GPUFlightAll performance labs

Matrix Multiplication 3: Share One 4 x 4 Tile

Every output C[row, col] needs one row of A and one column of B. That means neighboring output threads request many of the same A and B values. Use shared memory as a small workspace for one 4 x 4 matrix multiply: let the 16 threads load A and B together, synchronize, then reuse those tiles to calculate C.

Sign in to run this lab. GPU availability and credits are shown in the app.

Languages: CUDA C++, Python, Rust, HIP C++

Lesson

Global memory is large, but it is farther from the SM and its bandwidth is shared across the GPU. Repeatedly requesting the same values can therefore become a bottleneck. In this 4 x 4 multiply, 16 output threads each perform 4 products and request 2 inputs per product: 16 × 4 × 2 = 128 requested global reads. But the whole calculation contains only 16 A values and 16 B values. The block can cooperatively load those 32 values into its small, block-local shared memory, wait at __syncthreads(), and serve the repeated reads from there. That is 32 requested global reads instead of 128, a 4× reduction in algorithmic global-memory traffic. It does not promise an exact 4× runtime speedup: caches, synchronization, occupancy, and the GPU still affect measured performance.

The animated 2 by 2 arithmetic is illustrative. The assignment and its correctness checks use the 4 by 4 tile described by the runner.

What you will learn

  • Trace how one A row is reused across a row of C and one B column across a column of C.
  • Calculate why one 4 x 4 output tile requests 128 global-memory reads without shared-memory reuse.
  • Cooperatively load one A tile and one B tile into block-local shared memory.
  • Use __syncthreads() before any thread reads the shared tiles.
  • Reuse each shared value across multiple output elements.

Starter code

    __shared__ float aTile[4][4];
    __shared__ float bTile[4][4];

    const int tx = threadIdx.x;
    const int ty = threadIdx.y;

    // 1. Every thread loads one A value and one B value into shared memory.
    // TODO: aTile[ty][tx] = ...
    // TODO: bTile[ty][tx] = ...

    // 2. Make sure the whole block can see both complete tiles.
    // TODO: synchronize the block.

    // 3. Each thread computes one C value by reusing the shared tiles.
    float sum = 0.0f;
    for (int k = 0; k < N; ++k) {
        // TODO: accumulate one product from aTile and bTile.
    }

    // TODO: store the completed value in C.

Sign in to run