Computación GPU: por qué, cuándo y cómo. Además de algunas pruebas

Todos saben desde hace mucho tiempo que en las tarjetas de video no solo puedes jugar juguetes, sino también realizar cosas que no están relacionadas con los juegos, por ejemplo, entrenar una red neuronal, recordar criptomonedas o realizar cálculos científicos. Cómo sucedió, puede leerlo aquí , pero quería tocar el tema de por qué la GPU puede ser interesante para el programador promedio (no relacionado con GameDev) cómo abordar el desarrollo en la GPU sin perder mucho tiempo, decidir si mire en esta dirección y " descubra con los dedos" qué beneficio puede obtener. 



El artículo fue escrito en base a mi presentación en HighLoad ++. Discute principalmente las tecnologías ofrecidas por NVIDIA. No tengo ningún propósito de anunciar ningún producto, solo los doy a modo de ejemplo, y seguro que se puede encontrar algo similar en los fabricantes de la competencia.

¿Por qué contar con la GPU?


Se pueden comparar dos procesadores de acuerdo con diferentes criterios, probablemente los más populares son la frecuencia y el número de núcleos, el tamaño de los cachés, etc., pero al final, estamos interesados ​​en cuántas operaciones puede realizar un procesador por unidad de tiempo, qué tipo de operación es, pero una pregunta aparte Una métrica común es el número de operaciones de punto flotante por segundo - flops. Y cuando queremos comparar cálido con suave, y en nuestro caso GPU con CPU, esta métrica es útil.

El siguiente gráfico muestra el crecimiento de estos mismos fracasos a lo largo del tiempo para procesadores y tarjetas de video.


(Los datos se recopilan de fuentes abiertas, no hay datos para 2019-20 años, porque no todo es tan hermoso allí, pero las GPU aún ganan)

Bueno, es tentador, ¿no? ¡Cambiamos todos los cálculos de la CPU a la GPU y obtenemos ocho veces el mejor rendimiento!

Pero, por supuesto, no todo es tan simple. No puede simplemente tomar y transferir todo a la GPU, por eso, hablaremos más.

Arquitectura de GPU y su comparación con la CPU


Les traigo a muchos una imagen familiar con la arquitectura de la CPU y los elementos básicos:


CPU Core

¿Qué es tan especial? Un núcleo y un montón de bloques auxiliares.

Ahora echemos un vistazo a la arquitectura GPU:


La

tarjeta de video GPU Core A tiene muchos núcleos de procesamiento, generalmente varios miles, pero se combinan en bloques; para las tarjetas de video NVIDIA, generalmente 32 cada una, y tienen elementos comunes, incluidos y se registra. La arquitectura del núcleo de la GPU y los elementos lógicos es mucho más simple que en la CPU, es decir, no hay pretratadores, predictores de brunch y mucho más.

Bueno, estos son los puntos clave de la diferencia en la arquitectura de la CPU y la GPU y, de hecho, imponen restricciones o, por el contrario, abren las posibilidades de lo que podemos leer efectivamente en la GPU.

No mencioné un punto más importante, por lo general, la tarjeta de video y el procesador no "hurgan" entre ellos y escriben datos en la tarjeta de video y leen el resultado; estas son operaciones separadas y pueden resultar en un "cuello de botella" en su sistema, un gráfico del tiempo de bombeo versus el tamaño Los datos se dan más adelante en el artículo.

Limitaciones y características de la GPU


¿Qué limitaciones impone esta arquitectura a los algoritmos ejecutables?

  • Si estamos calculando en una GPU, entonces no podemos seleccionar solo un núcleo, se asignará un bloque completo de núcleos (32 para NVIDIA).
  • Todos los núcleos ejecutan las mismas instrucciones, pero con datos diferentes (hablaremos de esto más adelante), tales cálculos se denominan Datos de instrucciones múltiples o SIMD (aunque NVIDIA introduce su refinamiento). 
  • Debido al conjunto relativamente simple de bloques lógicos y registros generales, a la GPU realmente no le gusta la ramificación, y de hecho la lógica compleja en los algoritmos.

¿Qué oportunidades abre?

  • En realidad, la aceleración de esos mismos cálculos SIMD. El ejemplo más simple es la adición de matrices por elementos, y analicémoslo.

Reducción de algoritmos clásicos a representación SIMD


Transformación


Tenemos dos matrices, A y B, y queremos agregar un elemento de la matriz B a cada elemento de la matriz A. A continuación se muestra un ejemplo en C, aunque espero que quede claro para aquellos que no hablan este idioma:

void func(float *A, float *B, size)
{ 
   for (int i = 0; i < size; i++) 
   { 
       A[i] += B[i]
   } 
}

Bucle clásico de elementos en un bucle y tiempo de ejecución lineal.

Ahora veamos cómo se verá ese código para la GPU:

void func(float *A, float *B, size) 
{ 
   int i = threadIdx.x; 
   if (i < size) 
      A[i] += B[i] 
}

Y aquí ya es interesante, apareció la variable threadIdx, que no parecía declarar en ningún lado. Sí, su sistema nos proporciona. Imagine que en el ejemplo anterior la matriz consta de tres elementos y desea ejecutarla en tres hilos paralelos. Para hacer esto, necesitará agregar otro parámetro: el índice o el número de flujo. Esto es lo que la tarjeta de video hace por nosotros, aunque pasa el índice como una variable estática y puede funcionar con varias dimensiones a la vez: x, y, z.

Otro matiz, si va a iniciar una gran cantidad de transmisiones paralelas a la vez, entonces las transmisiones tendrán que dividirse en bloques (una característica arquitectónica de las tarjetas de video). El tamaño máximo de bloque depende de la tarjeta de video, y el índice del elemento para el que realizamos los cálculos deberá obtenerse de la siguiente manera:

int i = blockIdx.x * blockDim.x + threadIdx.x; // blockIdx –  , blockDim –  , threadIdx –    

Como resultado, lo que tenemos: muchos subprocesos de ejecución paralela que ejecutan el mismo código, pero con diferentes índices y, en consecuencia, datos, es decir. El mismo SIMD.

Este es el ejemplo más simple, pero si desea trabajar con la GPU, debe llevar su tarea a la misma forma. Desafortunadamente, esto no siempre es posible y, en algunos casos, puede ser objeto de una disertación doctoral, pero, sin embargo, los algoritmos clásicos aún pueden llevarse a esta forma.

Agregación


Veamos ahora cómo se verá la agregación agregada a la representación SIMD:
 

Tenemos una matriz de n elementos. En la primera etapa, comenzamos n / 2 hilos y cada hilo agrega dos elementos, es decir En una iteración, sumamos la mitad de los elementos en la matriz. Y luego, en el bucle, repetimos lo mismo para la matriz recién creada, hasta que agreguemos los dos últimos elementos. Como puede ver, cuanto más pequeño es el tamaño de la matriz, menos hilos paralelos podemos comenzar, es decir en una GPU, tiene sentido agregar matrices de un tamaño suficientemente grande. Tal algoritmo se puede usar para calcular la suma de elementos (por cierto, no se olvide del posible desbordamiento del tipo de datos con el que está trabajando), busque el máximo, mínimo o simplemente busque.

Clasificación


Pero la clasificación ya parece mucho más complicada.

Los dos algoritmos de clasificación más populares en la GPU son:

  • Tipo bitónico
  • Radix-sort

Pero la clasificación por radix todavía se usa con más frecuencia, y la implementación lista para producción se puede encontrar en algunas bibliotecas. No analizaré en detalle cómo funcionan estos algoritmos; aquellos que estén interesados ​​pueden encontrar una descripción de clasificación por radix en https://www.codeproject.com/Articles/543451/Parallel-Radix-Sort-on-the-GPU-using-Cplusplus- AMP y https://stackoverflow.com/a/26229897

Pero la idea es que incluso un algoritmo no lineal como la clasificación se puede reducir a una vista SIMD.

Y ahora, antes de ver los números reales que se pueden obtener de la GPU, descubramos cómo programar para este milagro de la tecnología.

Donde empezar


Las dos tecnologías más comunes que se pueden usar para programar en la GPU:

  • Opencl
  • Cuda

OpenCL es un estándar que es compatible con la mayoría de los fabricantes de tarjetas de video, incluyendo y en dispositivos móviles, también el código escrito en OpenCL puede ejecutarse en la CPU.

Puede usar OpenCL desde C / C ++, hay carpetas para otros idiomas.

Para OpenCL, me gustó más el libro OpenCL en acción . También describe diferentes algoritmos en la GPU, incluidos Bitonic-sort y Radix-sort.

CUDA es la tecnología patentada y SDK de NVIDIA. Puede escribir en C / C ++ o usar enlaces a otros idiomas.

Comparar OpenCL y CUDA es algo incorrecto, porque uno es el estándar, el otro es todo el SDK. Sin embargo, muchas personas eligen CUDA para el desarrollo de tarjetas de video, a pesar de que la tecnología es propietaria, aunque gratuita y solo funciona en tarjetas NVIDIA. Hay varias razones para esto:

  • API
  • , GPU, (host)
  • , ..

Las peculiaridades incluyen el hecho de que CUDA viene con su propio compilador, que también puede compilar código C / C ++ estándar.

El libro de CUDA más completo que encontré fue la Programación profesional de CUDA C , aunque ya está un poco desactualizado, sin embargo, analiza muchos matices técnicos de programación para tarjetas NVIDIA.

Pero, ¿qué pasa si no quiero pasar un par de meses leyendo estos libros, escribiendo mi propio programa para una tarjeta de video, probando y depurando, y luego descubriendo que esto no es para mí? 

Como dije, hay una gran cantidad de bibliotecas que ocultan la complejidad del desarrollo bajo la GPU: XGBoost, cuBLAS, TensorFlow, PyTorch y otras, consideraremos la biblioteca de empuje, ya que es menos especializado que las otras bibliotecas anteriores, pero al mismo tiempo implementa algoritmos básicos, por ejemplo, clasificación, búsqueda, agregación, y con alta probabilidad puede ser aplicable en sus tareas.

Thrust es una biblioteca de C ++ que tiene como objetivo "reemplazar" los algoritmos STL estándar con algoritmos basados ​​en GPU. Por ejemplo, ordenar una matriz de números usando esta biblioteca en una tarjeta de video se vería así:

thrust::host_vector<DataType> h_vec(size); //    
std::generate(h_vec.begin(), h_vec.end(), rand); //   
thrust::device_vector<DataType> d_vec = h_vec; //         
thrust::sort(d_vec.begin(), d_vec.end()); //    
thrust::copy(d_vec.begin(), d_vec.end(), h_vec.begin()); //   ,     

(no olvide que el ejemplo debe ser compilado por un compilador de NVIDIA)

Como puede ver, thrust :: sort es muy similar a un algoritmo similar de STL. Esta biblioteca esconde muchas dificultades, en particular el desarrollo de un subprograma (más precisamente, el núcleo), que se ejecutará en la tarjeta de video, pero al mismo tiempo priva de flexibilidad. Por ejemplo, si queremos clasificar varios gigabytes de datos, sería lógico enviar un dato a la tarjeta para comenzar a ordenar, y mientras la clasificación está en progreso, enviar más datos a la tarjeta. Este enfoque se llama ocultación de latencia y permite un uso más eficiente de los recursos del mapa del servidor, pero, desafortunadamente, cuando usamos bibliotecas de alto nivel, tales oportunidades permanecen ocultas. Pero para la creación de prototipos y la medición del rendimiento, son iguales, especialmente con el empuje que puede medir qué gastos generales proporciona la transferencia de datos.

Escribí un pequeño punto de referencia usando esta biblioteca, que ejecuta varios algoritmos populares con diferentes cantidades de datos en la GPU, veamos cuáles son los resultados.

Resultados del algoritmo de GPU


Para probar la GPU, tomé una instancia en AWS con una tarjeta gráfica Tesla k80, esta no es la tarjeta de servidor más potente hasta la fecha (la más potente Tesla v100), pero es la más asequible y tiene a bordo:

  • 4992 granos de CUDA
  • 24 GB de memoria
  • 480 Gb / s - ancho de banda de memoria 

Y para las pruebas en la CPU, tomé una instancia con un procesador Intel Xeon CPU E5-2686 v4 @ 2.30GHz

Transformación



Tiempo de ejecución de transformación en la GPU y la CPU en ms

Como puede ver, la transformación habitual de los elementos de la matriz es aproximadamente la misma en el tiempo, tanto en la GPU como en la CPU. ¿Y por qué? Debido a que la sobrecarga para enviar datos a la tarjeta y viceversa consume todo el aumento del rendimiento (hablaremos de la sobrecarga por separado), y hay relativamente pocos cálculos en la tarjeta. Además, no olvide que los procesadores también son compatibles con las instrucciones SIMD, y los compiladores en casos simples pueden usarlas de manera efectiva. 

Veamos ahora qué tan eficientemente se realiza la agregación en la GPU.

Agregación



Tiempo de ejecución de agregación en GPU y CPU en ms

En el ejemplo de agregación, ya vemos un aumento significativo en el rendimiento con un aumento en el volumen de datos. También vale la pena prestar atención al hecho de que bombeamos una gran cantidad de datos a la memoria de la tarjeta, y solo se recupera un valor agregado, es decir Los gastos generales para transferir datos de la tarjeta a la RAM son mínimos.

Pasemos al ejemplo más interesante: la clasificación.

Clasificación



Tiempo de clasificación para la GPU y la CPU en ms

A pesar de que enviamos toda la matriz de datos a la tarjeta de video y viceversa, la clasificación a la GPU de 800 MB de datos es aproximadamente 25 veces más rápida que en el procesador.

Sobrecarga de transferencia de datos


Como se puede ver en el ejemplo de transformación, no siempre es obvio si la GPU será efectiva incluso en aquellas tareas que funcionan bien en paralelo. La razón de esto es una sobrecarga para transferir datos desde la RAM de la computadora a la memoria de la tarjeta de video (en las consolas de juegos, por cierto, la memoria se comparte entre la CPU y la GPU, y no hay necesidad de transferir datos). Una de las características de una tarjeta de video es el ancho de banda de la memoria o el ancho de banda de la memoria, que determina el ancho de banda teórico de la tarjeta. Para Tesla k80 es de 480 GB / s, para Tesla v100 ya es de 900 GB / s. Además, la versión PCI Express y la implementación de cómo transferirá los datos a la tarjeta afectarán el rendimiento, por ejemplo, esto puede hacerse en varias transmisiones paralelas.

Veamos los resultados prácticos que se obtuvieron para la tarjeta gráfica Tesla k80 en la nube de Amazon:


Tiempo para transferir datos a la GPU, ordenar y transferir datos de vuelta a la RAM en ms

HtoD - transferir datos a la tarjeta de video

GPU Ejecución - ordenar en la tarjeta de video

DtoH - copiar datos de la tarjeta de video a la RAM


Lo primero a tener en cuenta es que leer los datos de la tarjeta de video es más rápido que escríbelos ahí abajo.

El segundo: cuando trabaja con una tarjeta de video, puede obtener una latencia de 350 microsegundos, y esto puede ser suficiente para algunas aplicaciones de baja latencia.

El siguiente cuadro muestra una sobrecarga para más datos:


Tiempo para transferir datos a la GPU, ordenar y transferir datos a la RAM en ms

Uso del servidor


La pregunta más común es ¿en qué se diferencia una tarjeta de video de juego de una de servidor? Según las características, son muy similares, pero los precios difieren significativamente.


Las principales diferencias entre el servidor (NVIDIA) y la tarjeta de juego:

  • Garantía del fabricante (la tarjeta de juego no está diseñada para uso del servidor)
  • Posibles problemas de virtualización para una tarjeta gráfica de consumidor
  • Disponibilidad del mecanismo de corrección de errores en la tarjeta del servidor
  • El número de subprocesos paralelos (no núcleos CUDA) o soporte para Hyper-Q, que le permite trabajar con la tarjeta desde varios subprocesos en la CPU, por ejemplo, cargar datos a la tarjeta desde un subproceso e iniciar cálculos desde otro

Estas son, quizás, las principales diferencias importantes que encontré.

Multithreading


Después de descubrir cómo ejecutar el algoritmo más simple en la tarjeta de video y qué resultados se pueden esperar, la siguiente pregunta lógica es cómo se comportará la tarjeta de video al procesar varias solicitudes paralelas. Como respuesta, tengo dos gráficos de computación en la GPU y un procesador con 4 y 32 núcleos:


El tiempo necesario para realizar cálculos matemáticos en la GPU y la CPU con matrices de 1000 x 60 en ms

. Este gráfico realiza cálculos con matrices de 1000 x 60 elementos. Los cálculos se inician a partir de varias secuencias de programas, se crea una secuencia separada para la GPU para cada secuencia de CPU (se utiliza el Hyper-Q). 

Como puede ver, el procesador hace frente a esta carga muy bien, mientras que la latencia para una solicitud por GPU aumenta significativamente con un aumento en el número de solicitudes paralelas.


El tiempo para realizar cálculos matemáticos en la GPU y la CPU con matrices de 10,000 x 60 en ms.

En el segundo gráfico, los mismos cálculos, pero con matrices 10 veces más largas, y la GPU se comporta mucho mejor bajo tal carga. Estos gráficos son muy indicativos, y podemos concluir: el comportamiento bajo carga depende de la naturaleza de la carga misma. Un procesador también puede manejar cálculos matriciales de manera bastante eficiente, pero hasta cierto punto. Para una tarjeta de video, es característico que para una carga informática pequeña, el rendimiento se reduzca aproximadamente linealmente. Con un aumento en la carga y el número de hilos paralelos, la tarjeta de video se adapta mejor. 

Es difícil plantear la hipótesis de cómo se comportará la GPU en diversas situaciones, pero como puede ver, bajo ciertas condiciones, una tarjeta de servidor puede procesar solicitudes de varios flujos paralelos de manera bastante eficiente.

Discutiremos algunas preguntas más que puede tener si aún decide usar la GPU en sus proyectos.

Límite de recursos


Como ya dijimos, los dos recursos principales de una tarjeta de video son los núcleos informáticos y la memoria.

Por ejemplo, tenemos varios procesos o contenedores que usan una tarjeta de video, y nos gustaría poder compartir la tarjeta de video entre ellos. Desafortunadamente, no hay una API simple para esto. NVIDIA ofrece tecnología vGPU , pero no encontré la tarjeta Tesla k80 en la lista de compatibles, y hasta donde puedo entender por la descripción, la tecnología está más enfocada en pantallas virtuales que en cálculos. Quizás AMD ofrece algo más adecuado.

Por lo tanto, si planea usar la GPU en sus proyectos, debe confiar en el hecho de que la aplicación usará la tarjeta de video exclusivamente, o controlará mediante programación la cantidad de memoria asignada y la cantidad de núcleos utilizados para los cálculos.

Contenedores y GPU


Si descubrió el límite de recursos, entonces la siguiente pregunta lógica: ¿qué pasa si hay varias tarjetas de video en el servidor?

Nuevamente, puede decidir a nivel de aplicación qué GPU usará.

Otra forma más conveniente son los contenedores Docker. Puede usar contenedores regulares, pero NVIDIA ofrece sus contenedores NGC , con versiones optimizadas de varios software, bibliotecas y controladores. Para un contenedor, puede limitar la cantidad de GPU utilizadas y su visibilidad al contenedor. Los gastos generales en el uso de contenedores son aproximadamente del 3%.

Trabajar en un clúster


Otra pregunta, ¿qué hacer si desea realizar una tarea en varias GPU dentro del mismo servidor o clúster?

Si elige una biblioteca similar a empuje o una solución de nivel inferior, entonces la tarea deberá resolverse manualmente. Los marcos de alto nivel, por ejemplo, para el aprendizaje automático o las redes neuronales, por lo general admiten la capacidad de usar varias tarjetas de fábrica.

Además, me gustaría señalar que, por ejemplo, NVIDIA ofrece una interfaz para el intercambio directo de datos entre tarjetas: NVLINK , que es significativamente más rápido que PCI Express. Y hay tecnología para el acceso directo a la memoria de la tarjeta desde otros dispositivos PCI Express: GPUDirect RDMA , incl. y red .

Recomendaciones


Si está considerando usar la GPU en sus proyectos, entonces la GPU es más adecuada para usted si:

  • Su tarea puede reducirse a una vista SIMD
  • Es posible cargar la mayoría de los datos en el mapa antes de los cálculos (caché)
  • El desafío implica la computación intensiva

También debe hacer preguntas por adelantado:

  • ¿Cuántas consultas paralelas serán 
  • ¿Qué latencia esperas?
  • ¿Necesita una tarjeta para su carga? ¿Necesita un servidor con varias tarjetas o un grupo de servidores GPU? 

¡Eso es todo, espero que el material te sea útil y te ayude a tomar la decisión correcta!

Referencias


Benchmark y resultados en github - https://github.com/tishden/gpu_benchmark/tree/master/cuda

Además del tema, una grabación del informe "Bases de datos de GPU - Arquitectura, rendimiento y perspectivas de uso"

Seminario web NVIDIA NGC Containers - http : //bit.ly/2UmVIVt o http://bit.ly/2x4vJKF

All Articles