调试自动中断 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后，再尝试询问第二个问题。一步一步的来，否则你会越来越乱。

玫瑰幻想 发表于 2013-11-6 15:26
楼主您好：

Grid launch failed, 90%是您的启动配置不当！ （例如kernel的形状，shared memory动态配置等 ...

程序是这样启动的：

1. float *matd;
   
2. cudaMalloc(&matd,m*n*sizeof(float));
   
3. cudaMemcpy(matd,mat,m*n*sizeof(float),cudaMemcpyHostToDevice);
   
4. filter<<<512,n>>>(matd);
   
5. cudaThreadSynchronize();
   
6. denoise<<<512,nq>>>(matd);

复制代码

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

翊君 发表于 2013-11-6 15:40
程序是这样启动的：
m=160000,n=128,nq=32
内部分配的shared memory的大小是固定大小的。

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，以辅助您定位访存越界。

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

ice 发表于 2013-11-6 19:10
LZ您好：

您这个线程的形状是没有问题的，所以此时是玫瑰斑竹说的90%可能性之外的剩余10%的可能中的情况 ...

我的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的时候会出现异常。
根据您的提示，可能是访存越界，但是越界一般是发生在最后的存储单元，而我现在无法在后面断下来。
请经验丰富的版主们指点迷津。

翊君 发表于 2013-11-6 21:19
我的Nsight3.0不知道算不算早期版本，memory check功能是一直都开着的。
我已经找到了问题所在的核函数， ...

LZ您好：

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

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

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

目前无其他建议。

祝您好运~