¿Cómo se lanza un kernel CUDA?

Question

He creado 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

Ejecuta una cuadrícula 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 subprocesos que ejecutan la misma instrucción o están inactivos (porque se ramificaron, por ejemplo, al salir de un ciclo antes que los vecinos dentro de la misma distorsión o sin tomar el if). En una GPU de Fermi, como máximo dos distorsiones se ejecutan en un multiprocesador a la vez.

Cuando hay latencia (es decir, paradas 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 bloques diferentes registros utilizados por cada hilo y la cantidad de memoria compartida utilizada por a/the block (s)).

Esta programación se realiza de forma 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 (¿cuán grande es un bloque?), blockIdx (¿dónde está mi bloque en la cuadrícula?) Y gridDim (qué tamaño 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 (todos los hay) por otro lanzamiento del kernel. Para la ejecución no secuencial del núcleo secuencial, no se necesita más sincronización.

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

Puede sincronizar todos los hilos dentro de un bloque con la instrucción intrínseca __syncthreads() (todos los hilos se activarán después, aunque, como siempre, como máximo dos distorsiones se pueden ejecutar en una GPU 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 una urdimbre están siempre "sincronizados", aunque algunos pueden 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 voto de warp que le pueden ayudar a tomar decisiones de sucursal o calcular sumas enteras (prefijo).

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

Answer 2

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

Answer 3

permite echar un ejemplo de adición de 4 * 4 matrices .. tiene dos matrices A y B, que tiene 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]; 
} 

permite echar un ejemplo de adición de 16 * 16 matrices .. tiene dos matrices A y B, que tienen la dimensión 16 * 16 ..

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

Ahora, una rejilla se inicia con una función de kernel .. Una rejilla 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 rejilla puede tener max 1024 no de hilos threads.Those también pueden estar dispuestos en 3 formas dimensionales (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.

así, la configuración de cuadrícula se dim3 Grid(1,1,1) es decir rejilla tendrá un solo bloque y configuración bloque será dim3 block(16,1,1) bloque es decir, tendrá 16 hilos de columna dispuestos sabia.

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! ¡Entonces dedique algo de tiempo a los libros cuda, diferentes algoritmos y lápiz de papel, por supuesto!