矩阵相乘在 CUDA 上给出错误的结果

Question

我编写了一小段代码，使用 CUDA 将 2 个方阵相乘。然而，事实证明大多数单元格计算错误。根据我用过的教程，一切都会好起来的。

__global__ void gpuMM(int *C, int *A, int *B, int N)
{
    int row = blockIdx.x*blockDim.x + threadIdx.x;
    int col = blockIdx.y*blockDim.y + threadIdx.y;
    int sum = 0;
    for (int n = 0; n < N; ++n)
        sum += A[row*N+n]*B[n*N+col];

    C[row*N+col] = sum;
}

#define ROW_SIZE 5
#define MATRIX_LENGTH ROW_SIZE*ROW_SIZE
#define BLOCK_SIZE 16

void MultiplyMatrixCUDA(int * pResult, int* pFactorA, int*pFactorB)
{
    int size = MATRIX_LENGTH*sizeof(int);
    int *dA,*dB,*dC;
    cudaMalloc(&dA,size);
    cudaMalloc(&dB,size);
    cudaMalloc(&dC,size);
    int K = 100;
    dim3 threadBlock(BLOCK_SIZE,BLOCK_SIZE);
    dim3 grid(K,K);

    printf("A:\n");
    DrawMatrix(pFactorA);
    printf("\n");

    printf("B:\n");
    DrawMatrix(pFactorB);
    printf("\n");

    // Copy matrices from the host to device
    cudaMemcpy(dA,pFactorA,size,cudaMemcpyHostToDevice);
    cudaMemcpy(dB,pFactorB,size,cudaMemcpyHostToDevice);

    //Execute the matrix multiplication kernel
    gpuMM<<<grid,threadBlock>>>(dC,dA,dB,ROW_SIZE);


    // Allocate memory to store the GPU answer on the host
    int *C;
    C = new int[MATRIX_LENGTH];

    // Now copy the GPU result back to CPU
    cudaMemcpy(C,dC,size,cudaMemcpyDeviceToHost);

    cudaFree(dA);
    cudaFree(dB);
    cudaFree(dC);

    printf("\nC from CUDA:\n");
    DrawMatrix(C);
    printf("\nC:\n");
    DrawMatrix(MultiplyWithCPU(pResult,pFactorA, pFactorB));  // the code of multiplying function is irrevelant, I'm sure it works fine (double-checked)

}

结果表明矩阵乘以标准 CPU 方法是正确的，但 CUDA 是错误的：

第一行总是正确的，但所有其他部分都是完全随机的。有时它们是消极的，有时不是。有时它们接近真实值，有时它们完全不同。

我的错误是什么？我看不出哪里失败了。算法看起来不错，变量似乎传递正确，但有些东西不起作用。

--- 编辑

所有变量（pResult 和两个 pFactors）都在代码的其他部分初始化（然后删除）。

Answer 1

由于每个块的线程数不等于输出矩阵中的元素数（您将 5x5 矩阵映射到 16x16 块上），因此一些线程正在访问/写入无效的内存位置。

解决方案包括双边界检查以解决问题。这将导致一些线程空闲。内核应该是这样的：

__global__ void gpuMM(int *C, int *A, int *B, int N)
{
    int row = blockIdx.x*blockDim.x + threadIdx.x;
    int col = blockIdx.y*blockDim.y + threadIdx.y;

    if( (row < N) && (col < N))
    {
        int sum = 0;
        for (int n = 0; n < N; ++n){
            sum += A[row*N+n]*B[n*N+col];       
        }
        C[row*N+col] = sum;
    }
}

另一个解决方案 - 更有效，实际上取决于您的设备 - 每个块启动更少的线程（在本例中为 25 个）。