H100 GPU’nun özellikleri
- 80GB HBM3 bellek ve 3TB/s bant genişliği sunar (gerçekte biraz daha düşüktür)
- 50MB L2 önbellek ve 12TB/s bant genişliği sunar. GPU içinde iki adet 25MB bölüme ayrılmıştır ve bir crossbar ile bağlanır (crossbar performans düşüşü etkenidir)
- 132 adet Streaming Multiprocessor (SM) içerir ve her SM şu yapıya sahiptir:
- 256KB L1 önbellek içinde en fazla 227KB paylaşımlı bellek sunar (birlikte yaklaşık 33TB/s bant genişliği)
- Asenkron adres üretimi ve bellek fetching yapabilen Tensor Memory Accelerator (TMA) sunar. On-chip bellek ağı desteği gibi işlevler de sağlar ancak bu yazıda ele alınmıyor
- 4 quadrant’a ayrılır ve her quadrant; warp scheduler, 512 vector register (her biri 32 adet 4-byte word içerir), matrix multiply için tensor core ve sum/multiply gibi paralel işlemleri destekleyen gömülü instruction’lardan oluşur
H100 GPU performansını optimize etme ipuçları
- Tensor Core’ları mümkün olduğunca etkin kullanmak önemlidir. H100, FP16 matris çarpımı için 989 TFLOPs performans sunduğundan, Tensor Core kullanım oranı toplam GPU kullanımını büyük ölçüde belirler
- Ancak Tensor Core’ları en üst düzeyde kullanmak için aşağıdakiler dikkate alınmalıdır:
- Warp Group Matrix Multiply Accumulate (WGMMA) instruction’larını kullanmak zorunludur, ancak kullanımları zordur
- Shared Memory düşünüldüğü kadar hızlı değildir ve kullanırken dikkatli olunmalıdır
- Adres üretiminin maliyeti yüksektir; bu yüzden Tensor Memory Accelerator (TMA) gibi yapılarla optimizasyon gerekir
- Occupancy’yi artırmak hâlâ faydalıdır ve register temel kaynak olmaya devam eder
- Bu özellikler yalnızca H100 için değil, başka GPU’lar için de bir ölçüde geçerlidir; ancak özellikle H100’de Tensor Core kullanımı hem önemli hem de zordur
ThunderKittens: H100 için optimize edilmiş CUDA embedded DSL
- NVIDIA H100 gibi modern GPU’ların performansını mümkün olan en üst seviyeye çıkarmak için geliştirilmiş, CUDA tabanlı bir embedded DSL
- Aşağıdaki 4 tile tabanlı şablon türünü sunar:
- Register Tile (register’da saklanan 2D tensör)
- Register Vector (register’da saklanan 1D tensör)
- Shared Tile (Shared Memory’de saklanan 2D tensör)
- Shared Vector (Shared Memory’de saklanan 1D tensör)
- Ayrıca tile işlemleri için çeşitli operatörleri (exp, mul, sum vb.) warp veya warp group seviyesinde sunar
- Mevcut Flash Attention ve Flash Attention 2 kernel’ları ThunderKittens ile uygulandığında kod önemli ölçüde sadeleşiyor ve H100 üzerinde en fazla %30 performans artışı sağlıyor
- Based Linear Attention kernel’ı da ThunderKittens ile uygulanarak 215 TFLOPs performansa ulaşıldı (algoritmanın yapısı gereği recompute dahil edildiğinde 300 TFLOPs’un üzerine çıkıyor)
Felsefi açıdan değerlendirme
- ThunderKittens’ın iyi çalışmasının nedeni, her şeyi desteklemeye çalışmaması; bunun yerine basit ama GPU mimarisiyle çok iyi örtüşen tile tabanlı bir soyutlama sunmasıdır
- Geleneksel 32-bit word yerine 1024-bit vector register kullanmak da bir ilerleme, ancak asıl gereken paradigma değişimi 16x16 tile’ları register’ın temel birimi olarak görmektir
- Yapay zeka iş yüklerinin sonuçta çoğunlukla matris çarpımı, reduction ve reshape etrafında dönmesi nedeniyle tile tabanlı yaklaşım mantıklıdır; donanım açısından da systolic array’in ötesinde küçük matris çarpımlarını destekleyecek yönde evrim beklenebilir
- Daha da ileri giderek, yapay zeka algoritmalarını donanıma optimize edilmiş biçimde tasarlama yönünde düşünce yapısını değiştirmek gerekir. Örneğin RNN durum boyutunu SM içine sığacak kadar sınırlamak ve işlem yoğunluğunu da donanımın gerektirdiği seviyeye göre ayarlamak gibi bir yaklaşım gereklidir
GN⁺ görüşü
- ThunderKittens, CUDA’ya aşina geliştiriciler için cazip bir seçenek olabilir. Bunun nedeni, mevcut kernel kodunu büyük ölçüde değiştirmeden kolayca performans artışı elde edilebilmesidir
- Ancak yeni başlayanlar için giriş engeli hâlâ yüksek olabilir. İleride daha çeşitli örnek kodlar ve öğrenme materyalleriyle desteklenmesi gerekiyor gibi görünüyor
- ThunderKittens desteğinin yalnızca H100 ile sınırlı kalmayıp AMD gibi diğer GPU’lara da genişletilmesinin planlanması ilgi çekici. Bunun vendor bağımlılığını azaltmaya katkı sağlayabileceği düşünülüyor
- Nihayetinde en önemli nokta, yapay zeka modelinin/algoritmasının kendisini donanıma optimize edilmiş biçimde tasarlamaktır. Bunun için önce donanım özelliklerinin derinlemesine anlaşılması gerekir; ThunderKittens geliştiricilerin bu tür içgörüler edinmesine yardımcı olabilir
- Hazy Research’ün sürekli Ar-Ge çalışmaları ve açık kaynak katkılarının CUDA ekosisteminin canlanmasına büyük katkı sağlaması bekleniyor. Ancak uzun vadede daha yüksek soyutlama seviyesine sahip framework’lerin ortaya çıkmasına da ihtiyaç var gibi görünüyor
1 yorum
Hacker News yorumu
Yapay zeka donanımının gereksinimleri giderek daha net hale geliyor. GPU'lar aslında başka amaçlar için tasarlanmıştı, ancak iyi matris çarpımı donanımına sahip oldukları için yapay zekada kullanılıyorlar. "AI GPU" gerçek bir GPU'nun bazı işlevlerini dışarıda bırakabilir ve eğilim daha kısa sayı biçimlerine (16 bit, 8 bit, 2 bit, 1 bit kayan nokta) doğru gidiyor. Bu makale, 16x16 tile seven donanımın birçok avantaja sahip olduğunu düşündürüyor.
Yapay zekaya özel yardımcı işlemcilere (NPU) ihtiyaç var. Özellikle geliştiriciler, uzmanlar, oyuncular ve benzerleri için prosumer düzeyi masaüstü sistemlerde gerekli. GPU'lar kurumsal tarafta işe yarıyor, ancak kişisel bilişim açısından yapay zeka için kullanımları elverişsiz. Özellikle VRAM sınırlamaları ve Vulkan dışında standart, açık bir API'nin olmaması sorun.
Yapay zeka araştırmalarını ilerletmek için sinirbilim ve psikolojinin daha iyi incelenmesi gerekiyor. Ayrıca sinir ağlarının grafik topolojisiyle ilgili unsurlar da bağlantılı olabilir.
Bu, CUTLASS'ın kullanıcı dostu hale getirilmiş bir versiyonu mu?
ThunderKittens maskotunda kedi/Sony Aibo havası var. Yapay zekayla iyi üretilmiş gibi görünüyor.
Evrensel Temel Hesaplama (UBC), Evrensel Temel Gelir'in ikamesi olarak düşünülürse, bu oldukça distopik bir gelecek olur. Nvidia gibi tek bir şirketin tüm hesaplamayı ürettiğini hayal edin.