CUDA编程——性能优化基本技巧

最新推荐文章于 2025-10-23 12:34:00 发布

原创

最新推荐文章于 2025-10-23 12:34:00 发布 · 1.7k 阅读

24 ·

CC 4.0 BY-SA版权

文章标签：

#性能优化 #矩阵 #CUDA

本文主要介绍下面三种技巧：

使用 __restrict__ 让编译器放心地优化指针访存
想办法让同一个 Warp 中的线程的访存 Pattern 尽可能连续，以利用 Memory coalescing
使用 Shared memory

0. 弄清Kernael函数是Compute-bound 还是 Memory-bound

先摆出一个知识点，一般来说，Compute-bound 的 Kernel 不太常见，常见的 Compute-bound 的 Kernel 可能只有矩阵乘法与卷积核比较大的卷积，大多数都是Memory-bound，所以下面我们主要关注如何优化访存。

在经典的冯诺依曼架构下，ALU (Arithmetic Logic Unit，计算逻辑单元，可以简单理解为加法器、乘法器等) 要从内存中取操作数，进行对应的计算（如乘法），并写回内存。所以，计算速度会受到两个因素的限制：ALU 进行计算的速度，与内存的存取速度。如果一个程序的运行速度瓶颈在于前者，那么称其为 Compute-bound 的；如果瓶颈在于后者，那么称其为 Memory-bound 的。

由于 CPU 中运算单元较少，且 CPU 具有多级缓存，所以空间连续性、时间连续性较好的程序在 CPU 上一般是 Compute-bound 的。而 GPU 则恰恰相反：GPU 的核心的规模一般很大，比如 RTX 4090 可以在一秒内做 82.58T 次 float16 运算（暂不考虑 Tensor core），但其内存带宽只有 1TB/s，每秒只能传输 0.5T 个 float16。这便导致 GPU 上的操作更可能会受到内存带宽的限制，成为 Memory-bound。

如何估测一个 CUDA Kernel 是 Compute-bound 还是 Memory-bound 呢？我们可以计算它的 “算存比”，也即，

算存比 = 计算次数/访存次数

并将其与 GPU 的每秒能做的运算次数/每秒能做的访存次数做比较。

例子：

__global__ void pointwise_add_kernel(int* C, const int* A, const int* B, int n) {
    for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < n; i += blockDim.x * gridDim.x)
        C[i] = A[i] + B[i];
}

对于上面的 pointwise_add_kernel，其需要访问 3N 次内存（读取A[i]、读取B[i]、写入C[i]），同时做 N次加法，所以其存算比为 N/3N=1/3

对于RTX 4090 可以在一秒内做 82.58T 次 float16 运算（暂不考虑 Tensor core），但其内存带宽只有 1TB/s，每秒只能传输 0.5T 个 float16。起存算比为 82.58/0.5=165.16

因此这个 pointwise_add_kernel为 Memory-bound。