Điều gì xảy ra bên trong khi chạy một kernel CUDA

 (fergusfinn.com)
2 điểm bởi GN⁺ 3 giờ trước | 1 bình luận | Chia sẻ qua WhatsApp
  • Ngay cả một chương trình CUDA cộng vector đơn giản cũng phải đi qua pipeline biên dịch, các lời gọi driver, hàng đợi lệnh GPU, lập lịch warp, phân cấp bộ nhớ và semaphore hoàn tất trước khi thu được kết quả 2.000000
  • nvcc tách mã host và mã device, tạo PTX bằng cicc, tạo SASS bằng ptxas, rồi đóng gói cubin và PTX vào fatbin và đưa vào trong tệp thực thi Linux
  • Cú pháp launch vadd<<<4096, 256>>> được chuyển thành host launch stub, còn các đối số da, db, dc, n được truyền tới driver thông qua CUDA runtime và libcuda.so.1
  • Việc thực thi trên GPU bắt đầu bằng QMD, pushbuffer, GPFIFO, GP_PUT, thao tác ghi MMIO doorbell; 128 SM của RTX 4090 thực thi cấu hình 4096 block và 256 thread theo đơn vị warp
  • Kernel này bị chi phối bởi băng thông bộ nhớ do mật độ tính toán thấp, cần truyền 12 byte cho mỗi phép cộng float; Nsight Compute ghi nhận 10.78μs, 79.65% đỉnh DRAM, warp issue 5.17%

Kernel ví dụ và phạm vi quan sát

  • Chương trình ví dụ dùng kernel CUDA vadd để cộng hai mảng float và lưu vào mảng thứ ba
    • Xử lý 1.048.576 float với n = 1 << 20
    • Cấu hình launch là vadd<<<4096, 256>>>(da, db, dc, n), dùng 4096 * 256 = n thread
  • Khi biên dịch cho RTX 4090 bằng nvcc -arch=sm_89 và chạy, chương trình in ra c[0]=2.000000 c[n-1]=2.000000
  • Ngay cả kết quả một dòng này cũng có sự tham gia của hàng chục triệu lệnh CPU, device file, khoảng 900 lời gọi ioctl và thanh ghi doorbell được memory-map

Cách nvcc tạo tệp thực thi

  • Dùng nvcc --keep có thể xem trực tiếp các artifact của pipeline biên dịch
    • vadd.ptx: PTX của mã device do cicc tạo
    • vadd.sm_89.cubin: SASS của mã device do ptxas tạo
    • vadd.fatbin: fatbin đóng gói cubin và PTX
    • vadd.cudafe1.stub.c: host launch stub và mã đăng ký kernel
    • vadd.o: object host cuối cùng có chứa fatbin
  • Mã host được xử lý bởi trình biên dịch host, còn kernel device vadd đi qua các bước cicc và ptxas
  • PTX là ISA ảo, dùng vô số thanh ghi ảo có kiểu và không phản ánh trực tiếp số thanh ghi phần cứng thực tế
    • PTX ví dụ bao gồm phép tính blockIdx.x * blockDim.x + threadIdx.x, kiểm tra biên, global load, float add và global store
    • Con trỏ CUDA mặc định là generic pointer, nên được chuyển sang global address bằng cvta.to.global rồi dùng ld.global
    • mul.wide.s32 chuyển index thành offset theo đơn vị 4 byte là sizeof(float) và mở rộng từ 32 bit lên 64 bit
  • SASS là lệnh thực tế theo từng kiến trúc; trong output cho RTX 4090, nó xuất hiện dưới dạng cô đọng hơn PTX
    • S2R sao chép các thanh ghi đặc biệt như SR_CTAID.X, SR_TID.X sang thanh ghi thông thường
    • Tổ hợp mul.wide và add của PTX được gộp thành IMAD.WIDE trong SASS
    • Phép chuyển đổi cvta được hấp thụ vào quá trình định địa chỉ
  • Toán hạng c[0x0][...] trỏ tới constant bank 0 do driver quản lý
    • Các con trỏ a, b, c nằm tại 0x160, 0x168, 0x170
    • n nằm tại 0x178
    • Launch geometry như blockDim.x và các giá trị ABI cũng nằm trong cùng bank
  • cubin là tệp ELF, cùng định dạng container với tệp thực thi Linux
    • fatbinary đóng gói cubin và PTX cùng nhau
    • Trên RTX 4090 này, SASS được thực thi thực tế, nhưng PTX được đưa vào làm fallback để driver có thể JIT compile trên kiến trúc khác
    • PTX là plain text dài dòng nên nvcc mặc định nén nó

Cách mã host chuẩn bị launch

  • Frontend trình biên dịch cudafe++ chèn một constructor ẩn chạy trước main
    • Constructor này đăng ký fatbinary được nhúng với CUDA runtime
    • Nó liên kết function pointer phía host vadd với tên device kernel đã bị mangle trong fatbin
  • Cú pháp vadd<<<4096, 256>>>(da, db, dc, n) được chuyển thành host launch stub được sinh ra
    • da, db, dc, n được căn chỉnh và đưa vào argument buffer trong host memory lần lượt ở các offset 0, 8, 16, 24
    • Các offset này tương ứng với vị trí 0x160, 0x168, 0x170, 0x178 mà SASS đọc trong constant bank 0
  • Stub gọi __cudaLaunch và truyền địa chỉ hàm dummy vadd phía host
    • Địa chỉ này không phải địa chỉ hàm để CPU thực thi, mà được dùng làm key để tra bảng đăng ký của runtime
    • Runtime tìm device symbol name tương ứng rồi chuyển sang libcuda.so.1, user-mode driver mã nguồn đóng
  • Ở lần gọi GPU đầu tiên, CUDA runtime mở động libcuda.so.1 và tạo context
    • Trong strace có thể thấy /lib/x86_64-linux-gnu/libcuda.so.1 được mở
    • Context bao gồm một channel để CPU giao tiếp với GPU
  • Từ CUDA 12.2, module loading mặc định là lazy
    • Việc upload SASS cubin được trì hoãn cho đến khi một kernel cụ thể được launch lần đầu
    • Có thể điều khiển bằng CUDA_MODULE_LOADING

Hàng đợi lệnh truyền công việc cho GPU

  • GPU không nhận lời gọi hàm rồi jump tới entry point như CPU
    • Nó đọc driver command stream trong host memory qua bus PCIe
    • cuLaunchKernel đưa launch command đã hoàn chỉnh vào stream này và báo cho GPU
  • Ở lần chạy đầu, driver sao chép SASS của kernel vào bộ nhớ GPU
    • Cấp phát code buffer và sao chép SASS
  • Channel có hai cấu trúc cốt lõi nằm trong host RAM
    • pushbuffer: vùng nhớ nơi driver ghi các method là lệnh GPU
    • GPFIFO: ring buffer con trỏ trỏ tới các span của pushbuffer
  • Một GPFIFO entry gồm hai word 32 bit biểu diễn (base, length) của span pushbuffer
  • GPU và driver theo dõi vị trí tiêu thụ và sản xuất công việc bằng hai cursor
    • GP_GET: cho biết GPU đã tiêu thụ đến đâu
    • GP_PUT: cho biết driver đã sản xuất đến đâu
    • Cả hai đều nằm trong cấu trúc per-channel gọi là USERD
  • Khi launch kernel, driver ghi các method vào span pushbuffer, để GPFIFO entry trỏ tới nó rồi tăng GP_PUT
  • Trên GPU hiện đại, host engine không liên tục giám sát cursor nên cần doorbell
    • GPU map một cửa sổ thanh ghi nhỏ vào process
    • Driver ghi work-submit token của channel vào thanh ghi doorbell
    • Sau khi nhận doorbell, host engine đọc GP_PUT rồi lấy GPFIFO entry và span pushbuffer bằng DMA

Thông tin thực thi chứa trong QMD

  • Launch bắt đầu bằng một burst method SET_INLINE_QMD_ADDRESS_A/B và LOAD_INLINE_QMD_DATA
  • QMD(Queue Meta Data) là launch descriptor của compute grid
    • Bao gồm kích thước grid và block là 4096, 256
    • Bao gồm số thanh ghi mỗi thread và yêu cầu shared memory
    • Bao gồm địa chỉ bắt đầu chương trình và địa chỉ constant bank chứa đối số kernel
    • Cũng bao gồm vị trí để báo hoàn tất
  • Các đối số do host stub đóng gói được driver sao chép vào constant bank, và địa chỉ bank đó được ghi trong QMD
  • QMD cho GPU biết vị trí SASS, cách cấu hình chương trình song song và vị trí tín hiệu hoàn tất
  • cuLaunchKernel trả về ngay khi doorbell được rung
    • Lời gọi là bất đồng bộ, nên CPU có thể tiếp tục chạy trong khi công việc GPU đang diễn ra

SM, warp và occupancy

  • Host engine chuyển QMD cho compute work distributor
    • Thành phần này tồn tại một lần trên toàn GPU
    • Nó phân phối linear SASS instruction stream tới các SM để thực thi như một chương trình song song
  • GPU mục tiêu GeForce RTX 4090 dùng 128 SM
    • Launch gồm 4096 block và 256 thread mỗi block
  • Mỗi SM có instruction cache cục bộ, còn active warp giữ program counter
    • Từ Volta trở đi có mô hình Independent Thread Scheduling với program counter và call stack theo từng thread
    • Việc issue vẫn diễn ra theo đơn vị warp
  • Trong kernel ví dụ, resource limit quyết định block residency
    • 256 threads = 8 warps mỗi block
    • ptxas dành 16 thanh ghi mỗi thread
    • Theo tiêu chí thanh ghi, có thể có 16 block mỗi SM
    • Dung lượng thread là 1.536 active thread mỗi SM, nên chỉ có 1536 / 256 = 6 block
    • Vì vậy tối đa 6 block mỗi SM, tức 48 warp ở trạng thái resident
  • SM được chia thành 4 processing block, tức sub-partition
    • 48 resident warp được phân phối đều cho 4 sub-partition
    • Mỗi warp scheduler quản lý 12 active warp khi đầy
    • Mỗi cycle, nó chọn một eligible warp và dispatch lệnh tiếp theo tới 32 lane

Điều kiện để warp ở trạng thái eligible

  • GPU không trích xuất nhiều phụ thuộc động từ một thread đơn như thực thi out-of-order của CPU
    • Nó đặt nhiều resident warp; khi xảy ra stall thì chuyển sang warp khác để che giấu latency
    • Compiler lập lịch timing có thể dự đoán, còn hardware scoreboard xử lý phần khó dự đoán
  • Lệnh SASS 128 bit chứa control-code payload do ptxas ghi
    • Lệnh có fixed latency chứa static stall count
    • Yield hint cho biết có nhường priority scheduler hay không
    • Thao tác variable latency dùng 6 physical scoreboard barrier theo từng warp
  • Trong đoạn SASS ví dụ, hai LDG.E set cùng scoreboard barrier B2
    • FADD có B2 là wait-on
    • Warp đó ở trạng thái ineligible cho đến khi hai load quay về và barrier được clear
    • Trong thời gian đó, scheduler chọn warp khác trong cùng sub-partition
  • Đoạn chuyển từ FADD sang STG.E được xử lý bằng fixed latency
    • FADD có stall=5, park warp trong vài cycle cho đến khi kết quả R9 sẵn sàng
    • Không cần barrier riêng
  • Control payload này bị ẩn trong output mặc định của nvdisasm
    • Nó nằm trong word 64 bit thứ hai của raw 128-bit encoding từ cuobjdump -sass
    • Layout không được tài liệu hóa mà được tái dựng bằng microbenchmarking

Truy cập bộ nhớ và đo hiệu năng

  • Khi warp thực thi LDG.E, 32 thread lần lượt tính địa chỉ của mình
    • Ví dụ truy cập mảng float liên tiếp, nên toàn bộ warp yêu cầu một khối liên tục 32 * 4 = 128 bytes
  • Load/store unit của SM thực hiện request coalescing
    • Gộp 32 yêu cầu 4 byte thành 4 sector request 32 byte
    • Nếu truy cập không liên tiếp, có thể phải đọc nhiều dữ liệu hơn mức cần thiết
  • Coalesced request trước hết kiểm tra L1 Data Cache cục bộ của SM
    • Nếu miss, nó đi qua crossbar interconnect tới slice L2 Cache 72MB
    • Nếu vẫn miss ở L2, nó đi qua memory controller và memory bus tới GDDR6X VRAM
  • Store STG.E về nguyên tắc đi theo cùng đường ngược lại
  • Các chỉ số Nsight Compute cho thấy kernel này bị 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 có mật độ tính toán rất thấp
    • Thực hiện 1 phép cộng float cho mỗi tổng cộng 12 byte truyền tải: hai load 4 byte và một store 4 byte
    • Xét riêng DRAM read, đọc 8,4MB trong 10.78μs, tương đương khoảng 780GB/s, xấp xỉ 4/5 mức đỉnh
    • Output c 4MB nằm vừa trong L2 72MB, nên không bị flush xuống DRAM cho đến khi device-to-host copy đọc nó

Quá trình kết quả quay về CPU

  • Kernel launch trả về CPU ngay khi doorbell được rung, nên GPU phải thông báo hoàn tất riêng
  • Khi cả 4096 block đều retire, GPU post completion semaphore chứa trong QMD
    • Fence field của QMD nằm ở các word 23–24
  • Trong default stream, cudaMemcpy(c, dc, ...) được đặt sau kernel
    • GPU copy engine bị gated cho đến khi semaphore được bật
    • Vì c vẫn đang dirty trong L2 72MB, lần đọc của copy engine được xử lý từ L2 mà không cần vòng đi-về DRAM
    • Dữ liệu di chuyển qua PCIe sang host memory
  • Khi copy kết thúc, copy engine post semaphore của chính nó
    • Việc chờ cudaMemcpy ở host kết thúc
    • c trở lại là host memory thông thường
    • printf đọc c[0] và c[n-1] từ RAM rồi in ra stdout

Cách nhìn vào bên trong launch

  • Chỉ đọc open kernel modules thì khó xác nhận trực tiếp một số hành vi vì libcuda là mã nguồn đóng
  • Method write không đi qua syscall mà được ghi trực tiếp vào write-combined buffer đã được mapping, nên muốn xem pushbuffer thì phải đọc memory
  • Có thể dùng shim LD_PRELOAD để bọc mmap và ghi lại các vùng được mapping từ /dev/nvidia*
    • Nếu test program gọi hàm dump của shim ngay sau launch, nó có thể in pushbuffer đã map
    • Dump tìm burst method tương ứng với SET_INLINE_QMD_ADDRESS_A
  • Header method của pushbuffer chứa opcode, payload count, subchannel index và register offset dưới dạng bit field
    • 0x0318 là SET_INLINE_QMD_ADDRESS_A
    • 0x0320 + i * 4 là LOAD_INLINE_QMD_DATA(i)
    • Trong dump có thể thấy burst increasing-method với count 66; hai address word và 64 QMD word, tổng cộng QMD 256 byte, được đưa inline
    • Word 12 trong QMD là 0x1000, word 18 là 0x100, tương ứng với 4096 và 256 của launch
  • Driver setup được thực hiện bằng ioctl
    • Với chương trình một kernel, strace ghi nhận 948 lời gọi ioctl
    • Phần lớn là setup một lần
    • Các file descriptor chính là /dev/nvidiactl và /dev/nvidia-uvm
    • Magic byte ioctl của NVIDIA resource manager là 0x46, tức 'F'
    • Command number 0x2A được diễn giải là NV_ESC_RM_CONTROL, còn 0x2B là NV_ESC_RM_ALLOC
  • Trong vadd.cudafe1.stub.c do nvcc --keep tạo ra cũng có thể xem mã đăng ký lúc khởi động
    • Hàm có gắn __attribute__((__constructor__)) chạy trước main
    • __cudaRegisterBinary và __cudaRegisterEntry liên kết host function pointer vadd với device entry point _Z4vaddPKfS0_Pfi

Bài viết liên quan

1 bình luận

 
GN⁺ 3 giờ trước
Các ý kiến trên Hacker News

  • Một bài viết thú vị, và phần giải thích về semaphore của stream mặc định cũng rất hay
    Thích ở chỗ CUDA ngầm xử lý việc đồng bộ hóa lệnh, còn các lệnh song song thì cho phép dùng có chọn lọc thông qua stream
    Điều này tương phản với Vulkan, vốn ngay từ đầu đã đẩy toàn bộ độ phức tạp của đồng bộ hóa cho người dùng

  • Về phía phần cứng thì có một số tài liệu công khai
    Không nhất thiết phải đọc mã nguồn kernel để tìm tài liệu về method hay định dạng QMD
    Xem https://github.com/NVIDIA/open-gpu-doc/blob/master/classes/c...

  • Rất hữu ích
    Đặc biệt phần về doorbell và QMD là hữu ích nhất, vì cho thấy cú pháp chạy CUDA thực sự được nối với nội dung gửi xuống GPU như thế nào
    Phần lớn các bài giải thích dừng lại quanh kernel, block, warp; còn bài này giúp lần theo đường đi CPU→driver→GPU dễ hơn nhiều

  • Mã điều khiển phức tạp hơn một chút so với mô tả trong bài
    Thực tế nó giống tra bảng hơn là một bit trong control word

  • Hiện nay có những công ty lấy việc tối ưu kernel để chạy nhanh hơn làm nghiệp vụ chính
    Tò mò liệu một ngày nào đó các công ty đó có bị một thư viện mã nguồn mở làm việc này cực tốt vượt qua hay không
    Với Nvidia thì có vẻ họ có thể tung ra thứ như vậy bất cứ lúc nào
    Hoặc cũng có thể mọi chuyện sẽ tốt hơn khi các nhà cung cấp lớn thâu tóm các công ty này để biến tốc độ suy luận thành moat

    • Trong ngắn hạn, thâu tóm kiểu acqui-hire có vẻ khá khả thi
      Tuy nhiên, nhìn vào việc các mô hình tiến bộ trên những benchmark liên quan như kernelbench, tôi nghĩ rốt cuộc các giải pháp tổng quát hơn cũng sẽ xuất hiện
      Vấn đề là mỗi thế hệ phần cứng mới thường lại có các ràng buộc hoặc tính năng mà mô hình cũ chưa từng thấy
      Ví dụ tcgen05 của Blackwell từng là một trường hợp ngoài phân phối
      Nếu các mô hình bắt đầu khái quát hóa tốt hơn thì đây có thể không phải rào cản chí mạng, nhưng ít nhất hiện tại nó vẫn là trở ngại
      [1] https://kernelbench.com/
    • Khi chạy CUDA ở quy mô lớn, thời gian kỹ sư phải bỏ ra để xử lý lỗi driver và thư viện của Nvidia nhiều đến mức phát ngán
      Tôi hiếm thấy ai mong muốn phụ thuộc nhiều hơn vào thư viện của Nvidia
    • Có lẽ là không
      Vì các chi tiết của workload — tức các tham số chính xác, cách biểu diễn dữ liệu trong bộ nhớ, phạm vi giá trị, v.v. — làm chiến lược tối ưu hóa khác nhau rất nhiều

  • Tôi vừa hoàn thành thạc sĩ HPC và đã học các môn CUDA, MPI+CUDA, OpenCL; nếu đọc được bài này trước khi vào học thì chắc đã giúp ích hơn nhiều
    Đặc biệt thích phần trước sau đoạn nói về việc một warp có thể thực thi nghĩa là gì

  • Trước hết, đây là một bài viết hay, đào khá sâu vào nhiều ngóc ngách
    Tuy nhiên, nếu không đi qua runtime API của CUDA thì rất nhiều phần mang màu sắc voodoo trong không gian người dùng sẽ biến mất
    Dùng driver API, nhận mã nguồn kernel dưới dạng chuỗi rồi biên dịch bằng trình biên dịch runtime của NVIDIA, bạn sẽ thấy rõ hơn chuyện gì đang diễn ra
    Không phải toàn bộ, nhưng khá nhiều phần trở nên minh bạch
    Phiên bản “thô sơ” hơn nằm ở đây:
    https://github.com/NVIDIA/cuda-samples/tree/master/cpp/0_Int...
    Nếu muốn xem cùng nội dung đó ở dạng API C++ hiện đại, dễ đọc hơn nhiều nhưng vẫn hoàn toàn minh bạch, hãy xem cái này:
    https://github.com/eyalroz/cuda-api-wrappers/blob/master/exa...
    Đây là chương trình ví dụ trong thư viện chỉ gồm header CUDA API wrappers của tôi

    • Driver API hay ở chỗ nó cho phép coi CUDA kernel như shader có thể hot-reload
      Có thể phát triển bằng cách thay đổi mã trong lúc đang chạy, khá thú vị

  • Trên bare metal à?