CuTe DSL

• Bug fixing and improvements
  • Fixed a failed kernel compilation issue when setmaxnreg is used together with specific warp-specialized patterns
    (!3420)
  • Fixed a leak with jit/kernel decorator (!3421)
  • Fixed bugs in cutlass.jax.cutlass_call with tensor aliasing and handling of optional tensors
  • Fixed unexpected compile cache-misses with cutlass.jax.cutlass_call
  • Reduced the protobuf version requirement of IKET profiler from 6.30 to 4.21
  • Fixed import time write to CUTE_DSL_LIBS (!3462)
  • Fixed export_to_c outputs i32 dynamic shapes in headers for i64 inputs (!3447)
  • Improved performance of constructing cutlass.Numeric from Python value (!3443)
  • Fixed TVM-FFI env stream detection for GPU tensors in tuple
    (!3444)
  • Fixed TVM-FFI error when passed a tensor with non-zero byte offset
    (!3450)

This release has been tested against the following packages:

• FlashAttention: main (a369df7)
• Quack: main (60d8808)
• FlashInfer: main (145b101)
• cuDNN-Frontend:deveop(66efedf)
  • Pytorch: main(cf30153)

CuTe DSL

• New features:

  • Introduced the Primitives API which provides a lower-level abstraction beneath CuTe enabling Tensor Core programming through SIMT. This provides a stable, thin wrapper over NVVM operations to use where CuTe abstractions reduce development velocity. Primitives are released as experimental and will evolve based on user feedback.

    NOTE: Primitives is a transitional API until a CUDA Python-like solution is available.

  • Introduced the Task Scheduling framework. This provides static analysis of execution schedules for warp-specialized kernels. Compilation stops when known concurrency issues are detected. Also provides tools for visualizing resource/task dependencies and analyzing a kernel's schedule.

  • Improved compiler diagnostics. Register spills and use of local memory can now be reported at compile time with source line numbers. Preliminary support for detecting classes of NVVM synchronization and execution hazards at compile-time when using the Primitives API. Better reporting of compiler errors that previously did not include source line numbers.

This release has been tested against the following packages:

• FlashAttention: main (c75d019)
• Quack: main (79517ae)
• FlashInfer: main (4b964ec)
• cuDNN-Frontend:deveop(5235c2b)
• Pytorch: main(cf30153)

CUTLASS Operator API

• Custom epilogue fusions enhancements:
  • Add support for scalar reductions.
  • Add support to specify data movement strategy for each operand being loaded/stored.

C++

• Add implementation of 2-kernel backward targeting at FP8 in FMHA example.
  • Added a backward fused multi-head attention benchmark with multi-precision (FP16/FP8), configurable batch/sequence/head sizes, variable-length and masking options, plus built-in correctness checks and runtime/throughput reporting.
  • The 2-kernel backward has approximately 25% improvement compared to 1-kernel implementation at FP8 without mask on Blackwell SM103 chip.
• Support CUDA 12.6 and newer structured bindings headers in NVRTC.
• Add fp32/fp16/bf16/e4m3/e5m2 -> e2m1 (FP4) in NumericArrayConverter.
• Optimize the E2M1 -> FP16 LUT decode helpers _e2m1_to_half_x2 and _e2m1_to_half_x4 by merging mask before prmt.
• Fix some issues:
  • Update streamk heuristic algorithm to optimize some kernels with mix cluster sizes.
  • Avoid integer-sequence get ambiguity in CuTe tuple algorithms.
  • Fix a TMA creation driver bug: detect if the first 128KiB is mapped in conservatively by checking if the tensor is compact, if so it is valid to flip the bit otherwise zero it.
• Various improvements and fixes from the community and CUTLASS team. Thanks to everyone who submitted PRs!
• Optimal code generation with CUDA toolkit versions 13.3.

CuTe DSL

• Bug fixing and improvements
  • Fixed a failed kernel compilation issue when setmaxnreg is used together with specific warp-specialized patterns
    (!3382)
  • Fixed a leak with jit/kernel decorator (!3421)
  • Fixed bugs in cutlass.jax.cutlass_call with tensor aliasing and handling of optional tensors
  • Reduced the protobuf version requirement of IKET profiler from 6.30 to 4.21
  • Fixed import time write to CUTE_DSL_LIBS (!3462)
  • Fixed export_to_c outputs i32 dynamic shapes in headers for i64 inputs (!3447)
  • Improved performance of constructing cutlass.Numeric from Python value (!3443)
  • Fixed TVM-FFI env stream detection for GPU tensors in tuple
    (!3444)
  • Fixed TVM-FFI error when passed a tensor with non-zero byte offset
    (!3450)

This release has been tested against the following packages:

• FlashAttention: main (a369df7)
• Quack: main (680ef82)
• FlashInfer: main (0263dc2)
• cuDNN-Frontend:deveop(ab9efe15)
• Pytorch: main(cf30153)

CuTe DSL

• Bug fixing and improvements
  • Reverted the TMA bulk copy elect_one change from 4.6.0 to reset the behavior to align with 4.5.x releases
  • Fixed a vectorized fp32->f8 conversion issue (!3382)
  • Fixed a CuTe DSL fp8 grouped_gemm_dglu kernel compilation issue (!397)
  • Fixed opt-level setting issue for ptxas (!3389)
  • set_name_prefix of kernel now supports full customization on compiled kernel name (!3389)
  • Fixed issue in (!3349)
  • Fixed issue in (!3351)
  • Reduced JIT compile overhead by ~50 ms per cute.compile
  • Speed up importing cutlass.cute by 3.8x with Torch installed, 1.25x without.

This release has been tested against the following packages:

• FlashAttention: main (00756db)
• Quack: main (79517ae)
• FlashInfer: main (766f94b)

CuTe DSL

• Bug fixing and improvements
  • Fixed following issues:
    • #3243
    • #3359
    • #3365
    • #3312
  • Fixed a compilation failure on Thor with 12.9 wheel
  • Fixed a per regression of flash attention v2 on Ampere
  • Allow custom and multiple FFI call registration for Jax

• Release documentation that explains how to accurately profiling GEMM performance.

CuTe DSL

• New features

  • New fine-grained compilation API: cute.compile_to that gives control over the what stage the compiler outputs. This feature allows customization of the path from compilation to runtime execution. cute.compile_to is considered experimental in 4.6.
  • Experimental Feature: Added the IKET (In-Kernel-Event-Tracing) profiler for instrumentation-based intra-kernel activities tracing. This enables fine-grained profiling and makes it easier to understand persistent, warp-specialized kernels' performance. This is a beta feature provided by CUTLASS Python until a NVIDIA DevTools product is released, there is no guarantee that this interface will remain stable!
  • Distribute compiler binaries to accompany cute.compile_to allowing users to build customized compile-execute pipelines outside of Python. Both static and shared compiler and executor/runtime libraries will be provided. Compiler binaries will be uploaded to GitHub with each release.
  • Supported AoT cross-compilation for aarch64-linux-gnu
  • Support for two launch attributes: launch completion events (cudaLaunchAttributeLaunchCompletionEvent), for recording an event once all thread blocks have begun executing, and launch programmatic events (cudaLaunchAttributeProgrammaticEvent), for PDL event-based synchronization
  • Supported auto calculating per-kernel shared memory carveout preference, or use new launch option preferred_smem_carveout to set manually.
  • Auto-deduced smem size for launching kernels
    • Launch config smem now defaults to None for auto-calculating kernel shared memory usage, which is recommended unless manual control is required.
    • Warnings will be raised when the manually set shared memory size is insufficient or exceeds the GPU maximum.
    • The default shared memory usage calculation aligns with CUDA C++ static shared memory behavior, i.e. summing all allocations additively.
    • An additional launch option smem_merge_branch_allocs is provided to merge shared memory allocations across mutually exclusive code branches, which is recommended for inlined mega-kernels to reduce total footprint.
  • SASS dumping in DSL is now supported in a self-contained manner - no CUDA toolkit installation required to get nvdisasm

• Bug fixing and improvements

  • Add the missing elect_one in cute.copy for bulk copy.
    • The elect_one required for async bulk copy was missing in cute.copy. It's now generated in cute.copy automatically.
    • Nesting elect_one will cause functionality issues. Please remove elect_one around cute.copy with async bulk copy.
    • Elect_one around direct async bulk copy instruction should be kept as it bypasses the cute.copy layer and will not be affected by this fix.
    • Affected copy atoms are CopyBulkG2SOp, CopyBulkG2SMulticastOp, CopyBulkS2GOp, CopyBulkS2GByteMaskOp, and CopyBulkS2SOp.
    • An Example showing changes to avoid nesting elect_one could be found in this PR
  • Improvements on linter support with more type ignores cleaned up
  • Improvements on tvm-ffi CUDA runtime error diagnostics
  • Improvements on dataclass support for TVM-FFI
  • Fixed a regression on compilation time
  • Enhancement on compile time checks to reject mis-aligned smem operand for TMA
  • Long-deprecated API clean-up, including:
    • cute.core.ThrMma, please use cute.ThrMma instead
    • cute.core.ThrCopy, please use cute.ThrCopy instead
    • cute.make_fragment, please use cute.make_rmem_tensor instead
  • Fixed following issues
    • #3132
    • #3170
    • #3259
    • #3290
    • #3298
    • #3312
    • #3327
    • #3329
    • #3342
    • #3348

CUTLASS Operator API

• CUTLASS Operator API is a new addition to the CUTLASS Python stack, providing easy interfaces
  to discover CUTLASS Python DSL kernels & integrate them in your code.
  • pip install nvidia-cutlass-operators to get started
• Operator API Overview
  • See currently supported kernels
• Basic GEMM tutorial
  • More tutorials here
• GitHub source
• API Reference

CUTLASS C++

• Add example 113 for Hopper GEMM with activation fusion.
  • Supports standard and gated activations (e.g., SiLu) with fp8 and fp16 inputs.
  • Covers both regular GEMM and grouped GEMM variants.
• Improve SM90 grouped/ptr-array GEMM with EVT support.
  • Adds the EVT (Epilogue Visitor Tree) plumbing required to do activation, bias, and auxiliary-tensor fusion inside SM90 grouped and ptr-array GEMM kernels.
• Add ptr-array TMA collective for tensor/token-scaled FP8 grouped GEMM Blackwell SM120/SM121 kernels.
  • Implement CollectiveMma and CollectiveBuilder specializations for MainloopSm120ArrayTmaWarpSpecialized, enabling ptr-array grouped GEMM (MoE expert dispatch) with tensor- and token-level FP8 scaling.
  • Corresponding unit test
• Add tileN = 8,16 for Blackwell SM120 blockscale GEMM kernels.
• Fix DescriptorIterator::operator+ in mma_traits_sm100.hpp to use 32-bit arithmetic on CUDA toolkit version <= 13.3, preserving the high half of the smem descriptor.
• Fix a CUDA structured bindings header issue.
• Various improvements and fixes from the community and CUTLASS team. Thanks to everyone who submitted PRs!
• Optimal code generation with CUDA toolkit versions 13.3.

CuTe DSL

• Bug fixing and improvements
  • Fixed a compilation time regression issue in 4.5.0. Compilation times now match those in the 4.4 and 4.6 branches.

CuTe DSL

• New features

  • Python 3.14t is now supported with GIL enabled

• Bug fixing and improvements

  • Fixed following issues: https://github.com/NVIDIA/cutlass/issues/3240 https://github.com/NVIDIA/cutlass/issues/3241

CUTLASS C++

• Fix missing convert fucntion in EVT for fp4 kernels.
• Avoid instantiate 2sm tma kernels where ctaN is none power of 64 when ctaN > 128 in profiler.

CuTe DSL

• Bug fixing and improvements
  • 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
  • Fixed Jax int64 stride divisibility issue
  • Fixed issues for SM120 blockscaled MMAs
    • added missing MXFP8MMAOP and MXF8F6F4MMAOP for sm120.

CUTLASS C++

• 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.

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

• 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

• 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 some kernel issues:
  • Fix l2_capacity=0 handling in Blackwell SM100/SM120 kernel templates
  • Fix CUTLASS clang build issues
  • Fix atomicCAS read-modify-write loop in ConstSubbyteReference
  • Replace __nv_atomic_load_n with volatile for CUDA 11.4 compatibility in subbyte reference
  • Remove PipelineStorage shadowing in SM100 complex epilogue
  • Fix build issue in SM90 epilogue fusion visitor TMA warpspecialized
• Fix some profiler issues:
  • Add missing reference kernels for blockwise GEMM profiler.

CuTe DSL

• New features
  • CuTe DSL now supports Python 3.14 for both x86_64 and aarch64
  • Runtime Pointer/Tensor/FakeTensor now supports cache_key, providing a stable, hashable representation that simplifies and improves compiled function caching.
• Bug fixing and improvements
  • Fixed Hopper FMHA causal attention performance regression on CUDA toolkit 13.1 by optimizing mbarrier synchronization to avoid unnecessary convergence barriers.
  • Fix kernel loading race condition when multiple GPU are present in the same process in JAX.

CUTLASS C++

• Enable Blackwell SM120f compilation of examples and exposes NVFP4/MX Grouped GEMM in the CUTLASS Profiler.

CuTe DSL

• Bug fixing and improvements
  • Fixed a segfault issue with tvm-ffi on aarch64

CuTe DSL

• New features

  • CuTe DSL now supports CUDA toolkit 13.1!
    • Set up with cutlass/python/CuTeDSL/setup.sh --cu13
    • Refer to https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/quick_start.html for more details
  • GB300 is now supported in CuTe DSL with CTK 13.1
    • Refer to SM103 batched 3xFP4 blockscaled GEMM kernel for example kernel
  • cute.experimental: introduce a higher-level, composable layer on top of existing CuTe DSL APIs (not a separate abstraction), which can be mixed with existing Cute DSL building blocks.
    • Fragment-free programming model: copy/dot APIs take memrefs directly instead of descriptors/fragments.
    • Automatic TMA descriptor generation and update insertion.
    • Automatic vectorization and predication for SIMT copies.
    • New pipeline abstraction with convenience wrappers
    • New Partition ops to simplify partitioning logic.
    • Device-side TMA descriptor allocation, initialization, and management
    • These examples can be found here https://github.com/NVIDIA/cutlass/tree/main/examples/python/CuTeDSL/experimental
  • Ahead of Time (AoT) compilation is now available!
    • Refer to files under https://github.com/NVIDIA/cutlass/tree/main/examples/python/CuTeDSL/cute/export for example usage
  • JAX support - you can now use CuTeDSL along with JAX
    • Refer to files under https://github.com/NVIDIA/cutlass/tree/main/examples/python/CuTeDSL/jax for example usage
  • Introduced versioning support in DSL:
    • cutlass.version for a string representation of DSL version
    • cutlass.CUDA_VERSION for a version class to tell the CUDA version used for DSL
  • Added CopyDsmemStoreOp to store data to distributed shared memory with explicit synchronization.
  • Grouped GEMM example now supports device-only problem shapes.
  • We allow grid carve-out without problem shapes being available on host.
  • Tma+LdMatrix features for loading+unpacking narrow-width types (refer to mixed_input_fmha_decode.py for example usage).
  • It is possible now to have customized epilogue fusion for persistent dense GEMM through a Python Epilogue Fusion Configuration (EFC) function, somewhat similar to CUTLASS C++ EVT. It also provides a PyTorch evaluator to compare the results.

• More examples of authorizing peak-performance kernels

  • SM103 batched 3xFP4 blockscaled GEMM kernel
  • Mixed input FMHA decode example with support for int4 KV (int8 KV supported in 4.3)
  • New acc_scale grouped mixed input gemm kernel variant is introduced to deliver better performance for decoding cases.
  • All mixed_input_gemm examples are moved into a separate folder mixed_input_gemm. Common utility functions are also extracted into mixed_input_host_utils.py under the same folder.

• Bug fixing and improvements

  • Fixed an issue that both branches of if are executed
  • Fixed cute.printf with f-string
  • Fixed an indexing issue of scalar tensor
  • Fixed small K reference check error for cta_tile_n = 256 case with overlapping accumulator optimization in Blackwell SM100 persistent dense blockscaled GEMM with static scheduling.

• API changes

  • Deprecate get_num_tmem_alloc_cols from blackwell_helpers.py. Use the one from tmem_allocator.py instead.
  • Deprecate SM100_TMEM_CAPACITY_COLUMNS and SM100_TMEM_MIN_ALLOC_COLUMNS.
  • LdMatrix16x16x8bOp and StMatrix16x8x8bOp now require explicit transpose=True when calling init, to avoid ambiguity in data transposition.
  • LdMatrix16x16x8bOp copy traits updated to be faithful to PTX without permutations. Permuted variant is renamed to LdMatrix16x8x8bOp.
  • Grouped GEMM example takes the argument --host_problem_shape_available. If the argument is provided, grid is carved out based upon the host problem shapes, otherwise, we launch maximum possible SMs.
  • hardware_info.get_max_active_cluster support pass in specific stream to query. Useful for green context based SM partition.
  • group_bulk_copy_modes in async bulk copy example is now deprecated, use group_modes directly instead.
  • Deprecate nvvm wrapper from using nvvm enum, use str instead.
  • cute.arch.calc_packed_f32x2_op default enable ftz to default disable ftz
  • In CuTe DSL with CTK 13.1, following APIs in cutlass.cute.arch now require string literal instead of enum as argument:
    • fence_proxy
    • fence_view_async_tmem_op
    • calc_packed_f32x2_op
    • warp_redux_sync
    • atomic_add
    • atomic_and
    • atomic_or
    • atomic_xor
    • atomic_max
    • atomic_min
    • atomic_exch
    • atomic_cas
    • store
    • load

• Use 'Advanced control file' for mixed input gemm examples for better performance.

  • Advanced control file is an experimental feature of CUDA compiler. The controls file contains internal compiler settings tuned for specific kernels with a specific version of CUDA toolkit to get better GPU kernel code. More details and documentation on how to create these controls files will be provided in future CUDA toolkit release. Note: The advanced compiler control file is not expected to work for kernels that it was not tuned for. There is no compatibility guarantee, and the controls file will not work for CUDA toolkit with a different version.

CUTLASS C++

• Add example 93 for Blackwell low latency generation phase GQA kernel.
  • Flash Decoding with cluster reduction.
  • Kernel design details please check Readme.
• Add Blackwell SM100 State Space Decomposition (SSD) kernel in example 112.
• Add Hopper SM90 State Space Decomposition (SSD) kernel in example 111.
• Add example 94 for Ada FP8xFP8 -> BF16 GEMM with blockwise dequantization of input matrices in the MMA loop with FP32 accumulation.
  • Generate additional device/kernel/threadblock files in CUTLASS include directory that add functionality to carry the scaling tensors + use them in MMA loop.
  • Add gemm_blockwise to include files in default_mma_core_sm80
• Add Hopper e2m1 to fp32 optimized conversion and e2m1 * TF32 tensor core GEMM.
  • Set MmaType to tfloat32_t for FP32 mode.
  • TF32 provides FP32 inputs with reduced precision (19-bit vs 32-bit)
  • Set TileShapeK=64 for TF32 (K must be multiple of 8)
  • Shuffle optimization enabled via compute_memory_reordering_atom<tfloat32_t>()
  • E2M1 -> FP32 -> TF32 TC path for mixed-precision GEMM
  • Enable example 55 with TF32 support
• Add support for arbitrary application-provided strides for block-scale tensors.
  • Users and applications now must pass valid block-scale strides in all cases, even when the tensor is packed.
• Support 4x blockscaled public ptx for CUDA 13.1.
• Allow non-static TmaGbasis in AuxTmaParams.
  • Some cases in attention kernel may require non-static tma_gbasis.
  • Relax the restriction on TmaGbasis parameter of AuxTmaParams and users are allowed to manually construct a dynamic gbasis.
• Fix some kernel issues:
  • Fix MSVC pre process issue.
  • Fix a self assign issue in GEMV kernel.
  • Fix a TMA descriptor bug where the CUDA driver is not properly setting the OOB address gen mode correctly.
  • Fix memory fence for clc scheduler in Blackwell SM120 pingpong kernel.
  • Fix missing SMEM alignment in Blackwell SM120 scale factors.
  • Fix a PDL issue for grouped gemm.
  • Fix divide-by-zero issue in canimplement for sm100 implicit gemm kernels.
  • Fix cluster swizzle for Grouped GEMMs.
    • Move host-side swizzling heuristics to device.
    • Apply swizzle per group based on problem shape and max swizzle size.
    • Improve examples and unit tests.
• Fix some profiler issues:
  • Fix a core dump issue for nvfp4 grouped GEMM kernel.
  • Fix inconsistent GEMM verification logic.
  • Rework grouped gemm verification logic for different types.
  • Fix api break change in using nvMatmulHeuristics.
• Fix some failed links under media/docs.
• Various improvements and fixes from the community and CUTLASS team. Thanks to everyone who submitted PRs!
• Optimal code generation with CUDA toolkit versions 13.1.

CuTe DSL

• Bug fixing and improvements
  • Fixed the unexpected CPU overhead issue introduced by 4.3.4
• Update copyright to 2026.

CUTLASS C++

• Update copyright to 2026.
• Use CUDA Driver Get Version Runtime APIs Rather than Driver APIs.

CuTe DSL

• New features

  • Added PDL support along with example Kernel launch with Programmatic Dependent Launch

• Bug fixing and improvements

  • Fixed a frame refcnt issue with cuda graph
  • Enhancement for tvm-ffi AoT case for earlier module unload
  • Fixed order issue in make_smem_layout_a in utils/hopper_helpers.py

CUTLASS C++

• Work around a driver TMA descriptor related bug which will cause occasionally errors on Blackwell when the tensor's backing memory allocation is less than 128KB and it is not a dense non-overlapping tensor.

CuTe DSL

• New features

  • Supported namedtuple and kwargs for JIT function arguments in tvm-ffi
  • Supported variadic tuples for JIT function argument in tvm-ffi

• Bug fixing and improvements

  • Fixed an issue when JIT function argument with union type annotation for tvm-ffi
  • Clearer error message for the case of runtime error cudaErrorInsufficientDriver

CuTe DSL

• New features

  • New env var CUTE_DSL_CACHE_DIR to specify the path for dumping caches

• Bug fixing and improvements

  • Fixed an issue of CUDA JitExecutor when unloading kernels
  • Fixed an issue of allocating max smem when there's statically allocated smem

CuTe DSL

• New features
  • Added Blackwell SM103 support
  • Multiple dependent DSOs in the wheel have been merged into one single DSO
• Bug fixing and improvements
  • Fixed device reset issue with tvm-ffi
  • Fixed tvm-ffi export compiled function

CUTLASS C++

• Support blockscaled variant of ragged contiguous grouped gemm with the new simplified MoE API in example 92.
  • The new example works for all microscaling types.

CuTe DSL

• New features:
  • Supported Apache TVM-FFI for further reduced host runtime overhead for JIT functions, better PyTorch and ML frameworks interopability
  • Added fake tensor and stream to decouple compile jit function with "from_dlpack" flow. Now we no longer require users to have real tensor when compile jit function.
  • Added FastDivmodDivisor with Python operator overloads, new APIs, Cute dialect integration, and optimized static tile scheduler performance for faster index mapping.
  • Added l2 cache evict priority for tma related ops. Users could do fine-grain l2 cache control.
• Debuggability improvements:
  • Supported source location tracking for DSL APIs (Allow tools like nsight profiling to correlate perf metrics with Python source code)
  • Supported dumping PTX and CUBIN code: Hello World Example
• More examples and notebooks to get started with CuTe DSL:
  • Improved performance of elementwise example:
    • Generalize code to handle list of input tensors
    • Generalize TV layout computation to handle different data types
  • Improved Blackwell SM100 persistent dense GEMM with static scheduling:
    • To demonstrate usage of new Pipeline APIs PipelineProducer and PipelineConsumer to simplify code without explicit pipeline state management (Exiting APIs are still maintained)
    • Separated epilogue code for non-TMA and TMA implementation
  • Tutorial for Blackwell GEMM: Basic Blackwell SM100 GEMM
    • Baseline Blackwell GEMM achieves 84% SOL performance with MNK 8K
    • More examples are coming for demo of optimization: Baseline + X
  • Tutorial for Async Pipeline API
  • Reworked elementwise add notebook with more details and detailed explanation about TV layout
    • Updated implementation to handle general data type and multiple inputs
    • Updated explanation for TV layout in simpler language
    • Added visualization of TV Layout with 3rd party utils
  • Benchmark and autotune demonstration
• More examples of authorizing peak-performance kernels:
  • Blackwell SM100 mixed-input GEMM
  • Blackwell SM100 persistent blockwise dense GEMM
  • Blackwell SM100 persistent blockwise contiguous grouped dense GEMM
  • Blackwell SM100 persistent blockwise masked grouped dense GEMM
  • Blackwell SM100 fmha bwd
  • Blackwell SM100 mla
  • Hopper SM90 persistent dense GEMM with static scheduling
  • Blackwell GeForce batched dense GEMM
  • Ampere HSTU Attention
• API updates:
  • Please refer to DSL API changelog for details
• Bug fixings and improvements
  • Add mma_tiler_n=64 and mma_tiler_n=192 support in Blackwell SM100 persistent dense blockscaled GEMM with static scheduling.
  • Fixed TensorSSA.reduce to support static value as initial value
  • Updated docstring for following APIs to be more concise and easier to understand:
    • make_layout_tv
    • is_static
    • PipelineAsync
    • SmemAllocator
  • Fixed documentation for pipeline, utils and cute.math
  • Added overlapping accumulator optimization for block tile N = 256 case for better epilogue latency hiding in Blackwell SM100 persistent dense blockscaled GEMM with static scheduling.
  • Fixed TensorSSA.getitem indexing to match CuTe's indexing convention
  • Fixed an issue with cutlass.max and cutlass.min
  • Fixed an issue with mark_compact_shape_dynamic

CUTLASS C++

• Further enhance Blackwell SM100 Attention kernels in example 77.
  • Add softmax skip correction.
  • Fix a shared memory allocation bug where it needs to opt in maximum dynamics shared memory explicitly once it exceeds 48KB.
  • Fix a dead hang issue caused by early return warp.
• Add support through cmdline argument lists for batch, no_verif, cluster_shape and cluster_shape_fallback in example 89.
• Add Ragged Contiguous Grouped gemm kernel in example 92.
  • This kernel uses a TMA 3D load to load the weights matrix and use the tensormap update method to load activations.
• Add 256x128 tile size support for Hopper SM90 deepgemm in example 67.
  • Performance is optimized to align with Deepseek implementation.
• Simplification of API for MoE gemms.
  • Instead of requiring users to call several cute utilities to set up the stride, API moe_stride_utils is introduced to help setup strides in the kernel.
  • Instead of requiring users to set vectors like problem_shapes_device and problem_shapes_hosts, a new problem shape struct called MoEProblemShape is introduced which takes in max_m, max_n, max_k and counts vector as input and deduce problem shapes internally whenever required.
• Enable GEMM_K = 0 in grouped gemm.
• Optimize group gemm kernels by enabling async TMA desc update.
• Support Blackwell SM100 convolution stream-K kernel.
  • Unit tests: fprop_streamK, dgrad_streamK, wgrad_streamK.
• Add Blackwell SM100 sparse gemm compressor unit tests.
  • Unit tests: compressor_fp16.
  • Add sub-bytes and runtime data type support in compressor unit test testbed.
• Add profiler support for:
  • Blackwell SM100 and SM120 blockscaled sparse kernels.
  • New MoE grouped gemm API.
  • Blackwell SM100 cpasync kernel.
• Fix some kernel issues:
  • Fix a race check issue of Blackwell SM103 kernels by adding missing elect one for prefetch barrier initialization.
  • Allow user to directly specify the number of stages for Hopper sm90 mixed input gemm.
  • Remove warnings caused by cuda vector type alignment setting in CUDA 13.
  • Remove problematic cutlass::int8_t and replace it with int8_t.
  • Fix a few bugs in distributed gemm API and examples.
  • Fix handling negative zero in sparse compressor.
  • Add missing wait_on_dependent_grids for PDL use case.
• Fix some profiler issues:
  • Add some missing reference kernels.
  • Support VoidC reference kernels.
  • Add calculation of scale factor A and B in function bytes_with_problem_shape of block scaled profiler.
  • Fix an issue when epilogue tile N is not divided by default subtile N.
• Various improvements and fixes from the community and CUTLASS team. Thanks to everyone who submitted PRs!
• Optimal code generation with CUDA toolkit versions 13.0U1.

CuTe DSL

• Bug fixings and improvements
  • Fixed an issue when running DSL codes with cuda-python 13.0
  • Fixed an issue when running inductor with DSL codes
  • Fixed an issue with unexpected logging when running DSL codes in FlashInfer
  • Fixed the issue reported in https://github.com/NVIDIA/cutlass/issues/2647
  • Fixed an issue when conditional define of variables outside of dynamic control flow

CUTLASS C++

• Bypass EVT for nosmem blockwise kernels on Blackwell.
• Rename cutlass/python/cutlass directory to cutlass/python/cutlass_cppgen.

CuTe DSL

• More Python versions are now supported for both x86-64 and aarch64, including
  • Python 3.10, 3.11, 3.12, and 3.13
• Added new example and updated notebook to get started with CuTe DSL
  • Call kernels with dlpack bypassed
  • Updates on TensorSSA demonstration
    • Added a section for introducing the broadcast
• API updates
  • Please refer to DSL API changelog for details
• Bug fixings and improvements
  • Fixed cute.print_tensor for coordinate tensor
  • Fixed cute.print for tuple of layouts
  • Fixed frozen object is not properly updated after fully assigned in dynamic control flow
  • Fixed assign tuple/list element in a dynamic control flow may cause compilation failure
  • Improved error message when CUDA context is not initialized
  • Improved docstring of congruent and weakly_congruent

CUTLASS C++

• Support for Blackwell SM103 kernels for B300 GPUs.
  • Collective mainloop codes: Blockscaled datatypes with support for dense GEMM mainloop
  • New GEMM and epilogue dispatch policies for collectives, kernel layers, and builders.
  • Kernel codes: Blockscaled datatypes with support for dense GEMM kernel.
• Set of examples that demonstrate the usage of the 3.x API for targeting Blackwell SM103 architecture:
  • Blockscaled ultra fp4 dense GEMM.
  • Blockscaled ultra fp4 dense grouped GEMM.
• Set of unit tests that demonstrate the usage of Blackwell SM103 blockscaled GEMM
  • Unit test files with prefix name of sm103_ under GEMM device unit tests.
• Support for Blackwell SM121 kernels for DGX Spark GPUs.
  • Share the major codes with Blackwell SM120 kernels.
• Add support for heuristics-based kernel filtering and autotuning using nvidia-matmul-heuristics to find the best kernels for a given scenario.
  • Details please refer to heuristics doc.
• Further enhance Blackwell SM100 Attention kernels in example 77.
  • Add fused reduction kernel support for cutlass MLA.
  • Add softmax skip correction.
  • Support for GQA in FMHA backward kernel.
  • Fix an issue where get_unmasked_trip_count may return a negative value.
  • Fix an issue where mbarriers are initialized with a zero arrival count.
  • Fix a corner case issue where the sequence length of q is not a multiple of tile_q.
  • Remove tma padding for forward kernel inputs.
• Add Blackwell SM100 kernels for MoEs (focusing on Low-Latency inference performance): example 92. It uses TMA (for weights) and CPASYNC (for tokens) to load input matrices and allow only one problem dimension to vary across groups/experts, unlike general Grouped GEMMs. Note: further API simplifications and kernel improvements are upcoming. Any feedback on API is welcome.
• Further enhance blockwise and groupwise GEMMs on Hopper and Blackwell
  • On Blackwell SM120, a blockwise gemm kernel is added: example 87.
  • On Hopper, add K major scale factor support for SM90 blockwise kernels.
  • On Hopper, relax the restriction that the k dimension of the problem size has to be the multiple of the k dimension of the tile size.
  • On Hopper, grouped version supports the case when k = 0.
• Support for Blackwell SM100 fp4 gemv kernels.
  • Kernel codes: Gemv kernel.
  • Example codes: example 91
• Support for Blackwell SM100 legacy mixed input GEMM kernels.
  • Collective mainloop codes: Mixed input mainloop.
  • Kernel codes: Mixed input kernel.
  • Example codes: example 86.
• Support for Blackwell SM100 cpasync kernel.
  • Collective mainloop codes: cpasync mainloop.
  • Kernel codes: cpasync kernel.
• Support Blackwell SM120 mixed input blockscaled grouped GEMM.
• Instantiating more Blackwell kernels in profiler.
  • Blackwell SM100 and SM103 kernels support CUTLASS_LIBRARY_INSTANTIATION_LEVEL to instantiate all possible combinations.
  • To use this feature, CUTLASS_LIBRARY_KERNELS must be non-empty. Profiler will combine CUTLASS_LIBRARY_KERNELS and CUTLASS_LIBRARY_INSTANTIATION_LEVEL to instantiate specific kernels.
  • Details please check Profiler Doc.
• Fix some profiler issues:
  • Modify default cluster callback values to none 0 to avoid profiler failure when these values are not set in command line.
  • Fix some no output and timeout issues.
  • Fix Pingpong Blockwise Hopper library generation.
• From CUDA 13.0, the Blackwell SM101 for Thor GPUs is renamed to SM110.
  • For CUDA toolkit version < 13.0, SM101 is still used for Thor GPUs.
  • For CUDA toolkit version >= 13.0, SM110 is used for Thor GPUs and SM101 is no longer valid.
• Rename legacy Python API package from cutlass to cutlass_cppgen and add Blackwell EVT support to legacy Python interface.
  • Restructuring the C++ Blackwell SM100 Collective Epilogue Builder to work with the Python interface's EpilogueDescriptors.
  • Added Blackwell SM100 EVT Emitter on the Python side and routed most emission through Hopper SM90 Emitter.
  • Added some support for running SM100 kernels via the Python interface.
• CuTe changes:
  • Fix inaccurate GridDim calculation under CuTe tutorial.
  • Add movmatrix support.
  • Fix smallest MMA-N allowed for Blackwell fp8 and fp16 gemm kernels.
  • Support fp16 accmulator for sm89 fp8 mma.
  • Shorten nullspace implementation.
  • Isolate and comment on cosize hacks.
  • Important documentation correction: E<0,1> == 1@0@1.
• Fix some kernel issues:
  • Fix Hopper SM90 group gemm kernel to only use the commit group and wait group instead of also waiting on mbarriers.
  • Fix a tiny bug when K is large for Blackwell SM103 fp4 grouped GEMM kernel.
• Add following unit tests:
  • fp16 accmulator for sm89 fp8 mma
  • movmatrix test
  • fp8 narrow mma n and fp16 narrow mma n
• Various improvements and fixes from the community and CUTLASS team. Thanks to everyone who submitted PRs!
• Optimal code generation with CUDA toolkit versions 13.0U1.

CuTe DSL

• Add aarch64 support, you can now pip install nvidia-cutlass-dsl on GB200 systems!
• More examples demonstrating how to use CuTe DSL to write peak-performance kernels
  • Blackwell Mamba2 SSD
  • Blackwell SM100 persistent dense blockscaled GEMM with static scheduling
• API updates
  • Please refer to FUNCTIONALITY.md for details

CUTLASS C++

• Further enhance Blackwell SM100 Attention kernels in example 77.
  • Add variable sequence length support for FMHA Backward kernel.
  • Add varlen test support to Backward runner.
  • Codes support empty batch sequences.
• Replace subbyte_iterator with cute::recast_ptr when constructing logical iterators/arrays.
• CuTe changes:
  • Rewrite ArithTuple and ScaledBasis for robustness and clarity.
  • Remove buggy and kludgy get_layoutA|B|C_MN and friends from Atoms/TiledX.
  • Factor out print_latex and friends and rewrite.
  • Factor out print_svg and friends and rewrite.
• Support Blackwell SM100 SIMT packed fp32x2 kernels.
• Support residual add for implicit gemm kernels.
• Various fixes for CUTLASS C++ Python interface's EVT tracer:
  • Add verifier for sm90 to report the invalid input.
  • When adding an edge to the graph, if the edge already exists, add an identity compute node to avoid having multiple parallel edges.
  • Register operations of tanh, sigmoid, exp, gelu to the python ast frontend.
  • Replace the NotImplemented Error by packing all nodes into a single topological visitor node as a fallback.
• Fix profiler bugs in exhaustive perf search.
  • Fix incorrect cluster shape output issue when doing exhaustive search.
  • Fix a bug in profiler grouped GEMM for setting tile scheduler swizzles, cluster shapes, and raster orders.
• Fix some profiler issues.
  • Complete the reference for Blackwell blockwise gemm kernels.
  • Fix incorrect regex logic for L1 test.

CuTe DSL

CuTe DSL is a Python DSL centered around CuTe's abstractions

• Enables authoring kernels in Python to reach peak performance on NVIDIA GPUs
• Core DSL implementation files
• DSL quick start
• DSL Overview
• Educational notebooks for getting started with CuTe DSL

CUTLASS C++

• Support Family Specific Architecture Features which was introduced in CUDA 12.9
• Further improved Blockwise and Groupwise GEMMs on Hopper and Blackwell
• Enhance Blackwell SM100 Attention kernels in example 77
• Add Blackwell SM100 implicit GEMM conv fprop/dgrad/wgrad unit tests
• New Hopper SM90 FMHA example, similar in design to the existing Blackwell FMHA
• Cute enhancements: CuTe C++ reduce op
• Other functional and performance enhancements