使用二维线程在 CUDA 中添加两个矩阵

Question

我是 CUDA 的新手。我试图添加两个向量，它工作正常。现在我想添加两个矩阵。我想使用二维线程（threadIdx.x 和 threadIdx.y）添加两个矩阵。我在 Internet 上找到了此代码，并进行了一些更改以显示结果。它编译。但是显示出意外的结果，它看起来像内存地址。请帮助我，提前谢谢你。

#include <stdio.h>
#include <stdlib.h>

#define N 5
#define BLOCK_DIM 10

__global__ void matrixAdd (int *a, int *b, int *c) {
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int row = blockIdx.y * blockDim.y + threadIdx.y;

    int index = col + row * N;

    if (col < N && row < N) {
        c[index] = a[index] + b[index];
    }

}

int main() {
    int a[N][N], b[N][N], c[N][N];
    int *dev_a, *dev_b, *dev_c;

    int size = N * N;

    for(int i=0; i<N; i++)
        for (int j=0; j<N; j++){
            a[i][j] = 1;
            b[i][j] = 2;
        }

    cudaMalloc((void**)&dev_a, size);
    cudaMalloc((void**)&dev_b, size);
    cudaMalloc((void**)&dev_c, size);

    cudaMemcpy(dev_a, a, size, cudaMemcpyHostToDevice);
    cudaMemcpy(dev_b, b, size, cudaMemcpyHostToDevice);

    dim3 dimBlock(BLOCK_DIM, BLOCK_DIM);
    dim3 dimGrid((int)ceil(N/dimBlock.x),(int)ceil(N/dimBlock.y));

    matrixAdd<<<dimGrid,dimBlock>>>(dev_a,dev_b,dev_c);
    cudaDeviceSynchronize();

    for(int i=0; i<N; i++){
        for (int j=0; j<N; j++){
            printf("%d\t", c[i][j] );
        }
        printf("\n");
    }

    cudaMemcpy(c, dev_c, size, cudaMemcpyDeviceToHost);

    cudaFree(dev_a); 
    cudaFree(dev_b); 
    cudaFree(dev_c);
}

输出是

0   0   -780197879  32659   1   
0   452489360   32764   6303208 0   
4198328 0   452489376   32764   4198181 
0   2   0   4198557 0   
4196864 0   0   0   4198480 

我的预期输出是元素 3 的 5x5 矩阵。请帮助我。

Answer 1

您应该始终使用proper cuda error checking 并使用cuda-memcheck 运行您的代码，任何时候您在使用CUDA 代码时遇到问题。我建议您在在此处寻求帮助之前进行此操作。即使您不理解错误输出，它也会对那些试图帮助您的人有用。

1. 如果您进行了正确的 CUDA 错误检查，您会在内核启动时收到“无效参数”错误的通知。它是由于这个计算而产生的：

   dim3 dimGrid((int)ceil(N/dimBlock.x),(int)ceil(N/dimBlock.y));
   

   如果您在计算后打印出实际计算值：

   printf("dimGrid.x = %d, dimGrid.y = %d\n", dimGrid.x, dimGrid.y);
   

   你会发现它们都是零。这是违法的。

   这个结构：

   N/dimBlock.x
   

   正在使用整数除法。当 N=5 和 dimBlock.x = 10 时，这两个数字的整数除法为零。使用ceil 不影响这个，你写的方式。

   有很多方法可以解决它。一种可能的方法是像这样进行算术运算：

   dim3 dimGrid((N+dimBlock.x-1)/dimBlock.x, (N+dimBlock.y-1)/dimBlock.y);
   

2. 下一个错误是在计算您的size 变量时：

   int size = N * N;
   

   cudaMalloc 和 cudaMemcpy，如 malloc 和 memcpy，期望大小以 字节 为单位。所以你应该这样做：

   int size = N * N * sizeof(int);
   

3. 最后，您希望将数据复制回主机在打印出来之前。所以这是不正确的：

   for(int i=0; i<N; i++){
       for (int j=0; j<N; j++){
           printf("%d\t", c[i][j] );
       }
       printf("\n");
   }
   
   cudaMemcpy(c, dev_c, size, cudaMemcpyDeviceToHost);
   

   你想这样做：

   cudaMemcpy(c, dev_c, size, cudaMemcpyDeviceToHost);
   for(int i=0; i<N; i++){
       for (int j=0; j<N; j++){
           printf("%d\t", c[i][j] );
       }
       printf("\n");
   }
   

这是一个工作示例，显示了对代码的上述修改：

$ cat t1058.cu
#include <stdio.h>
#include <stdlib.h>

#define N 5
#define BLOCK_DIM 10

__global__ void matrixAdd (int *a, int *b, int *c) {
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int row = blockIdx.y * blockDim.y + threadIdx.y;

    int index = col + row * N;

    if (col < N && row < N) {
        c[index] = a[index] + b[index];
    }

}

int main() {
    int a[N][N], b[N][N], c[N][N];
    int *dev_a, *dev_b, *dev_c;

    int size = N * N * sizeof(int);

    for(int i=0; i<N; i++)
        for (int j=0; j<N; j++){
            a[i][j] = 1;
            b[i][j] = 2;
        }

    cudaMalloc((void**)&dev_a, size);
    cudaMalloc((void**)&dev_b, size);
    cudaMalloc((void**)&dev_c, size);

    cudaMemcpy(dev_a, a, size, cudaMemcpyHostToDevice);
    cudaMemcpy(dev_b, b, size, cudaMemcpyHostToDevice);

    dim3 dimBlock(BLOCK_DIM, BLOCK_DIM);
    //dim3 dimGrid((int)ceil(N/dimBlock.x),(int)ceil(N/dimBlock.y));
    dim3 dimGrid((N+dimBlock.x-1)/dimBlock.x, (N+dimBlock.y-1)/dimBlock.y);
    printf("dimGrid.x = %d, dimGrid.y = %d\n", dimGrid.x, dimGrid.y);
    matrixAdd<<<dimGrid,dimBlock>>>(dev_a,dev_b,dev_c);
    cudaDeviceSynchronize();
    cudaMemcpy(c, dev_c, size, cudaMemcpyDeviceToHost);

    for(int i=0; i<N; i++){
        for (int j=0; j<N; j++){
            printf("%d\t", c[i][j] );
        }
        printf("\n");
    }


    cudaFree(dev_a);
    cudaFree(dev_b);
    cudaFree(dev_c);
}
$ nvcc -o t1058 t1058.cu
$ cuda-memcheck ./t1058
========= CUDA-MEMCHECK
dimGrid.x = 1, dimGrid.y = 1
3 3 3 3 3
3 3 3 3 3
3 3 3 3 3
3 3 3 3 3
3 3 3 3 3
========= ERROR SUMMARY: 0 errors
$