Matrix Multiplication 4: Reuse a Shared-Memory Tile

Apply the shared-memory tile pattern to N = 4096 on your selected GPU. Load each 16 x 16 A and B tile once, let the block reuse it, then compare the profile with the naive kernel from Stage 2.

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

Lesson

The naïve kernel fetches a value of A or B from global memory separately for many output threads. A tile changes that: the block loads each value once into shared memory, synchronizes, and reuses it for all output elements in the tile. The math is unchanged; only where the reuse happens changes. Use a square 16 x 16 block and reserve 2048 bytes for two float tiles.

Shared tiling reduces issued global loads. Cache hits can make physical device-memory traffic and runtime change by different factors.

What you will learn

Cooperatively load A and B tiles into dynamic shared memory.
Use __syncthreads() before consuming a tile and before overwriting it.
Compare the tiled kernel with the profiled naïve baseline.

Starter code

    // One 16 x 16 block produces one output tile. Dynamic shared memory holds
    // an A tile followed by a B tile: 2 * 16 * 16 * sizeof(float) = 2048 bytes.
    extern __shared__ float shared[];
    const int tile = blockDim.x;
    float* aTile = shared;
    float* bTile = shared + tile * tile;

    // Apply the four-step pattern from Stage 3 to every tile along K:
    // 1. cooperatively load A and B, 2. synchronize, 3. reuse the tiles,
    // 4. synchronize before loading the next pair.
    // TODO: derive coordinates, accumulate the full result, and store C[row, col].