HipKittens: Kernel AMD yang Cepat dan Kuat

 (hazyresearch.stanford.edu)
3 poin oleh GN⁺ 2025-11-16 | 1 komentar | Bagikan ke WhatsApp
  • HipKittens adalah proyek yang menyediakan kernel berperforma tinggi dan primitif pemrograman berbasis C++ untuk GPU AMD guna meningkatkan efisiensi komputasi AI
  • Komponen yang ada di ekosistem AMD seperti AITER, PyTorch, Triton, TileLang, Composable Kernel memperlihatkan keterbatasan karena performa yang tidak stabil dan dukungan yang belum matang
  • HipKittens berpusat pada abstraksi berbasis tile (tile abstraction), dengan antarmuka bersama antara NVIDIA dan AMD tetap dipertahankan sambil memisahkan implementasi per perangkat keras
  • Kernel yang ditulis dengan kurang lebih 500 baris kode atau kurang mencapai performa yang lebih cepat daripada kernel assembly AMD buatan tangan yang sudah ada
  • Proyek ini menawarkan landasan praktis untuk memperluas komputasi AI ke lingkungan multi-silikon serta menunjukkan kemungkinan transisi menuju ekosistem perangkat keras terbuka

Kondisi dan masalah dalam ekosistem GPU AMD

  • Komputasi AI hingga kini berkembang dengan fokus pada satu vendor perangkat keras, tetapi GPU AMD menawarkan performa komputasi tingkat tertinggi dan bandwidth memori sangat besar pada generasi terbaru
    • Contoh: AMD MI355X OAM memiliki BF16 2.5 PFLOPs, MXFP8 5.0 PFLOPs, memori 288GB, bandwidth 8.0TB/s
  • Namun, karena belum adanya software stack yang matang, performa ini tidak dapat dimanfaatkan sepenuhnya dalam workflow AI nyata
  • Komponen perangkat lunak utama AMD mencakup AITER, PyTorch, Triton, Mojo, TileLang, Composable Kernel (CK) dan lainnya
    • Kernel backpropagation SDPA Llama GQA milik AITER dan PyTorch masing-masing hanya mencapai 30% dan 24% dari performa SoTA
    • Mojo tertahan di sekitar 50% performa akibat bank conflict
    • TileLang hanya mendukung hingga CDNA3, dan kurangnya fitur inti serta ketergantungan pada CK menambah kompleksitas
    • Triton mengalami kesulitan dalam pelacakan masa hidup register dan optimasi akses memori
  • Akibatnya, kernel AMD dengan performa tertinggi masih harus ditulis langsung dalam assembly oleh para ahli, yang menjadi kendala besar untuk skalabilitas dan pemeliharaan
Iklan

Perbandingan dengan ekosistem yang berpusat pada CUDA

  • Di komunitas AI, sering dinilai ada hambatan masuk ekosistem CUDA (CUDA moat)
  • Di masa lalu, pengembangan kernel NVIDIA juga membutuhkan bertahun-tahun upaya berbasis CUDA/CUTLASS level rendah
  • Belakangan, perkembangan DSL (bahasa spesifik domain) dan alat bantu berbasis AI telah menyederhanakan pengembangan kernel NVIDIA
    • Contoh: ThunderKittens (TK), CuTe DSL, TinyGrad “tinykittens”, TileLang, Gluon dan lainnya
  • Berdasarkan tren ini, di AMD juga mulai dieksplorasi kebutuhan akan primitif pemrograman baru

Prinsip desain HipKittens

  • HipKittens adalah versi yang dikembangkan untuk AMD, menyusul ThunderKittens (NVIDIA) dan ThunderMittens (Apple Silicon)
  • Konsep inti: abstraksi berbasis tile (tile abstraction)
    1. Generalisasi abstraksi tile – model komputasi berbasis tile yang efektif di NVIDIA juga dapat diterapkan secara alami di AMD
    2. Implementasi backend yang spesifik terhadap arsitektur – pola akses memori dan penjadwalan register dirancang berbeda untuk tiap perangkat keras
    3. Adaptivitas strategi penjadwalan – pada AMD CDNA3·CDNA4, penjadwalan berbasis wave tidak efisien, tetapi penjadwalan per tile tetap mempertahankan kesederhanaan dan performa
    Iklan
  • Dengan memisahkan antarmuka (tile dan operasi) dari implementasi (pemetaan ke perangkat keras), HipKittens menawarkan model yang dapat diterapkan secara umum ke berbagai arsitektur GPU

Hasil performa

  • Kernel Attention Forward: ditulis dengan sekitar 500 baris kode, dan mencapai performa lebih baik daripada kernel assembly AITER
    • Unggul pada berbagai dimensi head (D) dan panjang sekuens (N), baik dalam konfigurasi Causal maupun Non-Causal
  • Kernel GEMM: terdiri dari loop inti kurang dari 100 baris, dan mencapai performa terbaik untuk komputasi BF16 maupun FP8
  • Peningkatan dibanding performa terbaik sebelumnya juga terkonfirmasi pada kernel Attention Backward, Rotary, dan Fused Dropout-Residual-LayerNorm
  • Semua kernel tetap menjaga keterbacaan dan kemudahan modifikasi sambil memperoleh performa setara assembly buatan tangan

Ekspansi menuju AI multi-silikon

  • Untuk mewujudkan potensi AI, diperlukan pemanfaatan perangkat keras yang beragam dan terbuka
  • HipKittens bertujuan menjadikan GPU AMD sebagai platform yang benar-benar dapat diakses oleh pengembang AI
  • Dikombinasikan dengan software stack open source milik AMD, proyek ini menunjukkan kemungkinan transisi menuju ekosistem AI berbasis multi-silikon
  • Tulisan lanjutan akan membahas detail struktur teknis HipKittens (pratinjau part two)

Bacaan terkait

1 komentar

 
GN⁺ 2025-11-16
Opini Hacker News

  • Kami punya kontrak dengan AMD dan sedang melatih Llama 405B di MI350X untuk MLPerf.
    Situasi AMD memang jelas membaik. Jika Anda punya GPU AMD, saya sarankan mencoba memasang PyTorch dari pytorch.org dengan memilih Linux+ROCm. Tiga tahun lalu situasinya terasa putus asa, tetapi sekarang sebagian besar fitur utama sudah berjalan dengan baik. Saya menjalankan nanochat di MI300X dan langsung jalan begitu saja. MI350X juga stabil dengan cara yang sama.
    Tentu masih tertinggal dari NVIDIA, dan masih butuh banyak investasi pada ekosistem software, compiler, dan driver. Tapi dibandingkan situasi putus asa dua tahun lalu, sekarang sudah terlihat harapan.
    Jika ingin melihat bagian yang masih kurang dari backend LLVM AMD, menarik untuk membandingkan kode HipKittens dengan CUDA Kittens.
    Untuk training, NVIDIA dan Google ada di posisi 1, AMD di posisi 2, dan sisanya nyaris tidak ada. Intel dan Tenstorrent masih jauh, Huawei contoh kodenya mati karena segfault, Groq sudah menyerah menjual chip, Cerebras bahkan tidak bisa didapatkan, dan saya kehilangan minat pada Trainium karena butuh 5 hari hanya untuk mendapatkan satu instance

    • Saya penasaran seberapa jauh tinygrad dari kemampuan mengekspresikan atau mengeksplorasi fitur seperti optimisasi memori atau yang spesifik ke warp seperti ini. Saya juga ingin tahu seberapa rumit menambahkan fitur semacam itu ke tinygrad
    • Saya penasaran apakah menjalankan ROCm + PyTorch di hardware konsumen (non-MI) juga memerlukan driver kernel proprietari
    • Ungkapan "Cerebras tidak bisa didapatkan di mana pun" justru terdengar seperti tanda kemenangan
    • Saya sudah bekerja sebagai CEO AMD NeoCloud selama 2 tahun terakhir. Melihat langsung turnaround AMD seperti ini sangat menggembirakan.
      Setup awalnya memang masih agak kasar, tetapi jauh lebih baik dibanding sebelumnya. Misalnya, bahkan sebulan lalu nanochat belum berjalan dengan baik, sementara sekarang sudah bisa. Yang penting adalah sekarang orang-orang mulai tertarik pada ekosistem AMD.
      AI membutuhkan keragaman hardware. Jika semua hardware dan software AI dimonopoli oleh satu perusahaan, itu mungkin bagus bagi pemegang saham, tetapi buruk bagi kemajuan teknologi

  • Saya tidak mengerti valuasi perusahaan NVIDIA. Saat ini hanya segelintir algoritme seperti Transformer yang menang, dan data menjadi lebih penting. Karena tidak lagi dibutuhkan beragam kode HPC seperti dulu, para pesaing sekarang hanya perlu mengoptimalkan himpunan algoritme yang sempit.
    Jika bisa menjalankan vLLM dan Transformer dengan efisien, pasar yang sangat besar akan terbuka. Kalau begitu, rasanya AMD atau Huawei juga bisa dengan mudah mengejar, jadi saya penasaran apa sebenarnya parit pertahanan (moat) NVIDIA. Apakah InfiniBand saja cukup?

    • Benar, moat NVIDIA makin melemah. Perusahaan-perusahaan raksasa seperti MS, Google, Amazon, dan Apple semuanya sedang mengembangkan chip sendiri untuk menghindari ketergantungan pada NVidia.
      Di GPU datacenter NVIDIA masih kuat, tetapi Google memakai TPU dalam skala besar dan OpenAI juga memesan hardware AMD.
      Ekosistem CUDA masih penting, tetapi semua pihak aktif berupaya keluar dari sana. AMD, Intel, Qualcomm, dan lainnya juga ikut dalam persaingan ini. HipKittens tampak seperti langkah penting menuju ekosistem software yang netral
    • Cara termudah mengimplementasikan “segelintir algoritme” adalah membuat hardware komputasi serbaguna. Karena siklus pengembangan hardware panjang, pendekatan ini lebih aman
    • InfiniBand sedang digantikan oleh UEC. Untuk inferensi (inference), InfiniBand tidak dibutuhkan. Jadi di pasar inferensi tidak ada moat. Menggunakan AMD atau Google TPU adalah pilihan yang masuk akal
    • Senjata utama NVIDIA yang sebenarnya adalah ekosistem library CUDA yang sangat luas. Ada kode untuk hampir semua bidang
    • Transformer bukan satu teknologi tunggal. Cara implementasinya sangat beragam. Karena itu vLLM atau TRL tidak sesederhana kelihatannya

  • AMD seharusnya bisa saja mendanai proyek seperti ini, tetapi menurut saya mereka mungkin melewatkan peluang itu lagi. Dalam urusan GPU dan AI, mereka memang selalu seperti itu

    • AMD hanya berinvestasi pada software seminimal mungkin agar produk bisa dirilis. Bahkan sistem pengujian performanya pun lemah, dan bug regresi langsung sampai ke pelanggan.
      HipKittens adalah perbaikan, tetapi di internal AMD masih kurang kemampuan untuk melacak performa kernel. DevOps dialihdayakan ke TCS di India, dan pihak sana tidak benar-benar memahami situasinya.
      Tim yang punya pemimpin bagus biasanya menjalankan tim shadow IT sendiri. ROCm tidak punya tim seperti itu, dan perbaikan baru mulai setelah pelanggan cloud besar mengeluh.
      Bahkan ketika merekrut talenta, mereka menawarkan gaji di bawah pasar. Perbandingannya karena mereka menyamakan diri dengan Qualcomm atau Walmart.
      Dalam 4 tahun terakhir, bonus tidak pernah dibayarkan penuh
    • Ungkapan “AMD tidak pernah melewatkan kesempatan untuk melewatkan kesempatan” sangat tepat. Hardware Instinct pada dasarnya sudah mampu bersaing dengan NVidia, tetapi dukungan softwarenya mengerikan.
      Seperti dulu dengan Radeon VII, mereka memutus dukungan atau terlalu sering mengacak ulang ekosistem, sehingga tidak pernah matang. Karena kurangnya kompatibilitas CUDA dan minim investasi, sebagian besar organisasi akhirnya memilih NVIDIA saja.
      Meski begitu, belakangan ini mereka terus berinvestasi di ekosistem ROCm dan Instinct, jadi sedikit demi sedikit membaik. Hanya saja, di bidang networking NVIDIA masih jauh lebih unggul
    • Kalau melihat tabel perbandingan performa, AMD bisa saja sekarang sudah setara NVIDIA. Tapi mereka gagal karena tidak ada software. Mendesain chip itu pekerjaan yang lebih sulit, jadi masalahnya justru mereka tidak memahami software
    • Dulu ada orang-orang yang berkontribusi kernel yang dioptimalkan untuk ROCm, tetapi AMD menutup PR dan tidak menggabungkannya. Perilaku itu benar-benar sulit dipahami
    • Sekarang pendanaannya ada dan proyek ini sedang berjalan dengan baik

  • Saya penasaran apakah ada proyek yang dibangun di atas ThunderKittens.
    Jika ini versi yang di-port ke HIP, rasanya nilai kepraktisannya jauh lebih besar daripada sekadar port CUDA biasa

  • Ungkapan “raw assembly vs cooked assembly” menarik.
    Dulu menulis kode assembly langsung untuk CPU itu hal yang umum. Untuk GPU pun tidak perlu terlalu takut. Pada akhirnya compiler juga yang akan menghasilkan kode itu

  • Secara matematis, inferensi (inference) adalah operasi aljabar linear/BLAS dasar.
    Saya penasaran apakah mungkin ada API inferensi yang ringkas yang mencakup 80% kasus penggunaan dengan mempertimbangkan dtype dan sparsity. Sepertinya tidak perlu serumit CUDA

  • composable-kernel adalah software yang paling sering menyebabkan OOM (kehabisan memori) di sistem Gentoo saya. Saat Clang mengompilasi CK, pemakaian memorinya lebih dari 2,5GB per thread

    • Saya meninjau paket CK untuk Debian, dan saat build dengan -j32 bahkan workstation 64GB pun kena OOM. Saat dijalankan dengan -j1, build akhirnya berhasil setelah 190 jam.
      Di server build resmi sepertinya jumlah arsitektur harus dikurangi. Saya dengar sebagian besar paket dependensi hanya membutuhkan header. Sedang ada pekerjaan perbaikan untuk mengurangi waktu build
    • Butuh lebih dari 10 menit untuk instansiasi template bagi satu kernel saja itu tingkat yang mengejutkan.
      device_grouped_conv2d_fwd_xdl_ngchw_gkcyx_ngkhw_f16_instance, entah apa sebenarnya yang dilakukan pada clang

  • Saya penasaran apakah perkembangan seperti ini akan membuat chip AMD konsumen juga bisa menjalankan LLM dengan baik.
    Misalnya saya sedang mempertimbangkan laptop HP OMEN MAX 16 dengan CPU/iGPU AMD dan RTX 5080, apakah sisi AMD bisa bersaing dengan RTX?

    • Orang bisa menganggap dGPU selalu lebih cepat, tetapi batas kapasitas memori bisa menjadi hambatan.
      Tulisan blog ini atau hasil Mac kelas atas cukup menarik untuk dilihat
    • Saya menjalankan Qwen3 Coder 30B di RTX7900XTX dengan Ollama. Berjalan cukup baik. Sebagian beban tampaknya dialihkan ke memori sistem dan CPU Ryzen 7.
      Sedikit lebih lambat daripada API Sonnet 4, tetapi untuk performa lokal sebesar ini saya sudah sangat puas.
      Kombinasi Opencode + Ollama + Qwen3 Coder adalah alternatif yang sangat bagus untuk ClaudeCode + Sonnet4.
      Tentu, kalau Anda berharap AI menggantikan seluruh pekerjaan coding, mungkin rasanya berbeda. Tetapi sebagai asisten pribadi, ini sempurna

  • Saya tidak mengerti kenapa AMD sepenuhnya mengabaikan B300