Every training run and every inference server on NVIDIA hardware eventually bottoms out in a matrix multiply, and for most of the industry that multiply comes from CUTLASS. Started in 2017 and cited by Vijay Thakkar and colleagues, the project at version 4.8.0 in this tree is NVIDIA’s collection of C++ template abstractions for high-performance GEMM and its relatives, BSD 3-Clause licensed, spanning FP64 through TF32, FP16, BF16, both FP8 flavors, the block-scaled NVFP4, MXFP4, MXFP6, and MXFP8 families, narrow integers, and even 1-bit types across Volta, Turing, Ampere, Ada, Hopper, Blackwell, and now the new Rubin SM107 generation. Version 4 added the CuTe DSL, Python-native kernel authoring with orders-of-magnitude faster compiles, and the README’s release notes name the company it keeps: every release is tested against FlashAttention, Quack, FlashInfer, cuDNN Frontend, PyTorch main, and TensorRT-LLM main. As always in this series, what follows is an educational tour of published source code.

The architecture is a hierarchy of named layers, each of which can be swapped without touching the others. At the bottom, CuTe defines layouts and tensors, the algebra of shapes, strides, and swizzles. Atoms bind that algebra to actual hardware instructions. Collectives assemble atoms into mainloops and epilogues, kernels schedule tiles across persistent blocks, and a thin device adapter exposes the whole assembly as a runnable GEMM. The Python DSL mirrors the same concepts so that a kernel written in CuTe C++ and one written with the cute.jit decorator lower to comparable PTX. Reading the overview from left to right walks one GEMM from a Layout object to a scheduled tensor-core instruction.

Overview of CUTLASS: CuTe layouts and tensors feed MMA and copy atoms bound to instruction wrappers, collective builders and warp-specialized kernels consume them through async stage machinery, the GemmUniversalAdapter exposes runnable GEMMs, the CuTe DSL mirrors it all in Python, and the Operator API plus examples and profiler close the loop.

Reading the overview from left to right:

• layout.hpp is the 1971-line heart of CuTe, defining hierarchically multidimensional layouts and the composition algebra of tiling and partitioning.
• tensor.hpp packages a layout with a data array, with the 1124-line tensor_impl.hpp doing the iterator bookkeeping.
• The atom folder holds instruction traits such as mma_traits_sm107.hpp at 565 lines and copy_traits_sm90.hpp for TMA loads.
• The arch folder wraps raw PTX across 35 files, led by mma_sm100_umma.hpp at 2041 lines for Blackwell’s tcgen05 instructions.
• The kernel layer stages warp-specialized GEMMs per generation, from sm90 cooperative and pingpong variants to sm100_gemm_tma_warpspecialized.hpp at 964 lines and the Rubin sm107 kernel.
• GemmUniversalAdapter at class line 123 of an 819-line file turns any kernel-plus-collectives assembly into a runnable device GEMM.
• Tile schedulers at 826 lines map output tiles to persistent blocks, with group and stream-K siblings for MoE workloads and tail smoothing.
• cute/core.py is the Python DSL heart at 6728 lines, with tile_to_shape at 5139 and local_partition at 5177.

Why You Need This

Writing a GEMM that hits hardware peak is the classic months-long systems project: shared-memory bank conflicts, register pressure, async copy staging, and instruction scheduling all interact, and every new GPU generation resets the board. CUTLASS exists so that nobody has to restart that project. Its README charts show the 3.x kernels operating at a large fraction of theoretical peak on Blackwell SM100 across every supported dtype, and because PyTorch, TensorRT-LLM, FlashAttention, and FlashInfer all build on these templates, the performance work NVIDIA ships here propagates through the entire ecosystem at the next dependency bump. If you have ever wondered where the speed of an fp8 or block-scaled fp4 matmul actually comes from, the answer is data movement planned by these headers: TMA bulk copies staged through multi-stage buffers, accumulators parked in tensor memory, and warpgroups kept busy with cooperative or pingpong overlap.

The second reason is portability of ideas. CuTe’s insight is that a kernel is mostly bookkeeping, and bookkeeping can be expressed as layout algebra: a partition, a tiling, a swizzle are values you can print, compose, and reason about without running anything. Once your mainloop is written against atoms and layouts rather than raw instructions, retargeting it is a matter of swapping the atom. That is exactly how this repository carries one conceptual design from Volta’s mma instructions through Hopper’s wgmma, Blackwell’s tcgen05 with 2-SM variants, and Rubin’s SM107 UMMA with FP8 at MMA_K=64 and FP4 at MMA_K=128, extended tensor memory of 576 columns, and 328 KB of shared memory. The compatibility table runs from a V100 on CUDA 11.4 to B300 and DGX Spark on CUDA 13.0, with the caveat that architecture-accelerated instructions require the sm_90a and sm_100a targets and that Rubin needs the R615 driver from CUDA 13.4 GA.

The third reason is the Python door. The CuTe DSL lets you write kernels with the cute.jit and cute.kernel decorators, traced into the same IR the C++ templates emit, with compile times in seconds instead of minutes and no glue code to call it from PyTorch. FlashAttention-4, which we toured earlier in this series, is written in exactly this DSL, and this release keeps polishing the path: the cute_ext extension compiler preview via CUTE_DSL_USE_EXTENSION_COMPILER=1, x86_64 Windows support, ahead-of-time host targets including QNX 8.0, JAX integration through cutlass.jax.cutlass_call, and an Operator API whose candidates can now be ranked by nvMatmulHeuristics on Blackwell. For framework authors the deal is simple: describe an operator in cutlass.operators, get a tuned kernel without writing C++.

Detail view: int_tuple and swizzle feed layout.hpp, tensor_impl glues iterators to layouts, atom traits wrap SM90 wgmma, SM100 tcgen05, and SM107 UMMA instructions, the adapter runs cooperative and warpspecialized kernels through builders and stage machinery, and the DSL and Operator API surfaces mirror the C++ stack.

The CuTe core is small enough to read in an afternoon and rewards it. int_tuple.hpp defines the hierarchical integer tuples that serve as shapes and strides, swizzle.hpp adds bit-swizzled layout wrappers that permute shared-memory addresses to dodge bank conflicts, and layout.hpp composes them into the 1971-line algebra at the center. tensor_impl.hpp, 1124 lines, glues an iterator to a layout and supplies local_tile and local_partition, the two calls every mainloop uses to slice a global problem into per-block tiles. algorithm/gemm.hpp then provides the 500-line generic MMA loop that consumes atom traits over partitioned tensors, with functional.hpp adding elementwise combinators for fused work. Print a layout with cute::print and the whole abstraction becomes concrete, which is why the tutorials lean on it so heavily.

The instruction floor shows three generations side by side. For Hopper, mma_traits_sm90.hpp, at 144 lines, declares wgmma atom traits, and copy_traits_sm90.hpp lets mainloops express TMA loads as ordinary CuTe copies. For Blackwell, mma_sm100_umma.hpp spans 2041 lines of tcgen05 UMMA wrappers with its 957-line descriptor builder in mma_sm100_desc.hpp encoding dtype, transpose, and scale-mode bitfields. For Rubin, mma_sm107_umma.hpp, 553 lines, pairs with mma_traits_sm107.hpp at 565 lines. Above them, the kernel layer files are named like a menu: sm90_gemm_tma_warpspecialized_cooperative.hpp runs 893 lines with a pingpong sibling that overlaps epilogues, sm100_gemm_tma_warpspecialized.hpp runs 964 lines parking accumulators in tensor memory, and sm107_gemm_tma_warpspecialized.hpp reuses the SM100 builder family against Rubin’s larger memories. The collective builders do the assembly, sm100_umma_builder.inl at 350 lines for mainloops and the 1841-line sm100 epilogue builder for TMEM epilogues with fusion callbacks, while the 1406-line Hopper stage-machinery header choreographs TMA barriers between producer and consumer warps.

The adapter and its predecessors complete the C++ story. GemmUniversalAdapter, defined at line 123 of gemm_universal_adapter.h, initializes workspace, picks occupancy-aware launch parameters, and runs any assembled kernel behind one interface, which is why a single example file can swap schedules without changing call sites. The 2.x world still lives beside it: default_gemm_universal.h and its siblings cover sparse, grouped, rank-2K, planar-complex, and split-K assemblies from the Volta and Ampere era, all resting on the warp-level mma.h at 90 lines. Tile scheduling matured alongside, from the sm90 scheduler with group and stream-K variants to the 826-line sm100 scheduler with cluster-aware and dynamic persistent modes that Blackwell MoE kernels lean on.

The Python stack mirrors all of it. cute/core.py, at 6728 lines, re-exports the layout algebra with tile_to_shape at 5139 and local_partition at 5177, tensor.py adds a 3309-line Tensor whose copy and partition calls lower to the same paths as C++, and cutlass_dsl/cutlass.py hosts the decorators users actually write, jit at 1878 and kernel at 1957, traced by the pyir engine under base_dsl. The working example is in the repo: examples/python/CuTeDSL/dsl_tutorials/call_from_jit.py imports cutlass and cutlass.cute, wraps a GEMM in an @cute.jit function at line 163, and converts PyTorch buffers through the JitArgument protocol. The Operator API sits a level higher in operators/cutlass/operators with base.py describing intent and manifest.py enumerating candidates, and operators/cutlass/kernels exposes standalone persistent GEMMs including a Rubin flavor, callable without any operator ceremony. The 100 example directories, from 88_hopper_fmha through 92_blackwell_moe_gemm and 93_blackwell_low_latency_gqa to the cute/rubin block-scaled samples, double as the test corpus for the profiler under tools.

From Install to a Tuned GEMM

Build from source with the architecture-specific flags, then sweep configurations with the bundled profiler:

git clone https://github.com/NVIDIA/cutlass.git
cd cutlass && mkdir build && cd build
cmake .. -DCUTLASS_NVCC_ARCHS="90a;100a" -DCMAKE_BUILD_TYPE=Release
make cutlass_profiler -j
./tools/profiler/cutlass_profiler --operation=GEMM --m=4096 --n=4096 --k=4096

For the Python DSL, install the wheel and trace a kernel with the same decorators the repo’s tutorials use:

import cutlass
import cutlass.cute as cute

@cute.jit
def tensor_op_gemm_wrapper(a_tma, b_tma, mC, mD, ...):
    # compose layouts, partition tiles, launch the kernel
    ...

# examples/python/CuTeDSL/dsl_tutorials/call_from_jit.py shows the full flow:
# PyTorch tensors enter through make_ptr, exit as a JIT-compiled GEMM result

pip install nvidia-cutlass-dsl gives you the DSL package that FlashAttention-4 and the Operator API build on, and the notebooks under examples/python/CuTeDSL/cute/notebooks walk from first layout to block-scaled GEMM.

Honest limits: this is template metaprogramming, and C++ compile times for instantiated kernels remain a real tax, which is precisely why the DSL exists. Windows builds for the 3.x C++ stack are currently broken per the README, the sm_90a and sm_100a targets trade forward compatibility for speed, and GeForce RTX 50 owners on SM120 cannot run SM100 datacenter kernels. The Operator API is young and its Rubin support is explicitly a preview pending performance tuning. None of that changes the core fact: if a matrix multiply matters to your workload, the most battle-tested implementations on the planet are open source in this single tree.

Enjoyed this post? Never miss out on future posts by following us