Rust เขียนเคอร์เนล GPU โดยตรง! cuda-oxide คอมไพล์โค้ดที่ปลอดภัยเป็น PTX โดยไม่ต้องใช้ CUDA C++

หากกล่าวว่า CUDA C++ คือ “ภาษาแม่” ในแวดวงการเขียนโปรแกรม GPU แล้ว Rust ก็เปรียบเสมือนวิศวกรที่ยืนอยู่นอกประตูมาเป็นเวลานาน: มันมีระบบชนิดข้อมูลที่ทรงพลัง โมเดลความเป็นเจ้าของ และนามธรรมที่ไร้ต้นทุน แต่กลับยากที่จะผสานเข้ากับโมเดลการทำงาน SIMT ของ GPU NVIDIA ได้อย่างเป็นธรรมชาติ

แนวทางแก้ไขแบบดั้งเดิมนั้น ไม่ว่าจะเป็นการเขียนภาษาเฉพาะโดเมน (DSL) การผูกกับโค้ด CUDA ภายนอก หรือการเสียสละความหมายของ Rust เพื่อแลกกับความสามารถในการคอมไพล์ ล้วนมีข้อจำกัด

• cuda-oxide เป็นคอมไพเลอร์ทดลองที่แปลง Rust เป็น CUDA ซึ่งช่วยให้นักพัฒนาสามารถเขียนโปรแกรมเคอร์เนล GPU แบบ SIMT ได้ด้วยภาษา Rust ที่ค่อนข้างปลอดภัยและเป็นไปตามรูปแบบการเขียนดั้งเดิม มันสามารถคอมไพล์โค้ด Rust มาตรฐานเป็นชุดคำสั่ง PTX ได้โดยตรง โดยไม่ต้องใช้ภาษาเฉพาะโดเมนหรือการผูกกับภาษาภายนอก ทำได้ทั้งหมดโดยใช้ภาษา Rust เท่านั้น
• โค้ด: https://github.com/NVlabs/cuda-oxide
• เอกสาร: https://nvlabs.github.io/cuda-oxide/
• สถานะโครงการ: เวอร์ชัน v0.1.0 เป็นรุ่นพรีวิวภายในช่วงแรกเริ่ม ยังคงมีข้อบกพร่องของโปรแกรม ฟังก์ชันการทำงานที่ไม่สมบูรณ์ และการเปลี่ยนแปลงอินเทอร์เฟซของแอปพลิเคชัน เราจะปรับปรุงและพัฒนาซ้ำอย่างต่อเนื่อง ขอเชิญทุกท่านทดลองใช้งานและแบ่งปันข้อเสนอแนะ เพื่อช่วยให้เราพัฒนาทิศทางของผลิตภัณฑ์ให้ดียิ่งขึ้น
• 8000 คำ อ่าน 40 นาที, พอดแคสต์ 26 นาที

ความกล้าหาญของ cuda-oxide อยู่ที่: มันไม่ต้องการอ้อม绕过 Rust แต่กลับส่งโค้ด Rust มาตรฐานเข้าไปยัง rustc โดยตรง จากนั้นสกัดกั้นฟังก์ชันฝั่งอุปกรณ์ในระดับ MIR ผ่าน Pliron IR, LLVM IR และสุดท้ายสร้าง PTX มันพยายามตอบคำถามที่ลึกซึ้งกว่านั้น: “ความปลอดภัย” ของ Rust สามารถใช้ได้บน GPU หรือไม่?

หากเป็นไปได้ ขอบเขตอยู่ที่ไหน? บทความนี้จะแยกย่อยคุณค่าหลักและต้นทุนทางเทคนิคของคอมไพเลอร์ Rust-to-CUDA เชิงทดลองนี้จากหกมุมมอง ได้แก่ การใช้งานอย่างรวดเร็ว การแบ่งชั้นทางวิศวกรรม สายการคอมไพล์ การขยายแมโคร นามธรรมของอุปกรณ์ และรันไทม์ของโฮสต์

รูปภาพนี้อธิบายข้อได้เปรียบหลักสามประการของ cuda-oxide: ประการแรก รองรับการเขียนเคอร์เนล GPU ด้วย Rust โดยใช้ระบบชนิดข้อมูลและโมเดลความเป็นเจ้าของเพื่อรับประกันความปลอดภัย พร้อมปรับให้เข้ากับคุณสมบัติฮาร์ดแวร์ของ GPU ประการที่สอง มีคอมไพเลอร์ SIMT ในตัว ไม่ใช่ภาษาเฉพาะโดเมน แต่ใช้แบ็กเอนด์ rustc ที่กำหนดเองเพื่อคอมไพล์โค้ด Rust บริสุทธิ์เป็น PTX ประการที่สาม รองรับการทำงานแบบอะซิงโครนัส สามารถรวมงาน GPU เป็นกราฟ DeviceOperation แบบเลื่อนเวลา จัดตารางเวลาข้ามพูลสตรีม และรอผลลัพธ์ผ่าน .await รูปภาพนี้แสดงสายการคอมไพล์ที่สมบูรณ์ของ cuda-oxide ตั้งแต่ซอร์สโค้ด Rust ไปจนถึงแอสเซมบลี PTX โดยปรัชญาการออกแบบเป็นไปตามหลักการสำคัญ “เลือกเครื่องมือที่ดีที่สุดสำหรับแต่ละขั้นตอน แต่ควบคุมสายการคอมไพล์ทั้งหมด” ส่วนหน้าจะใช้ rustc และ stable MIR ซ้ำ ใช้ประโยชน์จากความสามารถในการตรวจสอบชนิด การตรวจสอบการยืม และการปรับแต่ง MIR ที่สมบูรณ์ โดยไม่ต้องสร้างส่วนหน้าคอมไพเลอร์จากศูนย์ ส่วนกลางใช้ pliron (เฟรมเวิร์กคล้าย MLIR) ที่เขียนด้วย Rust ล้วนๆ กำหนดไดอะเล็กต์แบบกำหนดเองสามแบบ ได้แก่ MIR, LLVM และ NVIDIA GPU intrinsics เพื่อทำการแปลง MIR เป็น LLVM IR โดยไม่มี dependencies C++ ส่วนท้ายอาศัยแบ็กเอนด์ LLVM NVPTX ที่สมบูรณ์ สร้างแอสเซมบลี PTX โดยตรง ใช้ประโยชน์จากการปรับแต่งเชิงลึกสำหรับสถาปัตยกรรม GPU คอมไพเลอร์ทั้งหมด (ยกเว้นการเรียก llc ในขั้นตอนสุดท้าย) เขียนด้วย Rust สามารถดีบักและรันผ่านเครื่องมือ Rust มาตรฐาน ทำให้เกิดความสมดุลระหว่างการใช้เครื่องมือที่มีอยู่ซ้ำอย่างมีประสิทธิภาพและการควบคุมสายการคอมไพล์ที่ปรับแต่งได้อย่างสมบูรณ์ โดยคำนึงถึงทั้งประสิทธิภาพการพัฒนาและการปรับแต่งประสิทธิภาพ ตัวอย่างจริง: แผนภาพตัวอย่างการทำงานแบบขนานของการแพร่กระจายไปข้างหน้าสำหรับ MLP สี่ชั้น น้ำหนักที่ใช้ร่วมกันจะถูกอัปโหลดเพียงครั้งเดียวด้วยประเภทการนับอ้างอิงแบบอะตอมมิกของกล่องอุปกรณ์ และโคลนไปยังแต่ละแบตช์ด้วยค่าใช้จ่ายต่ำมาก เพื่อหลีกเลี่ยงการคัดลอกข้อมูลซ้ำซ้อน ตัวจัดตารางเวลาแบบ Round-robin จะจัดสรรแบตช์ไปยังสี่สตรีม โดยไปป์ไลน์ที่สลับกันจะสร้างการทับซ้อนของการดำเนินการบนไทม์ไลน์ของ GPU ฝั่งโฮสต์จะแปลงแต่ละแบตช์เป็นงานอิสระผ่าน Tokio ในขณะที่เฟสต่างๆ เช่น H2D, GEMM บน GPU จะถูกจัดเรียงสลับกัน ทำให้หลายแบตช์ทำงานพร้อมกัน เมื่อเทียบกับการดำเนินการแบบอนุกรมทีละแบตช์ โหมดขนานสามารถใช้ประโยชน์จาก SM ที่ว่างของ GPU ได้อย่างเต็มที่ เวลาทั้งหมดประมาณ 1.3 เท่าของแบตช์เดียวเท่านั้น ช่วยเพิ่มปริมาณงานได้อย่างมาก

unsetunsetสารบัญunsetunset

• หนึ่ง. เริ่มต้นใช้งานอย่างรวดเร็ว: ทำให้ Rust kernel ทำงานก่อน
• สอง. โปรเจกต์นี้ทำอะไรกันแน่: ไม่ใช่การผูก CUDA แต่เป็น rustc codegen backend
• 2.1 crate ฝั่งผู้ใช้และ crate ฝั่งคอมไพเลอร์
• สาม. cargo-oxide: ซ่อนแบ็กเอนด์ที่ซับซ้อนไว้เบื้องหลังคำสั่งย่อย cargo
• สี่. สายการคอมไพล์หลัก: host ไป LLVM, device ไป cuda-oxide
• 4.1 codegen_crate คือประตูแยกเส้นทาง
• 4.2 สะพาน stable_mir: แปลระหว่างชนิดภายในของ rustc กับไปป์ไลน์ของตัวเอง
• 4.3 mir-importer: จาก Rust MIR สู่โลก Pliron
• ห้า. มาโคร #[kernel]: ผู้ใช้เขียนฟังก์ชันปกติ คอมไพเลอร์เห็นสัญลักษณ์ที่สงวนไว้
• หก. โมเดลความปลอดภัยฝั่งอุปกรณ์: หัวใจสำคัญของ safe(ish) อยู่ที่ DisjointSlice และ ThreadIndex
• 6.1 ThreadIndex: สร้างได้จากพิกัดเธรดฮาร์ดแวร์เท่านั้น
• 6.2 DisjointSlice: การเข้าถึงแบบผันแปรได้โดยปริยายต้องมี ThreadIndex
• เจ็ด. รันไทม์โฮสต์: จากไฟล์ PTX ไปยัง cuLaunchKernel
• แปด. ทำไมการปรับแต่ง MIR ถึงส่งผลต่อความถูกต้องของ GPU: คำเตือนของ JumpThreading
• เก้า. Rust รองรับอะไรได้บ้างและไม่สามารถรองรับอะไรบน GPU
• สิบ. ความก้าวหน้าที่แท้จริงของ cuda-oxide: ยกระดับการเขียนโปรแกรม GPU จากปัญหา FFI สู่ปัญหาภาษาแบ็กเอนด์
• บทสรุป: มันไม่ใช่จุดสิ้นสุด แต่เป็นเส้นทางที่ควรจับตามอง

unsetunsetหนึ่ง. เริ่มต้นใช้งานอย่างรวดเร็ว: ทำให้ Rust kernel ทำงานก่อนunsetunset

ปัจจุบัน cuda-oxide อยู่ในสถานะโครงการ alpha ระยะแรก README ระบุชัดเจนว่ายังเป็นคอมไพเลอร์เชิงทดลอง API และฟังก์ชันอาจเปลี่ยนแปลง เส้นทางประสบการณ์ที่น้อยที่สุดไม่ใช่การเรียก rustc ด้วยตนเอง แต่ใช้คำสั่งย่อย cargo-oxide ที่โครงการจัดเตรียมไว้

ข้อกำหนดการกำหนดค่าสภาพแวดล้อม

การรันโปรเจกต์นี้ต้องเป็นไปตามข้อกำหนดสภาพแวดล้อมที่เข้มงวดหลายประการ: ต้องใช้ Rust nightly เวอร์ชัน (ในคลังกำหนดเป็น nightly-2026-04-03 และต้องติดตั้งคอมโพเนนต์ rust-src และ rustc-dev), CUDA Toolkit 12.x หรือสูงกว่า, LLVM 21+ ที่มีแบ็กเอนด์ NVPTX, ไฟล์ส่วนหัว Clang/libclang และระบบปฏิบัติการ Linux ต้องเน้นย้ำเป็นพิเศษว่า LLVM 21 ไม่ใช่เงื่อนไขตกแต่งที่ขาดไม่ได้: ตามเอกสาร README โปรเจกต์นี้จะปล่อย GPU intrinsic ที่ใหม่กว่า เช่น TMA, tcgen05, WGMMA ในขณะที่ LLVM 20 และเวอร์ชันก่อนหน้าไม่สามารถจัดการคำสั่งเหล่านี้ได้อย่างสมบูรณ์

# ที่มา: README.md  

# ติดตั้ง cargo-oxide ในโปรเจกต์ภายนอก  
cargo install --git https://github.com/NVlabs/cuda-oxide.git cargo-oxide  

# หากต้องการเตรียม Rust nightly ด้วยตนเอง  
rustup toolchain install nightly-2026-04-03  
rustup component add rust-src rustc-dev --toolchain nightly-2026-04-03  

# ตรวจสอบ CUDA และ LLVM  
export PATH="/usr/local/cuda/bin:$PATH"  
nvcc --version  
llc-21 --version | grep nvptx  

# ตรวจสอบสภาพแวดล้อม  
cargo oxide doctor  

# สร้างและรันตัวอย่างการบวกเวกเตอร์ที่คลาสสิกที่สุด  
cargo oxide run vecadd  

# ดูสายการคอมไพล์ทั้งหมด  
cargo oxide pipeline vecadd  

# ดีบักด้วย cuda-gdb  
cargo oxide debug vecadd --tui  

หากคุณต้องการสร้างโปรเจกต์อิสระของคุณเอง crates/cargo-oxide/README.md มีคำสั่งเทมเพลตที่ตรงไปตรงมามากกว่า:

# ที่มา: crates/cargo-oxide/README.md  

cargo oxide new my_kernel  
cd my_kernel  
cargo oxide run  

# เทมเพลตอะซิงโครนัส: tokio + cuda-async  
cargo oxide new my_project --async  

รายละเอียดเพิ่มเติมเกี่ยวกับ dependencies ก่อนหน้า เวอร์ชัน LLVM รายละเอียดการติดตั้ง CUDA/Clang สามารถดูได้จาก README ของโปรเจกต์และ cuda-oxide-book สำหรับผู้เริ่มต้น เกณฑ์ตัดสินที่สำคัญที่สุดคือ: cargo oxide doctor ผ่านหรือไม่ และ cargo oxide run vecadd แสดงผล ✓ SUCCESS: All 1024 elements correct! หรือไม่

สอง. โปรเจกต์นี้ทำอะไรกันแน่: ไม่ใช่การผูก CUDA แต่เป็น rustc codegen backend

README ของ cuda-oxide นิยามมันว่าเป็น “custom rustc backend for compiling GPU kernels in pure Rust” คำสำคัญในประโยคนี้ไม่ใช่ GPU หรือ Rust แต่เป็น rustc backend

นั่นหมายความว่ามันไม่ใช่การห่อหุ้ม CUDA C API ใน Rust หรือการแปล Rust AST เป็น DSL บางอย่าง มันเข้าไปแทรกแซงในขั้นตอน codegen ของ rustc โดยตรง ซอร์สโค้ดยังคงผ่านส่วนหน้า Rust มาตรฐาน: การแยกวิเคราะห์, HIR, การตรวจสอบชนิด, การสร้าง MIR, การปรับแต่ง MIR เมื่อถึงขั้นตอน codegen แบ็กเอนด์ของ cuda-oxide จะทำหน้าที่ระบุฟังก์ชันอุปกรณ์ที่ทำเครื่องหมายด้วย #[kernel] และดึงกราฟการเรียกฝั่งอุปกรณ์ที่เข้าถึงได้ออกมา ผ่านสายการคอมไพล์อุปกรณ์แยกต่างหาก MIR → Pliron → LLVM IR → PTX ส่วนโค้ดโฮสต์ทั่วไปจะยังคงถูกคอมไพล์เป็นไบนารีฝั่ง CPU โดย rustc_codegen_llvm มาตรฐาน

ไฟล์ Cargo.toml ที่รากของคลังได้เปิดเผยการแบ่งชั้นนี้แล้ว:

# ที่มา: Cargo.toml  

[workspace]  
members = [  
# Core crates  
"crates/cuda-device",  
"crates/cuda-host",  
"crates/cuda-macros",  
"crates/dialect-llvm",  
"crates/dialect-mir",  
"crates/dialect-nvvm",  
"crates/mir-importer",  
"crates/mir-lower",  
"crates/cargo-oxide",  

# FFI bindings  
"crates/cuda-bindings",  
"crates/cuda-core",  
"crates/cuda-async",  
"crates/libnvvm-sys",  
"crates/nvjitlink-sys",  
]  

มีรายละเอียดที่ควรสังเกต: crates/rustc-codegen-cuda ไม่ใช่สมาชิกของ workspace หมายเหตุใน Cargo.toml รากระบุว่ามันต้องการฟีเจอร์ nightly rustc พิเศษและขั้นตอนการสร้างที่แตกต่างกัน นี่แสดงให้เห็นว่า “เครื่องยนต์หลัก” ของ cuda-oxide ไม่ใช่แค่ Rust crate ทั่วไป แต่เป็นแบ็กเอนด์การคอมไพล์ .so ที่จะถูกโหลดแบบไดนามิกโดย rustc

2.1 crate ฝั่งผู้ใช้และ crate ฝั่งคอมไพเลอร์

จากภาพรวม crate ใน README โปรเจกต์แบ่งออกเป็นสามชั้นโดยประมาณ:

ชั้นแรกคือ API ที่ผู้ใช้สัมผัสโดยตรง:

• cuda-device: intrinsic ฝั่งอุปกรณ์, ดัชนีเธรด, warp, barrier, หน่วยความจำที่ใช้ร่วมกัน, TMA ฯลฯ
• cuda-host: มาโคร launch ฝั่งโฮสต์, การโหลดโมดูล, ตัวช่วย LTOIR
• cuda-core: การห่อหุ้ม RAII ที่ปลอดภัยของ CUDA Driver API เช่น CudaContext, CudaStream, DeviceBuffer<T>
• cuda-async: ห่อหุ้มการทำงานของ GPU เป็น DeviceOperation แบบขี้เกียจ, ประกอบกันได้ และสามารถ .await ได้
• cuda-macros: จัดหามาโครกระบวนการ เช่น #[kernel], #[device], cuda_launch!

ชั้นที่สองคือภายในคอมไพเลอร์:

• rustc-codegen-cuda: แบ็กเอนด์ codegen rustc ที่กำหนดเอง
• mir-importer: แปล Rust MIR เป็น dialect-mir และขับเคลื่อนไปป์ไลน์ที่ตามมา
• mir-lower: ลดระดับจาก dialect-mir ไปเป็น dialect-llvm
• dialect-mir, dialect-llvm, dialect-nvvm: ไดอะเล็กต์ IR ที่ใช้ Pliron

ชั้นที่สามคือเครื่องมือ:

• cargo-oxide: คำสั่งย่อย cargo ที่面向ผู้ใช้ รับผิดชอบการสร้างแบ็กเอนด์ การส่งผ่านพารามิเตอร์ การรันตัวอย่าง การดีบัก และการตรวจสอบสภาพแวดล้อม

สถาปัตยกรรมแบบแบ่งชั้นนี้本质上สร้าง “สะพาน” ขึ้นมา: ปลายด้านหนึ่งคือโค้ด #[kernel] fn vecadd(...) ที่นักพัฒนา Rust เขียน อีกด้านหนึ่งคือไฟล์ .ptx หรือ .cubin ที่ CUDA driver สามารถโหลดและดำเนินการได้ การเชื่อมต่อระหว่างกลางไม่ใช่การแทนที่ข้อความอย่างง่าย แต่เป็นแบ็กเอนด์คอมไพเลอร์ที่สมบูรณ์

สาม. cargo-oxide: ซ่อนแบ็กเอนด์ที่ซับซ้อนไว้เบื้องหลังคำสั่งย่อย cargo

ผู้ใช้ทั่วไปไม่จำเป็นต้องใช้ rustc -Z codegen-backend=... โดยตรง คุณค่าของ cargo-oxide อยู่ที่การรวมขั้นตอนการสร้างแบ็กเอนด์ การกำหนดค่าตัวแปรสภาพแวดล้อม การจัดการเส้นทางตัวอย่าง และคำสั่งรันทั้งหมดไว้ในจุดเข้าใช้งานเดียว

คำจำกัดความ CLI ใน crates/cargo-oxide/src/main.rs นั้นตรงไปตรงมามาก:

// ที่มา: crates/cargo-oxide/src/main.rs  

enum Commands {  
Run {  
example: Option<String>,  
#[arg(long)]  
dlto: bool,  
#[arg(long)]  
emit_nvvm_ir: bool,  
#[arg(long)]  
arch: Option<String>,  
#[arg(long)]  
features: Option<String>,  
},  
Build { /* compile only */ },  
Pipeline { example: String, /* show MIR → PTX pipeline */ },  
Debug { example: String, cgdb: bool, tui: bool },  
Fmt { check: bool },  
New { name: String, async_mode: bool },  
Doctor,  
Setup,  
}

ส่วนที่สำคัญที่สุดคือตรรกะการค้นหาแบ็กเอนด์ cargo-oxide จะค้นหา librustc_codegen_cuda.so ตามลำดับความสำคัญที่กำหนดไว้: ตรวจสอบตัวแปรสภาพแวดล้อม CUDA_OXIDE_BACKEND ก่อน จากนั้นตรวจสอบคลังในเครื่อง แล้วค้นหาไดเรกทอรีแคช และสุดท้ายดำเนินการโคลนคลังแบบตื้นและสร้างโดยอัตโนมัติ

// ที่มา: crates/cargo-oxide/src/backend.rs  

pub fn find_or_build_backend(workspace_root: &Path) -> PathBuf {  
if let Ok(path) = std::env::var("CUDA_OXIDE_BACKEND") {  
let p = PathBuf::from(&path);  
if p.exists() {  
return p;  
}  
}  

let codegen_crate = workspace_root.join("crates/rustc-codegen-cuda");  
if codegen_crate.is_dir() {  
let so_path = codegen_crate.join("target/debug/librustc_codegen_cuda.so");  
build_backend_from_source(&codegen_crate);  
return so_path;  
}  

if let Some(cache_dir) = cache_directory() {  
let cached_so = cache_dir.join("librustc_codegen_cuda.so");  
if cached_so.exists() {  
return cached_so;  
}  
}  

auto_fetch_and_build()  
}

เมื่อมองดูเผินๆ การออกแบบนี้ดูเหมือนเป็นเพียง “การปรับปรุงประสบการณ์การติดตั้ง” แต่สำหรับคอมไพเลอร์เชิงทดลองแล้ว มันมีความสำคัญอย่างยิ่ง เนื่องจาก cuda-oxide พึ่งพา API ส่วนตัวของ rustc, เครื่องมือ nightly, LLVM เวอร์ชันเฉพาะ และเครื่องมือ CUDA อย่างมาก หากให้ผู้ใช้ประกอบส่วนประกอบเหล่านี้ด้วยตนเอง ขั้นตอนแรกก็จะถูกขัดขวางด้วยการกำหนดค่าสภาพแวดล้อมที่ซับซ้อน ฟังก์ชัน cargo oxide doctor และการสร้างแบ็กเอนด์อัตโนมัติ本质上เป็นการจัดหาจุดเข้าใช้งานที่สามารถดำเนินการได้สำหรับระบบคอมไพเลอร์ที่ไม่เสถียรแต่มีฟังก์ชันซับซ้อน

สี่. สายการคอมไพล์หลัก: host ไป LLVM, device ไป cuda-oxide

ไฟล์หลักที่สุดของ cuda-oxide คือ crates/rustc-codegen-cuda/src/lib.rs เอกสารที่ด้านบนของไฟล์นี้ได้อธิบายสถาปัตยกรรมโดยรวมอย่างชัดเจนแล้ว:

• ส่วนหน้า rustc สร้าง MIR ที่ถูกปรับแต่งแล้ว
• แบ็กเอนด์ cuda-oxide สแกนหา kernel ที่จุดเข้า codegen_crate
• โค้ดอุปกรณ์เข้าสู่ไปป์ไลน์ของตัวเอง ส่วนโค้ดโฮสต์จะถูกจัดการโดย LLVM backend มาตรฐาน

4.1 codegen_crate ในฐานะประตูแยกเส้นทาง

โค้ดต่อไปนี้เป็นส่วนสำคัญที่ทำหน้าที่เป็น “ศูนย์กลางการจราจร” ในทั้งโปรเจกต์:

// ที่มา: crates/rustc-codegen-cuda/src/lib.rs  

fn codegen_crate(&self, tcx: TyCtxt<'_>, crate_info: &CrateInfo) -> Box<dyn Any> {  
with_no_trimmed_paths!({  
let mono_partitions = tcx.collect_and_partition_mono_items(());  
let kernel_count = collector::count_kernels_in_cgus(tcx, mono_partitions.codegen_units);  
let device_fn_count =  
collector::count_device_fns_in_cgus(tcx, mono_partitions.codegen_units);  

let has_device_code = kernel_count > 0 || device_fn_count > 0;  

if has_device_code {  
let collection_result = collector::collect_device_functions(  
tcx,  
mono_partitions.codegen_units,  
self.config.verbose,  
);  

let device_config = device_codegen::DeviceCodegenConfig {  
output_dir: self.config.ptx_output_dir.clone().unwrap_or_else(|| {  
std::env::current_dir().unwrap_or_else(|_| ".".into())  
}),  
output_name: tcx.crate_name(rustc_hir::def_id::LOCAL_CRATE).to_string(),  
verbose: self.config.verbose,  
dump_rustc_mir: self.config.dump_rustc_mir,  
dump_mir_dialect: self.config.dump_mir_dialect,  
dump_llvm_dialect: self.config.dump_llvm_dialect,  
};  

device_codegen::generate_device_code(  
tcx,  
&collection_result.functions,  
&collection_result.device_externs,  
&device_config,  
).unwrap_or_else(|e| {  
tcx.dcx().fatal(format!(  
"[rustc_codegen_cuda] Device codegen failed: {}", e  
));  
});  
}  

// Host code ถูกส่งต่อไปยัง LLVM backend มาตรฐานอย่างสมบูรณ์  
self.llvm_backend.codegen_crate(tcx, crate_info)  
})  
}  

ตรรกะส่วนนี้มีประเด็นการออกแบบหลักสามประการ

• ประการแรก มันพึ่งพาผลลัพธ์ของ monomorphization ของ rustc อย่างมาก kernel แบบเจนเนอริกไม่ได้สร้างเทมเพลตโดยตรงในระดับซอร์สโค้ด แต่รอให้ rustc สร้างอินสแตนซ์แบบ monomorphized ตามจุดเรียกจริง จากนั้นแบ็กเอนด์จึงรวบรวมและประมวลผลอย่างเป็นหนึ่งเดียว
• ประการที่สอง มันใช้การเข้าถึงได้ของ kernel เพื่อกำหนดขอบเขตระหว่างอุปกรณ์และโฮสต์ ไม่ใช่ทุกฟังก์ชันที่จะถูกรวมไว้ในขั้นตอนการคอมไพล์ GPU เฉพาะฟังก์ชันที่เริ่มต้นจาก kernel หรือ device function และสามารถเข้าถึงได้ผ่านห่วงโซ่การเรียกเท่านั้นที่จะถูกรวบรวม
• สุดท้าย มันไม่รบกวนกระบวนการคอมไพล์ฝั่งโฮสต์ โค้ดที่面向 CPU ทั้งหมดยังคงถูกจัดการโดย rustc_codegen_llvm กลยุทธ์การออกแบบแบบ “ห่อหุ้ม LLVM backend” นี้ช่วยลดความซับซ้อนในการดำเนินงานของโครงการได้อย่างมาก พร้อมทั้งหลีกเลี่ยงภาระงานมหาศาลในการต้องนำโฮสต์ codegen ของ Rust ทั้งหมดมาใช้ใหม่

4.2 สะพาน stable_mir: แปลระหว่างชนิดภายในของ rustc กับไปป์ไลน์ของตัวเอง

ตัวสร้างโค้ดอุปกรณ์: สะพานจากการแสดงผลภายในสู่อินเทอร์เฟซที่เสถียร

ไฟล์ device_codegen.rs เปิดเผยปัญหาทางวิศวกรรมคอมไพเลอร์ที่สมจริงมาก: แบ็กเอนด์การสร้างโค้ดรับชนิดภายใน rustc_middle แต่คอมโพเนนต์ mir-importer ที่มีอยู่ของ cuda-oxide ชอบใช้อินเทอร์เฟซ rustc_public (หรือ stable MIR) มากกว่า เพื่อแก้ไขข้อขัดแย้งนี้ ระบบจำเป็นต้องสร้างเลเยอร์การแปลงกลาง

// ที่มา: crates/rustc-codegen-cuda/src/device_codegen.rs  

let result = rustc_internal::run(tcx, || {  
let stable_functions: Vec<mir_importer::CollectedFunction> = functions  
.iter()  
.zip(export_names.iter())  
.filter_map(|(func, (export_name, is_kernel))| {  
let stable_instance = rustc_internal::stable(func.instance);  

Some(mir_importer::CollectedFunction {  
instance: stable_instance,  
is_kernel: *is_kernel,  
export_name: export_name.clone(),  
})  
})  
.collect();  

let pipeline_config = mir_importer::PipelineConfig {  
output_dir: output_dir.clone(),  
output_name: output_name.clone(),  
verbose,  
show_mir_dialect: show_mir,  
show_llvm_dialect: show_llvm,  
emit_ltoir,  
ltoir_arch: ltoir_arch.clone(),  
emit_nvvm_ir,  
};  

// Rust MIR → dialect-mir → mem2reg → dialect-llvm → LLVM IR → PTX  
mir_importer::run_pipeline(  
&stable_functions,  
&stable_device_externs,  
&pipeline_config,  
)  
});  

แนวคิดการออกแบบหลักของโค้ดนี้คือ: cuda-oxide ไม่ได้ดำเนินการลดระดับทั้งหมดบน MIR ภายในของ rustc โดยตรง แต่จะแปลงข้อมูลเป็นการแสดงผลกลางที่เสถียรกว่าก่อน จากนั้นจึงส่งต่อไปยัง mir-importer นี่เป็นทางเลือกทางวิศวกรรมที่ปฏิบัติได้จริงมาก แม้จะเสียค่าใช้จ่ายในการแปลงชนิดครั้งหนึ่ง แต่ก็แลกมาด้วยความสามารถในการนำไปใช้ซ้ำของไปป์ไลน์คอมไพล์ทั้งหมดและการบำรุงรักษาที่ดีขึ้น

mir-importer: จาก Rust MIR สู่โลก Pliron

เป้าหมายการออกแบบของ mir-importer คือการแปล Rust MIR เป็น dialect-mir จากนั้นขับเคลื่อนกระบวนการลดระดับที่ตามมา มันใช้รูปแบบ alloca + load/store เพื่อแสดงตัวแปรท้องถิ่น จากนั้นยกระดับกลับเป็นรูปแบบ SSA ผ่าน mem2reg กลยุทธ์นี้สอดคล้องกับแนวทางของส่วนหน้า LLVM แบบดั้งเดิมอย่างมาก: สร้างรูปแบบหน่วยความจำที่เรียบง่ายก่อน จากนั้นจึงคืนคุณภาพ SSA ผ่าน pass การปรับแต่ง

// ที่มา: crates/mir-importer/src/pipeline.rs  

pub fn run_pipeline(  
functions: &[CollectedFunction],  
device_externs: &[DeviceExternDecl],  
config: &PipelineConfig,  
) -> Result<CompilationResult, PipelineError> {  
let mut ctx = Context::new();  

crate::translator::register_dialects(&mut ctx);  

let module = pliron::builtin::ops::ModuleOp::new(&mut ctx, module_name);  
let module_op_ptr = module.get_operation();  

for func in functions {  
let body = func  
.instance  
.body()  
.ok_or_else(|| PipelineError::NoBody(func.export_name.clone()))?;  

let func_op_ptr = crate::translator::body::translate_body(  
&mut ctx,  
&body,  
&func.instance,  
func.is_kernel,  
Some(&func.export_name),  
&mut legaliser,  
)?;  

verify_operation(&ctx, func_op_ptr, &func.export_name)?;  
append_to_module(&ctx, module_op_ptr, func_op_ptr);  
}  

verify_operation(&ctx, module_op_ptr, "module")?;  

// จากนั้นรัน mem2reg แล้วลดระดับไปยัง dialect-llvm, ส่งออก LLVM IR, และเรียก llc เพื่อสร้าง PTX  
}  

หากเปรียบเทียบกระบวนการคอมไพล์ทั้งหมดกับการแปลงานวรรณกรรม ส่วนหน้า rustc จะแปลซอร์สโค้ด Rust เป็น “ภาษากลางที่มีความหมายชัดเจน” ที่เรียกว่า MIR ก่อน จากนั้น cuda-oxide จะแปล MIR เป็น Pliron dialect จากนั้น Pliron dialect จะลดระดับเป็น LLVM IR และสุดท้ายแบ็กเอนด์ LLVM NVPTX จะแปลงเป็น PTX แต่ละชั้นจะคงความหมายบางส่วนไว้ พร้อมกับค่อยๆ เข้าใกล้ฮาร์ดแวร์มากขึ้น

มาโคร #[kernel]: ผู้ใช้เขียนฟังก์ชันปกติ คอมไพเลอร์เห็นสัญลักษณ์ที่สงวนไว้

好的，以下是针对您提供的文章片段进行的深度重写与降重结果。我已严格遵守所有规则，确保原意不变，同时优化了表达和排版。

หก. โมเดลความปลอดภัยฝั่งอุปกรณ์: DisjointSlice และ ThreadIndex คือหัวใจสำคัญของความปลอดภัย (ish)

ในการเขียนโปรแกรม GPU การเขียนแบบขนานเป็นปัญหาด้านความปลอดภัยที่ใหญ่ที่สุด ความหมายของ &mut T ใน Rust ต้องการการอ้างอิงแบบผันแปรได้แบบเอกสิทธิ์ แต่เคอร์เนล CUDA หนึ่งตัวอาจรันเธรดนับพันหรือนับหมื่นพร้อมกัน หากแต่ละเธรดสามารถรับ &mut [T] องค์ประกอบใดๆ ได้ตามอำเภอใจ กฎ别名ของ Rust จะใช้ไม่ได้ทันที

cuda-oxide ไม่ได้อ้างว่า “GPU ปลอดภัยโดยธรรมชาติ” แต่แก้ปัญหาผ่านการออกแบบนามธรรมคู่หนึ่งคือ ThreadIndex และ DisjointSlice<T>

6.1 ThreadIndex: สร้างได้จากพิกัดเธรดฮาร์ดแวร์เท่านั้น

ThreadIndex เป็น newtype ที่โปร่งใส แต่คอนสตรัคเตอร์ของมันไม่ได้เปิดเผยต่อสาธารณะ โดยทั่วไปผู้ใช้จะได้รับอินสแตนซ์ผ่าน thread::index_1d() หรือ index_2d()

// ที่มา: crates/cuda-device/src/thread.rs

#[derive(Clone, Copy, Debug)]
#[repr(transparent)]
pub struct ThreadIndex(usize);

impl ThreadIndex {
#[inline(always)]
pub fn get(self) -> usize {
self.0
}
}

#[inline(always)]
pub fn index_1d() -> ThreadIndex {
let tid = threadIdx_x();
let bid = blockIdx_x();
let bdim = blockDim_x();

// bid * bdim + tid สร้างดัชนีเฉพาะสำหรับแต่ละเธรดในกริด 1D
ThreadIndex((bid * bdim + tid) as usize)
}

threadIdx_x(), blockIdx_x() และ blockDim_x()

---

⚠️ หมายเหตุ: เนื้อหาได้รับการแปลโดย AI และตรวจสอบโดยมนุษย์ หากมีข้อผิดพลาดโปรดแจ้ง

☕ สนับสนุนค่ากาแฟทีมงาน

หากคุณชอบบทความนี้ สามารถสนับสนุนเราได้ผ่าน PromptPay

SCAN TO PAY WITH ANY BANK

本文来自网络搜集，不代表คลื่นสร้างอนาคต立场，如有侵权，联系删除。转载请注明出处：https://www.itsolotime.com/th/archives/33571