为啥启动 Numba cuda 内核最多可使用 640 个线程，但在有大量可用 GPU 内存时却因 641 而失败？

Posted 2023-03-23

【中文标题】为啥启动 Numba cuda 内核最多可使用 640 个线程，但在有大量可用 GPU 内存时却因 641 而失败？

【英文标题】：Why launching a Numba cuda kernel works with up to 640 threads, but fails with 641 when there's plenty of GPU memory free?为什么启动 Numba cuda 内核最多可使用 640 个线程，但在有大量可用 GPU 内存时却因 641 而失败？

【发布时间】：2021-10-10 00:19:15

【问题描述】：

我有一个 Numba cuda 内核，我可以在 RTX 3090 上启动多达 640 个线程和 64 个块。

如果我尝试使用 641 个线程，则会失败：

Traceback (most recent call last):
  File "/home/stark/Work/mmr6/mmr/algos/company_analysis/_analysis_gpu_backup.py", line 905, in <module>
    load()
  File "/home/stark/Work/mmr6/mmr/algos/company_analysis/_analysis_gpu_backup.py", line 803, in load_markets
    run_simulations[algo_configs.BLOCK_COUNT, algo_configs.THREAD_COUNT, stream](
  File "/home/stark/anaconda3/envs/mmr-env/lib/python3.9/site-packages/numba/cuda/compiler.py", line 821, in __call__
    return self.dispatcher.call(args, self.griddim, self.blockdim,
  File "/home/stark/anaconda3/envs/mmr-env/lib/python3.9/site-packages/numba/cuda/compiler.py", line 966, in call
    kernel.launch(args, griddim, blockdim, stream, sharedmem)
  File "/home/stark/anaconda3/envs/mmr-env/lib/python3.9/site-packages/numba/cuda/compiler.py", line 693, in launch
    driver.launch_kernel(cufunc.handle,
  File "/home/stark/anaconda3/envs/mmr-env/lib/python3.9/site-packages/numba/cuda/cudadrv/driver.py", line 2094, in launch_kernel
    driver.cuLaunchKernel(cufunc_handle,
  File "/home/stark/anaconda3/envs/mmr-env/lib/python3.9/site-packages/numba/cuda/cudadrv/driver.py", line 300, in safe_cuda_api_call
    self._check_error(fname, retcode)
  File "/home/stark/anaconda3/envs/mmr-env/lib/python3.9/site-packages/numba/cuda/cudadrv/driver.py", line 335, in _check_error
    raise CudaAPIError(retcode, msg)
numba.cuda.cudadrv.driver.CudaAPIError: [701] Call to cuLaunchKernel results in CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES

但是当我查看 nvidia-smi 时，我发现运行 640 个线程只需要 2.9GB 内存。此 GPU 有 22GB 未使用。

在这种情况下还有什么问题？我在某处读到网格大小、块大小、寄存器使用和共享内存使用是考虑因素。我怎样才能知道我正在使用多少个寄存器和共享内存？

【参考方案1】：

通常是每个线程的寄存器问题 (CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES)。这在 SO cuda 标签上的许多问题中都有涉及，例如 this one。还有很多其他的，例如here。简而言之，每个线程块使用的总寄存器不能超过 GPU 的限制（见下文）。每个 adblock 使用的寄存器总数大约是每个线程的寄存器总数乘以每个块的线程数（可能会四舍五入以获得分配粒度）。

在 numba cuda 中解决此问题的主要方法是在您的 cuda.jit 装饰器中包含一个 maximum register usage parameter：

@cuda.jit( max_registers=40) 

您当然可以将其设置为其他值。一个简单的启发式方法是将每个 SM 的寄存器总数（或每个线程块，如果它较低）（可通过 CUDA deviceQuery 示例代码或the programming guide 的表 15 发现）除以每个块的线程总数希望启动。因此，如果您的 GPU SM 有 64K 寄存器，并且您希望每个块启动 1024 个线程，您将选择每个线程最多 64 个寄存器。这个数字应该适用于 RTX 3090。

【讨论】：

谢谢，就是这样！