作者 | kaiyuan 编辑 | 汽车人
原文链接：https://zhuanlan.zhihu.com/p/584501634
点击下方卡片，关注“自动驾驶之心”公众号
ADAS巨卷干货，即可获取
点击进入→自动驾驶之心【模型部署】技术交流群
后台回复【CUDA】获取CUDA实战书籍！
不管你是在学习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<<<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: 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<<<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 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>(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...
>. 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 ms5 提升存算重叠度
6 多用官方标准库（cuDLA/cudnn/cuFFT/cublas）
cudaMath：写算子前先看这个库里面有没有现成函数。<使用教程>
cuDNN: 深度学习相关的卷积/池化等运算优化，直接提速网络。<介绍与使用>。
cuFFT: 快速傅里叶变化，有FFT/FFTW。<使用教程>。
cuBLAS: 线性代数/矩阵运算 ，算子种类丰富。<使用教程>。
cuSPARS：稀疏运算API，涉及稀疏向量/矩阵/混合运算，稀疏操作优先考虑该库。<使用教程>
7 了解硬件上面的特殊单元
往期回顾
史上最全综述 | 3D目标检测算法汇总！（单目/双目/LiDAR/多模态/时序/半弱自监督）
【自动驾驶之心】全栈技术交流群
自动驾驶之心是首个自动驾驶开发者社区，聚焦目标检测、语义分割、全景分割、实例分割、关键点检测、车道线、目标跟踪、3D目标检测、BEV感知、多传感器融合、SLAM、光流估计、深度估计、轨迹预测、高精地图、NeRF、规划控制、模型部署落地、自动驾驶仿真测试、硬件配置、AI求职交流等方向；
添加汽车人助理微信邀请入群
备注：学校/公司+方向+昵称
875
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
