---
title: 'CuTe Layouts: Hierarchical Tensor Mapping'
url: https://www.emergentmind.com/topics/cute-layouts
type: topic
---

# CuTe Layouts: Hierarchical Tensor Mapping

Searching arXiv for recent papers on CuTe layouts and closely related formalizations.
CuTe layouts are abstract specifications of how multi-dimensional tensor data and threads are mapped to linear resources such as memory and registers. In the CuTe formulation, a layout is written as a shape–stride object, typically \(L = S : D\) or \(H = s : d\), and is equipped with a layout algebra including composition, coalescence, complementation, division, tiling, and inversion. This representation directly extends traditional flat shape/stride tensor layouts to hierarchical forms required by modern GPU tensor instructions, and it is used as the foundation of NVIDIA’s CUTLASS library and related efforts including CuTe DSL [2603.02298].

## 1. Hardware motivation and conceptual scope

CuTe was introduced for the regime in which tensor cores, hardware-optimized copy instructions, and thread-cooperative memory movement prescribe fixed and often complex data layouts. The cited treatment places this in the context of Volta, Turing, Ampere, Hopper, and Blackwell GPUs, where MMA/WMMA, TMA, and TMEM instructions require layouts that are not naturally expressible by ordinary flat row-major or column-major indexing [2603.02298].

In this setting, a layout is not restricted to a memory layout for tensor elements. The same formal object is used for data layouts, thread layouts, and instruction layouts. A data layout maps logical tensor coordinates to offsets; a thread layout maps thread indices to logical coordinates; an instruction layout maps architecturally specified coordinates to offsets or fragment positions. CuTe’s central claim is that these three cases should be manipulated by a single algebra rather than by unrelated ad hoc index formulas [2603.02298].

This scope distinguishes CuTe from conventional flat tensor descriptors. Traditional shape/stride representations can encode row-major, column-major, padded, and uniformly strided layouts, but they do not naturally encode hierarchical tiles, interleavings, bank-swizzling patterns, or thread/value layouts. CuTe generalizes the familiar “shape + stride” idiom rather than discarding it: flat layouts remain special cases, while hierarchical and instruction-specific layouts become first-class objects [2603.02298].

## 2. Hierarchical representation: shape, stride, coordinates, and codomain

A CuTe shape is an `HTuple` of positive integers. An `HTuple(T)` is either an element of \(T\) or a `Tuple(HTuple(T))`, so shapes may be arbitrarily nested. Two such tuples are **congruent** when they have the same tree shape, and **weakly congruent** when one is at least as refined as the other. The shape and stride of a layout must be congruent [2603.02298].

The size of a shape is the product of its leaf integers. For a flat layout \(H = s : d\) with \(s = (s_0,\dots,s_{n-1})\), the ISL-based formalization defines the integral coordinate space
$$
C = \{ [c] : 0 \le c < P \}, \quad P = \prod_{i=0}^{n-1} s_i,
$$
and the natural coordinate space
$$
C_n = \{ [c_0,\dots,c_{n-1}] : 0 \le c_i < s_i \text{ for all } i \}.
$$
The coordinate mapping \(M_C^H\) uses colexicographic ordering:
$$
M_C^H = \{ c \rightarrow [c \bmod s_0,\ \lfloor c/s_0 \rfloor \bmod s_1,\ \dots,\ \lfloor c/\prod_{i=0}^{n-2} s_i \rfloor ] : 0 \le c < P \},
$$
while the index mapping is affine:
$$
M_I^H = \{[c_0,\dots,c_{n-1}] \rightarrow [\sum_{i=0}^{n-1} c_i d_i] : \forall i,\ 0 \le c_i < s_i \}.
$$
Their composition,
$$
M_L^H = M_I^H \circ M_C^H,
$$
is the layout mapping itself [2511.10374].

The 2026 representation paper gives the same idea in semimodule form. A stride is an `HTuple(D)` congruent with the shape, where \(D\) may be \(\mathbb{Z}\), a coordinate space, or a bit-vector space such as \(\mathbb{F}_2^m\). The layout is then \(L = D \circ S\), with evaluation by an inner product on natural coordinates. This permits ordinary integer offsets, coordinate-valued layouts, and binary swizzles to be expressed within one framework [2603.02298].

Standard flat layouts remain immediate special cases. Row-major \((M,N):(N,1)\) yields \(l = N \cdot i + j\); column-major \((M,N):(1,M)\) yields \(l = i + M j\). Hierarchical layouts such as \((4,(2,2)):(2,(1,8))\) preserve the same underlying mapping while exposing a nested coordinate structure that can be folded, unfolded, and reinterpreted algebraically [2511.10374].

## 3. Layout algebra

The defining feature of CuTe is not only its representation but its algebra. Every layout can be written as a concatenation of sublayouts,
$$
L = (L_0,\dots,L_{n-1}),
$$
with evaluation
$$
L(c) = L_0(c_0) + \cdots + L_{n-1}(c_{n-1}).
$$
This exposes the per-mode decomposition on which “by-mode” operations are built [2603.02298].

**Coalescence** flattens or simplifies a layout while preserving its action on integral coordinates. The result has the same size and the same mapping \(i \mapsto\) offset, but reduced hierarchy or rank. The cited examples include \((2,(1,6)):(1,(6,2))\), which coalesces to \((2,6):(1,2)\), and \(((4,3),5):((15,1),3)\), which coalesces to \((4,15):(15,1)\) [2603.02298].

**Composition** is the central transformation operator. Given layouts \(A\) and \(B\), the composition \(A \circ B\) satisfies
$$
R(c) = A(B(c))
$$
when the codomain of \(B\) is compatible with the domain of \(A\). The representation paper makes explicit that composition is governed by stride divisibility and shape divisibility conditions; if these fail, no CuTe layout exists that exactly represents the requested composed mapping over the intended finite domain [2603.02298]. This is the algebraic mechanism behind transposes, reblocking, thread partitioning, and instruction-specific reinterpretations.

**Inversion** appears in right, left, and full forms. A right-pseudo-inverse \(L^\ddagger\) and a left-pseudo-inverse \(L^\dagger\) support reasoning about contiguity, injectivity, and membership of architecturally required offsets in a data layout. Full inverses exist for bijective or “compact” layouts [2603.02298].

**Complementation** constructs a layout whose image fills ordered holes outside the image of a given layout. CuTe then defines **logical product**
$$
A \otimes B = (A, A^* \circ B)
$$
and **logical divide**
$$
A \oslash B = A \circ (B, B^*_{\#A}).
$$
These operators are the basis of tiling and partitioning. The paper identifies `zipped_divide` as the practical operator that repacks the result into a tile coordinate and a grid coordinate, which is then used pervasively to partition matrices across blocks, warps, and threads [2603.02298].

## 4. Formal semantics: integer-set relations and categorical foundations

A major subsequent development is the formalization of CuTe layouts as integer set relations in ISL. In that account, coordinate spaces, index spaces, and layout mappings are all explicit Presburger-style relations. Composition, inversion, and complement are implemented directly by ISL operations such as relation composition, inverse, set difference, range, domain, and lexicographic minimization. This gives mathematically checkable semantics for basic layouts, hierarchical layouts, and swizzles, and it supports correctness verification and cross-system reasoning with Triton linear layouts [2511.10374].

This formalization treats hierarchical layouts by flattening them, computing the flat layout mapping, and then reconstructing index mappings for alternative compatible shapes. A key consequence is that some equivalent coordinate organizations yield affine index mappings, while others yield quasi-affine ones. Since CuTe syntax expects shape/stride-based layouts with strictly affine index mappings, whether a compatible shape is representable in CuTe becomes a precise semantic question rather than an informal judgment [2511.10374].

The same paper also models CuTe swizzles \(H_{b,m,s}\) as integer set relations. Scalar bitwise definitions are translated into binary coordinate mappings, binary swizzle mappings, and scalar layout mappings. The resulting relations are quasi-affine and precise enough to verify properties such as involution, while also showing that CuTe swizzles are a special case of the binary linear transformations associated with Triton layouts [2511.10374].

A different line of formalization gives categorical foundations for a class of **tractable layouts**. It defines two categories, `Tuple` and `Nest`, whose morphisms encode flat and nested layouts, respectively. The paper proves a one-to-one correspondence between non-degenerate tractable layouts and non-degenerate standard-form `Nest` morphisms, and establishes compatibility theorems such as
$$
L_{g \circ f} = L_g \circ L_f
$$
for composition, along with analogous results for coalesce, complement, logical division, and logical product [2601.05972].

These two formalisms are complementary. The ISL semantics emphasizes full relation-level analysis, quasi-affinity, swizzles, and cross-system unification; the categorical account emphasizes the internal algebra of tractable layouts and proves that the operations used in CuTe practice correspond exactly to categorical constructions on morphisms [2511.10374] [2601.05972].

## 5. Role in CUTLASS, thread/value layouts, and compile-time verification

CuTe layouts are described as the foundation of NVIDIA’s CUTLASS library and of CuTe DSL, and the 2026 representation paper states that they have been successfully deployed in production systems [2603.02298]. This deployment context matters because the layout algebra is not presented as a purely descriptive notation. It is the mechanism by which kernels derive, propagate, and statically verify the operand layouts prescribed by modern hardware.

The same layout type is used for memory tensors, thread arrangements, and instruction fragments. A representative example is the Ampere FP64 tensor core \(C\)-matrix partition, expressed as a thread/value layout
$$
((4,8),2) : ((16,1),8).
$$
When a data layout \(L_{\text{data}}\) is composed with this thread/value layout, the result is a mapping from `(thread, value)` to memory offset, and slicing that composition yields the values owned by a given thread [2603.02298].

The paper gives analogous instruction-layout examples for TMEM loads. For `tcgen05.ld.32x32b.x1`, the layout is represented as
$$
(1,128):(1,16384),
$$
and for `tcgen05.ld.16x256b.x1` as
$$
(8,(16,4)) : (1,(16384,32\cdot 16384)).
$$
Because these are ordinary CuTe layouts, one can use left inverses, compositions, and compatibility checks to determine whether a data layout matches the instruction’s required addressing pattern [2603.02298].

This is also the basis for generic kernel construction. The paper’s high-level GEMM is parameterized by layout types `ALayout`, `BLayout`, and `CLayout`; algorithmic logic is separated from the tensor-memory mapping. Tiling, partitioning, transposition, and thread assignment are introduced by layout operations such as `zipped_divide`, `logical_product`, and composition, rather than by hand-written index arithmetic [2603.02298].

The productivity claim is concrete. The representation paper contrasts CUTLASS v2’s roughly 300 separate layout types, amounting to roughly 55,000 lines of code, with CuTe’s core implementation of roughly 3,000 lines, while asserting that CuTe represents all of those prior cases and additional ones [2603.02298]. In this sense, CuTe layouts function as a compact intermediate language for architecturally constrained tensor indexing.

## 6. Scope, limitations, and related directions

The most important limitation is that not every useful layout formalism coincides with every other one. The categorical framework applies to a naturally occurring class of **tractable layouts** and explicitly does not cover arbitrary swizzles or Triton-like \(\mathbb{F}_2\)-linear layouts. Its closure properties and theorems therefore hold for a large practical subset rather than for all conceivable CuTe-style mappings [2601.05972].

Conversely, the ISL-based formalization is deliberately more expressive than CuTe’s native shape/stride layouts plus swizzles. It notes an **expressiveness gap**: there are integer relations, including modulo-based permutations such as
$$
S = \{ [c_0] \to [(-3 c_0) \bmod 16] : 0 \le c_0 \le 15\},
$$
that are naturally expressible in ISL but are not expressible as CuTe layouts or Triton linear layouts [2511.10374].

A second limitation is reconstruction. Given only a one-dimensional layout relation \(M_L\), inferring both shape and strides in full generality is left open. The ISL paper provides algorithms when either the shape or the strides are known, but not for the unrestricted inverse problem [2511.10374].

A third limitation concerns complexity. Presburger arithmetic and polyhedral operations can be expensive in the worst case, although the same paper argues that realistic deep-learning layouts remain within practical scale, citing ranks up to 6 and tiling depths around 4, corresponding to at most 24 dimensions in the polyhedral representation [2511.10374].

The forward direction is therefore twofold. One line seeks richer compiler integration and cross-system optimization, including MLIR- or FPL-oriented layout-aware compilation and unified reasoning across CuTe and Triton [2511.10374]. The other line seeks deeper internal understanding of CuTe’s existing algebra, as in the categorical account’s characterization of tractable layouts and its proof that composition, complement, logical division, and logical product on morphisms align with CUTLASS behavior [2601.05972]. Together, these developments position CuTe layouts as both a practical kernel-construction language and a subject of ongoing formalization.

Source: https://www.emergentmind.com/topics/cute-layouts