Comprender las dimensiones de la cuadrícula CUDA, las dimensiones del bloque y la organización de los hilos (explicación simple)

¿Cómo se organizan los hilos para ser ejecutados por una GPU?

Hardware

Si un dispositivo GPU tiene, por ejemplo, 4 unidades de multiprocesamiento, y pueden ejecutar 768 subprocesos cada una: entonces, en un momento dado, no más de 4 * 768 subprocesos realmente se ejecutarán en paralelo (si planificó más subprocesos, estarán esperando). su turno).

Software

los hilos están organizados en bloques. Un bloque es ejecutado por una unidad de multiprocesamiento. Los hilos de un bloque se pueden identificar (indexar) usando 1Dimension (x), 2Dimensions (x, y) o 3Dim indexes (x, y, z) pero en cualquier caso x y z < = 768 para nuestro ejemplo (se aplican otras restricciones) a x, y, z, consulte la guía y la capacidad de su dispositivo).

Obviamente, si necesita más de esos 4 * 768 hilos, necesita más de 4 bloques. Los bloques también pueden indexarse ​​1D, 2D o 3D. Hay una cola de bloques esperando ingresar a la GPU (porque, en nuestro ejemplo, la GPU tiene 4 multiprocesadores y solo 4 bloques se están ejecutando simultáneamente).

Ahora un caso simple: procesar una imagen de 512×512

Supongamos que queremos que un hilo procese un píxel (i, j).

Podemos usar bloques de 64 hilos cada uno. Entonces necesitamos 512 * 512/64 = 4096 bloques (para tener 512×512 hilos = 4096 * 64)

Es común organizar (para facilitar la indexación de la imagen) los hilos en bloques 2D con blockDim = 8 x 8 (los 64 hilos por bloque). Prefiero llamarlo threadsPerBlock.

dim3 threadsPerBlock(8, 8); // 64 threads 

y 2D gridDim = 64 x 64 bloques (se necesitan 4096 bloques). Prefiero llamarlo numBlocks.

 dim3 numBlocks(imageWidth/threadsPerBlock.x, /* for instance 512/8 = 64*/ imageHeight/threadsPerBlock.y); 

El núcleo se inicia así:

 myKernel < <>>( /* params for the kernel function */ ); 

Finalmente: habrá algo así como “una cola de 4096 bloques”, donde un bloque espera que se le asigne uno de los multiprocesadores de la GPU para ejecutar sus 64 subprocesos.

En el kernel, el píxel (i, j) que se procesa por un hilo se calcula de esta manera:

 uint i = (blockIdx.x * blockDim.x) + threadIdx.x; uint j = (blockIdx.y * blockDim.y) + threadIdx.y; 

supongamos que una GPU de 9800GT: 14 multiprocesadores, cada uno tiene 8 procesadores de hilo y warpsize es 32, lo que significa que cada subprocesador maneja hasta 32 subprocesos. 14 * 8 * 32 = 3584 es la cantidad máxima de hilos actuales cuncurrent.

si ejecuta este núcleo con más de 3584 hilos (digamos 4000 hilos y no es importante cómo define el bloque y la rejilla. gpu los tratará como si fueran iguales):

 func1(); __syncthreads(); func2(); __syncthreads(); 

entonces el orden de ejecución de esas dos funciones es el siguiente:

1.func1 se ejecuta para los primeros 3584 hilos

2.func2 se ejecuta para los primeros 3584 hilos

3.func1 se ejecuta para los hilos restantes

4.func2 se ejecuta para los hilos restantes