GPU Lab OPTIONAL LAB / GPU ARCHITECTURE
Sharing work within a block
A small change in layout can change bank conflicts.
Before you start: Chapters 5 and 8: warps, active lanes, shared memory and block barriers. This optional lab explores the layout of that shared workspace.
The mechanism
A tiled tensor operation can load data once and reuse it across threads in the same block. The layout of that reusable data matters.
New concepts
- Shared memory
- Fast, programmer-managed memory accessible to threads of the same block.
- Bank
- One independently serviced part of shared memory.
- Bank conflict
- Different words requested from the same bank need separate service rounds in this model.
- Broadcast
- Multiple lanes requesting the same shared-memory word can be served together on a read.
This model uses 32 banks. Successive four-byte words map to successive banks, wrapping around at bank zero. For word index w, the bank is w % 32, where % means remainder.
These bank rules describe reads; broadcasting a word does not make simultaneous writes safe.
Worked example
Worked example: column read, 8 lanes shown
Lanes read one column of a 32×32 tile. Word = lane × pitch. Bank = word % 32. With pitch 32, every lane hits bank 0: 32 distinct words in one bank need 32 service rounds. With pitch 33 (one padding word), the banks spread out: 1 round.
| Lane | L0 | L1 | L2 | L3 | L4 | L5 | L6 | L7 |
|---|---|---|---|---|---|---|---|---|
| Word (stride 32) | 0 | 32 | 64 | 96 | 128 | 160 | 192 | 224 |
| Bank (stride 32) | 0 | 0 | 0 | 0 | 0 | 0 | 0 | 0 |
| Word (stride 33) | 0 | 33 | 66 | 99 | 132 | 165 | 198 | 231 |
| Bank (stride 33) | 0 | 1 | 2 | 3 | 4 | 5 | 6 | 7 |
Experiment: try changing it
Imagine a square tile with 32 rows and 32 useful columns. Threads read one column, so successive lanes are separated by the row pitch: the number of allocated words per row. A pitch of 33 adds one unused padding word to each row.
Each round serves at most one distinct word from each bank. Banks work in parallel; the busiest bank determines the number of rounds. Try a row read, a column read, a padded column, and a broadcast.
Choose a configuration, then predict or inspect.
Requests grouped by bank
Select a round. Each bank card counts its distinct requested words; highlighted cards serve a word in this round. Select a lane to inspect its word and bank.
Model assumptions and limits
We model one warp instruction reading aligned FP32 words through 32 banks of four-byte width, one distinct word per bank per round, with read broadcast. Rounds express conflict multiplicity, not measured cycles. Global-memory traffic, shared-memory writes, barriers, races, allocation limits, 64-bit accesses and matrix-specific load instructions are not simulated.
CUDA Best Practices Guide: shared memory and memory banks ↗CONNECT TO CUDA
The programming connection
__shared__ places an array in shared memory. The two array dimensions allocate rows with a pitch of 33 words. Each thread fills one column across all rows. After the barrier, the threads read the first column, which thread zero filled. This fragment assumes a full tile, one block of 32 threads, and valid input and output allocations.
__syncthreads() is a block barrier: every thread in the block must reach it before any proceeds past it. Prior shared-memory writes become visible to the block. Here it ensures that thread zero's writes are ready before the other threads read them. A barrier inside a condition taken by only some block threads can hang or behave incorrectly. The bank model covers the final read; it does not simulate the loop or barrier.
This function is extracted from the complete self-checking CUDA program, which supplies its launch and host setup. Run it using the desktop appendix.
__global__ void readColumn(const float* input, float* output) {
// One block with 32 threads; a column read from a tile.
__shared__ float tile[32][33]; // 32 useful columns + padding
int lane = threadIdx.x;
for (int row = 0; row < 32; ++row) {
tile[row][lane] = input[row * 32 + lane];
}
__syncthreads();
output[lane] = tile[lane][0];
}TRY IT YOURSELF
Change one thing. Explain the result.
- Predict the rounds for a column read with pitch 32. Select lanes 0 and 1. They read different words: which banks receive them?
- Switch to a padded column with pitch 33. Explain why adding a padding word changes the bank mapping while preserving the useful column.
- Compare the unpadded column with a broadcast. Both use one bank. Why does only one of them need repeated service rounds?
- Choose every other word and reduce the active lanes to sixteen. Can this subset avoid the conflict that appears with a full warp?
Compare your explanation
Bank conflicts depend on both the bank and the word address. A count of requests per bank would incorrectly penalise broadcasts. Padding changes the row pitch, which changes the bank mapping for column reads.
Padding uses extra shared memory. That can affect how many blocks fit on an SM, connecting this layout choice to resident work in chapter 6. Revisit the cooperation chapter to apply block barriers to a computed tree and inspect an unsafe interleaving. The bank model on this page covers reads only.
References
Use these official references to check the concepts and code on this page.
- CUDA Best Practices: shared-memory banks (new tab)
Compare bank conflicts with the same-word broadcast exception.
- CUDA: synchronisation functions (new tab)
Check the scope and participation rules for __syncthreads().
These browser exercises do not run an installed PyTorch or CUDA toolkit. For optional local execution, choose documentation matching your installed tools. Links open online in a new tab; the questions and answer comparisons below work offline.
Use the documentation
Use the references to explain each answer, then apply it to this page’s example. No installation is needed.