Lecture 4 | CUDA Programming

[chufanchen]

Warp shuffles are a faster mechanism for moving data between threads in the same warp. There are 4 variants:

• __shfl_up_sync: copy from a lane with lower ID relative to caller
• __shfl_down_sync: copy from a lane with higher ID relative to caller
• __shfl_xor_sync: copy from a lane based on bitwise XOR of own lane ID
• __shfl_sync: copy from indexed lane ID

Here the lane ID is the position within the warp.

T __shfl_up_sync(unsigned mask, T var, unsigned int delta);
T __shfl_down_sync(unsigned mask, T var, unsigned int delta);

• mask controls which threads are involved - usually set to -1 or 0xffffffff, equivalent to all 1's.
• var is a local register variable(int, unsigned int, long long, unsigned long long, float or double)
• delta is the offset within the warp – if the appropriate thread does not exist (i.e. it’s off the end of the warp) then the value is taken from the current thread.

T __shfl_xor_sync(unsigned mask, T var, int laneMask);

• an XOR (exclusive or) operation is performed between laneMask and the calling thread’s laneID to determine the lane from which to copy the value(laneMask controls which bits of laneID are flipped)
• a “butterfly” type of addressing, very useful for reduction operations and FFTs

T __shfl_sync(unsigned mask, T var, int srcLane);

• copies data from srcLane

[chufanchen]

Key requirements for a reduction operator $\circ$ are:

• commutative: $a \circ b = b \circ a$
• associative: $a \circ (b \circ c) = (a \circ b) \circ c$

Together, they mean that the elements can be re-arranged and combined in any order.

Assuming each thread starts with one value, the approach is to

1. first add the values within each thread block, to form a partial sum. (Local reduction)
2. then add together the partial sums from all of the blocks. (Global reduction)

The first phase does parallel summation of $N$ values:

1. first sum them in pairs to get $N/2$ values
2. repeat the procedure until we have only one value

Are there any problems with warp divergence? Note that not all threads can be busy all of the time:

• $N/2$ operations in first phase
• $N/4$ in second
• $N/8$ in third
• etc.

For efficiency, we want to make sure that each warp is either fully active or fully inactive, as far as possible.

Where should data be held. Threads need to access results produced by other threads:

• global device arrays would be too slow, so use shared memory
• need to think about synchronisation

__global__ void sum(float *d_sum, float *d_data) {
  extern __shared__ float temp[];
  int tid = threadIdx.x;
  temp[tid] = d_data[tid+blockIdx.x*blockDim.x];
  for (int d=blockDim.x/2; d>0; d=d/2) {
    __syncthreads();
    if (tid<d) temp[tid] += temp[tid+d];
  }
  if (tid == 0) d_sum[blodkIdx.x] = temp[0];
}