SGEMM: 2D Register Tiling Visualizer

The Outer Product approach: Heavy math intensity with dual register arrays (regM and regN).

Global Reads
0
Shared Reads
0
Register Reads
0

Simulation Engine

Init Block
Animation Speed600ms
Current Focus: Block (0, 0)
// 2D Tiling: BM=4, BN=4, BK=4, TM=2, TN=2
// 4 Threads compute a 4x4 Output Block.
__shared__ float As[4][4];
__shared__ float Bs[4][4];
float threadResults[2 * 2] = {0.0f};
float regM[2] = {0.0f};
float regN[2] = {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 2D vectors to registers (Double Broadcast!)
for (int i = 0; i < 2; ++i) regM[i] = As[ty*2 + i][dotIdx];
for (int j = 0; j < 2; ++j) regN[j] = Bs[dotIdx][tx*2 + j];
// Pure Register Outer-Product
for (int i = 0; i < 2; ++i) {
  for (int j = 0; j < 2; ++j) {
  threadResults[i*2+j] += regM[i] * regN[j];
  }
}
}
__syncthreads();
}
// Write 2x2 grid from registers to global C
for (int i = 0; i < 2; ++i) {
  for (int j = 0; j < 2; ++j) {
    C[... + j] = threadResults[i*2+j];
  }
}

Global Memory (DRAM)

Matrix A (8 × 8)
Matrix B (8 × 8)

Shared Memory (SRAM)

As [4][4] (Loaded to regM)
Bs [4][4] (Loaded to regN)

Global Matrix C

Matrix C (8 × 8)
Each thread computes a 2×2 square grid

Thread Monitor (4 Threads | TM=2, TN=2)

Outer Product Regs
TID Action regM[0] regM[1] regN[0] regN[1] Res[0,0] Res[0,1] Res[1,0] Res[1,1]

Why 2D Register Tiling Dominates (The Outer Product)

Visualizing the Outer Product Mapping:

During the FMA Math step, watch the Shared Memory grids (As and Bs) carefully. The lightly colored column and row indicate the vectors loaded into private registers. The brightly highlighted pink and indigo cells show exactly which single scalar value is currently being used by the thread to compute the intersection.

Notice how as the math progresses, these highlighted operands "slide" across the loaded vectors to systematically calculate every intersection in the 2x2 output grid!

The Shift to Outer Products: Instead of loading 1 element to reuse along a 1D strip, each thread loads a column vector (regM) from A and a row vector (regN) from B.

The Math: When you multiply a column vector by a row vector, you get a full 2D matrix! The thread performs \(\text{TM} \times \text{TN}\) math operations using only those few registers.

Shattering the Memory Wall: Arithmetic Intensity skyrockets. The ratio of math operations to Shared Memory reads goes through the roof.

\(\text{Shared Reads: } (TM + TN)\)
\(\text{FMA Operations: } (TM \times TN)\)

The Problem with 1D Tiling (4×1 strip)

  • Load 1 value from Matrix B into a register.
  • Load 4 values from Matrix A.
  • Perform 4 Multiply-Accumulate (FMA) operations.
  • The Ratio: 5 Shared Memory reads for 4 math ops. The GPU's math cores wait for data!

The Magic of 2D Tiling (4×4 square)

  • Load a column vector from Matrix A (4 values).
  • Load a row vector from Matrix B (4 values).
  • Multiplying them gives a 4×4 grid. Perform 16 FMA operations using only those 8 registers.
  • The Ratio: 8 Shared Memory reads for 16 math ops! Memory latency is hidden.

Why this matters on real GPUs:

In production kernels, you usually see 8×8 or 16×16 2D register tiling. If a thread uses an 8×8 tile, it reads 8 values from A and 8 values from B (16 reads total), and uses them to perform 64 math operations. That means for every single time the GPU goes to Shared Memory, it executes 4 math operations. By the time it needs to read from memory again, the math cores have been kept so busy that the memory latency is completely hidden. This is how highly optimized matrix multipliers (like cuBLAS) reach 90%+ of the GPU's maximum theoretical speed!