Part I · the forcing function

Data layout and its algebra

The tensor core reads operands from a fixed set of bytes in a fixed arrangement. Before we can feed it, we need a precise language for "which logical element lives at which physical address" — and a fix for the bank conflicts that arrangement creates. That language is a shape : stride layout algebra, and this lesson is the whole grammar.

The same numbers, an order of magnitude apart

Start with the fact that motivates everything: the same numbers, in a different physical arrangement, can run an order of magnitude apart on the same GPU. A logical index — "give me element (i, j)" — says nothing about where that element physically sits. The hardware, on the other hand, is acutely sensitive to placement. Three questions decide whether a kernel flies or crawls, and all three are about arrangement, not values:

• Do 32 lanes' loads coalesce into one memory transaction, or scatter into 32?
• Do those 32 loads hit distinct shared-memory banks, or collide on one?
• Does the tile even match the byte arrangement the tensor core can read, so it lands in the right registers without a re-shuffle?

The coalescing question is covered in the basics (gpu_kernels/04). This lesson is about the other two — and to talk about them precisely we need a way to name an arrangement.

A layout is shape : stride

The entire algebra is one definition. A layout is a pair of tuples — a shape (the logical extents) and a tuple of strides (how far in memory you move per step of each coordinate). The map from a logical coordinate to a physical offset is just a dot product:

address(index) = index · strides = Σ_d index_d × stride_d

This is exactly CuTe's Layout (the shape:stride object at the heart of CUTLASS 3.x) and its coordinate-to-offset map. Take a 4×4 row-major matrix, written S[(4,4):(4,1)] — shape (4,4), strides (4,1):

The shape tells you the iteration space; the strides tell you the memory walk. Change the strides and you change the physical arrangement without touching a single value. That is the whole point: layout decouples the logical view from the physical bytes.

You already use this — it is PyTorch .stride()

If you have ever printed a tensor's strides, you have read a layout. The correspondence is exact:

>>> t = torch.arange(12).reshape(3, 4)
>>> t.stride()
(4, 1)                      # == CuTe  S[(3,4):(4,1)] , row-major

>>> tt = t.permute(1, 0)    # logical transpose → shape (4, 3)
>>> tt.stride()
(1, 4)                      # == CuTe  S[(4,3):(1,4)]
>>> tt.data_ptr() == t.data_ptr()
True                        # SAME bytes — nothing was copied

This is the cleanest demonstration of the algebra in a tool you already know. A transpose is a layout change, not a copy: permute swaps the entries of the stride tuple and leaves the data pointer untouched. The numbers never move; only the rule for reading them does. (This is also why a transposed tensor is non-contiguous — its strides are no longer descending — and why .contiguous() is the operation that actually pays for a copy.)

Tiles are just nested layouts

Tiling needs no special "tile" concept. To cut an 8×8 matrix into 2×4 tiles, you simply split each index into an outer coordinate (which tile) and an inner coordinate (where in the tile), and let the dot-product rule do the rest:

Read the strides as "how many elements does one step of this coordinate cost": stepping to the next tile-row jumps 16 elements (two full rows of 8), stepping a column-within-tile jumps 1. Tiling expressed with no special machinery — just split the index into outer and inner coordinates and compute strides for each. Every tensor-core kernel in this track is, underneath, a stack of these nested layouts: a grid of CTAs, a tile per CTA, a fragment per warp, a few registers per lane.

Naming the hardware axes

So far the coordinates have been ordinary matrix indices. The move that makes the algebra describe hardware is to let a coordinate name a physical resource. The course uses a small set of named axes; each maps to a CuTe thread-layout or value-layout coordinate:

A stride tagged with an axis says which physical thing moves when that coordinate increments. A tensor-core register fragment — the operand as it actually lives spread across a warp's registers — is then just a layout whose strides are tagged with lane and register axes:

The two @laneid strides say how the 8 rows and 4 columns of the fragment scatter across the 32 lanes; the @reg stride says which register slot inside each lane holds the second element. This is the same idea as the lane-to-element map you saw for ldmatrix in the classic tensor-core lesson (gpu_kernels/09) — only now it is written in the same notation as a plain matrix, so a kernel and its hardware fragment speak one language. In CuTe these are the thread-layout and value-layout halves of a TiledMMA.

The bank-conflict tension — and the swizzle that breaks it

Now the second of our three placement questions. Shared memory is physically split into 32 banks, each 4 bytes wide. The hardware serves one 4-byte word per bank per cycle, so the 32 lanes of a warp can read 32 words in one shot only if those words fall in 32 distinct banks. The address-to-bank rule is simply bank = (address / 4) mod 32; 32 consecutive words span all 32 banks, and the 128-byte stretch they occupy is one bank line.

Here is the tension that no plain affine layout can escape. A layout laid out for efficient row access — rows contiguous in memory — puts the elements of a single column at a fixed byte stride apart. If that stride is a multiple of the 128-byte bank line, every element of the column lands in the same bank, and the 32 accesses serialize. A layout tuned for columns has the mirror-image problem on rows. But a matrix tile must be read both ways — rows for one operand, columns for the other — so one affine arrangement always loses on one of the two.

The worked conflict

Take an fp16 (8, 64) tile in shared memory, row-major. One row is 64 × 2 bytes = 128 bytes = exactly one bank line. So row r begins at byte 128r, and column c of every row sits at the same offset-within-line. Walking down a column therefore returns to the same bank on every row → an 8-way conflict, serialized into 8 cycles instead of 1. The brief's example is the same shape at a different width — an fp16 (8, 64) tile — and it conflicts for exactly this reason.

The fix is a swizzle: a non-affine permutation of the address that scrambles columns differently on each row so that both row and column accesses spread across banks. The canonical form is XOR-with-row:

physical_col = logical_col XOR row

(Unicode ⊕ is the same operation.) Because each row XORs by a different value, the column that used to land in one bank on all 8 rows now lands in 8 different banks — conflict gone — while a row read, which XORs all its columns by the same constant, is merely permuted within the line and stays conflict-free. One cheap bit-twiddle satisfies both access patterns at once.

CuTe exposes this as Swizzle<B, M, S> and the named modes SWIZZLE_32B / SWIZZLE_64B / SWIZZLE_128B — each describes a swizzle over 8 atoms of N bytes. The rule of thumb: pick the largest atom the contiguous row dimension can fill. An 8 × 128-byte atom covers a full bank line, which is why SWIZZLE_128B gives conflict-free access to 8 rows and 8 columns at once in fp16 — the case the widget below makes concrete. Crucially, the swizzle is composed on top of the affine layout, not baked into it (CuTe composition(Swizzle<B,M,S>, layout)): the shape:stride part stays a clean dot product, and the swizzle is a separate address transform applied after it.

Widget · the bank-conflict visualizer

An fp16 8×N tile mapped onto the 32 shared-memory banks. The lit cells are the 8 elements of one column access — the pattern a tensor-core load performs constantly. Toggle the swizzle and watch where those 8 accesses land. With swizzle off, an aligned column repeats one bank 8× — an 8-way conflict, serialized. With SWIZZLE_128B (XOR-with-row), the 8 accesses fan out across 8 distinct banks — conflict-free. The conflict factor is computed live as max accesses to any single bank.

Two things to notice as you slide the column with swizzle off: the conflict factor stays pinned at 8× for every aligned column when N = 64 (one row = one bank line), and the cost in cycles tracks it one-for-one. Flip to SWIZZLE_128B and the factor collapses to 1× across the whole walk — the XOR has spread each column's 8 elements across 8 banks. Widen N and the row no longer divides the bank line evenly, which changes the unswizzled collision pattern; the swizzle keeps it conflict-free regardless. That is the entire payoff of the non-affine address transform, made visible.

TMA does the layout and the swizzle in one shot

The async copy engine we meet in lesson 04 (TMA) folds all of this into a single descriptor. A TMA tensor map is multi-dimensional: one 3D box can describe both the atom tiling (how the rectangular tile is cut) and the swizzle within each atom. So a single TMA load lays the tile out in shared memory and swizzles it as it writes — there is no separate swizzle pass, no extra kernel, no shuffle instructions. The arrangement we just spent the lesson naming becomes a field in a descriptor that hardware honors during the copy.

Layout is part of the instruction interface

Here is the punchline that echoes through the rest of the track. Layout is part of the instruction interface, not decoration. Three separate things must describe the same physical arrangement of the same bytes:

If any one of these disagrees with the others, the hardware still runs — there is no exception, no segfault, no compiler error. It simply reads the wrong bytes, or reads the right bytes slowly. A swizzle mismatch between the TMA descriptor and the MMA descriptor produces a kernel that finishes and returns a confidently wrong tensor. This is why layout bugs are the hardest class of bug in the whole paradigm, and why we built the algebra first: every async mechanism in Part II is a contract written in this language, and the contract only holds if all parties name the same layout.

The next bottleneck

We can now describe any arrangement of bytes precisely — affine or swizzled, in memory or scattered across a warp's registers. But describing a layout and knowing which layout the hardware demands are different things, and the demand has not been stable. Each tensor-core generation since Volta keeps the invariant D = A·B + C while changing how the instruction is issued, what operand layout it requires, and even where the accumulator lives. The next lesson asks: now that we can describe any layout, which layout does each generation's tensor core actually demand — and why did the answer keep changing?