4

CUDAカーネルが、ホスト側の呼び出し（たとえば、）なしでデバイスマップトメモリへの書き込みを同期することは可能cudaDeviceSynchronizeですか？次のプログラムを実行すると、カーネルの起動直後にページロックされたホストメモリを調べてもメモリの変更が表示されないため、カーネルはデバイスにマップされたメモリへの書き込みが完了するのを待ってから終了するようには見えません（遅延が挿入されるか、への呼び出しcudaDeviceSynchronizeがコメント化されていない場合を除きます）：

#include <stdio.h>
#include <cuda.h>

__global__ void func(int *a, int N) {
    int idx = threadIdx.x;

    if (idx < N) {
        a[idx] *= -1;
        __threadfence_system();
    }
}

int main(void) {
    int *a, *a_gpu;
    const int N = 8;
    size_t size = N*sizeof(int);

    cudaSetDeviceFlags(cudaDeviceMapHost);
    cudaHostAlloc((void **) &a, size, cudaHostAllocMapped);
    cudaHostGetDevicePointer((void **) &a_gpu, (void *) a, 0);

    for (int i = 0; i < N; i++) {
        a[i] = i;
    }
    for (int i = 0; i < N; i++) {
        printf("%i ", a[i]);
    }
    printf("\n");

    func<<<1, N>>>(a_gpu, N);
    // cudaDeviceSynchronize();

    for (int i = 0; i < N; i++) {
        printf("%i ", a[i]);
    }
    printf("\n");

    cudaFreeHost(a);
}

Linux上のCUDA4.2.9を使用してsm_20について上記をコンパイルし、Fermi GPU（S2050）で実行しています。

4

1 に答える 1

4

カーネルの起動は、カーネル アクティビティが発生する前にすぐにホスト コードに戻ります。このように、カーネルの実行はホストの実行とは非同期であり、ホストの実行をブロックしません。したがって、カーネルの結果を確認するには、少し待つか、バリア (cudaDeviceSynchronize() など) を使用する必要があります。

ここで説明されているように：

ホストとデバイス間の同時実行を容易にするために、一部の関数呼び出しは非同期です。デバイスが要求されたタスクを完了する前に、制御がホスト スレッドに返されます。これらは：

  • カーネルが起動します。
  • メモリは 2 つのアドレス間で同じデバイス メモリにコピーされます。
  • 64 KB 以下のメモリ ブロックのホストからデバイスへのメモリ コピー。
  • Async という接尾辞が付いた関数によって実行されるメモリ コピー。
  • メモリ セット関数呼び出し。

これはもちろん、GPU と CPU を同時に使用できるようにするための意図的なものです。この動作が望ましくない場合、すでに発見した簡単な解決策は、バリアを挿入することです。カーネルが生成するデータをすぐにホストにコピーする場合、別のバリアは必要ありません。カーネルの後の cudaMemcpy 呼び出しは、カーネルが完了するまで待ってからコピー操作を開始します。

あなたの質問に答えると思いますが、バリアを使用しなくてもカーネルの起動を同期させたいと考えています (なぜこれを行う必要があるのですか? cudaDeviceSynchronize() 呼び出しを追加するのは問題ですか?) これを行うことは可能です:

「プログラマーは、CUDA_LAUNCH_BLOCKING 環境変数を 1 に設定することにより、システムで実行されているすべての CUDA アプリケーションの非同期カーネル起動をグローバルに無効にすることができます。この機能はデバッグ目的でのみ提供されており、本番ソフトウェアを確実に実行する方法として決して使用しないでください。」

このsynchronous動作が必要な場合は、単にバリアを使用することをお勧めします (または、cudaMemcpy などの後続の別の cuda 呼び出しに依存します)。上記の方法を使用してそれに依存すると、他の誰かが環境変数を設定せずに実行しようとするとすぐにコードが壊れます。したがって、それは本当に良い考えではありません。

于 2012-12-05T01:40:06.293 に答える