SCALE - Executando CUDA em GPUs AMD sem modificações

• O SCALE é um toolkit de programação GPGPU que permite compilar nativamente aplicações CUDA para GPUs AMD
• Não é necessário modificar o programa CUDA nem o sistema de build, e o suporte a mais fabricantes de GPU e à API CUDA está em desenvolvimento

Como funciona?

• O SCALE traz algumas inovações importantes em comparação com outras soluções GPGPU multiplataforma
  • Aceita programas CUDA como estão. Não é necessário portar para outra linguagem. Isso também vale quando o programa usa asm PTX inline
  • O compilador do SCALE aceita as mesmas opções de linha de comando e o mesmo dialeto CUDA do nvcc, funcionando como um substituto drop-in
  • "Finge" ser uma instalação do NVIDIA CUDA Toolkit, para que ferramentas e scripts de build existentes continuem funcionando sem mudanças

Quais projetos foram testados?

• O SCALE foi validado compilando projetos CUDA de código aberto e executando os testes
• Atualmente, os seguintes projetos open source estão incluídos em testes automatizados noturnos e passam completamente
  • NVIDIA Thrust, Blender Cycles, AMGX, llama-cpp, faiss, xgboost, GOMC, stdgpu, hashcat

Quais GPUs são suportadas?

• Os seguintes alvos de GPU são suportados e incluídos nos testes Nightly
  • AMD gfx1030 (Navi 21, RDNA 2.0)
  • AMD gfx1100 (Navi 31, RDNA 3.0)
• Os seguintes alvos de GPU passaram por testes manuais ocasionais e "parecem funcionar"
  • AMD gfx1010
  • AMD gfx1101
• Há trabalho em andamento para o seguinte suporte de GPU
  • AMD gfx900 (Vega 10, GCN 5.0)
• Se você quiser rapidamente suporte para uma arquitetura específica de GPU AMD, entre em contato

Componentes do SCALE

• Compilador compatível com nvcc capaz de compilar CUDA no dialeto do nvcc para GPUs AMD, incluindo PTX asm
• Implementação da API de runtime e driver CUDA para GPUs AMD
• Biblioteca wrapper open source que fornece a API "CUDA-X" delegando para bibliotecas ROCm. Bibliotecas como cuBLAS e cuSOLVER são tratadas dessa forma

Diferenças entre o SCALE e outras soluções

• Em vez de oferecer uma nova forma de escrever software GPGPU, o SCALE permite compilar diretamente para GPUs AMD programas escritos na linguagem CUDA amplamente usada
• O SCALE tem como meta ser totalmente compatível com o NVIDIA CUDA. A ideia é que os usuários possam dar suporte a vários fabricantes de GPU sem manter múltiplas bases de código nem fazer concessões de desempenho
• A linguagem do SCALE é um superconjunto do NVIDIA CUDA e oferece extensões opcionais de linguagem que tornam a escrita de código GPU mais fácil e eficiente para usuários que queiram ir além do nvcc
• O SCALE é um trabalho em andamento. Se alguma API ausente impedir seu uso, entre em contato. As prioridades de desenvolvimento serão ajustadas

Resumo do GN⁺

• O SCALE é um toolkit importante que permite compilar nativamente aplicações CUDA para GPUs AMD
• Há uma grande vantagem para desenvolvedores, pois é possível executar em GPUs AMD sem precisar modificar programas CUDA existentes
• Como busca compatibilidade total com o NVIDIA CUDA, ele é vantajoso para dar suporte a vários fabricantes de GPU
• É um projeto em andamento, e se alguma API necessária estiver faltando, é possível contatar a equipe de desenvolvimento para ajustar prioridades
• Projetos com funcionalidade semelhante incluem ROCm e HIP

[GN⁺]

Comentários no Hacker News

• Muitas pessoas acham que a AMD deveria oferecer suporte a uma camada de tradução, mas há a opinião de que isso é uma má ideia

  • O CUDA não foi projetado para ser neutro em relação ao fornecedor, e a Nvidia poderia criar dificuldades técnicas e legais
  • Por exemplo, executar cuDNN ou cuBLAS por cima disso pode violar os contratos de licença
  • Essas bibliotecas da Nvidia passariam a fazer parte do limite de API que a AMD teria de reimplementar e oferecer suporte

• Há a opinião de que buscar compatibilidade com bugs é uma tolice

  • Os usuários importantes de CUDA são open source
  • A AMD poderia implementar suporte diretamente em projetos upstream como pytorch ou llama.cpp
  • Com suporte, a comunidade poderia fazer a manutenção

• Há a opinião de que é difícil entender como um código fortemente dependente de hardware poderia "simplesmente funcionar" na AMD

  • A maior parte do código CUDA sério considera o tamanho do arquivo de registradores e da memória compartilhada, instruções wgmma, layouts ideais de memória e registradores para tensor cores, instruções do tensor memory accelerator etc.

• Se for verdade, é impressionante, mas há a opinião de que não é open source e faltam detalhes precisos sobre como funciona

  • Não está claro por que hoje em dia se espera que projetos sejam open source ou pelo menos source-available

• Há a opinião de que a principal razão para a alta valorização da Nvidia é que a AMD não investe em tornar as GPUs úteis para ML

  • Pode ser que a AMD tema ações antitruste, ou que exista algo na sua abordagem de hardware que limite sua competitividade
  • A empresa parece ter deixado passar bilhões de dólares tanto no pico de demanda por GPUs para mineração de criptomoedas quanto no atual pico de demanda do boom de IA

• Há a opinião de que a AMD erra tanto que dá vontade de parabenizar projetos como este

  • Especialmente no Linux, é muito frustrante quando as capacidades existem fisicamente no notebook, mas não podem ser usadas

• Trabalhei na Spectral Compute alguns anos atrás

  • Era uma equipe técnica muito inteligente e competente
  • Na época, não só mirava AMD, como também superava o backend ptx básico do LLVM e o NVCC

• Já escrevi um pouco de CUDA

  • Fico curioso sobre qual é a configuração básica para escrever código para placas AMD

• Há a opinião de que este projeto é excelente

  • Estou ansioso para ver a AMD competir diretamente com a Nvidia

• É bom haver uma página sobre as limitações atuais, mas há a opinião de que a maior parte do que as pessoas descrevem como "CUDA" é uma pequena parte das capacidades reais do CUDA

  • Seria bom ter uma tabela comparativa de recursos avançados como warp shuffle, operações atômicas, DPX, TMA, MMA etc.
  • O ideal seria uma tabela mapeando instruções PTX para instruções equivalentes em RDNA, ou para a lista de instruções que as emulam

• Há a opinião de que isso pode ser real porque é tecnicamente possível

  • Analisar PTX inline e mapeá-lo para AMDGPU seria um grande sofrimento
  • Fazer o alvo AMDGPU funcionar a partir de código-fonte CUDA que não usa PTX inline é parecido com substituir por HIP
  • Alguns detalhes podem ser questionáveis, por exemplo, o modelo atômico pode não coincidir ou Volta pode ter um modelo diferente de ponteiro de instrução
  • Mas pode ser feito corretamente
  • A AMD não fará isso
  • O CUDA em geral não é algo tão bom assim, e a equipe jurídica criaria problemas
  • Mas outras pessoas certamente podem fazer isso