Peningkatan Performa GPU yang Eksplosif
(hazyresearch.stanford.edu)- Di tengah meningkatnya biaya komputasi AI, Hazy Research merangkum bahwa inti optimasi performa GPU adalah menjaga tensor core NVIDIA H100 tetap terus bekerja
- H100 menghasilkan 989 TFLOPs pada perkalian matriks half-precision, tetapi hanya sekitar 60 TFLOPs untuk komputasi umum, sehingga utilisasi turun tajam begitu tensor core berhenti
- Untuk mendekati performa maksimum, WGMMA, penempatan shared memory, pembuatan alamat, dan occupancy harus ditangani bersama; tanpa
wgmma.mma_async, microbenchmark hanya mencapai sekitar 63% dari puncak - DSL tertanam CUDA yang dirilis terbuka, ThunderKittens, menyederhanakan penulisan kernel keluarga FlashAttention dengan abstraksi tile dan vector yang membungkus kompleksitas seperti swizzling dan register layout
- Kernel forward FlashAttention-2 untuk H100 ditulis dalam sekitar 100 baris dan sekitar 30% lebih cepat daripada FlashAttention-2, sementara kernel Based linear attention berjalan pada 215 TFLOPs
Kondisi yang Menentukan Performa H100
- AI menggunakan banyak komputasi, dan dalam beberapa tahun terakhir Hazy Research mengerjakan cara agar AI memakai lebih sedikit komputasi atau berjalan lebih efisien pada komputasi yang tersedia
- Contoh penghematan komputasi mencakup Based, Monarch Mixer, H3, Hyena, dan S4
- Contoh eksekusi yang efisien mencakup FlashAttention, FlashAttention-2, dan FlashFFTConv
- Tujuan praktisnya adalah merangkum pelajaran dari upaya membuat GPU cepat, serta merilis ThunderKittens, sebuah DSL tertanam CUDA yang membantu penulisan kernel cepat
- Secara lebih luas, topik ini membahas bagaimana pemahaman hardware mengubah cara memandang komputasi AI
Struktur dan Bottleneck NVIDIA H100
- GPU H100 SXM dibahas berdasarkan konfigurasi berikut
- 80GB HBM3, bandwidth 3TB/s
- Cache L2 50MB, bandwidth 12TB/s, terbagi menjadi dua seksi 25MB di seluruh GPU dan terhubung dengan crossbar
-
132 SM
- Setiap SM memiliki cache L1 256KB yang mencakup shared memory hingga 227KB, dengan total bandwidth sekitar 33TB/s
- Hardware baru Hopper, Tensor Memory Accelerator(TMA), menangani pembuatan alamat asinkron dan memory fetch
- Setiap SM terdiri dari 4 quadrant, dan tiap quadrant memiliki warp scheduler, 512 vector register, tensor core untuk perkalian matriks, serta instruksi bawaan paralel
- Semua komputasi terjadi di SM, dan sebagian besar diproses di register
- Kunci performa pada H100 adalah menjaga tensor core terus dalam keadaan fed
- H100 menyediakan 989 TFLOPs untuk perkalian matriks half-precision, dan sekitar 60 TFLOPs untuk komputasi “lainnya”
- Pada cycle ketika tensor core digunakan, utilisasi hardware setidaknya mencapai 94%
- Pada cycle ketika tensor core tidak digunakan, utilisasi maksimum hanya sekitar 6%
WGMMA: Instruksi yang Diperlukan tetapi Rumit
- H100 memiliki instruksi warp group matrix multiply accumulate bernama
wgmma.mma_async- Di PTX:
wgmma.mma_async - Di SASS:
HGMMA/IGMMA/QGMMA/BGMMA
- Di PTX:
wmma.mma.syncdanmma.syncpada GPU sebelumnya bersifat sinkron: satu warp berisi 32 thread memasukkan data ke tensor core lalu menunggu hasilnyawgmma.mma_asyncmembuat 128 thread berurutan berkolaborasi dan tersinkronisasi di seluruh quadrant SM, lalu memulai perkalian matriks asinkron langsung dari shared memory- Selama perkalian matriks berjalan, warp dapat mengerjakan tugas lain dengan register
- Hasilnya dapat ditunggu pada waktu yang diinginkan
- Dalam microbenchmark, instruksi ini diperlukan untuk mengeluarkan seluruh compute H100
- Jika tidak digunakan, GPU teramati berhenti di sekitar 63% dari utilisasi puncak
- Ini mungkin karena tensor core membutuhkan pipeline hardware yang dalam bahkan untuk resource lokal
- Kesulitan terbesarnya adalah kompleksitas memory layout
- Layout shared memory yang tidak di-swizzle memiliki coalescing yang sangat buruk, sehingga membutuhkan banyak bandwidth L2
- Layout swizzled memerlukan waktu untuk dipahami karena dokumentasinya keliru
- Layout swizzled tampaknya hanya bekerja pada shape matriks tertentu dan tidak cocok dengan fitur lain dari
wgmma.mma_async - Hardware dapat melakukan sub-matrix transpose di jalur menuju tensor core, tetapi hanya ketika layout tidak di-swizzle
- Pada kernel seperti FlashAttention, TMA dan cache L2 cukup cepat untuk menyembunyikan masalah ini sampai batas tertentu
- Untuk memakai hardware sepenuhnya, request memory harus di-coalesce dan bank conflict harus dihindari, sehingga kontrol layout menjadi penting
Shared Memory dan Bank Conflict
- Latensi single-access shared memory tampaknya sekitar 30 cycle, dan selama waktu itu tensor core di SM dapat menjalankan hampir dua perkalian matriks persegi 32x32
- Pekerjaan sebelumnya seperti FlashAttention terutama berfokus pada bottleneck HBM-SRAM, dan dulu bottleneck ini memang penting
- Karena HBM makin cepat dan tensor core tumbuh lebih cepat daripada bagian lain chip, latensi kecil pada shared memory pun menjadi hal yang perlu dihilangkan atau disembunyikan
- Shared memory terbagi menjadi 32 bank, sehingga bank conflict dapat terjadi jika tidak berhati-hati
- Jika beberapa potongan memory berbeda pada memory bank yang sama diminta secara bersamaan, request akan diserialisasi
- Berdasarkan pengalaman, kernel dapat menjadi lambat secara tidak seimbang
- Register layout yang dibutuhkan instruksi WGMMA dan MMA dapat mengalami bank conflict jika digunakan secara sederhana
- Solusinya adalah menata ulang shared memory dengan berbagai pola swizzling untuk menghindari conflict
- Jika memungkinkan, lebih baik menghindari perpindahan antara register dan shared memory; saat diperlukan, sebaiknya gunakan hardware bawaan seperti WGMMA dan TMA untuk memindahkan data secara asinkron
- Perpindahan sinkron menggunakan warp nyata adalah cara paling umum, tetapi hampir seperti fallback terburuk
Pembuatan Alamat dan TMA
- Pada H100, tensor core dan memory sama-sama cepat, sehingga pekerjaan membuat memory address yang akan di-fetch itu sendiri memakan porsi signifikan dari resource chip
- Ini makin terlihat ketika pola interleaved atau swizzling yang kompleks ditambahkan
- Tensor Memory Accelerator(TMA) dari NVIDIA memungkinkan penentuan layout tensor multidimensi di global/shared memory, mengambil subtile dari tensor tersebut secara asinkron, lalu memicu barrier saat selesai
- TMA mengurangi biaya pembuatan alamat dan juga mempermudah penyusunan pipeline
- TMA dinilai esensial untuk mengeluarkan potensi H100, seperti
wgmma.mma_async- Berdasarkan pengalaman, TMA mungkin bahkan lebih penting daripada WGMMA
- TMA menghemat resource register dan instruction dispatch
- TMA juga memiliki kemampuan melakukan reduction asinkron ke global memory, yang berguna dalam backward kernel yang kompleks
- TMA juga memerlukan sebagian reverse engineering untuk memahami mode swizzling, tetapi tidak sesakit WGMMA
Biaya yang Disembunyikan Occupancy
- Dalam CUDA, occupancy berarti jumlah thread yang di-co-schedule pada hardware eksekusi yang sama
- Warp scheduler pada quadrant SM mencoba meng-issue instruction ke warp yang siap menerima instruksi pada setiap cycle
- H100 dalam beberapa hal kurang bergantung pada occupancy dibanding generasi sebelumnya
- Berkat fitur asinkron, satu instruction stream pun dapat secara bersamaan menyibukkan memory fetch, matrix multiply, shared memory reduction, dan register math
- Namun occupancy sangat berguna untuk menyembunyikan kesalahan dan biaya sinkronisasi
- Pipeline yang dirancang sempurna bisa cepat tanpa occupancy tambahan
- Dari pengamatan nyata, GPU NVIDIA tampak dirancang dengan mempertimbangkan occupancy
- Karena ada banyak kemungkinan sinkronisasi dan kesalahan, peningkatan occupancy sering memperbaiki utilisasi hardware yang terealisasi
- Pada H100, occupancy berguna, tetapi pada A100 dan RTX 4090 masing-masing dinilai lebih penting
- Disebutkan kemungkinan karena keduanya lebih bergantung pada instruction dispatch sinkron dibanding H100
ThunderKittens: DSL Kecil di Dalam CUDA
- ThunderKittens adalah DSL tertanam CUDA yang dibuat untuk mempermudah penulisan kernel cepat pada H100
- Awalnya dibuat untuk penggunaan internal lab, lalu kemudian dirilis publik
- Namanya dipilih karena kittens itu lucu dan mereka merasa lucu jika harus mengetik
kittens::di kode - ThunderKittens menargetkan kesederhanaan dan menyediakan empat templated type
- Register tiles: tensor 2D di atas register file
- Register vectors: tensor 1D di atas register file
- Shared tiles: tensor 2D di dalam shared memory
- Shared vectors: tensor 1D di dalam shared memory
- Tile diparameterisasi dengan height, width, dan layout
- Register vector diparameterisasi dengan length dan layout, sedangkan shared vector hanya memakai length
- shared vector umumnya tidak mengalami bank conflict
- Operasi yang disediakan memanipulasi tile dan vector pada level warp atau level warp group kolaboratif
- initializer: operasi seperti membuat shared vector menjadi zero
- unary op: seperti
exp - binary op: seperti
mul - row/column op: seperti
row_sum
- ThunderKittens tertanam di dalam CUDA, sehingga tidak seperti library semacam Triton, abstraksinya dijelaskan dapat gagal secara “gracefully”
- Jika ada fitur yang kurang, pengguna dapat memperluasnya dengan cara yang diinginkan
Contoh FlashAttention dan Performa
- Sebagai contoh ThunderKittens, disajikan kernel forward FlashAttention sederhana untuk RTX 4090
- Hanya menangani headdim=64
nharus kelipatan 256- Ditulis dalam sekitar 60 baris kode CUDA
- Utilisasi hardware 75%
- Sebagian besar kompleksitas ada pada algoritme itu sendiri, bukan pada pola swizzling atau register layout
- Forward pass FlashAttention-2 untuk H100 juga ditulis dengan ThunderKittens
- ThunderKittens membungkus kompleksitas TMA, WGMMA, mode swizzling, dan descriptor
- Kernelnya sekitar 100 baris
- Di H100, sekitar 30% lebih cepat daripada FlashAttention-2
- ThunderKittens membungkus layout dan instruction serta menyediakan primitive, seperti “mini-pytorch” yang dapat digunakan di GPU
- Kernel untuk Based linear attention dan architecture lain yang akan dirilis ke depan juga dipublikasikan bersama
- Kernel Based linear attention berjalan pada 215 TFLOPs
- Jika recompute dari algoritme itu sendiri diperhitungkan, angkanya melewati 300 TFLOPs
- Linear attention secara teori lebih efisien, tetapi secara historis efisiensinya jauh lebih rendah pada hardware nyata
- Hasil ini dipandang dapat memperluas cakupan aplikasi dengan throughput tinggi
Cara Berpikir Berpusat pada Tile
- Alasan ThunderKittens bekerja baik dinilai karena ia tidak mencoba melakukan semuanya
- CUDA jauh lebih ekspresif daripada ThunderKittens
- ThunderKittens adalah DSL kecil dan sederhana
- Abstraksi intinya adalah small tile, dan ini dianggap selaras dengan arah AI dan hardware
- ThunderKittens tidak mendukung dimensi yang lebih kecil dari 16
- Hardware juga dinilai tidak terlalu menginginkan dimensi sekecil itu
- Pertanyaannya diajukan kira-kira: “Jika matrix multiply lebih kecil dari 16x16, bisakah kita yakin itu AI?”
- Perspektif era CPU yang melihat 32-bit word sebagai register dinilai tidak cocok untuk hardware AI
- Vector register 1024-bit CUDA dilihat sebagai satu langkah ke arah yang benar
- Di sini, register adalah data dari tile 16x16
- AI masih berpusat pada matrix multiply, reduction, dan reshape, sehingga abstraksi tile dianggap cocok untuk AI maupun hardware
- Ke depan, ide-ide AI perlu disusun ulang agar terpetakan dengan baik ke hardware
- Ukuran recurrent state harus cukup besar agar dapat masuk ke SM
- Compute density tidak boleh lebih rendah daripada tingkat yang dituntut hardware
- Menyesuaikan pelajaran dari hardware ke desain AI merupakan arah penting ke depan
Rencana Dukungan AMD
- Dukungan AMD hardware untuk ThunderKittens akan segera hadir
1 komentar
Komentar Hacker News
Pertanyaan “kalau perkalian matriksnya lebih kecil dari 16x16, yakin itu benar-benar AI?” menarik
Kebutuhan hardware AI makin terlihat jelas. GPU awalnya dirancang untuk kegunaan yang sama sekali berbeda, tetapi dipakai untuk AI karena hardware perkalian matriksnya bagus, dan “AI GPU” bisa mengurangi sebagian fungsi yang ada pada GPU sungguhan
Representasi angka juga bergerak ke arah yang makin pendek, seperti floating point 16-bit, 8-bit, 2-bit, hingga 1-bit, dan suatu saat titik yang tepat akan ditentukan. Tulisan ini menunjukkan bahwa hardware yang menyukai tile 16x16 cukup masuk akal. Saat ini mungkin sudah ada orang yang menulis hal semacam ini dalam VHDL, atau kemungkinan besar akan segera melakukannya
Pada akhirnya, sepertinya akan muncul perangkat yang lebih sederhana, kurang serbaguna, dan murah, yang hanya menjalankan operasi “AI” semaksimal mungkin tanpa beban hardware yang tidak perlu
Nvidia mungkin juga sedang mengerjakannya, tetapi dari sisi bisnis mungkin lebih baik mempertahankan bentuk kartu video, yaitu perangkat yang menggabungkan gim/hiburan/cryptocurrency/AI
[1] https://github.com/hollance/neural-engine/blob/master/docs/a...
Ini mengingatkan pada masa ketika Nervana milik Naveen Rao membuat driver Nvidia Maxwell yang lebih cepat daripada driver Nvidia sendiri. Tidak semua kesalahan dokumentasi pada produk yang tumbuh cepat adalah taktik menghadapi pesaing, tetapi mengingat para peneliti butuh waktu lama untuk merekayasa balik wgmma, ditambah situasi politik AS-Tiongkok seputar H100, terlihat seperti Nvidia memakai trik lama untuk mempertahankan parit pertahanannya
Jadi, alih-alih terlalu mendalami keunikan H100, kita perlu melihat bahwa “hardware seperti apa yang diinginkan AI” juga mencakup situasi komersial
https://www.amd.com/en/products/accelerators/alveo/v80.html
XDNA Architecture
https://www.amd.com/en/technologies/xdna.html
Bagian “Kebohongan NVIDIA. Ini adalah representasi yang sangat menyesatkan tentang tata letak 128b swizzled wgmma yang sebenarnya. Karena diagram ini membuatku kehilangan 3 minggu hidup yang tak bisa dikembalikan, aku mempermalukannya di depan umum” terasa berkesan
Aku penasaran apakah ada orang yang akan terkejut mengetahui bahwa bagian yang sangat besar dari kemajuan AI ada pada rekayasa seperti optimasi perkalian matriks, dan bahwa sebagian besar rekayasa itu adalah rekayasa balik chip NVIDIA
Warp scheduler, 4 kuadran, tensor memory accelerator, tata letak unswizzled wgmma…
Batas antara istilah GPU dan technobabble ala Star Trek makin kabur
Saat melihat tulisan lain pun kadang terpikir hal yang sama. Aku membayangkan seperti apa rasanya bagi seseorang yang menerima tautan tulisan di sini lalu membacanya. Mungkin seperti masuk ke acara penggemar Trek yang sedang mendiskusikan warp core
Untuk mengurangi konsumsi daya dan meningkatkan kecepatan inferensi AI, sepertinya yang terbaik adalah beralih ke rangkaian aproksimasi analog
Yang dibutuhkan bukan perkalian dan penjumlahan floating-point yang sempurna, melainkan hanya perangkat yang menerima dua tegangan input dan menghasilkan tegangan output yang cukup mendekati hasil perkalian
Keuntungan besarnya adalah, alih-alih merepresentasikan float16 dengan 16 jalur, angka itu direpresentasikan lewat tegangan pada 1 jalur. Secara teori, presisi yang jauh lebih tinggi daripada float32 pun mungkin bisa dicapai. Selain itu, nilainya bisa dihubungkan langsung tanpa dimuat ke unit aritmetika-logika, sehingga penghematan area die dan daya berpotensi mencapai beberapa orde magnitudo
Misalnya, menerima bahwa satu dari sejuta bit output terbalik demi meningkatkan rasio performa/daya. Ini mungkin sulit pada float32, di mana satu nilai infinity saja bisa merusak semuanya, tetapi pada int8 sepertinya masih bisa ditoleransi jika yang diinginkan 0 namun sesekali keluar 128
[1] Saya tidak begitu yakin apakah unit floating-point matriks pada H100 benar-benar mematuhi IEEE 754
Jaringan saraf biologis tidak nyaris terhubung penuh seperti jaringan saraf buatan pada umumnya, dan koefisien koneksi input/output neuron kurang dari 10, jadi sangat lokal. Sepengetahuan kita, dalam biologi juga tidak ada backpropagation; yang ada justru umpan balik dan rekursi
Mungkin juga ada sel atau proses pendukung yang belum kita ketahui tetapi esensial bagi fungsi sistem saraf pusat. Bahkan pada tingkat tinggi, kemungkinan ada cukup banyak konektivitas yang "di-hardcode", dan sebagian sudah diketahui. Misalnya, neuron pendengaran di telinga saling terhubung dan terjadi sesuatu yang mirip konvolusi untuk menentukan lokasi suara. Ini bukan fenomena emergen, melainkan fungsi yang bisa terjadi tanpa pelatihan
Tidak mengherankan, karena kehidupan menemukannya melalui miliaran tahun dan jumlah generasi yang kira-kira sebanding. Secara teori ini juga mungkin dilakukan dengan perangkat lunak, tetapi mengingat otak primata/manusia memiliki lebih dari 1 triliun neuron, ini akan sangat sulit bahkan dengan mesin kelas seribu core saat ini. Bahkan "cloud" pun tidak akan memenuhi konektivitas dan latensi yang dibutuhkan
Akan keren jika pendekatan seperti ini berhasil memodelkan sesuatu setingkat cacing atau serangga
Tulisan ini mengingatkan saya lagi pada kesenangan yang saya rasakan di kelas CS 149 pemrograman paralel
Gaya tulisan artikel ini benar-benar mengesankan, dan saya menantikan melihat ini di AMD MI300x. Kalau ingin memakai waktu di perangkat saya, beri tahu saya
Saya juga penasaran seberapa baik ia benar-benar bekerja, atau apakah lebih baik menabung sedikit lagi untuk membeli XTX alih-alih 7900 XT, serta seberapa besar berkurangnya VRAM memengaruhi kegunaan nyata
Pembaca tidak seharusnya perlu mencari di knowyourmeme.com untuk memahami apa yang ingin dikatakan para penulis. Saya bahkan tidak tahu apa arti judul ini, dan menurut saya itu sangat meleset dari tujuannya
Saya penasaran harus mulai dari mana dan mengikuti roadmap seperti apa agar bisa memahami tulisan seperti ini sepenuhnya
Lalu sebaiknya coba menulis sendiri kernel CUDA yang melakukan perkalian vektor-matriks. Dengan pycuda, Anda bisa fokus pada kernel dan menulis sisanya dengan Python. Anda bisa bilang ke ChatGPT bahwa Anda ingin membuat sendiri implementasi untuk mengalikan vektor 4000 elemen dengan matriks 4000x12000, lalu minta dipandu melalui seluruh prosesnya
Untuk sewa GPU, Runpod bagus, dan saat ini tersedia dari GPU murah sampai H100. Mulailah dengan GPU kelas rendah dulu
Saya menghabiskan 2 bulan untuk mengimplementasikan dan mengoptimalkan kernel perkalian matriks dengan Spiral
Grafik di README GitHub (https://github.com/HazyResearch/ThunderKittens/blob/main/att...) terlalu memusingkan. Apakah diagram batang bergelombang seperti ini legal? :P
[1]: https://matplotlib.org/stable/gallery/showcase/xkcd.html#sph...
Nama ThunderKittens itu bagus. Saya ingin melihat ThunderKittens menangani backpropagation FlashAttention, yang satu orde magnitudo lebih sulit daripada forward pass
causal: https://github.com/HazyResearch/ThunderKittens/blob/main/exa...
non-causal: https://github.com/HazyResearch/ThunderKittens/blob/main/exa...
Bukankah riset seperti ini sudah dilakukan oleh tim-tim yang membuat NPU saat ini? Misalnya, chip Groq bisa mencapai performa seperti sekarang karena memakai arsitektur khusus AI. Di sisi konsumen, Apple Silicon juga cukup mumpuni
Saya bukan orang di bidang ini, tetapi rasanya ada batas jika hanya mengandalkan prosesor serbaguna yang berkomunikasi lewat jalur yang relatif lambat. Memikirkan ulang desain di tingkat hardware, lalu pada akhirnya menurunkan harga di pasar konsumen, tampaknya strategi jangka panjang yang lebih baik
Dengan beberapa ratus dolar orang bisa membeli GPU Nvidia, atau laptop gaming dengan 4050 dan VRAM 6GB seharga 900 dolar, jadi sulit menyebut AI berbasis CPU itu mumpuni
Di tempat kerja saya juga tidak ada GPU, jadi saya mencoba berbasis CPU, tetapi selain memakai model kecil dan menunggu, itu tidak realistis. Akhirnya saya meminta komputer dengan GPU
"Secara teknis bisa" berbeda dengan "benar-benar nyaman dipakai". Nvidia benar-benar enak digunakan, sedangkan CPU menyakitkan dan membuat frustrasi