def __call__(self, arr, size=None, res=None, init=0, stream=0):
        """Performs a full reduction.

        :param arr: A host or device array.
        :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]
Performs a full reduction.

        :param arr: A host or device array.
        :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.
        
numba/cuda/kernels/reduction.py