Shared-memory tiling and bank conflicts
Design a cooperative shared-memory tile and predict its bank mapping, barrier locations, capacity cost, and global-transaction change
Before you begin
Follow a linked prerequisite when you cannot explain it from memory. Background skills without links are stated as abilities rather than hidden terminology.
- Chapter 5 block barriers
- Two-dimensional matrix tiles
- Previous section's requested-versus-transferred byte calculation
Shared memory is storage owned by one thread block. Threads explicitly write its contents, synchronize when required, and read values written by other threads. This makes shared memory useful when a block must reuse an input tile, exchange partial results, or change the order in which global memory is accessed.
Begin with matrix multiplication. Several output elements use the same input values. Without staging, different threads can request the same values from global memory repeatedly. A block can instead load a tile once into shared memory and let its threads reuse the tile. This reduces global traffic only when the tile is reused enough to repay the load and synchronization cost.
Shared memory has its own physical access rules. The address space is divided into banks that can serve different addresses concurrently. A warp request that maps different lane addresses to the same bank may be split into several service steps. Shared memory is therefore much closer to execution than off-chip device memory, but an unfavorable bank mapping can still limit a kernel.
Use shared memory for explicit reuse or rearrangement
A block first assigns global loads to threads so that warp addresses are coalesced. Each thread stores its values into a shared tile. After the required barrier, threads can read the tile in a different order or read values loaded by neighboring threads. This two-step access is common in transposes, reductions, stencils, and matrix kernels.
Count reuse before adding the tile. If each loaded value is consumed once by the same thread, shared memory may only add instructions and a barrier. If each value feeds many outputs or converts scattered global accesses into coalesced transactions, the added on-chip work can be beneficial. Hardware caches may already supply some reuse, so compare total global traffic before and after.
Calculate the bank mapping
On NVIDIA devices of compute capability 5.x and later, the standard model uses 32 banks with successive 32-bit words assigned to successive banks. For a 32-bit word address w, bank = w mod 32 is a useful starting calculation. Consult the target documentation for wider accesses and architecture details.
For a row-major 32 × 32 FP32 tile, lanes that read one row use consecutive word addresses and therefore consecutive banks. Lanes that read one column use addresses separated by 32 words. Every address has the same remainder modulo 32, so the request maps many different words to one bank and is serialized.
Declaring the physical tile as 32 × 33 changes the row stride to 33 words. Column addresses now advance by 33, so their bank indices advance by one modulo 32. The extra column is padding; it does not change the logical matrix. When several lanes request exactly the same shared-memory word, hardware can broadcast it, so repeated bank number alone is not sufficient to claim a conflict. Compare the complete addresses.
Protect tile phases and account for capacity
A cooperative tile normally needs a barrier after its load because some threads will read values written by other warps. A loop that reuses the same shared buffer can also need a barrier after consumption so that fast producers do not overwrite the tile while slower consumers still read it. Draw the states load, ready, consume, and reusable for each buffer.
Double buffering allocates two tiles so one stage can load future data while another stage computes. Asynchronous copy features can separate arrival from waiting more precisely, but they do not eliminate the dependency. The consumer must still wait until its tile is ready, and a buffer cannot be reused before its readers finish.
Shared-memory capacity is allocated per resident block. Larger tiles and more stages can improve reuse and overlap while reducing the number of blocks on an SM. Recalculate residency after padding and after every stage-count change. A small increase can cross a capacity threshold.
Worked example: Removing conflicts from a tiled transpose
A block loads a 32×32 FP32 tile by rows and writes it by columns. Global accesses are coalesced, but the shared-memory read is slow.
- A row-major tile with stride 32 causes column-reading lanes to address words separated by 32 positions. With 32 banks, those positions repeatedly map to the same bank.
- Declare the tile as 32×33. The extra column changes the row stride, so successive column values rotate across banks.
- Keep global coordinates based on the logical 32×32 tile; padding exists only in shared storage. Synchronize after the cooperative load before transposed reads.
- Compare bank-conflict metrics, shared-memory throughput, occupancy, and end-to-end time. The extra 128 bytes per tile may affect residency for already-large blocks.
Conclusion: Padding changes the physical row stride from 32 to 33 words, so column accesses distribute across banks instead of repeatedly selecting one bank. The logical 32 × 32 contents remain unchanged. The measured gain should be compared with the small increase in shared-memory capacity.
Common errors
- Assuming every repeated bank is a conflict. Broadcast of a common address can be efficient; different addresses in one bank are the problematic pattern.
- Padding without recalculating shared capacity. A small per-row change can cross a block-residency threshold in a large multistage tile.
- Forgetting the second phase barrier. Preventing a read-before-write does not automatically prevent overwrite-before-last-read in a loop.
Reference summary
Shared memory lets one block load a tile through coalesced global accesses, reuse it across threads, or read it in a different order. A barrier is required when one thread or warp consumes values produced by another.
Shared memory is divided into banks. Requests to different banks can proceed concurrently. Requests for different words in the same bank can be serialized, while requests for one common word can use broadcast behavior. Calculate complete word addresses before classifying a conflict.
__shared__ float tile[T][T + 1];
// +1 padding changes the bank mapping for column access
tile[threadIdx.y][threadIdx.x] = in[row * width + col];
__syncthreads();
out[col * height + row] = tile[threadIdx.x][threadIdx.y];- Coalesce global reads and writes on the two sides of a layout transformation.
- Pad or swizzle tiles when column access creates bank conflicts.
- Double-buffer or pipeline tiles only after the simple tiled kernel is correct and measured.
- Count shared-memory bytes per block; larger tiles can reduce resident blocks.
Knowledge check
Why does adding one padding element to a square shared-memory tile change performance without changing its mathematical contents?
Show the answer guide
- A square shared-memory tile whose row stride is a multiple of the bank count can map many lanes in a column access to the same bank.
- Adding one padding element changes the row stride and therefore changes the bank index of successive rows without changing the logical matrix values.
- The padded layout can turn one highly conflicted warp request into fewer serialized bank transactions; shared-memory bank-conflict metrics and timing test the mechanism.
Expose a shared-memory bank conflict
Implement a tiled transpose or column-read microbenchmark with an unpadded square tile and with one extra shared-memory column. Keep global-memory input and output identical.
Evidence to keep: Kernel source, shared-memory address-to-bank calculation for one warp, correctness results, raw timings, and profiler conflict or transaction measurements.
Shared memory organizes cooperation across threads. Registers hold each thread's immediate state and determine how much independent work can remain live.
Sources and further reading
The lesson explanations synthesize primary papers and official references. Follow the originals for proofs, exact APIs, architecture details, and version-sensitive behavior.