CUDA Memory Coalescing: what to measure and why

The mechanism

What cuda memory coalescing means

When a warp executes a global load or store, the hardware examines the addresses requested by its threads and combines them into the transactions needed for those addresses. Current CUDA guidance discusses global memory in terms of 32-byte transactions. Consecutive threads accessing consecutive 4-byte words can make productive use of each fetched segment; a large stride can make each thread require a different segment while using only a small part of it.

• Coalescing is a property of the warp’s address pattern for a specific memory instruction.
• Useful bytes divided by transferred bytes is a practical way to reason about efficiency.
• Alignment, element size, active lanes, cache behavior, and compute capability affect the observed transactions.
• A kernel can have high occupancy and still be limited by inefficient global-memory access.

A minimal pattern

Change one thing at a time

The snippets isolate the address or launch pattern being discussed. A real optimization still has to preserve output and be measured in the complete kernel.

const int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < count) {
  output[i] = input[i * stride];
}

const int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < count) {
  output[i] = input[i];
}

Profiler evidence

What to measure

Check correctness and required data semantics before changing the layout. Then compare kernel duration, bytes requested, sectors or transactions, memory throughput, and global-load or global-store efficiency. Source correlation is especially valuable because a kernel may contain several memory instructions with very different access patterns.

Reason about the result

How to interpret the measurements

Fewer transactions for the same useful work usually indicate better coalescing. Runtime improves when global-memory service was important enough to limit the kernel. If transaction efficiency improves without a runtime change, another limit—such as instruction throughput, synchronization, latency elsewhere, or insufficient problem size—may dominate the measured run.

Avoid false conclusions

Common mistakes

• Changing an index expression without preserving the algorithm’s required data mapping.
• Assuming every non-unit stride has the same cost on every architecture and data type.
• Looking only at aggregate bandwidth instead of the specific source instruction causing excess transactions.
• Optimizing one load while introducing scattered stores or additional synchronization elsewhere.