タイルのフェーズをループする

A full row times a full column is too big for one tile. So you split that long sum into chunks of width TILE, one chunk per phase.

If the matrices are N wide and tiles are TILE wide, you need N / TILE phases to cover the whole inner dimension.

Wrap your load-sync-compute steps in a loop over phase. Each pass slides the tile window further along the row of A and column of B.

The local sum variable lives outside the loop, so it keeps growing. Each phase adds its slice of the dot product to the running total.

Each phase shifts the column you read from A and the row you read from B by phase * TILE. That is how the window advances.

Just like before, call __syncthreads() after the loads so the tile is complete before anyone computes on it.

The inner loop adds TILE products into sum, using only the freshly loaded tile. It contributes one chunk of the final dot product.

End each phase with another __syncthreads() so no thread overwrites the tile while a slower thread is still reading it. 🚧

One barrier guards reads-after-load, the other guards loads-after-read. Together they keep every thread in lockstep across phases.

After all phases finish, the running sum is the complete dot product. Store it into C once, guarded by a bounds check.

When N is not a clean multiple of TILE, load zero for out-of-range elements so the extra products add nothing to the sum.

Think about where the running total lives during the phase loop.

You looped over phases, shifting tiles, syncing twice, and accumulating partial sums into one total. Now let us measure how much faster it is. 📈