CUTLASS 4.8.0 dev

[Junkai-Wu]

CuTe DSL

• New features

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

    • CuTe DSL
      • Rubin new FP8 and FP4 Tensor Core support
      • B collector reuse
      • Extended TMEM size from 512 COL to 576 COL
      • Larger shared memory allocations (328KB)
      • Enhanced mixed precision throughput (FP8/FP4)
    • Primitives
      • Rubin new FP8 and FP4 Tensor Core support
      • Extended TMEM size from 512 COL to 576 COL

    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.

  • 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, but generated PTX/SASS may differ. The pipeline is planned to become the default in a future release.

  • 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):
      • FP4 blockscaled GEMM
      • 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
      • Attention
        • GQA Decode
      • Grouped GEMM
        • Unscaled and blockscaled grouped GEMM
      • Top-K
    • Ampere (CuTe extension):
      • SIMT GEMM
• 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)

This release has been tested against the following packages:

• FlashAttention: main (0251105)
• Quack: main (60d8808)
• FlashInfer: main (109d44f)
• cuDNN-Frontend: deveop (25b3d51)
• Pytorch: main (cf30153)
• TensorRT-LLM: main (1cef02e)

CUTLASS Operator API

• Operator API features and functionality:

  • 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, in addition to those exposed through the Operator interface in cutlass.operators. This allows kernels to be called directly without looking them up first.

  • Custom Epilogue fusions now support partial (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 dev.