cudaDeviceSynchronize() の代替関数の存在について

Question

質問事項

cuda 11.6からcudaDeviceSynchronize関数は非推奨になりましたが，ドキュメントを確認しても，代替関数の記載がありません。いずれ代替の関数が提供されるまで，従来の非推奨関数をつかっていくのか，他に代替方法があるのか，その点を教えていただきたく，質問しました。よろしくお願いします。

実現したいこと

複数のblockの処理結果をまとめるスレッドにおいて，同期させる必要があります。
device 上で　block の全処理を終了するまで停止させたいといことです。できれば，コンパイラからの警告抑制をせずに，同期処理を行いたいということです。

補足情報（FW/ツールのバージョンなど）

ubuntu 20.04
cuda 11.7

Answer 1

kernelを分けて二段階launchするとか、cooperative_groupのGrid Synchronization使うとか。

Comments

[jbe00214]

Kernelを分けてやってみました。kernelの間にcudaDeviceSynchronizeを挟んで。すると2倍遅くなりました。GridSynchronizeは使えるか調べてみます。アドバイスありがとうございます。

[jbe00214]

cooperative_group の sync()を使ってみました。おかげさまでうまくいきましたが，<<<,>>>の代わりに，cudaLaunchCooperativeKernelを呼ぶことになり，この関数はHostでしか呼び出せないので，結局はkernelを分けて実行するのと同じという結論になりました。しかし，新たな選択肢を習得できたので感謝しています。ありがとうございました。