Necesito para asignar dinámicamente algunas matrices en el interior del núcleo de la función. ¿Cómo puede un yo que?

Mi código es algo así como:

__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 que no va a funcionar. Si esto fue en el host de código podría utilizar malloc. cudaMalloc necesita un puntero en el host, y otros en el dispositivo. En el interior del núcleo de la función no tengo el host puntero.

Así que, ¿qué debo hacer?

Si toma demasiado tiempo (algunos segundos) para asignar todas las matrices (necesito unos 4 de tamaño n y 5 de tamaño nn), esto no será un problema. Dado que el núcleo probablemente una duración de 20 minutos, al menos.

  • Probablemente te interese leer la sección de asignación dinámica de memoria en el código del dispositivo en el CUDA C guía de programadores. Esta capacidad requiere de la capacidad de cálculo de 2.0 o mayor en la GPU.
  • ¿Cuál es la configuración (bloques de hilos) que va a ejecutar este kernel? ¿Cuáles son los rangos típicos de n y nn (para los tamaños pequeños puede exprimir en registros o memoria compartida).
InformationsquelleAutor Granada | 2012-11-20

5 Comentarios

  1. 28

    Asignación dinámica de memoria sólo se admite en capacidad de cálculo 2.x y hardware más reciente. Usted puede utilizar el C++ nueva palabra clave o malloc en el núcleo, por lo que su ejemplo podría ser:

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

    Este asigna memoria en un local de la memoria en tiempo de ejecución del montón que tiene la vida útil del contexto, así que asegúrese de liberar la memoria después de que el núcleo se termina de ejecutar si su intención no es la de utilizar la memoria de nuevo. También debe tener en cuenta que el tiempo de ejecución de la memoria heap no se puede acceder directamente desde el host Api, por lo que no puede pasar un puntero asignado dentro de un núcleo como un argumento para cudaMemcpy, por ejemplo.

    • Tengo una situación similar donde necesito tener asignada dinámicamente matrices. Dichas matrices se tiene que acceder por cada subproceso para escribir propósito. Estoy confundido que si declaro este dinámico proceso de asignación en el interior del núcleo, crearía 4 veces tales matrices, si las dimensiones de kernel (1,4), es decir, nThreads = 4 y nBlocks = 1.
    • Es free apropiado aquí, o hay otra función para la liberación del montón local dentro de un núcleo?
    • No acaba de uso libre o eliminar
  2. 13

    @talonmies contestado a su pregunta sobre cómo asignar memoria dinámicamente dentro de un núcleo. Esto está pensado como una forma complementaria de respuesta, abordando el rendimiento de __device__ malloc() y una alternativa usted puede ser que desee considerar.

    Asignar memoria dinámicamente en el núcleo puede ser tentador, porque permite a la GPU código parece más a la de la CPU de código. Pero puede afectar seriamente el rendimiento. Escribí un auto-contenido de la prueba y ha incluido a continuación. La prueba se inicia alrededor de 2,6 millones de hilos. Cada rosca rellena de 16 enteros de memoria global con algunos de los valores que se derivan de la rosca índice, a continuación, resume los valores y devuelve la suma.

    La prueba implementa dos enfoques. El primer método utiliza __device__ malloc() y el segundo enfoque utiliza la memoria que se asigna antes de que el kernel se ejecuta.

    En mi dispositivo 2.0, el kernel se ejecuta en 1500ms cuando se utiliza __device__ malloc() y 27 ms cuando el uso de pre-memoria asignada. En otras palabras, la prueba se lleva a 56 x no a ejecutar cuando la memoria se asigna de forma dinámica dentro del kernel. El tiempo incluye el bucle exterior cudaMalloc() /cudaFree(), que no es parte del núcleo. Si el mismo kernel es lanzado muchas veces con el mismo número de hilos, como es a menudo el caso, el costo de la cudaMalloc() /cudaFree() es amortizado a lo largo de todo el núcleo de su lanzamiento. Que trae la diferencia aún mayor, alrededor de 60x.

    Especular, creo que el impacto en el rendimiento es consecuencia en parte implícita de la serialización. La GPU debe probablemente serializar todas las llamadas simultáneas a __device__ malloc() con el fin de proporcionar separar fragmentos de memoria para cada persona que llama.

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

    El problema potencial con la asignación de memoria delante es que, si sólo algunos de los subprocesos necesidad de asignar la memoria, y no se sabe que los hilos son, será necesario asignar memoria para todos los subprocesos. Si no hay suficiente memoria para eso, puede ser más eficaz para reducir el número de hilos por núcleo de la llamada, a continuación, utilizando __device__ malloc(). Otras soluciones quizá termine reimplementing lo __device__ malloc() está haciendo en el fondo, y que iba a ver un similar rendimiento.

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

    De 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
    • Debe el tiempo el cudaMalloc en la segunda prueba. De lo contrario, estás comparando un coche listo para correr (segunda prueba) a un vehículo detenido en un garaje (primera prueba). Los dos núcleos necesitan los mismos requisitos de almacenamiento.
    • Además pQB objeción: su cudaMalloc asigna una gran variedad, y esto es en comparación a la asignación de 2.5 millones de pequeñas matrices (para cada subproceso uno). Un procedimiento de este tipo es, por supuesto, más lento, y una prueba de la CPU muestra, que tu informó 60x desaceleración es realmente un buen trabajo (me pongo 1000x tiempos de desaceleración, a condición de que el código no violación de segmento — asignador de necesidades para soportar tantas matrices). Feria de la prueba es: asignar el mismo (uno) de la matriz, (1) por cudaMalloc, (2) por kernel<<<1,1>>>. Veo la kernel asignación ser más lento ~3 veces. Así que este es el verdadero impacto en el rendimiento.
    • Gracias. Yo había dejado la cudaMalloc() fuera de la cronología, suponiendo que no iba a ser mensurables. Para mi sorpresa, agregando que en sí causó un cambio, pasando de 60x a 56 x. He actualizado la respuesta y se añade una propaganda acerca de las implicaciones de la inclusión de la cudaMalloc() / cudaFree() en el tiempo.
    • El propósito de la prueba era mostrar las implicaciones de rendimiento en el uso de __device__ malloc() y a mostrar una forma alternativa de llevar a cabo la tarea para la que muchos consideran __device__ malloc(). El propósito no era el de comparar el rendimiento de un solo cudaMalloc() con una sola __device__ malloc().
    • aseado prueba! Creo que el punto principal es mostrar la diferencia en la asignación de muchas pequeñas matrices, ya sea en el dispositivo o en el de acogida. Pero, independientemente, con el mismo número de llamadas a malloc. Me gustaría pensar que «por supuesto» una sola llamada malloc va a ser más rápido que muchas de las llamadas a malloc.
  3. 2

    Si el valor de n y nn eran conocidos antes de que el kernel se llama, entonces, ¿por qué no cudaMalloc la memoria en el lado del host y de paso en el dispositivo de puntero de memoria para el kernel?

    • Debido a que cada núcleo debe propios de una matriz.
    • Son ustedes la presentación de múltiples kenel simultáneamente? No podía asignar suficiente espacio y a cada núcleo sólo comparte parte de ella?
    • si me abre, por ejemplo, de 1000 granos y si necesito 10 matrices de tamaño n. El yo debe hacer 10 matrices de tamaño n*1000? Y compartir esta a través de los granos mediante threadid y blockid?
  4. 0

    Ejecutó un experimento basado en los conceptos de @rogerdahl del post. Supuestos:

    • 4MB de memoria asignada en 64B trozos.
    • 1 GPU bloque y 32 hilos de urdimbre en ese bloque
    • Ejecutar en un P100

    La malloc+llamadas gratis locales a la GPU, parecía ser mucho más rápido que el cudaMalloc + cudaFree llamadas. El programa de salida:

    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

    Voy a dejar el código para timer.h y timer.cpp, pero 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 usted encuentra errores, por favor, lmk en los comentarios, y voy a tratar de arreglar.

    Y me encontré de nuevo con los más grandes de todo:

    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 aún más lento por mucho:

    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
    • También de la nota, malloc + free tomó básicamente la misma cantidad de tiempo como new y delete.
  5. 0

    Tal vez debería probar

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

    lugar

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

Dejar respuesta

Please enter your comment!
Please enter your name here