SGEMM: 1D Register Tiling Visualizer

Increasing Arithmetic Intensity: Threads load into Registers to reuse data and slash Shared Memory access.

This visualizer uses a scaled-down teaching configuration (BM=8, BN=4, BK=4, TM=4) so the dataflow fits on screen. The CUDA kernel in this repo uses the same algorithmic structure with larger production constants.

Global Reads

Shared Reads

Simulation Engine

Init Block

Problem Size

M=8, N=8, K=8

M N K

Block tiling stays fixed at BM=8, BN=4, BK=4, TM=4.

Tiling Terms

BM = number of output rows each block computes

BN = number of output columns each block computes

BK = how much of the K dimension is loaded at a time

TM = number of output rows each thread computes

Selected Thread

T(0,0)

Follow one thread's 4-register output strip across all phases.

Animation Speed600ms

Current Focus: Block (0, 0)

// 1D Tiling: BM=8, BN=4, BK=4, TM=4
// 8 Threads per block compute an 8x4 C-Tile.

__shared__ float As[8][4];
__shared__ float Bs[4][4];

const int ty = threadIdx.y; // 0 to 1
const int tx = threadIdx.x; // 0 to 3
float threadResults[4] = {0.0f};

for (int bkIdx = 0; bkIdx < K; bkIdx += 4) {

// Cooperative load to Shared Memory
As[...] = A[...];
Bs[...] = B[...];

__syncthreads();

for (int dotIdx = 0; dotIdx < 4; ++dotIdx) {

// Load 1 element from Shared to Register
// Hardware Broadcasts to identical `tx` threads
float Btmp = Bs[dotIdx][tx];

#pragma unroll
for (int resIdx = 0; resIdx < 4; ++resIdx) {

// Reuse Btmp 4 times! Math Intensity Up!
threadResults[resIdx] += As[ty * 4 + resIdx][dotIdx] * Btmp;

}

__syncthreads();

}

// Write accumulated results from registers back to global memory C
#pragma unroll
for (int resIdx = 0; resIdx < 4; ++resIdx) {
C[... + resIdx*N + tx] = threadResults[resIdx];
}

Global Memory (DRAM)

Matrix A (8 × 8)

Matrix B (8 × 8)

Shared Memory (SRAM)

As [8][4]

Bs [4][4]

BTMP

T(0,0)

Shared tiles are empty until the cooperative load phase begins.

Highlight Legend

Amber outline: rows or columns used by the selected thread across the whole tile.

Blue glow: the exact As or Bs cell the selected thread is using right now.

Purple glow/fill: block-wide active FMA work for the current (dotIdx, resultIndex) step, including other threads.

Pink highlight: register-side activity, such as loading or reusing Btmp.

Blue strip on C: the 4-output column owned by the selected thread.

Solid blue write: a finished value is being stored back to global C.

Matrix C

Matrix C (8 × 8)

Each thread computes a 4×1 vertical strip

Thread Monitor (8 Threads | TM=4)

Registers & Math

Selected Thread Lens

T(0,0) -> C rows 0..3, col 0

As Rows Used By This Thread

Bs Column Used By This Thread

TID	Action	Btmp	threadResults[0]	threadResults[1]	threadResults[2]	threadResults[3]

Why 1D Register Tiling (Work-Per-Thread)?

Is this 1D Tiling? Yes! Each thread computes a 1D column vector (in our viz, a \(4 \times 1\) tile) of the output block.

The Core Idea: We load a single value from Matrix B in shared memory into a fast, private hardware Register (Btmp).

Increasing Arithmetic Intensity: Because the thread calculates multiple rows (TM), it can reuse Btmp to multiply against multiple elements from Matrix A.

\(\text{SRAM Reads for B} \div \text{TM}\)

The Impact on the "Memory Wall":

Standard Shared Memory tiling fixes the Global Memory bottleneck, but replaces it with a Shared Memory bottleneck (Shared Memory is fast, but Registers are vastly faster).

Without Register Tiling (TM=1): 1 Shared Read of A + 1 Shared Read of B = 1 FMA instruction. (Ratio 2:1)
With Register Tiling (TM=4): We read Btmp once. Then we read 4 values from A to perform 4 FMA instructions. That's 5 Shared Reads for 4 FMAs. (Ratio 1.25:1)
With TM=8: 9 Shared Reads for 8 FMAs. (Ratio ~1.1:1). The GPU is no longer starved for memory and can calculate at maximum theoretical speed!

Why is the Shared Tile for A (8×4) larger than B (4×4)?

It directly relates to the work each thread is doing! Our block has 8 threads arranged in a 2×4 grid. Since each thread computes a 4×1 vertical strip (TM = 4), the block collectively computes an 8×4 block of the output Matrix C.

To compute an 8×4 block of C, the threads need 8 rows from Matrix A but only 4 columns from Matrix B. Thus, the tile fetched from A (As) must be twice as tall as the tile fetched from B (Bs).