HipKittens: kernels rápidos y potentes para AMD

• HipKittens es un proyecto que ofrece kernels de alto rendimiento y primitivas de programación basadas en C++ para GPUs AMD, con el fin de mejorar la eficiencia de las cargas de trabajo de IA
• En el ecosistema actual de AMD, AITER, PyTorch, Triton, TileLang y Composable Kernel han mostrado limitaciones debido a rendimiento inestable y soporte inmaduro
• HipKittens se centra en la abstracción basada en tiles (tile abstraction), separando las implementaciones específicas de hardware mientras mantiene una interfaz común entre NVIDIA y AMD
• Kernels escritos con menos de unas 500 líneas de código logran un rendimiento superior al de los kernels manuales en ensamblador existentes para AMD
• Presenta una base práctica para expandir las cargas de trabajo de IA a entornos multi-silicio, y sugiere la posibilidad de avanzar hacia un ecosistema de hardware abierto

---

Estado actual y problemas del ecosistema de GPU de AMD

• Hasta ahora, las cargas de trabajo de IA han evolucionado principalmente alrededor de un solo proveedor de hardware, pero las GPUs AMD ofrecen rendimiento de cómputo de primer nivel y ancho de banda de memoria en las generaciones más recientes
  • Ejemplo: AMD MI355X OAM ofrece BF16 2.5 PFLOPs, MXFP8 5.0 PFLOPs, 288 GB de memoria y 8.0 TB/s de ancho de banda
• Sin embargo, debido a la falta de un stack de software maduro, ese rendimiento no se aprovecha en los flujos de trabajo reales de IA
• Los principales componentes de software de AMD incluyen AITER, PyTorch, Triton, Mojo, TileLang y Composable Kernel (CK)
  • Los kernels de retropropagación SDPA Llama GQA de AITER y PyTorch solo alcanzan 30% y 24% del rendimiento SoTA, respectivamente
  • Mojo se queda en alrededor del 50% del rendimiento debido a conflictos de banco (bank conflict)
  • TileLang solo da soporte hasta CDNA3 y añade complejidad por la falta de funciones clave y su dependencia de CK
  • Triton tiene dificultades con el seguimiento de la vida útil de registros y la optimización del acceso a memoria
• Como resultado, los kernels AMD de mayor rendimiento todavía deben ser escritos manualmente en ensamblador por expertos, lo que impone fuertes límites a la escalabilidad y el mantenimiento

Comparación con el ecosistema centrado en CUDA

• En la comunidad de IA, se considera que existe una barrera de entrada del ecosistema CUDA (CUDA moat)
• En el pasado, el desarrollo de kernels para NVIDIA también requería años de trabajo basados en CUDA/CUTLASS de bajo nivel
• Recientemente, los avances en DSL (lenguajes específicos de dominio) y en herramientas asistidas por IA han simplificado el desarrollo de kernels para NVIDIA
  • Ejemplos: ThunderKittens (TK), CuTe DSL, TinyGrad “tinykittens”, TileLang y Gluon
• A partir de esta tendencia, se explora también en AMD la necesidad de nuevas primitivas de programación

Principios de diseño de HipKittens

• HipKittens es la versión desarrollada para AMD, después de ThunderKittens (NVIDIA) y ThunderMittens (Apple Silicon)
• Concepto central: abstracción basada en tiles (tile abstraction)
  1. Generalidad de la abstracción por tiles – el modelo de cómputo basado en tiles, que fue efectivo en NVIDIA, también se aplica de forma natural en AMD
  2. Implementación backend especializada por arquitectura – los patrones de acceso a memoria y la planificación de registros se diseñan de forma distinta según el hardware
  3. Adaptabilidad de la estrategia de scheduling – en CDNA3 y CDNA4 de AMD, el scheduling basado en waves es ineficiente, pero el scheduling por tiles sigue manteniendo simplicidad y rendimiento
• Al separar la interfaz (tiles y operaciones) de la implementación (mapeo al hardware), se propone un modelo aplicable en común a diversas arquitecturas de GPU

Resultados de rendimiento

• Kernel Attention Forward: fue escrito con unas 500 líneas de código y logra mejor rendimiento que el kernel en ensamblador de AITER
  • Supera a la alternativa existente en configuraciones Causal/Non-Causal y con diversas dimensiones de cabeza (D) y longitudes de secuencia (N)
• Kernel GEMM: está compuesto por menos de 100 líneas en el loop principal y alcanza rendimiento líder tanto en BF16 como en FP8
• También se confirmaron mejoras frente al mejor rendimiento previo en los kernels Attention Backward, Rotary y Fused Dropout-Residual-LayerNorm
• Todos los kernels mantienen legibilidad y facilidad de modificación, al tiempo que alcanzan rendimiento a nivel de ensamblador escrito a mano

Expansión hacia una IA multi-silicio

• Para materializar el potencial de la IA, se necesita aprovechar hardware diverso y abierto
• HipKittens busca convertir a las GPUs AMD en una plataforma realmente accesible para desarrolladores de IA
• En combinación con el stack de software open source de AMD, plantea la posibilidad de una transición hacia un ecosistema de IA basado en múltiples silicios
• En una publicación posterior se abordará la estructura técnica detallada de HipKittens (avance de la parte dos)

[GN⁺]

Comentarios de Hacker News

• Tenemos un contrato con AMD y estamos entrenando Llama 405B en MI350X para MLPerf.
  La situación de AMD definitivamente está mejorando. Si tienes una GPU de AMD, recomiendo instalar PyTorch desde pytorch.org haciendo clic en Linux+ROCm. Hace 3 años era algo desolador, pero ahora la mayoría de las funciones principales funcionan bien. Ejecuté nanochat en MI300X y simplemente funcionó de inmediato. MI350X también es estable.
  Claro, todavía está por detrás de NVIDIA, y hace falta mucha inversión en el ecosistema de software, compiladores y drivers. Pero comparado con la situación desesperante de hace 2 años, ahora ya se ve esperanza.
  Si quieres ver qué le falta al backend LLVM de AMD, conviene comparar el código de HipKittens con CUDA Kittens.
  Para entrenamiento, NVIDIA y Google están en el puesto 1, AMD en el 2, y fuera de eso casi no hay nada. Intel y Tenstorrent todavía están lejos, los ejemplos de Huawei morían con segfault, Groq dejó de vender chips, Cerebras ni siquiera se puede conseguir. Trainium tardó 5 días en darme una instancia y perdí el interés.

  • Me pregunto qué tan lejos está Tinygrad de poder expresar o explorar funciones como estas de optimización de memoria o especialización por warp. También quisiera saber qué tan complejo sería agregar esto a tinygrad.
  • Me pregunto si para ejecutar ROCm + PyTorch en hardware de consumo (no-MI) también hace falta un driver de kernel propietario.
  • La frase “Cerebras no se puede conseguir en ningún lado” suena más bien como una señal de victoria.
  • Llevo 2 años trabajando como CEO de una NeoCloud de AMD. Me da mucho gusto ver este turnaround de AMD en primera fila.
    La configuración inicial todavía es un poco rústica, pero está muchísimo mejor que antes. Por ejemplo, hace apenas un mes nanochat no funcionaba bien, y ahora sí. Lo importante es que ahora la gente sí está prestando atención al ecosistema de AMD.
    La IA necesita diversidad de hardware. Que una sola empresa monopolice todo el hardware y software de IA quizá sea bueno para los accionistas, pero perjudica el avance tecnológico.

• No entiendo la valuación de NVIDIA. En este momento ganaron unos pocos algoritmos como Transformer, y los datos se volvieron más importantes. Ya no se necesita una gran variedad de código HPC como antes, así que ahora los competidores solo tienen que optimizar un conjunto reducido de algoritmos.
  Si puedes ejecutar vLLM y Transformer de forma eficiente, se abre un mercado enorme. Entonces parecería que AMD o Huawei también podrían alcanzarlo con relativa facilidad, así que me pregunto cuál es el verdadero foso defensivo (moat) de NVIDIA. ¿InfiniBand por sí solo basta?

  • Sí, el foso de NVIDIA se está debilitando cada vez más. Grandes empresas como MS, Google, Amazon y Apple están desarrollando sus propios chips para evitar depender de NVidia.
    En GPUs para centros de datos NVIDIA sigue siendo fuerte, pero Google ya usa TPU a gran escala y OpenAI también está pidiendo hardware de AMD.
    El ecosistema de CUDA sigue siendo importante, pero hay mucho movimiento para salir de ahí. AMD, Intel, Qualcomm y otros también se metieron a esta competencia. HipKittens parece un paso importante hacia un ecosistema de software neutral.
  • La forma más fácil de implementar “unos pocos algoritmos” es crear hardware de cómputo de propósito general. Como el ciclo de desarrollo de hardware es largo, ese enfoque es estable.
  • InfiniBand está siendo reemplazado por UEC. Para inferencia (inference) no hace falta InfiniBand. Así que en el mercado de inferencia no hay foso defensivo. Lo sensato es usar AMD o Google TPU.
  • El arma real de NVIDIA es su enorme ecosistema de bibliotecas CUDA. Hay código para prácticamente cualquier área.
  • Transformer no es una sola tecnología. Hay muchísimas formas de implementarlo. Por eso vLLM o TRL no son tan simples.

• AMD bien podría estar financiando este tipo de proyectos, pero creo que otra vez debe haber dejado pasar la oportunidad. Siempre ha sido así con GPU e IA.

  • AMD invierte solo en el software mínimo necesario para sacar productos. Incluso su sistema de pruebas de rendimiento es deficiente, y los bugs de regresión les llegan tal cual a los clientes.
    HipKittens es una mejora, pero dentro de AMD falta capacidad para seguir el rendimiento de los kernels. DevOps se tercerizó con TCS en India, y allá no entienden bien la situación.
    Los equipos con buenos líderes operan sus propios equipos de TI en la sombra. ROCm no tenía algo así, y al final las mejoras empezaron cuando los grandes clientes de nube se quejaron.
    Incluso cuando contratan talento, ofrecen salarios por debajo del mercado. Porque se comparan con niveles de Qualcomm o Walmart.
    En los últimos 4 años nunca se pagó el bono completo.
  • La frase “AMD no pierde oportunidades; pierde la oportunidad de aprovecharlas” le queda perfecta. El hardware Instinct en realidad está a un nivel capaz de competir con NVidia, pero el soporte de software fue terrible.
    Como pasó antes con Radeon VII, cortaban el soporte o reiniciaban el ecosistema una y otra vez, así que nunca maduraba. Por la falta de compatibilidad con CUDA y de inversión, la mayoría de las organizaciones simplemente elegían NVIDIA.
    Aun así, últimamente sí han estado invirtiendo de forma constante en ROCm y el ecosistema de Instinct, así que poco a poco está mejorando. Eso sí, en redes NVIDIA todavía está muy por delante.
  • Viendo las tablas comparativas de rendimiento, AMD podría haber estado al nivel de NVIDIA a estas alturas. Pero fracasó por la falta de software. Diseñar chips es más difícil, y aun así el problema fue no entender el software.
  • Antes hubo gente que contribuyó kernels optimizados para ROCm, pero AMD cerró los PR y no los fusionó. De verdad es una conducta incomprensible.
  • Ahora sí hay financiamiento y el proyecto está funcionando con normalidad.

• Me pregunto si hay proyectos construidos sobre ThunderKittens.
  Si esta es una versión porteada a HIP, parecería tener un valor mucho más práctico que un simple port de CUDA.

• La expresión “raw assembly vs cooked assembly” me parece interesante.
  Antes era común escribir código ensamblador directamente para CPU. En GPU tampoco habría que tenerle tanto miedo. Al final, es el compilador el que termina generando ese código.

• Viéndolo matemáticamente, la inferencia (inference) son operaciones básicas de álgebra lineal/BLAS.
  Me pregunto si sería posible una API de inferencia compacta que cubra el 80% de los casos de uso teniendo en cuenta dtype y sparsity. No parece que necesite ser tan compleja como CUDA.

• composable-kernel es el software que más veces me ha provocado OOM (falta de memoria) en mi sistema Gentoo. Clang usa más de 2.5 GB por hilo al compilar CK.

  • Revisé el paquete de CK para Debian y al compilar con -j32 daba OOM incluso en una workstation con 64 GB. Con -j1 terminó con éxito después de 190 horas.
    En los servidores oficiales de build probablemente habría que reducir la cantidad de arquitecturas. Tengo entendido que la mayoría de los paquetes dependientes solo necesita los headers. Se está trabajando en mejoras para reducir el tiempo de compilación.
  • Que la instanciación de templates para un solo kernel tarde más de 10 minutos es impactante.
    device_grouped_conv2d_fwd_xdl_ngchw_gkcyx_ngkhw_f16_instance, no sé qué demonios le están haciendo a clang.

• Me pregunto si este avance permitirá ejecutar bien LLMs también en chips AMD de consumo.
  Por ejemplo, estoy considerando una laptop HP OMEN MAX 16 con CPU/iGPU de AMD y una RTX 5080; ¿la parte de AMD podría competir con la RTX?

  • Uno podría pensar que una dGPU siempre será más rápida, pero el límite de capacidad de memoria termina frenándola.
    Esta entrada de blog o los resultados avanzados en Mac son interesantes.
  • Estoy ejecutando Qwen3 Coder 30B con Ollama en una RTX7900XTX. Funciona bastante bien. Parte de la carga parece irse a la memoria del sistema y a un CPU Ryzen 7.
    Es un poco más lento que la API de Sonnet 4, pero para este nivel de rendimiento en local estoy más que satisfecho.
    La combinación Opencode + Ollama + Qwen3 Coder es una excelente alternativa a ClaudeCode + Sonnet4.
    Claro, si necesitas que la IA haga absolutamente toda la programación, quizá lo verías distinto. Pero como asistente personal es perfecto.

• No entiendo por qué AMD ignoró por completo a B300.