C++ ve CUDA kullanarak sıfırdan bir LLM çıkarım motoru oluşturma

• C++ ve CUDA kullanarak hiçbir kütüphane olmadan bir LLM çıkarım motorunun nasıl kurulacağı
• Bu sayede LLM çıkarımının tüm katmanını anlayabilir ve farklı optimizasyonların çıkarım hızını nasıl etkilediğini doğrudan görebilirsiniz
• Hedef: tek CPU + GPU sunucusunda tek batch ile hızlı çıkarım yapabilecek şekilde modeli uygulamak ve llama.cpp'den daha hızlı token işleme hızına ulaşmak

1. LLM mimarisi ve çıkarıma genel bakış

• Başlıca LLM'lerin çoğu, ardışık transformer blokları kullanan aynı mimariyi izler.
• Model yükleme, özelleştirilebilir transformer blok sınıfları tanımlayıp bunları bir dizi halinde birleştirerek safetensors ağırlıklarıyla başlatmaktır.
• Çıkarım çoğunlukla tek batch ile yapılır ve yürütmenin büyük kısmını "decode aşaması" oluşturur.

1.1 Çıkarıma genel bakış

• Çıkarım, verilen prompt token'larını modele aktararak KV cache'i dolduran prefill aşaması ile modeli tekrar tekrar çalıştırarak token üreten decode aşamasına ayrılır
  • Prefill aşaması: prompt token'larını işler ve KV cache'i başlatır
  • Decode aşaması: her seferinde bir token üretir
• KV cache: önceki key/value çiftlerini saklayarak geçmiş bağlamla attention hesaplamasını hızlandırır
• Modelin forward pass'i, embedding tablosunu kullanarak token ID'lerini embedding vektörlerine eşler ve durumları transformer blok dizisi boyunca dönüştürür

1.2 Darboğazlar ve benchmark

• Darboğaz: modern donanımda sınırlayıcı unsur bellek bant genişliğidir
  • Model çıkarımında her token'ı üretmek için tüm modeli okumak gerekir ve bellek bant genişliği hesaplamadan daha büyük bir kısıt haline gelir
• Model quantization, çıkarım hızını iyileştirmede etkilidir
• Teorik maksimum token işleme kapasitesi donanıma göre değişir ve gerçek performans çeşitli inference engine'ler üzerinden doğrulanabilir
• Teorik hız sınırı:
  • AMD EPYC 7702P: en fazla 13.6 tok/s (FP16 bazında)
  • RTX 4090: en fazla 67.1 tok/s (FP16 bazında)
• Benchmark:
  • llama.cpp: CPU 8.7 tok/s, GPU 61 tok/s
  • calm: GPU 66 tok/s

2. CPU tabanlı çıkarım

• CPU üzerindeki ilk uygulama tek iş parçacıklıdır ve yalnızca FP32 ağırlıkları destekler
• Multithreading ile kodu paralelleştirmeye başlayabilir, SIMD kullanarak da performansı artırabilirsiniz

2.1 Multithreading

• OpenMP kullanarak matris-vektör çarpımı (matmul) ve multi-head attention paralelleştirilip performans iyileştirildi
• Optimizasyon sonucu: hız 0.6 tok/s → 4.4 tok/s

2.2 Ağırlık quantization ve SIMD optimizasyonu

• Quantization: FP32 ağırlıkları FP16'ya quantize ederek bellek kullanımı yarıya indirildi ve performans artırıldı
• SIMD: AVX2 kullanılarak 8 adet FP32 değeri aynı anda işlenecek şekilde optimize edildi
• Sonuç: 8.4 tok/s

3. GPU tabanlı çıkarım

• Model FP16'ya quantize edilerek RTX 4090'a yüklenebilir ve GPU çıkarımı uygulanmaya başlanabilir
• CUDA ile C++ fonksiyonları (kernel'ler) GPU üzerinde paralel çalıştırılabilir

3.1 CUDA'ya basit port etme

• CPU işlemleri bire bir CUDA kernel'lerine dönüştürülerek GPU backend'i uygulanabilir
• CUDA kernel'leri asenkron çalışır, ancak aynı stream içinde sıralı yürütülür
• Sorun: iş parçacığı verimsizliği nedeniyle GPU kaynakları yeterince kullanılamıyor → 2.9 tok/s ile yavaş

3.2 Daha iyi matris çarpımı (matmul)

• Matris çarpımı CPU tarafında çalışma süresinin büyük bölümünü kaplar ve OpenMP ile optimize edilebilir
• GPU'da blok başına 1 satır işlenecek şekilde düzenlenerek iş parçacığı kullanım oranı artırılabilir
• Optimizasyon yöntemi:
  1. Bir blok bir satırı işler, blok içindeki iş parçacıkları hesaplama için birlikte çalışır
  2. Warp reduction uygulanır
• Sonuç: hız 51.7 tok/s'ye çıktı

3.3 Kernel fusion ve ek optimizasyonlar

• Kernel'leri birleştirerek performans artırılabilir
  • Kernel fusion: ardışık işlemleri tek bir kernel altında birleştirerek bellek erişimi ve hesaplama süresini en aza indirir
• Bellek erişim desenlerini optimize etme ve alan yeniden kullanımı ile 56.1 tok/s elde edildi

3.4 Attention optimizasyonu ve uzun bağlam işleme

• Sorun: uzun bağlamlarda attention kernel'i performans darboğazı haline geliyor
• Çözüm:
  1. Bellek erişimi optimizasyonu: ardışık bellek blokları okunacak şekilde yeniden tasarlandı
  2. Kayıp kayan nokta değerleri sorununu çözmek için atomicAdd yerine paylaşımlı bellek kullanıldı
• Optimizasyon sonucu:
  • Kısa bağlam: 63.8 tok/s (llama.cpp'nin 61.0 tok/s değerinden daha hızlı)
  • Uzun bağlam: 58.8 tok/s

3.5 KV cache quantization ve derleyici optimizasyon sorunları

• KV cache'in FP16'ya quantize edilmesi performans düşüşüne yol açtı (derleyici optimizasyonu yetersizliği)
• Çözüm: döngüler elle unroll edildi ve bellek prefetching uygulandı
• Sonuç: FP32'ye kıyasla yaklaşık 2 kat hız artışı ve uzun bağlam performansında 58.8 tok/s korundu

4. Gelecekteki iyileştirme yönleri

• Prompt prefill optimizasyonu: birden çok token'ı aynı anda işleyerek ilk token üretim süresini azaltma
• Attention kernel fusion: FlashAttention benzeri optimizasyon tekniklerini uygulama
• Daha yüksek quantization: FP8, INT8, INT4 uygulama ve aktivasyon/cache quantization
• Kernel optimizasyonu: bellek bant genişliği ve hesaplama verimliliğini en üst düzeye çıkaran ileri tekniklerin eklenmesi
• Kütüphane kullanımı: cuDNN, cuBLAS gibi kütüphanelerden yararlanarak optimizasyon süresini kısaltma

Sonuç özeti:

• CPU ve GPU'da çeşitli optimizasyonlarla 63.8 tok/s hızına ulaşıldı
• llama.cpp ve calm'e yakın ya da daha iyi performans elde edildi
• Hiçbir kütüphane olmadan yalnızca C++ ve CUDA ile yüksek performanslı bir LLM çıkarım motoru uygulandı

[GN⁺]

Hacker News görüşleri

• Yazar, blog yazısının ilgi görmesinden memnun ve geri bildirim duymak istiyor
• Bir okur, yazının harika olduğunu söyleyip ne kadar sürede yazıldığını merak ediyor
  • GPGPU alanında çalışan biri olarak benzer bir yazı yazmak istiyor, ancak ne kadar zaman alacağının belirsizliği yüzünden tereddüt ediyor
• Başka bir okur, kodun tensor core'ları veya wgmma komutlarını kullanmadığını düşünüyor
  • Bu tür programlamanın aynı anda birçok işi yürütmeyi gerektirdiği için zor olduğunu açıklıyor
  • Bant genişliği kısıtı nedeniyle ek hesaplamaya gerek olmayabileceğini belirtiyor
  • Blogdaki kodun başka hızlandırıcılara taşınırken muhtemelen iyi çalışacağını değerlendiriyor
  • wgmma kullanmanın Nvidia'nın nesilleri arasında taşınabilirliği düşürebileceğinden endişe ediyor
• Bir başka okur, buna benzer Python materyalleri arıyor ve bunları ekibiyle paylaşmak istiyor
  • Performanstan ziyade kavramsal olarak eksiksiz ve öğretici tarzda, kısa materyaller istiyor
• Bir kullanıcı, kendi Mistral sürümünü ve saniye başına token performansını karşılaştırmak istiyor
  • README'deki kuantizasyon bölümüne bakılması öneriliyor
• __shfl_down için günümüzde warp senkronizasyon sorunları nedeniyle tavsiye edilmediği yönünde bir görüş var