Comprender las dimensiones de la cuadrícula CUDA, las dimensiones del bloque y la organización de subprocesos (explicación simple) [cerrado]

161

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

cuda nvidia cibercitizen1
fuente
La Guía de programación de CUDA debería ser un buen lugar para comenzar con esto. También recomendaría revisar la introducción de CUDA desde aquí .
Tom

Respuestas:

287

Hardware

Si un dispositivo GPU tiene, por ejemplo, 4 unidades de multiprocesamiento, y pueden ejecutar 768 subprocesos cada una: en un momento dado, no más de 4 * 768 subprocesos realmente se ejecutarán en paralelo (si planeó 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 pueden identificarse (indexarse) 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 hilos 4 * 768, necesita más de 4 bloques. Los bloques también pueden indexarse ​​en 1D, 2D o 3D. Hay una cola de bloques esperando para ingresar a la GPU (porque, en nuestro ejemplo, la GPU tiene 4 multiprocesadores y solo 4 bloques se ejecutan simultáneamente).

Ahora un caso simple: procesar una imagen de 512x512

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 hilos de 512x512 = 4096 * 64)

Es común organizar (para facilitar la indexación de la imagen) los hilos en bloques 2D que tienen 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 los 4096 bloques). Prefiero llamarlo numBlocks.

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

El kernel se inicia así:

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

Finalmente: habrá algo como "una cola de 4096 bloques", donde un bloque está esperando que se le asigne uno de los multiprocesadores de la GPU para ejecutar sus 64 subprocesos.

En el núcleo, el píxel (i, j) que procesará un subproceso se calcula de esta manera:

uint i = (blockIdx.x * blockDim.x) + threadIdx.x;
uint j = (blockIdx.y * blockDim.y) + threadIdx.y;
cibercitizen1
fuente
11
Si cada bloque puede ejecutar 768 hilos, ¿por qué usar solo 64? Si usa el límite máximo de 768, tendrá menos bloques y un mejor rendimiento.
Aliza
10
@Aliza: los bloques son lógicos , el límite de 768 hilos es para cada unidad de procesamiento físico . Utiliza bloques de acuerdo con las especificaciones de su problema para distribuir el trabajo a los hilos. No es probable que siempre pueda usar bloques de 768 hilos para cada problema que tenga. Imagine que tiene que procesar una imagen de 64x64 (4096 píxeles). 4096/768 = 5.333333 bloques?
cibercitizen1
1
Los bloques son lógicos, pero cada bloque está asignado a un núcleo. Si hay más bloques que núcleo, los bloques se ponen en cola hasta que los núcleos se liberan. En su ejemplo, puede usar 6 bloques y hacer que los hilos adicionales no hagan nada (2/3 de los hilos en el sexto bloque).
Aliza
3
@ cibercitizen1 - Creo que el punto de Aliza es bueno: si es posible, uno quiere usar tantos hilos por bloque como sea posible. Si hay una restricción que requiere menos subprocesos, es mejor explicar por qué ese podría ser el caso en un segundo ejemplo (pero aún explica el caso más simple y deseable, primero).
66
@ this Sí, tal vez. Pero el caso es que la cantidad de memoria que necesita cada subproceso depende de la aplicación. Por ejemplo, en mi último programa, cada hilo invoca una función de optimización de mínimos cuadrados, que requiere "mucha" memoria. Tanto, que los bloques no pueden ser más grandes que hilos 4x4. Aun así, la aceleración obtenida fue dramática, frente a la versión secuencial.
cibercitizen1
9

Supongamos una GPU 9800GT:

  • tiene 14 multiprocesadores (SM)
  • cada SM tiene 8 procesadores de hilo (procesadores de flujo AKA, SP o núcleos)
  • permite hasta 512 hilos por bloque
  • warpsize es 32 (lo que significa que cada uno de los procesadores de hilos 14x8 = 112 puede programar hasta 32 hilos)

https://www.tutorialspoint.com/cuda/cuda_threads.htm

Un bloque no puede tener más subprocesos activos que 512, por __syncthreadslo tanto , solo puede sincronizar un número limitado de subprocesos. es decir, si ejecuta lo siguiente con 600 hilos:

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

entonces el kernel debe ejecutarse dos veces y el orden de ejecución será:

  1. func1 se ejecuta para los primeros 512 hilos
  2. func2 se ejecuta para los primeros 512 hilos
  3. func1 se ejecuta para los hilos restantes
  4. func2 se ejecuta para los hilos restantes

Nota:

El punto principal __syncthreadses una operación de todo el bloque y no sincroniza todos los hilos.

No estoy seguro del número exacto de subprocesos que __syncthreadspueden sincronizarse, ya que puede crear un bloque con más de 512 subprocesos y dejar que la urdimbre se encargue de la programación. Según tengo entendido, es más exacto decir: func1 se ejecuta al menos para los primeros 512 hilos.

Antes de editar esta respuesta (en 2010) midí 14x8x32 hilos sincronizados usando __syncthreads .

Le agradecería mucho que alguien vuelva a probar esto para obtener una información más precisa.

Bizhan
fuente
¿Qué sucede si func2 () depende de los resultados de func1 ()? Creo que esto está mal
Chris
@Chris Escribí esto hace siete años, pero si recuerdo bien, hice una prueba al respecto y llegué a la conclusión de que los núcleos con más hilos que gpu se comportan de esta manera. Si pruebas este caso y alcanzas un resultado diferente, tendré que eliminar esta publicación.
Bizhan
Lo siento, creo que esto también está mal, que la GPU solo puede ejecutar 112 hilos simultáneamente.
Steven Lu
@StevenLu, ¿lo has probado? Tampoco creo que 112 hilos concurrentes tengan sentido para una GPU. 112 es el número de procesadores de flujo. Apenas puedo recordar CUDA ahora :)
Bizhan
1
@StevenLu el número máximo de hilos no es el problema aquí, __syncthreadses una operación de todo el bloque y el hecho de que en realidad no sincroniza todos los hilos es una molestia para los estudiantes de CUDA. Así que actualicé mi respuesta en función de la información que me diste. Realmente lo aprecio.
Bizhan