如何使用嵌套的 for 循环将两个 2d（倾斜）数组相加？

Posted 2023-02-23

【中文标题】如何使用嵌套的 for 循环将两个 2d（倾斜）数组相加？

【英文标题】：How can I add up two 2d (pitched) arrays using nested for loops?

【发布时间】：2011-09-02 11:50:20

【问题描述】：

我是 cuda 的新手。我想将两个二维数组加到第三个数组中。 我使用以下代码：

cudaMallocPitch((void**)&device_a, &pitch, 2*sizeof(int),2);
cudaMallocPitch((void**)&device_b, &pitch, 2*sizeof(int),2);
cudaMallocPitch((void**)&device_c, &pitch, 2*sizeof(int),2);

现在我的问题是我不想将这些数组用作扁平二维数组 我想在我的内核代码中使用两个 for 循环并将结果放在第三个数组中，例如

__global__ void add(int *dev_a ,int *dev_b,int* dec_c)

    for i=0;i<2;i++)
     
      for j=0;j<2;j++)
      
        dev_c[i][j]=dev_a[i][j]+dev_b[i][j];
      
    

我如何在 CUDA 中做到这一点？ 请告诉我如何以这种方式使用二维数组？

使用 2d-array 的内核调用应该是什么？ 如果可能，请使用代码示例进行说明。

【问题讨论】：

请格式化代码 - 缩进 4 个空格。

【参考方案1】：

简短的回答是，你不能。 cudaMallocPitch()函数正如它的名字所暗示的那样，它分配有间距的线性内存，其中间距被选择为GPU内存控制器和纹理硬件的最佳选择。

如果您想在内核中使用指针数组，内核代码必须如下所示：

__global___ void add(int *dev_a[] ,int *dev_b[], int* dec_c[])

    for i=0;i<2;i++)  
      for j=0;j<2;j++) 
        dev_c[i][j]=dev_a[i][j]+dev_b[i][j];
      
    

然后你需要在主机端嵌套cudaMalloc 调用来构造指针数组并将其复制到设备内存。对于您相当简单的 2x2 示例，分配单个数组的代码如下所示：

int ** h_a = (int **)malloc(2 * sizeof(int *));
cudaMalloc((void**)&h_a[0], 2*sizeof(int));
cudaMalloc((void**)&h_a[1], 2*sizeof(int));

int **d_a;
cudaMalloc((void ***)&d_a, 2 * sizeof(int *));
cudaMemcpy(d_a, h_a, 2*sizeof(int *), cudaMemcpyHostToDevice);

这会将分配的设备指针数组留在 d_a 中，然后您会将其传递给内核。

出于代码复杂性和性能原因，您真的不想这样做，在 CUDA 代码中使用指针数组比使用线性内存的替代方案更难和。

---

为了说明在 CUDA 中使用指针数组有什么愚蠢之处，这里有一个完整的示例问题，它结合了上述两个想法：

#include <cstdio>
__global__ void add(int * dev_a[], int * dev_b[], int * dev_c[])

    for(int i=0;i<2;i++)
     
        for(int j=0;j<2;j++)
        
            dev_c[i][j]=dev_a[i][j]+dev_b[i][j];
        
    


inline void GPUassert(cudaError_t code, char * file, int line, bool Abort=true)

    if (code != 0) 
        fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code),file,line);
        if (Abort) exit(code);
           


#define GPUerrchk(ans)  GPUassert((ans), __FILE__, __LINE__); 

int main(void)

    const int aa[2][2]=1,2,3,4;
    const int bb[2][2]=5,6,7,8;
    int cc[2][2];

    int ** h_a = (int **)malloc(2 * sizeof(int *));
    for(int i=0; i<2;i++)
        GPUerrchk(cudaMalloc((void**)&h_a[i], 2*sizeof(int)));
        GPUerrchk(cudaMemcpy(h_a[i], &aa[i][0], 2*sizeof(int), cudaMemcpyHostToDevice));
    

    int **d_a;
    GPUerrchk(cudaMalloc((void ***)&d_a, 2 * sizeof(int *)));
    GPUerrchk(cudaMemcpy(d_a, h_a, 2*sizeof(int *), cudaMemcpyHostToDevice));

    int ** h_b = (int **)malloc(2 * sizeof(int *));
    for(int i=0; i<2;i++)
        GPUerrchk(cudaMalloc((void**)&h_b[i], 2*sizeof(int)));
        GPUerrchk(cudaMemcpy(h_b[i], &bb[i][0], 2*sizeof(int), cudaMemcpyHostToDevice));
    

    int ** d_b;
    GPUerrchk(cudaMalloc((void ***)&d_b, 2 * sizeof(int *)));
    GPUerrchk(cudaMemcpy(d_b, h_b, 2*sizeof(int *), cudaMemcpyHostToDevice));

    int ** h_c = (int **)malloc(2 * sizeof(int *));
    for(int i=0; i<2;i++)
        GPUerrchk(cudaMalloc((void**)&h_c[i], 2*sizeof(int)));
    

    int ** d_c;
    GPUerrchk(cudaMalloc((void ***)&d_c, 2 * sizeof(int *)));
    GPUerrchk(cudaMemcpy(d_c, h_c, 2*sizeof(int *), cudaMemcpyHostToDevice));

    add<<<1,1>>>(d_a,d_b,d_c);
    GPUerrchk(cudaPeekAtLastError());

    for(int i=0; i<2;i++)
        GPUerrchk(cudaMemcpy(&cc[i][0], h_c[i], 2*sizeof(int), cudaMemcpyDeviceToHost));
    

    for(int i=0;i<2;i++) 
        for(int j=0;j<2;j++) 
            printf("(%d,%d):%d\n",i,j,cc[i][j]);
        
    

    return cudaThreadExit();

我建议你研究它，直到你理解它的作用，以及为什么它与使用线性内存相比是一个糟糕的主意。

【讨论】：

是的，你是对的。现在假设我这样做应该是我的内核调用

谢谢。是的，你是对的。现在假设我这样做应该是我的内核调用吗？我会说我使用 cudaMalloc((void ***)&d_a, 2 * sizeof(int *));但它显示错误； h_a 的另一件事为什么你使用 cuda malloc 请解释一下

【参考方案2】：

您不需要在设备内部使用 for 循环。试试这个代码。

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

#define N 800
__global__ void  matrixAdd(float* A, float* B, float* C)

int i = threadIdx.x;
int j = blockIdx.x;
C[N*j+i] = A[N*j+i] + B[N*j+i];


int main (void) 
clock_t start = clock();
float a[N][N], b[N][N], c[N][N];
float *dev_a, *dev_b, *dev_c;

cudaMalloc((void **)&dev_a, N * N * sizeof(float));
cudaMalloc((void **)&dev_b, N * N * sizeof(float));
cudaMalloc((void **)&dev_c, N * N * sizeof(float));

for (int i = 0; i < N; i++)
    for (int j = 0; j < N; j++)    
        a[i][j] = rand() % 10;
        b[i][j] = rand() % 10;
    


cudaMemcpy(dev_a, a, N * N * sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(dev_b, b, N * N * sizeof(float), cudaMemcpyHostToDevice);

matrixAdd <<<N,N>>> (dev_a, dev_b, dev_c);
cudaMemcpy(c, dev_c, N * N * sizeof(float), cudaMemcpyDeviceToHost);

for (int i = 0; i < N; i++)
    for (int j = 0; j < N; j++)
    printf("[%d, %d ]= %f + %f = %f\n",i,j, a[i][j], b[i][j], c[i][j]);
    

printf("Time elapsed: %f\n", ((double)clock() - start) / CLOCKS_PER_SEC);

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

return 0; 

【讨论】：

这仅适用于静态分配的数组，其中尺寸在编译时是已知的。对于问题中指出的任何类型的动态分配（例如cudaMalloc 等），这将不起作用。

顺便说一下，在将数组传递给内核之前，您仍在展平数组，这不是相关用户想要的。