parameters: (self, arr, size=None, res=None, init=0, stream=0)
docstring: Performs a full reduction.
  
          :param arr: A host or device array. If a device array is given, the
                      reduction is performed inplace and the values in the array
                      are overwritten. If a host array is given, it is copied to
                      the device automatically.
          :param size: Optional integer specifying the number of elements in
                      ``arr`` to reduce. If this parameter is not specified, the
                      entire array is reduced.
          :param res: Optional device array into which to write the reduction
                      result to. The result is written into the first element of
                      this array. If this parameter is specified, then no
                      communication of the reduction output takes place from the
                      device to the host.
          :param init: Optional initial value for the reduction, the type of which
                      must match ``arr.dtype``.
          :param stream: Optional CUDA stream in which to perform the reduction.
                      If no stream is specified, the default stream of 0 is
                      used.
          :return: If ``res`` is specified, ``None`` is returned. Otherwise, the
                  result of the reduction is returned.
docstring_summary: Performs a full reduction.
docstring_tokens: Performs a full reduction .
function: def __call__(self, arr, size=None, res=None, init=0, stream=0):
          """Performs a full reduction.
  
          :param arr: A host or device array. If a device array is given, the
                      reduction is performed inplace and the values in the array
                      are overwritten. If a host array is given, it is copied to
                      the device automatically.
          :param size: Optional integer specifying the number of elements in
                      ``arr`` to reduce. If this parameter is not specified, the
                      entire array is reduced.
          :param res: Optional device array into which to write the reduction
                      result to. The result is written into the first element of
                      this array. If this parameter is specified, then no
                      communication of the reduction output takes place from the
                      device to the host.
          :param init: Optional initial value for the reduction, the type of which
                      must match ``arr.dtype``.
          :param stream: Optional CUDA stream in which to perform the reduction.
                      If no stream is specified, the default stream of 0 is
                      used.
          :return: If ``res`` is specified, ``None`` is returned. Otherwise, the
                  result of the reduction is returned.
          """
          from numba import cuda
  
          # ensure 1d array
          if arr.ndim != 1:
              raise TypeError("only support 1D array")
  
          # adjust array size
          if size is not None:
              arr = arr[:size]
  
          init = arr.dtype.type(init)  # ensure the right type
  
          # return `init` if `arr` is empty
          if arr.size < 1:
              return init
  
          kernel = self._compile(arr.dtype)
  
          # Perform the reduction on the GPU
          blocksize = _NUMWARPS * _WARPSIZE
          size_full = (arr.size // blocksize) * blocksize
          size_partial = arr.size - size_full
          full_blockct = min(size_full // blocksize, _WARPSIZE * 2)
  
          # allocate size of partials array
          partials_size = full_blockct
          if size_partial:
              partials_size += 1
          partials = cuda.device_array(shape=partials_size, dtype=arr.dtype)
  
          if size_full:
              # kernel for the fully populated threadblocks
              kernel[full_blockct, blocksize, stream](arr[:size_full],
                                                      partials[:full_blockct],
                                                      init,
                                                      True)
  
          if size_partial:
              # kernel for partially populated threadblocks
              kernel[1, size_partial, stream](arr[size_full:],
                                              partials[full_blockct:],
                                              init,
                                              not full_blockct)
  
          if partials.size > 1:
              # finish up
              kernel[1, partials_size, stream](partials, partials, init, False)
  
          # handle return value
          if res is not None:
              res[:1].copy_to_device(partials[:1], stream=stream)
              return
          else:
              return partials[0]
function_tokens: def __call__ ( self , arr , size = None , res = None , init = 0 , stream = 0 ) : from numba import cuda # ensure 1d array if arr . ndim != 1 : raise TypeError ( "only support 1D array" ) # adjust array size if size is not None : arr = arr [ : size ] init = arr . dtype . type ( init ) # ensure the right type # return `init` if `arr` is empty if arr . size < 1 : return init kernel = self . _compile ( arr . dtype ) # Perform the reduction on the GPU blocksize = _NUMWARPS * _WARPSIZE size_full = ( arr . size // blocksize ) * blocksize size_partial = arr . size - size_full full_blockct = min ( size_full // blocksize , _WARPSIZE * 2 ) # allocate size of partials array partials_size = full_blockct if size_partial : partials_size += 1 partials = cuda . device_array ( shape = partials_size , dtype = arr . dtype ) if size_full : # kernel for the fully populated threadblocks kernel [ full_blockct , blocksize , stream ] ( arr [ : size_full ] , partials [ : full_blockct ] , init , True ) if size_partial : # kernel for partially populated threadblocks kernel [ 1 , size_partial , stream ] ( arr [ size_full : ] , partials [ full_blockct : ] , init , not full_blockct ) if partials . size > 1 : # finish up kernel [ 1 , partials_size , stream ] ( partials , partials , init , False ) # handle return value if res is not None : res [ : 1 ] . copy_to_device ( partials [ : 1 ] , stream = stream ) return else : return partials [ 0 ]