Tiling the Inner Product
Loading sub-tiles of A and B per phase.
Tiling the Inner Product is a free CUDA Academy lesson on CoddyKit — lesson 2 of 4. You can read the complete lesson below for free — then practise it hands-on in the browser with a built-in code editor and a 24/7 AI tutor. It is part of the CUDA Academy learning path, one of 4 lessons in the course, and your progress syncs across the web and the CoddyKit app.
The Tiling Idea
Tiling breaks the matrices into small square tiles that fit in fast on-chip memory. Threads cooperate to load a tile once and reuse it many times.
Why Shared Memory
A tile lives in __shared__ memory, visible to every thread in the block. Reading it is far faster than hitting global memory again and again. ⚡
__shared__ float As[TILE][TILE];
__shared__ float Bs[TILE][TILE];One Block, One Output Tile
Each block is responsible for one TILE-by-TILE patch of the output C. Its threads team up to compute that whole patch together.
Tile Size Matches Block
You pick TILE equal to the block's width, often 16 or 32. That way each thread loads exactly one element of each tile.
#define TILE 16
dim3 threads(TILE, TILE);Mapping Thread to Tile Slot
Inside the tile, a thread's slot is just its threadIdx. Its global row and col still come from the block and thread indices.
int ty = threadIdx.y, tx = threadIdx.x;
int row = blockIdx.y*TILE + ty;
int col = blockIdx.x*TILE + tx;Loading a Tile of A
Each thread copies one element of A's current tile into shared memory. Together the block stages a full TILE-by-TILE block of A.
As[ty][tx] = A[row*N + (phase*TILE + tx)];Loading a Tile of B
At the same time, each thread loads one element of B's tile. Now both tiles sit on chip, ready for fast repeated reads.
Bs[ty][tx] = B[(phase*TILE + ty)*N + col];Sync Before You Compute
Call __syncthreads() so every thread finishes loading before anyone reads the tile. Skipping this gives garbage results.
__syncthreads();Compute on the Tile
Now each thread does a short loop over the tile, reading only shared memory. These reads are dramatically cheaper than global ones.
for (int k = 0; k < TILE; ++k)
sum += As[ty][k] * Bs[k][tx];The Reuse Payoff
Every loaded value gets used by TILE threads instead of one. That data reuse is the whole reason tiled matmul flies.
One Tile Is Not Enough
A single tile only covers part of the dot product. You repeat the load-sync-compute steps across many phases, which the next lesson handles.
Quick Check
Recall the order of steps when working with a shared tile.
Recap
You staged tiles of A and B into shared memory, synced, then computed with cheap on-chip reads. Reuse is the win. Phases come next. 🧱
Frequently asked questions
Is the “Tiling the Inner Product” lesson free?
Yes — the full text of “Tiling the Inner Product” is free to read here on the web, and the CUDA Academy course includes 4 lessons in total. To practise it interactively (a built-in code editor and a 24/7 AI tutor) and unlock the rest of the CUDA Academy course, upgrade to CoddyKit PRO.
What will I learn in “Tiling the Inner Product”?
Loading sub-tiles of A and B per phase. You practise CUDA Academy with hands-on code you run directly in the browser, and a 24/7 AI tutor answers your questions as you work through the lesson.
Do I need any experience to start CUDA Academy?
No prior experience is required. CUDA Academy on CoddyKit is structured for beginners through advanced learners; this is — lesson 2 of 4, so you can start here or from the beginning and move at your own pace.
How long does the “Tiling the Inner Product” lesson take?
Most CoddyKit lessons take about 5–10 minutes. Each one is bite-sized and interactive, so you make steady progress and pick up exactly where you left off across the web and the app.
Can I write and run code in this CUDA Academy lesson?
Yes. Every CUDA Academy lesson includes a built-in code editor, so you write and run real code right in your browser and get instant AI feedback — no local setup required.
All lessons in this course
- The Naive Matmul Kernel
- Tiling the Inner Product
- Looping Over Tile Phases
- Measuring the Speedup