Tiling at Scale: Read the Traffic, Not the Clock

If you followed the Matrix Multiplication path you have already staged one small tile by hand. This is that same pattern at production scale, and it arrives with a trap: tile this N x N kernel and the wall clock will barely move. At these sizes the L2 cache is already absorbing the redundant reads for you, so elapsed time and SM throughput look nearly identical before and after. The SASS global-sector counter is what grades you here, because it measures the traffic your kernel actually issues rather than the traffic today's cache happens to hide. You will also meet the case a small hand-worked tile never shows you: the last partial tile, when N is not a multiple of the tile size.

Languages: CUDA C++, Python, Rust

Lesson

A naive matrix multiply is correct but memory-hungry: to produce one output element it reads N elements of A and N of B, and the block next door re-reads almost all of the same data. Tiling fixes the ratio: each block cooperatively loads a TILE x TILE block of A and of B into shared memory, every thread in the block reuses those TILE values, and only then does the block move to the next tile along K. Each input element is now read from global memory about N/TILE times instead of N.

Here is the modern-GPU twist this lab is really about: run the naive kernel under the profiler and it does NOT look starved - SM throughput sits near 95%, L1 hit rate near 90%, and at N=2048 the whole B matrix (16 MB) fits inside a desktop GPU's L2 cache, so the redundant reads never reach global-memory request and the wall clock barely moves when you tile. The caches are doing your reuse for you, silently, at this size. The SASS global-sector counter tells the truth the clock hides: the naive kernel ISSUES roughly ten times the global-memory traffic, and that redundancy comes straight back as slowdown the moment the working set outgrows the cache (bigger N, real workloads, shared GPUs). Your tiled kernel's L1 hit rate will crater toward zero - that is the signature of success, not failure: the reuse moved out of the cache and into shared memory, where it is explicit and guaranteed instead of a property of today's cache size.

Metric contract: algorithmic reads/writes count source-level sweeps. GLOBAL_MEMORY_SECTORS counts SASS-derived issued/requested sectors in the selected profiling scope, including repeated launches. It is cache-blind: it is neither L2 traffic nor device-memory bytes. Trace elapsed time is a separate measurement; a global load may hit cache.

Cache reuse can hide device-memory latency for this workload. Wall-clock improvement is an observation to measure, not a guarantee that it will remain small.

What you will learn

Read the SASS global-sector counter as the ground truth for memory traffic, independent of what the caches absorb.
See why wall time and SM throughput can look identical for naive and tiled matmul at cache-resident sizes - and why the traffic still matters.
Stage a TILE x TILE block of each input into shared memory and reuse it across the block.
Use __syncthreads() to make the whole block agree on the tile before and after computing on it.
Handle the last, partial tile when N is not a multiple of the tile size (zero-pad, do not read out of bounds).

Starter code

    // Naive: one thread per output element. Each thread walks a full row of A and
    // a full column of B straight from global memory - lots of redundant global-memory request reads.
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    if (row < N && col < N) {
        float sum = 0.0f;
        for (int k = 0; k < N; ++k) {
            sum += A[row * N + k] * B[k * N + col];
        }
        C[row * N + col] = sum;
    }