¿Cómo asignar matrices dinámicamente dentro de un kernel?

11 minutos de lectura

¿Como asignar matrices dinamicamente dentro de un kernel
Granada

Necesito asignar dinámicamente algunas matrices dentro de la función del núcleo. ¿Cómo puedo hacer eso?

Mi código es algo así:

__global__ func(float *grid_d,int n, int nn){  
    int i,j;  
    float x[n],y[nn];  
    //Do some really cool and heavy computations here that takes hours.  
}

Pero eso no funcionará. Si esto estuviera dentro del código del host, podría usar malloc. cudaMalloc necesita un puntero en el host y otro en el dispositivo. Dentro de la función del kernel no tengo el puntero del host.

¿Entonces qué debo hacer?

Si toma demasiado tiempo (algunos segundos) asignar todas las matrices (necesito alrededor de 4 de tamaño n y 5 de tamaño nn), esto no será un problema. Dado que el kernel probablemente se ejecutará durante 20 minutos, al menos.

  • Probablemente quiera leer la sección sobre asignación de memoria dinámica en el código del dispositivo en el Guía de programadores de CUDA C. Esta capacidad requiere capacidad de cómputo 2.0 o superior en su GPU.

    – Roberto Crovella

    20 de noviembre de 2012 a las 19:10

  • ¿Cuál es la configuración (bloques, subprocesos) en la que ejecutará este núcleo? ¿Cuáles son los rangos típicos de n y nn (para tamaños pequeños, puede comprimirlos en registros o memoria compartida).

    – P. Marecki

    22 de noviembre de 2012 a las 11:03


¿Como asignar matrices dinamicamente dentro de un kernel
garras

La asignación de memoria dinámica solo se admite en la capacidad de cómputo 2.x y hardware más reciente. Puede usar la palabra clave new de C++ o malloc en el kernel, por lo que su ejemplo podría convertirse en:

__global__ func(float *grid_d,int n, int nn){  
    int i,j;  
    float *x = new float[n], *y = new float[nn];   
}

Esto asigna memoria en un montón de tiempo de ejecución de memoria local que tiene la duración del contexto, así que asegúrese de liberar la memoria después de que el kernel termine de ejecutarse si su intención es no usar la memoria nuevamente. También debe tener en cuenta que no se puede acceder a la memoria del montón en tiempo de ejecución directamente desde las API del host, por lo que no puede pasar un puntero asignado dentro de un núcleo como argumento para cudaMemcpypor ejemplo.

  • Tengo una situación similar en la que necesito tener matrices asignadas dinámicamente. Cada subproceso debe acceder a esas matrices para fines de escritura. Estoy confundido de que si declaro este proceso de asignación dinámica dentro del kernel, ¿crearía 4 veces tales matrices si las dimensiones del kernel son (1,4), es decir, nThreads = 4 y nBlocks = 1?

    – skm

    29 de julio de 2015 a las 12:15

  • Es free apropiado aquí, o hay otra función para liberar del montón local dentro de un núcleo?

    – landó

    29/10/2015 a las 16:54

  • @landau No, solo usa gratis o elimina

    – garras

    29/10/2015 a las 16:56


1647551469 49 ¿Como asignar matrices dinamicamente dentro de un kernel
roger dahl

@talonmies respondió a su pregunta sobre cómo asignar memoria dinámicamente dentro de un núcleo. Esto pretende ser una respuesta complementaria, abordando el rendimiento de __device__ malloc() y una alternativa que quizás desee considerar.

Asignar memoria dinámicamente en el kernel puede ser tentador porque permite que el código de la GPU se parezca más al código de la CPU. Pero puede afectar seriamente el rendimiento. Escribí una prueba independiente y la he incluido a continuación. La prueba lanza unos 2,6 millones de subprocesos. Cada subproceso llena 16 enteros de la memoria global con algunos valores derivados del índice del subproceso, luego suma los valores y devuelve la suma.

La prueba implementa dos enfoques. El primer enfoque utiliza __device__ malloc() y el segundo enfoque utiliza la memoria que se asigna antes de que se ejecute el núcleo.

En mi dispositivo 2.0, el kernel se ejecuta en 1500ms cuando uso __device__ malloc() y 27 ms cuando se usa memoria preasignada. En otras palabras, la prueba toma 56 veces más para ejecutarse cuando la memoria se asigna dinámicamente dentro del kernel. El tiempo incluye el bucle exterior. cudaMalloc() / cudaFree(), que no es parte del kernel. Si el mismo núcleo se lanza muchas veces con el mismo número de subprocesos, como suele ser el caso, el costo del cudaMalloc() / cudaFree() se amortiza sobre todos los lanzamientos del kernel. Eso eleva la diferencia aún más, a alrededor de 60x.

Especulando, creo que el impacto en el rendimiento se debe en parte a la serialización implícita. La GPU probablemente debe serializar todas las llamadas simultáneas a __device__ malloc() para proporcionar fragmentos separados de memoria a cada llamador.

La versión que no usa __device__ malloc() asigna toda la memoria de la GPU antes de ejecutar el kernel. Un puntero a la memoria se pasa al kernel. Cada subproceso calcula un índice en la memoria previamente asignada en lugar de utilizar un __device__ malloc().

El problema potencial con la asignación de memoria por adelantado es que, si solo algunos subprocesos necesitan asignar memoria y no se sabe qué subprocesos son, será necesario asignar memoria para todos los subprocesos. Si no hay suficiente memoria para eso, podría ser más eficiente reducir la cantidad de subprocesos por llamada al kernel y luego usar __device__ malloc(). Otras soluciones probablemente terminarían reimplementando lo que __device__ malloc() está haciendo en segundo plano, y vería un impacto de rendimiento similar.

Probar el rendimiento de __device__ malloc():

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>

const int N_ITEMS(16);

#define USE_DYNAMIC_MALLOC

__global__ void test_malloc(int* totals)
{
  int tx(blockIdx.x * blockDim.x + threadIdx.x);

  int* s(new int[N_ITEMS]);

  for (int i(0); i < N_ITEMS; ++i) {
    s[i] = tx * i;
  }

  int total(0);
  for (int i(0); i < N_ITEMS; ++i) {
    total += s[i];
  }

  totals[tx] = total;

  delete[] s;
}

__global__ void test_malloc_2(int* items, int* totals)
{
  int tx(blockIdx.x * blockDim.x + threadIdx.x);

  int* s(items + tx * N_ITEMS);

  for (int i(0); i < N_ITEMS; ++i) {
    s[i] = tx * i;
  }

  int total(0);
  for (int i(0); i < N_ITEMS; ++i) {
    total += s[i];
  }

  totals[tx] = total;
}

int main()
{
  cudaError_t cuda_status;

  cudaSetDevice(0);

  int blocks_per_launch(1024 * 10);
  int threads_per_block(256);

  int threads_per_launch(blocks_per_launch * threads_per_block);

  int* totals_d;
  cudaMalloc((void**)&totals_d, threads_per_launch * sizeof(int));

  cudaEvent_t start, stop;
  cudaEventCreate(&start);
  cudaEventCreate(&stop);

  cudaDeviceSynchronize();
  cudaEventRecord(start, 0);

#ifdef USE_DYNAMIC_MALLOC
  cudaDeviceSetLimit(cudaLimitMallocHeapSize, threads_per_launch * N_ITEMS * sizeof(int));

  test_malloc<<<blocks_per_launch, threads_per_block>>>(totals_d);
#else
  int* items_d;
  cudaMalloc((void**)&items_d, threads_per_launch * sizeof(int) * N_ITEMS);

  test_malloc_2<<<blocks_per_launch, threads_per_block>>>(items_d, totals_d);

  cudaFree(items_d);
#endif

  cuda_status = cudaDeviceSynchronize();
  if (cuda_status != cudaSuccess) {
    printf("Error: %d\n", cuda_status);
    exit(1);
  }

  cudaEventRecord(stop, 0);
  cudaEventSynchronize(stop);
  float elapsedTime;
  cudaEventElapsedTime(&elapsedTime, start, stop);

  printf("Elapsed: %f\n", elapsedTime);

  int* totals_h(new int[threads_per_launch]);
  cuda_status = cudaMemcpy(totals_h, totals_d, threads_per_launch * sizeof(int), cudaMemcpyDeviceToHost);
  if (cuda_status != cudaSuccess) {
    printf("Error: %d\n", cuda_status);
    exit(1);
  }

  for (int i(0); i < 10; ++i) {
    printf("%d ", totals_h[i]);
  }
  printf("\n");

  cudaFree(totals_d);
  delete[] totals_h;

  return cuda_status;
}

Producción:

C:\rd\projects\test_cuda_malloc\Release>test_cuda_malloc.exe
Elapsed: 27.311169
0 120 240 360 480 600 720 840 960 1080

C:\rd\projects\test_cuda_malloc\Release>test_cuda_malloc.exe
Elapsed: 1516.711914
0 120 240 360 480 600 720 840 960 1080

  • Debe cronometrar el cudaMalloc en la segunda prueba. De lo contrario, está comparando un automóvil listo para funcionar (segunda prueba) con un automóvil detenido en un garaje (primera prueba). Ambos núcleos necesitan los mismos requisitos de almacenamiento.

    – pQB

    21 de noviembre de 2012 a las 7:16

  • Además de la objeción pQB: su cudaMalloc asigna una matriz grande, y esto se compara con la asignación de 2,5 millones de matrices pequeñas (para cada subproceso uno). Tal procedimiento es, por supuesto, más lento, y una prueba en la CPU muestra que su desaceleración de 60x informada es en realidad un buen trabajo (obtengo una desaceleración de 1000x veces, siempre que el código no se segfault: el asignador necesita manejar tantas matrices). La prueba justa es: asignar la misma (una) matriz, (1) por cudaMalloc(2) por kernel<<<1,1>>>. veo el kernel la asignación es más lenta ~ 3 veces. Así que este es el verdadero éxito de rendimiento.

    – P. Marecki

    21 de noviembre de 2012 a las 9:08

  • @pQB: Gracias. Había dejado el cudaMalloc() fuera del tiempo, asumiendo que no sería medible. Para mi sorpresa, agregarlo provocó un cambio, pasando de 60x a 56x. Actualicé la respuesta y agregué una reseña sobre las implicaciones de incluir cudaMalloc() / cudaFree() en el tiempo.

    –Roger Dahl

    21 de noviembre de 2012 a las 16:07

  • @PMarecki: El propósito de la prueba era mostrar las implicaciones de rendimiento del uso __device__ malloc() y para mostrar una forma alternativa de llevar a cabo la tarea para la que muchos considerarían __device__ malloc(). El propósito no era comparar el desempeño de un solo cudaMalloc() con un solo __device__ malloc().

    –Roger Dahl

    21 de noviembre de 2012 a las 16:14

  • ¡Prueba ordenada de @RogerDahl! Creo que el punto principal es mostrar la diferencia en la asignación de muchas matrices pequeñas, ya sea en el dispositivo o en el host. Pero, independientemente, con el mismo número de llamadas malloc. Creo que “por supuesto” una sola llamada malloc será más rápida que muchas llamadas individuales malloc.

    – parteinteresada333

    10 abr 2018 a las 20:02

Si el valor de n y nn se conocía antes de llamar al kernel, ¿por qué no cudaMalloc la memoria en el lado del host y pasa el puntero de memoria del dispositivo al kernel?

  • Porque cada kernel debe poseer una matriz.

    – Granada

    21 de noviembre de 2012 a las 9:56


  • ¿Está lanzando múltiples criaderos al mismo tiempo? ¿No podría asignar suficiente espacio y cada kernel solo comparte parte de él?

    – Hong Zhou

    22 de noviembre de 2012 a las 6:12

  • si lanzo, por ejemplo, 1000 kernels y si necesito 10 arreglos de tamaño n. ¿Debería hacer 10 arreglos de tamaño n*1000? ¿Y compartir esto entre los núcleos usando threadid y blockid?

    – Granada

    22 de noviembre de 2012 a las 22:47

1647551470 160 ¿Como asignar matrices dinamicamente dentro de un kernel
parteinteresada333

Realicé un experimento basado en los conceptos de la publicación de @rogerdahl. Suposiciones:

  • 4 MB de memoria asignada en fragmentos de 64B.
  • 1 bloque de GPU y 32 subprocesos warp en ese bloque
  • Ejecutar en un P100

Las llamadas malloc+free locales a la GPU parecían ser mucho más rápidas que las cudaMalloc + cudaFree llamadas La salida del programa:

Starting timer for cuda malloc timer
Stopping timer for cuda malloc timer
         timer for cuda malloc timer took 1.169631s
Starting timer for device malloc timer
Stopping timer for device malloc timer
         timer for device malloc timer took 0.029794s

Estoy omitiendo el código para timer.h y timer.cpppero aquí está el código para la prueba en sí:

#include "cuda_runtime.h"
#include <stdio.h>
#include <thrust/system/cuda/error.h>

#include "timer.h"

static void CheckCudaErrorAux (const char *, unsigned, const char *, cudaError_t);
#define CUDA_CHECK_RETURN(value) CheckCudaErrorAux(__FILE__,__LINE__, #value, value)

const int BLOCK_COUNT = 1;
const int THREADS_PER_BLOCK = 32;
const int ITERATIONS = 1 << 12;
const int ITERATIONS_PER_BLOCKTHREAD = ITERATIONS / (BLOCK_COUNT * THREADS_PER_BLOCK);

const int ARRAY_SIZE = 64;


void CheckCudaErrorAux (const char *file, unsigned line, const char *statement, cudaError_t err) {
    if (err == cudaSuccess)
        return;
    std::cerr << statement<<" returned " << cudaGetErrorString(err) << "("<<err<< ") at "<<file<<":"<<line << std::endl;
    exit (1);
}

__global__ void mallocai() {
    for (int i = 0; i < ITERATIONS_PER_BLOCKTHREAD; ++i) {
        int * foo;
        foo = (int *) malloc(sizeof(int) * ARRAY_SIZE);
        free(foo);
    }
}

int main() {

    Timer cuda_malloc_timer("cuda malloc timer");

    for (int i = 0; i < ITERATIONS; ++ i) {
        if (i == 1) cuda_malloc_timer.start(); // let it warm up one cycle
        int * foo;
        cudaMalloc(&foo, sizeof(int) * ARRAY_SIZE);
        cudaFree(foo);
    }
    cuda_malloc_timer.stop_and_report();
    CUDA_CHECK_RETURN(cudaDeviceSynchronize());

    Timer device_malloc_timer("device malloc timer");
    device_malloc_timer.start();
    mallocai<<<BLOCK_COUNT, THREADS_PER_BLOCK>>>();
    CUDA_CHECK_RETURN(cudaDeviceSynchronize());
    device_malloc_timer.stop_and_report();
}

Si encuentra errores, por favor déjelos en los comentarios e intentaré corregirlos.

Y los volví a ejecutar con todo más grande:

const int BLOCK_COUNT = 56;
const int THREADS_PER_BLOCK = 1024;
const int ITERATIONS = 1 << 18;
const int ITERATIONS_PER_BLOCKTHREAD = ITERATIONS / (BLOCK_COUNT * THREADS_PER_BLOCK);

const int ARRAY_SIZE = 1024;

Y cudaMalloc fue todavía mucho más lento:

Starting timer for cuda malloc timer
Stopping timer for cuda malloc timer
         timer for cuda malloc timer took 74.878016s
Starting timer for device malloc timer
Stopping timer for device malloc timer
         timer for device malloc timer took 0.167331s

1647551470 531 ¿Como asignar matrices dinamicamente dentro de un kernel
tirando

Tal vez deberías probar

cudaMalloc(&foo,sizeof(int) * ARRAY_SIZE * ITERATIONS);
cudaFree(foo);

en lugar de

for (int i = 0; i < ITERATIONS; ++ i) {
    if (i == 1) cuda_malloc_timer.start(); // let it warm up one cycle
    int * foo;
    cudaMalloc(&foo, sizeof(int) * ARRAY_SIZE);
    cudaFree(foo);
}

¿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