TL;DR: A warp's 32 threads execute a load together; the memory system services it in 32-byte sectors (four per 128-byte cache line). If the 32 threads read 32 consecutive floats (128 bytes), the request touches 4 sectors and every byte fetched is used: 100% efficiency. If thread i reads element i × 2 (stride 2), the same instruction touches 8 sectors and uses half of each byte fetched: 50%. At stride 32 floats (128 bytes) each thread's load lands in its own sector: 32 sectors for 128 useful bytes, 12.5% efficiency, and the kernel runs at an eighth of the bandwidth. The classic case is a row-major matrix accessed down a column, where consecutive threads step by the row length; swapping which index the thread id maps to fixes it. In Nsight Compute, the number to read is sectors per request (ideal 4 for 4-byte loads) and the L1 and L2 "bytes requested versus bytes transferred" ratio.
How to approach it
Describe the mechanism (a warp's load becomes sector requests) and give the three cases with numbers. Then the matrix example with code. Then the profiler metrics and what values mean. Close with the fixes and the reversal (when strided is unavoidable and what to do).
A strong answer
A typical situation: a kernel that reads a 4,096 × 4,096 float matrix column-wise runs at 400 GB/s on an H100 rated at 3.35 TB/s; the profiler shows 32 sectors per request; a one-line index swap brings it to 2.8 TB/s.
The mechanism, with the arithmetic:
a warp: 32 threads issue one 4-byte load each → 128 bytes wanted
memory: served in 32-byte sectors; the ideal is 4 sectors for the 128 bytes
case A, contiguous: thread t reads a[base + t]
addresses span 128 contiguous bytes (aligned) → 4 sectors → 128 bytes fetched, 128 used → 100%
case B, stride 2: thread t reads a[base + 2t]
addresses span 256 bytes → 8 sectors → 256 fetched, 128 used → 50%; twice the traffic per useful byte
case C, stride 32 (a column of a 32-wide row-major matrix, or any stride ≥ 8 floats): thread t reads a[base + 32t]
each address is in a different sector → 32 sectors → 1,024 fetched, 128 used → 12.5%; 8× the traffic
case D, misaligned contiguous: a[base + t] with base not a multiple of 32 bytes
128 bytes spanning 5 sectors → 160 fetched → 80%; a minor cost, worth aligning allocations to avoid
sanity: the bandwidth a strided kernel achieves is the peak times the efficiency: 3.35 TB/s × 12.5% ≈ 420 GB/s,
which is the scenario's number.
Memory Coalescing has the sector model and the cache-line details; GPU Memory Hierarchy explains where the sectors come from (L2 and HBM granularity).
The matrix example:
// C[i][j] = A[i][j] * 2 for an N×N row-major matrix (A[i*N + j])
// version 1: thread id maps to the ROW; consecutive threads in a warp read consecutive rows → stride N
__global__ void scale_bad(const float* A, float* C, int N) {
int i = blockIdx.x * blockDim.x + threadIdx.x; // row
if (i < N)
for (int j = 0; j < N; ++j) C[i * N + j] = 2.0f * A[i * N + j]; // warp: 32 rows, same j → 32 sectors
}
// version 2: thread id maps to the COLUMN; consecutive threads read consecutive elements of one row → contiguous
__global__ void scale_good(const float* A, float* C, int N) {
int j = blockIdx.x * blockDim.x + threadIdx.x; // column
if (j < N)
for (int i = 0; i < N; ++i) C[i * N + j] = 2.0f * A[i * N + j]; // warp: 32 columns, same i → 4 sectors
}
// version 3 (2D grid): both indices from thread ids, x fastest → contiguous, and no loop
__global__ void scale_2d(const float* A, float* C, int N) {
int j = blockIdx.x * blockDim.x + threadIdx.x;
int i = blockIdx.y * blockDim.y + threadIdx.y;
if (i < N && j < N) C[i * N + j] = 2.0f * A[i * N + j];
}
The rule behind it: the fastest-varying thread index (threadIdx.x) should map to the fastest-varying memory index (the last dimension in row-major layout). Version 1 has each warp touching 32 rows at the same column, stride N floats, 32 sectors per request. Version 2 has each warp on 32 consecutive columns of one row, 4 sectors. Version 3 is the idiomatic 2D form with a 32 × 8 block (256 threads) whose x dimension is a warp reading 128 contiguous bytes.
The profiler view:
Nsight Compute, Memory Workload Analysis section:
"L1/TEX: sectors per request" for global loads: 4.0 is ideal for 4-byte accesses (8.0 for 8-byte, 16 for float4);
version 1 shows 32.0
"Memory throughput" vs "DRAM throughput %": a kernel at 12% of peak DRAM bandwidth with high L2 sector traffic is
fetching bytes it does not use
"L2: bytes requested / bytes transferred" (or the "global load efficiency" derived metric in older tools): the
fraction of fetched bytes the warp actually used; 12.5% in version 1
the source view: per-line sectors per request, which points at the exact load
Nsight Systems shows only that the kernel is slow; the sector counts are Compute's job
Profiling with Nsight walks the sections; Roofline Model is where the 12.5% shows up as a kernel far below the bandwidth roof at low intensity.
The fixes, in order of preference: change the thread-to-data mapping (free); change the data layout (store the matrix transposed, or as structure-of-arrays instead of array-of-structures, so the access becomes contiguous); use shared memory as a staging tile (read a tile contiguously, then access it in the pattern the algorithm needs, which is the transpose kernel's trick; Shared Memory and Bank Conflicts covers what that introduces); vectorize with float4 when access is contiguous, so each thread asks for 16 bytes and a warp asks for 512 in 16 sectors per request, which reduces instruction count and helps reach peak on wide buses.
The reversal condition: a gather with random indices (an embedding lookup, a sparse operation) cannot be coalesced; its efficiency is set by the data, and the design responses are to sort or bucket indices before the gather, to make the gathered rows at least 128 bytes so each row is its own efficient request, and to accept that the kernel is bound by sector traffic rather than by bytes used.
What interviewers probe next
- "Does the same apply to stores?" Yes; a warp's store becomes sector writes, and partial sectors cost a read-modify-write at L2; strided stores are as bad as strided loads.
- "What about 8-byte and 16-byte loads?" A warp of 16-byte loads wants 512 bytes in 16 sectors; still 100% efficient if contiguous; the ideal sectors per request scales with the access width.
- "How does the L1 cache change this?" Reuse within a block can hit L1 and avoid L2 traffic, but the sector granularity is the same; a strided pattern with no reuse gets no help.
- "Can the compiler fix it?" No; the mapping from threads to addresses is the program's; the compiler can vectorize adjacent accesses within a thread, not reorganize across threads.
Common mistakes
- Mapping
threadIdx.xto the row of a row-major matrix. - Reading "GPU utilization" or SM activity and concluding the kernel is fine while it fetches 8× the bytes it uses.
- Fixing coalescing with shared memory and introducing bank conflicts of the same magnitude.
- Ignoring alignment on hand-computed offsets.
Key takeaways
- A warp's load becomes sector requests; 4 sectors per 128 bytes is ideal; efficiency = useful bytes ÷ fetched bytes.
- Stride 2 halves efficiency; stride 32 floats gives 12.5% and an 8× slowdown.
- Fastest thread index to fastest memory index; swap the mapping or the layout; stage through shared memory when the algorithm needs the other order.
- Read sectors per request and the requested-versus-transferred ratio in Nsight Compute.