100 DAYS OF GPU Course preview
Day 11Module 2
in-technical-review

CUDA memory coalescing, measured

Here is a benchmark result from the CUDA 120-day challenge, reported at https://github.com/AdepojuJeremy/CUDA-120-DAYS--CHALLENGE/issues/4 (checked 2026-08-29):

Coalesced Kernel Time:                0.195584 ms
Non-Coalesced Kernel Time (stride=2): 0.061440 ms

Strided access, three times faster. Every textbook says the opposite.

The benchmark is wrong in a useful way. Both kernels launched the same number of threads, but the strided one skipped elements. It touched half as many elements and moved half as many bytes.

This page explains coalescing and shows how to compare access patterns without changing the amount of work.

What a warp does to memory

A warp is 32 threads that issue instructions together. When those 32 threads execute a load, the hardware does not perform 32 separate loads. It looks at the 32 addresses and works out the smallest set of memory transactions that covers them.

Global memory is served in 32-byte sectors, and sectors are grouped into 128-byte cache lines. Each fetch of a sector is one transaction. A transaction fetches a sector.

When threads 0 through 31 read data[0] through data[31], those 32 floats occupy 128 contiguous bytes. The warp needs four sectors and four transactions, and it uses every byte fetched.

Now suppose thread 0 reads data[0], thread 1 reads data[8], and thread 2 reads data[16]. The addresses are 32 bytes apart, so each lane uses its own sector.

The warp needs 32 transactions instead of 4. Each transaction moves 32 bytes to provide the 4 bytes that one lane requested, so seven eighths of the bandwidth carries unused data.

Push the stride to 32 floats and the transaction count does not get worse, because 32 is already one per lane. What changes is that the lanes now land in 32 different 128-byte lines instead of eight, which costs you in the cache rather than at the bus.

Coalescing is not a compiler option. It describes how the addresses from 32 lanes combine into memory transactions.

The rule, and the exception people expect

The rule: consecutive threads should touch consecutive addresses. data[globalIndex] is right. data[globalIndex * stride] is not.

Giving each thread a contiguous chunk looks sensible because it often suits CPU threads. On a GPU, it makes one warp access strided addresses.

// The CPU-correct, GPU-wrong pattern
size_t base = threadId * chunkSize;
for (int k = 0; k < chunkSize; ++k) {
    out[base + k] = in[base + k];
}

On a CPU with a private cache per core, each core can stream its own region. On a GPU, the 32 warp lanes execute the same k, so their addresses are chunkSize elements apart. A chunkSize of 32 gives each lane its own sector.

The fix is the grid-stride loop from day 8, where consecutive threads stay adjacent on every iteration.

Measure the same amount of work

The full program is in code/day11-coalescing/coalescing.cu. Its checks make the comparison valid.

Three rules the program follows:

Every variant moves the same bytes. The strided kernel performs n reads and n writes, just like the coalesced kernel. The program computes bandwidth from one constant, not once per kernel.

Time with events, not the clock. Kernel launches are asynchronous. A host-side timer around a launch measures the launch. Day 9 covers this.

Warm up first. The first launch of each kernel pays a module load cost, because CUDA_MODULE_LOADING=LAZY has been the default since CUDA 12.2. Warm up every kernel you intend to time, not just the first one.

Here are the kernels:

// One kernel serves every stride, so the instruction mix is identical across
// rows and only the address pattern differs.
//
// The launch is a 2D grid: x walks `m = n / stride` columns, y walks `stride`
// rows. Element index j = col * stride + row. Sweeping col over [0, m) and
// row over [0, stride) hits every j in [0, n) exactly once, so this is a
// permutation, not a sampling.
//
// Consecutive threads in a warp have consecutive `col`, so their addresses
// are `stride` elements apart. At stride 1 that is 4 bytes apart, which is
// perfectly coalesced. At stride 32 it is 128 bytes apart, so every lane
// lands in its own 32-byte sector.
//
// There is no modulo and no division here. An earlier version used `% n` on a
// runtime size_t, which nvcc cannot strength-reduce to a mask, so the strided
// rows paid for a 64-bit division the coalesced row did not. That broke the
// claim that the access pattern is the only variable.
__global__ void copyPermuted(const float* __restrict__ in,
                             float* __restrict__ out, size_t m, int stride) {
    size_t col = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
    if (col < m) {
        size_t j = col * static_cast<size_t>(stride) + blockIdx.y;
        out[j] = in[j];
    }
}

The first version of that kernel exposed two benchmark errors.

The first index was j = (i * stride) % n. It stays in range, but it is not a permutation when n is a power of two.

At stride 16, it touches half the buffer's 32-byte sectors. At stride 32, it touches one quarter. Those rows moved half and one quarter of the traffic that the program reported.

This is the same error as the benchmark at the top of the page. The 2D launch above forms a permutation at every stride, so each row accesses every element exactly once.

The first version also ran a 64-bit % n that the coalesced kernel did not. Because n arrives as a runtime size_t, the compiler cannot replace the operation with a mask. The current kernel uses no division.

Results

Measured on a Tesla T4, driver 595.84, CUDA 12.6 (V12.6.85), 64 Mi floats, 256 MiB per buffer, 10 timed runs after 3 warm-ups. Repeat runs agree closely on every row except stride 1, which moves a few percent between runs as the clocks ramp on a kernel this short.

Pattern Time (ms) GB/s Fraction of stride 1
stride 1 (coalesced) 2.305 232.9 1
stride 2 6.158 87.2 1/2.7
stride 4 12.836 41.8 1/5.6
stride 8 26.701 20.1 1/11.6
stride 16 37.728 14.2 1/16.4
stride 32 56.037 9.6 1/24.3
chunked, 32 per thread 16.493 32.6 1/7.1

Every row moved 536,870,912 bytes. The program prints that total, and the --json output records it for each row.

Two things here are not what the theory predicts

The curve does not flatten at stride 8. A 32-byte sector holds eight floats, so from stride 8 onward each lane already uses one sector. The transaction count stops growing, but bandwidth still falls from 20.1 to 14.2 to 9.6 GB/s.

Past stride 8, the lanes spread across more 128-byte cache lines, so the L2 hit rate drops. At stride 32, they also spread across more DRAM pages, which adds row-buffer misses and TLB pressure.

The sector model predicts transactions, not bandwidth. Other parts of the memory system account for the rest of the curve.

The profiler settles it

Nsight Compute counts the two things that separate the halves of that claim. l1tex__t_requests_pipe_lsu_mem_global_op_ld.sum is how many warp-wide load instructions issued. l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum is how many 32-byte sectors those instructions actually pulled.

Stride Requests Sectors Sectors per request GB/s
1 2,097,152 8,388,608 4
2 2,097,152 16,777,216 8
4 2,097,152 33,554,432 16
8 2,097,152 67,108,864 32 20.1
16 2,097,152 67,108,864 32 14.2
32 2,097,152 67,108,864 32 9.6

The request count never moves. The access pattern cannot change how many instructions issue, only what each one costs.

At stride 1 a warp's 32 lanes span 128 bytes, which is four sectors, and the profiler says four. Sectors per request then doubles with the stride until stride 8, where it pins at 32 and stops. That is the saturation the sector model predicts, measured directly.

Now look at the bandwidth column across those last three rows. Sectors are identical at 67,108,864. Bandwidth falls from 20.1 to 14.2 to 9.6.

The profiler shows that transactions stop growing at stride 8, while the timer shows that the kernel keeps slowing down. A sector count cannot show the lost locality from extra 128-byte lines, DRAM pages, and TLB pressure.

The measured loss is worse than the model even where the model applies. At stride 2 the model predicts half; the T4 delivers a third. At stride 4 it predicts a quarter and delivers a sixth. Treat 1/min(s, 8) as a ceiling on efficiency, not an estimate of it.

Chunked access is faster than stride 32. When each thread walks 32 contiguous elements, the warp's lanes are 32 elements apart at any instant. Yet stride 32 measures 9.6 GB/s, while chunked access reaches 32.6 GB/s.

The chunked kernel reuses cache lines. A whole warp covers 1024 contiguous elements across its loop, so later iterations can use data from lines fetched earlier. The strided kernel does not revisit those lines.

Chunked access is still about seven times slower than coalesced access. Its cache reuse makes it less costly than the per-instruction stride suggests.

Do not memorise these numbers. The GPU, driver, and problem size will change them. The order should hold: coalesced access is fastest, each stride hurts, and bandwidth can keep falling after the transaction count stops rising.

Run it yourself

The full program allocates 512 MiB and runs seventy timed launches across seven patterns, which is too much for Compiler Explorer's 20 second cap. Use Colab or a local GPU for the full run.

A cut-down version, one coalesced and one strided kernel over 4 Mi elements, finishes in about two seconds and fits fine. Target sm_75 or lower: the Compiler Explorer runner is a Tesla T4.

Exercise

Measure the bandwidth curve across strides 1, 2, 4, 8, 16 and 32, then explain its shape in two sentences.

Check: the harness verifies that every measurement moved the same number of bytes, then checks that your stride-1 result is at least eight times your stride-32 result. It reports the ratio; on the reference T4 it is 24.3x, so the eight times bound has real headroom and a failure means something is wrong. If the bytes-moved assertion fails, you have reproduced the course bug.

Hint 1

Think about how many 32-byte sectors one warp's 32 addresses span at each stride. A sector holds 8 floats. At stride 1 the warp spans four sectors.

At stride 8 it already spans thirty-two.

Hint 2

Count transactions first, then ask whether transactions are the only thing that costs you. Where does the transaction count stop growing, and does your measured curve stop falling at the same place?

Solution

The curve has two regimes.

Up to stride 8 the transaction count grows with the stride. A 32-byte sector holds eight floats, so at stride s the warp uses at most 1/min(s, 8) of every byte fetched. Measured, the T4 does slightly worse than that ceiling at every point: a third at stride 2 where the model says half, a sixth at stride 4 where it says a quarter.

From stride 8 the transaction count is pinned at 32, one sector per lane, and cannot grow. Bandwidth still falls: 20.1, then 14.2, then 9.6 GB/s. That drop is not transactions, it is locality.

The lanes keep spreading across more 128-byte lines and then across more DRAM pages, so L2 hit rate falls and row-buffer misses rise.

A transaction model does not predict bandwidth. Measure both when performance matters.

Pitfalls

Comparing kernels that move different byte counts. This caused the bug at the top of the page. Check the byte count before comparing bandwidth, especially when a strided kernel appears faster.

Giving each thread a contiguous chunk. Correct on a CPU, worst case on a GPU. Consecutive lanes must hold consecutive addresses at each step of the loop.

Timing with a host clock. time ./a.out measures process start, context creation, allocation and copies. Day 9 covers what each of those costs.

Forgetting the warm-up. The first launch of each kernel pays a module load. A benchmark without warm-up mostly measures loading.

Assuming coalescing is always the bottleneck. It matters when a kernel is bandwidth-bound, which most simple kernels are. A compute-bound kernel can have sloppy access patterns and not care. Day 42 shows how to tell which one you have before optimising.

Go deeper

Next

Day 12 takes the pattern you just measured and applies it to the transpose problem, where reading rows means writing columns and one side is always strided. Day 13 fixes it with shared memory.