Using gridDim.x in reduction (max/min/sum) cuda kernels

[coreylowman]

Often, it is recommend to loop over items in a cuda kernel like:

for (unsigned int i = tid; i < n; i += blockDim.x * gridDim.x) {
   ...
}

This was recently done for most kernels in #787

However, this is not currently implemented for reductions.

Currently, each chunk to be reduced is spread across thread blocks. So the threads in a single block may be working on 1 or more chunks. How does this work with grid striding?

Additionally, how would shared memory be used?