GPU Challenges / CUDA

Parallel Sum Reduction: CUDA Exercise

hard Runs on a real NVIDIA T4 GPU in your browser

Parallel reduction turns an O(n) sequential sum into O(log n) parallel steps. It is a common GPU interview question and the building block for sums, maxima, norms and softmax. Getting it right means understanding shared memory, synchronization and how work is split across blocks.

The problem

Write a kernel that sums all elements of an array using parallel reduction. The array size is 1024 elements. Use shared memory for efficiency.

Build a block-level reduction in shared memory, then combine the partial sums on the host to recover the full array sum.

Your goal

  • Load one input value per thread into shared memory.
  • Reduce values inside the block using synchronized halving steps.
  • Write one partial sum per block into the output array.

How to approach it

  1. Load one element per thread from global memory into a __shared__ buffer.
  2. Loop with stride = blockDim.x / 2 down to 1, adding sdata[tid + stride] into sdata[tid] when tid < stride.
  3. Call __syncthreads() after the load and after every step of the loop.
  4. Write sdata[0] for each block to an output array and sum those partial results.

Common mistakes

  • Using the interleaved tid % (2 * stride) pattern, which works but causes warp divergence.
  • Reading shared memory before every thread has finished writing it.
  • Forgetting that the final answer needs the per-block sums combined, not just block 0.

What success looks like

  • The final host-side accumulated value should match the expected sum.
  • Shared memory reads and writes should stay inside the allocated block buffer.
  • Every reduction step should synchronize before the next one starts.

Starter code

Start from this scaffold. Fill in the TODOs, then run it against the hidden tests.

parallel-reduction.cu
123456789101112131415
#include <cuda_runtime.h>

// Kernel is launched as:  reductionSum<<<ceil(n/256), 256, 256*sizeof(float)>>>(input, output, n)
// gridDim.x  = ceil(n / 256)   blockDim.x = 256   shared mem = 256 floats
// tid = threadIdx.x              (lane within this block, 0..255)
// i   = blockIdx.x*256 + tid    (global element index)
__global__ void reductionSum(const float* input, float* output, int n) {
    extern __shared__ float sdata[];
    unsigned int tid = threadIdx.x;
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
    // TODO: Load input[i] into sdata[tid] (use 0.0f if i >= n)
    // TODO: __syncthreads()
    // TODO: Tree reduction: halve active threads each iteration
    // TODO: Thread 0 writes sdata[0] to output[blockIdx.x]
}
Solve it in the editor →Free account, no GPU or CUDA install needed.

Concepts used

  • Shared memory staging
  • Synchronization with __syncthreads()
  • Tree reduction patterns

Learn them first:

More CUDA challenges