CUDA 2D array nvidia

Question

我使用cudaMallocPitch和cudaMemcpy2D为2D数组。我不确定我已经编码正确，即使我不能得到正确的输出。有谁能帮忙吗?谁能调试我的错误吗?提前感谢。

#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("%dt",m1_c[i][j]);
        }
        printf("n");
    }
    printf("nn");
    for(i=1;i<=4;i++)
    {
        for(j=1;j<=4;j++)
        {
            printf("%dt",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("%dt",m1_c[i][j]);
        }
        printf("n");
    }
    err=cudaFree(m1_d);
    err=cudaFree(m2_d);
    err=cudaDeviceReset();      
}

Answer 1

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

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

下面的代码以一种或另一种方式解决了上述所有问题:

#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 - ABORTINGn"); 
            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("%dt",m1_c[(i*4)+j]);
        }
        printf("n");
    }
    printf("nn");
    for(i=0;i<4;i++)
    {
        for(j=0;j<4;j++)
        {
            printf("%dt",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("%dt",m1_c[(i*4)+j]);
        }
        printf("n");
    }
    err=cudaFree(m1_d);
    err=cudaFree(m2_d);
    err=cudaDeviceReset();
}