Struct ForgeRuntime 

pub struct ForgeRuntime { /* private fields */ }

A ready-to-use handle for running certified forge kernels on real data.

Owns the GPU compute context (one device/queue/slab pair) and, optionally, a ManifestCache directory plus this machine’s topology hash so tuned schedules can be looked up. Construct once and reuse across many calls; each topk / ternary_gemv / p64_project call allocates transiently, dispatches, reads back, and frees its transient allocations.

Implementations

impl ForgeRuntime

pub fn new( capacity_bytes: usize, cache_dir: Option<PathBuf>, ) -> Result<Self, ForgeError>

Build the GPU context, optionally attaching a manifest-cache directory for tuned schedules.

capacity_bytes sizes the device slab (inputs + outputs must fit within it per call). When cache_dir is Some, Self::tuned_schedule will consult the cache; when it is None, every kernel uses its documented default schedule. The topology hash used for cache keys is derived from the live adapter, so a cache produced on different hardware is simply never matched (it is not an error).

Example

let mut rt = ForgeRuntime::new(64 * 1024 * 1024, None)?;
let top = rt.topk(&[3.0, 1.0, 2.0, 0.5], 2)?; // largest-2 per block

pub fn tuned_schedule(&self, builtin: BuiltinKernel) -> Schedule

The tuned Schedule for builtin on this hardware.

If a cache is attached and holds a [TuningManifest] for (topology_hash, builtin), the winning schedule from that record is returned. Otherwise — no cache, no topology hash, no record for this kernel, or a cache read error — the per-kernel default is returned.

Default policy (matches what the evaluate_* oracle paths use): every built-in defaults to Schedule { workgroup_size: 64, .. } (items_per_invocation = 1, vector_width = 1). For top-k that 64 is also the per-block size (block_size == workgroup_size), so the default top-k processes the input in 64-element blocks. Populate the cache for this machine with shader auto-tune-all.

pub fn topk(&mut self, input: &[f32], k: usize) -> Result<Vec<f32>, ForgeError>

Real-data per-block top-k: returns, for each block_size-element block of input (with block_size = tuned_schedule.workgroup_size), the k largest values in descending order, concatenated block-by-block.

The tail block (when input.len() is not a multiple of block_size) is padded with f32::MIN by the kernel, exactly as the certified path does, so short blocks still emit k values (the padding sentinels sort last).

Buffer wiring is identical to evaluate_topk: binding 0 = input (storage-read), binding 1 = output (storage-read-write, num_blocks * k f32s), binding 2 = TopKParams (uniform); dispatch element_count = input.len(). The CALLER’s input is fed directly — no oracle, no test vectors, no comparison.

Example

let top2 = rt.topk(&[5.0, 1.0, 9.0, 2.0], 2)?; // one 4-elem tail block -> [9.0, 5.0]

pub fn ternary_gemv( &mut self, x: &[f32], packed_w: &[u32], scale: &[f32], m: usize, k: usize, ) -> Result<Vec<f32>, ForgeError>

Real-data ternary (BitNet-style) GEMV with on-the-fly dequant: out[o] = scale[o] * sum_{i<k} ternary(w[o][i]) * x[i] for m output rows.

packed_w holds the 2-bit ternary codes, 16 codes per u32 (0 -> 0.0, 1 -> +1.0, 2 -> -1.0, code 3 unused), laid out as m rows of ceil(k / 16) words each, row-major. scale has length m, x length k.

Buffer wiring is identical to evaluate_ternary_gemv: binding 0 = x, 1 = w_packed, 2 = scale (all storage-read), 3 = output (storage-read-write, m f32s), 4 = TernaryGemvParams (uniform); dispatch element_count = m. The CALLER’s tensors are fed directly — no oracle.

Example

// 2 rows x 4 cols, codes packed low-to-high (1=+1, 2=-1): row0=+1,+1,+1,+1; row1=-1,-1,-1,-1
let out = rt.ternary_gemv(&[1.0, 2.0, 3.0, 4.0], &[0x55, 0xAA], &[2.0, 10.0], 2, 4)?;
// -> [2*(1+2+3+4), 10*(-1-2-3-4)] = [20.0, -100.0]

pub fn p64_project( &mut self, records_bytes: &[u8], weights: &[f32], record_count: usize, ) -> Result<Vec<f32>, ForgeError>

Real-data P64 projection: out[r] = sum_{w<16} weights[w] * f32(p64[r].word[w]) for record_count records.

records_bytes is the packed P64 GPU words exactly as evaluate_p64 lays them out — a contiguous array of P64GpuWords64 (64 bytes / record, 16 u32 words / record), so records_bytes.len() must be record_count * 64. weights has length 16.

Buffer wiring is identical to evaluate_p64: binding 0 = input (P64 records, storage-read), 1 = weights (storage-read), 2 = output (storage-read-write, record_count f32s); dispatch element_count = record_count. The CALLER’s records are fed directly — no oracle.

Example

let recs = [P64GpuWords64::from_u64_fields([1, 0, 0, 0, 0, 0, 0, 0])];
let bytes: &[u8] = bytemuck::cast_slice(&recs);
let weights = [1.0f32; 16];
let out = rt.p64_project(bytes, &weights, 1)?; // out[0] = word[0] = 1.0

pub fn gemm( &mut self, a: &[f32], b: &[f32], m: usize, k: usize, n: usize, ) -> Result<Vec<f32>, ForgeError>

Real-data dense GEMM: row-major C[M×N] = A[M×K] · B[K×N], all f32, i.e. C[i][j] = sum_{k<K} a[i*K + k] * b[k*N + j], for m * n output elements.

a must have m * k elements and b must have k * n elements, both row-major. The returned vector has m * n elements, row-major.

Buffer wiring is identical to evaluate_gemm: binding 0 = a, 1 = b (both storage-read), 2 = c (storage-read-write, m*n f32s), 3 = GemmParams (uniform); dispatch element_count = m*n. The CALLER’s matrices are fed directly — no oracle, no test vectors.

Example

// A (2×3) · B (3×2): A=[[1,2,3],[4,5,6]], B=[[7,8],[9,10],[11,12]]
let a = [1.0, 2.0, 3.0, 4.0, 5.0, 6.0];
let b = [7.0, 8.0, 9.0, 10.0, 11.0, 12.0];
let c = rt.gemm(&a, &b, 2, 3, 2)?; // -> [58, 64, 139, 154]

pub fn gemv( &mut self, a: &[f32], x: &[f32], m: usize, n: usize, ) -> Result<Vec<f32>, ForgeError>

Real-data dense GEMV: row-major y[M] = A[M×N] · x[N], all f32, i.e. y[i] = sum_{j<N} a[i*N + j] * x[j], for m output rows.

a must have m * n elements (row-major) and x must have n elements. The returned vector has m elements.

Buffer wiring is identical to evaluate_gemv: binding 0 = a, 1 = x (both storage-read), 2 = y (storage-read-write, m f32s), 3 = GemvParams (uniform); dispatch element_count = m. The CALLER’s matrix/vector are fed directly — no oracle, no test vectors.

Example

// A (2×3) · x (3): A=[[1,2,3],[4,5,6]], x=[1,1,1]
let a = [1.0, 2.0, 3.0, 4.0, 5.0, 6.0];
let x = [1.0, 1.0, 1.0];
let y = rt.gemv(&a, &x, 2, 3)?; // -> [6, 15]

pub fn fft( &mut self, complex_interleaved: &[f32], ) -> Result<Vec<f32>, ForgeError>

Real-data forward FFT: out = DFT(in) of n = complex_interleaved.len()/2 complex points, computed by the workgroup-local radix-2 Decimation-In-Time kernel. The input and output are interleaved f32 — element j is (buf[2*j], buf[2*j+1]) = (real, imag) — so both have length 2*n.

Precondition: n must be a power of two and <= 1024 (the kernel runs ONE workgroup of n threads, so n is also the workgroup size, capped by the maximum workgroup size). Unlike the other runtime methods, the schedule here is pinned to workgroup_size = n (the transform length is the parallel width), not the tuned default; n ranges over distinct power-of-two sizes.

Buffer wiring is identical to evaluate_fft: binding 0 = input (storage-read, 2*n f32), 1 = output (storage-read-write, 2*n f32), 2 = FftParams (uniform); dispatch element_count = n so exactly one workgroup launches. The CALLER’s signal is fed directly — no oracle, no test vectors, no comparison.

Example

// A real unit impulse at index 0 (rest zero) has a flat spectrum: all ones.
let mut signal = vec![0.0f32; 2 * 8];
signal[0] = 1.0;
let spectrum = rt.fft(&signal)?; // every bin == (1, 0)