เกิดอะไรขึ้นภายในเมื่อคุณรัน CUDA kernel

 (fergusfinn.com)
2 คะแนน โดย GN⁺ 4 시간 전 | 1 ความคิดเห็น | แชร์ทาง WhatsApp
  • แม้แต่โปรแกรม CUDA สำหรับบวกเวกเตอร์อย่างง่ายก็ยังต้องผ่าน pipeline การคอมไพล์, การเรียกไดรเวอร์, คิวคำสั่งของ GPU, การจัดตาราง warp, ลำดับชั้นหน่วยความจำ และ semaphore สำหรับแจ้งเสร็จสิ้น กว่าจะได้ผลลัพธ์ 2.000000
  • nvcc แยกโค้ดฝั่งโฮสต์และโค้ดฝั่งดีไวซ์ แล้วสร้าง PTX ด้วย cicc, สร้าง SASS ด้วย ptxas, จากนั้นรวม cubin และ PTX เป็น fatbin แล้วใส่ไว้ในไฟล์ executable บน Linux
  • ไวยากรณ์ launch vadd<<<4096, 256>>> จะถูกแปลงเป็น host launch stub และอาร์กิวเมนต์ da, db, dc, n จะถูกส่งไปยังไดรเวอร์ผ่าน CUDA runtime และ libcuda.so.1
  • การรันบน GPU เริ่มจาก QMD, pushbuffer, GPFIFO, GP_PUT และการเขียน MMIO แบบ doorbell โดย RTX 4090 ที่มี 128 SM จะรันคอนฟิก 4096 บล็อกและ 256 เธรดในระดับ warp
  • kernel นี้มี arithmetic intensity ต่ำ เพราะต้องย้ายข้อมูล 12 ไบต์ต่อการบวก float 1 ครั้ง ทำให้ใน Nsight Compute วัดได้ 10.78μs, ใช้ DRAM peak 79.65% และมี warp issue 5.17% จึงถูกจำกัดด้วยแบนด์วิดท์หน่วยความจำ

ตัวอย่าง kernel และขอบเขตการสังเกต

  • โปรแกรมตัวอย่างใช้ CUDA kernel vadd เพื่อบวกอาร์เรย์ float สองชุดแล้วเก็บลงในอาร์เรย์ชุดที่สาม
    • ใช้ n = 1 << 20 เพื่อประมวลผล float 1,048,576 ค่า
    • คอนฟิก launch คือ vadd<<<4096, 256>>>(da, db, dc, n) และใช้เธรด 4096 * 256 = n ตัว
  • เมื่อคอมไพล์ด้วย nvcc -arch=sm_89 และรันบน RTX 4090 จะได้เอาต์พุต c[0]=2.000000 c[n-1]=2.000000
  • แม้ผลลัพธ์จะมีเพียงบรรทัดเดียว แต่เบื้องหลังก็มีคำสั่ง CPU หลายสิบล้านคำสั่ง, device file, ioctl ราว 900 ครั้ง และรีจิสเตอร์ doorbell ที่แมปด้วยหน่วยความจำเข้ามาเกี่ยวข้อง

กระบวนการที่ nvcc สร้างไฟล์ executable

  • หากใช้ nvcc --keep จะสามารถดูผลลัพธ์จาก pipeline การคอมไพล์ได้โดยตรง
    • vadd.ptx: PTX ของโค้ดดีไวซ์ที่สร้างโดย cicc
    • vadd.sm_89.cubin: SASS ของโค้ดดีไวซ์ที่สร้างโดย ptxas
    • vadd.fatbin: fatbin ที่รวม cubin และ PTX
    • vadd.cudafe1.stub.c: host launch stub และโค้ดลงทะเบียน kernel
    • vadd.o: host object สุดท้ายที่มี fatbin ฝังอยู่
  • โค้ดฝั่งโฮสต์จะถูกประมวลผลด้วย host compiler ส่วน kernel ฝั่งดีไวซ์ vadd จะผ่านขั้นตอน cicc และ ptxas
  • PTX เป็น ISA เสมือนที่ใช้ virtual register แบบมีชนิดข้อมูลและมีจำนวนไม่จำกัด จึงไม่ได้สะท้อนจำนวนรีจิสเตอร์ของฮาร์ดแวร์จริงโดยตรง
    • PTX ของตัวอย่างมีการคำนวณ blockIdx.x * blockDim.x + threadIdx.x, การตรวจขอบเขต, global load, การบวก float และ global store
    • pointer ของ CUDA โดยปริยายเป็น generic pointer จึงต้องแปลงเป็น global address ด้วย cvta.to.global ก่อนใช้ ld.global
    • mul.wide.s32 ใช้แปลง index ให้เป็น offset หน่วย 4 ไบต์ตาม sizeof(float) และขยายจาก 32 บิตเป็น 64 บิต
  • SASS คือชุดคำสั่งจริงที่ขึ้นกับสถาปัตยกรรม และในเอาต์พุตสำหรับ RTX 4090 จะมีรูปแบบที่กระชับกว่า PTX
    • S2R ใช้คัดลอก special register เช่น SR_CTAID.X, SR_TID.X ไปยังรีจิสเตอร์ทั่วไป
    • ชุด mul.wide กับ add ใน PTX จะถูกรวมเป็น IMAD.WIDE ใน SASS
    • การแปลง cvta จะถูกรวมเข้าไปในกระบวนการอ้างที่อยู่
  • โอเปอแรนด์ c[0x0][...] ชี้ไปยัง constant bank 0 ที่ไดรเวอร์จัดการ
    • pointer a, b, c อยู่ที่ 0x160, 0x168, 0x170
    • n อยู่ที่ 0x178
    • ค่า launch geometry เช่น blockDim.x และค่า ABI ก็อยู่ใน bank เดียวกัน
  • cubin เป็นไฟล์ ELF ซึ่งเป็นฟอร์แมตคอนเทนเนอร์แบบเดียวกับไฟล์ executable บน Linux
    • fatbinary จะรวมทั้ง cubin และ PTX ไว้ด้วยกัน
    • บน RTX 4090 นี้ SASS คือสิ่งที่ถูกรันจริง ส่วน PTX ถูกเก็บไว้เป็น fallback สำหรับให้ไดรเวอร์ทำ JIT คอมไพล์บนสถาปัตยกรรมอื่น
    • PTX เป็น plain text แบบ verbose จึงถูก nvcc บีบอัดโดยปริยาย

วิธีที่โค้ดฝั่งโฮสต์เตรียมการ launch

  • คอมไพเลอร์ฟรอนต์เอนด์ cudafe++ จะแทรก constructor ที่ซ่อนอยู่ซึ่งรันก่อน main
    • constructor นี้จะลงทะเบียน embedded fatbinary กับ CUDA runtime
    • และเชื่อม function pointer ฝั่งโฮสต์ vadd เข้ากับชื่อดีไวซ์ kernel แบบ mangled ภายใน fatbin
  • ไวยากรณ์ vadd<<<4096, 256>>>(da, db, dc, n) จะถูกแปลงเป็น host launch stub ที่สร้างขึ้นอัตโนมัติ
    • da, db, dc, n จะถูกจัดวางลงใน argument buffer บน host memory ที่ออฟเซ็ต 0, 8, 16, 24
    • ออฟเซ็ตเหล่านี้สอดคล้องกับตำแหน่ง 0x160, 0x168, 0x170, 0x178 ที่ SASS อ่านจาก constant bank 0
  • stub จะเรียก __cudaLaunch พร้อมส่งที่อยู่ของฟังก์ชัน dummy vadd ฝั่งโฮสต์
    • ที่อยู่นี้ไม่ใช่ที่อยู่ของฟังก์ชันที่จะรันบน CPU แต่ใช้เป็นคีย์สำหรับค้นหาในตารางลงทะเบียนของ runtime
    • runtime จะค้นหาชื่อ device symbol ที่ตรงกันก่อนส่งต่อไปยัง libcuda.so.1 ซึ่งเป็น user-mode driver แบบ closed source
  • เมื่อมีการเรียกใช้ GPU ครั้งแรก CUDA runtime จะเปิด libcuda.so.1 แบบไดนามิกและสร้าง context
    • ใน strace จะเห็นการเปิด /lib/x86_64-linux-gnu/libcuda.so.1
    • ภายใน context จะมี channel ที่ CPU ใช้สื่อสารกับ GPU
  • ตั้งแต่ CUDA 12.2 เป็นต้นไป การโหลดโมดูลเป็นแบบ lazy โดยปริยาย
    • การอัปโหลด SASS cubin จะถูกเลื่อนไปจนกว่าจะมีการ launch kernel นั้นครั้งแรก
    • สามารถควบคุมได้ด้วย CUDA_MODULE_LOADING

คิวคำสั่งที่ใช้ส่งงานไปยัง GPU

  • GPU ไม่ได้รับการเรียกฟังก์ชันแล้วกระโดดไปยัง entry point แบบ CPU
    • แต่จะอ่าน driver command stream ใน host memory ผ่านบัส PCIe
    • cuLaunchKernel จะใส่คำสั่ง launch ที่สมบูรณ์ลงใน stream นี้แล้วแจ้ง GPU
  • ในการรันครั้งแรก ไดรเวอร์จะคัดลอก SASS ของ kernel ไปยังหน่วยความจำของ GPU
    • มีการจอง code buffer แล้วคัดลอก SASS ลงไป
  • ภายใน channel มีโครงสร้างสำคัญสองส่วนที่อยู่ใน host RAM
    • pushbuffer: พื้นที่หน่วยความจำที่ไดรเวอร์ใช้เขียน method ซึ่งเป็นคำสั่งของ GPU
    • GPFIFO: ring buffer ของ pointer ที่ชี้ไปยังช่วงของ pushbuffer
  • รายการ GPFIFO หนึ่งรายการประกอบด้วย word 32 บิตสองคำ ซึ่งแทน (base, length) ของช่วงใน pushbuffer
  • GPU และไดรเวอร์ติดตามตำแหน่งการผลิตและการใช้ข้อมูลด้วยเคอร์เซอร์สองตัว
    • GP_GET: ระบุว่า GPU ใช้ข้อมูลไปถึงไหนแล้ว
    • GP_PUT: ระบุว่าไดรเวอร์ผลิตข้อมูลไปถึงไหนแล้ว
    • ทั้งคู่เก็บอยู่ในโครงสร้างต่อ channel ที่ชื่อ USERD
  • ตอน launch kernel ไดรเวอร์จะเขียน method ลงในช่วงของ pushbuffer, ทำให้รายการ GPFIFO ชี้มายังช่วงนั้น แล้วเลื่อน GP_PUT ไปข้างหน้า
  • บน GPU สมัยใหม่ host engine จะไม่คอยเฝ้าดูเคอร์เซอร์อย่างต่อเนื่อง จึงต้องใช้ doorbell
    • GPU จะแมปหน้าต่างรีจิสเตอร์ขนาดเล็กให้กับโปรเซส
    • ไดรเวอร์จะเขียน work-submit token ของ channel ลงใน doorbell register
    • host engine เมื่อได้รับ doorbell แล้วจึงค่อยอ่าน GP_PUT และดึงรายการ GPFIFO กับช่วง pushbuffer ผ่าน DMA

ข้อมูลการรันที่อยู่ใน QMD

  • การ launch เริ่มจากชุด method burst ของ SET_INLINE_QMD_ADDRESS_A/B และ LOAD_INLINE_QMD_DATA
  • QMD(Queue Meta Data) คือ launch descriptor ของ compute grid
    • มีขนาด grid และ block คือ 4096, 256
    • มีจำนวนรีจิสเตอร์ต่อเธรดและความต้องการ shared memory
    • มีที่อยู่เริ่มต้นของโปรแกรมและที่อยู่ constant bank ที่เก็บอาร์กิวเมนต์ของ kernel
    • มีตำแหน่งสำหรับแจ้งว่าเสร็จสิ้นด้วย
  • อาร์กิวเมนต์ที่ host stub แพ็กไว้จะถูกไดรเวอร์คัดลอกไปยัง constant bank และบันทึกที่อยู่ของ bank นี้ลงใน QMD
  • QMD ใช้บอก GPU ว่า SASS อยู่ที่ไหน, จะจัดโปรแกรมแบบขนานอย่างไร และต้องส่งสัญญาณเสร็จสิ้นที่จุดใด
  • cuLaunchKernel จะคืนค่าทันทีที่มีการกด doorbell
    • การเรียกนี้เป็นแบบ asynchronous ดังนั้น CPU จึงทำงานต่อได้ในขณะที่ GPU กำลังประมวลผลงาน

SM, warp และ occupancy

  • host engine จะส่ง QMD ต่อไปยัง compute work distributor
    • องค์ประกอบนี้มีเพียงหนึ่งเดียวทั้ง GPU
    • หน้าที่คือกระจาย linear SASS instruction stream ไปยัง SM ต่าง ๆ เพื่อให้รันเป็นโปรแกรมแบบขนาน
  • GPU เป้าหมาย GeForce RTX 4090 ใช้ 128 SM
    • การ launch นี้ประกอบด้วย 4096 บล็อก และ 256 เธรดต่อบล็อก
  • แต่ละ SM มี local instruction cache ของตัวเอง และ active warp จะเก็บ program counter เอาไว้
    • ตั้งแต่ Volta เป็นต้นมาใช้โมเดล Independent Thread Scheduling ที่มี program counter และ call stack ต่อเธรด
    • แต่การ issue คำสั่งยังคงทำในระดับ warp
  • ใน kernel ตัวอย่าง ข้อจำกัดด้านทรัพยากรเป็นตัวกำหนดจำนวน block ที่อยู่ resident ได้
    • ต่อบล็อกมี 256 threads = 8 warps
    • ptxas จองรีจิสเตอร์ไว้ 16 ตัวต่อเธรด
    • หากดูจากรีจิสเตอร์จะวางได้ 16 บล็อกต่อ SM
    • แต่เพดานเธรดคือ 1,536 active threads ต่อ SM ดังนั้น 1536 / 256 = 6 จึงวางได้เพียง 6 บล็อก
    • ดังนั้นสูงสุดต่อ SM คือ 6 บล็อก หรือ 48 warp ที่อยู่ในสถานะ resident
  • SM ถูกแบ่งเป็น 4 processing block หรือ sub-partition
    • 48 resident warp จะถูกกระจายเท่า ๆ กันไปยัง 4 sub-partition
    • warp scheduler แต่ละตัวจะดูแล active warp ได้ 12 ตัวเมื่ออยู่ในสถานะเต็ม
    • ในแต่ละ cycle จะเลือก eligible warp หนึ่งตัวแล้ว dispatch คำสั่งถัดไปไปยัง 32 lane

เงื่อนไขที่ทำให้ warp อยู่ในสถานะ eligible

  • GPU ไม่ได้ดึง dependency แบบไดนามิกจากเธรดเดี่ยวได้มากเหมือนการรัน out-of-order ของ CPU
    • แต่จะรักษา resident warp ไว้จำนวนมาก และเมื่อมี stall ก็สลับไปทำ warp อื่นเพื่อซ่อน latency
    • คอมไพเลอร์จะจัดตารางในส่วนที่ทำนายเวลาได้ ส่วน hardware scoreboard จะดูแลส่วนที่ทำนายได้ยาก
  • ใน SASS instruction ขนาด 128 บิต จะมี control-code payload ที่ ptxas เขียนไว้
    • สำหรับคำสั่งที่มี latency คงที่ จะมี static stall count ติดมาด้วย
    • yield hint ใช้บอกว่าจะยอมลดลำดับความสำคัญให้ scheduler หรือไม่
    • สำหรับ operation ที่มี latency ไม่แน่นอน จะใช้ physical scoreboard barrier ต่อ warp จำนวน 6 ตัว
  • ในช่วง SASS ของตัวอย่าง คำสั่ง LDG.E สองตัวจะ set scoreboard barrier เดียวกันคือ B2
    • ส่วน FADD จะมี B2 เป็น wait-on
    • จนกว่าการโหลดทั้งสองจะกลับมาและ barrier ถูก clear, warp นั้นจะยังอยู่ในสถานะ ineligible
    • ระหว่างนั้น scheduler จะเลือก warp อื่นใน sub-partition เดียวกัน
  • ช่วงจาก FADD ไป STG.E ใช้การจัดการแบบ fixed latency
    • FADD มี stall=5 และจะพัก warp ไว้หลาย cycle จนกว่าผลลัพธ์ใน R9 จะพร้อม
    • จึงไม่ต้องใช้ barrier แยกต่างหาก
  • control payload นี้ถูกซ่อนไว้ในเอาต์พุตปริยายของ nvdisasm
    • แต่จะอยู่ใน word 64 บิตตัวที่สองของ raw 128-bit encoding จาก cuobjdump -sass
    • layout นี้ไม่ได้มีเอกสารทางการรองรับ แต่ถูกสร้างภาพย้อนกลับขึ้นมาจาก microbenchmarking

การเข้าถึงหน่วยความจำและการวัดประสิทธิภาพ

  • เมื่อ warp รัน LDG.E เธรดทั้ง 32 จะคำนวณที่อยู่ของตัวเอง
    • ในตัวอย่างนี้เป็นการเข้าถึงอาร์เรย์ float แบบต่อเนื่อง ทำให้ทั้ง warp ขอข้อมูลเป็นบล็อกต่อเนื่องขนาด 32 * 4 = 128 bytes
  • load/store unit ของ SM จะทำ request coalescing
    • รวมคำขอ 4 ไบต์ 32 รายการให้เป็น sector request ขนาด 32 ไบต์จำนวน 4 รายการ
    • หากไม่ได้เข้าถึงต่อเนื่อง ก็อาจต้องอ่านข้อมูลมากกว่าที่จำเป็น
  • คำขอที่ coalesce แล้วจะตรวจที่ SM local L1 Data Cache ก่อน
    • หาก miss จะวิ่งผ่าน crossbar interconnect ไปยัง L2 Cache slice ขนาด 72MB
    • ถ้ายัง miss ที่ L2 อีก ก็จะผ่าน memory controller และ memory bus ไปยัง VRAM แบบ GDDR6X
  • store STG.E ก็จะใช้เส้นทางเดียวกันในทิศทางกลับกันโดยหลักการ
  • ค่าที่วัดจาก Nsight Compute แสดงให้เห็นว่า kernel นี้ถูกจำกัดด้วยหน่วยความจำ
    • 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 ต่ำมาก
    • ทำ float add 1 ครั้งต่อการย้ายข้อมูลรวม 12 ไบต์ จาก 4-byte load สองครั้งและ 4-byte store หนึ่งครั้ง
    • หากดูเฉพาะการอ่าน DRAM จะอ่าน 8.4MB ใน 10.78μs หรือประมาณ 780GB/s ซึ่งอยู่ที่ราว 4/5 ของค่าสูงสุด
    • เอาต์พุต c ขนาด 4MB สามารถอยู่ใน L2 ขนาด 72MB ได้ ดังนั้นจะยังไม่ถูก flush ลง DRAM จนกว่าการคัดลอกกลับไปยังโฮสต์จะมาอ่าน

กระบวนการที่ผลลัพธ์กลับมาสู่ CPU

  • การ launch kernel จะคืนกลับสู่ CPU ตั้งแต่ตอนกด doorbell ดังนั้น GPU ต้องมีวิธีแจ้งแยกต่างหากเมื่อทำเสร็จ
  • เมื่อ block ทั้ง 4096 ตัว retire ครบแล้ว GPU จะ post completion semaphore ที่บันทึกไว้ใน QMD
    • field สำหรับ fence ใน QMD อยู่ที่ words 23–24
  • ใน default stream คำสั่ง cudaMemcpy(c, dc, ...) จะถูกวางต่อจาก kernel
    • GPU copy engine จะอยู่ในสถานะ gated จนกว่า semaphore จะถูกยกขึ้น
    • เนื่องจาก c ยังอยู่ในสถานะ dirty ภายใน L2 ขนาด 72MB การอ่านของ copy engine จึงเสร็จจาก L2 ได้โดยไม่ต้องวิ่งไปกลับ DRAM
    • จากนั้นข้อมูลจะถูกย้ายผ่าน PCIe ไปยัง host memory
  • เมื่อการคัดลอกเสร็จ copy engine จะ post semaphore ของตัวเอง
    • การรอของ cudaMemcpy ฝั่งโฮสต์จึงสิ้นสุดลง
    • c กลับมาเป็น host memory ปกติอีกครั้ง
    • printf จะอ่าน c[0] และ c[n-1] จาก RAM แล้วพิมพ์ออกทาง stdout

วิธีส่องดูภายในของ launch

  • การอ่าน open kernel modules อย่างเดียวไม่เพียงพอ เพราะ libcuda เป็น closed source จึงยืนยันพฤติกรรมบางส่วนได้ยาก
  • การเขียน method ไม่ได้ผ่าน syscall แต่เขียนลง write-combined buffer ที่แมปไว้แล้วโดยตรง ดังนั้นหากอยากดู pushbuffer ต้องอ่านหน่วยความจำ
  • สามารถใช้ LD_PRELOAD shim ครอบ mmap เพื่อบันทึกพื้นที่ที่ถูกแมปจาก /dev/nvidia*
    • หากโปรแกรมทดสอบเรียกฟังก์ชัน dump ของ shim ทันทีหลัง launch ก็จะพิมพ์ pushbuffer ที่ถูกแมปออกมาได้
    • dump จะค้นหา method burst ที่ตรงกับ SET_INLINE_QMD_ADDRESS_A
  • header ของ method ใน pushbuffer จะเก็บ opcode, จำนวน payload, ดัชนี subchannel และ register offset ไว้ในรูป bit field
    • 0x0318 คือ SET_INLINE_QMD_ADDRESS_A
    • 0x0320 + i * 4 คือ LOAD_INLINE_QMD_DATA(i)
    • ใน dump จะเห็น increasing-method burst ที่มี count 66 ซึ่งประกอบด้วย address word 2 คำและ QMD 64 word รวมเป็น QMD แบบ inline ขนาด 256 ไบต์
    • ภายใน QMD, word 12 คือ 0x1000 และ word 18 คือ 0x100 ซึ่งตรงกับค่า 4096 และ 256 ของการ launch
  • การตั้งค่าไดรเวอร์เกิดขึ้นผ่าน ioctl
    • ในโปรแกรมที่มี kernel เดียว strace บันทึก ioctl ได้ 948 ครั้ง
    • ส่วนใหญ่เป็นการตั้งค่าแบบ one-time setup
    • file descriptor หลักคือ /dev/nvidiactl และ /dev/nvidia-uvm
    • ioctl magic byte ของ NVIDIA resource manager คือ 0x46 หรือ 'F'
    • หมายเลขคำสั่ง 0x2A ตีความได้เป็น NV_ESC_RM_CONTROL และ 0x2B เป็น NV_ESC_RM_ALLOC
  • ใน vadd.cudafe1.stub.c ที่สร้างจาก nvcc --keep จะมองเห็นโค้ด registration ตอนเริ่มต้นได้ด้วย
    • ฟังก์ชันที่มี __attribute__((__constructor__)) จะรันก่อน main
    • __cudaRegisterBinary และ __cudaRegisterEntry ใช้เชื่อม host function pointer vadd เข้ากับ device entry point _Z4vaddPKfS0_Pfi

บทความที่เกี่ยวข้อง

1 ความคิดเห็น

 
GN⁺ 4 시간 전
ความเห็นจาก Hacker News

  • เป็นบทความที่น่าสนใจ และคำอธิบายเรื่อง semaphore ของ default stream ก็น่าสนุกด้วย
    ชอบตรงที่ CUDA จัดการการซิงโครไนซ์คำสั่งให้แบบแฝง ๆ และเปิดให้ใช้คำสั่งแบบขนานผ่าน stream ได้ตามต้องการ
    ตรงข้ามกับ Vulkan ที่โยนความซับซ้อนของการซิงโครไนซ์ทั้งหมดให้ผู้ใช้ตั้งแต่แรก

  • ฝั่งฮาร์ดแวร์มี เอกสารที่เปิดเผยต่อสาธารณะ อยู่บ้าง
    ไม่จำเป็นต้องไปอ่านซอร์สเคอร์เนลเพื่อหาเอกสารเมธอดหรือฟอร์แมต QMD เสมอไป
    ดูได้ที่ https://github.com/NVIDIA/open-gpu-doc/blob/master/classes/c...

  • มีประโยชน์มาก
    โดยเฉพาะส่วน doorbell และ QMD ที่ช่วยให้เห็นว่ารูปแบบการรันของ CUDA เชื่อมต่อกับสิ่งที่ถูกส่งไปยัง GPU จริง ๆ อย่างไร ซึ่งเป็นส่วนที่ช่วยได้มากที่สุด
    คำอธิบายส่วนใหญ่มักหยุดอยู่แถว ๆ kernel, block, warp แต่บทความนี้ทำให้ตามเส้นทาง CPU→driver→GPU ได้ง่ายกว่ามาก

  • โค้ดควบคุมจริงซับซ้อนกว่าที่บทความอธิบายนิดหน่อย
    ในความเป็นจริงมันใกล้เคียงกับ การ lookup ตาราง มากกว่าจะเป็นบิตใน control word

  • ตอนนี้มีบริษัทที่มีงานหลักคือปรับแต่งเคอร์เนลให้รันได้เร็วขึ้น
    ก็สงสัยว่าบริษัทเหล่านั้นวันหนึ่งจะถูก ไลบรารีโอเพนซอร์ส ที่เก่งเรื่องนี้มาก ๆ แซงไปหรือไม่
    ถ้าเป็น Nvidia ก็ดูเหมือนจะปล่อยอะไรแบบนั้นออกมาได้ทุกเมื่อ
    หรืออีกทาง บริษัทเหล่านี้อาจไปได้ดีกว่าเดิม ถ้าผู้ให้บริการรายใหญ่เข้าซื้อเพื่อใช้เป็น moat สำหรับเร่งความเร็ว inference

    • ในระยะสั้น การเข้าซื้อแบบ acqui-hire ก็ดูมีความเป็นไปได้พอสมควร
      แต่พอดูจากการที่โมเดลพัฒนาในเบนช์มาร์กที่เกี่ยวข้องอย่าง kernelbench ก็คิดว่าสุดท้ายวิธีแก้ที่เป็นสากลมากขึ้นก็น่าจะต้องเกิดขึ้น
      ปัญหาคือในฮาร์ดแวร์แต่ละเจเนอเรชันใหม่ มักมีข้อจำกัดหรือฟีเจอร์ที่โมเดลเดิมไม่เคยเห็นมาก่อนโผล่มาอยู่เรื่อย ๆ
      ตัวอย่างเช่น tcgen05 ของ Blackwell ก็เคยเป็นกรณีนอกการกระจายมาก่อน
      ถ้าโมเดลเริ่ม generalize ได้ดีขึ้น มันอาจไม่ใช่อุปสรรคที่ร้ายแรงนัก แต่ตอนนี้อย่างน้อยก็ยังเป็นจุดสะดุดอยู่
      [1] https://kernelbench.com/
    • ถ้ารัน CUDA ในสเกลใหญ่ เวลาเชิงวิศวกรรมจะถูกกินไปกับการรับมือ บั๊กในไดรเวอร์และไลบรารีของ Nvidia มากจนน่าขยะแขยง
      ไม่ค่อยเห็นใครคาดหวังอยากพึ่งพาไลบรารีของ Nvidia มากกว่าเดิมเท่าไร
    • คิดว่าไม่น่าจะเป็นแบบนั้น
      เพราะรายละเอียดของ workload เช่น พารามิเตอร์ที่แน่นอน การแทนข้อมูลในหน่วยความจำ และช่วงของค่า ล้วนทำให้กลยุทธ์การปรับแต่งต่างกันมาก

  • เพิ่งเรียนจบปริญญาโทด้าน HPC และเคยเรียนวิชา CUDA, MPI+CUDA, OpenCL มาก่อน ถ้าได้อ่านบทความแบบนี้ก่อนเรียนคงช่วยได้มากกว่านี้มาก
    โดยเฉพาะบริบทก่อนหลังของส่วนที่พูดถึง ความหมายของการที่ warp พร้อมรันได้ นั้นดีมาก

  • ก่อนอื่นเลย เป็นบทความที่ดีและเจาะรายละเอียดหลายมุมมาก
    แต่ถ้าไม่ผ่าน runtime API ของ CUDA ส่วนที่ดูเหมือน มนตร์ดำ ใน user space หลายอย่างจะหายไป
    ถ้าใช้ driver API และรับซอร์สเคอร์เนลเป็นสตริงแล้วคอมไพล์ด้วยคอมไพเลอร์แบบรันไทม์ของ NVIDIA ก็จะเห็นได้ชัดขึ้นว่ามีอะไรเกิดขึ้น
    แม้จะไม่ทั้งหมด แต่หลายส่วนก็โปร่งใสขึ้นมาก
    เวอร์ชันที่ “ดิบ” กว่านี้อยู่ที่นี่:
    https://github.com/NVIDIA/cuda-samples/tree/master/cpp/0_Int...
    ถ้าอยากดูเนื้อหาเดียวกันในรูปแบบ modern C++ API ที่อ่านง่ายกว่ามากและยังโปร่งใสเต็มที่ ก็ดูอันนี้ได้:
    https://github.com/eyalroz/cuda-api-wrappers/blob/master/exa...
    เป็นโปรแกรมตัวอย่างจากไลบรารี header-only CUDA API wrappers ของผม

    • ชอบที่ driver API ทำให้จัดการ CUDA kernel ได้เหมือน shader ที่ hot-reload ได้
      สนุกดีที่สามารถเปลี่ยนโค้ดระหว่างรันเพื่อพัฒนาไปด้วยได้

  • บน bare metal เหรอ?