一个CUDA程序如果使用的寄存器数量过多,会导致在SM上同时驻留的线程和block数量减少,继而导致程序性能不足。

__launch_bounds__maxrregcount都可以用来限制cuda程序的寄存器数量,但是两者是不同的机制。

__launch_bounds__

__global__ void

__launch_bounds__(maxThreadsPerBlock, minBlocksPerMultiprocessor)
MyKernel(...)
{
 ...
}
  • maxThreadsPerBlock 指定了每个block最多可以包含多少个线程。它对应了PTX中的.maxntid指令。
  • minBlocksPerMultiprocessor 是可选的,它指定了每个SM上最少需要驻留多少block,它对应了PTX的.minnctapersm指令。

指定这两个参数以后,compiler会推出每个线程的寄存器使用上限值L。

compiler将通过下面的方法来优化寄存器使用:

  • 如果compiler第一次得出的寄存器使用量高于L,那么它将通过spill寄存器内容到local memory增加指令数量(通过更多的寄存器内容交换)来减少寄存器使用量,直到小于等于L。
  • 如果第一次得出的寄存器量低于L,
    • 如果指定了maxThreadsPerBlock但是没有指定minBlocksPerMultiprocessor,compiler则会尽量减少寄存器以便满足这个max。
  • 如果两个都指定了,则compiler会尽量增加寄存器数量(直到L)来减少指令数量隐藏指令延迟。

maxrregcount

maxrregcount是一个编译参数,简单暴力的将所有的kernel函数寄存器(除了使用launch_bound限制的kernel函数以外)使用的最大量限制为对应的数值。

它只关心寄存器的使用数量,我感觉它的优化策略,跟launch_bound第一条差不多。

  • 如果原来default编译出来寄存器数量超过了maxrregcount,就尽量减少寄存器数量。
  • 但是如果原来的寄存器数量很少,而maxrregcount限制值较大,寄存器数量可能减少、不变或者增加

Reference

Limiting register usage in CUDA: launch_bounds vs maxrregcount

CUDA Programming Guide

CUDA 编程之 launch bounds


文章版权归 FindHao 所有丨本站默认采用CC-BY-NC-SA 4.0协议进行授权|
转载必须包含本声明,并以超链接形式注明作者 FindHao 和本文原始地址:
https://findhao.net/easycoding/2468.html

Comments