动态并行

加速计算 CUDA
1

手头有个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左右的。 版主大人，这是为什么？谢谢！

2

（1）请发kernel部分的代码，谢谢。
（2）请提供您的测时部分的代码，谢谢。
（3）请简要描述您的实现意图，谢谢。
（4）如果可能，请用profiler的测时结果，而不要自己测试。

3

抱歉，我编辑过了帖子，可是依旧显示不全的。

4

您直接给出的为何：
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个函数名就能比较自动脑补出它们的具体实现。因此无法直接回答，您说呢？）

5

手头有个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左右的。 版主大人，这是为什么？谢谢！

6

编辑了帖子，可是显示的一直是我关机前自动保存的内容。抱歉~

7

楼主您好，看到您的代码了，以及您的测试结果（慢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的卡，因为无法替您测试一些东西，抱歉）

8

首先，我要自我检讨下，我程序里有个值弄错了。所以造成50倍的差距。我现在就按照版主的方法来给出一个测试结果的。

9

楼主您的大喘气已经超出了我的认知能力范围了。。。

呜呜。

10

版主，我按照您说的这两种方式即：
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倍，给版主造成的误解，真心抱歉了）

11

嗯嗯。谁都会搞错，没关系的。

楼主你那样要深刻分析了，以及，我建议这忏这样用，这样实际将启动kernel的交互操作从host上移动到了device上，避免了来回的设备-host交互，节省时间。以及，计时您应该只考虑这4个kernel的执行时间，这样才能突出比较的片段。

为何楼主您那样（将后三个放在kernel1的某个线程的最后启动）会导致5s，不好说。
我目前还未申请到3.5的卡，无法评测任何3.5相关内容，只能建议你这样做，而不能深刻分析为何不这样做会变得极度缓慢。请谅解。

以及，关于event计时出现混乱值往往是因为忘记了同步，建议检查下这个。

感谢来访。