# Parallel reduction by key implementations

I have an implementation of the reduction approach used in this document. Furthermore, I extended (crudely) this so I can reduce-by-key.

In my setup I can assume that a key array variable is constant through subsequent calls to a reduction kernel and that it is sorted in increasing order, and can thus pre-process and make any assumptions about the reduction necessary. I can further assume that each individual key will only ever be seen at the block level and as such I do not need to consider reduction across blocks.

To extend the approach used in the above document, I have padded data (with -1) such that each warp only deals with a single key value, and then the results of each individual warp's reduction is atomically added to the output array. However, this is inefficient as I have to launch many times more threads that are really necessary, and most are inactive.

An example key:

[0 0 0 0 0 0 0 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 ...


Below is the setup I have.

__global__ void reduce(int N, int* key, float* out)
{
int i = blockIdx.x * blockDim.x + ti;
if (i < N)
{
int li = ti % warpSize;

// code which computes a value per thread to be reduced,
// for threads where key[i] == -1, val = 0.f
float val = 123.45f;

if (li == 0)
out[key[i]] = 0.f;

// warp reduction
int offset = 0;
for (offset = warpSize / 2; offset > 0; offset /= 2)
val += __shfl_down(val, offset);

// atomically update out with each warp's result
if (li == 0)

What improvements could be appropriate for such a reduce-by-key implementation? For example, I think I should be using __ballot to only reduce those with keys matching a predicate, and could therefore remove my padding with -1s, although can't get my head around it to actually implement it.