Module oracle 

GPU-forge differential oracle: CPU references, the numeric comparator, the GpuEvaluation evidence bundle, and the per-kernel evaluators that certify emitted shaders against those references.

Split by concern into submodules ([params], [report], [reference], [evaluate], [kernels]); the full public surface is re-exported here so every existing crate::wgsl_forge::oracle::<Item> path resolves exactly as before.

Structs

AffineParams

ComparisonReport

FfnParams

16-byte uniform block for the fused-FFN kernel.

FftParams

16-byte uniform block for the radix-2 FFT kernel: n complex elements, log2n = log2(n). The kernel runs one workgroup of n threads.

GemmParams

16-byte uniform block for the dense GEMM kernel: row-major C[M×N] = A[M×K]·B[K×N].

GemvParams

16-byte uniform block for the dense GEMV kernel: row-major y[M] = A[M×N]·x[N].

GpuEvaluation

OracleCase

OracleTolerance

TernaryGemvParams

16-byte uniform block for the ternary-GEMV kernel (k_words == ceil(k/16)).

TopKParams

16-byte uniform block for the top-k kernel (block_size == workgroup size).

Constants

TERNARY_CODES_PER_WORD

Number of 2-bit ternary codes packed into one u32 word.

Functions

affine_cpu

candidate_evaluation

certify_builtin

compare_f32

dft_cpu

Naive O(N²) forward Discrete Fourier Transform, the reference the GPU radix-2 FFT is differentially checked against. Complex data is interleaved f32: element j is (input[2*j], input[2*j+1]) = (real, imag), so both the input slice and the returned vector hold 2*N f32.

evaluate_affine

Cross-backend differential-oracle evaluation of the affine kernel against affine_cpu (plan §7). Generic over OracleContext, so the same code runs on wgpu (via WgpuComputeContext) and CUDA (via CudaComputeContext); the backend only differs inside OracleContext::run_kernel. The CPU-reference vectors, bindings, dispatch sizing and tolerance are identical to what the wgpu-inline and evaluate_affine_cuda paths used before unification.

evaluate_affine_cuda

Cross-backend oracle (plan §7/§10): runs the affine kernel through the native CUDA backend (CUDA-C compiled to PTX by NVRTC) and checks it against the same CPU reference vectors used for the wgpu backend. Requires a CUDA device.

evaluate_builtin

evaluate_coopmat_loadstore

Diagnostic: cooperative-matrix load→store round-trip (no multiply). Loads a as a role-C fragment and stores it to c; c must equal a. This verifies coopLoadT/coopStoreT work on the adapter (they do — the coopMultiplyAdd path is the one currently blocked on the experimental backend).

evaluate_ffn

Cross-backend differential-oracle evaluation for the fused FFN against ffn_cpu (plan §7). Generic over OracleContext — the same code runs on wgpu and CUDA; only OracleContext::run_kernel differs. One workgroup-thread per output element; the output buffer is output_size. Tensors, bindings, dispatch sizing and tolerance are identical to the prior wgpu/evaluate_ffn_cuda paths.

evaluate_ffn_cuda

Cross-backend oracle for the fused FFN via the CUDA backend. Thin wrapper over the generic evaluate_ffn (warmups = 0, samples = 1, one dispatch as before).

evaluate_fft

Differential-oracle evaluation of the radix-2 FFT (out = forward DFT(in)) against dft_cpu. One workgroup of n = schedule.workgroup_size threads (one complex element per thread; n must be a power of two), mirroring the single-workgroup dispatch of [evaluate_topk]: element_count = n with workgroup_size = n launches exactly one workgroup. The input/output buffers hold 2*n interleaved f32.

evaluate_gemm

Differential-oracle evaluation for the dense GEMM against gemm_cpu. One workgroup-thread per output element; the output buffer is m * n elements. Row-major C[M×N] = A[M×K] · B[K×N], all f32.

evaluate_gemv

Differential-oracle evaluation for the dense GEMV against gemv_cpu. One workgroup-thread per output ROW; the output buffer is m elements. Row-major y[M] = A[M×N] · x[N], all f32.

evaluate_matmul_tc

Differential-oracle evaluation of the cooperative-matrix (tensor-core) 8x8 GEMM tile C = A * B against matmul_cpu. All-f32 — the only coopmat configuration wgpu/naga 29 implements (see crate::wgsl_forge::emit::coopmat). One subgroup (32-lane NVIDIA warp) cooperatively computes the tile; the row-major loads/store reproduce the row-major CPU reference, so agreement is to f32 precision (a tiny tolerance covers tensor-core accumulation order).

evaluate_matmul_tc_cuda

Tensor-core oracle: runs the genuine f16-input WMMA GEMM (C = A * B, 16x16x16) on the CUDA backend via the nvcuda::wmma fragment API, compiled by NVRTC for the device’s compute capability, and checks it against the row-major CPU reference. This is the reduced-precision tensor-core path (f16 A/B inputs, f32 accumulator) that wgpu/naga 29’s cooperative-matrix backend cannot express — 29 implements only all-f32 8x8x8, and even that multiply is non-functional on the 29.0.3 execution path (no published fix; see crate::wgsl_forge::emit::coopmat). Requires a CUDA device with compute capability >= 7.0 (Volta+).

evaluate_p64

Differential-oracle evaluation for the P64 projection against p64_project_cpu.

evaluate_rayprobe

Differential-oracle evaluation of the ray-query (ray-probe) kernel: builds a BLAS+TLAS for rayprobe_scene, dispatches the emitted ray_probe WGSL over rayprobe_rays on the GPU, and checks the committed hit distances against rayprobe_cpu. Requires a ray-query-capable adapter (RT cores).

evaluate_ternary_gemv

Differential-oracle evaluation for the ternary GEMV against ternary_gemv_cpu. One workgroup-thread per output row; the output buffer is m elements.

evaluate_topk

Cross-backend differential-oracle evaluation for the top-k kernel against topk_cpu (plan §7). Generic over OracleContext — the same code runs on wgpu and CUDA; only OracleContext::run_kernel differs.

evaluate_topk_cuda

Cross-backend oracle for top-k via the CUDA backend (CUDA-C __shared__). Thin wrapper over the generic evaluate_topk (warmups = 0, samples = 1, one dispatch as before; block_size is schedule.workgroup_size, i.e. 64).

ffn_cpu

CPU reference for the fused FFN, matching the emitted kernel’s op order exactly (hidden outer, input inner) so GPU/CPU agree within tolerance: out[o] = sum_h w2[o,h] * gelu(sum_i w1[h,i] * input[i]).

ffn_tensors

Deterministic FFN test tensors. Weights are scaled by 1/sqrt(fan_in) so the pre-activations stay O(1) and GPU/CPU agree within a modest tolerance.

fft_inputs

Deterministic complex test signal as interleaved f32 (2*n values), drawn from the same xorshift stream as every other oracle vector so it is reproducible. Both the real and imaginary parts land in [-1, 1].

gemm_cpu

Row-major general dense GEMM reference, the bit-for-bit mirror of the emitted gemm kernel: C[M×N] = A[M×K] · B[K×N], i.e. C[i][j] = sum_{k<K} A[i*K + k] * B[k*N + j]. The inner-sum order (k ascending) matches the kernel’s kk loop so GPU/CPU agree to f32 summation precision.

gemm_tensors

Deterministic GEMM test tensors: A (M×K) and B (K×N), both drawn from the same xorshift stream as every other oracle vector and scaled by 1/sqrt(K) so the length-K dot products stay O(1) and GPU/CPU agree within a tight tolerance.

gemv_cpu

Row-major dense GEMV reference, the bit-for-bit mirror of the emitted gemv kernel: y[M] = A[M×N] · x[N], i.e. y[i] = sum_{j<N} A[i*N + j] * x[j]. The inner-sum order (j ascending) matches the kernel’s j loop so GPU/CPU agree to f32 summation precision.

gemv_tensors

Deterministic GEMV test tensors: A (M×N) and x (N), both drawn from the same xorshift stream as every other oracle vector and scaled by 1/sqrt(N) so the length-N dot products stay O(1) and GPU/CPU agree within a tight tolerance.

matmul_cpu

Row-major n×n matrix multiply reference: c[i][j] = sum_k a[i][k] * b[k][j].

p64_project_cpu

CPU reference for the P64 projection: out[r] = sum_w weights[w] * f32(word_w), reading the 16 packed u32 words in the same lane order as the kernel.

p64_records

Deterministic P64 descriptors with small (f32-exact) u32 words.

rayprobe_cpu

CPU reference for the ray-probe kernel: for each ray, the nearest committed triangle hit t within [t_min, t_max], or -1.0 on a miss — matching the emitter’s hits[i] = committed.t else -1.0.

rayprobe_rays

The fixed ray set as the 8-float-per-ray layout the emitter expects (origin.xyz, dir.xyz, t_min, t_max). All rays originate at z = -1 and point along +z; hit rays target clear triangle interiors (away from edges) so GPU BVH traversal and the CPU reference agree, and miss rays point well outside.

rayprobe_scene

The fixed ray-probe scene: three world-space triangles as a flat f32 list (9 floats/triangle = 3 verts × xyz). Two coplanar triangles tile the quad [0,2]² at z = 2; a third sits behind them at z = 4, so rays through the lower-left region hit two triangles and must commit the nearer one (t at z=2).

ternary_gemv_cpu

CPU reference for the BitNet-style ternary GEMV, the bit-for-bit mirror of the emitted kernel: out[o] = scale[o] * sum_{i<K} ternary(w[o,i]) * x[i].

ternary_gemv_tensors

Deterministic ternary-GEMV test tensors: the activation vector x (length K), the 2-bit-packed ternary weights (M * ceil(K/16) words), and the per-row scales (length M). Codes are drawn from the xorshift stream and reduced into {0,1,2} so the weights only ever decode to {0, +1, -1} (never the unused 3), keeping the GPU and CPU paths bit-identical.

topk_cpu

CPU reference for the per-block top-k: the k largest values of each block_size-element block, in descending order. Blocks shorter than block_size (the tail) are padded with the sentinel, mirroring the GPU kernel’s out-of-range loads.

topk_inputs

Deterministic xorshift test vector in [-1, 1], matching the affine generator.