Peningkatan Performa GPU yang Eksplosif

Karakteristik GPU H100

• Menyediakan memori HBM3 80GB dan bandwidth 3TB/s (dalam praktiknya sedikit lebih rendah)
• Menyediakan cache L2 50MB dan bandwidth 12TB/s. Pada GPU, ini dibagi menjadi dua seksi 25MB yang terhubung melalui crossbar (crossbar menjadi faktor penurunan performa)
• Terdiri dari 132 Streaming Multiprocessor (SM), dan tiap SM memiliki konfigurasi berikut:
  • Menyediakan hingga 227KB shared memory di dalam cache L1 256KB (secara gabungan sekitar 33TB/s bandwidth)
  • Menyediakan Tensor Memory Accelerator (TMA) yang memungkinkan pembuatan alamat asinkron dan memory fetching. Fitur seperti dukungan jaringan memori on-chip juga tersedia, tetapi tidak dibahas dalam tulisan ini
  • Dibagi menjadi 4 kuadran, dan tiap kuadran terdiri dari warp scheduler, 512 vector register (masing-masing berisi 32 word 4-byte), tensor core untuk matrix multiply, serta instruction bawaan yang mendukung operasi paralel seperti sum/multiply

Tips optimasi performa GPU H100

• Penting untuk memanfaatkan Tensor Core semaksimal mungkin. H100 memiliki performa operasi perkalian matriks FP16 sebesar 989 TFLOPs, sehingga tingkat pemanfaatan Tensor Core sangat menentukan utilisasi keseluruhan GPU
• Namun, untuk memanfaatkan Tensor Core semaksimal mungkin, hal-hal berikut perlu dipertimbangkan:
  • Wajib memanfaatkan instruksi Warp Group Matrix Multiply Accumulate (WGMMA), tetapi penggunaannya tidak mudah
  • Shared Memory tidak secepat yang dibayangkan, dan penggunaannya memerlukan perhatian yang cermat
  • Pembuatan alamat memiliki biaya yang besar, sehingga perlu dioptimalkan dengan memanfaatkan Tensor Memory Accelerator (TMA) dan sejenisnya
  • Meningkatkan occupancy tetap membantu, dan register adalah resource utama
• Karakteristik ini sampai batas tertentu berlaku juga pada GPU lain, tetapi khususnya pada H100, pemanfaatan Tensor Core sangat penting sekaligus cukup rumit

ThunderKittens: CUDA embedded DSL yang dioptimalkan untuk H100

• Embedded DSL berbasis CUDA yang dikembangkan untuk memaksimalkan performa GPU modern seperti NVIDIA H100
• Menyediakan 4 tipe template berbasis tile berikut:
  • Register Tile (tensor 2D yang disimpan di register)
  • Register Vector (tensor 1D yang disimpan di register)
  • Shared Tile (tensor 2D yang disimpan di Shared Memory)
  • Shared Vector (tensor 1D yang disimpan di Shared Memory)
• Juga menyediakan berbagai operator untuk manipulasi tile (exp, mul, sum, dll.) pada level warp atau warp group
• Saat kernel Flash Attention dan Flash Attention 2 yang sudah ada diimplementasikan dengan ThunderKittens, kodenya menjadi jauh lebih ringkas dan performa di H100 meningkat hingga 30%
• Kernel Based Linear Attention juga diimplementasikan dengan ThunderKittens dan mencapai performa 215 TFLOPs (karena karakteristik algoritmenya, jika termasuk recompute maka lebih dari 300 TFLOPs)

Tinjauan dari sudut pandang filosofis

• Alasan ThunderKittens bekerja dengan baik adalah karena ia tidak berusaha mendukung semuanya, melainkan menyediakan abstraksi berbasis tile yang sederhana namun sangat cocok dengan arsitektur GPU
• Penggunaan vector register 1024-bit alih-alih word 32-bit tradisional juga merupakan kemajuan, tetapi yang lebih penting adalah pergeseran paradigma untuk melihat tile 16x16 sebagai satuan register
• Mengingat workload AI pada akhirnya didominasi oleh matrix multiply, reduction, dan reshape, pendekatan berbasis tile masuk akal, dan dari sudut pandang hardware juga akan berevolusi ke arah mendukung perkalian matriks kecil selain systolic array
• Lebih jauh lagi, perlu ada perubahan cara berpikir ke arah merancang algoritme AI dalam bentuk yang dioptimalkan untuk hardware. Misalnya, ukuran state RNN dibatasi agar muat di SM, dan densitas komputasinya juga disesuaikan dengan tingkat yang dibutuhkan hardware

Opini GN⁺

• ThunderKittens bisa menjadi opsi yang menarik bagi developer yang sudah akrab dengan CUDA. Alasannya, peningkatan performa bisa diperoleh dengan mudah tanpa perlu banyak mengubah kode kernel yang ada
• Namun, untuk pemula, hambatan masuknya masih bisa cukup tinggi. Ke depan, tampaknya perlu lebih banyak contoh kode dan materi pembelajaran yang mendukung
• Rencana untuk memperluas dukungan ThunderKittens bukan hanya ke H100 tetapi juga ke GPU lain seperti AMD terasa menarik. Ini tampaknya bisa membantu mengurangi vendor lock-in
• Pada akhirnya, poin yang sangat penting adalah merancang model/algoritme AI itu sendiri dalam bentuk yang dioptimalkan untuk hardware. Untuk itu, pemahaman mendalam tentang karakteristik hardware harus didahulukan, dan ThunderKittens dapat membantu developer memperoleh insight semacam itu
• Diharapkan riset dan pengembangan berkelanjutan serta kontribusi open source dari Hazy Research akan sangat membantu menghidupkan ekosistem CUDA. Namun, dalam jangka panjang, tampaknya juga dibutuhkan kemunculan framework dengan level abstraksi yang lebih tinggi

[GN⁺]

Komentar Hacker News

• Kebutuhan perangkat keras AI semakin jelas. GPU awalnya dirancang untuk tujuan lain, tetapi digunakan untuk AI karena memiliki perangkat keras perkalian matriks yang baik. "GPU AI" dapat menghilangkan sebagian fungsi GPU yang sebenarnya, dan trennya mengarah ke angka yang lebih pendek (floating point 16-bit, 8-bit, 2-bit, 1-bit). Makalah ini menyiratkan bahwa perangkat keras yang menyukai tile 16x16 memiliki banyak keuntungan.

• Diperlukan prosesor pendamping khusus AI (NPU). Ini terutama dibutuhkan pada sistem desktop kelas prosumer untuk pengembang, profesional, gamer, dan lainnya. GPU berfungsi di perusahaan, tetapi kurang praktis untuk penggunaan AI dari sisi komputasi pribadi. Terutama karena keterbatasan VRAM dan tidak adanya API terbuka standar selain Vulkan.

• Untuk memajukan riset AI, kita perlu mempelajari neurosains dan psikologi dengan lebih baik. Selain itu, hal-hal yang berkaitan dengan topologi graf pada jaringan saraf juga mungkin relevan.

• Apakah ini CUTLASS yang dibuat ramah pengguna?

• Maskot ThunderKittens punya nuansa kucing/Sony Aibo. Tampaknya dihasilkan dengan baik oleh AI.

• Jika Universal Basic Compute (UBC) dipertimbangkan sebagai pengganti Universal Basic Income, itu akan menjadi masa depan yang sangat distopis. Bayangkan satu perusahaan seperti Nvidia membuat semua komputasi.