4

同時実行が期待される CUDA カーネルでシリアライゼーションの問題に遭遇しました。カーネルの実行を追跡するためのマーカーとして cudaEvents を使用しています。

複数のストリームを持つ同時実行カーネルでの私の実験では、それぞれのストリームでイベントを使用すると、同時実行カーネルがシリアル化されることがわかりました。

以下のコードは、この問題を示しています。以下にリストされている同時カーネル実行機能を持つ 2 つの異なるデバイスでこれをテストしました。

  1. Tesla C2070、ドライバー バージョン 4.10、ランタイム バージョン 4.10、CUDA 機能 2.0
  2. Tesla M2090、ドライバー バージョン 4.10、ランタイム バージョン 4.10、CUDA 機能 2.0

USE_EVENTS マクロを変更することにより、イベントを使用してプログラムを実行したり、イベントを使用せずにプログラムを実行したりできます。同時実行とシリアル実行による違いを確認できます。

#include<cuda.h>
#include<pthread.h>
#include<stdio.h>
#include<stdlib.h>
#include<stdint.h>

#define CUDA_SAFE_CALL( call) do {                                        \
cudaError_t err = call;                                                    \
if( cudaSuccess != err) {                                                \
fprintf(stderr, "Cuda error in call at file '%s' in line %i : %s.\n", \
__FILE__, __LINE__, cudaGetErrorString( err) );              \
exit(-1);                                                     \
} } while (0)



// Device code
__global__ void VecAdd(uint64_t len)
{
    volatile int a;
    for(uint64_t n = 0 ; n < len ; n ++)
        a++; 
    return ;
}

#define USE_EVENTS

int
main(int argc, char *argv[])
{

    cudaStream_t stream[2];
    for(int i = 0 ; i < 2 ; i++) 
        CUDA_SAFE_CALL(cudaStreamCreate(&stream[i]));

#ifdef USE_EVENTS
    cudaEvent_t e[4];
    CUDA_SAFE_CALL(cudaEventCreate(&e[0]));
    CUDA_SAFE_CALL(cudaEventCreate(&e[1]));
    CUDA_SAFE_CALL(cudaEventRecord(e[0],stream[0]));
#endif
    VecAdd<<<1, 32, 0, stream[0]>>>(0xfffffff);

#ifdef USE_EVENTS
    CUDA_SAFE_CALL(cudaEventRecord(e[1],stream[0]));
#endif

#ifdef USE_EVENTS
    CUDA_SAFE_CALL(cudaEventCreate(&e[2]));
    CUDA_SAFE_CALL(cudaEventCreate(&e[3]));
    CUDA_SAFE_CALL(cudaEventRecord(e[2],stream[1]));
#endif
    VecAdd<<<1, 32, 0, stream[1]>>>(0xfffffff);

#ifdef USE_EVENTS
    CUDA_SAFE_CALL(cudaEventRecord(e[3],stream[1]));
#endif
    CUDA_SAFE_CALL(cudaDeviceSynchronize());

    for(int i = 0 ; i < 2 ; i++) 
        CUDA_SAFE_CALL(cudaStreamDestroy(stream[i]));

    return 0;

}

これが発生する理由と、このシリアライゼーションを回避する方法についての提案は役に立ちます。

4

2 に答える 2

3

上記の例の問題は、次の順序で機能します。

1 event record on stream A
2 launch on stream A
3 event record on Stream A
4 event record on stream B
5 launch on stream B
6 event record on stream B

同じストリームに対する CUDA 操作は発行順に実行されます。異なるストリームの CUDA 操作は同時に実行される場合があります。

プログラミング モデルの定義では、同時実行性が必要です。ただし、現在のデバイスでは、この作業は単一のプッシュ バッファーを介して GPU に発行されます。これにより、GPU は操作 2 が完了するのを待ってから操作 3 を発行し、操作 4 が完了するのを待ってから 5 を発行します。... イベント レコードが削除された場合、操作は次のようになります。

1 launch on stream A
2 launch on stream B

操作 1 と 2 は異なるストリーム上にあるため、GPU は 2 つの操作を同時に実行できます。

Parallel Nsight と CUDA コマンド ライン プロファイラー (v4.2) を使用して、同時操作の時間を計ることができます。コマンド ライン プロファイラー オプションは「conckerneltrace」です。この機能は、NVIDIA Visual Profiler の将来のバージョンで表示されるはずです。

于 2012-05-09T03:19:37.160 に答える
1

私は基本的に同じ問題をデバッグしていました。説明は完全ではないようですが、グレッグの答えは非常に役に立ちました。本当の問題は、4が発行されたときにop3が2を待機していることです。4が別のストリームにある場合でも、発行キューで待機中のカーネル/イベントがすでに存在する場合は、発行できません。これは、ストリームごとに複数のカーネルが連続して発行される場合と似ています。これは、次のようにストリーム終了イベントを遅らせることで解決できます。

  1. ストリームAのイベントレコード（開始タイマー）
  2. ストリームAで起動
  3. ストリームBのイベントレコード（開始タイマー）
  4. ストリームBで起動
  5. ストリームAのイベントレコード（終了タイマー）
  6. ストリームBのイベントレコード（終了タイマー）

起動は非同期であるため、ストリームの終了イベントは、そのストリームでの以前のカーネル起動が完了し、他のすべてのストリームのカーネルの問題が起動されるまで待機します。明らかに、これにより、特定のハードウェアで同時に発行できるよりも多くのストリームがある場合、エンドタイマーの発行が遅すぎます。

于 2012-06-14T13:20:54.263 に答える