同一個串流內kernel間的運行需要同步嗎?

加速计算 CUDA
1 

int main()
{
  ...
   initial<<<...>>>();
   cudaThreadSynchronize();
   
   for(double t=dt;t<dt*5;t=t+dt)
   {
   step1<<<grid,block>>>(d_Vx,d_Vy,d_VxT,d_VyT);
   cudaThreadSynchronize();
   bool tf=1;
   step2<<<grid,block>>>(d_rho,d_VxT,d_VyT);
   cudaThreadSynchronize();
   for(int step=0;step<4000;step=step+1)
   {
   if(tf)
   {poisson<<<grid,block>>>(d_p,d_pT,d_rho);}
   else
   {poisson<<<grid,block>>>(d_pT,d_p,d_rho);}

   tf= !(tf);
   /*host coda*/
   }
   cudaThreadSynchronize();
   step3<<<grid,block>>>(d_Vx,d_Vy,d_VxT,d_VyT,d_p);
   cudaThreadSynchronize();
   
   }
 ...
 return 0;
}

我的觀念可能不太清楚，想要請教一下同步的事情。

以上是我的程式碼，大概就是先配置好變數後其他事情通通都在GPU上運作

我看visual profiler中kernel的運行是依照順序進行的，而kernel的呼叫是早於運行(不加同步的話)

所以呼叫的時間被執行時給隱藏，即使我的每一步的資料有相依性，但因為kernel式串行運作所以保障資料同步

但如果我要加入host在kernel間且資料是要kernel算完的話，則需要先同步等待GPU算完，
然後再要求GPU等待直到CPU算完把資料丟回去

如此才確保資料的正確性。我想確認這樣的理解有沒有錯？

因為我發現我把同步拿掉後時間會省下很多，但計算的資料應該沒有錯誤，所以想確認一下。

另外，我看visual profiler在cudaMemcpy時有些會標[sync]有些則沒有，原因是什麼？

是說某些cudaMemcpy並沒有自動同步的效果嗎？

2

楼主您好，多个一次启动的kernel中间无需host和GPU进行同步，您可以在总工作完成后，执行一次同步（显式的或者隐式的。例如cudaMemcpy会自动等待之前的所有kernel完成)

关于您的第二个问题，某些cudaMemcpy在profiler里被标记上了Sync字样。这个我不知道是怎么回事。建议其他会员、版主、NVIDIA官方支持、总版主给予解答。

最后，一个建议：
cudaThreadSynchronize()已经不推荐使用了。这个是在3.2之前的context模型(每个host thread有自己的gpu context)时使用的。现在这个函数已经被重定向为cudaDeviceSynchronize()了。您可以使用这个。

3

感謝版主的建議，我看的書可能比較舊以後會改正使用cudaDeviceSynchronize()

4

LZ您好，您的问题简要答复如下：

1：同一个stream里面的kernel是顺序执行的，无需同步以“等待前一个kernel结束”，前一个kernel结束后一个kernel才运行是被保证的。
2：以及如果您的kernel处理的数据都在显存上的话，您也只需要反复发射kernel即可，无需其他处理。如果有需要传回host memory的数据，那么其间插入的cudaMemcpy是同步操作，会确保之前kernel计算都完成之后，才copy回来。
3：以及cudaMemcpy也有异步版本，详情您可以参考官方手册。

祝您编码顺利~