“CUDA GPU Hesaplamanın Assembly Dili Olarak PTX’in Anlaşılması”

Paralel iplik yürütme (PTX), CUDA’nın başlangıcından beri kullanılan sanal makine talimat seti mimarisidir. PTX, NVIDIA’nın CUDA GPU hesaplama platformunun assembly dilini temsil eder.

Bu yazıda, PTX’in ne anlama geldiğini, ne için kullanıldığını açıklayacağız ve CUDA uygulamalarınızdan en iyi şekilde yararlanabilmeniz için bilmeniz gerekenleri ele alacağız. İlk olarak, CUDA’nın kodu nasıl oluşturduğuna, depoladığına ve yüklediğine bakacağız. Ardından PTX’in ileriye dönük uyumluluğu nasıl sağladığı ve CUDA’ya hedeflemesi için alan spesifik ve diğer programlama dillerinin nasıl kullanılabileceğini göstereceğiz.

Talimat Seti Mimarisi

Bir talimat seti mimarisi (ISA), bir işlemcinin hangi talimatları yürüttüğünü, bunların formatını, bu talimatların davranışlarını ve ikili kodlamalarını belirten bir spesifikasyondur. Her işlemcinin bir ISA’sı vardır. Örneğin, x86_64 bir CPU ISA’sıdır. ARM64 başka bir örnektir. Bir GPU’nun da bir ISA’sı vardır. NVIDIA GPU’larının ISA’sı farklı nesil veya farklı ürün serilerinde değişiklik gösterebilir.

Sanal makine ISA’sı, bir sanal işlemcinin desteklediği talimatlar, formatlar ve davranışlar için bir spesifikasyondur. Bu, aslında üretilmeyen bir işlemci için yalnızca bir ISA olduğu anlamına gelir. Sanal makine ISAları, talimatların çalıştırılması için bir fiziksel işlemci üzerinde yalnızca gerekli olan ikili kodlamayı belirlemeyebilir.

PTX’in CUDA Platformundaki Rolü

PTX’in CUDA platformundaki yerine iyi bir örnek vermek için, basit bir CUDA dosyasının nasıl derlendiğini göstereceğiz. Bu kaynak dosya, klasik bir örnek olan iki vektörü paralel olarak toplayan tek bir çekirdek içermektedir ve uygulama ana fonksiyonunun iskeletini ve bir yardımcı fonksiyonu barındırmaktadır.

__global__ void vecAdd(float* a, float* b, float* c, int n)
{
    int index = threadIdx.x + blockIdx.x * blockDim.x;
    if (index 

Bu dosya NVIDIA’nın CUDA derleyici aracı NVCC ile derlendiğinde, kaynak kodu GPU için ve CPU için ayrı ayrı kodlara bölünür. GPU kodu GPU derleyicisine ve CPU kodu host derleyicisine gönderilir. Host derleyici NVCC’nin bir parçası değildir; NVCC, komut satırında verilen veya sistemdeki varsayılan derleyici olan bir host derleyiciye çağrıda bulunur.

GPU’da çalışacak fonksiyonlar ve çekirdekler için GPU derleyicisi CUDA platformunun assembly dilini olan PTX’i oluşturur. Bu assembly (PTX) daha sonra исполнитель (assembler) olan ptxas aracılığıyla GPU için yürütülebilir ikili kod oluşturur. GPU ikilisi CUDA binary’nin kısaltması olan *cubin* olarak adlandırılır.

GPU kodunun derlenmesi iki aşamadan oluşur: ilk olarak, yüksek seviyeli dil kodu (C++) PTX’e derlenir. Daha sonra PTX, *cubin*’e derlenir. İkili çıktı istendiğinde ptxas aracının çağrılması NVCC tarafından otomatik olarak yapılır.

Bu derleme yolu, clang gibi popüler derleyicilerin işleyişine benzerdir. clang ilk olarak kodu bir sanal makine ISA’sı olan **LLVM IR** (Low Level Virtual Machine, Ara Düzlem Temsili) şeklinde derler. Daha sonra, ikinci bir aşama veya arka uç derleyici, **LLVM** olarak adlandırılan, sanal makine temsilini belirli bir işlemci için yürütülebilir koda çevirir. Bu yapının avantajı, bir programın LLVM IR’inin, LLVM arka uç derleyicisi tarafından desteklenen herhangi bir donanım mimarisine derlenebilmesidir.

PTX, bir programın sanal makine temsilinin, çok çeşitli NVIDIA GPU’larına derlenebilmesi açısından LLVM IR’a benzer. Önemli olarak, belirli bir GPU için PTX’in derlenmesi uygulama çalışma zamanında anlık (JIT) olarak gerçekleşebilir. Şekil 1’de gösterildiği üzere, bir uygulamanın yürütülebilir hali hem GPU ikili kodları (cubin) hem de PTX kodunu içerebilir. PTX’in yürütülebilir hale gömülmesi, CUDA’nın PTX’i uygulama çalışma zamanında uygun cubin’e JIT derlemesine olanak tanır. PTX için JIT derleyici, NVIDIA GPU sürücüsü‘nde bulunur.

PTX’i uygulama içinde gömmek, uygulamanın derlendiği sırada ilk derleme aşaması olan yüksek seviyeli dilin PTX’e dönüşmesini sağlarken, uygulamanın çalışma zamanında PTX’ten cubin’e geçiş aşamasını geciktirmeyi mümkün kılar. Aşağıda, yukarıdaki örnekten gelen *vecAdd* çekirdeği için PTX kodu bulunmaktadır. Assembly dili görenler, PTX’in sözdizimi ve formatını tanıdık bulacaklardır. Kodun detaylarını anlamak zorunda değilsiniz. Bunun yerine bu kod, PTX’e bir göz atmanızı sağlamak ve PTX’in ne olduğu hakkında bir fikir vermek için sağlanmıştır: CUDA platformunun assembly dili.

.visible .entry _Z6vecAddPfS_S_j(
.param .u64 _Z6vecAddPfS_S_j_param_0,
.param .u64 _Z6vecAddPfS_S_j_param_1,
.param .u64 _Z6vecAddPfS_S_j_param_2,
.param .u32 _Z6vecAddPfS_S_j_param_3
)
{
.reg .pred %p;
.reg .f32 %f;
.reg .b32 %r;
.reg .b64 %rd;


ld.param.u64 %rd1, [_Z6vecAddPfS_S_j_param_0];
ld.param.u64 %rd2, [_Z6vecAddPfS_S_j_param_1];
ld.param.u64 %rd3, [_Z6vecAddPfS_S_j_param_2];
ld.param.u32 %r2, [_Z6vecAddPfS_S_j_param_3];
mov.u32 %r3, %tid.x;
mov.u32 %r4, %ntid.x;
mov.u32 %r5, %ctaid.x;
mad.lo.s32 %r1, %r5, %r4, %r3;
setp.ge.u32 %p1, %r1, %r2;
@%p1 bra $L__BB0_2;


cvta.to.global.u64 %rd4, %rd1;
mul.wide.u32 %rd5, %r1, 4;
add.s64 %rd6, %rd4, %rd5;
cvta.to.global.u64 %rd7, %rd2;
add.s64 %rd8, %rd7, %rd5;
ld.global.f32 %f1, [%rd8];
ld.global.f32 %f2, [%rd6];
add.f32 %f3, %f2, %f1;
cvta.to.global.u64 %rd9, %rd3;
add.s64 %rd10, %rd9, %rd5;
st.global.f32 [%rd10], %f3;


$L__BB0_2:
ret;
}

Hesaplama Kapasitesi ve NVIDIA GPU Donanım ISAları

Tüm NVIDIA GPU’ları, hesaplama kapasitesi olarak bilinen bir sürüm tanıtıcısına sahiptir. Her hesaplama kapasitesinin bir ana ve bir alt sürüm numarası vardır. Örneğin, hesaplama kapasitesi 8.6, 8 ana ve 6 alt sürümüne sahiptir.

Her işlemci gibi, NVIDIA GPU’larının da belirli bir ISA’sı vardır. Farklı nesillerdeki GPU’ların farklı ISAları vardır. Bu ISAlar, bir GPU’nun hesaplama kapasitesine karşılık gelen bir sürüm numarası ile tanımlanır. Bir ikili dosya (cubin) derlendiğinde, belirli bir hesaplama kapasitesi için derlenir.

Örneğin, NVIDIA Ampere neslinden GeForce ve RTX GPU’ları 8.6 hesaplama kapasitesine sahiptir ve bunların cubin sürümü sm_86‘dır. Tüm cubin sürümleri, hesaplama kapasitesinin ana ve alt numaralarına karşılık gelen sm_XY formatına sahiptir.

NVIDIA GPU’ları, farklı nesillerde ve aynı nesildeki farklı ürünlerde bile değişik ISAlara sahip olabilir. Bu nedenle PTX’e sahip olmak önemlidir.

PTX: Versiyonlu Sanal Mühendislik Assembly Dili

PTX, bir sanal makine ISA’sıdır. Daha önce de belirtildiği gibi, sanal makine ISA’sı belirli bir gerçek işlemci için değil, bir varsayımsal işlemci için desteklenen talimatlar kümesidir. Sanal makine ISAları, gerçek donanım ISAlarından biraz daha soyut olduğundan, CUDA assemble aracı ptxas geleneksel bir assemblerdan ziyade bir derleyici gibi davranır. PTX programlarını GPU ikili dosyalarına, yani cubin’lere derler.

GPU özellikleri, CUDA’nın ilk tanıtıldığı günden bu yana büyümüştür. Yeni GPU donanım nesilleri ile birlikte GPU’lara yeni özellikler eklenmektedir. PTX’in tanımladığı sanal makine de buna eşlik etmiştir. PTX spesifikasyonunda yapılan değişiklikler genellikle yeni talimatların eklenmesiyle ilgilidir.

Sonuç olarak, farklı talimat setlerini destekleyen farklı PTX sürümleri bulunmaktadır. PTX sürüm numarası, sanal mimaride hangi talimatların mevcut olduğunu gösterir. Cubin sürümleri gibi, bu sürüm numaraları bir GPU hesaplama kapasitesine karşılık gelir.

Örneğin, NVIDIA Ampere nesli GA100 GPU’su, 8.0 hesaplama kapasitesine sahiptir. *compute_80* olarak adlandırılan PTX versiyonu, GA100 tarafından desteklenen tüm talimatlara sahiptir. PTX sürümleri compute_XY olarak adlandırılır; burada X ve Y, hesaplama kapasitesinin ana ve alt numaralarına karşılık gelir.

GPU Kodu Uyumluluğu

CUDA, farklı GPU’lar arasında iki uyumluluk mekanizması sağlamaktadır: ikili uyumluluk ve PTX JIT uyumluluğu.

İkili Uyumluluk

NVIDIA GPU’ları, ana hesaplama kapasitesi sürümünde, alt sürüm aynı veya daha yüksek olduğu sürece ikili uyumludur. Bu, sm_86 için derlenmiş bir cubin’in, sm_8x‘in herhangi bir versiyonu için yüklenebileceği anlamına gelir; burada x, 6 ve üzerindeki değerlerdir.

Örneğin, *sm_86* için derlenmiş bir cubin (örneğin, NVIDIA RTX A4000), aynı zamanda *sm_89* (örneğin, NVIDIA RTX 4000 Ada Generation) üzerinde de çalışabilir. Ancak, cubin’in alt sürüm numarası, GPU’nunkinden daha düşük olan bir 8.0 hesaplama kapasitesine sahip bir cihazda çalışmaz.

NVIDIA GPU’ları, ana hesaplama kapasitesi sürümleri arasında ikili olarak uyumlu değildir. *sm_86* için derlenen bir cubin, 9.0 hesaplama kapasitesine sahip (örneğin, NVIDIA Hopper mimarisi) veya daha yeni bir GPU’da çalıştirilamaz.

PTX JIT Uyumluluğu

PTX’in bir çalıştırılabilir dosyaya gömülmesi, farklı ana sürümlere sahip GPU’lar arasında uyumluluk sağlamanın bir mekanizmasıdır. Şekil 1’de gösterildiği gibi, hem PTX hem de cubin son uygulama yürütülebilir dosyalarında depolanabilir. PTX ve cubin aynı zamanda kütüphanelerde de yer alabilir.

PTX kodu, bir uygulama veya kütüphane ikili dosyasında depolandığında, çalıştırıldığı GPU için JIT derlenebilir. Örneğin, uygulama veya kütüphane *compute_70* için hedeflenen PTX barındırıyorsa, bu PTX, hesaplama kapasitesi 7.0 veya daha yüksek olan herhangi bir GPU için JIT derlenebilir. Ancak, PTX, hedef sürümden daha düşük hesaplama kapasitelerine JIT derlenemez. Örneğin, *compute_70* için hedeflenen bir PTX, 5.x veya 6.x hesaplama kapasitesine sahip bir GPU için JIT derlenemez.

Fatbin’ler

CUDA uygulamaları veya kütüphaneleri derlenirken, çeşitli cubinler ve PTX sürümlerini içeren bir kapsayıcı oluşturulmaktadır; bu kapsayıcıya *fatbin* denir. Örneğin, Şekil 2’deki fatbin, *compute_70* için PTX’in yanı sıra *sm_70*, *sm_80* ve *sm_86* için cubin içerir. Bu, uygulamanın hesaplama kapasitesi 7.0, 8.0 ve 8.6 GPU’ları için ikili kodları önceden barındırdığı anlamına gelir. *sm_86* cubini, uygulama hesaplama kapasitesi 8.9 GPU’sunda çalıştırıldığında da yüklenebilir.

*compute_70* PTX, hesaplama kapasitesi 7.0 veya daha yüksek olan herhangi bir GPU için JIT derlenebilir; bu nedenle bu uygulama, cubin’lerin mevcut hedeflerinden daha yeni GPU’larda çalıştırılabilir. Örneğin, bu uygulama 9.0, 10.0 veya 12.0 hesaplama kapasitesine sahip bir GPU’da çalıştırılabilir. Tablo 1, Şekil 2’deki fatbin’de bulunan her bileşenin bu belirli örnekte nasıl uyumluluk sağladığını göstermektedir.

PTX, uygulama başladığında veya kod GPU üzerinde ilk kez kullanıldığında sürücü tarafından derlenir. JIT derlemenin ne zaman gerçekleşeceğini kontrol etme hakkında daha fazla bilgi için Tembel Yükleme bölümüne göz atın.

PTX’in Avantajları

PTX, bir ara kod formatı olarak kullanıldığında, CUDA platformu geliştiricilerin henüz yaratılmamış GPU’larda çalışacak uygulama ikili dosyaları geliştirmelerine olanak tanır. 2018 yılında NVIDIA Turing mimarisi (CC 7.5) için derlenen bir uygulama, 2025’te NVIDIA Blackwell (CC 12.0) GPU’sunda çalışabilir, ayrıca gelecekte çıkacak GPU’larda da çalışabilir.

Uygulamalara veya kütüphanelere GPU kodunun PTX temsilini gömerek, CUDA sürücüsü, PTX kodunu çalışma zamanında herhangi bir mimari için JIT derlemeyi sağlar. Uygulama veya kütüphane dağıtan geliştiriciler için bu durum, uygulamanın veya kütüphanenin gelecekteki GPU mimarilerinde, ikili dosyayı güncellemeye gerek kalmadan çalışabilmesi anlamına gelir.

PTX, CUDA GPU hesaplama platformunun assembly dilini temsil etmenin yanı sıra, başka herhangi bir dil için hedef alınabilir bir temsildir. Örneğin, bir alan spesifik dil (DSL) derleyicisi PTX kodu üretebilir; bu kod daha sonra NVIDIA GPU’larında çalıştırılabilir. OpenAI Triton, PTX üreten bir DSL örneğidir.

DSL oluşturma konusunda ilgilenen geliştiricilerin, NVVM IR ve libNVVM‘yi incelemeleri faydalı olacaktır, çünkü bu kaynaklar PTX oluşturma için özel bir uygulama geliştirmekten daha iyi alternatifler sunmaktadır.

PTX’i Doğru Bir Şekilde Yazmak

NVIDIA tarafından belgelenen PTX sanal makine ISA’sı bulunmaktadır. PTX kodunu elle yazmak mümkündür. Diğer assembly dilleri gibi, bu genelde büyük yazılım projeleri için mantıklı bir seçim değildir. Yüksek seviyeli diller, assembly dillerine veya sanal assembly dillerine göre daha iyi geliştirici verimliliği sağlar.

PTX’i elle yazmak, işlemcinin kodu nasıl çalıştıracağı konusunda içgörülere sahip uzmanların, nihai yürütülebilir dosyadaki talimatlara ince bir kontrol sağlarken kullanabileceği bir yöntemdir. Bu yöntem ile performans iyileştirmeleri elde etmek mümkün olsa da, çoğu geliştirici için bu genelde gerekli ya da tavsiye edilen bir durum değildir.

Yine de bazı geliştiriciler, milyarlarca kez çalıştırılacak iç döngü kodlarında doğrudan PTX yazarak optimizasyon yapmayı tercih edebilmektedir. Bu tür kodlarda, en küçük performans kazanımları bile yüksek tekrar sayısıyla çarpıldığında anlamlı hale gelir ve bu tür ince ayarın çabası değerli hale gelir. Gelecek bir yazıda bu tür optimizasyon örneklerine de yer vermeyi planlıyoruz.

libcu++, CUDA Toolkit’i içinde yer almaktadır ve cuda::ptx ad alanı, PTX talimatlarına doğrudan eşlenen işlevler sunmaktadır. Bu, C++ uygulamalarında belirli PTX talimatlarını kullanmayı kolaylaştırır. Daha fazla bilgi için libcu++ belgelerine göz atabilirsiniz. Ayrıca NVIDIA, C++ kodu içine PTX’yi doğrudan inleyerek kullanma ile ilgili belgeler de sunmaktadır.

Özet

PTX, CUDA GPU hesaplama platformu için assembly dili gibi düşünülebilen sanal makine ISA’sıdır. PTX, CUDA GPU hesaplama platformunun önemli bir parçasıdır. Daha yüksek seviyeli diller, PTX’e derlenir ve ardından PTX, derleme zamanında veya çalışma zamanında ikili koda dönüştürülür.

PTX kodunu ikili dosyalarına gömen uygulamalar ve kütüphaneler, çeşitli nesiller arası uyumluluk elde edebilir. Ayrıca, diğer programlama dilleri veya alan spesifik dil derleyicileri de PTX’e derleyebilir; bu PTX, ptxas veya JIT derlemesi yoluyla NVIDIA GPU’larında çalışabilecek ikili kodlar üretir.

Geliştirici olarak GPU kodunuzun uyumluluğunu maksimize etmek için uygulamanıza veya kütüphanenize PTX dahil etmeyi unutmayın.

Teşekkürler

Bu yazıya katkıda bulunan Rob Armstrong ve Jake Hemstad’a teşekkürler.