CUDA çekirdeğini çalıştırdığınızda içeride neler olur

 (fergusfinn.com)
2 puan yazan GN⁺ 4 시간 전 | 1 yorum | WhatsApp'ta paylaş
  • Basit bir vektör toplama CUDA programı bile 2.000000 sonucunu elde edene kadar derleme hattı, sürücü çağrıları, GPU komut kuyruğu, warp zamanlama, bellek hiyerarşisi ve tamamlama semaforlarından geçer
  • nvcc, host kodu ile device kodunu ayırır; cicc ile PTX, ptxas ile SASS üretir ve cubin ile PTX'i bir fatbin içinde paketleyip Linux yürütülebilir dosyasının içine yerleştirir
  • vadd<<<4096, 256>>> launch sözdizimi bir host launch stub'a dönüştürülür; da, db, dc, n argümanları CUDA runtime ve libcuda.so.1 üzerinden sürücüye iletilir
  • GPU yürütmesi QMD, pushbuffer, GPFIFO, GP_PUT ve doorbell MMIO yazımıyla başlar; RTX 4090'ın 128 SM'i 4096 blok ve 256 thread yapılandırmasını warp düzeyinde yürütür
  • Bu çekirdek, her float toplama başına 12 bayt aktarım gerektiren düşük aritmetik yoğunluk nedeniyle Nsight Compute'ta 10.78μs, DRAM tepesinin %79.65'i ve %5.17 warp issue ile bellek bant genişliğine bağımlıdır

Örnek çekirdek ve inceleme kapsamı

  • Örnek program, iki float dizisini toplayıp üçüncü diziye yazan vadd CUDA çekirdeğini kullanır
    • n = 1 << 20 ile 1.048.576 adet float işlenir
    • launch yapılandırması vadd<<<4096, 256>>>(da, db, dc, n) olup 4096 * 256 = n thread kullanır
  • RTX 4090 için nvcc -arch=sm_89 ile derlenip çalıştırıldığında c[0]=2.000000 c[n-1]=2.000000 çıktısı alınır
  • Bu tek satırlık sonucun arkasında bile on milyonlarca CPU komutu, device file, yaklaşık 900 adet ioctl ve belleğe eşlenmiş doorbell register'ı bulunur

nvcc yürütülebilir dosyayı nasıl oluşturur

  • nvcc --keep kullanıldığında derleme hattısının çıktıları doğrudan görülebilir
    • vadd.ptx: cicc tarafından üretilen device kodunun PTX'i
    • vadd.sm_89.cubin: ptxas tarafından üretilen device kodunun SASS'ı
    • vadd.fatbin: cubin ve PTX'i bir araya getiren fatbin
    • vadd.cudafe1.stub.c: host launch stub ve çekirdek kayıt kodu
    • vadd.o: fatbin'i içeren son host object
  • Host kodu host derleyici tarafından işlenirken, vadd device çekirdeği cicc ve ptxas aşamalarından geçer
  • PTX, tipli sonsuz sanal register kullanan bir sanal ISA'dır ve gerçek donanımdaki register sayısını doğrudan yansıtmaz
    • Örnek PTX; blockIdx.x * blockDim.x + threadIdx.x hesabı, sınır denetimi, global load, float add ve global store içerir
    • CUDA pointer'ları varsayılan olarak generic pointer olduğu için, ld.global kullanılmadan önce cvta.to.global ile global adrese dönüştürülür
    • mul.wide.s32, index'i sizeof(float) olan 4 baytlık ofsete çevirir ve 32 bitten 64 bite genişletir
  • SASS, mimariye özgü gerçek komutlardır ve RTX 4090 hedefli çıktıda PTX'ten daha sıkıştırılmış görünür
    • S2R, SR_CTAID.X, SR_TID.X gibi özel register'ları genel amaçlı register'lara kopyalar
    • PTX'teki mul.wide ile add birleşimi SASS'ta IMAD.WIDE olarak birleştirilir
    • cvta dönüşümü adresleme sürecine yedirilir
  • c[0x0][...] operandı, sürücü tarafından yönetilen constant bank 0'ı işaret eder
    • a, b, c pointer'ları 0x160, 0x168, 0x170 konumlarındadır
    • n, 0x178 konumundadır
    • blockDim.x gibi launch geometrisi ve ABI değerleri de aynı bankta yer alır
  • cubin, Linux yürütülebilir dosyalarıyla aynı kapsayıcı biçim olan bir ELF dosyasıdır
    • fatbinary, cubin ve PTX'i birlikte paketler
    • Bu RTX 4090 üzerinde gerçek yürütme SASS ile yapılır, ancak PTX başka mimarilerde sürücünün JIT derleme yapabilmesi için fallback olarak eklenir
    • PTX ayrıntılı bir düz metin olduğundan nvcc bunu varsayılan olarak sıkıştırır

Host kodu launch'ı nasıl hazırlar

  • Derleyici frontend'i cudafe++, main'den önce çalışan gizli bir constructor ekler
    • Bu constructor, gömülü fatbinary'yi CUDA runtime'a kaydeder
    • Host tarafındaki vadd fonksiyon pointer'ını, fatbin içindeki mangled device kernel name ile eşler
  • vadd<<<4096, 256>>>(da, db, dc, n) sözdizimi, üretilmiş bir host launch stub ile değiştirilir
    • da, db, dc, n; host belleğindeki argument buffer'a sırasıyla 0, 8, 16, 24 ofsetleriyle hizalanarak yerleştirilir
    • Bu ofsetler, SASS'ın constant bank 0'dan okuduğu 0x160, 0x168, 0x170, 0x178 konumlarıyla eşleşir
  • Stub, __cudaLaunch çağrısı yaparken host tarafındaki dummy vadd fonksiyon adresini iletir
    • Bu adres, CPU'da çalıştırılacak bir fonksiyon adresi değil; runtime kayıt tablosunda arama yapmak için kullanılan bir anahtardır
    • Runtime karşılık gelen device symbol name'i bulduktan sonra kapalı kaynak user-mode driver olan libcuda.so.1e geçer
  • İlk GPU çağrısında CUDA runtime, libcuda.so.1 dosyasını dinamik olarak açar ve bir context oluşturur
    • strace içinde /lib/x86_64-linux-gnu/libcuda.so.1 dosyasının açıldığını görmek mümkündür
    • Bu context, CPU'nun GPU ile haberleştiği channel'ı içerir
  • CUDA 12.2'den itibaren module loading varsayılan olarak lazy'dir
    • Belirli bir çekirdek ilk kez launch edilene kadar SASS cubin yüklemesi ertelenir
    • Bu davranış CUDA_MODULE_LOADING ile kontrol edilebilir

GPU'ya iş aktaran komut kuyruğu

  • GPU, CPU gibi fonksiyon çağrısı alıp bir entry point'e atlamaz
    • Bunun yerine, PCIe bus üzerinden host belleğindeki sürücü komut akışını okur
    • cuLaunchKernel, tamamlanmış launch komutunu bu akışa yerleştirir ve GPU'ya haber verir
  • İlk çalıştırmada sürücü, çekirdeğin SASS kodunu GPU belleğine kopyalar
    • Bir code buffer ayırır ve SASS'ı buraya kopyalar
  • Channel içinde host RAM'de duran iki temel yapı vardır
    • pushbuffer: sürücünün GPU komutu olan method'ları yazdığı bellek alanı
    • GPFIFO: pushbuffer span'lerini gösteren bir pointer ring buffer
  • Bir GPFIFO girdisi, pushbuffer span'inin (base, length) bilgisini taşıyan iki adet 32 bitlik word'den oluşur
  • GPU ve sürücü, tüketim ve üretim konumlarını iki cursor ile izler
    • GP_GET: GPU'nun ne kadarını tükettiğini gösterir
    • GP_PUT: sürücünün ne kadarını ürettiğini gösterir
    • Her ikisi de USERD adlı channel başına yapıda bulunur
  • Çekirdek launch edildiğinde sürücü, pushbuffer span'ine method'ları yazar, GPFIFO girdisini buna işaret ettirir ve ardından GP_PUT değerini ilerletir
  • Modern GPU'larda host engine cursor'ları sürekli izlemediği için doorbell gerekir
    • GPU, prosese küçük bir register window eşler
    • Sürücü, channel'ın work-submit token'ını doorbell register'ına yazar
    • Host engine, doorbell sinyalini aldıktan sonra GP_PUT değerini okur ve GPFIFO girdisiyle pushbuffer span'ini DMA üzerinden getirir

QMD'nin taşıdığı yürütme bilgisi

  • Launch, SET_INLINE_QMD_ADDRESS_A/B ve LOAD_INLINE_QMD_DATA method burst'ü ile başlar
  • QMD(Queue Meta Data), compute grid için kullanılan launch descriptor'dur
    • 4096 ve 256 olan grid ve block boyutlarını içerir
    • Thread başına register sayısı ile shared memory gereksinimini içerir
    • Program başlangıç adresi ile çekirdek argümanlarını taşıyan constant bank adresini içerir
    • Tamamlamanın nereye bildirileceğini de içerir
  • Host stub'ın paketlediği argümanlar sürücü tarafından constant bank'e kopyalanır ve bu bankın adresi QMD'ye yazılır
  • QMD, GPU'ya SASS'ın nerede olduğunu, paralel programın nasıl yapılandırılacağını ve completion signal'ın nereye yazılacağını söyler
  • cuLaunchKernel, doorbell çaldığı anda geri döner
    • Çağrı asenkron olduğu için CPU, GPU işi sürerken çalışmaya devam edebilir

SM, warp ve doluluk

  • Host engine, QMD'yi compute work distributor'a iletir
    • Bu bileşen tüm GPU üzerinde tektir
    • Doğrusal SASS instruction stream'ini SM'lere dağıtarak paralel program olarak yürütülmesini sağlar
  • Hedef GPU olan GeForce RTX 4090, 128 SM kullanır
    • Launch, 4096 bloktan ve blok başına 256 thread'den oluşur
  • Her SM'in yerel bir instruction cache'i vardır ve aktif warp'lar kendi program counter'larını korur
    • Volta sonrasında thread başına program counter ve call stack sunan Independent Thread Scheduling modeli vardır
    • Ancak issue işlemi hâlâ warp düzeyinde yapılır
  • Örnek çekirdekte block residency'yi kaynak sınırları belirler
    • Blok başına 256 threads = 8 warps
    • ptxas, thread başına 16 register ayırır
    • Register açısından SM başına 16 blok mümkündür
    • Ancak thread kapasitesi SM başına 1.536 aktif thread olduğundan 1536 / 256 = 6 blokla sınırlanır
    • Dolayısıyla SM başına en fazla 6 blok, yani 48 warp resident durumda olabilir
  • SM, 4 processing block yani sub-partition'a ayrılır
    • Resident 48 warp, 4 sub-partition arasında eşit dağıtılır
    • Her warp scheduler tam dolu durumda 12 aktif warp yönetir
    • Her cycle'da eligible olan bir warp seçilir ve sonraki komut 32 lane'e dispatch edilir

Bir warp'ın eligible duruma gelme koşulları

  • GPU, CPU'nun out-of-order yürütmesine benzer biçimde tek bir thread içinden büyük ölçüde dinamik bağımlılık çıkarmaz
    • Bunun yerine çok sayıda resident warp tutar ve biri stall olduğunda gecikmeyi gizlemek için diğerine geçer
    • Derleyici öngörülebilir zamanlamayı planlar, donanım scoreboard ise öngörülemez kısımları yönetir
  • 128 bitlik SASS instruction içinde ptxas tarafından yazılan control-code payload bulunur
    • Sabit gecikmeli instruction'larda statik stall count yer alır
    • Yield hint, scheduler önceliğinin bırakılıp bırakılmayacağını belirtir
    • Değişken gecikmeli işlemler için warp başına 6 fiziksel scoreboard barrier kullanılır
  • Örnek SASS bölümünde iki LDG.E, aynı scoreboard barrier olan B2yi set eder
    • FADD, B2 için wait-on bilgisi taşır
    • İki load geri dönüp barrier temizlenene kadar o warp ineligible durumda kalır
    • Bu sırada scheduler aynı sub-partition'daki başka bir warp'ı seçer
  • FADD ile STG.E arasındaki bölüm sabit gecikme olarak ele alınır
    • FADD, stall=5 taşır ve R9 sonucu hazır olana kadar warp'ı birkaç cycle bekletir
    • Ayrı bir barrier gerekmez
  • Bu control payload, nvdisasm varsayılan çıktısında gizlidir
    • cuobjdump -sass tarafından verilen ham 128 bit encoding'de ikinci 64 bitlik word içinde yer alır
    • Yerleşim belgelenmiş değildir; microbenchmarking ile yeniden oluşturulmuştur

Bellek erişimi ve performans ölçümü

  • Warp, LDG.E yürüttüğünde 32 thread'in her biri kendi adresini hesaplar
    • Örnekte ardışık float dizi erişimi olduğu için, tüm warp 32 * 4 = 128 bytes boyutunda ardışık bir blok ister
  • SM load/store unit, request coalescing yapar
    • 32 adet 4 baytlık isteği 4 adet 32 baytlık sector request'e birleştirir
    • Erişimler ardışık olmasaydı, gerekenden daha fazla veri okunabilirdi
  • Coalesced request önce SM içindeki yerel L1 Data Cache'i kontrol eder
    • Miss olursa crossbar interconnect üzerinden 72MB L2 Cache slice'a gider
    • L2'de de miss olursa memory controller ve memory bus üzerinden GDDR6X VRAM'e gider
  • STG.E store işlemi de ilke olarak bunun ters yönündeki aynı yolu izler
  • Nsight Compute ölçümleri, bu çekirdeğin memory-bound olduğunu gösterir
    • launch__grid_size: 4.096
    • launch__block_size: 256
    • launch__registers_per_thread: 16
    • launch__waves_per_multiprocessor: 5.33
    • sm__warps_active.avg.pct_of_peak: 82.77%
    • smsp__issue_active.avg.pct_of_peak: 5.17%
    • dram__throughput.avg.pct_of_peak: 79.65%
    • gpu__time_duration.sum: 10.78μs
  • Çekirdeğin aritmetik yoğunluğu çok düşüktür
    • İki adet 4 baytlık load ve bir adet 4 baytlık store ile toplam 12 bayt aktarım başına yalnızca bir float toplama yapılır
    • DRAM okuma açısından bakıldığında, 8.4MB veri 10.78μs'de okunarak yaklaşık 780GB/s elde edilir; bu, tepe değerin yaklaşık beşte dördüdür
    • 4MB boyutundaki c çıktısı 72MB L2'ye sığdığı için, device-to-host copy bunu okumadan önce DRAM'e flush edilmez

Sonucun CPU'ya geri dönmesi

  • Kernel launch, doorbell çaldığı anda CPU'ya döndüğünden, GPU tamamlandığını ayrıca bildirmek zorundadır
  • 4096 bloğun tamamı retire edildiğinde GPU, QMD'de belirtilen completion semaphore'u post eder
    • QMD'nin fence alanı 23–24. word'lerde yer alır
  • Varsayılan stream'de cudaMemcpy(c, dc, ...), çekirdeğin arkasına yerleştirilir
    • GPU copy engine, semaphore yükselene kadar gated durumda kalır
    • c hâlâ 72MB L2 içinde dirty durumda olduğundan, copy engine read işlemi DRAM'e gidip gelmeden L2 üzerinden tamamlanır
    • Veri, PCIe üzerinden host belleğine taşınır
  • Kopyalama bittiğinde copy engine kendi semaphore'unu post eder
    • Host tarafındaki cudaMemcpy bekleyişi sona erer
    • c yeniden sıradan host belleği olur
    • printf, c[0] ve c[n-1] değerlerini RAM'den okuyup stdout'a yazar

Launch'ın içini nasıl görebilirsiniz

  • Açık kernel modüllerini okumak tek başına yeterli değildir; çünkü libcuda kapalı kaynak olduğundan bazı davranışlar doğrudan doğrulanamaz
  • Method write işlemleri syscall üzerinden geçmez; önceden eşlenmiş write-combined buffer'a doğrudan yazılır, bu yüzden pushbuffer'ı görmek için belleğin kendisini okumak gerekir
  • LD_PRELOAD shim ile mmap sarmalanarak /dev/nvidia* üzerinde eşlenen bölgeler kaydedilebilir
    • Test programı launch'tan hemen sonra shim'in dump fonksiyonunu çağırırsa, eşlenen pushbuffer çıktılanabilir
    • Dump, SET_INLINE_QMD_ADDRESS_A ile ilişkili method burst'ü arar
  • Pushbuffer method header'ı, opcode, payload count, subchannel index ve register offset'i bit alanlarında taşır
    • 0x0318, SET_INLINE_QMD_ADDRESS_A'ya karşılık gelir
    • 0x0320 + i * 4, LOAD_INLINE_QMD_DATA(i) anlamına gelir
    • Dump içinde count değeri 66 olan increasing-method burst görülür; 2 adres word'ü ile 64 QMD word'ü, yani toplam 256 baytlık QMD inline biçimde taşınır
    • QMD içindeki 12. word 0x1000, 18. word ise 0x100 olup launch'taki 4096 ve 256 değerlerine karşılık gelir
  • Sürücü kurulumu ioctl ile yapılır
    • Tek çekirdekli bir programda strace, 948 adet ioctl kaydeder
    • Bunların çoğu tek seferlik kurulum işlemleridir
    • Başlıca file descriptor'lar /dev/nvidiactl ve /dev/nvidia-uvm'dir
    • NVIDIA resource manager ioctl magic byte'ı 0x46, yani 'F'tir
    • 0x2A komut numarası NV_ESC_RM_CONTROL, 0x2B ise NV_ESC_RM_ALLOC olarak yorumlanır
  • nvcc --keep ile üretilen vadd.cudafe1.stub.c içinde başlangıç kayıt kodu da görülebilir
    • __attribute__((__constructor__)) ile işaretlenen bir fonksiyon main'den önce çalışır
    • __cudaRegisterBinary ve __cudaRegisterEntry aracılığıyla host function pointer vadd ile device entry point _Z4vaddPKfS0_Pfi birbirine bağlanır

İlgili okumalar

1 yorum

 
GN⁺ 4 시간 전
Hacker News görüşleri

  • İlgi çekici bir yazıydı; varsayılan akışın semaforları açıklaması da güzeldi
    CUDA'nın komut senkronizasyonunu örtük olarak halletmesi ve paralel komutları akışlar üzerinden isteğe bağlı kullandırması hoş
    Senkronizasyon karmaşıklığının tamamını en baştan kullanıcıya yükleyen Vulkan ile tezat oluşturuyor

  • Donanım tarafında bazı açık belgeler var
    Metot dokümantasyonunu ya da QMD biçimini bulmak için mutlaka kernel kaynağını okumak gerekmiyor
    https://github.com/NVIDIA/open-gpu-doc/blob/master/classes/c... bkz.

  • Çok faydalıydı
    Özellikle doorbell ve QMD kısmı, CUDA çalıştırma söz diziminin GPU'ya gerçekten gönderilen içerikle nasıl bağlandığını gösterdiği için en yararlı bölümdü
    Çoğu açıklama kernel, blok ve warp civarında duruyor; bu yazıysa CPU→sürücü→GPU yolunu takip etmeyi çok daha kolaylaştırıyor

  • Kontrol kodu, yazıda anlatılandan biraz daha karmaşık
    Gerçekte kontrol sözcüğü içindeki bitlerden ziyade daha çok tablo araması gibi

  • Artık kernel'leri optimize edip daha hızlı çalıştırmayı ana işi yapan şirketler var
    Bu şirketlerin bir gün bunu çok iyi yapan bir açık kaynak kütüphaneye yenilip yenilmeyeceğini merak ediyorum
    Nvidia'nın böyle bir şeyi istediği anda çıkarabilecek gibi bir hali var
    Ya da büyük sağlayıcılar çıkarım hızını artıran bir moat olarak görmek için bu şirketleri satın alırsa daha da iyi durumda olabilirler

    • Kısa vadede yetenek kazanımı odaklı satın almalar oldukça olası görünüyor
      Yine de kernelbench gibi ilgili benchmark'larda modellerin gelişimine bakınca, daha genel amaçlı çözümler de sonunda ortaya çıkacaktır diye düşünüyorum
      Sorun şu ki her yeni donanım neslinde mevcut modellerin daha önce görmediği kısıtlar ya da özellikler sık sık çıkıyor
      Örneğin Blackwell'in tcgen05'i bir dönem dağılım dışı bir örnekti
      Modeller daha iyi genelleme yapmaya başladığında bu ölümcül bir engel olmayabilir ama en azından şu anda hâlâ bir pürüz
      [1] https://kernelbench.com/
    • CUDA'yı büyük ölçekte çalıştırdığınızda Nvidia sürücüsü ve kütüphane hatalarıyla uğraşmak için iğrenç denecek kadar fazla mühendislik zamanı harcanıyor
      Nvidia kütüphanelerine daha fazla bağımlı olmayı bekleyen pek kimse görmedim
    • Muhtemelen hayır
      Çünkü iş yükünün ayrıntıları; yani tam parametreler, bellekteki veri temsili, değer aralıkları vb. optimizasyon stratejilerini büyük ölçüde ayrıştırıyor

  • HPC yüksek lisansını yeni bitirdim; CUDA, MPI+CUDA ve OpenCL dersleri aldım, bu tür bir yazıyı derslerden önce okumuş olsam çok daha faydalı olurdu
    Özellikle bir warp'ın çalıştırılabilir olmasının ne anlama geldiğini ele alan bölümün öncesi ve sonrası çok iyiydi

  • Öncelikle, pek çok köşeyi bucağı iyi kurcalayan güzel bir yazı
    Ancak CUDA'nın runtime APIsi kullanılmadığında kullanıcı alanındaki büyü gibi görünen birçok kısım ortadan kalkıyor
    Sürücü API'sini kullanıp kernel kaynağını string olarak alarak NVIDIA'nın çalışma zamanı derleyicisiyle derlerseniz, neler olduğunu çok daha iyi görebilirsiniz
    Hepsi değil ama epey büyük bir kısmı şeffaflaşıyor
    Daha “ham” bir sürüm burada:
    https://github.com/NVIDIA/cuda-samples/tree/master/cpp/0_Int...
    Aynı şeyi çok daha okunabilir ama yine de tamamen şeffaf bir modern C++ API biçiminde görmek isterseniz buna bakın:
    https://github.com/eyalroz/cuda-api-wrappers/blob/master/exa...
    Bu, benim CUDA API wrappers başlık-dosyası-yalnız kütüphanemin örnek programı

    • Sürücü API'si güzel, çünkü CUDA kernel'lerini hot-reload edilebilir shader'lar gibi ele alabiliyorsunuz
      Çalışma sırasında kodu değiştirerek geliştirme yapabilmek eğlenceli

  • Bare metal'de mi?