-1

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

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

  • 我有一定数量的矩阵（比如matN）A，B和C，然后我将（H2D）一个矩阵转移给A，一个矩阵转移给B，然后将它们相乘，以转移回（D2H）一个C；

  • 我有A、B 和 C 的任何一个，但我在时间为 A 和 BmatN传输 >1（比如说）矩阵，执行精确的乘法，然后传回结果矩阵 C。chunkchunkchunk

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

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

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


/**** 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，在这两种情况下我都没有收到错误。

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

4

1 回答 1

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在结束之前开始复制A和B在设备上call 0:squareMatMulKernel，因此您可能会使用不正确的值A和/或B来计算您的第一次迭代。

  • call 1:squareMatMulKernelC在您从调用 0检索值之前开始，因此您可以C使用来自call 1.

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

  • 在您的示例中使用同步与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]); 
}

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

于 2019-06-19T09:35:47.303 回答