ACCELERATING BACKWARD DATA GRADIENT BY INCREASING TENSOR CORE UTILIZATION IN CUTLASS

Manish Gupta, March 21, 2021

致谢
议程
概述
CUTLASS 卷积性能
CUTLASS 跨步反向数据梯度 (DGRAD)
- 跨步 DGRAD (1D 理解)
- 跨步前向传播与跨步反向传播 (Strided Fprop vs. Strided Dgrad)
- 朴素跨步反向传播与 CUTLASS 跨步反向传播 (Naïve Strided Dgrad vs. CUTLASS Strided Dgrad)
- CUTLASS 跨步反向传播 (1D 卷积，步长 = 2，滤波器 = 3)
- CUTLASS 跨步反向传播 [小滤波器] (1D 卷积，步长 = 2，滤波器 = 1)
- CUTLASS 跨步反向传播 (步长 > 1)
- 2D 中的跨步反向传播 (隐式 GEMM 卷积)
- 4D 张量上的 2D 反向传播梯度 - 定义
隐式 GEMM 卷积 (Implicit GEMM Convolution)
- 反向数据梯度 (Dgrad)
- 朴素跨步反向传播 (隐式 GEMM) - CUTLASS 2.5
- CUTLASS 跨步反向传播 (隐式 GEMM) - CUTLASS 2.6+
- 实现细节与算法步骤
  - 加载输出梯度矩阵 (Dy) 和滤波器矩阵 (w) 的瓦片
  - 计算矩阵乘积累加
  - 迭代 KRS 维度
- 构建连贯完整的抽象
- 高效地组合复杂算法的抽象
- 隐式GEMM卷积 - CUTLASS 2.5 组件
- 隐式GEMM卷积 - STRIDED DGRAD (CUTLASS 2.6+ 组件)
更多地利用 CUTLASS
- 使用 Tensor Cores 加速单精度计算
- 分组 GEMM (Grouped GEMM)
- CUTLASS Python
结论
参考文献

致谢

幻灯片中包含致谢页，感谢了CUTLASS GitHub社区、CUTLASS团队的开发者和产品管理人员，以及众多贡献者。

议程

本次演讲的议程包括：
* 概述：介绍CUTLASS 2.6 (GTC March 2021) 到 CUTLASS 2.9 (GTC March 2022) 的发展。
* 卷积性能 (CUTLASS 2.9)：分析 Fprop, Dgrad 和 Wgrad 在 CUDA 11.6 上的性能。
* 跨步 Dgrad：探讨朴素的跨步 Dgrad（包含冗余 MMAs）和 CUTLASS 跨步 Dgrad（移除冗余 MMAs）。
* 隐式 Gemm 卷积：构建连贯完整的抽象，高效地组合复杂算法。
* 使用 CUTLASS 完成更多任务：介绍使用 Tensor Cores (3xTF32) 加速单精度算术运算，分组 GEMM，以及 CUTLASS Python 示例。

概述

CUTLASS 是用于深度学习和线性代数的 CUDA C++ 模板库。该库在不同GTC年份和CUDA版本下持续演进，并增加了新功能。

CUTLASS版本和功能时间线
Page 5

CUTLASS 的新进展 (自2021年7月起)

CUTLASS 2.6 - 2021年7月
- 新增并改进了跨步 Dgrad (分析迭代器)。
- 实现了 Epilogue 融合模式。
CUTLASS 2.7 - 2021年9月
- 新增并改进了跨步 Dgrad (优化迭代器)。
- 引入了半精度 GELU_taylor 激活函数。
- 对融合 GEMM + GEMM 示例进行了性能调优。
- 支持小于 128b 对齐的卷积。
CUTLASS 2.8 - 2021年11月
- 利用 Tensor Cores (3xTF32) 加速单精度算术运算。
- 每个批次支持可变大小的 GEMM 问题 (分组 GEMM)。
CUTLASS 2.9 - 即将发布的版本
- 第一层卷积。
- CUTLASS Python 示例。
- 更多功能...

CUTLASS 卷积性能

CUTLASS 2.8 在 NVIDIA A100 (CUDA 11.6) 上混合精度训练 (F16 <- F16 * F16 + F32) 相对于 cuDNN 的性能表现：

CUTLASS 2.8 卷积性能与cuDNN的比较
Page 7

图表显示了 ResNet50 不同层（2到20）的 fprop、dgrad 和 wgrad 相对于 cuDNN 的加速比。

fprop 在大多数层上表现出与 cuDNN 相近或略高的性能，最高达到约 110%。
dgrad 的性能波动较大，在某些层（如层7）表现出高达约 135% 的加速，但在其他层可能低于 cuDNN。
wgrad 的性能普遍低于 cuDNN，大部分在 80% 到 100% 之间。
为了获得最佳性能，建议使用 CUDA 11.4 或更高版本。

CUTLASS 跨步反向数据梯度 (DGRAD)

CUTLASS 在处理跨步 Dgrad 方面取得了显著进展。CUTLASS 2.8 的跨步 Dgrad 性能相较于 CUTLASS 2.5 有了巨大提升。

CUTLASS 2.5 vs CUTLASS 2.8 跨步Dgrad性能对比
Page 9

该图比较了 CUTLASS 2.5 和 CUTLASS 2.8 在 NVIDIA GA100 (1005Mhz) 上，使用 TensorOp (F16 <- F16*F16+F32) 进行跨步 Dgrad (步长为 2x2) 相对于 cuDNN 8.3.21 (CUDA 11.6 Toolkit) 的性能。
* CUTLASS 2.5 (黄色条) 的性能远低于 cuDNN，在所有测试层中都表现不佳，最低只有 20% 左右。
* CUTLASS 2.8 (绿色条) 则显示了显著的性能提升，在许多层中超过了 cuDNN，最高达到约 135% 的加速比 (例如层10和层14)。
* 测试的跨步层包括：
* 层 1-6：ResNet50，1x1 滤波器
* 层 7-9：RNXT JoC，3x3 滤波器
* 层 10-15：MaskRCNN MLPERF，1x1 滤波器

为了获得最佳性能，建议使用 CUDA 11.4 或更高版本。

跨步 DGRAD (1D 理解)

为了更好地理解跨步 DGRAD，我们首先回顾 1D 卷积的基本概念。

前向传播 | Y = FPROP(X, W)
对于步长为 1，滤波器大小为 3 的 1D 卷积：

1D 步长1，滤波器3的前向传播
Page 11

图示展示了激活 (x) 和滤波器 (w) 如何通过乘积累加操作生成输出 (y) 的过程。每个输出元素 (y) 由激活 (x) 和滤波器 (w) 的乘积累加贡献。

反向数据梯度 | DX = DGRAD(DY, W)
对于步长为 1，滤波器大小为 3 的 1D 卷积的反向数据梯度：

1D 步长1，滤波器3的反向数据梯度
Page 12

图示展示了输出梯度 (Dy) 和滤波器 (w) 如何通过乘积累加操作生成激活梯度 (Dx) 的过程。每个激活梯度元素 (Dx) 由输出梯度 (Dy) 和滤波器 (w) 的乘积累加贡献。

跨步前向传播 | Y = FPROP(X, W)
对于步长为 2，滤波器大小为 3 的 1D 卷积：

1D 步长2，滤波器3的跨步前向传播
Page 13

图示展示了当步长增加到 2 时，激活 (x) 和滤波器 (w) 如何通过乘积累加操作生成输出 (y)。此时，输出 (y) 中的元素间隔更远。

跨步反向数据梯度 | DX = DGRAD(DY, W)
对于步长为 2，滤波器大小为 3 的 1D 卷积的反向数据梯度：

1D 步长2，滤波器3的跨步反向数据梯度
Page 14

图示展示了当步长为 2 时，输出梯度 (Dy) 和滤波器 (w) 如何通过乘积累加操作生成激活梯度 (Dx)。

反向数据梯度 (DGRAD) [跨步] 朴素实现
对于步长为 2，滤波器大小为 3 的 1D 卷积的朴素实现：

1D 步长2，滤波器3的跨步反向数据梯度朴素实现
Page 15

朴素实现中，为了处理跨步情况，会在滤波器贡献中引入零填充 (用 'Z' 表示)。这可能导致冗余的乘积累加操作。每个激活梯度元素 (Dx) 仍由输出梯度 (Dy) 和滤波器 (w) 的乘积累加贡献。

跨步前向传播与跨步反向传播 (STRIDED FPROP VS. STRIDED DGRAD)

1D 卷积（步长 = 2，滤波器 = 3）

本节对比了步长为 2、滤波器大小为 3 的 1D 卷积中的跨步前向传播 (STRIDED FPROP) 和跨步梯度反向传播 (STRIDED DGRAD)。

STRIDED FPROP: 协作线程阵列 (CTA) 平铺在输出 (y) 的元素上。
STRIDED DGRAD: 协作线程阵列 (CTA) 平铺在激活梯度 (Dx) 的元素上。

STRIDED FPROP VS. STRIDED DGRAD
Page 16

其中 CTA 代表 Cooperative Thread Array 或 Thread block。

朴素跨步反向传播与 CUTLASS 跨步反向传播 (NAÏVE STRIDED DGRAD VS CUTLASS STRIDED DGRAD)

1D 卷积（步长 = 2，滤波器 = 3）

本节对比了朴素跨步反向传播（CUTLASS 2.5）和 CUTLASS 优化的跨步反向传播（CUTLASS 2.6）在 1D 卷积（步长 = 2，滤波器 = 3）中的实现。

朴素跨步反向传播 (CUTLASS 2.5)：
- 需要 3 次 GEMM-K 迭代。
- 处理中包含冗余的“Z”元素。
CUTLASS 跨步反向传播 (CUTLASS 2.6)：
- 需要 1-2 次 GEMM-K 迭代。
- 通过优化，直接处理有效数据，减少了冗余。

NAÏVE STRIDED DGRAD VS CUTLASS STRIDED DGRAD
Page 17

CUTLASS 跨步反向传播 (CUTLASS STRIDED DGRAD)

1D 卷积（步长 = 2，滤波器 = 3）

CUTLASS 跨步反向传播方法通过以下三个阶段进行优化：

序言 (Prologue)：映射 Dx 元素以跳过冗余的乘加操作。
主循环 (Mainloop)：只运行有用的 GEMM-K 迭代。
- CTA 0-1 计算 Dx[i]，其中 i={1, 3, 5, 7}（2 次乘加）。
- CTA 2-3 计算 Dx[i]，其中 i={0, 2, 4, 6}（1 次乘加）。
结语 (Epilogue)：映射 Dx 元素以存储到正确的目的位置。

CUTLASS STRIDED DGRAD
Page 18

CUTLASS 跨步反向传播 [小滤波器] (CUTLASS STRIDED DGRAD [SMALL FILTER])

1D 卷积（步长 = 2，滤波器 = 1）| 泛化（步长 > 滤波器）

本节比较了在小滤波器情况下（步长 = 2，滤波器 = 1）以及步长大于滤波器大小的泛化场景下，朴素跨步反向传播 (CUTLASS 2.5) 和 CUTLASS 跨步反向传播 (CUTLASS 2.6) 的表现。

朴素跨步反向传播 (CUTLASS 2.5)：在小滤波器情况下仍然可能包含冗余操作，如图中“Z”所示。
CUTLASS 跨步反向传播 (CUTLASS 2.6)：即使在小滤波器情况下也能进行有效优化，只处理必要的计算，减少了迭代次数。

CUTLASS STRIDED DGRAD [SMALL FILTER]
Page 19

CUTLASS 跨步反向传播 (CUTLASS STRIDED DGRAD)

(步长 > 1)

当步长大于 1 时，CUTLASS 跨步反向传播的伪代码结构如下：

// Prologue
id -> mapped_id // 将 CTA 映射到计算有效的 Dx 元素

// Mainloop
accumulators = 0 // 清零累加器
if (isMainloopRequired(blockIdx.x)) {
    accumulators = Dy * w
}

// Epilogue
id -> mapped_id // 映射回目标 Dx 元素位置
Dx_source = ((beta==0) ? 0 : Dx[mapped_id])
Dx[mapped_id] = alpha * (accumulators) + beta * Dx_source // 存储 Dx

CUTLASS STRIDED DGRAD (Stride > 1)
Page 20

2D 中的跨步反向传播 (STRIDED DGRAD IN 2D)

(隐式 GEMM 卷积)

本部分将探讨 2D 卷积中跨步反向传播的实现，特别是通过隐式 GEMM (General Matrix Multiply) 卷积的方式。

STRIDED DGRAD IN 2D (IMPLICIT GEMM CONVOLUTIONS)
Page 21

4D 张量上的 2D 反向传播梯度 - 定义 (2D DGRAD ON 4D TENSORS - DEFINITION)

反向数据传播 | Dx = CONV (Dy, w)

2D 反向传播梯度（DGRAD）在 4D 张量上的定义如下：

$$ \mathbf{Dx}[n, h, w, c] = \sum_{k=0}^{K-1} \sum_{r=0}^{R-1} \sum_{s=0}^{S-1} (\mathbf{Dy}[n, \bar{p}(h,r), \bar{q}(w,s), k] * \mathbf{w}[k,r,s,c]) $$

其中：

$$ \bar{p}(h,r) = (h + \text{pad\_h} - r * \text{dilation\_h}) / \text{stride\_h} $$

$$ \bar{q}(w,s) = (w + \text{pad\_w} - s * \text{dilation\_w}) / \text{stride\_w} $$

激活梯度张量 (Activation Gradient Tensor) (Dx)：NHWC = {1, 4, 4, 3}
输出梯度张量 (Output Gradient Tensor) (Dy)：NPQK = {1, 3, 3, 4}
滤波器张量 (Filter Tensor) (w)：KRSC = {4, 2, 2, 3}

2D DGRAD ON 4D TENSORS - DEFINITION
Page 22

隐式 GEMM 卷积 (IMPLICIT GEMM CONVOLUTION)

反向数据梯度 (Dgrad)

隐式 GEMM 卷积通过将 4D 张量映射到 2D 矩阵来执行 Dgrad 计算：

4D 张量:
- Dy[N,P,Q,K] : 4D 输出梯度张量
- w[K,R,S,C] : 4D 滤波器张量
- Dx[N,H,W,C] : 4D 激活梯度张量
GEMM M-by-N-by-K 维度:
- GEMM-M = NHW
- GEMM-N = C
- GEMM-K = KRS
2D 矩阵:
- GEMM-A : Dy 矩阵 (NHW, KRS)
- GEMM-B : w 矩阵 (KRS, C)
- GEMM-C : Dx 矩阵 (NHW, C)

IMPLICIT GEMM CONVOLUTION
Page 23

朴素跨步反向传播 (隐式 GEMM) - CUTLASS 2.5 (NAÏVE STRIDED DGRAD (IMPLICIT GEMM) - CUTLASS 2.5)

为了覆盖整个 GEMM-K (KRS) 维度：

通过遍历以下维度来处理 KRS 维度中的瓦片：
- 滤波器 s 位置
- 滤波器 r 位置
- Tile_K 输出通道 k 元素
启动平铺迭代以覆盖所有通道元素 (K) 和滤波器位置 (R-by-S)。

计算 gemm_k_iterations 的公式为：
gemm_k_iterations = R * S * ((K + Tile_K - 1) / Tile_K)

示例尺寸：
- Tile_M = 128
- Tile_N = 128
- Tile_K = 32
- R-by-S = 3-by-3
- Stride = 2-by-2

NAÏVE STRIDED DGRAD (IMPLICIT GEMM) - CUTLASS 2.5
Page 24

CUTLASS 跨步反向传播 (隐式 GEMM) - CUTLASS 2.6+ (CUTLASS STRIDED DGRAD (IMPLICIT GEMM) - CUTLASS 2.6+)

CUTLASS 2.6+ 版本的隐式 GEMM 跨步反向传播优化如下：

序言 (Prologue)：将 NHW 维度映射到有效的起始滤波器位置。
主循环 (Mainloop)：仅计算对激活梯度 (Dx) 有贡献的 MMAs（矩阵乘加）。
结语 (Epilogue)：映射累加器行以写入到映射的 NHW 行。

示例尺寸：
- Tile_M = 128
- Tile_N = 128
- Tile_K = 32
- R-by-S = 3-by-3
- Stride = 2-by-2

CUTLASS STRIDED DGRAD (IMPLICIT GEMM) - CUTLASS 2.6+
Page 25

实现细节与算法步骤

本节继续深入探讨隐式 GEMM 卷积的实现细节。

隐式 GEMM 卷积算法的主要步骤如下：

将 Dy 和 w 矩阵的瓦片加载到共享内存中。
在共享内存的操作数上计算矩阵乘加 (mma)。
迭代遍历 KRS 维度。

IMPLICIT GEMM CONVOLUTION
Page 27

加载输出梯度矩阵 (Dy) 和滤波器矩阵 (w) 的瓦片

算法的第一步是加载 Dy 和 w 矩阵的瓦片到共享内存中。

示例：
- Tile_M = 128
- Tile_N = 128
- Tile_K = 32
- 输入类型 = F16

IMPLICIT GEMM CONVOLUTION
Page 28

详细加载过程：

GEMM-A 输出梯度矩阵 (Dy) 的瓦片:
- 加载 128-by-32 元素。
- 每行映射到唯一的 (n, h, w) 坐标。
- 每列映射到唯一的滤波器 k 坐标。
GEMM-B 滤波器矩阵 (w) 的瓦片:
- 加载 32-by-128 元素。
- 每行映射到唯一的滤波器 k 坐标。
- 每列映射到唯一的通道 c 坐标。
每个线程从每个瓦片发出多个向量加载:
- 例如，T0 的第一次加载将 k=0..7 的 F16 操作数带入共享内存。

示例尺寸与 Page 28 相同。

IMPLICIT GEMM CONVOLUTION
Page 29

计算矩阵乘积累加

在隐式GEMM卷积中，warp级的矩阵乘积累加（MMA）操作器直接使用 cutlass::gemm::warp 模块来利用NVIDIA Tensor Cores。

上图展示了示例瓦片大小：Tile_M = 128，Tile_N = 128，Tile_K = 32。数据从全局内存（Global Memory）通过共享内存（Shared Memory）和寄存器文件（Register Files）流向Tensor Cores，使用 cutlass::gemm::warp::MmaTensorOp 进行处理。

迭代 KRS 维度

隐式GEMM卷积算法包含以下步骤：

将Dy（输出梯度矩阵）和滤波器w的瓦片加载到共享内存中。
在共享内存中的操作数上计算矩阵乘积累加（mma）。
遍历KRS维度：
a) 推进以在共享内存中加载下一个瓦片。
b) 确保所有滤波器位置 (r,s) 和输出通道K的累加。

示例瓦片大小：Tile_M = 128，Tile_N = 128，Tile_K = 32。输入类型为F16。

上图说明了滤波器矩阵 (w)、输出梯度矩阵 (Dy) 和激活梯度矩阵 (Dx) 的尺寸以及瓦片大小。

为了覆盖整个GEMM-K (KRS) 维度：
通过遍历以下内容，在KRS维度中处理瓦片：
* 滤波器 s 位置
* 滤波器 r 位置
* Tile_K 滤波器 k 元素

启动足够的迭代次数以覆盖所有输出通道元素 (K) 和滤波器位置 (R-by-S)：
num_tiled_iterations = R * S * ((K + Tile_K – 1)/Tile_K)

示例大小：Tile_M = 128，Tile_N = 128，Tile_K = 32，R-by-S = 3-by-3，Stride = 1-by-1。

上图展示了如何通过瓦片迭代来覆盖3x3滤波器位置和0..K-1通道元素。

构建连贯完整的抽象

cutlass::conv::threadblock::Iterators
CUTLASS卷积迭代器实现了以下抽象：
* advance(): 移动到GEMM-K中的下一个瓦片迭代。
* operator++(): 移动到线程的下一个加载位置。
* at(): 应用 p (h,r) 和 q (w,s) 函数将 hw 映射到 pq，并返回Dy张量 [n,p,q,k] 中的坐标。
* valid(): 检查全局内存中张量的越界访问。
* get(): 根据张量坐标从全局内存中获取指针。

上图详细说明了这些迭代器函数在遍历输出梯度矩阵 (Dy) 和滤波器矩阵 (w) 时的作用。

高效地组合复杂算法的抽象

幻灯片强调了抽象在高效组合复杂算法中的重要性。

隐式GEMM卷积 - CUTLASS 2.5 组件

CUTLASS 2.5 隐式GEMM卷积的组件包括 cutlass::conv::threadblock、cutlass::gemm::warp 和 cutlass::epilogue。
数据流经全局内存 -> 共享内存 -> 寄存器文件 / CUDA/Tensor Cores -> SMEM -> CUDA Cores -> 全局内存。
每个阶段都定义了其布局（Layout）、迭代器（IteratorsA/B）和矩阵乘法累加（Mma）操作。

上图展示了CUTLASS 2.5的组件架构及其数据流，以及一个 cutlass::conv::kernel::ImplicitGemmConvolution 的代码片段。

隐式GEMM卷积 - STRIDED DGRAD (CUTLASS 2.6+ 组件)

CUTLASS 2.6+ 对步进（strided）DGRAD隐式GEMM卷积的组件进行了更新。主要更新体现在 IteratorsA/B 中使用了 PredicatedTileIteratorStridedDgrad，并且内核名称变更为 cutlass::conv::kernel::ImplicitGemmConvolutionStridedDgrad。

上图展示了CUTLASS 2.6+的更新组件架构和对应的代码片段。

结论

最终结论总结了CUTLASS的各项优势和进展：

CUTLASS卷积
- CUTLASS卷积性能与cuDNN 8.3.3 (> 95%) 相当。
- CUTLASS 2.6中的步进Dgrad比CUTLASS 2.5提供高达4倍的加速。
- 支持小于128b的对齐。
分组GEMM
- 为混合专家自然语言模型实现了1.75倍的端到端性能加速。
- 将可变批大小的GEMM分组到一个单一内核中。
- 不需要预处理。
使用3xTF32在Tensor Cores上加速单精度计算
- 相比CUDA核心的单精度计算，性能提升2倍。
- 相比CUDA核心的单精度计算，精度更好。
CUTLASS Python示例
- 期待Python程序员提供更多反馈。
CUTLASS 2.9
- 即将发布的版本值得关注。

参考文献

GTC 2022 会议
- GTC 2022 [S41491] NVIDIA数学库的最新发展
- GTC 2022 [S41486] CUDA: 新特性与展望
- GTC 2022 [S41606] 使用CUTLASS融合多个GEMM以达到极致性能
- GTC 2022 [S41611] 使用Int4/Int8混合精度进行自动模型压缩和加速的通用框架
GTC 2022 与专家联系 (现场会议)
- [CWE41721] 与专家联系: NVIDIA数学库 (星期一, 3月21日, 2:00 PM - 2:50 PM PDT)
往年 (2018 - 2021) GTC CUTLASS 会议
- GTC 2018 [S8854] CUTLASS: CUDA中各级别和规模的稠密线性代数软件原语
- GTC 2019 [S9593] cuTENSOR: CUDA中的高性能张量操作 (与cuTENSOR联合演讲)
- GTC 2020 [S21745] 开发CUDA内核以将Tensor Cores推向NVIDIA A100的绝对极限
- GTC 2021 [S31883] CUTLASS中加速卷积与Tensor Cores

ACCELERATING BACKWARD DATA GRADIENT BY INCREASING TENSOR CORE UTILIZATION IN CUTLASS

目录

致谢

议程

概述

CUTLASS 卷积性能

CUTLASS 跨步反向数据梯度 (DGRAD)

跨步 DGRAD (1D 理解)

跨步前向传播与跨步反向传播 (STRIDED FPROP VS. STRIDED DGRAD)

1D 卷积（步长 = 2，滤波器 = 3）

朴素跨步反向传播与 CUTLASS 跨步反向传播 (NAÏVE STRIDED DGRAD VS CUTLASS STRIDED DGRAD)

1D 卷积（步长 = 2，滤波器 = 3）

CUTLASS 跨步反向传播 (CUTLASS STRIDED DGRAD)

1D 卷积（步长 = 2，滤波器 = 3）

CUTLASS 跨步反向传播 [小滤波器] (CUTLASS STRIDED DGRAD [SMALL FILTER])

1D 卷积（步长 = 2，滤波器 = 1）| 泛化（步长 > 滤波器）

CUTLASS 跨步反向传播 (CUTLASS STRIDED DGRAD)

(步长 > 1)

2D 中的跨步反向传播 (STRIDED DGRAD IN 2D)

(隐式 GEMM 卷积)

4D 张量上的 2D 反向传播梯度 - 定义 (2D DGRAD ON 4D TENSORS - DEFINITION)

反向数据传播 | Dx = CONV (Dy, w)

隐式 GEMM 卷积 (IMPLICIT GEMM CONVOLUTION)

反向数据梯度 (Dgrad)

朴素跨步反向传播 (隐式 GEMM) - CUTLASS 2.5 (NAÏVE STRIDED DGRAD (IMPLICIT GEMM) - CUTLASS 2.5)

CUTLASS 跨步反向传播 (隐式 GEMM) - CUTLASS 2.6+ (CUTLASS STRIDED DGRAD (IMPLICIT GEMM) - CUTLASS 2.6+)

实现细节与算法步骤

加载输出梯度矩阵 (Dy) 和滤波器矩阵 (w) 的瓦片

计算矩阵乘积累加

迭代 KRS 维度

构建连贯完整的抽象

高效地组合复杂算法的抽象

隐式GEMM卷积 - CUTLASS 2.5 组件

隐式GEMM卷积 - STRIDED DGRAD (CUTLASS 2.6+ 组件)

更多地利用 CUTLASS

使用 Tensor Cores 加速单精度计算

分组 GEMM (GROUPED GEMM)

CUTLASS Python

结论

参考文献

登录

注册

忘记密码

重发验证邮件