调试自动中断 grid lunch failure

今天遇到两个问题,有劳版主及各位过目。
1、调试一个比较大型的程序,调试的时候,在没有设断点的地方中断了,提示:
CUDA grid launch failed: CUcontext: 3231464 CUmodule: 322612552 Function: _Z7denoisePf
这是什么意思呢,什么原因造成的?

2、我的程序有八个kernel,kernel之间是串行的关系,没有设置流,我是将每个核函数调试好之后再通过main函数组装起来的,debug模式下运行的时候几秒可以完成。但是我在其中的某个kernel中设置断点,然后启动调试,要等十几分钟才能在断点处停下来,这样的话我就没有办法调试组装好之后的完整程序。这是什么原因造成的?

楼主您好:

Grid launch failed, 90%是您的启动配置不当! (例如kernel的形状,shared memory动态配置等)

请您立刻查看您的启动配置。

请您先尝试修正BUG后,再尝试询问第二个问题。一步一步的来,否则你会越来越乱。

程序是这样启动的:

float *matd;
 cudaMalloc(&matd,m*n*sizeof(float));
 cudaMemcpy(matd,mat,m*n*sizeof(float),cudaMemcpyHostToDevice);
 filter<<<512,n>>>(matd);
 cudaThreadSynchronize();
 denoise<<<512,nq>>>(matd);

m=160000,n=128,nq=32
内部分配的shared memory的大小是固定大小的。
shared float appx[n2+4scale],detl[n+4*scale];
程序在denoise内部某一条指令自动中断

LZ您好:

您这个线程的形状是没有问题的,所以此时是玫瑰斑竹说的90%可能性之外的剩余10%的可能中的情况,一般为kernel代码在很靠前的地方就访存越界了,此时nsight会提示CUDA grid lanuch failed。请您检查代码。

另外,部分早期版本的nsight在这种情况下会出现“卡着不动”的情况,机器亦不死机。经测试,nsight 3.2无此现象,可以立即报告lanuch failed。您也可以选nsight菜单——option——CUDA——CUDA Memory Checker——Enable Memory Checker选为 True,以辅助您定位访存越界。

大致如此,祝您编码顺利~

我的Nsight3.0不知道算不算早期版本,memory check功能是一直都开着的。
我已经找到了问题所在的核函数,但无法知道是什么问题,问题描述:
无论是DEBUG还是RELEASE模式,只要不调试,而是直接运行,就没有任何问题,并且运行结果是正确的。一旦启动调试,就会出现问题。具体现象为:
我在appx[tid]=matd[xr];指令处设置条件断点,其中xr为global memory地址索引,tid为shared memory地址索引。假如我给这个kernel每个block分配的是32线程,则当条件为“xr<32”时可以正常断下来,当断点条件为xr大于等于32时,例如xr=32,则不能正常断下来,也不会跳出,就像您说的停下来不动了。假如我给每个BLOCK分配64线程,则当xr<64正常断下,xr>=64不能断下。
不能断下我就等吧,希望它会断下来,等了好久不见动静,然后屏幕一黑,显示没有信号,但主机依然在运行。我只能重启电脑,我被这个程序深深地折服了。
我当时调试的时候只跟踪默认的tid=0,xr=0的数据,发现一切正常。但我没有想到当xr超过blockDim的时候会出现异常。
根据您的提示,可能是访存越界,但是越界一般是发生在最后的存储单元,而我现在无法在后面断下来。
请经验丰富的版主们指点迷津。

LZ您好:

前面说的启动线程形状的问题以及kernel早期越界的问题都是针对“lanuch failed”说的,不是针对您无故卡死的情况说的,虽然某些nsight版本在kernel早期越界的时候也会卡死。

刚才实机测试了nsight 3.2下,类似您说的情况,能正常在断点处停下来。

请您先升级到nsight 3.2(注意需要较新的驱动支持),然后再尝试一下。

目前无其他建议。

祝您好运~