使用 CUDA 流和 memCpyAsync 的错误结果，添加 cudaDeviceSynchronize 变得正确

Question

我正在开发 CUDA 矩阵乘法，但我做了一些修改以观察它们如何影响性能。

我试图观察一个简单矩阵乘法内核的行为（并且我正在测量 GPU 事件时间的变化）。但我正在两种特定的不同条件下对其进行测试：

我有一定数量的矩阵（比如 matN）用于 A、B 和 C，然后我传输 (H2D) 一个矩阵用于 A，一个用于 B，然后将它们相乘，转回 (D2H) 一个 C;
我对 A、B 和 C 都有 matN，但我在时间上为 A 和 B 传输 >1（比如 chunk）矩阵，准确执行 chunk乘法，传回chunk结果矩阵C.

在第一种情况下 (chunk = 1) 一切都按预期工作，但在第二种情况下 (chunk > 1) 我得到一些 C 是正确的，而另一些是错误的。

但是如果我在 cudaMemcpyAsync 之后加上 cudaDeviceSynchronize()，我得到的所有结果都是正确的。

这是执行我上面刚刚描述的代码的一部分：


/**** main.cpp ****/

    int chunk = matN/iters;    
    #ifdef LOWPAR
        GRIDx= 1;
        GRIDy= 1;
        label="LOW";
    #else
       int sizeX = M;
       int sizeY = N;
       GRIDx = ceil((sizeX)/BLOCK);
       GRIDy = ceil((sizeY)/BLOCK);
       label="";
    #endif

    const int bytesA = M*K*sizeof(float);
    const int bytesB = K*N*sizeof(float);
    const int bytesC = M*N*sizeof(float);

    //device mem allocation
    float *Ad, *Bd, *Cd;
    gpuErrchk( cudaMalloc((void **)&Ad, bytesA*chunk) );
    gpuErrchk( cudaMalloc((void **)&Bd, bytesB*chunk) );
    gpuErrchk( cudaMalloc((void **)&Cd, bytesC*chunk) );
    //host pinned mem allocation
    float *A, *B, *C;
    gpuErrchk( cudaMallocHost((void **)&A, bytesA*matN) );
    gpuErrchk( cudaMallocHost((void **)&B, bytesB*matN) );
    gpuErrchk( cudaMallocHost((void **)&C, bytesC*matN) );

    //host data init
    for(int i=0; i<matN; ++i){
        randomMatrix(M, K, A+(i*M*K));
        randomMatrix(K, N, B+(i*K*N));
    } 

    //event start
    createAndStartEvent(&startEvent, &stopEvent);

    if (square)
    {          
        label += "SQUARE";
        int size = N*N;
        for (int i = 0; i < iters; ++i) { 
            int j = i%nStream;            
            int idx = i*size*chunk;
            newSquareMatMulKer(A+idx, B+idx, C+idx, Ad, Bd, Cd, N, chunk, stream[j]); 
        }
    }
    else {
        ...
    } 

    msTot = endEvent(&startEvent, &stopEvent);
    #ifdef MEASURES          
        printMeasures(square, label, msTot, millis.count(), matN, iters, devId);
    #else
        float *_A, *_B, *_C, *tmpC;
        tmpC = (float *)calloc(1,bytesC*chunk);
        for (int s=0; s<matN; ++s)
        {
            _A = A+(s*M*K);
            _B = B+(s*K*N);
            _C = C+(s*M*N);
            memset(tmpC, 0, bytesC*chunk);

            hostMatMul(_A, _B, tmpC, M, K, N);
            checkMatEquality(_C, tmpC, M, N);
        }   
    #endif


/**** matmul.cu ****/

__global__ void squareMatMulKernel(float* A, float* B, float* C, int N, int chunk) {

    int ROW = blockIdx.x*blockDim.x+threadIdx.x;
    int COL = blockIdx.y*blockDim.y+threadIdx.y;


    if (ROW<N && COL<N) {
        int size=N*N;
        int offs = 0;
        float tmpSum=0.0f;

        for (int s=0; s<chunk; ++s)
        {
            offs = s*size;
            tmpSum = 0.0f;

            for (int i = 0; i < N; ++i) {
                tmpSum += A[offs+(ROW*N)+i] * B[offs+(i*N)+COL];
            }

            C[offs+(ROW*N)+COL] = tmpSum;
        }
    }
    return ;
}




void newSquareMatMulKer(float *A, float *B, float *C, float *Ad, float *Bd, float *Cd, 
            int n, int chunk, cudaStream_t strm)
{
    int size = n*n;
    int bytesMat = size*sizeof(float);

    dim3 dimBlock(BLOCK,BLOCK,1);
    dim3 dimGrid(GRIDx, GRIDy,1); 

    gpuErrchk( cudaMemcpyAsync(Ad, A, bytesMat*chunk, cudaMemcpyHostToDevice, strm) );    
    gpuErrchk( cudaMemcpyAsync(Bd, B, bytesMat*chunk, cudaMemcpyHostToDevice, strm) );   

    #ifdef LOWPAR
        squareMatMulGridStrideKer<<<dimGrid, dimBlock, 0, strm>>>(Ad, Bd, Cd, n, chunk);
    #else
        squareMatMulKernel<<<dimGrid, dimBlock, 0, strm>>>(Ad, Bd, Cd, n, chunk);
    #endif
    squareMatMulKernel<<<dimGrid, dimBlock, 0, strm>>>(Ad, Bd, Cd, n, chunk);

    gpuErrchk( cudaMemcpyAsync( C, Cd, bytesMat*chunk, cudaMemcpyDeviceToHost, strm) );

    cudaDeviceSynchronize();
        ^ ^ ^ ^ ^ ^
}

我尝试使用 cuda-gdb 进行调试，但没有出现任何异常，gpuErrchk 不会在 CUDA API 调用中引发任何错误。我运行代码也使用了 memcheck，在有和没有 cudaDeviceSynchronize 的情况下，在这两种情况下我都没有错误。

我想我可以说这是一个同步问题，但我不明白这背后的原因是什么。有人能发现我哪里出错了吗？也非常感谢其他代码风格建议。

Answer 1

如果您使用多个流，您可以在使用它们之前覆盖 Ad 和 Bd。

iters = 2 和 nStream = 2 的示例：

for (int i = 0; i < iters; ++i) { 
  int j = i%nStream;            
  int idx = i*size*chunk;
  newSquareMatMulKer(A+idx, B+idx, C+idx, Ad, Bd, Cd, N, chunk, stream[j]); 
}

在此循环中，您将调用

newSquareMatMulKer(A, B, C, Ad, Bd, Cd, N, chunk, stream[0]); // call 0
newSquareMatMulKer(A+idx, B+idx, C+idx, Ad, Bd, Cd, N, chunk, stream[1]); // call 1

由于您在设备上为这两个调用使用相同的内存区域，您可能会遇到几个同步问题：

call 1 在 call 0:squareMatMulKernel 结束前开始在设备上复制 A 和 B，因此您可能会使用不正确的 [=18= 值] and/or B 来计算你的第一次迭代。
call 1:squareMatMulKernel 在您从调用 0 中检索 C 的值之前开始，因此您可以使用 call 1 中的值覆盖 C。

要解决这个问题，我看到了两种方法：

使用 cudaDeviceSynchronize();.
例如，您可以在两个设备端分配更多内存（每个流一个工作空间）。

''

//device mem allocation
float *Ad, *Bd, *Cd;
gpuErrchk( cudaMalloc((void **)&Ad, bytesA*chunk*nStream) );
gpuErrchk( cudaMalloc((void **)&Bd, bytesB*chunk*nStream) );
gpuErrchk( cudaMalloc((void **)&Cd, bytesC*chunk*nStream) );

/* code here */

for (int i = 0; i < iters; ++i) { 
  int j = i%nStream;            
  int idx = i*size*chunk;
  int offset_stream = j*size*chunk;
  newSquareMatMulKer(A+idx, B+idx, C+idx, 
    Ad + offset_stream , 
    Bd + offset_stream , 
    Cd + offset_stream , N, chunk, stream[j]); 
}

在这种情况下，您不需要在循环结束前进行同步。

使用 CUDA 流和 memCpyAsync 的错误结果，添加 cudaDeviceSynchronize 变得正确

Wrong results using CUDA streams and memCpyAsync, become correct adding cudaDeviceSynchronize

cuda

cuda-streams