Asignación de memoria en GPU para matriz dinámica de estructuras

5 minutos de lectura

Asignacion de memoria en GPU para matriz dinamica de estructuras
Bakus123

Tengo un problema al pasar una matriz de estructura al kernel gpu. Me basé en este tema: falla de segmentación de cudaMemcpy y escribí algo así:

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

struct Test {
    char *array;
};

__global__ void kernel(Test *dev_test) {
    for(int i=0; i < 5; i++) {
        printf("Kernel[0][i]: %c \n", dev_test[0].array[i]);
    }
}

int main(void) {

    int n = 4, size = 5;
    Test *dev_test, *test;

    test = (Test*)malloc(sizeof(Test)*n);
    for(int i = 0; i < n; i++)
        test[i].array = (char*)malloc(size * sizeof(char));

    for(int i=0; i < n; i++) {
        char temp[] = { 'a', 'b', 'c', 'd' , 'e' };
        memcpy(test[i].array, temp, size * sizeof(char));
    }

    cudaMalloc((void**)&dev_test, n * sizeof(Test));
    cudaMemcpy(dev_test, test, n * sizeof(Test), cudaMemcpyHostToDevice);
    for(int i=0; i < n; i++) {
        cudaMalloc((void**)&(test[i].array), size * sizeof(char));
        cudaMemcpy(&(dev_test[i].array), &(test[i].array), size * sizeof(char), cudaMemcpyHostToDevice);
    }

    kernel<<<1, 1>>>(dev_test);
    cudaDeviceSynchronize();

    //  memory free
    return 0;
}

No hay ningún error, pero los valores mostrados en el kernel son incorrectos. ¿Qué estoy haciendo mal? Gracias de antemano por cualquier ayuda.

  • ¿Por qué es cudaMalloc((void**)&(test[i].array), size * sizeof(char)); y no cudaMalloc((void**)&(dev_test[i].array), size * sizeof(char)); ? Además, debería ser cudaMemcpy(dev_test[i].array, test[i].array, size * sizeof(char), cudaMemcpyHostToDevice);.

    – francisco

    6 mayo 2015 a las 17:02


  • @francis, no funciona (falla de segmentación (núcleo descargado)). En gpu no podemos asignar memoria de forma estándar.

    – Bakus123

    6 mayo 2015 a las 17:06


  • Consejo amistoso adicional: no elija el código de una pregunta, excepto si ha entendido el problema que enfrenta el autor de la pregunta… Lo siento si mi sugerencia no funcionó. Mi sugerencia fue asignar memoria para dev_test[i].arrayno para test[i].array que ya está asignado en la CPU por test[i].array = (char*)malloc(size * sizeof(char));.

    – francisco

    6 mayo 2015 a las 17:15


  • @francis, está bien, no hay problema. sí test[i].array ya está asignado pero solo en la CPU, no en la GPU. No podemos asignar memoria para dev_test[i].array, porque esta memoria solo es visible desde el dispositivo. Al menos yo lo entiendo así.

    – Bakus123

    6 mayo 2015 a las 17:26


1646955967 838 Asignacion de memoria en GPU para matriz dinamica de estructuras
Roberto Crovella

  1. Esto está asignando un nuevo puntero a la memoria del host:

     test[i].array = (char*)malloc(size * sizeof(char));
    
  2. Esto está copiando datos a esa región en la memoria del host:

     memcpy(test[i].array, temp, size * sizeof(char));
    
  3. Esto es sobrescribiendo el puntero previamente asignado a la memoria del host (del paso 1 anterior) con un nuevo puntero a la memoria del dispositivo:

     cudaMalloc((void**)&(test[i].array), size * sizeof(char));
    

Después del paso 3, los datos que configuró en el paso 2 se pierden por completo y ya no se puede acceder a ellos de ninguna manera. Refiriéndose a los pasos 3 y 4 en la pregunta/respuesta que vinculó:

3. Cree un puntero int separado en el host, llamémoslo myhostptr

4.cudaMalloc int almacenamiento en el dispositivo para myhostptr

No has hecho esto. No creó un puntero separado. Reutilizó (borró, sobrescribió) un puntero existente, que apuntaba a los datos que le interesaban en el host. Esta pregunta / respuesta, también vinculada a la respuesta que vinculó, brinda casi exactamente los pasos que debe seguir, en codigo.

Aquí hay una versión modificada de su código, que implementa correctamente los pasos faltantes 3 y 4 (y 5) que no implementó correctamente de acuerdo con la pregunta/respuesta que vinculó: (consulte los comentarios que describen los pasos 3,4,5)

$ cat t755.cu
#include <stdio.h>
#include <stdlib.h>

struct Test {
    char *array;
};

__global__ void kernel(Test *dev_test) {
    for(int i=0; i < 5; i++) {
        printf("Kernel[0][i]: %c \n", dev_test[0].array[i]);
    }
}

int main(void) {

    int n = 4, size = 5;
    Test *dev_test, *test;

    test = (Test*)malloc(sizeof(Test)*n);
    for(int i = 0; i < n; i++)
        test[i].array = (char*)malloc(size * sizeof(char));

    for(int i=0; i < n; i++) {
        char temp[] = { 'a', 'b', 'c', 'd' , 'e' };
        memcpy(test[i].array, temp, size * sizeof(char));
    }

    cudaMalloc((void**)&dev_test, n * sizeof(Test));
    cudaMemcpy(dev_test, test, n * sizeof(Test), cudaMemcpyHostToDevice);

    // Step 3:
    char *temp_data[n];
    // Step 4:
    for (int i=0; i < n; i++)
      cudaMalloc(&(temp_data[i]), size*sizeof(char));
    // Step 5:
    for (int i=0; i < n; i++)
      cudaMemcpy(&(dev_test[i].array), &(temp_data[i]), sizeof(char *), cudaMemcpyHostToDevice);
    // now copy the embedded data:
    for (int i=0; i < n; i++)
      cudaMemcpy(temp_data[i], test[i].array, size*sizeof(char), cudaMemcpyHostToDevice);

    kernel<<<1, 1>>>(dev_test);
    cudaDeviceSynchronize();

    //  memory free
    return 0;
}

$ nvcc -o t755 t755.cu
$ cuda-memcheck ./t755
========= CUDA-MEMCHECK
Kernel[0][i]: a
Kernel[0][i]: b
Kernel[0][i]: c
Kernel[0][i]: d
Kernel[0][i]: e
========= ERROR SUMMARY: 0 errors
$

Dado que la metodología anterior puede ser un desafío para los principiantes, el consejo habitual es no hacerlo, sino aplanar sus estructuras de datos. Aplanar generalmente significa reorganizar el almacenamiento de datos para eliminar los punteros incrustados que deben asignarse por separado.

Un ejemplo trivial de aplanar esta estructura de datos sería usar esto en su lugar:

struct Test {
    char array[5];
};

Se reconoce, por supuesto, que esto especial El enfoque no serviría para todos los propósitos, pero debería ilustrar la idea/intención general. Con esa modificación, como ejemplo, el código se vuelve mucho más simple:

$ cat t755.cu
#include <stdio.h>
#include <stdlib.h>

struct Test {
    char array[5];
};

__global__ void kernel(Test *dev_test) {
    for(int i=0; i < 5; i++) {
        printf("Kernel[0][i]: %c \n", dev_test[0].array[i]);
    }
}

int main(void) {

    int n = 4, size = 5;
    Test *dev_test, *test;

    test = (Test*)malloc(sizeof(Test)*n);

    for(int i=0; i < n; i++) {
        char temp[] = { 'a', 'b', 'c', 'd' , 'e' };
        memcpy(test[i].array, temp, size * sizeof(char));
    }

    cudaMalloc((void**)&dev_test, n * sizeof(Test));
    cudaMemcpy(dev_test, test, n * sizeof(Test), cudaMemcpyHostToDevice);

    kernel<<<1, 1>>>(dev_test);
    cudaDeviceSynchronize();

    //  memory free
    return 0;
}
$ nvcc -o t755 t755.cu
$ cuda-memcheck ./t755
========= CUDA-MEMCHECK
Kernel[0][i]: a
Kernel[0][i]: b
Kernel[0][i]: c
Kernel[0][i]: d
Kernel[0][i]: e
========= ERROR SUMMARY: 0 errors
$

  • muchas gracias. ¿Qué quiere decir con “aplanar sus estructuras de datos”?

    – Bakus123

    6 mayo 2015 a las 18:24


  • actualicé mi respuesta para responder a esta pregunta. Sin embargo, si busca en la etiqueta CUDA, encontrará muchas referencias y ejemplos de “aplanamiento”.

    – Roberto Crovella

    6 mayo 2015 a las 18:43

¿Ha sido útil esta solución?

Esta web utiliza cookies propias y de terceros para su correcto funcionamiento y para fines analíticos y para mostrarte publicidad relacionada con sus preferencias en base a un perfil elaborado a partir de tus hábitos de navegación. Al hacer clic en el botón Aceptar, acepta el uso de estas tecnologías y el procesamiento de tus datos para estos propósitos. Configurar y más información
Privacidad