CUTLASS Files

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)

This release has been tested against the following packages:

• FlashAttention: main (8d3a3b8)
• Quack: main (35266c3)
• FlashInfer: main (5d0c89e)
• cuDNN-Frontend: deveop (e0317d1)
• Pytorch: main (7d5f021)
• TensorRT-LLM: main (c295dd9)

CUTLASS Operator API

• Dense and blockscaled GEMMs in Operator API have preliminary Rubin support. These are provided as a preview and may need additional performance tuning.

• Updated GEMMs include:

  • Dense GEMMs: FP8xFP8
  • Blockscaled GEMM: {MXFP8}x{MXFP4, MXFP8} and {MXFP4, NVFP4}x{MXFP4, NVFP4} (including support for the new UE5M3 scale factor dtype for NVFP4).

• These kernels utilize the below new features in Rubin:

  • Higher SMEM (328KB) and TMEM capacity (288KB)
  • B-buffer reuse
  • Enhanced mixed precision throughput

• Operators can now be ranked by their estimated performance when nvMatmulHeuristics is available. See tutorial here. NOTE: This currently only supports Blackwell kernels as nvMatmulHeuristics does not yet support Rubin.

• Standalone kernel implementations are now exposed through cutlass.kernels. These kernels can be used directly, in addition to being discoverable and usable via the Operator interface in cutlass.operators.

• Custom Epilogue fusions now support per-row or per-column reductions.

• IndexPtrGroupedGemmArguments is now used to represent Grouped GEMM with contiguous-offset/index-pointers. Existing GroupedGemmArguments is deprecated and will be removed in a future release.

C++

• Added initial Rubin support (SM107) with CuTe C++ building blocks:
• Rubin Tensor Core MMA instructions and corresponding CuTe MMA traits.
• CuTe examples that demonstrate the use of Rubin SM107 Tensor Core instructions:
• Dense FP8 GEMM.
• Block-scaled FP8 GEMM.
• Mixed-precision block-scaled FP8/FP4 GEMM.
• Block-scaled FP4 GEMM.
• Adjusted shared-memory and tensor-memory capacity handling for Rubin SM107:
• Set the SM107 shared-memory capacity to 327 KiB and added launch support for oversized shared-memory configurations.
• Set the SM107 TMEM capacity to 576 columns per SM, updated the CuTe 1SM and 2SM TMEM allocators for Rubin's exclusive allocation path.
• Enabled the existing SM100-compatible GEMM and convolution for the new SM107 sm_107a and sm_107f targets:
• Set of unit tests for Rubin SM107 SIMT GEMM, dense FP8 GEMM, block-scaled FP8 GEMM, block-scaled FP4 GEMM, and mixed-precision, complex, and 9xBF16 GEMM
• Various improvements and fixes from the community and CUTLASS team. Thanks to everyone who submitted PRs!
• Optimal code generation with CUDA toolkit versions 13.4.

NOTE: Executing Rubin kernels (SM107) requires the R615 driver which will be released with CUDA Toolkit 13.4 GA. R610 from CUDA Toolkit 13.4 Developer Preview is not sufficient.