CUDA kernel चलाने पर अंदर क्या होता है

 (fergusfinn.com)
2 पॉइंट द्वारा GN⁺ 3 시간 전 | 1 टिप्पणियां | WhatsApp पर शेयर करें
  • एक साधारण vector addition CUDA program भी परिणाम 2.000000 पाने तक compile pipeline, driver calls, GPU command queue, warp scheduling, memory hierarchy और completion semaphore से गुजरता है
  • nvcc host code और device code को अलग करके cicc से PTX, ptxas से SASS बनाता है, और cubin व PTX को fatbin में बांधकर Linux executable file के अंदर डालता है
  • vadd<<<4096, 256>>> launch syntax host launch stub में बदल जाता है, और arguments da, db, dc, n CUDA runtime और libcuda.so.1 के जरिए driver तक पहुंचते हैं
  • GPU execution QMD, pushbuffer, GPFIFO, GP_PUT, doorbell MMIO write से शुरू होता है, और RTX 4090 के 128 SM, 4096 blocks और 256 threads की configuration को warp units में execute करते हैं
  • यह kernel float addition के प्रति 1 operation पर 12 bytes transfer की जरूरत वाली कम arithmetic intensity के कारण Nsight Compute में 10.78μs, DRAM peak का 79.65%, warp issue 5.17% दिखाता है और memory bandwidth पर निर्भर रहता है

Example kernel और observation scope

  • Example program vadd CUDA kernel से दो float arrays को जोड़कर तीसरे array में store करता है
    • n = 1 << 20 से 1,048,576 floats process किए जाते हैं
    • launch configuration vadd<<<4096, 256>>>(da, db, dc, n) है और 4096 * 256 = n threads का उपयोग करती है
  • RTX 4090 के लिए nvcc -arch=sm_89 से compile करके चलाने पर c[0]=2.000000 c[n-1]=2.000000 print होता है
  • इस एक line के result में भी CPU के करोड़ों instructions, device file, लगभग 900 ioctl, और memory-mapped doorbell register शामिल होते हैं

nvcc executable file कैसे बनाता है

  • nvcc --keep का उपयोग करने पर compile pipeline के outputs सीधे देखे जा सकते हैं
    • vadd.ptx: cicc द्वारा बनाया गया device code का PTX
    • vadd.sm_89.cubin: ptxas द्वारा बनाया गया device code का SASS
    • vadd.fatbin: cubin और PTX को बांधने वाला fatbin
    • vadd.cudafe1.stub.c: host launch stub और kernel registration code
    • vadd.o: fatbin शामिल करने वाला final host object
  • Host code को host compiler process करता है, और device kernel vadd cicc तथा ptxas stages से गुजरता है
  • PTX एक virtual ISA है, जो typed infinite virtual registers का उपयोग करता है और actual hardware register count को सीधे reflect नहीं करता
    • Example PTX में blockIdx.x * blockDim.x + threadIdx.x calculation, bounds check, global load, float add, global store शामिल हैं
    • CUDA pointers default रूप से generic pointer होते हैं, इसलिए cvta.to.global से global address में convert करने के बाद ld.global का उपयोग किया जाता है
    • mul.wide.s32 index को sizeof(float) यानी 4-byte unit offset में बदलता है और 32-bit से 64-bit में expand करता है
  • SASS architecture-specific actual instructions हैं, और RTX 4090 target output में PTX से अधिक compact रूप में दिखते हैं
    • S2R, SR_CTAID.X, SR_TID.X जैसे special registers को general registers में copy करता है
    • PTX के mul.wide और add combination को SASS में IMAD.WIDE में merge किया जाता है
    • cvta conversion address calculation process में absorb हो जाता है
  • c[0x0][...] operand driver-managed constant bank 0 को refer करता है
    • pointers a, b, c 0x160, 0x168, 0x170 पर स्थित होते हैं
    • n 0x178 पर स्थित होता है
    • blockDim.x जैसी launch geometry और ABI values भी उसी bank में होती हैं
  • cubin Linux executable file जैसे container format वाली ELF file है
    • fatbinary cubin और PTX को साथ में बांधता है
    • इस RTX 4090 पर SASS actual में execute होता है, लेकिन PTX को अन्य architectures पर driver द्वारा JIT compile किए जा सकने वाले fallback के रूप में शामिल किया जाता है
    • PTX verbose plain text है, इसलिए nvcc default रूप से इसे compress करता है

Host code launch कैसे prepare करता है

  • Compiler frontend cudafe++ main से पहले execute होने वाला hidden constructor insert करता है
    • यह constructor embedded fatbinary को CUDA runtime में register करता है
    • Host-side function pointer vadd और fatbin के अंदर के mangled device kernel name को connect करता है
  • vadd<<<4096, 256>>>(da, db, dc, n) syntax generated host launch stub में बदल जाता है
    • da, db, dc, n host memory के argument buffer में क्रमशः offsets 0, 8, 16, 24 पर align होकर जाते हैं
    • ये offsets SASS द्वारा constant bank 0 से पढ़े जाने वाले 0x160, 0x168, 0x170, 0x178 locations से correspond करते हैं
  • Stub __cudaLaunch call करते हुए host-side dummy vadd function address pass करता है
    • यह address CPU पर execute होने वाले function का address नहीं, बल्कि runtime registration table lookup के key की तरह use होता है
    • Runtime corresponding device symbol name ढूंढने के बाद closed-source user-mode driver libcuda.so.1 में चला जाता है
  • पहले GPU call पर CUDA runtime dynamically libcuda.so.1 खोलता है और context बनाता है
    • strace में /lib/x86_64-linux-gnu/libcuda.so.1 open होते हुए देखा जा सकता है
    • Context में वह channel शामिल होता है जिससे CPU, GPU के साथ communicate करता है
  • CUDA 12.2 से module loading default रूप से lazy है
    • किसी specific kernel के पहली बार launch होने तक SASS cubin upload को defer करता है
    • इसे CUDA_MODULE_LOADING से control किया जा सकता है

GPU को काम भेजने वाली command queue

  • GPU, CPU की तरह function call लेकर entry point पर jump नहीं करता
    • PCIe bus के पार host memory में मौजूद driver command stream को read करता है
    • cuLaunchKernel completed launch command को इस stream में डालता है और GPU को notify करता है
  • पहली execution में driver kernel SASS को GPU memory में copy करता है
    • code buffer allocate करके SASS copy करता है
  • Channel में host RAM में मौजूद दो core structures होते हैं
    • pushbuffer: memory area जहां driver GPU command यानी method लिखता है
    • GPFIFO: pushbuffer span को point करने वाला pointer ring buffer
  • GPFIFO entry pushbuffer span के (base, length) को दिखाने वाले दो 32-bit words से बनी होती है
  • GPU और driver दो cursors से work consumption और production positions track करते हैं
    • GP_GET: GPU कहां तक consume कर चुका है यह दिखाता है
    • GP_PUT: driver कहां तक produce कर चुका है यह दिखाता है
    • दोनों USERD नाम के per-channel structure में होते हैं
  • Kernel launch पर driver pushbuffer span में method लिखता है, GPFIFO entry को उस पर point कराता है, और फिर GP_PUT को आगे बढ़ाता है
  • Modern GPUs में host engine cursor को लगातार monitor नहीं करता, इसलिए doorbell जरूरी होता है
    • GPU process में एक small register window map करता है
    • Driver channel का work-submit token doorbell register में लिखता है
    • Doorbell मिलने के बाद host engine GP_PUT read करता है और GPFIFO entry व pushbuffer span को DMA से fetch करता है

QMD में क्या execution information होती है

  • Launch SET_INLINE_QMD_ADDRESS_A/B और LOAD_INLINE_QMD_DATA method burst से शुरू होता है
  • QMD(Queue Meta Data) compute grid का launch descriptor है
    • grid और block size 4096, 256 शामिल करता है
    • प्रति thread register count और shared memory requirement शामिल करता है
    • program start address और kernel arguments वाली constant bank address शामिल करता है
    • completion notify करने की location भी शामिल करता है
  • Host stub द्वारा packed arguments को driver constant bank में copy करता है, और QMD में उस bank address को record किया जाता है
  • QMD GPU को SASS location, parallel program configuration method, और completion signal location बताता है
  • cuLaunchKernel doorbell बजते ही return कर देता है
    • Call asynchronous है, इसलिए CPU GPU work चलते समय भी execution जारी रख सकता है

SM, warp, occupancy

  • Host engine QMD को compute work distributor को pass करता है
    • यह component पूरे GPU में एक होता है
    • linear SASS instruction stream को SMs में distribute करके parallel program के रूप में execute कराता है
  • Target GPU GeForce RTX 4090 128 SM का उपयोग करता है
    • Launch 4096 blocks और प्रति block 256 threads से बना है
  • हर SM के पास local instruction cache होता है, और active warp program counter maintain करता है
    • Volta के बाद Independent Thread Scheduling model है, जिसमें per-thread program counter और call stack होते हैं
    • Issue अब भी warp unit में होता है
  • Example kernel में resource limit block residency तय करता है
    • प्रति block 256 threads = 8 warps
    • ptxas प्रति thread 16 registers reserve करता है
    • Register के आधार पर प्रति SM 16 blocks संभव हैं
    • Thread capacity प्रति SM 1,536 active threads है, इसलिए 1536 / 256 = 6 blocks ही संभव हैं
    • इसलिए प्रति SM maximum 6 blocks, यानी 48 warps resident state में होते हैं
  • SM 4 processing blocks, यानी sub-partitions में बंटा है
    • 48 resident warps 4 sub-partitions में evenly distribute होते हैं
    • हर warp scheduler full state में 12 active warps manage करता है
    • हर cycle एक eligible warp चुनकर 32 lanes पर अगला instruction dispatch करता है

Warp eligible state में कब होता है

  • GPU CPU के out-of-order execution की तरह single thread से dynamic dependencies बहुत ज्यादा extract नहीं करता
    • कई resident warps रखता है और stall होने पर latency छिपाने के लिए दूसरे warp पर switch करता है
    • Compiler predictable timing को schedule करता है, और hardware scoreboard मुश्किल-से-predict हिस्सों को handle करता है
  • 128-bit SASS instruction में ptxas द्वारा लिखा गया control-code payload होता है
    • fixed-latency instruction में static stall count होता है
    • yield hint बताता है कि scheduler priority छोड़नी है या नहीं
    • variable-latency operation में per-warp physical scoreboard barriers 6 उपयोग किए जाते हैं
  • Example SASS segment में दो LDG.E वही scoreboard barrier B2 set करते हैं
    • FADD के पास B2 wait-on के रूप में होता है
    • जब तक दोनों loads लौटकर barrier clear नहीं करते, वह warp ineligible state में रहता है
    • इस दौरान scheduler उसी sub-partition के दूसरे warp को चुनता है
  • FADD से STG.E पर जाने वाला segment fixed latency से handle होता है
    • FADD में stall=5 होता है, और R9 result ready होने तक warp को कुछ cycles park करता है
    • अलग barrier की जरूरत नहीं होती
  • यह control payload nvdisasm default output में hidden रहता है
    • cuobjdump -sass के raw 128-bit encoding में दूसरे 64-bit word में शामिल होता है
    • layout documented नहीं है, बल्कि microbenchmarking से reconstruct किया गया है

Memory access और performance measurement

  • जब warp LDG.E execute करता है, तो 32 threads अपनी-अपनी address calculate करते हैं
    • Example में consecutive float array access है, इसलिए पूरा warp 32 * 4 = 128 bytes contiguous block request करता है
  • SM load/store unit request coalescing करता है
    • 32 four-byte requests को 4 thirty-two-byte sector requests में merge करता है
    • अगर access contiguous नहीं होता, तो जरूरत से ज्यादा data read हो सकता था
  • Coalesced request पहले SM local L1 Data Cache check करता है
    • miss होने पर crossbar interconnect से 72MB L2 Cache slice तक जाता है
    • L2 में भी miss होने पर memory controller और memory bus से होकर GDDR6X VRAM तक जाता है
  • STG.E store भी principle में उल्टी दिशा में वही path follow करता है
  • Nsight Compute measurements दिखाते हैं कि यह kernel memory-bound है
    • 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
  • Kernel की arithmetic intensity बहुत कम है
    • दो 4-byte loads और एक 4-byte store, यानी कुल 12 bytes transfer पर 1 float add करता है
    • DRAM read के हिसाब से 8.4MB को 10.78μs में read करके लगभग 780GB/s मिलता है, जो peak का करीब 4/5 है
    • 4MB output c 72MB L2 में fit हो जाता है, इसलिए device-to-host copy द्वारा read होने तक DRAM में flush नहीं होता

Result CPU पर वापस कैसे आता है

  • Kernel launch doorbell बजते ही CPU पर return हो जाता है, इसलिए GPU को completion अलग से बतानी पड़ती है
  • सभी 4096 blocks retire होने पर GPU QMD में मौजूद completion semaphore post करता है
    • QMD का fence field words 23–24 में होता है
  • Default stream में cudaMemcpy(c, dc, ...) kernel के बाद रखा जाता है
    • GPU copy engine semaphore set होने तक gated state में रहता है
    • क्योंकि c अभी 72MB L2 में dirty state में है, copy engine read DRAM round-trip के बिना L2 से handle होता है
    • Data PCIe पार करके host memory में move होता है
  • Copy खत्म होने पर copy engine अपना semaphore post करता है
    • Host का cudaMemcpy wait खत्म हो जाता है
    • c फिर से normal host memory बन जाता है
    • printf RAM से c[0] और c[n-1] read करके stdout पर print करता है

Launch के अंदर झांकने का तरीका

  • केवल open kernel modules पढ़कर कुछ behavior directly verify करना मुश्किल है, क्योंकि libcuda closed-source है
  • Method write syscall से नहीं गुजरता और पहले से mapped write-combined buffer में direct लिखा जाता है, इसलिए pushbuffer देखने के लिए memory read करनी होगी
  • LD_PRELOAD shim से mmap को wrap करके /dev/nvidia* से mapped regions record किए जा सकते हैं
    • test program launch के तुरंत बाद shim के dump function को call करे तो mapped pushbuffer print किया जा सकता है
    • dump SET_INLINE_QMD_ADDRESS_A से correspond करने वाला method burst ढूंढता है
  • Pushbuffer method header opcode, payload count, subchannel index, register offset को bit field में रखता है
    • 0x0318 SET_INLINE_QMD_ADDRESS_A है
    • 0x0320 + i * 4 LOAD_INLINE_QMD_DATA(i) है
    • dump में count 66 का increasing-method burst दिखता है, और दो address words व 64 QMD words, कुल 256-byte QMD inline load होता है
    • QMD के अंदर word 12 0x1000, word 18 0x100 है, जो launch के 4096 और 256 से correspond करते हैं
  • Driver setup ioctl से proceed होता है
    • one-kernel program में strace 948 ioctl record करता है
    • इनमें से अधिकतर one-time setup हैं
    • मुख्य file descriptors /dev/nvidiactl और /dev/nvidia-uvm हैं
    • NVIDIA resource manager ioctl magic byte 0x46, यानी 'F' है
    • command number 0x2A को NV_ESC_RM_CONTROL, 0x2B को NV_ESC_RM_ALLOC के रूप में interpret किया जाता है
  • nvcc --keep से generated vadd.cudafe1.stub.c में startup registration code भी देखा जा सकता है
    • __attribute__((__constructor__)) लगा function main से पहले execute होता है
    • __cudaRegisterBinary और __cudaRegisterEntry के जरिए host function pointer vadd और device entry point _Z4vaddPKfS0_Pfi connect होते हैं

संबंधित पढ़ाई

1 टिप्पणियां

 
GN⁺ 3 시간 전
Hacker News की राय

  • लेख दिलचस्प था, और default stream के semaphore की व्याख्या भी मजेदार लगी
    यह बात अच्छी है कि CUDA command synchronization को implicitly संभाल देता है, और parallel commands को streams के जरिए वैकल्पिक रूप से इस्तेमाल करने देता है
    यह Vulkan के उलट है, जो शुरुआत से ही synchronization की सारी जटिलता user पर डाल देता है

  • hardware वाली तरफ कुछ public documentation मौजूद है
    method docs या QMD format खोजने के लिए kernel source पढ़ना जरूरी नहीं है
    https://github.com/NVIDIA/open-gpu-doc/blob/master/classes/c... देखें

  • बहुत उपयोगी था
    खासकर doorbell और QMD वाला हिस्सा सबसे मददगार लगा, क्योंकि उसने दिखाया कि CUDA execution syntax असल में GPU को submit की जाने वाली चीजों से कैसे जुड़ता है
    ज्यादातर explanations kernel, block, warp के आसपास ही रुक जाती हैं, लेकिन यह लेख CPU→driver→GPU path को कहीं ज्यादा आसानी से follow करने लायक बनाता है

  • control code लेख में बताए गए से थोड़ा ज्यादा जटिल है
    असल में यह control word के अंदर bits जैसा कम, table lookup जैसा ज्यादा है

  • आजकल ऐसी कंपनियां हैं जिनका मुख्य काम kernels को optimize करके उन्हें तेज चलाना है
    सोचता हूं कि क्या ऐसी कंपनियां कभी किसी open-source library से पीछे छूट जाएंगी, जो यह काम बेहद अच्छी तरह करती हो
    Nvidia चाहे तो कभी भी ऐसी चीज जारी कर सकता है, ऐसा लगता है
    या फिर बड़े providers inference speed बढ़ाने के moat के रूप में इन कंपनियों को acquire करें, तो उनका भविष्य और बेहतर भी हो सकता है

    • short term में acqui-hire काफी संभव लगता है
      हालांकि kernelbench जैसे related benchmarks में models की प्रगति देखकर लगता है कि ज्यादा generalized solutions भी आखिरकार आएंगे ही
      समस्या यह है कि हर नई hardware generation में अक्सर ऐसी constraints या features आ जाते हैं जिन्हें पुराने models ने देखा नहीं होता
      उदाहरण के लिए Blackwell का tcgen05 भी कभी out-of-distribution case था
      अगर models बेहतर generalize करना शुरू कर दें तो यह घातक बाधा नहीं रह सकती, लेकिन कम से कम अभी यह अब भी अड़चन है
      [1] https://kernelbench.com/
    • CUDA को बड़े scale पर चलाने पर Nvidia driver और library bugs से निपटने में घिनौनी हद तक engineer time लगता है
      मैंने बहुत कम लोगों को Nvidia libraries पर और निर्भर होने की उम्मीद करते देखा है
    • शायद ऐसा नहीं होगा
      क्योंकि workload की details—यानी exact parameters, memory में data representation, values की range वगैरह—optimization strategy को काफी अलग-अलग कर देती हैं

  • अभी-अभी HPC में master's पूरा किया है और CUDA, MPI+CUDA, OpenCL classes ली थीं; classes से पहले ऐसा लेख पढ़ा होता तो बहुत ज्यादा मदद मिलती
    खासकर warp के executable होने का मतलब वाले हिस्से का context अच्छा था

  • सबसे पहले, यह कई कोनों में अच्छी तरह उतरने वाला बढ़िया लेख है
    हालांकि अगर CUDA के runtime API से होकर न जाएं, तो user space की बहुत-सी voodoo जैसी चीजें गायब हो जाती हैं
    driver API इस्तेमाल करें, kernel source को string के रूप में लेकर NVIDIA के runtime compiler से compile करें, तो क्या हो रहा है यह ज्यादा साफ दिखता है
    सब कुछ नहीं, लेकिन काफी कुछ transparent हो जाता है
    ज्यादा “primitive” version यहां है:
    https://github.com/NVIDIA/cuda-samples/tree/master/cpp/0_Int...
    इसी चीज को ज्यादा readable, फिर भी पूरी तरह transparent modern C++ API रूप में देखने के लिए यह देखें:
    https://github.com/eyalroz/cuda-api-wrappers/blob/master/exa...
    यह मेरी CUDA API wrappers header-only library का example program है

    • driver API अच्छी है क्योंकि इससे CUDA kernels को hot-reloadable shaders की तरह treat किया जा सकता है
      चलते समय code बदलते हुए develop कर सकते हैं, इसलिए मजेदार है

  • bare metal पर?