3
\$\begingroup\$

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 ti = threadIdx.x;
        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;

        // reset out to zero ready for atomic addition later
        if (li == 0)
          out[key[i]] = 0.f;
        __syncthreads();

        // 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)
          atomicAdd(&out[key[i]], val);
      }
   }

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.

\$\endgroup\$
  • \$\begingroup\$ I have actually gone through this myself now with time and am using warp shuffles. When I have time I will add an answer for future reference and others. \$\endgroup\$ – James Jan 28 '16 at 1:56

Your Answer

By clicking “Post Your Answer”, you agree to our terms of service, privacy policy and cookie policy

Browse other questions tagged or ask your own question.