作者 | kaiyuan  编辑 | 汽车人

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

点击下方卡片，关注“自动驾驶之心”公众号

ADAS巨卷干货，即可获取

点击进入→自动驾驶之心【模型部署】技术交流群

后台回复【CUDA】获取CUDA实战书籍！

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

1. 常用‘print’辅助理解

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

3. 性能提升找准瓶颈点

4. 减少数据的拷贝/换页

5. 提升存算重叠度

6. 多用官方标准库

7. 清楚硬件上面的特殊单元 全文涉及示例代码（欢迎star，后续不断更新）：

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

1 常用‘printf’辅助理解

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

• 线程数量相同情况下kernel<<<N, 1>>> 和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”>

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

Case0: the diff between <<<1, N>>> with <<<N, 1>>>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: 0blockIdx.x: 0 threadIdx.x: 1blockIdx.x: 0 threadIdx.x: 2blockIdx.x: 0 threadIdx.x: 3blockIdx.x: 0 threadIdx.x: 4blockIdx.x: 0 threadIdx.x: 5blockIdx.x: 0 threadIdx.x: 6blockIdx.x: 0 threadIdx.x: 7Kernel 1 invocation with N threads (N blocks, 1 thread/block) N =8blockIdx.x: 1 threadIdx.x: 0blockIdx.x: 6 threadIdx.x: 0blockIdx.x: 2 threadIdx.x: 0blockIdx.x: 5 threadIdx.x: 0blockIdx.x: 7 threadIdx.x: 0blockIdx.x: 3 threadIdx.x: 0blockIdx.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: 2blockIdx: x=1 y= 0 z=0 threadIdx x=0 y=0 z=0; offset= 2blockIdx: 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: 1blockIdx: x=0 y= 0 z=0 threadIdx x=0 y=0 z=0; offset= 0blockIdx: x=0 y= 0 z=0 threadIdx x=1 y=0 z=0; offset= 1Case2: 2 dimension, grid: 2 x 1  block: 2 x 2blockIdx: x=1 y= 0 z=0 threadIdx x=0 y=0 z=0; offset= 2blockIdx: x=1 y= 0 z=0 threadIdx x=1 y=0 z=0; offset= 3blockIdx: x=1 y= 0 z=0 threadIdx x=0 y=1 z=0; offset= 6blockIdx: 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: 1blockIdx: x=0 y= 0 z=0 threadIdx x=0 y=0 z=0; offset= 0blockIdx: x=0 y= 0 z=0 threadIdx x=1 y=0 z=0; offset= 1blockIdx: x=0 y= 0 z=0 threadIdx x=0 y=1 z=0; offset= 4blockIdx: 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<<<numBlocks, blockSize>>>(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  cudaMallocManaged1.9        6,989,834          1    6,989,834.0    6,989,834.0  6,989,834    6,989,834            0.0  cudaDeviceSynchronize0.2          790,933          2      395,466.5      395,466.5    360,910      430,023       48,870.3  cudaFree0.0           39,267          1       39,267.0       39,267.0     39,267       39,267            0.0  cudaLaunchKernel[5/7] Executing 'gpukernsum' stats reportCUDA 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 reportCUDA 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 reportCUDA 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自带的工具，使用简单，步骤如下

1. 编译的时候加上 -pg 参数

2. 运行程序

3. 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     totaltime   seconds   seconds    calls  ns/call  ns/call  name62.50      0.03     0.03  1048576    23.84    23.84  std::fmax(float, float)25.00      0.04     0.01                             main12.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_main0.00      0.04     0.00        1     0.00     0.00  cudaError cudaLaunchKernel<char>(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使用简单。使用的一般步骤：

1. 通过nvprof 导出记录文件；

2. 启动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...
>. Data tranfer via global memory.  VectorAdd throughput: 1.271375 GB/s
>. Data tranfer via  zero copy.     VectorAdd throughput: 714.285706 GB/s

共享内存：用户可直接使用的片上存储。对于需要反复使用的数据，将数据放到共享内存中，因为共享内存的速度与L1 cache相同，相比全局显存效率更高。

求和运算示例代码：shared_mem.cu 编译：“$ nvcc -lcuda -I../memory_opt/ shared_mem.cu -o smem_run”，运行“./smem_run” 结果：

[Shared Memory Application: Array Sum.] - Starting...
Sum array with shared memory.       Elapsed time: 0.007025 ms
Sum array without shared memory.    Elapsed time: 0.011110 ms

5 提升存算重叠度

6 多用官方标准库（cuDLA/cudnn/cuFFT/cublas）

cudaMath：写算子前先看这个库里面有没有现成函数。<使用教程>

cuDNN: 深度学习相关的卷积/池化等运算优化，直接提速网络。<介绍与使用>。

cuFFT: 快速傅里叶变化，有FFT/FFTW。<使用教程>。

cuBLAS: 线性代数/矩阵运算 ，算子种类丰富。<使用教程>。

cuSPARS：稀疏运算API，涉及稀疏向量/矩阵/混合运算，稀疏操作优先考虑该库。<使用教程>

7 了解硬件上面的特殊单元

往期回顾

史上最全综述 | 3D目标检测算法汇总！（单目/双目/LiDAR/多模态/时序/半弱自监督）

【自动驾驶之心】全栈技术交流群

自动驾驶之心是首个自动驾驶开发者社区，聚焦目标检测、语义分割、全景分割、实例分割、关键点检测、车道线、目标跟踪、3D目标检测、BEV感知、多传感器融合、SLAM、光流估计、深度估计、轨迹预测、高精地图、NeRF、规划控制、模型部署落地、自动驾驶仿真测试、硬件配置、AI求职交流等方向；

添加汽车人助理微信邀请入群

备注：学校/公司+方向+昵称