Yerel büyük dil modellerine (LLM) olan talep arttıkça, geliştiriciler PyTorch ve vLLM gibi standart kütüphanelerin ötesine geçerek son derece özelleştirilmiş çıkarım hatları oluşturuyor. Standart kütüphaneler mükemmel varsayılanlar sunsa da, taşınabilirlik ve kullanım kolaylığı için mikro-optimizasyonlardan genellikle vazgeçerler. Tüketici sınıfı GPU'lardan veya veri merkezi A100'lerinden her bir verimlilik birimini çıkarmak için CUDA çekirdek optimizasyonuna derinlemesine dalmanız gerekir. Bu yazı, dönüştürücü (transformer) çıkarım mimarileri için özel olarak tasarlanmış CUDA çekirdeklerini ince ayarlamak için kritik teknikleri incelemektedir.
Hafıza Hiyerarşisi Darboğazını Anlamak
LLM çıkarımındaki birincil darboğaz genellikle hesaplama ile sınırlı değildir; bellek ile sınırlıdır. Dikkat (attention) mekanizmaları ve beslemeli-ileri katmanlar, Yüksek Bant Genişlikli Bellek (HBM) ile GPU'nun yonga içi belleği arasında devasa miktarda veri hareketi gerektirir. İnce ayarın ilk adımı, bellek birleştirmeyi (memory coalescing) sağlamaktır. Bir warp içindeki bitişik iş parçacıkları, bitişik bellek adreslerine eriştiğinde donanım bu istekleri tek bir bellek işlemi olarak birleştirir. Veri yapılarının veya erişim desenlerinin hizalanmaması, performansı bir büyüklük mertebeye kadar düşürebilir.
Standart bir matris çarpımı çekirdeğini ele alalım. İş parçacığı dizinleri dikkatle eşleştirilmezse, maliyetli yayılma (broadcast) işlemleriyle karşılaşabilirsiniz. İş parçacığı bloğu boyutlarınızın warp boyutunun (genellikle 32) katları olduğundan ve global bellek erişimlerinin ardışık olduğundan emin olun.
Paylaşılan Belleğin Stratejik Kullanımı
Paylaşılan bellek, GPU yongası üzerinde bulunan ve global belleğe göre gecikme açısından birkaç mertebesi daha düşük olan yüksek hızlı SRAM'dir. Optimize edilmiş Flash Attention veya RoPE (Döndürücü Konumsal Gömme) uygulayanlar gibi özel LLM çekirdekleri için performans açısından paylaşılan bellek kullanımı tartışmasızdır.
Bununla birlikte, paylaşılan bellek kıt bir kaynaktır. Aşırı tahsis, performansı yok eden kayıt dökülmesine (register spilling) yol açar. Anahtar, döşeme (tiling) yapmaktır. Paralellik en üst düzeyde tutulurken verileri paylaşılan bellek sınırları içine sığdırmak için verileri döşemeniz gerekir. Aşağıda, döşeme tabanlı bir matris yüklemesi için paylaşılan belleğin nasıl bildirileceğine ve kullanılacağına dair basitleştirilmiş bir örnek bulunmaktadır:
// CUDA çekirdeği içinde
extern __shared__ float sdata[];
// Veriyi paylaşılan belleğe yükle
unsigned int tid = threadIdx.x;
unsigned int row = blockIdx.y * blockDim.y + threadIdx.y;
unsigned int col = blockIdx.x * blockDim.x + threadIdx.x;
if (row < M && col < N) {
sdata[tid] = globalMatrix[row * N + col];
}
__syncthreads(); // Tüm iş parçacıklarının verilerini yüklediğinden emin olun
// Paylaşılan belleği kullanarak hesaplama yap
float sum = 0.0f;
for (int i = 0; i < TILE_SIZE; ++i) {
sum += sdata[tid] * otherMatrix[i * N + col];
}
Bu kod parçacığında, __syncthreads() kritik öneme sahiptir. Bir bariyer olarak hareket eder ve herhangi bir iş parçacığı sdata içinden okumaya başlamadan önce bloktaki tüm iş parçacıklarının verileri yüklemeyi tamamladığını sağlar. Bu olmadan, yarış durumları (race conditions) hesaplamalarınızı bozacaktır.
İşgal ve Kayıt Baskısı
İşgal (Occupancy), çoklu işlemci başına aktif warp sayısının mümkün olan maksimum warp sayısına oranını ifade eder. Yüksek işgal, bir warp belleği beklerken planlayıcının başka bir warpa geçmesine izin vererek bellek gecikmesini gizlemeye yardımcı olur. Ancak, işgali artırmak bir denge oyunudur. İş parçacığı başına kullanılan her kayıt, bir çoklu işlemcide sığabilecek toplam warp sayısını azaltır.
Kayıt baskısını teşhis etmek için nvprof veya ncu (NVIDIA Nsight Compute) kullanın. Çekirdeğiniz, "kayıt sınırları" nedeniyle düşük işgal bildiriyorsa, kodunuzu yeniden yapılandırmanız gerekir. Yaygın stratejiler şunları içerir:
- Yerel değişkenleri
__shared__belleğe yükseltmek. - Anahtar alanı gerektiren dizilerin ve döngülerin boyutunu azaltmak.
- Kompilatörün işaretçi aritmetiğini optimize etmesine yardımcı olmak için
__restrict__işaretçilerini kullanmak.
Karma Hassasiyet için Tensor Çekirdeklerinden Yararlanma
Modern NVIDIA GPU'ları, karma hassasiyetli matris çarpımları (FP16, BF16, INT8) için tasarlanmış özel donanım birimleri olan Tensor Çekirdeklerine sahiptir. Hassasiyet kaybına genellikle tolerans gösteren LLM çıkarımı için, CUTLASS veya manuel __mma talimatları aracılığıyla Tensor Çekirdeklerini kullanacak şekilde matmul çekirdeklerinizi dönüştürmek, standart FP32 veya FP16 CUDA çekirdeklerine kıyasla 4x ila 10x hızlanma sağlayabilir.
Özel çekirdekler yazarken, veri düzeninizin NHWC (veya Tensor Çekirdekleri için uygun) olduğundan emin olun, NCHW değil ve matris boyutlarınızı belirli mimariye (Volta, Ampere veya Hopper) bağlı olarak 8 veya 16'nın katlarına hizalayın.
Sonuç
Yerel LLM çıkarımı için CUDA çekirdeklerini ince ayarlamak, korkaklar için değildir. GPU mimarisi, hafıza hiyerarşileri ve derleyici optimizasyonları hakkında derin bir anlayış gerektirir. Ancak, ödüller önemli ölçüde büyüktür. Bellek birleştirmeyi, paylaşılan bellek döşemesini, işgal yönetimini ve Tensor Çekirdeği kullanımını ustalaşarak, genel amaçlı çerçevelerin performansını aşan çıkarım hatları oluşturabilirsiniz. Darboğazınızı profillemeye başlayın, bu teknikleri yinelemeli olarak uygulayın ve Nsight Compute gibi araçlarla etkimi her zaman ölçün. Verimli yerel AI'nın geleceği, GPU dilini konuşabilenlerin elindedir.