bloques, hilos, warpSize

Question

Ha habido mucha discusión sobre cómo elegir #blocks & blockSize, pero aún me falta algo. Muchos de mis preocupaciones frente a esta cuestión: How CUDA Blocks/Warps/Threads map onto CUDA Cores? (Para simplificar la discusión, hay suficiente perThread & memoria perBlock límites de memoria no son un problema aquí..)

kernelA<<<nBlocks, nThreads>>>(varA,constB, nThreadsTotal); 

1) Mantener el SM tan ocupado como sea posible, Debería establecer nThreads en un múltiplo de warpSize. ¿Cierto?

2) Un SM solo puede ejecutar un kernel a la vez. Eso es todo HWcores de ese SM ejecutando solo kernelA. (No algunos HWcores ejecutan kernelA, mientras que otros ejecutan kernelB.) Así que si solo tengo un hilo para ejecutar, estoy "perdiendo" el otro HWcores. ¿Cierto?

3) Si los problemas del programador warp funcionan en unidades de warpSize (32 hilos), y cada SM tiene 32 HWcores, entonces el SM se utilizará por completo. ¿Qué sucede cuando el SM tiene 48 HWcores? ¿Cómo puedo mantener los 48 núcleos completos utilizados cuando el planificador está emitiendo trabajos en fragmentos de 32? (Si el párrafo anterior es verdadero, ¿no sería mejor si el planificador emitió trabajo en unidades de tamaño HWcore?)

4) Parece que el planificador warp pone en cola 2 tareas a la vez. De modo que cuando el kernel que se está ejecutando se detenga o se bloquee, se intercambie el segundo kernel. (No está claro, pero supongo que la cola aquí tiene más de 2 núcleos). ¿Es esto correcto?

5) Si mi HW tiene un límite superior de 512 hilos por bloque (nThreadsMax), eso no significa que el núcleo con 512 hilos se ejecutará más rápido en un bloque. (De nuevo, no es un problema.) Existe una buena posibilidad de que obtenga un mejor rendimiento si extiendo el kernel de 512 hilos a través de muchos bloques, no solo uno. El bloque se ejecuta en uno o muchos SM. ¿Cierto?

5a) Estoy pensando que cuanto más pequeño mejor, ¿pero importa lo pequeño que sea nBlocks? La pregunta es, ¿cómo elegir el valor de nBlocks que sea decente? (No necesariamente óptimo.) ¿Existe un enfoque matemático para elegir nBlocks, o simplemente es trial-n-err.

Answer 1

Déjame intentar responder tus preguntas una por una.

1. Eso es correcto.
2. ¿Qué quiere decir exactamente con "HWcores"? La primera parte de su declaración es correcta.

3. Según el NVIDIA Fermi Compute Architecture Whitepaper: "El SM programa hilos en grupos de 32 hilos paralelos llamados urdimbres. Cada SM presenta dos planificadores de deformación y dos unidades de despacho de instrucciones, permitiendo que se emitan y ejecuten dos urdimbres al mismo tiempo. dos warps y emite una instrucción de cada warp a un grupo de dieciséis núcleos, dieciséis unidades de carga/almacén o cuatro SFU. Debido a que los warps se ejecutan de forma independiente, el programador de Fermi no necesita verificar las dependencias desde el flujo de instrucciones ".

   Además, el NVIDIA Keppler Architecture Whitepaper establece: "El programador de deformación cuádruple de Kepler selecciona cuatro distorsiones, y se pueden enviar dos instrucciones independientes por urdimbre en cada ciclo".

   Por lo tanto, los núcleos "sobrantes" se utilizan al programar más de un salto a la vez.

4. El programador warp programa warps del mismo kernel, no kernels diferentes.

5. No del todo cierto: cada bloque está bloqueado en un único SM, ya que es donde reside su memoria compartida.
6. Eso es un problema complicado y depende de cómo se implemente su kernel. Es posible que desee echar un vistazo a nVidia Webinar Better Performance at Lower Occupancy por Vasily Volkov, que explica algunos de los problemas más importantes. Principalmente, sin embargo, le sugiero que elija su conteo de hilos para mejorar la ocupación, usando el CUDA Occupancy Calculator.

Answer 2

1) Sí.

2) Los dispositivos CC 2.0 - 3.0 pueden ejecutar hasta 16 redes al mismo tiempo. Cada SM está limitado a 8 bloques, por lo que para alcanzar una concurrencia total, el dispositivo debe tener al menos 2 SM.

3) Sí, los programadores warp seleccionan y emiten warps a la vez. Olvídese del concepto de núcleos CUDA que son irrelevantes. Para ocultar la latencia, necesita tener un alto paralelismo de nivel de instrucción o una alta ocupación. Se recomienda tener> 25% para CC 1.x y> 50% para CC> = 2.0. En general, CC 3.0 requiere una mayor ocupación que los dispositivos 2.0 debido a la duplicación de los programadores, pero solo un 33% de aumento en los urdidos por SM. El experimento Nsight VSE Issue Efficiency es la mejor manera de determinar si tuvo suficientes warps para ocultar la latencia de la instrucción y la memoria. Desafortunadamente, el Visual Profiler no tiene esta métrica.

4) El algoritmo del planificador warp no está documentado; sin embargo, no considera en qué cuadrícula se originó el bloque de hilos. Para los dispositivos CC 2.xy 3.0, el distribuidor de trabajo de CUDA distribuirá todos los bloques desde una cuadrícula antes de distribuir los bloques desde la siguiente cuadrícula; sin embargo, esto no está garantizado por el modelo de programación.

5) Para mantener el SM ocupado, debe tener suficientes bloques para llenar el dispositivo. Después de eso, debes asegurarte de tener suficientes urdimbres para alcanzar una ocupación razonable. Existen ventajas y desventajas para usar bloques de hilos grandes. En general, los bloques de hilos grandes usan menos memoria caché de instrucciones y tienen huellas más pequeñas en la memoria caché; sin embargo, los bloques de subprocesos grandes bloquean las conexiones sincrónicas (SM puede volverse menos eficiente ya que hay menos distorsiones entre las que elegir) y tienden a mantener las instrucciones ejecutándose en unidades de ejecución similares. Recomiendo probar 128 o 256 hilos por bloque de hilo para comenzar. Hay buenas razones para ambos bloques de hilos más grandes y más pequeños. 5a) Use la calculadora de ocupación. Elegir un tamaño de bloque de hilo demasiado grande a menudo hará que los registros lo limiten. Escoger un tamaño de bloque de hilo demasiado pequeño puede encontrarlo limitado por la memoria compartida o por los 8 bloques por límite de SM.