parallel processing - ultima - ¿Cómo se lanza un kernel CUDA?

tutorial cuda

Question

Creé una aplicación CUDA simple para agregar dos matrices. Está compilando bien. Quiero saber cómo se lanzará el kernel por todos los hilos y cuál será el flujo dentro de CUDA. Quiero decir, de qué manera cada hilo ejecutará cada elemento de las matrices.

Sé que este es un concepto muy básico, pero no sé esto. Estoy confundido con respecto al flujo.

Answer 1

Prueba ''Cuda-gdb'' , que es el depurador CUDA.

Answer 2

Lanzas una grilla de bloques.

Los bloques se asignan indivisiblemente a multiprocesadores (donde el número de bloques en el multiprocesador determina la cantidad de memoria compartida disponible).

Los bloques se dividen en urdimbres. Para una GPU de Fermi que tiene 32 hilos que ejecutan la misma instrucción o que están inactivos (porque se ramificaron, por ejemplo, saliendo de un ciclo antes que los vecinos dentro de la misma curva o no tomando la misma if lo hicieron). En una GPU de Fermi, como máximo dos distorsiones se ejecutan en un multiprocesador a la vez.

Siempre que haya latencia (es decir, bloqueos de ejecución para acceso a memoria o dependencias de datos para completar) se ejecuta otro warp (el número de warps que se ajustan a un multiprocesador - del mismo bloque o de bloques diferentes) está determinado por el número de registros utilizados por cada subproceso y la cantidad de memoria compartida utilizada por a / the block (s)).

Esta programación ocurre de manera transparente. Es decir, no tienes que pensar demasiado en eso. Sin embargo, es posible que desee utilizar los vectores enteros predefinidos threadIdx (¿dónde está mi hilo dentro del bloque?), blockDim (¿qué tan grande es un bloque?), blockIdx (¿dónde está mi bloque en la cuadrícula?) Y gridDim (¿qué tan grande es? ¿la cuadrícula?) para dividir el trabajo (leer: entrada y salida) entre los hilos. También es posible que desee leer cómo acceder efectivamente a los diferentes tipos de memoria (para que se puedan atender varios hilos dentro de una sola transacción), pero eso está alejando el tema.

NSight proporciona un depurador gráfico que le da una buena idea de lo que está sucediendo en el dispositivo una vez que pasó por la jungla de la jerga. Lo mismo ocurre con su generador de perfiles con respecto a las cosas que no verá en el depurador (por ejemplo, razones de bloqueo o presión de la memoria).

Puede sincronizar todos los hilos dentro de la grilla (todo lo que hay) por otro lanzamiento del kernel. Para la ejecución secuencial del kernel sin superposición, no se necesita más sincronización.

Los hilos dentro de una cuadrícula (o una kernel ejecutada, como quiera llamarla) pueden comunicarse a través de la memoria global usando operaciones atómicas (para aritmética) o vallas de memoria apropiadas (para acceso de carga o almacenamiento).

Puede sincronizar todos los subprocesos dentro de un bloque con la instrucción intrínseca __syncthreads() (todos los subprocesos estarán activos después, aunque, como siempre, se pueden ejecutar como máximo dos distorsiones en una GPU de Fermi). Los hilos dentro de un bloque pueden comunicarse a través de la memoria compartida o global usando operaciones atómicas (para aritmética) o vallas de memoria apropiadas (para acceso de carga o tienda).

Como se mencionó anteriormente, todos los hilos dentro de un warp siempre están "sincronizados", aunque algunos podrían estar inactivos. Se pueden comunicar a través de la memoria compartida o global (o "intercambio de carril" en el próximo hardware con capacidad de cálculo 3). Puede usar operaciones atómicas (para aritmética) y variables compartidas o globales volátiles (el acceso a la carga o a la tienda se realiza secuencialmente dentro de la misma urdimbre). El calificador volátil le dice al compilador que siempre acceda a la memoria y nunca se registra cuyo estado no puede ser visto por otros hilos.

Además, existen funciones de votación en toda la disformidad que pueden ayudarlo a tomar decisiones de sucursal o calcular sumas de enteros (prefijos).

OK, eso es básicamente eso. Espero que ayude. Tenía un buen flujo de escritura :-).

Answer 3

Tomemos un ejemplo de suma de matrices 4 * 4 ... usted tiene dos matrices A y B, que tienen una dimensión 4 * 4.

int main()
{
 int *a, *b, *c;            //To store your matrix A & B in RAM. Result will be stored in matrix C
 int *ad, *bd, *cd;         // To store matrices into GPU''s RAM. 
 int N =4;                 //No of rows and columns.

 size_t size=sizeof(float)* N * N;

 a=(float*)malloc(size);     //Allocate space of RAM for matrix A
 b=(float*)malloc(size);     //Allocate space of RAM for matrix B

//allocate memory on device
  cudaMalloc(&ad,size);
  cudaMalloc(&bd,size);
  cudaMalloc(&cd,size);

//initialize host memory with its own indices
    for(i=0;i<N;i++)
      {
    for(j=0;j<N;j++)
         {
            a[i * N + j]=(float)(i * N + j);
            b[i * N + j]= -(float)(i * N + j);
         }
      }

//copy data from host memory to device memory
     cudaMemcpy(ad, a, size, cudaMemcpyHostToDevice);
     cudaMemcpy(bd, b, size, cudaMemcpyHostToDevice);

//calculate execution configuration 
   dim3 grid (1, 1, 1); 
   dim3 block (16, 1, 1);

//each block contains N * N threads, each thread calculates 1 data element

    add_matrices<<<grid, block>>>(ad, bd, cd, N);

   cudaMemcpy(c,cd,size,cudaMemcpyDeviceToHost);  
   printf("Matrix A was---/n");
    for(i=0;i<N;i++)
    {
        for(j=0;j<N;j++)
            printf("%f ",a[i*N+j]);
        printf("/n");
    }

   printf("/nMatrix B was---/n");
   for(i=0;i<N;i++)
    {
        for(j=0;j<N;j++)
            printf("%f ",b[i*N+j]);
        printf("/n");
    }

    printf("/nAddition of A and B gives C----/n");
    for(i=0;i<N;i++)
    {
        for(j=0;j<N;j++)
            printf("%f ",c[i*N+j]);   //if correctly evaluated, all values will be 0
        printf("/n");
    }



    //deallocate host and device memories
    cudaFree(ad); 
    cudaFree(bd); 
    cudaFree (cd);

    free(a);
    free(b);
    free(c);

    getch();
    return 1;
}

/////Kernel Part

__global__ void add_matrices(float *ad,float *bd,float *cd,int N)
{
  int index;
  index = blockIDx.x * blockDim.x + threadIDx.x            

  cd[index] = ad[index] + bd[index];
}

Tomemos un ejemplo de suma de 16 * 16 matrices ... tiene dos matrices A y B, que tienen una dimensión 16 * 16.

En primer lugar, debe decidir la configuración de su hilo. Se supone que debes iniciar una función de kernel, que realizará el cálculo paralelo de tu adición de matriz, que se ejecutará en tu dispositivo GPU.

Ahora, se lanza una cuadrícula con una función de kernel. Una cuadrícula puede tener un máximo de 65.535 no de bloques que se pueden organizar de 3 formas tridimensionales. (65535 * 65535 * 65535).

Cada bloque en la cuadrícula puede tener un máximo de 1024 no de hilos. Estos hilos también se pueden organizar en 3 formas tridimensionales (1024 * 1024 * 64)

Ahora nuestro problema es la adición de 16 * 16 matrices ...

A | 1  2  3  4 |        B | 1  2  3  4 |      C| 1  2  3  4 |
  | 5  6  7  8 |   +      | 5  6  7  8 |   =   | 5  6  7  8 | 
  | 9 10 11 12 |          | 9 10 11 12 |       | 9 10 11 12 |  
  | 13 14 15 16|          | 13 14 15 16|       | 13 14 15 16|

Necesitamos 16 hilos para realizar el cálculo.

i.e. A(1,1) + B (1,1) = C(1,1)
     A(1,2) + B (1,2) = C(1,2) 
     .        .          .
     .        .          . 
     A(4,4) + B (4,4) = C(4,4) 

Todos estos hilos se ejecutarán simultáneamente. Entonces necesitamos un bloque con 16 hilos. Para nuestra conveniencia, arreglaremos los hilos en forma (16 * 1 * 1) en un bloque, ya que ninguno de los hilos tiene 16, por lo que solo necesitamos un bloque para almacenar esos 16 hilos.

entonces, la configuración de la grilla será dim3 Grid(1,1,1) es decir, la grilla tendrá solo un bloque y la configuración del bloque será dim3 block(16,1,1) es decir, el bloque tendrá 16 hilos dispuestos en forma de columna.

El siguiente programa le dará una idea clara de su ejecución. Comprender la parte de indexación (es decir, threadIDs, blockDim, blockID) es la parte importante. Necesitas revisar la literatura de CUDA. ¡Una vez que tenga una idea clara sobre la indexación, ganará la mitad de la batalla! ¡Así que dedique algo de tiempo a los libros cuda, a los diferentes algoritmos y al lápiz de papel, por supuesto!