GPU Challenges / CUDA

Dot Product: CUDA Exercise

medium Runs on a real NVIDIA T4 GPU in your browser

A dot product combines two GPU patterns: an element-wise map (multiply a[i] by b[i]) and a reduction (sum everything). It is the simplest exercise where threads must cooperate, which introduces shared memory and __syncthreads().

The problem

Compute the dot product of two vectors on the GPU. The dot product is the sum of element-wise products: result = sum(a[i] * b[i]). Use parallel reduction with shared memory for efficiency. Array size N = 4096.

Compute a parallel dot product by multiplying vector elements in the kernel and reducing the partial products efficiently inside each block.

Your goal

  • Multiply a[i] * b[i] before loading into shared memory.
  • Reduce each block's partial products down to a single value.
  • Return block-level partial sums for final accumulation on the host.

How to approach it

  1. Each thread multiplies its pair of elements, using a grid-stride loop if there are more elements than threads.
  2. Store each thread's partial product in a __shared__ array sized to the block.
  3. Reduce the shared array in halving steps, calling __syncthreads() between steps.
  4. Have thread 0 of each block write the block's sum, then add the block sums on the host or with atomicAdd.

Common mistakes

  • Missing __syncthreads() between reduction steps, which causes intermittent wrong totals.
  • Placing __syncthreads() inside a branch that not all threads reach, which can hang the block.
  • Assuming a block size that is not a power of two in the halving loop.

What success looks like

  • The combined result should match the expected dot product printed by the host.
  • Threads beyond n should contribute zero, not garbage values.
  • The partial sum array should contain one valid result per launched block.

Starter code

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

dot-product.cu
123456789101112131415
#include <cuda_runtime.h>
#define THREADS_PER_BLOCK 256

// Kernel is launched as:  dotProduct<<<ceil(n/256), 256, 256*sizeof(float)>>>(a, b, partialSums, 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 dotProduct(const float* a, const float* b, float* partialSums, int n) {
    extern __shared__ float sdata[];
    unsigned int tid = threadIdx.x;
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
    // TODO: Load a[i]*b[i] into sdata[tid] (0.0f if i >= n)
    // TODO: __syncthreads() then tree reduction
    // TODO: Thread 0 writes sdata[0] to partialSums[blockIdx.x]
}
Solve it in the editor →Free account, no GPU or CUDA install needed.

Concepts used

  • Element-wise parallel work
  • Reduction reuse from previous challenge
  • Numerical verification against CPU expectations

Learn them first:

More CUDA challenges