loops - two - multidimensional array c#

¿Cómo puedo agregar dos arrays 2d(pitch) usando anidados para bucles?

Question

Soy nuevo en cuda. Quiero agregar dos arreglos de 2d en una tercera matriz. Yo uso el siguiente código:

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);

ahora mi problema es que no quiero usar estos arreglos como una matriz plana de 2 d todo en mi código kernel quiero di es usar dos para bucle y poner el resultado en la tercera matriz como

__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];
      }
    }
}

¿Cómo puedo hacer esto en CUDA? por favor dime cómo usar una matriz de 2 d de esta manera?

¿Cuál debería ser la llamada del kernel para usar 2d-array? Si es posible, explique usando muestras de código.

Answer 1

La respuesta corta es que no puedes. La función cudaMallocPitch() hace exactamente lo que su nombre implica, asigna memoria lineal inclinada, donde se elige el tono óptimo para el controlador de memoria GPU y el hardware de textura.

Si desea utilizar matrices de punteros en el kernel, el código del kernel debería verse así:

__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];
      }
    }
}

y luego necesitaría llamadas cudaMalloc anidadas en el lado del host para construir la matriz de punteros y copiarla en la memoria del dispositivo. Para su ejemplo bastante trivial de 2x2, el código para asignar una única matriz se vería así:

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);

Lo cual dejaría el conjunto de punteros de dispositivos asignados en d_a, y lo pasaría a su kernel.

Por razones de complejidad y rendimiento del código, realmente no desea hacer eso, usar matrices de punteros en el código CUDA es más difícil y más lento que la alternativa que utiliza la memoria lineal.

Para mostrar cuán loco es el uso de matrices de punteros en CUDA, aquí hay un ejemplo completo de trabajo de su problema de muestra que combina las dos ideas anteriores:

#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();
}

Te recomiendo que lo estudies hasta que entiendas lo que hace, y por qué es una idea tan pobre en comparación con el uso de la memoria lineal.

Answer 2

No necesita usar bucles para dentro del dispositivo. Prueba este código

#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; 
}