解释“启动资源太多”

考虑以下Python代码：

from numpy import float64
from pycuda import compiler, gpuarray
import pycuda.autoinit

# N > 960 is crucial!
N = 961
code = """
__global__ void kern(double *v)
{
    double a = v[0]*v[2];
    double lmax = fmax(0.0, a), lmin = fmax(0.0, -a);
    double smax = sqrt(lmax),   smin = sqrt(lmin);

    if(smax > 0.2) {
        smax = fmin(smax, 0.2)/smax ;
        smin = (smin > 0.0) ? fmin(smin, 0.2)/smin : 0.0;
        smin = lmin + smin*a;

        v[0] = v[0]*smin + smax*lmax;
        v[2] = v[2]*smin + smax*lmax;
    }
}
"""
kernel_func = compiler.SourceModule(code).get_function("kern")
kernel_func(gpuarray.zeros(3, float64), block=(N,1,1))

执行此操作将得到：

Traceback (most recent call last):
  File "test.py", line 25, in <module>
    kernel_func(gpuarray.zeros(3, float64), block=(N,1,1))
  File "/usr/lib/python3.5/site-packages/pycuda/driver.py", line 402, in function_call
    func._launch_kernel(grid, block, arg_buf, shared, None)
pycuda._driver.LaunchError: cuLaunchKernel failed: too many resources requested for launch

我的设置：在Ubuntu 16.04.1(64位),内核4.4.0,nvcc V7.5.17上具有pycuda == 2016.1.2和numpy == 1.11.1的Python v3.5.2.显卡是Nvidia GeForce GTX 480.

您可以在机器上重现吗？您有什么主意,是什么导致此错误消息？

备注：我知道原则上存在竞争条件,因为所有内核都尝试更改v [0]和v [2].但是无论如何,内核都不应该到达if块的内部！而且,我可以在没有竞争条件的情况下重现错误,但是它要复杂得多.

解决方法:

几乎可以肯定,您正在达到每个块的寄存器限制.

读取relevant documentation,您的设备每个块最多只能有32k个32位寄存器.当块大小大于960个线程(30个扭曲)时,内核启动需要太多寄存器,启动失败. NVIDIA提供了一个Excel电子表格和advice,用于确定如何确定每个线程的内核注册要求以及可用于在设备上启动内核的限制块大小.

标签：python,numpy,cuda,pycuda
来源： https://codeday.me/bug/20191012/1899868.html