NVIDIA Blackwell ve CUDA 12.9 ile Aileye Özel Mimari Özellikler Tanıtıldı

CUDA platformunun design süreçlerinden biri, NVIDIA GPU’larının kod uyumluluğuna destek sağlamaktı. Bu tasarım, yeni GPU’ların, önceki GPU’lar için yazılmış programları herhangi bir değişiklik yapmadan çalıştırabilmesi anlamına geliyor. Bu, CUDA’nın iki temel özelliği ile sağlanıyor:

• NVIDIA Paralel Thread Execution (PTX) sanal talimat seti mimarisi (ISA)
• NVIDIA sürücüsü, PTX kodunu çalışma zamanında just-in-time (JIT) derliyor.

PTX, NVIDIA GPU’larını hedefleyen sanal ISA’dır. Bunu, belirli bir fiziksel çip donanım mimarisiyle sınırlı olmayan bir montaj kodu gibi düşünebilirsiniz; gelecekteki GPU mimarileriyle uyumlu olabilmesi için genel bir tasarıma sahiptir.

NVIDIA, CUDA platformunu genel amaçlı program yazımını mümkün kılmak için oluşturduğundan beri, PTX bu süreçte önemli bir rol oynamaktadır. Önceki GPU’lar için oluşturulan PTX kodu, mevcut sürücüler tarafından JIT derlenerek günümüzdeki GPU’larda çalıştırılabilir.

Şimdi basit bir örnek düşünelim. Bu örnek, GPU adını ve hesaplama kabiliyetini ekrana yazdıran, ayrıca GPU çekirdek içinden “Merhaba” yazan bir kod parçasıdır.

#include <stdio.h>
#include <iostream>

__global__ void printfKernel()
{
    printf(">>>>>>>>>>>>>>>>>>>>n" );
    printf("THREAD %d' den MERHABAn", threadIdx.x );
    printf(">>>>>>>>>>>>>>>>>>>>n" );
}

int main(int argc, char** argv)
{
    // Cihaz özelliklerini kontrol et ve göster
    cudaDeviceProp deviceProp;
    cudaGetDeviceProperties(&deviceProp, 0);

    std::cout << deviceProp.name << std::endl;
    std::cout << "Hesaplama Kabiliyeti: " << deviceProp.major
        << "." << deviceProp.minor << std::endl;
    printfKernel<<<1,1>>>();
    cudaDeviceSynchronize();

    std::cout << "Program Sonu" << std::endl;
    return 0;
}

Bu kodu CUDA 12.8 ile derlediğimizde ve NVIDIA RTX 4000 Ada sistemimizde çalıştırdığımızda şu sonuca ulaşıyoruz:

$ nvcc -o x.device_info device_info.cu

$ ./x.device_info
NVIDIA RTX 4000 Ada Nesli
Hesaplama Kabiliyeti: 8.9
>>>>>>>>>>>>>>>>>>>>
THREAD 0' den MERHABA
>>>>>>>>>>>>>>>>>>>>
Program Sonu

Compiler bayraklarını belirtmediğimiz için NVCC, bu derleme sürümünde desteklenen en düşük PTX hedefine göre çalışır. Hangi PTX mimarisi ve hangi CUDA ikili (cubin) mimarisinin kodunuzda olduğunu cuobjdump kullanarak görebiliriz (çıktı kısaltılmıştır):

$ cuobjdump x.device_info

Fatbin elf kodu:
================
mimari = sm_52

>>>> kesildi <<<

Fatbin ptx kodu:
================
mimari = sm_52

>>>> kesildi <<<

Ayrıca ELF ve PTX görünüyorsa, bu, hem cubin hem de PTX’nin nesne dosyasında gömülü olduğu anlamına gelir. Mimarisi sm_52‘dir; bu hesaplama kabiliyeti (CC) 5.2’dir. CC, X.Y şeklinde temsil edilir; burada X büyük revizyon numarasını ve Y küçük revizyon numarasını ifade eder.

JIT Derleme ile Uyumlu Çalışma

Örneğimize geri dönersek, GPU 8.9 CC’sindeyken bu kod nasıl oluyor da çalışabiliyor? İşte burada JIT derlemenin önemi devreye giriyor. CUDA sürücüsü, PTX’yi CC 8.9 GPU’sunda çalıştırmak için derler. Kodunuz, PTX’nin, GPU’nuzun mimarisinin eşit veya önceki bir mimariden üretilmiş olması durumunda sağlıklı bir şekilde çalışacaktır.

Bunu doğrulamak için bazı derleyici bayrakları ekeleyelim. -gencode arch=compute_75,code=compute_75 argümanını ekleyerek NVCC’ye uygulamanız için compute_75 (hesaplama kabiliyeti 7.5) ile PTX oluşturmak istediğinizi belirtiriz. Bunu, cuobjdump ile de doğrulayabilirsiniz. NVCC’nin nasıl PTX oluşturduğu ve ikili kod ürettiği hakkında daha fazla bilgi için, Understanding PTX, the Assembly Language of CUDA GPU Computing‘deki Şekil 1’e bakabilirsiniz.

Bu yapı çalışıyor.

$ nvcc -gencode arch=compute_75,code=compute_75 -o x.device_info device_info.cu

$ ./x.device_info
NVIDIA RTX 4000 Ada Nesli
Hesaplama Kabiliyeti: 8.9
>>>>>>>>>>>>>>>>>>>>
THREAD 0' den MERHABA
>>>>>>>>>>>>>>>>>>>>
Program Sonu

Şimdi code=compute_75‘i code=sm_75 olarak değiştirirsek, bu durumda NVCC’nin PTX’yi daha önce belirtildiği gibi bırakması gerektiği yerine sm_75 için bir cubin’e derlemesine ve bu cubin’i yürütülebilir dosyaya yerleştirmesine neden oluyoruz. Bunu cuobjdump ile kontrol edelim. Çıktı şöyle olsun:

$ nvcc -gencode arch=compute_75,code=sm_75 -o x.device_info device_info.cu

$ ./x.device_info
NVIDIA RTX 4000 Ada Nesli
Hesaplama Kabiliyeti: 8.9
Program Sonu

Dikkatli bakarsanız, “THREAD 0' den MERHABA” ifadesinin basılmadığını göreceksiniz. Tüm hata kontrolü kodlarını temiz tutmak için çıkardık.

Hata Kontrolü ve Uyum Kuralları

Eğer hata kontrolü ekleseydik, ki gerçek kodda her zaman yapmalısınız, GPU çekirdeğinin çalıştırılmadığını ve aldığımız hata mesajının “Cihazda çalıştırmak için kernel görüntüsü mevcut değil.” olduğunu görecektiniz. Bu, uygulamada, bu CC 8.9 cihazla uyumlu kernel bulunmadığı anlamına gelir ve bu nedenle kernel çalıştırılmamıştır.

CUDA uyumlu GPU’lar için (Tegra hariç) geçerli olan kural şunlardır:

• PTX uyumluluğu: Belirli bir CC’ye sahip PTX ile oluşturulan herhangi bir kod, o CC’ye ve sonraki CC’ye sahip GPU’larda çalışır.
• Cubin uyumluluğu: Belirli bir CC’ye sahip bir cubin ile oluşturulan herhangi bir kod, o CC’ye ve daha sonra gelen aynı büyük kabiliyete sahip GPU’larda çalışır. Örneğin, CC 8.6’ya sahip bir GPU, CC 8.0 için oluşturulan bir cubini çalıştırabilir. Tam tersi doğru değildir. CC 8.6 için bir cubin derlerseniz, yalnızca CC 8.6 ve üzeri ile çalışır, CC 8.0 ile çalışmaz.

NVIDIA Hopper ile Gelen Özellikler

NVIDIA Hopper mimarisi (CC 9.0) ile birlikte, yalnızca belirli bir hedef mimaride var olma garantisi bulunan mimari özellik setleri adında yeni ve özel bir özellik seti tanıtıldı. Bu özellikler çoğunlukla Tensor Core’ların kullanımı ile ilgilidir.

Bu özellikleri kullanabilmek için, PTX veya cubin kodunu uygulamanıza gömmek zorundasınız; bunu compute_90a bayrağı ile PTX veya sm_90a bayrağı ile cubin ile yapabilirsiniz. Mimari spesifik hedef oluşturduğunuzda, bu kod gelecekteki GPU mimarileri ile ileriye dönük uyumlu değildir.

Örneğin, CUDA çekirdeklerinizi şu NVCC ile derleyebilirsiniz:

$ nvcc -gencode arch=compute_90a,code=sm_90a -c kernel.cu

Bu durumda, kodunuz yalnızca CC 9.0 cihazlarında yüklenir ve çalıştırılır; ileriye dönük uyum yoktur.

Aile-Bazlı Özellik Seti – NVIDIA Blackwell

NVIDIA Blackwell mimarisi ve CUDA 12.9 ile birlikte, yeni bir özellik kategorisi olarak aile spesifik özellikler tanıtılmıştır.

Aile spesifik özellikler, farklı küçük hesaplama kabiliyetlerine sahip cihazlar tarafından desteklenmesine rağmen birden fazla küçük hesaplama kabiliyetine sahip cihazlarla uyumludur. Tüm aile içindeki cihazlar, aynı büyük hesaplama kabiliyeti sürümünde ortak özellikler taşır. Aile spesifik özellikler belirli bir ailedeki cihazlarda mevcut olma garantisi taşır; bu cihazlar, daha yüksek küçük hesaplama kabiliyetlerine sahip daha sonraki GPU’larla birlikte çalışabilir.

Aile spesifik derleyici hedefi, mimari spesifik hedefe benzer; ancak a yerine f eki kullanılarak tanımlanır.

Şu şekilde bir NVCC ile CUDA çekirdeklerinizi derleyebilirsiniz:

$ nvcc -gencode arch=compute_100f,code=sm_100 -c kernel.cu

Bu durumda, kodunuz yalnızca 10.x hesaplama kabiliyetine sahip cihazlarda çalışır. Gelecekte 10.x hesaplama kabiliyetine sahip yeni GPU’lar tanıtıldığında, kodunuz bu GPU’larda da uyumlu olacaktır çünkü bunlar sm_100f ailesine dahildir.

Geliştirici Yönergeleri

Artık mimari ve aile spesifik kod hedeflerinin nasıl inşa edildiğini açıkladığımız için, uygulamalarınızı inşa ederken bazı öneriler sunmak istiyoruz.

Genel olarak, kodunuzu olabildiği kadar çok mimaride çalışacak şekilde inşa etmelisiniz. Mimari veya aile spesifik özellikler kullanmıyorsanız, uygulamanızı her zamanki gibi inşa etmeniz gerekir. Eğer bu özellikleri kullanıyorsanız, a veya f bayraklarını kullanarak kodunuzu derlemeniz gerektiğini bilin.

Portatiflik istiyorsanız, kodunuzda uygun guardlar eklemsiniz. Aşağıdaki makroları kullanarak, kullandığınız aile veya mimari spesifik kodlara bağlı olarak kod yollarını kontrol edebilirsiniz:

• __CUDA_ARCH_FAMILY_SPECIFIC__
• __CUDA_ARCH_SPECIFIC__

Bu makrolar, __CUDA_ARCH__ ile benzer şekilde tanımlıdır. Daha fazla bilgi için CUDA Programlama Kılavuzu‘na başvurabilirsiniz.

Eğer uygulamanız CUTLASS gibi bir kütüphane kullanıyorsa veya CUTLASS’ın bir kütüphanesini kullanıyorsanız, yalnızca CC 9.0 veya üzerindeki cihazlarda çalışmakta iseniz, mimari spesifik hedeflerle uygulamanızı inşa etmelisiniz. CUTLASS, yüksek performans için özel olarak tasarlanmıştır ve mimari spesifik özelliklerin maksimum performansla çalışacağı özel kod yollarını içerir. Bu kütüphaneler, diğer mimarilerle tam uyum sağlamak için zaten içlerinde geri dönüş yolları bulundurmaktadır.

Uygulamada Uygulama

Bu mimari ve aile spesifik hedeflerin nasıl kullanıldığını tartıştıktan sonra, her şeyin bir araya geldiği noktaya gelelim.

Genel Durum

Öncelikle, kodunuzun mimari veya aile spesifik özellikleri kullanıp kullanmadığını belirlemelisiniz. Eğer bu özellikleri kullanmıyorsanız, kodunuzu her zamanki gibi derlemeye devam edebilirsiniz.

En iyi performansı ve gelecekteki uyumluluğu sağlamak için, çalışacağınız her mimari için ikili kod inşa etmek tipik kılavuzluktur. Bu size en iyi performansı sağlayacaktır. En yeni mimariye PTX gömerek, gelecekte büyük uyumluluk kendine sağlamış olursunuz. Örneğin, kodunuzun CC 8.0, 9.0 ve 10.0 cihazlarında çalışacağını biliyorsanız, CUDA çekirdeklerinizi şu şekilde derleyebilirsiniz:

$ nvcc -gencode arch=compute_80,code=sm_80 
       -gencode arch=compute_90,code=sm_90
       -gencode arch=compute_100,code=sm_100 
       -gencode arch=compute_100,code=compute_100 -c kernels.cu

Aile-Spesifik Özellikler

Özel bir optimizasyon yapmak üzere, özellikle farklı mimariler arasında taşınabilirliğini istiyorsanız, bu özellikleri kullanmaya karar vermelisiniz. Öncelikle, bu özelliklerin aile spesifik özellik setinde yer alıp almadığını belirlemelisiniz. Eğer öyleyse, hedeflerinizi f eki ile derleyebilirsiniz, böylece o aile içindeki GPU’lar için uyumluluğunuz olur. Ancak, aile dışındaki GPU’lara taşınabilirlik istiyorsanız, aile-spesifik özellikleri kullanan herhangi bir kod için rezerv pistleri eklemelisiniz.

Tipik olarak, aile spesifik kodu uygulamanız boyunca bu makroları kullanarak koruyabilirsiniz. Örneğin, mimari spesifik özellikleri kullanmayacak şekilde optimize edilen kodunuzun düzeni şu şekilde olabilir :

$ nvcc -gencode arch=compute_100f,code=sm_100 -c kernels.cu

Bu, kodunuza CC 8.0, 8.9 ve 10.0’da çalışma yeteneği katacak ve gömülü PTX ile gelecekteki cihazlarda da çalışmanıza olanak tanıyacaktır.

Diğer bir senaryo, yalnızca bu ailedeki cihazlara özgü özellikleri kullanan bir uygulama tasarlamanız gerektiğidir. Bu durumda, 100f ailesindeki özelliklerden yararlandığınızı varsayalım, o zaman yukarıdaki gibi bir kod düzenlemesi olmalıdır:

$ nvcc  -gencode arch=compute_100f,code=sm_100 -c kernels.cu

Bu, yalnızca o aile içindeki cihazlarda taşınabilirlik sağlayacaktır.

Mimari-Özgü Özellikler

Eğer kodunuzun aile spesifik özellikler dışındaki özellikler gerektirdiğini düşünüyorsanız ve bu özellikleri kullanıyorsanız a bayrağı ile oluşturmalısınız.

Aile spesifik durumda olduğu gibi, kod taşınabilirliği sağlamanıza yardımcı olacak bir dizi kod yolu oluşturmalısınız. Aşağıdaki örnekte, bu yapıyı şu şekilde görebilirsiniz:

$ nvcc -gencode arch=compute_80,code=sm_80 
       -gencode arch=compute_90,code=sm_90
       -gencode arch=compute_100a,code=sm_100a 
       -gencode arch=compute_100,code=compute_100 -c kernels.cu

Bu, CC 8.0, CC 9.0, CC 10.0 ve sonrasında çalışmanıza olanak tanır.

Eğer belirli bir mimaride optimize edilmiş özel bir uygulama tasarlıyorsanız, bu durumda şu şekilde inşa edebilirsiniz:

$ nvcc  -gencode arch=compute_100a,code=sm_100a -c kernels.cu

Bu durumda uygulamanız, yalnızca CC 10.0 üzerinde çalışacak ve diğer GPU’larla uyumsuz olacaktır.

Sonuçlandırma

Özetle, kodunuzu düşünürken şu basit akışı izleyebilirsiniz:

• Doğrudan PTX mi yazıyorsunuz, yoksa CUTLASS gibi bir kütüphane mı çağırıyorsunuz? Eğer hayırsa, f veya a bayraklarını dahil etmelisiniz.
• Eğer PTX kodu yazıyorsanız? Eğer PTX kodunuzun kullanıldığı bir kütüphane varsa, bu kütüphanenin belgelerine başvurun ve en iyi mimariye göre nasıl derleyeceğinizi belirleyin.

Bu yazıda, kodları mimari ve aileye özgü özellikler ile kullanma konusunda bilgi verdik. Bu özellikleri kullanmak için, PTX yazıyor olmanız veya bu tür bir kütüphaneyi çağırmanız gerekmektedir.

Teşekkürler

Aşağıdaki NVIDIA katkılarına teşekkürler: Cody Addison, Vyas Venkataraman, Rob Armstrong, Girish Bharambe ve Mridula Prakash.