Hilos en CUDA (II)

La ocupación es el ratio entre el número de hilos o warps que se ejecutan sobre un SM para un kernel concreto y el número máximo de hilos o warps que se podrían ejecutar potencialmente sobre dicho SM para ese kernel. Es decir: occ = threadsPerSM / maxThreadsPerSM.

Los hilos en ejecución (reales) en un SM dependen de muchos factores, entre ellos: el tamaño de bloque elegido, el número de registros en uso por cada hilo, la memoria compartida necesaria y su configuración e incluso la compute capability de la GPU. El denominador, el máximo número de hilos teórico por SM es únicamente dependiente de la generación de la tarjeta y su compute capability.

Así pues, por ejemplo para una tarjeta con compute capability 3.7 tenemos un máximo de 2048 hilos por multiprocesador y una limitación de 16 bloques por multiprocesador. De esa forma, si ejecutamos un kernel con una configuración de 16 hilos por bloque, podremos tener un total de 256 hilos sin exceder la limitación de 16 bloques por multiprocesador. Esto significa que como mucho ejecutaremos 256 hilos a la vez, mientras que el máximo teórico es de 2048 por lo que la ocupación será de 256/2048 = 0.125 o lo que es lo mismo del 12.5%, una utilización bastante pobre de los recursos.

Por norma general, una mayor ocupación implica una mejor utilización de los recursos de cómputo de la tarjeta y por lo tanto un mayor rendimiento. No obstante, el rendimiento final depende de varios factores y no solo de la ocupación por lo que en ocasiones merece más la pena ejecutar un kernel con una ocupación baja que aproveche mejor otros factores (Un gran ejemplo de este fenómeno es la presentación ed Vasily Volkov: Better Performance at Lower Occupancy que puede ser consultada aquí http://www.nvidia.com/content/gtc-2010/pdfs/2238_gtc2010.pdf).

Problema 1

Si necesitamos tantos hilos como elementos en dos vectores para llevar a cabo la suma de los mismos, ¿qué expresión de las siguientes sería la correcta para llevar a cabo el mapeo entre hilos y elementos?

idx = threadIdx.x + threadIdx.y;
idx = blockIdx.x + threadIdx.x;
idx = blockIdx.x * blockDim.x + threadIdx.x;
idx = blockIdx.x * threadIdx.x;

Problema 2

Continuando con el ejercicio anterior, si además queremos que cada hilo calcule dos posiciones adyacentes del vector de elementos, ¿qué expresión de las siguientes sería la correcta para llevar a cabo el mapeo entre hilos y elementos?

idx = blockIdx.x * blockDim.x + threadIdx.x + 2;
idx = blockIdx.x * threadIdx.x * 2;
idx = (blockIdx.x * blockDim.x + threadIdx.x) * 2;
idx = blockIdx.x * blockDim.x * 2 + threadIdx.x;

Problema 3

Para una operación suma de vectores, asume que el tamaño de los vectores a sumar es 2000, cada hilo calcula 1 elemento de salida, y el tamaño del bloque es de 512 hilos, ¿cuántos hilos habrá en total en el grid ejecutado?

2000
2024
2048
2096

Problema 4

Si un SM de un dispositivo CUDA puede ejecutar hasta 1536 hilos y hasta 4 bloques a la vez, ¿cuál de las siguientes configuraciones obtendría mayor rendimiento (en este caso mayor número de hilos en ejecución)? Asumimos que en nuestro problema, el rendimiento viene determinado principalmente por la ocupación.

128 hilos por bloque
256 hilos por bloque
512 hilos por bloque
1024 hilos por bloque

Problema 5

Necesitamos escribir un kernel que opere sobre una imagen de tamaño 400x900 píxeles. Queremos asignar un hilo para cada pixel. El número de hilos por bloque tiene que ser cuadrado y utilizar el máximo número de hilos por bloque posible en tu dispositivo (arquitectura Fermi con Compute Capability 2.0).

¿Qué tamaño de grid y bloque elegirías?
¿Cuántos hilos no realizarán ningún cómputo?

Problemas 3: Hilos en CUDA (II)