هسته های GPU بومی در فضاهای کاری Cargo

از NVIDIA cuda-rust زنجیر ابزار همراه با nvptx64-nvidia-cuda هدف، لایه‌های میانی را که مدت‌ها در توسعه پردازنده‌های گرافیکی Rust با مشکل مواجه بودند، حذف می‌کند. این مقاله به ساخت یک فضای کار کامل بار با دو جعبه می پردازد: الف gpu-kernel جعبه ای که یک هسته Rust را به بایت کد NVPTX کامپایل می کند و a host-runner جعبه ای که این هسته را در یک رشته دستگاه CUDA اجرا می کند و نتایج را تأیید می کند.

محتوا

وضعیت: را cuda-rust اکوسیستم جعبه (cuda_std، cuda_kernel، cuda_driverتوضیح داده شده در این مقاله، هنوز به صورت عمومی در crates.io منتشر نشده است. نمونه‌های کد، گردش کار مورد نظر را تشریح می‌کنند. قبل از تلاش برای بازتولید این مراحل، در دسترس بودن جعبه را در crates.io یا مخازن رسمی GitHub NVIDIA بررسی کنید. ساخت کیس دستگاه نیاز به یک هر شب زنگ بزن زنجیر ابزار با rust-src جزء؛ در nvptx64-nvidia-cuda هدف است سطح-3 و از طریق قابل دسترسی نیست rustup target add روی زنگ پایدار

چرا هسته های GPU بومی در Rust اکنون مهم هستند

برای سال‌ها، توسعه‌دهندگان Rust که به دنبال بارهای کاری با سرعت پردازش گرافیکی هستند، با واقعیتی ناراحت‌کننده مواجه شده‌اند. نوشتن هسته های GPU بومی در Rust به معنای مونتاژ زنجیره های ابزار شکننده است: کد چسبنده C++/CUDA که از طریق NVCC کامپایل شده است، cc هک برای ساخت جعبه هایی که کامپایلرهای خارجی، بسته بندی های خطرناک FFI را در اطراف PTX کامپایل شده دستی یا پروژه های آزمایشی مانند rust-gpu که هرگز برای محاسبات عمومی به طور کامل تثبیت نمی شوند. هر رویکرد اصطکاک، شکاف های ایمنی یا هر دو را معرفی کرد.

این وضعیت در حال تغییر است. از NVIDIA cuda-rust زنجیر ابزار همراه با nvptx64-nvidia-cuda target، این لایه های میانی را در نمای فعلی خود حذف می کند. توسعه دهندگان Rust استانداردی را تدوین می کنند no_std با استفاده از بایت کد NVPTX مستقیما زنگ بزنید rustcباطن اصلی LLVM، آن بایت کد را از طریق اتصالات درایور Rust native CUDA بارگیری کنید و هسته ها را روی سخت افزار GPU بدون نوشتن یا کامپایل کردن یک خط C++ اجرا کنید.

این مقاله به ساخت یک فضای کار کامل بار با دو جعبه می پردازد: الف gpu-kernel جعبه ای که یک هسته Rust را به بایت کد NVPTX کامپایل می کند و a host-runner جعبه ای که این هسته را در یک رشته دستگاه CUDA اجرا می کند و نتایج را تأیید می کند. هر مرحله از ابزارهای Cargo و بومی Rust API استفاده می کند.

پیش نیازها:

• زنگ زدگی هر شب زنجیره ابزار (مثلا nightly-2025-06-01 یا بعد از آن)، ص rust-src جزء نصب شده
• CUDA Toolkit 12.x (آخرین نسخه را در آرشیو CUDA Toolkit بررسی کنید)
• پردازنده گرافیکی NVIDIA با Compute 7.0 یا بالاتر
• لینوکس (یا WSL2)
• درایور کارت گرافیک NVIDIA سازگار (یادداشت های انتشار CUDA Toolkit را برای حداقل نسخه های درایور بررسی کنید)

چگونه nvptx64-nvidia-cuda با خطوط لوله سنتی NVCC متفاوت است

مسیر ساخت اساساً با توسعه CUDA قدیمی در Rust متفاوت است. اس nvptx64-nvidia-cuda هدف، جریان به این صورت عمل می کند: ارسال کد منبع Rust rustc با استفاده از پشتیبان NVPTX LLVM که متن اسمبلی PTX را مستقیماً منتشر می کند. سپس درایور CUDA این PTX را در زمان اجرا بارگذاری می کند و JIT-را از طریق معماری GPU خاص کامپایل می کند. cuModuleLoadData یا API های معادل.

این را با مسیر سنتی مقایسه کنید: کد میزبانی که Rust استفاده می کند bindgen یا در cc جعبه پرداخت NVCC که کامپایل می شود .cu فایل ها در .cubin یا .ptx مصنوعات سپس Rust از طریق فراخوانی های FFI ناامن که API درایور CUDA را بسته بندی می کند، به این مصنوعات دسترسی پیدا می کند. هر عبور از مرز، از Rust تا سیستم ساخت ++C تا NVCC تا لایه FFI، پتانسیل انواع ناهماهنگ، رفتار تعریف‌نشده و آسیب‌پذیری سیستم را معرفی می‌کند.

کد دستگاه باید باشد no_std زیرا زمینه های اجرای GPU فاقد سیستم عامل، تخصیص دهنده استاندارد و دسترسی به تماس های سیستم سمت میزبان نیستند. را no_std محدودیت یک محدودیت زنجیره ابزار نیست، بلکه بازتابی از مدل اجرای سخت افزار است.

موارد و اجزای کلیدی

از NVIDIA cuda-rust اکوسیستم جعبه سه جزء اصلی را فراهم می کند. cuda_std کتابخانه استاندارد سمت دستگاه را تامین می‌کند: موارد اولیه نمایه‌سازی رشته، اعلان‌های حافظه مشترک، موانع همگام‌سازی، و اجزای داخلی ریاضی. cuda_kernel فراهم می کند #[cuda_kernel] ماکرو ویژگی که توابع کامپایل سمت دستگاه را علامت گذاری می کند و ترجمه ABI مورد نیاز برای نقاط ورودی PTX را مدیریت می کند. cuda_driver اتصالات بومی Rust را به API درایور CUDA برای بارگذاری ماژول، مدیریت حافظه، راه‌اندازی هسته و همگام‌سازی نخ سمت میزبان نشان می‌دهد.

توجه: ممکن است این جعبه‌ها هنوز در crates.io در دسترس نباشند. مخازن رسمی NVIDIA را برای URL های منبع بررسی کنید و آنها را به عنوان اضافه کنید git وابستگی ها با تجدید نظر خاص در صورت نیاز:

cuda_std = { git = "https://github.com/NVIDIA/cuda-rust", rev = "..." }

را nvptx64-nvidia-cuda دستورالعمل های سه گانه را هدف قرار دهید rustc برای انتشار NVPTX IR از طریق LLVM به جای کد ماشین بومی. CUDA Toolkit توسط سیستم میزبان برای سازگاری درایور و دسترسی به کتابخانه زمان اجرا مورد نیاز است. خط لوله ساخت Rust NVCC را صدا نمی کند، اما جعبه ابزار CUDA کتابخانه درایور را فراهم می کند (libcuda.so) برای اتصال لازم است.

چرا یک فضای کاری بار با جعبه های متعدد؟

کد دستگاه و کد میزبان نمی توانند در یک جعبه همزمان وجود داشته باشند. الف no_std با هدف قرار دادن صندوق ها nvptx64-nvidia-cuda اساساً با الف ناسازگار است std پرداخت با هدف قرار دادن سه قلو بومی میزبان. فضاهای کاری بار این موضوع را به وضوح حل می کنند: جعبه دستگاه در یک PTX تحت یک هدف کامپایل می شود، جعبه میزبان در باینری خود تحت دیگری کامپایل می شود و هر دو یک ریشه فضای کاری مشترک دارند. یک جعبه مشترک نازک می تواند در صورت نیاز حاوی انواع یا ثابت های عمومی باشد.

توسعه دهندگان Rust استانداردی را تدوین می کنند no_std با استفاده از بایت کد NVPTX مستقیما زنگ بزنید rustcباطن اصلی LLVM، آن بایت کد را از طریق اتصالات درایور Rust native CUDA بارگیری کنید و هسته ها را روی سخت افزار GPU بدون نوشتن یا کامپایل کردن یک خط C++ اجرا کنید.

راه اندازی فضای کاری بار

سطح فضای کاری Cargo.toml

ریشه فضای کاری هر دو را به عنوان عضو تعریف می کند و از حل کننده نسخه 2 برای ادغام صحیح توابع بین اهداف استفاده می کند:

[workspace]
members = ["gpu-kernel", "host-runner"]
resolver = "2"

[workspace.dependencies]
cuda_std = "0.3"
cuda_kernel = "0.3"
cuda_driver = "0.3"

مهم: اگر این جعبه‌ها در crates.io منتشر نشده‌اند، رشته‌های نسخه را جایگزین کنید git حاشیه نویسی به منبع که به مخزن رسمی و یک ویرایش یا برچسب خاص اشاره دارد.

این ساختار تضمین می‌کند که نسخه‌های وابستگی در هر دو مخزن همگام می‌مانند، در حالی که به هر کدام اجازه می‌دهد پیکربندی هدف خاص خود را مشخص کند.

CUDA-Rust Toolchain و NVPTX Target را نصب کنید

قبل از نوشتن Rust، محیط ساخت باید پیکربندی شود:

rustup toolchain install nightly


rustup component add rust-src --toolchain nightly


rustup run nightly rustc --version


nvidia-smi


ls /usr/local/cuda/lib64/libcuda*

را nvptx64-nvidia-cuda هدف یک است سطح-3 هدف از طریق در دسترس نیست rustup target add روی هر زنجیره ساز در عوض، کتابخانه استاندارد (core) باید با استفاده از منبع ساخته شود -Z build-std=core در عصر این از طریق پیکربندی شده است .cargo/config.toml در جعبه دستگاه (به بخش بعدی مراجعه کنید).

پیکربندی .cargo/config.toml

را gpu-kernel جعبه نیاز به یک .cargo/config.toml به هدف سطح 3 اشاره کرده و آن را فعال کنید build-std:

[build]
target = "nvptx64-nvidia-cuda"




[unstable]
build-std = ["core"]

بدون این فایل، cargo build زیرا جعبه دستگاه با خطاهای از دست رفته از کار می افتد core مصنوعات

توجه: را panic = "abort" تنظیمات پروفایل متعلق به Cargo.toml، نه در .cargo/config.toml. را .cargo/config.toml فایل پشتیبانی نمی کند [profile.*] جداول; چنین بخش هایی در سکوت نادیده گرفته می شوند. نگاه کن gpu-kernel/Cargo.toml پیکربندی زیر برای قرارگیری مناسب

سنجاق کردن Nightly Toolchain

a اضافه کنید rust-toolchain.toml در ریشه فضای کاری برای اطمینان از ساخت های قابل تکرار:

[toolchain]
channel = "nightly"
components = ["rust-src"]

ساختار دایرکتوری پروژه

طرح کامل فضای کاری:

cuda-rust-workspace/
├── Cargo.toml                  # Workspace root
├── rust-toolchain.toml         # Pins nightly toolchain
├── gpu-kernel/
│   ├── .cargo/
│   │   └── config.toml         # NVPTX target + build-std config
│   ├── Cargo.toml              # Device crate config
│   └── src/
│       └── lib.rs              # GPU kernel source (no_std)
└── host-runner/
    ├── Cargo.toml              # Host crate config
    ├── build.rs                # Build script cross-compiling gpu-kernel
    ├── src/
    │   └── main.rs             # Host-side launch code
    └── tests/
        └── integration.rs      # Kernel output validation tests

هر فایل در این درخت در بخش های زیر ایجاد می شود. را build.rs ج host-runner چسبی است که آن را می سازد cargo build از ریشه فضای کاری، کد دستگاه و میزبان را در یک تماس کامپایل کنید.

نوشتن جعبه کرنل GPU

gpu-kernel/Cargo.toml پیکربندی

جعبه هسته برای تولید خروجی PTX به جای دودویی بومی به پیکربندی خاصی نیاز دارد:

[package]
name = "gpu-kernel"
version = "0.1.0"
edition = "2021"

[lib]
crate-type = ["rlib"]

[dependencies]
cuda_std = { workspace = true }
cuda_kernel = { workspace = true }

[profile.release]
opt-level = 3
lto = true
codegen-units = 1
panic = "abort"

متن PTX با پیمایش استخراج می شود RUSTFLAGS="--emit=asm" به rustc در حین ساخت دستگاه؛ هیچ کتابخانه مشترکی ایجاد نمی شود. را rlib نوع جعبه مناسب برای a no_std جعبه کتابخانه را panic = "abort" تنظیم مورد نیاز است زیرا هدف NVPTX از باز کردن پشته پشتیبانی نمی کند. تنظیمات پروفایل انتشار بر کیفیت PTX منتشر شده تأثیر می‌گذارد: LTO و ماژول تک کد می‌توانند اندازه PTX را کاهش دهند و قابلیت پیش‌بینی تخصیص ثبت را بهبود بخشند، که مستقیماً بر اشغال GPU تأثیر می‌گذارد.

نوشتن الف no_std وکتور-افزودن هسته

هسته خود Rust استاندارد با ویژگی‌ها و ویژگی‌های درونی دستگاه است:

#![no_std]

use cuda_kernel::cuda_kernel;
use cuda_std::{thread_idx, block_idx, block_dim};

#[panic_handler]
fn panic(_: &core::panic::PanicInfo) -> ! {
    
    
    cuda_std::trap()
}





#[cuda_kernel]
pub unsafe fn vector_add(a: *const f32, b: *const f32, out: *mut f32, len: u32) {
    let idx = (block_idx().x as u32) * (block_dim().x as u32) + (thread_idx().x as u32);

    if idx  len {
        let i = idx as usize;
        *out.add(i) = *a.add(i) + *b.add(i);
    }
}

را #[panic_handler] برای هر مورد لازم است no_std جعبه بدون آن، ساخت شکست می خورد #[panic_handler] function required. دستکاری کننده استفاده می کند cuda_std::trap() برای ایجاد یک استثنا در سمت دستگاه که compute-sanitizer می تواند به جای آویزان کردن پایه در یک حلقه بی نهایت گزارش دهد.

را #[cuda_kernel] ویژگی ماکرو این تابع را به یک نقطه ورودی PTX با قرارداد فراخوانی صحیح تبدیل می کند. امضای تابع از نشانگرهای خام استفاده می کند زیرا حافظه دستگاه از مدل مرجع امن Rust قابل آدرس دهی نیست. GPU در یک فضای حافظه جداگانه اجرا می شود که در آن ضمانت های طول عمر و اشغال سمت میزبان اعمال نمی شود.

بررسی مرز if idx is critical. Unlike CPU code where out-of-bounds access triggers a panic or segfault, GPU kernels that access memory beyond allocated bounds may silently read garbage or corrupt unrelated memory regions. This pattern, computing a global thread index and guarding with a length check, is the fundamental safety idiom for GPU kernels regardless of language.

درک ایندکس رشته و حافظه مشترک در Rust

محاسبه شاخص جهانی block_idx().x * block_dim().x + thread_idx().x هر رشته را به یک عنصر منحصر به فرد در آرایه های ورودی نگاشت می کند. block_dim().x تعداد رشته ها در هر بلوک (تعیین شده در هنگام راه اندازی) را برمی گرداند. block_idx().x مشخص می کند که نخ به کدام بلوک تعلق دارد و thread_idx().x موقعیت نخ در بلوک آن است.

برای هسته هایی که نیاز به استفاده مجدد از داده ها در یک بلوک از رشته ها دارند، حافظه مشترک ضروری است. مثال زیر یک الگوی دسترسی کاشی شده با همگام سازی را نشان می دهد:

use cuda_std::{shared_memory, syncthreads};

const TILE_SIZE: usize = 256;

#[cuda_kernel]
pub unsafe fn vector_add_shared(a: *const f32, b: *const f32, out: *mut f32, len: u32) {
    
    
    
    #[shared_memory]
    static mut TILE_A: [f32; TILE_SIZE] = [0.0; TILE_SIZE];
    #[shared_memory]
    static mut TILE_B: [f32; TILE_SIZE] = [0.0; TILE_SIZE];

    let tid = thread_idx().x as usize;
    let gid = (block_idx().x as u32) * (block_dim().x as u32) + (thread_idx().x as u32);

    
    
    if gid  len {
        TILE_A[tid] = *a.add(gid as usize);
        TILE_B[tid] = *b.add(gid as usize);
    } else {
        TILE_A[tid] = 0.0;
        TILE_B[tid] = 0.0;
    }

    
    
    syncthreads();

    if gid  len {
        *out.add(gid as usize) = TILE_A[tid] + TILE_B[tid];
    }
}

آرایه‌هایی از حافظه مشترک در SRAM روی تراشه زندگی می‌کنند، که در معماری‌های مدرن NVIDIA تقریباً 5-30 برابر تأخیر کمتری نسبت به حافظه جهانی دستگاه ارائه می‌کند (راهنمای برنامه‌نویسی CUDA C++، بخش 5.3.2 برای ارقام خاص معماری را ببینید). را syncthreads() تماس معادل Rust با CUDA C++ است __syncthreads()، اطمینان حاصل شود که همه رشته های موجود در بلوک قبل از خواندن هر رشته ای از آرایه های مشترک، نوشته های خود را کامل کرده اند. به این توجه کنید syncthreads() باید توسط تمام رشته های بلوک به طور یکنواخت به دست آید. قرار دادن آن در یک شاخه شرطی که در آن برخی از رشته ها واگرا می شوند، رفتار تعریف نشده در CUDA است. موضوعات خارج از محدوده بنویسید 0.0 برای اطمینان از اینکه اسلات های حافظه مشترک قبل از مانع مقداردهی اولیه شده اند.

برای یک بردار ساده، اضافه کردن این الگوی ترتیبی آموزشی است. برای عملیات ضرب ماتریس یا الگو، ذخیره حافظه مشترک مکانیسم اولیه برای دستیابی به حداکثر استفاده از پهنای باند حافظه است.

ساختن PTX از Rust: The Automated Build Script

host-runner/build.rs: باکس کرنل را متقاطع کامپایل کنید

اسکریپت ساخت در جعبه میزبان، ساخت دستگاه را خودکار می کند، بنابراین یک واحد cargo build از ریشه فضای کاری هم آرتیفکت PTX و هم باینری میزبان را تولید می کند:

use std::env;
use std::path::PathBuf;
use std::process::Command;

fn main() {
    println!("cargo:rerun-if-changed=../gpu-kernel/src/");
    println!("cargo:rerun-if-changed=../gpu-kernel/Cargo.toml");

    
    let cargo = env::var("CARGO").unwrap_or_else(|_| "cargo".into());

    
    let existing_rustflags = env::var("RUSTFLAGS").unwrap_or_default();
    let rustflags = format!("{} --emit=asm", existing_rustflags);

    let output = Command::new(&cargo)
        .args([
            "build",
            "--manifest-path", "../gpu-kernel/Cargo.toml",
            "--target", "nvptx64-nvidia-cuda",
            "--release",
        ])
        .env("RUSTFLAGS", rustflags.trim())
        
        .env(
            "CARGO_MAKEFLAGS",
            env::var("CARGO_MAKEFLAGS").unwrap_or_default(),
        )
        .output()
        .expect("Failed to spawn cargo for gpu-kernel");

    if !output.status.success() {
        panic!(
            "gpu-kernel compilation failed:
stdout: {}
stderr: {}",
            String::from_utf8_lossy(&output.stdout),
            String::from_utf8_lossy(&output.stderr)
        );
    }

    
    let manifest_dir = env::var("CARGO_MANIFEST_DIR").unwrap();
    let workspace_root = PathBuf::from(&manifest_dir)
        .parent()
        .expect("host-runner has no parent directory")
        .to_path_buf();

    let target_dir = env::var("CARGO_TARGET_DIR")
        .map(PathBuf::from)
        .unwrap_or_else(|_| workspace_root.join("target"));

    
    let candidates = [
        target_dir.join("nvptx64-nvidia-cuda/release/gpu_kernel.s"),
        target_dir.join("nvptx64-nvidia-cuda/release/gpu_kernel.ptx"),
    ];

    let ptx_path = candidates
        .iter()
        .find(|p| p.exists())
        .unwrap_or_else(|| {
            panic!(
                "PTX file not found. Searched:
{}
{}
\
                 Run: find {:?} -name 'gpu_kernel*' to locate the artifact.",
                candidates[0].display(),
                candidates[1].display(),
                target_dir.join("nvptx64-nvidia-cuda/release/"),
            )
        });

    let out_dir = env::var("OUT_DIR").unwrap();
    let dest = PathBuf::from(&out_dir).join("gpu_kernel.ptx");
    std::fs::copy(ptx_path, &dest).expect("Failed to copy PTX to OUT_DIR");

    println!("cargo:rustc-env=GPU_KERNEL_PTX={}", dest.display());
}

را cargo:rerun-if-changed دستورالعمل ها تضمین می کنند که جعبه دستگاه تنها زمانی دوباره کامپایل می شود که منبع یا پیکربندی آن تغییر کند. را --emit=asm پرچم دستور می دهد rustc برای پخش متن اسمبلی PTX (باطن NVPTX PTX را از طریق این مکانیسم تولید می کند، نه از طریق یک کتابخانه مشترک). فایل PTX کامپایل شده در فایل کپی می شود OUT_DIR بنابراین می توان آن را در زمان کامپایل توسط جعبه میزبان ارسال کرد. را cargo:rustc-env خط مسیر PTX را به عنوان یک متغیر محیطی در دسترس در زمان کامپایل میزبان نمایش می دهد.

توجه: استفاده env::var("CARGO") به جای هاردکد "cargo" تضمین می کند که اسکریپت ساخت همان Cargo باینری را که بیلد والد را اجرا می کند فراخوانی می کند. این در تنظیمات متقابل کامپایل، محیط های Nix و هنگام استفاده از رابط های جایگزین Cargo مهم است cargo-nextest.

را پیدا و جاسازی کنید .ptx فایل

کد میزبان PTX را در زمان کامپایل با استفاده از include_str!:

const PTX_SOURCE: &str = include_str!(env!("GPU_KERNEL_PTX"));

کامپایل زمان تعبیه PTX با include_str! یک باینری مستقل با PTX تعبیه شده ایجاد می کند که برای ساخت های قابل تکرار و سناریوهای استقرار که در آن طرح سیستم فایل تضمین نشده است ایده آل است. جایگزین، بارگیری PTX از مسیر فایل در زمان اجرا، انعطاف‌پذیری را برای توسعه تکراری ارائه می‌دهد که در آن کامپایل مجدد باینری میزبان برای هر تغییر هسته نامطلوب است. برای پیاده‌سازی‌هایی که پیوند هسته/دودویی قابل قبول است، PTX داخلی یک کلاس از خطاهای مرتبط با فایل‌های PTX گمشده یا ناهماهنگ را به قیمت نیاز به بازگشت کامل تغییر هسته حذف می‌کند.

راه اندازی هسته سمت میزبان و مدیریت حافظه

host-runner وابستگی صندوق پول

[package]
name = "host-runner"
version = "0.1.0"
edition = "2021"
build = "build.rs"

[dependencies]
cuda_driver = { workspace = true }

[features]
gpu-tests = []

را cuda_driver crate API های سمت میزبان را برای مدیریت زمینه، تخصیص حافظه، بارگذاری ماژول و ارسال هسته فراهم می کند. هیچ وابستگی ساخت خارج از کتابخانه استاندارد مورد نیاز نیست زیرا اسکریپت ساخت فقط از آن استفاده می کند std::process::Command.

زمینه CUDA را راه اندازی کنید و ماژول PTX را بارگذاری کنید

use cuda_driver::{CudaDevice, CudaContext, CuModule, CuFunction};

const PTX_SOURCE: &str = include_str!(env!("GPU_KERNEL_PTX"));

fn main() {
    
    cuda_driver::init().expect("Failed to initialize CUDA driver");

    
    let device = CudaDevice::get(0).expect("No CUDA device found");
    let ctx = CudaContext::new(&device).expect("Failed to create CUDA context");

    
    let module = CuModule::load_data(PTX_SOURCE).expect("Failed to load PTX module");
    let kernel: CuFunction = module
        .get_function("vector_add")
        .expect("Kernel function 'vector_add' not found in PTX module");

را CuModule::load_data call رشته PTX را به درایور CUDA ارسال می کند، که JIT آن را به مجموعه دستورات بومی GPU در معرض کامپایل می کند. را get_function تماس یک کنترل کننده را با نام بازیابی می کند. نام باید با نماد صادر شده مطابقت داشته باشد #[cuda_kernel] ماکرو در جعبه دستگاه

نکته: را #[cuda_kernel] ماکرو ممکن است نام تابع PTX صادر شده را خراب یا تزئین کند. اگر get_function("vector_add") در زمان اجرا خراب می شود، PTX را بررسی کنید تا نام نقطه ورودی واقعی را پیدا کنید:

grep "\.entry" target/nvptx64-nvidia-cuda/release/gpu_kernel.s

از نام موجود در آن استفاده کنید get_function یک تماس

عدم تطابق نام در اینجا یک خطای زمان اجرا ایجاد می کند، نه یک خطای زمان کامپایل، که یکی از لبه های ناهموار باقی مانده در زنجیره ابزار است.

تخصیص حافظه دستگاه و انتقال داده

    
    let n: usize = 1_048_576;
    assert!(
        n  u32::MAX as usize,
        "n={} exceeds u32::MAX; update kernel signature to u64 first", n
    );
    let n32: u32 = n as u32;

    let host_a: Vecf32> = (0..n).map(|i| i as f32).collect();
    let host_b: Vecf32> = (0..n).map(|i| (i as f32) * 2.0).collect();
    let mut host_out: Vecf32> = vec![0.0f32; n];

    
    let dev_a = cuda_driver::DeviceBuffer::from_slice(&host_a)
        .expect("Failed to allocate device buffer for a");
    let dev_b = cuda_driver::DeviceBuffer::from_slice(&host_b)
        .expect("Failed to allocate device buffer for b");
    let mut dev_out = cuda_driver::DeviceBuffer::f32>::zeroed(n)
        .expect("Failed to allocate device buffer for out");

DeviceBuffer::from_slice تخصیص میزبان به دستگاه و کپی را در یک تماس انجام می دهد. پارامترسازی نوع (DeviceBuffer::) تضمین می کند که محاسبات اندازه سطح بایت و محاسبات نشانگر توسط سیستم نوع به جای دستی انجام می شود. size_of::() * n محاسبات این حذف می کند void* قالب‌هایی که کد CUDA C را آزار می‌دهند و منبع ثابتی از خطاهای مخفی خرابی حافظه هستند.

شروع کرنل در یک موضوع CUDA

    
    let threads_per_block: u32 = 256;
    let blocks: u32 = (n32 + threads_per_block - 1) / threads_per_block;

    
    let stream = cuda_driver::CudaStream::new().expect("Failed to create CUDA stream");

    
    
    
    let ptr_a = dev_a.as_device_ptr();
    let ptr_b = dev_b.as_device_ptr();
    let ptr_out = dev_out.as_device_ptr();

    unsafe {
        cuda_driver::launch_kernel(
            &kernel,
            (blocks, 1, 1),       
            (threads_per_block, 1, 1), 
            0,                     
            &stream,
            &[
                &ptr_a   as *const _ as *const std::ffi::c_void,
                &ptr_b   as *const _ as *const std::ffi::c_void,
                &ptr_out as *const _ as *const std::ffi::c_void,
                &n32     as *const _ as *const std::ffi::c_void,
            ],
        )
        .expect("Kernel launch failed");
    }

    
    stream.synchronize().expect("Stream synchronization failed");
    dev_out.copy_to(&mut host_out).expect("Device-to-host copy failed");

    
    let epsilon: f32 = 1e-5;
    for i in 0..n {
        let expected = host_a[i] + host_b[i];
        assert!(
            (host_out[i] - expected).abs()  epsilon,
            "Mismatch at index {}: got {}, expected {}",
            i,
            host_out[i],
            expected
        );
    }
    println!("All {} elements verified correct.", n);
}

هشدار: یک اشتباه رایج گذر است &dev_a as *const _ as *const c_void، که یک اشاره گر را به کشور میزبان ارسال می کند DeviceBuffer struct – نه اشاره گر دستگاهی که هسته انتظار دارد. این باعث دسترسی غیرقانونی به حافظه یا خرابی داده های GPU بی صدا می شود. همیشه ابتدا نشانگر دستگاه خام را واکشی کنید (مثلاً از طریق as_device_ptr()) و یک اشاره گر را به که ارزش

را launch_kernel تماس است unsafe زیرا درایور CUDA نمی تواند در زمان کامپایل بررسی کند که انواع پارامترهای هسته با امضای تابع مطابقت دارند.

را launch_kernel تماس است unsafe زیرا درایور CUDA نمی تواند در زمان کامپایل بررسی کند که انواع پارامترهای هسته با امضای تابع مطابقت دارند. در هسته len یک پارامتر است u32، و میزبان عبور می کند n32: u32 – این مسابقه حیاتی است. میزبان عبوری usize (8 بایت در x86-64) برای پارامتر هسته اعلام شده به عنوان u32 (4 بایت) به دلیل عدم تطابق اندازه ABI، پارامترهای مجاور را در آرایه پارامتر خراب می کند. اندازه شبکه و بلوک از مدل استاندارد CUDA پیروی می کند: blocks * threads_per_block باید حداقل باشد n برای پوشاندن تمام عناصر، و بررسی کرنل داخلی در برابر رشته های دنباله ای که از طول آرایه بیشتر است محافظت می کند.

همگام سازی جریان از طریق stream.synchronize() تا زمانی که تمام عملیات روی آن رشته به پایان برسد، رشته میزبان را مسدود می کند. پرش از این تماس قبل از بازخوانی نتایج، یک منبع رایج خطاهای غیر قطعی است، زیرا میزبان حافظه دستگاه بدون مقدار اولیه یا نیمه نوشته شده را می‌خواند.

آزمایش و اعتبارسنجی خروجی هسته

نوشتن تست های ادغام سمت میزبان

صحت هسته GPU را می توان در استاندارد تایید کرد cargo test. از آنجایی که آزمایش‌ها به یک GPU با قابلیت CUDA نیاز دارند، آن‌ها را پشت یک پرچم ویژگی قرار دهید تا محیط‌های CI بدون دسترسی به GPU بتوانند از آن‌ها عبور کنند:

use cuda_driver::{CudaDevice, CudaContext, CuModule, CudaStream, DeviceBuffer};

const PTX_SOURCE: &str = include_str!(env!("GPU_KERNEL_PTX"));

#[cfg(feature = "gpu-tests")]
#[test]
fn test_vector_add_correctness() {
    cuda_driver::init().unwrap();
    let device = CudaDevice::get(0).unwrap();
    let _context = CudaContext::new(&device).unwrap();

    let module = CuModule::load_data(PTX_SOURCE).unwrap();
    let kernel = module.get_function("vector_add").unwrap();

    let n: usize = 4096;
    assert!(n  u32::MAX as usize);
    let n32: u32 = n as u32;

    let host_a: Vecf32> = (0..n).map(|i| i as f32).collect();
    let host_b: Vecf32> = (0..n).map(|i| (i as f32) * 0.5).collect();
    let mut host_out: Vecf32> = vec![0.0f32; n];

    let dev_a = DeviceBuffer::from_slice(&host_a).unwrap();
    let dev_b = DeviceBuffer::from_slice(&host_b).unwrap();
    let mut dev_out = DeviceBuffer::f32>::zeroed(n).unwrap();

    let stream = CudaStream::new().unwrap();
    let threads_per_block: u32 = 256;
    let blocks: u32 = (n32 + threads_per_block - 1) / threads_per_block;

    let ptr_a = dev_a.as_device_ptr();
    let ptr_b = dev_b.as_device_ptr();
    let ptr_out = dev_out.as_device_ptr();

    unsafe {
        cuda_driver::launch_kernel(
            &kernel,
            (blocks, 1, 1),
            (threads_per_block, 1, 1),
            0,
            &stream,
            &[
                &ptr_a   as *const _ as *const std::ffi::c_void,
                &ptr_b   as *const _ as *const std::ffi::c_void,
                &ptr_out as *const _ as *const std::ffi::c_void,
                &n32     as *const _ as *const std::ffi::c_void,
            ],
        )
        .unwrap();
    }

    stream.synchronize().unwrap();
    dev_out.copy_to(&mut host_out).unwrap();

    let epsilon: f32 = 1e-5;
    for i in 0..n {
        let expected = host_a[i] + host_b[i];
        assert!(
            (host_out[i] - expected).abs()  epsilon,
            "Mismatch at index {}: got {}, expected {}",
            i,
            host_out[i],
            expected
        );
    }
}

تست های GPU را به طور صریح با cargo test --features gpu-tests. محیط های CI بدون دسترسی به GPU می توانند اجرا شوند cargo test بدون پرچم برای رد شدن از این تست ها.

تحمل اپسیلون تفاوت در عملکرد ممیز شناور بین حساب میزبان و دستگاه را نشان می دهد.

مشکلات رایج و نکات عیب یابی

اگر blocks * threads_per_block کمتر از طول آرایه است، عناصر پایانی هرگز محاسبه نمی شوند. هسته خطا نمی دهد. این فقط هرگز برای آن شاخص ها رشته نمی شود. همیشه بلوک های دارای جداسازی سقف را محاسبه کنید: (n + threads_per_block - 1) / threads_per_block. این رایج ترین منبع خطاهای صحت بی صدا در کد GPU است.

درج تصادفی std انواع موجود در جعبه دستگاه خطاهای مرموز LLVM را در طول کامپایل PTX ایجاد می کند. هر گونه وابستگی اضافه شده به gpu-kernel باید باشد no_std سازگار وابستگی های گذرا را با دقت حسابرسی کنید.

شما همچنین ممکن است ضربه بزنید cuModuleLoadData زمانی که نسخه PTX با درایور مطابقت نداشته باشد خراب می شود. CUDA Toolkit PTX سازگار با درایورهای همان خانواده نسخه اصلی را تولید می کند، اما کار با درایورهای بسیار قدیمی تر شکست خواهد خورد. آن را بررسی کنید nvidia-smi خروجی در برابر یادداشت های انتشار CUDA Compatibility Toolkit.

اشکال زدایی رفتار سمت دستگاه با کمک NVIDIA امکان پذیر است compute-sanitizer ابزاری که می‌تواند دسترسی‌های خارج از محدوده، شرایط مسابقه و خواندن حافظه اولیه را در PTX تولید شده توسط Rust شناسایی کند، درست مانند کدهای تولید شده توسط NVCC:

compute-sanitizer --tool memcheck ./target/release/host-runner

به این توجه کنید compute-sanitizer همراه با جعبه ابزار CUDA و باید روی شما باشد PATH (معمولا در /usr/local/cuda/bin/compute-sanitizer).

ملاحظات عملکرد و مراحل بعدی

زمان استفاده از این زنجیره ابزار در مقابل جایگزین ها

را cuda-rust toolchain به طور خاص برای سخت افزار NVIDIA ساخته شده است و PTX را تولید می کند که مدل اجرای کامل CUDA را هدف قرار می دهد. برای بارهای کاری که چندین فروشنده GPU را هدف قرار می دهد، wgpu با شیدرهای محاسباتی WGSL یا nagaخطوط لوله مبتنی بر پلتفرم قابلیت حمل و نقل بین پلتفرمی را به قیمت از دست دادن دسترسی به اولیه‌های سطح پیچ (درهم‌آمیزی، بولتن‌ها)، سفارش حافظه مشترک صریح، و تنظیم مستقیم مشغولیت فراهم می‌کنند. را rust-gpu این پروژه، در حالی که پیشگام است، عمدتاً بر روی کامپایل سایه زن های گرافیکی از طریق SPIR-V تمرکز دارد و تا سال 2024 به یک نسخه پایدار برای محاسبات عمومی نرسیده است. اتصالات OpenCL از طریق ocl گزینه ای برای محیط های ناهمگن باقی می ماند، اما فاقد مزایای مشخص است cuda-rust فراهم می کند: پارامترهای اندازه بافر زمان کامپایل، بافرهای دستگاه تایپ شده که حذف می شوند void* ریخته گری و ادغام فضای کاری باری ممتاز.

را cuda-rust Toolchain برای بارهای کاری انحصاری NVIDIA مناسب است، جایی که ایمنی نوع Rust، مدل مالکیت و اکوسیستم Cargo قفل شدن پلت فرم را توجیه می کند.

گسترش فضای کاری

افزودن a shared-types جعبه به فضای کاری اجازه می دهد تا ساختارهای داده و ثابت هایی را که هم دستگاه و هم کد میزبان استفاده می کنند، تعریف کنند، مشروط بر اینکه این انواع no_std سازگار عملکرد هسته نمایه با NVIDIA nsys (ردیابی سطح سیستم) و ncu (تحلیل سطح هسته)؛ هر دو به طور یکسان با کد تولید شده توسط Rust PTX و کد تولید شده توسط NVCC کار می کنند. برای مقیاس بندی به خطوط لوله چند هسته ای و چند رشته ای، اضافه کنید #[cuda_kernel] ویژگی های موجود در جعبه دستگاه، تماس بگیرید get_function برای هر کدام در هاست و از وابستگی های رشته صریح برای دستور اجرا استفاده کنید.

تصویر کامل

فضای کاری ساخته شده از طریق این مقاله الف gpu-kernel جعبه کامپایل Rust به کد بایت NVPTX و a host-runner اجرای یک کارتریج و اعتبارسنجی این هسته‌ها از طریق اتصالات درایور CUDA بومی Rust، یک گردش کار کامل توسعه GPU بدون C++ و فراخوانی NVCC را نشان می‌دهد. اشاره گر خام پرتاب می کند launch_kernel گفتگوها باقی می ماند unsafe، زیرا درایور CUDA نمی تواند انواع پارامترها را در زمان کامپایل بررسی کند. این شکاف بین نوع سیستم میزبان و آرایه پارامترهای تایپ نشده راننده مهمترین مرز ایمنی است که هنوز حل نشده است.

این شکاف بین نوع سیستم میزبان و آرایه پارامترهای تایپ نشده راننده مهمترین مرز ایمنی است که هنوز حل نشده است.

در دسترس بودن کادر را علامت بزنید، هسته افزودن برداری را با محاسبات مخصوص دامنه جایگزین کنید و با پیاده سازی های معادل CUDA C++ مقایسه کنید تا زنجیره ابزار را برای بارهای کاری واقعی ارزیابی کنید.