You cannot buy 4 bytes. The shop only sells 32.

Lesson 4 · memory coalescing · the hardware bill for the strides you chose in Lesson 3

Lesson 3 ended with a scattered traversal and a promise to put a number on the waste. Here is the number. This is the single highest-leverage idea in memory-bound kernel work, and it is why two kernels that read "the same data" can differ by 8×.

The one fact the whole lesson hangs on

DRAM is not byte-addressable from the GPU's point of view. Memory is delivered in fixed-size sectors of 32 bytes. Ask for one float32 — 4 bytes — and the hardware moves 32 bytes across the bus and throws 28 of them away.

That would be catastrophic, except for the second fact: a warp of 32 threads issues its load as one instruction. The memory system looks at all 32 addresses together and issues the minimum number of sectors that covers them. So the question is never "how many bytes do I need?" It is:

How many 32-byte sectors do my 32 lanes land in? Four is perfect. Thirty-two is a disaster. Everything interesting lives between them.

Access pattern (32 lanes × fp32)	Sectors	Bytes moved	Useful
Stride 1, aligned — the happy path	4	128 B	100%
Stride 1, misaligned by 4 bytes	5	160 B	80%
Stride 2	8	256 B	50%
Stride 8 or more	32	1024 B	12.5%

Read the last row again. Same 128 useful bytes. Eight times the traffic. On a memory-bound kernel — which, per Lesson 2, is most of them — that is an 8× slowdown that no amount of arithmetic cleverness recovers.

Why a stride-8 read costs the same as stride-64

Once your lanes are far enough apart that no two share a sector, each lane needs its own. For float32 that happens at stride 8 — 8 × 4 = 32 bytes — and it never gets worse, because 32 sectors for 32 lanes is already one each. Efficiency bottoms out at element_size / 32: 12.5% for fp32, 6.25% for fp16.

This is exactly the transposed read from Lesson 3. A row of x.T with an inner stride of 1024 is not "slightly worse" than a contiguous row — it is at the floor.

What Triton actually does for you here

The Lesson 1 offsets, re-read as an access pattern rather than as indices.

Listen first: the vector-add kernel built its offsets as block_start plus tl.arange from zero to BLOCK_SIZE. That is a run of consecutive integers, so consecutive lanes get consecutive addresses, so the loads coalesce perfectly — and you never wrote the word coalescing. Compare the second form, where the offsets are multiplied by a stride of eight. Same number of elements, same amount of arithmetic, and the compiler can do nothing about it: the addresses genuinely are scattered, so the hardware genuinely must fetch thirty-two sectors.

offsets = block_start + tl.arange(0, BLOCK_SIZE)        # consecutive → COALESCED
x = tl.load(x_ptr + offsets)                            # 4 sectors per warp

offsets = block_start + tl.arange(0, BLOCK_SIZE) * 8    # stride 8 → SCATTERED
x = tl.load(x_ptr + offsets)                            # 32 sectors per warp, 12.5% useful

The Triton docs promise "automatic coalescing". That promise is real but narrow: the compiler assigns your tile's elements to lanes in the order that produces the fewest transactions for the addresses you described. It will not — cannot — invent a different access pattern. Coalescing is a property of your indexing, not a compiler feature.

Micro-world: count the sectors

One warp, 32 lanes, one load instruction

Each block below is one 32-byte sector. Orange means the hardware had to fetch it. The useful bytes never change — only the bill does.

Element type

Stride (elements): 1

Start offset (bytes): 0

32-byte sectors on the memory bus

fetched (you pay) skipped

Bytes that crossed the bus

Start at stride 1, offset 0. Four sectors, 128 bytes fetched, 128 useful. This is the only configuration you should ever be happy with.
Nudge the start offset to 4 bytes. Nothing about your data changed — but a fifth sector appears and you just lost 20% of your bandwidth to alignment alone.
Walk the stride up: 1 → 2 → 4 → 8. Watch the fetched bytes double each time while the useful bytes sit still. This is the whole lesson in one drag.
Keep going past 8. Nothing gets worse. You have hit the floor — one sector per lane. Stride 8 and stride 16 cost exactly the same, which is why "slightly strided" is a myth for fp32.
Switch to float16 at stride 1, then stride 16. The floor is lower for smaller types — 6.25%. Narrow dtypes make coalescing matter more, not less. Remember that in Lesson 5.

Check yourself

Cold-recall defense (one breath):
Memory moves in 32-byte sectors, and a warp of 32 lanes issues one load whose addresses the hardware covers with the fewest sectors it can. Stride 1 aligned = 4 sectors, 100% useful. Stride 8 for fp32 = 32 sectors, 12.5% useful — an 8× bandwidth tax on a memory-bound kernel, and the floor is element_size / 32. Misalignment alone costs one extra sector. A transposed read from Lesson 3 is the strided case. Triton coalesces the lane mapping, not your indexing — the access pattern is yours to choose.

What this means for Experiment 0001

Vector add reads two float32s and writes one per element: 12 bytes of traffic per element, for one add. Arithmetic intensity ≈ 0.083 FLOP/byte against a machine balance of 153 — about as memory-bound as code gets. So the only honest way to report that kernel is GB/s, not TFLOP/s, and the ceiling on an A100 is 2039 / 12 ≈ 170 billion elements per second. If your T4 run reports a bandwidth well under the card's spec, the first suspect is not Triton — it is the access pattern, and this micro-world is where you diagnose it.

Primary source

Read the CUDA C++ Best Practices Guide § Coalesced Access to Global Memory (~10 min). It is the source of record, with the same sector arithmetic and the misaligned-access figures. Read it after playing with the micro-world above — the diagrams will land instead of blur.

💬 I'm your teacher for this — ask me followups any time. The natural next question: "if strided access is this bad, how does matmul ever go fast?" That is Lesson 6, and the answer is reuse.