GPU Challenges / CUDA

SAXPY: Single Precision A*X Plus Y: CUDA Exercise

easy Runs on a real NVIDIA T4 GPU in your browser

SAXPY (Single-precision A times X Plus Y) is a level-1 BLAS routine and a standard GPU benchmark. It is memory-bound: each element needs two loads and one store for only two floating-point operations, which makes it a good first look at why memory bandwidth, not arithmetic, limits many kernels.

The problem

Implement the BLAS SAXPY operation: y = a*x + y, where a is a scalar and x, y are vectors. This is one of the most fundamental GPU operations. Vector size N = 8192. Each thread handles one element.

Implement the BLAS SAXPY pattern y = a*x + y, one element per thread, with a clean 1D launch and minimal kernel logic.

Your goal

  • Compute the global index once and reuse it.
  • Apply the SAXPY formula only when the index is valid.
  • Leave the host code's launch configuration and data flow intact.

How to approach it

  1. Pass the scalar a by value; only x and y need device memory.
  2. Compute one global index per thread and update y[i] = a * x[i] + y[i] in place.
  3. Cover N = 8192 elements with a 1D grid and keep the bounds check.
  4. Verify a few outputs on the host against the same formula computed on the CPU.

Common mistakes

  • Copying a to device memory and passing a pointer, which adds work for no benefit.
  • Writing the result to a new array but printing the old y.
  • Comparing floats for exact equality instead of within a small tolerance.

What success looks like

  • The first printed values should match the expected SAXPY math.
  • The kernel should handle N values without skipping the tail.
  • No extra temporaries are needed beyond the index and formula.

Starter code

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

saxpy.cu
123456789
#include <cuda_runtime.h>

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

Concepts used

  • Embarrassingly parallel kernels
  • Scalar-plus-vector operations
  • Simple correctness validation from printed results

Learn them first:

More CUDA challenges