2 of 4

CUTLASS, CuTe, and CUDA Tile

CUTLASS is NVIDIA’s own answer to the same problem, aimed squarely at one computation: it is “a collection of abstractions for implementing high-performance matrix-matrix multiplication (GEMM) and related computations at all levels and scales within CUDA,” decomposing that work into reusable, modular components. Through CUTLASS 2.x, that decomposition mirrored the GPU’s own hierarchy (thread, warp, threadblock) directly. CUTLASS 3.0 broke that mirroring on purpose: Hopper’s warp-group-wide instructions do not correspond to any single warp or thread-level concept, so tying the library’s structure to one generation’s hardware layout kept breaking on the next generation.

What replaced it is CuTe, described as “a collection of C++ CUDA template abstractions for defining and operating on hierarchically multidimensional layouts of threads and data.” A CuTe layout is, at bottom, a function from integers to integers, and CuTe builds a full layout algebra on top of that idea: composition, product, and divide let a kernel author build a thread-to-data mapping for an entire GEMM tile out of a handful of primitive layouts, instead of hand-writing a new iterator type for every architecture-specific access pattern. CUTLASS 3.0 replaced most of its 2.x-era named iterator types with this one vocabulary type, on the argument that a mapping expressed as an algebra can be checked at compile time: “if the code compiles, it’s probably correct.”

The shape:stride notation reads mechanically. In the layout algebra documentation’s worked example B = (4,3):(3,1), the shape (4,3) says coordinates run over a 4 by 3 grid, and the stride (3,1) says a step in the first coordinate advances the underlying index by 3 while a step in the second advances it by 1: the doc’s own table evaluates B(1,0) to 3 and B(0,1) to 1. Because a layout is just a function from integers to integers, the same function can often be written more than one way; the documentation notes that column-major layouts like (_2,_4):(_1,_2) “act identically to” _8:_1 for 1-D coordinates, and CuTe’s coalesce operation is the simplifier that finds the smaller spelling without changing the function:

coalesce, from the CuTe layout algebra documentation

auto layout = Layout<Shape <_2,Shape <_1,_6>>,
                     Stride<_1,Stride<_6,_2>>>{};
auto result = coalesce(layout);    // _12:_1

The algebra’s central operation is functional composition, R := A ∘ B with R(c) = A(B(c)), and the documentation calls it “the core of CuTe,” used “in just about every higher-level operation.” Its worked example composes A = (6,2):(8,2) with B = (4,3):(3,1), evaluates all twelve outputs by hand, and lands on the observation the whole library rests on: the result is itself a layout, ((2,2),3):((24,2),8), and it is compatible with B, meaning every coordinate of B is also a valid coordinate of the result. Selecting every third element of one layout through another never leaves the algebra, which is what lets the product and divide operations, and ultimately a whole thread-to-data partition for a GEMM tile, be built out of compositions and checked at compile time.

Composition and coalesce are the algebra’s easy half, and on their own they only rewrite a mapping into a different spelling of itself. The operations that split one, so that a thread or a threadblock can be handed its share of a tensor, are the subject of the next few pages, and the rest of this section climbs from there: Tile IR, which drops the thread-to-data mapping altogether, and then the CUTLASS layers that turn any of this into a kernel you can launch.

Partitioning: complement and divide#

Composition selects coordinates. Partitioning also needs the ones it did not select, and supplying them is the job of complement, an operation that “attempts to find another layout that represents the ‘rest’,” the elements the first layout never touches. The complement of a layout A under a size M is bounded by that size, has a codomain disjoint from A’s, and is ordered, its strides positive and increasing, which is what makes it unique. Two of the documentation’s examples cover the two jobs it does: complement(4:1, 24) is 6:4, where the layout 4:1 is “effectively repeated 6 times with 6:4,” and complement(6:4, 24) is 4:1, where “the ‘hole’ in 6:4 is filled with 4:1.”

Divide is complement and composition put together, and it is the operation that actually partitions data. Informally, logical_divide(A, B) “splits a layout A into two modes”: “in the first mode are all elements pointed to by B and in the second mode are all elements not pointed to by B.” Formally it is A composed with the concatenation of B and B’s complement, and the documentation is careful to note that this introduces nothing new, being “defined only in terms of concatenation, composition, and complement”:

logical_divide, from the CuTe layout algebra documentation

template <class LShape, class LStride,
          class TShape, class TStride>
auto logical_divide(Layout<LShape,LStride> const& layout,
                    Layout<TShape,TStride> const& tiler)
{
  return composition(layout, make_layout(tiler, complement(tiler, size(layout))));
}

The one-dimensional worked example tiles A = (4,2,3):(2,1,8) with the tiler B = 4:2: out of a 24-element vector in A’s storage order, extract tiles of 4 elements strided by 2. Three steps follow the definition. The complement of 4:2 under 24 is (2,3):(1,8); concatenating gives (4,(2,3)):(2,(1,8)); composing A with that gives ((2,2),(2,3)):((4,1),(2,8)), in which “the first mode of the result is the tile of data and the second mode of the result iterates over each tile,” six tiles in this case. The same thing generalizes to two dimensions by handing divide a tiler per mode, which the documentation describes as “a kind of gather operation or as simply a permutation on the rows and cols.”

That result is awkward to index, so CuTe ships rearrangements of it. For a layout of shape (M, N, L, ...) under a tiler <TileM, TileN>, logical_divide leaves ((TileM,RestM), (TileN,RestN), L, ...), preserving which mode is which, while zipped_divide gathers the subtiles into one mode and the leftovers into another, ((TileM,TileN), (RestM,RestN,L,...)); tiled_divide and flat_divide unpack that second mode by degrees. The zipped form is the one you index. With zd the zipped divide, “the offset to the 3rd tile is zd(0,3),” the (1,2)th tile is zd(0,make_coord(1,2)), and the tile itself is always layout<0>(zd), which the documentation observes is always equal to composition(a, b).

One wrapper up is the call a kernel author actually writes. In the CuTe GEMM tutorial the threadblock tile is just a shape, make_shape(bM, bN, bK), and each threadblock’s view of A, B, and C comes from local_tile, which the tutorial defines as exactly two steps: apply the tiler with zipped_divide, then slice the “Rest” mode with this threadblock’s coordinate. The Step<_1, X,_1> argument projects away the modes that do not apply to that operand, and the underscore in the coordinate keeps the k mode whole instead of selecting one value of it, so gA comes back rank-3 as (BLK_M,BLK_K,k): the first two modes are this threadblock’s tile, and the last indexes every k-tile it will reduce, which is what the mainloop iterates over.

local_tile and the two steps it wraps, from the CuTe GEMM tutorial

auto cta_coord = make_coord(blockIdx.x, blockIdx.y, _);              // (m,n,k)
Tensor gA = local_tile(mA, cta_tiler, cta_coord, Step<_1, X,_1>{});  // (BLK_M,BLK_K,k)

// ((BLK_M,BLK_K),(m,k))
Tensor gA_mk = zipped_divide(mA, select<0,2>(cta_tiler));
// (BLK_M,BLK_K,k)
Tensor gA = gA_mk(make_coord(_,_), select<0,2>(cta_coord));

This is the payoff of insisting that a layout be a function and that the operations on it close over layouts. A tiling that a CUDA programmer would write as index arithmetic, and would have to rewrite for every operand and every architecture, is here three named operations over two layouts, checked at compile time, with the same call producing a threadblock’s tile of A, of B, and of C by changing only which modes it projects.

CUDA Tile IR#

CUDA Tile IR, introduced in CUDA 13.1, takes the tile idea a level higher than either CUTLASS or Triton do on their own. Where CUTLASS and CuTe still hand a C++ (or now Python, via CuTe DSL) programmer explicit control over the thread-to-data layout, a tile block in Tile IR is written with no thread-to-data mapping at all; the compiler decides how a tile block’s work lands on real SM threads, the memory hierarchy, and tensor cores. Tile IR is positioned as a compilation target beneath higher-level DSLs, not a replacement for CUTLASS or Triton themselves, which is exactly the role it plays for Triton’s Tile IR backend described in the previous section.

The programming model documentation shows what that looks like as actual code. A Tile IR program is a module of tile kernels declared with entry, and every value in a kernel is a tensor whose rank, shape, and element type are statically known; rank-0 tensors are scalars, and global memory is only ever reached through tensors built from pointer parameters. Its first worked example, a 128-element vector addition, spends most of its lines constructing a tile of 128 pointers from a scalar base pointer: an iota builds the offset vector 0 through 127, a reshape and broadcast replicate the base pointer across the tile, and an offset adds the two. The arithmetic itself is then three statements:

Tile IR vector addition, from the Tile IR programming model documentation

%a_val, %token_a = load_ptr_tko weak %a_tensor : tile<128xptr<f32>> -> tile<128xf32>, token
%b_val, %token_b = load_ptr_tko weak %b_tensor : tile<128xptr<f32>> -> tile<128xf32>, token
%c_val = addf %a_val, %b_val rounding<nearest_even> : tile<128xf32>
store_ptr_tko weak %c_tensor, %c_val : tile<128xptr<f32>>, tile<128xf32> -> token

The documentation’s summary of that kernel is the model in one sentence: “this code is written from a single thread of control, but its level of parallelism will be determined by the compiler.” Scaling past one tile reuses CUDA’s launch shape rather than replacing it: tile blocks group into a 1-d, 2-d, or 3-d tile grid, the grid size set at launch determines how many tile blocks run, and each block queries its position with get_tile_block_id and the grid’s dimensions with get_num_tile_blocks, the role blockIdx plays in the kernels of Foundations, one level of hierarchy up.

Vector addition is where the Triton section said the interesting part starts, and Tile IR’s documentation agrees: it lays out a GEMM progression that computes a two-dimensional GEMM for a single block first, then generalizes “step by step to a full GEMM by utilizing the tile grid and introducing control flow and manual tiling.” The first step is the vector-addition kernel with square tiles substituted for flat ones. The iota now builds 4096 offsets reshaped to 64 by 64, the same reshape and broadcast turn each base pointer into a 64 by 64 tile of pointers, and the multiply is a single mmaf, the floating-point matrix-multiply-accumulate operation, with mmai its integer counterpart:

A single-block 64 by 64 GEMM, from the Tile IR programming model documentation

%offset_flat = iota : tile<4096xi32>
%offset = reshape %offset_flat : tile<4096xi32> -> tile<64x64xi32>

%A_block, %token_a = load_ptr_tko weak %a_tensor : tile<64x64xptr<f32>> -> tile<64x64xf32>, token
%B_block, %token_b = load_ptr_tko weak %b_tensor : tile<64x64xptr<f32>> -> tile<64x64xf32>, token
%C_block = mmaf %A_block, %B_block, %init_accum: tile<64x64xf32>, tile<64x64xf32>, tile<64x64xf32>
store_ptr_tko weak %c_tensor, %C_block : tile<64x64xptr<f32>>, tile<64x64xf32> -> token

Set that beside the CuTe partitioning above and the difference in level is the whole point. There is no thread layout, no shared memory, no tensor-core instruction selected, and no divide: one mmaf over a 64 by 64 tile is the entire inner computation, and the documentation’s own gloss is that “basic matrix multiplication is achieved by constructing 2D tiles from pointers and performing matrix-multiply-accumulate (MMA) operations within a single block.” What this version cannot do is scale, since it is written against one static block with simplifying assumptions about the input pointers, and the documentation asks the question those assumptions raise, what happens “if we want to run the tile kernel over a large input problem size in parallel.” The tile grid, control flow, and manual tiling are the answer, and they are the same three concerns the Triton matmul handled with a program-id grid, a k loop, and block-size constants.

Tile IR is also something a reader can run rather than only read about. NVIDIA ships it as open source aligned with the CUDA Toolkit 13.1 release: an “MLIR-based intermediate representation and compiler infrastructure for CUDA kernel optimization,” made of a CUDA Tile MLIR dialect, Python bindings for programmatic IR construction and manipulation, a bytecode format with serialization both ways, and a conformance test suite checked against the specification. The path from text to a running kernel is short. The cuda-tile-translate tool turns an MLIR program into Tile IR bytecode, which the CUDA driver API can load and JIT compile directly, or tileiras from the CUDA Toolkit can compile ahead of time into a cubin for a particular GPU target.

Assembling a CUTLASS kernel#

CuTe is the vocabulary; the library above it is a stack CUTLASS documents as Device, Kernel, Collective, Tiled MMA and Copy, and Atom, ordered highest to lowest. An Atom is “the smallest collection of threads and data that must participate in the execution of a hardware-accelerated math or copy operation,” indivisible in space rather than in time. A Collective is its mirror image, “the largest collection of threads onto which mma atoms and copy atoms are tiled,” meaning the largest group that can still cooperate through hardware features such as asynchronous array copy, MMA on tiles resident in shared memory, cluster and block synchronization, and hardware barriers for asynchronous data dependencies. The tiled layer between them turns one atom into a block-wide operation.

The producer-consumer pipeline. CUTLASS's asynchronous pipeline. Producer threads load tiles into a circular list of shared-memory stages while consumer threads compute on tiles one slot behind, the two sides synchronized through acquire, commit, wait, and release barrier operations. In the highlighted slot the load of tile k+1 runs while tile k is being computed: the software pipelining CUTLASS calls critical to hiding the latency of global memory loads.Illustrative numbers

Assembling a kernel therefore runs bottom-up and stops at the host. The documentation states the order plainly: compose the required collective mainloop and epilogue, compose those into a kernel type, and wrap the kernel in a device-layer adapter. Its Hopper example is four declarations long, three of which appear below, and the entire configuration of the mainloop sits in the argument list of a single builder:

Kernel assembly, from the CUTLASS 3.0 GEMM API documentation

// Step 1: Generate the required collective layer mainloop specialization
using CollectiveMainloop = typename cutlass::gemm::collective::CollectiveBuilder<
    ArchTag, OperatorClass,
    ElementA, LayoutA, AlignmentA,
    ElementB, LayoutB, AlignmentB,
    ElementAccumulator,
    TilesShape, ClusterShape,
    cutlass::gemm::collective::StageCountAuto,
    cutlass::gemm::collective::KernelScheduleAuto
  >::CollectiveOp;

// Step 3: Compose the mainloop and epilogue together at the kernel layer
using GemmKernel = cutlass::gemm::kernel::GemmUniversal<
    cute::Shape<int,int,int,int>, // ProblemShape [M,N,K,L]
    CollectiveMainloop,
    CollectiveEpilogue
>;

// Step 4: Wrap up the kernel::GemmUniversal kernel class
// with the device adapter to obtain a host-side handle to the kernel
using GemmHandle = cutlass::gemm::device::GemmUniversalAdapter<GemmKernel>;

The builder exists because the collective it produces is an expert interface nobody wants to fill in by hand. CollectiveMma allows “full control over all the properties of the collective’s GPU micro-kernel” across fifteen template parameters, from strides to shared-memory layout atoms to per-operand transforms, while CollectiveBuilder “accepts CUTLASS 2.x equivalent input template arguments, and attempts to build the best performing CollectiveMma from the given parameters.” The last two arguments are where the decisions get delegated: StageCountAuto “allows the collective builder to compute the size of a single stage’s size in shared memory and maximize the shared memory usage assuming 1 threadblock / multiprocessor occupancy,” so the pipeline depth of the figure above is not a number the programmer picks but whatever fits, and KernelScheduleAuto picks “the best kernel schedule available for the given set of parameters” unless an explicit tag overrides it.

Those tags are the library’s answer to the combinatorial problem underneath. Collectives “are not generic. Instead, they must be specialized for each algorithm and GPU architecture,” so CUTLASS 3.0 dispatches on a policy type: Hopper’s MainloopSm90TmaGmmaWarpSpecialized carries a stage count and a cluster shape “over which TMA multicast will take place,” and pairs with a schedule chosen from a small enumerated set including KernelTmaWarpSpecializedPingpong and KernelTmaWarpSpecializedCooperative. The rule that makes the indirection worth its cost is stated both ways: “a single kernel schedule can support multiple mainloop implementations,” and “a single mainloop can be composed with multiple possible kernel schedules.” A new schedule does not fork the mainloop, and the Blackwell dispatch policies of Current hardware are the same naming scheme one architecture on.

Above the collective, the kernel layer is “a collection of all clusters in the grid” that orders the collectives, marshals the threads of a warp-specialized schedule into their roles, swizzles the grid, and tiles the input tensors before invoking the collectives. GemmUniversal is stateless, and 3.x promotes the problem shape to a top-level template argument so a fully static instantiation is possible when the shapes are known at compile time. Above it, GemmUniversalAdapter is “a stateful, reusable GEMM handle” and deliberately “does not specify the grid shape. The kernel controls the grid shape and other kernel-specific launch parameters,” which is what lets every 3.x kernel share one launch path.

CUTLASS 4.x adds a second front door to the same machinery. NVIDIA describes CUTLASS Python as “the Python kernel-authoring stack in CUTLASS 4.x,” for writing kernels “with Python syntax while retaining explicit control over the GPU memory, thread, and data hierarchy,” and the vocabulary carries over intact: tensors, layouts, MMA and copy atoms, tiled operations, and pipeline APIs that “coordinate asynchronous copies, barriers, compute work, and producer/consumer warp groups.” Two decorators do the work, @jit for host-side functions and @kernel for GPU kernels, and the launch parameters read like the C++ ones made dynamic: grid, block, cluster, a smem that defaults to computing shared-memory usage automatically, and opt-ins such as a fallback_cluster that turns the requested cluster into a preference allowing “graceful degradation when hardware cannot satisfy the preferred dimensions.” Compilation is just-in-time, with IR modules cached across calls and DLPack handing tensors to and from PyTorch and JAX, and CUTE_DSL_LINEINFO=1 emits line information so tools “can correlate generated PTX/SASS back to the Python source that produced it.”

The documentation is careful about what that buys. NVIDIA states that the Python DSL “is optimized for the NVIDIA Blackwell Architecture and achieves performance within 2% of handwritten C++ implementations,” with Hopper and Ampere support “available as an experimental feature at launch,” a vendor claim rather than anything measured here. It also draws the boundary: CUTLASS Python “is not a replacement for the CUTLASS C++ library or its 2.x and 3.x APIs,” it “focuses on authoring and tuning individual kernel instances,” and it “does not currently provide the full GEMM/Conv profiler or library interface available in CUTLASS C++.”