为什么启动一个Numba Cuda内核可以处理多达640个线程，但在有足够的GPU内存可用时却会在641个线程上失败？

问题描述

我有一个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时，我发现它只需要2.9 GB的内存就可以运行640个线程。此GPU有22 GB未使用的空间。

在这种情况下还会有什么问题？我在某处读到网格大小、挡路大小、寄存器使用率和共享内存使用率都是需要考虑的因素。如何才能知道我正在使用多少寄存器和共享内存？

---

解决方案

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

在Numba Cuda中解决此问题的主要方法是在cuda.jit装饰符中包含maximum register usage parameter：

@cuda.jit( max_registers=40) 

您当然可以将其设置为其他值。一个简单的启发式方法是将每个SM(或每个标题挡路，如果它更低)的寄存器总数(可通过CUDAdeviceQuery示例代码或the programming guide的表15发现)除以您希望启动的每个挡路的线程总数。因此，如果您的GPUSM有64K寄存器，并且您希望每个挡路启动1024线程，那么您可以选择每个线程最多64个寄存器。该号码应适用于腾讯通3090。