yiran103的专栏

前一篇文章介绍了 Stream-K: Work-centric Parallel Decomposition for Dense Matrix-Matrix Multiplication on the GPU 论文，下面对其代码实现进行分析。cutlass 的 examples/47_ampere_gemm_universal_streamk 展示了 GEMM Stream-K 算法在 Ampere 架构上的使用。对比了普通 Gemm 以及 Split-K 算法和 Stream-K 的性能

在 NVIDIA GTC22 秋季会议上，CUTLASS: Python API, Enhancements, and NVIDIA Hopper 介绍了 CUTLASS~2.11 中引入的 Stream-K 分解：在这里插入图片描述几个月后公开的 Stream-K: Work-centric Parallel Decomposition for Dense Matrix-Matrix Multiplication on the GPU 论文对其进行了更详细的介绍。CUTLASS 的 GEMM 实现由三大

表 1.1 比较了一些处理器上乘法和除法的时间。这张表展示了乘法和除法时间差距的增长趋势。因此，中提出了使用整数乘法进行任意非零整数常数和运行时不变量之间除法的算法。文档中记录了更广泛的处理指令性能，其中 Intel IceLake 处理器的乘除法指令延迟和吞吐倒数如下表所示：可以看出，在现代 CPU 处理器上除法开销大的情况并未发生改变。NVIDIA 和 AMD GPU 均不支持整数除法指令，CUDA C++ Programming Guide。

CUTLASS 是 CUDA C++ 模板抽象的集合，用于在 CUDA 内的所有级别和规模上实现高性能矩阵-矩阵乘法 (GEMM) 和相关计算。它采用了类似于 cuBLAS 和 cuDNN 中实现的分层分解和数据移动策略。CUTLASS 最新版本为3.3，相比1.3.3变动较大。然而重温一下1.3.3仍然是有意义的。因为它更易于理解： 与 PROGRAMMING TENSOR CORES: NATIVE VOLTA TENSOR CORES WITH CUTLASS 中介绍的内容相匹配；

PROGRAMMING TENSOR CORES: NATIVE VOLTA TENSOR CORES WITH CUTLASS 源自于 GTC Silicon Valley-2019: cuTENSOR: High-performance Tensor Operations in CUDA，介绍了 CUTLASS 1.3 中基于 Volta Tensor Core 实现高效矩阵乘法计算的策略。主要内容为以下三点： CUDA 10.1中mma.sync指令介绍； Global Memory

Modeling Deep Learning Accelerator Enabled GPUs 发表在 ISPASS 2019 上。文章研究了 NVIDIA 的 Volta 和 Turing 架构中张量核的设计，并提出了 Volta 中张量核的架构模型。基于实现该模型，并且支持 CUTLASS 运行。发现其性能与硬件非常吻合，与 Titan V GPU 相比，获得了99.6%的 IPC 相关性。文中还展示了 Turing 架构中张量核的操作数矩阵元素到线程的映射，并发现它们与 Volta 张量核的行为不同。

Implicit GEMM 是将卷积操作表述为 GEMM (广义矩阵-矩阵积)。卷积接受激活张量并对其应用滑动滤波器以产生输出张量。二维卷积可以映射到矩阵乘：组建一个包含激活张量元素的卷积矩阵，然后由滤波张量形成的矩阵乘以该矩阵。该算法的最早形式通过通常称为 im2col 的操作显式构造卷积矩阵。生成的矩阵按照滤波器大小复制每个激活元素，消耗额外的存储容量和内存带宽。隐式 GEMM 算法是 CUDA 中分块、分层 GEMM 计算的一种变体：当数据从全局内存加载到共享内存时，通过周密地更新指针和谓词，它

CUTLASS 实现了 CUTLASS: Fast Linear Algebra in CUDA C++ 和 CUTLASS GTC2018 talk 中描述的分层分块结构。基本的三重嵌套循环计算矩阵乘法可以应用分块和拼贴，以匹配硬件、内存局部性和并行编程模型中的并发性。CUTLASS 中 GEMM 映射到 NVIDIA GPU 的结构如以下嵌套循环所示。