3

私は休眠中のコードでホストスレッドでビジーウェイトを取り除く方法を探しています（そのコードをコピーしないでください、それは私の問題のアイデアを示すだけで、多くの基本的なバグがあります）：

cudaStream_t steams[S_N];
for (int i = 0; i < S_N; i++) {
    cudaStreamCreate(streams[i]);
}
int sid = 0;
for (int d = 0; d < DATA_SIZE; d+=DATA_STEP) {
     while (true) {
         if (cudaStreamQuery(streams[sid])) == cudaSuccess) { //BUSY WAITING !!!!
             cudaMemcpyAssync(d_data, h_data + d, DATA_STEP, cudaMemcpyHostToDevice, streams[sid]);
             kernel<<<gridDim, blockDim, smSize streams[sid]>>>(d_data, DATA_STEP);
             break;
         }
         sid = ++sid % S_N;
     }

}

ホストスレッドをアイドル状態にして、何らかのストリームが終了するのを待ってから、別のストリームを準備して実行する方法はありますか？

編集：ビジーウェイトを強調するために、コードにwhile（true）を追加しました。ここで、すべてのストリームを実行し、そのうちのどれが別の新しいストリームの実行を終了したかを確認します。cudaStreamSynchronize特定のストリームが終了するのを待ちますが、最初にジョブを終了したストリームのいずれかを待ちたいです。

EDIT2：私は忙しいのをやめました-休眠の方法で待っています：

cudaStream_t steams[S_N];
for (int i = 0; i < S_N; i++) {
    cudaStreamCreate(streams[i]);
}
int sid = 0;
for (int d = 0; d < DATA_SIZE; d+=DATA_STEP) {
    cudaMemcpyAssync(d_data, h_data + d, DATA_STEP, cudaMemcpyHostToDevice, streams[sid]);
    kernel<<<gridDim, blockDim, smSize streams[sid]>>>(d_data, DATA_STEP);
    sid = ++sid % S_N;
}
for (int i = 0; i < S_N; i++) {
    cudaStreamSynchronize(streams[i]);
    cudaStreamDestroy(streams[i]);
}

ただし、ホストスレッドでビジーウェイトを使用しているバージョンよりも少し遅いようです。これは、現在、ストリームにジョブを静的に分散しているためだと思います。そのため、1つのストリームが作業を終了すると、各ストリームが作業を終了するまでアイドル状態になります。以前のバージョンでは、作業が最初のアイドルストリームに動的に分散されたため、より効率的でしたが、ホストスレッドでビジーウェイトが発生していました。

4

5 に答える 5

6

本当の答えは、cudaThreadSynchronizeを使用して以前のすべての起動が完了するのを待ち、 cudaStreamSynchronizeを使用して特定のストリームのすべての起動が完了するのを待ち、 cudaEventSynchronizeを使用して特定のストリームの特定のイベントのみが記録されるのを待つことです。

ただし、コードでストリームと同期を使用するには、ストリームと同期がどのように機能するかを理解する必要があります。

ストリームをまったく使用しないとどうなりますか？次のコードを検討してください。

kernel <<< gridDim, blockDim >>> (d_data, DATA_STEP);
host_func1();
cudaThreadSynchronize();
host_func2();

カーネルが起動され、ホストはhost_func1とカーネルを同時に実行するために移動します。次に、ホストとデバイスが同期されます。つまり、ホストはカーネルが終了するのを待ってから、host_func2（）に移動します。

では、2つの異なるカーネルがある場合はどうなりますか？

kernel1 <<<gridDim, blockDim >>> (d_data + d1, DATA_STEP);
kernel2 <<<gridDim, blockDim >>> (d_data + d2, DATA_STEP);

kernel1は非同期で起動されます！ホストが移動し、kernel1が終了する前にkernel2が起動されます。ただし、kernel2は両方ともストリーム0（デフォルトのストリーム）で起動されているため、kernel1が終了するまで実行されません。次の代替案を検討してください。

kernel1 <<<gridDim, blockDim>>> (d_data + d1, DATA_STEP);
cudaThreadSynchronize();
kernel2 <<<gridDim, blockDim>>> (d_data + d2, DATA_STEP);

デバイスは同じストリームで起動されたカーネルをすでに同期しているため、これを行う必要はまったくありません。

ですから、あなたが探している機能はすでに存在していると思います...カーネルは常に同じストリーム内の以前の起動が終了するのを待ってから開始します（ホストが通過したとしても）。つまり、以前の起動が終了するのを待ちたい場合は、単にストリームを使用しないでください。このコードは正常に機能します。

for (int d = 0; d < DATA_SIZE; d+=DATA_STEP) {
    cudaMemcpyAsync(d_data, h_data + d, DATA_STEP, cudaMemcpyHostToDevice, 0);
    kernel<<<gridDim, blockDim, smSize, 0>>>(d_data, DATA_STEP);
 }

さて、ストリームに移りましょう。ストリームを使用して、デバイスの同時実行を管理できます。

ストリームをキューと考えてください。さまざまなmemcpy呼び出しとカーネル起動をさまざまなキューに入れることができます。次に、ストリーム1のカーネルとストリーム2の起動は非同期です。それらは同時に、または任意の順序で実行できます。デバイスで一度に1つのmemcpy/kernelのみが実行されていることを確認したい場合は、ストリームを使用しないでください。同様に、カーネルを特定の順序で実行する場合は、ストリームを使用しないでください。

とはいえ、ストリーム1に入れられたものはすべて順番に実行されるので、わざわざ同期しないでください。同期は、ホストとデバイスの呼び出しを同期するためのものであり、2つの異なるデバイスの呼び出しではありません。したがって、異なるデバイスメモリを使用し、相互に影響を与えないために複数のカーネルを同時に実行する場合は、ストリームを使用します。何かのようなもの...

cudaStream_t steams[S_N];
for (int i = 0; i < S_N; i++) {
    cudaStreamCreate(streams[i]);
}

int sid = 0;
for (int d = 0; d < DATA_SIZE; d+=DATA_STEP) {
     cudaMemcpyAsync(d_data, h_data + d, DATA_STEP, cudaMemcpyHostToDevice, streams[sid]);
     kernel<<<gridDim, blockDim, smSize streams[sid]>>>(d_data, DATA_STEP);
     sid = ++sid % S_N;
 }

明示的なデバイス同期は必要ありません。

于 2011-02-25T13:02:46.097 に答える
4

その問題を解決するための私の考えは、1つのストリームごとに1つのホストスレッドを持つことです。そのホストスレッドはcudaStreamSynchronizeを呼び出して、ストリームコマンドが完了するまで待機します。残念ながら、CUDA 3.2では1つのCUDAコンテキストを処理できるホストスレッドは1つだけであるため、これは不可能です。つまり、1つのCUDA対応GPUごとに1つのホストスレッドを意味します。

うまくいけば、CUDA 4.0でそれが可能になるでしょう：CUDA4.0RCニュース

編集：私はオープンmpを使用してCUDA4.0RCでテストしました。cudaストリームごとに1つのホストスレッドを作成しました。そしてそれは働き始めました。

于 2011-03-01T09:04:33.203 に答える
3

あります：cudaEventRecord(event, stream)とcudaEventSynchronize(event)。リファレンスマニュアルhttp://developer.download.nvidia.com/compute/cuda/3_2/toolkit/docs/CUDA_Toolkit_Reference_Manual.pdfにすべての詳細があります。

編集：ところで、ストリームはカーネルとメモリ転送の同時実行に便利です。現在のストリームが終了するのを待って実行をシリアル化するのはなぜですか？

于 2011-02-24T18:02:45.357 に答える
2

cudaStreamQueryの代わりに、 cudaStreamSynchronizeが必要です

int sid = 0;
for (int d = 0; d < DATA_SIZE; d+=DATA_STEP) {
     cudaStreamSynchronize(streams[sid]);
     cudaMemcpyAssync(d_data, h_data + d, DATA_STEP, cudaMemcpyHostToDevice, streams[sid]);
     kernel<<<gridDim, blockDim, smSize streams[sid]>>>(d_data, DATA_STEP);
     sid = ++sid % S_N;
}

（cudaThreadSynchronizeを使用して、すべてのストリームでの起動を待機し、cudaEventSynchronizeを使用したイベントを使用して、より高度なホスト/デバイスの同期を行うこともできます。）

これらの同期機能で発生する待機の種類をさらに制御できます。cudaDeviceBlockingSyncフラグなどのリファレンスマニュアルを参照してください。ただし、デフォルトはおそらくあなたが望むものです。

于 2011-02-24T18:18:13.867 に答える
1

データチャンクをコピーし、そのデータチャンクで別のforループでカーネルを実行する必要があります。それはより効率的になります。

このような：

size = N*sizeof(float)/nStreams;

for (i=0; i<nStreams; i++){
    offset = i*N/nStreams;
    cudaMemcpyAsync(a_d+offset, a_h+offset, size, cudaMemcpyHostToDevice, stream[i]);
}


for (i=0; i<nStreams; i++){
    offset = i*N/nStreams;
    kernel<<<N(nThreads*nStreams), nThreads, 0, stream[i]>>> (a_d+offset);
}

このようにして、メモリコピーは前のストリームのカーネル実行を待つ必要はなく、その逆も同様です。

于 2011-02-25T05:43:06.417 に答える