GPU Challenges / CUDA

Matrix Multiplication: CUDA Exercise

medium Runs on a real NVIDIA T4 GPU in your browser

Matrix multiplication is the workload GPUs are built around: it underlies every neural network layer. This exercise moves from 1D to 2D launches, mapping each thread to one row and column of the output matrix and computing a dot product of a row of A with a column of B.

The problem

Implement a simple 2D matrix multiplication kernel. Multiply two NxN matrices where N = 64. Each thread computes one output element.

Map a 2D CUDA launch onto matrix rows and columns so each thread computes exactly one output element of C.

Your goal

  • Derive row and column coordinates from blockIdx, blockDim, and threadIdx.
  • Accumulate the dot product for one output cell across k = 0..N-1.
  • Write the result into the correct flattened row-major location in C.

How to approach it

  1. Launch a 2D grid with dim3 block(16, 16) and enough blocks to cover the 64x64 output.
  2. Derive row from the y coordinates and col from the x coordinates of the block and thread.
  3. Loop over k, accumulating A[row * N + k] * B[k * N + col] in a register.
  4. Write the sum to C[row * N + col] only when both row and col are inside the matrix.

Common mistakes

  • Swapping row and column, which still compiles but produces a transposed or wrong result.
  • Accumulating directly into global memory inside the loop instead of a local variable.
  • Launching a 1D grid and trying to recover 2D coordinates by hand.

What success looks like

  • For the provided input, each output element should equal N * 2.0f.
  • Threads outside the matrix bounds should do no work.
  • The output buffer should be filled consistently across different rows and columns.

Starter code

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

matrix-multiply.cu
123456789101112
#include <cuda_runtime.h>
#define N 64

// Kernel is launched as:  matMul<<<dim3(ceil(N/16),ceil(N/16)), dim3(16,16)>>>(A, B, C)
// gridDim  = (ceil(N/16), ceil(N/16))   blockDim = (16, 16)
// Row:  int row = blockIdx.y * blockDim.y + threadIdx.y;
// Col:  int col = blockIdx.x * blockDim.x + threadIdx.x;
__global__ void matMul(const float* A, const float* B, float* C) {
    // C[row][col] = sum_k A[row][k] * B[k][col]
    // TODO: Use blockIdx.x/y and threadIdx.x/y to assign one output element per thread
    // TODO: Accumulate the dot product and write to C[row * N + col]
}
Solve it in the editor →Free account, no GPU or CUDA install needed.

Concepts used

  • 2D block geometry
  • Row-major indexing math
  • Nested accumulation inside a kernel

Learn them first:

More CUDA challenges