AMD GPU’nun ‘brrr’ hızında çalışmasını sağlamanın yolu

• HipKittens, AMD GPU’ların potansiyel performansını ortaya çıkarmak için tasarlanmış bir programlama primitifi koleksiyonu olup bellek erişimi, zamanlama ve önbellek yeniden kullanımını optimize eder
• AMD MI355X GPU, 256 işlem birimi ve 8 chiplet (XCD) yapısına sahiptir; büyük bir register dosyası ve ince taneli matris çekirdeği komutları sunar
• NVIDIA’dan farklı olarak AMD’de register yeniden tahsisi, asenkron matris komutları ve mbarrier bulunmadığından, wave specialization yerine 8-wave ping-pong ve 4-wave interleave zamanlaması daha etkilidir
• HipKittens, chiplet farkındalıklı (grid) zamanlama ile L2 ve LLC önbellek yerelliğini iyileştirerek GEMM ve Attention işlemlerinde daha yüksek bant genişliği ve TFLOPS artışı sağlar
• Bu yaklaşım, AMD GPU ekosistemindeki yazılım olgunluğu eksikliğini telafi eder ve çeşitli donanım temelli yapay zeka hesaplamalarının ölçeklenebilirliğini artıracak bir temel sunar

---

AMD CDNA GPU mimarisi ve performans özellikleri

• AMD MI355X GPU, 256 işlem birimi (CU) içerir ve her CU 4 SIMD’den oluşur
  • Her SIMD, 64 iş parçacığından oluşan bir wave çalıştırır; bu da NVIDIA’nın 32 iş parçacıklı warp yapısıyla karşılaştırılır
• MI355X, B200’e kıyasla %70 seviyesinde SRAM’e (165KB) sahiptir ve asenkron matris çarpımı komutları, register yeniden tahsisi, tensor memory acceleration, mbarrier özelliklerinden yoksundur
• Buna karşılık, 2 kat daha büyük bir register dosyası ve %60 daha fazla işlemci sayısı (256 CU’ya karşı 160 SM) sunar
  • Küçük ve ince taneli matris çekirdeği komutlarını destekler ve doğrudan global→shared memory yükleme (TMA benzeri) yeteneğine sahiptir
• AMD, 8 chiplet’ten (XCD) oluşan bir chiplet mimarisi kullanır; her XCD bağımsız bir L2 önbelleğe sahiptir ve üst seviyede LLC önbelleği bulunur
• Tabloya göre MI355X; BF16 2.5 PFLOPs, MXFP8 5.0 PFLOPs, MXFP6 10.1 PFLOPs hesaplama performansı ile 288GB bellek kapasitesi ve 8TB/s bant genişliği sunar

AMD için kernel tasarımındaki zorluklar

• Bellek erişimi optimizasyonu: HIPCC derleyici kısıtları ve kamuya açık olmayan I/O davranışları nedeniyle veri yerleşimi ve swizzle desenlerinin tasarımı önemlidir
• İşlemci içi zamanlama: AMD’de shared memory yerine register dosyası ve küçük matris komutlarından yararlanmak gerekir
• İşlemciler arası zamanlama: Chiplet tabanlı yapı nedeniyle önbellek seviyesindeki NUMA etkilerini dikkate alan iş dağıtımı gerekir

HipKittens’in bellek erişim desenleri

• HipKittens (HK), temel veri birimi olarak tile kullanır ve PyTorch benzeri işlem fonksiyonları sunar
  • Tile; veri tipi, boyut ve düzen ile tanımlanır ve C++ template metaprogramming sayesinde çeşitli girdilere uyarlanır
• Register zamanlaması: HIPCC belirli register’ları MFMA girdisi olarak kullanamadığı için HK, açık register sabitleme özelliği sunar
  • Geliştirici, en yüksek performanslı kernel’i yazabilmek için register’ları doğrudan belirleyebilir
• Register düzeni: AMD’de veri tipi ve matris biçimine göre düzen değiştiğinden tek bir swizzle deseni kullanılamaz
  • Örneğin 16×16 bf16 tile ile 16×32 bf16 tile farklı swizzle desenleri gerektirir
• Komut fazı yapısı: AMD’nin shared memory komutları süreksiz faz gruplarına sahiptir ve iç dokümantasyonu yetersizdir
  • HK bunun için tersine mühendislikle geliştirilmiş bir çözücü sunar
• Adres üretimi: AMD, asenkron HBM→shared memory yüklemeyi destekler ve HBM adres swizzle’ı ile optimizasyon yapar

İşlemci içi zamanlama: Wave desenleri

• Wave specialization, NVIDIA’da etkili olsa da AMD’de register yeniden tahsisi olmaması nedeniyle performansı düşürür
  • Producer wave gereksiz register’ları işgal ederken, Consumer wave yetersiz register yüzünden spill yaşar
• HK’nin deney sonuçlarına göre wave specialization, AMD’de aritmetik yoğunluğun düşmesine ve bellek darboğazına yol açar
  • Örnek: GEMM’de HK 0/8 yapılandırması 1605 TFLOPs, CUTLASS ise 1570 TFLOPs üretir
• Alternatif zamanlama desenleri
  • 8-wave ping-pong: İki wave sırayla bellek/hesaplama kümelerini yürütür
  • 4-wave interleave: Bir wave, bellek ve hesaplamayı ince taneli biçimde iç içe yürütür
  • 8-wave yaklaşımı daha sade koda sahipken, 4-wave daha ince denetim sunar ancak kodu uzatır
  • GEMM ve Attention Forward’da 8-wave, SoTA düzeyinde performans sağlar

İşlemciler arası zamanlama: Chiplet farkındalıklı yaklaşım

• AMD MI355X, 8 XCD chiplet’ine sahiptir ve her chiplet bağımsız bir L2 önbellek taşır
  • Thread block’lar round-robin yöntemiyle chiplet’lere atandığı için grid sırası önbellek yeniden kullanım verimini doğrudan etkiler
• Basit bir row-major yerleşim, düşük L2 önbellek yeniden kullanımına neden olarak bant genişliği kaybı yaratır
  • Örnek: L2 %55, LLC %95, 15.1 TB/s, 1113 TFLOPs
• HK, chiplet farkındalıklı (grid) zamanlama getirerek L2 ve LLC önbellek yerelliğini birlikte kullanır
  • Thread block’ları çıktı matrisinin bitişik bölgeleri bazında gruplandırarak girdi verisinin yeniden kullanımını en üst düzeye çıkarır

Gerçek kernel örnekleri

• Attention Forward ve BF16 GEMM kernel’lerinin hot loop kısmı, HK’nin 8-wave ping-pong zamanlamasını kullanır
  • Her döngüde Compute–Memory kümeleri sırayla çalıştırılır ve schedule barrier ile senkronize edilir
  • Kod örneğinde mma_AtB, load, exp2, col_sum gibi HK işlemleri tekrar tekrar kullanılır

Sonuç: Çoklu silikonlu yapay zeka çağında AMD

• HipKittens, AMD CDNA3 ve CDNA4’te rekabetçi performans elde eder
  • Üç temel unsur: optimize edilmiş bellek erişimi, AMD odaklı wave zamanlaması, chiplet farkındalıklı grid zamanlaması
• HK kernel’leri AMD tarafında en yüksek performansı elde ederken, NVIDIA Blackwell kernel’leriyle de rekabet edecek seviyeye ulaşır
• Yapay zeka hesaplamalarında çeşitlilik için AMD GPU erişilebilirliğinin artırılması gerekir; HipKittens bunun için temel bir yazılım zemini sunar
• AMD’nin HIPCC register zamanlamasını iyileştirmesi, gelecekte önemli bir gelişim alanı olarak öne çıkıyor

[GN⁺]

Hacker News görüşleri

• HipKittens hakkındaki tartışmaya göz atmanızı öneririm
• Aynı araştırmayı ele alan bir HipKittens: Fast and furious AMD kernels yazısı da var. George Hotz ve AMD çalışanlarının yorumları bulunuyor
• Akademide bu tür sorunların ele alınması sevindirici, ama bunun sonuçta AMD içinde çözülmesi gereken bir sorun olduğunu düşünüyorum
  • Bana göre donanım şirketleri sadece donanım yapmalı. Böylece teşvikler saf kalır. Performans %20 daha düşük olsa bile bunun daha iyi olduğunu düşünüyorum
  • Kesinlikle katılıyorum. AMD bu meseleyi 10 yıl önce erteledi ve ancak şimdi yetişmeye çalışıyor. Donanım harika ama firmware yazma becerisi eksikliği yüzünden potansiyelini kullanamıyor
  • Ama bu araştırma ekibi daha önce Nvidia GPU'lar için de benzer bir yazılım geliştirmişti. Görünüşe göre çok iyi araştırmacılar kendi uzmanlıklarını ortaya koyuyor
  • Bildiğim kadarıyla AMD bu konuyu zaten çeşitli seviyelerde ele alıyor ve tinycorp ile de iş birliği yapıyor
• Yazı, AMD GPU'ların mimari karmaşıklığı nedeniyle optimizasyonun zor olduğu izlenimini veriyor. Ancak uzun vadede AMD'nin yaklaşımı daha iyi ölçeklenebilir olabilir. Nvidia 2 chiplet kullanırken AMD 8 chiplet'li bir yapıya sahip, bu yüzden bellek yerelliği sorunları var. Gelecekte chiplet sayısı daha da artacağı için, bugünkü karmaşıklıkla başa çıkma deneyimi uzun vadede faydalı olabilir
  • AMD'de yüksek performans için warp specialization gerekmediğinden programlama daha basit
• Birçok geliştirici AMD GPU'ları genel geliştiriciler için “go brrr” hale getirmeye çalıştı ama başarısız oldu. AMD'nin neden kendi başına yazılım sorunlarını çözmediğini anlamıyorum. Artık yeterince paraları da var; geliştirici işe almamak mazeret olamaz. Veri merkezi GPU'ları fena değil ama bireysel olarak ML ve yapay zeka denemeleri yapmak istediğinizde Nvidia hâlâ çok daha iyi. Beş yıllık RTX 3090'ımın şimdiye kadar çıkan tüm AMD tüketici GPU'larından daha iyi olduğunu düşünüyorum
  • AMD geliştirici deneyimi berbat. Sürücü çökmesi hata raporlarını bile kabul etmiyorlar
  • Kısa süre önce çıkarım sunucumu NVidia 5090'dan iki adet AMD R9700 32GB'a geçirdim ve deneyim tamamen olumluydu. Fedora çekirdeğinde DKMS ayarı yapmadan doğrudan çalıştı, ROCm ile konteyner bağlamak da kolaydı. Sadece Ollama ve Storyteller ayarlarını değiştirmem yetti. CUDA'dan çok daha rahat bir deneyimdi
  • Nvidia, Unreal Engine fork'unu bile doğrudan kendisi sürdürüyor. AMD'nin seviyesi bununla rekabet bile edemiyor
  • Nvidia, donanım şirketleri arasında yazılım mühendislerine rekabetçi ücret sunan tek şirket. AMD'de ise yazılımı hâlâ “gerçek iş” olarak görmeyen bir kültür var ve bu tür bir ataleti değiştirmek zor
• Mojo'nun AMD GPU'larda geliştirici deneyimini (devX) iyileştirmeye yönelik fikirleri vardı; bunun ne durumda olduğunu merak ediyorum
• AMD'nin yazılımı iyileştirmek için milyarlarca dolar yatırım yapmamasını anlamıyorum. Nvidia dünyanın en değerli şirketi ve AMD onun tek rakibi
  • AMD de çabalıyor ama her yıl donanım yenilemeye alışmış bir kurum kültürünü yazılım merkezli bir kültüre dönüştürmek zor. Yazılım, donanım gibi doğrudan gelir üretmediği için yönetim bunu daha düşük önceliğe koyma eğiliminde oluyor. Ayrıca dış tedarikçilerin açık kaynak olarak kod sağlaması kısa vadede iyi görünse de uzun vadeli kaliteye zarar verebiliyor. Donanım trendlerinden birini bile kaçırırsanız rakibin gerisine düşme riski büyük
  • Birden fazla GPU üreticisinde çalıştım; yazılımı varlık (asset) olarak görüp yatırım yapan tek şirket Nvidia. Diğerleri bunu sadece maliyet olarak görüyor
• Kişisel olarak “go brr” memesini pek sevmiyorum ama Stanford gibi yerlerde kullanıldığını görmek eğlenceli
  • Aslında daha bir yıl önce ThunderKittens duyurusunda da “go brr” kullanılmıştı
  • Böyle bir meme üniversitenin resmî kanallarında görünüyorsa, bu belki de trendin bittiğinin işaretidir
• Projenin kendisi harika ama AMD'nin neden bunu doğrudan kendisinin yapmadığı soru işareti. AMD hâlâ olgun bir yazılım yığınının önemini kavrayamamış gibi görünüyor. CUDA gibi, tüm kartlarda kullanılabilen birleşik bir yığına ihtiyaç var. Bir zamanlar AMD'nin bir gün yetişeceğine inanıyordum ama artık neredeyse umudu kestim
• Proje iyi ama yazının kendisi garip biçimde yazılmış gibi hissettiriyor
  • Metin çok tuhaf. Ya yapay zekaya aşırı yaslanılmış ya da yapay zeka üslubu taklit edilmiş gibi. “part one'a bakın” ya da “AMD GPU'ları go brr yapma yolu” gibi ifadeler tekrar edip duruyor. Özellikle teknik kısımlarda, grafikle anlatılması gereken şeylerin 100 satırlık kodla açıklanmış olması hayal kırıklığı yaratıyor