为什么启动 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 内核,我可以在 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.9GB 的内存就可以 运行 640 个线程。此 GPU 有 22GB 未使用。

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

这通常是每个线程的寄存器问题 (CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES)。这在 SO cuda 标签(例如 this one. There are many others also such as here 上的许多问题中都有涉及。简而言之,每个线程块使用的总寄存器不能超过 GPU 的限制(见下文)。每个 theadblock 使用的总寄存器大约是每个线程的寄存器总数乘以每个块的线程数(可能为分配粒度四舍五入)。

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

@cuda.jit( max_registers=40) 

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