概述

CUTLASS 4.7.0

CUTLASS 4.7.0 - 2026年8月

CUTLASS 是一套用于在 CUDA 中实现高性能矩阵乘法（GEMM）及相关计算的抽象集合，涵盖各级别与各规模。它融合了层次化分解和数据移动策略，将这些“动态组件”分解为可重用、模块化的软件组件与抽象。

概念并行化层次结构中不同级别的原语可通过自定义分块大小、数据类型和其他算法策略进行特化和调优。由此产生的灵活性简化了它们在自定义内核和应用程序中作为构建块的使用。

自2017年以来，CUTLASS 一直提供用于高性能线性代数的 CUDA C++ 模板抽象。这些抽象广泛支持各种计算，包括混合精度计算、专用数据移动（异步复制）以及针对 FP64、FP32、TF32、FP16、BF16 的乘积累加抽象， 通过张量核心指令进行的 FP32 仿真， 8位浮点类型（e5m2 和 e4m3）， 块缩放数据类型（NVIDIA NVFP4 以及 OCP 标准 MXFP4、MXFP6、MXFP8）， 窄整数类型（4位和8位有符号及无符号整数）， 以及二进制1位数据类型（在架构支持此类数据类型的原生操作时），覆盖 NVIDIA 的 Volta、Turing、Ampere、Ada、Hopper 和 Blackwell 架构。

在这个丰富的基于 C++ 的内核编程抽象生态系统基础上，CUTLASS 4 引入了 CUTLASS DSL。这是一种 Python 原生接口，用于基于核心 CUTLASS 和 CuTe 概念编写高性能 CUDA 内核，且不损失任何性能。这提供了更平缓的学习曲线、快几个数量级的编译时间、无需编写粘合代码即可与深度学习框架原生集成，以及更直观的元编程，无需深厚的 C++ 专业知识。

总体而言，我们将 CUTLASS DSL 视为一系列领域特定语言（DSL）。随着 4.0 版本的发布，我们推出了其中首个 DSL——CuTe DSL。这是一种低级编程模型，与 CuTe C++ 抽象完全一致，公开了诸如布局、张量、硬件原子等核心概念，并允许对硬件线程和数据层次结构进行完全控制。

CuTe DSL 展示了针对 NVIDIA Ampere、Hopper 和 Blackwell 架构所实现的可编程、高吞吐量 Tensor Cores 的优化矩阵乘法及其他线性代数运算。

我们相信，它将成为学生、研究人员和性能工程师不可或缺的工具——降低 GPU 编程的学习门槛，快速原型化内核设计，并将优化的解决方案投入生产。

CuTe DSL 目前处于公开测试阶段，将于 2026 年夏末结束测试阶段。

如需快速入门，请参考：

• CUTLASS C++ 快速入门指南。
  • CuTe DSL 快速入门指南。

CUTLASS 4.7 新特性

CuTe DSL

• 新特性：

  • 引入了 Primitives API，它在 CuTe 之下提供了更低级别的抽象，支持通过 SIMT 进行 Tensor Core 编程。这为 NVVM 操作提供了一个稳定、轻量的封装，可用于 CuTe 抽象降低开发效率的场景。Primitives 目前为实验性版本，将根据用户反馈进行演进。

    注意：Primitives 是一个过渡性 API，直至 CUDA Python 类解决方案可用。

  • 引入了任务调度框架。该框架可对 warp 专用内核的执行调度进行静态分析。当检测到已知的并发问题时，编译会停止。同时提供了用于可视化资源/任务依赖关系和分析内核调度的工具。

  • 改进了编译器诊断功能。现在可以在编译时报告寄存器溢出和本地内存使用情况，并附带源代码行号。在使用 Primitives API 时，初步支持在编译时检测各类 NVVM 同步和执行风险。对以前不包含源代码行号的编译器错误提供了更好的报告。

此版本已针对以下软件包进行测试：

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

CUTLASS 算子 API

• 自定义 epilogue 融合增强：
  • 增加对标量约简的支持。
  • 增加对为每个加载/存储的操作数指定数据移动策略的支持。

C++

• 在 FMHA 示例 中添加面向 FP8 的双内核反向传播实现。
  • 添加了一个支持多精度（FP16/FP8）的反向融合多头注意力基准测试，可配置批次/序列/头大小、变长和掩码选项，并内置正确性检查以及运行时/吞吐量报告。
  • 在 Blackwell SM103 芯片上，无掩码的 FP8 双内核反向传播实现相比单内核实现大约有 25% 的性能提升。
• 支持 CUDA 12.6 及更新版本的 NVRTC 结构化绑定头文件。
• 在 NumericArrayConverter 中添加 fp32/fp16/bf16/e4m3/e5m2 到 e2m1（FP4）的转换。
• 通过在 prmt 之前合并掩码，优化 E2M1 到 FP16 的 LUT 解码辅助函数 _e2m1_to_half_x2 和 _e2m1_to_half_x4。
• 修复部分问题：
  • 更新 streamk 启发式算法，以优化某些具有混合簇大小的内核。
  • 避免 CuTe 元组算法中的 integer-sequence get 歧义。
  • 修复 TMA 创建驱动程序 bug：通过检查张量是否紧凑来保守地检测前 128KiB 是否被映射，如果是，则可以翻转该位，否则将其置零。
• 来自社区和 CUTLASS 团队的各种改进和修复。感谢所有提交 PR 的人！
• 使用 CUDA 工具包 13.3 版本生成优化代码。

注意：已知在所有 CUDA 工具包版本下，Windows 平台上的 CUTLASS 4.x 构建均存在问题。CUTLASS 团队正在努力修复。

有关所有过往版本和更新的详细信息，请参阅 变更日志。

性能

CUTLASS 原语效率极高。当用于构建设备级 GEMM 内核时，它们几乎能达到理论峰值吞吐量的最佳利用率。下图展示了 CUTLASS 3.8 在 NVIDIA Blackwell SM100 架构 GPU 上运行时，针对各种输入和输出数据类型的性能占理论峰值利用率的百分比。

以下两张图展示了自 CUTLASS 3.1 以来，在 NVIDIA H100（NVIDIA Hopper 架构）上 CUTLASS 性能的持续提升。CUTLASS 3.5.1 是使用 CUDA 12.5u1 工具包 编译的。张量核心操作使用 CUDA 的 mma 和 wgmma 指令实现。

CuTe

CUTLASS 3.0 引入了一个新的核心库 CuTe，用于描述和操作线程张量与数据张量。 CuTe 是一组 C++ CUDA 模板抽象，用于定义和操作线程与数据的分层多维布局。 CuTe 提供 Layout 和 Tensor 对象，可紧凑地封装数据的类型、形状、内存空间和布局，同时为用户执行复杂的索引计算。 这使程序员能够专注于算法的逻辑描述，而 CuTe 负责处理繁琐的记录工作。借助这些工具，我们可以快速设计、实现和修改所有稠密线性代数运算。

CuTe 的核心抽象是分层多维布局，这些布局可与数据数组组合以表示张量。 布局的表示能力足以涵盖我们实现高效稠密线性代数所需的几乎所有内容。 布局还可以通过函数组合进行合并和操作，我们基于此构建了大量常见操作，如分块和分区。

CUTLASS 3.0 及更高版本在其模板的 GEMM 层次结构中全面采用了 CuTe。 这极大地简化了设计，并提高了代码的可组合性和可读性。 有关 CuTe 的更多特定文档，请参见其专用文档目录。

兼容性

最低要求：

• 架构：Volta（计算能力 7.0）
• 编译器：必须至少支持 C++17
• CUDA 工具包版本：11.4

CUTLASS 需要 C++17 主机编译器，并且使用CUDA 12.8 工具包构建时性能最佳。 它还与 CUDA 11.4、CUDA 11.5、CUDA 11.6、CUDA 11.7、CUDA 11.8 以及所有其他 CUDA 12.x 版本兼容。

操作系统

我们已在以下环境中进行了测试。

操作系统	编译器
Ubuntu 18.04	GCC 7.5.0
Ubuntu 20.04	GCC 10.3.0
Ubuntu 22.04	GCC 11.2.0

注意：GCC 8.5.0 在折叠表达式和重载运算符方面存在已知的回归问题。建议使用 GCC 7.5.0 或（首选）GCC >= 9。

注意：已知在所有 CUDA 工具包下，Windows 平台上的 CUTLASS 3.x 构建均无法正常工作。 CUTLASS 团队正在努力修复此问题。

硬件

CUTLASS 已在以下 NVIDIA GPU 上成功运行，并且有望在基于 Volta、Turing、Ampere、Ada 和 Hopper 架构的 NVIDIA GPU 上高效运行。

GPU	CUDA 计算能力	CUTLASS-3 所需的最低 CUDA 工具包版本
NVIDIA V100 Tensor Core GPU	7.0	11.4
NVIDIA TitanV	7.0	11.4
NVIDIA GeForce RTX 20x0 系列	7.5	11.4
NVIDIA T4	7.5	11.4
NVIDIA A100 Tensor Core GPU	8.0	11.4
NVIDIA A10	8.6	11.4
NVIDIA GeForce RTX 30x0 系列	8.6	11.4
NVIDIA GeForce RTX 40x0 系列	8.9	11.8
NVIDIA L40	8.9	11.8
NVIDIA H100 Tensor Core GPU	9.0	11.8
NVIDIA H200 Tensor Core GPU	9.0	11.8
NVIDIA B200 Tensor Core GPU	10.0	12.8
NVIDIA B300 Tensor Core GPU	10.3	13.0
NVIDIA DRIVE Thor	11.0	13.0
NVIDIA GeForce RTX 50x0 系列	12.0	12.8
NVIDIA DGX Spark	12.1	13.0

目标架构

通常，为一种目标架构生成的 PTX 代码可以在未来的架构上运行（即具有前向兼容性）。 然而，CUDA 12.0 引入了“架构加速特性”的概念，其 PTX 不保证前向兼容性。 一些 Hopper 和 Blackwell PTX 指令属于此类架构加速特性，因此需要 sm_90a 或 sm100a 目标架构（注意末尾附加的“a”）。有关此特性和其他架构加速指令的更多详细信息，请参阅 CUDA 文档。

目标架构信息通过 cmake 标志 CUTLASS_NVCC_ARCHS 传递给 CUTLASS。为了在 Hopper GH100 上实现最佳性能，用户需要将 CUTLASS 构建为以 90a 为目标架构。 如果用户不慎构建了使用 SM90a 特性（例如 Hopper 张量核心指令）的内核，而使用的是 SM90 目标（注意缺少“a”），无论是使用 CUDA Toolkit 12 还是 11.8，该内核都可能在运行时出错。

cmake .. -DCUTLASS_NVCC_ARCHS="90a"

或者

cmake .. -DCUTLASS_NVCC_ARCHS="100a"

注意：数据中心产品中使用的 NVIDIA Blackwell SM100 架构与支持 NVIDIA Blackwell GeForce RTX 50 系列 GPU 的架构（SM120）具有不同的计算能力。因此，针对 Blackwell SM100 架构使用架构条件特性（通过 sm100a）编译的内核与 RTX 50 系列 GPU 不兼容。

有关哪些内核需要哪些目标架构的详细信息，请参考功能文档。

文档

CUTLASS 在以下文档和随附的 Doxygen 文档中进行了描述。

• 快速入门指南 - 构建和运行 CUTLASS 的基础知识
• 功能特性 - 总结 CUTLASS 中可用的功能
• CUDA 中的高效 GEMM - 描述如何在 CUDA 中高效实现 GEMM 内核
• CUTLASS 3.x 设计 - 描述 CUTLASS 3.x 的设计、其优势以及 CuTe 如何使我们能够编写更具组合性的组件
• GEMM API 3.x - 描述 CUTLASS 3.x GEMM 模型和 C++ 模板概念
• GEMM API 2.x - 描述 CUTLASS 2.x GEMM 模型和 C++ 模板概念
• 隐式 GEMM 卷积 - 描述 CUTLASS 中的二维和三维卷积
• 代码组织 - 描述 CUTLASS 项目的组织和内容
• 术语表 - 描述代码中使用的术语
• 编程指南 - 编写高效现代 CUDA C++ 的指南
• 基本类型 - 描述 CUTLASS 中用于表示数值量和数组的基本 C++ 类
• 布局 - 描述内存中矩阵和张量的布局
• 瓦片迭代器 - 描述用于迭代内存中矩阵瓦片的 C++ 概念
• CUTLASS 性能分析器 - 命令行驱动的性能分析应用程序
• CUTLASS 实用工具 - 用于促进快速开发的其他模板
• 依赖内核启动 - 描述 Hopper 中的一项新功能，该功能允许在同一流中重叠依赖内核，以及它在 CUTLASS 中的使用方式。

资源

我们在 2018 年 GPU 技术大会的演讲中也介绍了高效 GEMM 的结构，相关内容可查看 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

构建 CUTLASS

CUTLASS 是一个仅含头文件的模板库，其他项目使用时无需对其进行构建。客户端应用程序只需在其包含路径中指定 CUTLASS 的 include/ 目录即可。

CUTLASS 的单元测试、示例和工具可通过 CMake 进行构建。CMake 的最低版本要求在 快速入门指南 中有说明。请确保 CUDACXX 环境变量指向系统中已安装的 CUDA 工具包中的 NVCC。

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

在 CUTLASS 项目内创建一个构建目录，然后运行 CMake。默认情况下，CUTLASS 将为 CUDA 架构版本 5.0、6.0、6.1、7.0、7.5、8.0、8.6、8.9 和 9.0 构建内核。若要缩短编译时间，您可以通过修改 CMake 配置设置 CUTLASS_NVCC_ARCHS 来指定要为其构建 CUTLASS 的架构。

$ mkdir build && cd build

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

从 build/ 目录中，通过使用 make 构建目标 test_unit 来编译并运行 CUTLASS 单元测试。

单元测试被组织为多个二进制文件，这些文件与 CUTLASS 的顶级命名空间相对应，并且可以通过 make 的 -j 命令行参数并行执行。

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

所有测试应在受支持的平台上通过，不过测试的确切数量可能会随时间变化。

项目结构

CUTLASS 是一个仅包含头文件的库，同时配有工具、实用程序、示例和单元测试。 Doxygen 文档提供了 CUTLASS 项目中定义的所有文件、类和模板概念的完整列表。

有关源代码组织的详细说明，请参见 CUTLASS 文档，下面概述了几个主要组件。

CUTLASS 模板库

include/                     # client applications should target this directory in their build's include paths

  cutlass/                   # CUDA Templates for Linear Algebra Subroutines and Solvers - headers only

    arch/                    # direct exposure of architecture features (including instruction-level GEMMs)

    conv/                    # code specialized for convolution

    epilogue/                # code specialized for the epilogue of gemm/convolution

    gemm/                    # code specialized for general matrix product computations

    layout/                  # layout definitions for matrices, tensors, and other mathematical objects in memory

    platform/                # CUDA-capable Standard Library components

    reduction/               # bandwidth-limited reduction kernels that do not fit the "gemm" model

    thread/                  # simt code that can be performed within a CUDA thread

    transform/               # code specialized for layout, type, and domain transformations

    *                        # core vocabulary types, containers, and basic numeric operations

  cute/                      # CuTe Layout, layout algebra, MMA/Copy atoms, tiled MMA/Copy

    algorithm/               # Definitions of core operations such as copy, gemm, and operations on cute::tuples

    arch/                    # Bare bones PTX wrapper structs for copy and math instructions

    atom/                    # Meta-information either link to or built from arch/ operators

      mma_atom.hpp           # cute::Mma_Atom and cute::TiledMma

      copy_atom.hpp          # cute::Copy_Atom and cute::TiledCopy

      *sm*.hpp               # Arch specific meta-information for copy and math operations

    *                        # Core library types such as Shape, Stride, Layout, Tensor, and associated operations

CUTLASS SDK 示例

CUTLASS SDK 示例 应用 CUTLASS 模板来实现基本计算。

工具

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/unit/ 目录包含使用 Google Test 实现的单元测试，这些测试展示了 Core API 组件的基本用法以及对 CUTLASS GEMM 计算的完整测试。

有关构建和运行单元测试的说明，请参见 快速入门指南。

性能分析

tools/profiler/ 目录包含一个命令行工具，用于启动各个 GEMM 内核。其构建方法如下：

$ make cutlass_profiler -j16

构建所有 GEMM 和卷积内核（构建时间较长）

默认情况下，每种数据类型、数学指令和布局仅实例化一种分块大小。 要实例化所有分块大小，请在从空的 build/ 目录运行 CMake 时设置以下环境变量。 请注意，这将生成数万个内核，导致构建时间很长。 这还会导致二进制文件体积过大，在某些平台上链接器可能无法成功构建库。 因此，强烈建议仅生成一部分内核，如下文小节所示。

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

构建 GEMM 和卷积核的子集（缩短构建时间）

若要严格编译单个核或一小组核，可使用带通配符的逗号分隔核名称列表来缩减核集合。以下示例展示了为 NVIDIA Ampere 和 Turing 架构构建单个核或核子集的方法：

构建 Tensor Core GEMM 核子集

要为 NVIDIA Ampere 和 Turing 架构编译具有 FP32 累加和 FP16 输入的 Tensor Core GEMM 核子集，请使用以下 cmake 命令行：

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

用于分析部分 Tensor Core GEMM 内核的示例命令行如下：

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



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

构建一个 CUDA Core GEMM 内核

若要编译一个面向 NVIDIA Ampere 和 Turing 架构的 SGEMM 内核，请使用以下 cmake 命令行：

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

用于分析单个 SGEMM CUDA 内核的命令行示例如下：

$ ./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

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

构建部分 Tensor Core 卷积核

若要编译部分 Tensor Core 卷积核（这些卷积核实现前向传播（fprop），采用 FP32 累加和 FP16 输入，目标为 NVIDIA Ampere 和 Turing 架构），请使用以下 cmake 命令行：

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

用于分析部分 Tensor Core 卷积核的命令行示例如下：

$ ./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

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

构建一个卷积 CUDA 内核

要编译和运行一个实现前向传播（fprop）的 CUDA Core 卷积内核，该内核采用 F32 累加和 FP32 输入，针对 NVIDIA Ampere 和 Turing 架构，请使用以下 cmake 命令行：

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

用于分析一个 CUDA Core 卷积核的示例命令行：

$ ./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


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

编译 CUTLASS 内核和 CUTLASS 性能分析器的更多详细信息

• 请通过以下链接获取更多关于选择性编译 CUTLASS 内核的 CMake 示例：
  • GEMM CMake 示例
  • 隐式 GEMM 卷积 CMake 示例
• 此处描述了关于 CUTLASS 性能分析器的更多详细信息。

关于

CUTLASS 由 NVIDIA Corporation 作为开源软件发布，采用 3 条款“新”BSD 许可证。

贡献者

CUTLASS 开发人员和贡献者的官方列表可在此处获取：贡献者。

版权

版权所有 (c) 2017 - 2026 NVIDIA CORPORATION & AFFILIATES。保留所有权利。 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.