CUDA 2D 阵列英伟达

Question

我将cudaMallocPitch 和cudaMemcpy2D 用于二维数组。即使我无法正确获得输出，我也不确定我的编码是否正确。有人可以帮忙吗？任何人都可以调试我的错误吗？提前致谢。

#include<stdio.h>
#include<cuda.h>
#define siz 4*sizeof(int)
__global__ void addmatrix(int *m1,int *m2,size_t pitch)
{
    int r=threadIdx.x;
    int *r1=m1+r*pitch;
    int *r2=m2+r*pitch;
    int c;
    for(c=1;c<=4;c++)
    {
        r1[c]+=r2[c];
    }
}
int main()
{
    int i,j;
    int **m1_c,**m2_c;
    int *m1_d,*m2_d;
    size_t pitch;
    cudaError_t err;
    m1_c=(int **)malloc(4*sizeof(int *));
    for(i=1;i<=4;i++)
    {
        m1_c[i]=(int *)malloc(siz);
    }
    m2_c=(int **)malloc(4*sizeof(int *));
    for(i=1;i<=4;i++)
    {
        m2_c[i]=(int *)malloc(siz);
    }
    for(i=1;i<=4;i++)
    {
        for(j=1;j<=4;j++)
        {
            m1_c[i][j]=rand()%10;
            m2_c[i][j]=rand()%10;
        }
    }
    for(i=1;i<=4;i++)
    {
        for(j=1;j<=4;j++)
        {
            printf("%d\t",m1_c[i][j]);
        }
        printf("\n");
    }
    printf("\n\n");
    for(i=1;i<=4;i++)
    {
        for(j=1;j<=4;j++)
        {
            printf("%d\t",m2_c[i][j]);
        }
        printf("\n");
    }
    err=cudaMallocPitch((void **)&m1_d,&pitch,siz,siz);
    err=cudaMallocPitch((void **)&m2_d,&pitch,siz,siz);
    err=cudaMemcpy2D(m1_d,pitch,m1_c,siz,siz,4,cudaMemcpyHostToDevice);
    err=cudaMemcpy2D(m2_d,pitch,m2_c,siz,siz,4,cudaMemcpyHostToDevice);
    dim3 grid(1);
    dim3 block(16);
    addmatrix<<<grid,block>>>(m1_d,m2_d,siz);
    cudaMemcpy2D(m1_c,siz,m1_d,pitch,siz,4,cudaMemcpyDeviceToHost);

    for(i=1;i<=4;i++)
    {
        for(j=1;j<=4;j++)
        {
            printf("%d\t",m1_c[i][j]);
        }
        printf("\n");
    }
    err=cudaFree(m1_d);
    err=cudaFree(m2_d);
    err=cudaDeviceReset();      
}

Answer 1

所以这段代码有几个问题。排名不分先后：

1. 您正在对从 1 到 4 的各种数组进行索引，但这在 C 中是不正确的。C 索引从零开始，并且比维度小一。这与 CUDA 无关。
2. cudaMemcpy2D 需要两个指针（src 和 dst），它们都是指向内存中线性数组的指针。我意识到这很令人困惑，因为 2D 出现在整个描述中，但是这两个指针参数基本上都是相同类型的（指向内存的指针），并且您正在传递两种不同类型的指针（一个是指向内存的指针，另一个是指向内存的指针）是指向内存指针的指针）。从 cudaMemcpy2D 的定义很清楚，你的用法是不正确的。有很多关于如何使用 cudaMemcpy2D 的示例的已回答问题，我建议您搜索并查看其中的一些。请注意，修复此问题可能会导致您从根本上重新考虑如何将数据存储在主机矩阵上。有很多关于处理多维矩阵的问题，例如this one——如果可能的话，你应该将它们展平。请注意，在您当前的代码中，使用 cudaMemcpy2D 的此错误正在破坏主机矩阵上的指针数组，这会在您尝试打印结果时导致段错误。
3. 您传递给cudaMallocPitch 的参数不太正确。对于width 和height 参数，您将传递siz 这是矩阵维度以字节为单位。但是您应该只传递 width 参数的字节维度。对于 height 参数，您应该传递行数，即在您的情况下为 4。对 cudaMemcpy2D 的调用也有类似的要求，但您已经做到了。
4. 现在让我们看看您的内核。在调用中，您将启动一个由 16 个线程组成的块的网格。由于您的矩阵有 16 个元素，这似乎是明智的。这意味着线程策略，其中每个线程将负责结果的单个元素。但是看看你的内核代码，你让每个线程计算一整行的结果，即 4 个元素。有两种方法可以解决这个问题：您可以将网格减少到 4 个线程而不是 16 个线程（从代码修改的角度来看可能更简单），或者您可以重新编写内核（消除 for 循环）并拥有每个线程计算单个输出元素（可能会并行执行更多工作）。
5. 此外，在您的内核中，您在基于指针算术的索引中使用pitch 参数。但请记住，音高以 bytes 为单位，对于指针算术索引，编译器希望参数在 elements 中 - 它会根据数据类型。同样，这确实是一个 C 问题，而不是特定于 CUDA。您可以通过在内核中使用 pitch 的任何位置使用 (pitch/sizeof(int)) 来解决此问题。
6. 您将siz 传递给您的内核。您应该将 pitch 传递给 pitch 参数。 siz 实际上是主机数据存储上的“间距”，但 pitch 是设备上存储的间距。内核在设备存储上运行，因此需要正确的音高。
7. 作为建议，对所有 cuda API 调用和内核调用执行 cuda error checking。

这里有一些代码以一种或另一种方式解决了上述所有问题：

#include<stdio.h>
#define siz (4*sizeof(int))

#define cudaCheckErrors(msg) \
    do { \
        cudaError_t __err = cudaGetLastError(); \
        if (__err != cudaSuccess) { \
            fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
                msg, cudaGetErrorString(__err), \
                __FILE__, __LINE__); \
            fprintf(stderr, "*** FAILED - ABORTING\n"); \
            exit(1); \
        } \
    } while (0)

__global__ void addmatrix(int *m1,int *m2,size_t pitch)
{
    int r=threadIdx.x;
    int *r1=m1+r*(pitch/sizeof(int));
    int *r2=m2+r*(pitch/sizeof(int));
    int c;
    for(c=0;c<4;c++)
    {
        r1[c]+=r2[c];
    }
}
int main()
{
    int i,j;
    int *m1_c,*m2_c;
    int *m1_d,*m2_d;
    size_t pitch;
    cudaError_t err;
    m1_c=(int *)malloc(16*sizeof(int));
    m2_c=(int *)malloc(16*sizeof(int));
    for(i=0;i<4;i++)
    {
        for(j=0;j<4;j++)
        {
            m1_c[(i*4)+j]=rand()%10;
            m2_c[(i*4)+j]=rand()%10;
        }
    }
    for(i=0;i<4;i++)
    {
        for(j=0;j<4;j++)
        {
            printf("%d\t",m1_c[(i*4)+j]);
        }
        printf("\n");
    }
    printf("\n\n");
    for(i=0;i<4;i++)
    {
        for(j=0;j<4;j++)
        {
            printf("%d\t",m2_c[(i*4)+j]);
        }
        printf("\n");
    }
    err=cudaMallocPitch((void **)&m1_d,&pitch,siz,4);
    cudaCheckErrors("cm1");
    err=cudaMallocPitch((void **)&m2_d,&pitch,siz,4);
    cudaCheckErrors("cm2");
    err=cudaMemcpy2D(m1_d,pitch,m1_c,siz,siz,4,cudaMemcpyHostToDevice);
    cudaCheckErrors("cm3");
    err=cudaMemcpy2D(m2_d,pitch,m2_c,siz,siz,4,cudaMemcpyHostToDevice);
    cudaCheckErrors("cm4");
    dim3 grid(1);
    dim3 block(4);
    addmatrix<<<grid,block>>>(m1_d,m2_d,pitch);
    cudaMemcpy2D(m1_c,siz,m1_d,pitch,siz,4,cudaMemcpyDeviceToHost);
    cudaCheckErrors("cm5");

    for(i=0;i<4;i++)
    {
        for(j=0;j<4;j++)
        {
            printf("%d\t",m1_c[(i*4)+j]);
        }
        printf("\n");
    }
    err=cudaFree(m1_d);
    err=cudaFree(m2_d);
    err=cudaDeviceReset();
}