CUDA記錄並等待事件不起作用？

Question

遵循我的設置的骨架。這樣執行它不會給出正確的結果。這很可能是由於內核使用它們的異步數據傳輸尚未完成。我使用預處理器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); 

    } 

請不要發佈建議重新安排的答案。就像，把你的內核分成幾個，然後分開發送。

Answer 1

，你在做什麼 - 或者至少使用一個事件的兩個流進行同步 - 應該工作。它基本上是不可能說爲什麼你的實際代碼不起作用，因爲您已選擇不發佈它，魔鬼總是在細節。

然而，這裏是一個完整的，可運行的例子，我認爲類似你正在嘗試做的，它的方式是使用流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; 
} 

內核是「融合向量添加/減少「，只是無稽之談，但它依賴於在內核執行之前將三個輸入中的最後一個置零以產生正確的答案，其應該簡單地是輸入數據點數的兩倍。正如在你的例子中，內核執行和異步輸入數組複製是在不同的流中，所以複製和執行可能會重疊。沒有健全的原因，前兩個大投入在每次迭代在這種情況下複製，其他的最後一個副本（這是關鍵的一年）之前完成，而不是引入延遲，並提高它會不正確地與內核重疊的機會。這可能是您出錯的地方，因爲我不相信CUDA內存模型可以保證異步修改正在運行的內核正在訪問的內存是安全的。 如果這是你正在嘗試做的，那麼期望它失敗。但是如果沒有看到真實的代碼，就不可能說更多。

由此可以看出，如果沒有cudaStreamWaitEvent在內核啓動之前同步兩個流，內核將不會產生正確的結果。您的僞代碼和此示例之間唯一的區別是執行流上的cudaStreamSynchronize的位置。在這裏，我將它放在內核啓動後，以確保內核在傳輸前完成以將結果收集回主機。這可能是關鍵的區別，但是也沒有真正的代碼等於沒有真正的代碼分析....

所有我可以建議你是用這個例子玩得到它是如何工作的感覺。我知道有可能在最近的Nsight for Windows版本中對異步代碼進行概要分析而不用人工序列化執行流。如果你無法從這個例子或你自己的代碼中找出問題，那麼這可能能夠幫助你診斷你的問題。