Büyük dil modellerini (LLM’ler) üretimde kullanırken, yüksek hız gereksinimlerini karşılamak için dikkatli olmak gerekir. Bu süreçteki kritik aşamalardan biri LLM kod çözümü‘dür; burada bir sonraki token’a zaman vermek, optimize edilmesi gereken önemli bir ölçüttür. Çalışma süresi gecikmelerini en aza indirmek için, genellikle tensor paralelliği uygulanarak MLP ve projeksiyon GEMM katmanları üzerinde çoklu GPU’lara dağıtım yapılır.
Kod çözüm aşamasında, iletilen verilerin boyutları ve hesaplama ihtiyaçları görece küçüktür. Derleyiciden kaynaklanan statik gecikmeler, örneğin kernel çağrısı veya iletişim yapılandırması, iletişim ve hesaplamalardaki süreleri baskın hale getirir. Bu makalede, toplam kod çözüm gecikmesine katkıda bulunan bu tür gecikmeleri en aza indirmek için geliştirdiğimiz bazı teknikleri paylaşacağız.
8 yollu tensor paralelGemma2 LLM modeli üzerinde, 8 adet NVIDIA H100 Tensor Core GPU ile tek bir düğümde kod çözüm aşamasını çalıştırdık. Tensor paralel katmanlarda tüm azaltma işlemi, toplam kod çözüm gecikmelerinde önemli bir darboğaz oluşturdu. Bu kolektif işlemi, yaklaşık 30 KB/s büyüklüğündeki mesajlarla çalışarak toplam kod çözüm gecikmesinin yaklaşık %23’ünü oluşturdu. İletim ve hesaplama kernel’leri arasındaki veri bağımlılıkları nedeniyle, iletim süreci hesaplama ile üst üste gelemiyordu.
Geleneksel Çözüm ve Tasarlanan Yenilikçi Yöntem
Geleneksel tüm azaltma yöntemleri, N GPU etrafında halka şeklinde veri değişimi yapmayı içeren halka algoritması kullanıyor. Halka algoritması, orta ve büyük boyutlu mesajlar ile bant genişliği açısından optimal olsa da, küçük mesaj boyutları için veri değişimleri ve tekrarlanan inter-GPU senkronizasyon engelleri önemli gecikmelere yol açmaktadır (yaklaşık 2 kat). Bu nedenle, biz bir özel tek aşamalı tüm azaltma algoritması geliştirdik. Bu algoritma, her bir cihazın, düğüm üzerindeki akışlarını toplamak ve aynı anda azaltma yapmasını sağlar; diğer bir deyişle, tüm toplama işlemi ve ardından lokal bir azaltma işlemi gerçekleştirir.
Bu yöntem, veri değişimlerini arttırsa da, NVLink iletişimi sayesinde değişimler eş zamanlı gerçekleştirildiğinden, toplam iletişim gecikmesi azaltılmaktadır. Ayrıca, bir cudaDeviceEnablePeerAccess kullanarak ek bellek kopyalama yükünü de ortadan kaldırdık; bu sayede, komşu GPU’larda kaydedilmiş tamponlara doğrudan erişebiliyoruz. Bu uygulama, bir süreçte çalışan tek düğüm çoklu GPU yapılandırmasında oldukça yararlıdır, çünkü paylaşılan bir CUDA bağlamı sayesinde, akış birimleri arasındaki cihaz belleği adreslerine erişim kolaylaşmaktadır.
Kod ve Uygulama Entegrasyonu
Aşağıda bir örnekle kodumuzun uygulanış şekli gösterilmektedir:
// Fused Tek aşamalı Tüm Azaltma + Kök Ortalama Normalizasyon çekirdek fonksiyonu
__global__ void OneShotARNormKernel(std::vector<T*> peer_comm_buffer_ptrs, T* sum_vec, T* weight_buffer, float eps, int hidden_size)
{
for (int ii = 0; ii < NUM_RANKS; ++ii)
{
// Tek aşamalı Tüm Azaltma
sum_vec = add(sum_vec, peer_comm_buffer_ptrs[ii][thread_offset]);
}
// Normalizasyonun Infüzyonu
squares = compute_square(sum_vec);
summed_squares = block_reduce_sum(squares);
float denom = __fsqrt_rn(__fdividef(summed_squares, hidden_size) + eps);
// Gerekirse her öğe için öğrenilebilir parametreyi yükle
if learnable_affine
{
weight_vec = weight_buffer + thread_offset;
}
sum_vec = rms_norm(denom, sum_vec, weight_vec);
}
Yukarıdaki kod örneği, mahalle iletişim tamponlarını GPU’dan okuma işlemini gerçekleştiren ve toplam azaltma ve normalizasyonu tek bir işlemde birleştiren çekirdek fonksiyonunu göstermektedir. Bu çekirdek, JAX uygulamalarında bir özel çağrı olarak entegre edilmiştir.
Verimlilik ve Gecikme İyileştirmeleri
Bu özelleştirilmiş çekirdek, komşu katman normalizasyonu ve noktasal toplama işlemleri ile birleştirilmiştir. Böylelikle kernel çağrı gecikmeleri ve bellek hareketleri en aza indirilmiştir. Sağladığımız bu birliktelik sayesinde, bağımsız tüm azaltma çekirdeğine göre yaklaşık 3 kat kernel süresi hız artışı elde ettik ve kod çözüm süresinde %27’lik bir iyileşme sağladık. Özel çekirdek ile birlikte, modeldeki diğer hesaplama çekirdekleri de tek bir CUDA Grafiği olarak gruplanmış ve başlatılmıştır; bu da ek %5’lik bir iyileşme sağlamıştır.
Sonuç olarak, küçük mesaj boyutları için iletişim gecikmelerini azaltmak, özellikle hesaplama-iletim üst üste gelmediği durumlarda önemlidir. Daha yüksek verimlilik için özelleştirilmiş uygulama çekirdekleri, iletişim blokları ve hesaplamaların birlikte başlatılmasına olanak sağlar. Daha da önemlisi, Mosaic-GPU DSL, veri iletişimi ve toplama işlemlerini bir araya getiren dağıtım desenlerini ifade etme imkanı sunar. Bu tür uygulamalarda geliştirilen fırlatıcı ve kıyaslayıcı yapılar, verimsiz iletişimi minimize ederek önemli kazançlar sağlayabilir.
