OxiCUDA 0.4.0 Release

[0.4.0] - 2026-07-01

This release is an on-device validation pass: for the first time, hand-written PTX kernels across more than 60 crates were JIT-compiled and executed on real NVIDIA GPU hardware (an RTX A4000, sm_86, CUDA 12.4) rather than only checked for CPU-logic parity. This surfaced and fixed dozens of genuine bugs that no amount of CPU-side testing could have caught — kernels that never compiled (ptxas rejected them outright), kernels that compiled but computed the wrong thing (register shadowing, base-2/base-e mixups, races), and kernels that were still bare stubs behind a real, tested CPU reference implementation. Every fix was verified fail→revert→pass on the actual device. Alongside the validation sweep: several algorithms went from partial/proxy PTX implementations to full ones (P1 FEM assembly, NODE tree inference, soft-MoE dispatch, 3D Gaussian splatting projection/SH), a few new modules landed (variable-depth NODE trees, TabR-style retrieval, preconditioned CG, GPT-NeoX RoPE), and analytic test coverage was added for dozens of previously-untested modules.

Added

• On-device GPU validation harness: a feature-gated gpu-tests Cargo feature plus a src/gpu_tests.rs module per crate, JIT-compiling each crate's hand-written PTX via Module::from_ptx, launching it on a live CUDA device, and asserting numerical equivalence to a CPU oracle. Rolled out workspace-wide (more than 60 crates); every test skips gracefully when no device is present. See Fixed below for what it caught.
• Crates that ran clean on the very first on-device pass (JIT-loaded and matched their CPU oracle with zero bugs found): oxicuda-pinn (7 kernels), oxicuda-bayes (7), oxicuda-federated (7), oxicuda-continual (7), oxicuda-peft (7), oxicuda-meta (7), oxicuda-tn (7), oxicuda-sketch (7), oxicuda-graphalg (7), oxicuda-cvx (7), oxicuda-gen (6), oxicuda-adversarial (7), oxicuda-hdc (7) — 42+ kernels across 13 crates with no defects on first real-hardware execution. oxicuda-pde (7) and oxicuda-numeric (7) likewise matched cleanly but each surfaced one honestly-documented (not fixed) caveat: pde's fem_assemble_kernel (pre-completion, see below) was confirmed to do only a signed-triangle-area scatter; numeric's bessel_recurrence aliases the next point's J_0 when multiple points share a launch (documented, validation scoped to the single-point calling convention).
• oxicuda-pde: fem_assemble_kernel completed from a partial stub (per-element signed-area scatter into one matrix entry) to a full unconstrained dense P1 stiffness assembly — the 3×3 local K_ij = (1/(4·Area))·(b_i·b_j + c_i·c_j) per element, atomically scattered into the dense global matrix. Validated element-wise against the crate's own p1_local_stiffness to 1e-4 rel / 1e-5 abs.
• oxicuda-tabular: sparsemax_kernel replaced a dead threshold pass with the exact O(D²) Martins & Astudillo largest-support search; quantile_norm_kernel replaced a 2-bucket heuristic with true empirical-CDF linear-scan + interpolation; node_tree_eval_kernel (previously hardcoded to 2 leaves, ignoring depth) now runs the full multi-level NODE tree with per-level entmax-1.5 bisection and a 2^depth-leaf mixture. Validated against sparsemax/QuantileTransformer::transform/NodeTree::forward at 1e-4–1e-5. Also new: VarObliviousLayer (variable-depth NODE oblivious trees via outer-product leaf gating) and TabRecordLayer (TabR-style retrieval: encode → scaled −L2 similarity → entmax attention → convex combination), reusing the crate's entmax/entmoid/sparsemax simplex code (19 tests).
• oxicuda-moe: soft_moe_dispatch_kernel completed from a first-slot-only proxy to the real 3-pass slot-softmax dispatch matrix D[t,s] = softmax(x·Φ/√d) over all slots. Validated against SoftMoeRouter::dispatch_weights to 5e-4, every output row confirmed to sum to 1.
• oxicuda-geometry3d: project_kernel now emits the full EWA 2D covariance Σ_2d = J·R·Σ_3d·Rᵀ·Jᵀ + 0.3·I (previously never written); sh_eval_kernel now evaluates all 9 L=0..2 spherical-harmonic terms per RGB channel (previously a reduced 5-term basis). Validated against project_gaussian/Gaussian3d::sh_color on the A4000.
• oxicuda-recsys: implemented 4 previously-empty-loop stub kernels — embedding_lookup, dot_score, bpr_gradient, lightgcn_propagate — now real, validated bit-exact / to ~1e-4 against Bpr/LightGcn CPU references. The remaining PTX surface (softmax_topk, negsample_uniform) is documented as still-stub with loud STUB/PTX-BUG doc comments designed to fail the day each is implemented for real.
• oxicuda-webgpu: naga_tests.rs — real WGSL parse+validate (naga::front::wgsl::parse_str + valid::Validator) across all 15 shader generators (31 tests), replacing prior substring-only shader checks.
• oxicuda-pinn: hand-written pure-Rust FFT (iterative radix-2 Cooley–Tukey + Bluestein chirp-z for arbitrary/prime N) replacing the FNO spectral path's O(N²) brute-force DFT; wired into 1D + separable-2D spectral_conv, zero rustfft/oxicuda-fft dependency (fno_3d's DFT is a noted follow-up).
• oxicuda-nas: LatencyLut::to_bytes/from_bytes — a dependency-free little-endian persistence format (magic LLUT + version + stable OpKind discriminants), 7 tests including round-trip identity across all 8 OpKind variants.
• oxicuda-cvx: preconditioned conjugate gradient — pcg_solve/pcg_solve_counted/cg_solve_counted plus a Preconditioner trait with IdentityPrecond/JacobiPrecond (Jacobi cuts a κ=1e4 diagonal system from 6 CG iterations to 1).
• oxicuda-blas, oxicuda-ptx: new PTX kernel templates for broadcast bias-add and a numerically-stable causal (masked) softmax, F32/F64.
• oxicuda-dnn: GPT-NeoX half-split partial-rotary RoPE (NeoXRopeConfig, apply_rope_neox_half_split) alongside the existing GPT-J/RoFormer interleaved Rope, plus a RoPE-NeoX attention integration.
• oxicuda-ptx, oxicuda-metal, oxicuda-memory: new f64 math-intrinsic codegen module (body_builder/math_f64.rs), Metal backend function/type additions, and device_buffer helpers.
• Large-scale analytic test-coverage expansion for previously zero-coverage modules (property-based/closed-form assertions, not smoke tests): oxicuda-rlhf (kl_control, alignment metrics, PPO GAE rollout, SFT/reward/preference losses — 54 tests), oxicuda-recsys (popularity_neg + a 116-test suite across 16 models: BERT4Rec, SASRec, GRU4Rec, PLE, MMoE, ESMM, DeepFM, AutoInt, Wide&Deep, ALS, NMF, NGCF, LightGCN, NCF, Two-Tower, hard-neg sampling), oxicuda-peft (merge/p-tuning-v2 + a 67-test suite across 9 adapter variants), oxicuda-meta (few-shot/linear-head + a 69-test suite across the MAML family), oxicuda-numeric (8 Gauss-Patterson exactness tests to degree 46), oxicuda-evol (8 CMA-ES Jacobi-eigensolver spectral-identity tests), oxicuda-solver (16 PDE/ODE tests with measured O(h²) convergence), oxicuda-ann (43 tests across HNSW/kNN-graph/IVF/IVFPQ).
• Test suite expanded to 38,093 passing tests (workspace-wide, --all-features; 37,166 with default features), up from 36,984 at 0.3.0.

Changed

• oxicuda-solver: syevd (symmetric eigensolver) and the blocked Householder QR / one-sided-Jacobi SVD device paths previously launched an incomplete GPU kernel and read back fabricated (never-computed) values. Replaced with an explicit, documented exact-CPU host fallback — no GPU acceleration yet, pending on-device follow-up — rather than silently returning wrong results.
• oxicuda-ssl: barlow_cross_corr_wgmma, nt_xent_softmax_warp, and gather_features_bulk are now documented as intentionally Hopper/Blackwell-only PTX (wgmma, redux.sync, TMA); each has an on-device-confirmed portable scalar fallback for Ampere and older.

Fixed

• Register-shadowing of CUDA's built-in special registers (.reg declarations literally named %tid/%ntid/%ctaid/%warpid, clobbering the special registers like %tid.x actually read from) — the single most common defect class this pass found, affecting oxicuda-primitives, oxicuda-train (all 9 optimizer kernels), oxicuda-ann (hnsw_neighbor_eval/ivf_assign/topk_select), oxicuda-rl (all 5 kernels), oxicuda-dist-infer (all 5 kernels), and oxicuda-timeseries (all 7 kernels, plus a special register used directly as a mad operand). All renamed to non-colliding register names.
• Base-2 (ex2.approx/lg2.approx) used where the math needs base-e — silently plausible, genuinely wrong: oxicuda-survival (Cox risk/score/info, ~18–30% off), oxicuda-seq (HMM forward log-sum-exp, ~30% off), oxicuda-ot (Sinkhorn/unbalanced log-sum-exp, ~20% off), oxicuda-rlhf (BT/DPO/KTO losses), oxicuda-nerf (volume_render's alpha compositing), oxicuda-gnn (softmax_edge, masked by softmax's own scale-invariance — only a base-e CPU oracle caught it), and oxicuda-audio (ctc_alpha_kernel, bundled with two other defects below). All fixed with the correct log2(e)/ln(2) scaling.
• Kernels that never compiled at all — invalid PTX rejected outright by ptxas: undeclared/out-of-range registers (oxicuda-multimodal's bilinear_pool/temporal_pool, oxicuda-audio's ctc_alpha_kernel, oxicuda-quant's all 5 kernels, oxicuda-moe's expert_dispatch_kernel), a duplicate register declaration (oxicuda-geom2d's point_in_aabb), mid-function .reg declarations (oxicuda-distill's at_pool_kernel/gram_matrix_kernel), a missing .reg .pred entirely (oxicuda-recsys's als_update_step), illegal scaled-register shared-memory addressing (oxicuda-ann's topk_select, oxicuda-infer's logits_softmax, oxicuda-causal's expm_pade_kernel across dozens of sites — the latter also stored immediate literals directly via st.shared.f32, which st cannot take as a source operand), the [smem] bracket form used as arithmetic instead of load/store (oxicuda-lm's rms_norm/causal_attn_softmax), a non-existent atom.exch.s32 (oxicuda-tda's boundary_reduce, needs .b32), a non-existent cos.approx.f64/lg2.approx.f64 (oxicuda-fft's precompute_window, oxicuda-evol's gaussian_mutate_kernel), an unsupported 4-byte cp.async.cg transaction (oxicuda-cs's iht_step_cp_async, needs .ca for sub-16-byte transactions — a separate deadlock from out-of-range threads skipping bar.sync was fixed in the same kernel), a malformed branch label (oxicuda-sparse's spmv_bsr), and invalid braced predication plus nonexistent f64 SFU forms (oxicuda-privacy, 6 of 7 differential-privacy kernels). All now compile and are ptxas-verified on sm_86.
• Kernels that compiled but computed the wrong thing: oxicuda-manifold's knn_topk was correct only for k=1 (missing the ascending bubble-up pass for k>1); oxicuda-vision's bilinear_interp/roi_align used the non-existent floor.f32 (replaced with cvt.rmi.f32.f32); oxicuda-stats's mean_var/rank_assign were missing the .rn rounding qualifier on cvt.f32.u32; oxicuda-quantum's statevector simulator had 8 stacked defects (partial 4×4 gate matrix-vector products, wrong bit-insertion masks, an unguarded swap race, divergent-lane shfl.sync, wrong Taylor-series hex-float constants) — essentially every operation was wrong; oxicuda-anomaly's lof_reach_dist_kernel had a loop-index register clobber that collapsed every reach_dist to kd_j; oxicuda-mamba's parallel_scan used the wrong shuffle direction (100% wrong) and wkv_forward used the wrong softmax pivot (42.6% error); oxicuda-audio's rel_pos_bias_kernel had an unsigned-underflow clamp bug and stats_pool_kernel had a 32-lane write race; oxicuda-nas's gumbel_softmax_kernel used log2(e) where ln(2) was needed (its reciprocal, ~2.08× error); oxicuda-geometry3d's sh_eval_kernel referenced a register past its declared bank and dropped a dx factor from every channel's c1·Y11 term; oxicuda-moe's expert_ffn_kernel GELU tanh approximation was missing a factor of 2 in its exponent; oxicuda-infer's flagship paged_attention had 5 stacked defects (invalid 64-bit mul.wide.u32, a partial dot product instead of the full Σ_d, V read through the K pointer, base-2 softmax, wrong GQA head mapping) and rope_apply dropped a sign term plus used an imprecise log2(10000) constant.
• oxicuda-solver: lu.rs's launch_gemm_update swapped the grid X/Y axes against the actual row/column mapping (dropping part of the GEMM update on non-square trailing tiles), and its Padé matrix-exponential kernel always loaded f64 coefficients even in the f32 variant — both fixed. Separately, the LU (panel_lu/trsm_unit_lower/gemm_update/pivot_swap) and Cholesky (panel_cholesky) kernels were literal ret; stub bodies performing no factorization at all — implemented with real bar.sync-staged panel/trailing updates, validated by two new on-device test suites including an independent splitmix64-seeded adversarial harness.
• oxicuda-rand: AES round-key words needed .swap_bytes() for the device's endianness; mrg32k3a/philox/xorwow's Box-Muller kernels shared one f32/f64 register pool, corrupting Gaussian sampling on all three engines; 4 philox_optimized branch targets were missing the $ label sigil.
• oxicuda-signal: dct2_permute/dct3_pretwiddle/dct3_unpermute/fir_direct used illegal brace predication; fir_direct's bounds guard was an always-false src > u64::MAX (silently zeroing all output); dct3_pretwiddle/dct4_postscale emitted f32 immediates into f64 instructions.
• oxicuda-ssl: random_mask_kernel had a spurious ×0.5 roughly doubling the effective drop probability.
• oxicuda-tabular: feature_tokenize_kernel addressed its weight/bias rows with stride 1 instead of feat*embed_dim and never looped over the embedding dimension.
• oxicuda-rlhf: dpo_loss_kernel's grid-stride accumulation clobbered the register holding beta with an atomic's return value; ppo_rlhf/ppo_step.rs's per-step value-loss loop (a CPU-side bug, found by the new test suite) read values[0] for every step instead of the current step's value.
• oxicuda-ann: ivf/ivf.rs search (CPU-side, found by the new test suite) indexed insertion-ordered vectors with a list-traversal counter instead of per-list storage, scoring against the wrong stored vector whenever add() calls interleaved across coarse lists.
• oxicuda-recsys: (CPU-side, found by the new test suite) multitask/ple.rs fed shared experts at layer>0 the wrong-length input from the previous layer; factorization/als.rs's gauss_jordan had a redundant row-swap that skipped exchanging column 0 whenever the pivot was off-diagonal.
• oxicuda-blas (GEMM), oxicuda-dnn (LayerNorm, and separately implicit_gemm/conv1x1/depthwise which were comment-only stubs with no arithmetic), oxicuda-sparse (f64 CSR SpMV, including an illegal 64-bit shfl.sync.down.b64 split into .b32 halves, and the mixed-precision SpMV FP64 path): a shared invalid-PTX bug class — an .f32-declared register bank and/or single-precision zero literal used in .f64 instructions. Root-caused once in oxicuda-ptx's PtxType (precision-correct zero-literal encoding + correctly-rounded cvt selection) and applied across all affected sites; LayerNorm's .maxntid directive was also misplaced inside the kernel body with a stray semicolon.
• oxicuda-snn: the atan surrogate-gradient kernel was off by π² (α·π/(1+x²) instead of α/(π·(1+x²))), affecting every SM target sm_75–sm_100.
• oxicuda-webgpu: conv2d WGSL codegen named a buffer filter, a reserved WGSL keyword, failing naga parsing outright — renamed to kernel_w.
• oxicuda-metal: cleared 2 clippy warnings in memory.rs (missing dead_code cfg-gate, a redundant explicit drop).

Full Changelog: v0.3.0...v0.4.0