GPU Challenges / CUDA

Vector Addition: CUDA Exercise

easy Runs on a real NVIDIA T4 GPU in your browser

Vector addition is the "hello world" of CUDA. It teaches the single most important pattern in GPU programming: turning a thread's block and thread coordinates into a global index, so thousands of threads each handle exactly one array element in parallel.

The problem

Write a CUDA kernel that adds two arrays element-wise. Each thread should handle one element. The arrays are of size N = 1024.

Write the classic CUDA warm-up kernel: each thread computes one element of the output vector and guards against running past the end of the array.

Your goal

  • Compute the global thread index with blockIdx.x, blockDim.x, and threadIdx.x.
  • Only write to c[i] when the computed index is inside the vector length.
  • Keep the host launch configuration aligned with the problem size N = 1024.

How to approach it

  1. Compute the global index as blockIdx.x * blockDim.x + threadIdx.x.
  2. Guard the write with if (i < n), because the last block usually has more threads than remaining elements.
  3. Launch ceil(n / 256) blocks of 256 threads so every element is covered.
  4. Copy inputs to the device, launch the kernel, copy the result back and print a few values to check them.

Common mistakes

  • Leaving out the bounds check, which writes past the end of the array when n is not a multiple of the block size.
  • Using threadIdx.x alone as the index, so every block overwrites the same first 256 elements.
  • Forgetting cudaMemcpy back to the host and printing uninitialised host memory.

What success looks like

  • The first few printed results should equal a[i] + b[i].
  • No out-of-bounds writes should happen for the last block.
  • The kernel should work for any n, not only exact multiples of 256.

Starter code

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

vector-add.cu
12345678910
#include <cuda_runtime.h>

// Kernel is launched as:  vectorAdd<<<ceil(n/256), 256>>>(a, b, c, n)
// gridDim.x  = ceil(n / 256)   (one block per 256 elements)
// blockDim.x = 256
// Thread ID:  int i = blockIdx.x * blockDim.x + threadIdx.x;
__global__ void vectorAdd(const float* a, const float* b, float* c, int n) {
    // TODO: Compute global thread index i
    // TODO: If i < n, set c[i] = a[i] + b[i]
}
Solve it in the editor →Free account, no GPU or CUDA install needed.

Concepts used

  • 1D grid and block indexing
  • Boundary checks in simple kernels
  • Host-to-device and device-to-host memory flow

Learn them first:

More CUDA challenges