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

Question

我是CUDA的新手。我试图添加两个 vector ，它工作正常。现在我要添加两个矩阵。我想使用二维线程(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

每当遇到CUDA代码问题时，都应始终使用proper cuda error checking并使用cuda-memcheck运行代码。建议您先在寻求帮助，然后再执行。即使您不理解错误输出，它对于试图帮助您的人也会很有用。

如果您已经执行了正确的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);

下一个错误是您的size变量的计算:

int size = N * N;

cudaMalloc和cudaMemcpy与malloc和memcpy一样，期望大小为字节。因此，您应该这样做:

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

最后，您要在打印出之前将数据复制回主机。所以这是不正确的:

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
$