2

私のセットアップのスケルトンに従ってください。このように実行すると、正しい結果が得られません。これは、カーネルが使用するときに非同期データ転送が完了していないことが原因である可能性が最も高いです。if-elseプリプロセッサステートメントを使用して「フェイルセーフ」バージョンを実装しました。部分を翻訳するとelse、プログラムは正常に実行されます。理解できません。どうして？

、、... は単なるプレースホルダーですin1。out1もちろん、それらは for ループの反復ごとに異なるコンテナーを指します。非同期転送ができるようにします。ただし、反復内ではout1、転送によって使用されるものとカーネルによって使用されるものは同じです。

  cudaStream_t streams[2];
  cudaEvent_t  evCopied;

  cudaStreamCreate(&streams[0]); // TRANSFER
  cudaStreamCreate(&streams[1]); // KERNEL

  cudaEventCreate(&evCopied);

  // many iterations
  for () {

    // Here I want overlapping of transfers with previous kernel
    cudaMemcpyAsync( out1, in1, size1, cudaMemcpyDefault, streams[0] );
    cudaMemcpyAsync( out2, in2, size2, cudaMemcpyDefault, streams[0] );
    cudaMemcpyAsync( out3, in3, size3, cudaMemcpyDefault, streams[0] );

#if 1
    // make sure host thread doesn't "run away"
    cudaStreamSynchronize( streams[1] );
    cudaEventRecord( evCopied , streams[0] );
    cudaStreamWaitEvent( streams[1] , evCopied , 0);
#else
    // this gives the correct results
    cudaStreamSynchronize( streams[0] );
    cudaStreamSynchronize( streams[1] );
#endif

    kernel<<< grid , sh_mem , streams[1] >>>(out1,out2,out3);

  }

セットアップの再配置を示唆する回答を投稿しないでください。たとえば、カーネルをいくつかに分割し、別々のストリームで発行します。

4

1 に答える 1

3

あなたがしていること、または少なくともイベントを使用して2つのストリームを同期させることはうまくいくはずです。投稿しないことを選択したために実際のコードが機能しない理由を言うことは基本的に不可能であり、悪魔は常に詳細にあります。

ただし、これは完全で実行可能な例であり、あなたがやろうとしていることと同様の方法でストリーム API を使用しており、正しく動作していると思います。

#include <cstdio>

typedef unsigned int uint;

template<uint bsz>
__global__ void kernel(uint * a, uint * b, uint * c, const uint N)
{
    __shared__ volatile uint buf[bsz];
    uint tid = threadIdx.x + blockIdx.x * blockDim.x;
    uint stride = blockDim.x * gridDim.x;
    uint val = 0;
    for(uint i=tid; i<N; i+=stride) {
        val += a[i] + b[i];
    }
    buf[threadIdx.x] = val; __syncthreads();

#pragma unroll
    for(uint i=(threadIdx.x+warpSize); (threadIdx.x<warpSize)&&(i<bsz); i+=warpSize)
        buf[threadIdx.x] += buf[i];

    if (threadIdx.x < 16) buf[threadIdx.x] += buf[threadIdx.x+16];
    if (threadIdx.x < 8)  buf[threadIdx.x] += buf[threadIdx.x+8];
    if (threadIdx.x < 4)  buf[threadIdx.x] += buf[threadIdx.x+4];
    if (threadIdx.x < 2)  buf[threadIdx.x] += buf[threadIdx.x+2];
    if (threadIdx.x == 0) c[blockIdx.x] += buf[0] + buf[1];

}

#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true)
{
   if (code != cudaSuccess) 
   {
      fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
      if (abort) exit(code);
   }
}

int main(void)
{
    const int nruns = 16, ntransfers = 3;
    const int Nb = 32, Nt = 192, Nr = 3000, N = Nr * Nb * Nt;
    const size_t szNb = Nb * sizeof(uint), szN = size_t(N) * sizeof(uint);
    size_t sz[4] = { szN, szN, szNb, szNb };

    uint * d[ntransfers+1];
    for(int i=0; i<ntransfers+1; i++)
        gpuErrchk(cudaMallocHost((void **)&d[i], sz[i]));
    uint * a = d[0], * b = d[1], * c = d[2], * out = d[3];

    for(uint i=0; i<N; i++) {
        a[i] = b[i] = 1; 
        if (i<Nb) c[i] = 0;
    }

    uint * _d[3];
    for(int i=0; i<ntransfers; i++)
        gpuErrchk(cudaMalloc((void **)&_d[i], sz[i])); 
    uint * _a = _d[0], * _b = _d[1], * _c = _d[2];

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

    cudaEvent_t sync_event;
    gpuErrchk(cudaEventCreate(&sync_event)); 

    uint results[nruns];
    for(int j=0; j<nruns; j++) {
        for(int i=0; i<ntransfers; i++)
            gpuErrchk(cudaMemcpyAsync(_d[i], d[i], sz[i], cudaMemcpyHostToDevice, stream[0]));

        gpuErrchk(cudaEventRecord(sync_event, stream[0]));
        gpuErrchk(cudaStreamWaitEvent(stream[1], sync_event, 0));

        kernel<Nt><<<Nb, Nt, 0, stream[1]>>>(_a, _b, _c, N);
        gpuErrchk(cudaPeekAtLastError());

        gpuErrchk(cudaMemcpyAsync(out, _c, szNb, cudaMemcpyDeviceToHost, stream[1]));
        gpuErrchk(cudaStreamSynchronize(stream[1]));

        results[j] = uint(0);
        for(int i=0; i<Nb; i++) results[j]+= out[i];
    }

    for(int j=0; j<nruns; j++) 
        fprintf(stdout, "%3d: ans = %u\n", j, results[j]);

    gpuErrchk(cudaDeviceReset());
    return 0;
}

カーネルは「融合ベクトル加算/削減」であり、まったくナンセンスですが、カーネル実行前に 3 つの入力の最後のゼロを使用して、正しい答えを生成します。これは、入力データ ポイントの数の 2 倍である必要があります。あなたの例のように、カーネルの実行と非同期入力配列のコピーは異なるストリームにあるため、コピーと実行が重複する可能性があります。この場合、反復ごとに最初の 2 つの大きな入力をコピーする正当な理由はありません。最後のコピー (重要なコピー) が完了する前に遅延が発生し、カーネルと誤ってオーバーラップする可能性が高くなるからです。CUDA メモリ モデルでは、実行中のカーネルがアクセスするメモリを非同期的に変更しても安全であるとは思えないため、ここが間違っている可能性があります。それがあなたがやろうとしていることなら、それが失敗することを期待してください。しかし、実際のコードを見なければ、それ以上のことは言えません。

cudaStreamWaitEventこれで、カーネルが起動する前に 2 つのストリームを同期しないと、カーネルが正しい結果を生成しないことがわかります。疑似コードとこの例の唯一の違いはcudaStreamSynchronize、実行ストリーム上の の場所です。ここでは、転送前にカーネルが終了して結果をホストに戻すようにするために、カーネルの起動後に配置しました。それが重要な違いかもしれませんが、実際のコードは実際のコード分析と同じではありません....

私が提案できるのは、この例で遊んで、それがどのように機能するかを感じることだけです. Nsight for Windows のごく最近のバージョンでは、プロファイリングによって実行ストリームを人為的にシリアル化することなく、非同期コードをプロファイリングできる可能性があることを理解しています。この例または独自のコードから問題を解決できない場合は、問題の診断に役立つ可能性があります。

于 2012-06-02T15:15:27.780 に答える