Third post in the Foundations series (previously: GPU architecture history and What is CUTLASS). The CUTLASS post ended by noting that CUTLASS 3.0 rebuilt its GEMM kernels on a new core library, CuTe, without changing the underlying device → threadblock → warp → thread hierarchy — only how each level's tile gets described. This post is about that description: what a CuTe Layout actually is.

A Layout is a pair: Shape and Stride

Strip away everything else, and a CuTe Layout is just two tuples:

  • Shape — the logical extent of each dimension (how many elements)
  • Stride — how many elements to skip in linear memory for each step along that dimension

A Layout's job is to convert a logical coordinate — (i, j) for a 2-D array — into a single linear memory offset. It does that with an inner product: multiply each coordinate component by its matching stride component, and sum them.

That's genuinely the whole idea. Nothing about "row-major" or "column-major" is a special case CuTe has to handle separately — it's just a different Stride paired with the same Shape.

Same Shape, different Stride

Take a 4×3 matrix — Shape (4, 3) either way. Store it row-major, and each step in the row dimension jumps 3 elements (the row length); each step in the column dimension jumps 1 element. Store it column-major, and it's reversed:

Two 4x3 grids showing the same logical matrix laid out in memory two different ways — row-major with Stride (3,1) and column-major with Stride (1,4) — each cell labeled with its resulting linear memory offset

// Row-major: Shape (4,3), Stride (3,1)
auto layout_row = make_layout(make_shape(4, 3), make_stride(3, 1));
// layout_row(i, j) == i*3 + j
 
// Column-major: Shape (4,3), Stride (1,4)
auto layout_col = make_layout(make_shape(4, 3), make_stride(1, 4));
// layout_col(i, j) == i + j*4

Same logical matrix, same Shape, two different memory layouts — and the code that indexes into it doesn't need to know which one it's looking at. That decoupling (the Layout is a first-class value, not a compile-time-fixed choice baked into every access) is most of the point.

Shapes nest

The examples above are flat — one number per dimension. CuTe Shapes are allowed to be nested tuples instead, and that's where layouts stop being "just strided arrays" and start being able to describe tiling.

A Shape like (4, (2, 2)) describes 4 outer elements, each containing a 2×2 sub-structure — the same data addressable either as one flat coordinate or as a hierarchical one, both resolving to the same memory offsets. This is precisely the mechanism CUTLASS 3.x uses to express "a threadblock's tile within the full matrix," or "a thread's fragment within a warp's tile" — the exact hierarchy from the previous post: device → threadblock → warp → thread is, underneath, nested Layouts, not four separate systems.

Composing layouts to tile a GEMM

Here's the piece that ties the abstraction back to something concrete. Extracting one threadblock's tile out of a full matrix in a CUTLASS 3.x kernel looks like this:

auto cta_tiler = make_shape(bM, bN, bK);              // e.g. 128x128x8
auto cta_coord = make_coord(blockIdx.x, blockIdx.y, _);
 
Tensor gA = local_tile(mA, cta_tiler, cta_coord, Step<_1, X, _1>{});

local_tile does two things at once: it reinterprets the full matrix's layout as (tile_shape, num_tiles) — the nested-shape idea above, applied automatically — and then indexes into it using this threadblock's coordinate (blockIdx.x, blockIdx.y). The result, gA, is a Tensor — a Layout plus the actual data pointer — that addresses only the slice of global memory this threadblock owns.

Going one level deeper, local_partition does the same kind of thing for threads instead of threadblocks — assigning each thread its own sub-tensor out of a shared-memory tile:

auto tA = make_layout(make_shape(Int<32>{}, Int<8>{}));  // 32x8 thread layout
Tensor tAsA = local_partition(sA, tA, threadIdx.x);       // this thread's slice

Same idea, one layer down: a Layout describing where threads sit relative to a tile, used to carve out exactly the data one thread is responsible for.

Why bother with all this instead of index math

You could write i*3 + j by hand — nobody needs CuTe to compute a single strided offset. What you can't easily do by hand, at least not without a lot of bug-prone repetition, is keep four or five layers of tiling (global memory → threadblock tile → warp tile → thread fragment) consistent with each other as you change tile sizes, data types, or target architecture. Composing Layouts — dividing them, partitioning them, nesting them — lets CUTLASS guarantee that consistency structurally: if the shapes don't compose, it's a compile error, not a kernel that silently reads the wrong address.

That composability is also what let CUTLASS 3.x collapse a large set of Ampere/Turing-era, generation-specific tiling classes into one general mechanism that extends cleanly to Hopper's warp-group MMA and beyond — which is where this series goes next.

References

Layout diagram above is original artwork made for this post.