CUDA C++ Best Practices Guide 概要-CSDN博客

本文链接：https://blog.csdn.net/m0_61864577/article/details/142596250

CUDA C++ Best Practices Guide 概要

一.整体优化策略
二.优先级
三.其它知识点

CUDA C++ Best Practices Guide 概要

一.整体优化策略

1.最大化并行度

算法的并行化
Kernel的并行化
多Stream并行
HOST和DEVICE并行
计算指令与IO指令的并行
指令并行

2.最大化内存带宽

减少HOST跟DEVICE之间的数据传输
减少Kernel对DRAM的访问,有时甚至采用重计算
合并访问
减少共享内存的Bank冲突

3.最大化指令吞吐

使用高吞吐的指令
在精度和性能之前平衡,在不影响最终结果的情况下,选择性能更高的指令
优先使用intrinsic指令或单精度指令
注意控制流指令

二.优先级

1.高优先级

为了最大限度地提高开发人员的工作效率，分析应用程序以确定热点和瓶颈
首先应该关注如何并行化串行代码
最大限度地减少主机和设备之间的数据传输，即使某些Kernel相比CPU没有性能增益
尽可能对DRAM合并访问
最小化DRAM的使用,尽可能使用共享内存访问
避免在同一个线程束中使用不同的执行路径
在衡量性能和优化效益时，使用计算的有效带宽作为衡量指标。

2.中优先级

使用共享内存以避免从全局内存进行冗余传输
为了隐藏由寄存器依赖性引起的延迟，每个SM保持足够数量的活动线程(即占用率)
每个块的线程数应该是32个线程的倍数，因为这可以提供最佳的计算效率并促进合并
当速度胜过精度时，使用快速数学库
尽可能选择更快、更专业的数学函数，而不是更慢、更通用的函数

3.中低优先级

使用有符号整数而不是无符号整数作为循环计数器

4.低优先级

使用移位操作来避免昂贵的除法和取模计算
避免编译器将double自动转换为float
使编译器更容易使用分支预测代替循环或控制语句
在CUDA工具包2.2及更高版本的集成gpu上使用零拷贝操作

三.其它知识点

建议首次顺序阅读
按评估(理论性能)、并行化改造、优化、部署的顺序,多次迭代
评估性能瓶颈,及优化后的加速比
并行化:使用cublas cudnn等库;算法的并行化改造;cuda c实现Kernel
优化是个迭代的过程,了解需求很重要,不必掌握所有的优化技巧后再开始
部署:一次迭代后就可部署,可尽早受益,也能降低风险
本文中的建议根据影响范围给出优先级,在做低优先的优化之前,确保已完成高优先级的优化
现在处理器越来越依靠多核来提高性能(多个计算单元),这要求代码能并行
先定位热点,再评估此处是否能并行,并了解当前及未来的计算规模
HOST支持只能有限数量的限程,然而GPU一个warp就32个线程,一个80个SM的GPU,甚至能达到160000个线程
GPU的线程切换开销远于CPU。CPU为优化少量线程下的延迟而设计,GPU为大量轻量级并发线程而设计,吞吐优先
CPU和GPU各有所长,CPU适合串行的任务而GPU适合并行的任务,在它们组成的混合架构中,需发挥各自的专长,以使整体性能最优
CUDA的轻量级线程模型使用我们可以在GPU上运行大量的线程,运行大量的线程也是GPU能获取高性能的前提
HOST和GPU都用于拥有自己的内存,在运行GPU程序之前需将数据从HOST拷贝到Device,应减少它们之间的数据传输
需要执行的计算量应能匹配上数据搬运的成本,如果搬移的成本大于计算的耗时,可能没有收益
比如矩阵的加法,计算跟跟传输的比值为(1:3),这种情况下收益并不明显,如果是乘加则为1:N,收益会随规模增加而增加大
数据尽可能的驻留在设备内存,减少HOST跟DEVICE之间的数据传输,即使某些计算GPU比HOST计算慢
合理的访问模式,可以让访存合并和提高Cache命中率,避免随机访问.完全随机的情况下,GPU能通过隐藏延迟,比CPU更有优势
确保性能分析时的用数据是真实数据,来自真实场景
使用gprof等工具得到热点函数,我们先关注TOP1的函数,将TOP2及之后的热点函数放在下一轮迭代
GPU的加速比取决于代码的并行度,如果并行度不高,建议放在HOST上执行,除非会导致大量传输成本
Amdahl’s Law即Strong scaling(S=1/((1-P)+(P/N))) P表示可并行化的耗时占比,N是处理器数据量
描述了在固定规模下,加速比跟处理器数量的关系。从公式上看来,优化P是重点
Weak Scaling and Gustafson’s Law：S = N + (1 − P)(1 − N) 描述了随着处理器增加到固定的问题,问题总量变大
S是此时的加速比，如果并行部分的占比是100%,则加速比为N;如果并行部分的占比越低,加速比越低
可能通过并行库(cublas,cudnn),OpenACC实现加速,也可以通过CUDA C++实现,如果没有明显的热点函数,建议对工程进行重构,这对其它平台也有益
所有的优化,结果正确是前提
并行执行跟串行流程不一样,特别是浮点运算,理论上数值不能保存完全一致,需要定义一套精度评价机制(比如判断误差的阈值),每次优化后,跟目标进行比较
除了比较精度以外,也可以建立性能比较标准,用来评价不同优化方案对性能的影响
采用多个__device__函数,而不是一个揉在一起的__global__函数,方便对每个功能进行单元测试。比如,写一个函数用来处理寻址逻辑,对其进行单独调试,可以简化调试工作
如果一段设备代码最终没有写DRAM,CUDA编译器会将其优化掉
可以用__host__ device 修饰函数,该函数在host和device都可以运行,通过比较二者的差异,方便提前发现问题。该方法还能减少代码重复(特别是大量需要在host和device上都运行的函数)
采用CUDA-GDB调试
基于相同精度类型进行比较,允许存在误差,不能期望浮点值完全一样;浮点运算没有结合性(A+B)+C不等于A+(B+C);并不是所有的CUDA指令都遵循IEEE 754标准;x86处理可以使用80bit double扩展,跟CUDA纯64bit有差异
在用CPU统计Kernel耗时时别忘了cudaDeviceSynchronize
GPU会交错调度不同stream的任务,统计的时间也会包含在内
sync操作会导致GPU处理管线停顿,会影响性能,谨慎使用
几乎所有的代码更改都应该在它们如何影响带宽的背景下进行。带宽可能会受到存储数据的内存、数据的布局方式和访问数据的顺序以及其他因素的显著影响
根据内存理论带宽和实测带宽的差异,设计优化方案,逐步优化
理论带宽可能通过GPU数据手册或NCU中的DRAM的几峰值性能获取.速度等于主频位宽/82(双边沿采样).有时1G=1024^3 有些是10^9.请保持计算规则一致
开启ECC后GDDR有效容量减少6.25%,有效带宽减少20%。HBM2用于专用的ECC资源,不影响性能
实际带宽可以通过实际传输的数据量除以时间.或通过NCU获取.最小的内存事务为32字节,有可能比kernel实际需要的多
可以比较硬件峰值带宽、实际使用的带宽以及有效带宽来定位瓶颈
每次host到device的传输都存在开销.因此,建议将多个小块传输合并成一个大块进行传输(尽管拼接存在开销)
host内存建议pin住(用cudaHostAlloc分配,或者malloc之后调用cudaHostRegister)。但过度使用pin内存会影响系统的整体性能
通过cudaMemcpyAsync让host计算和拷贝并行起来,也能让Kernel计算与拷贝并行起来
通过异步拷贝,可以将一块计算拆成多个小块.整体耗时从(拷贝耗时+计算耗时)优化到(计算耗时+拷贝耗时除以拆分次数)
在集成GPU上使用零拷贝cudaHostAlloc+cudaHostGetDevicePointer
从compute capability 2.0开始,支持UVA,这使得cudaHostAlloc分配的地址可同时指向HOST和DEVICE,不需要再调用cudaHostGetDevicePointer,但malloc的内存仍需要,UVA也是P2P的前提条件
合并访问以warp为单位,一个warp需要的事务数等于这32个线程需要的字节总数(按32对齐以膈)除以32
在计算能力为6.0或更高版本的设备上，L1缓存是默认开启，但无论全局加载是否缓存在L1中，数据访问单元都是32字节。
ECC打开时,合并访问更加重要,分散的访问会增加ECC内存传输开销，特别是在将数据写入全局时
warp内线程访问地址只要落在同一个32字节的对齐区间，整个32字节的数据就会被读取(不论顺序,也不管是否全部访问)
如果一个warp的访问地址不是按32字节对齐,则会多加载一个32字节.即4(有效的sector数)/5(实际加载的sector)。如果数据已经在cache中,差距可能不会这么明显
如果按步长为2进行访问,一个warp的访问效率为50%.最差的情况是每个线程访问一个sector,4/32=1/8
将L2 accessPolicyWindow中的 num_bytes固定为X字节并调整 hitRatio,使总持久性数据中的随机X字节驻留在L2预留缓存部分中(不要超过)。其余部分将使用streaming属性访问,有助于减少缓存抖动
共享内存每个Bank每个cycle具有32位的带宽,连续的32位字被分配给连续的bank.warp大小为32个线程，bank数量也为32
当一个块中的多个线程使用相同的DRAM数据时,可以用共享内存复用数据,减少对DRAM的访问。此外还可用来避免非合并内存访问
从DRAM中以合并模式加载和存储数据，然后在共享内存中对其进行重新排序。除了BANK冲突外，共享内存中的线程束不会对非顺序或非对齐访问造成任何损失
__pipeline_memcpy_async+__pipeline_commit()+__pipeline_wait_prior(0);可以绕过寄存器直接从dram加载数据到共享内存
当同步拷贝的数量是4的倍数时,编译器会进行优化,性能最佳。但异步拷贝没有这个约束,都很好。总体上使用异步方式拷贝8或16字节时,性能最佳
Local Memory用于存放自动变量(局部变量)，当编译器发现寄存器空间不足时，会将大型结构体、数组或可能动态索引的数组放入Local Memory
查看ptx是是否有.local声明或是否有ld.local,st.local指令,可以知道是否使用了local memory.也可以让nvcc在编译时加–ptxas-options=-v进行报告
Constant Memory会被Constant Cache缓存,warp内的线程对不同地址的访问是串行的,如果一个warp中的所有线程都访问Constant Memory同一个位置，那么速度跟寄存器访问一样快
编译器和硬件线程调度器会尽可能优化指令调度以避免寄存器银行冲突，但应用程序无法直接控制这些冲突，因此将数据打包成浮点数4（float4）或整型4（int4）等向量数据类型并没有寄存器层面的实际必要性。
可以使用 -maxrregcount=N 编译器命令行选项或 launch bounds限定符来控制每个线程分配的最大寄存器数
使用 cudaMalloc() 和 cudaFree() 进行设备内存分配和释放是耗时操作，建议使用 cudaMallocAsync() 和 cudaFreeAsync() 这些有序池分配器来更高效地管理设备内存。
一些新版本的Linux发行版默认启用自动NUMA平衡，但在某些情况下，这可能会降低运行在NVIDIA GPU上的应用性能，因此用户应手动调整应用的NUMA特性以实现最佳性能。如 numactl --membind=0,8
尽量让设备上的SM保持忙碌，SM之间不均衡会导致性能不佳，因此需优化线程和块的使用，提升硬件利用率，避免阻碍任务自由分配的做法；通过让多个独立Kernel并发执行，可进一步提高硬件利用率
线程指令是顺序执行的，一个warp暂停或阻塞时，通过执行其他warps是隐藏延迟、保持硬件忙碌的唯一方式，活跃warps的数量被称为占用率（occupancy），是评估硬件利用效率的重要指标。
更occupancy并不总是等同于更高性能，但低occupancy会影响隐藏内存延迟的能力，从而导致性能下降。
使用 launch_bounds(maxThreadsPerBlock,minBlocksPerMultiprocessor) 可以帮助编译器优化内核的资源使用，提高性能,有助于避免因线程块过大而引起的寄存器溢出或共享内存不足问题
在计算占用率时，线程使用的寄存器数量是关键因素，但需考虑寄存器分配的粒度对线程块和占用率的实际影响
寄存器依赖在指令使用先前指令写入寄存器结果时产生，尽管算术指令有约4个周期的延迟，但其他warp线程的执行可以完全隐藏这种延迟
每个块的线程数应该是warpsize的倍数，以避免浪费计算资源和更好的访存合并。
每个块至少应该使用64个线程，并且只有在每个SM有多个并发块的情况下才应该使用
每个块128到256个线程之间是一个很好的初始范围，可以用于不同块大小的实验
如果延迟会影响性能，请在每个sm上使用几个较小的线程块，而不是一个大的线程块。这对经常调用 __syncthreads 的Kernel特别有用。
通过调整执行配置中的动态分配共享内存参数来实验观察占用率对性能的影响，这样能在不修改内核的情况下降低占用率并评估性能变化。
同一个GPU上，可以同时分配多个上下文及其相关资源（如全局内存），但在任意时刻只有一个上下文可以执行任务；多个上下文会在同一GPU上以时间切片的方式分时运行。
创建额外的上下文会导致每个上下文的数据占用额外的内存，并增加上下文切换的时间开销。
此外，当来自多个上下文的任务本可以并发执行时，上下文切换的需求会降低使用效率。简单来说，多个上下文的存在可能会导致资源使用不充分和性能下降。
为了确保符合 IEEE-754 标准，单精度时应显式调用 rsqrtf()，双精度时应调用 rsqrt() 来计算倒数平方根。
编译器只有在不违反 IEEE-754 语义的情况下，才会自动将 1.0f/sqrtf(x) 优化为 rsqrtf()。因此，显式调用可以避免不确定的优化行为，确保精度和性能。
在 char 或 short 上操作的函数，编译器会插入转换指令将操作数转换为 int,将双精度的常量(不带f后缀)用于单精度复点指令。
对于某些分数指数，通过使用平方根、立方根及其倒数进行幂运算，相较于直接使用 pow() 函数，计算速度可以显著提高。且能获得更准确的结果
带下划线前缀的函数（如 __functionName()) 直接映射到硬件级别，速度更快但精度较低；
没有下划线的函数(如 functionName())速度较慢但精度更高，尤其当需要对参数大小进行调整时，后者会显著变慢并可能使用高延迟的本地内存，从而影响性能。
-use_fast_math 编译选项将所有 functionName() 调用转换为 __functionName() 调用，降低精度并禁用单精度非规格化支持，提高性能但可能减少数值准确性
建议仅在可接受的情况下选择性使用此优化，该选项仅影响单精度浮点计算。
对于较小的整数幂（如 (x^2) 或 (x^3)），显式乘法速度通常比使用 pow() 更快，尤其当需要多次计算相同基数的幂时，显式乘法有助于编译器进行公共子表达式消除优化，从而显著提升性能。
对于以2或10为底的幂运算，使用 exp2()、exp2f()、exp10() 或 exp10f() 函数更好，它们性能与 exp()、expf() 类似，可以比 pow()、powf() 快十倍，因为后者由于处理多种特殊情况导致寄存器压力和指令计数较高。
对于以 (1/3) 为指数的幂运算，使用 cbrt() 或 cbrtf() 函数，相比通用的 pow() 或 powf()，它们速度快得多；同样地，对于以 (-1/3) 为指数的幂运算，使用 rcbrt() 或 rcbrtf()。
将(sin(π*)) 替换为 sinpi(expr)可以提高准确性和性能，因为 sinpi() 简化了参数缩减并使用了无限精确的数学 (π) 运算；类似地，对于 cos() 和 sincos() 替换为 cospi() 和 sincospi() 也有相同优势。
默认情况下， nvcc 编译器生成符合IEEE标准的代码，但它也提供了一些选项来生成不太准确但更快的代码：
-ftz=true （反规格化的数字被刷新为零） -prec-div=false （不太精确的除法） -prec-sqrt=false （不太精确的平方根）
流程控制指令（如 if、switch、do、for、while）可能导致同一个warp中的线程路径分歧，从而需要分别执行不同路径，增加指令总数，影响指令吞吐量。
建议用线程ID做判断条件,比如threadIdx / WSIZE.按warp对齐,可以保证不存在warp分歧
对于只包含少量指令的分支，warp分歧通常仅导致较小的性能损失，因为编译器可能使用谓词化技术，通过为每个线程设定条件码或谓词来避免实际分支，使得不满足条件的线程不写入结果，也不计算地址或读取操作数。
从Volta架构开始，独立线程调度允许warp在数据依赖的条件块之外保持分歧，可以使用显式的__syncwarp()确保warp在随后指令执行前重新汇聚。
在使用分支预测时，所有指令都会被调度执行，但只有谓词为真的指令会实际执行和写入结果，谓词为假的指令则不会写入结果或读取操作数。
编译器仅在由分支条件控制的指令数量小于或等于某个阈值时，才将分支指令替换为谓词指令
在完成应用程序部分组件的GPU加速后，可以将结果与原始预期进行比较，尽快将部分并行化实现投入生产，以便尽早获利并通过渐进式改进降低风险。
每代NVIDIA处理器都会添加CUDA新特性，程序员需了解架构特性，包括计算能力和CUDA运行时及驱动API的版本号。
计算能力描述硬件特性，反映设备支持的指令集及其他规格，如每个块的最大线程数和每个多处理器的寄存器数量；更高版本的计算能力是较低版本的超集，因此向后兼容。
某些硬件特性不在计算能力描述中，如异步数据传输与内核执行重叠的能力可通过调用cudaGetDeviceProperties()来确定
例如asyncEngineCount字段显示是否可以重叠执行及并行传输次数，而canMapHostMemory字段指示是否支持零拷贝数据传输。
为了针对特定版本的NVIDIA硬件和CUDA软件，使用nvcc的-arch、-code和-gencode选项，例如，使用warp shuffle操作的代码必须使用-arch=sm_30（或更高计算能力）进行编译。
与CUDA DriverAPI相比，CUDA Runtime通过提供隐式初始化、上下文管理和设备代码模块管理大大简化了设备管理，且nvcc生成的C++主机代码依赖CUDA Runtime
因此链接到此代码的应用程序会依赖CUDA Runtime，使用cuBLAS、cuFFT等CUDA工具包库的代码也同样依赖CUDA Runtime。
CUDA Runtime在内核启动前负责内核加载、设置内核参数和启动配置，并执行隐式驱动版本检查、代码初始化、CUDA上下文管理、CUDA模块管理（从cubin到函数映射）、内核配置及参数传递。
CUDA编译器（nvcc）通过分离和引导编译处理CUDA和非CUDA代码，作为CUDA编译工具链的一部分，CUDA运行时API为开发者提供简化设备管理和内核执行的高级C++接口，而CUDA驱动API提供面向NVIDIA硬件的低级编程接口。
从CUDA 11开始，工具包版本基于行业标准的语义版本化方案：.X.Y.Z
X代表主要版本-API已更改且二进制兼容性已被破坏。
Y代表次要版本-新API的引入，旧API的弃用，以及源代码兼容性可能会被破坏，但二进制兼容性仍保持不变
Z代表发布/补丁版本-新的更新和补丁将增加此版本
源代码兼容性指的是库提供的一组保证，当使用特定版本库构建的应用程序（使用SDK）在安装新版本SDK后仍能继续构建和运行。
然而，CUDA驱动和运行时在不同SDK版本之间未保证源代码兼容性，可能导致API弃用或移除，需对旧版本编译成功的应用进行修改才能在新版本下编译。
即便如此，CUDA保证CUDA驱动API接口的二进制兼容性，因此旧版工具包编译的应用程序二进制仍受支持。
二进制兼容性指的是应用程序在动态链接到库的不同版本时仍能正常运行的一组保证。
CUDA驱动API具有版本化的C风格ABI，确保即使应用程序在旧驱动上运行，仍能在现代驱动上正常工作，但源代码可能需要修改以使用新特性。
CUDA驱动API是二进制兼容的，但不是源代码兼容的，因为每个CUDA工具包版本需要最低NVIDIA驱动版本支持。
cubin在硬件架构的次要版本更新中保持兼容，但在主要版本更新或逆向次要版本更新时不保证兼容。对于未来的GPU架构，为了便携性，应用程序应加载PTX代码，以便由NVIDIA驱动进行及时编译。
PTX定义了一种用于并行线程执行的虚拟机和指令集架构。PTX程序在加载时通过CUDA驱动的JIT编译器转译为目标硬件指令集。为了在旧版本驱动上支持运行，代码需使用静态ptxjitcompiler库或NVRTC生成特定架构代码。
静态链接推荐：尽量使用静态链接CUDA运行时库，以减少对外部依赖的需求，从而避免库之间的接口或者命名空间冲突。
语义版本控制：维护和更新你的库接口时要遵循语义版本控制，以确保兼容性不被破坏。重要的接口变化（比如弃用或修改API）需要增加主版本号，而添加新功能则增加次版本号。
兼容性措施：通过条件性使用新特性和提供后备方案，保证代码在较旧的驱动程序上仍能工作。
避免暴露动态结构：不要公开可能会改变的ABI结构，使用嵌入大小的指针结构可以更好地保证稳定性。
动态链接注意事项：当使用动态链接库时，确保使用的库版本不低于所需的最低版本，以防止因版本不一致导致的问题。
在多GPU系统中，建议在同一应用中使用相同类型的GPU而不是混合不同代，以确保一致性，可用cudaChooseDevice()函数选择最符合需求的设备。
用cudaGetDeviceProperties检测硬件和软件配置
由于某些CUDA API调用和所有内核启动与主机代码是异步的，错误可能在主机与设备同步时才报告，因此在调用cudaMemcpy()或cudaDeviceSynchronize()时要特别注意。
始终检查所有CUDA API函数的错误返回值，包括不容易出错的函数，以便尽快检测和处理错误，同时在内核启动后立即调用cudaGetLastError()检查内核启动中可能发生的错误，以确保程序的完整性和数据的正确性。
每代支持CUDA的设备都有一个计算能力版本，指示设备所支持的功能集，编译时指定目标GPU的计算能力，有助于应用程序内核实现最佳性能并充分利用GPU特性。
构建多个计算能力版本时，CUDA驱动在运行时根据设备的计算能力选择合适的二进制执行；如果缺乏对应本地二进制但有中间的PTX代码，驱动将即时编译PTX为本地二进制，而如果PTX也不可用，内核启动将失败。
CUDA运行时静态链接时，同一应用进程内可共存多个版本的CUDA运行时，例如应用程序和其插件分别静态链接不同版本的运行时是可以接受的，只要安装的NVIDIA驱动程序能够支持这两者即可。
默认是静态链接,编译时加上–cudart=shared改为动态链接
nvidia-smi可以查询ECC,GPU利用率,活动的进程 ,时钟，温度，风扇,电源,各种ID。并修改ECC模式,复位ECC,复位GPU，修改计算模式,持续模式
nvidia-smi和NVML的ID是按PCIE枚举的顺序,且运行时和CUDA_VISIBLE_DEVICES采用的是逻辑ID(排序后的ID)。通过通过UUID或总线号进行关联