CUTLASS 3.5.1

CUTLASS 3.5.1 - 2024年7月

CUTLASS是一组CUDA C++模板抽象，用于在CUDA的所有级别和规模内实现高性能矩阵-矩阵乘法（GEMM）和相关计算。它采用了类似于实现cuBLAS和cuDNN的层次分解和数据移动策略。CUTLASS将这些“移动部分”分解成可重用的、模块化的软件组件，通过C++模板类进行抽象。可以通过自定义瓦片大小、数据类型和其他算法策略对概念并行化层次结构的不同级别进行特化和调整。这种灵活性简化了它们作为自定义内核和应用程序构建块的使用。

为了支持广泛的各类应用，CUTLASS提供了对混合精度计算的广泛支持，为半精度浮点（FP16）、BFloat16（BF16）、张量浮点32（TF32）、单精度浮点（FP32）、《a href="./examples/27_ampere_3xtf32_fast_accurate_tensorop_gemm" rel="nofollow">FP32通过张量核心指令仿真）、双精度浮点（FP64）类型、整数数据类型（4b和8b）和二进制数据类型（1b）提供了专门的数据移动和乘累加抽象。CUTLASS演示了针对NVIDIA Volta、Turing、Ampere和Hopper架构实现的可编程、高吞吐量张量核心的矩阵乘法操作。

查看快速入门指南以快速开始。

查看功能列表，以获取执行模型层次结构每层支持的操作列表。

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

CuTe 的核心抽象是分层多维布局，可以与数据数组组合来表示张量。布局的表示足以表示我们需要实现高效稠密线性代数的大部分内容。布局也可以通过函数组合进行组合和处理，我们在此之上构建了大量常见操作，如分块和分区。

CUTLASS 3.0 及以后的版本在其模板的 GEMM 层次中采用了 CuTe。这极大地简化了设计，并提高了代码的可组合性和可读性。更多关于 CuTe 的特定文档可以在其 专用文档目录 中找到。

除了 GEMM 之外，CUTLASS 通过隐式 GEMM 算法实现了高性能卷积。隐式 GEMM 将卷积操作表示为 GEMM，从而利用了 CUTLASS 的模块化 GEMM 管道。这允许 CUTLASS 通过重用高度优化的 GEMM 组件来构建卷积。

CUTLASS 3.5 的新特性

CUTLASS 3.5.1 是 CUTLASS 的一个更新，增加了

• 100 行代码中的最小 SM90 WGMMA + TMA GEMM 示例.
• 在 TMA 复制原子中暴露 L2 cache_hint
• 在 CUTLASS 库分析器 和 示例 48 中暴露光栅顺序和分块打乱范围。
• 基于 TMA 存储且支持 EVT 的后缀 用于 Hopper 指针数组批量内核。
• 为 CUTLASS 2.x Ampere 内核提供新的 GemmSparseUniversal API 以启用串行和并行 split-k，以更好地支持 LLM 推理和新的小型分块大小。
• CUDA 主机适配器 扩展以支持 TMA 描述符构建驱动程序 API。
• 将更多 Hopper fprop、dgrad 和 wgrad 卷积内核 包含到 CUTLASS 库和分析器中。
• 在卷积内核中支持残差加（beta != 0）。
• 为 CUTLASS 2.x 提供新的卷积 后缀 以支持非打包 NHWC 输出。
• 重构 CUTLASS 核心目录中的 包含文件 以减少循环依赖，并添加 测试 以防止它们。
• 设置 VSCode 以与 CUTLASS 一起良好工作的指南 和 扩展代码风格指南。
• 更好地支持 MSVC 作为主机编译器。
• 许多性能优化、改进和错误修复，包括 FlashAttention-2 的修复。
• 使用 CUDA 工具包版本 12.4 和 12.5u1 生成最优代码。
• 通知
  • 即将推出的 CUTLASS 3.6 版本将包含对 CUTLASS 3.x 卷积 kernel::ConvUniversal API 的破坏性重构，以使其与 gemm::GemmUniversal 保持一致。在此之后，3.x 卷积 API 将不再被视为测试版 API。
  • 即将推出的 CUTLASS 3.6 版本将包含对 Hopper TMA 指针数组批量后缀的破坏性重构，以支持分组 GEMM。

CUTLASS 3.5.0 是 CUTLASS 的一个更新，增加了

• 针对 Hopper SM90A 的隐式 GEMM 卷积，通过 WGMMA + TMA im2col。
  • CUTLASS 3.x 中使用 CuTe 的本地实现，与 GEMM 的设计层次结构相同。
  • 以秩无关的方式支持 1D、2D 和 3D 卷积。
  • 支持 Fprop、Dgrad 和 Wgrad 算法。
  • 支持通过 3.x API 实现的 2D 和 3D 卷积的 CUTLASS 分析器。
  • 注意：这是一个测试版。CUTLASS 的进一步更新将包括重大性能改进、功能启用和可能的 API 破坏性更改，直到 3.7 版本发布。请对设计提出您的反馈！
• 支持通过 2.x API 的 Ada (SM89) FP8 矩阵核心。需要 CUDA 12.4 或更高版本。
• CuTe 和 CUTLASS 3.x 中的 Ampere 收集/散射卷积示例。
  • 展示如何使用 CUTLASS 3.x 和 CuTe 编写和优化自定义内核，以及将卷积作为 GETT 特殊化的实现策略。
  • 实现了一个粗粒度稀疏收集/散射内核，在 Ampere 类矩阵核心上达到峰值性能。
• 为提高窄高和宽短矩阵的性能，CUTLASS 2.x 添加了 32x 和 16x 瓦片大小。
• 更新 CuTe 文档以包括 cute::Tensor<>、MMA 原子，以及全新的 CuTe GEMM 教程系列。
• 扩展 CuTe 以支持 L2 预取 和 TMA 存储+归约。
• 删除对一些 CUTLASS 2.x API 头文件中 C++11 的要求。所有 CUTLASS 文件现在都需要 C++17。
• 修复以大大减少构建警告。
• 来自社区的更新和错误修复（谢谢！）
• CUTLASS 3.5.1 是 CUTLASS 的小更新，包括一些错误修复和改进，包括 FlashAttention-2 构建修复。

最低要求

• 架构：Volta
• 编译器：必须支持至少 C++17
• CUDA 工具包版本：11.4

从 CUTLASS 3.0 开始，CUTLASS 删除了对以下内容的支持

• Maxwell 和 Pascal GPU 架构
• Ubuntu 16.04
• CUDA 10.2
• C++ 语言版本低于 17。

有关版本和更新的详细信息，请参阅 CHANGELOG。

性能

CUTLASS 原语非常高效。当用于构建全局 GEMM 内核时，它们的表现与 cuBLAS 的标量 GEMM 计算相当。上图显示了自 CUTLASS 3.1 以来在 NVIDIA H100（NVIDIA Hopper 架构）上的持续 CUTLASS 性能改进。CUTLASS 3.5.1 使用 CUDA 12.5u1 工具包 编译。使用 CUDA 的 mma 和 wgmma 指令实现矩阵核心操作。

当使用 CUTLASS 块构建全局隐式 GEMM（Fprop、Dgrad 和 Wgrad）内核时，CUTLASS 性能也类似于在 NVIDIA A100 上运行 Resnet-50 层时的 cuDNN，如上图所示。使用 CUDA 的 mma 指令 实现矩阵核心操作。

兼容性

CUTLASS 需要一个 C++17 主机编译器，并且当使用 CUDA 12.4 工具包 构建时表现最佳。它也与 CUDA 11.4、CUDA 11.5、CUDA 11.6、CUDA 11.7、CUDA 11.8、CUDA 12.0、CUDA 12.1、CUDA 12.2.2、CUDA 12.3.1 和 CUDA 12.3.2 兼容。

操作系统

我们已经测试了以下环境。

操作系统	编译器
Ubuntu 18.04	GCC 7.5.0
Ubuntu 20.04	GCC 10.3.0
Ubuntu 22.04	GCC 11.2.0
Ubuntu 22.04	Clang 10.0.0
Ubuntu 22.04	Clang 14.0.6
Ubuntu 22.04	Clang 17.0.6
Windows 10.0	Visual Studio 2019 v16.11.27

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

硬件

CUTLASS 在以下 NVIDIA GPU 上运行成功，预计在 Volta、Turing、Ampere、Ada 和 Hopper 架构的 NVIDIA GPU 上效率较高。

GPU	CUDA 计算能力	CUTLASS-3 需要的最小 CUDA 工具包版本
NVIDIA V100 张量核心 GPU	7.0	11.4
NVIDIA TitanV	7.0	11.4
NVIDIA GeForce RTX 2080 TI, 2080, 2070	7.5	11.4
NVIDIA T4	7.5	11.4
NVIDIA A100 张量核心 GPU	8.0	11.4
NVIDIA A10	8.6	11.4
NVIDIA GeForce RTX 3090	8.6	11.4
NVIDIA GeForce RTX 4090	8.9	11.8
NVIDIA L40	8.9	11.8
NVIDIA H100 张量核心 GPU	9.0	11.8

目标架构

一般来说，为某个目标架构生成的 PTX 代码可以在未来的架构上运行（即，它是向前兼容的）。然而，CUDA 12.0 引入了“架构加速功能”的概念，其 PTX 没有向前兼容的保证。一些 Hopper PTX 指令属于此类架构加速功能，因此需要 sm_90a 目标架构（注意附加的“a”）。有关此和其他架构加速指令的更多详细信息，请参阅 CUDA 文档。

目标架构信息通过 cmake 标志 CUTLASS_NVCC_ARCHS 传递给 CUTLASS。为了在 Hopper GH100 上最大化性能，用户需要将目标架构设置为 90a 来构建 CUTLASS。如果用户意外构建了一个使用 SM90a 功能的内核（例如，Hopper 张量核心指令），则使用带有 CTK 12 或 11.8 的 SM90 目标（注意没有“a”），内核可能会因运行时错误而失败。

cmake .. -DCUTLASS_NVCC_ARCHS="90a" 

请参阅 功能文档 了解哪些内核需要哪些目标架构的详细信息。

文档

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 中的 2D 和 3D 卷积
• 代码组织 - 描述 CUTLASS 项目的组织结构和内容
• 术语 - 描述代码中使用的术语
• 编程指南 - 写作高效现代 CUDA C++ 的指南
• 基本类型 - 描述 CUTLASS 中用于表示数值和数组的 C++ 基本类
• 布局 - 描述内存中矩阵和张量的布局
• Tile 迭代器 - 描述用于迭代内存中矩阵块的 C++ 概念
• CUTLASS Profiler - 基于命令行的性能分析应用程序
• CUTLASS 工具 - 用于加速快速开发的额外模板

资源

我们在2018年GPU技术大会上介绍了高效GEMM的结构，请参阅此处。

• CUTLASS：CUDA中所有级别和规模的密集线性代数软件原语
• 在NVIDIA A100上开发CUDA内核以将张量核心性能发挥到极致
• 在CUTLASS中使用张量核心加速卷积
• 通过增加张量核心利用率在CUTLASS中加速反向数据梯度
• CUTLASS：Python API，增强功能，以及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/目录开始，通过构建目标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/                 #                            manging 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和卷积内核（长构建时间）

默认情况下，对于每种数据类型、数学指令和布局，仅实例化一个瓷砖大小。要实例化所有，请运行CMake时设置以下环境变量。请注意，这将导致数万个内核和漫长的构建时间。这也会导致二进制文件大小过大，并且在某些平台上可能会在构建库时链接失败。因此，强烈建议仅生成如以下子部分中所示的部分内核。

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

构建GEMM和卷积内核的子集（减少构建时间）

为了严格编译一个内核或一组小的内核，可以使用带有通配符字符的内核名逗号分隔列表来减少内核集。以下示例展示了为NVIDIA Ampere和Turing架构构建精确的一个或一组内核。

构建Tensor Core GEMM内核的子集

要编译针对NVIDIA Ampere和Turing架构的Tensor Core GEMM内核子集，并使用FP32累加和FP16输入，请使用以下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卷积内核的子集

要编译实现前向传播（fprop）的Tensor core卷积内核子集，并使用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内核

要编译和运行一个CUDA Core卷积内核，实现前向传播（fprop）使用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公司作为开源软件发布，遵循3-clause "New" BSD许可。

贡献者

CUTLASS的开发者和贡献者的官方列表可在此处找到：CONTRIBUTORS。

版权

版权（c）2017 - 2024 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.