Stride, Misalignment, and the Cost of a Bad Layout

Threads are cheap and memory is not. A kernel can add two numbers per thread per cycle without noticing, but the same kernel can lose most of its bandwidth to an access pattern that looks, in the source, like ordinary array indexing. The pattern is a property of how the loop index maps onto the array index, and it is decided long before the kernel runs: in the choice of whether a structure is stored as an array of records or a record of arrays, in whether a matrix is kept row-major, and in where the alignment of a buffer happens to fall.

What a stride costs, in sectors

The memory system moves 32-byte sectors and forms a warp request from the addresses that thirty-two threads supply at one instruction. The number of sectors touched therefore decides the cost, and the stride decides the sectors. A warp loading one 4-byte word per thread at stride s bytes touches a number of sectors that grows with s until the stride exceeds the array footprint, at which point every thread is in its own sector.

Stride between neighbouring lanes	Sectors touched per warp	Bytes moved per 4 useful bytes
4 bytes, contiguous	4	1
8 bytes	8	2
16 bytes	16	4
32 bytes or more	32	8

The right column is the over-fetch factor. It is a lower bound: it counts only what a single request costs, and a kernel whose accesses are strided usually also has fewer requests in flight per unit of bandwidth, which compounds the loss.

Two observations follow. First, the unit matters more than the index arithmetic: a stride of two floats and a stride of two doubles are different patterns in bytes, so the table above counts bytes because the hardware does. Second, the table stops mattering once the stride reaches a full sector. Beyond 32 bytes per lane, the added cost comes from row activations and from the address distribution across memory partitions rather than from sector count, so doubling the stride beyond that point is roughly free in transaction terms.

Misalignment splits a run across lines

A fully contiguous warp access of 128 bytes costs four sectors provided the run begins on a 128-byte boundary. If it begins at an offset that is not a multiple of the sector size, the run spans five sectors and two cache lines, because the first and last sectors are each only partly used. The over-fetch is 25 percent on a fetch that covers exactly one line, and it grows worse for smaller requests: a warp reading 32 bytes starting at an offset of 4 bytes spans two sectors instead of one, doubling the cost of an access that had no stride problem at all.

Alignment on the device is usually not the programmer’s fault. The allocation call returns addresses aligned far beyond any scalar requirement, and the offsets introduced by pointer arithmetic are the common cause. Slicing a buffer, passing base + offset to a kernel, viewing a struct field as an array, or packing two arrays into one allocation with a hand-computed offset are all ways to start a warp access at an address that is not a multiple of the sector size. The same effect appears in strided multi-dimensional indexing, where the row length in bytes is not a multiple of the 32-byte sector.

The fix is arithmetic rather than architectural: make the row length, the slice offset, and the packed allocation sizes multiples of the alignment the access pattern wants. When the layout is fixed by an external format, the remedy is to copy into a padded staging buffer once rather than pay the misalignment on every access.

The row-major trap

Consider a matrix of R rows and C columns stored row-major, so that element (r, c) lives at r * C + c. A kernel that assigns one thread per element and maps threadIdx.x to the column index produces perfect coalescing: adjacent lanes read adjacent columns of one row. A kernel that instead maps threadIdx.x to the row index reads down a column, where the stride between lanes is C elements. If C is 64 and the elements are 4 bytes, every lane is in a different 256-byte region, so a warp of 32 lanes touches 32 sectors for 128 useful bytes — an eight-fold over-fetch from an index swap that changes nothing about the arithmetic.

The mistake is easy to make because it is invisible. The kernel computes the right answer in both cases; only the addresses differ. It appears whenever the natural parallel axis of the algorithm is orthogonal to the storage axis: reading a column of a row-major matrix, gathering the same field from an array of records, or accessing one node of an array-of-structures per lane.

There are three ways out, and they differ in where the cost lands.

The cheapest, when it is available, is to change the traversal order. If the algorithm permits assigning lanes to the contiguous axis and iterating over the other one inside the kernel, the stride disappears with no data movement at all. This is why matrix kernels usually parallelize over columns and loop over rows rather than the reverse.

The second is to change the layout once, on the host or in a separate preparation step, to struct-of-arrays. Storing each field in its own dense array converts an interleaved gather into a contiguous read, at the price of a transformation pass over the data and of any algorithm that benefits from record locality. If the same field is read many times, the pass amortizes. If it is read once, the transformation costs more than it saves.

The third is to transpose through a per-block scratchpad, which absorbs the strided read at shared-memory bandwidth and emits contiguous global writes.

Transposing through a scratchpad

The shared-memory transpose exploits the fact that a strided read from global memory costs sectors, while a strided read from shared memory costs bank conflicts, and the two have different units. Shared memory is banked into 32 four-byte-wide banks, and a warp reading 32 consecutive words touches all 32 banks exactly once, which is a conflict-free access. A warp reading down a column of a padded tile is also conflict-free if the padding makes the column stride odd with respect to the bank count.

#define TS 32
#define PAD 1

// Launched with blockDim = (32, 8) over a square n by n matrix.
__global__ void transpose(const float *in, float *out, int n) {
    __shared__ float tile[TS][TS + PAD];   // +1 shifts each row by one bank
    int x = blockIdx.x * TS + threadIdx.x;
    int y = blockIdx.y * TS + threadIdx.y;

    for (int j = 0; j < TS; j += 8)        // lanes vary x: coalesced read
        tile[threadIdx.y + j][threadIdx.x] = in[(y + j) * n + x];
    __syncthreads();

    // Swap the block coordinates so lanes again run along the contiguous axis.
    x = blockIdx.y * TS + threadIdx.x;
    y = blockIdx.x * TS + threadIdx.y;
    for (int j = 0; j < TS; j += 8)        // lanes vary x: coalesced write
        out[(y + j) * n + x] = tile[threadIdx.x][threadIdx.y + j];
}

The read phase writes a 32 by 32 tile into shared memory with lanes varying along x, which is the contiguous axis of the input, so global reads coalesce. The write phase reads the tile in transposed order and writes with lanes varying along x of the output, having swapped which block coordinate indexes the outer axis. The padding of one extra column spreads the transposed read across all 32 banks; without it every lane in a step would land on a single bank and the warp would serialize into 32 passes. The exact conflict mechanics belong with the banked scratchpad itself; what matters here is that the scratchpad is used as a staging buffer whose access cost is measured in conflicts rather than sectors.

What the transformation itself costs

Every layout fix buys bandwidth with something else. A change in traversal order costs nothing but is limited to algorithms whose parallel axis can be reassigned. A host-side or kernel-side repack costs one pass over the data, which for a read-once workload is a net loss and for a read-many workload is a net win that grows with the reuse factor. A shared-memory transpose costs a scratchpad sized to the tile — a few kilobytes for a 32 by 32 float tile, tens of kilobytes for a 128 by 128 one — which reduces the number of blocks an SM can host and therefore occupancy; it also costs a barrier, which serializes the block, and shared memory traffic that competes with whatever else the kernel does with the scratchpad.

The honest summary is that a stride problem should be solved at the lowest level that will accept it. Fix the loop mapping first, then consider a one-time repack if the reuse justifies it, and reach for a shared-memory staging tile when the data must be consumed in a different order than it is stored and the reuse within a block is high enough to fund the barrier. Reaching for the transpose first can cost more than the stride it removes.