DEVELOPING CUDA KERNELS TO PUSH TENSOR CORES TO THE ABSOLUTE LIMIT ON NVIDIA A100

Andrew Kerr, May 21, 2020

致谢 (ACKNOWLEDGEMENTS)

• CUTLASS 团队: Andrew Kerr, Haicheng Wu, Manish Gupta, Duane Merrill, Pradeep Ramani
• 贡献者: Mostafa Hagog, Timothy Costa, Alan Kaatz, John Tran, Stephen Jones, Kyrylo Perelygin, Luke Durant, Piotr Majcher, Paul Springer, Markus Hohnerbach
• 鸣谢: Joel McCormack, Julien Demouth, Olivier Giroux, Bryce Lelbach, Cris Cecka

议程 (AGENDA)

1. 概述
   • NVIDIA Ampere 架构与 CUTLASS 2.2
2. NVIDIA Ampere 架构上的 Tensor Cores
   • 加速矩阵运算
3. 为 Tensor Cores 实现高效数据移动
   • 最大化性能的策略
4. NVIDIA A100 上的 CUTLASS
   • 用于 Tensor Cores 的最优 CUDA C++ 模板

概述 (OVERVIEW)

NVIDIA AMPERE 架构：NVIDIA A100

• 全新且更快的 Tensor Core 操作
  • 浮点 Tensor Core 操作比 F32 CUDA Cores 快 8x 和 16x
  • 整数 Tensor Core 操作比 F32 CUDA Cores 快 32x 和 64x
  • 全新的 IEEE 双精度 Tensor Cores 比 F64 CUDA Cores 快 2x
• 额外的数据类型和模式
  • Bfloat16, double, Tensor Float 32
• 异步复制 (Asynchronous copy)
  • 直接复制到共享内存中 - 实现深度软件流水线
• 更多新特性请参阅 "Inside NVIDIA Ampere Architecture"

编程 NVIDIA AMPERE 架构

可以通过以下两种方式利用 Tensor Cores 进行深度学习和数学库编程：

1. 高级库: 使用 cuDNN, cuBLAS, cuTENSOR 等 CUDA 加速的数学库，这些库在底层调用了 CUDA kernels。
2. CUDA C++ 设备代码: 直接使用 CUTLASS, CUDA Math API, CUB, Thrust, libcu++ 等进行底层编程。

本次演讲主要面向 CUDA 程序员，重点关注如何使用 CUDA C++ 和 CUTLASS 库进行编程。

CUTLASS

CUTLASS 是一个用于深度学习和线性代数的 CUDA C++ 模板库。其发展历程如下：

• CUTLASS 预览版 (CUDA 9.1)
• CUTLASS 1.0 (CUDA 9.2)
• CUTLASS 1.3 (CUDA 10.1) - 原生支持 NVIDIA V100 Tensor Cores
• CUTLASS 2.0 (CUDA 10.2) - 原生支持 NVIDIA Turing Tensor Cores
• CUTLASS 2.2 (CUDA 11) - 针对 NVIDIA A100

CUTLASS 新特性

• CUTLASS 2.2: 在 NVIDIA Ampere 架构上实现最优性能
  • Tensor Cores 吞吐量更高：所有数据类型提速超过2倍
  • 新的浮点类型：bfloat16, Tensor Float 32, double
  • 使用 cp.async 实现深度软件流水线：高效且容忍延迟
• CUTLASS 2.1:
  • 平面复数（complex-valued）GEMM，支持批处理，面向 Volta 和 Turing Tensor Cores
  • BLAS 风格的主机端 API
• CUTLASS 2.0:
  • 使用现代 C++11 编程进行重要重构
  • 高效，尤其适用于 Turing Tensor Cores
  • Tensor Core 编程模型：CUDA 中线性代数核的可复用组件
  • 文档、性能分析工具、参考实现、SDK 示例等

CUTLASS 在 NVIDIA AMPERE 架构上的性能

CUTLASS 2.2 - CUDA 11 Toolkit - NVIDIA A100
下图展示了在 m=3456, n=4096 条件下，不同精度下 Tensor Core 相对于 CUDA Core 的性能提升。
- 混合精度浮点: BF16 相比 F32 提升 13x，TF32 相比 F32 提升 5.7x。
- 双精度浮点: F64 Tensor Core 相比 F64 CUDA Core 提升 2x。
- 混合精度整数: INT4 相比 INT8 (CUDA Core) 提升 13.8x，INT8 (Tensor Core) 提升 7.7x。

NVIDIA AMPERE 架构上的 TENSOR CORES

什么是 TENSOR CORES？

Tensor Cores 执行 D = op(A, B) + C 形式的矩阵运算。

• 矩阵操作:
  • 矩阵乘加 (Matrix multiply-add)
  • XOR-POPC
• 输入数据类型 (A, B):
  • half, bfloat16, Tensor Float 32, double, int8, int4, bin1
• 累加数据类型 (C, D):
  • half, float, int32_t, double

这是一个 M x N x K 的矩阵操作，它是一个 warp 同步的集合操作（warp-synchronous, collective operation）。一个 warp 内的 32 个线程共同持有 A、B、C 和 D 操作数。

NVIDIA AMPERE 架构 - TENSOR CORE 操作

下表总结了 Ampere 架构上 Tensor Core 的各种 mma.sync（矩阵乘加）指令、支持的数据类型、形状以及相比于 F32 CUDA Cores 和前代架构的加速比。例如，对于 F16/BF16 数据类型，其在 A100 上的速度是 F32 CUDA Cores 的 16x。对于 INT4 数据，加速比可达 64x。

TENSOR CORE 操作：基本形状

Tensor Core 的操作是 warp 级别的。下图展示了一个 warp 内的32个线程（T0-T31）如何协同工作。以一个 8x8x128b 的 warp 级 Tensor Core 操作为例，展示了线程如何持有和处理数据片段。

S8 * S8 + S32 矩阵乘加运算 (8-by-8-by-16)

此页展示了针对 Tensor Core 的 8-by-8-by-16 形状的 S8 * S8 + S32 矩阵乘加操作。左侧图示说明了32个线程 (T0-T31) 如何处理输入数据。每个线程处理的数据片段，以及它们如何组合成最终的64位结果 (r0, r1)。

右侧提供了通过内联 PTX (Parallel Thread Execution) 指令调用此操作的示例代码。mma.sync.aligned.m8n8k16.row.col.s8.s8.s32 指令明确了操作的形状 (m8n8k16)、数据布局 (row.col) 以及操作数的数据类型（S8, S8, S32）。

扩展 M 维度

通过组合多个基础的 Tensor Core 操作，可以扩展矩阵运算的维度。此图展示了如何将两个 8-by-8 的操作在 M 维度上堆叠，从而实现一个 warp 级别的 16-by-8-by-8-128b Tensor Core 操作。这种扩展利用了 warp 中所有线程的计算能力来处理更大的矩阵块。

F16 * F16 + F32 矩阵乘加运算 (16-by-8-by-8)

此页介绍了 16-by-8-by-8 形状的 F16 * F16 + F32 操作。该操作使用半精度浮点数 (F16) 作为输入，并累加到单精度浮点数 (F32) 的累加器中。左侧图示说明了数据的分布和处理流程。

右侧的内联 PTX 代码示例展示了如何调用 mma.sync.aligned.m16n8k8.row.col.f16.f16.f32 指令来执行此操作。

扩展 K 维度

与扩展 M 维度类似，也可以在 K 维度上进行扩展以处理更复杂的矩阵乘法。此图演示了如何通过组合操作来扩展 K 维度，形成一个 warp 级别的 16-by-8-by-8-256b Tensor Core 操作。这对于深度学习中常见的卷积等运算至关重要。

F16 * F16 + F32 矩阵乘加运算 (16-by-8-by-16)

此页展示了另一种 F16 * F16 + F32 操作，其形状为 16-by-8-by-16。这表示 K 维度增加了一倍，允许每个操作处理更多的数据。

右侧的 PTX 代码示例 mma.sync.aligned.m16n8k16.row.col.f32.f16.f32 旨在执行此操作。

S8 * S8 + S32 矩阵乘加运算 (16-by-8-by-32)

本页介绍了 16-by-8-by-32 形状的 S8 * S8 + S32 整数矩阵乘加操作。这种操作在推理任务中非常常见，其中输入和权重通常被量化为8位整数。

右侧的 PTX 代码示例 mma.sync.aligned.m16n8k32.row.col.s32.s8.s8.s32 展示了如何调用此高吞吐量的整数运算指令。

半精度 (Half-Precision) 运算: F16 * F16 + F16 (16-by-8-by-16)

此页介绍了一种纯半精度的 Tensor Core 操作，其中输入、输出和累加器均为 F16 类型。与使用 F32 累加器相比，这种 F16 * F16 + F16 的操作可以减少寄存器的使用量。例如，累加器 D 和 C 只需要两个寄存器，而 F32 累加器则需要四个。这可以提高寄存器文件的利用效率。

对应的 PTX 指令为 mma.sync.aligned.m16n8k16.row.col.f16.f16.f16。

双精度 (Double-Precision) 运算: F64 * F64 + F64 (8-by-8-by-4)

NVIDIA Tensor Cores 也支持双精度浮点运算，这对于科学计算和 HPC 应用至关重要。此页展示了 8-by-8-by-4 形状的 F64 * F64 + F64 操作。

该操作使用 uint64_t 类型的寄存器来处理操作数，并产生128位的累加结果。相应的 PTX 指令是 mma.sync.aligned.m8n8k4.row.col.f64.f64.f64。

CUTLASS: 将 PTX 封装在模板中

为了简化 Tensor Core 编程，NVIDIA 提供了 CUTLASS 库。CUTLASS 是一个 C++ 模板库，它将底层的、复杂的 PTX 指令封装在易于使用的模板化组件中。

此页展示了 cutlass::arch::Mma 模板的结构。开发者可以通过指定模板参数（如矩阵形状、线程数、数据类型、布局等）来定义所需的矩阵乘加操作，而无需直接编写内联汇编。

CUTLASS 示例: 16-by-8-by-16

此页提供了一个使用 CUTLASS 的具体代码示例，用于执行 16-by-8-by-16 的矩阵乘加操作。

通过 arch::Mma<GemmShape<16, 8, 16>, 32, ...> 定义一个矩阵操作 mma，然后在 CUDA kernel 中直接调用 mma(C, A, B, C) 即可执行原位矩阵乘加。这种高级抽象极大地简化了代码，提高了可读性和可维护性。

为 Tensor Cores 提供高效的数据移动

Tensor Cores 的 "Hello World" 示例

使用 Tensor Cores 的基本编程流程包括以下步骤：
1. 将每个线程映射到矩阵操作的坐标。
2. 从内存加载输入数据。
3. 执行矩阵操作。
4. 将结果存储回内存。

右侧的 CUDA 示例代码展示了这一流程。代码首先计算每个线程访问矩阵 A 和 B 的坐标以及累加器矩阵的坐标，然后计算线性内存偏移量，最后通过 asm 块发出 mma.sync.aligned PTX 指令来执行 Tensor Core 操作。

性能影响分析

直接从全局内存加载数据给 Tensor Core 会带来性能瓶颈。
- 计算分析：
- 从内存加载 A 和 B 输入：每个线程 2 x 4B = 8B。一个 warp (32线程) 加载 256B。
- 执行一次 Tensor Core 操作：对于 m8n8k16 S8 操作，一个 warp 执行 2 * 8 * 8 * 16 = 2048 次浮点操作 (flops)。
- 算术强度：2048 flops / 256 B = 8 flops/byte。
- 硬件规格 (NVIDIA A100)：
- 峰值算力：624 TFLOP/s (INT8)
- 内存带宽：1.6 TB/s (HBM2)
- 机器平衡点：624 / 1.6 ≈ 400 flops/byte。
- 结论：由于该内核的算术强度 (8 flops/byte) 远低于机器的平衡点 (400 flops/byte)，因此它的性能受限于全局内存带宽，计算单元（Tensor Cores）大部分时间在等待数据。

为数据路径提供数据：通过共享内存高效存取

为了解决内存带宽瓶颈，需要采用分层、分块的内存模型，以最大化数据重用。

该图展示了一个高效的通用矩阵乘法 (GEMM) 的数据流：
1. Blocked GEMM: 将大的矩阵从全局内存 (Global Memory) 分块加载到共享内存 (Shared Memory)。
2. Thread Block Tile: 线程块内的线程协作将数据从共享内存加载到寄存器文件 (Register File)。
3. Warp Tile: 一个 warp 内的线程将寄存器中的数据提供给 CUDA/Tensor Cores 进行计算。
4. 计算结果写回寄存器，通过 Epilogue Functor 进行后续处理（如激活函数），然后写回全局内存。

核心思想是在共享内存和寄存器中重用数据，减少对高延迟、低带宽的全局内存的访问。

为数据路径提供数据：优化从全局内存到 Tensor Cores 的数据移动

为了尽可能高效地将数据从全局内存移动到 Tensor Cores，需要关注以下几点：
- 延迟容忍的流水线: 从全局内存加载数据时，使用流水线技术来隐藏内存访问延迟。
- 无冲突的共享内存存储: 设计数据布局和访问模式，避免写入共享内存时发生 bank conflicts。
- 无冲突的共享内存加载: 同样，在从共享内存读取数据到寄存器时，也要避免 bank conflicts。

图中详细展示了从全局内存到 Tensor Cores 的数据路径，强调了在每个阶段（Blocked GEMM -> Thread Block Tile -> Warp Tile -> Tensor Op）的数据组织和流动。

异步复制：高效流水线

NVIDIA Ampere 架构的一项新特性是 cp.async 指令，它支持从全局内存 (Global Memory) 直接异步复制到共享内存 (Shared Memory)。这一特性使得构建高效的软件流水线成为可能。

• 最小化数据移动：数据路径从 L2 → L1 → RF → SMEM 简化为 L2 → SMEM。
• 节省寄存器：不再需要寄存器文件 (RF) 来保存长延迟加载指令的结果。
• 间接性：可以提前获取多个阶段的数据，以更好地容忍延迟。

通过在共享内存中使用循环缓冲区（Circular buffer），可以实现数据加载和计算的流水线操作，如下图所示，其中 cp.async 指令负责将数据块写入缓冲区，而计算核心则通过 ld.shared 读取已就绪的数据。

馈送数据路径

为了最大化 Tensor Core 的利用率，必须尽可能高效地将数据从全局内存移动到 Tensor Core。这主要通过以下三个关键策略实现：

1. 延迟容忍的流水线：从全局内存开始构建流水线，以隐藏数据访问延迟。
2. 无冲突的共享内存存储：确保将数据从全局内存写入共享内存时不会发生存储体冲突 (bank conflict)。
3. 无冲突的共享内存加载：确保从共享内存加载数据到寄存器时不会发生存储体冲突。

整个数据流可以概括为：数据首先以分块 GEMM (Blocked GEMM) 的形式存在于全局内存中，然后被加载到线程块对应的共享内存瓦片 (Thread Block Tile)，接着加载到 Warp 对应的寄存器文件瓦片 (Warp Tile)，最终被送入 Tensor Core 进行运算。

从全局内存到 Tensor Cores

下图展示了数据从全局内存到共享内存，再到 Tensor Core 的映射过程。cp.async 指令负责将数据从全局内存复制到共享内存。随后，这些数据被重新组织并加载，以匹配 Tensor Core 操作所需的输入格式。

LDMATRIX: 获取 Tensor Core 操作数

ldmatrix 是一条 PTX 指令，专门用于从共享内存中加载一个矩阵到寄存器中，以供 Tensor Core 使用。

• 每个线程提供一个指向共享内存中 128 位数据行的指针。
• 每条 128 位的数据行会被广播给一组四个线程（这组线程可能与提供指针的线程不同）。
• 加载后的数据排列方式与 Tensor Core 操作的输入要求完全匹配。

下图清晰地展示了线程（如 T0, T8, T16, T24）如何提供指针，以及 ldmatrix 指令如何根据这些指针从共享内存中抓取数据并将其组织成 Tensor Core 所需的矩阵格式。

ldmatrix PTX 指令详解

ldmatrix 指令通过内联 PTX 汇编在 CUDA C++ 中使用。其基本工作原理是，每个线程提供一个指针，指令会加载 128 位（即 4 个 32 位值）的数据，并广播给一个四线程组。

以下是 ldmatrix 指令的内联 PTX 汇编示例：

// Inline PTX assembly for ldmatrix
uint32_t R[4];
uint32_t smem_ptr;

asm volatile (
  "ldmatrix.sync.aligned.x4.m8n8.shared.b16 "
  "{%0, %1, %2, %3}, [%4];"
  : "=r"(R[0]), "=r"(R[1]), "=r"(R[2]), "=r"(R[3])
  : "r"(smem_ptr)
);

下图整合了 cp.async 和 ldmatrix 两个步骤，完整地展示了数据从全局内存到共享内存，再到 Tensor Cores 的全过程。

NVIDIA Ampere 架构 - 共享内存存储体时序

在 Ampere 架构中，共享内存的访问是以多个阶段（phase）进行的，这对于理解和避免存储体冲突至关重要。

• 一个 Warp 中的线程分为多个访问阶段。
• 访问 128 位（16B）的数据需要 4 个阶段。
• 每个阶段处理一个 Warp 中 8 个连续的线程。
  • Phase 0: T0 .. T7
  • Phase 1: T8 .. T15
  • Phase 2: T16 .. T23
  • Phase 3: T24 .. T31

共享内存存储体冲突问题

如果将数据从全局内存直接按线程顺序连续存放到共享内存中，会导致严重的存储体冲突。例如，在 Phase 0 中，线程 T0 到 T7 会同时访问共享内存。如果它们访问的数据位于同一个存储体（bank），就会发生冲突，导致访问操作被串行化，严重影响性能。

解决方案：置换共享内存布局 (Permuted Shared Memory Layout)

为了解决存储体冲突问题，可以采用一种置换（或称 swizzling）的共享内存布局。通过一个 XOR 函数将线程索引映射到共享内存地址，从而确保同一访问阶段内的线程访问的是不同的存储体。

• 从全局内存加载：每个线程从全局内存中加载 128 位数据。
• 存储到共享内存：数据在存入共享内存时，其地址根据线程 ID 进行了置换。

这种置换布局可以有效地避免存储和加载过程中的存储体冲突。以下图示（Page 40-43）分阶段展示了置换布局如何避免冲突：

• Phase 0 (T0..T7): 线程 T0-T7 的数据被存储到共享内存的第一行，由于地址置换，它们分布在不同的存储体上，因此无冲突。
• Phase 1 (T8..T15): 线程 T8-T15 的数据被存储到第二行，同样无冲突。
• Phase 2 (T16..T23): 线程 T16-T23 的数据被存储到第三行，无冲突。
• Phase 3 (T24..T31): 线程 T24-T31 的数据被存储到第四行，无冲突。

通过这种方式，数据从全局内存到共享内存的存储过程是无冲突的。

重新审视数据路径

回顾数据路径，我们已经解决了从全局内存到共享内存的流水线延迟和存储冲突问题。接下来需要解决从共享内存加载到寄存器时的冲突问题。

从共享内存加载到寄存器

尽管数据在共享内存中是以置换后的布局存储的，ldmatrix 指令能够高效地处理这种布局。每个线程提供指向其在置換布局中对应数据的指针，ldmatrix 指令会负责将这些分散的数据重新组合成 Tensor Core 所需的逻辑上连续的矩阵。这样，从共享内存到寄存器的加载过程也是无冲突的，从而完成了整个高效的数据馈送路径。

本节展示了数据如何从共享内存加载到线程块（threadblock）中各个线程的寄存器。

• 线程块瓦片（Threadblock Tile）的逻辑视图：该视图显示了线程（T0-T31）在逻辑上的排列。
• 从共享内存加载矩阵：该图展示了每个线程负责加载的数据块。例如，线程T8负责加载特定位置的数据块，线程T16负责加载另一块，以此类推。
• 共享内存指针：图中右侧详细说明了每个线程（T0-T31）的共享内存指针如何指向并加载其对应的数据。这个过程是分阶段进行的，以确保高效和无冲突的数据加载。

下图以线程T8、T16、T24和T31为例，逐步展示了它们的加载过程。

前进到下一个K组

在矩阵乘法（GEMM）中，计算是沿着K维度分块进行的。当一个K维度的块（例如，K=0..15）计算完成后，需要移动到下一个K维度的块（例如，K=16..31）。

为了高效地处理下一个数据块，共享内存的指针会进行更新。如下图所示，通过对共享内存指针进行简单的异或操作（smem_ptr ^= 2），可以快速切换到用于下一个K组数据的缓冲区，这是一种双缓冲（double buffering）技术的实现。

为下一个K组从共享内存加载到寄存器

当计算进行到下一个K组（例如 K=16..31）时，会重复与之前类似的数据加载过程。线程会从共享内存的另一部分（由更新后的指针指向）加载新的数据块到寄存器中。这个过程同样是分阶段（Phase 0, 1, 2, 3）进行的，确保数据流的连续性和高效性。

CUTLASS: 作为Tensor Core最优抽象层的CUDA C++模板库

CUTLASS 旨在为 Tensor Cores 提供一个最优的抽象层，其核心特性包括：
- 延迟容忍的流水线：从全局内存（Global Memory）开始，构建高效的数据流水线以隐藏内存访问延迟。
- 无冲突的共享内存存储：优化数据在共享内存中的布局和存储方式，避免 bank conflict。
- 无冲突的共享内存加载：优化从共享内存加载数据到寄存器的方式，同样避免 bank conflict。

下图展示了 CUTLASS 如何将分块的 GEMM 计算任务在不同的内存层次（全局内存、共享内存、寄存器文件）和计算单元（Tensor Cores）之间进行映射和调度。

CUTLASS: Tensor Core 的最优抽象

CUTLASS 通过 C++ 模板提供了一种高级编程模型，将复杂的底层硬件细节抽象出来。开发者可以通过定义 GEMM 的形状、数据类型和布局来实例化一个高效的矩阵乘法操作。

下图展示了从共享内存到 Warp 级别的矩阵乘法，再到 Tensor Core 操作的映射，并给出了相应的 CUTLASS 代码示例。

CUTLASS 编程模型详解

CUTLASS 的核心组件包括：

• Tile Iterator Constructors (瓦片迭代器构造函数)：初始化指向经过重排的共享内存缓冲区的指针。
• Fragments (片段)：基于寄存器的数据数组，用于存放每个线程的数据。
• Tile Iterator (瓦片迭代器)：
  • load()：从经过重排的共享内存缓冲区中获取数据。
  • operator++()：迭代器前进到共享内存中的下一个逻辑矩阵。
• Warp-level matrix multiply (Warp级矩阵乘法)：将一个大的矩阵乘法分解为多个 Tensor Core 操作。

CUTLASS 在 NVIDIA A100 上的表现

CUTLASS 相对于 cuBLAS 的性能

下图展示了 CUTLASS 2.2 在 CUDA 11 Toolkit 和 NVIDIA A100 GPU 上，相对于高度优化的 cuBLAS 库的性能表现。

• 测试涵盖了不同的 GEMM 类型（DGEMM, IGEMM, SGEMM）和 Tensor Core 操作（f16, f32, TF32），以及不同的矩阵布局（NN, NT, TN, TT）。
• 结果显示，CUTLASS 的性能非常接近 cuBLAS，在多种情况下达到了 cuBLAS 性能的 95% 以上，甚至 99%。对于 TensorOp (TF32) 等操作，性能也能达到 80%-83%。

这证明了 CUTLASS 作为一个高级抽象层，在提供灵活性的同时，几乎没有性能损失。

跨三代 GPU 架构的性能对比

下图进一步比较了 CUTLASS 2.2 在三代不同的 GPU 架构（TitanV, 2080Ti, A100）上相对于 cuBLAS 的性能。

• 结果表明，CUTLASS 在不同架构上都保持了与 cuBLAS 相媲美的高性能。
• 在 A100 上的性能通常是最高的，这得益于新架构的改进。
• 对于 TensorOp (TF32)，由于该特性是 A100 架构引入的，因此只有 A100 的数据。

任意问题规模

CUTLASS 模板覆盖了整个设计空间。下图展示了在NVIDIA A100上使用Tensor Cores（F16 * F16 + F32）的CUTLASS 2.2的性能表现。图中显示，随着内存对齐（alignment）的改善，性能（以GFLOP/s计）显著提升。128b对齐实现了接近峰值的性能。即使是对于较差的对齐（如16b），其性能也远超CUDA 10.2及更早的版本。这表明CUTLASS能够为各种GEMM K维度提供高效的内核。

结论

结论：NVIDIA A100 快速且可编程

CUDA中的NVIDIA A100 Tensor Cores
* 矩阵计算实现数量级的加速。
* 可通过 mma.sync 在CUDA中进行编程，无额外开销。
* 内核设计可避免内存瓶颈。
* CUDA 11工具套件能够达到接近峰值的性能。

CUTLASS 2.2：2020年5月发布
* 用于CUDA开发的开源CUDA C++模板库。
* 提供可重用的构建模块，用于在NVIDIA GPU上利用Tensor Cores。
* 在NVIDIA Ampere架构上实现接近最优的性能。

立即尝试！ https://github.com/NVIDIA/cutlass

参考文献

NVIDIA Ampere 架构:
* "Inside the NVIDIA Ampere Architecture" (GTC 2020 - S21730)
* "NVIDIA Ampere Architecture In-Depth" (博客文章)
* "CUDA New Features and Beyond" (GTC 2020 - S21760)
* "Tensor Core Performance on NVIDIA GPUs" (GTC 2020 - S21929)
* "Inside the Compilers, Libraries and Tools for Accelerated Computing" (GTC 2020 - S21766)

CUTLASS:
* https://github.com/NVIDIA/cutlass (开源软件，New BSD 许可证)
* GTC 2018 和 GTC 2019 的演讲：GEMM 结构和 Volta Tensor Cores
* CUTLASS Parallel For All 博客文章