CUTLASS + CuTe — the industrial kernel language

When to invest: After you have finished Boehm’s SGEMM ladder and Triton tutorials 01–06. Not before. CUTLASS/CuTe is what you reach for when Triton has hit its ceiling and you need Hopper/Blackwell primitives with full control.

Reality check: For inference-engineering job scope, you need to read CUTLASS/CuTe fluently and modify it, not necessarily author from scratch. FA3 and FA4 lean on CuTe (FA4 is literally written in CuTe-DSL). Being able to open include/cutlass/gemm/kernel/sm90_gemm_tma_warpspecialized_pingpong.hpp and understand what it does is the minimum bar.

The layered architecture

CUTLASS is layered from high-level GEMM API down to individual PTX instructions:

cutlass::gemm::device::Gemm            <— top-level, user-facing
    │
cutlass::gemm::kernel::Gemm            <— kernel-level (grid/block/warp iteration)
    │
cutlass::gemm::threadblock::Mma        <— block-tile MMA loop
    │
cutlass::gemm::warp::Mma               <— warp-tile MMA
    │
cutlass::arch::mma_sync / wgmma        <— PTX-level tensor-core instructions

On Hopper (SM90), the kernel-level classes use CuTe everywhere — layouts, tensors, tiles, tiled MMA — to express the algorithm.

CuTe — the math primitive

CuTe is the layout algebra library that ships with CUTLASS 3.x. It replaces the older “iterators” with a single unified abstraction: Layout = Shape × Stride.

The one insight: a Layout is a function from a coordinate to an offset. Composition of layouts (composition, logical_divide, zipped_divide, tiled_product) lets you build complex indexing patterns algebraically instead of writing nested loops with hand-computed strides.

Once you internalize this, HBM/SMEM/register tiling becomes function composition, not stride bookkeeping.

Minimum vocabulary

  • Shape: tuple describing tensor dimensions, potentially hierarchical (e.g., ((4,32),(2,16)) for a warp-tiled layout).

  • Stride: matching-shape tuple of strides (in element units).

  • Layout: a (Shape, Stride) pair, callable as L(coord) → offset.

  • Tensor: a Layout + a pointer.

  • tiled_divide / logical_divide: reshape a layout into (outer_tile, inner_tile).

  • local_partition / local_tile: partition a tensor across threads/warps.

  • TiledMMA: describes how a warp-group’s MMAs tile a bigger MMA operation.

Example (mental): partition a global A tile among 128 threads so each thread owns 4 rows × 16 cols of the tile. In old CUTLASS you’d write nested loops with stride math. In CuTe you compose the block layout with a thread layout — one line.

Learning path (in order)

Phase A — read (2–3 evenings)

  1. CuTe Quickstart (NVIDIA official): https://docs.nvidia.com/cutlass/latest/media/docs/cpp/cute/00_quickstart.html

  2. CuTe Layout (NVIDIA official): https://docs.nvidia.com/cutlass/latest/media/docs/cpp/cute/01_layout.html

  3. CuTe Layout Algebra (NVIDIA official, rigorous): https://docs.nvidia.com/cutlass/latest/media/docs/cpp/cute/02_layout_algebra.html

These three are dense but load-bearing. Read them once, then move on.

Phase B — outside perspectives

  1. Lei Mao — CuTe Layout Algebra: https://leimao.github.io/article/CuTe-Layout-Algebra/ — the best mathematical treatment on the open web. Read after NVIDIA docs; it will click.

  2. GPU MODE Lecture 15 — CUTLASS deep dive (Andreas Köpf). Notes: https://christianjmills.com/posts/cuda-mode-notes/lecture-015 — excellent conceptual overview.

  3. Colfax Research — https://research.colfax-intl.com/category/papers/deep-learning — technical notes on Hopper/Blackwell GEMM. Their FA3 note and their WGMMA introduction are must-reads.

Phase C — hands on

  1. CutlassAcademy — curated exercises: https://github.com/MekkCyber/CutlassAcademy — progressive problems with solutions.

  2. mlai.blog CuTe basics + attention — http://mlai.blog/2025-05-10-cute-basics — implements FA-style attention using CUTLASS/CuTe. Read alongside FA3 paper.

  3. CUTLASS example kernels — in the CUTLASS repo, examples/48_hopper_warp_specialized_gemm and examples/62_hopper_sparse_gemm are compact reference kernels. Compile them, run under Nsight, tweak block sizes.

When to invest — the decision tree

Do I need Blackwell-specific features (TMEM, tcgen05, 2-CTA MMA)?
  └ YES  → CuTe (via CUTLASS C++ or CuTe-DSL Python)
  └ NO   → Is Triton fast enough (within 10–20% of hand-optimum)?
              └ YES → Triton
              └ NO  → CUTLASS/CuTe (Hopper)

Signals CUTLASS is warranted:

  • You need explicit control of TMA descriptors + wgmma pipelining.

  • You’re writing a fused GEMM+epilogue with unusual epilogue (e.g., custom quant, split-K reductions).

  • Triton’s autoinference is choosing a bad layout and you can’t nudge it.

  • You’re targeting frontier throughput (top 5% of possible).

Signals Triton is fine:

  • You’re within 10–20% of cuBLAS/cuDNN and that’s good enough for your service.

  • You need to iterate rapidly.

  • You’re comfortable with torch.compile handling most of the fusion story.

CuTe-DSL — the Python front-end (2025+)

NVIDIA has released CuTe-DSL, a Python DSL that maps to the same CuTe primitives. Two big deals:

  1. FA4 is written in Python via CuTe-DSL (arXiv <phone_number_or_numberic_id_or_random_id_46>). Frontier kernels in Python is now viable.

  2. Integrates with JAX via cutlass.jax.cutlass_call, allowing CuTe kernels to be called from JAX code paths.

Positioning: CuTe-DSL is the “Triton for CuTe.” You keep the Python velocity, but you get explicit CuTe semantics + Blackwell primitives. When to reach for it: writing Blackwell-native kernels where Triton’s Blackwell path (via 3.6/3.7) doesn’t cover your use case.

Learning:

What to actually build in this file’s scope

Don’t try to build a full CUTLASS-competitive kernel in Phase 2. Instead:

  1. Read CuTe quickstart + layout + layout-algebra (Phase A above).

  2. Compile and run CUTLASS example 48_hopper_warp_specialized_gemm on H100 (or an Ampere equivalent on Ampere hardware).

  3. Modify one thing — change the block tile size or the pingpong stages — and observe the perf delta.

  4. Read one Colfax note end-to-end.

  5. Read the FA3 paper (arXiv <phone_number_or_numberic_id_or_random_id_47>) with the intention of understanding its CuTe primitives, not just its algorithm.

That’s Phase-2-appropriate depth. Real CUTLASS authoring is a Phase 4+ activity.

Common pitfalls

  • Trying to memorize the API surface. CUTLASS has thousands of template parameters. Understand the layered abstraction and grep for the specific class you need.

  • Reading arch/ before collective/. The arch/ layer is thin PTX wrappers; the interesting design is in collective/ and kernel/.

  • Fighting the templates. If your compile error is 200 lines long, you’re using it wrong. Start from an example that works and modify incrementally.

  • Ignoring CuTe printers. cute::print(layout); is your best debugging tool. Use it.

References