### 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 * 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](https://github.com/NVIDIA/cutlass/tree/main/examples/python/CuTeDSL/ampere/call_bypass_dlpack.py) - Updates on [TensorSSA demonstration](https://github.com/NVIDIA/cutlass/tree/main/examples/python/CuTeDSL/notebooks/tensorssa.ipynb) + Added a section for introducing the broadcast * API updates - Please refer to [DSL API changelog](https://docs.nvidia.com/cutlass/media/docs/pythonDSL/cute_dsl_api/changelog.html) 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](https://github.com/NVIDIA/cutlass/tree/main/include/cutlass/gemm/collective/sm103_blockscaled_mma_warpspecialized.hpp) - New [GEMM](https://github.com/NVIDIA/cutlass/tree/main/include/cutlass/gemm/dispatch_policy.hpp) and [epilogue](https://github.com/NVIDIA/cutlass/tree/main/include/cutlass/epilogue/dispatch_policy.hpp) dispatch policies for collectives, kernel layers, and builders. - Kernel codes: [Blockscaled datatypes with support for dense GEMM kernel](https://github.com/NVIDIA/cutlass/tree/main/include/cutlass/gemm/kernel/sm103_blockscaled_gemm_tma_warpspecialized.hpp). * Set of examples that demonstrate the usage of the 3.x API for targeting Blackwell SM103 architecture: - [Blockscaled ultra fp4 dense GEMM](https://github.com/NVIDIA/cutlass/tree/main/examples/89_sm103_fp4_ultra_gemm/). - [Blockscaled ultra fp4 dense grouped GEMM](https://github.com/NVIDIA/cutlass/tree/main/examples/90_sm103_fp4_ultra_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](https://github.com/NVIDIA/cutlass/tree/main/test/unit/gemm/device/). * 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](https://github.com/NVIDIA/cutlass/tree/main/media/docs/cpp/heuristics.md). * Further enhance Blackwell SM100 Attention kernels in [example 77](https://github.com/NVIDIA/cutlass/tree/main/examples/77_blackwell_fmha/). - 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](https://github.com/NVIDIA/cutlass/tree/main/examples/92_blackwell_moe_gemm/). 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](https://github.com/NVIDIA/cutlass/tree/main/examples/87_blackwell_geforce_gemm_blockwise/). - 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](https://github.com/NVIDIA/cutlass/tree/main/include/cutlass/gemm/kernel/gemv_blockscaled.h). - Example codes: [example 91](https://github.com/NVIDIA/cutlass/tree/main/examples/91_fp4_gemv/) * Support for Blackwell SM100 legacy mixed input GEMM kernels. - Collective mainloop codes: [Mixed input mainloop](https://github.com/NVIDIA/cutlass/tree/main/include/cutlass/gemm/collective/sm100_mma_warpspecialized_mixed_input.hpp). - Kernel codes: [Mixed input kernel](https://github.com/NVIDIA/cutlass/tree/main/include/cutlass/gemm/kernel/sm100_gemm_tma_warpspecialized_mixed_input_transform.hpp). - Example codes: [example 86](https://github.com/NVIDIA/cutlass/tree/main/examples/86_blackwell_mixed_dtype_gemm/). * Support for Blackwell SM100 cpasync kernel. - Collective mainloop codes: [cpasync mainloop](https://github.com/NVIDIA/cutlass/tree/main/include/cutlass/gemm/collective/sm100_mma_cpasync_warpspecialized.hpp). - Kernel codes: [cpasync kernel](https://github.com/NVIDIA/cutlass/tree/main/include/cutlass/gemm/kernel/sm100_gemm_cpasync_warpspecialized.hpp). * 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](https://github.com/NVIDIA/cutlass/tree/main/media/docs/cpp/profiler.md). * 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](https://github.com/NVIDIA/cutlass/tree/main/examples/cute/tutorial/blackwell/). - Add [movmatrix](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-instructions-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](https://github.com/NVIDIA/cutlass/tree/main/test/unit/cute/ampere/cooperative_gemm.cu) - [movmatrix test](https://github.com/NVIDIA/cutlass/tree/main/test/unit/cute/turing/movm.cu) - [fp8 narrow mma n](https://github.com/NVIDIA/cutlass/tree/main/test/unit/gemm/device/sm100_tensorop_gemm/f16_f16_void_f32_narrow_mma_n.cu) and [fp16 narrow mma n](test/unit/gemm/device/sm100_tensorop_gemm/f8_f8_void_bf16_narrow_mma_n.cu) * 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](https://github.com/NVIDIA/cutlass/tree/main/examples/python/CuTeDSL/blackwell/mamba2_ssd/mamba2_ssd.py) - [Blackwell SM100 persistent dense blockscaled GEMM with static scheduling](https://github.com/NVIDIA/cutlass/tree/main/examples/python/CuTeDSL/blackwell/dense_blockscaled_gemm_persistent.py) * API updates - Please refer to [FUNCTIONALITY.md](https://github.com/NVIDIA/cutlass/blob/main/FUNCTIONALITY.md) for details **CUTLASS C++** * Further enhance Blackwell SM100 Attention kernels in [example 77](https://github.com/NVIDIA/cutlass/tree/main/examples/77_blackwell_fmha/). - 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** * 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](https://github.com/NVIDIA/cutlass/tree/main/examples/python/CuTeDSL/blackwell/mamba2_ssd/mamba2_ssd.py) - [Blackwell SM100 persistent dense blockscaled GEMM with static scheduling](https://github.com/NVIDIA/cutlass/tree/main/examples/python/CuTeDSL/blackwell/dense_blockscaled_gemm_persistent.py) * API updates - Please refer to [FUNCTIONALITY.md](https://github.com/NVIDIA/cutlass/blob/main/FUNCTIONALITY.md) for details **CUTLASS C++** * Further enhance Blackwell SM100 Attention kernels in [example 77](https://github.com/NVIDIA/cutlass/tree/main/examples/77_blackwell_fmha/). - 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.

* Fixed [Blockwise](./examples/67_hopper_fp8_warp_specialized_gemm_with_blockwise_scaling/67_hopper_fp8_warp_specialized_gemm_with_blockwise_scaling.cu) and [Groupwise](./examples/67_hopper_fp8_warp_specialized_gemm_with_blockwise_scaling/67_hopper_fp8_warp_specialized_gemm_with_groupwise_scaling.cu) GEMM hang issue when problem size K is 128. * Optimal code generation with CUDA toolkit versions 12.9.

* Fixed Group Gemm hang issue in CUTLASS 3.x * Improved Hopper [Blockwise](./examples/67_hopper_fp8_warp_specialized_gemm_with_blockwise_scaling/67_hopper_fp8_warp_specialized_gemm_with_blockwise_scaling.cu) and [Groupwise](./examples/67_hopper_fp8_warp_specialized_gemm_with_blockwise_scaling/67_hopper_fp8_warp_specialized_gemm_with_groupwise_scaling.cu) GEMM performance.

* Support for Blackwell SM120 kernels for GeForce GPUs in CUTLASS 3.x API: - Collective mainloops that target for: * [Blockscaled datatypes with support for dense GEMM](./include/cutlass/gemm/collective/sm120_blockscaled_mma_tma.hpp) * [Blockscaled datatypes with support for sparse GEMM](./include/cutlass/gemm/collective/sm120_blockscaled_sparse_mma_tma.hpp) - New [GEMM](./include/cutlass/gemm/dispatch_policy.hpp) and [epilogue](./include/cutlass/epilogue/dispatch_policy.hpp) dispatch policies for collectives, kernel layers, and builders. - [Blackwell SM120 epilogue](./include/cutlass/epilogue/fusion/sm120_visitor_store_tma_warpspecialized.hpp) and [full set of EVT fusions](./include/cutlass/epilogue/fusion/sm120_callbacks_tma_warpspecialized.hpp). * Set of examples that demonstrate the usage of the 3.x API for targeting Blackwell SM120 architecture: - [Blockscaled GEMM with NVFP4 input datatype and BF16 output tensor](./examples/79_blackwell_geforce_gemm/79a_blackwell_geforce_nvfp4_bf16_gemm.cu). - [Blockscaled GEMM with NVFP4 input datatype and NVFP4 output tensor with scale factor generation](./examples/79_blackwell_geforce_gemm/79b_blackwell_geforce_nvfp4_nvfp4_gemm.cu). - [Blockscaled GEMM with mixed input datatype (MXFP8 and MXFP6) and BF16 output tensor](./examples/79_blackwell_geforce_gemm/79c_blackwell_geforce_mixed_mxfp8_mxfp6_bf16_gemm.cu). - [Grouped GEMM with nvfp4 datatype](./examples/79_blackwell_geforce_gemm/79d_blackwell_geforce_nvfp4_grouped_gemm.cu). - [Sparse Blockscaled GEMM with mxfp8 input datatype and BF16 output tensor](./examples/80_blackwell_geforce_sparse_gemm/80a_blackwell_geforce_mxfp8_bf16_sparse_gemm.cu). - [Sparse Blockscaled GEMM with NVFP4 input datatype and NVFP4 output tensor](./examples/80_blackwell_geforce_sparse_gemm/80b_blackwell_geforce_nvfp4_nvfp4_sparse_gemm.cu). * Set of unit tests that demonstrate the usage of both [sparse](./test/unit/gemm/device/sm120_blockscaled_sparse_tensorop_gemm/) and [dense](./test/unit/gemm/device/sm120_blockscaled_tensorop_gemm/) Blackwell SM120 blockscaled GEMM. * Support for Blackwell SM100 Sparse kernels: - Collective mainloop that target for * [SM100 Sparse GEMM](./include/cutlass/gemm/collective/sm100_sparse_mma_warpspecialized.hpp) * Set of example that demonstrate the usage of the 3.x API for targeting Blackwell SM100 Sparse GEMM: - [Sparse GEMM](./examples/83_blackwell_sparse_gemm/83_blackwell_sparse_gemm.cu) - [Blockscaled Sparse GEMM with NVFP4 input data type](./examples/84_blackwell_narrow_precision_sparse_gemm/84a_blackwell_nvfp4_bf16_sparse_gemm.cu) - [Blockscaled Sparse GEMM with mixed input data type (MXFP8 and MXFP4)](./examples/84_blackwell_narrow_precision_sparse_gemm/84b_blackwell_mixed_mxfp8_bf16_sparse_gemm.cu) * Set of unit tests that demonstrate the usage of [sparse](./test/unit/gemm/device/sm100_sparse_tensorop_gemm) and [blockscaled sparse](./test/unit/gemm/device/sm100_blockscaled_sparse_tensorop_gemm) Blackwell SM100 GEMM. * A new Multi-head Latent Attention (MLA) for SM100 Blackwell architecture in CUTLASS [example](./examples/77_blackwell_fmha/) covers the flashMLA-like weight-absorbed decoding use-case. * A new FMHA Backward kernel for SM100 Blackwell architecture extends CUTLASS [example](./examples/77_blackwell_fmha/) to show how the five backward pass MMAs can be fused into a single kernel to achieve high performance. * A new [distributed GEMM example](./examples/82_blackwell_distributed_gemm/82_blackwell_distributed_gemm.cu) for SM100 Blackwell architecture. * Enhancement and new support of block-wise and group-wise GEMM for Hopper and Blackwell architectures: - Enhancement of [blockwise GEMM](./examples/67_hopper_fp8_warp_specialized_gemm_with_blockwise_scaling/67_hopper_fp8_warp_specialized_gemm_with_blockwise_scaling.cu) for Hopper architecture. - Enhancement of [groupwise GEMM](./examples/67_hopper_fp8_warp_specialized_gemm_with_blockwise_scaling/67_hopper_fp8_warp_specialized_gemm_with_groupwise_scaling.cu) for Hopper architecture. - Support for [grouped GEMM with blockwise and groupwise scaling](./examples/68_hopper_fp8_warp_specialized_grouped_gemm_with_blockwise_scaling/) for Hopper architecture. - Support for [grouped-wise GEMM](./tools/profiler/src/blockwise_gemm_operation_profiler.cu) in CUTLASS profiler. - Support for [blockwise GEMM](./examples/81_blackwell_gemm_blockwise/81_blackwell_gemm_blockwise.cu) for Blackwell architecture. - Support for [groupwise GEMM](./examples/81_blackwell_gemm_blockwise/81_blackwell_gemm_groupwise.cu) for Blackwell architecture. - Support for [grouped GEMM with blockwise](./examples/81_blackwell_gemm_blockwise/81_blackwell_grouped_gemm_blockwise.cu) and [groupwise scaling](./examples/81_blackwell_gemm_blockwise/81_blackwell_grouped_gemm_groupwise.cu) for Blackwell architecture. * Added support for enhanced kernel performance search (auto-tuning) in CUTLASS profiler: - Sorting performance results by GFLOPs/second: Users can now sort the final performance report based on GFLOPs/second, making it easier to identify the most efficient kernels. - Exhaustive search for best kernel performance in GFLOPs/second: The profiler now searches for the best-performing kernel across a range of problem sizes, swizzle sizes, rasterization orders, and dynamic cluster configurations to maximize performance. - Performance search under a fixed GEMM shape: Enables exhaustive tuning within a fixed GEMM shape, exploring various kernel parameters to find the best configuration. - More detailed introductions and examples to leverage this feature can be found in [profiler.md](./media/docs/cpp/profiler.md#exhaustive-search-mode-and-top-k-output-ranking-according-to-performance-in-gflopss). * Support `void` as the D element in sm100 kernel epilogues.