Tiny GPU: la GPU mínima implementada en Verilog

• tiny-gpu es una implementación mínima de GPU basada en Verilog para aprender desde cero cómo funciona una GPU a nivel de hardware, y se enfoca en los principios comunes entre GPGPU y los aceleradores de ML más que en hardware dedicado a gráficos
• La implementación está compuesta por menos de 15 archivos Verilog documentados, documentación de la arquitectura y la ISA, kernels de suma y multiplicación de matrices, y soporte para simulación de kernels y trazas de ejecución
• La GPU ejecuta un solo kernel a la vez; se cargan la memoria de programa y la memoria de datos, se configura thread_count y luego se activa la señal start para iniciar el kernel
• Para simplificar, cada core procesa un solo block a la vez y cada hilo tiene su propia ALU, LSU, PC y archivo de registros, pero se asume que todos los hilos convergen al mismo PC después de cada instrucción
• Se excluyen deliberadamente la mayoría de las funciones de las GPU modernas, como cachés multinivel, memoria compartida, memory coalescing, pipelining, warp scheduling, branch divergence y barriers, para priorizar una estructura orientada al aprendizaje

---

El problema que tiny-gpu intenta resolver

• Hay mucho material para aprender CPU desde la arquitectura hasta las señales de control, pero los detalles técnicos de bajo nivel de las GPU modernas siguen siendo en su mayoría propietarios debido a lo competitivo del mercado
• Hay muchos recursos sobre programación de GPU, pero casi no hay material para aprender cómo funciona una GPU a nivel de hardware
• Implementaciones open source de GPU como Miaow y VeriGPU buscan completar funcionalidad y comportamiento, por lo que su estructura es compleja
• tiny-gpu elimina gran parte de la complejidad de una tarjeta gráfica de nivel producción y se concentra en los elementos clave comunes a los aceleradores de hardware modernos
  • Componentes importantes de la arquitectura de GPU
  • Cómo se implementa en hardware el modelo de programación SIMD
  • Cómo una GPU maneja el ancho de banda de memoria limitado

Arquitectura general

• tiny-gpu está diseñada para ejecutar un solo kernel a la vez
• El procedimiento de ejecución del kernel es el siguiente
  • Cargar el código del kernel en la memoria global de programa
  • Cargar los datos necesarios en la memoria de datos
  • Indicar en los registros de control del dispositivo cuántos hilos se van a ejecutar
  • Poner la señal start en high para ejecutar el kernel
• La GPU está compuesta por las siguientes unidades
  • Registros de control del dispositivo
  • Dispatcher
  • Un número variable de cores de cómputo
  • Controladores de memoria para la memoria de datos y la memoria de programa
  • Caché

Ejecución del kernel y distribución de hilos

• Los registros de control del dispositivo almacenan los metadatos de ejecución del kernel; en tiny-gpu solo guardan thread_count, que es el número total de hilos a ejecutar
• El dispatcher distribuye los hilos entre varios cores de cómputo cuando comienza el kernel
  • Agrupa los hilos que pueden ejecutarse en paralelo en blocks
  • Envía los blocks a los cores disponibles para que los procesen
  • Cuando termina el procesamiento de todos los blocks, indica que la ejecución del kernel finalizó
• El core simplificado procesa un solo block a la vez
• Cada hilo tiene una ALU, LSU, PC y archivo de registros dedicados
• Gestionar la ejecución de instrucciones de los hilos sobre estos recursos es uno de los problemas difíciles en una GPU

Estructura de memoria y controladores

• La GPU está hecha para interactuar con memoria global externa y, para simplificar, separa la memoria de datos de la memoria de programa
• Especificaciones de la memoria de datos
  • Direccionamiento de 8 bits
  • 256 filas en total
  • Datos de 8 bits
  • Cada fila almacena valores menores a 256
• Especificaciones de la memoria de programa
  • Direccionamiento de 8 bits
  • 256 filas en total
  • Datos de 16 bits
  • Según la ISA, cada instrucción ocupa 16 bits
• El controlador de memoria rastrea las solicitudes de memoria que llegan desde los cores, limita las solicitudes de acuerdo con el ancho de banda real de la memoria externa y entrega las respuestas al recurso correcto
• Cada controlador de memoria tiene un número fijo de canales según el ancho de banda de la memoria global
• La caché es una función en desarrollo que guarda en la SRAM del dispositivo los datos traídos de memoria externa para recuperarlos más rápido en solicitudes posteriores y dejar el ancho de banda de memoria disponible para datos nuevos

Estructura interna del core

• Cada core tiene un único scheduler que administra la ejecución de los hilos
• El scheduler de tiny-gpu ejecuta hasta el final las instrucciones de un block y luego toma un nuevo block, ejecutando las instrucciones de todos los hilos en orden sincronizado
• En schedulers más avanzados se puede mejorar el aprovechamiento de recursos con pipelining y warp scheduling
• La principal limitación del scheduler es la latencia que se produce al cargar y almacenar datos en la memoria global
  • La mayoría de las instrucciones se pueden ejecutar de manera síncrona
  • Operaciones load-store como LDR y STR son asíncronas, así que la ejecución de instrucciones debe organizarse alrededor de esos largos tiempos de espera
• El Fetcher obtiene de manera asíncrona desde la memoria de programa la instrucción del contador de programa actual
• El Decoder decodifica la instrucción obtenida en señales de control para la ejecución de los hilos
• El archivo de registros de cada hilo guarda los datos con los que se está calculando y hace posible el patrón SIMD
  • Los registros de solo lectura incluyen %blockIdx, %blockDim, %threadIdx
  • El kernel puede ejecutarse con datos distintos según el ID local del hilo
• La ALU de cada hilo procesa las instrucciones aritméticas ADD, SUB, MUL, DIV
• CMP produce si el resultado de la diferencia entre dos registros es negativo, cero o positivo, y guarda el resultado en el registro NZP de la unidad PC
• La LSU de cada hilo accede a la memoria global de datos y maneja LDR, STR y la latencia asíncrona de memoria
• El PC de cada hilo determina la siguiente instrucción a ejecutar
  • Por defecto aumenta en 1 con cada instrucción
  • BRnzp salta a una fila específica de la memoria de programa si se cumple la condición del registro NZP configurado por el CMP anterior
  • Los loops y condicionales se implementan de esta manera
• Para simplificar, tiny-gpu asume que todos los hilos convergen al mismo PC después de cada instrucción
• En una GPU real, hilos individuales pueden bifurcarse hacia distintos PC, y en ese caso el grupo de hilos que se estaba procesando junto se divide en varios flujos de ejecución, lo que produce branch divergence

ISA

• tiny-gpu implementa una ISA de 11 instrucciones para ejecutar kernels simples de prueba de concepto, como suma y multiplicación de matrices
• Instrucciones soportadas
  • BRnzp: salta a otra fila de la memoria de programa si se cumple la condición NZP
  • CMP: compara los valores de dos registros y guarda el resultado en el registro NZP
  • ADD, SUB, MUL, DIV: operaciones aritméticas básicas para matemáticas de tensores
  • LDR: carga datos desde la memoria global
  • STR: guarda datos en la memoria global
  • CONST: carga un valor constante en un registro
  • RET: señala el fin de la ejecución del hilo actual
• Cada registro se especifica con 4 bits, por lo que hay un total de 16 registros
  • 13 de ellos, de R0 a R12, son registros de propósito general de lectura y escritura
  • Los últimos 3 son registros especiales de solo lectura que proporcionan %blockIdx, %blockDim, %threadIdx, necesarios para SIMD

Flujo de ejecución

• Cada core sigue el siguiente flujo de control por etapas al ejecutar instrucciones
  • FETCH: obtener la siguiente instrucción del PC actual
  • DECODE: decodificar la instrucción en señales de control
  • REQUEST: solicitar datos a la memoria global si se requiere LDR o STR
  • WAIT: esperar la respuesta de la memoria global si es necesario
  • EXECUTE: realizar el cálculo sobre los datos
  • UPDATE: actualizar el archivo de registros y el registro NZP
• Este flujo de control está organizado para favorecer la simplicidad y la comprensión
• En una implementación real, algunas etapas podrían comprimirse para optimizar el tiempo de procesamiento, o coordinar la ejecución de varias instrucciones en los recursos del core mediante pipelining
• Cada hilo realiza sus cálculos siguiendo la misma ruta de ejecución sobre los datos de su archivo de registros dedicado
• Es similar a un diagrama de CPU, pero se diferencia en que %blockIdx, %blockDim, %threadIdx están en registros de solo lectura para habilitar la funcionalidad SIMD

Kernels de ejemplo

• Se escribieron kernels de suma y multiplicación de matrices como prueba de concepto de la ISA
• Los archivos de prueba del repositorio pueden simular completamente estos kernels en la GPU y generar el estado de la memoria de datos y una traza completa de ejecución

• Suma de matrices

  • matadd.asm suma dos matrices de 1 x 8
  • Las 8 sumas por elemento se ejecutan cada una en un hilo distinto
  • Muestra la programación SIMD usando los registros %blockIdx, %blockDim, %threadIdx
  • Incluye gestión de memoria asíncrona mediante las instrucciones LDR y STR

• Multiplicación de matrices

  • matmul.asm multiplica dos matrices de 2 x 2
  • Calcula por elemento el producto punto entre la fila y la columna correspondientes
  • Muestra ramificación dentro del hilo usando CMP y BRnzp
  • Como todas las bifurcaciones vuelven a converger, funciona con la implementación actual de tiny-gpu

Simulación

• Para ejecutar la simulación de kernels se necesita iverilog y cocotb
• Procedimiento de preparación
  • Instalar el compilador Verilog y cocotb con brew install icarus-verilog y pip3 install cocotb
  • Descargar y descomprimir la última versión de sv2v y agregar el binario a $PATH
  • Ejecutar mkdir build en la raíz del repositorio
• La simulación de kernels se ejecuta con make test_matadd y make test_matmul
• Los resultados de ejecución se escriben en archivos de log dentro de test/logs
  • Estado inicial de la memoria de datos
  • Traza completa de ejecución del kernel
  • Estado final de la memoria de datos
• Al inicio de cada archivo de log se ven las matrices de entrada, y al final en la memoria de datos final se ve la matriz resultado
• La traza de ejecución incluye el estado de ejecución de todos los hilos de todos los cores en cada ciclo
  • Instrucción actual
  • PC
  • Valores de registros
  • Información de estado

Funciones avanzadas de GPU omitidas intencionalmente

• tiny-gpu deja fuera la mayoría de los elementos de rendimiento y funcionalidad de las GPU modernas para simplificar

• Cachés multinivel y memoria compartida

  • Las GPU modernas usan varias capas de caché para reducir el acceso a la memoria global
  • tiny-gpu solo implementa una única capa de caché que guarda datos recientes entre los recursos solicitantes y el controlador de memoria
  • Las cachés multinivel almacenan datos de uso frecuente más cerca del lugar donde se usan para reducir el tiempo de carga
  • Las GPU también pueden usar memoria compartida para que los hilos del mismo block intercambien resultados compartidos

• Memory coalescing

  • Durante la ejecución paralela, varios hilos suelen acceder a direcciones contiguas, como elementos adyacentes de una matriz
  • El memory coalescing analiza las solicitudes de memoria en cola y combina solicitudes contiguas en una sola transacción
  • El objetivo es reducir el tiempo dedicado al direccionamiento y procesar las solicitudes en conjunto

• Pipelining

  • Los cores de tiny-gpu no comienzan la siguiente instrucción hasta que termina la ejecución de una instrucción de un grupo de hilos
  • Las GPU modernas hacen streaming de varias ejecuciones de instrucciones secuenciales mientras garantizan ejecución ordenada para instrucciones con dependencias
  • Esto mejora el aprovechamiento de recursos para que el core no quede ocioso, por ejemplo mientras espera solicitudes de memoria asíncronas

• Warp scheduling

  • Divide un block en warps, que son lotes de hilos que pueden ejecutarse juntos
  • Cuando un warp está en espera, ejecuta instrucciones de otro warp para procesar varios warps a la vez en un solo core
  • Es parecido al pipelining, pero trabaja con instrucciones de hilos diferentes

• Branch divergence

  • tiny-gpu asume que todos los hilos de un mismo lote están en el mismo PC después de cada instrucción
  • En la práctica, hilos individuales pueden bifurcarse a distintas líneas según los datos
  • Los hilos con distintos PC se separan en flujos de ejecución independientes y también hay que gestionar el momento en que vuelven a converger

• Sincronización y barriers

  • Las GPU modernas pueden establecer barriers para que un grupo de hilos del mismo block espere hasta que todos lleguen a un punto específico
  • Esto es útil para garantizar que el procesamiento de datos se haya completado cuando los hilos necesitan intercambiar datos compartidos

Próximos pasos

• Las mejoras futuras contempladas son las siguientes
  • Agregar una caché de instrucciones simple
  • Construir un adaptador para poder usar la GPU en Tiny Tapeout 7
  • Agregar branch divergence básica
  • Agregar memory coalescing básico
  • Agregar pipelining básico
  • Optimizar el flujo de control y el uso de registros para mejorar el tiempo de ciclo
  • Escribir kernels gráficos básicos o agregar hardware gráfico simple para mostrar capacidades gráficas
• Quienes quieran mejorar el repositorio pueden contribuir mediante un PR

[GN⁺]

Opiniones de Hacker News

• El mercado de las GPU es tan competitivo que la mayoría de los detalles técnicos de bajo nivel de las arquitecturas modernas siguen siendo privados.
  Como excepción, Intel publica mucha documentación técnica sobre sus GPU: https://kiwitree.net/~lina/intel-gfx-docs/prm/
  También se pueden encontrar en línea los manuales de i810/815 y, salvo un extraño vacío en el período anterior a 965 donde faltan 855/910/915/945, la documentación ha sido bastante constante.
  • AMD también publica bastante documentación: https://www.amd.com/en/developer/browse-by-resource-type/documentation.html
    Incluye incluso documentos de arquitectura del conjunto de instrucciones de productos actuales y pasados, pero parece más orientada a implementadores que a explicaciones de alto nivel para aficionados interesados.
  • El driver de Linux de Intel también es de buena calidad y está en mainline.
    Ojalá todas las compañías siguieran este enfoque.
  • Es material de 2018, pero tiene cierta relación: The Thirty Million Line Problem - Casey Muratori
• Es un proyecto realmente genial, y da gusto ver este tipo de proyectos de hardware desarrollarse en público.
  Dicho eso, creo que esto se parece más a un coprocesador SIMD.
  Para llamarlo GPU, creo que debería tener al menos algún tipo de salida de pantalla.
  Sé que el término se ha vuelto bastante flexible últimamente, con Nvidia y otros vendiendo como GPU variantes de arquitecturas gráficas solo para servidores, pero la parte gráfica del diseño de una GPU sigue representando una porción considerable de la complejidad.
  • Si procesa gráficos, creo que puede considerarse una GPU aunque no tenga salida.
    Una GPU que no genera salida sigue siendo útil.
    En mi trabajo hay unas 75 estaciones de trabajo con Quadro de gama media; las tarjetas solo tienen mini-DisplayPort y la empresa solo compra cables HDMI, así que todas están conectadas a la gráfica integrada.
    Aun así, esas tarjetas aceleran software y procesan gráficos; simplemente no muestran nada en pantalla.
• Excelente. Apoyo mucho el trabajo en GPU de núcleo abierto.
  Hay otro ejemplo: https://github.com/jbush001/NyuziProcessor
  • Sería bueno tener una implementación mínima de CUDA para alguno de estos procesadores de núcleo abierto.
    ¿Qué volumen haría falta para que TSMC u otra fundición pudiera producir económicamente este tipo de procesador?
• Es un proyecto realmente excelente.
  Quiero probar con FPGA, pero sinceramente es difícil siquiera saber por dónde empezar, y todo el campo se siente bastante intimidante.
  Mi objetivo final es hacer una tarjeta aceleradora para LLM, que aunque es una meta totalmente arbitraria, parece tener mucho en común con este proyecto; probablemente solo diferiría en la parte de offloading de memoria para cargar modelos más grandes.
  • Hay que cambiar el marco mental.
    La introducción a FPGA debe dividirse en varias subhabilidades, y también hay que ajustar las expectativas.
    No esperaríamos que un ingeniero de software empezara construyendo una computadora completa desde los principios básicos, escribiera una arquitectura de conjunto de instrucciones, entendiera código máquina, lo convirtiera en ensamblador y luego desarrollara un lenguaje de programación para crear aplicaciones en Python.
    Lo correcto es empezar arriba e ir bajando por la pila.
    Si abstraes la complejidad y te enfocas en construir sistemas con IP prefabricada, el diseño con FPGA es bastante fácil.
    Normalmente recomendaría algo como MATLAB, porque con HDL Coder puedes crear una aplicación inicial en un DevKit que tenga un reference design.
    De lo contrario, aparece una enorme carga de aprender arquitectura de cómputo digital, Verilog, timing, transceptores/E/S, planificación de pines, Quartus/Vivado, simulación/verificación, sistemas embebidos, etc.
    En resumen, empieza con diseño a nivel de sistema, aprende a tomar IP plug-and-play y conectarla en el nivel superior, y coloca ese módulo dentro de un diseño de referencia ya preparado.
    Después puedes ir quitando capas gradualmente para exponer la complejidad que hay debajo.
  • Estoy en la misma situación, y mi plan es este:
    1. Leer Digital Design and Computer Architecture, de Harris y Harris. (2022). Elsevier: https://doi.org/10.1016/c2019-0-00213-0
    2. Seguir el curso RVFpga de los autores y construir una CPU RISC-V real sobre una FPGA: https://www.youtube.com/watch?v=ePv3xD3ZmnY
  • Recomiendo esta ruta:
    1. Clonar el repositorio educativo https://github.com/yuri-panchul/basics-graphics-music. Es una colección de prácticas simples para quienes aprenden Verilog desde cero, escrita por Yuri Panchul, quien trabajó en Imagination en desarrollo de GPU.
    2. Conseguir una de las decenas de placas FPGA compatibles y accesorios como botones y LED.
    3. Instalar Yosys y las herramientas relacionadas.
    4. Empezar por lab01 DeMorgan y hacer tantas prácticas del repositorio como sea posible.
       Puedes combinar las prácticas con la lectura de Harris&Harris.
       Cuando termines las prácticas y el libro, será momento de iniciar tu propio proyecto.
       Como referencia, en HackerMojo también hay reuniones semanales, y puedes participar por Zoom aunque no estés en el Valley.
  • No sé en qué etapa estás, pero estos recursos me ayudaron a entender mejor la lógica digital y la arquitectura de CPU/GPU.
    1. https://learn.saylor.org/course/CS301
    2. https://www.coursera.org/learn/comparch

3. https://hdlbits.01xz.net/wiki/Main_Page

• Si quieres acelerar LLMs, primero tienes que entender la arquitectura.
  Puedes empezar por ahí.
  El hardware en realidad es la parte fácil, y también la parte difícil desde el punto de vista de la fabricación.
• ¿Hay alguna razón para mezclar operadores de asignación non-blocking y asignación blocking en este bloque always secuencial?
  • Eso parece una variable local.
  • Si no te obsesiona demasiado que coincidan los resultados de simulación y síntesis, está bien hacerlo así.
• Hace mucho hice algo parecido en VHDL.
  Había un sitio llamado opencores que reunía varios proyectos HDL open source.
  Me pregunto si hoy existen buenos simuladores HDL distribuidos a gran escala, de nivel HPC.
  Parece razonable usar GPUs modernas para simulación a nivel RTL.
  • No es que “existía”; todavía existe: https://opencores.org/projects?language=VHDL
    ¿O era otro sitio parecido y no el mismo?
• ¿La ALU implementa la instrucción DIV tal cual a nivel de hardware?
  ¿Es normal que algo como un core CUDA moderno tenga división como instrucción real, o normalmente se emula por software?
  Como un circuito de división en hardware ocupa muchísimo espacio, no esperaba que estuviera dentro de una ALU de GPU.
  En Verilog es demasiado fácil escribir una sola línea como DIV: begin alu_out_reg <= rs / rt; end, pero esa línea consume mucho silicio.
  Si solo simulas Verilog, quizá no te des cuenta de eso.
  • Esto es simplemente un proyecto de alguien aprendiendo Verilog.
    El proyecto se queda en la simulación, y para convertirlo en hardware real haría falta mucho más trabajo.
• Además, es una “GPU” sin funciones gráficas.
  Personalmente creo que a esto habría que llamarlo de otra forma.
  • La primera pregunta es por qué CPU y GPU se separaron en primer lugar.
    La brecha entre ambas se está reduciendo y las dos partes están agregando funciones de la otra, pero todavía hay diferencias considerables.
    En mi opinión tiene que ver con la ley de Amdahl [0].
    En ese sentido, se podría llamar a la CPU un procesador optimizado para latencia, y a la GPU un procesador optimizado para throughput.
    Más concretamente, [1] también se podría llamar a la CPU un procesador de dependencias de datos largas y profundas, y a la GPU uno de dependencias de datos anchas y planas.
    [0]: https://en.wikipedia.org/wiki/Amdahl%27s_law
    [1]: https://en.wikipedia.org/wiki/Data_dependency
  • Se podría llamar TPU, es decir, unidad de procesamiento de tensores.
    Un tensor es simplemente un arreglo n-dimensional.
    Encima se puede poner software o firmware para hacerlo comportarse como una GPU.
  • He estado pensando en iniciar un proyecto para hacer un “adaptador de pantalla”, pero antes de empezar me quedé trabado porque no logré entender el protocolo de comunicación entre el driver GOP de UEFI y el adaptador de pantalla.
    Intenté armar las piezas a partir del código fuente de EDK2, pero no queda claro cuánto de eso es específico de QEMU.
  • Se le puede llamar MPU, es decir, unidad de procesamiento de matrices.
  • Creo que el término que se está imponiendo es AIA, es decir, acelerador de IA.
• Que tiny-gpu asuma que todos los hilos “convergen” al mismo contador de programa después de cada instrucción es una simplificación demasiado ingenua.
  En una GPU real, los hilos individuales pueden ramificarse a distintos PC, y aparece divergencia de ramas, donde un grupo de hilos que al principio se procesaba junto se separa en ejecuciones distintas.
  Habría sido mejor que hiciera programación de GPU antes de intentar hacer una GPU en silicio.
  Además, tampoco parece correcto llamarlo SIMD.
  Esta es la misma persona que antes conectó circuitos de otros para hacer parpadear un LED y dijo que había hecho una CPU.
  • ¿Lo primero no equivale a llamar __syncthreads() en cada ejecución?