CUTLASS 4.8.0

[Junkai-Wu]

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.

---

This discussion was created from the release CUTLASS 4.8.0.