TL;DR: For scalar four-byte shared-memory accesses on the stated 32-bank model, the bank is the word index modulo 32. Reading a column of a 32-word-wide tile sends 32 distinct words to one bank. A 33-word row stride spreads that column across all 32 banks. Padding costs 128 bytes per 32-row FP32 tile, or 3.125%; it removes this conflict, not necessarily a 32× whole-kernel slowdown.
How to approach it
First state the scalar bank model and distinguish different words in one bank from a same-word broadcast. Derive the column mapping before and after padding. Then inspect a small kernel with that exact access, and explain what the profiler and resource budget must confirm.
A strong answer
A transpose gains coalesced global reads by staging a tile in shared memory, but still spends time on the shared read. The addresses now need a different analysis: global sectors group nearby byte addresses into requests, while shared-memory banks determine which local resources serve the warp’s distinct words.
Derive the layout change
Shared memory is split into 32 banks, and each bank can deliver one 4-byte word per cycle. When the lanes of a warp read words that sit in different banks, they are served together. When several lanes need different words from the same bank, those reads take turns, which is a bank conflict.
NVIDIA's shared-memory guidance describes 32 banks with successive 32-bit words assigned to successive banks on compute capability 5.x and later. Multiple lanes reading the same word can be served by broadcast. Different words in one bank create a conflict under this scalar model.
row-major tile element [r][c], row stride P words:
word index = P × r + c; bank = (P × r + c) mod 32
column 0, lanes vary r from 0 to 31:
P=32 → bank 0 for every lane, with 32 distinct words
P=33 → bank r, all 32 banks used
storage: 32 × 32 × 4 = 4,096 bytes; 32 × 33 × 4 = 4,224 bytes
extra = 128 bytes / 4,096 = 3.125%
sanity: adding one word per row changes addresses, not the logical matrix values.
Shared Memory and Bank Conflicts explains the resource model. GPU Memory Hierarchy distinguishes it from cache-line alignment. A conflict degree predicts how this access must be split into conflict-free work; it is not an instruction latency in clock cycles or a whole-kernel speedup ratio.
A small kernel that exposes the access
The following example reads a fixed 32×32 input tile and outputs its first column. Its launch contract is exactly one one-dimensional block of 32 threads. The input has 1,024 floats, the output has 32 floats, and they do not overlap. Instantiate PITCH=32 for the conflicting layout or PITCH=33 for the padded one.
template<int PITCH>
__global__ void readColumn(const float* __restrict__ input,
float* __restrict__ output) {
static_assert(PITCH == 32 || PITCH == 33);
__shared__ float tile[32][PITCH];
const int lane = threadIdx.x;
for (int row = 0; row < 32; ++row)
tile[row][lane] = input[row * 32 + lane];
__syncthreads();
output[lane] = tile[lane][0];
}
Each staging instruction reads one contiguous input row across the warp. All threads reach the barrier before any reads from the staged column. The output stores are contiguous. Only the shared column-read mapping changes between the two template instances. In a full transpose, the same idea appears when exchanging row and column roles; rectangular dimensions and partial tiles also need bounds checks that do not make barrier participation divergent.
The example isolates a layout mechanism. A compiler can transform code, and the small launch is not intended to fill a GPU or establish throughput. Inspect generated instructions and the actual shared-memory requests when profiling. Verification scope: both template bodies were executed with a CPU thread/barrier harness and checked against the expected column. This does not substitute for native CUDA compilation, sanitizer checks or device timing.
Confirm the intended change, then measure its value
Nsight Compute's shared-memory table exposes requests, wavefronts and bank-conflict information. Compare actual versus ideal work for the same instruction width and active mask, then attribute excess work to the source line. Use the available metrics for the installed GPU/tool version rather than assuming one hardcoded metric name is portable.
Wider vector operations can require multiple ideal wavefronts even without a conflict. Do not call that extra necessary work a bank conflict, and do not claim that using an eight-byte type physically changes modern banks into eight-byte banks. Account for the generated instruction and its service groups. Likewise, same-address reads can broadcast, but concurrent non-atomic writes to one destination are a correctness problem.
Padding uses shared capacity and may change alignment. A one-word pad is useful for this scalar column access; it is not a universal layout for vector loads, asynchronous copies or tensor-core instructions. Swizzling can preserve a denser layout while rearranging bank ownership, but must satisfy the consuming instruction's contract. CUTLASS and Tensor Core Kernels discusses those specialized layouts.
For scalar full-warp distinct-word strides, the conflict pattern follows the stride modulo 32: stride 34 has the same bank increment as two and revisits banks after 16 lanes. Odd strides visit all banks in this model. That statement is about row stride, not a rule that every padding amount must be odd for every tile.
Use padding when this conflict is on the critical path and the new layout preserves correctness, alignment and useful occupancy. Reverse it when accesses are already conflict-free, the added shared memory reduces useful residency, or another supported layout wins the full-kernel measurement. Source arithmetic identifies the mechanism; profiling decides its importance.
What interviewers probe next
- "Why does 34 differ from 33?" Its increment modulo 32 is two, so it uses only 16 banks for 32 distinct words.
- "Is a broadcast free?" It avoids that bank conflict, but the load still performs work and has latency.
- "Why not pad every tile?" Padding can waste capacity or violate another instruction's layout constraints.
Common mistakes
- Equating a 32-way scalar conflict with 32× end-to-end runtime.
- Treating global-sector coalescing and shared-bank conflicts as the same mechanism.
- Adding bounds checks that let some threads skip a required block barrier.
Key takeaways
- For the stated scalar model, bank = word index modulo 32.
- A 32→33 word row stride spreads a column across 32 banks.
- The 128-byte pad costs 3.125%; keep it only when the measured benefit survives other resource costs.