Overview

CUTLASS 4.8.1

CUTLASS 4.8.1 - Oct. 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 released 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 2026.

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

What's New in CUTLASS 4.8

CuTe DSL

• New features

• Initial Rubin support to accelerate dense GEMMs. The following features are available:

  • CuTe DSL and CuTe extensions
  • Support for higher-throughput FP8 (MMA_K=64) and FP4 (MMA_K=128) Tensor Core MMA instructions
  • B collector reuse
  • Extended TMEM size from 512 COL to 576 COL
  • Larger shared memory allocations (328KB)
  • Enhanced mixed precision throughput (FP8/FP4)
  • Softmax acceleration related features
  • Primitives
  • Support for higher-throughput FP8 (MMA_K=64) and FP4 (MMA_K=128) Tensor Core MMA instructions
  • B collector reuse
  • Extended TMEM size from 512 COL to 576 COL
  • Larger shared memory allocations (328KB)
  • Enhanced mixed precision throughput (FP8/FP4)
  • Softmax acceleration related feature
  • 2:4 sparsity support for FP4

• CuTe DSL extensions has several new features:

  • CTA-V maps are now inferred automatically for cute_ext TMA load, store, multicast, and reduce-store operations. Explicit CTA-V maps remain supported as overrides.
  • Added asynchronous atomic TMA reduce-store and sparse MMA operations.
  • Added reusable cute_ext GEMM mainloop and TMA epilogue helpers.
  • Added opt-in TMEM accumulator-buffer planning, including overlapping ping-pong storage for capacity-constrained kernels.
  • Improved device-side TMA descriptor updates and grouped GEMM performance through SMEM-staged updates, workspace reuse, and reduced prologue and synchronization overhead.
• This release includes an opt-in preview of the CuTe DSL extensions (cute_ext) compiler pipeline for ordinary Cute DSL kernels. This pipeline lets user mix cute_ext APIs directly into @cute.jit and @cute.kernel code and is required for kernels that mix the two API surfaces. You may test this feature with the following: CUTE_DSL_USE_EXTENSION_COMPILER=1 python your_program.py The pipeline is expected to preserve program behavior and performance, but generated PTX/SASS may differ. Note that this pipeline will become the default in the future, no earlier than 4.10.
• Added examples for better control over Primitives' compiler warnings/errors introduced in 4.7.0. See the CuTeDSL/experimental/compiler_diagnostic/ directory.
• IKET Profiler Tool
  • Rubin kernels (sm107) can now be profiled.
  • It is now possible to only dump timing data for a specific cluster to reduce profiling overhead. Previously all clusters were profiled.
  • Task Scheduling can instrument the schedule with IKET ranges when constructing TaskManager objects (iket_enable_profiling=True). Task execution will generate an IKET range and individual pipeline stages in a schedule may generate separate ranges (iket_profiling_stages).
• A number of new examples were added in this release:
  • Rubin (CuTe):
  • Dense GEMM for legacy data type with B collector reuse as applicable
  • Grouped GEMM with B collector reuse
  • Dense blockscaled GEMM with FP4/FP6/FP8 mixed precision and UE5M3 / block-32 scale-factor support
  • Grouped blockscaled GEMM with B collector reuse as applicable
  • Blockwise GEMM
  • Rubin (CuTe extension):
  • Support for higher-throughput FP8 (MMA_K=64) and FP4 (MMA_K=128) blockscaled GEMM with UE5M3 scale-factor
  • Grouped GEMM with B collector reuse
  • Blackwell (CuTe extension):
  • Dense GEMMs
    • Back-to-back GEMM
    • Blockscaled GEMM
    • Persistent GEMM with alpha/beta scaling
    • CLC scheduler/dynamic persistent GEMM
    • GLU GEMM
    • Mixed input GEMM
    • Planar complex GEMM
    • Input transform GEMM
    • GeForce pingpong dense GEMM
    • Blackwell Ultra blockscaled GEMM
  • Dense Convolutions
    • Implicit-Gemm Fprop Conv
    • Blocksclaed Implicit-Gemm Fprop Conv
    • GeForce Implicit-Gemm Fprop Conv
    • GeForce Blockscaled Implicit-Gemm Fporp Conv
  • Attention
    • GQA Decode
  • Grouped GEMM
    • Unscaled and blockscaled grouped GEMM
  • Top-K
  • Ampere (CuTe extension):
  • SIMT GEMM
• CuTe DSL now supports x86_64 Windows
• CuTe DSL AoT now supports new host target: QNX8.0
• Notebooks are restructured under examples/python/CuTeDSL/cute/notebooks and new notebooks for primitives will be added under examples/python/CuTeDSL/notebooks

• Numpy is now not a default dependency

• Bug fixes and improvements:

• nvidia-cuda-nvdisasm is now an optional dependency of nvidia-cutlass-dsl via the optional [sass] extra. SASS dumping (CUTE_DSL_KEEP=sass / KeepSASS) now resolves nvdisasm from the bundled wheel (recommended since its version matches the DSL toolchain) or from a local CUDA Toolkit (CUDA_HOME/CUDA_PATH). A locally-provided nvdisasm must come from a CUDA Toolkit at least as new as the toolchain that produced the CUBIN. Installations that never dump SASS are unaffected.
• Reduced the protobuf version requirement of IKET profiler from 6.30 to 4.21. This should make protobuf an easier requirement to satisfy in preparation for transioning IKET to an optional extra.
• Improved JAX PyTree input/output aliasing for cutlass.jax.cutlass_call
• Fixed a regression from 4.6.0 where cute.autovec_copy emitted per-element instead of vectorized instructions for tensors with a dynamic stride (!3463)
• Fixed TVM-FFI env stream detection for GPU tensors in tuple (!3444)
• Fixed GPU link-libraries compile-option order so it is stable across processes (!3564)
• Fixed preprocessor IndexError on staged bool() with no arguments (!3506)
• Rejected cute.compile on @cute.kernel with a user error instead of an ICE (!3429)
• Fixed CuTe DSL crashing the Python interpreter when used in a REPL (!3413)
• Fixed a cuDNN Frontend FROST SDPA backward compilation failure issue (!3594)
• Fixed SIGABRTs in TVM-FFI launch for cuDNN Frontend SM100 ragged SDPA kernel (!3595)
• Fixed an issue with cute.extern where functions with zero arguments would generate invalid type signatures. Compiling such functions a second time would fail.
• Clear error/warning are now emitted for CopyUniversalOp with memory attributes specified (!3711).
• CUTLASS Primitives' tensormap_replace now keeps device-side TMA descriptor updates safe when a descriptor is shrunk below its original extent (e.g. grouped GEMM) and fixes a possible illegal memory access on Blackwell/Rubin. This behavior can be opted out by specifying keep_oob_fetch_window=True. See https://docs.nvidia.com/datacenter/tesla/tesla-release-notes-595-91-07/index.html#known-issues for more details on the known cuTensorMapEncodeTiled bug.

• Fixed unintended sys.path modification that could break imports in some environments where Torch was not installed e.g. JAX.

• Known issues

• Due to a bug in CTK, the use of some packed instructions result in invalid code generation. Given the nature of the instructions and how they are explicitly used, we have chosen to not silently rewrite these, but rather are flagging until the up