Second post in the "Foundations" series (the first covered GPU architecture generations, Volta to Blackwell). This one is about a question that doesn't have an obvious answer the first time you hit it: NVIDIA already ships cuBLAS, a fast, well-tested matrix multiply library — so what is CUTLASS for, and why would anyone reach for it instead?

The problem cuBLAS doesn't solve

cublasSgemm() is fast. For a plain matrix multiply, it's close to the best you'll get out of the hardware, and you don't have to think about tiling, shared memory, or Tensor Core instructions to use it. That's exactly the problem: it's a black box. You call it, it computes `C = αAB

  • βC`, and that's the entire interface.

Real workloads usually want more than that:

  • Fusion — bias add, activation function, or a scale-and-clamp applied to the GEMM output, without writing the result to global memory and launching a second kernel to read it back.
  • Non-standard data types — mixed precision (FP16 in, FP32 accumulate), INT8/INT4 for quantized inference, or newer formats like FP8, none of which cuBLAS exposes with full flexibility.
  • Problem-specific specialization — grouped or batched GEMMs with irregular shapes, or a tile size tuned for one specific, known problem size instead of a generic one.

None of this is available by calling into cuBLAS. To get it, you'd normally have to write the kernel yourself — and a hand-written GEMM that actually reaches Tensor Core peak throughput is a genuinely hard piece of software: multi-level tiling, double-buffered shared memory, avoiding bank conflicts, and using the right matrix-multiply-accumulate instruction for the target architecture (wmma, mma.sync, or wgmma, depending on generation). Most of that work has nothing to do with your actual problem — it's the same scaffolding every fast GEMM kernel needs.

CUTLASS (CUDA Templates for Linear Algebra Subroutines) is that scaffolding, factored out into reusable, composable C++ templates — open source, maintained by NVIDIA, and used internally to build parts of cuBLAS itself.

The idea: match the abstraction to the hardware

CUTLASS's core design decision is to decompose a GEMM into layers that correspond directly to the GPU's own execution and memory hierarchy, rather than treating "matrix multiply" as one opaque operation:

Three-column diagram showing how CUTLASS's abstraction layers (device, threadblock, warp, thread) map onto the GPU's execution hierarchy and memory hierarchy

Each row is doing the same conceptual job — computing a tile of the output — just at a different scale, backed by a different part of the memory system:

  • At the device level, the whole problem is split into tiles, one per thread block.
  • Each thread block cooperatively loads its slice of A and B into shared memory, since re-reading from global memory (HBM) for every reuse would be far too slow.
  • Each warp takes a sub-tile of that shared-memory data and drives the Tensor Core, using matrix fragments held in registers.
  • At the thread level, individual lanes participate in the actual MMA (multiply-accumulate) instruction issued by the warp.

This is exactly the tiling discipline any hand-written high-performance GEMM has to implement — CUTLASS just gives each layer a name, a C++ type, and a well-defined interface, so you can swap out one layer (say, the data type or the epilogue) without re-deriving the other three.

What using it actually looks like

At the simplest level — a plain FP32 GEMM with no fusion — declaring and running a CUTLASS kernel looks like this:

#include "cutlass/gemm/device/gemm.h"
 
using ColumnMajor = cutlass::layout::ColumnMajor;
 
using CutlassGemm = cutlass::gemm::device::Gemm<
    float,        // data type of A
    ColumnMajor,  // layout of A
    float,        // data type of B
    ColumnMajor,  // layout of B
    float,        // data type of C
    ColumnMajor>; // layout of C
 
CutlassGemm gemm_operator;
 
CutlassGemm::Arguments args(
    {M, N, K},       // problem size
    {A, lda},        // source matrix A
    {B, ldb},        // source matrix B
    {C, ldc},        // source matrix C
    {C, ldc},         // destination matrix D
    {alpha, beta});   // epilogue scalars
 
cutlass::Status status = gemm_operator(args);

CutlassGemm here is a C++ type, not a function — the compiler generates a specialized kernel for exactly this combination of data types and layouts at compile time. That's the mechanism behind the customization CUTLASS is built for: change float to cutlass::half_t and int8_t, add an epilogue that fuses a ReLU, and the compiler produces a different, still-specialized kernel — no runtime dispatch, no black box.

This particular shape of example (plain data types, no fusion) is intentionally the simplest possible one. It's also, not coincidentally, close to what cuBLAS already gives you — the difference only shows up once you start changing the template parameters to things cuBLAS can't express.

CUTLASS 2.x vs. 3.x: the CuTe rewrite

CUTLASS has existed since 2017, and the API above is from its original (2.x) design. Starting with CUTLASS 3.0, released alongside Hopper, NVIDIA introduced CuTe — a new core library for describing hierarchically multi-dimensional layouts of threads and data — and rebuilt CUTLASS's GEMM kernels on top of it, adding direct support for Hopper's wgmma instruction and the Tensor Memory Accelerator (TMA) for asynchronous bulk data movement.

The hierarchy in the diagram above didn't change between 2.x and 3.x — device, threadblock, warp, and thread are still the levels a GEMM gets decomposed into. What changed is how each level's tile is described: CuTe replaces a set of fairly rigid, generation-specific template classes with a small, composable algebra of shapes and strides that works uniformly across architectures. That's specific enough to be its own post — next in this series.

Why this is worth knowing before writing any kernels

CUTLASS is a big, actively-developed codebase, and it's easy to treat it as something to copy-paste from rather than understand. The one idea worth taking away from this post: every abstraction layer in CUTLASS exists because it corresponds to a real, physical layer of the GPU. When something about CUTLASS's API feels arbitrary, it usually isn't — it's tracking a hardware boundary (a memory space, a synchronization scope, an instruction granularity) that any fast kernel has to respect, CUTLASS or not.

References

Hierarchy diagram above is original artwork made for this post.