det-infer

Architecture-invariant bitwise LLM inference, loadable through kernels. The model-as-a-kernel engine with every architecture-dependent operation removed: logits and greedy token streams are bit-identical on every CUDA architecture, so the same weights and prompt reproduce the same bits on an L4, an H200, or a Blackwell workstation, indefinitely. An entire greedy generation, prompt consumption included, is one kernel launch. Batched decode carries the guarantee per sequence: a sequence's logits are identical whether it runs alone or alongside others in a batch, so a batch is reproducible on every card and independent of how sequences are grouped.

Usage

from kernels import get_kernel

di = get_kernel("phanerozoic/det-infer", version=1, trust_remote_code=True)

mm = di.MegaModel.from_pretrained("HuggingFaceTB/SmolLM2-135M")
tokens = mm.generate(prompt_ids, max_new=128)   # one kernel launch, total
logits = mm.decode_step(token_id, pos)          # one launch, fp32 [V] logits

streams = mm.generate_batch([ids_a, ids_b, ids_c], max_new=128)  # batched

version selects the release branch; trust_remote_code is required by kernels for publishers without the trusted-publisher mark.

The API is identical to model-as-a-kernel: decode_step, prefill, generate(single_launch=True), decode_step_timed, the batched generate_batch / decode_batch (up to batch_max() sequences, 16 on the llama-family path), and the raw ops.mak_* program launches. Supported: llama-family architectures with GQA, optional qk-norm, and plain rope (attention scaling 1), plus the gemma-4-E2B architecture (per-layer embeddings, shared KV, sliding windows, gelu gating, softcapped head; batch 1 on gemma), single device, positions to one million (the Cody-Waite rope reduction bound). Weights are dense bf16 or bitsandbytes nf4; nf4 weights are stored packed and dequantized in the kernel. A GGUF (llama.cpp) checkpoint loads with MegaModel.from_gguf, dequantized to bf16 at load through the gguf reader, so any GGUF quant type runs with the same guarantee.

API

How it works

GPU floating-point add, multiply, fma, divide, and sqrt are IEEE 754 bit-exact on every architecture; what varies between cards is reduction order and the vendor math library. The engine already fixes reduction order by construction (fixed trees, lane-striped GEMV accumulation, launch geometry that never touches arithmetic — the cp.async weight streaming and the attention K/V staging are pure scheduling). det-infer removes the second source: exp, tanh, cos, sin, and rsqrt run through IEEE-only implementations instead of libdevice. exp runs entirely in double-float fp32 pairs (Dekker exact products and Knuth exact sums on the full-rate fma units, an exactly representable three-term Cody-Waite ln2 split, one rounding at the collapse); tanh is built on that exp; sin/cos use an fp64 two-term reduction with double-float polynomials; rsqrt runs through fp64 sqrt and divide. The sum-of-squares accumulations are written with explicit round-per-op intrinsics so ptxas cannot contract them differently per target. For llama-family ropes the frequency table is computed by a fixed 50-digit decimal recipe (exp and ln series, llama3 wavelength scaling in fp64 rational arithmetic) rather than taken from any host math library; gemma-4 tables come from the model's own initialization. No operation in the forward pass depends on the architecture or the CUDA toolchain.

nf4-quantized weights are stored packed and dequantized inside the GEMV: each 4-bit index selects a codebook entry, multiplied by its block's absmax and rounded to bf16, with the products accumulated in the same order as the dense path. The dequantization is integer indexing and IEEE multiplication and does not depend on the architecture.

Certification is a pinned digest: the test suite hashes the fp32 logits of a fixed teacher-forced sequence and a greedy generation for each covered model. The SHA-256 digests are constants in the test file; passing on a new card, torch release, or CUDA toolchain is the proof of bit-identity with every other combination.

Correctness

Verified against transformers eager: argmax agreement and mean logit deviation at the same level as model-as-a-kernel (the IEEE-only transcendentals are at least as accurate as libdevice; Llama-3.2-1B teacher-forced argmax agrees 64/64). Fused, phase-per-launch, and repeated runs are mutually bitwise identical, and the single-launch generation emits the same tokens as one launch per token.

Invariance: pinned logits and greedy-token digests for SmolLM2-135M, Qwen3-0.6B, Llama-3.2-1B, Llama-3.2-3B, zephyr-7b-beta (the Mistral architecture), and gemma-4-E2B-it, minted on RTX 6000 Ada (sm89), reproduce byte-for-byte on L4 (sm89), H200 (sm90), and RTX PRO 6000 Blackwell (sm120), under torch 2.11, 2.12, and 2.13, across compiled variants spanning the cu126, cu128, cu130, and cu132 toolchains: nine card-torch-toolchain combinations against one digest set (sm120 runs cu128 and newer; older toolchains do not target it). The digests are further invariant to launch geometry (grid sizes 64 to 284 blocks, weight-ring depths 1 to 8), to concurrent device load, and along 1500-token generations; the attention chunk length sets the softmax partial split and is fixed at 128. Batch invariance is exact: a sequence's fp32 logits are bitwise identical whether it decodes alone or inside a batch of 16, so the per-sequence digest is the batched certificate as well. The digests are the constants in tests/test_det_infer.py; reproducing them on a new combination is the certificate of bit-identity with every other.

Measured

The deterministic math costs about 8 percent of decode throughput on SmolLM2-135M on RTX 6000 Ada (824 vs 898 tok/s closed loop, medians of three runs on the current builds of both engines). Consumer cards issue fp64 at 1/64 rate, so the hot transcendentals run in double-float arithmetic on the full-rate fp32 units; fp64 remains only in the sincos range reduction and the per-row rsqrt. Speedups over transformers eager match model-as-a-kernel's order of magnitude.

Batched decode amortizes the weight reads across sequences. Steady-state throughput on Llama-3.2-1B (RTX 6000 Ada) rises from 256 tokens per second at batch 1 to 1673 at batch 16, a 6.5x aggregate; SmolLM2-135M reaches 5.2x. Sequences up to 8 keep their input staged in shared memory; wider batches route the large projections through a global scratch, since the persistent kernel holds one occupancy across all phases and a shared panel large enough for 16 sequences would halve it.

License

Apache-2.0.