我有个工程中使用了多个核函数,其他核函数里面下断点,采用Nsight调试运行时都能在断点处停下来,但是有一个核函数比较奇怪,明明里面有断点,但是采用Nsight调试运行时在断点处却停不下来,这可能是什么原因呢?
显卡GT420,cuda4.2,Vc2008,Nsight好像是2.0还是2.2来着
我有个工程中使用了多个核函数,其他核函数里面下断点,采用Nsight调试运行时都能在断点处停下来,但是有一个核函数比较奇怪,明明里面有断点,但是采用Nsight调试运行时在断点处却停不下来,这可能是什么原因呢?
显卡GT420,cuda4.2,Vc2008,Nsight好像是2.0还是2.2来着
能否发下代码?
很可能此断点无法命中, 或者此断点处没有生成任何指令。
断点肯定可以命中,因为我是在核函数第一句下的
建议您发下代码。
以及楼主如果不方便发代码也没关系,您还可以排查这点:
如果您的kernel是分散于多个.cu中的,您看看您这个无法调试的kernel所在的.cu, 是否没有使用-G编译?而其他能调试的可能您选择了-G。
这个也是有可能的。
我是在工程设置中定义的cuda rules,那个里面是-G无疑,单个cu文件可以搞成不一样的cuda编译规则?
可以的。但最好建议不要。单个文件继承整个项目的的值最好。
今天尝试将核函数一部分语句注释掉,然后发现Nsight调试运行时能在断点处停下来,我现在怀疑是不是核函数的语句太多了,或者是占用的寄存器太多了,或者调用__device__函数是占用的堆栈太多了。
Nsight调试运行时不能在断点处停下来,是不是也不能说明核函数没有正常运行。
要判断核函数是否正常运行,除了getlasterror,有没有什么好用的宏?类似CUDA_SAFE_CALL之类的
核函数内语句数量,寄存器使用数量,调用__device__函数,一般皆不影响。
如果您打算调试检查您的代码,也可以使用检查返回值的方法。
cudaGetLastError()函数使用有一定限制,在之前的帖子里面有详尽的讨论,您可以找来看看。
大致建议为,对于同步操作,请直接检查返回值;对于异步操作,请在之后添加cudaDeviceSynchronize()并检查此函数的返回值。
cudaGetLastError()因为有自己维护的error记录缓冲,可能会造成一定的误解,当然可能用起来稍微简洁一点。
大致建议如上。
请问贵论坛在哪里提供帖子搜索啊?
本论坛目前无搜索功能,已向维护团队反映过,还没改好。
您可以使用google搜索暂时代替。
祝您好运!
针对该问题,我虽然还没有搞清楚原因,但有个变通的方法:
首先把核函数给拆了,然后把__device__函数的调用给注释掉,直接把函数体贴到核函数里面,结果现在可以断下来了。
我想追问3个问题:
1.核函数里面的堆栈使用是不是有限制?比如堆栈使用量不能超过多少?
2.__forceinline__是不是与直接把函数体贴到核函数里面等效?如果等效的话,为啥统计出的局部内存使用量是不一样的?
3.Nsight2.2是不是有在某种情况下,断点断不下来的bug?
LZ您好,您的问题大致答复如下:
1:您说的堆栈是指调用__device__函数时的?我不清楚此限制的大小,但是一般使用似乎没有出过问题。
2:您只能通过__forceinline__来暗示编译器内联,但是无法强迫编译器如此,编译器可以不听您的。
3:一个已知的问题是,在CUDA5和nsight 2.2联用的时候,__device__函数中的断点无法停下来,您可以尝试nsight 3.0。
大致如上,供您参考。
祝您好运~
继续追问:
1.加上__forceinline__之后,编译器输出的当地内存使用信息和堆栈使用信息都与不加不同,如果没有内联的话,这是为啥?
2.假设内联成功的话,进行内联是不是与直接把函数体贴到核函数里面等效?如果等效的话,为啥统计出的局部内存使用量是不一样的?
3.在CUDA5和nsight 2.2联用的时候,__device__函数中的断点无法停下来,是所有__device__函数都无法断下来,还是只有一部分无法断下来吗?
4.我是cuda4.2与nsight 2.2联用,有一些gobal函数中断点无法断下来,而另外的gobal函数和device函数都可以断下来。把那个gobal函数给拆了,然后把__device__函数的调用给注释掉,直接把函数体贴到核函数里面,结果现在可以断下来了。针对这种情况,能否做个简单的推断呢?
1&2:__forceinline__只是暗示编译器将下述函数内联,但是编译器的行为永远是他自己认为最合适的操作,他可能内联该函数也可能不内联。编译器的具体编译结果,您可以dump出汇编代码然后予以分析,但是我无法告知您编译器的具体行为。
3:nsight 2.2的release note里面只表示支持 CUDA Toolkit 4.2 。因此CUDA 5本身不在支持范围内,也不保证所有功能可用。只是经过实践尝试,是基本可用的,只是无法进入__device__函数的断点。由于我没有详细测试过这个问题,所以无法回答您是所有的__device__函数都不能进入断点还是只有一部分。只能建议您如果真要使用CUDA 5+nsight 2.2,那么这是一个已知问题,请设法避免,或者升级到nsight 3。
4:您好,无法推断。
大致如上,供您参考,欢迎讨论和补充,祝您好运!