CUTLASS

Part I: The mental model

one GEMM:  D = alpha * (A @ B) + beta * C        A is MxK, B is KxN, D is MxN

decomposed onto the hardware, level by level:

GLOBAL MEMORY (HBM)         the whole A, B, C, D live here
      |   tile it
      v
THREADBLOCK / CTA tile      one thread block owns an Mt x Nt output tile
      |                     it streams K in steps, staging A/B tiles in SHARED MEMORY
      v
WARP tile                   one warp (or warpgroup) owns an Mw x Nw sub-tile
      |                     reads operands from SHARED MEMORY into REGISTERS
      v
INSTRUCTION / MMA atom      one tensor-core instruction does a small Mi x Ni x Ki
                            multiply-accumulate straight out of registers/shared

the loop that ties it together:

mainloop over K:   load next A/B tile (cp.async / TMA)  ||  MMA on current tile
                   double/multi-buffered shared memory keeps loads and math overlapped
      |
      v
epilogue:          accumulators in registers -> alpha/beta, bias, activation,
                   arbitrary fused elementwise (EVT) -> store D to global memory

The one-sentence identity. CUTLASS is a template library that expresses one matrix multiply as a nested hierarchy of tiles, one per level of the GPU memory system, so that each level reuses data enough times to keep the next-faster level, and ultimately the tensor cores, busy. A matrix multiply is almost pure arithmetic sitting on top of a data-movement problem. The tensor cores can do arithmetic far faster than memory can supply operands, so the only way to reach peak is to load each element of A and B once from slow memory and then reuse it many times out of fast memory. Tiling is how you buy that reuse, and doing it at every level, global to shared, shared to register, register to tensor core, is the whole game. Part V makes this quantitative with the arithmetic-intensity story.

The second load-bearing idea is CuTe, short for CUDA Tensors. A Layout in CuTe is a pair of a Shape and a Stride, and it is nothing more than a function from a logical coordinate to a linear offset. A Tensor is a pointer plus a Layout. That sounds trivial, but layouts can be hierarchical, they nest, and they support a real algebra of composition, division, and product. With that algebra a single object can describe a tile of a matrix, the partition of that tile across the threads of a warp, or the fragment a tensor-core instruction expects in registers, and all of these compose by the same rules. Once the mapping from threads to data is itself a value you can compute with, tiling a GEMM stops being index arithmetic scattered through the kernel and becomes a few layout operations you can read, print, and check. CUTLASS 3.x rebuilt the whole library on CuTe for exactly this reason.

Two consequences follow. First, CUTLASS is a toolkit, not a drop-in library. cuBLAS gives you a function to call. CUTLASS gives you the templates to assemble a kernel with the exact data types, tile shapes, pipeline depth, and fused epilogue you want, and then get near cuBLAS-class performance on the newest hardware, often before a tuned closed-source path exists. Second, because the building blocks are open and composable, CUTLASS became the substrate other projects build on. FlashAttention and the fused attention kernels inside inference engines like vLLM, SGLang, and TensorRT-LLM lean on CuTe atoms and CUTLASS collectives for their matrix multiplies rather than reinventing them. Everything here is written against the CUTLASS 3.x and 4.x line as of mid 2026. The API surface moves quickly between releases, so where an exact template signature is likely to have shifted I say so and stay at the level of the idea.

Part II: Using it

CUTLASS is a header-only library for the parts you include, plus a CMake project for the examples, tests, and profiler. You need an NVIDIA GPU, a recent CUDA toolkit, and a C++17 compiler. To use the headers in your own code you only have to put include/ on the include path. To build the examples and the profiler you clone and run CMake, telling it which GPU architectures to target:

git clone https://github.com/NVIDIA/cutlass
cd cutlass
mkdir build && cd build

# target one architecture: 80 = Ampere (A100), 90a = Hopper (H100),
# 100a = Blackwell. The 'a' suffix enables arch-accelerated features.
cmake .. -DCUTLASS_NVCC_ARCHS=90a

# building every kernel is enormous and slow; scope it down while learning
make cutlass_profiler -j 16

The first thing to know is that CUTLASS can generate an overwhelming number of kernel instances, one per combination of data type, tile shape, and layout, and building all of them can take hours and many gigabytes. CUTLASS_NVCC_ARCHS restricts the target architectures and CUTLASS_LIBRARY_KERNELS filters which instances get compiled into the instance library. Keep both narrow while you are learning. The single most useful build target is cutlass_profiler, a command-line tool that enumerates the compiled kernels, runs any of them against a problem size, and reports throughput. It is the fastest way to see what the library can do without writing a line of C++:

# run every applicable GEMM instance for a 4096-cube problem and rank them
./tools/profiler/cutlass_profiler --operation=Gemm \
    --m=4096 --n=4096 --k=4096

# pin the data types and see which tile shapes win
./tools/profiler/cutlass_profiler --operation=Gemm \
    --A=f16:row --B=f16:col --C=f16:row --accum=f32 \
    --m=8192 --n=8192 --k=8192

Now the smallest amount of C++ that runs a real tensor-core GEMM. The 2.x device API is the gentlest entry point, and it still ships and works. You pick element types and layouts for A, B, and C, an accumulator type, an operator class that says use the tensor cores, and a target architecture, and the template picks sensible tile shapes for you:

#include <cutlass/gemm/device/gemm.h>
#include <cutlass/numeric_types.h>

// A is row-major fp16, B is column-major fp16, C/D are row-major fp16,
// accumulation happens in fp32, on Ampere tensor cores.
using Gemm = cutlass::gemm::device::Gemm<
    cutlass::half_t, cutlass::layout::RowMajor,      // A: element, layout
    cutlass::half_t, cutlass::layout::ColumnMajor,   // B
    cutlass::half_t, cutlass::layout::RowMajor,      // C and D
    float,                                           // accumulator element
    cutlass::arch::OpClassTensorOp,                  // use tensor cores
    cutlass::arch::Sm80>;                            // target Ampere

Gemm gemm_op;
cutlass::Status status = gemm_op({
    {M, N, K},                 // problem size
    {ptrA, lda}, {ptrB, ldb},  // A and B as {pointer, leading dim}
    {ptrC, ldc}, {ptrD, ldd},  // source C and destination D
    {alpha, beta}              // epilogue scalars for alpha*AB + beta*C
});
if (status != cutlass::Status::kSuccess) { /* handle */ }

The default epilogue there is LinearCombination, which computes alpha * accumulator + beta * C and casts to the output type. Swapping it for a fused variant, adding a bias vector or a ReLU, is a matter of naming a different epilogue functor in the template. That is the first taste of the design, the kernel is assembled from parts, and the epilogue is one of the parts you get to choose.

On Hopper and Blackwell the recommended path is the 3.x collective builder, which asks you for the problem shape, a threadblock tile, and a cluster shape, then auto-selects the pipeline schedule, stage count, and copy atoms appropriate to the architecture. The signature has changed across releases, so treat this as schematic and cross-check it against the version's examples/ directory:

#include <cutlass/gemm/collective/collective_builder.hpp>
#include <cutlass/epilogue/collective/collective_builder.hpp>
#include <cutlass/gemm/kernel/gemm_universal.hpp>
#include <cutlass/gemm/device/gemm_universal_adapter.h>
using namespace cute;

using TileShape    = Shape<_128, _256, _64>;  // CTA tile: M, N, K
using ClusterShape = Shape<_2,   _1,   _1>;    // Hopper thread block cluster

// the mainloop: how A and B are staged and multiplied
using Mainloop = typename cutlass::gemm::collective::CollectiveBuilder<
    cutlass::arch::Sm90, cutlass::arch::OpClassTensorOp,
    cutlass::half_t, cutlass::layout::RowMajor,    8,   // A + alignment
    cutlass::half_t, cutlass::layout::ColumnMajor, 8,   // B + alignment
    float,                                              // accumulator
    TileShape, ClusterShape,
    cutlass::gemm::collective::StageCountAuto,          // pick shared-mem depth
    cutlass::gemm::collective::KernelScheduleAuto       // pick warp schedule
>::CollectiveOp;

// the epilogue: how accumulators become D (fusions plug in here)
using Epilogue = typename cutlass::epilogue::collective::CollectiveBuilder<
    cutlass::arch::Sm90, cutlass::arch::OpClassTensorOp,
    TileShape, ClusterShape,
    cutlass::epilogue::collective::EpilogueTileAuto,
    float, float,                                       // compute, accumulator
    cutlass::half_t, cutlass::layout::RowMajor, 8,      // C
    cutlass::half_t, cutlass::layout::RowMajor, 8,      // D
    cutlass::epilogue::collective::EpilogueScheduleAuto
>::CollectiveOp;

using GemmKernel = cutlass::gemm::kernel::GemmUniversal<
    Shape<int,int,int,int>, Mainloop, Epilogue>;     // problem is M,N,K,L
using Gemm = cutlass::gemm::device::GemmUniversalAdapter<GemmKernel>;

The mistakes beginners make. First, trying to build everything. Scope the architectures and the kernel filter, or the first make will run for an hour. Second, confusing element and accumulator types. Tensor-core GEMM typically multiplies fp16 or bf16 operands but accumulates in fp32, and getting the accumulator wrong quietly costs accuracy or fails to compile. Third, ignoring alignment. The alignment number in the collective builder, and the vectorized access CUTLASS wants, mean your leading dimensions and pointers usually must be aligned to 8 or 16 elements, and a misaligned tensor either falls back to a slow path or refuses to run. Fourth, expecting the newest tile shapes to run on old hardware. A Hopper wgmma schedule needs -DCUTLASS_NVCC_ARCHS=90a and an H100. Compile for the architecture you will actually run on.

Part III: When it is the right tool

CUTLASS is the right tool when you need a matrix multiply or convolution that cuBLAS cannot express, and you are willing to work in C++ templates to get it. The common triggers are a fused epilogue cuBLAS does not offer, a quantized or mixed-input GEMM such as fp8 by fp8 or int4 weights by fp16 activations, a grouped or batched GEMM with irregular shapes as in mixture-of-experts, an unusual data layout, or simply wanting an open kernel you can read, modify, and specialize. It is also the place where support for a brand-new architecture lands first, so if you need to use Hopper or Blackwell tensor cores at peak the week the GPU ships, CUTLASS is often the only game in town.

The honest cases for alternatives. cuBLAS and cuBLASLt when a plain or lightly-fused GEMM is all you need, since they are a tuned closed-source library you call rather than a toolkit you assemble, and cuBLASLt already fuses common bias and activation epilogues. Triton when you want to write a fused kernel in Python with a block-level programming model and a JIT, trading some peak performance and some control over the exact instruction stream for a much shorter path from idea to working kernel. Plain hand-written CUDA with the mma intrinsics when your kernel is small and you want no abstraction at all. On AMD hardware the closest analogue is Composable Kernel, since CUTLASS is NVIDIA-only. And for the layers below GEMM, the scan, sort, and reduce primitives, reach for CUB and Thrust rather than CUTLASS.

There is a real learning-cost warning here. CUTLASS is famous for template error messages that run to hundreds of lines, because the entire kernel configuration lives in the type system and a single mismatched tile shape or unsupported type combination surfaces as a wall of substitution failures. The payoff is that a configuration that compiles is a configuration the library believes is valid, and much of the correctness is checked before the kernel ever runs. But budget time for reading types, not just code. The newer CuTe DSL, a Python front end over the same layout algebra that recent releases ship, exists partly to soften this, letting you author kernels against CuTe abstractions in Python while keeping the semantics of the C++ library.

The architecture-shaped caution is about matching the kernel to the problem shape. A GEMM that is large and square is compute-bound and almost any reasonable tiling reaches peak. A GEMM that is tall and skinny, a large K with small M and N as in a single-token decode step, has too few output tiles to fill the GPU, and the arithmetic intensity is low enough that memory bandwidth, not the tensor cores, sets the ceiling. This is the regime where split-K and Stream-K schedules matter, where the choice between them is not cosmetic, and where blindly using the tile shape that won on a square problem gives you a fraction of the achievable throughput. The right tile and the right scheduler depend on the shape, which is exactly why CUTLASS parameterizes both.

Part IV: The full life of one GEMM

The specimen is one tensor-core GEMM, D = alpha * A @ B + beta * C, with fp16 operands and fp32 accumulation, running the 3.x collective path on a Hopper GPU. The Ampere path is the same story with cp.async in place of TMA and mma.sync in place of wgmma, and I note the fork where it matters.

Stage 1: host setup and launch

On the host you instantiate the Gemm type from Part II, build an Arguments struct holding the problem shape, the device pointers and strides for A, B, C, and D, and the epilogue parameters such as alpha and beta. You then ask the adapter how much scratch it needs with Gemm::get_workspace_size(args), allocate that workspace, call can_implement(args) to validate the configuration against the problem, then initialize and run. Under the hood run computes a grid shape from the problem and the CTA tile, sets the dynamic shared-memory size, and launches the single kernel. Everything after this is on the device.

Stage 2: tile scheduling and the grid

The output matrix is covered by a grid of Mt-by-Nt tiles, and a tile scheduler assigns tiles to thread blocks. The naive scheduler hands block (i, j) the tile at row block i, column block j, but CUTLASS also swizzles this mapping so that blocks running at the same time touch nearby regions of A and B, which improves L2 cache reuse. For skewed shapes the Stream-K scheduler instead splits the K dimension across blocks so every streaming multiprocessor gets a balanced share of work, then reduces the partial results. On Hopper the grid is organized into thread block clusters, small groups of blocks that can read each other's shared memory, which the cluster shape in Part II selects. Each block now knows which output tile it owns.

Stage 3: the prologue and CuTe partitioning

Inside the kernel, before any math, the block sets up its view of the data with CuTe. It wraps the global A, B, C, D pointers in tensors with their strides, then uses local_tile to slice out just this block's Mt-by-K slab of A and K-by-Nt slab of B. It allocates the shared-memory staging buffers, which are multi-buffered so several K-tiles can be in flight at once, and it builds a TiledMMA from the architecture's MMA atom and asks it, through partition_C, for this thread's slice of the accumulator, a small register fragment that starts zeroed. All of the ugly index arithmetic that used to fill a GEMM kernel is now a handful of layout operations over these tensors.

Stage 4: the mainloop, load and compute overlapped

This is the heart of the kernel and where most of the time is spent. The block walks the K dimension in tile-sized steps, and at every step two things happen concurrently. Loads: a copy atom brings the next A-tile and B-tile from global memory into the next free shared-memory buffer. On Hopper this is a single TMA bulk copy per tile, issued by one thread and driven by the Tensor Memory Accelerator, which knows the tile's shape from a descriptor and signals completion through an mbarrier. On Ampere it is a set of cp.async instructions that copy global to shared without staging through registers. Compute: the tensor cores run MMA instructions over the shared-memory buffer that a previous iteration already filled, accumulating into the register fragment. On Ampere the warp first uses ldmatrix to pull operand fragments from shared memory into registers in the exact layout the mma.sync instruction wants. On Hopper the wgmma instruction is issued by a whole warpgroup of 128 threads and reads its operands straight from shared memory through descriptors, so no explicit register load is needed. Because the shared memory is multi-buffered, the loads for tile k+1 overlap the math on tile k, and the pipeline barriers make a producer set of warps and a consumer set of warps hand buffers back and forth. This warp specialization, producers doing TMA and consumers doing wgmma, is the defining shape of a Hopper CUTLASS kernel.

multistage software pipeline over K (shared memory is a ring of buffers):

buffer:   b0     b1     b2     b0     b1  ...
load:    [k0]->[k1]->[k2]->[k3]->[k4]           producer warps: TMA / cp.async
mma:           [k0]->[k1]->[k2]->[k3]->[k4]     consumer warps: wgmma / mma.sync
                ^ compute on k0 overlaps load of k2
result:  register accumulator grows by A[:,k] @ B[k,:] each step

Stage 5: the epilogue and fusion

When the K loop finishes, each thread holds its slice of the raw A @ B accumulator in registers, in whatever scattered layout the tensor-core instruction produced. The epilogue turns that into the stored result. In the simple case it loads the matching piece of C, computes alpha * accumulator + beta * C, casts to the output type, and stores. Because the MMA output layout is not friendly for coalesced global writes, the epilogue usually routes the data through shared memory once to reorganize it into wide, aligned store transactions, and on Hopper the store itself can be a TMA copy back to global. The powerful case is fusion. CUTLASS 3.x expresses the epilogue as an Epilogue Visitor Tree, a small compute graph whose nodes are elementwise ops, broadcasts of a bias vector, activation functions, and reductions. You can fuse a bias add, a GELU, a per-channel scale for dequantization, and a reduction for the next layer's normalization, all into the same kernel, so the accumulator is consumed while it is still hot in registers and never makes an extra trip to memory.

Stage 6: the store, and any split-K reduction

The epilogue writes D to global memory and the kernel exits. If the configuration used split-K or Stream-K, several blocks each computed a partial sum over part of K, and a short second phase reduces those partials into the final D, either in a separate reduction kernel or through atomics into the output. That closes the loop of one GEMM, operands streamed in through the memory hierarchy, multiplied on the tensor cores with loads hidden under math, fused in the epilogue, and the result written back once. The same six stages describe a convolution too, because CUTLASS implements convolution as an implicit GEMM, forming the im2col matrix on the fly inside the mainloop rather than materializing it in memory.

Part V: Internals deep dives

Deep dive: the arithmetic-intensity story

Everything about CUTLASS follows from one number. Arithmetic intensity is the ratio of floating-point operations performed to bytes moved from memory, in FLOP per byte. The roofline model says the throughput you can attain is min(peak_flops, intensity * peak_bandwidth). Below a threshold you are memory-bound and bandwidth caps you. Above it you are compute-bound and the tensor cores cap you. That threshold, the ridge point, is peak_flops / peak_bandwidth. For a modern data-center GPU the tensor cores are so fast relative to memory that this ridge point is high. An A100, for instance, delivers roughly 300 TFLOP/s of dense fp16 tensor-core throughput against roughly 2 TB/s of HBM bandwidth, which puts the ridge point near 150 FLOP per byte. To be compute-bound you must do about 150 floating-point operations for every byte you read.

Now measure a GEMM against that bar. A matrix multiply of shape M-by-N-by-K does 2*M*N*K FLOP. The minimal data it must touch is M*K + K*N + M*N elements. A naive kernel that recomputes reads gets nowhere near the minimum and lands far to the memory-bound left of the ridge. Tiling fixes this. Suppose a block holds an Mt-by-Nt output tile in fast storage and streams K. At each K-step it reads Mt + Nt new elements and does 2 * Mt * Nt FLOP, so the intensity of the tile is:

intensity  =  (2 * Mt * Nt FLOP) / ((Mt + Nt) elements * bytes_per_element)

for a square tile Mt = Nt = T, with b bytes per element:

intensity  =  2 * T^2 / (2 * T * b)  =  T / b

so intensity grows LINEARLY with tile size T.

fp16, b = 2 bytes, ridge point ~150 FLOP/byte  =>  need T / 2 > 150  =>  T > 300
                                                    per level of reuse combined

That single line, intensity equals tile size over bytes per element, is the reason the whole hierarchy exists. A bigger tile reuses each loaded element more times, which raises arithmetic intensity, which is the only way to cross the ridge point and become compute-bound. But no single level of storage is both large enough and fast enough to hold a tile big enough on its own. Registers are fastest but tiny, shared memory is bigger but slower, L2 and HBM bigger and slower still. So CUTLASS multiplies the reuse across levels. The threadblock tile, often 128-by-256, gets its reuse against slow global memory by living in shared memory. The warp tile, often 64-by-64 or 64-by-128, gets its reuse against shared memory by living in registers. The instruction tile, the 16-by-8-by-16 shape of one Ampere MMA or the 64-wide shape of a Hopper wgmma, gets its reuse right at the tensor-core inputs. Each level's job is to raise the effective intensity seen by the level below it, and the product of the per-level reuse is what lifts a bandwidth-bound problem onto the compute roofline. Read the media/docs/efficient_gemm.md note alongside this section, it walks the same hierarchy from NVIDIA's own framing.

Deep dive: CuTe layout algebra

A CuTe Layout is a Shape paired with a Stride, and it is a function from a coordinate to an integer offset. The layout (4,2):(1,4) has shape four by two and strides one and four, so coordinate (i,j) maps to i*1 + j*4. That is a column-major four-by-two tile. Change the strides to (2,1) and the same shape becomes row-major. Shapes and strides can be nested tuples, which is what makes layouts hierarchical, a shape like ((2,2),2) describes a tile-of-tiles and the algebra handles the nesting automatically. You build them with make_layout and static integers like _4{}, and you can print_layout them to see the coordinate-to-offset table directly:

#include <cute/tensor.hpp>
using namespace cute;

// a 4x2 column-major layout: coord (i,j) -> i*1 + j*4
auto layout = make_layout(make_shape (_4{}, _2{}),
                          make_stride(_1{}, _4{}));
print_layout(layout);            // prints the 2D index table

// a tensor is just a pointer wearing a layout
auto A = make_tensor(ptrA, make_layout(make_shape(M, K), make_stride(K, _1{})));

// slice out one CTA's tile without copying anything
auto gA = local_tile(A, make_shape(_128{}, _64{}), make_coord(block_m, k));

The reason this is an algebra and not just a struct is the set of operations that combine layouts. composition feeds one layout's output into another's input, which is how you express slicing and reshaping. complement fills in the strides a layout does not cover. The products, logical_product, blocked_product, and raked_product, tile one layout by another, and the divides, logical_divide and zipped_divide, partition a layout into tiles. The two you meet constantly in kernels are local_tile, which gives a block its tile of a big tensor, and local_partition, which gives a thread its elements of a tile according to a thread layout. Because the thread-to-data mapping is itself a layout, the same value can describe a matrix tile and the way threads carve it up, and partitioning becomes one algebraic operation instead of a page of index math. The trap to internalize early is that a layout is bookkeeping, not data. Two very different layouts can point at the same bytes, and a correct-looking kernel with the wrong thread layout produces bank conflicts or wrong answers with no type error, because every layout is a valid layout. The CuTe docs under media/docs/cute/ build this up carefully and are worth reading in order.

Deep dive: MMA atoms, copy atoms, and tiling them

An atom in CuTe is the smallest indivisible unit of a hardware operation, wrapped so the layout algebra can reason about it. An MMA_Atom corresponds to exactly one tensor-core instruction and carries the thread-value layouts that say which thread holds which element of the operands and the result. The atom names encode the instruction directly. An Ampere fp16 atom reads like SM80_16x8x16_F32F16F16F32_TN, which parses as architecture SM80, MMA shape 16-by-8-by-16, output and operand and accumulator types F32, F16, F16, F32, and the TN operand arrangement the instruction requires. A Hopper warpgroup atom reads like SM90_64x128x16_F32F16F16_SS, a 64-by-N-by-16 wgmma that takes both operands from shared memory, which is what the SS suffix means. You rarely name these by hand on the 3.x path, the collective builder picks them, but knowing the naming convention lets you read any CUTLASS kernel.

A single atom is tiny, one instruction over a handful of threads. A TiledMMA replicates an atom across the warps of a block to cover the whole warp tile, and it exposes partition_A, partition_B, and partition_C so each thread can find its fragments. The mirror image on the data-movement side is the Copy_Atom, one copy instruction such as a cp.async transfer, an ldmatrix shared-to- register load, or a TMA bulk copy, and TiledCopy tiles it across threads the same way. The whole mainloop is then two tiled objects working together, a TiledCopy feeding shared memory and a TiledMMA consuming it, and the actual math is one call to cute::gemm(tiled_mma, fragA, fragB, fragC). Expressing both the compute and the data movement as tiled atoms means the same algebra that lays out the matrices also lays out the instructions, so a kernel is a composition of a few well-typed pieces rather than a monolith. One more layer worth naming is the shared-memory swizzle. Tensor cores read shared memory in patterns that would collide on the memory banks if the layout were naive, so CuTe composes a Swizzle onto the shared-memory layout to scatter addresses across banks and avoid conflicts, and that swizzle is, again, just another layout in the composition.

Deep dive: the collective builder and kernel schedules

The 3.x library is organized around collectives. A CollectiveMainloop owns everything about streaming A and B and running the MMAs, and a CollectiveEpilogue owns everything about turning accumulators into output. Above them a GemmUniversal kernel glues a mainloop, an epilogue, and a tile scheduler, and a GemmUniversalAdapter wraps the kernel with the host-side launch machinery. You could assemble a collective by hand, but the CollectiveBuilder exists to choose the hard parts for you. Given the architecture, the types, a tile shape, and a cluster shape, it selects the copy atoms, the MMA atom, the number of pipeline stages that fit in shared memory, and the warp schedule. The schedule is the interesting choice on Hopper. KernelScheduleAuto resolves to a warp-specialized schedule, and the named variants trade off differently. A cooperative schedule has all consumer warps collaborate on one output tile, while a pingpong schedule has two consumer warpgroups alternate on two tiles so one can run the epilogue while the other runs math. Which wins depends on the shape and the epilogue cost, which is why both exist and why the profiler is the way to decide. The exact spelling of these schedule tags and builder parameters has moved between 3.x releases, so learn the roles and check the headers under include/cutlass/gemm/collective/ for the version you have.

Deep dive: data types and quantization

CUTLASS defines its own numeric types in cutlass/numeric_types.h so the templates can carry precision in the type system. Beyond half_t and bfloat16_t there is tfloat32_t, the reduced-mantissa TF32 format the tensor cores use for fast fp32-ish math, the fp8 formats float_e4m3_t and float_e5m2_t, and narrow integer types down to int4. The point of first-class narrow types is that they raise arithmetic intensity twice over. Fewer bytes per element means the ridge-point math from the first deep dive needs a smaller tile to become compute-bound, and the tensor cores run narrow types at higher FLOP rates, so quantized GEMM is where a great deal of inference speed comes from. This is exactly the machinery inference engines reach for. A vLLM or TensorRT-LLM fp8 or int8 matmul is a CUTLASS GEMM with a narrow input type and an epilogue that fuses the dequantization scale, and a mixture-of-experts layer becomes a grouped GEMM where each expert is one problem in a batch of differently-shaped problems. The newest Blackwell tensor cores push this further with block-scaled four-bit formats, where a small group of elements shares a scale factor, and CUTLASS models those as paired data-and-scale operands in the same collective framework.

Part VI: Reading the repository

The repository is large but well-partitioned, and you can read the parts that matter without touching the rest. Paths below reflect the 3.x and 4.x layout as of mid 2026, and a few may have shifted.

Stage 0, orientation. Read the top-level README.md for the feature and architecture matrix, then media/docs/efficient_gemm.md for the hierarchical GEMM picture in NVIDIA's words, then browse media/docs/cute/ for the CuTe introduction. Questions to hold, what is the difference between the 2.x and 3.x APIs, what does CuTe replace, and which architectures does your checkout support.

Stage 1, CuTe first. Everything else rests on it. Read include/cute/layout.hpp and include/cute/tensor.hpp for the core types, then the algorithms in include/cute/algorithm/ (notably gemm.hpp and copy.hpp), then the atoms in include/cute/atom/ (mma_atom.hpp and copy_atom.hpp). Run the CuTe examples under examples/cute/ and print layouts as you go. Questions, what does a layout actually compute, what does local_partition do to a tile, and what is a thread-value layout.

Stage 2, one GEMM end to end. Start with a basic example such as examples/00_basic_gemm to see the 2.x device API, then a Hopper example such as the warp-specialized GEMM examples in the higher-numbered examples/ directories to see the 3.x collective path. Follow the include chain from the example into include/cutlass/gemm/device/, then gemm/kernel/, then gemm/collective/. Questions, where is the mainloop, where is the tile scheduler chosen, and where does the epilogue attach.

Stage 3, the collective internals. Read the collective mainloop and its pipeline in include/cutlass/gemm/collective/, the epilogue and its fusion nodes in include/cutlass/epilogue/ including epilogue/fusion/ for the visitor trees, and the architecture-specific instruction wrappers in include/cutlass/arch/ and the CuTe include/cute/arch/. Questions, how does the multistage pipeline synchronize producers and consumers, and how does an Epilogue Visitor Tree turn a fusion into code.

Stage 4, the tools. Read tools/profiler/ to understand how kernels are enumerated and benchmarked, tools/library/ for how instances are generated and registered, and tools/util/ for the host-side reference GEMM and tensor helpers you use when checking your own kernel against a ground truth. Questions, how does the profiler pick which kernels apply to a problem, and how would you add a new instance to the library.

Stage 5, the frontier. The python/ tree holds the Python interface and the CuTe DSL, the include/cutlass/conv/ tree holds the implicit-GEMM convolutions, and the grouped and mixture-of-experts GEMM examples show the irregular-shape machinery. Where not to start, the architecture-specific Blackwell paths and the block-scaled narrow-precision types are the most in-flux corner of the codebase and make more sense once the dense Hopper story is solid.

Part VII: Hands-on labs

Labs 1 and 2 need only a checkout and a compiler-visible GPU. Labs 3 through 5 want a tensor-core GPU, Ampere or newer. Log and menu formats vary with the release.

Lab 1: print a layout. Concept: CuTe layouts as functions.

// layout_lab.cu  -- compile with nvcc -I include, run on host or device
#include <cute/tensor.hpp>
using namespace cute;
int main() {
  auto col = make_layout(make_shape(_4{}, _2{}), make_stride(_1{}, _4{}));
  auto row = make_layout(make_shape(_4{}, _2{}), make_stride(_2{}, _1{}));
  print_layout(col);   // column-major table
  print_layout(row);   // row-major table
  // predict the offset of coord (2,1) for each BEFORE running
  print(col(make_coord(2,1))); print("\n");
  print(row(make_coord(2,1))); print("\n");
}

Predict both offsets from the stride before you run, then confirm. Change one shape to a nested tuple such as make_shape(make_shape(_2{},_2{}), _2{}) and read how the hierarchical layout prints.

Lab 2: enumerate and benchmark with the profiler. Concept: the instance library and tile-shape sensitivity.

# after building cutlass_profiler for your arch
./tools/profiler/cutlass_profiler --operation=Gemm \
    --A=f16:row --B=f16:col --C=f16:row --accum=f32 \
    --m=4096 --n=4096 --k=4096

# now make it skinny and watch the winning kernel change
./tools/profiler/cutlass_profiler --operation=Gemm \
    --A=f16:row --B=f16:col --C=f16:row --accum=f32 \
    --m=16 --n=16 --k=32768

Note which tile shape and which kernel win the square problem, then note how the winner changes for the tall-skinny K-heavy problem, and connect that to the arithmetic-intensity argument and to split-K.

Lab 3: a first device GEMM. Concept: the 2.x device API and correctness checking.

// build the fp16 Gemm from Part II, fill A and B with known values,
// run it, then compare against the host reference in tools/util:
#include <cutlass/util/reference/host/gemm.h>
// compute D_ref on the host, then assert max_abs_error(D, D_ref) is small.
// deliberately misalign lda by 1 and observe can_implement() reject it.

The lesson is the workflow, instantiate, run, and always validate against tools/util's reference before trusting a kernel. The alignment failure teaches why leading dimensions matter.

Lab 4: fuse an activation into the epilogue. Concept: epilogue as a swappable part.

// swap the default LinearCombination epilogue for a variant that applies
// a ReLU or GELU after alpha*AB + beta*C, e.g. LinearCombinationRelu, and
// confirm the fused kernel matches "GEMM then activation" done separately,
// at lower total time because the accumulator never round-trips memory.

Time the fused kernel against running the GEMM and the activation as two passes. The gap is the extra memory traffic you removed, which is the whole point of epilogue fusion.

Lab 5: quantized GEMM. Concept: narrow types and dequantization fusion.

# find an fp8 or int8 GEMM instance in the profiler and benchmark it
./tools/profiler/cutlass_profiler --operation=Gemm \
    --A=e4m3:row --B=e4m3:col --C=f16:row --accum=f32 \
    --m=8192 --n=8192 --k=8192

Compare the fp8 throughput against the fp16 result from Lab 2 on the same shape, and reason about why narrower operands move the ridge point and let a smaller tile become compute-bound. This is the exact kernel shape inference engines use for quantized matmuls.

Part VIII: Questions and model answers

Understanding checks. Answer aloud before reading.

1. What is CUTLASS in one sentence?

A header library of CUDA C++ templates, built on the CuTe layout algebra, for assembling matrix-multiply and convolution kernels that decompose one GEMM into a hierarchy of tiles matching the GPU memory levels and run at close to peak tensor-core throughput.

2. Why is a GEMM tiled at three levels instead of one?

Because no single storage level is both large enough to hold a tile big enough for high arithmetic intensity and fast enough to feed the tensor cores. The threadblock tile buys reuse against global memory by living in shared memory, the warp tile buys reuse against shared memory by living in registers, and the instruction tile sits at the tensor-core inputs. The product of the per-level reuse is what lifts the kernel onto the compute roofline.

3. What is a CuTe Layout?

A shape paired with a stride, understood as a function from a logical coordinate to a linear offset. Layouts can nest, so one value can describe a tile, a tile-of-tiles, or the mapping of threads onto data, and they compose under a real algebra of composition, product, and division.

4. What is an MMA atom and how does a TiledMMA use it?

An MMA atom wraps exactly one tensor-core instruction together with the thread-value layouts describing which thread holds which operand and result element. A TiledMMA replicates that atom across the warps of a block to cover the warp tile and exposes partition methods so each thread finds its fragments, so the block-level GEMM is the atom tiled by the algebra.

5. What happens in the mainloop, and why is it a pipeline?

The block walks K in tile steps, and at each step it loads the next A and B tiles into a free shared-memory buffer while the tensor cores multiply a buffer a previous step already filled. Multi-buffering and pipeline barriers overlap the loads with the math so the tensor cores rarely wait on memory, which is the difference between reaching peak and stalling.

6. How does Hopper's mainloop differ from Ampere's?

Ampere uses cp.async to copy global to shared without staging through registers, then ldmatrix to load operand fragments into registers for mma.sync. Hopper uses the TMA engine for bulk asynchronous tile copies and the warpgroup-wide wgmma instruction that reads operands directly from shared memory, and it typically warp-specializes producers doing TMA against consumers doing wgmma.

7. What is epilogue fusion and why does it help?

The epilogue turns the raw accumulator into the stored output, and fusion lets you express a graph of extra elementwise work there, a bias add, an activation, a dequant scale, a reduction, so it runs while the accumulator is still in registers. It helps because it removes the extra global-memory round trips a separate pass would cost, and for memory-bound epilogues that traffic is the dominant cost.

8. State the roofline condition for a GEMM to be compute-bound.

Its arithmetic intensity, FLOP divided by bytes moved, must exceed the ridge point, which is peak FLOP/s divided by peak bandwidth. For a square tile of side T with b bytes per element the tile intensity is about T over b, so larger tiles and narrower types both push the kernel above the ridge and onto the compute roofline.

9. When would you choose Triton or cuBLAS over CUTLASS?

cuBLAS or cuBLASLt when a plain or lightly-fused GEMM suffices and you want a tuned function to call rather than a toolkit to assemble. Triton when you want to write a fused kernel in Python quickly and can accept slightly less peak and less control over the exact instructions. CUTLASS when you need an unusual fusion, a quantized or mixed-input type, an irregular grouped shape, or day-one performance on the newest hardware.

10. How do FlashAttention and inference engines use CUTLASS?

They build their fused kernels on CuTe atoms and CUTLASS collectives rather than reinventing tensor-core plumbing. FlashAttention's Hopper kernels use CuTe MMA and copy atoms, TMA, and warp specialization, and engines like vLLM, SGLang, and TensorRT-LLM use CUTLASS GEMMs for their quantized fp8 and int8 matmuls and for grouped mixture-of-experts layers, with the dequantization fused into the epilogue.

11. Why can the same template machinery do convolution?

CUTLASS implements convolution as an implicit GEMM. The im2col matrix that a convolution is mathematically equivalent to is never materialized in memory. Instead the mainloop forms the needed tiles on the fly from the input tensor, so the same tiled mainloop and epilogue that serve a GEMM serve a convolution.

12. A tall-skinny GEMM, large K with tiny M and N, runs far below peak. Why, and what fixes it?

There are too few output tiles to fill the GPU, so most streaming multiprocessors sit idle, and the intensity is low enough that bandwidth caps you regardless. Splitting the K dimension across blocks, split-K or the load-balanced Stream-K scheduler, spreads the work over the machine and then reduces the partial sums, restoring occupancy.

Part IX: Design lessons

Make the layout a value you compute with. CuTe turns the thread-to-data mapping into a first-class object with an algebra, so tiling and partitioning become operations you can read, print, and compose instead of index arithmetic smeared through the kernel. The same instinct shows up wherever a hard invariant is lifted into data you can manipulate, shapes in an array language, layouts in a database planner, types in a compiler.

Match the abstraction to the memory hierarchy. The three tile levels are not arbitrary, each one exists to buy reuse against exactly one level of storage. When your abstraction layers mirror the physical cost structure, performance reasoning becomes local, you can look at one level and know what it is responsible for. This is the same discipline as designing cache-oblivious algorithms or blocking a numerical kernel to the L1 and L2 sizes.

Put the configuration in the type system. CUTLASS encodes tile shapes, data types, and schedules as template parameters, so an invalid kernel usually fails to compile rather than failing at runtime. The cost is famously long error messages, and the benefit is that a kernel that builds is a kernel the library has already largely validated. It is the same trade a strongly-typed builder or a session-typed protocol makes, move errors left, pay in up-front strictness.

Separate the parts of a kernel. Splitting a GEMM into a collective mainloop, a collective epilogue, and a tile scheduler means you can swap the epilogue for a fused one, swap the scheduler for Stream-K, or retarget the mainloop to a new tensor-core instruction without rewriting the others. Decomposing along the axes that actually vary is what lets one library cover thousands of kernel variants from a handful of composable pieces.

Ship the measurement tool with the library. The profiler is not an afterthought, it is how you discover that the best tile shape depends on the problem shape, which is a fact the library cannot know for you. Building enumeration and benchmarking into the toolkit makes tuning a first-class activity rather than a bespoke script every user has to write.

Part X: Memorization framework

The one-sentence summary. CUTLASS decomposes one matrix multiply into threadblock, warp, and instruction tiles, streams the operands through shared memory and registers with a pipelined mainloop so loads hide under tensor-core math, and fuses the finishing work into the epilogue, all described by the CuTe layout algebra so the whole kernel is a composition of a few well-typed pieces.

D = alpha*A@B + beta*C
  -> tile scheduler assigns each block an Mt x Nt output tile (swizzle / Stream-K)
  -> prologue: CuTe local_tile the slabs of A,B, build TiledMMA, zero accumulator
  -> mainloop over K:  load next tile (TMA / cp.async)  ||  MMA (wgmma / mma.sync)
                       multi-buffered shared memory overlaps load and compute
  -> epilogue: accumulators -> alpha/beta, bias, activation, fused EVT -> store D
  -> optional split-K / Stream-K reduction of partial sums

The pieces mapped to the tree:

layout algebra    include/cute/            (layout.hpp, tensor.hpp, atom/, algorithm/)
GEMM building     include/cutlass/gemm/    (device/, kernel/, collective/, warp/, threadblock/)
epilogue+fusion   include/cutlass/epilogue/ (thread/, collective/, fusion/ for EVT)
instructions      include/cutlass/arch/ and include/cute/arch/  (mma, cp.async, tma)
convolution       include/cutlass/conv/    (implicit GEMM)
tools             tools/profiler/, tools/library/, tools/util/ (reference + benchmark)
docs              media/docs/ (efficient_gemm.md, cute/, gemm api notes)

Memorize these blocks:

• Tile hierarchy: threadblock tile lives in shared memory, warp tile in registers, instruction tile at the tensor-core inputs. Each level buys reuse against the next-slower memory.
• Roofline: attainable = min(peak_flops, intensity * bandwidth). Ridge point = peak_flops / bandwidth. Square-tile intensity is about T/b, so bigger tiles and narrower types both push toward compute-bound.
• CuTe: Layout = Shape:Stride is a coord-to-offset function. Atoms wrap one instruction, TiledMMA and TiledCopy tile atoms across threads, the mainloop is a TiledCopy feeding a TiledMMA.
• Mainloop: multistage pipeline over K, loads (TMA on Hopper, cp.async on Ampere) overlap MMAs (wgmma on Hopper, mma.sync on Ampere) via multi-buffered shared memory and pipeline barriers.
• Epilogue: alpha*AB + beta*C by default, or an Epilogue Visitor Tree fusing bias, activation, dequant, and reduction while the accumulator is still in registers.

Part XI: Papers and further reading

The ideas in this walkthrough trace back to a short list of papers and NVIDIA documents, and each one rewards a direct read. Where this site derives the same idea in depth, the companion link points there.

1. Kerr et al., CUTLASS, Fast Linear Algebra in CUDA C++, 2017. The article that introduced the library and the threadblock-warp-thread tile hierarchy this chapter walks.
2. NVIDIA, Efficient GEMM in CUDA, maintained in-tree. The library's own account of the tiled mainloop, software pipelining, split-K and sliced-K, and warp specialization, the best single companion to Parts IV and V.
3. Williams et al., Roofline, An Insightful Visual Performance Model for Multicore Architectures, CACM 2009. The source of the arithmetic-intensity argument that explains why every tile level exists. The parallel computing class on this site builds the same model from scratch.
4. Markidis et al., NVIDIA Tensor Core Programmability, Performance and Precision, 2018. An early measured look at what the tensor cores actually deliver and what feeding them takes. The hardware context lives in the advanced systems architecture class.
5. Andersch et al., NVIDIA Hopper Architecture In-Depth, 2022. The authoritative description of TMA, thread block clusters, and the H100 tensor cores behind the wgmma mainloop of Part IV.
6. NVIDIA, NVIDIA Blackwell Architecture Technical Overview. The vendor overview of the newest tensor cores and the block-scaled narrow formats Part V touches.
7. Osama et al., Stream-K, Work-centric Parallel Decomposition for Dense Matrix-Matrix Multiplication on the GPU, 2023. The load-balanced tile scheduler that fixes the tall-skinny GEMM regime, written by CUTLASS authors and shipped in the library.
8. Dao et al., FlashAttention, Fast and Memory-Efficient Exact Attention with IO-Awareness, 2022. The most famous kernel built on the ideas here, traced in the FlashAttention walkthrough and the online softmax note.
9. Shah et al., FlashAttention-3, Fast and Accurate Attention with Asynchrony and Low-precision, 2024. The Hopper rewrite that leans directly on CuTe atoms, TMA, and warp specialization, a case study in using CUTLASS as a substrate.
10. Tillet et al., Triton, An Intermediate Language and Compiler for Tiled Neural Network Computations, MAPL 2019. The main alternative programming model from Part III, covered in the Triton walkthrough.
11. Micikevicius et al., FP8 Formats for Deep Learning, 2022. Defines the e4m3 and e5m2 types behind the quantized GEMMs of Part V. The precision tradeoffs are worked in the mixed precision note.

Part XII: Final takeaway

If the memory-hierarchy and roofline ideas underneath all of this are the gap, the parallel computing class builds them from the ground up, and the attention-kernel story that leans hardest on CuTe is traced in the FlashAttention chapter and the online softmax derivation. Then come back and read one Hopper GEMM example top to bottom. Once the tile hierarchy, the pipelined mainloop, and the layout algebra are in your head, the kernel reads like a direct transcription of the roofline argument, which is the entire point.