HipKittens: kernel AMD rápidos e poderosos

• HipKittens é um projeto que fornece kernels de alto desempenho e primitivas de programação baseadas em C++ para GPUs AMD, aumentando a eficiência de operações de IA
• No ecossistema AMD atual, AITER, PyTorch, Triton, TileLang e Composable Kernel expõem limitações por causa de desempenho instável e suporte ainda imaturo
• O HipKittens é centrado em uma abstração baseada em tiles (tile abstraction), separando implementações específicas de hardware enquanto mantém uma interface comum entre NVIDIA e AMD
• Kernels escritos com cerca de 500 linhas de código ou menos alcançam desempenho superior ao dos kernels AMD escritos manualmente em assembly
• O projeto apresenta uma base prática para expandir operações de IA para ambientes com múltiplos silícios e aponta para uma possível transição a um ecossistema de hardware aberto

---

Situação atual e problemas do ecossistema de GPU da AMD

• As operações de IA evoluíram até agora de forma centrada em um único fornecedor de hardware, mas as GPUs AMD oferecem desempenho computacional de nível máximo e largura de banda de memória nas gerações mais recentes
  • Ex.: AMD MI355X OAM com BF16 2.5 PFLOPs, MXFP8 5.0 PFLOPs, memória de 288GB e largura de banda de 8.0TB/s
• Porém, por causa da ausência de uma stack de software madura, esse desempenho não vem sendo aproveitado em fluxos de trabalho reais de IA
• Os principais componentes de software da AMD incluem AITER, PyTorch, Triton, Mojo, TileLang e Composable Kernel (CK)
  • Os kernels de retropropagação SDPA Llama GQA de AITER e PyTorch atingem apenas 30% e 24% do desempenho SoTA, respectivamente
  • O Mojo fica em cerca de 50% de desempenho devido a conflitos de banco (bank conflict)
  • O TileLang suporta apenas até CDNA3 e aumenta a complexidade por causa da falta de recursos centrais e da dependência de CK
  • O Triton enfrenta dificuldades no rastreamento do tempo de vida de registradores e na otimização de acessos à memória
• Como resultado, os kernels AMD de maior desempenho ainda precisam ser escritos manualmente em assembly por especialistas, o que impõe grandes limitações de escalabilidade e manutenção

Comparação com o ecossistema centrado em CUDA

• Na comunidade de IA, considera-se que existe uma barreira de entrada do ecossistema CUDA (CUDA moat)
• No passado, o desenvolvimento de kernels para a NVIDIA também exigia anos de esforço com base em CUDA/CUTLASS de baixo nível
• Mais recentemente, o desenvolvimento de kernels para NVIDIA foi simplificado com o avanço de DSLs (linguagens específicas de domínio) e ferramentas assistidas por IA
  • Ex.: ThunderKittens (TK), CuTe DSL, TinyGrad “tinykittens”, TileLang e Gluon
• Com base nessa tendência, explora-se na AMD a necessidade de novas primitivas de programação

Princípios de design do HipKittens

• O HipKittens é a versão desenvolvida para AMD, na sequência de ThunderKittens (NVIDIA) e ThunderMittens (Apple Silicon)
• Conceito central: abstração baseada em tiles (tile abstraction)
  1. Universalidade da abstração de tiles – o modelo de operações baseado em tiles, que foi eficaz na NVIDIA, também se aplica naturalmente à AMD
  2. Implementação de backend específica por arquitetura – padrões de acesso à memória e escalonamento de registradores são projetados de forma diferente para cada hardware
  3. Adaptabilidade da estratégia de escalonamento – em CDNA3 e CDNA4 da AMD, o escalonamento baseado em wave é ineficiente, mas o escalonamento por tile ainda preserva simplicidade e desempenho
• Ao separar interface (tiles e operações) e implementação (mapeamento para hardware), o projeto apresenta um modelo aplicável em comum a várias arquiteturas de GPU

Resultados de desempenho

• Kernel Attention Forward: escrito com cerca de 500 linhas de código e com desempenho superior ao kernel em assembly da AITER
  • Supera a alternativa em várias dimensões de head (D) e comprimentos de sequência (N), tanto em configuração Causal quanto Non-Causal
• Kernel GEMM: composto por um loop principal com menos de 100 linhas, alcançando desempenho máximo tanto em BF16 quanto em FP8
• Também foram confirmadas melhorias em relação ao melhor desempenho anterior nos kernels Attention Backward, Rotary e Fused Dropout-Residual-LayerNorm
• Todos os kernels mantêm legibilidade e facilidade de modificação ao mesmo tempo em que garantem desempenho em nível de assembly escrito manualmente

Expansão para IA em múltiplos silícios

• Para concretizar o potencial da IA, é necessário usar hardware diverso e aberto
• O HipKittens tem como objetivo tornar as GPUs AMD uma plataforma realmente acessível para desenvolvedores de IA
• Em combinação com a stack de software open source da AMD, o projeto aponta para uma possível transição a um ecossistema de IA baseado em múltiplos silícios
• Um texto posterior abordará a estrutura técnica detalhada do HipKittens (prévia da parte dois)

[GN⁺]

Comentários no Hacker News

• Fechamos um contrato com a AMD e estamos treinando o Llama 405B no MI350X para o MLPerf.
  A situação da AMD claramente está melhorando. Se você tem uma GPU AMD, recomendo instalar o PyTorch clicando em Linux+ROCm em pytorch.org. Há 3 anos era desesperador, mas agora a maior parte dos recursos principais funciona bem. Rodei o nanochat no MI300X e simplesmente funcionou de cara. O MI350X também está estável.
  Claro, ainda está atrás da NVIDIA, e precisa de muito investimento no ecossistema de software, compiladores e drivers. Mas, comparado à situação desesperadora de 2 anos atrás, agora dá para ter esperança.
  Se quiser ver onde o backend LLVM da AMD ainda deixa a desejar, vale comparar o código do HipKittens com o do CUDA Kittens.
  Para treinamento, NVIDIA e Google estão em 1º lugar, AMD em 2º, e fora isso quase não há mais ninguém. Intel e Tenstorrent ainda estão longe, os exemplos da Huawei morriam com segfault, a Groq desistiu de vender chips, a Cerebras nem dá para conseguir, e o Trainium demorou 5 dias para me liberar uma instância, então perdi o interesse

  • Fico curioso para saber quão longe o Tinygrad está de conseguir expressar ou explorar recursos como essas otimizações de memória e especialização por warp. Também queria entender o quão complexo seria adicionar isso ao tinygrad
  • Fico curioso se, para rodar ROCm + PyTorch em hardware de consumo (não MI), ainda é preciso um driver de kernel proprietário
  • A frase “não dá para conseguir Cerebras em lugar nenhum” quase soa como um sinal de vitória
  • Sou CEO de uma NeoCloud AMD há 2 anos. É muito bom ver essa virada da AMD de perto.
    A configuração inicial ainda é um pouco áspera, mas melhorou muito em relação ao passado. Por exemplo, um mês atrás o nanochat ainda não funcionava direito, e agora funciona. O importante é que agora as pessoas estão interessadas no ecossistema da AMD.
    A IA precisa de diversidade de hardware. Uma única empresa monopolizar todo o hardware e software de IA pode até ser bom para os acionistas, mas faz mal para o avanço tecnológico

• Não entendo o valor de mercado da NVIDIA. No momento, alguns poucos algoritmos, como Transformer, venceram, e os dados ficaram mais importantes. Já não é mais preciso ter a variedade de códigos HPC de antes, então os concorrentes agora só precisam otimizar um conjunto estreito de algoritmos.
  Se conseguir rodar vLLM e Transformer de forma eficiente, abre-se um mercado gigantesco. Nesse caso, parece que AMD ou Huawei também poderiam alcançar fácil. Então qual é o verdadeiro moat da NVIDIA? Só o InfiniBand basta?

  • Sim, o moat da NVIDIA está enfraquecendo aos poucos. Gigantes como MS, Google, Amazon e Apple estão todos desenvolvendo seus próprios chips para evitar dependência da NVIDIA.
    Em GPUs para datacenter, a NVIDIA ainda é forte, mas o Google já usa TPU em larga escala, e a OpenAI também está encomendando hardware da AMD.
    O ecossistema CUDA ainda é importante, mas há um movimento forte de todo mundo tentando sair dele. AMD, Intel, Qualcomm e outras também entraram nessa disputa. O HipKittens parece um passo importante rumo a um ecossistema de software neutro
  • A forma mais fácil de implementar “alguns poucos algoritmos” é criar hardware de computação de propósito geral. Como o ciclo de desenvolvimento de hardware é longo, essa abordagem é estável
  • O InfiniBand está sendo substituído pelo UEC. Para inferência, não é preciso InfiniBand. Então não há moat no mercado de inferência. O sensato é usar AMD ou Google TPU
  • A verdadeira arma da NVIDIA é seu vasto ecossistema de bibliotecas CUDA. Há código para praticamente qualquer área
  • Transformer não é uma única tecnologia. Há muitas variações de implementação. Por isso vLLM e TRL não são tão simples assim

• Parece o tipo de projeto que a AMD deveria financiar, mas acho que ela de novo deve ter perdido a oportunidade. Com GPU e IA, sempre foi assim

  • A AMD investe só o mínimo de software necessário para lançar um produto. Até a estrutura de testes de desempenho é fraca, e bugs de regressão acabam chegando ao cliente.
    O HipKittens é uma melhora, mas dentro da AMD falta capacidade de acompanhar desempenho de kernel. O DevOps foi terceirizado para a TCS na Índia, e eles não entendem bem a situação.
    Equipes com boa liderança acabam tocando seu próprio time paralelo de TI. O ROCm não tinha isso, e as melhorias só começaram quando grandes clientes de nuvem reclamaram.
    Mesmo quando contratam talentos, oferecem salários abaixo do mercado. Fazem comparação com Qualcomm ou Walmart.
    Nos últimos 4 anos, o bônus nunca foi pago integralmente
  • A frase “a AMD não perde oportunidades sem antes desperdiçá-las” cai como uma luva. O hardware Instinct já está praticamente em nível competitivo com a NVidia, mas o suporte de software era terrível.
    Como no passado com a Radeon VII, a AMD cortava suporte ou reformulava o ecossistema com frequência, impedindo maturidade. Por falta de compatibilidade com CUDA e de investimento, a maioria das organizações simplesmente escolhia NVIDIA.
    Ainda assim, recentemente a empresa vem investindo de forma consistente no ecossistema ROCm e Instinct, então as coisas estão melhorando aos poucos. Mas na área de redes, a NVIDIA ainda está bem à frente
  • Pelas tabelas comparativas de desempenho, a AMD já poderia estar no nível da NVIDIA a esta altura. Mas fracassou por falta de software. Projetar chips é mais difícil, e mesmo assim o problema foi não entender software
  • No passado houve gente contribuindo kernels otimizados para ROCm, mas a AMD fechava os PRs e não fazia merge. É um comportamento realmente incompreensível
  • Agora o projeto está recebendo financiamento e funcionando normalmente

• Fico curioso se existe algum projeto construído em cima do ThunderKittens.
  Se isto for uma versão portada para HIP, parece ter um valor muito mais prático do que um simples port de CUDA

• A expressão “raw assembly vs cooked assembly” é interessante.
  Antigamente era comum escrever código assembly manualmente para CPU. Em GPU também não precisa haver tanto medo. No fim, é isso que o compilador acaba gerando

• Do ponto de vista matemático, inferência é basicamente álgebra linear/BLAS.
  Fico pensando se não daria para ter uma API de inferência enxuta cobrindo 80% dos casos de uso com dtype e sparsity em mente. Não parece que precisaria ser tão complexa quanto CUDA

• O composable-kernel foi o software que mais frequentemente causou OOM (falta de memória) no meu sistema Gentoo. O Clang usa mais de 2,5 GB por thread ao compilar o CK

  • Revisei o pacote CK para Debian, e ao compilar com -j32 ele deu OOM até numa workstation com 64 GB. Com -j1, terminou com sucesso em 190 horas.
    Nos servidores oficiais de build, provavelmente vai ser preciso reduzir a quantidade de arquiteturas. Ouvi dizer que a maioria dos pacotes dependentes só precisa dos headers. Há trabalho em andamento para melhorar isso e reduzir o tempo de compilação
  • Levar mais de 10 minutos para instanciar templates de um único kernel é um nível impressionante.
    device_grouped_conv2d_fwd_xdl_ngchw_gkcyx_ngkhw_f16_instance, não faço ideia do que estão fazendo com o clang

• Fico curioso se esse avanço vai permitir rodar LLMs bem também em chips AMD de consumo.
  Por exemplo, estou considerando um notebook HP OMEN MAX 16 com CPU/iGPU AMD e RTX 5080, mas será que o lado AMD consegue competir com a RTX?

  • Dá para achar que uma dGPU sempre será mais rápida, mas a limitação de capacidade de memória pesa contra isso.
    Este post de blog e os resultados avançados do Mac são interessantes
  • Estou rodando Qwen3 Coder 30B no Ollama com uma RTX7900XTX. Funciona muito bem. Parte da carga parece ir para a memória do sistema e para um Ryzen 7 CPU.
    É um pouco mais lento que a API do Sonnet 4, mas, para esse nível de desempenho local, estou bem satisfeito.
    A combinação Opencode + Ollama + Qwen3 Coder é uma excelente alternativa a ClaudeCode + Sonnet4.
    Claro, se você precisa que a IA faça tudo sozinha na programação, talvez sinta diferente. Mas como assistente pessoal, é perfeito

• Não entendo por que a AMD ignorou completamente o B300