Paralel İş Parçacığı Yürütme - Parallel Thread Execution

Paralel İş Parçacığı Yürütme (PTXveya NVPTX[1]) düşük seviyelidir paralel Konu icra sanal makine ve komut seti mimarisi kullanılan Nvidia 's CUDA programlama ortamı. nvcc derleyici, CUDA'da yazılmış kodu çevirir, bir C ++ benzeri bir dil, PTX komutlarına ve grafik sürücüsü, PTX talimatlarını ikili koda çeviren bir derleyici içerir[2] işlem çekirdekleri üzerinde çalıştırılabilir Nvidia GPU'lar. GNU Derleyici Koleksiyonu ayrıca PTX üretimi için temel beceriye sahiptir. OpenMP boşaltma.[3] Satır içi PTX montajı CUDA'da kullanılabilir.[4]

Kayıtlar

PTX, keyfi olarak büyük bir yazmaç seti kullanır; derleyicinin çıktısı neredeyse saf tek atama formu, genellikle ardışık kayıtlara atıfta bulunan ardışık satırlarla. Programlar, form beyanları ile başlar

.reg .u32 % r<335>; // işaretsiz 32 bitlik tam sayı türünde% r0,% r1, ...,% r334 335 kayıt bildir

Üç bağımsız değişkenli bir derleme dilidir ve hemen hemen tüm talimatlar üzerinde çalıştıkları veri türünü (işaret ve genişlik açısından) açıkça listeler. Kayıt adlarının önünde bir% karakteri vardır ve sabitler değişmezdir, örneğin:

shr.u64 % rd14, % rd12, 32; // işaretsiz 64 bitlik bir tamsayıyı% rd12'den 32 sıra sağa kaydır,% rd14 ile sonuçlanırcvt.u64.u32 % rd142, % r112; // işaretsiz 32 bitlik bir tamsayıyı 64 bit'e dönüştür

Koşul kayıtları vardır, ancak gölgelendirici 1.0 modelinde derlenmiş kod bunları yalnızca dallanma komutlarıyla birlikte kullanır; şartlı şube

@%s 14 sutyen $ etiket; // $ etikete dal

setp.cc.type komut, uygun türdeki iki kaydı karşılaştırmanın sonucuna bir yüklem kaydı ayarlar, ayrıca bir Ayarlamak talimat, nerede set.le.u32.u64 % r101, % rd12, % rd28 32 bitlik kaydı ayarlar % r101 -e 0xffffffff 64 bitlik kayıt % rd12 64 bitlik sicilden küçük veya ona eşittir % rd28. Aksi takdirde % r101 ayarlandı 0x00000000.

Sözde kayıtları gösteren önceden tanımlanmış birkaç tanımlayıcı vardır. Diğerlerinin yanı sıra, % tid,% ntid,% ctaid, ve % nctaid sırasıyla diş indeksleri, blok boyutları, blok indeksleri ve ızgara boyutlarını içerir.[5]

Durum uzayları

Yük (ld) ve sakla (st) komutlar, birkaç farklı durum alanından (bellek bankaları) birini ifade eder, örn. ld.paramSekiz durum alanı vardır:[5]

  • .reg : kayıtlar
  • .sreg : özel, salt okunur, platforma özgü kayıtlar
  • .const : paylaşılan, salt okunur bellek
  • .global : tüm iş parçacıkları tarafından paylaşılan global bellek
  • .yerel : yerel bellek, her iş parçacığına özel
  • .param : çekirdeğe geçirilen parametreler
  • .shared : bir bloktaki iş parçacıkları arasında paylaşılan bellek
  • .tex : genel doku belleği (kullanımdan kaldırıldı)

Paylaşılan bellek, PTX dosyasında formun başındaki satırlar aracılığıyla bildirilir:

.shared .align 8 .b8 pbatch_cache[15744]; // 8 baytlık bir sınıra hizalanmış 15.744 bayt tanımlayın

PTX'de çekirdek yazmak, PTX modüllerini CUDA Sürücü API'si aracılığıyla açıkça kaydetmeyi gerektirir; bu, tipik olarak CUDA Runtime API ve NVIDIA'nın CUDA derleyicisi nvcc'yi kullanmaktan daha zahmetlidir. GPU Ocelot projesi, GPU Ocelot artık aktif olarak korunmasa da, CUDA Runtime API çekirdek çağrılarının yanı sıra PTX modüllerini kaydetmek için bir API sağladı.[6]

Ayrıca bakınız

Referanslar

  1. ^ "NVPTX Arka Uç için Kullanıcı Kılavuzu - LLVM 7 belgeleri". llvm.org.
  2. ^ "CUDA Binary Utilities". docs.nvidia.com. Alındı 2019-10-19.
  3. ^ "nvptx". GCC Wiki.
  4. ^ "CUDA'da Satır İçi PTX Montajı". docs.nvidia.com. Alındı 2019-11-03.
  5. ^ a b "PTX ISA Sürüm 2.3" (PDF).
  6. ^ "GPUOCelot: PTX için dinamik bir derleme çerçevesi". github.com.

Dış bağlantılar