diff --git a/ACTIVE_DEVELOPERS.md b/ACTIVE_DEVELOPERS.md
new file mode 100644
index 00000000..6ae47b43
--- /dev/null
+++ b/ACTIVE_DEVELOPERS.md
@@ -0,0 +1,73 @@
+
+
+[README](./README.md#documentation) > **Active Developers**
+
+# CUTLASS Developers **
+
+Andrew Kerr (CUTLASS founding member)
+Dustyn Blasig
+Albert Xu
+Junkai Wu
+Xiuxia Zhang
+Haicheng Wu (CUTLASS founding member)
+Jack Yang
+Pradeep Ramani (CUTLASS 3.x founding member)
+Aditya Atluri
+Han Li
+Nick Zhao
+Ivan Yin
+Yu-Jung Chen
+Markus Hoehnerbach
+Honghao Lu
+Mihir Awatramani
+Hao Sheng
+Zekun Fan
+Aniket Shivam
+Siyu Liu
+Richard Cai
+Vikas Gupta
+Ethan Yan
+Vijay Thakkar (CUTLASS 3.x founding member)
+Cris Cecka (CuTe and CUTLASS 3.x founding member)
+Lawrence Ryan
+Qun Song
+Daniel Ricketts
+dePaul Miller
+Yuhan Li
+Saman Ashkiani
+Jack Chen
+Shang Zhang
+Petrick Liu
+Questa Wang
+Pramod Shenoy
+Jack Kosaian
+Yujia Zhai
+Zhaodong Chen
+Manas Sahni
+Shunfan Shao
+Fengqi Qiao
+Serif Yesil
+Aragorn Guan
+Heidi He
+Xiao Song
+Sergey Klevtsov
+Jiang Shao
+Ruqing Xu
+Mengyu Guo
+Tao Xie
+Linfeng Zheng
+Harrison Barclay
+Wenfei Tang
+Diksha Gohlyan
+Alexander Zhurkevich
+Siyuan Fu
+Hua Huang
+Xiufan Liang
+Ian Tramble
+Ali Hassani
+Shreya Gaur
+
+** _The list is sorted in order of the author's first contribution to the CUTLASS project._
+
+# CUTLASS Product Manager
+Matthew Nicely
diff --git a/CHANGELOG.md b/CHANGELOG.md
index 76f6bf6a..95419bcb 100644
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@@ -1,8 +1,59 @@
# NVIDIA CUTLASS Changelog
+
+## [3.8.0](https://github.com/NVIDIA/cutlass/releases/tag/v3.8.0) (2025-01-25)
+
+* Support for new CuTe building blocks specifically for Blackwell SM100 architecture:
+ - [5th generation Blackwell Tensor Core instructions (TCGen05)](./include/cute/atom/mma_traits_sm100.hpp) via CuTe MMA atoms.
+ - Extensions to [Tensor Memory Accelerator](./include/cute/atom/copy_traits_sm100_tma.hpp) via CuTe Copy atoms.
+ - Exposure of Blackwell's new tensor memory (note: distinct from TMA) as [`tmem`](./include/cute/pointer.hpp) across CuTe as a first class data locale.
+ - Exposure of [`tmem->rmem`, `rmem->tmem` and `smem->tmem data movement instructions`](./include/cute/atom/copy_traits_sm100.hpp) as copy atoms in CuTe.
+ - [`make_tmem_copy()`](./include/cute/atom/copy_traits_sm100.hpp) utility method to ease creation of tiled copies for tmem copy atoms.
+ - Support for [new variants of LDSM on Blackwell](./include/cute/atom/copy_traits_sm100.hpp) via CuTe Copy atoms.
+* Support for new CUTLASS building blocks specifically for Blackwell SM100 architecture:
+ - Various narrow precision [FP4, FP6, and FP8](./include/cutlass/exmy_base.h) formats as well as their [block-scaled variants NVFP4, MXFP4, MXFP6, and MXFP8](./include/cutlass/float_subbyte.h)
+ - [Pipelines that implement Blackwell specific synchronization](./include/cutlass/pipeline/sm100_pipeline.hpp).
+ - [Cluster launch control API supporting preferred and fallback cluster shapes](./include/cutlass/cluster_launch.hpp).
+ - Data types including NVFP4, MXFP4, MXFP6, and MXFP8 and all their supported element and scale factor types.
+ - Tile schedulers using [Blackwell's Cluster Launch Control (CLC) feature](./cutlass/media/docs/blackwell_cluster_launch_control.md) to implement dynamic persistence scheduling for [GEMMs](./include/cutlass/gemm/kernel/sm100_tile_scheduler.hpp), and [stream-K](./include/cutlass/gemm/kernel/sm100_tile_scheduler_stream_k.hpp).
+ - Extensions to testbeds and reference check code for unit tests and CUTLASS profiler.
+* Full support for Blackwell SM100 kernels in CUTLASS 3.x API:
+ - [Blackwell specific kernel layers](./include/cutlass/gemm/kernel/sm100_gemm_tma_warpspecialized.hpp) that
+ + Implement a new warp-specialization recipe tuned specifically for Blackwell SM100 architecture.
+ + Leverage all the new features such as CLC based tile scheduling, preferred cluster, and TMEM based double buffering of accumulators.
+ + Support stream-K load balancing for all kernel types everywhere via composable scheduler support.
+ - Blackwell collective mainloops that target the TCGen05 MMA instructions (both SS and TS) for
+ * [Non-block scaled data types without support for pointer array and grouped GEMM with TMA](./include/cutlass/gemm/collective/sm100_mma_warpspecialized.hpp)
+ * [Non-block scaled data types with support for pointer array and grouped GEMM with TMA](./include/cutlass/gemm/collective/sm100_mma_array_warpspecialized.hpp)
+ * [Block scaled data types without support for pointer array and grouped GEMM with TMA](./include/cutlass/gemm/collective/sm100_blockscaled_mma_warpspecialized.hpp)
+ * [Block scaled data types with support for pointer array and grouped GEMM with TMA](./include/cutlass/gemm/collective/sm100_blockscaled_mma_array_warpspecialized.hpp)
+ - Blackwell [collective mainloop for convolution kernels](./include/cutlass/conv/collective/sm100_implicit_gemm_umma_warpspecialized.hpp) supporting non-block scaled data types for fprop, dgrad, and wgrad.
+ - New [GEMM](./include/cutlass/gemm/dispatch_policy.hpp), [convolution](./include/cutlass/conv/dispatch_policy.hpp), and [epilogue](./include/cutlass/epilogue/dispatch_policy.hpp) dispatch policies for collectives, kernel layers, and builders.
+ - [Blackwell epilogue that supports loading accumulators from `tmem`](./include/cutlass/epilogue/collective/sm100_epilogue_tma_warpspecialized.hpp) and [full set of EVT fusions]().
+* CUTLASS library and profiler integration for block scaled data types for kernel emission, profiling, and verification.
+ - Support for preferred and fallback cluster shapes via profiler command line arguments parsing to set dynamic cluster shapes.
+ - Support for dynamic datatypes by parsing profiler via profiler command line arguments parsing to set dynamic datatype setting in TCGen05 MMA instruction descriptors.
+* Set of examples that demonstrate the usage of the 3.x API for targeting Blackwell SM100 architecture:
+ - [Basic FP16 and FP8 GEMMs with minimal changes from Hopper examples](./examples/70_blackwell_gemm/), demonstrating ease of migration for off the shelf kernels using the 3.x collective builder API.
+ - GEMM with [opt-in collective builder schedules showcasing available recipes](./examples/71_blackwell_gemm_with_collective_builder/71_blackwell_gemm_with_collective_builder.cu) for Blackwell.
+ - Block scaled data type GEMMs targeting Blackwell's native block scaled Tensor Cores:
+ + [NVFP4 inputs with BF16 output](./examples/72_blackwell_narrow_precision_gemm/72a_blackwell_nvfp4_bf16_gemm.cu)
+ + [NVFP4 inputs with NVFP4 output](./examples/72_blackwell_narrow_precision_gemm/72b_blackwell_nvfp4_nvfp4_gemm.cu)
+ + [Mixed MXFP8 and MXFP6 inputs with BF16 output](./examples/72_blackwell_narrow_precision_gemm/72c_blackwell_mixed_mxfp8_bf16_gemm.cu)
+ - GEMM example demonstrating [Blackwell's new preferred cluster support via dynamic cluster shapes](./examples/73_blackwell_gemm_preferred_cluster/blackwell_gemm_preferred_cluster.cu) for increased occupancy.
+ - [GEMM with CLC based StreamK scheduler for load balancing](./examples/74_blackwell_gemm_streamk/blackwell_gemm_streamk.cu).
+ - Grouped GEMM for [vanilla FP8 data inputs](./examples/75_blackwell_grouped_gemm/75_blackwell_grouped_gemm.cu) and [NVFP4 block scaled inputs](./examples/75_blackwell_grouped_gemm/75_blackwell_grouped_gemm_block_scaled.cu).
+ - Convolution kernels for [fprop](./examples/76_blackwell_conv/76_blackwell_conv_fprop.cu), [dgrad](./examples/76_blackwell_conv/76_blackwell_conv_dgrad.cu), and [wgrad](./examples/76_blackwell_conv/76_blackwell_conv_wgrad.cu).
+ - [Fused multi-head attention fprop kernel](./examples/77_blackwell_fmha/77_blackwell_fmha.cu) supporting fp16/bf16/fp8 data types across head dims of 32,64, and 128.
+* Documentation updates:
+ - [Quickstart - instantiating a Blackwell block-scaled GEMM](./media/docs/quickstart.md#instantiating-a-blackwell-gemm-kernel).
+ - Detailed [Blackwell block-scaled GEMM functionality documentation](./media/docs/narrow_and_mixed_precision_gemms.md)
+ - A new [functionality documentation](./media/docs/functionality.md) specifically for 3.x API comprehensively documenting all supported kernel types, data types, kernel features, minimum CUDA tookit support etc for 3.x supported architectures.
+ - Updates to [compatibility](./README.md#compatibility) section regarding supported compilers, operating systems, CUDA Toolkits, Hardware Architectures, and [Target Architecture](./README.md#Target-Architecture).
+
## [3.7.0](https://github.com/NVIDIA/cutlass/releases/tag/v3.7.0) (2025-01-11)
- [Hopper blockwise scaling FP8 GEMM](./examples/67_hopper_fp8_warp_specialized_gemm_with_blockwise_scaling/67_hopper_fp8_warp_specialized_gemm_with_blockwise_scaling.cu) uses 2D scaling tensor, assigning one value per threadblock. This allows a finer-grained scaling to be applied for each output tile per gemm-k iteration. The operands and scaling tensors are loaded from global memory to shared memory using TMA and cp_async, respectively. The scaling is applied inside the mainloop. Details with figures are [here](https://github.com/NVIDIA/cutlass/pull/1932#issue-2645398439).
- [Distributed GEMM](./examples/65_distributed_gemm/65_distributed_gemm.cu) is a new (experimental) API which can turn existing CUTLASS GEMM kernels into pipelined Tensor Parallel GEMMs that run efficiently on NVLink-based network of GPUs. Its pipelining schedules can hide most of the communication behind computation, and relies on point-to-point communication, which can simply use CUDA runtime's peer device access feature. It also utilizes remote TMA loads and memcopies with CUDA graphs to handle communication primarily through the Copy Engine, leaving all SMs free for Hopper's persistent kernels. For more details you can refer to the [DistGEMM blog post](https://blog.shi-labs.com/distributed-gemm-88be6a481e2b).
-- Improved persistent grid launch for Hopper kernels with large cluster sizes (>= size of 4) using the new `make_kernel_hardware_info` API as shown in [example 48](./examples/48_hopper_warp_specialized_gemm/48_hopper_warp_specialized_gemm.cu).
+- Improved persistent grid launch for Hopper kernels with large cluster sizes (>= size of 4) using the new `make_kernel_hardware_info` API as shown in [example 48](./examples/48_hopper_warp_specialized_gemm/48_hopper_warp_specialized_gemm.cu).
- Enabled high precision accumulation for Hopper FP8 Sparse GEMM.
- Potential API breaking changes:
+ Fix `cute::UniversalCopy` for type safety.
@@ -22,12 +73,7 @@
+ [INT8](./test/unit/gemm/device/sm90_sparse_gemm_s8_s8_s32_tensor_op_s32.cu)
+ [TF32](./test/unit/gemm/device/sm90_sparse_gemm_tf32_tf32_f32_tensor_op_f32.cu)
- A refactor to the CUTLASS 3.x convolution `kernel::ConvUniversal` [API](./include/cutlass/conv/kernel/sm90_implicit_gemm_tma_warpspecialized.hpp) to bring it in line with `gemm::GemmUniversal`. Now the 3.x convolution API is no longer considered as a beta API.
-- Improve [mixed input GEMM](./examples/55_hopper_mixed_dtype_gemm/README.md).
- + Added a [lookup table implementation](./examples/55_hopper_mixed_dtype_gemm/55_hopper_int4_fp8_gemm.cu) for `INT4`x`FP8` scale-only mode.
- + Added [layout pre-shuffling](./examples/55_hopper_mixed_dtype_gemm/55_hopper_int4_fp8_gemm.cu#L50-55) to optimize memory loading.
- + Added [interleaved conversion](./examples/55_hopper_mixed_dtype_gemm/55_hopper_int4_bf16_gemm.cu#L50-52) for `{INT4, UINT4, INT8}` x `{FP16, BF16}`.
- + Other general optimizations.
-- The suffixes of the mixed input kernel schedules have been removed. Use `KernelTmaWarpSpecialized`, `KernelTmaWarpSpecializedPingpong` and `KernelTmaWarpSpecializedCooperative` instead.
+- [An improved mixed input GEMM](./examples/55_hopper_mixed_dtype_gemm/README.md) and a [lookup table implementation](./examples/55_hopper_mixed_dtype_gemm/55_hopper_int4_fp8_gemm.cu) for `INT4`x`FP8` scale-only mode.
- [EVT nodes for Top-K selection and softmax](./include/cutlass/epilogue/fusion/sm90_visitor_topk_softmax.hpp) and [GEMM example using those](./examples/61_hopper_gemm_with_topk_and_softmax/61_hopper_gemm_with_topk_and_softmax.cu).
- [Programmatic Dependent Launch](./include/cutlass/arch/grid_dependency_control.h) (PDL) that leverages a new Hopper feature to speedup two back-to-back kernels, and its corresponding [documentations](./media/docs/dependent_kernel_launch.md).
- [A new debugging tool, synclog](./include/cutlass/arch/synclog.hpp), for dumping out all synchronization events from within a kernel to a file. Please see [synclog documentation](./media/docs/utilities.md#debugging-asynchronous-kernels-with-cutlasss-built-in-synclog-tool) for details.
diff --git a/CMakeLists.txt b/CMakeLists.txt
index e50fd76e..9892f067 100755
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -164,6 +164,11 @@ endif()
if (CUDA_VERSION VERSION_GREATER_EQUAL 12.0)
list(APPEND CUTLASS_NVCC_ARCHS_SUPPORTED 90a)
endif()
+
+if (CUDA_VERSION VERSION_GREATER_EQUAL 12.8)
+ list(APPEND CUTLASS_NVCC_ARCHS_SUPPORTED 100 100a)
+endif()
+
set(CUTLASS_NVCC_ARCHS ${CUTLASS_NVCC_ARCHS_SUPPORTED} CACHE STRING "The SM architectures requested.")
set(CUTLASS_NVCC_ARCHS_ENABLED ${CUTLASS_NVCC_ARCHS} CACHE STRING "The SM architectures to build code for.")
@@ -383,6 +388,21 @@ endif()
+
+###################################################################################################
+#
+# Blackwell features
+#
+###################################################################################################
+
+if (CUDA_VERSION VERSION_GREATER_EQUAL 12.8)
+ list(APPEND CUTLASS_CUDA_NVCC_FLAGS -DCUDA_BLACKWELL_TMA_SWIZZLE_ENABLED=1)
+
+ list(APPEND CUTLASS_CUDA_NVCC_FLAGS -DCUDA_ENABLE_PREFERRED_CLUSTER=1)
+endif()
+
+
+
# Warnings-as-error exceptions and warning suppressions for Clang builds
if (CUTLASS_CLANG_HOST_COMPILE)
diff --git a/CONTRIBUTORS.md b/CONTRIBUTORS.md
deleted file mode 100644
index 538bb658..00000000
--- a/CONTRIBUTORS.md
+++ /dev/null
@@ -1,87 +0,0 @@
-
-
-[README](./README.md#documentation) > **Contributors**
-
-# CUTLASS Developers and Contributors
-
-This is the official list of CUTLASS developers and contributors.
-
-## DEVELOPERS
-Vijay Thakkar
-Pradeep Ramani
-Cris Cecka
-Aniket Shivam
-Jack Kosaian
-Mark Hoemmen
-Richard Cai
-Honghao Lu
-Ethan Yan
-Haicheng Wu
-Andrew Kerr
-Dustyn Blasig
-Fengqi Qiao
-Duane Merrill
-Yujia Zhai
-Rawn Henry
-Sergey Klevtsov
-Shang Zhang
-Piotr Majcher
-Paul Springer
-Markus Hohnerbach
-Jin Wang
-Aditya Atluri
-
-## CuTe
-Cris Cecka
-Vijay Thakkar
-
-## CUTLASS Product Manager
-Matthew Nicely
-
-## Former CUTLASS Developers
-Manish Gupta
-Naila Farooqui
-David Tanner
-Manikandan Ananth
-Zhaodong Chen
-Chinmay Talegaonkar
-
-## CONTRIBUTORS
-Timothy Costa
-Julien Demouth
-Brian Fahs
-Michael Garland
-Michael Goldfarb
-Mostafa Hagog
-Fei Hu
-Alan Kaatz
-Tina Li
-Timmy Liu
-Wei Liu
-Tim Martin
-Duane Merrill
-Kevin Siu
-Markus Tavenrath
-John Tran
-Vicki Wang
-Junkai Wu
-Fung Xie
-Albert Xu
-Yang Xu
-Jack Yang
-Scott Yokim
-Xiuxia Zhang
-Nick Zhao
-
-## ACKNOWLEDGEMENTS
-
-Girish Bharambe
-Luke Durant
-Carter Edwards
-Olivier Giroux
-Stephen Jones
-Rishkul Kulkarni
-Bryce Lelbach
-Joel McCormack
-Kyrylo Perelygin
-Sean Treichler
diff --git a/README.md b/README.md
index 56637806..a74ac114 100644
--- a/README.md
+++ b/README.md
@@ -1,8 +1,8 @@

-# CUTLASS 3.7.0
+# CUTLASS 3.8.0
-_CUTLASS 3.7.0 - January 2025_
+_CUTLASS 3.8.0 - January 2025_
CUTLASS is a collection of CUDA C++ template abstractions for implementing
high-performance matrix-matrix multiplication (GEMM) and related computations at all levels
@@ -16,71 +16,96 @@ as building blocks within custom kernels and applications.
To support a wide variety of applications, CUTLASS provides extensive support for
mixed-precision computations, providing specialized data-movement and
-multiply-accumulate abstractions for half-precision floating
-point (FP16), BFloat16 (BF16), Tensor Float 32 (TF32),
-single-precision floating point (FP32),
-[FP32 emulation via tensor core instruction](./examples/27_ampere_3xtf32_fast_accurate_tensorop_gemm),
-double-precision floating
-point (FP64) types, integer data types (4b and 8b), and binary data types (1b).
-CUTLASS demonstrates warp-synchronous matrix multiply operations
+multiply-accumulate abstractions for FP64, FP32, TF32, FP16, BF16,
+[FP32 emulation via tensor core instruction](./examples/27_ampere_3xtf32_fast_accurate_tensorop_gemm),
+ 8b floating point types (e5m2 and e4m3),
+ block scaled data types (NVIDIA NVFP4 and OCP standard MXFP4, MXFP6, MXFP8),
+ narrow integer types (4 and 8b signed and unsigned integers),
+ and binary 1b data types (where architectures allow for the
+native support of such data types).
+CUTLASS demonstrates optimal matrix multiply operations
targeting the programmable, high-throughput _Tensor Cores_ implemented by
-NVIDIA's Volta, Turing, Ampere, and Hopper architectures.
+NVIDIA's Volta, Turing, Ampere, Ada, Hopper, and Blackwell architectures.
+
+In addition to GEMMs, CUTLASS implements high-performance convolution via
+the implicit GEMM algorithm. Implicit GEMM is the formulation of a convolution
+operation as a GEMM thereby taking advantage of CUTLASS's modular GEMM pipeline.
+This allows CUTLASS to build convolutions by reusing highly-optimized GEMM components.
See the [Quick Start Guide](./media/docs/quickstart.md) to get started quickly.
-See the [functionality listing](./media/docs/functionality.md) for the list of operations
-supported at each level of the execution model hierarchy.
+See the [functionality docs](./media/docs/functionality.md) for a more comprehensive
+list of kernel level features, data types, instructions, and minimum supported by CUTLASS on each GPU
+architecture.
-CUTLASS 3.0 introduced a new core library, CuTe, to describe and manipulate tensors of threads and data.
-CuTe is a collection of C++ CUDA template abstractions for defining and operating on hierarchically multidimensional layouts of threads and data. CuTe provides `Layout` and `Tensor` objects that compactly package the type, shape, memory space, and layout of data, while performing the complicated indexing for the user. This lets programmers focus on the logical descriptions of their algorithms while CuTe does the mechanical bookkeeping for them. With these tools, we can quickly design, implement, and modify all dense linear algebra operations.
+# What's New in CUTLASS 3.8
-The core abstractions of CuTe are hierarchically multidimensional layouts which can be composed with data arrays to represent tensors. The representation of layouts is powerful enough to represent nearly everything we need to implement efficient dense linear algebra. Layouts can also be combined and manipulated via functional composition, on which we build a large set of common operations such as tiling and partitioning.
+CUTLASS 3.8 is the first release that supports the NVIDIA Blackwell SM100 architecture.
+For a background on Blackwell's new features, please consult the PTX documentation for CUDA 12.8.
-CUTLASS 3.0 and beyond adopts CuTe throughout the GEMM hierarchy in its templates. This greatly simplifies the design
-and improves code composability and readability. More documentation specific to CuTe can be found in its [dedicated documentation directory](./media/docs/cute/00_quickstart.md).
+* Support for new CuTe building blocks specifically for Blackwell architecture:
+ - [5th generation Blackwell Tensor Core instructions (TCGen05)](./include/cute/atom/mma_traits_sm100.hpp) via CuTe MMA atoms.
+ - Extensions to [Tensor Memory Accelerator](./include/cute/atom/copy_traits_sm100_tma.hpp) via CuTe Copy atoms.
+ - Exposure of Blackwell's new tensor memory (note: distinct from TMA) as [`tmem`](./include/cute/pointer.hpp#L290) across CuTe as a first class data locale.
+ - Exposure of [`tmem->rmem`, `rmem->tmem` and `smem->tmem data movement instructions`](./include/cute/atom/copy_traits_sm100.hpp) as copy atoms in CuTe.
+ - [`make_tmem_copy()`](./include/cute/atom/copy_traits_sm100.hpp) utility method to ease creation of tiled copies for tmem copy atoms.
+ - Support for [new variants of LDSM on Blackwell](./include/cute/atom/copy_traits_sm100.hpp) via CuTe Copy atoms.
+* Support for new CUTLASS building blocks specifically for Blackwell architecture:
+ - Various narrow precision [FP4, FP6, and FP8](./include/cutlass/exmy_base.h) formats as well as their [block-scaled variants NVFP4, MXFP4, MXFP6, and MXFP8](./include/cutlass/float_subbyte.h)
+ - [Pipelines that implement Blackwell specific synchronization](./include/cutlass/pipeline/sm100_pipeline.hpp).
+ - [Cluster launch control API supporting preferred and fallback cluster shapes](./include/cutlass/cluster_launch.hpp).
+ - Data types including NVFP4, MXFP4, MXFP6, and MXFP8 and all their supported element and scale factor types.
+ - Tile schedulers using [Blackwell's Cluster Launch Control (CLC) feature](./cutlass/media/docs/blackwell_cluster_launch_control.md) to implement dynamic persistence scheduling for [GEMMs](./include/cutlass/gemm/kernel/sm100_tile_scheduler.hpp), and [stream-K](./include/cutlass/gemm/kernel/sm100_tile_scheduler_stream_k.hpp).
+ - Extensions to testbeds and reference check code for unit tests and CUTLASS profiler.
+* Full support for Blackwell kernels in CUTLASS 3.x API:
+ - [Blackwell specific kernel layers](./include/cutlass/gemm/kernel/sm100_gemm_tma_warpspecialized.hpp) that
+ + Implement a new warp-specialization recipe tuned specifically for Blackwell.
+ + Leverage all the new features such as CLC based tile scheduling, preferred cluster, and TMEM based double buffering of accumulators.
+ + Support stream-K load balancing for all kernel types everywhere via composable scheduler support.
+ - Blackwell collective mainloops that target the TCGen05 MMA instructions (both SS and TS) for
+ * [Non-block scaled data types without support for pointer array and grouped GEMM with TMA](./include/cutlass/gemm/collective/sm100_mma_warpspecialized.hpp)
+ * [Non-block scaled data types with support for pointer array and grouped GEMM with TMA](./include/cutlass/gemm/collective/sm100_mma_array_warpspecialized.hpp)
+ * [Block scaled data types without support for pointer array and grouped GEMM with TMA](./include/cutlass/gemm/collective/sm100_blockscaled_mma_warpspecialized.hpp)
+ * [Block scaled data types with support for pointer array and grouped GEMM with TMA](./include/cutlass/gemm/collective/sm100_blockscaled_mma_array_warpspecialized.hpp)
+ - Blackwell [collective mainloop for convolution kernels](./include/cutlass/conv/collective/sm100_implicit_gemm_umma_warpspecialized.hpp) supporting non-block scaled data types for fprop, dgrad, and wgrad.
+ - New [GEMM](./include/cutlass/gemm/dispatch_policy.hpp), [convolution](./include/cutlass/conv/dispatch_policy.hpp), and [epilogue](./include/cutlass/epilogue/dispatch_policy.hpp) dispatch policies for collectives, kernel layers, and builders.
+ - [Blackwell epilogue that supports loading accumulators from `tmem`](./include/cutlass/epilogue/collective/sm100_epilogue_tma_warpspecialized.hpp) and [full set of EVT fusions]().
+* CUTLASS library and profiler integration for block scaled data types for kernel emission, profiling, and verification.
+ - Support for preferred and fallback cluster shapes via profiler command line arguments parsing to set dynamic cluster shapes.
+ - Support for dynamic datatypes by parsing profiler via profiler command line arguments parsing to set dynamic datatype setting in TCGen05 MMA instruction descriptors.
+* Set of examples that demonstrate the usage of the 3.x API for targeting Blackwell
+ - [Basic FP16 and FP8 GEMMs with minimal changes from Hopper examples](./examples/70_blackwell_gemm/), demonstrating ease of migration for off the shelf kernels using the 3.x collective builder API.
+ - GEMM with [opt-in collective builder schedules showcasing available recipes](./examples/71_blackwell_gemm_with_collective_builder/71_blackwell_gemm_with_collective_builder.cu) for Blackwell.
+ - Block scaled data type GEMMs targeting Blackwell's native block scaled Tensor Cores:
+ + [NVFP4 inputs with BF16 output](./examples/72_blackwell_narrow_precision_gemm/72a_blackwell_nvfp4_bf16_gemm.cu)
+ + [NVFP4 inputs with NVFP4 output](./examples/72_blackwell_narrow_precision_gemm/72b_blackwell_nvfp4_nvfp4_gemm.cu)
+ + [Mixed MXFP8 and MXFP6 inputs with BF16 output](./examples/72_blackwell_narrow_precision_gemm/72c_blackwell_mixed_mxfp8_bf16_gemm.cu)
+ - GEMM example demonstrating [Blackwell's new preferred cluster support via dynamic cluster shapes](./examples/73_blackwell_gemm_preferred_cluster/blackwell_gemm_preferred_cluster.cu) for increased occupancy.
+ - [GEMM with CLC based StreamK scheduler for load balancing](./examples/74_blackwell_gemm_streamk/blackwell_gemm_streamk.cu).
+ - Grouped GEMM for [vanilla FP8 data inputs](./examples/75_blackwell_grouped_gemm/75_blackwell_grouped_gemm.cu) and [NVFP4 block scaled inputs](./examples/75_blackwell_grouped_gemm/75_blackwell_grouped_gemm_block_scaled.cu).
+ - Convolution kernels for [fprop](./examples/76_blackwell_conv/76_blackwell_conv_fprop.cu), [dgrad](./examples/76_blackwell_conv/76_blackwell_conv_dgrad.cu), and [wgrad](./examples/76_blackwell_conv/76_blackwell_conv_wgrad.cu).
+ - [Fused multi-head attention fprop kernel](./examples/77_blackwell_fmha/77_blackwell_fmha.cu) supporting fp16/bf16/fp8 data types across head dims of 32,64, and 128.
+* Documentation updates:
+ - [Quickstart - instantiating a Blackwell block-scaled GEMM](./media/docs/quickstart.md#instantiating-a-blackwell-gemm-kernel).
+ - Detailed [Blackwell block-scaled GEMM functionality documentation](./media/docs/narrow_and_mixed_precision_gemms.md)
+ - A new [functionality documentation](./media/docs/functionality.md) specifically for 3.x API comprehensively documenting all supported kernel types, data types, kernel features, minimum CUDA tookit support etc for 3.x supported architectures.
+ - Updates to [compatibility](./README.md#compatibility) section regarding supported compilers, operating systems, CUDA Toolkits, Hardware Architectures, and [Target Architecture](./README.md#Target-Architecture).
-In addition to GEMMs, CUTLASS implements high-performance convolution via the implicit GEMM algorithm. Implicit GEMM is the formulation of a convolution operation as a GEMM thereby taking advantage of CUTLASS's modular GEMM pipeline. This allows CUTLASS to build convolutions by reusing highly-optimized GEMM components.
+Note: CUTLASS 3.x builds are known to be broken on Windows platforms for all CUDA toolkits.
+CUTLASS team is working on a fix.
-# What's New in CUTLASS 3.7
-
-CUTLASS 3.7.0 is an update to CUTLASS adding:
-
-- A new [Hopper blockwise scaling FP8 GEMM](./examples/67_hopper_fp8_warp_specialized_gemm_with_blockwise_scaling/67_hopper_fp8_warp_specialized_gemm_with_blockwise_scaling.cu) where the operands and block scaling tensor are staged via shared memory.
-- [Distributed GEMM](./examples/65_distributed_gemm/65_distributed_gemm.cu) is an experimental pipelined Tensor Parallelism implementation utilizing existing CUTLASS kernels and CUDA runtime features, which can hide the most of communication behind computation.
-- Improved persistent grid launch for Hopper kernels with large cluster sizes (>= size of 4) using the new `make_kernel_hardware_info` API as shown in [example 48](./examples/48_hopper_warp_specialized_gemm/48_hopper_warp_specialized_gemm.cu).
-- Enabled high precision accumulation for Hopper FP8 Sparse GEMM.
-- Potential API breaking changes:
- + Fix `cute::UniversalCopy` for type safety.
- + No longer implicitly select `cute::SM80_CP_ASYNC_*` based on input tensors. This avoids implicit downstream synchronization requirements. To use `SM80_CP_ASYNC`, users must explicitly select the appropriate CopyAtom.
- + Fix `cute::SM80_CP_ASYNC_CACHEALWAYS`, `cute::SM80_CP_ASYNC_CACHEGLOBAL`, `cute::SM80_CP_ASYNC_CACHEALWAYS_ZFILL`, `cute::SM80_CP_ASYNC_CACHEGLOBAL_ZFILL` to avoid implicitly selecting `ZFILL` behavior on predication.
- + Remove `cute::copy_vec` in favor of `cute::copy_aligned` and `cute::copy(AutoVectorizingCopyWithAssumedAlignment,...)`.
- + A refactor of default epilogue struct `DefaultEpilogue` [API](./include/cutlass/epilogue/collective/default_epilogue.hpp) to avoid reading non-void `ElementC` value for `ElementC = void` kernel.
-- New CUTLASS profiler flags: `profiling-duration`, `min-iterations`, and `kernels-file` documented in [profiler.md](./media/docs/profiler.md#cutlass-profiler).
-- Various improvements and fixes from the community and CUTLASS team. Thanks to everyone who submitted PRs!
-
-Minimum requirements:
-
-- Architecture: Volta
-- Compiler: Must support at least C++17
-- CUDA Toolkit version: 11.4
-
-Starting from CUTLASS 3.0, CUTLASS removed support for the following:
-
-- Maxwell and Pascal GPU architectures
-- Ubuntu 16.04
-- CUDA 10.2
-- C++ language versions less than 17.
-
-**See the [CHANGELOG](CHANGELOG.md) for a detailed listing of releases and updates.**
+**See the [CHANGELOG](CHANGELOG.md) for details of all past releases and updates.**
# Performance
-
-
-
CUTLASS primitives are very efficient. When used to construct device-wide GEMM kernels,
-they exhibit peak performance comparable to cuBLAS for scalar GEMM
-computations. The above figure shows the continual CUTLASS performance improvements
+they exhibit nearly optimal utilization of peak theoretical throughput. The figure below
+shows CUTLASS 3.8's performance as a % of theoretical peak utilization
+on various input and output data types when run on NVIDIA Blackwell SM100 architecture GPU.
+
+
+
+The two figures below show the continual CUTLASS performance improvements
on an [NVIDIA H100](https://www.nvidia.com/en-us/data-center/h100/) (NVIDIA Hopper architecture) since
CUTLASS 3.1.
CUTLASS 3.5.1 was compiled with the [CUDA 12.5u1 Toolkit](https://developer.nvidia.com/cuda-downloads).
@@ -88,20 +113,45 @@ Tensor Core operations are implemented using CUDA's
[mma](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-instructions-mma) and
[wgmma](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-matrix-instructions) instructions.
-
+
+
-When using CUTLASS building blocks to construct device-wide implicit gemm (Fprop, Dgrad, and Wgrad)
-kernels, CUTLASS performance is also comparable to cuDNN when running Resnet-50 layers on an [NVIDIA A100](https://www.nvidia.com/en-us/data-center/a100/)
-as shown in the above figure. Tensor Core operations are implemented using CUDA's
-[mma instruction](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-instructions-mma).
+# CuTe
+
+CUTLASS 3.0 introduced a new core library, CuTe, to describe and manipulate tensors of threads and data.
+CuTe is a collection of C++ CUDA template abstractions for
+defining and operating on hierarchically multidimensional layouts of threads and data.
+CuTe provides `Layout` and `Tensor` objects that compactly package the type,
+shape, memory space, and layout of data, while performing the complicated indexing for the user.
+This lets programmers focus on the logical descriptions of their algorithms while
+CuTe does the mechanical bookkeeping for them. With these tools, we can quickly design,
+implement, and modify all dense linear algebra operations.
+
+The core abstractions of CuTe are hierarchically multidimensional layouts
+which can be composed with data arrays to represent tensors.
+The representation of layouts is powerful enough to represent nearly
+everything we need to implement efficient dense linear algebra.
+Layouts can also be combined and manipulated via functional composition, on which we build a large set of common operations such as tiling and partitioning.
+
+CUTLASS 3.0 and beyond adopts CuTe throughout the GEMM hierarchy in its templates.
+This greatly simplifies the design and improves code composability and readability.
+More documentation specific to CuTe can be found in its
+[dedicated documentation directory](./media/docs/cute/00_quickstart.md).
# Compatibility
+Minimum requirements:
+
+- Architecture: Volta (compute capability 7.0)
+- Compiler: Must support at least C++17
+- CUDA Toolkit version: 11.4
+
CUTLASS requires a C++17 host compiler and
-performs best when built with the [**CUDA 12.4 Toolkit**](https://developer.nvidia.com/cuda-downloads).
-It is also compatible with CUDA 11.4, CUDA 11.5, CUDA 11.6, CUDA 11.7, CUDA 11.8, CUDA 12.0, CUDA 12.1, CUDA 12.2.2, CUDA 12.3.1 and CUDA 12.3.2.
+performs best when built with the [**CUDA 12.8 Toolkit**](https://developer.nvidia.com/cuda-downloads).
+It is also compatible with CUDA 11.4, CUDA 11.5, CUDA 11.6, CUDA 11.7, CUDA 11.8, and all other CUDA 12.x versions.
## Operating Systems
+
We have tested the following environments.
|**Operating System** | **Compiler** |
@@ -109,47 +159,74 @@ We have tested the following environments.
| Ubuntu 18.04 | GCC 7.5.0 |
| Ubuntu 20.04 | GCC 10.3.0 |
| Ubuntu 22.04 | GCC 11.2.0 |
-| Ubuntu 22.04 | Clang 10.0.0 |
-| Ubuntu 22.04 | Clang 14.0.6 |
-| Ubuntu 22.04 | Clang 17.0.6 |
-| Windows 10.0 | Visual Studio 2019 v16.11.27 |
Note: GCC 8.5.0 has known regressions regarding fold expressions and overloaded operators. Using GCC 7.5.0 or (preferred) GCC >= 9 is recommended.
+Note: CUTLASS 3.x builds are known to be broken on Windows platforms for all CUDA toolkits.
+CUTLASS team is working on a fix.
+
## Hardware
+
CUTLASS runs successfully on the following NVIDIA GPUs, and it is expected to be efficient on Volta, Turing, Ampere, Ada, and Hopper architecture based NVIDIA GPUs.
|**GPU**|**CUDA Compute Capability**|**Minimum CUDA Toolkit Required by CUTLASS-3**|
|---|---|---|
|NVIDIA V100 Tensor Core GPU |7.0|11.4|
|NVIDIA TitanV |7.0|11.4|
-|NVIDIA GeForce RTX 2080 TI, 2080, 2070 |7.5|11.4|
+|NVIDIA GeForce RTX 20x0 series |7.5|11.4|
|NVIDIA T4 |7.5|11.4|
|NVIDIA A100 Tensor Core GPU |8.0|11.4|
|NVIDIA A10 |8.6|11.4|
-|NVIDIA GeForce RTX 3090 |8.6|11.4|
-|NVIDIA GeForce RTX 4090 |8.9|11.8|
+|NVIDIA GeForce RTX 30x0 series |8.6|11.4|
+|NVIDIA GeForce RTX 40x0 series |8.9|11.8|
|NVIDIA L40 |8.9|11.8|
|NVIDIA H100 Tensor Core GPU |9.0|11.8|
+|NVIDIA H200 Tensor Core GPU |9.0|11.8|
+|NVIDIA B200 Tensor Core GPU |10.0|12.8|
## Target Architecture
-In general, PTX code generated for one target architecture can be run on future architectures (i.e., it is forward compatible). However, CUDA 12.0 introduced the concept of "architecture-accelerated features" whose PTX does not have forward compatibility guarantees. Several Hopper PTX instructions fall under this category of architecture-accelerated features, and thus require a `sm_90a` target architecture (note the "a" appended). For more details on this and other architecture-accelerated instructions, please refer to the [CUDA Documentation](https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#feature-availability).
+In general, PTX code generated for one target architecture can be run on future architectures
+(i.e., it is forward compatible).
+However, CUDA 12.0 introduced the concept of "architecture-accelerated features" whose
+PTX does not have forward compatibility guarantees.
+Several Hopper and Blackwell PTX instructions fall under this category of
+architecture-accelerated features, and thus require a `sm_90a` or `sm100a` target architecture
+(note the "a" appended). For more details on this and other architecture-accelerated instructions,
+please refer to the [CUDA Documentation](https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#feature-availability).
-The target architecture information is passed on to CUTLASS via the cmake flag `CUTLASS_NVCC_ARCHS`. In order to maximize performance on Hopper GH100, users are required to build CUTLASS with `90a` as the target architecture. If a user accidentally builds a kernel which uses SM90a features (e.g. Hopper Tensor Core Instructions), using the SM90 target (note the lack of "a"), with either CUDA Toolkit 12 or 11.8, the kernel is expected to fail with a runtime error.
+The target architecture information is passed on to CUTLASS via the cmake flag
+`CUTLASS_NVCC_ARCHS`. In order to maximize performance on Hopper GH100,
+users are required to build CUTLASS with `90a` as the target architecture.
+If a user accidentally builds a kernel which uses SM90a features
+(e.g. Hopper Tensor Core Instructions), using the SM90 target
+(note the lack of "a"), with either CUDA Toolkit 12 or 11.8,
+the kernel is expected to fail with a runtime error.
```
-cmake .. -DCUTLASS_NVCC_ARCHS="90a"
+cmake .. -DCUTLASS_NVCC_ARCHS="90a"
+```
+Or
+
+```
+cmake .. -DCUTLASS_NVCC_ARCHS="100a"
```
-Please refer to the [functionality documentation](./media/docs/functionality.md) for details on which kernels require which target architectures.
+Note: The NVIDIA Blackwell SM100 architecture used in the datacenter
+products has a different compute capability than the one underpinning
+NVIDIA Blackwell GeForce RTX 50 series GPUs. As a result, kernels
+compiled for Blackwell SM100 architecture with arch conditional features
+(using `sm100a`) are not compatible with RTX 50 series GPUs.
+
+Please refer to the [functionality documentation](./media/docs/functionality.md)
+for details on which kernels require which target architectures.
# Documentation
CUTLASS is described in the following documents and the accompanying
[Doxygen documentation](https://nvidia.github.io/cutlass).
-- [Quick Start Guide](./media/docs/quickstart.md) - build and run CUTLASS
+- [Quick Start Guide](./media/docs/quickstart.md) - basics of building and running CUTLASS
- [Functionality](./media/docs/functionality.md) - summarizes functionality available in CUTLASS
- [Efficient GEMM in CUDA](./media/docs/efficient_gemm.md) - describes how GEMM kernels may be implemented efficiently in CUDA
- [CUTLASS 3.x Design](./media/docs/cutlass_3x_design.md) - describes the CUTLASS 3.x design, its benefits, and how CuTe enables us to write much more composable components
@@ -163,7 +240,7 @@ CUTLASS is described in the following documents and the accompanying
- [Layouts](./media/docs/layout.md) - describes layouts of matrices and tensors in memory
- [Tile Iterators](./media/docs/tile_iterator_concept.md) - describes C++ concepts for iterating over tiles of matrices in memory
- [CUTLASS Profiler](./media/docs/profiler.md) - command-line driven profiling application
-- [CUTLASS Utilities](./media/docs/utilities.md) - additional templates used to facilate rapid development
+- [CUTLASS Utilities](./media/docs/utilities.md) - additional templates used to facilitate rapid development
- [Dependent kernel launch](./media/docs/dependent_kernel_launch.md) - describes a new feature in Hopper which allows overlapping dependent
kernels in the same stream, and how it is used in CUTLASS.
@@ -171,11 +248,11 @@ kernels in the same stream, and how it is used in CUTLASS.
We have also described the structure of an efficient GEMM in our talk at the
[GPU Technology Conference 2018](http://on-demand.gputechconf.com/gtc/2018/presentation/s8854-cutlass-software-primitives-for-dense-linear-algebra-at-all-levels-and-scales-within-cuda.pdf).
- - [CUTLASS: Software Primitives for Dense Linear Algebra at All Levels and Scales within CUDA](https://www.nvidia.com/en-us/on-demand/session/gtcsiliconvalley2018-s8854/)
- - [Developing CUDA Kernels to Push Tensor Cores to the Absolute Limit on NVIDIA A100](https://www.nvidia.com/en-us/on-demand/session/gtcsj20-s21745/)
- - [Accelerating Convolution with Tensor Cores in CUTLASS](https://www.nvidia.com/en-us/on-demand/session/gtcspring21-s31883/)
- - [Accelerating Backward Data Gradient by Increasing Tensor Core Utilization in CUTLASS](https://www.nvidia.com/en-us/on-demand/session/gtcspring22-s41996/)
- - [CUTLASS: Python API, Enhancements, and NVIDIA Hopper](https://www.nvidia.com/en-us/on-demand/session/gtcfall22-a41131/)
+- [CUTLASS: Software Primitives for Dense Linear Algebra at All Levels and Scales within CUDA](https://www.nvidia.com/en-us/on-demand/session/gtcsiliconvalley2018-s8854/)
+- [Developing CUDA Kernels to Push Tensor Cores to the Absolute Limit on NVIDIA A100](https://www.nvidia.com/en-us/on-demand/session/gtcsj20-s21745/)
+- [Accelerating Convolution with Tensor Cores in CUTLASS](https://www.nvidia.com/en-us/on-demand/session/gtcspring21-s31883/)
+- [Accelerating Backward Data Gradient by Increasing Tensor Core Utilization in CUTLASS](https://www.nvidia.com/en-us/on-demand/session/gtcspring22-s41996/)
+- [CUTLASS: Python API, Enhancements, and NVIDIA Hopper](https://www.nvidia.com/en-us/on-demand/session/gtcfall22-a41131/)
# Building CUTLASS
diff --git a/examples/48_hopper_warp_specialized_gemm/48_hopper_warp_specialized_gemm.cu b/examples/48_hopper_warp_specialized_gemm/48_hopper_warp_specialized_gemm.cu
index 7ed4593c..97e9061e 100644
--- a/examples/48_hopper_warp_specialized_gemm/48_hopper_warp_specialized_gemm.cu
+++ b/examples/48_hopper_warp_specialized_gemm/48_hopper_warp_specialized_gemm.cu
@@ -489,6 +489,14 @@ int main(int argc, char const **args) {
<< "later (compute capability 90 or greater).\n";
return 0;
}
+
+ if (props.major != 9 || props.minor != 0) {
+ std::cerr
+ << "This example requires a GPU of NVIDIA's Hopper Architecture (compute capability 90).\n";
+ return 0;
+ }
+
+
//
// Parse options
//
diff --git a/examples/49_hopper_gemm_with_collective_builder/49_collective_builder.cu b/examples/49_hopper_gemm_with_collective_builder/49_collective_builder.cu
index 5852cd8d..ee841526 100644
--- a/examples/49_hopper_gemm_with_collective_builder/49_collective_builder.cu
+++ b/examples/49_hopper_gemm_with_collective_builder/49_collective_builder.cu
@@ -540,6 +540,15 @@ int main(int argc, char const **args) {
<< "later (compute capability 90 or greater) and CUDA 12.0 or greater.\n";
return 0;
}
+
+ else if (__CUDACC_VER_MAJOR__ < 12 || props.major != 9 || props.minor != 0) {
+ std::cout
+ << "This example requires a GPU of NVIDIA's Hopper Architecture "
+ << "(compute capability 90) and CUDA 12.0 or greater.\n";
+ return 0;
+ }
+
+
//
// Parse options
//
diff --git a/examples/50_hopper_gemm_with_epilogue_swizzle/50_hopper_gemm_with_epilogue_swizzle.cu b/examples/50_hopper_gemm_with_epilogue_swizzle/50_hopper_gemm_with_epilogue_swizzle.cu
index 69a3c030..6e91a3ba 100644
--- a/examples/50_hopper_gemm_with_epilogue_swizzle/50_hopper_gemm_with_epilogue_swizzle.cu
+++ b/examples/50_hopper_gemm_with_epilogue_swizzle/50_hopper_gemm_with_epilogue_swizzle.cu
@@ -356,6 +356,15 @@ int main(int argc, char const **args) {
<< "later (compute capability 90 or greater) and CUDA 12.0 or greater.\n";
return 0;
}
+
+ else if (__CUDACC_VER_MAJOR__ < 12 || props.major != 9 || props.minor != 0) {
+ std::cout
+ << "This example requires a GPU of NVIDIA's Hopper Architecture "
+ << "(compute capability 90) and CUDA 12.0 or greater.\n";
+ return 0;
+ }
+
+
//
// Parse options
//
diff --git a/examples/52_hopper_gather_scatter_fusion/52_hopper_gather_scatter_fusion.cu b/examples/52_hopper_gather_scatter_fusion/52_hopper_gather_scatter_fusion.cu
index 8a198420..49505284 100644
--- a/examples/52_hopper_gather_scatter_fusion/52_hopper_gather_scatter_fusion.cu
+++ b/examples/52_hopper_gather_scatter_fusion/52_hopper_gather_scatter_fusion.cu
@@ -626,6 +626,13 @@ int main(int argc, const char ** argv) {
std::cerr << "This example requires a device with compute capability 90 or higher.\n";
notSupported = true;
}
+
+ else if (props.major != 9 || props.minor != 0) {
+ std::cerr << "This example requires a GPU of NVIDIA's Hopper Architecture (compute capability 90).\n";
+ notSupported = true;
+ }
+
+
if (notSupported) {
return EXIT_SUCCESS; // Do not fail CI checks on unsupported systems
}
diff --git a/examples/53_hopper_gemm_permute/53_hopper_gemm_permute.cu b/examples/53_hopper_gemm_permute/53_hopper_gemm_permute.cu
index d8096d9e..66ab9b24 100644
--- a/examples/53_hopper_gemm_permute/53_hopper_gemm_permute.cu
+++ b/examples/53_hopper_gemm_permute/53_hopper_gemm_permute.cu
@@ -750,6 +750,13 @@ int main(int argc, char const **argv)
std::cerr << "This example requires a device with compute capability 90 or higher.\n";
notSupported = true;
}
+
+ else if (props.major != 9 || props.minor != 0) {
+ std::cerr << "This example requires a GPU of NVIDIA's Hopper Architecture (compute capability 90).\n";
+ notSupported = true;
+ }
+
+
if (notSupported) {
return EXIT_SUCCESS; // Do not fail CI checks on unsupported systems
}
diff --git a/examples/54_hopper_fp8_warp_specialized_gemm/54_hopper_fp8_warp_specialized_gemm.cu b/examples/54_hopper_fp8_warp_specialized_gemm/54_hopper_fp8_warp_specialized_gemm.cu
index a6f33d1c..f250e4b9 100644
--- a/examples/54_hopper_fp8_warp_specialized_gemm/54_hopper_fp8_warp_specialized_gemm.cu
+++ b/examples/54_hopper_fp8_warp_specialized_gemm/54_hopper_fp8_warp_specialized_gemm.cu
@@ -572,6 +572,13 @@ int main(int argc, char const **args) {
<< "later (compute capability 90 or greater).\n";
return 0;
}
+
+ else if (props.major != 9 || props.minor != 0) {
+ std::cerr << "This example requires a GPU of NVIDIA's Hopper Architecture (compute capability 90).\n";
+ return 0;
+ }
+
+
//
// Parse options
//
diff --git a/examples/55_hopper_mixed_dtype_gemm/55_hopper_int4_bf16_gemm.cu b/examples/55_hopper_mixed_dtype_gemm/55_hopper_int4_bf16_gemm.cu
index 8bca0a35..bdb59bfd 100644
--- a/examples/55_hopper_mixed_dtype_gemm/55_hopper_int4_bf16_gemm.cu
+++ b/examples/55_hopper_mixed_dtype_gemm/55_hopper_int4_bf16_gemm.cu
@@ -619,6 +619,13 @@ int main(int argc, char const **args) {
<< "later (compute capability 90 or greater).\n";
return 0;
}
+
+ else if (props.major != 9 || props.minor != 0) {
+ std::cerr << "This example requires a GPU of NVIDIA's Hopper Architecture (compute capability 90).\n";
+ return 0;
+ }
+
+
//
// Parse options
//
diff --git a/examples/55_hopper_mixed_dtype_gemm/55_hopper_int4_fp8_gemm.cu b/examples/55_hopper_mixed_dtype_gemm/55_hopper_int4_fp8_gemm.cu
index 7bc65f9b..581ccf88 100644
--- a/examples/55_hopper_mixed_dtype_gemm/55_hopper_int4_fp8_gemm.cu
+++ b/examples/55_hopper_mixed_dtype_gemm/55_hopper_int4_fp8_gemm.cu
@@ -524,6 +524,13 @@ int main(int argc, char const **args) {
<< "later (compute capability 90 or greater).\n";
return 0;
}
+
+ else if (props.major != 9 || props.minor != 0) {
+ std::cerr << "This example requires a GPU of NVIDIA's Hopper Architecture (compute capability 90).\n";
+ return 0;
+ }
+
+
//
// Parse options
//
diff --git a/examples/55_hopper_mixed_dtype_gemm/55_hopper_mixed_dtype_gemm.cu b/examples/55_hopper_mixed_dtype_gemm/55_hopper_mixed_dtype_gemm.cu
index cf64f37b..2629ee33 100644
--- a/examples/55_hopper_mixed_dtype_gemm/55_hopper_mixed_dtype_gemm.cu
+++ b/examples/55_hopper_mixed_dtype_gemm/55_hopper_mixed_dtype_gemm.cu
@@ -489,6 +489,13 @@ int main(int argc, char const **args) {
<< "later (compute capability 90 or greater).\n";
return 0;
}
+
+ else if (props.major != 9 || props.minor != 0) {
+ std::cerr << "This example requires a GPU of NVIDIA's Hopper Architecture (compute capability 90).\n";
+ return 0;
+ }
+
+
//
// Parse options
//
diff --git a/examples/55_hopper_mixed_dtype_gemm/README.md b/examples/55_hopper_mixed_dtype_gemm/README.md
index 48eca35c..ca64c901 100644
--- a/examples/55_hopper_mixed_dtype_gemm/README.md
+++ b/examples/55_hopper_mixed_dtype_gemm/README.md
@@ -7,14 +7,18 @@ When relying on `KernelScheduleAuto`, the main loop supporting different A and B
This first version only supports mixed type GEMMs using TMA.
+
## Performance
-While the example offers a harness for straightforward benchmarking, this initial implementation isn't optimized for performance in the majority of scenarios. We expect this implementation to be performant for `{fp16, bf16} x {int8, int4}` and `{fp8} x {int4}` for problems that are compute bound. Additionally, we expect good performance for `fp16, bf16` or `fp32` scales and zero-points. For best performance, it is ideal to have the scales and zero-points be the same type.
+While the example offers a harness for straightforward benchmarking, this initial implementation isn't optimized for performance in the majority of scenarios. We expect this implementation to be performant for `{fp16, bf16} x {int8, int4, int2}` and `{fp8} x {int4}` for problems that are compute bound. Additionally, we expect good performance for `fp16`, `bf16` or `fp32` scales and zero-points. For best performance, it is ideal to have the scales and zero-points be the same type as mma's type.
The scale only mode for `fp8 x int4` is significantly slower than direct conversion mode. There is a lookup-table workaround targeting this mode, as shown in `55_hopper_int4_fp8_gemm.cu`. To use this feature, use `cutlass::Array` as the scale type in the collective builder. However, it requires modifications to the encoding of quantized weights and scale factors. Also, scale with zero point mode is not supported for now.
+Additionally, it's recommended to reorder the narrow data type tensor such that elements read into register file by the same thread are contiguous in global and shared memory. The user can use the helper function `compute_memory_reordering_atom` and `reorder_tensor` to achieve this. See `55_hopper_int4_fp8_gemm.cu` and `55_hopper_int4_bf16_gemm.cu` for more details.
+
We are currently optimizing the following cases:
1. Memory bound cases for all types
+2. `fp8 x {int2, uint2}` case
## Limitations
diff --git a/examples/55_hopper_mixed_dtype_gemm/mixed_dtype_utils.hpp b/examples/55_hopper_mixed_dtype_gemm/mixed_dtype_utils.hpp
index 56382d83..724adabc 100644
--- a/examples/55_hopper_mixed_dtype_gemm/mixed_dtype_utils.hpp
+++ b/examples/55_hopper_mixed_dtype_gemm/mixed_dtype_utils.hpp
@@ -151,16 +151,16 @@ void mixed_dtype_profiling(
runtimes.reserve(options.iterations);
for (int iter = 0; iter < options.warmup + options.iterations; ++iter) {
- cudaEventRecord(start);
- CUTLASS_CHECK(gemm.run());
- cudaEventRecord(stop);
- cudaEventSynchronize(stop);
+ cudaEventRecord(start);
+ CUTLASS_CHECK(gemm.run());
+ cudaEventRecord(stop);
+ cudaEventSynchronize(stop);
- if (iter >= options.warmup) {
- float milliseconds = 0;
- cudaEventElapsedTime(&milliseconds, start, stop);
- runtimes.push_back(milliseconds);
- }
+ if (iter >= options.warmup) {
+ float milliseconds = 0;
+ cudaEventElapsedTime(&milliseconds, start, stop);
+ runtimes.push_back(milliseconds);
+ }
}
cudaEventDestroy(start);
diff --git a/examples/55_hopper_mixed_dtype_gemm/packed_scale.hpp b/examples/55_hopper_mixed_dtype_gemm/packed_scale.hpp
index 98d6df55..a595ca72 100644
--- a/examples/55_hopper_mixed_dtype_gemm/packed_scale.hpp
+++ b/examples/55_hopper_mixed_dtype_gemm/packed_scale.hpp
@@ -33,6 +33,8 @@
#include
+#include "cutlass/util/device_memory.h"
+#include "cutlass/integer_subbyte.h"
#include "cutlass/float8.h"
#include "cutlass/util/reference/device/tensor_fill.h"
@@ -197,7 +199,6 @@ bool initialize_packed_scale(
{
cutlass::packed_scale_t tmp(data_in[i]);
data_out[i] = reinterpret_cast const&>(tmp);
- // std::cout << data_in[i] << ":" << std::hex << static_cast(data_in[i].storage) << ",\t" << -data_in[i] << ":" << std::hex << static_cast((-data_in[i]).storage) << std::endl;
}
try {
block_out.copy_from_host(data_out.data());
diff --git a/examples/56_hopper_ptr_array_batched_gemm/56_hopper_ptr_array_batched_gemm.cu b/examples/56_hopper_ptr_array_batched_gemm/56_hopper_ptr_array_batched_gemm.cu
index 886d39a2..ec29bc05 100644
--- a/examples/56_hopper_ptr_array_batched_gemm/56_hopper_ptr_array_batched_gemm.cu
+++ b/examples/56_hopper_ptr_array_batched_gemm/56_hopper_ptr_array_batched_gemm.cu
@@ -519,6 +519,13 @@ int main(int argc, char const **args) {
<< "later (compute capability 90 or greater).\n";
return 0;
}
+
+ else if (props.major != 9 || props.minor != 0) {
+ std::cerr << "This example requires a GPU of NVIDIA's Hopper Architecture (compute capability 90).\n";
+ return 0;
+ }
+
+
//
// Parse options
//
diff --git a/examples/57_hopper_grouped_gemm/57_hopper_grouped_gemm.cu b/examples/57_hopper_grouped_gemm/57_hopper_grouped_gemm.cu
index 3aeafb4d..0f014cc7 100644
--- a/examples/57_hopper_grouped_gemm/57_hopper_grouped_gemm.cu
+++ b/examples/57_hopper_grouped_gemm/57_hopper_grouped_gemm.cu
@@ -737,6 +737,13 @@ int main(int argc, char const **args) {
<< "later (compute capability 90 or greater).\n";
return 0;
}
+
+ else if (props.major != 9 || props.minor != 0) {
+ std::cerr << "This example requires a GPU of NVIDIA's Hopper Architecture (compute capability 90).\n";
+ return 0;
+ }
+
+
//
// Parse options
//
diff --git a/examples/61_hopper_gemm_with_topk_and_softmax/61_hopper_gemm_with_topk_and_softmax.cu b/examples/61_hopper_gemm_with_topk_and_softmax/61_hopper_gemm_with_topk_and_softmax.cu
index a71a63eb..ac21697e 100644
--- a/examples/61_hopper_gemm_with_topk_and_softmax/61_hopper_gemm_with_topk_and_softmax.cu
+++ b/examples/61_hopper_gemm_with_topk_and_softmax/61_hopper_gemm_with_topk_and_softmax.cu
@@ -507,6 +507,13 @@ int main(int argc, char const **args) {
<< "later (compute capability 90 or greater).\n";
return 0;
}
+
+ else if (props.major != 9 || props.minor != 0) {
+ std::cerr << "This example requires a GPU of NVIDIA's Hopper Architecture (compute capability 90).\n";
+ return 0;
+ }
+
+
//
// Parse options
//
diff --git a/examples/62_hopper_sparse_gemm/62_hopper_sparse_gemm.cu b/examples/62_hopper_sparse_gemm/62_hopper_sparse_gemm.cu
index 01a10046..708d1db6 100644
--- a/examples/62_hopper_sparse_gemm/62_hopper_sparse_gemm.cu
+++ b/examples/62_hopper_sparse_gemm/62_hopper_sparse_gemm.cu
@@ -576,6 +576,14 @@ int main(int argc, char const **args) {
<< "later (compute capability 90 or greater).\n";
return 0;
}
+
+ if (props.major != 9 || props.minor != 0) {
+ std::cerr
+ << "This example requires a GPU of NVIDIA's Hopper Architecture (compute capability 90).\n";
+ return 0;
+ }
+
+
//
// Parse options
//
diff --git a/examples/63_hopper_gemm_with_weight_prefetch/63_hopper_gemm_with_weight_prefetch.cu b/examples/63_hopper_gemm_with_weight_prefetch/63_hopper_gemm_with_weight_prefetch.cu
index d1db304b..03b54f3e 100644
--- a/examples/63_hopper_gemm_with_weight_prefetch/63_hopper_gemm_with_weight_prefetch.cu
+++ b/examples/63_hopper_gemm_with_weight_prefetch/63_hopper_gemm_with_weight_prefetch.cu
@@ -475,6 +475,13 @@ int main(int argc, char const **args) {
<< "later (compute capability 90 or greater).\n";
return 0;
}
+
+ else if (props.major != 9 || props.minor != 0) {
+ std::cerr << "This example requires a GPU of NVIDIA's Hopper Architecture (compute capability 90).\n";
+ return 0;
+ }
+
+
//
// Parse options
//
diff --git a/examples/65_distributed_gemm/65_distributed_gemm.cu b/examples/65_distributed_gemm/65_distributed_gemm.cu
index f0b59ca3..2289d62a 100644
--- a/examples/65_distributed_gemm/65_distributed_gemm.cu
+++ b/examples/65_distributed_gemm/65_distributed_gemm.cu
@@ -133,7 +133,8 @@ using namespace cute;
using TP = _8;
static constexpr int TP_ = TP{};
-#if (defined(CUTLASS_ARCH_MMA_SM90_SUPPORTED) && (__CUDACC_VER_MAJOR__ >= 12) && (__CUDACC_VER_MINOR__ >= 4))
+#if defined(CUTLASS_ARCH_MMA_SM90_SUPPORTED) && \
+ (__CUDACC_VER_MAJOR__ > 12 || (__CUDACC_VER_MAJOR__ == 12 && __CUDACC_VER_MINOR__ >= 4))
// Distributed GEMM tiling/sharding schedule
// Choices:
@@ -344,7 +345,8 @@ struct Result {
};
-#if (defined(CUTLASS_ARCH_MMA_SM90_SUPPORTED) && (__CUDACC_VER_MAJOR__ >= 12) && (__CUDACC_VER_MINOR__ >= 4))
+#if defined(CUTLASS_ARCH_MMA_SM90_SUPPORTED) && \
+ (__CUDACC_VER_MAJOR__ > 12 || (__CUDACC_VER_MAJOR__ == 12 && __CUDACC_VER_MINOR__ >= 4))
/////////////////////////////////////////////////////////////////////////////////////////////////
/// GEMM setup and evaluation
diff --git a/examples/70_blackwell_gemm/70_blackwell_fp16_gemm.cu b/examples/70_blackwell_gemm/70_blackwell_fp16_gemm.cu
new file mode 100644
index 00000000..39123cac
--- /dev/null
+++ b/examples/70_blackwell_gemm/70_blackwell_fp16_gemm.cu
@@ -0,0 +1,483 @@
+/***************************************************************************************************
+ * Copyright (c) 2024 - 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: BSD-3-Clause
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ *
+ * 1. Redistributions of source code must retain the above copyright notice, this
+ * list of conditions and the following disclaimer.
+ *
+ * 2. Redistributions in binary form must reproduce the above copyright notice,
+ * this list of conditions and the following disclaimer in the documentation
+ * and/or other materials provided with the distribution.
+ *
+ * 3. Neither the name of the copyright holder nor the names of its
+ * contributors may be used to endorse or promote products derived from
+ * this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+ * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ * DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+ * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+ * DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+ * SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+ * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+ * OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+ * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ *
+ **************************************************************************************************/
+
+/*! \file
+ \brief A FP16 dense GEMM example for the NVIDIA Blackwell SM100 architecture using CUTLASS.
+
+ This example demonstrates minimal set of changes needed to transition from a Hopper CUTLASS 3.x
+ GEMM kernel (see example 48_hopper_warp_specialized_gemm) to a Blackwell 3.x CUTLASS GEMM kernel.
+
+ The Blackwell SM100 CUTLASS kernel uses of the following Blackwell SM100 features:
+
+ 1. New series of Tensor Core MMA Instructions (tcgen05) introduced on the Blackwell architecture (sm100a)
+ which have 2x throughput compared to Hopper Tensor Core MMA instructions (WGMMA).
+
+ Note that Hopper WGMMA Tensor Core MMA instructions are not compatible on Blackwell (See https://docs.nvidia.com/cuda/parallel-thread-execution).
+
+ 2. A new per-SM memory called Tensor Memory (TMEM) introduced on the Blackwell architecture (sm100a).
+ Blackwell SM100 Tensor Core MMA instructions store their accumulation results in TMEM instead of the
+ Register File. (Please refer to CUDA 12.8 docs on https://docs.nvidia.com/cuda/).
+
+ 3. An extended flavor of the warp-specialized kernel design introduced in Hopper enabled by use of TMEM
+ which allows us to decouple the execution of MMA and epilogue into separate warps.
+
+ 4. A new SW controlled dynamic scheduler based on cluster launch control (See https://docs.nvidia.com/cuda/parallel-thread-execution).
+
+ Usage:
+ $ ./examples/70_blackwell_gemm/70_blackwell_fp16_gemm --m=8192 --n=8192 --k=8192
+*/
+
+
+
+#include
+
+#include "cutlass/cutlass.h"
+
+#include "cute/tensor.hpp"
+#include "cutlass/tensor_ref.h"
+#include "cutlass/epilogue/thread/linear_combination.h"
+#include "cutlass/gemm/dispatch_policy.hpp"
+#include "cutlass/gemm/collective/collective_builder.hpp"
+#include "cutlass/epilogue/collective/collective_builder.hpp"
+#include "cutlass/gemm/device/gemm_universal_adapter.h"
+#include "cutlass/gemm/kernel/gemm_universal.hpp"
+#include "cutlass/gemm/kernel/tile_scheduler_params.h"
+
+#include "cutlass/util/command_line.h"
+#include "cutlass/util/distribution.h"
+#include "cutlass/util/host_tensor.h"
+#include "cutlass/util/packed_stride.hpp"
+#include "cutlass/util/tensor_view_io.h"
+#include "cutlass/util/reference/device/gemm.h"
+#include "cutlass/util/reference/device/tensor_compare.h"
+#include "cutlass/util/reference/device/tensor_fill.h"
+
+#include "helper.h"
+
+using namespace cute;
+
+#if defined(CUTLASS_ARCH_MMA_SM100_SUPPORTED)
+
+/////////////////////////////////////////////////////////////////////////////////////////////////
+/// GEMM kernel configurations
+/////////////////////////////////////////////////////////////////////////////////////////////////
+
+// A matrix configuration
+using ElementA = half_t; // Element type for A matrix operand
+using LayoutA = cutlass::layout::RowMajor; // Layout type for A matrix operand
+constexpr int AlignmentA = 128 / cutlass::sizeof_bits::value; // Memory access granularity/alignment of A matrix in units of elements (up to 16 bytes)
+
+// B matrix configuration
+using ElementB = half_t; // Element type for B matrix operand
+using LayoutB = cutlass::layout::ColumnMajor; // Layout type for B matrix operand
+constexpr int AlignmentB = 128 / cutlass::sizeof_bits::value; // Memory access granularity/alignment of B matrix in units of elements (up to 16 bytes)
+
+// C/D matrix configuration
+using ElementC = float; // Element type for C and D matrix operands
+using LayoutC = cutlass::layout::ColumnMajor; // Layout type for C and D matrix operands
+constexpr int AlignmentC = 128 / cutlass::sizeof_bits::value; // Memory access granularity/alignment of C matrix in units of elements (up to 16 bytes)
+
+// Kernel functional config
+using ElementAccumulator = float; // Element type for internal accumulation
+using ArchTag = cutlass::arch::Sm100; // Tag indicating the minimum SM that supports the intended feature
+using OperatorClass = cutlass::arch::OpClassTensorOp; // Operator class tag
+
+// MMA and Cluster Tile Shapes
+// Shape of the tile computed by tcgen05 MMA, could be across 2 SMs if Cluster Shape %2 == 0
+using MmaTileShape_MNK = Shape<_256,_128,_64>;
+// Shape of the threadblocks in a cluster
+using ClusterShape_MNK = Shape<_2,_2,_1>;
+// Shape of the threadblocks participating in a tcgen05 MMA. <1, 1, 1> for cta_group = 1, <2, 1, 1> for cta_group = 2
+using AtomThrShape_MNK = Shape<_2, _1, _1>;
+// Shape of the tile computed by each SM
+using PerSmTileShape_MNK = decltype(shape_div(MmaTileShape_MNK{}, AtomThrShape_MNK{}));
+
+// Build the epilogue
+using CollectiveEpilogue = typename cutlass::epilogue::collective::CollectiveBuilder<
+ ArchTag, OperatorClass,
+ PerSmTileShape_MNK, ClusterShape_MNK,
+ cutlass::epilogue::collective::EpilogueTileAuto,
+ ElementAccumulator, ElementAccumulator,
+ ElementC, LayoutC, AlignmentC,
+ ElementC, LayoutC, AlignmentC,
+ cutlass::epilogue::collective::EpilogueScheduleAuto
+ >::CollectiveOp;
+
+// Build the mainloop
+using CollectiveMainloop = typename cutlass::gemm::collective::CollectiveBuilder<
+ ArchTag, OperatorClass,
+ ElementA, LayoutA, AlignmentA,
+ ElementB, LayoutB, AlignmentB,
+ ElementAccumulator,
+ MmaTileShape_MNK, ClusterShape_MNK,
+ cutlass::gemm::collective::StageCountAutoCarveout(sizeof(typename CollectiveEpilogue::SharedStorage))>,
+ cutlass::gemm::collective::KernelScheduleAuto
+ >::CollectiveOp;
+
+// Compose into a kernel
+using GemmKernel = cutlass::gemm::kernel::GemmUniversal<
+ Shape, // Indicates ProblemShape
+ CollectiveMainloop,
+ CollectiveEpilogue,
+ void>; // Default to ClusterLaunchControl (CLC) based tile scheduler
+
+using Gemm = cutlass::gemm::device::GemmUniversalAdapter;
+
+// Reference device GEMM implementation type
+using DeviceGemmReference = cutlass::reference::device::Gemm<
+ ElementA,
+ LayoutA,
+ ElementB,
+ LayoutB,
+ ElementC,
+ LayoutC,
+ ElementAccumulator,
+ ElementAccumulator>;
+
+using StrideA = typename Gemm::GemmKernel::StrideA;
+using StrideB = typename Gemm::GemmKernel::StrideB;
+using StrideC = typename Gemm::GemmKernel::StrideC;
+using StrideD = typename Gemm::GemmKernel::StrideD;
+
+//
+// Data members
+//
+
+/// Initialization
+StrideA stride_A;
+StrideB stride_B;
+StrideC stride_C;
+StrideD stride_D;
+uint64_t seed;
+
+cutlass::DeviceAllocation block_A;
+cutlass::DeviceAllocation block_B;
+cutlass::DeviceAllocation block_C;
+cutlass::DeviceAllocation block_D;
+cutlass::DeviceAllocation block_ref_D;
+
+#endif // defined(CUTLASS_ARCH_MMA_SM100_SUPPORTED)
+
+/////////////////////////////////////////////////////////////////////////////////////////////////
+/// Testbed utility types
+/////////////////////////////////////////////////////////////////////////////////////////////////
+
+// Command line options parsing
+struct Options {
+
+ bool help;
+
+ float alpha, beta;
+ int iterations;
+ int m, n, k;
+
+ Options():
+ help(false),
+ m(8192), n(8192), k(8192),
+ alpha(1.f), beta(0.f),
+ iterations(10)
+ { }
+
+ // Parses the command line
+ void parse(int argc, char const **args) {
+ cutlass::CommandLine cmd(argc, args);
+
+ if (cmd.check_cmd_line_flag("help")) {
+ help = true;
+ return;
+ }
+
+ cmd.get_cmd_line_argument("m", m);
+ cmd.get_cmd_line_argument("n", n);
+ cmd.get_cmd_line_argument("k", k);
+ cmd.get_cmd_line_argument("alpha", alpha, 1.f);
+ cmd.get_cmd_line_argument("beta", beta, 0.f);
+ cmd.get_cmd_line_argument("iterations", iterations);
+ }
+
+ /// Prints the usage statement.
+ std::ostream & print_usage(std::ostream &out) const {
+
+ out << "70_blackwell_fp16_gemm\n\n"
+ << " Blackwell FP16 GEMM using a Warp Specialized kernel.\n\n"
+ << "Options:\n\n"
+ << " --help If specified, displays this usage statement\n\n"
+ << " --m= Sets the M extent of the GEMM\n"
+ << " --n= Sets the N extent of the GEMM\n"
+ << " --k= Sets the K extent of the GEMM\n"
+ << " --alpha= Epilogue scalar alpha\n"
+ << " --beta= Epilogue scalar beta\n\n"
+ << " --iterations= Number of profiling iterations to perform.\n\n";
+
+ out
+ << "\n\nExamples:\n\n"
+ << "$ " << "70_blackwell_fp16_gemm" << " --m=1024 --n=512 --k=1024 --alpha=2 --beta=0.707 \n\n";
+
+ return out;
+ }
+
+ /// Compute performance in GFLOP/s
+ double gflops(double runtime_s) const
+ {
+ // Two flops per multiply-add
+ uint64_t flop = uint64_t(2) * m * n * k;
+ double gflop = double(flop) / double(1.0e9);
+ return gflop / runtime_s;
+ }
+};
+
+/// Result structure
+struct Result
+{
+ double avg_runtime_ms;
+ double gflops;
+ cutlass::Status status;
+ cudaError_t error;
+ bool passed;
+
+ Result(
+ double avg_runtime_ms = 0,
+ double gflops = 0,
+ cutlass::Status status = cutlass::Status::kSuccess,
+ cudaError_t error = cudaSuccess)
+ :
+ avg_runtime_ms(avg_runtime_ms), gflops(gflops), status(status), error(error), passed(false)
+ {}
+
+};
+
+#if defined(CUTLASS_ARCH_MMA_SM100_SUPPORTED)
+
+/////////////////////////////////////////////////////////////////////////////////////////////////
+/// GEMM setup and evaluation
+/////////////////////////////////////////////////////////////////////////////////////////////////
+
+/// Helper to initialize a block of device data
+template
+bool initialize_block(
+ cutlass::DeviceAllocation& block,
+ uint64_t seed=2023) {
+
+ Element scope_max, scope_min;
+ int bits_input = cutlass::sizeof_bits::value;
+
+ if (bits_input == 1) {
+ scope_max = Element(2);
+ scope_min = Element(0);
+ } else if (bits_input <= 8) {
+ scope_max = Element(2);
+ scope_min = Element(-2);
+ } else {
+ scope_max = Element(8);
+ scope_min = Element(-8);
+ }
+
+ cutlass::reference::device::BlockFillRandomUniform(
+ block.get(), block.size(), seed, scope_max, scope_min, 0);
+
+ return true;
+}
+
+/// Initialize operands to be used in the GEMM and reference GEMM
+void initialize(const Options &options) {
+
+ stride_A = cutlass::make_cute_packed_stride(StrideA{}, {options.m, options.k, 1});
+ stride_B = cutlass::make_cute_packed_stride(StrideB{}, {options.n, options.k, 1});
+ stride_C = cutlass::make_cute_packed_stride(StrideC{}, {options.m, options.n, 1});
+ stride_D = cutlass::make_cute_packed_stride(StrideD{}, {options.m, options.n, 1});
+
+ block_A.reset(options.m * options.k);
+ block_B.reset(options.k * options.n);
+ block_C.reset(options.m * options.n);
+ block_D.reset(options.m * options.n);
+ block_ref_D.reset(options.m * options.n);
+
+ initialize_block(block_A, seed + 2023);
+ initialize_block(block_B, seed + 2022);
+ initialize_block(block_C, seed + 2021);
+}
+
+/// Populates a Gemm::Arguments structure from the given commandline options
+typename Gemm::Arguments args_from_options(const Options &options)
+{
+ typename Gemm::Arguments arguments{
+ cutlass::gemm::GemmUniversalMode::kGemm,
+ {options.m, options.n, options.k, 1},
+ {block_A.get(), stride_A, block_B.get(), stride_B},
+ {{options.alpha, options.beta}, block_C.get(), stride_C, block_D.get(), stride_D}
+ };
+
+ return arguments;
+}
+
+bool verify(const Options &options) {
+ cutlass::TensorRef ref_A(block_A.get(), Gemm::LayoutA::packed({options.m, options.k}));
+ cutlass::TensorRef ref_B(block_B.get(), Gemm::LayoutB::packed({options.k, options.n}));
+ cutlass::TensorRef ref_C(block_C.get(), Gemm::LayoutC::packed({options.m, options.n}));
+ cutlass::TensorRef ref_D(block_ref_D.get(), Gemm::LayoutD::packed({options.m, options.n}));
+
+ //
+ // Compute reference output
+ //
+
+ // Create instantiation for device reference gemm kernel
+ DeviceGemmReference gemm_reference;
+
+ // Launch device reference gemm kernel
+ gemm_reference(
+ {options.m, options.n, options.k},
+ ElementAccumulator(options.alpha),
+ ref_A,
+ ref_B,
+ ElementAccumulator(options.beta),
+ ref_C,
+ ref_D);
+
+ // Wait for kernel to finish
+ CUDA_CHECK(cudaDeviceSynchronize());
+
+ // Check if output from CUTLASS kernel and reference kernel are equal or not
+ bool passed = cutlass::reference::device::BlockCompareEqual(block_ref_D.get(), block_D.get(), block_D.size());
+
+ return passed;
+}
+
+/// Execute a given example GEMM computation
+template
+int run(Options &options)
+{
+ initialize(options);
+
+ // Instantiate CUTLASS kernel depending on templates
+ Gemm gemm;
+
+ // Create a structure of gemm kernel arguments suitable for invoking an instance of Gemm
+ auto arguments = args_from_options(options);
+
+ // Using the arguments, query for extra workspace required for matrix multiplication computation
+ size_t workspace_size = Gemm::get_workspace_size(arguments);
+
+ // Allocate workspace memory
+ cutlass::device_memory::allocation workspace(workspace_size);
+
+ // Check if the problem size is supported or not
+ CUTLASS_CHECK(gemm.can_implement(arguments));
+
+ // Initialize CUTLASS kernel with arguments and workspace pointer
+ CUTLASS_CHECK(gemm.initialize(arguments, workspace.get()));
+
+ // Correctness / Warmup iteration
+ CUTLASS_CHECK(gemm.run());
+
+ // Check if output from CUTLASS kernel and reference kernel are equal or not
+ Result result;
+ result.passed = verify(options);
+
+ std::cout << " Disposition: " << (result.passed ? "Passed" : "Failed") << std::endl;
+
+ if (!result.passed) {
+ exit(-1);
+ }
+
+ // Run profiling loop
+ if (options.iterations > 0)
+ {
+ GpuTimer timer;
+ timer.start();
+ for (int iter = 0; iter < options.iterations; ++iter) {
+ CUTLASS_CHECK(gemm.initialize(arguments, workspace.get()));
+ CUTLASS_CHECK(gemm.run());
+ }
+ timer.stop();
+
+ // Compute average runtime and GFLOPs.
+ float elapsed_ms = timer.elapsed_millis();
+ result.avg_runtime_ms = double(elapsed_ms) / double(options.iterations);
+ result.gflops = options.gflops(result.avg_runtime_ms / 1000.0);
+
+
+ std::cout << " Problem Size: " << options.m << 'x' << options.n << 'x' << options.k << std::endl;
+ std::cout << " Avg runtime: " << result.avg_runtime_ms << " ms" << std::endl;
+ std::cout << " GFLOPS: " << result.gflops << std::endl;
+ }
+
+ return 0;
+}
+
+#endif // defined(CUTLASS_ARCH_MMA_SM100_SUPPORTED)
+
+///////////////////////////////////////////////////////////////////////////////////////////////////
+
+int main(int argc, char const **args) {
+
+ // CUTLASS must be compiled with CUDA 12.0 Toolkit to run this example
+ // and must have compute capability at least 100a.
+
+ if (__CUDACC_VER_MAJOR__ < 12 || (__CUDACC_VER_MAJOR__ == 12 && __CUDACC_VER_MINOR__ < 8)) {
+ std::cerr << "This example requires CUDA 12.8 or newer." << std::endl;
+ // Returning zero so this test passes on older Toolkits. Its actions are no-op.
+ return 0;
+ }
+
+ cudaDeviceProp props;
+ int current_device_id;
+ CUDA_CHECK(cudaGetDevice(¤t_device_id));
+ CUDA_CHECK(cudaGetDeviceProperties(&props, current_device_id));
+ cudaError_t error = cudaGetDeviceProperties(&props, 0);
+ if (props.major != 10 || props.minor != 0) {
+ std::cerr << "This example requires a GPU with compute capability 100a)." << std::endl;
+ return 0;
+ }
+
+ //
+ // Parse options
+ //
+
+ Options options;
+
+ options.parse(argc, args);
+
+ if (options.help) {
+ options.print_usage(std::cout) << std::endl;
+ return 0;
+ }
+
+ //
+ // Evaluate CUTLASS kernels
+ //
+#if defined(CUTLASS_ARCH_MMA_SM100_SUPPORTED)
+ run(options);
+#endif // defined(CUTLASS_ARCH_MMA_SM100_SUPPORTED)
+
+ return 0;
+}
+
+/////////////////////////////////////////////////////////////////////////////////////////////////
diff --git a/examples/70_blackwell_gemm/70_blackwell_fp8_gemm.cu b/examples/70_blackwell_gemm/70_blackwell_fp8_gemm.cu
new file mode 100644
index 00000000..0b1758b9
--- /dev/null
+++ b/examples/70_blackwell_gemm/70_blackwell_fp8_gemm.cu
@@ -0,0 +1,671 @@
+/***************************************************************************************************
+ * Copyright (c) 2024 - 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: BSD-3-Clause
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ *
+ * 1. Redistributions of source code must retain the above copyright notice, this
+ * list of conditions and the following disclaimer.
+ *
+ * 2. Redistributions in binary form must reproduce the above copyright notice,
+ * this list of conditions and the following disclaimer in the documentation
+ * and/or other materials provided with the distribution.
+ *
+ * 3. Neither the name of the copyright holder nor the names of its
+ * contributors may be used to endorse or promote products derived from
+ * this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+ * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ * DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+ * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+ * DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+ * SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+ * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+ * OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+ * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ *
+ **************************************************************************************************/
+
+/*! \file
+ \brief A FP8 dense GEMM example for the NVIDIA Blackwell SM100 architecture using CUTLASS.
+
+ This example demonstrates minimal set of changes needed to transition from a Hopper CUTLASS 3.x
+ FP8 GEMM kernel (see example 54_hopper_fp8_warp_specialized_gemm) to a Blackwell SM100 FP8 GEMM kernel.
+
+ This example shows all important fusions used by FP8 gemm kernels,
+ i.e., scale factor for A, B, C, D tensor, the abs_max value of D tensor.
+
+ The Blackwell SM100 CUTLASS kernel uses of the following Blackwell SM100 features:
+
+ 1. New series of Tensor Core MMA Instructions (tcgen05) introduced on the Blackwell architecture (sm100a)
+ which have 2x throughput compared to Hopper Tensor Core MMA instructions (WGMMA).
+
+ Note that Hopper WGMMA Tensor Core MMA instructions are not compatible on Blackwell (See https://docs.nvidia.com/cuda/parallel-thread-execution).
+
+ 2. A new per-SM memory called Tensor Memory (TMEM) introduced on the Blackwell architecture (sm100a).
+ Blackwell SM100 Tensor Core MMA instructions store their accumulation results in TMEM instead of the
+ Register File. (Please refer to CUDA 12.8 docs on https://docs.nvidia.com/cuda/).
+
+ 3. An extended flavor of the warp-specialized kernel design introduced in Hopper enabled by use of TMEM
+ which allows us to decouple the execution of MMA and epilogue into separate warps.
+
+ 4. A new SW controlled dynamic scheduler based on cluster launch control (See https://docs.nvidia.com/cuda/parallel-thread-execution).
+
+ Usage:
+ $ ./examples/70_blackwell_gemm/70_blackwell_fp8_gemm --m=8192 --n=8192 --k=8192
+*/
+
+
+
+#include
+
+#include "cutlass/cutlass.h"
+
+#include "cute/tensor.hpp"
+#include "cutlass/tensor_ref.h"
+#include "cutlass/epilogue/thread/activation.h"
+#include "cutlass/gemm/dispatch_policy.hpp"
+#include "cutlass/gemm/collective/collective_builder.hpp"
+#include "cutlass/epilogue/dispatch_policy.hpp"
+#include "cutlass/epilogue/collective/collective_builder.hpp"
+#include "cutlass/gemm/device/gemm_universal_adapter.h"
+#include "cutlass/gemm/kernel/gemm_universal.hpp"
+#include "cutlass/gemm/kernel/tile_scheduler_params.h"
+
+#include "cutlass/util/command_line.h"
+#include "cutlass/util/distribution.h"
+#include "cutlass/util/host_tensor.h"
+#include "cutlass/util/packed_stride.hpp"
+#include "cutlass/util/tensor_view_io.h"
+#include "cutlass/util/reference/host/tensor_fill.h"
+#include "cutlass/util/reference/host/tensor_copy.h"
+#include "cutlass/util/reference/host/tensor_compare.h"
+#include "cutlass/util/reference/host/tensor_norm.h"
+#include "cutlass/util/reference/host/gett.hpp"
+
+
+#include "helper.h"
+
+using namespace cute;
+
+#if defined(CUTLASS_ARCH_MMA_SM100_SUPPORTED)
+
+
+/////////////////////////////////////////////////////////////////////////////////////////////////
+/// GEMM kernel configurations
+/////////////////////////////////////////////////////////////////////////////////////////////////
+// A matrix configuration
+using ElementA = cutlass::float_e4m3_t; // Element type for A matrix operand
+using LayoutA = cutlass::layout::RowMajor; // Layout type for A matrix operand
+constexpr int AlignmentA = 128 / cutlass::sizeof_bits::value; // Memory access granularity/alignment of A matrix in units of elements (up to 16 bytes)
+
+// B matrix configuration
+using ElementB = cutlass::float_e4m3_t; // Element type for B matrix operand
+using LayoutB = cutlass::layout::ColumnMajor; // Layout type for B matrix operand
+constexpr int AlignmentB = 128 / cutlass::sizeof_bits::value; // Memory access granularity/alignment of A matrix in units of elements (up to 16 bytes)
+
+// C/D matrix configuration
+using ElementC = cutlass::float_e4m3_t; // Element type for C and D matrix operands
+using LayoutC = cutlass::layout::ColumnMajor; // Layout type for C and D matrix operands
+constexpr int AlignmentC = 128 / cutlass::sizeof_bits::value; // Memory access granularity/alignment of A matrix in units of elements (up to 16 bytes)
+
+using ElementD = ElementC;
+using LayoutD = LayoutC;
+constexpr int AlignmentD = AlignmentC;
+
+// MMA type
+using ElementAccumulator = float;
+
+// Epilogue types
+using ElementBias = cutlass::half_t;
+using ElementCompute = float;
+using ElementAux = ElementC;
+using LayoutAux = LayoutC;
+using ElementAmax = float;
+
+// MMA and Cluster Tile Shapes
+// Shape of the tile computed by tcgen05 MMA, could be across 2 SMs if Cluster Shape %2 == 0
+using MmaTileShape_MNK = Shape<_256,_128,_64>;
+// Shape of the threadblocks in a cluster
+using ClusterShape_MNK = Shape<_2,_2,_1>;
+// Shape of the threadblocks participating in a tcgen05 MMA. <1, 1, 1> for cta_group = 1, <2, 1, 1> for cta_group = 2
+using AtomThrShape_MNK = Shape<_2, _1, _1>;
+// Shape of the tile computed by each SM
+using PerSmTileShape_MNK = decltype(shape_div(MmaTileShape_MNK{}, AtomThrShape_MNK{}));
+
+using FusionOp = cutlass::epilogue::fusion::ScaledLinCombPerRowBiasEltActAmaxAux<
+ LayoutC, cutlass::epilogue::thread::ReLU, ElementD, ElementCompute, ElementAux, ElementAmax, ElementBias>;
+
+using CollectiveEpilogue = typename cutlass::epilogue::collective::CollectiveBuilder<
+ cutlass::arch::Sm100, cutlass::arch::OpClassTensorOp,
+ PerSmTileShape_MNK, ClusterShape_MNK,
+ cutlass::epilogue::collective::EpilogueTileAuto,
+ ElementAccumulator, ElementCompute,
+ ElementC, LayoutC, AlignmentC,
+ ElementD, LayoutC, AlignmentD,
+ cutlass::epilogue::collective::EpilogueScheduleAuto,
+ FusionOp
+ >::CollectiveOp;
+
+using CollectiveMainloop = typename cutlass::gemm::collective::CollectiveBuilder<
+ cutlass::arch::Sm100, cutlass::arch::OpClassTensorOp,
+ ElementA, LayoutA, AlignmentA,
+ ElementB, LayoutB, AlignmentB,
+ ElementAccumulator,
+ MmaTileShape_MNK, ClusterShape_MNK,
+ cutlass::gemm::collective::StageCountAutoCarveout(sizeof(typename CollectiveEpilogue::SharedStorage))>,
+ cutlass::gemm::collective::KernelScheduleAuto
+ >::CollectiveOp;
+
+using GemmKernel = cutlass::gemm::kernel::GemmUniversal<
+ Shape,
+ CollectiveMainloop,
+ CollectiveEpilogue,
+ void>; // Default to ClusterLaunchControl (CLC) based tile scheduler
+
+using Gemm = cutlass::gemm::device::GemmUniversalAdapter;
+
+// Extract information from Gemm kernel.
+using EpilogueOutputOp = typename Gemm::EpilogueOutputOp;
+using ElementScalar = typename EpilogueOutputOp::ElementScalar;
+using ElementAmax = typename EpilogueOutputOp::ElementAmax;
+using ActivationFunctor = typename EpilogueOutputOp::ActivationFn;
+
+using StrideA = typename Gemm::GemmKernel::StrideA;
+using StrideB = typename Gemm::GemmKernel::StrideB;
+using StrideC = typename Gemm::GemmKernel::StrideC;
+using StrideD = typename Gemm::GemmKernel::StrideD;
+using StrideAux = StrideC;
+
+constexpr bool IsDFp8 =
+ cute::is_same_v or
+ cute::is_same_v;
+
+constexpr bool IsAuxFp8 =
+ cute::is_same_v or
+ cute::is_same_v;
+
+/// Initialization
+StrideA stride_A;
+StrideB stride_B;
+StrideC stride_C;
+StrideD stride_D;
+StrideAux stride_aux;
+uint64_t seed;
+
+cutlass::HostTensor tensor_A;
+cutlass::HostTensor tensor_B;
+cutlass::HostTensor tensor_C;
+cutlass::HostTensor tensor_D;
+cutlass::HostTensor tensor_ref_D;
+cutlass::HostTensor tensor_aux;
+cutlass::HostTensor tensor_ref_aux;
+
+using LayoutScalar = cutlass::layout::PackedVectorLayout;
+cutlass::HostTensor scalar_alpha;
+cutlass::HostTensor scalar_beta;
+cutlass::HostTensor scale_A;
+cutlass::HostTensor scale_B;
+cutlass::HostTensor scale_C;
+cutlass::HostTensor scale_D;
+cutlass::HostTensor scale_aux;
+cutlass::HostTensor abs_max_D;
+cutlass::HostTensor reference_abs_max_D;
+cutlass::HostTensor abs_max_aux;
+cutlass::HostTensor reference_abs_max_aux;
+
+#endif // defined(CUTLASS_ARCH_MMA_SM100_SUPPORTED)
+
+/////////////////////////////////////////////////////////////////////////////////////////////////
+/// Testbed utility types
+/////////////////////////////////////////////////////////////////////////////////////////////////
+
+// Command line options parsing
+struct Options {
+
+ bool help = false;
+
+ float alpha = 1.f, beta = 0.f;
+ float scale_a = 1.f, scale_b = 1.f, scale_c = 1.f, scale_d = 1.f, scale_aux = 1.f;
+ bool device_scale = false;
+ bool save_aux = true;
+ bool save_amax = true;
+ int iterations = 1000;
+ int m = 1024, n = 512, k = 1024, l = 1;
+
+ // Parses the command line
+ void parse(int argc, char const **args) {
+ cutlass::CommandLine cmd(argc, args);
+
+ if (cmd.check_cmd_line_flag("help")) {
+ help = true;
+ return;
+ }
+
+ cmd.get_cmd_line_argument("m", m);
+ cmd.get_cmd_line_argument("n", n);
+ cmd.get_cmd_line_argument("k", k);
+ cmd.get_cmd_line_argument("l", l);
+ cmd.get_cmd_line_argument("alpha", alpha, 1.f);
+ cmd.get_cmd_line_argument("beta", beta, 0.f);
+ cmd.get_cmd_line_argument("scale_a", scale_a, 1.f);
+ cmd.get_cmd_line_argument("scale_b", scale_b, 1.f);
+ cmd.get_cmd_line_argument("scale_c", scale_c, 1.f);
+ cmd.get_cmd_line_argument("scale_d", scale_d, 1.f);
+ cmd.get_cmd_line_argument("scale_aux", scale_aux, 1.f);
+ cmd.get_cmd_line_argument("device_scale", device_scale, false);
+ cmd.get_cmd_line_argument("save_aux", save_aux, true);
+ cmd.get_cmd_line_argument("save_amax", save_amax, true);
+ cmd.get_cmd_line_argument("iterations", iterations);
+ }
+
+ /// Prints the usage statement.
+ std::ostream & print_usage(std::ostream &out) const {
+
+ out << "70_blackwell_fp8_gemm\n\n"
+ << " Blackwell FP8 GEMM using a Warp Specialized kernel.\n\n"
+ << "Options:\n\n"
+ << " --help If specified, displays this usage statement\n\n"
+ << " --m= Sets the M extent of the GEMM\n"
+ << " --n= Sets the N extent of the GEMM\n"
+ << " --k= Sets the K extent of the GEMM\n"
+ << " --l= Sets the l extent (batch) of the GEMM\n"
+ << " --alpha= Epilogue scalar alpha\n"
+ << " --beta= Epilogue scalar beta\n"
+ << " --scale_a= Scaling factor for A\n"
+ << " --scale_b= Scaling factor for B\n"
+ << " --scale_c= Scaling factor for C\n"
+ << " --scale_d= Scaling factor for D (ignored for non-fp8 D)\n"
+ << " --scale_aux= Scaling factor for the auxiliary tensor (ignored for non-fp8 aux)\n"
+ << " --device_scale= Copy scalars to device memory before kernel launch (default: false)\n"
+ << " --save_aux= Save the pre-activation as an auxiliary tensor (default: true)\n"
+ << " --save_amax= Save the pre-scaled max absolute value of any fp8 outputs (aux and/or D) (default: true)\n"
+ << " --iterations= Number of profiling iterations to perform.\n\n";
+
+ out
+ << "\n\nExamples:\n\n"
+ << "$ " << "70_blackwell_fp8_gemm" << " --m=1024 --n=512 --k=1024 --alpha=2 --beta=0.707 \n\n";
+
+ return out;
+ }
+
+ /// Compute performance in GFLOP/s
+ double gflops(double runtime_s) const
+ {
+ // Two flops per multiply-add
+ uint64_t flop = uint64_t(2) * m * n * k;
+ double gflop = double(flop) / double(1.0e9);
+ return gflop / runtime_s;
+ }
+};
+
+/// Result structure
+struct Result
+{
+ double avg_runtime_ms;
+ double gflops;
+ cutlass::Status status;
+ cudaError_t error;
+ bool passed;
+
+ Result(
+ double avg_runtime_ms = 0,
+ double gflops = 0,
+ cutlass::Status status = cutlass::Status::kSuccess,
+ cudaError_t error = cudaSuccess)
+ :
+ avg_runtime_ms(avg_runtime_ms), gflops(gflops), status(status), error(error), passed(false)
+ {}
+
+};
+
+#if defined(CUTLASS_ARCH_MMA_SM100_SUPPORTED)
+
+/////////////////////////////////////////////////////////////////////////////////////////////////
+/// GEMM setup and evaluation
+/////////////////////////////////////////////////////////////////////////////////////////////////
+
+/// Helper to initialize a block of device data
+template
+bool initialize_tensor(
+ cutlass::TensorView view,
+ uint64_t seed) {
+
+ double scope_max, scope_min;
+ int bits_input = cutlass::sizeof_bits::value;
+ int bits_output = cutlass::sizeof_bits::value;
+
+ if (bits_input == 1) {
+ scope_max = 2;
+ scope_min = 0;
+ }
+ else if (bits_input <= 8) {
+ scope_max = 2;
+ scope_min = -2;
+ }
+ else if (bits_output == 16) {
+ scope_max = 5;
+ scope_min = -5;
+ }
+ else {
+ scope_max = 8;
+ scope_min = -8;
+ }
+ cutlass::reference::host::TensorFillRandomUniform(
+ view, seed, scope_max, scope_min, 0);
+
+ return true;
+}
+
+/// Initialize operands to be used in the GEMM and reference GEMM
+void initialize(const Options &options) {
+
+ stride_A = cutlass::make_cute_packed_stride(StrideA{}, cute::make_shape(options.m, options.k, options.l));
+ stride_B = cutlass::make_cute_packed_stride(StrideB{}, cute::make_shape(options.n, options.k, options.l));
+ stride_C = cutlass::make_cute_packed_stride(StrideC{}, cute::make_shape(options.m, options.n, options.l));
+ stride_D = cutlass::make_cute_packed_stride(StrideD{}, cute::make_shape(options.m, options.n, options.l));
+ stride_aux = stride_D;
+
+ auto a_coord = cutlass::make_Coord(options.m * options.l, options.k);
+ auto c_coord = cutlass::make_Coord(options.m * options.l, options.n);
+ auto b_coord = cutlass::make_Coord(options.k, options.n * options.l);
+
+ tensor_A.resize(a_coord);
+ tensor_B.resize(b_coord);
+ tensor_C.resize(c_coord);
+ tensor_D.resize(c_coord);
+ tensor_ref_D.resize(c_coord);
+
+ initialize_tensor(tensor_A.host_view(), seed + 2022);
+ initialize_tensor(tensor_B.host_view(), seed + 2023);
+ initialize_tensor(tensor_C.host_view(), seed + 2024);
+
+ tensor_A.sync_device();
+ tensor_B.sync_device();
+ tensor_C.sync_device();
+ tensor_D.sync_device();
+
+ if (options.save_aux) {
+ tensor_aux.resize(c_coord);
+ tensor_aux.sync_device();
+ tensor_ref_aux.resize(c_coord);
+ }
+
+ if (options.device_scale) {
+ scalar_alpha.resize(cutlass::make_Coord(1));
+ scalar_beta.resize(cutlass::make_Coord(1));
+ scale_A.resize(cutlass::make_Coord(1));
+ scale_B.resize(cutlass::make_Coord(1));
+ scale_C.resize(cutlass::make_Coord(1));
+ scale_D.resize(cutlass::make_Coord(1));
+ scale_aux.resize(cutlass::make_Coord(1));
+
+ cutlass::reference::host::TensorFill(scalar_alpha.host_view(), options.alpha);
+ cutlass::reference::host::TensorFill(scalar_beta.host_view(), options.beta);
+ cutlass::reference::host::TensorFill(scale_A.host_view(), options.scale_a);
+ cutlass::reference::host::TensorFill(scale_B.host_view(), options.scale_b);
+ cutlass::reference::host::TensorFill(scale_C.host_view(), options.scale_c);
+ cutlass::reference::host::TensorFill(scale_D.host_view(), options.scale_d);
+ cutlass::reference::host::TensorFill(scale_aux.host_view(), options.scale_aux);
+
+ scalar_alpha.sync_device();
+ scalar_beta.sync_device();
+ scale_A.sync_device();
+ scale_B.sync_device();
+ scale_C.sync_device();
+ scale_D.sync_device();
+ scale_aux.sync_device();
+ }
+
+ if (IsDFp8 && options.save_amax) {
+ abs_max_D.resize(cutlass::make_Coord(1));
+ abs_max_D.sync_device();
+ reference_abs_max_D.resize(cutlass::make_Coord(1));
+ }
+
+ if (IsAuxFp8 && options.save_aux && options.save_amax) {
+ abs_max_aux.resize(cutlass::make_Coord(1));
+ abs_max_aux.sync_device();
+ reference_abs_max_aux.resize(cutlass::make_Coord(1));
+ }
+}
+
+/// Populates a Gemm::Arguments structure from the given commandline options
+typename Gemm::Arguments args_from_options(const Options &options)
+{
+ typename Gemm::Arguments arguments{
+ cutlass::gemm::GemmUniversalMode::kGemm,
+ {options.m, options.n, options.k, options.l},
+ {tensor_A.device_data(), stride_A, tensor_B.device_data(), stride_B},
+ {
+ {}, // epilogue.thread
+ tensor_C.device_data(), stride_C,
+ tensor_D.device_data(), stride_D
+ }
+ };
+
+ auto &fusion_args = arguments.epilogue.thread;
+ fusion_args.alpha = options.alpha;
+ fusion_args.beta = options.beta;
+ fusion_args.alpha_ptr = scalar_alpha.device_data();
+ fusion_args.beta_ptr = scalar_beta.device_data();
+ fusion_args.scale_a = options.scale_a;
+ fusion_args.scale_b = options.scale_b;
+ fusion_args.scale_c = options.scale_c;
+ fusion_args.scale_a_ptr = scale_A.device_data();
+ fusion_args.scale_b_ptr = scale_B.device_data();
+ fusion_args.scale_c_ptr = scale_C.device_data();
+
+ // ignored if tensor types are not fp8
+ fusion_args.scale_d = options.scale_d;
+ fusion_args.scale_aux = options.scale_aux;
+ fusion_args.scale_d_ptr = scale_D.device_data();
+ fusion_args.scale_aux_ptr = scale_aux.device_data();
+
+ // leaving/setting these as nullptr disables the fusion at runtime
+ fusion_args.bias_ptr = nullptr;
+
+ if (options.save_aux) {
+ fusion_args.aux_ptr = tensor_aux.device_data();
+ fusion_args.dAux = stride_aux;
+ if (options.save_amax) {
+ fusion_args.amax_aux_ptr = abs_max_aux.device_data();
+ }
+ }
+
+ if (options.save_amax) {
+ fusion_args.amax_D_ptr = abs_max_D.device_data();
+ }
+
+ return arguments;
+}
+
+bool verify(const Options &options) {
+ //
+ // Compute reference output
+ //
+
+ // Create instantiation for device reference gemm kernel
+ auto A = cute::make_tensor(tensor_A.host_data(),
+ cute::make_layout(cute::make_shape(options.m, options.k, options.l), stride_A));
+ auto B = cute::make_tensor(tensor_B.host_data(),
+ cute::make_layout(cute::make_shape(options.n, options.k, options.l), stride_B));
+ auto C = cute::make_tensor(tensor_C.host_data(),
+ cute::make_layout(cute::make_shape(options.m, options.n, options.l), stride_C));
+ auto D = cute::make_tensor(tensor_ref_D.host_data(),
+ cute::make_layout(cute::make_shape(options.m, options.n, options.l), stride_D));
+ auto Aux = cute::make_tensor(tensor_ref_aux.host_data(),
+ cute::make_layout(cute::make_shape(options.m, options.n, options.l), stride_aux));
+ using unused_t = decltype(D);
+
+ cutlass::reference::host::GettMainloopParams mainloop_params{A, B};
+
+ cutlass::reference::host::GettEpilogueParams<
+ ElementScalar,
+ ElementScalar,
+ ElementAccumulator,
+ ElementCompute,
+ decltype(C),
+ decltype(D),
+ unused_t, // bias
+ decltype(Aux),
+ unused_t, // valpha
+ unused_t, // vbeta
+ ActivationFunctor
+ > epilogue_params;
+
+ epilogue_params.C = C;
+ epilogue_params.D = D;
+ epilogue_params.Aux = Aux;
+ epilogue_params.alpha = options.alpha;
+ epilogue_params.beta = options.beta;
+ epilogue_params.scale_a = options.scale_a;
+ epilogue_params.scale_b = options.scale_b;
+ epilogue_params.scale_c = options.scale_c;
+ epilogue_params.scale_d = options.scale_d;
+ epilogue_params.scale_aux = options.scale_aux;
+ epilogue_params.abs_max_D = reference_abs_max_D.host_data();
+ epilogue_params.abs_max_Aux = reference_abs_max_aux.host_data();
+
+ // get reference result
+ cutlass::reference::host::Gemm3x(mainloop_params, epilogue_params);
+
+ // compare_reference
+ tensor_D.sync_host();
+ bool passed = cutlass::reference::host::TensorEquals(tensor_ref_D.host_view(), tensor_D.host_view());
+
+ if (IsDFp8 && options.save_amax) {
+ abs_max_D.sync_host();
+ passed &= abs_max_D.at(cutlass::make_Coord(0)) == reference_abs_max_D.at(cutlass::make_Coord(0));
+ }
+
+ if (options.save_aux) {
+ tensor_aux.sync_host();
+ passed &= cutlass::reference::host::TensorEquals(tensor_ref_aux.host_view(), tensor_aux.host_view());
+ if (IsAuxFp8 && options.save_amax) {
+ abs_max_aux.sync_host();
+ passed &= abs_max_aux.at(cutlass::make_Coord(0)) == reference_abs_max_aux.at(cutlass::make_Coord(0));
+ }
+ }
+
+ return passed;
+}
+
+/// Execute a given example GEMM computation
+template
+int run(Options &options)
+{
+ initialize(options);
+
+
+ // Instantiate CUTLASS kernel depending on templates
+ Gemm gemm;
+
+ // Create a structure of gemm kernel arguments suitable for invoking an instance of Gemm
+ auto arguments = args_from_options(options);
+
+ // Using the arguments, query for extra workspace required for matrix multiplication computation
+ size_t workspace_size = Gemm::get_workspace_size(arguments);
+
+ // Allocate workspace memory
+ cutlass::device_memory::allocation workspace(workspace_size);
+
+
+ // Check if the problem size is supported or not
+ CUTLASS_CHECK(gemm.can_implement(arguments));
+
+
+ // Initialize CUTLASS kernel with arguments and workspace pointer
+ CUTLASS_CHECK(gemm.initialize(arguments, workspace.get()));
+
+
+ // Correctness / Warmup iteration
+ CUTLASS_CHECK(gemm.run());
+
+
+ // Check if output from CUTLASS kernel and reference kernel are equal or not
+ Result result;
+ result.passed = verify(options);
+
+ std::cout << " Disposition: " << (result.passed ? "Passed" : "Failed") << std::endl;
+
+ if (!result.passed) {
+ exit(-1);
+ }
+
+ // Run profiling loop
+ if (options.iterations > 0)
+ {
+ GpuTimer timer;
+ timer.start();
+ for (int iter = 0; iter < options.iterations; ++iter) {
+ CUTLASS_CHECK(gemm.run());
+ }
+ timer.stop();
+
+ // Compute average runtime and GFLOPs.
+ float elapsed_ms = timer.elapsed_millis();
+ result.avg_runtime_ms = double(elapsed_ms) / double(options.iterations);
+ result.gflops = options.gflops(result.avg_runtime_ms / 1000.0);
+
+ std::cout << " Problem Size: " << options.m << 'x' << options.n << 'x' << options.k << 'x' << options.l << std::endl;
+ std::cout << " Avg runtime: " << result.avg_runtime_ms << " ms" << std::endl;
+ std::cout << " GFLOPS: " << result.gflops << std::endl;
+ }
+
+ return 0;
+}
+
+#endif // defined(CUTLASS_ARCH_MMA_SM100_SUPPORTED)
+
+///////////////////////////////////////////////////////////////////////////////////////////////////
+
+int main(int argc, char const **args) {
+
+ // CUTLASS must be compiled with CUDA 12.0 Toolkit to run this example
+ // and must have compute capability at least sm100a.
+
+ if (__CUDACC_VER_MAJOR__ < 12) {
+ std::cerr << "This example requires CUDA 12 or newer.\n";
+ // Returning zero so this test passes on older Toolkits. Its actions are no-op.
+ return 0;
+ }
+
+ cudaDeviceProp props;
+ int current_device_id;
+ CUDA_CHECK(cudaGetDevice(¤t_device_id));
+ CUDA_CHECK(cudaGetDeviceProperties(&props, current_device_id));
+ cudaError_t error = cudaGetDeviceProperties(&props, 0);
+ if (props.major != 10 || props.minor != 0) {
+ std::cerr << "This example requires a GPU with compute capability 100a)." << std::endl;
+ return 0;
+ }
+
+
+ //
+ // Parse options
+ //
+
+ Options options;
+
+ options.parse(argc, args);
+
+ if (options.help) {
+ options.print_usage(std::cout) << std::endl;
+ return 0;
+ }
+
+ //
+ // Run
+ //
+#if defined(CUTLASS_ARCH_MMA_SM100_SUPPORTED)
+ run(options);
+#endif // defined(CUTLASS_ARCH_MMA_SM100_SUPPORTED)
+
+ return 0;
+}
+
+/////////////////////////////////////////////////////////////////////////////////////////////////
diff --git a/examples/70_blackwell_gemm/CMakeLists.txt b/examples/70_blackwell_gemm/CMakeLists.txt
new file mode 100644
index 00000000..d88a8c56
--- /dev/null
+++ b/examples/70_blackwell_gemm/CMakeLists.txt
@@ -0,0 +1,41 @@
+
+# Copyright (c) 2025 - 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-License-Identifier: BSD-3-Clause
+#
+# Redistribution and use in source and binary forms, with or without
+# modification, are permitted provided that the following conditions are met:
+#
+# 1. Redistributions of source code must retain the above copyright notice, this
+# list of conditions and the following disclaimer.
+#
+# 2. Redistributions in binary form must reproduce the above copyright notice,
+# this list of conditions and the following disclaimer in the documentation
+# and/or other materials provided with the distribution.
+#
+# 3. Neither the name of the copyright holder nor the names of its
+# contributors may be used to endorse or promote products derived from
+# this software without specific prior written permission.
+#
+# THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+# AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+# IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+# DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+# FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+# DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+# SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+# CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+# OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+# OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+
+if(NOT CUTLASS_NVCC_ARCHS STREQUAL "100")
+cutlass_example_add_executable(
+ 70_blackwell_fp16_gemm
+ 70_blackwell_fp16_gemm.cu
+)
+
+cutlass_example_add_executable(
+ 70_blackwell_fp8_gemm
+ 70_blackwell_fp8_gemm.cu
+)
+endif()
diff --git a/examples/71_blackwell_gemm_with_collective_builder/71_blackwell_gemm_with_collective_builder.cu b/examples/71_blackwell_gemm_with_collective_builder/71_blackwell_gemm_with_collective_builder.cu
new file mode 100644
index 00000000..6712d7a9
--- /dev/null
+++ b/examples/71_blackwell_gemm_with_collective_builder/71_blackwell_gemm_with_collective_builder.cu
@@ -0,0 +1,570 @@
+/***************************************************************************************************
+ * Copyright (c) 2025 - 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: BSD-3-Clause
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ *
+ * 1. Redistributions of source code must retain the above copyright notice, this
+ * list of conditions and the following disclaimer.
+ *
+ * 2. Redistributions in binary form must reproduce the above copyright notice,
+ * this list of conditions and the following disclaimer in the documentation
+ * and/or other materials provided with the distribution.
+ *
+ * 3. Neither the name of the copyright holder nor the names of its
+ * contributors may be used to endorse or promote products derived from
+ * this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+ * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ * DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+ * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+ * DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+ * SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+ * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+ * OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+ * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ *
+ **************************************************************************************************/
+
+/*! \file
+ \brief Blackwell SM100 GEMM example demonstrating compatible mainloop+epilogue builder schedules
+ and epilogue visitor tree (EVT) construction
+
+ Example usage:
+ $ ./examples/71_blackwell_gemm_with_collective_builder/71_blackwell_gemm_with_collective_builder \
+ --m=2048 --n=2048 --k=2048 --l=2
+*/
+
+#include
+
+#include "cute/tensor.hpp"
+
+#include "cutlass/cutlass.h"
+#include "cutlass/tensor_ref.h"
+#include "cutlass/epilogue/collective/default_epilogue.hpp"
+#include "cutlass/epilogue/thread/linear_combination.h"
+#include "cutlass/gemm/dispatch_policy.hpp"
+#include "cutlass/gemm/collective/collective_builder.hpp"
+#include "cutlass/epilogue/collective/collective_builder.hpp"
+#include "cutlass/gemm/device/gemm_universal_adapter.h"
+#include "cutlass/gemm/kernel/gemm_universal.hpp"
+#include "cutlass/gemm/kernel/tile_scheduler.hpp"
+
+#include "cutlass/util/command_line.h"
+#include "cutlass/util/distribution.h"
+#include "cutlass/util/host_tensor.h"
+#include "cutlass/util/packed_stride.hpp"
+#include "cutlass/util/tensor_view_io.h"
+#include "cutlass/util/reference/device/gemm_complex.h"
+#include "cutlass/util/reference/device/tensor_compare.h"
+#include "cutlass/util/reference/device/tensor_fill.h"
+
+using namespace cute;
+
+///////////////////////////////////////////////////////////////////////////////////////////////////
+
+/// Command line options parsing
+struct Options {
+
+ bool help;
+ bool error;
+
+ int m, n, k, l;
+ float alpha, beta;
+
+ Options():
+ help(false),
+ error(false),
+ m(2048), n(2048), k(2048), l(1),
+ alpha(1.f), beta(0.f)
+ { }
+
+ // Parses the command line
+ void parse(int argc, char const **args) {
+ cutlass::CommandLine cmd(argc, args);
+
+ if (cmd.check_cmd_line_flag("help")) {
+ help = true;
+ return;
+ }
+
+ cmd.get_cmd_line_argument("m", m, 2048);
+ cmd.get_cmd_line_argument("n", n, 2048);
+ cmd.get_cmd_line_argument("k", k, 2048);
+ cmd.get_cmd_line_argument("l", l, 1);
+ cmd.get_cmd_line_argument("alpha", alpha, 1.f);
+ cmd.get_cmd_line_argument("beta", beta, 0.f);
+ }
+
+ /// Prints the usage statement.
+ std::ostream & print_usage(std::ostream &out) const {
+
+ out << "71_blackwell_gemm_with_collective_builder\n\n"
+ << " This example showcases the use of CUTLASS's collective operation builders to easily construct\n"
+ << " performant kernels targeting NVIDIA's Blackwell architecture.\n\n"
+ << "Options:\n\n"
+ << " --help If specified, displays this usage statement\n\n"
+ << " --m= Sets the M extent of the GEMM\n"
+ << " --n= Sets the N extent of the GEMM\n"
+ << " --k= Sets the K extent of the GEMM\n"
+ << " --l= Sets the L extent (batch count) of the GEMM\n"
+ << " --alpha= Epilogue scalar alpha\n"
+ << " --beta= Epilogue scalar beta\n\n";
+
+ return out;
+ }
+};
+
+///////////////////////////////////////////////////////////////////////////////////////////////////
+
+/// Helper to initialize a block of device data
+template
+bool initialize_block(
+ cutlass::DeviceAllocation& block,
+ uint64_t seed=2023) {
+
+ Element scope_max, scope_min;
+ int bits_input = cutlass::sizeof_bits::value;
+
+ if (bits_input == 1) {
+ scope_max = 2;
+ scope_min = 0;
+ } else if (bits_input <= 8) {
+ scope_max = 2;
+ scope_min = -2;
+ } else {
+ scope_max = 8;
+ scope_min = -8;
+ }
+
+ cutlass::reference::device::BlockFillRandomUniform(
+ block.get(), block.size(), seed, scope_max, scope_min, 0);
+
+ return true;
+}
+
+///////////////////////////////////////////////////////////////////////////////////////////////////
+
+#if defined(CUTLASS_ARCH_MMA_SM100_SUPPORTED)
+
+// Wrapper to construct, run, and verify a GEMM. This example showcases CUTLASS's collective
+// operation builders by specializing the GEMM on the kernel+epilogue schedule it will use and the
+// number of pipeline stages.
+template <
+ // Type of kernel schedule to generate
+ class MainloopScheduleType = cutlass::gemm::collective::KernelScheduleAuto,
+ // Type of epilogue schedule to generate
+ class EpilogueScheduleType = cutlass::epilogue::collective::EpilogueScheduleAuto,
+ // Number of pipeline stages to use
+ class StageCountType = cutlass::gemm::collective::StageCountAuto,
+ // Do we use custom epilogue visitor tree (EVT) fusion
+ bool UseCustomEVT = false
+>
+struct ExampleRunner {
+
+ using LayoutA = cutlass::layout::RowMajor;
+ using LayoutB = cutlass::layout::ColumnMajor;
+ using LayoutC = cutlass::layout::ColumnMajor;
+ using LayoutD = cutlass::layout::ColumnMajor;
+
+ using ElementA = cutlass::half_t;
+ using ElementB = cutlass::half_t;
+ using ElementC = cutlass::half_t;
+ using ElementD = cutlass::half_t;
+ using ElementAccumulator = float;
+ using ElementCompute = float;
+ using ElementScalar = float;
+
+ using ClusterShapeMNK = Shape<_2,_2,_1>;
+ static constexpr bool Use2SmMma =
+ // Manually specified 2sm cluster MMA schedule, will error if cluster M is not a multiple of 2
+ std::is_same_v ||
+ // Auto schedule will try to select 2sm cluster MMA based on cluster M
+ std::is_same_v && size<0>(ClusterShapeMNK{}) % 2 == 0;
+ // The MNK layout of CTAs within a cluster MMA
+ using AtomThrMNK = std::conditional_t, Shape<_1,_1,_1>>;
+ // The MMA tile used by the mainloop collective. Blackwell 1sm MMA supports up to MMA tile M = 128, 2sm MMA supports up to MMA tile M = 256
+ using MmaTileMNK = std::conditional_t, Shape<_128,_128,_64>>;
+ // The Output tile used by the epilogue collective
+ using OutputTileMNK = decltype(shape_div(MmaTileMNK{}, AtomThrMNK{}));
+
+ // 16B alignment lets us use TMA
+ static constexpr int AlignmentA = 128 / cutlass::sizeof_bits::value;
+ static constexpr int AlignmentB = 128 / cutlass::sizeof_bits::value;
+ static constexpr int AlignmentC = 128 / cutlass::sizeof_bits::value;
+ static constexpr int AlignmentD = 128 / cutlass::sizeof_bits::value;
+
+ static constexpr auto RoundStyle = cutlass::FloatRoundStyle::round_to_nearest;
+
+ // Blackwell fusions for the most part use the same EVT nodes used in Hopper. Most Blackwell EVTs will alias to their Hopper counterparts.
+ // EVT nodes new to Blackwell mainly relate to narrow precision scale factor generation and are contained in include/cutlass/epilogue/fusion/sm100_visitor_*.hpp
+ // See include/cutlass/epilogue/fusion/sm100_callbacks_tma_warpspecialized.hpp for EVT construction using these new nodes
+ // Fusions relating to narrow-precision scale factor generation are demonstrated in example 72b and can only be used in blackwell kernels
+ using CustomEVT = // alpha * acc + beta * C
+ cutlass::epilogue::fusion::Sm90EVT, // beta * C + (alpha * acc)
+ cutlass::epilogue::fusion::Sm90ScalarBroadcast, // beta
+ cutlass::epilogue::fusion::Sm90SrcFetch, // C
+ cutlass::epilogue::fusion::Sm90EVT, // alpha * acc
+ cutlass::epilogue::fusion::Sm90ScalarBroadcast, // alpha
+ cutlass::epilogue::fusion::Sm90AccFetch // acc
+ >
+ >;
+
+ // As in Hopper, a predefined set of fusion operations are provided in include/cutlass/epilogue/fusion/operations.hpp and can be passed to the epilogue builder
+ // Fusions operations supported by the Hopper TMA epilogue will also be supported by the Blackwell TMA epilogue
+ // Fusions relating to narrow-precision scale factor generation are demonstrated in example 72b and can only be used in blackwell kernels
+ using DefaultOperation = cutlass::epilogue::fusion::LinearCombination;
+
+ using CollectiveEpilogue = typename cutlass::epilogue::collective::CollectiveBuilder<
+ cutlass::arch::Sm100, cutlass::arch::OpClassTensorOp,
+ OutputTileMNK, ClusterShapeMNK,
+ cutlass::epilogue::collective::EpilogueTileAuto,
+ ElementAccumulator, ElementCompute,
+ ElementC, LayoutC, AlignmentC,
+ ElementD, LayoutD, AlignmentD,
+ EpilogueScheduleType,
+ cute::conditional_t
+ >::CollectiveOp;
+
+ using CollectiveMainloop = typename cutlass::gemm::collective::CollectiveBuilder<
+ cutlass::arch::Sm100, cutlass::arch::OpClassTensorOp,
+ ElementA, LayoutA, AlignmentA,
+ ElementB, LayoutB, AlignmentB,
+ ElementAccumulator,
+ MmaTileMNK, ClusterShapeMNK,
+ cute::conditional_t,
+ cutlass::gemm::collective::StageCountAutoCarveout(sizeof(typename CollectiveEpilogue::SharedStorage))>,
+ StageCountType>,
+ MainloopScheduleType
+ >::CollectiveOp;
+
+ using GemmKernel = cutlass::gemm::kernel::GemmUniversal<
+ Shape,
+ CollectiveMainloop,
+ CollectiveEpilogue
+ >;
+
+ using Gemm = cutlass::gemm::device::GemmUniversalAdapter;
+
+ using ProblemShapeType = typename Gemm::GemmKernel::ProblemShape;
+
+ using StrideA = typename Gemm::GemmKernel::StrideA;
+ using StrideB = typename Gemm::GemmKernel::StrideB;
+ using StrideC = typename Gemm::GemmKernel::StrideC;
+ using StrideD = typename Gemm::GemmKernel::StrideD;
+
+ using LayoutTagA = cutlass::gemm::detail::StrideToLayoutTagA_t;
+ using LayoutTagB = cutlass::gemm::detail::StrideToLayoutTagB_t;
+ using LayoutTagC = cutlass::gemm::detail::StrideToLayoutTagC_t;
+ using LayoutTagD = cutlass::gemm::detail::StrideToLayoutTagC_t;
+
+ //
+ // Data members
+ //
+
+ /// Initialization
+ StrideA stride_A;
+ StrideB stride_B;
+ StrideC stride_C;
+ StrideD stride_D;
+ uint64_t seed = 0;
+
+ cutlass::DeviceAllocation block_A;
+ cutlass::DeviceAllocation block_B;
+ cutlass::DeviceAllocation block_C;
+ cutlass::DeviceAllocation block_D;
+ cutlass::DeviceAllocation block_ref_D;
+
+ //
+ // Methods
+ //
+
+ bool verify(const ProblemShapeType& problem_size, float alpha, float beta) {
+ auto [M, N, K, L] = problem_size;
+
+ cutlass::TensorRef ref_A(block_A.get(), Gemm::LayoutA::packed({M, K}));
+ cutlass::TensorRef ref_B(block_B.get(), Gemm::LayoutB::packed({K, N}));
+ cutlass::TensorRef ref_C(block_C.get(), Gemm::LayoutC::packed({M, N}));
+ cutlass::TensorRef ref_D(block_ref_D.get(), Gemm::LayoutD::packed({M, N}));
+
+ cutlass::reference::device::GemmComplex(
+ {M, N, K},
+ ElementScalar(alpha),
+ ref_A,
+ cutlass::ComplexTransform::kNone,
+ ref_B,
+ cutlass::ComplexTransform::kNone,
+ ElementScalar(beta),
+ ref_C,
+ ref_D,
+ ElementAccumulator(0),
+ L, // batch_count
+ M * K, // batch_stride_A
+ K * N, // batch_stride_B
+ M * N, // batch_stride_C
+ M * N // batch_stride_D
+ );
+
+ cudaError_t result = cudaDeviceSynchronize();
+ if (result != cudaSuccess) {
+ std::cerr << "Reference kernel failed. Last CUDA error: "
+ << cudaGetErrorString(result) << std::endl;
+ return false;
+ }
+
+ // Check if output from CUTLASS kernel and reference kernel are equal or not
+ bool passed = cutlass::reference::device::BlockCompareEqual(block_ref_D.get(), block_D.get(), block_D.size());
+
+ return passed;
+ }
+
+ /// Initialize operands to be used in the GEMM and reference GEMM
+ void initialize(const ProblemShapeType& problem_size) {
+ auto problem_shape_MNKL = cute::append<4>(problem_size, 1);
+ auto [M, N, K, L] = problem_shape_MNKL;
+
+ stride_A = cutlass::make_cute_packed_stride(StrideA{}, cute::make_shape(M, K, L));
+ stride_B = cutlass::make_cute_packed_stride(StrideB{}, cute::make_shape(N, K, L));
+ stride_C = cutlass::make_cute_packed_stride(StrideC{}, cute::make_shape(M, N, L));
+ stride_D = cutlass::make_cute_packed_stride(StrideD{}, cute::make_shape(M, N, L));
+
+ block_A.reset(M * K * L);
+ block_B.reset(K * N * L);
+ block_C.reset(M * N * L);
+ block_D.reset(M * N * L);
+ block_ref_D.reset(M * N * L);
+
+ initialize_block(block_A, seed + 2023);
+ initialize_block(block_B, seed + 2022);
+ initialize_block(block_C, seed + 2021);
+ }
+
+ bool run(const Options& options, const cutlass::KernelHardwareInfo& hw_info) {
+ ProblemShapeType problem_size = ProblemShapeType{options.m, options.n, options.k, options.l};
+
+ initialize(problem_size);
+
+ typename Gemm::Arguments arguments{
+ cutlass::gemm::GemmUniversalMode::kGemm,
+ problem_size,
+ {block_A.get(), stride_A, block_B.get(), stride_B},
+ {{}, // epilogue.thread
+ block_C.get(), stride_C, block_D.get(), stride_D},
+ hw_info
+ };
+
+ // See example 48 for details on custom EVT construction
+ if constexpr (UseCustomEVT) {
+ arguments.epilogue.thread =
+ { // ternary op : beta * C + (alpha * acc)
+ {{options.beta}}, // leaf op+args : beta
+ {}, // leaf op+args : C
+ { // binary op : alpha * acc
+ {{options.alpha}}, // leaf op+args : alpha
+ {}, // leaf op+args : acc
+ {} // binary args : multiplies
+ }, // end binary op
+ {} // ternary args : multiply_add
+ }; // end ternary op
+ }
+ // Pre-defined fusions will have flat, named args for user-friendlyness
+ else {
+ arguments.epilogue.thread.alpha = options.alpha;
+ arguments.epilogue.thread.beta = options.beta;
+ }
+
+ Gemm gemm_op;
+
+ size_t workspace_size = Gemm::get_workspace_size(arguments);
+ cutlass::device_memory::allocation workspace(workspace_size);
+
+ cutlass::Status status = gemm_op.can_implement(arguments);
+ if (status != cutlass::Status::kSuccess) {
+ std::cerr << "This kernel is not supported. Last CUDA error is: "
+ << cudaGetErrorString(cudaGetLastError()) << std::endl;
+ return false;
+ }
+
+ status = gemm_op.initialize(arguments, workspace.get());
+ if (status != cutlass::Status::kSuccess) {
+ std::cerr << "Failed to initialize the CUTLASS kernel. Last CUDA error is: "
+ << cudaGetErrorString(cudaGetLastError()) << std::endl;
+ return false;
+ }
+
+ // Run the GEMM
+ status = gemm_op.run();
+ if (status != cutlass::Status::kSuccess) {
+ std::cerr << "Failed to launch the CUTLASS kernel. Last CUDA error is: "
+ << cudaGetErrorString(cudaGetLastError()) << std::endl;
+ return false;
+ }
+
+ cudaError_t result = cudaDeviceSynchronize();
+ if (result != cudaSuccess) {
+ std::cerr << "Error running the CUTLASS kernel. Last CUDA error is: "
+ << cudaGetErrorString(result) << std::endl;
+ return false;
+ }
+
+ // Verify that the result is correct
+ bool passed = verify(problem_size, options.alpha, options.beta);
+ if (!passed) {
+ std::cerr << "Reference check failed" << std::endl;
+ }
+
+ return passed;
+ }
+
+};
+
+#endif // defined(CUTLASS_ARCH_MMA_SM100_SUPPORTED)
+
+///////////////////////////////////////////////////////////////////////////////////////////////////
+
+/// Helper to print a description of the example run and its result
+void print_result(const std::string& description, bool passed) {
+ std::cout << description << ": " << (passed ? "Passed" : "Failed") << std::endl;
+}
+
+///////////////////////////////////////////////////////////////////////////////////////////////////
+
+int main(int argc, char const **args) {
+
+ cudaDeviceProp props;
+
+ cudaError_t error = cudaGetDeviceProperties(&props, 0);
+ if (error != cudaSuccess) {
+ std::cerr << "cudaGetDeviceProperties() returned an error: " << cudaGetErrorString(error) << std::endl;
+ return -1;
+ }
+
+if (__CUDACC_VER_MAJOR__ < 12 || (__CUDACC_VER_MAJOR__ == 12 && __CUDACC_VER_MINOR__ < 8)) {
+ std::cerr << "This example requires CUDA 12.8 or newer." << std::endl;
+ // Returning zero so this test passes on older Toolkits. Its actions are no-op.
+ return 0;
+ }
+
+ if (!(props.major == 10 && props.minor == 0)) {
+ std::cerr << "This example requires a GPU of NVIDIA's Blackwell architecture (compute capability 100)." << std::endl;
+ return 0;
+ }
+
+ //
+ // Parse options
+ //
+
+ Options options;
+
+ options.parse(argc, args);
+
+ if (options.help) {
+ options.print_usage(std::cout) << std::endl;
+ return 0;
+ }
+
+ if (options.error) {
+ std::cerr << "Aborting execution." << std::endl;
+ return -1;
+ }
+
+#if defined(CUTLASS_ARCH_MMA_SM100_SUPPORTED)
+
+ //
+ // Run examples
+ //
+
+ // The KernelHardwareInfo struct holds the number of SMs on the GPU with a given device ID. This
+ // information is used by the underlying kernel.
+ cutlass::KernelHardwareInfo hw_info;
+
+ // Change device_id to another value if you are running on a machine with multiple GPUs and wish
+ // to use a GPU other than that with device ID 0.
+ hw_info.device_id = 0;
+ hw_info.sm_count = cutlass::KernelHardwareInfo::query_device_multiprocessor_count(hw_info.device_id);
+
+ bool passed;
+
+ // Auto mainloop and epilogue schedules must be used together to guarantee functionality
+ ExampleRunner<> runner_0;
+ passed = runner_0.run(options, hw_info);
+ print_result("KernelScheduleAuto mainloop schedule with EpilogueScheduleAuto epilogue schedule", passed);
+
+ // Mainloop stage counts can be specified manually
+ // It is the user's responsibility to ensure there is enough device smem to allocate manual stage counts
+ ExampleRunner<
+ cutlass::gemm::collective::KernelScheduleAuto,
+ cutlass::epilogue::collective::EpilogueScheduleAuto,
+ _3> runner_1;
+ passed = runner_1.run(options, hw_info);
+ print_result("KernelScheduleAuto mainloop schedule with EpilogueScheduleAuto epilogue schedule and 3 mainloop stages", passed);
+
+ // 1SM cluster MMA mainloop schedules can be used with direct store ("no-smem") epilogue schedules
+ ExampleRunner runner_2;
+ passed = runner_2.run(options, hw_info);
+ print_result("KernelTmaWarpSpecialized1SmSm100 mainloop schedule with NoSmemWarpSpecialized epilogue schedule", passed);
+
+ // 1SM cluster MMA mainloop schedules can also be used with 1SM TMA epilogue schedules
+ // 1SM cluster MMA mainloop schedules will not work with 2SM TMA epilogue schedules
+ ExampleRunner runner_3;
+ passed = runner_3.run(options, hw_info);
+ print_result("KernelTmaWarpSpecialized1SmSm100 mainloop schedule with NoSmemWarpSpecialized epilogue schedule", passed);
+
+ // 2SM cluster MMA mainloop schedules can be used with direct store ("no-smem") epilogue schedules
+ ExampleRunner runner_4;
+ passed = runner_4.run(options, hw_info);
+ print_result("KernelTmaWarpSpecialized2SmSm100 mainloop schedule with NoSmemWarpSpecialized epilogue schedule", passed);
+
+ // 2SM cluster MMA mainloop schedules can also be used with 2SM TMA epilogue schedules
+ // 2SM cluster MMA mainloop schedules will not work with SM TMA epilogue schedules
+ ExampleRunner runner_5;
+ passed = runner_5.run(options, hw_info);
+ print_result("KernelTmaWarpSpecialized2SmSm100 mainloop schedule with TmaWarpSpecialized2Sm epilogue schedule", passed);
+
+ // Blackwell Auto schedule supports custom EVT fusions
+ constexpr bool UseCustomEVT = true;
+ ExampleRunner<
+ cutlass::gemm::collective::KernelScheduleAuto,
+ cutlass::epilogue::collective::EpilogueScheduleAuto,
+ cutlass::gemm::collective::StageCountAuto,
+ UseCustomEVT> runner_6;
+ passed = runner_6.run(options, hw_info);
+ print_result("KernelScheduleAuto mainloop schedule with EpilogueScheduleAuto epilogue schedule and custom EVT", passed);
+
+ // 1SM TMA epilogue schedules support custom EVT fusions
+ ExampleRunner<
+ cutlass::gemm::KernelTmaWarpSpecialized1SmSm100,
+ cutlass::epilogue::TmaWarpSpecialized1Sm,
+ cutlass::gemm::collective::StageCountAuto,
+ UseCustomEVT> runner_7;
+ passed = runner_7.run(options, hw_info);
+ print_result("KernelTmaWarpSpecialized1SmSm100 mainloop schedule with TmaWarpSpecialized1Sm epilogue and custom EVT", passed);
+
+ // 2SM TMA epilogue schedules support custom EVT fusions
+ ExampleRunner<
+ cutlass::gemm::KernelTmaWarpSpecialized2SmSm100,
+ cutlass::epilogue::TmaWarpSpecialized2Sm,
+ cutlass::gemm::collective::StageCountAuto,
+ UseCustomEVT> runner_8;
+ passed = runner_8.run(options, hw_info);
+ print_result("KernelTmaWarpSpecialized2SmSm100 mainloop schedule with TmaWarpSpecialized2Sm epilogue and custom EVT", passed);
+
+
+ // Blackwell direct store epilogue schedule supports custom EVTs and named fusion operations as well (not supported for pre-Blackwell kernels)
+ ExampleRunner<
+ cutlass::gemm::KernelTmaWarpSpecialized1SmSm100,
+ cutlass::epilogue::NoSmemWarpSpecialized,
+ cutlass::gemm::collective::StageCountAuto,
+ UseCustomEVT> runner_9;
+ passed = runner_9.run(options, hw_info);
+ print_result("KernelTmaWarpSpecialized1SmSm100 mainloop schedule with NoSmemWarpSpecialized epilogue and custom EVT", passed);
+
+#endif
+
+ return 0;
+}
+
+/////////////////////////////////////////////////////////////////////////////////////////////////
diff --git a/examples/71_blackwell_gemm_with_collective_builder/CMakeLists.txt b/examples/71_blackwell_gemm_with_collective_builder/CMakeLists.txt
new file mode 100644
index 00000000..5bac6494
--- /dev/null
+++ b/examples/71_blackwell_gemm_with_collective_builder/CMakeLists.txt
@@ -0,0 +1,35 @@
+# Copyright (c) 2025 - 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-License-Identifier: BSD-3-Clause
+#
+# Redistribution and use in source and binary forms, with or without
+# modification, are permitted provided that the following conditions are met:
+#
+# 1. Redistributions of source code must retain the above copyright notice, this
+# list of conditions and the following disclaimer.
+#
+# 2. Redistributions in binary form must reproduce the above copyright notice,
+# this list of conditions and the following disclaimer in the documentation
+# and/or other materials provided with the distribution.
+#
+# 3. Neither the name of the copyright holder nor the names of its
+# contributors may be used to endorse or promote products derived from
+# this software without specific prior written permission.
+#
+# THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+# AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+# IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+# DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+# FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+# DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+# SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+# CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+# OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+# OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+# Both filenames are shorter to avoid MAX_PATH issues on Windows.
+if(NOT CUTLASS_NVCC_ARCHS STREQUAL "100")
+cutlass_example_add_executable(
+ 71_blackwell_gemm_with_collective_builder
+ 71_blackwell_gemm_with_collective_builder.cu
+ )
+endif()
diff --git a/examples/72_blackwell_narrow_precision_gemm/72a_blackwell_nvfp4_bf16_gemm.cu b/examples/72_blackwell_narrow_precision_gemm/72a_blackwell_nvfp4_bf16_gemm.cu
new file mode 100644
index 00000000..ec597966
--- /dev/null
+++ b/examples/72_blackwell_narrow_precision_gemm/72a_blackwell_nvfp4_bf16_gemm.cu
@@ -0,0 +1,544 @@
+/***************************************************************************************************
+ * Copyright (c) 2025 - 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: BSD-3-Clause
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ *
+ * 1. Redistributions of source code must retain the above copyright notice, this
+ * list of conditions and the following disclaimer.
+ *
+ * 2. Redistributions in binary form must reproduce the above copyright notice,
+ * this list of conditions and the following disclaimer in the documentation
+ * and/or other materials provided with the distribution.
+ *
+ * 3. Neither the name of the copyright holder nor the names of its
+ * contributors may be used to endorse or promote products derived from
+ * this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+ * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ * DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+ * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+ * DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+ * SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+ * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+ * OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+ * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ *
+ **************************************************************************************************/
+
+/*! \file
+ \brief A GEMM example using CUTLASS for the NVIDIA Blackwell SM100 architecture.
+
+ This example demonstrates a simple way to instantiate and run a blockscaled NVFP4 GEMM on the NVIDIA Blackwell SM100 architecture.
+
+ The Blackwell SM100 CUTLASS kernel uses the new Block Scaled Tensor Core MMA Instructions (tcgen05.mma.blockscaled) introduced
+ on the Blackwell architecture (sm100a) which have 2x throughput compared to fp8 Tensor Core MMA instructions (tcgen05.mma)
+ and 4x throughput compared to fp8 Hopper Tensor Core MMA Instructions (WGMMA) (See https://docs.nvidia.com/cuda/parallel-thread-execution).
+
+ Similar to 70_blackwell_gemm, this kernel leverages:
+ 1. Per-SM memory called Tensor Memory (TMEM) (Please refer to CUDA 12.8 docs on https://docs.nvidia.com/cuda/).
+
+ 2. The extended warp-specialized kernel design introduced in Hopper enabled by use of TMEM
+ which allows us to decouple the execution of MMA and epilogue into separate warps.
+
+ 3. A new SW controlled dynamic scheduler based on cluster launch control (See https://docs.nvidia.com/cuda/parallel-thread-execution).
+
+ Usage:
+
+ $ ./examples/72_blackwell_narrow_precision_gemm/72a_blackwell_nvfp4_bf16_gemm --m=2048 --n=2048 --k=2048
+*/
+
+#include
+
+#include "cutlass/cutlass.h"
+
+#include "cute/tensor.hpp"
+#include "cutlass/tensor_ref.h"
+#include "cutlass/epilogue/thread/linear_combination.h"
+#include "cutlass/gemm/dispatch_policy.hpp"
+#include "cutlass/gemm/collective/collective_builder.hpp"
+#include "cutlass/epilogue/collective/collective_builder.hpp"
+#include "cutlass/detail/sm100_blockscaled_layout.hpp"
+#include "cutlass/gemm/device/gemm_universal_adapter.h"
+#include "cutlass/gemm/kernel/gemm_universal.hpp"
+#include "cutlass/gemm/kernel/tile_scheduler_params.h"
+
+#include "cutlass/util/command_line.h"
+#include "cutlass/util/distribution.h"
+#include "cutlass/util/host_tensor.h"
+#include "cutlass/util/packed_stride.hpp"
+#include "cutlass/util/tensor_view_io.h"
+#include "cutlass/util/reference/device/gemm.h"
+#include "cutlass/util/reference/device/tensor_compare.h"
+#include "cutlass/util/reference/host/tensor_fill.h"
+#include "cutlass/util/reference/host/gett.hpp"
+#include "cutlass/util/reference/host/tensor_norm.h"
+#include "cutlass/util/reference/host/tensor_compare.h"
+
+
+#include
+
+#include "helper.h"
+
+using namespace cute;
+
+#if defined(CUTLASS_ARCH_MMA_SM100_SUPPORTED)
+
+
+/////////////////////////////////////////////////////////////////////////////////////////////////
+/// GEMM kernel configurations
+/////////////////////////////////////////////////////////////////////////////////////////////////
+
+// A matrix configuration
+using ElementA = cutlass::nv_float4_t; // Element type for A matrix operand
+using LayoutATag = cutlass::layout::RowMajor; // Layout type for A matrix operand
+constexpr int AlignmentA = 32; // Memory access granularity/alignment of A matrix in units of elements (up to 16 bytes)
+
+// B matrix configuration
+using ElementB = cutlass::nv_float4_t; // Element type for A matrix operand
+using LayoutBTag = cutlass::layout::ColumnMajor; // Layout type for B matrix operand
+constexpr int AlignmentB = 32; // Memory access granularity/alignment of B matrix in units of elements (up to 16 bytes)
+
+// C/D matrix configuration
+using ElementD = cutlass::bfloat16_t; // Element type for D matrix operand
+using ElementC = cutlass::bfloat16_t; // Element type for C matrix operand
+using LayoutCTag = cutlass::layout::RowMajor; // Layout type for C matrix operand
+using LayoutDTag = cutlass::layout::RowMajor; // Layout type for D matrix operand
+constexpr int AlignmentD = 128 / cutlass::sizeof_bits::value; // Memory access granularity/alignment of C matrix in units of elements (up to 16 bytes)
+constexpr int AlignmentC = 128 / cutlass::sizeof_bits::value; // Memory access granularity/alignment of C matrix in units of elements (up to 16 bytes)
+// Kernel functional config
+using ElementAccumulator = float; // Element type for internal accumulation
+using ArchTag = cutlass::arch::Sm100; // Tag indicating the minimum SM that supports the intended feature
+using OperatorClass = cutlass::arch::OpClassBlockScaledTensorOp; // Operator class tag
+
+// Kernel Perf config
+using MmaTileShape = Shape<_256,_256,_256>; // MMA's tile size
+using ClusterShape = Shape<_4,_4,_1>; // Shape of the threadblocks in a cluster
+using PerSmTileShape_MNK = Shape<_128,_256,_256>; // Threadblock-level tile size
+
+using CollectiveEpilogue = typename cutlass::epilogue::collective::CollectiveBuilder<
+ ArchTag, OperatorClass,
+ PerSmTileShape_MNK, ClusterShape,
+ cutlass::epilogue::collective::EpilogueTileAuto,
+ ElementAccumulator, ElementAccumulator,
+ ElementC, LayoutCTag, AlignmentC,
+ ElementD, LayoutDTag, AlignmentD,
+ cutlass::epilogue::collective::EpilogueScheduleAuto // Epilogue schedule policy
+ >::CollectiveOp;
+
+using CollectiveMainloop = typename cutlass::gemm::collective::CollectiveBuilder<
+ ArchTag, OperatorClass,
+ ElementA, LayoutATag, AlignmentA,
+ ElementB, LayoutBTag, AlignmentB,
+ ElementAccumulator,
+ MmaTileShape, ClusterShape,
+ cutlass::gemm::collective::StageCountAutoCarveout(sizeof(typename CollectiveEpilogue::SharedStorage))>,
+ cutlass::gemm::collective::KernelScheduleAuto // Kernel schedule policy. Auto or using targeted scheduling policy
+ >::CollectiveOp;
+
+using GemmKernel = cutlass::gemm::kernel::GemmUniversal<
+ Shape, // Indicates ProblemShape
+ CollectiveMainloop,
+ CollectiveEpilogue,
+ void>;
+
+using Gemm = cutlass::gemm::device::GemmUniversalAdapter;
+
+// Reference device GEMM implementation type
+using StrideA = typename Gemm::GemmKernel::StrideA;
+using LayoutA = decltype(cute::make_layout(make_shape(0,0,0), StrideA{}));
+using LayoutSFA = typename Gemm::GemmKernel::CollectiveMainloop::LayoutSFA; // Scale Factor tensors have an interleaved layout. Bring Layout instead of stride.
+using StrideB = typename Gemm::GemmKernel::StrideB;
+using LayoutB = decltype(cute::make_layout(make_shape(0,0,0), StrideB{}));
+using LayoutSFB = typename Gemm::GemmKernel::CollectiveMainloop::LayoutSFB; // Scale Factor tensors have an interleaved layout. Bring Layout instead of stride.
+using StrideC = typename Gemm::GemmKernel::StrideC;
+using LayoutC = decltype(cute::make_layout(make_shape(0,0,0), StrideC{}));
+using StrideD = typename Gemm::GemmKernel::StrideD;
+using LayoutD = decltype(cute::make_layout(make_shape(0,0,0), StrideD{}));
+
+//
+// Data members
+//
+
+/// Initialization
+StrideA stride_A;
+LayoutA layout_A;
+LayoutSFA layout_SFA;
+StrideB stride_B;
+LayoutB layout_B;
+LayoutSFB layout_SFB;
+StrideC stride_C;
+LayoutC layout_C;
+StrideD stride_D;
+LayoutD layout_D;
+uint64_t seed;
+
+// The HostTensors are only used for allocating memory on host and device, and transferring data between host and device
+// Use cute::Tensor and cute::Layout for iterating thru the matrix elements
+cutlass::HostTensor block_A;
+cutlass::HostTensor block_SFA;
+cutlass::HostTensor block_B;
+cutlass::HostTensor block_SFB;
+cutlass::HostTensor block_C;
+// Output Tensor
+cutlass::HostTensor block_D;
+// Reference Output Tensor
+cutlass::HostTensor block_reference_D;
+#endif // defined(CUTLASS_ARCH_MMA_SM100_SUPPORTED)
+
+template
+auto make_iterator(T* ptr) {
+ using namespace cute;
+ if constexpr (cute::is_subbyte_v) {
+ return subbyte_iterator(ptr);
+ }
+ else {
+ return ptr;
+ }
+}
+
+/////////////////////////////////////////////////////////////////////////////////////////////////
+/// Testbed utility types
+/////////////////////////////////////////////////////////////////////////////////////////////////
+
+// Command line options parsing
+struct Options {
+
+ bool help;
+
+ float alpha, beta;
+ int iterations;
+ int m, n, k;
+
+ Options():
+ help(false),
+ m(1024), n(1024), k(1024),
+ alpha(1.f), beta(0.f),
+ iterations(10)
+ { }
+
+ // Parses the command line
+ void parse(int argc, char const **args) {
+ cutlass::CommandLine cmd(argc, args);
+
+ if (cmd.check_cmd_line_flag("help")) {
+ help = true;
+ return;
+ }
+
+ cmd.get_cmd_line_argument("m", m);
+ cmd.get_cmd_line_argument("n", n);
+ cmd.get_cmd_line_argument("k", k);
+ cmd.get_cmd_line_argument("alpha", alpha, 1.f);
+ cmd.get_cmd_line_argument("beta", beta, 0.f);
+ cmd.get_cmd_line_argument("iterations", iterations);
+ }
+
+ /// Prints the usage statement.
+ std::ostream & print_usage(std::ostream &out) const {
+
+ out << "72a_blackwell_nvfp4_bf16_gemm\n\n"
+ << " Blackwell NVFP4 GEMM using a Warp Specialized kernel.\n\n"
+ << "Options:\n\n"
+ << " --help If specified, displays this usage statement\n\n"
+ << " --m= Sets the M extent of the GEMM\n"
+ << " --n= Sets the N extent of the GEMM\n"
+ << " --k= Sets the K extent of the GEMM\n"
+ << " --alpha= Epilogue scalar alpha\n"
+ << " --beta= Epilogue scalar beta\n\n"
+ << " --iterations= Number of profiling iterations to perform.\n\n";
+
+ out << "\n\nExamples:\n\n"
+ << "$ " << "./examples/72_blackwell_narrow_precision_gemm/72a_blackwell_nvfp4_bf16_gemm" << " --m=1024 --n=512 --k=1024 --alpha=2 --beta=0.707 \n\n";
+
+ return out;
+ }
+
+ /// Compute performance in GFLOP/s
+ double gflops(double runtime_s) const
+ {
+ // Two flops per multiply-add
+ uint64_t flop = uint64_t(2) * m * n * k;
+ double gflop = double(flop) / double(1.0e9);
+ return gflop / runtime_s;
+ }
+};
+
+/// Result structure
+struct Result
+{
+ double avg_runtime_ms;
+ double gflops;
+ cutlass::Status status;
+ cudaError_t error;
+ bool passed;
+
+ Result(
+ double avg_runtime_ms = 0,
+ double gflops = 0,
+ cutlass::Status status = cutlass::Status::kSuccess,
+ cudaError_t error = cudaSuccess)
+ :
+ avg_runtime_ms(avg_runtime_ms), gflops(gflops), status(status), error(error), passed(false)
+ {}
+
+};
+
+#if defined(CUTLASS_ARCH_MMA_SM100_SUPPORTED)
+
+/////////////////////////////////////////////////////////////////////////////////////////////////
+/// GEMM setup and evaluation
+/////////////////////////////////////////////////////////////////////////////////////////////////
+
+/// Helper to initialize a block of device data
+template
+bool initialize_block(
+ cutlass::TensorView view,
+ uint64_t seed) {
+
+ double scope_max, scope_min;
+ constexpr int bits_input = cutlass::sizeof_bits::value;
+
+ if constexpr (bits_input == 1) {
+ scope_max = 2;
+ scope_min = 0;
+ }
+ else if constexpr (bits_input <= 6) {
+ scope_max = 2;
+ scope_min = -2;
+ }
+ else if constexpr (bits_input <= 8) {
+ if constexpr (cute::is_same_v) {
+ scope_max = 4;
+ scope_min = 1;
+ }
+ else {
+ scope_max = 1;
+ scope_min = -1;
+ }
+ }
+ else{
+ scope_max = 4;
+ scope_min = -4;
+ }
+ cutlass::reference::host::TensorFillRandomUniform(
+ view, seed, scope_max, scope_min, 0);
+
+ return true;
+}
+
+/// Initialize operands to be used in the GEMM and reference GEMM
+void initialize(const Options &options) {
+ using namespace cute;
+ // For SFA and SFB tensors layouts
+ using Sm100BlkScaledConfig = typename Gemm::GemmKernel::CollectiveMainloop::Sm100BlkScaledConfig;
+
+ stride_A = cutlass::make_cute_packed_stride(StrideA{}, {options.m, options.k, 1});
+ stride_B = cutlass::make_cute_packed_stride(StrideB{}, {options.n, options.k, 1});
+ stride_C = cutlass::make_cute_packed_stride(StrideC{}, {options.m, options.n, 1});
+ stride_D = cutlass::make_cute_packed_stride(StrideD{}, {options.m, options.n, 1});
+
+ layout_A = make_layout(make_shape(options.m, options.k, 1), stride_A);
+ layout_B = make_layout(make_shape(options.n, options.k, 1), stride_B);
+ layout_C = make_layout(make_shape(options.m, options.n, 1), stride_C);
+ layout_D = make_layout(make_shape(options.m, options.n, 1), stride_D);
+ layout_SFA = Sm100BlkScaledConfig::tile_atom_to_shape_SFA(cute::make_shape(options.m, options.n, options.k, 1));
+ layout_SFB = Sm100BlkScaledConfig::tile_atom_to_shape_SFB(cute::make_shape(options.m, options.n, options.k, 1));
+
+ block_A.reset(cutlass::make_Coord(size(layout_A)));
+ block_B.reset(cutlass::make_Coord(size(layout_B)));
+ block_C.reset(cutlass::make_Coord(size(layout_C)));
+ block_D.reset(cutlass::make_Coord(size(layout_D)));
+ block_reference_D.reset(cutlass::make_Coord(size(layout_D)));
+ block_SFA.reset(cutlass::make_Coord(size(filter_zeros(layout_SFA))));
+ block_SFB.reset(cutlass::make_Coord(size(filter_zeros(layout_SFB))));
+
+ initialize_block(block_A.host_view(), seed + 2021);
+ initialize_block(block_B.host_view(), seed + 2022);
+ initialize_block(block_C.host_view(), seed + 2023);
+ initialize_block(block_SFA.host_view(), seed + 2024);
+ initialize_block(block_SFB.host_view(), seed + 2025);
+
+ block_A.sync_device();
+ block_B.sync_device();
+ block_C.sync_device();
+ block_SFA.sync_device();
+ block_SFB.sync_device();
+}
+
+// Populates a Gemm::Arguments structure from the given commandline options
+typename Gemm::Arguments args_from_options(const Options &options)
+{
+ typename Gemm::Arguments arguments {
+ cutlass::gemm::GemmUniversalMode::kGemm,
+ {options.m, options.n, options.k, 1},
+ { // Mainloop arguments
+ block_A.device_data(), stride_A,
+ block_B.device_data(), stride_B,
+ block_SFA.device_data(), layout_SFA,
+ block_SFB.device_data(), layout_SFB
+ },
+ { // Epilogue arguments
+ {options.alpha, options.beta},
+ block_C.device_data(), stride_C,
+ block_D.device_data(), stride_D
+ }
+ };
+
+ return arguments;
+}
+
+bool verify(const Options &options) {
+ using namespace cute;
+ // Create the arguments for host reference implementation
+ Tensor tensor_A = make_tensor(make_iterator(block_A.host_data()), layout_A);
+ Tensor tensor_SFA = make_tensor(block_SFA.host_data(), layout_SFA);
+ Tensor tensor_B = make_tensor(make_iterator(block_B.host_data()), layout_B);
+ Tensor tensor_SFB = make_tensor(block_SFB.host_data(), layout_SFB);
+
+ cutlass::reference::host::GettBlockScalingMainloopParams<
+ ElementAccumulator, // ElementAccumulator
+ decltype(tensor_A), // TensorA
+ decltype(tensor_SFA), // TensorSfA
+ decltype(tensor_B), // TensorB
+ decltype(tensor_SFB) // TensorSfB
+ > mainloop_params{tensor_A, tensor_SFA, tensor_B, tensor_SFB};
+
+ auto tensor_C = cute::make_tensor(make_iterator(block_C.host_data()), layout_C);
+ auto tensor_D = cute::make_tensor(make_iterator(block_reference_D.host_data()), layout_D);
+
+ cutlass::reference::host::GettBlockScalingEpilogueParams<
+ ElementAccumulator, // ElementScalar
+ ElementAccumulator, // ElementAccumulator
+ ElementAccumulator, // ElementCompute
+ decltype(tensor_C), // TensorC
+ decltype(tensor_D) // TensorD
+ > epilogue_params{options.alpha, options.beta, tensor_C, tensor_D};
+
+ cutlass::reference::host::Gemm3x(mainloop_params, epilogue_params);
+
+ // Comparison
+ block_D.sync_host();
+ bool passed = cutlass::reference::host::TensorEquals(block_reference_D.host_view(), block_D.host_view());
+ passed &= (cutlass::reference::host::TensorNorm(block_reference_D.host_view()) > 0);
+ passed &= (cutlass::reference::host::TensorNorm(block_D.host_view()) > 0);
+
+ return passed;
+}
+
+/// Execute a given example GEMM computation
+template
+int run(Options &options)
+{
+ initialize(options);
+
+ // Instantiate CUTLASS kernel depending on templates
+ Gemm gemm;
+
+ // Create a structure of gemm kernel arguments suitable for invoking an instance of Gemm
+ auto arguments = args_from_options(options);
+
+ // Using the arguments, query for extra workspace required for matrix multiplication computation
+ size_t workspace_size = Gemm::get_workspace_size(arguments);
+
+ // Allocate workspace memory
+ cutlass::device_memory::allocation workspace(workspace_size);
+
+ // Check if the problem size is supported or not
+ CUTLASS_CHECK(gemm.can_implement(arguments));
+
+ // Initialize CUTLASS kernel with arguments and workspace pointer
+ CUTLASS_CHECK(gemm.initialize(arguments, workspace.get()));
+
+ // Correctness / Warmup iteration
+ CUTLASS_CHECK(gemm.run());
+
+ cudaDeviceSynchronize();
+
+ // Check if output from CUTLASS kernel and reference kernel are equal or not
+ Result result;
+ result.passed = verify(options);
+
+ std::cout << " Disposition: " << (result.passed ? "Passed" : "Failed") << std::endl;
+
+ if (!result.passed) {
+ exit(-1);
+ }
+
+ // Run profiling loop
+ if (options.iterations > 0)
+ {
+ GpuTimer timer;
+ timer.start();
+ for (int iter = 0; iter < options.iterations; ++iter) {
+ CUTLASS_CHECK(gemm.initialize(arguments, workspace.get()));
+ CUTLASS_CHECK(gemm.run());
+ }
+ timer.stop();
+
+ // Compute average runtime and GFLOPs.
+ float elapsed_ms = timer.elapsed_millis();
+ result.avg_runtime_ms = double(elapsed_ms) / double(options.iterations);
+ result.gflops = options.gflops(result.avg_runtime_ms / 1000.0);
+
+
+ std::cout << " Problem Size: " << options.m << 'x' << options.n << 'x' << options.k << std::endl;
+ std::cout << " Avg runtime: " << result.avg_runtime_ms << " ms" << std::endl;
+ std::cout << " GFLOPS: " << result.gflops << std::endl;
+ }
+
+ return 0;
+}
+
+#endif // defined(CUTLASS_ARCH_MMA_SM100_SUPPORTED)
+
+///////////////////////////////////////////////////////////////////////////////////////////////////
+
+int main(int argc, char const **args) {
+
+ // CUTLASS must be compiled with CUDA 12.8 or higher Toolkit to run this example
+ // and must have compute capability at least 100.
+ if (__CUDACC_VER_MAJOR__ < 12 || (__CUDACC_VER_MAJOR__ == 12 && __CUDACC_VER_MINOR__ < 8)) {
+ std::cerr << "This example requires CUDA 12.8 or newer." << std::endl;
+ // Returning zero so this test passes on older Toolkits. Its actions are no-op.
+ return 0;
+ }
+
+ cudaDeviceProp props;
+ int current_device_id;
+ CUDA_CHECK(cudaGetDevice(¤t_device_id));
+
+ CUDA_CHECK(cudaGetDeviceProperties(&props, current_device_id));
+
+ if (!(props.major == 10 && props.minor == 0)) {
+ std::cerr << "This example requires a GPU of NVIDIA's Blackwell architecture (compute capability 100)." << std::endl;
+ return 0;
+ }
+
+ //
+ // Parse options
+ //
+
+ Options options;
+
+ options.parse(argc, args);
+
+ if (options.help) {
+ options.print_usage(std::cout) << std::endl;
+ return 0;
+ }
+
+ //
+ // Evaluate CUTLASS kernels
+ //
+#if defined(CUTLASS_ARCH_MMA_SM100_SUPPORTED)
+ run(options);
+#endif // defined(CUTLASS_ARCH_MMA_SM100_SUPPORTED)
+
+ return 0;
+}
+
+/////////////////////////////////////////////////////////////////////////////////////////////////
diff --git a/examples/72_blackwell_narrow_precision_gemm/72b_blackwell_nvfp4_nvfp4_gemm.cu b/examples/72_blackwell_narrow_precision_gemm/72b_blackwell_nvfp4_nvfp4_gemm.cu
new file mode 100644
index 00000000..cefa3e92
--- /dev/null
+++ b/examples/72_blackwell_narrow_precision_gemm/72b_blackwell_nvfp4_nvfp4_gemm.cu
@@ -0,0 +1,594 @@
+/***************************************************************************************************
+ * Copyright (c) 2025 - 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: BSD-3-Clause
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ *
+ * 1. Redistributions of source code must retain the above copyright notice, this
+ * list of conditions and the following disclaimer.
+ *
+ * 2. Redistributions in binary form must reproduce the above copyright notice,
+ * this list of conditions and the following disclaimer in the documentation
+ * and/or other materials provided with the distribution.
+ *
+ * 3. Neither the name of the copyright holder nor the names of its
+ * contributors may be used to endorse or promote products derived from
+ * this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+ * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ * DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+ * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+ * DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+ * SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+ * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+ * OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+ * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ *
+ **************************************************************************************************/
+
+/*! \file
+ \brief A GEMM example using CUTLASS for the NVIDIA Blackwell SM100 architecture.
+
+ This example demonstrate a simple way to instantiate and run a blockscaled NVFP4 GEMM on the NVIDIA Blackwell SM100 architecture
+ on NVIDIA Blackwell SM100 architecture. The kernel outputs quantized fp4 values with scale factors that be the input of another GEMM.
+
+ Similar to 72a_blackwell_nvfp4_bf16_gemm, this kernel leverages:
+ 1. Blockscaled tcgen05.mma instructions.
+
+ 2. Per-SM memory called Tensor Memory (TMEM)
+
+ 3. The extended warp-specialized kernel design introduced in Hopper enabled by use of TMEM
+ which allows us to decouple the execution of MMA and epilogue into separate warps.
+
+ 4. A new SW controlled dynamic scheduler based on cluster launch control (See https://docs.nvidia.com/cuda/parallel-thread-execution).
+
+ Usage:
+
+ $ ./examples/72_blackwell_narrow_precision_gemm/72b_blackwell_nvfp4_nvfp4_gemm --m=2048 --n=2048 --k=2048
+*/
+
+#include
+
+#include "cutlass/cutlass.h"
+
+#include "cute/tensor.hpp"
+#include "cutlass/tensor_ref.h"
+#include "cutlass/epilogue/thread/linear_combination.h"
+#include "cutlass/gemm/dispatch_policy.hpp"
+#include "cutlass/gemm/collective/collective_builder.hpp"
+#include "cutlass/epilogue/collective/collective_builder.hpp"
+#include "cutlass/detail/sm100_blockscaled_layout.hpp"
+#include "cutlass/gemm/device/gemm_universal_adapter.h"
+#include "cutlass/gemm/kernel/gemm_universal.hpp"
+#include "cutlass/gemm/kernel/tile_scheduler_params.h"
+
+#include "cutlass/util/command_line.h"
+#include "cutlass/util/distribution.h"
+#include "cutlass/util/host_tensor.h"
+#include "cutlass/util/packed_stride.hpp"
+#include "cutlass/util/tensor_view_io.h"
+#include "cutlass/util/reference/device/gemm.h"
+#include "cutlass/util/reference/device/tensor_compare.h"
+#include "cutlass/util/reference/host/tensor_fill.h"
+#include "cutlass/util/reference/host/gett.hpp"
+#include "cutlass/util/reference/host/tensor_norm.h"
+#include "cutlass/util/reference/host/tensor_compare.h"
+
+
+#include
+
+#include "helper.h"
+
+using namespace cute;
+
+#if defined(CUTLASS_ARCH_MMA_SM100_SUPPORTED)
+
+
+/////////////////////////////////////////////////////////////////////////////////////////////////
+/// GEMM kernel configurations
+/////////////////////////////////////////////////////////////////////////////////////////////////
+
+// A matrix configuration
+using ElementA = cutlass::nv_float4_t; // Element type for A matrix operand
+using LayoutATag = cutlass::layout::RowMajor; // Layout type for A matrix operand
+constexpr int AlignmentA = 32; // Memory access granularity/alignment of A matrix in units of elements (up to 16 bytes)
+
+// B matrix configuration
+using ElementB = cutlass::nv_float4_t; // Element type for A matrix operand
+using LayoutBTag = cutlass::layout::ColumnMajor; // Layout type for B matrix operand
+constexpr int AlignmentB = 32; // Memory access granularity/alignment of B matrix in units of elements (up to 16 bytes)
+
+// C/D matrix configuration
+using ElementD = cutlass::float_e2m1_t; // Element type for D matrix operand
+using ElementSFD = cutlass::float_ue8m0_t; // Element type for SFB matrix operand
+using ElementC = float; // Element type for C matrix operand
+using LayoutCTag = cutlass::layout::RowMajor; // Layout type for C matrix operand
+using LayoutDTag = cutlass::layout::RowMajor; // Layout type for D matrix operand
+using LayoutSFDTag = LayoutDTag; // Layout type for SFD should be same as D matrix operand
+
+constexpr int AlignmentD = 128 / cutlass::sizeof_bits::value; // Memory access granularity/alignment of C matrix in units of elements (up to 16 bytes)
+constexpr int AlignmentC = 128 / cutlass::sizeof_bits::value; // Memory access granularity/alignment of C matrix in units of elements (up to 16 bytes)
+
+// Kernel functional config
+using ElementAccumulator = float; // Element type for internal accumulation
+using ElementCompute = float; // Element type for internal accumulation
+using ArchTag = cutlass::arch::Sm100; // Tag indicating the minimum SM that supports the intended feature
+using OperatorClass = cutlass::arch::OpClassBlockScaledTensorOp; // Operator class tag
+
+// Kernel Perf config
+using MmaTileShape = Shape<_128,_128,_256>; // MMA's tile size
+using ClusterShape = Shape<_1,_1,_1>; // Shape of the threadblocks in a cluster
+using PerSmTileShape_MNK = Shape<_128,_128,_256>; // Threadblock-level tile size
+
+constexpr int InputSFVectorSize = 16;
+constexpr int OutputSFVectorSize = InputSFVectorSize;
+
+// D = alpha * acc + beta * C
+// With BlockScaleFactor generation.
+using FusionOperation = cutlass::epilogue::fusion::LinCombBlockScaleFactor<
+ OutputSFVectorSize,
+ ElementD,
+ ElementCompute,
+ ElementSFD, LayoutSFDTag,
+ ElementC>;
+
+using CollectiveEpilogue = typename cutlass::epilogue::collective::CollectiveBuilder<
+ ArchTag, OperatorClass,
+ PerSmTileShape_MNK, ClusterShape,
+ cutlass::epilogue::collective::EpilogueTileAuto,
+ ElementAccumulator, ElementAccumulator,
+ ElementC, LayoutCTag, AlignmentC,
+ ElementD, LayoutDTag, AlignmentD,
+ cutlass::epilogue::collective::EpilogueScheduleAuto, // Epilogue schedule policy
+ FusionOperation
+ >::CollectiveOp;
+
+using CollectiveMainloop = typename cutlass::gemm::collective::CollectiveBuilder<
+ ArchTag, OperatorClass,
+ ElementA, LayoutATag, AlignmentA,
+ ElementB, LayoutBTag, AlignmentB,
+ ElementAccumulator,
+ MmaTileShape, ClusterShape,
+ cutlass::gemm::collective::StageCountAutoCarveout(sizeof(typename CollectiveEpilogue::SharedStorage))>,
+ cutlass::gemm::collective::KernelScheduleAuto // Kernel schedule policy. Auto or using targeted scheduling policy
+ >::CollectiveOp;
+
+using GemmKernel = cutlass::gemm::kernel::GemmUniversal<
+ Shape, // Indicates ProblemShape
+ CollectiveMainloop,
+ CollectiveEpilogue,
+ void>;
+
+using Gemm = cutlass::gemm::device::GemmUniversalAdapter;
+
+// Reference device GEMM implementation type
+using StrideA = typename Gemm::GemmKernel::StrideA;
+using LayoutA = decltype(cute::make_layout(make_shape(0,0,0), StrideA{}));
+using LayoutSFA = typename Gemm::GemmKernel::CollectiveMainloop::LayoutSFA; // Scale Factor tensors have an interleaved layout. Bring Layout instead of stride.
+using StrideB = typename Gemm::GemmKernel::StrideB;
+using LayoutB = decltype(cute::make_layout(make_shape(0,0,0), StrideB{}));
+using LayoutSFB = typename Gemm::GemmKernel::CollectiveMainloop::LayoutSFB; // Scale Factor tensors have an interleaved layout. Bring Layout instead of stride.
+using StrideC = typename Gemm::GemmKernel::StrideC;
+using LayoutC = decltype(cute::make_layout(make_shape(0,0,0), StrideC{}));
+using StrideD = typename Gemm::GemmKernel::StrideD;
+using LayoutD = decltype(cute::make_layout(make_shape(0,0,0), StrideD{}));
+
+using FusionOp = typename Gemm::EpilogueOutputOp;
+constexpr bool IsBlockScaleSupported = FusionOp::IsBlockScaleSupported;
+using SfdOutputCfg = cutlass::detail::Sm100BlockScaledOutputConfig;
+using LayoutSFD = typename SfdOutputCfg::LayoutSF;
+
+//
+// Data members
+//
+
+/// Initialization
+StrideA stride_A;
+LayoutA layout_A;
+LayoutSFA layout_SFA;
+StrideB stride_B;
+LayoutB layout_B;
+LayoutSFB layout_SFB;
+StrideC stride_C;
+LayoutC layout_C;
+StrideD stride_D;
+LayoutD layout_D;
+LayoutSFD layout_SFD;
+
+uint64_t seed;
+
+// The HostTensors are only used for allocating memory on host and device, and transferring data between host and device
+// Use cute::Tensor and cute::Layout for iterating thru the matrix elements
+cutlass::HostTensor block_A;
+cutlass::HostTensor block_SFA;
+cutlass::HostTensor block_B;
+cutlass::HostTensor block_SFB;
+cutlass::HostTensor block_C;
+// Output Tensors
+cutlass::HostTensor block_D;
+cutlass::HostTensor block_SFD;
+// Reference Output Tensors
+cutlass::HostTensor block_reference_D;
+cutlass::HostTensor block_reference_SFD;
+// Matrix-wide normalization constant
+cutlass::HostTensor block_Normconst;
+
+#endif // defined(CUTLASS_ARCH_MMA_SM100_SUPPORTED)
+
+template
+auto make_iterator(T* ptr) {
+ using namespace cute;
+ if constexpr (cute::is_subbyte_v) {
+ return subbyte_iterator(ptr);
+ }
+ else {
+ return ptr;
+ }
+}
+
+/////////////////////////////////////////////////////////////////////////////////////////////////
+/// Testbed utility types
+/////////////////////////////////////////////////////////////////////////////////////////////////
+
+// Command line options parsing
+struct Options {
+
+ bool help;
+
+ float alpha, beta;
+ int iterations;
+ int m, n, k;
+
+ Options():
+ help(false),
+ m(1024), n(1024), k(1024),
+ alpha(1.f), beta(0.f),
+ iterations(10)
+ { }
+
+ // Parses the command line
+ void parse(int argc, char const **args) {
+ cutlass::CommandLine cmd(argc, args);
+
+ if (cmd.check_cmd_line_flag("help")) {
+ help = true;
+ return;
+ }
+
+ cmd.get_cmd_line_argument("m", m);
+ cmd.get_cmd_line_argument("n", n);
+ cmd.get_cmd_line_argument("k", k);
+ cmd.get_cmd_line_argument("alpha", alpha, 1.f);
+ cmd.get_cmd_line_argument("beta", beta, 0.f);
+ cmd.get_cmd_line_argument("iterations", iterations);
+ }
+
+ /// Prints the usage statement.
+ std::ostream & print_usage(std::ostream &out) const {
+
+ out << "72b_blackwell_nvfp4_nvfp4_gemm\n\n"
+ << " Blackwell NVFP4 GEMM using a Warp Specialized kernel.\n\n"
+ << "Options:\n\n"
+ << " --help If specified, displays this usage statement\n\n"
+ << " --m= Sets the M extent of the GEMM\n"
+ << " --n= Sets the N extent of the GEMM\n"
+ << " --k= Sets the K extent of the GEMM\n"
+ << " --alpha= Epilogue scalar alpha\n"
+ << " --beta= Epilogue scalar beta\n\n"
+ << " --iterations= Number of profiling iterations to perform.\n\n";
+
+ out << "\n\nExamples:\n\n"
+ << "$ " << "./examples/72_blackwell_narrow_precision_gemm/72b_blackwell_nvfp4_nvfp4_gemm" << " --m=1024 --n=512 --k=1024 --alpha=2 --beta=0.707 \n\n";
+
+ return out;
+ }
+
+ /// Compute performance in GFLOP/s
+ double gflops(double runtime_s) const
+ {
+ // Two flops per multiply-add
+ uint64_t flop = uint64_t(2) * m * n * k;
+ double gflop = double(flop) / double(1.0e9);
+ return gflop / runtime_s;
+ }
+};
+
+/// Result structure
+struct Result
+{
+ double avg_runtime_ms;
+ double gflops;
+ cutlass::Status status;
+ cudaError_t error;
+ bool passed;
+
+ Result(
+ double avg_runtime_ms = 0,
+ double gflops = 0,
+ cutlass::Status status = cutlass::Status::kSuccess,
+ cudaError_t error = cudaSuccess)
+ :
+ avg_runtime_ms(avg_runtime_ms), gflops(gflops), status(status), error(error), passed(false)
+ {}
+
+};
+
+#if defined(CUTLASS_ARCH_MMA_SM100_SUPPORTED)
+
+/////////////////////////////////////////////////////////////////////////////////////////////////
+/// GEMM setup and evaluation
+/////////////////////////////////////////////////////////////////////////////////////////////////
+
+/// Helper to initialize a block of device data
+template
+bool initialize_block(
+ cutlass::TensorView view,
+ uint64_t seed) {
+
+ double scope_max, scope_min;
+ constexpr int bits_input = cutlass::sizeof_bits::value;
+
+ if constexpr (bits_input == 1) {
+ scope_max = 2;
+ scope_min = 0;
+ }
+ else if constexpr (bits_input <= 6) {
+ scope_max = 2;
+ scope_min = -2;
+ }
+ else if constexpr (bits_input <= 8) {
+ if constexpr (cute::is_same_v) {
+ scope_max = 4;
+ scope_min = 1;
+ }
+ else {
+ scope_max = 1;
+ scope_min = -1;
+ }
+ }
+ else{
+ scope_max = 4;
+ scope_min = -4;
+ }
+ cutlass::reference::host::TensorFillRandomUniform(
+ view, seed, scope_max, scope_min, 0);
+
+ return true;
+}
+
+/// Initialize operands to be used in the GEMM and reference GEMM
+void initialize(const Options &options) {
+ using namespace cute;
+ // For SFA and SFB tensors layouts
+ using Sm100BlkScaledConfig = typename Gemm::GemmKernel::CollectiveMainloop::Sm100BlkScaledConfig;
+ // For SFD tensor layout
+ using Sm100BlockScaledOutputConfig = typename Gemm::GemmKernel::CollectiveMainloop::Sm100BlkScaledConfig;
+
+ stride_A = cutlass::make_cute_packed_stride(StrideA{}, {options.m, options.k, 1});
+ stride_B = cutlass::make_cute_packed_stride(StrideB{}, {options.n, options.k, 1});
+ stride_C = cutlass::make_cute_packed_stride(StrideC{}, {options.m, options.n, 1});
+ stride_D = cutlass::make_cute_packed_stride(StrideD{}, {options.m, options.n, 1});
+
+ layout_A = make_layout(make_shape(options.m, options.k, 1), stride_A);
+ layout_B = make_layout(make_shape(options.n, options.k, 1), stride_B);
+ layout_C = make_layout(make_shape(options.m, options.n, 1), stride_C);
+ layout_D = make_layout(make_shape(options.m, options.n, 1), stride_D);
+ layout_SFA = Sm100BlkScaledConfig::tile_atom_to_shape_SFA(cute::make_shape(options.m, options.n, options.k, 1));
+ layout_SFB = Sm100BlkScaledConfig::tile_atom_to_shape_SFB(cute::make_shape(options.m, options.n, options.k, 1));
+ layout_SFD = SfdOutputCfg::tile_atom_to_shape_SFD(cute::make_shape(options.m, options.n, options.k, 1));
+
+ block_A.reset(cutlass::make_Coord(size(layout_A)));
+ block_B.reset(cutlass::make_Coord(size(layout_B)));
+ block_C.reset(cutlass::make_Coord(size(layout_C)));
+ block_D.reset(cutlass::make_Coord(size(layout_D)));
+ block_reference_D.reset(cutlass::make_Coord(size(layout_D)));
+ block_reference_SFD.reset(cutlass::make_Coord(size(filter_zeros(layout_SFD))));
+ block_Normconst.reset(cutlass::make_Coord(1));
+
+ block_SFA.reset(cutlass::make_Coord(size(filter_zeros(layout_SFA))));
+ block_SFB.reset(cutlass::make_Coord(size(filter_zeros(layout_SFB))));
+ block_SFD.reset(cutlass::make_Coord(size(filter_zeros(layout_SFD))));
+
+ initialize_block(block_A.host_view(), seed + 2021);
+ initialize_block(block_B.host_view(), seed + 2022);
+ initialize_block(block_C.host_view(), seed + 2023);
+ initialize_block(block_SFA.host_view(), seed + 2024);
+ initialize_block(block_SFB.host_view(), seed + 2025);
+ block_Normconst.at(cutlass::make_Coord(0)) = 2;
+
+ block_A.sync_device();
+ block_B.sync_device();
+ block_C.sync_device();
+ block_D.sync_device();
+ block_SFA.sync_device();
+ block_SFB.sync_device();
+ block_SFD.sync_device();
+ block_Normconst.sync_device();
+
+}
+
+// Populates a Gemm::Arguments structure from the given commandline options
+typename Gemm::Arguments args_from_options(const Options &options)
+{
+ typename Gemm::Arguments arguments {
+ cutlass::gemm::GemmUniversalMode::kGemm,
+ {options.m, options.n, options.k, 1},
+ { // Mainloop arguments
+ block_A.device_data(), stride_A,
+ block_B.device_data(), stride_B,
+ block_SFA.device_data(), layout_SFA,
+ block_SFB.device_data(), layout_SFB
+ },
+ { // Epilogue arguments
+ { options.alpha, options.beta },
+ block_C.device_data(), stride_C,
+ block_D.device_data(), stride_D}
+ };
+
+ if constexpr (IsBlockScaleSupported) {
+ arguments.epilogue.thread.block_scale_factor_ptr = block_SFD.device_data();
+ arguments.epilogue.thread.norm_constant_ptr = block_Normconst.device_data();
+ }
+
+ return arguments;
+}
+
+bool verify(const Options &options) {
+ using namespace cute;
+ // Create the arguments for host reference implementation
+ Tensor tensor_A = make_tensor(make_iterator(block_A.host_data()), layout_A);
+ Tensor tensor_SFA = make_tensor(block_SFA.host_data(), layout_SFA);
+ Tensor tensor_B = make_tensor(make_iterator(block_B.host_data()), layout_B);
+ Tensor tensor_SFB = make_tensor(block_SFB.host_data(), layout_SFB);
+
+ // think about how to simplify the gemm3x interface.
+ cutlass::reference::host::GettBlockScalingMainloopParams<
+ ElementAccumulator, // ElementAccumulator
+ decltype(tensor_A), // TensorA
+ decltype(tensor_SFA), // TensorSfA
+ decltype(tensor_B), // TensorB
+ decltype(tensor_SFB) // TensorSfB
+ > mainloop_params{tensor_A, tensor_SFA, tensor_B, tensor_SFB};
+
+ Tensor tensor_C = cute::make_tensor(make_iterator(block_C.host_data()), layout_C);
+ Tensor tensor_D = cute::make_tensor(make_iterator(block_reference_D.host_data()), layout_D);
+ Tensor tensor_SFD = make_tensor(block_reference_SFD.host_data(), layout_SFD);
+
+ cutlass::reference::host::GettBlockScalingEpilogueParams<
+ ElementCompute, // ElementScalar
+ ElementAccumulator, // ElementAccumulator
+ ElementCompute, // ElementCompute
+ decltype(tensor_C), // TensorC
+ decltype(tensor_D), // TensorD
+ decltype(tensor_SFD), // TensorSfD
+ cute::Int,
+ cutlass::reference::host::SfStrategy::SfDGen
+ > epilogue_params {options.alpha, options.beta, tensor_C, tensor_D, tensor_SFD, block_Normconst.at(cutlass::make_Coord(0))};
+
+ cutlass::reference::host::Gemm3x(mainloop_params, epilogue_params);
+
+ // Comparison
+ block_D.sync_host();
+ bool passed = cutlass::reference::host::TensorEquals(block_reference_D.host_view(), block_D.host_view());
+ passed &= (cutlass::reference::host::TensorNorm(block_reference_D.host_view()) > 0);
+ passed &= (cutlass::reference::host::TensorNorm(block_D.host_view()) > 0);
+
+ return passed;
+}
+
+/// Execute a given example GEMM computation
+template
+int run(Options &options)
+{
+ initialize(options);
+
+ // Instantiate CUTLASS kernel depending on templates
+ Gemm gemm;
+
+ // Create a structure of gemm kernel arguments suitable for invoking an instance of Gemm
+ auto arguments = args_from_options(options);
+
+ // Using the arguments, query for extra workspace required for matrix multiplication computation
+ size_t workspace_size = Gemm::get_workspace_size(arguments);
+
+ // Allocate workspace memory
+ cutlass::device_memory::allocation workspace(workspace_size);
+
+ // Check if the problem size is supported or not
+ CUTLASS_CHECK(gemm.can_implement(arguments));
+
+ // Initialize CUTLASS kernel with arguments and workspace pointer
+ CUTLASS_CHECK(gemm.initialize(arguments, workspace.get()));
+
+ // Correctness / Warmup iteration
+ CUTLASS_CHECK(gemm.run());
+
+ cudaDeviceSynchronize();
+
+ // Check if output from CUTLASS kernel and reference kernel are equal or not
+ Result result;
+ result.passed = verify(options);
+
+ std::cout << " Disposition: " << (result.passed ? "Passed" : "Failed") << std::endl;
+
+ if (!result.passed) {
+ exit(-1);
+ }
+
+ // Run profiling loop
+ if (options.iterations > 0)
+ {
+ GpuTimer timer;
+ timer.start();
+ for (int iter = 0; iter < options.iterations; ++iter) {
+ CUTLASS_CHECK(gemm.initialize(arguments, workspace.get()));
+ CUTLASS_CHECK(gemm.run());
+ }
+ timer.stop();
+
+ // Compute average runtime and GFLOPs.
+ float elapsed_ms = timer.elapsed_millis();
+ result.avg_runtime_ms = double(elapsed_ms) / double(options.iterations);
+ result.gflops = options.gflops(result.avg_runtime_ms / 1000.0);
+
+
+ std::cout << " Problem Size: " << options.m << 'x' << options.n << 'x' << options.k << std::endl;
+ std::cout << " Avg runtime: " << result.avg_runtime_ms << " ms" << std::endl;
+ std::cout << " GFLOPS: " << result.gflops << std::endl;
+ }
+
+ return 0;
+}
+
+#endif // defined(CUTLASS_ARCH_MMA_SM100_SUPPORTED)
+
+///////////////////////////////////////////////////////////////////////////////////////////////////
+
+int main(int argc, char const **args) {
+
+ // CUTLASS must be compiled with CUDA 12.8 or higher Toolkit to run this example
+ // and must have compute capability at least 100.
+ if (__CUDACC_VER_MAJOR__ < 12 || (__CUDACC_VER_MAJOR__ == 12 && __CUDACC_VER_MINOR__ < 8)) {
+ std::cerr << "This example requires CUDA 12.8 or newer." << std::endl;
+ // Returning zero so this test passes on older Toolkits. Its actions are no-op.
+ return 0;
+ }
+
+ cudaDeviceProp props;
+ int current_device_id;
+ CUDA_CHECK(cudaGetDevice(¤t_device_id));
+
+ CUDA_CHECK(cudaGetDeviceProperties(&props, current_device_id));
+
+ if (!(props.major == 10 && props.minor == 0)) {
+ std::cerr << "This example requires a GPU of NVIDIA's Blackwell architecture (compute capability 100)." << std::endl;
+ return 0;
+ }
+
+ //
+ // Parse options
+ //
+
+ Options options;
+
+ options.parse(argc, args);
+
+ if (options.help) {
+ options.print_usage(std::cout) << std::endl;
+ return 0;
+ }
+
+ //
+ // Evaluate CUTLASS kernels
+ //
+#if defined(CUTLASS_ARCH_MMA_SM100_SUPPORTED)
+ run(options);
+#endif // defined(CUTLASS_ARCH_MMA_SM100_SUPPORTED)
+
+ return 0;
+}
+
+/////////////////////////////////////////////////////////////////////////////////////////////////
diff --git a/examples/72_blackwell_narrow_precision_gemm/72c_blackwell_mixed_mxfp8_bf16_gemm.cu b/examples/72_blackwell_narrow_precision_gemm/72c_blackwell_mixed_mxfp8_bf16_gemm.cu
new file mode 100644
index 00000000..b73f2c94
--- /dev/null
+++ b/examples/72_blackwell_narrow_precision_gemm/72c_blackwell_mixed_mxfp8_bf16_gemm.cu
@@ -0,0 +1,545 @@
+/***************************************************************************************************
+ * Copyright (c) 2025 - 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: BSD-3-Clause
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ *
+ * 1. Redistributions of source code must retain the above copyright notice, this
+ * list of conditions and the following disclaimer.
+ *
+ * 2. Redistributions in binary form must reproduce the above copyright notice,
+ * this list of conditions and the following disclaimer in the documentation
+ * and/or other materials provided with the distribution.
+ *
+ * 3. Neither the name of the copyright holder nor the names of its
+ * contributors may be used to endorse or promote products derived from
+ * this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+ * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ * DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+ * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+ * DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+ * SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+ * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+ * OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+ * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ *
+ **************************************************************************************************/
+
+/*! \file
+ \brief A GEMM example using CUTLASS for the NVIDIA Blackwell SM100 architecture.
+
+ This example demonstrates a simple way to instantiate and run a mixed precision blockscaled GEMM on the NVIDIA Blackwell SM100 architecture.
+ This Blackwell SM100 CUTLASS kernel uses the new Block Scaled Tensor Core MMA Instructions (tcgen05.mma.blockscaled) introduced
+ on the Blackwell architecture (sm100a) which have the same throughput compared to fp8 Tensor Core MMA instructions (tcgen05.mma)
+ and 2x throughput compared to fp8 Hopper Tensor Core MMA Instructions (WGMMA) (See https://docs.nvidia.com/cuda/parallel-thread-execution).
+
+ Similar to 72a_blackwell_nvfp4_fp32_gemm, this kernel leverages:
+ 1. Blockscaled tcgen05.mma instructions.
+
+ 2. Per-SM memory called Tensor Memory (TMEM) (Please refer to CUDA 12.8 docs on https://docs.nvidia.com/cuda/).
+
+ 3. The extended warp-specialized kernel design introduced in Hopper enabled by use of TMEM
+ which allows us to decouple the execution of MMA and epilogue into separate warps.
+
+ 4. A new SW controlled dynamic scheduler based on cluster launch control (See https://docs.nvidia.com/cuda/parallel-thread-execution).
+
+ Usage:
+
+ $ ./examples/72_blackwell_narrow_precision_gemm/72c_blackwell_mixed_mxfp8_bf16_gemm --m=2048 --n=2048 --k=2048
+*/
+
+#include
+
+#include "cutlass/cutlass.h"
+
+#include "cute/tensor.hpp"
+#include "cutlass/tensor_ref.h"
+#include "cutlass/epilogue/thread/linear_combination.h"
+#include "cutlass/gemm/dispatch_policy.hpp"
+#include "cutlass/gemm/collective/collective_builder.hpp"
+#include "cutlass/epilogue/collective/collective_builder.hpp"
+#include "cutlass/detail/sm100_blockscaled_layout.hpp"
+#include "cutlass/gemm/device/gemm_universal_adapter.h"
+#include "cutlass/gemm/kernel/gemm_universal.hpp"
+#include "cutlass/gemm/kernel/tile_scheduler_params.h"
+
+#include "cutlass/util/command_line.h"
+#include "cutlass/util/distribution.h"
+#include "cutlass/util/host_tensor.h"
+#include "cutlass/util/packed_stride.hpp"
+#include "cutlass/util/tensor_view_io.h"
+#include "cutlass/util/reference/device/gemm.h"
+#include "cutlass/util/reference/device/tensor_compare.h"
+#include "cutlass/util/reference/host/tensor_fill.h"
+#include "cutlass/util/reference/host/gett.hpp"
+#include "cutlass/util/reference/host/tensor_norm.h"
+#include "cutlass/util/reference/host/tensor_compare.h"
+
+
+#include
+
+#include "helper.h"
+
+using namespace cute;
+
+#if defined(CUTLASS_ARCH_MMA_SM100_SUPPORTED)
+
+
+/////////////////////////////////////////////////////////////////////////////////////////////////
+/// GEMM kernel configurations
+/////////////////////////////////////////////////////////////////////////////////////////////////
+
+// A matrix configuration
+using ElementA = cutlass::mx_float8_t; // Element type for A matrix operand
+using LayoutATag = cutlass::layout::RowMajor; // Layout type for A matrix operand
+constexpr int AlignmentA = 16; // Memory access granularity/alignment of A matrix in units of elements (up to 16 bytes)
+
+// B matrix configuration
+using ElementB = cutlass::mx_float4_t; // Element type for A matrix operand
+using LayoutBTag = cutlass::layout::ColumnMajor; // Layout type for B matrix operand
+constexpr int AlignmentB = 128; // Memory access granularity/alignment of B matrix in units of elements (up to 16 bytes)
+
+// C/D matrix configuration
+using ElementD = cutlass::bfloat16_t; // Element type for D matrix operand
+using ElementC = cutlass::bfloat16_t; // Element type for C matrix operand
+using LayoutCTag = cutlass::layout::RowMajor; // Layout type for C matrix operand
+using LayoutDTag = cutlass::layout::RowMajor; // Layout type for D matrix operand
+constexpr int AlignmentD = 128 / cutlass::sizeof_bits::value; // Memory access granularity/alignment of C matrix in units of elements (up to 16 bytes)
+constexpr int AlignmentC = 128 / cutlass::sizeof_bits::value; // Memory access granularity/alignment of C matrix in units of elements (up to 16 bytes)
+// Kernel functional config
+using ElementAccumulator = float; // Element type for internal accumulation
+using ArchTag = cutlass::arch::Sm100; // Tag indicating the minimum SM that supports the intended feature
+using OperatorClass = cutlass::arch::OpClassBlockScaledTensorOp; // Operator class tag
+
+// Kernel Perf config
+using MmaTileShape = Shape<_256,_256,_256>; // MMA's tile size
+using ClusterShape = Shape<_4,_4,_1>; // Shape of the threadblocks in a cluster
+using PerSmTileShape_MNK = Shape<_128,_256,_256>; // Threadblock-level tile size
+
+using CollectiveEpilogue = typename cutlass::epilogue::collective::CollectiveBuilder<
+ ArchTag, OperatorClass,
+ PerSmTileShape_MNK, ClusterShape,
+ cutlass::epilogue::collective::EpilogueTileAuto,
+ ElementAccumulator, ElementAccumulator,
+ ElementC, LayoutCTag, AlignmentC,
+ ElementD, LayoutDTag, AlignmentD,
+ cutlass::epilogue::collective::EpilogueScheduleAuto // Epilogue schedule policy
+ >::CollectiveOp;
+
+using CollectiveMainloop = typename cutlass::gemm::collective::CollectiveBuilder<
+ ArchTag, OperatorClass,
+ ElementA, LayoutATag, AlignmentA,
+ ElementB, LayoutBTag, AlignmentB,
+ ElementAccumulator,
+ MmaTileShape, ClusterShape,
+ cutlass::gemm::collective::StageCountAutoCarveout(sizeof(typename CollectiveEpilogue::SharedStorage))>,
+ cutlass::gemm::collective::KernelScheduleAuto // Kernel schedule policy. Auto or using targeted scheduling policy
+ >::CollectiveOp;
+
+using GemmKernel = cutlass::gemm::kernel::GemmUniversal<
+ Shape, // Indicates ProblemShape
+ CollectiveMainloop,
+ CollectiveEpilogue,
+ void>;
+
+using Gemm = cutlass::gemm::device::GemmUniversalAdapter;
+
+// Reference device GEMM implementation type
+using StrideA = typename Gemm::GemmKernel::StrideA;
+using LayoutA = decltype(cute::make_layout(make_shape(0,0,0), StrideA{}));
+using LayoutSFA = typename Gemm::GemmKernel::CollectiveMainloop::LayoutSFA; // Scale Factor tensors have an interleaved layout. Bring Layout instead of stride.
+using StrideB = typename Gemm::GemmKernel::StrideB;
+using LayoutB = decltype(cute::make_layout(make_shape(0,0,0), StrideB{}));
+using LayoutSFB = typename Gemm::GemmKernel::CollectiveMainloop::LayoutSFB; // Scale Factor tensors have an interleaved layout. Bring Layout instead of stride.
+using StrideC = typename Gemm::GemmKernel::StrideC;
+using LayoutC = decltype(cute::make_layout(make_shape(0,0,0), StrideC{}));
+using StrideD = typename Gemm::GemmKernel::StrideD;
+using LayoutD = decltype(cute::make_layout(make_shape(0,0,0), StrideD{}));
+
+//
+// Data members
+//
+
+/// Initialization
+StrideA stride_A;
+LayoutA layout_A;
+LayoutSFA layout_SFA;
+StrideB stride_B;
+LayoutB layout_B;
+LayoutSFB layout_SFB;
+StrideC stride_C;
+LayoutC layout_C;
+StrideD stride_D;
+LayoutD layout_D;
+uint64_t seed;
+
+// The HostTensors are only used for allocating memory on host and device, and transferring data between host and device
+// Use cute::Tensor and cute::Layout for iterating thru the matrix elements
+cutlass::HostTensor block_A;
+cutlass::HostTensor block_SFA;
+cutlass::HostTensor block_B;
+cutlass::HostTensor block_SFB;
+cutlass::HostTensor block_C;
+// Output Tensor
+cutlass::HostTensor block_D;
+// Reference Output Tensor
+cutlass::HostTensor block_reference_D;
+#endif // defined(CUTLASS_ARCH_MMA_SM100_SUPPORTED)
+
+template
+auto make_iterator(T* ptr) {
+ using namespace cute;
+ if constexpr (cute::is_subbyte_v) {
+ return subbyte_iterator(ptr);
+ }
+ else {
+ return ptr;
+ }
+}
+
+/////////////////////////////////////////////////////////////////////////////////////////////////
+/// Testbed utility types
+/////////////////////////////////////////////////////////////////////////////////////////////////
+
+// Command line options parsing
+struct Options {
+
+ bool help;
+
+ float alpha, beta;
+ int iterations;
+ int m, n, k;
+
+ Options():
+ help(false),
+ m(1024), n(1024), k(1024),
+ alpha(1.f), beta(0.f),
+ iterations(10)
+ { }
+
+ // Parses the command line
+ void parse(int argc, char const **args) {
+ cutlass::CommandLine cmd(argc, args);
+
+ if (cmd.check_cmd_line_flag("help")) {
+ help = true;
+ return;
+ }
+
+ cmd.get_cmd_line_argument("m", m);
+ cmd.get_cmd_line_argument("n", n);
+ cmd.get_cmd_line_argument("k", k);
+ cmd.get_cmd_line_argument("alpha", alpha, 1.f);
+ cmd.get_cmd_line_argument("beta", beta, 0.f);
+ cmd.get_cmd_line_argument("iterations", iterations);
+ }
+
+ /// Prints the usage statement.
+ std::ostream & print_usage(std::ostream &out) const {
+
+ out << "72c_blackwell_mixed_mxfp8_bf16_gemm\n\n"
+ << " Blackwell Mxfp8 x Mxfp4 GEMM using a Warp Specialized kernel.\n\n"
+ << "Options:\n\n"
+ << " --help If specified, displays this usage statement\n\n"
+ << " --m= Sets the M extent of the GEMM\n"
+ << " --n= Sets the N extent of the GEMM\n"
+ << " --k= Sets the K extent of the GEMM\n"
+ << " --alpha= Epilogue scalar alpha\n"
+ << " --beta= Epilogue scalar beta\n\n"
+ << " --iterations= Number of profiling iterations to perform.\n\n";
+
+ out << "\n\nExamples:\n\n"
+ << "$ " << "/examples/72_blackwell_narrow_precision_gemm/72c_blackwell_mixed_mxfp8_bf16_gemm" << " --m=1024 --n=512 --k=1024 --alpha=2 --beta=0.707 \n\n";
+
+ return out;
+ }
+
+ /// Compute performance in GFLOP/s
+ double gflops(double runtime_s) const
+ {
+ // Two flops per multiply-add
+ uint64_t flop = uint64_t(2) * m * n * k;
+ double gflop = double(flop) / double(1.0e9);
+ return gflop / runtime_s;
+ }
+};
+
+/// Result structure
+struct Result
+{
+ double avg_runtime_ms;
+ double gflops;
+ cutlass::Status status;
+ cudaError_t error;
+ bool passed;
+
+ Result(
+ double avg_runtime_ms = 0,
+ double gflops = 0,
+ cutlass::Status status = cutlass::Status::kSuccess,
+ cudaError_t error = cudaSuccess)
+ :
+ avg_runtime_ms(avg_runtime_ms), gflops(gflops), status(status), error(error), passed(false)
+ {}
+
+};
+
+#if defined(CUTLASS_ARCH_MMA_SM100_SUPPORTED)
+
+/////////////////////////////////////////////////////////////////////////////////////////////////
+/// GEMM setup and evaluation
+/////////////////////////////////////////////////////////////////////////////////////////////////
+
+/// Helper to initialize a block of device data
+template
+bool initialize_block(
+ cutlass::TensorView view,
+ uint64_t seed) {
+
+ double scope_max, scope_min;
+ constexpr int bits_input = cutlass::sizeof_bits::value;
+
+ if constexpr (bits_input == 1) {
+ scope_max = 2;
+ scope_min = 0;
+ }
+ else if constexpr (bits_input <= 6) {
+ scope_max = 2;
+ scope_min = -2;
+ }
+ else if constexpr (bits_input <= 8) {
+ if constexpr (cute::is_same_v) {
+ scope_max = 4;
+ scope_min = 1;
+ }
+ else {
+ scope_max = 1;
+ scope_min = -1;
+ }
+ }
+ else{
+ scope_max = 4;
+ scope_min = -4;
+ }
+ cutlass::reference::host::TensorFillRandomUniform(
+ view, seed, scope_max, scope_min, 0);
+
+ return true;
+}
+
+/// Initialize operands to be used in the GEMM and reference GEMM
+void initialize(const Options &options) {
+ using namespace cute;
+ // For SFA and SFB tensors layouts
+ using Sm100BlkScaledConfig = typename Gemm::GemmKernel::CollectiveMainloop::Sm100BlkScaledConfig;
+
+ stride_A = cutlass::make_cute_packed_stride(StrideA{}, {options.m, options.k, 1});
+ stride_B = cutlass::make_cute_packed_stride(StrideB{}, {options.n, options.k, 1});
+ stride_C = cutlass::make_cute_packed_stride(StrideC{}, {options.m, options.n, 1});
+ stride_D = cutlass::make_cute_packed_stride(StrideD{}, {options.m, options.n, 1});
+
+ layout_A = make_layout(make_shape(options.m, options.k, 1), stride_A);
+ layout_B = make_layout(make_shape(options.n, options.k, 1), stride_B);
+ layout_C = make_layout(make_shape(options.m, options.n, 1), stride_C);
+ layout_D = make_layout(make_shape(options.m, options.n, 1), stride_D);
+ layout_SFA = Sm100BlkScaledConfig::tile_atom_to_shape_SFA(cute::make_shape(options.m, options.n, options.k, 1));
+ layout_SFB = Sm100BlkScaledConfig::tile_atom_to_shape_SFB(cute::make_shape(options.m, options.n, options.k, 1));
+
+ block_A.reset(cutlass::make_Coord(size(layout_A)));
+ block_B.reset(cutlass::make_Coord(size(layout_B)));
+ block_C.reset(cutlass::make_Coord(size(layout_C)));
+ block_D.reset(cutlass::make_Coord(size(layout_D)));
+ block_reference_D.reset(cutlass::make_Coord(size(layout_D)));
+ block_SFA.reset(cutlass::make_Coord(size(filter_zeros(layout_SFA))));
+ block_SFB.reset(cutlass::make_Coord(size(filter_zeros(layout_SFB))));
+
+ initialize_block(block_A.host_view(), seed + 2021);
+ initialize_block(block_B.host_view(), seed + 2022);
+ initialize_block(block_C.host_view(), seed + 2023);
+ initialize_block(block_SFA.host_view(), seed + 2024);
+ initialize_block(block_SFB.host_view(), seed + 2025);
+
+ block_A.sync_device();
+ block_B.sync_device();
+ block_C.sync_device();
+ block_SFA.sync_device();
+ block_SFB.sync_device();
+}
+
+// Populates a Gemm::Arguments structure from the given commandline options
+typename Gemm::Arguments args_from_options(const Options &options)
+{
+ typename Gemm::Arguments arguments {
+ cutlass::gemm::GemmUniversalMode::kGemm,
+ {options.m, options.n, options.k, 1},
+ { // Mainloop arguments
+ block_A.device_data(), stride_A,
+ block_B.device_data(), stride_B,
+ block_SFA.device_data(), layout_SFA,
+ block_SFB.device_data(), layout_SFB
+ },
+ { // Epilogue arguments
+ {options.alpha, options.beta},
+ block_C.device_data(), stride_C,
+ block_D.device_data(), stride_D
+ }
+ };
+
+ return arguments;
+}
+
+bool verify(const Options &options) {
+ using namespace cute;
+ // Create the arguments for host reference implementation
+ Tensor tensor_A = make_tensor(make_iterator(block_A.host_data()), layout_A);
+ Tensor tensor_SFA = make_tensor(block_SFA.host_data(), layout_SFA);
+ Tensor tensor_B = make_tensor(make_iterator(block_B.host_data()), layout_B);
+ Tensor tensor_SFB = make_tensor(block_SFB.host_data(), layout_SFB);
+
+ cutlass::reference::host::GettBlockScalingMainloopParams<
+ ElementAccumulator, // ElementAccumulator
+ decltype(tensor_A), // TensorA
+ decltype(tensor_SFA), // TensorSfA
+ decltype(tensor_B), // TensorB
+ decltype(tensor_SFB) // TensorSfB
+ > mainloop_params{tensor_A, tensor_SFA, tensor_B, tensor_SFB};
+
+ auto tensor_C = cute::make_tensor(make_iterator(block_C.host_data()), layout_C);
+ auto tensor_D = cute::make_tensor(make_iterator(block_reference_D.host_data()), layout_D);
+
+ cutlass::reference::host::GettBlockScalingEpilogueParams<
+ ElementAccumulator, // ElementScalar
+ ElementAccumulator, // ElementAccumulator
+ ElementAccumulator, // ElementCompute
+ decltype(tensor_C), // TensorC
+ decltype(tensor_D) // TensorD
+ > epilogue_params{options.alpha, options.beta, tensor_C, tensor_D};
+
+ cutlass::reference::host::Gemm3x(mainloop_params, epilogue_params);
+
+ // Comparison
+ block_D.sync_host();
+ bool passed = cutlass::reference::host::TensorEquals(block_reference_D.host_view(), block_D.host_view());
+ passed &= (cutlass::reference::host::TensorNorm(block_reference_D.host_view()) > 0);
+ passed &= (cutlass::reference::host::TensorNorm(block_D.host_view()) > 0);
+
+ return passed;
+}
+
+/// Execute a given example GEMM computation
+template
+int run(Options &options)
+{
+ initialize(options);
+
+ // Instantiate CUTLASS kernel depending on templates
+ Gemm gemm;
+
+ // Create a structure of gemm kernel arguments suitable for invoking an instance of Gemm
+ auto arguments = args_from_options(options);
+
+ // Using the arguments, query for extra workspace required for matrix multiplication computation
+ size_t workspace_size = Gemm::get_workspace_size(arguments);
+
+ // Allocate workspace memory
+ cutlass::device_memory::allocation workspace(workspace_size);
+
+ // Check if the problem size is supported or not
+ CUTLASS_CHECK(gemm.can_implement(arguments));
+
+ // Initialize CUTLASS kernel with arguments and workspace pointer
+ CUTLASS_CHECK(gemm.initialize(arguments, workspace.get()));
+
+ // Correctness / Warmup iteration
+ CUTLASS_CHECK(gemm.run());
+
+ cudaDeviceSynchronize();
+
+ // Check if output from CUTLASS kernel and reference kernel are equal or not
+ Result result;
+ result.passed = verify(options);
+
+ std::cout << " Disposition: " << (result.passed ? "Passed" : "Failed") << std::endl;
+
+ if (!result.passed) {
+ exit(-1);
+ }
+
+ // Run profiling loop
+ if (options.iterations > 0)
+ {
+ GpuTimer timer;
+ timer.start();
+ for (int iter = 0; iter < options.iterations; ++iter) {
+ CUTLASS_CHECK(gemm.initialize(arguments, workspace.get()));
+ CUTLASS_CHECK(gemm.run());
+ }
+ timer.stop();
+
+ // Compute average runtime and GFLOPs.
+ float elapsed_ms = timer.elapsed_millis();
+ result.avg_runtime_ms = double(elapsed_ms) / double(options.iterations);
+ result.gflops = options.gflops(result.avg_runtime_ms / 1000.0);
+
+
+ std::cout << " Problem Size: " << options.m << 'x' << options.n << 'x' << options.k << std::endl;
+ std::cout << " Avg runtime: " << result.avg_runtime_ms << " ms" << std::endl;
+ std::cout << " GFLOPS: " << result.gflops << std::endl;
+ }
+
+ return 0;
+}
+
+#endif // defined(CUTLASS_ARCH_MMA_SM100_SUPPORTED)
+
+///////////////////////////////////////////////////////////////////////////////////////////////////
+
+int main(int argc, char const **args) {
+
+ // CUTLASS must be compiled with CUDA 12.8 or higher Toolkit to run this example
+ // and must have compute capability at least 100.
+ if (__CUDACC_VER_MAJOR__ < 12 || (__CUDACC_VER_MAJOR__ == 12 && __CUDACC_VER_MINOR__ < 8)) {
+ std::cerr << "This example requires CUDA 12.8 or newer." << std::endl;
+ // Returning zero so this test passes on older Toolkits. Its actions are no-op.
+ return 0;
+ }
+
+ cudaDeviceProp props;
+ int current_device_id;
+ CUDA_CHECK(cudaGetDevice(¤t_device_id));
+
+ CUDA_CHECK(cudaGetDeviceProperties(&props, current_device_id));
+
+ if (!(props.major == 10 && props.minor == 0)) {
+ std::cerr << "This example requires a GPU of NVIDIA's Blackwell architecture (compute capability 100)." << std::endl;
+ return 0;
+ }
+
+ //
+ // Parse options
+ //
+
+ Options options;
+
+ options.parse(argc, args);
+
+ if (options.help) {
+ options.print_usage(std::cout) << std::endl;
+ return 0;
+ }
+
+ //
+ // Evaluate CUTLASS kernels
+ //
+#if defined(CUTLASS_ARCH_MMA_SM100_SUPPORTED)
+ run(options);
+#endif // defined(CUTLASS_ARCH_MMA_SM100_SUPPORTED)
+
+ return 0;
+}
+
+/////////////////////////////////////////////////////////////////////////////////////////////////
diff --git a/examples/72_blackwell_narrow_precision_gemm/CMakeLists.txt b/examples/72_blackwell_narrow_precision_gemm/CMakeLists.txt
new file mode 100644
index 00000000..fa80c184
--- /dev/null
+++ b/examples/72_blackwell_narrow_precision_gemm/CMakeLists.txt
@@ -0,0 +1,46 @@
+
+# Copyright (c) 2025 - 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-License-Identifier: BSD-3-Clause
+#
+# Redistribution and use in source and binary forms, with or without
+# modification, are permitted provided that the following conditions are met:
+#
+# 1. Redistributions of source code must retain the above copyright notice, this
+# list of conditions and the following disclaimer.
+#
+# 2. Redistributions in binary form must reproduce the above copyright notice,
+# this list of conditions and the following disclaimer in the documentation
+# and/or other materials provided with the distribution.
+#
+# 3. Neither the name of the copyright holder nor the names of its
+# contributors may be used to endorse or promote products derived from
+# this software without specific prior written permission.
+#
+# THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+# AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+# IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+# DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+# FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+# DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+# SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+# CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+# OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+# OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+
+if(NOT CUTLASS_NVCC_ARCHS STREQUAL "100")
+cutlass_example_add_executable(
+ 72a_blackwell_nvfp4_bf16_gemm
+ 72a_blackwell_nvfp4_bf16_gemm.cu
+ )
+
+cutlass_example_add_executable(
+ 72b_blackwell_nvfp4_nvfp4_gemm
+ 72b_blackwell_nvfp4_nvfp4_gemm.cu
+ )
+
+cutlass_example_add_executable(
+ 72c_blackwell_mixed_mxfp8_bf16_gemm
+ 72c_blackwell_mixed_mxfp8_bf16_gemm.cu
+ )
+endif()
diff --git a/examples/73_blackwell_gemm_preferred_cluster/CMakeLists.txt b/examples/73_blackwell_gemm_preferred_cluster/CMakeLists.txt
new file mode 100644
index 00000000..0d0f7757
--- /dev/null
+++ b/examples/73_blackwell_gemm_preferred_cluster/CMakeLists.txt
@@ -0,0 +1,36 @@
+# Copyright (c) 2025 - 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-License-Identifier: BSD-3-Clause
+#
+# Redistribution and use in source and binary forms, with or without
+# modification, are permitted provided that the following conditions are met:
+#
+# 1. Redistributions of source code must retain the above copyright notice, this
+# list of conditions and the following disclaimer.
+#
+# 2. Redistributions in binary form must reproduce the above copyright notice,
+# this list of conditions and the following disclaimer in the documentation
+# and/or other materials provided with the distribution.
+#
+# 3. Neither the name of the copyright holder nor the names of its
+# contributors may be used to endorse or promote products derived from
+# this software without specific prior written permission.
+#
+# THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+# AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+# IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+# DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+# FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+# DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+# SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+# CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+# OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+# OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+
+
+if(NOT CUTLASS_NVCC_ARCHS STREQUAL "100")
+cutlass_example_add_executable(
+ 73_blackwell_gemm_preferred_cluster
+ blackwell_gemm_preferred_cluster.cu
+ )
+endif()
diff --git a/examples/73_blackwell_gemm_preferred_cluster/blackwell_gemm_preferred_cluster.cu b/examples/73_blackwell_gemm_preferred_cluster/blackwell_gemm_preferred_cluster.cu
new file mode 100644
index 00000000..fb62e844
--- /dev/null
+++ b/examples/73_blackwell_gemm_preferred_cluster/blackwell_gemm_preferred_cluster.cu
@@ -0,0 +1,541 @@
+/***************************************************************************************************
+ * Copyright (c) 2025 - 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: BSD-3-Clause
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ *
+ * 1. Redistributions of source code must retain the above copyright notice, this
+ * list of conditions and the following disclaimer.
+ *
+ * 2. Redistributions in binary form must reproduce the above copyright notice,
+ * this list of conditions and the following disclaimer in the documentation
+ * and/or other materials provided with the distribution.
+ *
+ * 3. Neither the name of the copyright holder nor the names of its
+ * contributors may be used to endorse or promote products derived from
+ * this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+ * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ * DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+ * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+ * DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+ * SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+ * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+ * OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+ * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ *
+ **************************************************************************************************/
+
+/*! \file
+ \brief A GEMM example using CUTLASS for the NVIDIA Blackwell SM100 architecture with preferred cluster.
+
+ With the introduction of NVIDIA Compute Capability 9.0, the CUDA programming model introduced
+ an optional hierarchy level known as Thread Block Clusters, which consist of multiple Thread Blocks.
+ While the CUDA programming model has supported the specification of cluster shapes at runtime
+ (Dynamic Clusters) since the Hopper architecture, CUTLASS has only provided support for Static
+ Clusters, meaning that cluster shapes must be defined at compile time.
+
+ Larger cluster shapes can achieve higher TMA multicast but may result in poor SM occupancy due
+ to quantization. For instance, a 2x2 cluster on an 18 SM GPU would only utilize 16 SMs, leaving
+ 2 SMs idle.
+
+ Starting with Compute Capability 10.0, the CUDA programming model adds the ability to specify
+ two clusters: preferred cluster and fallback cluster. For brevity, we refer to this as
+ Preferred Clusters. In the previous example, users can now launch an additional 2x1 cluster to
+ utilize the 2 idle SMs.
+
+ With CUTLASS 3.8, in addition to Dynamic Clusters, CUTLASS adds support for Preferred Dynamic Cluster,
+ the ability for users to specify two clusters shapes at runtime.
+
+ Terminology
+ * Static cluster: cluster shape is specified at compile time.
+ * Dynamic cluster: cluster shape is specified at runtime and set by the host.
+ * Preferred cluster: Kernel can be launched with two cluster shapes (preferred and fallback).
+
+ Preferred and fallback cluster shapes are subject to several constraints.
+ * Preferred cluster depth (Z dimension) must be the same as that of fallback cluster.
+ * Fallback cluster shape must evenly divide the preferred cluster shape.
+ * Preferred cluster shape must evenly divide the kernel launch grid shape.
+
+ This example demonstrates how to use the Dynamic Clusters and Preferred Clusters features in
+ CUTLASS 3.x Blackwell SM100 kernels. Users can specify preferred and fallback cluster shapes via GEMM arguments.
+
+ # Example:
+ ./73_blackwell_gemm_preferred_cluster" --m=4096 --n=4096 --k=4096 --preferred_cluster_m=4 --preferred_cluster_n=4 --fallback_cluster_m=2 --fallback_cluster_m=1
+*/
+
+
+
+#include
+#include
+#include
+#include
+
+#include "cutlass/cutlass.h"
+
+#include "cute/tensor.hpp"
+#include "cutlass/tensor_ref.h"
+#include "cutlass/epilogue/thread/linear_combination.h"
+#include "cutlass/gemm/dispatch_policy.hpp"
+#include "cutlass/gemm/collective/collective_builder.hpp"
+#include "cutlass/epilogue/collective/collective_builder.hpp"
+#include "cutlass/gemm/device/gemm_universal_adapter.h"
+#include "cutlass/gemm/kernel/gemm_universal.hpp"
+#include "cutlass/gemm/kernel/tile_scheduler_params.h"
+
+#include "cutlass/util/command_line.h"
+#include "cutlass/util/distribution.h"
+#include "cutlass/util/host_tensor.h"
+#include "cutlass/util/packed_stride.hpp"
+#include "cutlass/util/tensor_view_io.h"
+#include "cutlass/util/reference/device/gemm.h"
+#include "cutlass/util/reference/device/tensor_compare.h"
+#include "cutlass/util/reference/device/tensor_fill.h"
+
+#include "helper.h"
+
+using namespace cute;
+
+#if defined(CUTLASS_ARCH_MMA_SM100_SUPPORTED)
+
+
+/////////////////////////////////////////////////////////////////////////////////////////////////
+/// GEMM kernel configurations
+/////////////////////////////////////////////////////////////////////////////////////////////////
+
+// A matrix configuration
+using ElementA = half_t; // Element type for A matrix operand
+using LayoutA = cutlass::layout::RowMajor; // Layout type for A matrix operand
+constexpr int AlignmentA = 128 / cutlass::sizeof_bits::value; // Memory access granularity/alignment of A matrix in units of elements (up to 16 bytes)
+
+// B matrix configuration
+using ElementB = half_t; // Element type for B matrix operand
+using LayoutB = cutlass::layout::ColumnMajor; // Layout type for B matrix operand
+constexpr int AlignmentB = 128 / cutlass::sizeof_bits::value; // Memory access granularity/alignment of B matrix in units of elements (up to 16 bytes)
+
+// C/D matrix configuration
+using ElementC = float; // Element type for C and D matrix operands
+using LayoutC = cutlass::layout::ColumnMajor; // Layout type for C and D matrix operands
+constexpr int AlignmentC = 128 / cutlass::sizeof_bits::value; // Memory access granularity/alignment of C matrix in units of elements (up to 16 bytes)
+
+// Kernel functional config
+using ElementAccumulator = float; // Element type for internal accumulation
+using ArchTag = cutlass::arch::Sm100; // Tag indicating the minimum SM that supports the intended feature
+using OperatorClass = cutlass::arch::OpClassTensorOp; // Operator class tag
+
+// MMA and Cluster Tile Shapes
+// Shape of the tile computed by tcgen05 MMA, could be across 2 SMs if Cluster Shape % 2 == 0
+using MmaTileShape_MNK = Shape<_256,_128,_64>;
+// Shape of the threadblocks participating in a tcgen05 MMA. <1, 1, 1> for cta_group = 1, <2, 1, 1> for cta_group = 2
+using AtomThrShape_MNK = Shape<_2, _1, _1>;
+// Shape of the tile computed by each SM
+using PerSmTileShape_MNK = decltype(shape_div(MmaTileShape_MNK{}, AtomThrShape_MNK{}));
+// Shape of the cluster set to to indicate dynamic cluster shape
+using ClusterShape_MNK = Shape;
+// When dynamic cluster is used, KernelScheduleAuto always selects mainloop dispatch policy that
+// lowers to tcgen05 MMA cta_group = 1 as we don't know if the dynamic cluster M dimension will be a multiple of 2
+// To use KernelScheduleAuto, users need to set AtomThrShape_MNK to Shape<1, 1, 1>
+using KernelSchedule = cute::conditional_t;
+
+using CollectiveEpilogue = typename cutlass::epilogue::collective::CollectiveBuilder<
+ ArchTag, OperatorClass,
+ PerSmTileShape_MNK, ClusterShape_MNK,
+ cutlass::epilogue::collective::EpilogueTileAuto,
+ ElementAccumulator, ElementAccumulator,
+ ElementC, LayoutC, AlignmentC,
+ ElementC, LayoutC, AlignmentC,
+ cutlass::epilogue::collective::EpilogueScheduleAuto
+ >::CollectiveOp;
+
+using CollectiveMainloop = typename cutlass::gemm::collective::CollectiveBuilder<
+ ArchTag, OperatorClass,
+ ElementA, LayoutA, AlignmentA,
+ ElementB, LayoutB, AlignmentB,
+ ElementAccumulator,
+ MmaTileShape_MNK, ClusterShape_MNK,
+ cutlass::gemm::collective::StageCountAutoCarveout(sizeof(typename CollectiveEpilogue::SharedStorage))>,
+ KernelSchedule
+ >::CollectiveOp;
+
+using GemmKernel = cutlass::gemm::kernel::GemmUniversal<
+ Shape, // Indicates ProblemShape
+ CollectiveMainloop,
+ CollectiveEpilogue,
+ void // <--- Default to cluster launch control (CLC) scheduler
+>;
+
+using Gemm = cutlass::gemm::device::GemmUniversalAdapter;
+
+// Reference device GEMM implementation type
+using DeviceGemmReference = cutlass::reference::device::Gemm<
+ ElementA,
+ LayoutA,
+ ElementB,
+ LayoutB,
+ ElementC,
+ LayoutC,
+ ElementAccumulator,
+ ElementAccumulator>;
+
+using StrideA = typename Gemm::GemmKernel::StrideA;
+using StrideB = typename Gemm::GemmKernel::StrideB;
+using StrideC = typename Gemm::GemmKernel::StrideC;
+using StrideD = typename Gemm::GemmKernel::StrideD;
+
+//
+// Data members
+//
+
+/// Initialization
+StrideA stride_A;
+StrideB stride_B;
+StrideC stride_C;
+StrideD stride_D;
+uint64_t seed;
+
+cutlass::DeviceAllocation block_A;
+cutlass::DeviceAllocation block_B;
+cutlass::DeviceAllocation block_C;
+cutlass::DeviceAllocation block_D;
+cutlass::DeviceAllocation block_ref_D;
+
+#endif // defined(CUTLASS_ARCH_MMA_SM100_SUPPORTED)
+
+/////////////////////////////////////////////////////////////////////////////////////////////////
+/// Testbed utility types
+/////////////////////////////////////////////////////////////////////////////////////////////////
+
+// Command line options parsing
+struct Options {
+
+ bool help;
+
+ float alpha, beta;
+ int iterations;
+ int m, n, k;
+ int preferred_cluster_m, preferred_cluster_n, fallback_cluster_m, fallback_cluster_n;
+
+ Options():
+ help(false),
+ m(4096), n(4096), k(4096),
+ alpha(1.f), beta(0.f),
+ iterations(10),
+ preferred_cluster_m(4),
+ preferred_cluster_n(4),
+ fallback_cluster_m(2),
+ fallback_cluster_n(1)
+ { }
+
+ // Parses the command line
+ void parse(int argc, char const **args) {
+ cutlass::CommandLine cmd(argc, args);
+
+ if (cmd.check_cmd_line_flag("help")) {
+ help = true;
+ return;
+ }
+
+ cmd.get_cmd_line_argument("m", m);
+ cmd.get_cmd_line_argument("n", n);
+ cmd.get_cmd_line_argument("k", k);
+ cmd.get_cmd_line_argument("alpha", alpha, 1.f);
+ cmd.get_cmd_line_argument("beta", beta, 0.f);
+ cmd.get_cmd_line_argument("iterations", iterations);
+ cmd.get_cmd_line_argument("preferred_cluster_m", preferred_cluster_m, 4);
+ cmd.get_cmd_line_argument("preferred_cluster_n", preferred_cluster_n, 4);
+ cmd.get_cmd_line_argument("fallback_cluster_m", fallback_cluster_m, 2);
+ cmd.get_cmd_line_argument("fallback_cluster_n", fallback_cluster_n, 1);
+
+ if (!validate_cluster_shape()){
+ std::cout << "--Invalid cluster shapes" << std::endl;
+ help = true;
+ return;
+ }
+ }
+
+ /// Prints the usage statement.
+ std::ostream & print_usage(std::ostream &out) const {
+
+ out << "73_blackwell_gemm_preferred_cluster\n\n"
+ << " Blackwell FP16 GEMM using preferred cluster.\n\n"
+ << "Options:\n\n"
+ << " --help If specified, displays this usage statement\n\n"
+ << " --m= Sets the M extent of the GEMM\n"
+ << " --n= Sets the N extent of the GEMM\n"
+ << " --k= Sets the K extent of the GEMM\n"
+ << " --alpha= Epilogue scalar alpha\n"
+ << " --beta= Epilogue scalar beta\n"
+ << " --preferred_cluster_m= Sets the M extent of preferred cluster shape\n"
+ << " --preferred_cluster_n= Sets the N extent of preferred cluster shape\n"
+ << " --fallback_cluster_m= Sets the M extent of fallback cluster shape\n"
+ << " --fallback_cluster_n=