避免存储体冲突

Shared memory is split into 32 banks, one for each thread in a warp. Spreading accesses across them lets all 32 threads read at full speed.

Consecutive 4-byte words land in consecutive banks, wrapping around after 32. So word 0 is bank 0, word 32 is bank 0 again.

A bank conflict happens when two threads in a warp hit the same bank but different words. The hardware must serialize those accesses.

An N-way conflict turns one fast access into N slow ones. A 32-way conflict can make shared memory as slow as if it were single-file. 🐢

When each thread reads tile[threadIdx.x], every thread hits a distinct bank. That is conflict-free and runs at full bandwidth.

Reading with a stride of 32 sends every thread to the same bank. That is the worst case, a full 32-way conflict.

Any even stride that shares a factor with 32 causes partial conflicts. A stride of 2, for example, gives a 2-way conflict across the warp.

Good news: if all threads read the same address, the hardware broadcasts it in one shot. Same word is fine, only same bank with different words conflicts.

For 2D tiles, add one extra column with +1 padding. This shifts each row so column accesses no longer all land in one bank.

The extra column makes the row length coprime with 32. Now stepping down a column visits a different bank each time, killing the conflict.

Bank conflicts are invisible in source code. Use Nsight Compute to measure shared-memory conflicts before spending effort fixing them.

Let us check what makes shared access fast.

You learned that shared memory has 32 banks, that same-bank different-word access serializes, and that +1 padding fixes column conflicts. Next: dynamic shared memory. 🎯