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

recorrido de una transformacion lineal

Question

Necesito asignar dinámicamente algunas matrices dentro de la función kernel. ¿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 de host, podría usar malloc. cudaMalloc necesita un puntero en el host y otro en el dispositivo. Dentro de la función kernel no tengo el puntero del host.

¿Entonces qué debo hacer?

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

Answer 1

@talonmies respondió su pregunta sobre cómo asignar dinámicamente la memoria dentro de un kernel. Esto tiene la intención de 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 autocontenida y la he incluido a continuación. La prueba lanza unos 2.6 millones de hilos. Cada hilo llena 16 enteros de la memoria global con algunos valores derivados del índice de subprocesos, luego suma los valores y devuelve la suma.

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

En mi dispositivo 2.0, el kernel se ejecuta en 1500ms cuando usa __device__ malloc() y 27ms cuando usa memoria preasignada. En otras palabras, la prueba tarda 56 veces más en ejecutarse cuando la memoria se asigna dinámicamente dentro del kernel. El tiempo incluye el bucle externo cudaMalloc() / cudaFree() , que no es parte del kernel. Si el mismo núcleo se inicia muchas veces con el mismo número de subprocesos, como suele ser el caso, el costo de cudaMalloc() / cudaFree() se amortiza en todos los cudaFree() del kernel. Eso trae la diferencia aún más alta, 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 trozos de memoria separados a cada llamante.

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

El problema potencial con la asignación anticipada de memoria 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 hilos 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ían un rendimiento similar.

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

Salida:

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

Answer 2

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

Answer 3

La asignación dinámica de memoria solo se admite en la capacidad informática 2.x y el hardware más nuevo. Puede usar la palabra clave nueva de C ++ o malloc en el kernel, para que su ejemplo se convierta 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 la memoria local que tiene la vida útil del contexto, así que asegúrese de liberar la memoria después de que el núcleo termine de ejecutarse si su intención es no utilizar la memoria nuevamente. También debe tener en cuenta que no se puede acceder a la memoria dinámica del tiempo de ejecución directamente desde las API de host, por lo que no puede pasar un puntero asignado dentro de un kernel como argumento a cudaMemcpy , por ejemplo.