动态并行

手头有个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计时出现混乱值往往是因为忘记了同步,建议检查下这个。

感谢来访。