Native GPU Kernels in Cargo Workspaces

Writing Native GPU Kernels in Rust with the CUDA-Rust Toolchain

SitePoint TeamPublished inComputing·Programming·
October 1, 2026

The AI briefing for Developers

Stay up to date with AI tools, model releases, and developer workflows that matter.

Weekly. Free. One click to leave.

SitePoint Premium
Stay Relevant and Grow Your Career in Tech

• Premium Results
• Publish articles on SitePoint
• Daily curated jobs
• Learning Paths
• Discounts to dev tools

7 Day Free Trial. Cancel Anytime.

NVIDIA’s cuda-rust toolchain, paired with the nvptx64-nvidia-cuda target, removes the intermediate layers that have long plagued Rust GPU development. This article walks through building a complete two-crate Cargo workspace: a gpu-kernel crate that compiles a Rust kernel to NVPTX bytecode and a host-runner crate that launches that kernel on a CUDA device stream and validates results.

Status: The cuda-rust crate ecosystem (cuda_std, cuda_kernel, cuda_driver) described in this article has not yet been publicly released on crates.io. Code examples are illustrative of the intended workflow. Before attempting to reproduce these steps, verify crate availability on crates.io or NVIDIA’s official GitHub repositories. Building the device crate requires a Rust nightly toolchain with the rust-src component; the nvptx64-nvidia-cuda target is tier-3 and is not available via rustup target add on stable Rust.

Why Native GPU Kernels in Rust Matter Now

For years, Rust developers pursuing GPU-accelerated workloads faced an uncomfortable reality. Writing native GPU kernels in Rust meant cobbling together fragile toolchains: C++/CUDA glue code compiled through NVCC, cc crate build hacks that shelled out to external compilers, unsafe FFI wrappers around hand-compiled PTX, or experimental projects like rust-gpu that never fully stabilized for general-purpose compute. Each approach introduced friction, safety gaps, or both.

That situation is beginning to change. NVIDIA’s cuda-rust toolchain, paired with the nvptx64-nvidia-cuda target, removes these intermediate layers in its current preview. Rust developers compile standard no_std Rust directly to NVPTX bytecode using rustc‘s native LLVM backend, load that bytecode through Rust-native CUDA driver bindings, and execute kernels on GPU hardware without writing or compiling a single line of C++.

This article walks through building a complete two-crate Cargo workspace: a gpu-kernel crate that compiles a Rust kernel to NVPTX bytecode and a host-runner crate that launches that kernel on a CUDA device stream and validates results. Every step uses native Cargo tooling and Rust-native APIs.

• Rust nightly toolchain (e.g., nightly-2025-06-01 or later), with the rust-src component installed
• CUDA Toolkit 12.x (verify the latest release at the CUDA Toolkit Archive)
• An NVIDIA GPU with compute capability 7.0 or higher
• Linux (or WSL2)
• A compatible NVIDIA GPU driver (check the CUDA Toolkit release notes for minimum driver versions)

Architecture of the CUDA-Rust Toolchain

How nvptx64-nvidia-cuda Differs from Traditional NVCC Pipelines

The compilation path differs fundamentally from legacy CUDA development in Rust. With the nvptx64-nvidia-cuda target, the flow works like this: Rusth emits PTX assembly text directly. The CUDA driver then loads that PTX at runtime and JIT-compiles it to the specific GPU architecture

Contrast this with the traditional path: Rust host code uses bindgen or the cc crate to shell out to NVCC, which compiles .cu files into .cubin or .ptx artifacts. Rust then accesses those artifacts through unsafe FFI calls wrapping the CUDA driver API. Every boundary crossing, from Rust to the C++ build system to NVCC to the FFI layer, introduced potential for mismatched types, undefined behavior, and build system fragility.

Device code must be no_std because GPU execution contexts have no operating system, no standard allocator, and no access to host-side system calls. The no_std constraint is not a limitation of the toolchain but a reflection of the hardware execution model.

Key Crates and Components

NVIDIA’s cuda-rust crate ecosystem provides three primary components. cuda_std supplies the device-side standard library: thread indexing primitives, shared memory declarations, synchronization barriers, and math intrinsics. cuda_kernel provides the #[cuda_kernel] attribute macro that marks functions for device-side compilation and handles the ABI translation required for PTX entry points. cuda_driver exposes Rust-native bindings to the CUDA driver API for module loading, memory management, kernel launching, and stream synchronization on the host side.

Note: These crates may not yet be available on crates.io. Check NVIDIA’s official repositories foron if needed:

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

The nvptx64-nvidia-cuda target triple instructs rustc to emit NVPTX IR through LLVM rather than native machine code. The CUDA Toolkit is required on the host system for driver compatibility and runtime library access. The Rust compilation pipeline does not invoke NVCC, but the CUDA Toolkit provides the driver library (libcuda.so) needed for linking.

Why a Multi-Crate Cargo Workspace

Device code and host code cannot coexist in a single crate. A no_std crate targeting nvptx64-nvidia-cuda is fundamentally incompatible with a std crate targeting the host’s native triple. Cargo workspaces solve this cleanly: the device crate compiles to PTX under one target, the host crate compiles to a native binary under another, and both share a workspace root. A thin shared crate can hold common types or constants if needed.

Rust developers compile standard no_std Rust directly to NVPTX bytecode using rustc‘s native LLVM backend, load that bytecode through Rust-native CUDA driver bindings, and execute kernels on GPU hardware without writing or compiling a single line of C++.

Setting Up the Cargo Workspace

Workspace-Level Cargo.toml

The workspace root defines both crates as members and uses resolver version 2 for proper feature unification across targets:

[workspace]members=["gpu-kernel","host-runner"]resolver="2"[workspace.dependencies]cuda_std="0.3"cuda_kernel="0.3"cuda_driver="0.3"

Important: If these crates are not published on crates.io, replace the version strings with gitic revision or tag

This structure ensures that dependency versions stay synchronized across both crates while allowing each to specify its own target-specific configuration.

Installing the CUDA-Rust Toolchain and NVPTX Target

Before writing any Rust, the build environment must be configured:

rustup toolchain install nightlyrustup component add rust-src --toolchain nightlyrustup run nightly rustc --versionnvidia-smils /usr/local/cuda/lib64/libcuda*

The nvptx64-nvidia-cuda target is a tier-3 target. It is not availablebrary (core) must be built fromcargo/config.toml in the device crate (see next section)

Configuring .cargo/config.toml

The gpu-kernel crate requires a .cargo/config.toml to specify the tier-3 target and enable build-std:

[build]target="nvptx64-nvidia-cuda"[unstable]build-std=["core"]

Without this file, cargo build for the device crate will fail with errors about missing core artifacts.

Note: The panic = "abort" profile setting belongs in Cargo.toml, not in .cargo/config.toml. The .cargo/config.toml file does not support [profile.*] tables; such sections will be silently ignored. See the gpu-kernel/Cargo.toml configuration below for the correct placement.

Pinning the Nightly Toolchain

Add a rust-toolchain.toml at the workspace root to ensure reproducible builds:

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

Project Directory Structure

The complete workspace layout:

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

Every file in this tree is created over the following sections. The build.rs in host-runner is the glue that makes cargo build from the workspace root compile both device and host code in a single invocation.

Writing the GPU Kernel Crate

gpu-kernel/Cargo.toml Configuration

The kernel crate requires specific configuration to produce PTX output rather than a native binary:

[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=3lto=truecodegen-units=1panic="abort"

PTX text is extracted by passing RUSTFLAGS="--emit=asm" to rustc during the device build; no shared library is produced. The rlib crate type is appropriate for a no_std library crate. The panic = "abort" setting is required because the NVPTX target does not support stack unwinding. The release profile settings affect the quality of the emitted PTX: LTO and single codegen-unit can reduce PTX size and improve register allocation predictability, which directly affects occupancy on the GPU.

Writing a no_std Vector-Add Kernel

The kernel itself is standard Rust with device-specific attributes and intrinsics:

#![no_std]usecuda_kernel::cuda_kernel;usecuda_std::{thread_idx, block_idx, block_dim};#[panic_handler]fnpanic(_:&core::panic::PanicInfo)->!{cuda_std::trap()}#[cuda_kernel]pubunsafefnvector_add(a:*constf32, b:*constf32, out:*mutf32, len:u32){let idx =(block_idx().x asu32)*(block_dim().x asu32)+(thread_idx().x asu32);if idx < len {let i = idx asusize;*out.add(i)=*a.add(i)+*b.add(i);}}

The #[panic_handler] is required for any no_std crate; without it, compilation fails with #[panic_handler] function required. The handler uses cuda_std::trap() to raise a device-side exception that compute-sanitizer can report, rather than hanging the warp in an infinite loop.

The #[cuda_kernel] attribute macro transforms this function into a PTX entry point with the correct calling convention. The function signature uses raw pointers because device memory is not addressable from Rust’s safe reference model; the GPU operates on a separate memory space where host-side lifetime and borrow guarantees do not apply.

The bounds check if idx < len 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.

Understanding Thread Indexing and Shared Memory in Rust

The global index calculation block_idx().x * block_dim().x + thread_idx().x maps each thread to a unique element in the input arrays. block_dim().x returns the number of threads per block (set at launch time), block_idx().x identifies which block the thread belongs to, and thread_idx().x is the thread’s position within its block.

For kernels that require data reuse within a thread block, shared memory is essential. The following variant demonstrates a tiled access pattern with synchronization:

usecuda_std::{shared_memory, syncthreads};constTILE_SIZE:usize=256;#[cuda_kernel]pubunsafefnvector_add_shared(a:*constf32, b:*constf32, out:*mutf32, len:u32){#[shared_memory]staticmutTILE_A:[f32;TILE_SIZE]=[0.0;TILE_SIZE];#[shared_memory]staticmutTILE_B:[f32;TILE_SIZE]=[0.0;TILE_SIZE];let tid =thread_idx().x asusize;let gid =(block_idx().x asu32)*(block_dim().x asu32)+(thread_idx().x asu32);if gid < len {TILE_A[tid]=*a.add(gid asusize);TILE_B[tid]=*b.add(gid asusize);}else{TILE_A[tid]=0.0;TILE_B[tid]=0.0;}syncthreads();if gid < len {*out.add(gid asusize)=TILE_A[tid]+TILE_B[tid];}}

The shared memory arrays live in on-chip SRAM, which on modern NVIDIA architectures offers roughly 5-30x lower latency than global device memory (see the CUDA C++ Programming Guide, Section 5.3.2 for architecture-specific figures). The syncthreads() call is the Rust equivalent of CUDA C++’s __syncthreads(), ensuring all threads in the block have completed their writes before any thread reads from the shared arrays. Note that syncthreads() must be reached by all threads in the block uniformly; placing it inside a conditional branch where some threads diverge is undefined behavior in CUDA. The out-of-bounds threads write 0.0 to ensure the shared memory slots are initialized before the barrier.

For a simple vector add this tiling pattern is pedagogical; for matrix multiplication or stencil operations, shared memory tiling is the primary mechanism for achieving near-peak memory bandwidth utilization.

Building PTX from Rust: The Automated Build Script

host-runner/build.rs: Cross-Compiling the Kernel Crate

The build script in the host crate automates the device compilation so that a single cargo build from the workspace root produces both the PTX artifact and the host binary:

usestd::env;usestd::path::PathBuf;usestd::process::Command;fnmain(){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());}

The cargo:rerun-if-changed directives ensure that the device crate is only recompiled when itsc to emit PTX assembly text (the NVPTX backend produces PTX through this mechanism rather than a shared library). The compiled PTX file is copied into OUT_DIR so it can be referenced at compile time by the host crate. The cargo:rustc-env line exposes the PTX path as an environment variable available during host compilation

Note: Using env::var("CARGO") instead of hardcoding "cargo" ensures the build script invokes the same Cargo binary that is running the parent build. This is important in cross-compilation setups, Nix environments, and when using alternative Cargo frontends like cargo-nextest.

Locating and Embedding the .ptx File

The host code loads the PTX at compile time using include_str!:

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

Embedding the PTX at compile time with include_str! produces a self-contained binary with the PTX baked in, which is ideal for reproducible builds and deployment scenarios where file system layout is not guaranteed. The alternative, loading PTX from a file path at runtime, offers flexibility for iterative development where recompiling the host binary for every kernel change is undesirable. For deployments where kernel/binary version coupling is acceptable, embedded PTX eliminates a class of errors related to missing or mismatched PTX files at the cost of requiring a full rebuild per kernel change.

Host-Side Kernel Launch and Memory Management

host-runner Crate Dependencies

[package]name="host-runner"version="0.1.0"edition="2021"build="build.rs"[dependencies]cuda_driver={workspace=true}[features]gpu-tests=[]

The cuda_driver crate provides the host-side API for context management, memory allocation, module loading, and kernel dispatch. No build dependencies beyond the standard library are required since the build script uses only std::process::Command.

Initializing the CUDA Context and Loading the PTX Module

usecuda_driver::{CudaDevice,CudaContext,CuModule,CuFunction};constPTX_SOURCE:&str=include_str!(env!("GPU_KERNEL_PTX"));fnmain(){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");

The CuModule::load_data call passes the PTX string to the CUDA driver, which JIT-compiles it to the native instruction set of the detected GPU. The get_function call retrieves a handle by name; the name must match the symbol exported by the #[cuda_kernel] macro in the device crate.

Tip: The #[cuda_kernel] macro may mangle or decorate the exported PTX function name. If get_function("vector_add") fails at runtime, inspect the PTX to find the actual entry point name:

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

Use the name found there in the get_function call.

A name mismatch here produces a runtime error, not a compile-time error, which is one of the remaining rough edges in the toolchain.

Allocating Device Memory and Transferring Data

let n:usize=1_048_576;assert!(n <=u32::MAXasusize,"n={} exceeds u32::MAX; update kernel signature to u64 first", n);let n32:u32= n asu32;let host_a:Vec<f32>=(0..n).map(|i| i asf32).collect();let host_b:Vec<f32>=(0..n).map(|i|(i asf32)*2.0).collect();letmut host_out:Vec<f32>=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");letmut dev_out =cuda_driver::DeviceBuffer::<f32>::zeroed(n).expect("Failed to allocate device buffer for out");

DeviceBuffer::from_slice performs allocation and host-to-device copy in a single call. The type parameterization (DeviceBuffer::<f32>) ensures that the byte-level size calculations and pointer arithmetic are handled by the type system rather than by manual size_of::<T>() * n computations. This eliminates the void* casts that plague CUDA C code and are a persistent

Launching the Kernel on a CUDA Stream

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*conststd::ffi::c_void,&ptr_b   as*const _ as*conststd::ffi::c_void,&ptr_out as*const _ as*conststd::ffi::c_void,&n32     as*const _ as*conststd::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 in0..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);}

Warning: A common mistake is passing &dev_a as *const _ as *const c_void, which passes a pointer to the host-side DeviceBuffer struct — not the device pointer the kernel expects. This causes illegal memory accesses or silent data corruption on the GPU. Always extract the raw device pointer first (e.g.,) and pass a pointer to that value

The launch_kernel call is unsafe because the CUDA driver cannot verify at compile time that the kernel parameter types match the function signature.

The launch_kernel call is unsafe because the CUDA driver cannot verify at compile time that the kernel parameter types match the function signature. The kernel’s len parameter is u32, and the host passes n32: u32 — this match is critical. Passing a host usize (8 bytes on x86-64) for a kernel parameter declared as u32 (4 bytes) would corrupt adjacent parameters in the parameter array due to ABI size mismatch. Grid and block dimensions follow the standard CUDA model: blocks * threads_per_block must be at least n to cover all elements, with the kernel’s internal bounds check guarding against the tail threads that exceed the array length.

Stream synchronization blocks the host thread until all operations on that stream have completed. Omitting this call before reading results back is a commoned or partially written device memory

Testing and Validating Kernel Output

Writing Host-Side Integration Tests

GPU kernel correctness can be validated within standard cargo test. Since tests require a CUDA-capable GPU, gate them behind a feature flag so CI environments without GPU access can skip them:

usecuda_driver::{CudaDevice,CudaContext,CuModule,CudaStream,DeviceBuffer};constPTX_SOURCE:&str=include_str!(env!("GPU_KERNEL_PTX"));#[cfg(feature = "gpu-tests")]#[test]fntest_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::MAXasusize);let n32:u32= n asu32;let host_a:Vec<f32>=(0..n).map(|i| i asf32).collect();let host_b:Vec<f32>=(0..n).map(|i|(i asf32)*0.5).collect();letmut host_out:Vec<f32>=vec![0.0f32; n];let dev_a =DeviceBuffer::from_slice(&host_a).unwrap();let dev_b =DeviceBuffer::from_slice(&host_b).unwrap();letmut 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*conststd::ffi::c_void,&ptr_b   as*const _ as*conststd::ffi::c_void,&ptr_out as*const _ as*conststd::ffi::c_void,&n32     as*const _ as*conststd::ffi::c_void,],).unwrap();}stream.synchronize().unwrap();dev_out.copy_to(&mut host_out).unwrap();let epsilon:f32=1e-5;for i in0..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);}}

Run GPU tests explicitly with cargo test --features gpu-tests. CI environments without GPU access can run cargo test without the flag to skip these tests.

The epsilon tolerance accounts for floating-point representation differences between host and device arithmetic.

Common Pitfalls and Debugging Tips

If blocks * threads_per_block is less than the array length, trailing elements are never computed. The kernel does not error; it simply never executes threads for those indices. Always compute blocks with ceiling division: (n + threads_per_block – 1) / threads_per_block. This is the most frequent

Accidentally pulling in std types in the device crate produces cryptic LLVM errors during PTX compilation. Any dependency added to gpu-kernel must itself be no_std compatible. Audit transitive dependencies carefully.

You may also hit cuModuleLoadData failures when the PTX version does not match the driver. The CUDA Toolkit generates PTX compatible with drivers from the same major version family, but running against significantly older drivers will fail. Check nvidia-smi output against the CUDA Toolkit release notes for compatibility.

Debugging device-side behavior is possible using NVIDIA’s compute-sanitizer tool, which can detect out-of-bounds accesses, race conditions, and uninitialized memory reads in Rust-generated PTX just as it does with NVCC-generated code:

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

Note that compute-sanitizer ships with the CUDA Toolkit and must be on your PATH (typically at /usr/local/cuda/bin/compute-sanitizer).

Performance Considerations and Next Steps

When to Use This Toolchain vs. Alternatives

The cuda-rust toolchain is purpose-built for NVIDIA hardware and produces PTX that targets the full CUDA execution model. For workloads targeting multiple GPU vendors, wgpu with WGSL compute shaders or naga-based pipelines provide cross-platform portability at the cost of losing access to warp-level primitives (shuffles, ballots), explicit shared memory tiling, and direct occupancy tuning. The rust-gpu project, while pioneering, has focused primarily on graphics shader compilation through SPIR-V and, as of 2024, has not reached a stable release for general-purpose compute. OpenCL bindings via ocl remain an option for heterogeneous environments but lack the concrete benefits cuda-rust provides: compile-time buffer size parameterization, typed device buffers that eliminate void* casts, and first-class Cargo workspace integration.

The cuda-rust toolchain is the strongest fit for NVIDIA-exclusive workloads where Rust’s type safety, ownership model, and Cargo ecosystem justify the platform lock-in.

Extending the Workspace

Adding a shared-types crate to the workspace enables defining data structures and constants that both device and host code use, provided those types are no_std compatible. Profile kernel performance with NVIDIA’s nsys (system-level tracing) and ncu (kernel-level analysis); both work identically with Rust-generated PTX and NVCC-generated code. To scale to multi-kernel, multi-stream pipelines, add #[cuda_kernel] functions in the device crate, call get_function for each on the host, and use explicit stream dependencies for execution ordering.

The Full Picture

The workspace built through this article, a gpu-kernel crate compiling Rust to NVPTX bytecode and a host-runner crate launching and validating those kernels through Rust-native CUDA driver bindings, demonstrates a complete GPU development workflow with no C++ and no NVCC invocation. The raw pointer casts in launch_kernel calls remain unsafe, as the CUDA driver cannot verify parameter types at compile time. That gap between the host’s type system and the driver’s untyped parameter array is the most significant safety boundary still unsolved.

That gap between the host’s type system and the driver’s untyped parameter array is the most significant safety boundary still unsolved.

Verify crate availability, replace the vector-add kernel with a domain-specific computation, and benchmark against equivalent CUDA C++ implementations to evaluate the toolchain for real workloads.

Sharing our passion for building incredible internet things.