layout: true class: typo, typo-selection --- count: false class: nord-dark, middle, center # 🎮 GPU-Accelerated Randomized Vertex Cover ### Porting Python Numba CUDA → Rust cudarc @luk036 👨‍💻 · 2026 📅 · Network Optimization --- ### 📋 Agenda .pull-left[ #### 1️⃣ Background - Pitt's randomized algorithm 🎲 - Why GPU acceleration? 🚀 - Monte Carlo trials 🎰 #### 2️⃣ Python GPU Code - Numba CUDA decorator 🐍 - Kernel structure 🔬 - Host-side orchestration ⚙️ ] .pull-right[ #### 3️⃣ Rust cudarc Port - cudarc architecture 🦀 - CUDA C kernel compilation 📜 - Three-tier fallback 🔄 #### 4️⃣ Build & Profile - Feature flags ⚙️ - ncu profiling 🔍 - Diagnostic logging 📊 #### 5️⃣ Lessons Learned - CUDA version features ⚠️ - ERR_NVGPUCTRPERM 🛡️ - v12 API compatibility 🔧 ] --- class: nord-light, middle, center ## 🎲 Pitt's Randomized Vertex Cover --- ### 🎯 The Algorithm — Quick Recap - For each uncovered edge $(u, v)$, pick an endpoint with probability **inversely proportional to weight**: $$\Pr(\text{pick } u) = \frac{w(v)}{w(u) + w(v)}, \qquad \Pr(\text{pick } v) = \frac{w(u)}{w(u) + w(v)}$$ - **Lighter** vertices are more likely to be selected ✅ - Achieves **2-approximation** in expectation - **Trivially parallel**: each trial completely independent 🎰 - Post-processing: reverse-delete removes redundant vertices - 🎰 Monte Carlo Approach Instead of running **1 trial**, run **N trials** and pick the best: $$\text{Best cover} = \arg\min_{i=1}^{N} \sum_{v \in \text{cover}_i} w(v)$$ More trials → **better expected approximation** 📈 --- ### 🔬 Why GPU? The Parallelism Map .mermaid[ <pre> graph LR subgraph Graph["Input Graph"] direction LR V["Vertices V"] E["Edges E"] W["Weights w"] end Graph --> CONVERT["Convert to flat arrays\nedges[n_edges×2] i32\nweights[n_vertices] f32\nseeds[n_trials] i32"] CONVERT --> DEVICE["Upload to GPU"] subgraph GPU["GPU Execution"] direction LR T0["Thread 0\nTrial 0"] T1["Thread 1\nTrial 1"] T2["..."] TN["Thread N-1\nTrial N-1"] end DEVICE --> T0 DEVICE --> T1 DEVICE --> T2 DEVICE --> TN T0 --> MERGE["argmin costs"] T1 --> MERGE T2 --> MERGE TN --> MERGE MERGE --> RESULT["Best Cover"] style Graph fill:#4caf50 style GPU fill:#9c27b0 style MERGE fill:#ff9800 style RESULT fill:#4caf50 </pre> ] --- ### 🔢 Cover as Bitmask Instead of storing sets of vertex names, use **bitmask words** per trial: ``` word 0 word 1 word 2 Trial 0: |10110011...| |00011010...| |10000001...| Trial 1: |01101001...| |11000011...| |00110100...| ... Trial N: |00111100...| |01010110...| |11100010...| ``` - 1 bit per vertex → **32 vertices per uint32 word** 💪 - `num_words = (n_vertices + 31) / 32` - `cover[trial * num_words + word] & bit_mask` to check if vertex is covered - Fits entire cover state in **L1 cache** for small graphs 📊 --- class: nord-light, middle, center ## 🐍 Python: Numba CUDA --- ### 🎯 Numba CUDA Decorator ```python import numpy as np from numba import cuda THREADS_PER_BLOCK = 64 @cuda.jit def _pitt_kernel(edges, num_edges, weights, num_vertices, cover_words, costs, seeds): tid = cuda.grid(1) if tid >= costs.shape[0]: return seed = seeds[tid] for i in range(num_edges): u, v = edges[i, 0], edges[i, 1] # ... Pitt's rule with per-thread LCG RNG # Compute cost from bitmask costs[tid] = cost ``` - `@cuda.jit` — auto-compiles Python → PTX 🪄 - `cuda.grid(1)` — 1D thread index - `cuda.to_device()` / `copy_to_host()` — data transfer --- ### ⚙️ Host-Side Orchestration ```python def rand_vertex_cover_gpu(ugraph, weight, coverset=None, num_trials=1024, seed=None): # 1. NetworkX graph → flat NumPy arrays edges_np = np.array(edge_list, dtype=np.int32) weights_np = np.array(weight_list, dtype=np.float32) seeds_np = np.array(seed_list, dtype=np.int32) # 2. Allocate device memory d_edges, d_weights, d_covers, d_costs, d_seeds = [ cuda.to_device(arr) for arr in arrays ] # 3. Launch kernel blocks = ceil(num_trials / THREADS_PER_BLOCK) _pitt_kernel[blocks, THREADS_PER_BLOCK]( d_edges, n_edges, d_weights, n_vertices, d_covers, d_costs, d_seeds ) cuda.synchronize() # 4. Copy results back, find best trial return best_cover, best_cost ``` --- ### ✅ Pros & ❌ Cons of Numba CUDA > 🎯 **Goal**: Same algorithm power, but in **Rust** — zero runtime overhead, compile-time safety, graceful fallback. .pull-left[ #### ✅ Pros - **Write Python, run on GPU** — zero boilerplate 🪄 - `@cuda.jit` is a simple decorator - NumPy arrays auto-mapped to CUDA pointers - Python RNG → LCG in kernel - Perfect for prototyping ] .pull-right[ #### ❌ Cons - **Numba/Numba-CUDA** heavy dependency - **Python GIL** constraints on host side - No compile-time type checking 🐛 - Requires CUDA Toolkit + compatible Python - Single-GPU only - `RuntimeError` if no GPU found ] --- class: nord-light, middle, center ## 🦀 Rust: cudarc Port --- ### 🏗️ The cudarc Crate (v0.12) .mermaid[ <pre> graph TB subgraph cudarc["cudarc v0.12.1"] DRV["driver\nCudaDevice, CudaSlice\nCudaFunction, LaunchConfig"] NVRTC["nvrtc\ncompile_ptx()\nCUDA C → PTX"] end subgraph Our["rand_cover_gpu.rs"] KERN["PITT_KERNEL_CUDA\n(const &str)"] EXEC["try_gpu_execution()\n-#[cfg(feature = cuda)]"] HOST["rand_vertex_cover_gpu()\npublic API"] end KERN --> NVRTC NVRTC --> EXEC DRV --> EXEC HOST --> EXEC style cudarc fill:#2196f3 style Our fill:#ff9800 style KERN fill:#9c27b0 </pre> ] --- ### 📜 CUDA C Kernel — Embedded as Rust `&str` ```c const PITT_KERNEL_CUDA: &str = r#" extern "C" __global__ void pitt_kernel( const int* edges, int num_edges, const float* weights, int num_vertices, unsigned int* covers, float* costs, unsigned long long* seeds, int num_trials ) { int tid = blockIdx.x * blockDim.x + threadIdx.x; if (tid >= num_trials) return; unsigned long long seed = seeds[tid]; unsigned int* cover = &covers[tid * ((num_vertices + 31) / 32)]; for (int i = 0; i < num_edges; i++) { int u = edges[i * 2], v = edges[i * 2 + 1]; // ... Pitt's rule with LCG RNG } // Compute cost from bitmask costs[tid] = cost; } "#; ``` --- ### 🔄 Compilation & Execution Flow .pull-left[ .mermaid[ <pre> sequenceDiagram participant R as Rust Host participant N as NVRTC participant D as CUDA Driver participant G as GPU R->>R: Convert Graph → flat arrays R->>D: CudaDevice::new(0) R->>N: compile_ptx(PITT_KERNEL_CUDA) N->>R: Ptx object R->>D: load_ptx(ptx, "mod", &["pitt_kernel"]) R->>D: get_func("mod", "pitt_kernel") D->>R: CudaFunction R->>D: htod_copy(edges, weights, covers, seeds) D->>G: Copy H→D R->>G: launch(func, config, args) loop Each Thread G->>G: Pitt's algorithm end R->>D: synchronize() R->>D: dtoh_sync_copy(covers, costs) D->>R: Copy D→H R->>R: Find best trial (argmin) </pre> ] ] --- ### 🛡️ Three-Tier Fallback Architecture .font-sm.mb-xs[ ```rust pub fn rand_vertex_cover_gpu<W>(...) -> (HashSet<String>, W) { // Tier 1: Try GPU (only compiled with --features cuda) #[cfg(feature = "cuda")] { match try_gpu_execution(...) { Ok(result) => { eprintln!("✅ GPU execution succeeded"); return result; } Err(e) => eprintln!("⚠️ CUDA failed: {}", e), } } // Tier 2/3: CPU fallback cpu_multi_trial(...) // rayon parallel OR sequential } ``` | Tier | Backend | Feature | When | |------|---------|---------|------| | 🥇 | CUDA GPU | `cuda` | Kernel compiles + launches | | 🥈 | CPU Rayon | `rayon` | `cuda` disabled or GPU fails | | 🥉 | CPU Sequential | (default) | Neither feature, or tiny graphs | ] --- ### ⚠️ Key Fix: CUDA Version Feature `cudarc` v0.12 requires an **explicit CUDA version feature**: ```toml # Cargo.toml [dependencies.cudarc] version = "0.12" optional = true default-features = false features = ["driver", "nvrtc", "cuda-12030"] # ← MUST match installed CUDA [features] cuda = ["dep:cudarc"] ``` **Without** the version feature → **build panics** 💥: ``` Must specify one of: [cuda-version-from-build-system, cuda-12060, cuda-12050, cuda-12040, cuda-12030, ...] ``` Available version features map to your installed CUDA (12.3 → `cuda-12030`). --- ### 🔧 Build Commands .pull-left[ ```powershell # CPU-only (default) cargo build cargo test # GPU enabled cargo build --features cuda cargo test --features cuda # GPU + CPU parallel cargo build --features cuda,rayon ``` ] .pull-right[ .mermaid[ <pre> graph LR subgraph Features D["default = []"] --> CPU["CPU sequential"] CUDA["cuda"] --> GPU["GPU via cudarc"] R["rayon"] --> CPU end GPU --> CUDA_COMB["cuda + rayon"] CPU --> CUDA_COMB style D fill:#666 style CUDA fill:#9c27b0 style R fill:#2196f3 style GPU fill:#4caf50 style CUDA_COMB fill:#ff9800 </pre> ] ] --- class: nord-light, middle, center ## 🔍 Profiling & Verification --- ### 🛠️ Profiling Tools | Tool | Status | Use | |------|--------|-----| | `nvprof` | ⛔ **Deprecated** (CUDA 11+) — dead stub in v12 | Do NOT use | | `ncu` | ✅ **Nsight Compute** (modern) | `ncu --kernel-name regex:pitt` | | `nsys` | ✅ **Nsight Systems** (timeline) | `nsys profile` | | `nvidia-smi` | ✅ **Real-time monitoring** | `nvidia-smi pmon -i 0` | --- ### 🔬 ncu Profiling Command ```powershell # 1. Build with CUDA + no-run to get binary path cargo test --features cuda --lib --no-run # 2. Profile specific kernel with Nsight Compute ncu --kernel-name regex:pitt --set full ` target\debug\deps\netlistx_rs-d98804f73a79131b.exe ` rand_cover_gpu::tests::test_gpu_triangle --nocapture ``` Expected output with **full GPU trace**: ``` ==PROF== Connected to process ... test rand_cover_gpu::tests::test_gpu_triangle ... ok [rand_cover_gpu] ✅ GPU execution succeeded (64 trials) ``` --- ### 🛡️ ERR_NVGPUCTRPERM — The Permission Bug On Windows, `ncu` fails with: ``` ==ERROR== ERR_NVGPUCTRPERM - The user does not have permission to access NVIDIA GPU Performance Counters on the target device 0. ``` **Fix** — NVIDIA Control Panel (one-time): .mermaid[ <pre> graph LR NV["NVIDIA Control Panel"] --> DEV["Developer\nsection"] DEV --> COUNTERS["Manage GPU\nPerformance Counters"] COUNTERS --> ALLOW["Allow access for\nall users"] ALLOW --> REBOOT["Reboot 🔄"] style NV fill:#9c27b0 style DEV fill:#2196f3 style COUNTERS fill:#ff9800 style ALLOW fill:#4caf50 style REBOOT fill:#f44336 </pre> ] --- ### 📊 Diagnostic Logging The code emits clear diagnostic messages to stderr: | Message | Meaning | Action | |---------|---------|--------| | `✅ GPU execution succeeded (N trials)` | CUDA kernel ran 🎮 | Working — profile with `ncu` | | `⚠️ CUDA failed: {reason}` | Error — check reason 🔧 | Fix CUDA config, retry | | (silent) | CPU mode (no `cuda` feature) 💻 | Add `--features cuda` | ```rust #[cfg(debug_assertions)] eprintln!("✅ GPU execution succeeded ({} trials)", num_trials); ``` > 🎯 Always checked: you can **prove** the GPU path ran just by watching stderr! --- class: nord-light, middle, center ## 🏗️ Architecture Deep-Dive --- ### 🧬 Data Layout: Host → GPU .mermaid[ <pre> graph TB subgraph Host["Host Memory"] HF["edges_flat: Vec<i32>\n[n_edges × 2]"] HFW["weights_flat: Vec<f32>\n[n_vertices]"] HFS["seeds: Vec<u64>\n[num_trials]"] HFC["covers_host: Vec<u32>\n[num_trials × num_words]"] HFC2["costs_host: Vec<f32>\n[num_trials]"] end subgraph Device["Device Memory"] DF["d_edges: CudaSlice<i32>"] DFW["d_weights: CudaSlice<f32>"] DFS["d_seeds: CudaSlice<u64>"] DFC["d_covers: CudaSlice<u32>"] DFC2["d_costs: CudaSlice<f32>"] end HF -->|"htod_copy()"| DF HFW -->|"htod_copy()"| DFW HFS -->|"htod_copy()"| DFS HFC -->|"htod_copy()"| DFC DFC2 -.->|"alloc_zeros()"| Device DFC -->|"dtoh_sync_copy()"| HFC DFC2 -->|"dtoh_sync_copy()"| HFC2 style Host fill:#4caf50 style Device fill:#9c27b0 </pre> ] --- ### 🔪 Kernel Cutting Details ```cuda // Each CUDA thread = 1 independent Pitt trial __global__ void pitt_kernel( const int* edges, int num_edges, // graph topology const float* weights, int num_vertices, // vertex weights unsigned int* covers, // output: bitmask per trial float* costs, // output: cost per trial unsigned long long* seeds, // per-trial RNG state int num_trials // guard: tid < num_trials ) ``` **Kernel Config**: - `THREADS_PER_BLOCK = 64` (arbitrary choice, tunable) - `LaunchConfig::for_num_elems(grid_blocks * THREADS_PER_BLOCK)` - **Note**: cudarc's `for_num_elems` uses **1024 threads/block** internally — extra threads exit via `tid >= num_trials` guard --- ### 🔧 cudarc v0.12 API — Key Types .font-sm.mb-xs[ | cudarc Type | CPU Equivalent | Purpose | |------------|---------------|---------| | `CudaDevice` | `std::alloc::Global` | GPU memory allocator | | `CudaSlice<T>` | `Vec<T>` | Owned GPU memory | | `CudaView<T>` | `&[T]` | Immutable GPU slice | | `CudaViewMut<T>` | `&mut [T]` | Mutable GPU slice | | `CudaFunction` | `fn(...)` | Compiled GPU kernel | | `LaunchConfig` | — | Grid/block dimensions | | `LaunchAsync` | — | Trait for kernel launch | ] ```rust let dev = CudaDevice::new(0).unwrap(); // GPU 0 let d_data = dev.htod_copy(vec![1,2,3]).unwrap(); // H → D let result = dev.sync_reclaim(d_data).unwrap(); // D → H + free ``` --- class: nord-light, middle, center ## 📦 Putting It All Together --- .pull-left[ ### 🗂️ Source Files Overview .font-sm.mb-xs[ | File | Purpose | |------|---------| | `src/rand_cover_gpu.rs` | GPU module: kernel, host, fallback | | `src/rand_cover.rs` | CPU random cover + hypergraph multi-trial | | `src/lib.rs` | Module declaration + re-export | | `Cargo.toml` | Feature flags: `cuda`, `rayon` | ] ] .pull-right[ ### 🎮 Module API .font-sm.mb-xs[ ```rust // Public entry point — auto-selects GPU/CPU pub fn rand_vertex_cover_gpu<W>( grph: &Graph<String, (), Undirected>, weight: &HashMap<String, W>, coverset: &HashSet<String>, num_trials: usize, // default: 1024 seed: u64, ) -> (HashSet<String>, W) // Convenience wrapper pub fn rand_vertex_cover_gpu_default<W>( grph: &Graph<String, (), Undirected>, weight: &HashMap<String, W>, ) -> (HashSet<String>, W) ``` ] ] --- ### 🧪 Test Suite (6 GPU-specific tests) ```rust #[test] fn test_gpu_triangle() // 3-node cycle — valid cover #[test] fn test_gpu_line() // 3-node path #[test] fn test_gpu_weighted_edge() // Heavy vs light — picks light #[test] fn test_gpu_empty_graph() // Edge case #[test] fn test_gpu_with_initial_coverset() // Pre-seeded coverset #[test] fn test_gpu_deterministic_seed() // Reproducibility ``` ```powershell cargo test --features cuda --lib rand_cover_gpu -- --nocapture ``` ``` running 6 tests test test_gpu_triangle ... ok test test_gpu_line ... ok test test_gpu_weighted_edge ... ok test test_gpu_empty_graph ... ok test test_gpu_with_initial_coverset ... ok test test_gpu_deterministic_seed ... ok [rand_cover_gpu] ✅ GPU execution succeeded (64 trials) ← GPU confirmed! ``` --- class: nord-light, middle, center ## 📊 Results & Insights --- ### 📈 What We Achieved .font-sm.mb-xs[ | Metric | Value | |--------|-------| | **Python CUDA port** | ✅ Full feature parity | | **CUDA kernel** | `pitt_kernel` — embedded as `&str` | | **Compilation** | Runtime NVRTC from CUDA C source | | **API surface** | `rand_vertex_cover_gpu()` matches Python | | **Fallback** | 3-tier: GPU → Rayon CPU → sequential | | **Tests** | 6 GPU-specific + 188 previously = **194 total** | | **Profiling** | `ncu` confirmed GPU execution | | **Diagnostics** | `eprintln!` shows which path is active | ] --- ### 🐛 Bugs Found & Fixed .font-sm.mb-xs[ | Bug | Root Cause | Fix | |-----|-----------|-----| | `cudarc` build panic | Missing `cuda-12030` version feature | Added to `Cargo.toml` | | `driver` module not found | Missing `"driver"` feature flag | Added to `features = [...]` | | `DriverError` formatting | No `Display` impl → `{}` fails | Use `{:?}` everywhere | | `DeviceRepr` not satisfied | `&i32` not supported for kernel args | Pass scalars by value | | GPU path never ran | `#[cfg(feature = "cuda")]` not enabled by default | `cargo test --features cuda` | ] --- ### 🔮 Future Improvements .pull-left[ #### 🚀 Performance - Tune `THREADS_PER_BLOCK` for occupancy - Shared memory for frequently-accessed weights - Streams for concurrent graph+GPU overlap - Multi-GPU support for larger trials ] .pull-right[ #### 🛠️ Usability - Auto-detect CUDA version at build time - Benchmark vs Python Numba CUDA - Pre-compiled PTX in build script (`build.rs`) - `#[cfg(feature = "cuda")]` for hypergraph too - CI integration with GPU runners ] --- count: false class: nord-dark, middle, center ### ✅ GPU Port: Complete & Verified .mermaid[ <pre> graph LR P["🐍 Python\nNumba CUDA\nrand_cover_gpu.py"] R["🦀 Rust\ncudarc\nrand_cover_gpu.rs"] CUDA["🎮 CUDA C\npitt_kernel\n(embedded &str)"] NVRTC["📜 NVRTC\ncompile_ptx()"] PROFILE["🔍 ncu\nNsight Compute"] TEST["✅ 6 GPU tests\npassing"] P -->|"port"| R R --> CUDA CUDA --> NVRTC NVRTC --> PROFILE PROFILE --> TEST style P fill:#4caf50,color:#fff style R fill:#2196f3,color:#fff style CUDA fill:#9c27b0,color:#fff style NVRTC fill:#ff9800,color:#fff style PROFILE fill:#00bcd4,color:#fff style TEST fill:#4caf50,color:#fff </pre> ] ### 🎯 Key Takeaway **Rust + CUDA = zero-cost GPU acceleration with compile-time safety.** `cudarc` + NVRTC compiles CUDA C from a Rust `&str` at runtime — no external build system, no `.cu` files, no `nvcc`. @luk036 👨‍💻 · 2026 📅