CUDA-MODE课程笔记｜CUDA性能检查清单

我的课程笔记，欢迎关注：https://github.com/BBuf/how-to-optim-algorithm-in-cuda/tree/master/cuda-mode

CUDA-MODE课程笔记 第8课: CUDA性能检查清单

课程笔记

这节课实际上算是 CUDA-MODE 课程笔记 第一课: 如何在 PyTorch 中 profile CUDA kernels 这节课更细节的讲解。另外关于nsight compute相关指标细节解释可以参考 CUDA-MODE 第一课课后实战（上） ， CUDA-MODE 第一课课后实战（下） 这两篇笔记。

将GPU用于计算，我们最关心的肯定是性能。比较幸运的是，当我们掌握一些性能优化技巧之后它往往会经常被用到。这节课将会系统的介绍这些性能优化技巧。

本节课的课件和代码都在 https://github.com/cuda-mode/lectures 开源，我们可以用nvcc编译lecture8下面的cu文件，然后使用ncu进行profile。此外，这里的方法遵循了  https://arxiv.org/pdf/1804.06826.pdf 这篇论文的风格。up主也非常推荐大家阅读下这篇论文，用claude 3.5问了一下paper的主要内容，如下截图：

可以看到这篇论文主要是对Volta架构的GPU的架构细节进行分析，对性能优化是很重要的。

这张Slides从物理学的角度分析了下SRAM和DRAM的区别：

• DRAM由1个晶体管和1个电容器构成；SRAM: 由6个晶体管构成
• SRAM 比 DRAM 更快，但也更贵；SRAM 占用更多空间且发热更多；实际上SRAM就对应了GPU的Shared Memory，而DRAM对应的则是Shared Memory。

这里的youtube链接作者Bill是NVIDIA的首席科学家，他解释了很多为什么GPU设计成现在这个样子，并且由浅入深，基础细节讲的非常清楚。

这里的"性能检查清单"（Performance checklist），列出了一系列优化GPU程序性能的策略和技巧：

• 合并全局内存访问（Coalesced Global Memory Access）
• 最大化占用率（Maximize occupancy）
• 理解是内存受限还是计算受限（Understand if memory or compute bound）
• 最小化线程分化（Minimize control divergence）
• Tiling以更好的重用数据（Tiling of reused data）
• 私有化（Privatization）
• Thread Coarsening
• 使用更好的数学方法重写算法（Rewrite your algorithm using better math）

这里的Privatization指的应该就是Shared Memory/寄存器优化全局内存读取，而Coarsening大概指的就是一个线程应该完成多少任务，一般情况下我们让一个线程完成的任务尽量少，但是在Compute Bound情况下，让一个线程执行更多的工作可以让程序运行得更快。最后一点更好的数学方法重写算法的经典例子就是Flash Attention。

这张Slides讲述了GPU内存访问延迟的相关内容，下面的Figure3和表格都来自 https://arxiv.org/pdf/2208.11174 ，这个表格（Table IV），列出了不同类型内存的访问延迟（以时钟周期为单位）：

• 全局内存（Global memory）: 290 cycles
• L2 缓存: 200 cycles
• L1 缓存: 33 cycles
• 共享内存（Shared Memory）: 读取23 cycles，写入19 cycles

我后面也找到了这个paper里面做micro benchmark的代码：https://www.stuffedcow.net/research/cudabmk?q=research/cudabmk ，后面如果有空继续阅读下这篇 paper 以及测试代码。

这张Slides讲述了延迟（latency）在计算机系统中的重要性和一些相关概念。

• 标题 "It's the latency stupid" 强调了延迟的重要性。

• 吞吐量（Throughput）和延迟（Latency）的对比：

• 吞吐量容易提高，但延迟却很难降低。
• 举例说明：即使你可以并行使用80条电话线，每条线传输一个比特，但100毫秒的延迟仍然存在。

• 量化（Quantization）技术：

• 用于减少数据包大小的一种方法。
• 例如，Bolo（可能是某个系统或协议）尽可能使用字节（byte）而不是16位或32位字来减少数据包大小。

• 底部提供了一个网址链接，包含更多关于这个话题的详细讨论。

这张Slides开始介绍内存合并（Memory Coalescing）的概念。我们无法减少延迟，但可以通过读取连续的内存元素来隐藏延迟。Slides建议在进行案例研究时要关注以下三个方面：

• DRAM Throughput（DRAM吞吐量）
• Duration（持续时间）
• L1 cache throughput（L1缓存吞吐量）

这里说的内存合并的案例就是 https://github.com/cuda-mode/lectures/blob/main/lecture_008/coalesce.cu 这里所展示的。代码如下：

#include   
#include   
  
__global__ void copyDataNonCoalesced(float *in, float *out, int n) {  
    int index = blockIdx.x * blockDim.x + threadIdx.x;  
    if (index         out[index] = in[(index * 2) % n];  
    }  
}  
  
__global__ void copyDataCoalesced(float *in, float *out, int n) {  
    int index = blockIdx.x * blockDim.x + threadIdx.x;  
    if (index         out[index] = in[index];  
    }  
}  
  
void initializeArray(float *arr, int n) {  
    for(int i = 0; i         arr[i] = static_cast<float>(i);  
    }  
}  
  
int main() {  
    const int n = 1 <    float *in, *out;  
  
    cudaMallocManaged(&in, n * sizeof(float));  
    cudaMallocManaged(&out, n * sizeof(float));  
  
    initializeArray(in, n);  
  
    int blockSize = 128; // Define block size  
    // int blockSize = 1024; // change this when talking about occupancy  
    int numBlocks = (n + blockSize - 1) / blockSize; // Ensure there are enough blocks to cover all elements  
  
    // Launch non-coalesced kernel  
    copyDataNonCoalesced<<>>(in, out, n);  
    cudaDeviceSynchronize();  
  
    initializeArray(out, n); // Reset output array  
  
    // Launch coalesced kernel  
    copyDataCoalesced<<>>(in, out, n);  
    cudaDeviceSynchronize();  
  
    cudaFree(in);  
    cudaFree(out);  
  
    return 0;  
}  

这里段程序比较简单，用于演示内存合并（Memory Coalescing）的概念和其对性能的影响。它主要做了以下事情：

• 定义了两个CUDA kernel：

• copyDataNonCoalesced kernel：非合并内存访问模式，以非连续的方式读取输入数组（使用 (index * 2) % n 作为索引），这种访问模式会导致非合并的内存访问，可能降低性能。
• copyDataCoalesced kernel：合并内存访问模式，以连续的方式读取输入数组（直接使用 index 作为索引），这种访问模式允许合并内存访问，可以提高性能。

• 主函数：

• 分配统一内存（Unified Memory）用于输入和输出数组，初始化输入数组。
• 设置CUDA网格和块的大小，分别运行非合并和合并的kernel，在每次kernel执行后使用 cudaDeviceSynchronize() 确保GPU操作完成。

接着使用 nvcc -o benchmark coalesce.cu 来编译程序，然后执行 ncu benchmark 来Profile程序。

对于copyDataNonCoalesced kernel来说，DRAM内存吞吐量大约是89%，L1 Cache的吞吐量是30%，kernel的执行时间是764us。

对于copyDataCoalesced kernel来说，L1 Cache的吞吐量大约是37%，DRAM内存吞吐量是82%，执行时间是558us。

我们可以看到合并内存访问的kernel是有明显的性能提升的。可以预见，随着输入数据量的增大合并内存访问的优势会更明显。ncu的结果里面还提示计算的理论occupancy（100.0%）和实测的实际occupancy占用（77%）之间的差异可能是由于 kernel 执行期间的warp调度开销或工作负载不平衡导致的。在同一kernel 的不同块之间以及块内的不同 warps 之间都可能发生负载不平衡。把上面程序中的 int blockSize = 128 改成 int blockSize = 1024 再次用ncu profile，可以发现occupancy提升到了85.94%。

这张Slides讨论了GPU中的占用率（Occupancy）问题，主要内容如下：

• 两种quantization问题：

• a) Tile quantization：矩阵维度不能被线程块Tile大小整除。
• b) Wave quantization：Tile总数不能被GPU上的SM（流多处理器）数量整除。

• 性能图表比较和分析：

• 左图(a)：cuBLAS v10 上 NN GEMM 的性能
• 右图(b)：cuBLAS v11 上 NN GEMM 的性能
• 两图都是在 M = 1024, N = 1024 的矩阵维度下进行的测试
• 左图(a)显示性能呈现明显的阶梯状，有大幅波动。
• 右图(b)显示性能波动较小，整体更加平滑。我们可以看到cuBLAS v11 可能采用了更好的调度策略或优化技术，减少了由于Tile和Wave Quantization 导致的性能波动。

这张Slides讲解了在PyTorch中使用padding（填充）来解决Tensor Core矩阵乘法维度要求的问题。具体内容如下：

• 在PyTorch环境中，使用padding是解决某些问题的方法。

• 表格展示了不同cuBLAS和cuDNN版本下，使用Tensor Core的数据精度要求。这些要求适用于矩阵维度M、N和K。

• 版本区分：

• 左列：cuBLAS < 11.0 和 cuDNN < 7.6.3 的旧版本
• 右列：cuBLAS ≥ 11.0 和 cuDNN ≥ 7.6.3 的新版本

• 数据类型的要求：

• INT8：旧版本要求16的倍数；新版本总是可用，但16的倍数最高效，在A100上128的倍数最佳。
• FP16：旧版本要求8的倍数；新版本总是可用，但8的倍数最高效，在A100上64的倍数最佳。
• TF32：旧版本不适用；新版本总是可用，但4的倍数最高效，在A100上32的倍数最佳。
• FP64：旧版本不适用；新版本总是可用，但2的倍数最高效，在A100上16的倍数最佳。

新版本的cuBLAS和cuDNN提供了更灵活的Tensor Core使用条件。而A100 GPU可能需要更大的倍数来获得最佳性能。Padding可以用来将矩阵维度调整为这些推荐的倍数，以提高性能。

在CUDA中提升Occupancy的一个方法是修改kernel。

CUDA Occupancy calculator工具可以帮我们自动计算达到更好Occupancy的kernel启动参数，在上一节合并访存的.cu中调用这个Api结果显示，对于T4 GPU，最优的配置是网格大小为40，块大小为1024。代码见：https://github.com/cuda-mode/lectures/blob/main/lecture_008/occupancy.cu

在对这个程序进行ncu的时候有新的问题，那就是下面所展示的：

警告（WRN）：内存利用率高于计算利用率：请查看内存工作负载分析部分以识别DRAM瓶颈。检查内存重放（合并）指标，以确保您正在有效利用传输的字节。同时考虑是否可以通过每次内存访问执行更多工作（kernel融合）或是否有可以（重新）计算的值。

接下来开始讨论这个问题

讨论之前需要先了解一下这张Slides展示的Roofline模型，它决定了一个cuda kernel是compute bound还是memory bound。

这张Slides讲解了算术强度（Arithmetic intensity）的概念及其在处理器性能分析中的应用。这个slides来自gtc2019的一个讲解。

左侧指标是数学运算和内存操作的算法混合，称为算术强度。右侧指标是处理器的ops/byte比率。例如，V100 GPU可以执行125/0.9=139 FLOPS/B。比较算术强度和ops/byte比率可以指出算法受什么因素限制。

下面还给出了操作类型及其算术强度表格：

• Residual addition（残差加法）：0.166，受内存限制
• ReLU activation（ReLU激活）：0.25，受内存限制
• Batch normalization（批量归一化）：O(10)，受内存限制
• Convolution（卷积）：1-10000+（假设FP16数据），可能受内存或数学运算限制

链接：https://developer.download.nvidia.com/video/gputechconf/gtc/2019/presentation/s9926-tensor-core-performance-the-ultimate-guide.pdf

这张slides讲解了ReLU（Rectified Linear Unit）函数的算术强度分析：

• ReLU函数定义：f(x) = max(0, x)，应用于向量的每个元素。

• 操作描述：对每个元素进行1次读取、1次比较操作，可能还有1次写入。

• 数据类型：假设使用float32，即每个数占4字节（32位）。

• 计算分析：

• 操作数（Ops）：1（每个元素一次比较操作）
• 字节数（Byte）：2 * 4 = 8（读取和可能的写入，每次4字节）

• 算术强度计算：

• 最坏情况：1/8（当每个元素都需要写入时）
• 最好情况：1/4（当不需要写入时，只有读取操作） 结论：1/4 < 1，表明ReLU操作受内存带宽限制（Memory bound）

这张Slides对Float16的ReLU进行了算术强度分析，可以看打这种情况下最坏的算术强度是1/4，而不是Float32时的1/8，因此量化是可以提高计算强度的。

这张Slides讲解了矩阵乘法（Matmul）的算术强度分析。其中：

• FLOPS（浮点运算次数）计算：

• 对C中的每个输出元素，需要A的一行和B的一列做点积
• 需要N次乘法和N次加法
• 总FLOPS = M * K * 2N

• 字节数计算：

• 加载矩阵A和B：MN + NK
• 写入输出矩阵C：MK
• 总字节数 = MN + NK + MK

• 算术强度（AI）计算：

• AI = 2MNK / (MN + NK + MK)

• 结论：

• 对于大型矩阵，计算受限（Compute bound）
• 否则，带宽受限（Bandwidth bound）

这张Slides总结了如何优化不同类型的kernels：

• 带宽受限的kernel（Bandwidth Bound Kernels）优化策略：

• Fuse（融合）：合并多个操作以减少内存访问
• Quantize（量化）：使用更小的数据类型来减少内存传输
• Compile（编译）：可能指使用特定的编译技术来优化内存访问模式

• 计算受限的kernel（Compute Bound Kernels）优化策略：

• Write a better algorithm（编写更好的算法）：这意味着需要从算法层面进行优化

关于矩阵乘法Tiling减少全局内存访问请查看以前的 CUDA-MODE 课程笔记 第四课: PMPP 书的第4-5章笔记 。

这张Slides对应这里的代码：https://github.com/cuda-mode/lectures/blob/main/lecture_008/divergence.cu ，主要是对下面2个kernel进行分析：

__global__ void processArrayWithDivergence(int *data, int N) {  
    int idx = blockIdx.x * blockDim.x + threadIdx.x;  
    if (idx         if (data[idx] % 2 == 0) {  
            data[idx] = data[idx] * 2; // 注意这个分支比下面的分支要慢，可能一个Warp里执行这个分支的线程会落后，Warp里的其它线程必须等待这些线程计算完成  
        } else {  
            data[idx] = data[idx] + 1;  
        }  
    }  
}  
  
__global__ void processArrayWithoutDivergence(int *data, int N) {  
    int idx = blockIdx.x * blockDim.x + threadIdx.x;  
    if (idx         int isEven = !(data[idx] % 2); // 这里做的事情和上面相同，但是规避了线程分化问题  
        data[idx] = isEven * (data[idx] * 2) + (!isEven) * (data[idx] + 1);  
    }  
}  

• 控制分歧（control divergence）与占用率（occupancy）有关，但如果条件语句导致大量线程闲置，这是不好的。
• processArrayWithDivergence 耗时 0.074272 毫秒; processArrayWithoutDivergence 耗时 0.024704 毫秒;这表明去除control divergence可以显著提高性能（约3倍）。
• "ncu --set full divergence" 用这行命令来设置线程control divergence分析。

对于compute bound的kernel，让线程可以做更多工作，可能会更快。

• 性能比较：

• 运行命令：main ~/lecturex ./benchmark
• VecAdd 执行时间：0.245600 ms
• VecAddCoarsened 执行时间：0.015264 ms

• 关键观察：

• VecAddCoarsened启动了一半的线程数量
• 尽管线程数减少，但执行速度显著提高（约16倍）

这里的代码在 https://github.com/cuda-mode/lectures/blob/main/lecture_008/coarsening.cu 。

这也许可以解释Lecture 7中为什么对于Int4 Weight Only量化的高效kernel实现比普通的fp16的Kernel跑得更快。

这张Slides讨论了在GPU编程中的"私有化"（Privatization）技术。要点为：

• 将部分更新应用到数据的私有副本上，然后再写回全局或共享内存。

• 示例：

• 滑动窗口算法（Sliding window algorithm）
• 图示：1 2 [3] [4] [5] 6 7
• 这表明算法在一个局部窗口内进行操作。

• Privatization的优势：

• 更高的占用率（Higher occupancy）
• 更高的计算SM吞吐量（Higher compute SM throughput）