专栏名称: 小白学视觉

本公众号主要介绍机器视觉基础知识和新闻，以及在学习机器视觉时遇到的各种纠结和坑的心路历程。

CUDA编程之常用技巧与方法

小白学视觉 · 公众号 · · 2024-03-01 13:40

正文

点击上方 “小白学视觉”，选择加"星标"或“置顶”

重磅干货，第一时间送达

原文链接： https://zhuanlan.zhihu.com/p/584501634

不管你是在学习CUDA，还是在优化算子，掌握一些CUDA编程技巧，能够提升你的工作效率，甚至找到更优解。本文主要是介绍一些常用的技巧/方法，并配上实践code，希望对读者有所帮助。

常用‘print’辅助理解
使用统一内存降低编写难度
性能提升找准瓶颈点
减少数据的拷贝/换页
提升存算重叠度
多用官方标准库
清楚硬件上面的特殊单元全文涉及示例代码（欢迎star，后续不断更新）：

CUDA编程常用方法示例:https://github.com/CalvinXKY/BasicCUDA/tree/master/common_methods

1 常用‘printf’辅助理解

print函数不仅仅是编程中利器，在CUDA编程中我们同样需要常用print来获得过程信息。尤其是在很多debug场景下，我们需要进行数据索引和线程（thread）索引的计算校对，单纯读代码不一定能发现问题，这个时候不妨将这些数据全部打印出来。比如在"CUDA GUIDE" 第一章里面解释了grid、block、thread含义，初次接触只能有个大概的印象，但对于一些关联问题，不一定能够理解到位，比如：

线程数量相同情况下kernel<< >> 和kernel<<<1, N>>> 的配置有什么区别？
kernel里面定义的threadIdx 、blockIdx、blockDim、gridDim如何与线程对应？
一维线程与二维线程的坐标如何计算，以及计算是否正确？

针对问题1，2，我们可以直接在kernel里面加打印，如下：

__global__ void kernel(int mark)
{
    if (blockIdx.x == 0 && threadIdx.x == 0) 
        printf("=== kernel %d run info: gridDim.x: %d, blockDim.x: %d ===\n", \
              mark, gridDim.x, blockDim.x);
    __syncthreads();
    printf("blockIdx.x: %d threadIdx.x: %d\n", blockIdx.x,  threadIdx.x);
}

<示例代码：print_any.cu 编译方式“nvcc -lcuda print_any.cu -o print_any”运行“./print_any”>

通过打印我们可以直接看出<< >>与<<<1, N>>>的差异：

Case0: the diff between <<<1, N>>> with <<>>
 Kernel 0 invocation with N threads (1 blocks, N thread/block) N =8
  === kernel 0 run info: gridDim.x: 1, blockDim.x: 8 ===
    blockIdx.x: 0 threadIdx.x: 0
    blockIdx.x: 0 threadIdx.x: 1
    blockIdx.x: 0 threadIdx.x: 2
    blockIdx.x: 0 threadIdx.x: 3
    blockIdx.x: 0 threadIdx.x: 4
    blockIdx.x: 0 threadIdx.x: 5
    blockIdx.x: 0 threadIdx.x: 6
    blockIdx.x: 0 threadIdx.x: 7
 Kernel 1 invocation with N threads (N blocks, 1 thread/block) N =8
    blockIdx.x: 1 threadIdx.x: 0
    blockIdx.x: 6 threadIdx.x: 0
    blockIdx.x: 2 threadIdx.x: 0
    blockIdx.x: 5 threadIdx.x: 0
    blockIdx.x: 7 threadIdx.x: 0
    blockIdx.x: 3 threadIdx.x: 0
    blockIdx.x: 4 threadIdx.x: 0
  === kernel 1 run info: gridDim.x: 8, blockDim.x: 1 ===
    blockIdx.x: 0 threadIdx.x: 0

对于thread的坐标计算有1D/2D/3D三种情况，比如一个1d的坐标计算如下图所示：

在计算时，可以借助print来打印坐标的关系：

printf("    blockIdx: x=%d y= %d z=%d threadIdx x=%d y=%d z=%d; offset= %d\n",\
    blockIdx.x, blockIdx.y, blockIdx.z,  threadIdx.x, threadIdx.y, threadIdx.z, offset);

其中offset值（索引数据的偏移量）是保证每个线程的索引数据唯一，1D、2D、3D的计算不同。具体我们通过打印可看到其中的索引关系（示例代码：print_any.cu）：

Case1: 1 dimension, grid: 2  block: 2
    blockIdx: x=1 y= 0 z=0 threadIdx x=0 y=0 z=0; offset= 2
    blockIdx: x=1 y= 0 z=0 threadIdx x=1 y=0 z=0; offset= 3
============= The grid shape: gridDim.x: 2 gridDim.y: 1 gridDim.z: 1
============= The block shape: blockDim.x: 2 blockDim.y: 1 blockDim.z: 1
    blockIdx: x=0 y= 0 z=0 threadIdx x=0 y=0 z=0; offset= 0
    blockIdx: x=0 y= 0 z=0 threadIdx x=1 y=0 z=0; offset= 1

Case2: 2 dimension, grid: 2 x 1  block: 2 x 2
    blockIdx: x=1 y= 0 z=0 threadIdx x=0 y=0 z=0; offset= 2
    blockIdx: x=1 y= 0 z=0 threadIdx x=1 y=0 z=0; offset= 3
    blockIdx: x=1 y= 0 z=0 threadIdx x=0 y=1 z=0; offset= 6
    blockIdx: x=1 y= 0 z=0 threadIdx x=1 y=1 z=0; offset= 7
============= The grid shape: gridDim.x: 2 gridDim.y: 1 gridDim.z: 1
============= The block shape: blockDim.x: 2 blockDim.y: 2 blockDim.z: 1
    blockIdx: x=0 y= 0 z=0 threadIdx x=0 y=0 z=0; offset= 0
    blockIdx: x=0 y= 0 z=0 threadIdx x=1 y=0 z=0; offset= 1
    blockIdx: x=0 y= 0 z=0 threadIdx x=0 y=1 z=0; offset= 4
    blockIdx: x=0 y= 0 z=0 threadIdx x=1 y=1 z=0; offset= 5
....

从打印中我们可以知道：

不管是传入1d、2d、3d的数据，在函数里面的 gridDim、blockDim、blockIdx、threadIdx 格式一样，都包含了三个量（x, y, z）。
Dim中没有使用的维度，设置为：1；Idx中没有使用的维度设置为：0。

2. 使用统一内存降低编写难度

在code编写的初期，可以使用统一内存来降低编写与阅读难度。避免了GPUToHost、HostToGPU的操作，从而快速验证算法（kernel）的正确性，比如：

  float *x, *y;
  cudaMallocManaged(&x, N*sizeof(float));
  cudaMallocManaged(&y, N*sizeof(float));
  for (int i = 0; i < N; i++) {
    x[i] = 1.0f;
    y[i] = 2.0f;
  }
  int blockSize = 256;
  int numBlocks = (N + blockSize - 1) / blockSize;
  add<<>>(N, x, y);

代码中在给x y赋值时可直接在主机上进行操作，然后直接把数据代入add kernel中计算。示例代码：um_demo.cu 编译方式“nvcc -lcuda um_demo.cu -o um_demo”运行“./um_demo”。

3 性能提升找准瓶颈点

CUDA程序的性能不仅取决于GPU本身运算速度，也取决于主机机器的运算速度，我们需要借助一些工具来查看性能的瓶颈点，如果卡点在CPU的运算上，则优化CPU代码，如果在GPU运算，就优化GPU代码。常用工具：

nvprof：CUDA API计算时间统计工具
gprof：linux函数耗时统计
nvvp：运算过程可视化工具
events：CUDA API过程耗时统计
nsight/cupit: 工具套件

3.1 nvprof

nvprof 的使用方式非常简洁，只要安装了CUDA，直接在shell里面输入命令即可。如上面提到统一内存的例子中，我们可以通过nvprof查看各个过程的耗时：

$ nvprof um_demo
CUDA API Statistics:

 Time (%)  Total Time (ns)  Num Calls    Avg (ns)       Med (ns)     Min (ns)    Max (ns)     StdDev (ns)           Name
 --------  ---------------  ---------  -------------  -------------  ---------  -----------  -------------  ---------------------
     97.9      367,348,423          2  183,674,211.5  183,674,211.5     13,035  367,335,388  259,736,126.7  cudaMallocManaged
      1.9        6,989,834          1    6,989,834.0    6,989,834.0  6,989,834    6,989,834            0.0  cudaDeviceSynchronize
      0.2          790,933          2      395,466.5      395,466.5    360,910      430,023       48,870.3  cudaFree
      0.0           39,267          1       39,267.0       39,267.0     39,267       39,267            0.0  cudaLaunchKernel

[5/7] Executing 'gpukernsum' stats report

CUDA Kernel Statistics:

 Time (%)  Total Time (ns)  Instances   Avg (ns)     Med (ns)    Min (ns)   Max (ns)   StdDev (ns)             Name
 --------  ---------------  ---------  -----------  -----------  ---------  ---------  -----------  --------------------------
    100.0        6,655,089          1  6,655,089.0  6,655,089.0  6,655,089  6,655,089          0.0  add(int, float *, float *)

[6/7] Executing 'gpumemtimesum' stats report

CUDA Memory Operation Statistics (by time):

 Time (%)  Total Time (ns)  Count  Avg (ns)  Med (ns)  Min (ns)  Max (ns)  StdDev (ns)              Operation
 --------  ---------------  -----  --------  --------  --------  --------  -----------  ---------------------------------
     87.9        1,555,999    376   4,138.3   3,519.0     3,167    42,048      3,178.2  [CUDA Unified Memory memcpy HtoD]
     12.1          214,933     24   8,955.5   3,583.5     2,207    42,176     11,645.0  [CUDA Unified Memory memcpy DtoH]

[7/7] Executing 'gpumemsizesum' stats report

CUDA Memory Operation Statistics (by size):

 Total (MB)  Count  Avg (MB)  Med (MB)  Min (MB)  Max (MB)  StdDev (MB)              Operation
 ----------  -----  --------  --------  --------  --------  -----------  ---------------------------------
      8.389    376     0.022     0.004     0.004     0.971        0.079  [CUDA Unified Memory memcpy HtoD]
      4.194     24     0.175     0.033     0.004     1.044        0.307  [CUDA Unified Memory memcpy DtoH]

3.2 gprof

在优化CPU计算时，充分利用gprof工具。gprof 可以分析出在主机上运算的函数/API的耗时时间。由于gprof是linux自带的工具，使用简单，步骤如下

编译的时候加上 -pg 参数
运行程序
gprof查看运行结果

$ nvcc -pg -lcuda um_demo.cu -o um_demo
$ ./um_demo
$ gprof ./um_demo

这里给了一个参考示例gprof_readme，大家可以运行测试，获得的打印结果：

Flat profile:

Each sample counts as 0.01 seconds.
  %   cumulative   self              self     total
 time   seconds   seconds    calls  ns/call  ns/call  name
 62.50      0.03     0.03  1048576    23.84    23.84  std::fmax(float, float)
 25.00      0.04     0.01                             main
 12.50      0.04     0.01  1048576     4.77     4.77  std::fabs(float)
  0.00      0.04     0.00        2     0.00     0.00  cudaError cudaMallocManaged<float>(float**, unsigned long, unsigned int)
  0.00      0.04     0.00        2     0.00     0.00  dim3::dim3(unsigned int, unsigned int, unsigned int)
  0.00      0.04     0.00        1     0.00     0.00  _GLOBAL__sub_I_main
  0.00      0.04     0.00        1     0.00     0.00  cudaError cudaLaunchKernel(char const*, dim3, dim3, void**, unsigned long, CUstream_st*)
  0.00      0.04     0.00        1     0.00     0.00  __device_stub__Z3addiPfS_(int, float*, float*)
  0.00      0.04     0.00        1     0.00     0.00  add(int, float*, float*)
  0.00      0.04     0.00        1     0.00     0.00  __static_initialization_and_destruction_0(int, int)
  0.00      0.04     0.00        1     0.00     0.00  ____nv_dummy_param_ref(void*)
  0.00      0.04     0.00        1     0.00     0.00  __sti____cudaRegisterAll()
  0.00      0.04     0.00        1     0.00     0.00  __nv_cudaEntityRegisterCallback(void**)
  0.00      0.04     0.00        1     0.00     0.00  __nv_save_fatbinhandle_for_managed_rt(void**)

3.3 nvvp

nvvp是一个可视化UI工具，能够方便的看到算子的各个操作在运算周期内的情况，nvvp相比Nsight使用简单。使用的一般步骤：

通过nvprof 导出记录文件；
启动nvvp加载该文件；

$ nvprof -o output.%p ./um_demo
$ nvvp

启动nvvp界面工具导入output文件即可看到profile情况，e.g.:

更多可以参看nvvp详细教程。

3.4 event

在编写kernel函数时，我们一般需要知道kernel在GPU端的运行时间，通常使用event来统计时间，而不是使用cpu的timer（统计时间不准确！）。event使用示例如下，其中func为待统计的运算函数：

#define TIME_ELAPSE(func, elapsedTime, start, stop)  \
    cudaEventCreate(&start);                         \
    cudaEventCreate(&stop);                          \
    cudaEventRecord(start, 0);                       \
    (func);                                          \
    cudaEventRecord(stop, 0);                        \
    cudaEventSynchronize(stop);                      \
    cudaEventElapsedTime(&elapsedTime, start, stop); \
    cudaEventDestroy(start);                         \
    cudaEventDestroy(stop);

event的使用具体可以参看：定义：memory_opt 30Line ，使用示例zero_copy.cu

3.5 nsight/cupit/nvtx

nsight/cupit/nvtx使用成本相对更高，但功能更强大。

Nsight ：用于GPU资源/数据/性能分析，是一个CUDA编程的综合UI工具，可视化易操作，使用教程;

CUPTI （TheCUDA Profiling Tools Interface）CUDA调优专用API级工具，使用教程;

NVTX （The NVIDIA Tools Extension SDK ）主要是针对C语言的编程API，相对cupit简单点的API, 使用教程；

4 减少数据的拷贝/换页

如果运算时间主要消耗在数据传输/拷贝（通过工具能检查出来），可以通过 共享内存、零拷贝、页锁内存 等降低数据传输成本。

零拷贝 ：当数据保存在主机上，且GPU 只需要使用一次 时，我们借助零拷贝来实现数据传输。可以避免数据从全局显存的进出，从而提供效率。

例如向量加法运算，当使用零拷贝时，数据吞吐能够极大提高。

示例代码：zero_copy.cu 编译：“$ nvcc -lcuda -I../memory_opt/ zero_copy.cu -o zero_run”，运行“./zero_run”，结果：

[Zero Copy Opt Vector Add] - Starting...