Tiny GPU: a menor GPU implementada em Verilog

 (github.com/adam-maj)
2 pontos por GN⁺ 2024-04-27 | 1 comentários | Compartilhar no WhatsApp
  • tiny-gpu é uma implementação mínima de GPU baseada em Verilog para aprender, do zero, como uma GPU funciona no nível de hardware, com foco nos princípios em comum entre GPGPU e aceleradores de ML, e não em hardware dedicado a gráficos
  • A implementação é composta por menos de 15 arquivos Verilog documentados, documentação da arquitetura e da ISA, kernels de soma e multiplicação de matrizes, além de suporte para simulação de kernels e rastreamento de execução
  • A GPU executa um único kernel por vez; ela carrega a memória de programa e a memória de dados, define thread_count e então aciona o sinal start para iniciar o kernel
  • Para simplificar, cada core processa um block por vez e cada thread tem ALU, LSU, PC e arquivo de registradores próprios, mas assume-se que todas as threads convergem para o mesmo PC após cada instrução
  • Recursos como cache multinível, memória compartilhada, coalescimento de memória, pipelining, escalonamento de warps, divergência de ramificação e barreiras, comuns em GPUs modernas, foram em grande parte deixados de fora para priorizar uma estrutura voltada ao aprendizado

O problema que o tiny-gpu quer resolver

  • Há muito material para aprender CPUs, da arquitetura aos sinais de controle, mas os detalhes técnicos de baixo nível das GPUs modernas continuam em sua maioria proprietários por causa de um mercado competitivo
  • Existe bastante material sobre programação para GPU, mas quase nada para aprender como uma GPU funciona no nível de hardware
  • Implementações open source de GPU como Miaow e VeriGPU buscam completude funcional e execução real, o que torna sua estrutura complexa
  • tiny-gpu remove grande parte da complexidade de placas gráficas de nível de produção e se concentra nos elementos centrais comuns aos aceleradores de hardware modernos
    • Componentes importantes da arquitetura de GPU
    • Como o modelo de programação SIMD é implementado no hardware
    • Como a GPU lida com largura de banda de memória limitada

Arquitetura geral

  • O tiny-gpu foi projetado para executar apenas um kernel por vez
  • O procedimento de execução de um kernel é o seguinte
    • Carregar o código do kernel na memória global de programa
    • Carregar os dados necessários na memória de dados
    • Definir no registrador de controle do dispositivo o número de threads a executar
    • Colocar o sinal start em nível alto para executar o kernel
  • A GPU é composta pelas seguintes unidades
    • Registrador de controle do dispositivo
    • Dispatcher
    • Um número variável de cores de computação
    • Controladores de memória para a memória de dados e a memória de programa
    • Cache

Execução de kernel e distribuição de threads

  • O registrador de controle do dispositivo armazena metadados da execução do kernel e, no tiny-gpu, guarda apenas thread_count, o número total de threads a executar
  • O dispatcher distribui as threads entre vários cores de computação quando o kernel começa
    • Agrupa threads que podem ser executadas em paralelo em um block
    • Envia os blocks para os cores disponíveis
    • Sinaliza a conclusão da execução quando todos os blocks terminam
  • O core simplificado processa um block por vez
  • Cada thread possui ALU, LSU, PC e arquivo de registradores dedicados
  • Gerenciar a execução das instruções das threads sobre esses recursos é um dos problemas difíceis de uma GPU

Estrutura de memória e controladores

  • A GPU foi feita para se conectar à memória global externa e, para simplificar, separa memória de dados e memória de programa
  • Especificação da memória de dados
    • Endereçamento de 8 bits
    • 256 linhas no total
    • Dados de 8 bits
    • Cada linha armazena valores menores que 256
  • Especificação da memória de programa
    • Endereçamento de 8 bits
    • 256 linhas no total
    • Dados de 16 bits
    • Cada instrução tem 16 bits, de acordo com a ISA
  • O controlador de memória rastreia as requisições vindas dos cores, limita os pedidos de acordo com a largura de banda real da memória externa e entrega as respostas ao recurso correto
  • Cada controlador de memória tem um número fixo de canais de acordo com a largura de banda da memória global
  • O cache está em desenvolvimento e armazena em SRAM interna os dados trazidos da memória externa, permitindo recuperá-los mais rápido em requisições futuras e usar a largura de banda para dados novos

Estrutura interna do core

  • Cada core tem um único scheduler para gerenciar a execução das threads
  • O scheduler do tiny-gpu executa até o fim as instruções de um block antes de buscar um novo block e executa as instruções de todas as threads em ordem sincronizada
  • Em schedulers mais avançados, pipelining e escalonamento de warps podem aumentar a taxa de uso dos recursos
  • A principal limitação do scheduler é a latência ao carregar e armazenar dados na memória global
    • A maioria das instruções pode ser executada de forma síncrona
    • Operações load-store como LDR e STR são assíncronas, então a execução precisa ser organizada em torno desses tempos de espera longos
  • O Fetcher busca de forma assíncrona, na memória de programa, a instrução correspondente ao contador de programa atual
  • O Decoder decodifica a instrução buscada em sinais de controle para a execução das threads
  • O arquivo de registradores de cada thread guarda os dados em cálculo e viabiliza o padrão SIMD
    • Os registradores somente leitura contêm %blockIdx, %blockDim e %threadIdx
    • O kernel pode ser executado com dados diferentes conforme o ID local da thread
  • A ALU de cada thread processa as instruções aritméticas ADD, SUB, MUL e DIV
  • CMP informa se o resultado da diferença entre dois registradores é negativo, zero ou positivo e salva esse resultado no registrador NZP da unidade de PC
  • A LSU de cada thread acessa a memória global de dados e lida com LDR, STR e com a latência assíncrona da memória
  • O PC de cada thread determina a próxima instrução a ser executada
    • Por padrão, ele é incrementado em 1 a cada instrução
    • BRnzp desvia para uma linha específica da memória de programa se a condição do registrador NZP, definida pelo CMP anterior, for satisfeita
    • Loops e condicionais são implementados dessa forma
  • Para simplificar, o tiny-gpu assume que todas as threads convergem para o mesmo PC após cada instrução
  • Em GPUs reais, threads individuais podem desviar para PCs diferentes e, nesse caso, ocorre divergência de ramificação, em que um grupo de threads que era processado junto se divide em vários fluxos de execução

ISA

  • O tiny-gpu implementa uma ISA de 11 instruções para executar kernels simples de prova de conceito, como soma e multiplicação de matrizes
  • Instruções suportadas
    • BRnzp: salta para outra linha da memória de programa se a condição NZP for satisfeita
    • CMP: compara os valores de dois registradores e salva o resultado no registrador NZP
    • ADD, SUB, MUL, DIV: operações aritméticas básicas para matemática de tensores
    • LDR: carrega dados da memória global
    • STR: armazena dados na memória global
    • CONST: carrega um valor constante em um registrador
    • RET: sinaliza o fim da execução da thread atual
  • Cada registrador é especificado com 4 bits, totalizando 16 registradores
    • De R0 a R12, 13 são registradores livres de leitura e escrita
    • Os 3 últimos são registradores especiais somente leitura que fornecem %blockIdx, %blockDim e %threadIdx, necessários para SIMD

Fluxo de execução

  • Ao executar instruções, cada core segue o seguinte fluxo de controle
    • FETCH: buscar a próxima instrução no PC atual
    • DECODE: decodificar a instrução em sinais de controle
    • REQUEST: solicitar dados à memória global quando LDR ou STR forem necessários
    • WAIT: aguardar a resposta da memória global quando necessário
    • EXECUTE: realizar o cálculo sobre os dados
    • UPDATE: atualizar o arquivo de registradores e o registrador NZP
  • Esse fluxo de controle foi montado para favorecer simplicidade e compreensão
  • Em implementações reais, algumas etapas podem ser comprimidas para otimizar o tempo de processamento, ou coordenadas por pipelining para executar várias instruções nos recursos do core
  • Cada thread realiza o cálculo seguindo o mesmo caminho de execução sobre os dados do seu arquivo de registradores dedicado
  • O diagrama lembra o de uma CPU, mas se diferencia por ter %blockIdx, %blockDim e %threadIdx em registradores somente leitura, o que viabiliza a funcionalidade SIMD

Kernels de exemplo

  • Para provar os conceitos da ISA, foram escritos kernels de soma e multiplicação de matrizes
  • Os arquivos de teste do repositório conseguem simular completamente esses kernels na GPU e gerar o estado da memória de dados e o rastreamento completo da execução

  • Soma de matrizes

    • matadd.asm soma duas matrizes 1 x 8
    • Cada uma das 8 somas por elemento é executada em uma thread separada
    • Usa os registradores %blockIdx, %blockDim e %threadIdx para demonstrar programação SIMD
    • Inclui gerenciamento assíncrono de memória com instruções LDR e STR

  • Multiplicação de matrizes

    • matmul.asm multiplica duas matrizes 2 x 2
    • Calcula o produto interno das linhas e colunas correspondentes, elemento por elemento
    • Usa CMP e BRnzp para mostrar desvios dentro da thread
    • Como todos os desvios convergem novamente, isso funciona na implementação atual do tiny-gpu

Simulação

  • Para rodar a simulação dos kernels, são necessários iverilog e cocotb
  • Etapas de preparação
    • Instalar o compilador Verilog e o cocotb com brew install icarus-verilog e pip3 install cocotb
    • Baixar a versão mais recente de sv2v, descompactar e adicionar o binário ao $PATH
    • Executar mkdir build na raiz do repositório
  • A simulação dos kernels é executada com make test_matadd e make test_matmul
  • O resultado da execução é gravado em arquivos de log em test/logs
    • Estado inicial da memória de dados
    • Rastreamento completo da execução do kernel
    • Estado final da memória de dados
  • No início de cada arquivo de log aparecem as matrizes de entrada, e no fim, na memória de dados final, aparece a matriz resultante
  • O rastreamento da execução inclui, a cada ciclo, o estado de execução de todas as threads em todos os cores
    • Instrução atual
    • PC
    • Valores dos registradores
    • Informações de estado

Recursos avançados de GPU deixados de fora de propósito

  • O tiny-gpu deixa de fora a maior parte dos elementos de desempenho e funcionalidade das GPUs modernas para simplificar

  • Cache multinível e memória compartilhada

    • GPUs modernas usam várias camadas de cache para reduzir acessos à memória global
    • O tiny-gpu implementa apenas uma única camada de cache entre os recursos solicitantes e o controlador de memória, guardando dados recentes
    • Caches em várias camadas armazenam dados usados com frequência mais perto do local de uso para reduzir o tempo de carregamento
    • GPUs também podem usar memória compartilhada para permitir que threads do mesmo block troquem resultados compartilhados

  • Coalescimento de memória

    • Várias threads executadas em paralelo frequentemente acessam endereços contíguos, como elementos adjacentes de uma matriz
    • O coalescimento de memória analisa as requisições enfileiradas e combina requisições vizinhas em uma única transação
    • O objetivo é reduzir o tempo gasto com endereçamento e processar pedidos em conjunto

  • Pipelining

    • No core do tiny-gpu, a próxima instrução só começa depois que a execução de uma instrução termina para um conjunto de threads
    • GPUs modernas fazem streaming da execução de várias instruções sequenciais, garantindo execução em ordem quando há dependências
    • Isso aumenta o aproveitamento dos recursos e evita que eles fiquem ociosos durante esperas como as de requisições assíncronas à memória

  • Escalonamento de warps

    • Um block é dividido em warps, lotes de threads que podem ser executadas juntas
    • Quando um warp está esperando, instruções de outro warp podem ser executadas, permitindo tratar vários warps simultaneamente em um único core
    • É semelhante ao pipelining, mas lida com instruções de threads diferentes

  • Divergência de ramificação

    • O tiny-gpu assume que todas as threads de um mesmo lote têm o mesmo PC após cada instrução
    • Na prática, threads individuais podem desviar para linhas diferentes dependendo dos dados
    • Threads com PCs diferentes são divididas em fluxos de execução separados, e também é preciso gerenciar o ponto em que elas convergem novamente

  • Sincronização e barreiras

    • GPUs modernas podem definir barreiras para fazer grupos de threads do mesmo block esperarem até que todas cheguem a um ponto específico
    • Isso é útil para garantir que o processamento foi concluído quando as threads precisam trocar dados compartilhados

Próximos passos

  • As melhorias futuras incluem
    • Adicionar um cache de instruções simples
    • Construir um adaptador para usar a GPU no Tiny Tapeout 7
    • Adicionar divergência de ramificação básica
    • Adicionar coalescimento de memória básico
    • Adicionar pipelining básico
    • Otimizar o fluxo de controle e o uso de registradores para melhorar o tempo de ciclo
    • Escrever kernels gráficos básicos ou adicionar hardware gráfico simples para demonstrar recursos gráficos
  • Quem quiser melhorar o repositório pode contribuir com PRs

Leituras relacionadas

1 comentários

 
GN⁺ 2024-04-27
Opiniões no Hacker News
  • Como o mercado de GPUs é extremamente competitivo, a maioria dos detalhes técnicos de baixo nível das arquiteturas modernas permanece fechada
    Como exceção, a Intel publica bastante documentação técnica sobre GPUs: https://kiwitree.net/~lina/intel-gfx-docs/prm/
    Também é possível encontrar online os manuais do i810/815 e, tirando uma lacuna estranha no período anterior ao 965, sem o 855/910/915/945, a documentação foi bastante consistente
  • É um projeto realmente muito legal, e é bom ver projetos de hardware assim sendo desenvolvidos abertamente
    Ainda assim, eu diria que isto está mais para um coprocessador SIMD
    Para chamar de GPU, acho que deveria haver pelo menos algum tipo de saída de vídeo
    Sei que o termo ficou bem mais flexível recentemente, com Nvidia e outras vendendo também variantes de arquiteturas gráficas só para servidores como GPUs, mas a parte gráfica no projeto de uma GPU ainda responde por uma parcela considerável da complexidade
    • Se processa gráficos, acho que pode ser considerado uma GPU mesmo sem saída
      GPUs sem saída ainda são úteis
      No meu trabalho há cerca de 75 workstations com Quadro intermediárias, mas as placas só têm mini-DisplayPort e a empresa só compra cabos HDMI, então todas estão conectadas ao vídeo integrado
      Mesmo assim, essas placas aceleram software e processam gráficos; elas simplesmente não fazem a saída para a tela
  • Legal. Apoio muito o trabalho em GPUs open core
    Há outro exemplo também: https://github.com/jbush001/NyuziProcessor
    • Seria bom haver uma implementação mínima de CUDA para um desses processadores open core
      Qual volume seria necessário para fabricar economicamente um processador desses na TSMC ou em outra foundry?
  • Projeto realmente excelente
    Tenho vontade de mexer com FPGA, mas, sinceramente, é difícil até saber por onde começar, e a área inteira parece bastante intimidadora
    O objetivo final seria criar uma placa aceleradora para LLMs; embora seja uma meta totalmente arbitrária, acho que ela teria bastante sobreposição com este projeto, e talvez só diferisse na parte de offloading de memória para carregar modelos maiores
    • É preciso mudar o enquadramento mental
      A introdução a FPGA precisa ser dividida em várias subáreas, e as expectativas também precisam ser ajustadas
      Ninguém espera que um engenheiro de software, logo de início, construa um computador inteiro a partir dos princípios básicos, projete uma arquitetura de conjunto de instruções, entenda código de máquina, transforme isso em assembly e ainda desenvolva uma linguagem de programação para criar uma aplicação em Python
      O certo é começar por cima e ir descendo pela stack
      Se você abstrair a complexidade e se concentrar em construir sistemas com IPs prontos, o projeto em FPGA fica bastante fácil
      Normalmente eu recomendaria algo como MATLAB, porque em um DevKit com reference design é possível criar a aplicação inicial com HDL Coder
      Caso contrário, surge um fardo enorme de aprender arquitetura de computação digital, Verilog, timing, transceptores/I/O, planejamento de pinos, Quartus/Vivado, simulação/verificação, sistemas embarcados etc.
      Em resumo, comece pelo projeto em nível de sistema, aprenda a pegar IPs plug-and-play e conectá-los no nível superior, e depois coloque esse módulo em um projeto de referência pronto
      A partir daí, vá removendo as camadas gradualmente para expor a complexidade por baixo
    • Estou na mesma situação, e meu plano é este
      1. Ler Digital Design and Computer Architecture, de Harris e Harris. (2022). Elsevier: https://doi.org/10.1016/c2019-0-00213-0
      2. Seguir o curso RVFpga dos autores e construir uma CPU RISC-V real em um FPGA: https://www.youtube.com/watch?v=ePv3xD3ZmnY
    • Eu recomendaria este caminho
      1. Clonar o repositório educacional https://github.com/yuri-panchul/basics-graphics-music. É uma coleção de exercícios simples para quem está aprendendo Verilog do zero, escrita por Yuri Panchul, que trabalhou na Imagination em desenvolvimento de GPU
      2. Obter uma das dezenas de placas FPGA compatíveis e acessórios como botões e LEDs
      3. Instalar o Yosys e ferramentas relacionadas
      4. Começar pelo lab01 DeMorgan e fazer o máximo possível dos exercícios do repositório
        Dá para fazer os exercícios em paralelo com a leitura de Harris&Harris
        Depois de terminar os exercícios e o livro, será o momento de começar seu próprio projeto
        Como referência, também há encontros semanais no HackerMojo, e dá para participar por Zoom mesmo não estando no Valley
    • Não sei em que etapa você está, mas estes materiais me ajudaram a entender melhor lógica digital e arquitetura de CPU/GPU
      1. https://learn.saylor.org/course/CS301
      2. https://www.coursera.org/learn/comparch
  1. https://hdlbits.01xz.net/wiki/Main_Page
  • Se você quer acelerar LLMs, primeiro precisa entender arquitetura
    É por aí que se começa
    Hardware é, de fato, a parte fácil, e também a parte difícil do ponto de vista de fabricação
  • Há algum motivo para misturar operadores de atribuição non-blocking e atribuição blocking neste bloco always sequencial aqui?
    • Aquilo parece uma variável local
    • Se você não se preocupa tanto em fazer a simulação e o resultado da síntese baterem, dá para fazer assim
  • Muito tempo atrás, fiz algo parecido em VHDL
    Havia um site chamado opencores, que reunia vários projetos HDL open source
    Fico curioso se hoje existe algum simulador HDL distribuído em larga escala, no nível de HPC, que seja bom
    Parece fazer sentido usar GPUs modernas para simulação em nível RTL
  • A ALU implementa a instrução DIV diretamente no nível de hardware?
    Em algo como núcleos CUDA modernos, é comum haver divisão como instrução real, ou normalmente ela é emulada em software?
    Um circuito real de divisão em hardware ocupa uma área enorme, então eu não esperava que fosse entrar numa ALU de GPU
    Em Verilog é fácil demais escrever uma linha como DIV: begin alu_out_reg <= rs / rt; end, mas essa linha consome muito silício
    Se você estiver apenas simulando Verilog, talvez nem perceba isso
    • Isto é só um projeto de alguém aprendendo Verilog
      O projeto para na simulação; para transformá-lo em hardware real, seria necessário muito mais trabalho
  • Mais uma “GPU” sem funções gráficas
    Pessoalmente, acho que esse tipo de coisa deveria receber outro nome
    • A primeira pergunta é por que CPU e GPU se separaram em primeiro lugar
      A distância entre as duas está diminuindo, e ambas estão adicionando recursos uma da outra, mas ainda há diferenças consideráveis
      Na minha opinião, isso tem relação com a lei de Amdahl [0]
      Nesse sentido, uma CPU pode ser chamada de processador otimizado para latência, e uma GPU, de processador otimizado para throughput
      Mais especificamente, [1] também dá para chamar a CPU de processador de dependências de dados longas e profundas, e a GPU de processador de dependências de dados largas e rasas
      [0]: https://en.wikipedia.org/wiki/Amdahl%27s_law
      [1]: https://en.wikipedia.org/wiki/Data_dependency
    • Dá para chamar de TPU, ou unidade de processamento de tensores
      Um tensor é simplesmente um array n-dimensional
      Dá para colocar software ou firmware em cima disso e fazê-lo se comportar como uma GPU
    • Eu vinha pensando em começar um projeto para criar um “adaptador de vídeo”, mas travei antes mesmo de começar porque não consegui entender o protocolo de comunicação entre o driver GOP da UEFI e o adaptador de vídeo
      Tentei juntar as peças a partir do código-fonte do EDK2, mas não fica claro até que ponto aquilo é específico do QEMU
    • Dá para chamar de MPU, ou unidade de processamento de matrizes
    • Acho que o termo que está se consolidando é AIA, ou acelerador de IA
  • É uma simplificação ingênua demais o tiny-gpu presumir que todas as threads “convergem” para o mesmo contador de programa depois de cada instrução
    Em GPUs reais, threads individuais podem desviar para PCs diferentes, e ocorre divergência de ramificação, em que um grupo de threads que inicialmente era processado em conjunto se divide em execuções separadas
    Teria sido melhor fazer programação de GPU antes de tentar criar uma GPU em silício
    Além disso, nem parece muito correto chamar isso de SIMD
    Essa é a mesma pessoa que, no passado, conectou circuitos de outras pessoas para piscar um LED e disse que tinha feito uma CPU
    • O primeiro caso não seria equivalente a chamar __syncthreads() a cada execução?