Introducing cu-Jev (CUDA-Jev)

Almost everything we call an LLM serving engine is built to decode. Grow a KV cache, sample the next token, repeat. That is the right shape for a chatbot. It is the wrong shape for a decision.

A System One request is one long shared context and a handful of short questions about it. At the end of each question you do not generate anything. You read a few logits and stop. TypeSafe shipped that contract as Jev: a state, a map of typed questions (noul / choice / score), and typed answers with probabilities. Nothing is decoded, so output tokens are free, because there are none. The official model is closed. The open copies I could find wrap a chat engine and scrape logprob_token_ids. That works. It is also the wrong machine.

I wanted the contract on an open checkpoint, running as a CUDA engine rather than as a wrapper around a chat server. That is cu-Jev (CUDA-Jev, the same naming family as cuBLAS). Open Qwen3.5 checkpoints, 0.8B through 9B, become a decision function with Jev's wire format. The runtime is C and CUDA. Python is a ctypes binding, a prompt layout, and a FastAPI server. Point TYPESAFE_BASE_URL at it and the official SDK keeps working.

The workload

A request has to do three things:

1. ingest the state once,
2. run every question against that state without letting the questions see each other,
3. at the last token of each question, read the logits of a few vocabulary ids — the option labels.

There is no sampling and no KV growth. It is batched prefill with a shared prefix and a tiny gather at the end.

Each question is a short branch. It attends to the state's KV cache, and every linear-attention layer starts from the state's recurrent Gated-DeltaNet state: one copy, read-only, never duplicated per question. Branches never see each other. I tested that with an adversarial sibling that tries to pull the answer the other way. It does not move a branch past the cuBLAS reduction-order noise, which the isolation test bounds at about 0.03 logits.

The prompt is one shared prefix — system rules, then STATE:, then the state as text or pretty JSON — and one suffix per question, ending at the assistant's first token, where the labels are read. choice uses verified single-token labels A…Z, AA…; noul reads yes / no; score reads the digits 0…K-1. Softmax over those logits is the base model's next-token distribution, renormalized to the options. Confidence is $1-H(p)/\ln K$.

Why not vLLM

jevfire, openjev-sglang and simple-jev all sit on a chat engine and read logprobs. That is a reasonable first cut, and I would have stopped there if the prefix and the read-out were cheap enough. They are not, once you look at what the request is actually doing.

In cu-Jev the prefix is one KV copy plus one recurrent state that every branch reads. There are no prefix-cache blocks that have to land on the right alignment to hit. The read-out is a gather-dot over at most 255 rows of the embedding, not a 248k-wide LM head at every position. There is no scheduler, no tokenizer, and no Python between the request and the GEMMs: the whole eval is one synchronous C call on one stream. The price is generality. It runs Qwen3.5 and nothing else.

Qwen3.5 is a hybrid. Every fourth layer is gated GQA; the other three are Gated DeltaNet. Full attention has an output gate, zero-centred RMSNorm on q and k, and partial RoPE. The linear layers hold a 128×128 fp32 recurrent state per head. I wrote both mixers. Attention is a FlashAttention-2-style kernel on mma.sync.m16n8k16, bf16 in, fp32 accumulate. The GDN scan keeps the recurrent state in registers, one column split over two lanes, with no spills. GEMMs go through cuBLAS. Projections are row-fused at load, so each layer is four GEMMs. The residual stream stays fp32; GEMM inputs and outputs are bf16.

Does it match the reference

tests/python/test_oracle.py runs real prompts through cu-Jev and through transformers, then compares the full-vocabulary logits at the read-out position. Identical argmax and top-5 on every probe for 0.8B, 2B, 4B and 9B. Against an fp32 run of the reference, cu-Jev's next-token distributions sit closer to fp32 than Hugging Face's own bf16 run, because the residual stream never leaves fp32 while the GEMMs still run in bf16. On 4B the L1 distances are 0.006 / 0.001 / 0.009 for cu-Jev against 0.012 / 0.008 / 0.060 for HF-bf16.

The edge-case suite sweeps 55 (state, question) length pairs across every 16-token query tile and 64-key tile boundary, packs 40 branches into one call, checks 256 candidates against full logits, and covers state-reuse determinism plus every capacity error. 59 tests, all matching HF, with a synchronize-and-error-check after every kernel.

Shared prefix is the whole point

The state is the expensive part. If you pay for it once, sixteen short questions should be cheap. If you re-prefill the state with every question, you are running a chat engine in disguise. Same tokens, same model, three ways to answer 16 questions of 40 tokens on a 1 024-token state. First the RTX 5090:

Same workload on an RTX 3090:

Sharing the state is 13× faster per request, and 35× once the state is already cached. Against Hugging Face doing the same batched read-out, the 4B is 10× faster, 27× cached. The ratios match on both GPUs. Argmax agreement with the naive path and with HF is 100%.

The raw engine, cujev-bench, on CUDA 13.3, bf16:

On the 4B, nsys puts about 60% of the time in cuBLAS GEMMs at the bf16 tensor-core ceiling, about 210 TFLOPS on the 5090 and about 70 on the 3090, about 15% in the recurrent GDN scan, and about 10% in tensor-core attention. A 40-token question on the 4B costs about 2 ms on the 5090 and about 5 ms on the 3090, no matter how long the state is. The scan is serial over the state. A chunked (WY) prefill is the next kernel.

Accuracy, without a trained head

cu-Jev does not train a decision model. It reads what the base checkpoint already knows. On LocalLLaMA/typed-decisions, test split, 400 cases × 5 questions, zero-shot, one prompt for all four enterprise workflows:

For scale, these are numbers other projects publish on the same split, theirs rather than mine: TF-IDF plus logistic regression trained on the train split at 66.1%, TypeSafe Jev 1.13 at 72.7%, Laya at 76.6% and openJev-verdict-2.0 at 77.1%, the last two both trained encoders. Jev claims calibrated probabilities from RLCD. cu-Jev's probabilities are the base model's, renormalized over the options. They are not calibrated that way. Temperature or isotonic calibration on this split, reported per question type, is next, and so is prompt tuning per workflow. Right now the number is what the open checkpoint knows with a generic prompt.

The 4B is the one I would actually serve. 9B is a small bump in accuracy for a large bump in time. 27B is not in yet: six query heads per KV head needs a kernel variant.

Watch it fly

The quickest way to see what "state in, typed answers out" feels like at 11 decisions a second is to let it play a game. examples/starfighter is a browser lane shooter with no scripted pilot. Every tick the game writes what it sees as text —

Ship: lane 3 of 5, hull 2/3, shield ready, gun ready.
Lane 2: asteroid wall 34u (risky).
Lane 3 (yours): enemy bullet 18u, 5 credits 52u (bonus) (DANGER: hit within a second).
Lane 4: clear (safe).

— and asks cu-Jev four typed questions in one request: move (a choice between staying and stepping one lane left or right, each option described with what is in that lane), fire and shield (two nouls), and threat (a score that drives the meter). The answers pass through the game's own rules — one lane per tick, gun cooldown, shield charge — and the probabilities are drawn live next to the canvas. Asteroid walls with a single gap, credit clusters and fighters that drift between lanes while still far away make the lane choice matter every second.

That is Qwen3.5-4B on the RTX 3090 at about 95 ms per decision, roughly 530 tokens of state and questions per tick, prefilled fresh every tick because the state changes. Three consecutive autoplayed 40-second takes scored 1724, 1051 and 335 with zero, one and one hull losses; the recording is the first take, and examples/starfighter/record.py reproduces it. The pilot is not perfect. It dies about once a minute, mostly to a bullet fired from a lane it is stepping through, which is an honest picture of a 4B reading a text state at that rate. You can rewrite the pilot's orders while it plays, or tick "I fly" and take the keys yourself. (mp4)

bashCopy
uv run cujev serve --model models/Qwen3.5-4B
# open http://127.0.0.1:8080/play

The API is Jev's

pythonCopy
import os
os.environ["TYPESAFE_BASE_URL"] = "http://127.0.0.1:8080"
os.environ["TYPESAFE_API_KEY"] = "local"
from typesafe_sdk import Choice, Noul, Score, TypeSafeClient

r = TypeSafeClient().system_one(
    state="Hi, I've been trying to connect my Stripe account for 3 days and the "
          "integration keeps failing. I'm losing sales. Please help ASAP.",
    questions={
        "department": Choice(instructions="Which team should handle this",
                             criteria={"billing": "Payment or subscription issues",
                                       "technical": "Bugs or integration problems",
                                       "sales": "Pricing or account questions"}),
        "frustration": Score(instructions="How frustrated the customer appears",
                             criteria=["Calm, just stating facts", "Frustrated but civil",
                                       "Very angry, strong language"]),
        "is_urgent": Noul(instructions="The message conveys urgency or time-sensitivity"),
    })
r.answers["department"].choice     # "technical"  (Qwen3.5-4B: p=0.87)
r.answers["frustration"].score     # 1.11
r.answers["is_urgent"].noul        # 0.99

Validation errors are 422 with Jev's {detail: [{loc, msg, type}]} list. A bad key is 401. GET /v1/models and GET /health exist. The C API is three calls: cujev_model_open, cujev_state_prefill, cujev_eval.

bashCopy
git clone https://github.com/dtunai/cu-Jev && cd cu-Jev
CUJEV_CUDA_ARCH=120 uv sync          # RTX 5090
CUJEV_CUDA_ARCH=86 uv sync           # RTX 3090
uv run scripts/download_model.py Qwen3.5-4B
uv run cujev serve --model models/Qwen3.5-4B --port 8080

What it is not

It is not a chat model. It will not write the reply to the ticket. It will tell you which team owns it, how urgent it is, and whether the customer is angry, as distributions, in one request, with the state paid once.

It is not Jev. Jev is a trained, calibrated System One model. cu-Jev is an engine that turns an open hybrid LM into that API. What I care about is the shape. For a large class of production decisions, "state in, typed answers out" is a better description of the job than "call a chat model and parse the JSON." Once you stop decoding, the CUDA for it is a shared prefix, isolated branches, and a gather at the end.