手头有个K20的卡,就把之前的一个程序试着用动态并行去实现。
程序的逻辑是:
cudaMallloc(d_kernel2,d_kernel3,d_kernel4, d_result )
kernel1<<grid1,block1>>(d_kernel2);
kernel2<<grid2,block2>>(d_kernel2,d_kernel3);
最后把d_result传回主机端。程序执行的时间在0.01s左右。
改成动态并行的逻辑是:
cudaMallloc(d_kernel2,d_kernel3,d_kernel4, d_result )
kernel1<<grid1,block1>>(d_kernel2,d_kernel3,d_kernel4,d_result)
同样把d_result传回主机端。程序执行的时间却在0.5s左右的。 版主大人,这是为什么?谢谢!
(1)请发kernel部分的代码,谢谢。
(2)请提供您的测时部分的代码,谢谢。
(3)请简要描述您的实现意图,谢谢。
(4)如果可能,请用profiler的测时结果,而不要自己测试。
您直接给出的为何:
kernel1<<grid1,block1>>(d_kernel2,d_kernel3,d_kernel4,,d_result)
比
kernel1<<grid1,block1>>(d_kernel2);
kernel2<<grid2,block2>>(d_kernel2,d_kernel3);
要慢50倍,
无法回答。
(版主无超能力,无法直接给出2个函数名就能比较自动脑补出它们的具体实现。因此无法直接回答,您说呢?)
手头有个K20的卡,就把之前的一个程序试着用动态并行去实现。
程序的逻辑是:
cudaMallloc(d_kernel2,d_kernel3,d_kernel4, d_result )
kernel1<<grid1,block1>>(d_kernel2);
kernel2<<grid2,block2>>(d_kernel2,d_kernel3);
kernel3<<grid3,block3>>(d_kernel3,d_kernel4);
kernel4<<grid4,block4>>(d_kernel4, d_result );
最后把d_result传回主机端。程序执行的时间在0.01s左右。
改成动态并行的逻辑是:
cudaMallloc(d_kernel2,d_kernel3,d_kernel4, d_result )
kernel1<<grid1,block1>>(d_kernel2,d_kernel3,d_kernel4,d_result)
global void kernel1(d_kernel2,d_kernel3,d_kernel4,d_result)
{
…//获得d_kernel2的值
if(0==threadIdx.x){
kernel2<<grid2,block2>>(d_kernel2,d_kernel3);
kernel3<<grid3,block3>>(d_kernel3,d_kernel4);
kernel4<<grid4,block4>>(d_kernel4, d_result );
}
}
同样把d_result传回主机端。程序执行的时间却在0.5s左右的。 版主大人,这是为什么?谢谢!
编辑了帖子,可是显示的一直是我关机前自动保存的内容。抱歉~
楼主您好,看到您的代码了,以及您的测试结果(慢50倍)超出了目前我对Dynamic Parallelism的认识,
您能否按照之前的要求,发一下您的测时部分的代码?(您可以简单的用profiler看下)
以及,为了公平起见,您否则测时下贵代码这样写的时间:
原来host上:
kernel1<<grid1,block1>>(d_kernel2);
kernel2<<grid2,block2>>(d_kernel2,d_kernel3);
kernel3<<grid3,block3>>(d_kernel3,d_kernel4);
kernel4<<grid4,block4>>(d_kernel4, d_result );
改成:
host 上: kernel0<<<1,1>>>(…);
device 上:
global void kernel0(…)
{
kernel1<<grid1,block1>>(d_kernel2);
kernel2<<grid2,block2>>(d_kernel2,d_kernel3);
kernel3<<grid3,block3>>(d_kernel3,d_kernel4);
kernel4<<grid4,block4>>(d_kernel4, d_result );
}
希望您能配合下给出贵代码的测时结果。
(我还没有申请到3.5的卡,因为无法替您测试一些东西,抱歉)
首先,我要自我检讨下,我程序里有个值弄错了。所以造成50倍的差距。我现在就按照版主的方法来给出一个测试结果的。
版主,我按照您说的这两种方式即:
1、
kernel1<<grid1,block1>>(d_kernel2);
kernel2<<grid2,block2>>(d_kernel2,d_kernel3);
kernel3<<grid3,block3>>(d_kernel3,d_kernel4);
kernel4<<grid4,block4>>(d_kernel4, d_result );
2、
host 上: kernel0<<<1,1>>>(…);
device 上:
global void kernel0(…)
{
kernel1<<grid1,block1>>(d_kernel2);
kernel2<<grid2,block2>>(d_kernel2,d_kernel3);
kernel3<<grid3,block3>>(d_kernel3,d_kernel4);
kernel4<<grid4,block4>>(d_kernel4, d_result );
}
由于我用cudaEvent计时时,出来的结果是乱码,所以我用的clock()计时,循环了10次。计时部分包含了数据分配,传输,释放等所有操作,记录了十次。时间是10次的平均。
对于第一种时间是:0.007810s
对于第二种时间是:0.0045s。
用profile查看,第一种情况,默认流的执行时间为:356.68ms;第二种情况,默认流的执行时间为:173.683ms.(我不知道怎么传图,抱歉)
以及我自己动态并行写法,默认流的执行时间为5.375s。为什么会有这么大差别呢?我没想明白。
(之前结果由于自己程序中有个错误,造成慢了50倍,给版主造成的误解,真心抱歉了)
嗯嗯。谁都会搞错,没关系的。
楼主你那样要深刻分析了,以及,我建议这忏这样用,这样实际将启动kernel的交互操作从host上移动到了device上,避免了来回的设备-host交互,节省时间。以及,计时您应该只考虑这4个kernel的执行时间,这样才能突出比较的片段。
为何楼主您那样(将后三个放在kernel1的某个线程的最后启动)会导致5s,不好说。
我目前还未申请到3.5的卡,无法评测任何3.5相关内容,只能建议你这样做,而不能深刻分析为何不这样做会变得极度缓慢。请谅解。
以及,关于event计时出现混乱值往往是因为忘记了同步,建议检查下这个。
感谢来访。