too many resources required for launch

背景

当我们在使用一些低端的jetson设备的时候，比如nano, 偶尔会出现报错，报错显示：

too many resources required for launch

查资料可以发现，一般遇到这种情况就是两个问题，第一个就是寄存器不足，第二个就是共享内存不足。那么问题来了，我实现的一个kernel压根一点没用共享内存，那么就是寄存器不足，可是寄存器不足不是说可以使用显存的吗？我显存几个G为啥还说不够？以下内容仅仅适合高阶cuda玩家，普通小白直接看解决方案就行。

分析

1. 什么时候kernel会由于寄存器不足而溢出到显存里？
   这个是由编译器控制的(也就是编译时期决定的)，它不是由于运行时决定的，这也无法动态决定（换句话说就是不会因为你一个block中的线程太多就是改变由显存来充当寄存器）。当你写的kernel寄存器不够的时候，你的编译器会插入下面类似的指令：

STL  [R0], R1

在这里，R1将会被存在local memory（也就是显存中），R0中存放的是显存的相关地址，这种情况就是spill store 寄存器不足溢出到显存，可以通过-Xptxas=-v查看具体信息如下图，在这里我尝试去写个函数把寄存器用完，但是编译器会优先考虑性能，会自动配置寄存器，把数组转移到了stack中(也就是local memory)。

2. 当寄存器不足的时候，cuda运行时如何决定抛出『too many resources required for launch』？多少寄存器才能满足kernel launch.
   当在编译期间的时候，编译器是不知道kernel会被怎样调用，也不知道会有多少block per grid 和thread per block以及动态部署的shared memory, 所以在编译的时候其实就是在编译单个线程thread. 这个时候每个线程要使用的寄存器数量已经被确定了，已经是一个可执行代码，在运行的时候不会也不可能实时调整了。
   当到运行时的时候，cuda 就会知道每个线程的寄存器，launch的配置参数(block, thread), 机器对寄存器的限制数。当在kernel launch之前会计算：

registers_per_thread*threads_per_block <= max_registers_per_multiprocessor
上面这个解释一下，因为一个block不能拆分，所以个sm上至少也要有一个block, 所以计算方式如上, 此外寄存器个数经常会被调整到2或者4的倍数，所以上述公式左边还要有一个round_up。

如果上面的条件不符合，就会报错，不执行该kernel

3. 既然寄存器有溢出到显存机制，当寄存器不足的时候，为啥不溢出到显存里？

再次来解释这个问题，因为每个线程的寄存器个数在编译的时候就定死了，所以当你每一个block的thread个数特别多的时候，cuda runtime的时候就不会执行，也不会说这个时候挪用显存，因为代码已经写死了，要想该用显存，就必须重新编译代码。

解决

1. 在写kernel的时候限制

__global__ void
__launch_bounds__(maxThreadsPerBlock, minBlocksPerMultiprocessor)
MyKernel(...)

2. 在编译的时候限制寄存器的使用数量，-maxrregcount=18，这样寄存器在编译的时候每个线程就只能使用18个，不足就去用显存，spill store/load.
3. 在config配置的时候，每一个block的线程数减少

参考

链接