¿Por qué molestarse en saber sobre CUDA Warps?

gpu gpu-warp

Question

Tengo GeForce GTX460 SE, por lo que es: 6 SM x 48 núcleos CUDA = 288 núcleos CUDA. Se sabe que en uno Warp contiene 32 hilos, y que en un bloque simultáneamente (a la vez) se puede ejecutar solo un Warp. Es decir, en un solo multiprocesador (SM), puede ejecutar simultáneamente solo un bloque, un Warp y solo 32 hilos, incluso si hay 48 núcleos disponibles.

Y además, se puede usar un ejemplo para distribuir Hilo y Bloque de hormigón threadIdx.x y blockIdx.x. Para asignarlos use kernel <<< Blocks, Threads >>> (). Pero ¿cómo asignar un número específico de Warp-s y distribuirlos, y si no es posible, entonces por qué molestarse en saber sobre Warps?

Answer 1

¿Los hilos agrupados en Warps están necesariamente en orden, 1 - 32, 33 - 64 ...?

Sí, el modelo de programación garantiza que los hilos se agrupen en urdimbres en ese orden específico.

Como ejemplo simple de optimización de las rutas de código divergentes se puede utilizar la separación de todos los hilos en el bloque en grupos de 32 hilos? Por ejemplo: switch (threadIdx.s / 32) {case 0: / * 1 warp * / break; caso 1: / * 2 warp * / break; / * Etc * /}

Exactamente :)

¿Cuántos bytes se deben leer a la vez para Warp individual: 4 bytes * 32 hilos, 8 bytes * 32 hilos o 16 bytes * 32 hilos? Por lo que yo sé, la única transacción a la memoria global al mismo tiempo recibe 128 bytes.

Sí, las transacciones a la memoria global son 128 bytes. Por lo tanto, si cada subproceso lee una palabra de 32 bits de direcciones consecutivas (probablemente también deben estar alineadas en 128 bytes), todos los subprocesos de la disformidad se pueden atender con una sola transacción (4 bytes * 32 subprocesos = 128 bytes ) Si cada hilo lee más bytes, o si las direcciones no son consecutivas, se deben emitir más transacciones (con transacciones separadas para cada línea separada de 128 bytes que se toca).

Esto se describe en el Manual de programación CUDA 4.2, sección F.4.2, "Memoria global". También hay una propaganda que dice que la situación es diferente con los datos que están en caché solo en L2, ya que la caché L2 tiene líneas de caché de 32 bytes. No sé cómo hacer para que los datos se almacenen solo en L2 o con cuántas transacciones uno termina.

Answer 2

La situación es bastante más complicada de lo que describes.

Las unidades ALU (núcleos), carga / almacenamiento (LD / ST) y unidades de funciones especiales (SFU) (verde en la imagen) son unidades canalizadas. Mantienen los resultados de muchos cálculos u operaciones al mismo tiempo, en varias etapas de finalización. Por lo tanto, en un ciclo pueden aceptar una nueva operación y proporcionar los resultados de otra operación que se inició hace mucho tiempo (alrededor de 20 ciclos para las ALU, si no recuerdo mal). Entonces, un único SM en teoría tiene recursos para procesar 48 * 20 ciclos = 960 operaciones de ALU al mismo tiempo, que son 960/32 hilos por urdimbre = 30 urdimbres. Además, puede procesar operaciones de LD / ST y operaciones de SFU en cualquier momento de latencia y rendimiento.

Los programadores de deformación (amarillo en la imagen) pueden programar 2 * 32 hilos por urdimbre = 64 hilos en las tuberías por ciclo. Entonces esa es la cantidad de resultados que se pueden obtener por reloj. Entonces, dado que hay una combinación de recursos informáticos, 48 ​​núcleos, 16 LD / ST, 8 SFU, cada uno con latencias diferentes, se está procesando una combinación de warps al mismo tiempo. En cualquier ciclo dado, los planificadores de warp intentan "emparejar" dos warps para programar, para maximizar la utilización del SM.

Los planificadores de warp pueden emitir distorsiones desde diferentes bloques o desde diferentes lugares en el mismo bloque, si las instrucciones son independientes. Por lo tanto, las deformaciones de varios bloques se pueden procesar al mismo tiempo.

Además de la complejidad, las urdimbres que están ejecutando instrucciones para las cuales hay menos de 32 recursos, se deben emitir varias veces para que se reparen todos los subprocesos. Por ejemplo, hay 8 SFU, por lo que significa que una urdimbre que contiene una instrucción que requiere las SFU se debe programar 4 veces.

Esta descripción se simplifica. Existen otras restricciones que también entran en juego que determinan cómo la GPU programa el trabajo. Puede encontrar más información buscando en la web "fermi architecture".

Entonces, llegando a tu pregunta real,

¿Por qué molestarse en saber sobre Warps?

Saber el número de hilos en una urdimbre y tomarlo en consideración es importante cuando intenta maximizar el rendimiento de su algoritmo. Si no sigue estas reglas, perderá rendimiento:

• En la invocación del kernel, <<<Blocks, Threads>>> , intente elegir una cantidad de hilos que se dividan de manera pareja con el número de hilos en un warp. Si no lo haces, terminas lanzando un bloque que contiene hilos inactivos.

• En su kernel, intente que cada hilo en un warp siga la misma ruta de código. Si no lo haces, obtienes lo que se llama divergencia warp. Esto sucede porque la GPU debe ejecutar toda la distorsión a través de cada una de las rutas de código divergentes.

• En tu kernel, intenta tener cada hilo en una carga warp y almacenar datos en patrones específicos. Por ejemplo, haga que los hilos en una urdimbre accedan a palabras consecutivas de 32 bits en la memoria global.