CUDA优化总结

总体的CUDA优化分以下几个策略:

  1. 最大化并行执行来获得最大的利用率
  2. 优化内存使用方法来获得最大的内存吞吐量
  3. 优化指令使用方式来获得最大的指令吞吐量

具体再可以从以下几个方面入手:

CUDA配置优化:指的是在CUDA核函数执行之前,其参数的设置,根据显卡的参数,结合核函数的特点,需要注意以下几个参数:

  1. 占用率:每个多核(个人觉得也可以理解成CUDA流处理)上正在执行的线程数,除以多核上允许执行的最大线程数,对于gtx 1060,其多核上允许同时执行的最大线程数是2048,其块内允许最大线程数是1024,每个多核内运行的最大块数是32(这个还不是很确定),在这些参数下,当线程块的个数是32时,那么每个线程块内的线程数是64。高的线程占用率并不总是意味着高性能:存在一个临界点,当线程占用率超过某个值以后并不能够提升性能。但是,低的线程占用率却总是会导致性能降低。
    a 计算占用率:决定占用率的几个因素之一是可获取的寄存器个数。寄存器运行线程中的局部变量快速访问,但是,在多核中的所有线程都需要共享这些寄存器,并且寄存器的数量是有限的,并且寄存器是一次性分配给整个线程块,所以,如果线程块使用了大量的寄存器,那么多核中的线程块的个数就会降低,这就会导致低的占用率(个人理解,应该是降低线程块内每个线程的寄存器使用数量,这样可以增大占用率)。gtx 1060中每个块内的共享寄存器大小是48k,每个多核中的共享寄存器是64k,每个线程可以使用的最大寄存器个数是255个(每个寄存器是32位),每个线程块中的线程数最好是64的倍数,这样才能使性能最优
  2. 核并发执行
  3. 多上下文,在一个CUDA应用,中,在同一个GPU上,应该避免使用多个上下文。为了辅助这种情况,CUDA驱动API提供了一种方法去获取和管理每个GPU上的特殊的上下文,也叫首要上下文(primary context)。如果对于一个线程没有现成的上下文可以用,那么CUDA运行时将隐式的使用这个上下文
  4. 隐藏寄存器独立性,由于访问寄存器并不消耗额外的时钟周期,但是由于寄存器的读写具有独立性和存储的bank冲突,这会导致24个时钟周期的延迟
  5. 线程与线程块选择的启发性:由于有许多因素参与到块的大小选择,因此不可避免的需要一些经验。所以下面的几条应该遵循:
    a 线程块里面的线程数应该是线程warp(一般是32)的倍数
    b 当一个多核上有多个线程块的时候,线程块内的线程数最小应该是64(或者是其倍数)
    c 当你尝试不同的线程块内的线程数的时候,128到256之间的数是一个好的选择
    d 如果寄存器的latency影响性能的时候,应该使用多个线程块(3到4个),而不是使用一个大的线程块。当核函数频繁调用__syncthread()函数的时候更加有用
    需要注意的是,当线程块分配的寄存器个数超过了多核上现有的寄存器数量的时候,核函数将不能够启动,这对共享寄存器也是一样
  6. 共享内存的影响:在前面举的例子中(使用共享存储器和不实用共享存储器计算矩阵相乘),使用共享存储器能够减少对全局内存的访问,提高系统性能。然而他也同样能够对占用率产生限制,在许多情况下,一个核函数需要的共享存储器是与其线程块的块数相关的,比如:
    一个核函数可能需要一个32X32的共享存储器,但是每个线程块能够达到的最大线程数是512,那么,就不太可能为每个块分配一个32X32的共享存储器,在这种情况下,为每个块分配32X16或32X8的线程数,每个线程处理2个或4个共享存储器元素

CUDA指令优化
算术指令:单精度提供了最好的性能,高度推荐使用单精度浮点
单个算术操作的吞吐量在 CUDA_C_Programming_Guide中有详细讲解

  1. 除法取模操作,使用移位操作,替换昂贵的除法和取模操作:整数除法和取模操作特别耗时,并应该尽量避免或用位操作来替换:如果n是2的指数,那么,
    (i/n)的结果和(i>>log2(n))一样,(i%n)和(i&(n-1))结果一样
  2. 倒数平方根:单精度的倒数平方根应该显式的调用rsqrtf(),双精度的应该调用rsqrt(),
  3. 其他算术指令:避免自动的双精度到单精度的转换。在以下两种情况下,编译器将时不时的插入转换指令,增加额外的执行周期:
    a 函数操作char或者short时,其值通常要转换到int
    b 双精度浮点常量(没有后缀的浮点数如:1.0是双精度,1.0f是单精度),作为输入进行单精度浮点计算
    对于b,性能的影响外还有精度的影响,注意到这个不同之处在计算能力为2.x的设备上特别明显
  4. 小分数的指数运算:参见CUDA_C_Best_Practices_Guide,11.1.4
  5. 数学库:当速度要求大于精度要求时,建议使用快速数学库:有两种类型的运行时数学操作,可以根据其名字来区分,前面有下划线的和前面没有下划线的:
    比如__functionName(),functionaNmae().有下划线的函数将会被直接映射到硬件级别去执行,其速度快,但是精度较低。并且,其吞吐量大。相反,没有下划线的速度慢,但是精度高,吞吐量小。
  6. 与编译相关的nvcc的编译选项:
    -use_fast_math会将每一个遇到的没有下划线的函数替换成有下划线的函数
    -ftz 英文愿意:denormalized numbers are flushed to zero
    -prec-div=false(较低精度的除法)
    -prec-sqrt=false(较低精度的平方根)

CUDA内存指令优化
高度推荐:将全局内存的访问降到最低:转而尽可能的使用共享内存
内存指令包括了任何的读写,共享、局部、全局内存的操作。
推荐使用restrict关键字修饰参数,这样可以减少指令的个数,从而提高性能

CUDA控制指令优化
CUDA控制指令:if switch do for while,高度推荐:避免在一个warp(32线程)内执行不同的执行路径
任何流控制指令能够显著的影响指令吞吐量:因为其(不执行同一个执行路径)会导致同一个warp内的线程不收敛
当不执行同一个执行路径时,不同的执行路径江北串行化,因为线程束中的所有线程共享同一个程序计数器,这将增加这个warp内的总的指令数

  1. 分支预测

  2. 循环计数器:有符号vs无符号:在循环中,使用有符号的计数器而不是无符号的计数器

  3. 循环中发散线程的同步:高度推荐:避免在一个发散代码中使用__syncthreads(),即,要保证循环中的所有线程都会执行到这个函数,要是有某个线程执行不到这个函数,那就将出错

©️2020 CSDN 皮肤主题: 编程工作室 设计师:CSDN官方博客 返回首页