dim3换线程组织形式导致无法调试的问题

  • 8 replies
  • 869 views
dim3换线程组织形式导致无法调试的问题
« 于: 十一月 22, 2019, 10:12:31 pm »
kernel<<<32,32 >>>();
这种可以调试;
但改一下block里thread的组织结构后:
dim3 blockSize(32, 32);
kernel<<<32,blockSize >>> ();
这种无法调试,也就是核函数里的断点停不下来。

现在用的CUDA10.0,显卡是RTX2080ti,显卡驱动430.86。

Re: dim3换线程组织形式导致无法调试的问题
« 回复 #1 于: 十一月 23, 2019, 04:26:27 pm »
调用的时候就这三行(感觉就没跑这个核函数),如果只注销第二行,就可以跑进来。if(threadIdx.y==0)这种判断加过了(这里.x.y没改,不要在意),并没有效果。

Re: dim3换线程组织形式导致无法调试的问题
« 回复 #2 于: 十一月 25, 2019, 01:12:45 pm »
调用的时候就这三行(感觉就没跑这个核函数),如果只注销第二行,就可以跑进来。if(threadIdx.y==0)这种判断加过了(这里.x.y没改,不要在意),并没有效果。

if(threadIdx.y == 0)这种肯定能抵达的。那么如果多加了这一句话,if的body永远无法抵达的话,则不妨先看看kernel能否启动,有的时候因为debug编译不能正常启动kernel(资源超了,没优化),你需要手工添加:
__launch_bounds__(需要保证的block中的线程数)来修饰你的kernel。

建议的解决方案:
(1)检查kernel是否能够启动。
(2)如果不能,添加资源使用修饰,以便尝试让它能够启动。
(3)如果依然不能,报告论坛你遇到的具体情况。(例如:我能抵达if的上一句,或者kernel的开头,但如果继续单步就飞了,类似这种的具体情况描述)。

Re: dim3换线程组织形式导致无法调试的问题
« 回复 #3 于: 十一月 25, 2019, 01:14:31 pm »
if(threadIdx.y == 0)这种肯定能抵达的。那么如果多加了这一句话,if的body永远无法抵达的话,则不妨先看看kernel能否启动,有的时候因为debug编译不能正常启动kernel(资源超了,没优化),你需要手工添加:
__launch_bounds__(需要保证的block中的线程数)来修饰你的kernel。

建议的解决方案:
(1)检查kernel是否能够启动。
(2)如果不能,添加资源使用修饰,以便尝试让它能够启动。
(3)如果依然不能,报告论坛你遇到的具体情况。(例如:我能抵达if的上一句,或者kernel的开头,但如果继续单步就飞了,类似这种的具体情况描述)。

注意这个修饰是指的你的block中的总线程数量!
例如(32,32)的block,这里应该写1024。
这是一个很常见的debug编译不能启动kernel,而release编译(或者开启了优化后),能够启动kernel的重要原因。

Re: dim3换线程组织形式导致无法调试的问题
« 回复 #4 于: 十一月 25, 2019, 11:05:13 pm »
注意这个修饰是指的你的block中的总线程数量!
例如(32,32)的block,这里应该写1024。
这是一个很常见的debug编译不能启动kernel,而release编译(或者开启了优化后),能够启动kernel的重要原因。
用launch_bounds,在device code这里复现了手册上写的例子,没编译过,我是用法错了么。

Re: dim3换线程组织形式导致无法调试的问题
« 回复 #5 于: 十一月 25, 2019, 11:16:59 pm »
额。。。用法看了半天没看懂。

Re: dim3换线程组织形式导致无法调试的问题
« 回复 #6 于: 十一月 26, 2019, 12:53:37 pm »
用launch_bounds,在device code这里复现了手册上写的例子,没编译过,我是用法错了么。

你写重复了,直接写一次在函数最开头就可以了。

Re: dim3换线程组织形式导致无法调试的问题
« 回复 #7 于: 十一月 26, 2019, 12:58:57 pm »
定位到问题了, 是blockSize决定的,之前这就是我资源调用的瓶颈。
一下几个情况可以跑进核函数,可以正常调试:
1.kernel<<<32,32>>>>();
2.
dim3 gridSize(4,4,1);
dim3 blockSize(8,8,1);
kernel<<<gridSize,blockSize>>>>();
3.
dim3 gridSize(64,64,1);
dim3 blockSize(16,16,1);
kernel<<<gridSize,blockSize>>>>();
但就是blockSize写不成(32,32,1)
我想写成用32个thread组成一个任务单位,分配一定的资源,主线用threadIdx.x计算,需要大量计算的时候,用剩下的31个thread,辅助计算。
但是block数量虽然可以写上去,block里面的thread上不去。

Re: dim3换线程组织形式导致无法调试的问题
« 回复 #8 于: 十一月 26, 2019, 01:32:09 pm »
定位到问题了, 是blockSize决定的,之前这就是我资源调用的瓶颈。
一下几个情况可以跑进核函数,可以正常调试:
1.kernel<<<32,32>>>>();
2.
dim3 gridSize(4,4,1);
dim3 blockSize(8,8,1);
kernel<<<gridSize,blockSize>>>>();
3.
dim3 gridSize(64,64,1);
dim3 blockSize(16,16,1);
kernel<<<gridSize,blockSize>>>>();
但就是blockSize写不成(32,32,1)
我想写成用32个thread组成一个任务单位,分配一定的资源,主线用threadIdx.x计算,需要大量计算的时候,用剩下的31个thread,辅助计算。
但是block数量虽然可以写上去,block里面的thread上不去。

没太看懂你的中文。使用32 * 32线程的block, 还是32 * 1或者其他形状的block,主要取决于你block内部共享数据的需求。如果你需要在32个warps间有效的交换数据(例如走shared memory),则你应当上(32,32),启动不了就上launch bounds;反过来,如果你不需要交换数据,则你可以维持(32, 2), 甚至(32, 1)的大小(注意这种可能会影响occupancy)。

这个不是我能决定的,而是取决于你的内心。