5

私は、2 つの cuda ストリームが進行し、イベントによって管理され、互いに待機する小さなデモ プログラムを作成しようとしています。これまでのところ、このプログラムは次のようになります。

// event.cu
#include <iostream>
#include <cstdio>
#include <cuda_runtime.h>
#include <cuda.h>

using namespace std;

__global__ void k_A1() { printf("\tHi! I am Kernel A1.\n"); }
__global__ void k_B1() { printf("\tHi! I am Kernel B1.\n"); }
__global__ void k_A2() { printf("\tHi! I am Kernel A2.\n"); }
__global__ void k_B2() { printf("\tHi! I am Kernel B2.\n"); }

int main()
{
  cudaStream_t streamA, streamB;
  cudaEvent_t halfA, halfB;
  cudaStreamCreate(&streamA);
  cudaStreamCreate(&streamB);
  cudaEventCreate(&halfA);
  cudaEventCreate(&halfB);

  cout << "Here is the plan:" << endl <<
    "Stream A: A1, launch 'HalfA', wait for 'HalfB', A2." << endl <<
    "Stream B: Wait for 'HalfA', B1, launch 'HalfB', B2." << endl <<
    "I would expect: A1,B1, (A2 and B2 running concurrently)." << endl;

  k_A1<<<1,1,0,streamA>>>(); // A1!
  cudaEventRecord(halfA,streamA); // StreamA triggers halfA!
  cudaStreamWaitEvent(streamA,halfB,0); // StreamA waits for halfB.
  k_A2<<<1,1,0,streamA>>>(); // A2!

  cudaStreamWaitEvent(streamB,halfA,0); // StreamB waits, for halfA.
  k_B1<<<1,1,0,streamB>>>(); // B1!
  cudaEventRecord(halfB,streamB); // StreamB triggers halfB!
  k_B2<<<1,1,0,streamB>>>(); // B2!

  cudaEventDestroy(halfB);
  cudaEventDestroy(halfA);
  cudaStreamDestroy(streamB);
  cudaStreamDestroy(streamA);

  cout << "All has been started. Synchronize!" << endl;
  cudaDeviceSynchronize();
  return 0;
}

CUDA ストリームに関する私の理解は次のとおりです。ストリームは、タスクを追加できる一種のリストです。これらのタスクは連続して取り組まれます。したがって、私のプログラムでは、streamA が正常に動作することを確認できます。

  1. カーネル k_A1 を呼び出す
  2. トリガーハーフA
  3. 誰かが halfB をトリガーするのを待ちます
  4. カーネル k_A2 を呼び出す

そしてstreamBは

  1. 誰かが halfA をトリガーするのを待ちます
  2. カーネル k_B1 を呼び出す
  3. トリガーハーフB
  4. カーネル k_B2 を呼び出す

通常、両方のストリームは互いに非同期で実行される可能性があります。ただし、A1 が完了するまで streamB をブロックし、B1 が完了するまで streamA をブロックしたいと思います。

これはそれほど単純ではないようです。Tesla M2090 (CC 2.0) を搭載した私の Ubuntu では、

nvcc -arch=sm_20 event.cu && ./a.out

は

Here is the plan:
Stream A: A1, launch 'HalfA', wait for 'HalfB', A2.
Stream B: Wait for 'HalfA', B1, launch 'HalfB', B2.
I would expect: A1,B1, (A2 and B2 running concurrently).
All has been started. Synchronize!
        Hi! I am Kernel A1.
        Hi! I am Kernel A2.
        Hi! I am Kernel B1.
        Hi! I am Kernel B2.

そして、cudaEventRecord(halfB,streamB) の前に B1 が完了することを本当に期待していたでしょう。それにもかかわらず、ストリーム A は明らかに B1 の完了を待たないため、halfB の記録も待ちません。

さらに、cudaEventRecord コマンドを完全に削除すると、プログラムが cudaStreamWait コマンドでロックダウンすることが予想されます。しかし、そうではなく、同じ出力を生成します。私はここで何を見落としていますか？

4

1 に答える 1

7

これは、「halfB」が記録される前に「cudaStreamWaitEvent(streamA,halfB,0);」が呼び出されたためだと思います(cudaEventRecord(halfB,streamB);)。cudaStreamWaitEvent 呼び出しが、その前に閉じられた「halfB」を検索していた可能性があります。見つからなかったので、静かに前進しました。次のドキュメントを参照してください。

ストリームは、 onstreamへの最新のホスト呼び出しが完了するまで待機します。この呼び出しが戻ると、任意の関数 (およびを含む) を再度呼び出すことができ、その後の呼び出しは に影響しません。cudaEventRecord()eventcudaEventRecord()cudaEventDestroy()eventstream

深さ優先のコーディングを行う必要がある場合、解決策が見つかりませんでした。ただし、次のコードは、必要なものにつながる可能性があります。

  k_A1<<<1,1,0,streamA>>>(d); // A1!
  cudaEventRecord(halfA,streamA); // StreamA triggers halfA!
  cudaStreamWaitEvent(streamB,halfA,0); // StreamB waits, for halfA.
  k_B1<<<1,1,0,streamB>>>(d); // B1!
  cudaEventRecord(halfB,streamB); // StreamB triggers halfB!
  cudaStreamWaitEvent(streamA,halfB,0); // StreamA waits for halfB.
  k_A2<<<1,1,0,streamA>>>(d); // A2!
  k_B2<<<1,1,0,streamB>>>(d); // B2!

これはプロファイリングによって確認されます:

ここに画像の説明を入力

カーネル インターフェイスを変更したことに注意してください。

于 2013-03-23T22:17:49.280 に答える