コアレスアクセスとストライドアクセス

Coalescing depends on which address each thread touches. The map from thread to address decides whether one transaction serves the whole warp.

When neighboring threads read neighboring addresses, the warp covers one contiguous run. That tidy layout is a coalesced access.

Thread 0 hits index 0, thread 1 hits index 1, and so on. All 32 addresses fall in one aligned line, so the GPU needs a single transaction.

A stride is a fixed gap between the addresses successive threads touch. The moment that gap grows past one element, coalescing starts to break.

With a stride of 2, thread 0 reads 0, thread 1 reads 2, thread 2 reads 4. The warp now spans twice the memory, so it needs more transactions.

Double the stride and you roughly double the lines touched. Big strides can push a warp toward many transactions while using only a sliver of each line.

Giving each thread a whole column of a row-major matrix creates a huge stride. It looks neat in code but quietly scatters every warp.

Often the fix is just swapping which index varies fastest with the thread. Let consecutive threads walk consecutive memory and the reads coalesce again.

Even a coalesced pattern suffers if it starts mid-line. A misaligned offset shifts the warp across a boundary, splitting one read into two.

To judge a pattern, do not picture one thread. Picture all 32 lanes at once and ask how many lines their addresses cover together. 🔍

Make the fastest-changing index follow threadIdx.x. That single habit keeps most of your global reads coalesced for free.

Pick the access pattern that coalesces best.

You saw that coalesced reads keep neighbors together while a stride scatters them across lines. Keep threadIdx.x driving the fastest index. 🎉