CUTLASS crash course for beginners.

CUTLASS is NVIDIA's open-source C++ (and now Python) template library for writing GPU matrix-multiply kernels that run close to the speed of hand-tuned libraries like cuBLAS. This is a from-zero walkthrough for someone who has written CUDA before but has never opened CUTLASS: why matrix multiplication is deceptively hard to make fast on a GPU, the tiling hierarchy CUTLASS is built around, how to instantiate your first GEMM kernel, what changed in the CUTLASS 3.x rewrite (CuTe, collective builders), and how to actually build and run something from the repo.

1. What CUTLASS actually is

Start with the official one-line description, then unpack every word in it: CUTLASS is "a collection of CUDA C++ template abstractions and Python domain-specific languages designed to enable high-performance matrix-matrix multiplication (GEMM) and related computations."

• "Collection of C++ template abstractions" — CUTLASS isn't a black-box function you call. It's a toolbox of C++ template classes, each representing one layer of how a matrix multiply gets decomposed and executed on a GPU. You compose these templates together, and the compiler generates a specialized kernel for your exact combination of data types, matrix shapes, and target GPU.
• "GEMM" — General Matrix Multiply, the operation \( C = \alpha \cdot A \times B + \beta \cdot C \). It sounds narrow, but GEMM is the single most important compute pattern in modern deep learning — every linear layer, every attention projection, and (via a reformulation covered later) every convolution boils down to it.
• "High-performance" — the entire reason CUTLASS exists. A GEMM kernel you write yourself in an afternoon will typically run at 10-30% of a GPU's peak throughput. cuBLAS (NVIDIA's closed-source, hand-tuned GEMM library) routinely hits 80-95%+. CUTLASS's goal is to get you into that same range, but with source code you can read, modify, and specialize — which cuBLAS, being a compiled black box, never lets you do.

Put differently: CUTLASS exists because "write a fast GEMM kernel" is genuinely one of the hardest one-function programming problems in all of systems software, and NVIDIA decided to open-source their own internal solution to it rather than let every ML framework and inference engine reinvent it badly.

2. Why is matrix multiplication hard to make fast on a GPU?

This is the question that makes the rest of CUTLASS's design make sense. If you've never asked it, the natural assumption is "just write three nested loops" — and that assumption is exactly what makes GEMM a trap for beginners.

// The "obviously correct" GEMM. Also catastrophically slow on a GPU.
for (int i = 0; i < M; i++)
  for (int j = 0; j < N; j++)
    for (int k = 0; k < K; k++)
      C[i][j] += A[i][k] * B[k][j];

This is correct. It is also, if run naively with each thread reading directly from global memory (GPU DRAM) on every access, roughly an order of magnitude or more slower than a well-tuned kernel. The reason is a mismatch between how much arithmetic GEMM does versus how much data movement it needs — a relationship engineers call arithmetic intensity.

For an \(M \times K\) matrix times a \(K \times N\) matrix, the naive loop performs \(2MNK\) floating-point operations (one multiply and one add per inner-loop iteration) but re-reads the same elements of \(A\) and \(B\) from memory over and over — element \(A[i][k]\) gets re-read for every one of the \(N\) values of \(j\); element \(B[k][j]\) gets re-read for every one of the \(M\) values of \(i\). GPU global memory bandwidth is enormous compared to a CPU's, but it is still orders of magnitude slower than the GPU's arithmetic throughput — modern GPUs can do far more floating-point operations per second than they can move bytes per second. A kernel that re-fetches data instead of reusing it will spend nearly all its time waiting on memory, not computing — this is what "memory-bound" means, and it's the default failure mode for a naive GEMM.

The standard fix — and the organizing principle behind CUTLASS's entire design — is tiling: instead of one thread reading directly from global memory for every single multiply-add, cooperatively load a block (a "tile") of \(A\) and a tile of \(B\) into fast on-chip shared memory once, and then have many threads reuse that same on-chip copy for many multiply-adds before moving to the next tile. This single idea — load once, reuse many times, at every level of the hardware hierarchy — is what the rest of this article is really about.

3. A 60-second refresher on the GPU memory hierarchy

Tiling only makes sense once you can picture the hardware it's designed around. A GPU has (at least) three tiers of memory, each roughly an order of magnitude faster — and smaller — than the one below it.

CUTLASS's tiling hierarchy is essentially a direct mirror of this memory hierarchy — a tile of work sized to fit in shared memory gets staged there once and reused by many threads; a smaller tile sized to fit in registers gets staged there and reused by many arithmetic instructions. Every level of tiling described in the next section exists to match a level of this pyramid.

4. The CUTLASS 2.x hierarchy: device, kernel, threadblock, warp, thread

CUTLASS structures a GEMM as five nested levels, each owned by a different unit of GPU execution, each responsible for tiling the problem down to what the level below can handle. This is the "classic" CUTLASS design (versions 2.x, targeting Volta through Ampere/Turing) — the newer 3.x model in §6 restructures the lower levels, but this hierarchy is still the right mental model to learn first because every later concept is defined relative to it.

Level	CUTLASS namespace	What it does
Device	cutlass::gemm::device::Gemm	The host-callable entry point. Takes matrix pointers, shapes, and launches the CUDA kernel grid.
Kernel	cutlass::gemm::kernel::*	Runs once per threadblock. Computes which output tile this threadblock owns and drives the main loop over the K dimension.
Threadblock	cutlass::gemm::threadblock::Mma	Cooperatively loads a tile of A and a tile of B from global memory into shared memory, once per K-loop iteration, for the whole threadblock to reuse.
Warp	cutlass::gemm::warp::Mma	A group of 32 threads reads a sub-tile out of shared memory into registers and issues Tensor Core (or CUDA core) matrix-multiply-accumulate instructions.
Thread / Instruction	cutlass::arch::Mma	The literal hardware instruction — a single Tensor Core MMA op, or a scalar fused-multiply-add on older architectures without Tensor Cores.

The tile size chosen at each level — usually written as ThreadblockShape::{kM, kN, kK} and WarpShape::{kM, kN, kK} in code — is one of the biggest levers for performance, and it's a genuine tradeoff, not a free parameter to just maximize: a larger threadblock tile means fewer, larger reads from global memory (good — less bandwidth pressure), but it also means each thread needs to hold more of the accumulator in registers. Registers are the scarcest resource in the whole hierarchy (a few hundred per thread), and accumulator elements alone can eat up half a thread's entire register budget at large tile sizes — leaving too few registers active at once hurts the GPU's ability to hide memory latency by switching between threads. Picking tile sizes is exactly the kind of problem CUTLASS's template system is built to make explorable: change one template parameter, recompile, and get a completely different — but still correct — kernel.

4.1 Software pipelining: hiding latency, not just avoiding it

Tiling alone gets data close to the compute units, but there's still a real trip to global memory required to bring in the next tile of A and B for every iteration of the K-loop. CUTLASS overlaps that fetch with the current tile's computation using double buffering: at the threadblock level, two shared-memory tiles exist side by side — while warps compute against tile 0, the hardware's asynchronous copy engines are already fetching tile 1's data from global memory in the background. The same double-buffering trick is applied one level down, with register fragments, at the warp level. The effect is that the (unavoidable) memory latency for tile N+1 is hidden entirely behind the (already-happening) compute for tile N, rather than being paid as dead time between iterations.

5. Your first GEMM: reading a CUTLASS 2.x template instantiation

Everything above was theory. Here's what it looks like as actual code — the classic CUTLASS 2.x pattern for declaring a half-precision GEMM that runs on Tensor Cores on an SM75 (Turing) GPU:

#include <cutlass/gemm/device/gemm.h>
#include <cutlass/util/host_tensor.h>

using Gemm = cutlass::gemm::device::Gemm<
  cutlass::half_t,                    // ElementA — fp16 input
  cutlass::layout::ColumnMajor,       // LayoutA
  cutlass::half_t,                    // ElementB — fp16 input
  cutlass::layout::ColumnMajor,       // LayoutB
  cutlass::half_t,                    // ElementOutput — fp16 output
  cutlass::layout::ColumnMajor,       // LayoutOutput
  float,                              // ElementAccumulator — fp32 accumulation
  cutlass::arch::OpClassTensorOp,     // use Tensor Cores
  cutlass::arch::Sm75                 // target: Turing
>;

Gemm gemm_op;
cutlass::Status status = gemm_op({
  {M, N, K},
  A.device_ref(), B.device_ref(),
  C.device_ref(), C.device_ref(),
  {alpha, beta}
});

Every one of those template parameters is a decision CUTLASS is handing you explicitly, rather than hiding inside a compiled library like cuBLAS would:

• Element types can differ per matrix. Notice ElementA/ElementB are fp16 but ElementAccumulator is fp32 — this is mixed-precision accumulation, and it's not an edge case, it's the default recipe for every serious low-precision kernel: multiply in a narrow, fast format, but accumulate the running sum in a wider format so you don't lose precision over thousands of additions. The same pattern extends to fp8, and to the block-scaled MXFP4/MXFP6/MXFP8 formats CUTLASS also supports for the newest hardware.
• Layout is explicit, per matrix, and matters. ColumnMajor vs RowMajor for A and B individually change which memory-access pattern is contiguous — get this wrong and you silently get correct-but-slow code, since a mismatched layout defeats the coalesced-access patterns the kernel was tuned around.
• OpClassTensorOp is an architecture decision, not a performance hint. It tells CUTLASS to target the GPU's dedicated matrix-multiply hardware (Tensor Cores) instead of doing scalar fused-multiply-adds on ordinary CUDA cores — Tensor Cores are what actually gets you into cuBLAS's performance range on any GPU built since Volta.
• Sm75 pins the target architecture at compile time — CUTLASS is a template library, so the actual specialized kernel is a different, separately-compiled binary for every architecture you target; there's no runtime dispatch magic hiding underneath.

This is the essential trade CUTLASS makes versus calling cuBLAS: you write (or generate) more code, and you compile once per configuration, but every one of these decisions is visible and changeable — which is exactly what you need if your matrix shapes, data types, or fusion requirements don't match one of cuBLAS's pre-baked kernel variants.

6. CUTLASS 3.x: CuTe, collectives, and why the redesign happened

CUTLASS 3.0 (targeting Hopper and newer) is a substantial rewrite, not an incremental update — and understanding why it changed is more useful than memorizing the new API by itself. Two things drove it:

1. New hardware features had no clean place in the 2.x hierarchy. Hopper introduced thread-block clusters (groups of threadblocks that can cooperate and communicate directly, not just independently), the Tensor Memory Accelerator (TMA, a dedicated async copy engine for bulk tensor transfers), and warp-group-level MMA instructions. None of these map cleanly onto "threadblock → warp → thread" — they need a level between kernel and threadblock, and a fundamentally more flexible way to describe data movement than the fixed layout templates 2.x used.
2. The 2.x layout system didn't generalize. 2.x described tensor layouts (row-major, column-major, and a handful of special swizzled formats for shared memory) as a fixed enumeration of cases. Every new hardware feature or data format tended to need a new hand-written layout class. CuTe replaces that with an actual algebra — a small set of composable operations that can express essentially any layout, including the exotic swizzled and hierarchical layouts modern Tensor Core and TMA hardware need.

6.1 CuTe: layouts as shape + stride, composed algebraically

CuTe's central object is a Layout, and a Layout is just two pieces of information glued together: a Shape (how many valid coordinates exist in each dimension — the tensor's logical size) and a Stride (how many elements to jump in linear memory to move by one step in each dimension). A simple example: a 4×8 row-major matrix has shape (4,8) and stride (8,1) — moving one step in the row direction jumps 8 elements in memory; moving one step in the column direction jumps 1 element. That's it — that's a Layout.

What makes CuTe powerful isn't that basic definition, it's that shapes and strides can be arbitrarily nested (a shape can itself contain a tuple of shapes), and CuTe defines an algebra of operations — composition, complement, logical division, logical product — that let you combine Layouts to express tiling, partitioning across threads, and memory swizzling, all through the same small vocabulary instead of a new hand-written class per case. In practice this is what lets a single CuTe-based kernel description cleanly target wildly different physical memory patterns (a swizzled shared-memory tile, a TMA-friendly global-memory descriptor, a per-thread register fragment) without those being conceptually different kinds of objects — they're all just Layouts, composed differently.

6.2 The 3.x hierarchy: Device → Kernel → Collective → Tiled MMA/Copy → Atom

The 3.x hierarchy keeps Device and Kernel at the top (same idea as before — host entry point, then grid-level scheduling) but replaces "threadblock → warp → thread" with a new middle layer built around the word Collective — defined in the docs as "the largest collection of threads onto which mma atoms and copy atoms are tiled." A Collective can span a full thread-block cluster, not just one threadblock, which is precisely the hardware capability 2.x had no way to describe.

3.x concept	Lives in	Role
CollectiveMainloop	cutlass::gemm::collective::CollectiveMma	The main K-dimension loop — the direct successor to 2.x's threadblock-level Mma, but cluster-aware and TMA-capable.
CollectiveEpilogue	cutlass::epilogue::collective::	Post-matmul work: applying alpha/beta scaling, writing the result, optionally fusing in an activation function.
GemmUniversal (Kernel layer)	cutlass::gemm::kernel::	Orchestrates the collectives across the grid. Deliberately "stateless" and "universal" — it can dispatch to either 3.x or 2.x-style collectives via template introspection.
GemmUniversalAdapter (Device layer)	cutlass::gemm::device::	Host-side interface — marshals user-facing Arguments into kernel Params and handles cluster-aware kernel launches.
CollectiveBuilder	cutlass::gemm::collective::	A convenience layer that takes familiar 2.x-style parameters (element types, tile shape) and auto-generates an optimized CollectiveMma, so you don't have to hand-assemble every low-level template argument yourself.

// CUTLASS 3.x: building a collective mainloop via CollectiveBuilder
using CollectiveMainloop = typename cutlass::gemm::collective::CollectiveBuilder<
  cutlass::arch::Sm90,
  cutlass::arch::OpClassTensorOp,
  ElementA, LayoutA, AlignmentA,
  ElementB, LayoutB, AlignmentB,
  ElementAccumulator,
  TileShape, ClusterShape,
  cutlass::gemm::collective::StageCountAuto,
  cutlass::gemm::collective::KernelScheduleAuto
>::CollectiveOp;

Notice the shift in spirit versus the 2.x example in §5: you're no longer hand-picking a fixed Mma implementation — you're describing requirements (architecture, tile shape, cluster shape) and letting StageCountAuto/KernelScheduleAuto pick good defaults, with the option to override any of them explicitly once you know what you're tuning. This is CUTLASS's dispatch-policy pattern: a given mainloop implementation declares which kernel schedule it requires (for example, one named MainloopSm90TmaGmmaWarpSpecialized pairs with a schedule called KernelTmaWarpSpecializedCooperative), which lets many different kernel schedules compose with the same underlying mainloop implementation instead of every combination needing its own hand-written specialization — the same "avoid combinatorial duplication" motivation that drove CuTe's layout algebra shows up again here at the scheduling level.

7. Beyond GEMM: convolution as implicit GEMM

CUTLASS's scope isn't limited to literal matrix multiplication — it also supports convolution, via a reformulation called implicit GEMM. The insight: a convolution's sliding-window dot products can be re-expressed as a single (larger) matrix multiply, by unfolding overlapping input patches into rows of an implicit matrix and treating the convolution filters as the other operand — without ever materializing that unfolded matrix in memory (which is what makes it "implicit," as opposed to the classic and much more memory-hungry "im2col" approach). This lets convolution reuse essentially all of the same tiling, pipelining, and Tensor Core machinery already built for GEMM, rather than needing an entirely separate optimized code path.

8. Building CUTLASS and running your first kernel

CUTLASS is header-heavy template code plus a CMake build for its tests, examples, and profiler tool. Prerequisites: CUDA Toolkit 11.4+ (12.0+ recommended), CMake 3.18+, and a C++17-capable compiler (g++ 7.5.0 minimum).

# point CMake at nvcc, then configure for your target GPU architecture
export CUDACXX=${CUDA_INSTALL_PATH}/bin/nvcc
mkdir build && cd build
cmake .. -DCUTLASS_NVCC_ARCHS=90a    # Hopper
# cmake .. -DCUTLASS_NVCC_ARCHS=100a   # Blackwell
# cmake .. -DCUTLASS_NVCC_ARCHS=80     # Ampere

A full build (all tests, all example kernels, every data-type combination) takes a genuinely long time — CUTLASS's template instantiation surface is huge. For a first pass, restricting the build is worth doing:

# faster: build just the profiler, skip the full test suite
cmake .. -DCUTLASS_NVCC_ARCHS=90a -DCUTLASS_ENABLE_TESTS=OFF \
  -DCUTLASS_UNITY_BUILD_ENABLED=ON

# or filter by operation / kernel name if you know what you want
cmake .. -DCUTLASS_NVCC_ARCHS=90a -DCUTLASS_LIBRARY_OPERATIONS=conv2d
cmake .. -DCUTLASS_NVCC_ARCHS=90a -DCUTLASS_LIBRARY_KERNELS=cutlass3x*

Once built, the profiler tool is the fastest way to see CUTLASS actually run and benchmark itself against a specific problem shape — no need to write any code yet:

make cutlass_profiler -j12
./tools/profiler/cutlass_profiler --kernels=sgemm --m=4352 --n=4096 --k=4096

This prints the matched kernel, timing, achieved memory bandwidth, and GFLOP/s — a genuinely useful first exercise is running the same M/N/K shape across a few different kernel names and comparing throughput, to get a feel for how much tile-size and data-type choice actually move the needle in practice, before writing a single line of your own template code.

9. A suggested learning path from here

1. Run the profiler (§8) against a few problem shapes and kernel names — build intuition for what "fast" looks like before reading more theory.
2. Read (don't skim) the examples/ directory in the CUTLASS repo, starting from the lowest-numbered basic GEMM examples before jumping to fused-epilogue or grouped-GEMM examples.
3. Learn CuTe's Layout and layout-algebra docs directly — §6.1 above is deliberately a simplified entry point; the real payoff of CuTe only shows up once you've worked through composing a few layouts by hand.
4. Instantiate your own 2.x GEMM first, changing one template parameter at a time (data type, then layout, then tile shape) and re-running the profiler to see the effect — before attempting a 3.x collective-builder kernel, since 3.x assumes you already have the 2.x mental model as a baseline.
5. Only then move to 3.x collectives and CuTe kernels if you're targeting Hopper/Blackwell specifically — the extra conceptual machinery (clusters, TMA, dispatch policies) pays for itself only once you actually need those hardware features.

FAQ

Do I need CUTLASS if I can just call cuBLAS?

For a standard, well-shaped GEMM on a supported data type, cuBLAS is usually the right first choice — it's a mature, hand-tuned closed-source library and will often match or beat a first attempt at a CUTLASS kernel. CUTLASS earns its complexity when you need something cuBLAS doesn't offer as a pre-baked kernel: a fused epilogue (activation function, bias add, or another matmul fused directly into the same kernel), an unusual data type or mixed-precision combination, or when you're building an inference engine or kernel-generation tool that needs source-level control over the kernel rather than a fixed library call.

Is CUTLASS 3.x a strict replacement for 2.x?

Not in the sense of deprecating it outright — 2.x is still the right model for Volta/Turing/Ampere-class GPUs that don't have thread-block clusters or TMA, and the GemmUniversal kernel layer in 3.x is explicitly designed to still dispatch to 2.x-style collectives via template introspection. But for new development on Hopper or newer, 3.x's collective-builder pattern is the recommended entry point.

What's the practical difference between CUTLASS's Python DSL and the C++ templates?

The C++ templates are compile-time specialized — every configuration is its own separately-compiled kernel, which gives maximum performance but a slow edit-compile-run loop. The Python DSL layer uses JIT compilation for a much faster iteration cycle while still generating kernels reported to run within about 2% of hand-written C++ on Blackwell — a meaningfully smaller gap than Python bindings around compute kernels have historically carried.

Why does accumulation happen in a wider type than the inputs?

Summing thousands of low-precision products (as any GEMM's K-dimension reduction does) accumulates rounding error faster in a narrow format like fp16 or fp8 than in fp32. Multiplying in the narrow, fast format but accumulating the running sum in fp32 keeps the speed benefit of low-precision inputs while avoiding the accuracy collapse that pure low-precision accumulation would cause over a long reduction — this is why ElementAccumulator is a separate template parameter from ElementA/ElementB in every CUTLASS GEMM declaration.

Takeaways

• The entire reason CUTLASS is structured the way it is comes from one fact: GPUs can do far more arithmetic per second than they can move bytes per second, so a fast GEMM kernel is fundamentally a data-reuse problem, not a raw-compute problem.
• Tiling is the single unifying idea — device → kernel → threadblock → warp → thread (2.x) or device → kernel → collective → tiled-MMA → atom (3.x) is the same strategy applied recursively: stage a smaller chunk of data one memory tier closer to the compute units, and reuse it as many times as possible before moving on.
• Software pipelining (double buffering) is what turns "tiling reduces memory traffic" into "the kernel never actually waits on memory" — by prefetching the next tile while computing on the current one.
• CUTLASS 3.x's redesign wasn't a rewrite for its own sake — thread-block clusters and TMA on Hopper genuinely had no clean place in the old threadblock/warp/thread hierarchy, and CuTe's layout algebra genuinely replaces an unbounded set of hand-written layout special cases with one small composable vocabulary.
• Mixed-precision accumulation (narrow inputs, wide accumulator) is the default recipe, not a special case — expect to see it in essentially every low-precision GEMM template you read, in CUTLASS or elsewhere.
• Start with the profiler and 2.x templates before touching CuTe and collectives — the 3.x machinery is genuinely more powerful, but it's built as a superset of concepts the 2.x hierarchy teaches first.

References & further reading

• NVIDIA — CUTLASS documentation (official, latest) — the primary source this article is grounded in throughout.
• CUTLASS C++ Quickstart — build instructions and the 2.x GEMM instantiation pattern referenced in §5 and §8.
• CUTLASS 3.x GEMM API — CollectiveMainloop, CollectiveEpilogue, and the dispatch-policy pattern covered in §6.2.
• Efficient GEMM in CUDA — the tiling and double-buffering rationale behind §2 and §4.1.
• CuTe Layouts — the shape/stride formalism behind §6.1.
• CuTe Layout Algebra — composition, complement, logical division/product; the operations that make CuTe more than a data structure.
• NVIDIA/cutlass on GitHub — source, examples directory, and issue tracker.
• Colfax Research — CUTLASS Tutorial: GEMM with Thread Block Clusters on Blackwell — a hands-on walkthrough of the cluster/collective concepts from §6.2 on the newest hardware.
• cvam.sight — 26 LLM Architecture Concepts, Explained — many of those concepts (MLA, MoE, attention) are ultimately served by GEMM kernels like the ones this article covers.