cutlass Documentation

Technical Documentation: NVIDIA/cutlass

ℹ️ Provenance: Multi-Tier Ingestion · CodeWiki Reference

Overview (README.md) Copy Section

Overview

CUTLASS 4.8.1

CUTLASS 4.8.1 - Oct. 2026

CUTLASS is a collection of abstractions for implementing high-performance matrix-matrix multiplication (GEMM)
and related computations at all levels and scales within CUDA. It incorporates strategies for
hierarchical decomposition and data movement. CUTLASS decomposes these "moving parts" into reusable, modular
software components and abstractions.

Primitives for different levels of a conceptual parallelization hierarchy can be specialized and tuned
via custom tiling sizes, data types, and other algorithmic policy. The resulting flexibility simplifies
their use as building blocks within custom kernels and applications.

CUTLASS has been providing CUDA C++ template abstractions for high-performance linear algebra since 2017 and
these abstractions provide extensive support for a wide range of computations including
mixed-precision computations, specialized data-movement (async copy) and
multiply-accumulate abstractions for FP64, FP32, TF32, FP16, BF16,
FP32 emulation via tensor core instruction,
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) across NVIDIA's Volta, Turing, Ampere, Ada, Hopper, and Blackwell architectures.

To this rich ecosystem of C++ based kernel programming abstractions, CUTLASS 4 adds CUTLASS DSLs. These are Python native interfaces for writing high-performance CUDA kernels based on core CUTLASS and CuTe concepts without any performance compromises. This allows for a much smoother learning curve, orders of magnitude faster compile times, native integration with DL frameworks without writing glue code, and much more intuitive metaprogramming that does not require deep C++ expertise.

Overall, we envision CUTLASS DSLs as a family of domain-specific languages (DSLs). With the release of 4.0, we released the first of these in CuTe DSL. This is a low-level programming model that is fully consistent with CuTe C++ abstractions - exposing core concepts such as layouts, tensors, hardware atoms, and full control over the hardware thread and data hierarchy.

CuTe DSL demonstrates optimal matrix multiply and other linear algebra operations
targeting the programmable, high-throughput Tensor Cores implemented by
NVIDIA's Ampere, Hopper, and Blackwell architectures.

We believe it will become an indispensable tool for students, researchers, and performance
engineers alike - flattening the learning curve of GPU programming, rapidly prototyping kernel
designs, and bringing optimized solutions into production.

CuTe DSL is currently in public beta and will graduate out of beta by end of summer 2026.

To get started quickly - please refer :

• CUTLASS C++ Quick Start Guide.
• CuTe DSL Quick Start Guide.

What's New in CUTLASS 4.8

CuTe DSL Copy Chapter

• New features

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

    • CuTe DSL and CuTe extensions
      • Support for higher-throughput FP8 (MMA_K=64) and FP4 (MMA_K=128) Tensor Core MMA instructions
      • B collector reuse
      • Extended TMEM size from 512 COL to 576 COL
      • Larger shared memory allocations (328KB)
      • Enhanced mixed precision throughput (FP8/FP4)
      • Softmax acceleration related features
    • Primitives
      • Support for higher-throughput FP8 (MMA_K=64) and FP4 (MMA_K=128) Tensor Core MMA instructions
      • B collector reuse
      • Extended TMEM size from 512 COL to 576 COL
      • Larger shared memory allocations (328KB)
      • Enhanced mixed precision throughput (FP8/FP4)
      • Softmax acceleration related feature
      • 2:4 sparsity support for FP4

  • CuTe DSL extensions has several new features:

    • CTA-V maps are now inferred automatically for cute_ext TMA load, store, multicast, and reduce-store operations. Explicit CTA-V maps remain supported as overrides.
    • Added asynchronous atomic TMA reduce-store and sparse MMA operations.
    • Added reusable cute_ext GEMM mainloop and TMA epilogue helpers.
    • Added opt-in TMEM accumulator-buffer planning, including overlapping ping-pong storage for capacity-constrained kernels.
    • Improved device-side TMA descriptor updates and grouped GEMM performance through SMEM-staged updates, workspace reuse, and reduced prologue and synchronization overhead.

  • This release includes an opt-in preview of the CuTe DSL extensions (cute_ext) compiler pipeline for ordinary Cute DSL kernels. This pipeline lets user mix cute_ext APIs directly into @cute.jit and @cute.kernel code and is required for kernels that mix the two API surfaces. You may test this feature with the following:
    CUTE_DSL_USE_EXTENSION_COMPILER=1 python your_program.py
    The pipeline is expected to preserve program behavior and performance, but generated PTX/SASS may differ. Note that this pipeline will become the default in the future, no earlier than 4.10.

  • Added examples for better control over Primitives' compiler warnings/errors introduced in 4.7.0. See the CuTeDSL/experimental/compiler_diagnostic/ directory.

  • IKET Profiler Tool

    • Rubin kernels (sm107) can now be profiled.
    • It is now possible to only dump timing data for a specific cluster to reduce profiling overhead. Previously all clusters were profiled.
    • Task Scheduling can instrument the schedule with IKET ranges when constructing TaskManager objects (iket_enable_profiling=True). Task execution will generate an IKET range and individual pipeline stages in a schedule may generate separate ranges (iket_profiling_stages).

  • A number of new examples were added in this release:

    • Rubin (CuTe):
      • Dense GEMM for legacy data type with B collector reuse as applicable
      • Grouped GEMM with B collector reuse
      • Dense blockscaled GEMM with FP4/FP6/FP8 mixed precision and UE5M3 / block-32 scale-factor support
      • Grouped blockscaled GEMM with B collector reuse as applicable
      • Blockwise GEMM
    • Rubin (CuTe extension):
      • Support for higher-throughput FP8 (MMA_K=64) and FP4 (MMA_K=128) blockscaled GEMM with UE5M3 scale-factor
      • Grouped GEMM with B collector reuse
    • Blackwell (CuTe extension):
      • Dense GEMMs
        • Back-to-back GEMM
        • Blockscaled GEMM
        • Persistent GEMM with alpha/beta scaling
        • CLC scheduler/dynamic persistent GEMM
        • GLU GEMM
        • Mixed input GEMM
        • Planar complex GEMM
        • Input transform GEMM
        • GeForce pingpong dense GEMM
        • Blackwell Ultra blockscaled GEMM
      • Dense Convolutions
        • Implicit-Gemm Fprop Conv
        • Blocksclaed Implicit-Gemm Fprop Conv
        • GeForce Implicit-Gemm Fprop Conv
        • GeForce Blockscaled Implicit-Gemm Fporp Conv
      • Attention
        • GQA Decode
      • Grouped GEMM
        • Unscaled and blockscaled grouped GEMM
      • Top-K
    • Ampere (CuTe extension):
      • SIMT GEMM

  • CuTe DSL now supports x86_64 Windows

  • CuTe DSL AoT now supports new host target: QNX8.0

  • Notebooks are restructured under examples/python/CuTeDSL/cute/notebooks and new notebooks for primitives will be added under examples/python/CuTeDSL/notebooks

  • Numpy is now not a default dependency

• Bug fixes and improvements:

  • nvidia-cuda-nvdisasm is now an optional dependency of nvidia-cutlass-dsl via the optional [sass] extra. SASS dumping (CUTE_DSL_KEEP=sass / KeepSASS) now resolves nvdisasm from the bundled wheel (recommended since its version matches the DSL toolchain) or from a local CUDA Toolkit (CUDA_HOME/CUDA_PATH). A locally-provided nvdisasm must come from a CUDA Toolkit at least as new as the toolchain that produced the CUBIN. Installations that never dump SASS are unaffected.
  • Reduced the protobuf version requirement of IKET profiler from 6.30 to 4.21. This should make protobuf an easier requirement to satisfy in preparation for transioning IKET to an optional extra.
  • Improved JAX PyTree input/output aliasing for cutlass.jax.cutlass_call
  • Fixed a regression from 4.6.0 where cute.autovec_copy emitted per-element instead of
    vectorized instructions for tensors with a dynamic stride (!3463)
  • Fixed TVM-FFI env stream detection for GPU tensors in tuple
    (!3444)
  • Fixed GPU link-libraries compile-option order so it is stable across processes
    (!3564)
  • Fixed preprocessor IndexError on staged bool() with no arguments
    (!3506)
  • Rejected cute.compile on @cute.kernel with a user error instead of an ICE
    (!3429)
  • Fixed CuTe DSL crashing the Python interpreter when used in a REPL
    (!3413)
  • Fixed a cuDNN Frontend FROST SDPA backward compilation failure issue (!3594)
  • Fixed SIGABRTs in TVM-FFI launch for cuDNN Frontend SM100 ragged SDPA kernel (!3595)
  • Fixed an issue with cute.extern where functions with zero arguments would generate invalid type signatures. Compiling such functions a second time would fail.
  • Clear error/warning are now emitted for CopyUniversalOp with memory attributes specified (!3711).
  • CUTLASS Primitives' tensormap_replace now keeps device-side TMA descriptor updates safe when a descriptor is shrunk below its original extent (e.g. grouped GEMM) and fixes a possible illegal memory access on Blackwell/Rubin. This behavior can be opted out by specifying keep_oob_fetch_window=True. See https://docs.nvidia.com/datacenter/tesla/tesla-release-notes-595-91-07/index.html#known-issues for more details on the known cuTensorMapEncodeTiled bug.
  • Fixed unintended sys.path modification that could break imports in some environments where Torch was not installed e.g. JAX.

• Known issues

  • Due to a bug in CTK, the use of some packed instructions result in invalid code generation. Given the nature of the instructions and how they are explicitly used, we have chosen to not silently rewrite these, but rather are flagging until the upstream compiler fixes are available.
  • Avoid packed higher-throughput FP8 (MMA_K=128) subtraction via inline PTX sub.rn.e4m3x4.e4m3x4 or sub.rn.e5m2x4.e5m2x4, as incorrect code generation can produce wrong results.
  • Some CUTLASS Primitives packed mixed-input ops compile incorrectly on Rubin (sm107a/sm107f), resulting in wrong code generation. The specific patterns include:
    • Packed mul, sub, and add ops with bf16x2 or fp16x2 result type and f32x2 input type when result is converted to fp32 later. APIs affected: add_packed_f16x2_f32x2_f32x2, add_packed_bf16x2_f32x2_f32x2, mul_packed_f16x2_f32x2_f32x2, mul_packed_bf16x2_f32x2_f32x2, sub_packed_f16x2_f32x2_f32x2, sub_packed_bf16x2_f32x2_f32x2
      • For example:
        text Copy

        x = prims.mul_packed_bf16x2_f32x2_f32x2(a, b)
        out[0], out[1] = f32(x[0]), f32(x[1])  # out[0], out[1] are wrong

    • Packed mul op with bf16x2 or fp16x2 result type and mixed bf16x2 and fp16x2 input types when the first operand is a compilation-time constant. APIs affected: mul_packed_bf16x2_bf16x2_f16x2, mul_packed_f16x2_f16x2_bf16x2
      • For example:
        text Copy

        x = prims.mul_packed_bf16x2_bf16x2_f16x2(2.0, b) # x is wrong
        x = prims.mul_packed_f16x2_f16x2_bf16x2(2.0, b)  # x is wrong

This release has been tested against the following packages:

• FlashAttention: main (1fc0221)
• Quack: main (35266c32)
• FlashInfer: main (399b6c1)
• cuDNN-Frontend: deveop (986ef4d)
• Pytorch: main (ca883d5)
• TensorRT-LLM: main (cc63e66)

CUTLASS Operator API Copy Chapter

• Dense and blockscaled GEMMs in Operator API have preliminary Rubin support. These are provided as a preview and may need additional performance tuning.

  • Updated GEMMs include:

    • Dense GEMMs: FP8xFP8
    • Blockscaled GEMM: {MXFP8}x{MXFP4, MXFP8} and {MXFP4, NVFP4}x{MXFP4, NVFP4} (including support for the new UE5M3 scale factor dtype for NVFP4).

  • These kernels utilize the below new features in Rubin:

    • Higher SMEM (328KB) and TMEM capacity (288KB)
    • B-buffer reuse
    • Enhanced mixed precision throughput

• Operators can now be ranked by their estimated performance when nvMatmulHeuristics is available. See tutorial here. NOTE: This currently only supports Blackwell kernels as nvMatmulHeuristics does not yet support Rubin.

• Standalone kernel implementations are now exposed through cutlass.kernels. These kernels can be used directly, in addition to being discoverable and usable via the Operator interface in cutlass.operators.

• Custom Epilogue fusions now support per-row or per-column reductions.

• IndexPtrGroupedGemmArguments is now used to represent Grouped GEMM with contiguous-offset/index-pointers. Existing GroupedGemmArguments is deprecated and will be removed in a future release.

C++ Copy Chapter

• Added initial Rubin support (SM107) with CuTe C++ building blocks:

• Rubin Tensor Core MMA instructions and corresponding CuTe MMA traits.

• CuTe examples that demonstrate the use of Rubin SM107 Tensor Core instructions:

• Dense FP8 GEMM.
• Block-scaled FP8 GEMM.
• Mixed-precision block-scaled FP8/FP4 GEMM.
• Block-scaled FP4 GEMM.

• Adjusted shared-memory and tensor-memory capacity handling for Rubin SM107:

• Set the SM107 shared-memory capacity to 327 KiB and added launch support for oversized shared-memory configurations.
• Set the SM107 TMEM capacity to 576 columns per SM, updated the CuTe 1SM and 2SM TMEM allocators for Rubin's exclusive allocation path.

• Enabled the existing SM100-compatible GEMM and convolution for the new SM107 sm_107a and sm_107f targets:

• Set of unit tests for Rubin SM107 SIMT GEMM, dense FP8 GEMM, block-scaled FP8 GEMM, block-scaled FP4 GEMM, and mixed-precision, complex, and 9xBF16 GEMM

Note: Executing Rubin kernels (SM107) requires the R615 driver which will be released
with CUDA Toolkit 13.4 GA. R610 from CUDA Toolkit 13.4 Developer Preview is not sufficient.

See the CHANGELOG for details of all past releases and updates.

Performance

CUTLASS primitives are very efficient. When used to construct device-wide GEMM kernels,
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 (NVIDIA Hopper architecture) since
CUTLASS 3.1.
CUTLASS 3.5.1 was compiled with the CUDA 12.5u1 Toolkit.
Tensor Core operations are implemented using CUDA's
mma and
wgmma instructions.

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.

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.8 Toolkit.
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 Copy Chapter

We have tested the following environments.

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 down on Windows platforms for all CUDA toolkits.
CUTLASS team is working on a fix.

Hardware Copy Chapter

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.

Target Architecture Copy Chapter

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.

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"

Or

cmake .. -DCUTLASS_NVCC_ARCHS="100a"

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 (SM120). 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
for details on which kernels require which target architectures.

Documentation

CUTLASS is described in the following documents and the accompanying
Doxygen documentation.

• Quick Start Guide - basics of building and running CUTLASS
• Functionality - summarizes functionality available in CUTLASS
• Efficient GEMM in CUDA - describes how GEMM kernels may be implemented efficiently in CUDA
• CUTLASS 3.x Design - describes the CUTLASS 3.x design, its benefits, and how CuTe enables us to write much more composable components
• GEMM API 3.x - describes the CUTLASS 3.x GEMM model and C++ template concepts
• GEMM API 2.x - describes the CUTLASS 2.x GEMM model and C++ template concepts
• Implicit GEMM Convolution - describes 2-D and 3-D convolution in CUTLASS
• Code Organization - describes the organization and contents of the CUTLASS project
• Terminology - describes terms used in the code
• Programming Guidelines - guidelines for writing efficient modern CUDA C++
• Fundamental types - describes basic C++ classes used in CUTLASS to represent numeric quantities and arrays
• Layouts - describes layouts of matrices and tensors in memory
• Tile Iterators - describes C++ concepts for iterating over tiles of matrices in memory
• CUTLASS Profiler - command-line driven profiling application
• CUTLASS Utilities - additional templates used to facilitate rapid development
• Dependent kernel launch - describes a new feature in Hopper which allows overlapping dependent
  kernels in the same stream, and how it is used in CUTLASS.

Resources

We have also described the structure of an efficient GEMM in our talk at the
GPU Technology Conference 2018.

• CUTLASS: Software Primitives for Dense Linear Algebra at All Levels and Scales within CUDA
• Developing CUDA Kernels to Push Tensor Cores to the Absolute Limit on NVIDIA A100
• Accelerating Convolution with Tensor Cores in CUTLASS
• Accelerating Backward Data Gradient by Increasing Tensor Core Utilization in CUTLASS
• CUTLASS: Python API, Enhancements, and NVIDIA Hopper

Building CUTLASS

CUTLASS is a header-only template library and does not need to be built to be used by other
projects. Client applications should target CUTLASS's include/ directory in their include
paths.

CUTLASS unit tests, examples, and utilities can be build with CMake.
The minimum version of CMake is given in the Quickstart guide.
Make sure the CUDACXX environment variable points to NVCC in the CUDA Toolkit installed
on your system.

$ export CUDACXX=${CUDA_INSTALL_PATH}/bin/nvcc

Create a build directory within the CUTLASS project, then run CMake. By default CUTLASS will build kernels
for CUDA architecture versions 5.0, 6.0, 6.1, 7.0, 7.5, 8.0, 8.6, 8.9, and 9.0.
To reduce compile time you can specify
the architectures to build CUTLASS for by changing the CMake configuration setting
CUTLASS_NVCC_ARCHS.

$ mkdir build && cd build

$ cmake .. -DCUTLASS_NVCC_ARCHS=80               # compiles for NVIDIA's Ampere Architecture

From the build/ directory, compile and run the CUTLASS unit tests by building the target test_unit with make.

The unit tests are organized as several binaries mirroring the top-level namespaces of CUTLASS,
and they may be executed in parallel via make's -j command line argument.

$ make test_unit -j
...
...
...
[----------] Global test environment tear-down
[==========] 946 tests from 57 test cases ran. (10812 ms total)
[  PASSED  ] 946 tests.

All tests should pass on supported platforms, though the exact number of tests may vary over time.

Project Structure

CUTLASS is arranged as a header-only library along with Utilities, Tools, Examples, and unit tests.
Doxygen documentation provides a complete list of files, classes,
and template concepts defined in the CUTLASS project.

A detailed explanation of the source code organization may be found in the
CUTLASS documentation, but several main components are summarized below.

CUTLASS Template Library Copy Chapter

/* Detailed source-code truncated for AI context efficiency. */

CUTLASS SDK Examples Copy Section

CUTLASS SDK examples apply CUTLASS templates to implement basic computations.

Tools Copy Section

tools/
  library/                   # CUTLASS Instance Library - contains instantiations of all supported CUTLASS templates
    include/
      cutlass/
        library/

  profiler/                  # CUTLASS Profiler         - command-line utility for executing operations in the
                             #                            CUTLASS Library

  util/                      # CUTLASS Utilities        - contains numerous helper classes for
    include/                 #                            managing tensors in device memory, reference
      cutlass/               #                            implementations for GEMM, random initialization
        util/                #                            of tensors, and I/O.

Test Copy Section

The test/unit/ directory consist of unit tests implemented with Google Test that demonstrate
basic usage of Core API components and complete tests of the CUTLASS GEMM computations.

Instructions for building and running the Unit tests are described in the Quickstart guide.

Performance Profiling

The tools/profiler/ directory contains a command-line utility for launching each of the GEMM kernels.
It can be built as follows:

$ make cutlass_profiler -j16

Building all GEMM and Convolution kernels (_long_ build times) Copy Chapter

By default, only one tile size is instantiated for each data type, math instruction, and layout.
To instantiate all, set the following environment variable when running CMake from an empty build/ directory.
Beware, this results in tens of thousands of kernels and long build times.
This would also result in a large binary size and on some platforms linker to fail on building the library.
Therefore, it's highly recommended to generate only a subset of kernels as demonstrated in the sub-section below.

$ cmake .. -DCUTLASS_NVCC_ARCHS=90a -DCUTLASS_LIBRARY_KERNELS=all
...
$ make cutlass_profiler -j16

Building a subset of GEMM and Convolution kernels (_reduced_ build times) Copy Chapter

To compile strictly one kernel or a small set of kernels, a comma-delimited list of kernel names with
wildcard characters may be used to reduce the set of kernels. The following examples show building exactly one
or a subset of kernels for NVIDIA Ampere and Turing architecture:

Building a subset Tensor Core GEMM kernels Copy Section

To compile a subset of Tensor Core GEMM kernels with FP32 accumulation and FP16 input targeting NVIDIA Ampere and Turing architecture,
use the below cmake command line:

$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_tensorop_s*gemm_f16_*_nt_align8
...
$ make cutlass_profiler -j16

Example command line for profiling a subset of Tensor Core GEMM kernels is as follows:

./tools/profiler/cutlass_profiler --kernels=cutlass_tensorop_s*gemm_f16_*_nt_align8 --m=3456 --n=4096 --k=4096

...
=============================
  Problem ID: 1

        Provider: CUTLASS
   OperationKind: gemm
       Operation: cutlass_tensorop_s1688gemm_f16_256x128_32x2_nt_align8

          Status: Success
    Verification: ON
     Disposition: Passed

reference_device: Passed
          cuBLAS: Passed

       Arguments: --gemm_kind=universal --m=3456 --n=4096 --k=4096 --A=f16:column --B=f16:row --C=f32:column --alpha=1  \
                  --beta=0 --split_k_slices=1 --batch_count=1 --op_class=tensorop --accum=f32 --cta_m=256 --cta_n=128  \
                  --cta_k=32 --stages=2 --warps_m=4 --warps_n=2 --warps_k=1 --inst_m=16 --inst_n=8 --inst_k=8 --min_cc=75  \
                  --max_cc=1024

           Bytes: 118489088  bytes
           FLOPs: 115992428544  flops

         Runtime: 1.55948  ms
          Memory: 70.7616 GiB/s

            Math: 74378.8 GFLOP/s



=============================
...

Building one CUDA Core GEMM kernel Copy Section

To compile one SGEMM kernel targeting NVIDIA Ampere and Turing architecture, use the below cmake command line:

$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_simt_sgemm_128x128_8x2_nn_align1
...
$ make cutlass_profiler -j16

Example command line for profiling single SGEMM CUDA kernel is as follows:

$ ./tools/profiler/cutlass_profiler --kernels=sgemm --m=3456 --n=4096 --k=4096

=============================
  Problem ID: 1

        Provider: CUTLASS
   OperationKind: gemm
       Operation: cutlass_simt_sgemm_128x128_8x2_nn_align1

          Status: Success
    Verification: ON
     Disposition: Passed

          cuBLAS: Passed

       Arguments: --m=3456 --n=4096 --k=4096 --A=f32:column --B=f32:column --C=f32:column --alpha=1 --beta=0 --split_k_slices=1  \
                  --batch_count=1 --op_class=simt --accum=f32 --cta_m=128 --cta_n=128 --cta_k=8 --stages=2 --warps_m=4  \
                  --warps_n=2 --warps_k=1 --inst_m=1 --inst_n=1 --inst_k=1 --min_cc=50 --max_cc=1024

           Bytes: 180355072  bytes
           FLOPs: 115992428544  flops

         Runtime: 6.73655  ms
          Memory: 24.934 GiB/s

            Math: 17218.4 GFLOP/s

=============================

Building a subset of Tensor Core Convolution kernels Copy Section

To compile a subset of Tensor core convolution kernels implementing forward propagation (fprop) with FP32 accumulation
and FP16 input targeting NVIDIA Ampere and Turing architecture, use the below cmake command line:

$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_tensorop_s*fprop_optimized_f16
...
$ make cutlass_profiler -j16

Example command line for profiling a subset of Tensor Core convolution kernels is as follows:

$ ./tools/profiler/cutlass_profiler --kernels=cutlass_tensorop_s*fprop_optimized_f16 --n=8 --h=224 --w=224 --c=128 --k=128 --r=3 --s=3

...
=============================
  Problem ID: 1

        Provider: CUTLASS
   OperationKind: conv2d
       Operation: cutlass_tensorop_s16816fprop_optimized_f16_128x128_32x5_nhwc

          Status: Success
    Verification: ON
     Disposition: Passed

reference_device: Passed

       Arguments: --conv_kind=fprop --n=8 --h=224 --w=224 --c=128 --k=128 --r=3 --s=3 --p=224 --q=224 --pad_h=1 --pad_w=1  \
                  --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1 --Activation=f16:nhwc --Filter=f16:nhwc --Output=f32:nhwc  \
                  --conv_mode=cross --iterator_algorithm=optimized --alpha=1 --beta=0 --split_k_mode=serial --split_k_slices=1  \
                  --eq_gemm_provider=none --op_class=tensorop --accum=f32 --cta_m=128 --cta_n=128 --cta_k=32 --stages=5  \
                  --warps_m=2 --warps_n=2 --warps_k=1 --inst_m=16 --inst_n=8 --inst_k=16 --min_cc=80 --max_cc=1024

           Bytes: 1130659840  bytes
           FLOPs: 118482796544  flops

         Runtime: 0.711496  ms
          Memory: 1479.99 GiB/s

            Math: 166526 GFLOP/s

=============================
...

Building one Convolution CUDA kernel Copy Section

To compile and run one CUDA Core convolution kernel implementing forward propagation (fprop) with F32 accumulation
and FP32 input targeting NVIDIA Ampere and Turing architecture, use the below cmake command line:

$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_simt_sfprop_optimized_128x128_8x2_nhwc
...
$ make cutlass_profiler -j16

Example command line for profiling one CUDA Core convolution kernel:

$ ./tools/profiler/cutlass_profiler --kernels=cutlass_simt_sfprop_optimized_128x128_8x2_nhwc --n=8 --h=224 --w=224 --c=128 --k=128 --r=3 --s=3


=============================
  Problem ID: 1

        Provider: CUTLASS
   OperationKind: conv2d
       Operation: cutlass_simt_sfprop_optimized_128x128_8x2_nhwc

          Status: Success
    Verification: ON
     Disposition: Passed

reference_device: Passed

       Arguments: --conv_kind=fprop --n=8 --h=224 --w=224 --c=128 --k=128 --r=3 --s=3 --p=224 --q=224 --pad_h=1 --pad_w=1  \
                  --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1 --Activation=f32:nhwc --Filter=f32:nhwc --Output=f32:nhwc  \
                  --conv_mode=cross --iterator_algorithm=optimized --alpha=1 --beta=0 --split_k_mode=serial --split_k_slices=1  \
                  --eq_gemm_provider=none --op_class=simt --accum=f32 --cta_m=128 --cta_n=128 --cta_k=8 --stages=2 --warps_m=4  \
                  --warps_n=2 --warps_k=1 --inst_m=1 --inst_n=1 --inst_k=1 --min_cc=50 --max_cc=1024

           Bytes: 2055798784  bytes
           FLOPs: 118482796544  flops

         Runtime: 7.34266  ms
          Memory: 260.752 GiB/s

            Math: 16136.2 GFLOP/s


=============================

More Details on Compiling CUTLASS Kernels and CUTLASS Profiler Copy Chapter

• Please follow the links for more CMake examples on selectively compiling CUTLASS kernels:
  • GEMM CMake Examples
  • Implicit GEMM convolution CMake Examples
• Further details about the CUTLASS Profiler are described here.

About

CUTLASS is released by NVIDIA Corporation as Open Source software under the
3-clause "New" BSD license.

Contributors

The official list of CUTLASS developers and contributors is available here: CONTRIBUTORS.

Copyright

Copyright (c) 2017 - 2026 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.

Media Docs Cpp Quickstart (`media/docs/cpp/quickstart.md`) Copy Section

Quickstart

Prerequisites Copy Chapter

CUTLASS requires:

• NVIDIA CUDA Toolkit (11.4 or later required, 12.0 recommended)
• CMake 3.18+
• host compiler supporting C++17 or greater (minimum g++ 7.5.0)
• Python 3.6+

CUTLASS may be optionally compiled and linked with

• cuBLAS
• cuDNN v7.6 or later

Initial build steps Copy Chapter

Construct a build directory and run CMake.

$ export CUDACXX=${CUDA_INSTALL_PATH}/bin/nvcc

$ mkdir build && cd build

$ cmake .. -DCUTLASS_NVCC_ARCHS=90a            # compiles for NVIDIA Hopper GPU architecture
$ cmake .. -DCUTLASS_NVCC_ARCHS=100a           # compiles for NVIDIA Blackwell SM100 GPU architecture

If your goal is strictly to build only the CUTLASS Profiler and to minimize compilation time, we suggest
executing the following CMake command in an empty build/ directory.

$ cmake .. -DCUTLASS_NVCC_ARCHS=90a -DCUTLASS_ENABLE_TESTS=OFF -DCUTLASS_UNITY_BUILD_ENABLED=ON

This reduces overall compilation time by excluding unit tests and enabling the unity build.

You may reduce build times by compiling only certain operations by setting the CUTLASS_LIBRARY_OPERATIONS flag as shown below,
executed from an empty build/ directory. This only compiles 2-D convolution kernels.

$ cmake .. -DCUTLASS_NVCC_ARCHS=90a -DCUTLASS_LIBRARY_OPERATIONS=conv2d

You may also filter kernels by name by supplying a filter string with flag CUTLASS_LIBRARY_KERNELS. For example the below command selects only CUTLASS-3 kernels.

$ cmake .. -DCUTLASS_NVCC_ARCHS=90a -DCUTLASS_LIBRARY_KERNELS=cutlass3x*

See more examples on selectively compiling CUTLASS GEMM and convolution kernels here.

You may explicitly exclude cuBLAS and cuDNN as dependencies with the following CMake flags.

• -DCUTLASS_ENABLE_CUBLAS=OFF
• -DCUTLASS_ENABLE_CUDNN=OFF

Build and run the CUTLASS Profiler Copy Chapter

From the build/ directory created above, compile the CUTLASS Profiler.

$ make cutlass_profiler -j12

Then execute the CUTLASS Profiler computing GEMM, execute the following command.

$ ./tools/profiler/cutlass_profiler --kernels=sgemm --m=4352 --n=4096 --k=4096

=============================
  Problem ID: 1

    Provider: CUTLASS
   Operation: cutlass_simt_sgemm_128x128_nn

 Disposition: Passed
      Status: Success

   Arguments:  --m=4352 --n=4096 --k=4096 --A=f32:column --B=f32:column --C=f32:column --alpha=1 --beta=0  \
               --split_k_slices=1 --batch_count=1 --op_class=simt --accum=f32 --cta_m=128 --cta_n=128 --cta_k=8  \
               --stages=2 --warps_m=2 --warps_n=2 --warps_k=1 --inst_m=1 --inst_n=1 --inst_k=1 --min_cc=50  \
               --max_cc=1024

       Bytes: 52428800  bytes
       FLOPs: 146064539648  flops

     Runtime: 10.5424  ms
      Memory: 4.63158 GiB/s

        Math: 13854.9 GFLOP/s

To execute the CUTLASS Profiler for convolution, run the following example.

$ ./tools/profiler/cutlass_profiler --kernels=s1688fprop --n=8 --h=224 --w=224 --c=128 --k=128 --r=3 --s=3 --pad_h=1 --pad_w=1

To execute all CUTLASS 2-D convolution operators, execute the following.

$ ./tools/profiler/cutlass_profiler --operation=conv2d --n=8 --h=224 --w=224 --c=128 --k=128 --r=3 --s=3


=============================
  Problem ID: 1

        Provider: CUTLASS
   OperationKind: conv2d
       Operation: cutlass_simt_sfprop_optimized_128x128_8x2_nhwc

          Status: Success
    Verification: ON
     Disposition: Passed

reference_device: Passed

       Arguments: --conv_kind=fprop --n=8 --h=224 --w=224 --c=128 --k=128 --r=3 --s=3 --p=224 --q=224 --pad_h=1 --pad_w=1  \
                  --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1 --Activation=f32:nhwc --Filter=f32:nhwc --Output=f32:nhwc  \
                  --conv_mode=cross --iterator_algorithm=optimized --alpha=1 --beta=0 --split_k_mode=serial --split_k_slices=1  \
                  --eq_gemm_provider=none --op_class=simt --accum=f32 --cta_m=128 --cta_n=128 --cta_k=8 --stages=2 --warps_m=4  \
                  --warps_n=2 --warps_k=1 --inst_m=1 --inst_n=1 --inst_k=1 --min_cc=50 --max_cc=1024

           Bytes: 2055798784  bytes
           FLOPs: 118482796544  flops

         Runtime: 8.13237  ms
          Memory: 235.431 GiB/s

            Math: 14569.3 GFLOP/s

See documentation for the CUTLASS Profiler for more details.

Build and run CUTLASS Unit Tests Copy Chapter

From the build/ directory created above, simply build the target test_unit to compile and run
all unit tests.

$ make test_unit -j
...
...
...
[----------] Global test environment tear-down
[==========] 946 tests from 57 test cases ran. (10812 ms total)
[  PASSED  ] 946 tests.
$

The exact number of tests run is subject to change as we add more functionality.

No tests should fail. Unit tests automatically construct the appropriate runtime filters
to avoid executing on architectures that do not support all features under test.

The unit tests are arranged hierarchically mirroring the CUTLASS Template Library. This enables
parallelism in building and running tests as well as reducing compilation times when a specific
set of tests are desired.

For example, the following executes strictly the warp-level GEMM tests.

$ make test_unit_gemm_warp -j
...
...
[----------] 3 tests from SM75_warp_gemm_tensor_op_congruous_f16
[ RUN      ] SM75_warp_gemm_tensor_op_congruous_f16.128x128x8_32x128x8_16x8x8
[       OK ] SM75_warp_gemm_tensor_op_congruous_f16.128x128x8_32x128x8_16x8x8 (0 ms)
[ RUN      ] SM75_warp_gemm_tensor_op_congruous_f16.128x128x32_64x64x32_16x8x8
[       OK ] SM75_warp_gemm_tensor_op_congruous_f16.128x128x32_64x64x32_16x8x8 (2 ms)
[ RUN      ] SM75_warp_gemm_tensor_op_congruous_f16.128x128x32_32x32x32_16x8x8
[       OK ] SM75_warp_gemm_tensor_op_congruous_f16.128x128x32_32x32x32_16x8x8 (1 ms)
[----------] 3 tests from SM75_warp_gemm_tensor_op_congruous_f16 (3 ms total)
...
...
[----------] Global test environment tear-down
[==========] 104 tests from 32 test cases ran. (294 ms total)
[  PASSED  ] 104 tests.
[100%] Built target test_unit_gemm_warp

Building for Multiple Architectures Copy Chapter

To minimize compilation time, specific GPU architectures can be enabled via the CMake command,
selected by CUDA Compute Capability.

NVIDIA Blackwell Architecture.

$ cmake .. -DCUTLASS_NVCC_ARCHS=100a              # compiles for NVIDIA Blackwell GPU architecture

NVIDIA Hopper Architecture.

$ cmake .. -DCUTLASS_NVCC_ARCHS=90a              # compiles for NVIDIA Hopper GPU architecture

NVIDIA Ampere Architecture.

$ cmake .. -DCUTLASS_NVCC_ARCHS=80               # compiles for NVIDIA Ampere GPU architecture

NVIDIA Turing Architecture.

$ cmake .. -DCUTLASS_NVCC_ARCHS=75               # compiles for NVIDIA Turing GPU architecture

NVIDIA Volta Architecture.

$ cmake .. -DCUTLASS_NVCC_ARCHS=70               # compiles for NVIDIA Volta GPU architecture

NVIDIA Pascal Architecture.

$ cmake .. -DCUTLASS_NVCC_ARCHS="60;61"          # compiles for NVIDIA Pascal GPU architecture

NVIDIA Maxwell Architecture.

$ cmake .. -DCUTLASS_NVCC_ARCHS="50;53"          # compiles for NVIDIA Maxwell GPU architecture

Using CUTLASS within other applications Copy Chapter

Applications should list /include within their include paths. They must be
compiled as C++17 or greater.

Example: print the contents of a variable storing half-precision data.

#include <iostream>
#include <cutlass/cutlass.h>
#include <cutlass/numeric_types.h>
#include <cutlass/core_io.h>

int main() {

  cutlass::half_t x = 2.25_hf;

  std::cout << x << std::endl;

  return 0;
}

Launching a GEMM kernel in CUDA Copy Chapter

Example: launch a mixed-precision GEMM targeting Turing Tensor Cores.

Note, this example uses CUTLASS Utilities. Be sure tools/util/include is listed as an include path.

/* Detailed source-code truncated for AI context efficiency. */

Note, the above could be simplified as follows using helper methods defined in HostTensor.

cutlass::HostTensor<cutlass::half_t, cutlass::layout::ColumnMajor> A({M, K});
  cutlass::HostTensor<cutlass::half_t, cutlass::layout::ColumnMajor> B({K, N});
  cutlass::HostTensor<cutlass::half_t, cutlass::layout::ColumnMajor> C({M, N});

  //
  // Use the TensorRef returned by HostTensor::device_ref().
  //

  status = gemm_op({
    {M, N, K},
    A.device_ref(),            // TensorRef to A device tensor
    B.device_ref(),            // TensorRef to B device tensor
    C.device_ref(),            // TensorRef to C device tensor
    C.device_ref(),            // TensorRef to D device tensor - may be the same as C
    {alpha, beta}              // epilogue operation arguments
  });

Launching a GEMM kernel using CUTLASS 3.0 or newer Copy Chapter

Example: launch a mixed-precision GEMM targeting Hopper Tensor Cores.

/* Detailed source-code truncated for AI context efficiency. */

CUTLASS Library Copy Chapter

The CUTLASS Library defines an API for managing and executing collections of compiled
kernel instances and launching them from host code without template instantiations in client code.

The host-side launch API is designed to be analogous to BLAS implementations for convenience, though its
kernel selection procedure is intended only to be functionally sufficient. It may not launch the
optimal tile size for a given problem. It chooses the first available kernel whose data types,
layouts, and alignment constraints satisfy the given problem. Kernel instances and a data structure
describing them are completely available to client applications which may choose to implement their
own selection logic.

cuBLAS offers the best performance and functional coverage
for dense matrix computations on NVIDIA GPUs.

The CUTLASS Library is used by the CUTLASS Profiler to manage kernel instances, and it is also used
by several SDK examples.

• 10_planar_complex
• 11_planar_complex_array

The CUTLASS Library defines enumerated types describing numeric data types, matrix and tensor
layouts, math operation classes, complex transformations, and more.

Client applications should specify tools/library/include in their
include paths and link against libcutlas_lib.so.

The CUTLASS SDK example 10_planar_complex specifies
its dependency on the CUTLASS Library with the following CMake command.

target_link_libraries(
  10_planar_complex
  PRIVATE
  cutlass_lib
  cutlass_tools_util_includes
)

A sample kernel launch from host-side C++ is shown as follows.

/* Detailed source-code truncated for AI context efficiency. */

Example CMake Commands Copy Chapter

To instantiate all operations supporting all tile sizes, data types, and alignment constraints, specify
-DCUTLASS_LIBRARY_KERNELS=all when running cmake.

$ cmake .. -DCUTLASS_NVCC_ARCHS='70;75;80' -DCUTLASS_LIBRARY_KERNELS=all

The above command line generates about twenty thousand kernels targeting NVIDIA Ampere, Turing, and Volta architectures.
Compiling thousands of kernels for three different architectures is time-consuming. Additionally, this would also result
in a large binary size and on some platforms linker to fail on building the library.

Enabling the "unity build" instantiates multiple kernel instances in each compilation unit, thereby reducing binary size
and avoiding linker limitations on some platforms.

$ cmake .. -DCUTLASS_NVCC_ARCHS="70;75;80" -DCUTLASS_LIBRARY_KERNELS=all -DCUTLASS_UNITY_BUILD_ENABLED=ON

It is advised to only compile CUTLASS kernels for NVIDIA architectures one plans on running. Furthermore, kernels
can be selectively included in the CUTLASS Library by specifying filter strings and wildcard characters when executing CMake.

Several examples are defined below for convenience. They may be combined as a comma-delimited list.
Compling only the kernels desired reduces compilation time.

GEMM CMake Examples Copy Chapter

Example. All GEMM kernels targeting NVIDIA Ampere Tensor Cores.

$ cmake .. -DCUTLASS_NVCC_ARCHS=80 -DCUTLASS_LIBRARY_KERNELS=tensorop*gemm

Example. All GEMM kernels targeting NVIDIA Turing Tensor Cores.

$ cmake .. -DCUTLASS_NVCC_ARCHS=75 -DCUTLASS_LIBRARY_KERNELS=tensorop*gemm

Example. All GEMM kernels with FP32 accumulation targeting NVIDIA Ampere, Turing, and Volta architectures.

$ cmake .. -DCUTLASS_NVCC_ARCHS="70;75;80" -DCUTLASS_LIBRARY_KERNELS=s*gemm

Example. All kernels which expect A and B to be column-major or row-major targeting NVIDIA Ampere, Turing, and Volta architectures.

$ cmake .. -DCUTLASS_NVCC_ARCHS="70;75;80" -DCUTLASS_LIBRARY_KERNELS=gemm*nn,gemm*tt

Example. All planar complex GEMM variants targeting NVIDIA Ampere, Turing, and Volta architectures.

$ cmake .. -DCUTLASS_NVCC_ARCHS="70;75;80" -DCUTLASS_LIBRARY_KERNELS=planar_complex

Convolution CMake Examples Copy Chapter

Example. All convolution kernels targeting NVIDIA Ampere's 16816 Tensor Core operation

$ cmake .. -DCUTLASS_NVCC_ARCHS='80' -DCUTLASS_LIBRARY_KERNELS=s16816fprop,s16816dgrad,s16816wgrad

Example. All forward propagation (fprop) convolution kernels targeting CUDA Cores for multiple NVIDIA architectures

$ cmake .. -DCUTLASS_NVCC_ARCHS='50;60;61;70;75;80' -DCUTLASS_LIBRARY_KERNELS=sfprop

Example. All forward propagation (fprop) convolution kernels with FP32 accumulation and FP16 input targeting NVIDIA Ampere's 16816 Tensor Core operation

$ cmake .. -DCUTLASS_NVCC_ARCHS='80' -DCUTLASS_LIBRARY_KERNELS=s16816fprop_*_f16

Example. All backward weight gradient (wgrad) convolution kernels with FP32 accumulation, FP16 input, and optimized global memory iterator
targeting NVIDIA Ampere, Turing, and Volta Tensor Core operations

$ cmake .. -DCUTLASS_NVCC_ARCHS='70;75;80' -DCUTLASS_LIBRARY_KERNELS=tensorop*s*wgrad_optimized_f16

Instantiating a Blackwell SM100 GEMM kernel Copy Chapter

Blackwell SM100 kernels are instantiated very similarly to Hopper kernels. Let us start with an
FP8 GEMM without blockscaling
as an example.

The kernel starts with setting up datatypes and cluster shapes.

using LayoutA = cutlass::layout::RowMajor;
  using LayoutB = cutlass::layout::ColumnMajor;
  using LayoutC = cutlass::layout::ColumnMajor;
  using ElementA = cutlass::float_e4m3_t;
  using ElementB = cutlass::float_e4m3_t;
  using ElementC = cutlass::float_e4m3_t;
  using ElementD = cutlass::float_e4m3_t;
  using ElementAccumulator = float;
  using ElementCompute = float;
  using ElementBias = cutlass::half_t;
  using MmaTileShape = cute::Shape<_128,_64,Int<128 / sizeof(ElementA)>>;
  using ClusterShape = cute::Shape<_1,_1,_1>;

The epilogue needs to be instantiated first as the mainloop collective builder takes the shared memory budget of epilogue in the template parameter list. The 3.x epilogue collective builder API has not changed
for Blackwell, so the epilogue fusion is built in a same way as an SM90 epilogue.

using EpilogueSchedule = cutlass::epilogue::TmaWarpSpecialized1Sm;

  using FusionOperation = cutlass::epilogue::fusion::LinearCombination<
    ElementD,
    ElementCompute,
    ElementC
  >;

  using CollectiveEpilogue = typename cutlass::epilogue::collective::CollectiveBuilder<
      cutlass::arch::Sm100, cutlass::arch::OpClassTensorOp,
      MmaTileShape, ClusterShape,
      cutlass::epilogue::collective::EpilogueTileAuto,
      ElementAccumulator, ElementCompute,
      ElementC, LayoutC, 16 / sizeof(ElementC),
      ElementD, LayoutC, 16 / sizeof(ElementD),
      EpilogueSchedule,
      FusionOperation
    >::CollectiveOp;

One can refer to our Sm100 unit tests as examples of how to correctly
choose mainloop schedules. All of our dispatch policies can be found in dispatch_policy.hpp
and more comprehensive Blackwell specific documentation for valid
dispatch policies can be in blackwell_functionality.md.

using MainloopSchedule = cutlass::gemm::KernelTmaWarpSpecialized1SmSm100;
  using CollectiveMainloop = typename cutlass::gemm::collective::CollectiveBuilder<
      cutlass::arch::Sm100, cutlass::arch::OpClassTensorOp,
      ElementA, LayoutA, 16 / sizeof(ElementA),
      ElementB, LayoutB, 16 / sizeof(ElementB),
      ElementAccumulator,
      MmaTileShape, ClusterShape,
      cutlass::gemm::collective::StageCountAutoCarveout<static_cast<int>(sizeof(typename CollectiveEpilogue::SharedStorage))>,
      MainloopSchedule
    >::CollectiveOp;

  using GemmKernel = cutlass::gemm::kernel::GemmUniversal<
      Shape<int,int,int,int>,
      CollectiveMainloop,
      CollectiveEpilogue
  >;

Instantiating a blockscaled GEMM kernel is slightly different. Referring to an MXFP8 GEMM sample unit test, it takes a different tensor operation class:

using ElementA = cutlass::mx_float8_t<cutlass::float_e4m3_t>;
  using ElementB = cutlass::mx_float8_t<cutlass::float_e4m3_t>;

are needed in the mainloop builder:

using CollectiveMainloop = typename cutlass::gemm::collective::CollectiveBuilder<
      cutlass::arch::Sm100, cutlass::arch::OpClassTensorOp,
      ElementA, LayoutA, 16,
      ElementB, LayoutB, 16,
      ElementAccumulator,
      MmaTileShape, ClusterShape,
      cutlass::gemm::collective::StageCountAutoCarveout<static_cast<int>(sizeof(typename CollectiveEpilogue::SharedStorage))>,
      cutlass::gemm::KernelScheduleAuto
    >::CollectiveOp;

We encourage a user to refer to Sm100 unit tests and the generated profiler-based kernels as more comprehensive samples.

Copyright Copy Section

Copyright (c) 2017 - 2026 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.

Media Docs Cpp Cute 00 Quickstart (`media/docs/cpp/cute/00_quickstart.md`) Copy Section

Getting Started With CuTe

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 packages 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 abstraction of CuTe are the 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.

System Requirements Copy Chapter

CuTe shares CUTLASS 3.x's software requirements,
including NVCC with a C++17 host compiler.

Knowledge prerequisites Copy Chapter

CuTe is a CUDA C++ header-only library. It requires C++17
(the revision of the C++ Standard that was released in 2017).

Throughout this tutorial, we assume intermediate C++ experience.
For example, we assume that readers know
how to read and write templated functions and classes, and
how to use the auto keyword to deduce a function's return type.
We will be gentle with C++ and explain some things
that you might already know.

We also assume intermediate CUDA experience.
For example, readers must know
the difference between device and host code,
and how to launch kernels.

Building Tests and Examples Copy Chapter

CuTe's tests and examples build and run as part of CUTLASS's normal build process.

CuTe's unit tests live in the test/unit/cute subdirectory.

CuTe's examples live in the examples/cute subdirectory.

Library Organization Copy Chapter

CuTe is a header-only C++ library, so there is no source code that needs building. Library headers are contained within the top level include/cute directory, with components of the library grouped by directories that represent their semantics.

Tutorial Copy Chapter

This directory contains a CuTe tutorial in Markdown format.
The file
0x_gemm_tutorial.md
explains how to implement dense matrix-matrix multiply using CuTe components.
It gives a broad overview of CuTe and thus would be a good place to start.

Other files in this directory discuss specific parts of CuTe.

• 01_layout.md describes Layout, CuTe's core abstraction.

• 02_layout_algebra.md describes more advanced Layout operations and the CuTe layout algebra.

• 03_tensor.md describes Tensor,
  a multidimensional array abstraction which composes Layout
  with an array of data.

• 04_algorithms.md summarizes CuTe's
  generic algorithms that operate on Tensors.

• 0t_mma_atom.md demonstrates CuTe's meta-information and interface to our GPUs'
  architecture-specific Matrix Multiply-Accumulate (MMA) instructions.

• 0x_gemm_tutorial.md walks through building a GEMM from scratch using CuTe.

• 0y_predication.md explains what to do
  if a tiling doesn't fit evenly into a matrix.

• 0z_tma_tensors.md explains an advanced Tensor type that CuTe uses to support TMA loads and stores.

Quick Tips Copy Chapter

How do I print CuTe objects on host or device? Copy Section

The cute::print function has overloads for almost all CuTe types, including Pointers, Integers, Strides, Shapes, Layouts, and Tensors. When in doubt, try calling print on it.

CuTe's print functions work on either host or device.
Note that on device, printing is expensive.
Even just leaving print code in place on device,
even if it is never called
(e.g., printing in an if branch that is not taken at run time),
may generate slower code.
Thus, be sure to remove code that prints on device after debugging.

You might also only want to print on thread 0 of each threadblock, or threadblock 0 of the grid. The thread0() function returns true only for global thread 0 of the kernel, that is, for thread 0 of threadblock 0. A common idiom for printing CuTe objects to print only on global thread 0.

if (thread0()) {
  print(some_cute_object);
}

Some algorithms depend on some thread or threadblock,
so you may need to print on threads or threadblocks other than zero.
The header file
cute/util/debug.hpp,
among other utilities,
includes the function bool thread(int tid, int bid)
that returns true if running on thread tid and threadblock bid.

Other output formats

Some CuTe types have special printing functions that use a different output format.

The cute::print_layout function will display any rank-2 layout in a plain text table. This is excellent for visualizing the map from coordinates to indices.

The cute::print_tensor function will display any rank-1, rank-2, rank-3, or rank-4 tensor in a plain text multidimensional table. The values of the tensor are printed so you can verify the tile of data is what you expect after a copy, for example.

The cute::print_latex function will print LaTeX commands that you can use to build a nicely formatted and colored tables via pdflatex. This work for Layout, TiledCopy, and TiledMMA, which can be very useful to get a sense of layout patterns and partitioning patterns within CuTe.

Copyright Copy Chapter

Copyright (c) 2017 - 2026 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.

Media Docs Cpp Blackwell (`media/docs/cpp/blackwell.rst`) Copy Section

.. _blackwell:

Blackwell Specific

.. toctree::
:maxdepth: 2

Blackwell SM100/SM120 GEMMs
Blackwell Cluster Launch Control

Media Docs Cpp Blackwell Cluster Launch Control (`media/docs/cpp/blackwell_cluster_launch_control.md`) Copy Section

Blackwell Cluster Launch Control

Overview Copy Chapter

A GEMM workload usually consists of three phases: prologue, mainloop and epilogue. Each SM will process multiple output tiles in series if the number of output tiles are much more than the number of SMs, completely exposing the overhead of prologue and epilogue.

Consider a GEMM that has 20x20x1 output tiles, running on a GPU with 100 SMs. There is another kernel occupying all the resources of 20 SMs so only 80 SMs can be used. Assume cluster shape is 1x1x1. The following diagram shows how the schedule would look like for such a kernel.

Static Scheduler Copy Section

CUTLASS has adopted a software technique named persistent kernels. Persistent clusters, or Workers, can stay on the GPU throughout kernel execution and process multiple tiles, hiding prologue and epilogue costs. The tile scheduler statically determines the next output tile to process with zero overhead.

However, static scheduler is susceptible to workload imbalance if the resources of some SMs are unavailable. The following diagram illustrates this issue.

Dynamic Scheduler with Cluster Launch Control Copy Section

A fundamental limitation of persistent scheduling is that the number of SMs this kernel can utilize is unknown in real time. Some SMs might be occupied by another kernel and thus their resources are unavailable. This makes it challenging to load-balance work across SMs.

Blackwell introduces cluster launch control (CLC) for dynamic scheduling. (See https://docs.nvidia.com/cuda/parallel-thread-execution). With this feature, the kernel launches a grid containing as many threadblocks as there are output tiles to compute in the kernel -- just like one would in a non-persistent kernel. Here we define ClcID to be a coordinate from the 3D grid launched on GPU.

Cluster launch control follows the below rules:

1. A ClcID will be launched as a Worker when there are available resources.
2. A ClcID can be queried by an existing Worker via clusterlaunchcontrol.try_cancel instruction.
3. Every ClcID is guaranteed to be processed by either (1) or (2).
4. Each worker uses the {blockIdx.x, blockIdx.y, blockIdx.z} coordinate as the first output tile to process and uses the CLC query for subsequent processing of output tiles.
5. clusterlaunchcontrol.try_cancel instruction returns either a success signal with a ClcID or a decline signal. The most common reason of a decline is that all ClcIDs have been processed.
6. Cluster launch control works on the granularity of clusters. For example, a 2x2 persistent worker cluster's query will consume 2x2 ClcIDs at once.

The following diagram shows how the schedule would look like with cluster launch control.

Programming Model Copy Chapter

Pseudo Code Copy Section

Non-persistent kernel

// Non-persistent kernel
__device__ non_persistent_kernel(...) {
  setup_common_data_structures();
  dim3 workCoordinates = blockIdx;
  coordinate_specific_compute(workCoordinates);
}

Static Persistent Kernel

// Static Persistent Kernel
__device__ static_persistent_kernel(...) {
  setup_common_data_structures(...);
  dim3 workCoordinates = blockIdx;
  bool isValidId;
  do {
    coordinate_specific_compute(workCoordinates);
    std::tie(isValidId, workCoordinates) = staticTileScheduler.fetch_next_work();
  } while (isValidId);
}

Blackwell Dynamic Persistent Kernel

// Dynamic Persistent Kernel
__device__ clc_dynamic_persistent_kernel(...) {
  setup_common_data_structures(...);
  dim3 workCoordinates = blockIdx;
  dim3 newClcID;
  bool isValidId;
  do {
    coordinate_specific_compute(workCoordinates);
    std::tie(isValidId, newClcID) = clcTileScheduler.fetch_next_work();
    workCoordinates = newClcID;
  } while (isValidId);
}

Cluster Launch Control Pipeline Class Copy Section

Please refer to the PipelineCLCFetchAsync pipeline class defined in Cluster launch control pipeline class. Cluster launch control queries can be pipelined and managed by an asynchronous pipeline with producer-consumer relationship (See
pipeline document). The producer is the scheduler warp of the 0th CTA in the cluster and the consumers are all warps that need ClcIDs.

To setup a CLC pipeline correctly, we need to make sure the params are set to the right values:

• transaction_bytes is 16 as CLC will return a 16B response and store it in the specified shared memory address.
• consumer_arv_count is the thread count of all the consumer warps in the cluster.
• producer_arv_count is 1 because only one thread from scheduler warp will be elected to issue clusterlaunchcontrol.try_cancel.
• producer_blockid is 0 to denote that the first CTA in the cluster is producing.

Dynamic tile scheduler class Copy Section

Please refer to PersistentTileSchedulerSm100 class defined in sm100 dynamic persistent tile scheduler.

There are two important methods of the CLC scheduler class. The first is advance_to_next_work, which is intended to be executed by one elected thread from the scheduler warp. It effectively sends out the CLC query to the CLC. A CLC query response will be broadcast to the same shared memory address of all CTAs in the cluster.

The other method is named get_current_work. It simply loads the CLC response from the shared memory buffer indexed by a pipeline state.

The CLC pipeline and scheduler classes are used together to ensure correct functionality and necessary synchronization of CLC feature. Please refer to cluster launch control pipeline unit test.

Blackwell Warp-specialized Persistent Kernel Copy Chapter

Now, let's take a look at how CLC feature is used in our Blackwell dense GEMM kernel.

This particular warp-specialized kernel has the following warp assignment:

Scheduler warp is the producer of the CLC pipeline. The consumers are the MMA, Mainloop Load, Epilogue Load and Epilogue warps. In addition, the scheduler warp is its own consumer! This is because it needs the success information from the query to terminate the persistent loop on end-of-grid.

The CLC pipeline has a depth of 3 to overlap the CLC operations of multiple waves for latency hiding. The first ClcID is the preloaded blockIdx, which does not require CLC query and is fully static.

Copyright Copy Section

Copyright (c) 2025 - 2026 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.

Media Docs Cpp Blackwell Functionality (`media/docs/cpp/blackwell_functionality.md`) Copy Section

Blackwell SM100 GEMMs

TLDR; jump to block scaled GEMM example

Blackwell SM100 introduces tcgen05.mma instructions. tcgen05.mma instructions support all legacy types (tfloat32_t, half_t, bfloat16_t, int8_t, uint8_t) and
the new 4, 6, and 8-bits floating point datatypes with and without scale factors.
This document explains the new tcgen05.mma instructions supported by CUTLASS and how one can leverage CUTLASS to create
efficient SM100 GEMM kernels targeting these new mma instructions.

Blackwell SM100 has 7 new tcgen05.mma instructions. These instructions are 2x to 4x faster then Hopper Architecture's WGMMA instructions.

For more detailed information see tcgen05.mma PTX documentation.

New in Blackwell SM100 Copy Chapter

Block Scaled GEMMs Copy Section

Instructions with kind modifiers mxf8f6f4, mxf4, and nvf4mxf4 perform matrix multiplication operations with scale
factors of the form $D = C +( A \times SFA) * (B \times SFB)$. Scale factors are applied to GEMM-K dimension such that
every 16 or 32 elements of $A$ and $B$ matrices in K dimension have an associated scale factor (32 or 64 elements for sparse as sparse gemm compress 2x along k-dim). For example, an $M\times K$,
$A$ matrix has an associated $M \times \lceil K/32 \rceil$ SFA matrix; and an $N\times K$ $B$, matrix has an associated
$N \times \lceil K/32 \rceil$ SFB matrix. For block scaled GEMMs, an entry of output D matrix is
$D_{ij} = C_{ij} + \sum_{k} (A_{i,k} \times SFA_{i,k/SV}) \times (B_{j,k}\times SFB_{j,k/SV})$, in index notation, we SV is the scale factor vector size (16 or 32).
Further details can be found in
PTX documentation on block scaling.

Blackwell Narrow Precision Data Types Copy Section

Narrow-precision tcgen05.mma instructions can operate on several 4, 6, and 8-bit data types. Blackwell MMAs can operate
on five different 8-bit floating point values, of which only two (float_ue8m0_t and float_ue4m3_t) can be used as scale factor data types.
There are two 6-bit floating point types and one 4-bit floating point data type.
See PTX documentation for narrow precision data types for details.

Blackwell Narrow Precision Data Types

[^1]: Only valid as scale factor data types.

Block scaled MMAs use mx and nv types which are a pair of float8_t, float6_t, float4_t with 2 of the scale factor data types with a predetermined scale factor vector size. mx types follow OCP specification (see OCP Specification). The following types provided by CUTLASS can be used as inputs to collective builders to generate the block scaled kernels:

Blackwell Block Scaled Narrow Precision Data Types

Layouts, Tensor Alignment Requirements to Target `tcgen05.mma` Instructions Copy Chapter

Tables below list valid data type, and AB layout combinations. Note that the alignment is reported as number of elements. A and B matrix layouts are
represented with T and N. T represents row-major layouts, and N represents column-major layouts. For instance, TN is
row-major A matrix with column-major B matrix.

For legacy types (tf32, f16, bf16, i8 and u8) alignment requirements for A and B matrices are the same as in Hopper.
All four layouts (TT, NN, NT, TT) are supported for all legacy data types.

Table 1: Valid Data Type, Alignment, and Layout Combinations For MMAs with Legacy Types

For narrow precision Mmas, not all A/B type, and A/B layout combinations are supported by every tcgen05.mma instructions.
Furthermore, tensor copy instructions for subbyte types impose additional alignment requirements while loading narrow-precision
tensors from global memory to shared memory
(see PTX doc for details).

Below tables list valid layout, and alignment values for each A and B data type combination and their target tcgen05.mma
instructions supported by CUTLASS.

Table 2: Valid Data Type, Alignment, and Layout Combinations For Narrow Precision MMAs Without Block Scaling

Table 3: Valid Data Type, Alignment, and Layout Combinations for Block Scaled Narrow Precision MMAs

MMA tile shapes supported Copy Chapter

The alignment restrictions also limit the options for Mma Tile Shapes. Tables below list the supported/valid MmaTileShape,
Layout, and Dispatch Policy combinations for each row of Table 1, Table 2, and Table 3.

Table 4: Valid Tile Shapes and Dispatch Policies for legacy types (All rows of Table 1)

Table 5: Valid Tile Shapes and Dispatch Policies for {float4_t, float6_t} x {float4_t, float6_t} (Rows 1,2,3,6,10,11,12,and 15 of Table 2)

Table 6: Valid Tile Shapes and Dispatch Policies for float8_t x {float4_t, float6_t} (Rows 5,8,14,and 17 of Table 2)

Table 7: Valid Tile Shapes and Dispatch Policies for {float4_t, float6_t} x float8_t (Rows 4,7,13,and 16 of Table 2)

Table 8: Valid Tile Shapes and Dispatch Policies for float8_t x float8_t (Row 9,18 of Table 2)

Table 9: Valid Tile Shapes for nv_float4_t x nv_float4_t (Row 1 and 12 of Table 3)

Table 10: Valid Tile Shapes and Dispatch Policies for mx_float4_t x mx_float4_t (Row 2 and 13 of Table 3)

Table 11: Valid Tile Shapes and Dispatch Policies for mx_float4_t x mx_float4_t (Row 3 and 14 of Table 3)

Table 12: Valid Tile Shapes and Dispatch Policies for {mx_float4_t, mx_float6_t, mx_float8_t} x {mx_float4_t, mx_float6_t} (Rows 4, 5, 7, 8, 10, 15, 16, 18, 19, and 21 of Table 3)

Table 13: Valid Tile Shapes and Dispatch Policies for {mx_float4_t, mx_float6_t, mx_float8_t} x mx_float8_t (Rows 6, 9, 11, 17, 20, and 22 of Table 3)

Epilogue config supported Copy Chapter

Table 14: Epilogue Dispatch Policy

Table 15: Epilogue PerSmTileShape_MNK

MMA_TileShape_K is is generally 4 * MMA-Instruction-K. It depends on the config we defined in MMA tile shapes supported section.

Auto Kernel Dispatch Policies Copy Section

In addition to direct dispatch policies listed above, the user can also use auto policies for both non-block scaled narrow-precision
GEMMs (both sparse and dense), and block scaled narrow-precision GEMMs (only dense).

CUTLASS will do its best to find the most efficient kernel for given parameters, however, the preferred method for building
these kernels is to use direct kernel dispatch policies shown in the above tables.

• cutlass::gemm::collective::KernelScheduleAuto: For a given Mma Tile Size, data type and layout combinations choose instr kind (mxf8f6f4, mxf4, nvf4mxf4) and 1/2 SM tcgen05.mma(.sp).
• KernelTmaWarpSpecialized1SmBlockScaledSm100: Use 1 SM tcgen05.mma instruction and choose instr kind (mxf8f6f4, mxf4, nvf4mxf4) automatically.
• KernelTmaWarpSpecialized2SmBlockScaledSm100: Use 2 SM tcgen05.mma instruction and choose instr kind (mxf8f6f4, mxf4, nvf4mxf4) automatically.
• KernelSparseTmaWarpSpecialized1SmBlockScaledSm100: Use 1 SM tcgen05.mma.sp instruction and choose instr kind (mxf8f6f4, mxf4, nvf4mxf4) automatically.
• KernelSparseTmaWarpSpecialized2SmBlockScaledSm100: Use 2 SM tcgen05.mma.sp instruction and choose instr kind (mxf8f6f4, mxf4, nvf4mxf4) automatically.

Similarly for epilogues, we can use cutlass::epilogue::collective::EpilogueScheduleAuto.

Building a Block Scaled Kernel Copy Chapter

For non-blockscaled dense GEMM refer to quick start page. An example dense GEMM can be found:

1. Blackwell FP16 GEMM example.

An example sparse GEMM can be found:

1. Blackwell FP16 Sparse GEMM example.

Narrow precision and block scaled narrow precision kernels can be built using CUTLASS 3.x collective builder interface
(as described in CUTLASS 3.0 GEMM API). However, special attention needs to be given to
A and B matrix layouts, alignment requirements, and dispatch policies to obtain a functionally correct and performant kernel
which are listed above.

Several examples of block scaled dense GEMM kernels can be found in examples/72_blackwell_narrow_precision_gemm directory:

1. NVF4 Gemm with block scaling
2. NVF4 Gemm with block scaling and NVF4 output matrix
3. Mixed precision Mxf8 x Mxf8 GEMM with block scaling

Several examples of block scaled sparse GEMM kernels can be found in examples/84_blackwell_narrow_precision_sparse_gemm directory:

1. NVF4 Gemm with block scaling
2. Mixed precision Mxf4 x Mxf8 GEMM with block scaling

Collective builder interface expects the same arguments as any other CUTLASS 3.x kernels as described
here with a small difference for Collective MMA builder interface.
As in all Blackwell kernels, the TileShape_MNK argument expects the MmaTileShape_MNK which is the tile shape needed
by 1 or 2 SM tcgen05.mma instructions.

Let's consider building a block scaled GEMM where the A matrix is of type mx_float4_t and column-major (N), and the
B matrix is of type mx_float4_t and row-major (T). We first need to describe the A and B tensors, and find the
instruction that can support the selected A and B type and layout pair. Then, we will choose the performance parameters.

The skeleton C++ code is shown below:

/* Detailed source-code truncated for AI context efficiency. */

From the valid type and layout combinations Table 3, we see that only row 3 can support mx_float4_txmx_float4_t
combination with NT layout. As a result, we need to use the tcgen05.mma.kind:mxf8f6f4 instruction. Additionally, in order
to use tcgen05.mma.kind:mxf8f6f4, we see that A and B tensors both should be 128-element aligned.
Thus, we can describe A and B tensors as follows:

///////////////////////////////////////////////////////////
  //                Mainloop Builder Setup
  ///////////////////////////////////////////////////////////
  
  ///////////////////////////////////////////
  // 1. Describe A and B tensors
  ///////////////////////////////////////////
  using ElementA       = mx_float4_t;
  constexpr int AlignA = 128;
  using GmemLayoutA    = cutlass::layout::ColumnMajor;
  using ElementB       = mx_float4_t;
  constexpr int AlignB = 128;
  using GmemLayoutB    = cutlass::layout::RowMajor;

Next, we need to choose the performance parameters such as MmaTileShape_MNK, KernelMainloopPolicy,
and ClusterShape_MNK.

MmaTileShape_MNK supported for mx_float4_txmx_float4_t with mxf8f6f4 are listed in Table 11.
For NT layout, we see that 3 MmaTileShape_MNK are supported: 128x128x128, and 128x256x128 with 1SM instruction;
and 256x256x128 with 2SM instruction. Let's say, we expect to get the best performance with 256x256x128 MMA tile shape
for our GEMM problem. Then, we need to set the KernelMainloopPolicy to KernelTmaWarpSpecialized2SmMxf8f6f4Sm100.
Now, we need to choose the ClusterShape_MNK. Since we have selected a 2SM mma instruction, ClusterShape_MNK should be
compatible and its first mode should be a multiple of 2. ClusterShape_MNK = cute::Shape<_2, [_1|_2|_4], _1> or
ClusterShape_MNK = cute::Shape<_4, [_1|_2|_4], _1> would be valid options. Let's choose cute::Shape<_4,_4,_1>.
Our performance parameters looks like below:

//////////////////////////////////////////
  // 2. Choose Performance Parameters
  //////////////////////////////////////////

  // Tile and cluster shapes
  // Collective MMA takes tile shape of the MMA operation as input
  using KernelMainloopPolicy     = cutlass::gemm::KernelTmaWarpSpecialized2SmMxf8f6f4Sm100;
  using MmaTileShape_MNK         = cute::Shape<_256,_256,_128>;
  using ClusterShape_MNK         = cute::Shape<_4,_4,_1>;

After we config the main-loop, let's setup the epilogue.
A normal epilogue looks like below, we need to specify the output layout, datatype, alignment and PerSmTileShape_MNK, and let others to be default/auto.

PerSmTileShape_MNK should be deduced from the mainloop setup. For example, in above mainloop setup, the MmaTileShape_MNK is
256x256x128 and the KernelMainloopPolicy is 2sm policy.
It means each CTA is doing (256 / 2sm) x 256 x 128 output, so the PerSmTileShape_MNK is 128x256x128. The possible PerSmTileShape_MNK
is listed in Table 15

The epilogue scheduling policy is configurable, and it is common to set cutlass::epilogue::collective::EpilogueScheduleAuto
to allow the epilogue builder to automatically select the appropriate policy. However, it can also be explicitly defined to
use other policies based on the 1sm or 2sm MMA instruction. The available policies are listed in Table 14.

// Describe C and D tensors
  using ElementC = cutlass::half_t;
  constexpr int AlignC = 8;
  using GmemLayoutC = cutlass::layout::RowMajor;
  using ElementD = cutlass::float_e2m1_t;
  constexpr int AlignD = 32;
  using GmemLayoutD = cutlass::layout::RowMajor;
  // Mma's accumulator type
  using ElementAccumulator = float;
  // Epilogue computation's precision type
  using ElementCompute = float;
  
  //
  // Construct CollectiveEpilogue
  //

  using CollectiveEpilogue = typename cutlass::epilogue::collective::CollectiveBuilder<
      cutlass::arch::Sm100, cutlass::arch::OpClassBlockScaledTensorOp,      // Arch and Tensorop spec
      MmaTileShape_MNK, ClusterShape_MNK,                                   // MMA tile shape, and cluster shape
      cutlass::epilogue::collective::EpilogueTileAuto,                      // Epilogue subtile shape. Auto will find a suitable tile shape
      ElementAccumulator, ElementCompute,                                   // Mma instr's accumulator type and compute precision for epilogue
      ElementC, GmemLayoutC, AlignC,                                        // C tensor description
      ElementD, GmemLayoutD, AlignD,                                        // D tensor description
      cutlass::epilogue::TmaWarpSpecialized2Sm                              // Epilogue schedule policy
    >::CollectiveOp;

If we want to let the epilogue generate mxf4/nvf4/mxf6/mxf8 (i.e. elements + block-scalefactor), we need to setup the epilogue fusion into the builder.
First, we need to choose a SFDVectorSize indicates how many elements sharing the same block-scalefactor.
Then, we need to choose ElementSFD and GmemLayoutSFD which indicates the output datatype and which output-dim is used to generate the block-scalefactor.
Typically, GmemLayoutSFD would be same as the GmemLayoutD.

//
  // Construct FusionOperation
  //
  constexpr int SFDVectorSize = 16;
  // Define the fusion operation applied during epilogue
  using FusionOperation = cutlass::epilogue::fusion::LinCombBlockScaleFactor<
      SFDVectorSize,
      ElementD, ElementCompute, 
      ElementSFD, GmemLayoutSFD,
      ElementC
    >;

  using CollectiveEpilogue = typename cutlass::epilogue::collective::CollectiveBuilder<
      cutlass::arch::Sm100, cutlass::arch::OpClassBlockScaledTensorOp,      // Arch and Tensorop spec
      MmaTileShape_MNK, ClusterShape_MNK,                                   // MMA tile shape, and cluster shape
      cutlass::epilogue::collective::EpilogueTileAuto,                      // Epilogue subtile shape. Auto will find a suitable tile shape
      ElementAccumulator, ElementCompute,                                   // Mma instr's accumulator type and compute precision for epilogue
      ElementC, GmemLayoutC, AlignC,                                        // C tensor description
      ElementD, GmemLayoutD, AlignD,                                        // D tensor description
      cutlass::epilogue::TmaWarpSpecialized2Sm                              // Epilogue schedule policy
      FusionOperation                                                       // <================================== Pass the fusion config into epilogue builder.
    >::CollectiveOp;

Above example made a gentle introduction to using the fusion operations in the epilogue. For more detailed example, see
Blackwell GEMM with collective builder

Note that we have first discussed the CollectiveMainloop, then the CollectiveEpilogue for clarity.
However, the CollectiveMainloop needs to know the SMEM utilization of the epilogue. Therefore, it needs to be setup before the CollectiveMainloop. See examples/72_blackwell_narrow_precision_gemm directory for full kernel and run setup.

Scale Factor Layouts Copy Section

The scale factor layout consists of a 512B basic-block structure, as illustrated in the diagram below. Each block contains 128 M/N dimension and 4 scale factors (SF) along the K dimension.
The byte order of the basic storage chunk is row-major, meaning that M0SF0 to M0SF3, M32SF0 to M32SF3, M64SF0 to M64SF3, and M96SF0 to M96SF3 are stored consecutively in GMEM.

If the scale factor tensor exceeds M128xSF4, it indicates that there are multiple basic blocks along both the M and SFK dimensions. The arrangement of these basic blocks follows a K-major order. Here is a diagram illustrating the scenario where M equals 512 and the SFK is 16.

The creation of scale factor tensors' layouts are tedious. CUTLASS provides Sm1xxBlockScaledConfig to create these layouts easily
(See sm100_blockscaled_layout.hpp).
The interface to create SFA and SFB tensor layouts is as follows:

auto problem_shape = make_shape(M, N, K, L);
using SfConfig = Sm1xxBlockScaledConfig<SFVecSize>;

// SFA shape: ((32,4), ceil(M/128)), ((SFVecSize,4), ceil(K/4), L)
auto layout_sfa = SfConfig::tile_atom_to_shape_SFA(problem_shape);
// SFB shape: ((32,4), ceil(N/128)), ((SFVecSize,4), ceil(K/4), L)
auto layout_sfb = SfConfig::tile_atom_to_shape_SFB(problem_shape);

auto tensor_sfa = make_tensor(aptr, layout_sfa);
auto tensor_sfb = make_tensor(bptr, layout_sfb);
// Access SF for for element m,k of A tensor
auto val_a_mk = tensor_sfa(make_coord(m,k,0));

Blackwell SM120 GEMMs

The NVIDIA RTX 5000 Series GPUs introduce support for new narrow precision (4bit and 6bit) block-scaled and non-block-scaled tensor cores. The PTX ISA has extended the mma instructions to support these data formats which are 1x to 4x faster than Ada architecture's fp8 tensor cores. For more detailed information see mma PTX documentation.

CUTLASS 4.0 has added support for these newly introduced narrow precision GEMMs. Similar to the Blackwell SM100 GEMMs, the SM120 GEMMs can be built using the collective builder interface. See examples in examples/79_blackwell_geforce_gemm/ and unit tests listed below.

The data types supported and tensor alignment requirements are the same as the Blackwell SM100 GEMMs. The scale factor layout is also the same as SM100 mentioned above. OpClassTensorOp is used for non-blockscaled narrow precision GEMMs and OpClassBlockScaledTensorOp is used for blockscaled narrow precision GEMMs.

Besides the similarities, there are some key differences from the Blackwell SM100 GEMMs:

Cluster Size Copy Chapter

On Geforce series graphics card, there is no multicast feature therefore the cluster shape is fixed to 1x1x1.

Tensor Layout Copy Chapter

Only TN layout is supported. Matrix A is row major and matrix B is column major.

Pingpong v.s. cooperative kernel schedule Copy Chapter

Similar to Hopper's warp-group GEMM, SM120 GEMMs support both pingpong and cooperative kernel schedules. Pingpong kernel schedule has two groups of 4 MMA warps working on different output tiles, overlapping the mainloop and epilogue, while the cooperative kernel schedule has only one group of 8 MMA warps working on the same output tile. If KernelScheduleAuto is specified, KernelTmaWarpSpecializedCooperative will be selected by default.

Epilogue schedule: Copy Chapter

EpilogueScheduleAuto must be used.

Tile size: Copy Chapter

Below are tables that summarize the valid tile shapes and dispatch policies for SM120 GEMMs. If the output is float_6_t, the tile size in the leading dimension of output tensor must be 128.

Table 16: Valid Tile Shapes and Dispatch Policies for {float8_t, float_6_t, float_4_t} x {float8_t, float_6_t, float_4_t} of SM120 GEMMs

Table 17: Valid Tile Shapes for nv_float4_t x nv_float4_t of SM120 GEMMs

Table 18: Valid Tile Shapes and Dispatch Policies for mx_float4_t x mx_float4_t of SM120 GEMMs

Table 19: Valid Tile Shapes and Dispatch Policies for mx_float4_t x mx_float4_t of SM120 GEMMs

Specialized policies must be used to generate mixed-input-datatype mx_float4_t kernels.

Table 20: Valid Tile Shapes and Dispatch Policies for {mx_float4_t, mx_float6_t, mx_float8_t} x {mx_float4_t, mx_float6_t, mx_float8_t}

Copyright Copy Section

Copyright (c) 2025 - 2026 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.

Media Docs Cpp Code Organization (`media/docs/cpp/code_organization.md`) Copy Section

Code Organization

This document describes the layout of the CUTLASS repository. The main components are:

• CUTLASS Template Library - CUDA Templates for Linear Algebra Subroutines and Solvers (header only)
• CuTe Template Library - CUTLASS's core vocabulary layout type and associated algebra (header only)
• CUTLASS Utilities - Additional templates
• CUTLASS Instance Library - instantiations of CUTLASS templates covering the design space
• CUTLASS Profiler - CUTLASS Library, Profiler, and Utilities
• Examples - SDK examples of CUTLASS Template Library and components
• Media - supporting documentation and media content
• Tests - test components for CUTLASS Template Library and tools

CUTLASS Template Library Copy Chapter

CUDA Templates for Linear Algebra Subroutines and Solvers is a library of CUDA C++ template classes for
performing efficient matrix computations on NVIDIA GPUs.

Like NVIDIA CUB, the components of CUTLASS are organized hierarchically based on the scope of cooperative
elements. For example, warp-level GEMM components perform a matrix multiply collectively by the
set of threads within a warp. The following figure illustrates each layer.

Components are designed to be usable by client applications accessing functionailty at each scope.

CUTLASS Templates are implemented by header files in the following directory structure:

/* Detailed source-code truncated for AI context efficiency. */

See Programming Guidelines for further details about
conventions and design patterns used throughout CUTLASS.

CuTe Copy Chapter

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 packages 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. More documentation
for CuTe can be found in cute/.

Tools Copy Chapter

The tools/ directory contains clients of the CUTLASS Template library and includes the following.

CUTLASS Instance Library Copy Chapter

The CUTLASS Instance Library contains instantiations of the above CUTLASS templates covering supported configurations,
data types, block structure, and tile sizes. These instantiations are procedurally generated using a set of
scripts to span the design space.

tools/
  library/                   # static/dynamic library containing all kernel instantiations of interest
                             # (with some build-level filter switches to compile specific subsets)

    include/
      cutlass/
        library/             # header files for CUTLASS Deliverables Library (in cutlass::library:: namespace)

          handle.h           # implements a host-side API for launching kernels, similar to cuBLAS
          library.h          # defines enums and structs to describe the tiled structure of operator instances          
          manifest.h         # collection of all instances

    src/

python/
    cutlass_library/       # scripts to procedurally generate CUTLASS template instances

      gemm_operations.py
      library.py
      generator.py            # entry point of procedural generation scripts - invoked by cmake
      manifest.py

When CMake is executed, the CUTLASS Instance Library generator scripts are executed to construct a set of
instantiations in build/tools/library/generated/.

CUTLASS Profiler Copy Section

The CUTLASS Profiler is designed to load the CUTLASS Instance Library and execute all operations contained therein.
This command-line driven application constructs an execution environment for evaluating functionality and performance.
It is implemented in

tools/
  profiler/

and may be built as follows.

$ make cutlass_profiler -j

Further details about the CUTLASS Profiler are described here.

CUTLASS Utilities Copy Section

tools/util/ defines a companion library of headers and sources that support the CUTLASS test programs, examples, and other client applications. Its structure is as follows:

tools/
  util/
    include/
      cutlass/
        util/                   # CUTLASS Utility companion library

          reference/            #  functional reference implementation of CUTLASS operators
                                #    (minimal consideration for performance)
            
            detail/
              *

            device/             #  device-side reference implementations of CUTLASS operators
              thread/
              kernel/
                *
            host/               #  host-side reference implementations of CUTLASS operators
              *
          *

More details about CUTLASS Utilities may be found here.

Examples Copy Chapter

To demonstrate CUTLASS components, several SDK examples are implemented in examples/.

CUTLASS SDK examples apply CUTLASS templates to implement basic computations.

examples/
  00_basic_gemm/             # launches a basic GEMM with single precision inputs and outputs

  01_cutlass_utilities/      # demonstrates CUTLASS Utilities for allocating and initializing tensors
  
  02_dump_reg_smem/          # debugging utilities for printing register and shared memory contents
  
  03_visualize_layout/       # utility for visualizing all layout functions in CUTLASS

  04_tile_iterator/          # example demonstrating an iterator over tiles in memory

  05_batched_gemm/           # example demonstrating CUTLASS's batched strided GEMM operation

  06_splitK_gemm/            # exmaple demonstrating CUTLASS's Split-K parallel reduction kernel

  07_volta_tensorop_gemm/    # example demonstrating mixed precision GEMM using Volta Tensor Cores

  08_turing_tensorop_gemm/   # example demonstrating integer GEMM using Turing Tensor Cores

  10_planar_complex/         # example demonstrating planar complex GEMM kernels

  11_planar_complex_array/   # example demonstrating planar complex kernels with batch-specific problem sizes

  12_gemm_bias_relu/         # example demonstrating GEMM fused with bias and relu activation function

  13_fused_two_gemms/        # example demonstrating two GEMMs fused into one kernel

Media Copy Chapter

This directory contains documentation, images, and performance result data which accompanies the CUTLASS library and components.

Tests Copy Chapter

Test programs for CUTLASS. Tests are organized hierarchically, mirroring the organization of source files.

test/                        # unit tests for CUTLASS Template Library
  unit/
    arch/
    core/
    gemm/
      device/
      kernel/
      thread/
      threadblock/
      warp/
    reduction/
      kernel/
      thread/
    transform/
      threadblock/
      *

Tests can be built and run at the top level scope by invoking make test_unit or by building
and explicitly executing each individual target, e.g. cutlass_test_unit_gemm_device.

Tests are configured to specify appropriate GTest filter strings to avoid running except on
architectures where they are expected to pass. Thus, no tests should fail. The actual number
of tests run may vary over time as more are added.

Copyright

Copyright (c) 2017 - 2026 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.

Media Docs Cpp Cutlass 2x (`media/docs/cpp/cutlass_2x.rst`) Copy Section

.. _cutlass_2_x:

CUTLASS 2.x

.. toctree::
:maxdepth: 2

Layouts and Tensors<layout.md>
GEMM API<gemm_api.md>
Tile Iterator Concepts<tile_iterator_concept.md>
Utilities

Media Docs Cpp Cutlass 3x (`media/docs/cpp/cutlass_3x.rst`) Copy Section

.. _cutlass_3_x:

CUTLASS 3.x

.. toctree::
:maxdepth: 2

Design <cutlass_3x_design.md>
GEMM Backwards Compatibility <cutlass_3x_backwards_compatibility.md>
GEMM API <gemm_api_3x.md>

Media Docs Cpp Cutlass 3x Backwards Compatibility (`media/docs/cpp/cutlass_3x_backwards_compatibility.md`) Copy Section

CUTLASS 3.0 GEMM Backwards Compatibility

Although CUTLASS 3.0 restructures the GEMM hierarchy and introduces new types for the
threadblock layer and below, we intend the entire source code to be usable in user applications.
We expect users to be able to #include any source file from CUTLASS 3.0, whether
they implement the 2.x or the 3.x API, without breaking user builds. This means that a single
translation unit should be able to contain any valid kernel regardless of its API version. The
sections below discuss how device and kernel layer type names are made compatible across the
two API versions, and what the users can expect out of the threadblock layer API going forward.

Compatible Device API Copy Chapter

The entry point for CUTLASS's Device GEMM API
is the class
cutlass::gemm::device::GemmUniversalAdapter.
This class lives in the header file
include/cutlass/gemm/device/gemm_universal_adapter.h.

GemmUniversalAdapter is a "universal adapter"
and serves as a common device interface
for both CUTLASS 3.x and CUTLASS 2.x kernels.
Its template parameter GemmKernel,
the GEMM kernel type, can be any of the following:

• cutlass::gemm::kernel::GemmUniversal,
  implementing CUTLASS 3.x API kernels;
• cutlass::gemm::kernel::GemmUniversal,
  implementing CUTLASS 2.x API kernels;
• Any valid CUTLASS 2.x kernel layer GEMM that
  was previously composable with device::GemmUniversalAdapter

Users implementing new kernels in either API should prefer
using kernel::GemmUniversal as the kernel type
and compose it with device::GemmUniversalAdapter.
Users with existing kernel::Gemm kernels
can continue to use them as template arguments
of device::GemmUniversalAdapter. They can adopt
GemmUniversal as a gradual migration path,
since GemmUniversal accepts either 3.0 or 2.x collectives.
Please see the next section for kernel::GemmUniversal for details.

GemmUniversalAdapter presents a single
host-side interface to both 3.0 and 2.x kernels.
CUTLASS accomplishes this by
specializing GemmUniversalAdapter's implementation
on either 2.x API implementing kernel layer GEMMs, or 3.x API
implementing kernel layer GEMMs (as detected by gemm::detail::IsCutlass3GemmKernel
discussed below). As a result, GemmUniversalAdapter's behavior
might differ between the two specializations.

Device API design differences Copy Section

In CUTLASS 2.x, the Device API was more closely tied
to the Kernel API. In CUTLASS 3.0, the Device API
accepts any kernel type that meets the Kernel API
interface requirements. CUTLASS 3.0's Device API code is
parameterized by the kernel type, but this code
is generic; the same code works for any kernel type.

The device layer compatibility interface, device::GemmUniversalAdapter,
also provides reflective mappings from 3.0-specific types
back to the closest possible 2.x equivalent types. This is discussed further in the section below.

CUTLASS 3.0's device::GemmUniversalAdapter also exposes some new APIs that the 2.x device::GemmUniversalAdapter implementation does not. Most notably, this includes the ability to bypass the GemmKernel::Arguments to GemmKernel::Params lowering.

// Primary run() entry point API that is static allowing users to create and manage their own params.
static Status
run(Params& params, cudaStream_t stream = nullptr);

This new API is useful for the following scenarios.

• Running again does not require reinvoking GemmKernel::to_underlying_arguments()
• Manual control over construction of GemmKernel::Params for custom kernels with custom stride types
• Fully static problem shapes and strides for bespoke kernels where no argument mapping needs to take place

Compatible Kernel API Copy Chapter

CUTLASS 3.x API shares the kernel layer API with CUTLASS 2.x
through the single entry point type cutlass::gemm::kernel::GemmUniversal.
All kernel layer GEMMs are viewed as a composition of a collective mainloop
and a collective epilogue.

kernel::GemmUniversal implements both 2.x and 3.x APIs

The entry point for CUTLASS's kernel API is the class
cutlass::gemm::kernel::GemmUniversal.
This class' declaration lives in the header file
include/cutlass/gemm/kernel/gemm_universal.hpp.

/*
 * Stateless universal device GEMM kernel type that treats GEMM as
 * a composition of a collective mainloop and a collective epilogue.
 * SFIANE shims both 2.x and 3.0 API kernels based on ProblemShapeOrThreadblockMma_.
**/
template <
  class ProblemShapeOrThreadblockMma_,
  class CollectiveMainloopOrEpilogue_,
  class CollectiveEpilogueOrThreadblockSwizzle_,
  class TileScheduler_ = void,
  class Enable = void
>
class GemmUniversal;

We call this class "universal" because it can be built
using either the CUTLASS 3.0 or the 2.x mainloops and epilogues.
If GemmUniversal's first template argument
(ProblemShapeOrThreadblockMma_) is a cute::tuple,
then GemmUniversal assumes that
the remaining three template arguments
(the mainloop, epilogue, and grid swizzle)
implement the 3.0 APIs.
Otherwise, GemmUniversal assumes that
the remaining three template arguments
implement the 2.x APIs.
All the template arguments must be either
CUTLASS 3.0 or CUTLASS 2.x types. For example,
GemmUniversal does not permit using
a 2.x mainloop with a 3.0 collective epilogue.

CUTLASS 3.x implements various embodiments of kernel::GemmUniversal.
Each kernel layer schedule is specialized
for a GEMM scheduling algorithm and GPU architecture.
Specializations of kernel::GemmUniversal for 3.0 APIs live in
any of various gemm_*.hpp files in the directory
include/cutlass/gemm/kernel/.
The specialization to which to dispatch is decided through the dispatch policy's Schedule type.

Specializations for 2.x APIs live in the header file
include/cutlass/gemm/kernel/gemm_universal.h.

Kernel API design differences Copy Section

The CUTLASS 2.x Kernel API was more closely tied
to the Device API, as we mentioned above.
In particular, the 2.x Device API specified the grid shape
used to launch the Kernel API.
In CUTLASS 3.0, the Kernel API controls its own grid shape,
while the device adapter simply queries the kernel with which it needs to be launched.

This change is required to support various kernel schedules
that may need their own schedule specific grid planning logic.
For example, persistent kernel schedules generally only launch with
as many threadblocks as the number of multiprocessors on the GPU.

All CUTLASS 3 kernel::GemmUniversal specializations expose the following (static) API:

// Returns true if the kernel can execute the provided GEMM arguments.
static bool
can_implement(Arguments const& args);

// Returns a dim3 representing the threadblock shape.
static dim3
get_block_shape();

// Returns a dim3 representing the grid shape in terms of threadblocks.
static dim3
get_grid_shape(Params const& params);

The device adapter simply queries the kernel for these three before launching it on the device.
CUTLASS 3.0 provides a meta-function to detect whether a cutlass::gemm::kernel::* implements
the 3.x API or 2.x API:

// include/cutlass/gemm/gemm.h

namespace cutlass:gemm::detail {

// The following metafunction is used to detect whether a
// `kernel::Gemm` or `kernel::GemmUniversal` implements the CUTLASS 3.x API,
// by checking whether the problem shape type is aliased within.
template <class GemmKernel, class = void>
struct IsCutlass3GemmKernel;

} // namespace cutlass:gemm::detail

Users can dispatch their generic code against 2.x and 3.x specializations with
this as a type trait for the kernel API version.

Threadblock API and Inner Loops Copy Chapter

Much of the CUTLASS 3 GEMM hierarchy for mainloops and inner loops diverges
from that of CUTLASS 2.x. With that also comes the introduction of the
cutlass::gemm::collective layer as a direct replacement and a superset
of the 2.x cutlass::gemm::threadblock layer. Going forward,
CUTLASS 3.x will discontinue new developments in the following namespaces.

• cutlass::*::threadblock::*
• cutlass::*::warp::*
• cutlass::gemm::thread::*
• cutlass::arch::* (except barrier.h)

cutlass::gemm::collectives are a superset of the threadblock layer where
all new mainloops will be developed. Users should look to the CollectiveMma type
if they wish to author custom mainloop code in the 3.x API.

Similarly, for the GEMM inner loops, cute::MMA_Atoms replace the
gemm::warp and gemm::thread layer code. Going forward, all new PTX instructions
and associated metadata development will occur directly inside cute/arch/*.hpp and cute/atom/*.hpp.

The desired inner loop MMA iteration order and tiling can be achieved through careful
selection of the atom layout, value layout, and permutations of the cute::TiledMma.

For epilogues, the cutlass::epilogue::collective layer replaces cutlass::threadblock::collective. However, the thread-level epilogue elementwise operations
in cutlass::epilogue::thread will continue to be used in 3.x kernels as well, albeit, with
a more idiomatic epilogue vectorization strategy.
Example 50
shows how to use 2.x epilogue thread operators with 3.0 API kernels.

Porting from 2.x to 3.0 API Copy Chapter

CUTLASS 2.x layout tags and CUTLASS 3.0 major modes Copy Section

CUTLASS 2.x and CUTLASS 3.0 use both
different wording and different types
to describe the permitted layouts
of GEMM's input matrices A and B.

CUTLASS 3.0 does not use the terms "column major"
or "row major" to describe matrix layouts.
Starting with CUTLASS 3.0, adoption of CuTe allows us to decouple

• the coordinate mode order (logical shape) of layouts from

• the index space stride order of the backing storage.

In line with our switch to a conceptual GEMM hierarchy, we view the major modes not from a BLAS-3 perspective.
Rather, we divide the modes into two categories.

• "Inner modes" or "K-modes" are contracted over during the GEMM.
  Therefore, they are not present in the output tensor.

• "Outer modes" or "MN-modes" are preserved in the output.

Now, instead of RowMajor or ColumnMajor, whose major stride depends on whether we are referring to the
A or the B matrix, we uniformly employ the "K major" or "MN major" terminology and enforce the convention of all tensors having the shape [M/N, K, L] regardless of which mode is major. That is,

• the input matrix A has shape M x K,
• the input matrix B has shape N x K, and
• the input/output matrices C/D have shape M x N.

Note that this convention for B
differs from the BLAS's GEMM interface,
which specifies that B has shape K x N.

CUTLASS 3.0 uses these names of the modes
to specify which mode of a matrix has stride 1.
For the matrix A,

• "M major" means that the matrix is stride 1
  in the M mode, and
• "K major" means that the matrix is stride 1
  in the K mode.

For the matrix B,

• "N major" means that the matrix is stride 1
  in the N mode (which for B is mode 0,
  because the convention is that B is N x K); and
• "K major" means that the matrix is stride 1
  in the K mode (which for B is mode 1).

CUTLASS 2.x defines "layout tag" classes
cutlass::layout::ColumnMajor and cutlass::layout::RowMajor,
that live in the header file
cutlass/layout/matrix.h.
The interpretation of these layouts in GEMM
depends on whether they are applied
to the input matrix A or B. For the matrix A, "column major" means
that mode corresponding to M extent has stride 1,
and "row major" means that mode corresponding to K extent has stride 1.
This is the usual computer science definition
of column major and row major for a rank-2 array.
For the matrix B, the opposite holds:
"column major" means that mode corresponding to N extent has stride 1,
and "row major" means that mode corresponding to K extent has stride 1.

Using the convention of [outer, inner, batch] mode order for tensor logical shapes
avoids potential confusion with the meaning of column major and row major
changing depending on whether they are applied to A or B.

The table below summarizes our mode order convention and
mapping of 2.x layout tags to corresponding M-major, N-major, or K-major strides.

Notice that in CUTLASS 3.0, interpretation of layouts no longer changes based on
whether we are talking about the A or B matrix. M and N major inputs always have a
static size-1 stride in their 0th (outer) mode. Similarly, K major inputs
always contain the static size-1 stride in their 1st mode. This uniformity in stride order
allows us to represent tensor layouts much more cleanly and treat both A and B equally in our interfaces.
See for example the following snippet from our kernel/sm70_gemm.hpp
for Ampere kernel schedules.

// Represent the full tensors
Tensor mA_mkl = make_tensor(make_gmem_ptr(params.mainloop.ptr_A), make_shape(M,K,L), params.mainloop.dA); // (m,k,l)
Tensor mB_nkl = make_tensor(make_gmem_ptr(params.mainloop.ptr_B), make_shape(N,K,L), params.mainloop.dB); // (n,k,l)

// Get batch slice
Tensor mA_mk = mA_mkl(_,_,get<3>(blk_coord_mnkl)); // (m,k)
Tensor mB_nk = mB_nkl(_,_,get<3>(blk_coord_mnkl)); // (n,k)

// Slice to get the tiles for which this thread block is responsible
Tensor gA = local_tile(mA_mk, blk_shape, take<0,3>(blk_coord_mnkl), Step<_1, X,_1>{}); // (BLK_M,BLK_K,k)
Tensor gB = local_tile(mB_nk, blk_shape, take<0,3>(blk_coord_mnkl), Step< X,_1,_1>{}); // (BLK_N,BLK_K,k)

As seem in this snippet, all input tensors have the logical shape [outer, inner, batch],
and the strides could represent either outer or inner
(or any other complex hierarchical stride) major storage.
CuTe layouts always maintain the logical consistency of the coordinate spaces regardless of the strides.

By convention, in CUTLASS 3.0, we treat the M and N mode as the 0th mode,
and K mode as the 1st mode of the stride.

Conversions between 2.x tags and 3.0 types Copy Section

Starting with CUTLASS 3.0, all layouts are described using
cute::Shape and cute::Stride which compose into a cute::Layout<Shape, Stride>.
In CUTLASS 2.x, various layout tags such as cutlass::layout::RowMajor are used to specialize
template implementations. These tag types only encode information about the tensor strides,
as 2.x layouts did not incorporate any concept of tensor shape in the layout tags themselves.
Users may find a need to convert between CUTLASS 2.x layout tags, and 3.0
CuTe stride types. CUTLASS 3.0 gemm::collective::CollectiveBuilder interfaces
also accept these 2.x layout tags as input parameters in their template API as a convenience for users.
At every entry point into CUTLASS 3.0, these tags get converted to their corresponding CuTe Stride type with
metafunctions that best approximate their corresponding cute::Stride.

• cutlass::gemm::detail::TagToStrideA_t<LayoutTag>
• cutlass::gemm::detail::TagToStrideB_t<LayoutTag>
• cutlass::gemm::detail::TagToStrideC_t<LayoutTag>

By convention, and to match user expectations, the cute::Stride types that these
map onto always contain one static mode corresponding to the layout tag, and two 64-bit
dynamic stride modes corresponding to the minor mode and the batch mode. Batch
mode is included by default as all CUTLASS 3.0 kernels support packed batch-mode GEMMs
out of the box.

The cutlass/gemm/gemm.h#440
header file includes functions
that can be useful for converting
from CUTLASS 3.0 cute::Strides back to CUTLASS 2.x layout tags.

• cutlass::gemm::detail::StrideToLayoutTagA_t<CuteStride>
• cutlass::gemm::detail::StrideToLayoutTagB_t<CuteStride>
• cutlass::gemm::detail::StrideToLayoutTagC_t<CuteStride>

These metafunctions take the CuTe Stride as a template parameter and
attempt to find the size-1 stride in the idiomatic M, N, or K modes
to best approximate a corresponding 2.x layout tag type.
Note that this may not work in general for any cute::Stride
as the mapping between the stride and tag type is not bijective.

These mapping utilities are kept in a detail namespace
as we do not guarantee stability of their implementation.
Their behavior may change in future releases as we add new features.
However, we do expect these type names to remain stable. For users who want
these 2.x reflective types from an assembled kernel with a more stable API,
the specialization of cutlass::gemm::device::GemmUniversalAdapter
for CUTLASS 3.0 kernel provides all aliases for all 2.x type aliases
in addition to the layout tags. You can see how they are used in the header file
cutlass/gemm/device/gemm_universal_adapter.h.
Here is an excerpt.

/* Detailed source-code truncated for AI context efficiency. */

CUTLASS's library and profiler use these reflective interfaces to
obtain the kernel's configuration parameters. Users can use these to approximate the CUTLASS 2.x types
for 3.0 API kernels. However, the reflective interfaces cannot always match the types exactly,
as the mappings are not always bijective.

Copyright Copy Section

Copyright (c) 2023 - 2026 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.

Media Docs Cpp Cutlass 3x Design (`media/docs/cpp/cutlass_3x_design.md`) Copy Section

CUTLASS 3.0 Design

CUTLASS 3.0 is a major enhancement over the abstractions of CUTLASS 2.x
and aims to make usage of all layers of the GEMM hierarchy easier and more composable
while still achieving peak performance on Hardware.

CUTLASS 3.0 design goals Copy Chapter

CUTLASS 3.0 has the following design goals, in no particular order.

• Simplify expressing and manipulating data and thread layouts across
  the GEMM hierarchy with CuTe layouts and layout algebra.

• Improve code readability and learning curve by
  reducing the number of named types.

• Functional correctness by default,
  actionable static asserts otherwise.

• Single, clear points of performance tuning and custom kernel extensions.

• Support for NVIDIA Hopper GPUs with great performance using
  features such as Tensor Cores, tensor memory accelerator, and thread block clusters.

A new Conceptual GEMM Hierarchy Copy Chapter

CUTLASS 2.x decomposes the moving parts of a GEMM operation
across a hierarchy that closely mirrors the organization of GPU
architectures. This discussed in detail within the
CUTLASS 2.x GEMM API documentation.
This design, however, sometimes results in a coupling that is too tight
to extend to newer GPU features that might not fit into the same architectural
hierarchy. For instance, Hopper's warp-group wide instructions do not naturally
fit into any warp or thread layer GEMM concept in CUTLASS 2.x. Even for Volta tensor cores,
instructions that atomically exist at the quad-pair granularity are first tiled at
the warp level before use. This hints at the brittleness of the abstraction power.

CUTLASS 3.0 detaches its interface layers from the hardware,
centering them instead around the natural structure of GEMM algorithms
not tied to any particular GPU generation.
This makes CUTLASS's code more robust to GPU architecture evolution,
less prone to implementation detail leakage, and provides users
with a consistent interface to hardware acceleration regardless of
the architecture specific details.

The new conceptual GEMM hierarchy is discussed in detail in the dedicated
CUTLASS 3.0 GEMM API documentation readme,
along with code examples of the core concepts and types.

Adoption of CuTe Layout and Tensors Copy Chapter

CUTLASS 3.0 introduces 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 packages the type, shape, memory space, and layout of data, while performing the complicated indexing for the user.

CUTLASS 3.0 adopts CuTe throughout the GEMM hierarchy in its templates, greatly simplifying the design,
improving code composability, and readability. More documentation specific to CuTe can be found in its dedicated documentation directory.

Programming massively parallel systems with various layers of logical thread and data hierarchies is not a trivial task.

• cute::Layouts always maintain logical consistency of their coordinates,
  allowing us to check pre- and post-conditions at compile time for all static inner loops.
• Explicit thread to data mapping allows users and kernel authors to inspect and reason about operations
  from a single point in the source code.
• Layouts provide a single point of performance tuning, as most optimizations can be done by careful
  selection of thread and data layouts.
• Formalized algebra makes manipulation of and reasoning about thread->data mapping explicit in source code.
• Single vocabulary type (cute::Layout) subsumes every iterator and layout in CUTLASS 2.x CUTLASS 2.x uses many bespoke thread maps, iterators, and data layouts. Iterators are fundamentally 1-D, whereas most layouts we encounter in the GPU hierarchy are fundamentally n-D.

Reducing the number of named types and iterator concepts Copy Chapter

CUTLASS 2.x design preferred introducing bespoke named types for each
architecture specific thread and data layout. For instance, gemm::treadblock namespace
contains implementation for MmaMultistage, MmaPlanarComplexMultistage, MmaPipelined etc.
despite them providing mainloops for GEMMs. To spell these types the same way in generic code,
CUTLASS 2.x provides aliases through its default_x_configuration.h files, however,
these aliases make the code much harder to read as the user has to perform type substitution
mentally in order to understand the codebase.

CUTLASS 3.0 greatly reduces the number of named types used throughout by

• Replacing all iterator concepts for all memory domains with cute::Tensors
• Dispatching mainloop and epilogue implementations on tag-dispatch policies rather than naming new types
• Dispatching kernel layer schedules on tag-dispatch policies rather than naming new types

Reducing the number of named types has many benefits:

• It makes writing generic code easier, as the primary type names share the same lexical
  without aliasing through configuration providers.
• It flattens the learning curve of CUTLASS by greatly reducing the mental context required
  as the library only exposes a handful of named types.
• It provides a clear, singular extension point for users to plug in their customizations
  through the dispatch policies.

Correctness by default, Performance through clear, individual points of tuning Copy Chapter

CUTLASS 2.x maintained its thread layouts as implicit indexing math implemented
as a part of 1D iterators. This meant that the thread to data layout mapping
was implicit in the imperative structure of the C++ code itself and did not have
a formal algebra we could use to manipulate these mappings. Each iterator
had to re-implement its indexing and mapping logic. This made it hard to learn
how this mapping was performed for existing iterators, and even harder to
implement custom layout functions for the core inner loops of a GEMM.

CUTLASS 3.0 replaces all iterator concepts from CUTLASS 2.x
with a single layout type for thread and data tensors.
CuTe's formalized layout algebra is then used at every layer of
the GEMM hierarchy to manipulate the mapping between the two.
CuTe layouts always maintain logical consistency, and for fully static layouts
(such as in the core unrolled inner loops), provide
compile time checks that break builds if this consistency is violated.
In this way, CuTe reifies the thread-to-data-layout mapping,
makes it easier to write code that is "correct by construction".
If the code compiles, it's probably correct.

Copyright Copy Section

Copyright (c) 2017 - 2026 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.

Media Docs Cpp Dependent Kernel Launch (`media/docs/cpp/dependent_kernel_launch.md`) Copy Section

Dependent kernel launches

The Hopper and Blackwell architectures supports a new feature through which two kernels in the same stream can
overlap their execution, named
Programmatic Dependent Launch (PDL).
This allows kernels with conflict in global memory to programmatically and safely overlap portions
of their execution. Primary kernel can signal it is about to finish execution, and the next kernel is expected to
programmatically wait on the previous kernel to finish flushing its memory.

We enable PDL by setting a flag through the extended CUDA launch APIs. All CUTLASS kernels with PDL support
will wait on the prior kernel to flush its output to memory and signal the next kernel to start. This means
they can safely be dropped in with any other set of kernels using PDL as long as they also adhere to waiting on
the prior to flush its memory as well.

For more information, we refer you to the PDL section in the CUDA Programming Guide.

Using dependent launch in CUTLASS Copy Chapter

When building CUTLASS, you can use the CUTLASS_ENABLE_GDC_FOR_SM90 and CUTLASS_ENABLE_GDC_FOR_SM100 macro
respectively to enable PDL-related instructions:

cmake . -DCUTLASS_ENABLE_GDC_FOR_SM90=1

Note that this only adds PDL-related instructions to the kernels, but to actually allow a dependent
launch, you must also run your GEMM kernel with PDL:

gemm.run(
  /* stream = */ stream,
  /* cuda_adapter = */ nullptr,
  /* launch_with_pdl = */ true
);_

Model-Aware Optimizations with PDL Copy Chapter

In example 63, we use PDL to explicitly optimize for
performance of kernels where we know that one of the input matrices (our weights) will not be produced by a prior
kernel. In that case, we only need to wait on the prior kernels memory flush in order to load the other input matrix
(our activations). During our prologue, we can prefetch our weights to improve performance for memory bandwidth-bound
problem sizes. For more information, we refer the reader to the example.

Copyright Copy Chapter

Copyright (c) 2017 - 2026 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.

Media Docs Cpp Doxygen Mainpage (`media/docs/cpp/doxygen_mainpage.md`) Copy Section

CUTLASS 3.0

CUTLASS 3.0 - January 2023

CUTLASS is a collection of CUDA C++ template abstractions for implementing
high-performance matrix-multiplication (GEMM) at all levels and scales within CUDA.
It incorporates strategies for hierarchical decomposition and data movement similar
to those used to implement cuBLAS. CUTLASS decomposes these "moving parts" into
reusable, modular software components abstracted by C++ template classes. These
components can be specialized
and tuned via custom tiling sizes, data types, and other algorithmic policies. The
resulting flexibility simplifies their use 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 8-bit integer, half-precision floating
point (FP16), single-precision floating point (FP32), and double-precision floating
point (FP64) types. Furthermore, CUTLASS exploits the Tensor Cores and asynchronous
memory copy operations of the latest NVIDIA GPU architectures.

What's New in CUTLASS 3.0

For an overview of CUTLASS 3.0's GEMM interface levels,
please refer to the
CUTLASS 3.0 GEMM API document.
To learn how to migrate code using CUTLASS 2.x's interface
to CUTLASS 3.0, please refer to the
backwards compatibility document.

GEMM examples

For a code example showing how to define
a GEMM kernel using CUTLASS, please refer to
the quickstart guide.
The examples directory
has a variety of examples.

Copyright

Copyright (c) 2017 - 2026 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.

Media Docs Cpp Efficient Gemm (`media/docs/cpp/efficient_gemm.md`) Copy Section

Efficient GEMM in CUDA

CUTLASS implements the hierarchically blocked structure described in
CUTLASS: Fast Linear Algebra in CUDA C++
and the CUTLASS GTC2018 talk.

Hierarchical Structure Copy Chapter

The basic triple loop nest computing matrix multiply may be blocked and tiled to match
concurrency in hardware, memory locality, and parallel programming models. In CUTLASS,
GEMM is mapped to NVIDIA GPUs with the structure illustrated by the following loop nest.

/* Detailed source-code truncated for AI context efficiency. */

This tiled loop nest targets concurrency among

• threadblocks,
• warps, and
• CUDA and Tensor Cores.

It takes advantage of memory locality within

• shared memory and
• registers.

The figure below illustrates the flow of data within this structure.
This is the hierarchical GEMM computation embodied by CUTLASS. Each stage depicts a
nested level of tiling which corresponds to a layer of concurrency within the CUDA execution model and to a
level within the memory hierarchy, becoming increasingly finer moving left to right.

Threadblock-level GEMM Copy Section

Each threadblock computes its portion of the output GEMM by iteratively loading tiles of input
matrices and computing an accumulated matrix product. At the threadblock level, data are loaded from
global memory. The blocking strategy in general is key to achieving efficiency. However, the programmer
must balance multiple conflicting goals. A
larger threadblock means fewer fetches from global memory, thereby ensuring that DRAM bandwidth
does not become a bottleneck.
However, large threadblock tiles may not match the dimensions of the problem well. If either the
GEMM M or N dimension is small, some threads within the threadblock may not perform meaningful
work, as the threadblock may be partially outside the bounds of the problem. If both M and N
are small while K is large, this scheme may launch relatively few threadblocks and fail to
make full use of all multiprocessors within the GPU. Strategies to optimize performance for this case,
as described in the section Parallelized Reductions,
partition the GEMM K dimension across multiple threadblocks or multiple warps. These threadblocks
or warps compute matrix products in parallel; the products are then reduced to compute the result.

In CUTLASS, the dimensions of the threadblock tile are specified as ThreadblockShape::{kM, kN, kK}
and may be tuned to specialize the GEMM computation for the target processor and dimensions of
the GEMM problem.

Warp-level GEMM Copy Section

The warp-level GEMM maps to the warp-level parallelism within the CUDA execution model. Multiple
warps within a threadblock fetch data from shared memory into registers and perform computations.
Warp-level GEMMs may be implemented either by TensorCores issuing
mma.sync
or wmma
instructions, or by thread-level matrix computations issued to CUDA cores.
For maximum performance, access to shared memory should be bank conflict free. To maximize data
reuse within the warp, a large warp-level GEMM tile should be chosen.

Thread-level GEMM Copy Section

At the lowest level of blocking, each thread is responsible for processing a certain number of
elements. Threads cannot access each other's registers, so we choose an organization that enables
reuse of values held in registers for multiple math instructions. This results in a 2D tiled
structure within a thread, in which each thread issues a sequence of independent math instructions
to the CUDA cores and computes an accumulated outer product.

SGEMM, IGEMM, HGEMM, and DGEMM are computed by SIMT math instructions issued by thread-level matrix multiply
procedures.

Epilogue Copy Chapter

The above code focuses only on the matrix multiply computation C = AB whose result is
held in the registers of each thread within the threadblock. The mapping of logical elements
in the output tile to each thread is chosen to maximize performance of the matrix multiply
computation but does not result in efficient, coalesced loads and stores to global memory.

The epilogue is a separate phase in which threads exchange data through shared memory then
cooperatively access global memory using efficient striped access patterns. It is also
the phase in which linear scaling and other elementwise operations may be conveniently
computed using the matrix product results as inputs.

CUTLASS defines several typical epilogue operations such as linear scaling and clamping,
but other device-side function call operators may be used to perform custom operations.

Optimizations Copy Chapter

The hierarchical structure described above yields an efficient mapping to the CUDA execution model and
CUDA/TensorCores in NVIDIA GPUs. The following sections describe strategies for obtaining peak performance
for all corners of the design space, maximizing parallelism and exploiting data locality wherever possible.

Pipelining Copy Section

The blocked structure demands a large storage allocation within the registers of each CUDA thread. The
accumulator elements typically occupy at least half a thread's total register budget. Consequently,
occupancy -- the number of concurrent threads, warps, and threadblocks -- is relatively low compared
to other classes of GPU workloads. This limits the GPU's ability to hide memory latency and other stalls
by context switching to other concurrent threads within an SM.

To mitigate the effects of memory latency, CUTLASS uses software pipelining to overlap memory accesses
with other computation within a thread. CUTLASS accomplishes this by double buffering at the
following scopes.

• Threadblock-scoped shared memory tiles: two tiles are allocated in shared memory.
  One is used to load data for the current matrix operation,
  while the other tile is used to buffer data loaded from global memory
  for the next mainloop iteration.

• Warp-scoped matrix fragments: two fragments are allocated within registers.
  One fragment is passed to CUDA and TensorCores during the current matrix computation,
  while the other is used to receive shared memory fetch returns
  for the next warp-level matrix operation.

The following diagram illustrates the efficient, pipelined mainloop body used in CUTLASS GEMMs.

Threadblock Rasterization Copy Section

To maximize reuse of data held in the last level cache, CUTLASS defines several functions to
affect the mapping of threadblocks to logical partitions of the GEMM problem. These map
consecutively launched threadblocks to packed two-dimensional regions of the partitioned GEMM
problem to increase the probability that these will access the same tiles of global memory at
approximately the same time.

Several functions are defined in cutlass/gemm/threadblock_swizzle.h.

Parallelized Reductions Copy Section

Split K - reduction across threadblocks

Matrix product computations expose parallelism among O(MN) independent inner product
computations. For sufficiently large problem sizes, a GEMM kernel in CUTLASS may approach
the theoretical maximum computational throughput. For small problems, however, there are
too few threadblocks to efficiently occupy the entire GPU.

As a recourse, parallelizing the reduction performed during the inner product computation
enables more threadblocks to execute concurrently while still taking advantage of the throughput
benefits of large threadblock-level GEMM tiles.

CUTLASS implements parallel reductions across threadblocks by partitioning the GEMM K dimension
and launching an additional set of threadblocks for each partition. Consequently, we refer to
this strategy within CUTLASS as "parallel reduction splitK." The "parallel reduction splitK" strategy
requires the execution of 2 kernels: partitionedK GEMM, and batched reduction.

PartitionedK GEMM resembles one flavor of batched strided GEMM. Instead of requiring users
to specify the problem size of each batch, partitionedK GEMM asks for the overall problem size and the
number of partitions that will be applied along the K dimension for operands A and B. For example,
parameters of m=128, n=128, k=4096 and partition=16 will result in 16 batched strided GEMMs
with each batch of m=128, n=128, k=256. PartitionedK also allows scenario where k is not divisible
by the partition count.

For example, parameters of m=128, n=128, k=4096 and partition=20
will result in 20 batched strided GEMMs.
The first 19 batches will have m=128, n=128, and k=4096/20=204,
and the last batch will have m=128, n=128, and k=220.

The batched reduction kernel takes as input the output (C) of partitionedK GEMM,
and performs a reduction along the K-dimension.
Users must manage workspace memory to store this intermediate result.

Sliced K - reduction across warps

Similar to the split-k scenario, sliced-k aims at improving the efficiency of kernels
with smaller M and N dimensions, but large K dimension.
At the thread-block level, the parameters CtaTileN and CtaTileM expose parallelism
by partitioning the work among warps.
Larger warpTiles expose better instruction-level parallelism (ILP) and reuse,
but also limit the number of warps running per threadblock, which reduces efficiency.

In order to improve efficiency in such scenarios, partitioning the warpTiles also along ctaTileK
helps use the hardware more efficiently by allowing more warps to run concurrently in a CTA.
Sliced-k kernels break down a threadblock's computation among participating warps
not just among the CtaTileN, CtaTileM dimension, but also the CtaTileK dimension.
Thus, sliced-k entails a small cost in form of a reduction
which has to happen at the end among the participating warps.
This is because each warp computes using only a "slice" of CtaTileK,
so each warp only has a partial sum before the reduction.

Hopper Warp Specialization Copy Section

Note: the following section on warp-specialization contains details that are specific
to the Hopper kernel design. Blackwell SM100 kernels have a substantially different warp-specialization structure,
however, the concept of separating out producer and consumer agents still applies.

Starting with Hopper, CUTLASS 3.0 incorporates the concept of Warp Specialization
as part of the kernel design. A thread block is partitioned into two sets of warps, producer warp group and consumer warp group. The producer warp group loads data from global memory into shared memory buffers using the new Tensor Memory Accelerator (TMA).

Producer warp group (DMA) waits for the shared memory buffers to be signaled as empty by the consumer warp group using the newly added Async Pipeline class (refer). Once the data is written into the shared memory, TMA is also updates the barrier associated with that stage to notify affected threads that the buffer has been filled. The Consumer warp group (MMA) on the other hand waits for the producer warp group to signal that the buffer is filled and then launches tensor core MMA operations. Finally, the consumer warp group releases the buffers for the next set of TMA loads to happens.

Warp-Specialized Persistent Cooperative kernel design

Another flavor of Warp-Specialized kernel design being introduced starting with Hopper is the Warp-Specialized Persistent Cooperative kernel. Like the Warp-Specialized kernel, the concepts of warp groups and barrier synchronization between warp groups remain the same in the cooperative design.
The distinctive feature of the Warp-Specialized Persistent Cooperative kernel are the following :

• Persistent thread blocks launched to occupy as many SMs as mentioned in the KernelHardwareInfo struct. These persistent thread blocks are used to tile the output and thus (potentially) compute multiple output tiles through their lifetime. The main benefit this adds is amortization of the thread-block launch and kernel prologue overheads which are typical of all kernels.
• Presence of two consumer warp groups cooperating on the same output tile by splitting the tile in half across the M dimension. This allows for larger tile sizes to be enabled - since the register pressure per consumer warp group is reduced - and hence improving performance.

Since each thread block now computes multiple output tiles, the shape of the grid launch and the scheduling of tiles to the thread blocks is managed using the new Tile Scheduler. The Tile Scheduler considers the shape of the clusters as well as the available number of available SMs to compute a valid scheduling of the output tiles to launched thread blocks.

Warp-Specialized Persistent Ping-Pong kernel design

The third kernel design is the Warp-Specialized Persistent Ping-Pong kernel.
Like the Warp-Specialized Persistent Cooperative, kernel the concepts of warp groups, barrier synchronization between warp groups, and the shape of the grid launch remain the same in the persistent ping-pong design.
The distinctive feature of the Warp-Specialized Persistent Ping-Pong kernel is the following :

• The two consumer warp groups are assigned a different output tile using the Tile Scheduler. This allows for epilogue of one consumer warp group to be overlapped with the math operations of the other consumer warp group - thus maximizing tensor core utilization.
• The producer warp group synchronizes using the Ordered Sequence Barrier to fill buffers of the two consumer warp groups one after the other in order.

Resources

The following additional resources describe design and implementation details of GEMMs
targeting NVIDIA GPUs.

• Developing CUDA Kernels to Push Tensor Cores to the Absolute Limit on NVIDIA A100. (SR 21745)
• CUTLASS: Fast Linear Algebra in CUDA C++
• CUTLASS: SOFTWARE PRIMITIVES FOR DENSE LINEAR ALGEBRA AT ALL LEVELS AND SCALES WITHIN CUDA
• Programming Tensor Cores: NATIVE VOLTA TENSOR CORES WITH CUTLASS
• CUDA Programming Guide: warp matrix functions
• Matrix Multiply Accumulate Instructions

Copyright

Copyright (c) 2017 - 2026 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.

Media Docs Cpp Functionality (`media/docs/cpp/functionality.md`) Copy Section

Functionality

Note : CUTLASS-3 requires users to use CUDA 11.4 or newer, and SM70 or newer, for the target toolkit and architecture, respectively.

• N - Column Major Matrix
• T - Row Major matrix
• {N,T} x {N,T} - All combinations, i.e., NN, NT, TN, TT
• NHWC - 4 dimension tensor used for convolution
• NCxHWx - Interleaved 4 dimension tensor used for convolution
• f - floating point
• s - signed int
• b - bit
• cf - complex float
• bf16 - bfloat16
• tf32 - tfloat32
• Simt - Use Simt CUDA Core MMA
• TensorOp - Use Tensor Core MMA
• SpTensorOp - Use Sparse Tensor Core MMA
• WmmaTensorOp - Use WMMA abstraction to use Tensor Core MMA

Device-level GEMM Copy Chapter

The following tables summarize device-level GEMM kernels in CUTLASS, organized by opcode class, data type, and layout.
Hyperlinks to relevant unit tests demonstrate how specific template instances may be defined.

CUTLASS 3.x Kernels Copy Section

CUTLASS 2.x Kernels Copy Section

Device-level Implicit GEMM convolution Copy Chapter

The following table summarizes device-level implicit GEMM convolution kernels in CUTLASS, organized by opcode class, data type, and layout.
Hyperlinks to relevant conv2d fprop unit tests demonstrate how specific template instances may be defined.
One can find and/or create equivalent dgrad and wgrad convolutional operators.

Warp-level Matrix Multiply with Tensor Cores Copy Chapter

The following table summarizes supported warp level shapes for each TensorOp instruction.

TensorOp instructions depend on a permuted shared memory layout that can be efficiently
loaded from. The following tables summarize the destination shared memory layout that
can be targeted by matrix operands. It is assumed that each thread loads 128b vectors
from global memory with layout specified in the column "GMEM Layout."

TensorOp 8-by-8-by-4.

TensorOp 16-by-8-by-8.

TensorOp 16-by-8-by-8.

TensorOp 16-by-8-by-16.

TensorOp 8-by-8-by-4.

TensorOp 8-by-8-by-16.

TensorOp 16-by-8-by-32.

TensorOp 8-by-8-by-32.

TensorOp 16-by-8-by-64.

TensorOp 8-by-8-by-128.

SpTensorOp 16-by-8-by-16.

SpTensorOp 16-by-8-by-32.

SpTensorOp 16-by-8-by-64.

SpTensorOp 16-by-8-by-128.

Warp-level Matrix Multiply with CUDA WMMA API Copy Chapter

The following table summarizes supported warp level shapes for each WmmaTensorOp instruction.

CUDA exposes warp-level matrix operations in the CUDA C++ WMMA API. The CUDA C++ WMMA API exposes Tensor Cores via a set of functions and types in the nvcuda::wmma namespace. The functions and types in nvcuda::wmma provide target-independent APIs and implement architecture-specific tensor operation using TensorOp instruction underneath. CUTLASS exposes WMMA API through WmmaTensorOp. The WmmaTensorOp supports canonical shared memory layouts. The following table summarizes the destination shared memory layout that can be targeted by matrix operands. The WMMA API expects that matrices in shared memory loaded by nvcuda::wmma::load_matrix_sync() satisfy 128 bit alignment.

WmmaTensorOp (all matrix sizes and data types).

Copyright Copy Section

Copyright (c) 2017 - 2026 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.

Media Docs Cpp Fundamental Types (`media/docs/cpp/fundamental_types.md`) Copy Section

Fundamental Types

CUTLASS defies several fundamental numeric and container classes upon which computations and
algorithms algorithms for linear algebra computations are implemented.

Where possible, CUTLASS fundamental types mirror the C++ Standard Library. However, there are circumstances that necessitate divergence from the Standard Library's specification. In such cases, the CUTLASS implementation adopts unique capitalization to distinguish that standard vocabulary types may not be safely substituted in all cases.

Most types in CUTLASS are usable in both host code and device code. Moreover, they are functional regardless of compute capability, but they may only be efficient when hardware support is present.

Numeric Types Copy Chapter

CUTLASS defines classes for the following numeric data types.

• half_t: IEEE half-precision floating point (exponent: 5b, mantissa: 10b; literal suffix _hf)
• bfloat16_t: BFloat16 data type (exponent: 8b, mantissa: 7b; literal suffix _bf16)
• tfloat32_t: Tensor Float 32 data type (exponent: 8b, mantissa: 10b; literal suffix _tf32)
• int4_t, uint4_t: 4b signed and unsigned integer (literal suffx _s4, _u4)
• bin1_t: 1b binary numeric type (literal suffix _b1)
• float_e5m2_t: 8bits signed float (exponent: 5 bits, mantissa: 2 bits)
• float_e4m3_t: 8bits signed float (exponent: 4 bits, mantissa: 3 bits)
• float_ue4m3_t: 8bits unsigned float (exponent: 4 bits, mantissa: 3 bits)
• float_ue8m0_t: 8bits unsigned float (exponent: 8 bits, mantissa: 0 bits)
• float_e3m2_t: 6bits signed float (exponent: 3 bits, mantissa: 2 bits)
• float_e2m3_t: 6bits signed float (exponent: 2 bits, mantissa: 3 bits)
• float_e2m1_t: 4bits signed float (exponent: 2 bits, mantissa: 1 bits)
• type_erased_dynamic_float8_t: Type agnostic 8 bits signed float allowing the user to provide a specific datatype as runtime argument.
• type_erased_dynamic_float6_t: Type agnostic 6 bits signed float allowing the user to provide a specific datatype as runtime argument.
• type_erased_dynamic_float4_t: Type agnostic 4 bits signed float allowing the user to provide a specific datatype as runtime argument.
• mx_float8_t<float_e5m2_t> or mx_float8_t<float_e4m3_t> : Block scaled data type with fp8 element type and float_ue8m0_t scale factor and vector size of 32.
• mx_float6_t<float_e3m2_t> or mx_float6_t<float_e2m3_t> : Block scaled data type with fp6 element type and float_ue8m0_t scale factor and vector size of 32.
• mx_float4_t<float_e2m1_t> : Block scaled data type with signed e2m1 element type and float_ue8m0_t scale factor and vector size of 32.
• nv_float4_t<float_e2m1_t> : Block scaled data type with signed e2m1 element type and float_ue4m3_t scale factor and vector size of 16.
• complex<T>: defines complex-valued data type based on the supplied real-valued numeric type

Numeric types in CUTLASS may be used in both host and device code and are intended to function
like any other plain-old-data type.

If CUTLASS is compiled with CUTLASS_F16C_ENABLED, then hardware conversion is used for
half-precision types in host code. Regardless, cutlass::half_t uses the most efficient
NVIDIA GPU hardware instructions available in device code.

Example:

#include <iostream>
#include <cutlass/numeric_types.h>

__global__ void kernel(cutlass::half_t x) {
  printf("Device: %f\n", float(x * 2.0_hf));
}

int main() {

  cutlass::half_t x = 0.5_hf;

  std::cin >> x;

  std::cout << "Host: " << 2.0_hf * x << std::endl;

  kernel<<< dim3(1,1), dim3(1,1,1) >>>(x);

  return 0;
}

Containers Copy Chapter

CUTLASS uses the following containers extensively for implementing efficient CUDA kernels.

Array Copy Section