Overview

CUTLASS 4.5.2

CUTLASS 4.5.2 - May 2026

CUTLASS is a collection of abstractions for implementing high-performance matrix-matrix multiplication (GEMM) and related computations at all levels and scales within CUDA. It incorporates strategies for hierarchical decomposition and data movement. CUTLASS decomposes these "moving parts" into reusable, modular software components and abstractions.

Primitives for different levels of a conceptual parallelization hierarchy can be specialized and tuned via custom tiling sizes, data types, and other algorithmic policy. The resulting flexibility simplifies their use as building blocks within custom kernels and applications.

CUTLASS has been providing CUDA C++ template abstractions for high-performance linear algebra since 2017 and these abstractions provide extensive support for a wide range of computations including mixed-precision computations, specialized data-movement (async copy) and multiply-accumulate abstractions for FP64, FP32, TF32, FP16, BF16, FP32 emulation via tensor core instruction, 8b floating point types (e5m2 and e4m3), block scaled data types (NVIDIA NVFP4 and OCP standard MXFP4, MXFP6, MXFP8), narrow integer types (4 and 8b signed and unsigned integers), and binary 1b data types (where architectures allow for the native support of such data types) across NVIDIA's Volta, Turing, Ampere, Ada, Hopper, and Blackwell architectures.

To this rich ecosystem of C++ based kernel programming abstractions, CUTLASS 4 adds CUTLASS DSLs. These are Python native interfaces for writing high-performance CUDA kernels based on core CUTLASS and CuTe concepts without any performance compromises. This allows for a much smoother learning curve, orders of magnitude faster compile times, native integration with DL frameworks without writing glue code, and much more intuitive metaprogramming that does not require deep C++ expertise.

Overall we envision CUTLASS DSLs as a family of domain-specific languages (DSLs). With the release of 4.0, we are releasing the first of these in CuTe DSL. This is a low level programming model that is fully consistent with CuTe C++ abstractions — exposing core concepts such as layouts, tensors, hardware atoms, and full control over the hardware thread and data hierarchy.

CuTe DSL demonstrates optimal matrix multiply and other linear algebra operations targeting the programmable, high-throughput Tensor Cores implemented by NVIDIA's Ampere, Hopper, and Blackwell architectures.

We believe it will become an indispensable tool for students, researchers, and performance engineers alike — flattening the learning curve of GPU programming, rapidly prototyping kernel designs, and bringing optimized solutions into production.

CuTe DSL is currently in public beta and will graduate out of beta by end of summer 2025.

To get started quickly - please refer : - CUTLASS C++ Quick Start Guide. - CuTe DSL Quick Start Guide.

What's New in CUTLASS 4.5

CuTe DSL

• New features
• New Block API block_copy() to simplify TMA and S2T copy. Users can ignore detail about multicast and 2CTA partition for TMA by block_copy() and need not to invoke tma_partition(). And users can remove bulk of S2T initialization to simplify S2T copy.
• MXF8F6F4 mixed precision support
  • BlockScaled MMA now supports MXF8MXF4 or MXF8MXF6
• Block Scaled MMA for SM120 now works on Spark
• EFC broadcast semantics support
  • EFC epilogue functions can now broadcast and remap tensor modes via C.remap_modes[:, 0, 1] subscript syntax (where : marks a broadcast dimension and integers select source mode indices). Covers scalar broadcast, row/column broadcast, and arbitrary mode permutations (e.g. transpose). The PyTorch reference evaluator mirrors the same transformations.
• Initial linter support: Improved type hints on CuTe DSL APIs to support static type checkers like MyPy
• dataclasses.dataclass is now supported for JIT compilaton and cute.compile for both plain and tvm-ffi path
• cute.copy now supports user specified loop unrolling

• Python 3.14t is now supported with GIL enabled

• Bug fixing and improvements

• Improved source code correlation for profiling/debugging
• Fixed an aarch64 segfault issue with tvm-ffi
• Re-organization for CuTe DSL examples/tutorials for better discoverability
• Fixed following issues: https://github.com/NVIDIA/cutlass/issues/3219 https://github.com/NVIDIA/cutlass/issues/3218 https://github.com/NVIDIA/cutlass/issues/3212 https://github.com/NVIDIA/cutlass/issues/3210 https://github.com/NVIDIA/cutlass/issues/3208 https://github.com/NVIDIA/cutlass/issues/3201 https://github.com/NVIDIA/cutlass/issues/3227 https://github.com/NVIDIA/cutlass/issues/3240 https://github.com/NVIDIA/cutlass/issues/3241
• Fixed Jax int64 stride divisibility issue

• Fixed issues for SM120 blockscaled MMAs

  • added missing MXFP8MMAOP and MXF8F6F4MMAOP for sm120.

• More examples of authorizing peak-performance kernels

• MOE examles

  • A new style of grouped-gemm that aligns to torch's grouped_mm and scaled_groued_mm interface.
  • Expert-wise tensormap descriptor setup by a cheap helper kernel (~2us) to avoid long latency in tile switching, kernel structure is much more closer to a normal GEMM.
  • Compared to torch_210_cu13, very few problem has worse perf in B200.
    • mxfp8_2dx3d: avg 1.29 speedup;
    • mxfp8_2dx2d: avg 1.41 speedup;
      • nvfp4_2dx3d: avg 1.11 speedup;
    • nvfp4_2dx2d: avg 1.12 speedup (worst case 0.98)
    • bf16_2dx3d: avg 1.15 speedup (worst case 0.98)
    • bf16_2dx2d: avg 1.17 speedup (worst case 0.96)
    • Note: The perf is measured from torch profiler, this impl includes the helper kernel + main kernel, while torch's includes its setup kernel and cutlass_cpp main kernel.

• API changes

• ab_dtype is deprecated in make_trivial_tiled_mma and make_blockscaled_trivial_tiled_mma from blackwell_helpers.py. Please specify a_dtype and b_dtype separately instead.

CUTLASS C++

• Add 2SM MMA instruction support to mixed TMA+CpAsync SM100 vanilla GEMM kernels.
• Mixed TMA+CpAsync can now accept static, but non trivial cluster shapes.
• Uses TMA multicast for A tile when using non-trivial cluster size along N mode.
• Uses an additional barrier (mma_trampoline_barrier) to track cp.async arrivals in both CTAs.
• Changes included in example 92.
• Add support for 128x32xK and 128x64xK tile sizes for SM120 blockscaled MMA collective builders, yielding up to 30% performance improvement on Blackwell SM121 related kernels.
• Add static load to tensor memory support, included in example 77.
• Use 64-bit adds for SM100 MMA descriptor offsets and reduce move instructions for improved code generation.
• Add example 95 to support green context SM partition
• Enables launching GEMM on stream with partial SM allocation.
• Add Snake activation functor for EVT.
• Fix SM100 F8F6F4 SS MMA (1SM and 2SM) traits to use typed op templates.
• Add UE8M0 (uniform exponent distribution) initialization support in tensor fill utilities.
• Add cvt.rn.bf16x2.e4m3x2 conversion instruction support to numeric_conversion.h.
• Update example 93 with paged KV cache support for Blackwell low-latency GQA.
• Fix some kernel issues:
• Fix l2_capacity=0 handling in Blackwell SM100/SM120 kernel templates
• Fix CUTLASS clang build issues
• Remove PipelineStorage shadowing in SM100 complex epilogue
• Fix build issue in SM90 epilogue fusion visitor TMA warpspecialized
• Fix missing convert fucntion in EVT for fp4 kernels
• Fix some profiler issues:
• Add missing reference kernels for blockwise GEMM profiler.
• Avoid instantiate 2sm tma kernels where ctaN is none power of 64 when ctaN > 128 in profiler.

Note: CUTLASS 4.x builds are known to be down on Windows platforms for all CUDA toolkits. CUTLASS team is working on a fix.

See the CHANGELOG for details of all past releases and updates.

Performance

CUTLASS primitives are very efficient. When used to construct device-wide GEMM kernels, they exhibit nearly optimal utilization of peak theoretical throughput. The figure below shows CUTLASS 3.8's performance as a % of theoretical peak utilization on various input and output data types when run on NVIDIA Blackwell SM100 architecture GPU.

The two figures below show the continual CUTLASS performance improvements on an NVIDIA H100 (NVIDIA Hopper architecture) since CUTLASS 3.1. CUTLASS 3.5.1 was compiled with the CUDA 12.5u1 Toolkit. Tensor Core operations are implemented using CUDA's mma and wgmma instructions.

CuTe

CUTLASS 3.0 introduced a new core library, CuTe, to describe and manipulate tensors of threads and data. CuTe is a collection of C++ CUDA template abstractions for defining and operating on hierarchically multidimensional layouts of threads and data. CuTe provides Layout and Tensor objects that compactly package the type, shape, memory space, and layout of data, while performing the complicated indexing for the user. This lets programmers focus on the logical descriptions of their algorithms while CuTe does the mechanical bookkeeping for them. With these tools, we can quickly design, implement, and modify all dense linear algebra operations.

The core abstractions of CuTe are hierarchically multidimensional layouts which can be composed with data arrays to represent tensors. The representation of layouts is powerful enough to represent nearly everything we need to implement efficient dense linear algebra. Layouts can also be combined and manipulated via functional composition, on which we build a large set of common operations such as tiling and partitioning.

CUTLASS 3.0 and beyond adopts CuTe throughout the GEMM hierarchy in its templates. This greatly simplifies the design and improves code composability and readability. More documentation specific to CuTe can be found in its dedicated documentation directory.

Compatibility

Minimum requirements:

• Architecture: Volta (compute capability 7.0)
• Compiler: Must support at least C++17
• CUDA Toolkit version: 11.4

CUTLASS requires a C++17 host compiler and performs best when built with the CUDA 12.8 Toolkit. It is also compatible with CUDA 11.4, CUDA 11.5, CUDA 11.6, CUDA 11.7, CUDA 11.8, and all other CUDA 12.x versions.

Operating Systems

We have tested the following environments.

Note: GCC 8.5.0 has known regressions regarding fold expressions and overloaded operators. Using GCC 7.5.0 or (preferred) GCC >= 9 is recommended.

Note: CUTLASS 3.x builds are known to be down on Windows platforms for all CUDA toolkits. CUTLASS team is working on a fix.

Hardware

CUTLASS runs successfully on the following NVIDIA GPUs, and it is expected to be efficient on Volta, Turing, Ampere, Ada, and Hopper architecture based NVIDIA GPUs.

Target Architecture

In general, PTX code generated for one target architecture can be run on future architectures (i.e., it is forward compatible). However, CUDA 12.0 introduced the concept of "architecture-accelerated features" whose PTX does not have forward compatibility guarantees. Several Hopper and Blackwell PTX instructions fall under this category of architecture-accelerated features, and thus require a sm_90a or sm100a target architecture (note the "a" appended). For more details on this and other architecture-accel