cudaMemcpyとcudaMemcpyAsyncの違い

<スパン cudaMemcpyが同期で、cudaMemcpyAsyncが非同期であることは簡単に理解できます。具体的な理解には、以下の概念を明確にする必要があります。

1. CUDA ストリーム

cudaにおけるストリームとは、ホストコードから発行されデバイス上で実行される一連の操作であり、順次実行を保証する必要があります。異なるストリーム内のオペレーションは、クロス実行または同時実行することができます。

2. デフォルトのストリーム

デバイス操作には、データ転送とカーネルがあり、cudaでは、すべてのデバイス操作は、ストリームで実行されます。デフォルトストリームはデバイス操作の同期ストリームで、デフォルトストリーム内の操作の実行は、デバイス上の以前の任意のストリーム内のすべての操作が完了したときにのみ開始され、デフォルトストリーム内の他の操作が開始できる前にデフォルトストリーム内の1つの操作が完了しなければならないことを意味します。

例えば、次のようなコードです。

cudaMemcpy(d_a, a, numBytes, cudaMemcpyHostToDevice); increment<<<1,N>>>(d_a) cudaMemcpy(a, d_a, numBytes, cudaMemcpyDeviceToHost);

デバイス側からは、3つの操作はすべてデフォルトストリームで、順次実行され、ホスト側からは、データ転送はブロッキングまたは同期で、カーネルは非同期である。最初のホストからデバイスへのデータ転送は同期で、ホストからデバイスへのデータ転送が完了するまで、CPUスレッドは2行目に到達できない。カーネルの処理が終わると、CPUスレッドは3行目に移動しますが、デバイス側が2行目を実行しているため、2行目への転送を開始することができません。

<スパン 3. デフォルト以外のストリーム

デフォルト以外のストリームでのデータ転送は、ホスト側でノンブロッキングであり、転送が処理された直後にホストスレッドに制御を戻す関数cudaMemcpyAsync()を使用します。

参考：https://devblogs.nvidia.com/parallelforall/how-overlap-data-transfers-cuda-cc/