system
1
请问一下dynamic parallel中gpu能不能控制循环次数，我想把如下代码全部迁移到gpu上去，怎么改动呢？
while(iterSech<maxSech)
{
if(iterSech != 0)
{
step = step * 0.5;
kernle1<<<dimGrid,dimBlock>>> (…, step_d+iterSech ,…);
}
else
{
step = 100;
kernle2<<<dimGrid,dimBlock>>>(…);
}
gpu_Costfunc_f(…,step_d+iterSech, f_d…);
CudaSafeCall(cudaMemcpy(f,f_d, sizeof(double),cudaMemcpyDeviceToHost));
ftemp=old_f+0.0001 * step * ftemp;
if(*f<=ftemp)
{
CublasSafeCall( cublasDaxpy(m_cublasHandle, Zs, step_d+iterSech, dk_d, 1,x_d, 1) );
break;
}
iterSech=iterSech+1;
}
system
2
楼主您好：
关于是否能控制启动kernel的次数，这个显然是可以的。（继续用普通的for/while/if循环都行）。
但是您后者的cudaMemcpy和cublas可能有麻烦，他们可能无法在Device端使用。
感谢来访。
system
3
话说您可以考虑下cudaMemcpyAsync的，这个是支持的。
cublas刚才看了下，5.0起提供了device版本的，您可以查看手册适合改成使用Device版本的。
以及，您的 gpu_Costfunc_f()函数可能需要改写，以便让它适合device上运行。
大致如此吧。
system
4
gpu_Costfunc_f()函数内部是kernel和cufft的调用
我想把下面这部分cpu代码移到gpu上执行，可是这里break怎么跳出while循环呢?
ftemp=old_f+0.0001 * step * ftemp;
if(*f<=ftemp)
{
CublasSafeCall( cublasDaxpy(m_cublasHandle, Zs, step_d+iterSech, dk_d, 1,x_d, 1) );
break;
}
iterSech=iterSech+1;
system
5
如果您是说：
while(…)
{
您的含有break的这段代码（您可能需要去掉您的CublasSafeCall宏）
}
那么您原本的break具有跳出while的功能，现在的break也依然有跳出while的功能，这个是不变的。
您说呢？
system
6
是不是需要修改while内部已经gpu_costfun_f()函数内部的所有kernel，让每个kernel的最后调用下一个kernel，把原来的cpu调用kernel全部交个gpu调用kernel？
然后把gpu_Costfunc_f()函数内部是kernel和cufft的调用
我想把下面这部分cpu代码移到gpu上执行，可是这里break怎么跳出while循环呢?
ftemp=old_f+0.0001 * step * ftemp;
if(*f<=ftemp)
{
CublasSafeCall( cublasDaxpy(m_cublasHandle, Zs, step_d+iterSech, dk_d, 1,x_d, 1) );
break;
}
iterSech=iterSech+1;
这个代码修改成kernel，在这个内部判断（*f<=ftemp）如果满足条件就返回，不满足条件调用while开始的第一个kernel，这样对吗？
system
7
稍微补充一下5#的玫瑰斑竹：
LZ在4#表示“gpu_Costfunc_f()函数内部是kernel和cufft的调用”，以及根据1#的内容是打算从device端调用，那么cufft有可能不能在device端调用的，请LZ研究下最新的cufft手册，看看有没有明确指出这一点，或者简单测试一下。
祝LZ编码顺利~
system
8
是不是需要修改while内部以及gpu_costfun_f()函数内部的所有kernel，让每个kernel的最后调用下一个kernel，把原来的cpu调用kernel全部改成gpu调用kernel？
我想把下面这部分cpu代码移到写成一个kernel
ftemp=old_f+0.0001 * step * ftemp;
if(*f<=ftemp)
{
CublasSafeCall( cublasDaxpy(m_cublasHandle, Zs, step_d+iterSech, dk_d, 1,x_d, 1) );
break;
}
iterSech=iterSech+1;
在这个kernel内部判断（*f<=ftemp）如果满足条件就返回（作为最后一个kernel退出），如果不满足条件再次调用while开始的第一个kernel，这样对吗？
system
9
LZ您好：
1：各种分支判断和循环，如while，if，for以及break，continue等等都是和C/C++一模一样的，含义绝无修改。
2：不用修改kernel，原先计算的kernel保持原样即可。
3：您直接维护一个小的kernel，实现上述循环和分支判断的逻辑，而将那些大的计算kenrel分别被这个小的控制kernel调用即可，一个大的计算kernel执行完毕，逻辑上会回到原先调用这个kenrel的地方，然后继续向下执行，和host端调用十分类似，此时您可以容易地使用逻辑判断。
4：cufft可能不能在device端调用，请您测试一下再考虑使用。
祝您好运~
system
11
ice你好，你的意思是不是这样的
global kernel_f
{
kernel1<<<>>>();
kernel2<<<>>>();
//下面代码是gpu_costfun_f的简写，假设kernel可以调用cufft
kernel3<<<>>>();
cufft();
kernel4<<<>>>();
if（f_d<=ftemp_d）
break;
else
kernel_f<<<>>>(); //这里再递归调用自己
}
system
12
LZ您好：
我不是这个意思。
您原来的代码中并无递归调用，您为何要在这里加上呢？
您只需要写一个做控制用的小kernel，里面该循环循环，该判断判断，该break就break就完了。
你可以把这个小的kernel想象成device端一个具备kenrel启动能力的特殊的kenrel，或者设想成一个小型的嵌入在GPU里面干活的CPU来执行这些都行（虽然实际原理不是这样）。DP实现起来是非常顺畅和直接的，不知为何您的思路一直如此拧巴。
祝您编码顺利~
system
13
global kernel1()
{
if（f_d<=ftemp_d）
break;
}
while（）
{
kernel1<<<>>>();
}
function();
我的问题是kernel1调用之后如果条件满足执行了break之后，cpu是不是就会跳出while循环执行function？
在kernel内部调用break 怎么跳出while循环，这点怎么理解呢？
system
14
LZ您好：
无法回答更多了，没有人告诉您可以用break结束函数的，您这样越写越离谱了，请回顾下基础知识。
以及，前面的回帖中已经提供了对您的问题的较完备的建议，请您参考。
一般将不再回复本帖，敬请谅解。
祝您好运~
system
15
是写一个kernel，在这个kernel中调用其他kernel以及while循环和break吗？
global kernel
{
while()
{
kerenl1<<<>>>();
kernel2<<<>>>();
if（f_d<=ftemp_d）
break;
}
}
是这样吗?
system
16
我查了一下 cufft好像还不能在设备端调用，那又该怎么办呢？是不是要自己写fft的kernel呢？
system
17
LZ您好：
如果您确实打算DP实现的话，那么您需要自己实现fft了，或者等等看下一版cufft能否提供device端的版本。
祝您好运~
system
18
有没有测试过host调用kernel的延时和device调用kernel的延时的差别呢
system
19
LZ您好：
并无测过这个比较，但是考虑到不用DP的话，经常需要传回host端数据，然后再启动kernel，这个整个事件绝对比device端就地启动kernel更慢。
祝您编码顺利~