Yang Terjadi di Dalam Saat Menjalankan Kernel CUDA

 (fergusfinn.com)
2 poin oleh GN⁺ 3 jam lalu | 1 komentar | Bagikan ke WhatsApp
  • Bahkan program CUDA penjumlahan vektor sederhana pun, sebelum mendapatkan hasil 2.000000, melewati pipeline kompilasi, pemanggilan driver, antrean perintah GPU, penjadwalan warp, hierarki memori, dan semaphore penyelesaian
  • nvcc memisahkan kode host dan kode device, membuat PTX dengan cicc, SASS dengan ptxas, lalu membundel cubin dan PTX ke dalam fatbin dan memasukkannya ke file executable Linux
  • Sintaks launch vadd<<<4096, 256>>> diubah menjadi host launch stub, dan argumen da, db, dc, n diteruskan ke driver melalui CUDA runtime dan libcuda.so.1
  • Eksekusi GPU dimulai dengan QMD, pushbuffer, GPFIFO, GP_PUT, serta penulisan MMIO doorbell; 128 SM pada RTX 4090 menjalankan konfigurasi 4096 blok dan 256 thread dalam satuan warp
  • Karena kernel ini memiliki intensitas aritmetika rendah, yaitu membutuhkan transfer 12 byte untuk setiap 1 operasi penjumlahan float, di Nsight Compute ia bergantung pada bandwidth memori: 10,78 μs, 79,65% dari puncak DRAM, dan warp issue 5,17%

Kernel contoh dan cakupan pengamatan

  • Program contoh menggunakan kernel CUDA vadd untuk menjumlahkan dua array float dan menyimpan hasilnya ke array ketiga
    • Dengan n = 1 << 20, program memproses 1.048.576 nilai float
    • Konfigurasi launch adalah vadd<<<4096, 256>>>(da, db, dc, n), menggunakan 4096 * 256 = n thread
  • Jika dikompilasi untuk RTX 4090 dengan nvcc -arch=sm_89 dan dijalankan, outputnya adalah c[0]=2.000000 c[n-1]=2.000000
  • Bahkan untuk satu baris hasil ini, puluhan juta instruksi CPU, device file, sekitar 900 ioctl, dan register doorbell yang dipetakan ke memori ikut terlibat

Proses nvcc membuat file executable

  • Dengan nvcc --keep, artefak dari pipeline kompilasi dapat diperiksa langsung
    • vadd.ptx: PTX dari kode device yang dibuat oleh cicc
    • vadd.sm_89.cubin: SASS dari kode device yang dibuat oleh ptxas
    • vadd.fatbin: fatbin yang membundel cubin dan PTX
    • vadd.cudafe1.stub.c: host launch stub dan kode registrasi kernel
    • vadd.o: objek host akhir yang berisi fatbin
  • Kode host diproses oleh compiler host, sedangkan kernel device vadd melewati tahap cicc dan ptxas
  • PTX adalah ISA virtual, menggunakan register virtual tak terbatas yang bertipe, dan tidak merefleksikan langsung jumlah register hardware yang sebenarnya
    • PTX pada contoh mencakup perhitungan blockIdx.x * blockDim.x + threadIdx.x, pemeriksaan batas, global load, float add, dan global store
    • Pointer CUDA secara default adalah generic pointer, sehingga dikonversi terlebih dahulu menjadi global address dengan cvta.to.global, lalu menggunakan ld.global
    • mul.wide.s32 mengubah index menjadi offset satuan 4 byte, yaitu sizeof(float), dan memperluasnya dari 32 bit ke 64 bit
  • SASS adalah instruksi nyata yang spesifik untuk arsitektur, dan pada output untuk RTX 4090 tampil lebih ringkas dibanding PTX
    • S2R menyalin special register seperti SR_CTAID.X dan SR_TID.X ke register umum
    • Kombinasi mul.wide dan add pada PTX digabungkan menjadi IMAD.WIDE di SASS
    • Konversi cvta diserap ke dalam proses pengalamatan
  • Operand c[0x0][...] menunjuk ke constant bank 0 yang dikelola driver
    • Pointer a, b, c berada di 0x160, 0x168, 0x170
    • n berada di 0x178
    • Launch geometry seperti blockDim.x dan nilai ABI juga berada di bank yang sama
  • cubin adalah file ELF, format kontainer yang sama dengan file executable Linux
    • fatbinary membundel cubin dan PTX bersama-sama
    • Pada RTX 4090 ini, SASS yang benar-benar dijalankan, tetapi PTX disertakan sebagai fallback agar driver dapat melakukan kompilasi JIT di arsitektur lain
    • Karena PTX berupa plain text yang verbose, nvcc mengompresinya secara default

Cara kode host menyiapkan launch

  • Frontend compiler cudafe++ menyisipkan constructor tersembunyi yang dijalankan sebelum main
    • Constructor ini mendaftarkan embedded fatbinary ke CUDA runtime
    • Ia menghubungkan function pointer vadd di sisi host dengan nama device kernel yang sudah di-mangle di dalam fatbin
  • Sintaks vadd<<<4096, 256>>>(da, db, dc, n) diubah menjadi host launch stub yang dihasilkan
    • da, db, dc, n ditempatkan di argument buffer pada host memory, masing-masing disejajarkan pada offset 0, 8, 16, 24
    • Offset ini bersesuaian dengan lokasi 0x160, 0x168, 0x170, 0x178 di constant bank 0 yang dibaca oleh SASS
  • Stub memanggil __cudaLaunch sambil meneruskan alamat fungsi dummy vadd di sisi host
    • Alamat ini bukan alamat fungsi yang akan dijalankan di CPU, melainkan dipakai sebagai key untuk mencari tabel registrasi runtime
    • Runtime menemukan device symbol name yang sesuai, lalu berpindah ke libcuda.so.1, user-mode driver closed-source
  • Pada pemanggilan GPU pertama, CUDA runtime membuka libcuda.so.1 secara dinamis dan membuat context
    • Di strace, dapat terlihat bahwa /lib/x86_64-linux-gnu/libcuda.so.1 dibuka
    • Context mencakup channel yang digunakan CPU untuk berkomunikasi dengan GPU
  • Sejak CUDA 12.2, module loading secara default bersifat lazy
    • Upload cubin SASS ditunda sampai kernel tertentu pertama kali di-launch
    • Perilaku ini dapat dikendalikan dengan CUDA_MODULE_LOADING

Antrean perintah yang mengirim pekerjaan ke GPU

  • GPU tidak menerima function call seperti CPU lalu melakukan jump ke entry point
    • Ia membaca driver command stream di dalam host memory melalui bus PCIe
    • cuLaunchKernel memasukkan launch command yang sudah lengkap ke stream ini dan memberi tahu GPU
  • Pada eksekusi pertama, driver menyalin SASS kernel ke memori GPU
    • Driver mengalokasikan code buffer dan menyalin SASS ke sana
  • Channel memiliki dua struktur inti yang berada di host RAM
    • pushbuffer: area memori tempat driver menulis method, yaitu command GPU
    • GPFIFO: pointer ring buffer yang menunjuk ke rentang pushbuffer
  • Entry GPFIFO terdiri dari dua word 32 bit yang menunjukkan (base, length) dari rentang pushbuffer
  • GPU dan driver melacak posisi konsumsi dan produksi pekerjaan dengan dua cursor
    • GP_GET: menunjukkan sampai mana GPU sudah mengonsumsi pekerjaan
    • GP_PUT: menunjukkan sampai mana driver sudah memproduksi pekerjaan
    • Keduanya berada dalam struktur per-channel bernama USERD
  • Saat kernel launch, driver menulis method ke rentang pushbuffer, membuat entry GPFIFO menunjuk ke sana, lalu memajukan GP_PUT
  • Pada GPU modern, host engine tidak terus-menerus memantau cursor, sehingga doorbell diperlukan
    • GPU memetakan sebuah register window kecil ke process
    • Driver menulis work-submit token milik channel ke register doorbell
    • Setelah menerima doorbell, host engine membaca GP_PUT dan mengambil entry GPFIFO serta rentang pushbuffer melalui DMA

Informasi eksekusi yang dimuat QMD

  • Launch dimulai dengan burst method SET_INLINE_QMD_ADDRESS_A/B dan LOAD_INLINE_QMD_DATA
  • QMD(Queue Meta Data) adalah launch descriptor untuk compute grid
    • Memuat ukuran grid dan block, yaitu 4096 dan 256
    • Memuat jumlah register per thread dan kebutuhan shared memory
    • Memuat alamat awal program dan alamat constant bank yang berisi argumen kernel
    • Memuat juga lokasi untuk memberi tahu penyelesaian
  • Argumen yang dipaketkan oleh host stub disalin driver ke constant bank, dan alamat bank tersebut dicatat di QMD
  • QMD memberi tahu GPU lokasi SASS, cara membentuk program paralel, dan lokasi signal penyelesaian
  • cuLaunchKernel kembali begitu doorbell berbunyi
    • Karena pemanggilannya asynchronous, CPU dapat terus berjalan sementara pekerjaan GPU berlangsung

SM, warp, dan occupancy

  • Host engine meneruskan QMD ke compute work distributor
    • Komponen ini ada satu untuk seluruh GPU
    • Ia mendistribusikan linear SASS instruction stream ke SM agar dijalankan sebagai program paralel
  • GPU target, GeForce RTX 4090, menggunakan 128 SM
    • Launch terdiri dari 4096 block dan 256 thread per block
  • Setiap SM memiliki local instruction cache, dan active warp mempertahankan program counter
    • Sejak Volta, ada model Independent Thread Scheduling dengan program counter dan call stack per thread
    • Namun issue tetap dilakukan dalam satuan warp
  • Pada kernel contoh, resource limit menentukan block residency
    • 256 threads = 8 warps per block
    • ptxas mencadangkan 16 register per thread
    • Berdasarkan register, bisa ada 16 block per SM
    • Kapasitas thread adalah 1.536 active thread per SM, sehingga hanya 1536 / 256 = 6 block yang mungkin
    • Karena itu, maksimum 6 block per SM, atau 48 warp, berada dalam keadaan resident
  • SM dibagi menjadi 4 processing block, yaitu sub-partition
    • 48 resident warp dibagi merata ke 4 sub-partition
    • Dalam kondisi penuh, setiap warp scheduler mengelola 12 active warp
    • Setiap cycle, scheduler memilih satu eligible warp dan men-dispatch instruksi berikutnya ke 32 lane

Syarat sebuah warp menjadi eligible

  • GPU tidak mengekstraksi dependensi dinamis secara besar-besaran dari satu thread seperti eksekusi out-of-order pada CPU
    • Ia menyembunyikan latency dengan menyediakan banyak resident warp dan beralih ke warp lain saat terjadi stall
    • Compiler menjadwalkan timing yang dapat diprediksi, sedangkan hardware scoreboard menangani bagian yang sulit diprediksi
  • Instruksi SASS 128 bit berisi control-code payload yang ditulis oleh ptxas
    • Untuk instruksi fixed-latency, terdapat static stall count
    • Yield hint memberi tahu apakah scheduler harus menyerahkan prioritas
    • Untuk operasi variable-latency, digunakan 6 per-warp physical scoreboard barrier
  • Pada bagian SASS contoh, dua LDG.E menyetel scoreboard barrier yang sama, B2
    • FADD memiliki B2 sebagai wait-on
    • Sampai kedua load kembali dan barrier tersebut clear, warp itu berada dalam keadaan ineligible
    • Selama itu, scheduler memilih warp lain dari sub-partition yang sama
  • Perpindahan dari FADD ke STG.E ditangani sebagai fixed latency
    • FADD memiliki stall=5, sehingga warp diparkir beberapa cycle sampai hasil R9 siap
    • Barrier terpisah tidak diperlukan
  • Control payload ini disembunyikan dalam output default nvdisasm
    • Dalam raw 128-bit encoding dari cuobjdump -sass, ia berada di word 64 bit kedua
    • Layout-nya tidak terdokumentasi, melainkan direkonstruksi melalui microbenchmarking

Akses memori dan pengukuran performa

  • Saat warp menjalankan LDG.E, 32 thread masing-masing menghitung alamat
    • Contoh ini mengakses array float berurutan, sehingga seluruh warp meminta blok kontigu 32 * 4 = 128 bytes
  • Unit load/store SM melakukan request coalescing
    • Ia menggabungkan 32 request 4 byte menjadi 4 sector request berukuran 32 byte
    • Jika aksesnya tidak berurutan, data yang dibaca bisa lebih banyak dari yang diperlukan
  • Request yang sudah di-coalesce terlebih dahulu memeriksa SM local L1 Data Cache
    • Jika miss, request melewati crossbar interconnect menuju slice L2 Cache 72 MB
    • Jika di L2 juga miss, request melewati memory controller dan memory bus menuju VRAM GDDR6X
  • Store STG.E pada prinsipnya juga mengikuti jalur yang sama dengan arah berlawanan
  • Hasil pengukuran Nsight Compute menunjukkan kernel ini 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 ini memiliki intensitas aritmetika yang sangat rendah
    • Ia melakukan 1 operasi float add untuk total transfer 12 byte: dua load 4 byte dan satu store 4 byte
    • Dari sisi DRAM read, 8,4 MB dibaca dalam 10,78 μs, sekitar 780 GB/s, kira-kira 4/5 dari puncak
    • Output c sebesar 4 MB muat di L2 72 MB, sehingga tidak di-flush ke DRAM sampai device-to-host copy membacanya

Proses hasil kembali ke CPU

  • Karena kernel launch kembali ke CPU begitu doorbell dibunyikan, GPU harus memberitahukan penyelesaian secara terpisah
  • Setelah semua 4096 block retire, GPU melakukan post ke completion semaphore yang ada di QMD
    • Field fence pada QMD berada di word 23–24
  • Pada default stream, cudaMemcpy(c, dc, ...) ditempatkan setelah kernel
    • GPU copy engine berada dalam keadaan gated sampai semaphore naik
    • Karena c masih dirty di L2 72 MB, pembacaan oleh copy engine dilayani dari L2 tanpa perjalanan pulang-pergi ke DRAM
    • Data berpindah melalui PCIe ke host memory
  • Setelah copy selesai, copy engine melakukan post ke semaphore-nya sendiri
    • Penantian cudaMemcpy di host berakhir
    • c kembali menjadi host memory biasa
    • printf membaca c[0] dan c[n-1] dari RAM lalu mencetaknya ke stdout

Cara mengintip bagian dalam launch

  • Membaca open kernel modules saja tidak cukup untuk memastikan sebagian perilaku secara langsung karena libcuda bersifat closed-source
  • Method write tidak melewati syscall dan ditulis langsung ke write-combined buffer yang sudah dipetakan, sehingga untuk melihat pushbuffer perlu membaca memori
  • Dengan LD_PRELOAD shim untuk membungkus mmap, area yang dipetakan dari /dev/nvidia* dapat dicatat
    • Jika test program memanggil fungsi dump milik shim tepat setelah launch, pushbuffer yang dipetakan dapat dicetak
    • Dump mencari method burst yang sesuai dengan SET_INLINE_QMD_ADDRESS_A
  • Header method pushbuffer menyimpan opcode, payload count, subchannel index, dan register offset sebagai bit field
    • 0x0318 adalah SET_INLINE_QMD_ADDRESS_A
    • 0x0320 + i * 4 adalah LOAD_INLINE_QMD_DATA(i)
    • Dalam dump terlihat increasing-method burst dengan count 66; dua address word dan 64 QMD word, total QMD 256 byte, dimuat secara inline
    • Word 12 di dalam QMD adalah 0x1000, dan word 18 adalah 0x100, yang sesuai dengan 4096 dan 256 pada launch
  • Setup driver dilakukan dengan ioctl
    • Pada program one-kernel, strace mencatat 948 ioctl
    • Sebagian besar adalah setup sekali jalan
    • File descriptor utama adalah /dev/nvidiactl dan /dev/nvidia-uvm
    • Magic byte ioctl resource manager NVIDIA adalah 0x46, yaitu 'F'
    • Command number 0x2A ditafsirkan sebagai NV_ESC_RM_CONTROL, dan 0x2B sebagai NV_ESC_RM_ALLOC
  • Di vadd.cudafe1.stub.c yang dihasilkan oleh nvcc --keep, kode registrasi startup juga dapat dilihat
    • Fungsi dengan __attribute__((__constructor__)) dijalankan sebelum main
    • Melalui __cudaRegisterBinary dan __cudaRegisterEntry, function pointer host vadd dihubungkan dengan device entry point _Z4vaddPKfS0_Pfi

Bacaan terkait

1 komentar

 
GN⁺ 3 jam lalu
Komentar Hacker News

  • Tulisan yang menarik, dan penjelasan tentang semaphore pada stream default juga seru
    Saya suka bahwa CUDA menangani sinkronisasi perintah secara implisit, dan memungkinkan perintah paralel dipakai secara selektif lewat stream
    Ini kontras dengan Vulkan, yang sejak awal melempar seluruh kompleksitas sinkronisasi kepada pengguna

  • Di sisi hardware ada beberapa dokumen publik
    Tidak perlu selalu membaca source kernel untuk mencari dokumentasi method atau format QMD
    Lihat https://github.com/NVIDIA/open-gpu-doc/blob/master/classes/c...

  • Sangat berguna
    Terutama bagian doorbell dan QMD paling membantu karena menunjukkan bagaimana sintaks eksekusi CUDA terhubung dengan apa yang benar-benar dikirim ke GPU
    Kebanyakan penjelasan berhenti di sekitar kernel, block, dan warp, tetapi tulisan ini membuat jalur CPU→driver→GPU jauh lebih mudah diikuti

  • Kode kontrolnya sedikit lebih rumit daripada yang dijelaskan di tulisan
    Dalam praktiknya, ini lebih mirip lookup tabel daripada bit di dalam control word

  • Sekarang ada perusahaan-perusahaan yang pekerjaan utamanya mengoptimalkan kernel agar berjalan lebih cepat
    Saya penasaran apakah suatu hari perusahaan seperti itu akan tersisih oleh library open source yang melakukannya dengan sangat baik
    Nvidia tampaknya bisa saja merilis hal seperti itu kapan saja
    Atau bisa juga hasilnya justru lebih baik jika penyedia besar mengakuisisi perusahaan-perusahaan ini untuk menjadikannya moat dalam mempercepat inferensi

    • Dalam jangka pendek, acqui-hire tampaknya cukup mungkin
      Namun melihat perkembangan model pada benchmark terkait seperti kernelbench, saya rasa solusi yang lebih general pada akhirnya pasti akan muncul juga
      Masalahnya, setiap generasi hardware baru sering menghadirkan batasan atau fitur yang belum pernah dilihat model lama
      Misalnya, tcgen05 di Blackwell dulunya merupakan kasus di luar distribusi
      Jika model mulai bisa menggeneralisasi dengan lebih baik, ini mungkin bukan penghalang fatal, tetapi setidaknya untuk saat ini masih menjadi ganjalan
      [1] https://kernelbench.com/
    • Menjalankan CUDA dalam skala besar menghabiskan waktu engineer dalam jumlah yang menjijikkan banyaknya untuk menangani bug driver dan library Nvidia
      Saya jarang melihat orang berharap menjadi lebih bergantung pada library Nvidia
    • Sepertinya tidak
      Karena detail workload—yakni parameter yang tepat, representasi data di memori, rentang nilai, dan sebagainya—sangat memengaruhi strategi optimasi

  • Saya baru saja menyelesaikan S2 HPC dan mengambil kelas CUDA, MPI+CUDA, serta OpenCL; rasanya tulisan seperti ini akan jauh lebih membantu kalau saya membacanya sebelum kelas
    Saya terutama suka konteks di sekitar bagian tentang apa artinya sebuah warp bisa dieksekusi

  • Pertama-tama, ini tulisan bagus yang menggali banyak sudut secara mendalam
    Namun jika tidak melewati runtime API CUDA, banyak bagian voodoo di user space akan hilang
    Dengan memakai driver API dan mengompilasi source kernel yang diterima sebagai string menggunakan runtime compiler NVIDIA, kita bisa melihat lebih jelas apa yang terjadi
    Tidak semuanya, tetapi cukup banyak bagian menjadi transparan
    Versi yang lebih “mentah” ada di sini:
    https://github.com/NVIDIA/cuda-samples/tree/master/cpp/0_Int...
    Jika ingin melihat hal yang sama dalam bentuk API C++ modern yang jauh lebih mudah dibaca tetapi tetap sepenuhnya transparan, lihat ini:
    https://github.com/eyalroz/cuda-api-wrappers/blob/master/exa...
    Ini adalah program contoh dari library header-only CUDA API wrappers saya

    • Driver API bagus karena memungkinkan CUDA kernel diperlakukan seperti shader yang bisa di-hot reload
      Menyenangkan karena kita bisa mengembangkan sambil mengubah kode saat program berjalan

  • Di bare metal?