GPUFlight CUDA coordinates, strides, and layouts
Twelve field notes on how launched threads map to logical data and physical addresses. Each note names the question to ask and a common source of mistakes.
Contents · 12 topics
Launch arithmetic
Which elements does one thread own?
Mark the start of the block, then step by one thread, one block, or one whole grid as the schedule requires.
Watch for: For a nonempty input, the last valid index is N − 1. A rounded-up launch can have extra threads.
Coordinates & ranks
When do you flatten the block and when do you flatten the image?
In a CUDA block, x changes fastest. Divide a flat rank to remove its fastest-changing axis; take the remainder to recover that axis.
Watch for: A block-major enumeration of threads and the packed row-major order of image pixels are different enumerations.
Memory layouts
How far from the base pointer is this element?
Start at the base and add coordinate × stride for each axis. Write the unit next to every stride before multiplying.
Watch for: A byte pitch belongs on a byte pointer. Adding it to a float pointer scales it again. Widen before multiplying large dimensions.
Tensor strides
Which axis actually changes the address?
Name each logical axis, attach its physical stride, then add the products. A broadcast axis contributes zero offset.
Watch for: Permuting a view changes its strides; it does not necessarily move any data.
Warp mapping
Which lanes participate, and where does a partner live?
First compute the thread rank within its block. Split it into a warp index and lane, then apply the instruction’s group and participation rules.
Watch for: The lane is not generally threadIdx.x. In a short row, a warp can span several rows.
Memory access patterns
Which addresses share a sector or a bank?
List the addresses an instruction requests. Group byte addresses by sector, or word addresses by shared bank under the stated model.
Watch for: Lanes that read one shared word may be broadcast. Distinct words in the same bank are a different case.
Tiles & shared memory
How do local tile coordinates become global coordinates?
Track the tile origin, the thread’s local coordinates, and the leading dimension separately for the input, shared tile, and output.
Watch for: Padding a shared tile changes its physical stride. It does not change the logical transpose shape.
Parallel collectives
Which partner, rank, or output slot belongs to a thread?
Draw one round of the algorithm. For placement, count qualifying items before the item, then add the group’s global base.
Watch for: Inclusive and exclusive scans differ at the current element. Compaction uses the count before a kept element.
Stencils & convolution
Which input coordinate supplies this output tap?
Map an output back through stride and padding, then add the tap displacement. Apply the specified boundary policy last.
Watch for: Clamping, wrapping and zero padding are different operations. A negative candidate is not yet a legal global address.
Sparse & irregular data
Who owns an entry when rows have different lengths?
Read each pair of adjacent offsets as a half-open interval. Separate a slot in a neighbor list from the vertex ID stored there.
Watch for: Equal adjacent offsets denote an empty row. A scatter can send several threads to one destination.
Permutations & partners
Where does a butterfly, compare pair, or packed pair go?
Write the stage rule first. A bit can select a partner, the lower member of a pair, or a sorting direction; these are separate decisions.
Watch for: A valid partner within a subgroup may need the subgroup base added back to become a global index.
Production kernels
How do multiple mappings compose into one address?
Work one space at a time: identify the owner, translate to local coordinates, resolve any indirection, then apply storage strides.
Watch for: Capacity is not a live count. Padding, circular buffers and split workspaces each have their own bounds.