HipKittens: Hızlı ve Güçlü AMD Çekirdekleri

 (hazyresearch.stanford.edu)
3 puan yazan GN⁺ 2025-11-16 | 1 yorum | WhatsApp'ta paylaş
  • HipKittens, AMD GPU'lar için yüksek performanslı çekirdekler ve C++ tabanlı programlama primitifleri sunarak yapay zeka hesaplamalarının verimliliğini artıran bir proje
  • Mevcut AMD ekosistemindeki AITER, PyTorch, Triton, TileLang, Composable Kernel gibi bileşenler, kararsız performans ve olgunlaşmamış destek nedeniyle sınırlarını ortaya koyuyor
  • HipKittens, tile abstraction (karo tabanlı soyutlama) etrafında şekilleniyor; NVIDIA ile AMD arasında ortak bir arayüzü korurken donanıma özgü uygulamaları ayırıyor
  • Yaklaşık 500 satırdan az kodla yazılmış çekirdekler, AMD'nin mevcut elle yazılmış assembly çekirdeklerinden daha yüksek performansa ulaşıyor
  • Yapay zeka hesaplamalarını çoklu silikon ortamlarına genişletmek için pratik bir temel sunuyor ve açık donanım ekosistemine geçiş olasılığını gösteriyor

AMD GPU ekosisteminin mevcut durumu ve sorunları

  • Yapay zeka hesaplamaları bugüne kadar tek bir donanım üreticisi merkezli gelişti, ancak AMD GPU'lar son nesilde en üst düzey hesaplama performansı ve bellek bant genişliği sunuyor
    • Örnek: AMD MI355X OAM, BF16 2.5 PFLOPs, MXFP8 5.0 PFLOPs, 288GB bellek, 8.0TB/s bant genişliği sunuyor
  • Ancak olgun bir yazılım yığınının olmaması nedeniyle bu performans gerçek yapay zeka iş akışlarında kullanılamıyor
  • AMD'nin başlıca yazılım bileşenleri AITER, PyTorch, Triton, Mojo, TileLang, Composable Kernel (CK) gibi araçlardan oluşuyor
    • AITER ve PyTorch içindeki SDPA Llama GQA backward çekirdekleri, SoTA'nın sırasıyla yalnızca %30 ve %24 düzeyinde performansına ulaşabiliyor
    • Mojo, bank conflict nedeniyle performansın %50 seviyesinde kalmasına neden oluyor
    • TileLang, yalnızca CDNA3'e kadar destek veriyor ve çekirdek özellik eksikliği ile CK bağımlılığı yüzünden karmaşıklığı artırıyor
    • Triton, register yaşam süresi takibi ve bellek erişimi optimizasyonunda zorluk yaşıyor
  • Sonuç olarak, AMD'de en yüksek performanslı çekirdeklerin uzmanlar tarafından assembly ile doğrudan yazılması gerekiyor; bu da ölçeklenebilirlik ve bakım açısından büyük kısıtlar yaratıyor

CUDA merkezli ekosistemle karşılaştırma

  • Yapay zeka topluluğu, CUDA ekosisteminin giriş bariyerine (CUDA moat) dikkat çekiyor
  • Geçmişte NVIDIA çekirdek geliştirme de düşük seviyeli CUDA/CUTLASS tabanında yıllar süren çaba gerektiriyordu
  • Son dönemde DSL (alan özel dil) ve yapay zeka destekli araçların gelişmesiyle NVIDIA çekirdek geliştirme basitleşti
    • Örnek: ThunderKittens (TK), CuTe DSL, TinyGrad “tinykittens”, TileLang, Gluon
  • Bu eğilimin üzerine AMD tarafında da yeni programlama primitiflerine duyulan ihtiyaç araştırılıyor

HipKittens'in tasarım ilkeleri

  • HipKittens, ThunderKittens (NVIDIA) ve ThunderMittens (Apple Silicon) sonrasında AMD için geliştirilen sürüm
  • Temel kavram: tile abstraction (karo tabanlı soyutlama)
    1. Tile soyutlamasının genelliği – NVIDIA'da etkili olan karo tabanlı işlem modeli AMD'de de doğal biçimde uygulanabiliyor
    2. Backend uygulamasının mimariye özgü olması – Bellek erişim desenleri ve register zamanlaması donanıma göre farklı tasarlanıyor
    3. Zamanlama stratejisinin uyarlanabilirliği – AMD'nin CDNA3·CDNA4 mimarilerinde wave tabanlı zamanlama verimsiz olsa da karo düzeyindeki zamanlama sadelik ve performansı koruyor
  • Arayüzü (karolar ve işlemler) ile uygulamayı (donanım eşlemesi) ayırarak, farklı GPU mimarilerine ortak biçimde uygulanabilecek bir model sunuyor

Performans sonuçları

  • Attention Forward çekirdeği: Yaklaşık 500 satır kodla yazıldı ve AITER'in assembly çekirdeğinden daha iyi performans elde etti
    • Farklı head dimension (D) ve sequence length (N) değerlerinde, hem Causal hem de Non-Causal ayarlarda üstünlük gösterdi
  • GEMM çekirdeği: 100 satırdan kısa bir çekirdek döngü ile oluşturuldu ve BF16 ile FP8 hesaplamalarının ikisinde de en yüksek performansı elde etti
  • Attention Backward, Rotary, Fused Dropout-Residual-LayerNorm çekirdeklerinde de mevcut en iyi performansın üzerine çıkıldığı doğrulandı
  • Tüm çekirdekler okunabilirlik ve düzenleme kolaylığını korurken aynı zamanda elle yazılmış assembly düzeyinde performans sağladı

Çoklu silikon yapay zekaya doğru genişleme

  • Yapay zekanın potansiyelini gerçekleştirmek için çeşitli ve açık donanımların kullanımı gerekiyor
  • HipKittens, AMD GPU'ları yapay zeka geliştiricileri için pratikte erişilebilir bir platforma dönüştürmeyi hedefliyor
  • AMD'nin açık kaynak yazılım yığınıyla birleştiğinde, çoklu silikon tabanlı yapay zeka ekosistemine geçiş ihtimalini gösteriyor
  • Sonraki yazıda HipKittens'in teknik ayrıntılı yapısı ele alınacak (part two duyurusu)

İlgili okumalar

1 yorum

 
GN⁺ 2025-11-16
Hacker News görüşleri

  • AMD ile anlaşmamız var ve Llama 405B modelini MI350X üzerinde MLPerf için eğitiyoruz.
    AMD'nin durumu kesinlikle iyileşiyor. AMD GPU'nuz varsa pytorch.org üzerinden Linux+ROCm seçeneğine tıklayıp PyTorch kurmanızı öneririm. 3 yıl önce durum umutsuzdu ama şimdi ana akım özelliklerin çoğu düzgün çalışıyor. MI300X'te nanochat çalıştırdım ve doğrudan sorunsuz açıldı. MI350X de aynı şekilde kararlı.
    Elbette NVIDIA'nın gerisinde ve yazılım ekosistemi, derleyici ve sürücüler için daha çok yatırım gerekiyor. Ama 2 yıl önceki umutsuz tabloya kıyasla artık umut görünüyor.
    AMD'nin LLVM arka ucunda eksik kalan noktaları görmek istiyorsanız HipKittens kodunu CUDA Kittens ile karşılaştırabilirsiniz.
    Eğitim tarafında NVIDIA ve Google birinci ligde, AMD ikinci sırada, geri kalanlar ise neredeyse yok. Intel ve Tenstorrent hâlâ çok geride, Huawei'de örnekler segfault ile çöküyordu, Groq çip satmayı bıraktı, Cerebras'a ise zaten erişilemiyor. Trainium için tek bir instance almak 5 gün sürünce ilgimi kaybettim

    • Tinygrad'ın bu tür bellek optimizasyonu ve warp odaklı özellikleri ifade edebilmesi ya da arayabilmesi için daha ne kadar yolu var merak ediyorum. Bu özellikleri tinygrad'a eklemenin ne kadar karmaşık olduğunu da merak ediyorum
    • Tüketici sınıfı donanımda (MI dışı) ROCm + PyTorch çalıştırmak için tescilli çekirdek sürücüsü gerekip gerekmediğini merak ediyorum
    • “Cerebras'a hiçbir yerde ulaşılamıyor” sözü kulağıma daha çok zafer işareti gibi geliyor
    • Son 2 yıldır bir AMD NeoCloud'un CEO'su olarak çalışıyorum. AMD'nin bu geri dönüşüne bizzat tanıklık etmek gerçekten sevindirici.
      İlk kurulum hâlâ biraz pürüzlü ama geçmişe kıyasla çok daha iyi. Örneğin bir ay önce nanochat düzgün çalışmıyordu, şimdi çalışıyor. Asıl önemli olan, artık insanların AMD ekosistemine ilgi duyması.
      Yapay zekanın donanım çeşitliliğine ihtiyacı var. Tüm yapay zeka donanımı ve yazılımının tek bir şirket tarafından tekelleştirilmesi hissedarlar için iyi olabilir ama teknolojik ilerleme için zararlı

  • NVIDIA'nın piyasa değerini anlamıyorum. Şu anda Transformer gibi az sayıdaki algoritma kazandı ve veri daha önemli hâle geldi. Eskisi gibi çok çeşitli HPC kodlarına ihtiyaç olmadığı için artık rakiplerin yalnızca dar bir algoritma kümesini optimize etmesi yeterli.
    vLLM ve Transformer'ı verimli şekilde çalıştırabiliyorsanız çok büyük bir pazar açılıyor. O zaman AMD ya da Huawei'nin de kolayca yetişebilmesi gerekir gibi geliyor; bu durumda NVIDIA'nın gerçek hendeği (moat) nedir merak ediyorum. Sadece InfiniBand yeterli mi?

    • Evet, NVIDIA'nın hendeği giderek zayıflıyor. MS, Google, Amazon ve Apple gibi dev şirketlerin hepsi NVIDIA'ya bağımlılıktan kaçınmak için kendi çiplerini geliştiriyor.
      Veri merkezi GPU'larında NVIDIA hâlâ güçlü ama Google TPU'ları büyük ölçekte kullanıyor ve OpenAI de AMD donanımı sipariş ediyor.
      CUDA ekosistemi hâlâ önemli ama herkes oradan çıkmanın yollarını arıyor. AMD, Intel, Qualcomm ve diğerleri de bu yarışa girmiş durumda. HipKittens, tarafsız bir yazılım ekosistemine giden yolda önemli bir adım gibi görünüyor
    • “Az sayıdaki algoritmayı” uygulamanın en kolay yolu genel amaçlı işlem donanımı üretmektir. Donanımın geliştirme döngüsü uzun olduğu için bu yaklaşım daha güvenlidir
    • InfiniBand yerini UEC'ye bırakıyor. Çıkarım (inference) için InfiniBand gerekmiyor. Bu yüzden çıkarım pazarında bir hendek yok. AMD ya da Google TPU kullanmak daha mantıklı
    • NVIDIA'nın asıl silahı devasa CUDA kütüphane ekosistemi. Neredeyse her alana karşılık gelen kod mevcut
    • Transformer tek bir teknoloji değil. Uygulama biçimleri çok çeşitli. Bu yüzden vLLM ya da TRL o kadar da basit değil

  • AMD'nin bu tür projelere fon sağlaması beklenirdi ama bence yine fırsatı kaçırmış olabilir. GPU ve yapay zeka tarafında hep böyle oldu

    • AMD, ürün çıkarabilmek için gereken asgari yazılıma yatırım yapıyor. Performans test altyapısı bile zayıf ve regresyon hataları olduğu gibi müşteriye ulaşıyor.
      HipKittens bir iyileştirme ama AMD içinde çekirdek performansını takip edebilecek yeterlilik eksik. DevOps Hindistan'daki TCS'ye dış kaynak verilmiş ve onlar da duruma pek hâkim değil.
      İyi liderleri olan ekipler kendi gölge BT ekiplerini kuruyor. ROCm'de böyle bir ekip yoktu; iyileştirmeler ancak büyük bulut müşterileri itiraz etmeye başlayınca geldi.
      Yetenekli insanları işe alsalar bile piyasa altı maaş teklif ediyorlar. Kıyas noktaları Qualcomm ya da Walmart seviyesinde kalıyor.
      Son 4 yılda primler hiç tam ödenmedi
    • “AMD fırsatı kaçırmaz, kaçırmayı da kaçırmaz” sözü tam oturuyor. Instinct donanımı aslında NVIDIA ile rekabet edebilecek düzeyde ama yazılım desteği korkunçtu.
      Geçmişte Radeon VII'de olduğu gibi desteği kesiyor ya da ekosistemi sık sık baştan kurarak olgunlaşmasını engelliyordu. CUDA uyumluluğunun zayıf olması ve yatırım eksikliği nedeniyle çoğu kuruluş gidip NVIDIA'yı seçti.
      Yine de son dönemde ROCm ve Instinct ekosistemine düzenli yatırım yapılıyor, bu yüzden yavaş yavaş düzeliyor. Ama ağ tarafında NVIDIA hâlâ çok ileride
    • Performans karşılaştırma tablolarına bakınca AMD bugün NVIDIA seviyesinde olabilirmiş gibi görünüyor. Ama yazılım eksikliği yüzünden başarısız oldu. Aslında daha zor olan iş çip tasarımıyken, asıl sorun yazılımı anlamamalarıydı
    • Geçmişte ROCm için optimize çekirdekler katkılayan kişiler vardı ama AMD PR'leri kapatıp birleştirmedi. Gerçekten anlaşılmaz bir davranış
    • Şu anda fon sağlanıyor ve proje düzgün çalışıyor

  • ThunderKittens üzerine inşa edilmiş projeler olup olmadığını merak ediyorum.
    Eğer bu HIP'e taşınmış bir sürümse, sıradan bir CUDA portundan çok daha fazla pratik değer taşıyabilir gibi görünüyor

  • “raw assembly vs cooked assembly” ifadesi ilginç.
    Eskiden CPU için assembly kodunu doğrudan yazmak yaygındı. GPU tarafında da bundan fazla korkmamak gerekir. Sonuçta o kodu derleyici üretiyor

  • Matematiksel olarak çıkarım (inference) temel doğrusal cebir/BLAS işlemlerinden ibaret.
    dtype ve sparsity dikkate alınarak kullanım senaryolarının %80'ini kapsayan yalın bir çıkarım API'si mümkün olabilir mi merak ediyorum. CUDA kadar karmaşık olması gerekmiyor gibi

  • composable-kernel, Gentoo sistemimde en sık OOM (bellek yetersizliği) yaşatan yazılım oldu. Clang, CK derlerken iş parçacığı başına 2,5 GB'den fazla kullanıyor

    • Debian için CK paketine baktım; -j32 ile derlenirken 64 GB'lık iş istasyonunda bile OOM oldu. -j1 ile çalıştırınca 190 saat sonra başarıyla tamamlandı.
      Resmî derleme sunucularında mimari sayısını azaltmak gerekebilir. Bağımlı paketlerin çoğunda yalnızca başlıklara ihtiyaç olduğu söyleniyor. Derleme süresini kısaltmak için iyileştirme çalışmaları sürüyor
    • Tek bir çekirdek için şablon örneklemesinin 10 dakikadan fazla sürmesi inanılmaz bir seviye.
      device_grouped_conv2d_fwd_xdl_ngchw_gkcyx_ngkhw_f16_instance, clang'a tam olarak ne yaptırıldığını gerçekten anlamıyorum

  • Bu gelişmelerin tüketici sınıfı AMD çiplerinde de LLM'leri iyi çalıştırmayı mümkün kılıp kılmayacağını merak ediyorum.
    Örneğin AMD CPU/iGPU ile RTX 5080'i birlikte barındıran HP OMEN MAX 16 dizüstünü düşünüyorum; AMD tarafı RTX ile rekabet edebilir mi?

    • dGPU'nun her zaman daha hızlı olduğunu düşünebilirsiniz ama bellek kapasitesi sınırı ciddi engel oluyor.
      Bu blog yazısı ya da üst düzey Mac sonuçları bu açıdan ilginç
    • Ben RTX7900XTX üzerinde Ollama ile Qwen3 Coder 30B çalıştırıyorum. Gayet iyi çalışıyor. Yükün bir kısmı sistem belleğine ve Ryzen 7 CPU'ya kayıyor gibi görünüyor.
      Sonnet 4 API'sinden biraz daha yavaş ama yerelde bu seviyede performans benim için fazlasıyla yeterli.
      Opencode + Ollama + Qwen3 Coder birleşimi, ClaudeCode + Sonnet4 için harika bir alternatif.
      Tabii yapay zekanın kodlamanın tamamını üstlenmesi gerekiyorsa farklı düşünebilirsiniz. Ama kişisel yardımcı olarak mükemmel

  • AMD'nin B300'ü neden tamamen görmezden geldiğini anlamıyorum