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

您正在做的事情——或者至少使用一个事件来同步两个流——应该可以工作。基本上不可能说为什么您的实际代码不起作用，因为您选择不发布它，而魔鬼总是在细节中。

但是，这是一个完整的、可运行的示例，我认为它以类似于您正在尝试做的方式使用流 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 中，可以在不人工序列化执行流的情况下分析异步代码。如果您无法从本示例或您自己的代码中解决问题，这可能会帮助您诊断问题。

于 2012-06-02T15:15:27.780 回答