reduceN.cuh

// Block-wide reduction where each thread locally reduces N
// values before letting a single warp take over
template <typename T, typename ReduceOp, int N>
__device__ T reduceBlockN(T *smem,
                         int numVals,
                         ReduceOp reduceOp,
                         T init) {
  T local = threadIdx.x < numVals ? smem[threadIdx.x] : init;

#pragma unroll
  for (int i = 1; i < N; ++i) {
    int index = threadIdx.x + (i * blockDim.x);
    T next = index < numVals ? smem[index] : init;
    local = reduceOp(local, next);
  }

  return reduceBlock<T, ReduceOp>(smem, blockDim.x < numVals ? blockDim.x : numVals, local, reduceOp, init);
}