'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.'

def __call__(self, arr, size=None, res=None, init=0, stream=0):

    from numba import cuda
    if (arr.ndim != 1):
        raise TypeError('only    support    1D    array')
    if (size is not None):
        arr = arr[:size]
    init = arr.dtype.type(init)
    if (arr.size < 1):
        return init
    kernel = self._compile(arr.dtype)
    blocksize = (_NUMWARPS * _WARPSIZE)
    size_full = ((arr.size // blocksize) * blocksize)
    size_partial = (arr.size - size_full)
    full_blockct = min((size_full // blocksize), (_WARPSIZE * 2))
    partials_size = full_blockct
    if size_partial:
        partials_size += 1
    partials = cuda.device_array(shape=partials_size, dtype=arr.dtype)
    if size_full:
        kernel[(full_blockct, blocksize, stream)](arr[:size_full], partials[:full_blockct], init, True)
    if size_partial:
        kernel[(1, size_partial, stream)](arr[size_full:], partials[full_blockct:], init, (not full_blockct))
    if (partials.size > 1):
        kernel[(1, partials_size, stream)](partials, partials, init, False)
    if (res is not None):
        res[:1].copy_to_device(partials[:1], stream=stream)
        return
    else:
        return partials[0]
