我的程序大致如下：
while(iterSech<maxSech)
{
kernel1<<<dimGrid,dimBlock>>>(…);
kernel2<<<dimGrid,dimBlock>>>(…);
fun(f);
if(*f<=ftemp)
break;
iterSech=iterSech+1;
}
fun(…,…,f)
{
kernel3<<<dimGrid,dimBlock>>>(…);
kernel4<<<dimGrid,dimBlock>>>(…);
}
profiler while循环发现，两次循环之间间隔时间有500us，在computer行上是空闲的，而在Runtime API行上为cudaLaunch、 cudaFunSetCacheConfig、cudaBindTexture、cudaUnbindTexture、cudaSetupArgument等（上传了几次profiler截图都没成功只能这样描述了）
我对profiler分析报的时间轴有一些不理解的地方[list=1]
[]computer行的kernel的执行和 Runtime API行的 cudaFree是重叠的。
[]computer行 一段空闲时间段和 Runtime API行的cudaLaunch、 cudaFunSetCacheConfig、cudaBindTexture、cudaUnbindTexture、cudaSetupArgument等是重合的。
[/list]以上两点怎么理解呢？难道所有的kernel都配置完成了再执行？
楼主您确定您的fun(f)里，
没有调用任何的纹理绑定、解除绑定、设置L1/shared比例么？
您给出的2次循环间的5个API调用，cudaLaunch和setupArgument是自动的，可以无视。
但其他的三个纹理绑定，解除绑定，设置Shared memory比例，这个不会自动进行的。
所以最大的可能是，
您的fun()里写的不好，每次都重新绑定了纹理，重新设置了L1和shared memory的比例，这个可能故你试图封装一些kernel调用为普通的C函数时候，不小心每次都重复执行的。
关于kernel执行和cudaFree重叠。这个您先无视。这应该是实际发出命令的时间，和您的kernel在GPU上延迟一段时间后执行造成的错觉。可以无视。
感谢您的深夜来访。
再说下上文最后一点，
因为您启动kernel后，kernel不会在0ns内立刻执行，这有一定的延迟的，而此时，如果您下面立刻来个cudaFree(), 那么会造成从时间线上看上去您的cudaFree()提前于kernel执行了。
但此时实际上是安全的：
您的cudaFree()在您的kernel执行完毕之前，都会处于挂起状态，不会实际执行的（但会占据时间线，甚至会造成您的cudaFree和您的kernel在“同时”执行的假象）。
而当您的kernel执行完毕的瞬间，cudaFree才开始被实际的执行。
（这叫隐式同步。不仅仅是cudaFree。cudaMalloc也是如此，您的kernel完成之前，实际上不会为您执行cudaMalloc的）
所以大可放心，不用因为看上去时间线很接近或者完全重叠，而担心缓冲区半路被释放。
感谢您的夜晚来访。
谢谢玫瑰版主的提醒，
我的fun（）中调用了cufft，所以纹理绑定，解除绑定，设置Shared memory比例等应该是cufft内部调用的。
我现在想不通的是，Runtime API行中对kernel参数的设置，kernel的启动等与kernel的执行好像错开了一段时间，我们平时测量kernel的执行时间并没有包括kernel参数设置启动的时间,这部分时间好像很长，这难道是你所说的kernel在GPU上延迟执行的错觉吗?
compute行所显示的每个kernel的开始与结束时间应该的真实的吧？
我的两次循环间隔500us难道是cpu中断gpu的时间？如果是这个时间怎么回这么长呢，因为我的cpu每次循环只做了判断。
你想测试你的代码，就不要引入cufft, 单独抽取出来你的kernel测试。
避免cufft的神马纹理绑定，Shared memory大小设定之类的影响。
关于你的第二个问题，从你<<<>>>到你的kernel函数实际开始执行的延迟，实际上并无妨碍，
先不说几百us的执行延迟是否大或者小。但凡考虑一点：只要你连续的<<<>>>, kernel也会实际上连续的执行的。所以你不用担心一次启动的中间的延时。
你说呢？
玫瑰版主您好，我的问题的原因找到了。我上面的伪代码漏了一行
cudamemcpy(f_h,f_d,sizeof(double),cudaMemcpyDeviceToHost)
补充后是这样的：
while(iterSech<maxSech)
{
kernel1<<<dimGrid,dimBlock>>>(…);
kernel2<<<dimGrid,dimBlock>>>(…);
fun(f_d);
cudamemcpy(f_h,f_d,sizeof(double),cudaMemcpyDeviceToHost);
if(*f_h<=ftemp)
break;
iterSech=iterSech+1;
}
fun(…,…,f)
{
kernel3<<<dimGrid,dimBlock>>>(…);
kernel4<<<dimGrid,dimBlock>>>(…);
}
两次循环的500us延时是由cudamemcpy引起的。
我把cudamemcpy行注释之后，在timeline中两次循环间隔很短，几乎就是两个kernel的调用间隔时间。
所以又有了三个问题：
（1）为什么cudamemcpy 耗时这样严重？
（2）想把整个代码都交给GPU完成，包括如下的判断。
if(*f_h<=ftemp)
break;
kepler架构的GPU不是可以实现 Dynamic Parallelism吗，不知道可不可以完成循环的控制（跳出循环）？
（3）我的GPU为GTX680 是不是不支持 Dynamic Parallelism，可以推荐几款支持 Dynamic Parallelism的显卡吗？
LZ您好：
1：cudaMemcpy()有这样的延迟开销很正常，考虑到他需要做很多后台的工作，还要通过pci-e总线建立通信等等。
2：我觉得是可以的，您可以上一个小的kernel专门做控制，然后这个kernel通过DP发布的kernel才是真正计算用的。
3：GTX680是SM3.0的卡，不支持DP，如果您需要以较低的成本评估DP，那么可以考虑GK28核心的GT630显卡。如果您需要计算效能较强的，那么推荐tesla K20/K20X 以及geforce TITAN/780/780Ti。
大致如此，感谢您深夜来访，祝您晚安~