IQuestRT: Faster IQuest-Q1 Inference for Coding Agents

IQuestRT is a dedicated inference runtime for IQuest-Q1 on NVIDIA GPUs, designed to shorten the generation cycles of coding agents.

It coordinates kernels, communication, sampling, and data movement across each speculative round. The GPU chooses how far to draft and launches the next round itself, while dependent kernels prepare inputs early to reduce waiting.

Performance and measurement scope

The sweep covers seven context lengths, from 1K to 192K. It measures the same five-draft computation at each length.

Across the sweep, increasing context from 1K to 192K raises median step latency by 16%, from 3.23 ms to 3.76 ms.

Real-request timing covers verification, sampling, recursive drafting, cache settlement, and the final readback of the device loop’s result log. Prefill, model loading, graph capture, and graph upload are excluded.

Why single-request decode speed matters

A coding agent reads code, generates a change, calls a tool, and reasons about the result. Tool outputs extend its context; generation determines when its next action can begin. Decode latency accumulates across these sequential steps.

IQuest-Q1 is a sparse mixture-of-experts model for CLI agents. Its inference path combines sliding-window and full-attention layers with recursive multi-token prediction (MTP). The same physical draft layer runs repeatedly to propose candidates, which the target model verifies together. Later draft depths generally consume just one live hidden-state row.

Small computation shapes and repeated synchronization make scheduling central to performance. We measure one decode step as a complete speculative round, tracking its latency, output tokens, and time per output token.

IQuestRT execution

IQuestRT captures a small family of graphs for four to eight live verification rows. The GPU chooses additional draft depths, settles the caches, and launches the next step without a host round trip. Within each graph, dependent kernels prepare independent inputs while their producers finish.

Tokens, acceptance counts, draft decisions, and cache lengths remain on the device throughout the decode loop. The host launches its first step and reads the result log once the request reaches a stop token or its output budget. The included server streams its output, so it reads the log about every 32 tokens to emit them and check for cancellation; the reported runs decode without callbacks.

Q8 describes the eight-row kernel width. A step verifies one committed token and three to seven drafts; any remaining rows are padding. The main MoE computes only live rows. Each recursive draft depth preserves its required KV writes, then continues with one seed row.

A profiled trace step shows where the GPU time goes. Each block links to the section that shapes it.

Use the same cross-GPU read for readiness and reduction

Small-batch TP8 exchanges relatively little activation data at each boundary, but the main model has 88 layers. Synchronization after attention and the MLP therefore repeats many times. IQuestRT fuses all-reduce with residual processing and normalization, and uses the data itself to indicate readiness.

Each rank's output buffer is initially filled with NaNs. Producers replace them with computed values. Consumers read the switch-reduced result through NVLink multicast addresses. If any rank has not written an element, the reduction still contains a NaN. Once an entire vector is ready, its consumer can proceed.

The same switch read checks readiness and obtains the reduced value. Buffer rotation and carefully ordered NaN rearming make reuse safe across layers and steps. Reads that remain unready have a failure exit.

The tails now separate data readers from normalization work. Main MLP readers issue their first switch read at entry, while other threads stage normalization weights. A second reader group issues a staggered read; the first NaN-free result claims the vector and continues. This reduces the cost of an unlucky read landing just before the last producer. Buffer re-arming and vector ownership preserve the reduction result.

Read only the KV tiles needed by the current positions

Under TP8, each GPU owns six query heads and one KV head. IQuestRT combines eight query rows and six heads into 48 matrix rows, reuses K/V within a block, and partitions work along the sequence.

Full-attention reads beyond 4,096 positions use a streaming kernel built on tcgen05 instructions. TMA transfers K/V, Tensor Cores perform QK and PV, separate warps handle softmax, and score and output accumulators reside in Tensor Memory. A four-stage buffer overlaps transfers with computation. K and V complete separate barriers, so QK begins as soon as K arrives; PV waits for V. Tensor Memory is freed after the output accumulator has been read, before the partial stores finish. These final scheduling changes preserve every MMA operand and operation order.

The current query positions determine the read extent. Full attention streams the valid KV history up to those positions and applies a causal mask at the boundary. Preallocated capacity beyond that range does not add reads. The learned attention sink has a zero value, so it contributes to the softmax denominator without requiring a concatenated K/V tensor.

A separate combine kernel merges the splits. For streaming attention it issues the partial loads together and uses four independent FP32 accumulation chains, shortening a serial dependency while retaining the weight and denominator calculations. Through PDL, combine also starts the output projection’s weight preparation early.

Sliding-window layers use a separate path. Their physical ring holds one additional 64-slot tile beyond the 4,096-position window. When verification writes several KV rows before reading them, those writes cannot evict history still needed by an earlier query. Every live verification row sees its complete window.

Fold routing into the attention tail

The main MoE stack no longer launches a separate router. O projection also computes partial router logits from the attention output and residual base. The attention tail reduces those FP32 partials, selects the eight experts, and publishes their weights beside the normalized expert input.

This moves routing into work the layer already performs and removes a repeated kernel boundary. The initial fold saved 88.5 μs per step on the trace replay. Later refinements staged shared attention rows once, ranked logits before the final normalization scale arrived, and assigned one softmax exponential to each lane.

The fold uses the unrounded FP32 expression rather than routing from the rounded BF16 activation, so it changes numerical results. The draft retains its separate BF16 router. Expert prefetch remains a cache hint and never overrides either routing decision.

Prepare expert weights before routing finishes

MoE experts are selected dynamically, so their weight addresses depend on the router. Fetching those weights is a substantial cost at small batch sizes.

IQuestRT consumes the checkpoint's E4M3 weights with 128×128 block scales. After matrix computation, Gate/Up applies SiLU and multiplication, then writes the E4M3 activations needed by Down. Down uses eight warps to process a token's eight routed experts in parallel. PDL lets it prepare weights while Gate is still running.

Prefetch runs only on sliding-window layers with four or five live verification rows. A side branch predicts experts from the incoming hidden state and warms their Gate/Up weights in L2. At six to eight rows the branch ends too late and costs more than it saves, so those graphs launch no prefetch. The predictor uses a checkpoint-specific refit of the folded router weights; the final expert choice still comes from the attention tail. The reported runs used this refit, the same prefetch file the README downloads.

Include sampling and acceptance in the execution path

IQuestRT uses the checkpoint's defaults: temperature=1, top-k=20, top-p=0.95. The draft also samples and saves its actual proposal distribution.

The 160,000-token vocabulary is distributed over eight ranks. Each rank extracts a local Top-20 from 20,000 logits. After exchanging candidates, the sampler selects the global Top-20 from 160 candidates, applies top-p, and draws a token. All ranks use the same candidates and random numbers determined by the request seed and token position, avoiding a separate broadcast of the chosen token.

Verification uses standard rejection sampling. A candidate d is accepted with probability min(1, p(d)/q(d)), where p is the target distribution and q is the draft's actual proposal distribution. At the first rejection, a replacement is sampled from normalized max(0, p−q). If every candidate is accepted, the target’s extra row supplies a bonus token: seven accepted drafts plus one bonus give a maximum acceptance length of eight. Only the drafts actually generated are eligible for acceptance. Acceptance, next-round input preparation, and KV rollback are included in decode timing.

This implements the sampling rule relative to the current verifier distribution. Numerical differences from FP8 activations or changed reduction orders still require separate qualification against a higher-precision reference.

Choose draft depth on the device

Three recursive depths always run. After each subsequent decision, a lightweight control kernel either launches the next depth or ends drafting, up to seven depths. It uses features of the drafts already generated, including their served probabilities and entropy, to estimate whether another depth is worth its cost. A draft is never discarded because of its own sampled probability.

The first three depths belong to the verification graph; later depths have separate device-launchable graphs. A skipped depth therefore executes no kernels. Against a fixed draft count, the stop rule reduced measured time per output token by 2.94% on ten held-out seeds in a paired replay.

Draft entry is compact too. At depth 0, one launch prepares geometry and the embedding–hidden operand before EH projection. At later depths, the EH kernel builds that operand itself, while an extra CTA prepares positions, RoPE, and the mask. The one-row computation avoids a separate entry launch without introducing a grid-wide barrier.

Keep the request loop on the GPU

The final control settles both caches, records the step’s result, and tail-launches the graph for the next verification width. A stop token or an exhausted token budget ends the chain. Launching the next graph before the boundary’s final writes is safe because the graph starts only after the current graph completes.

This removes the host read, bookkeeping, and next launch between steps. Paired replay attributed roughly 22–25 μs per step to eliminating that gap, after paying for the device launch; a leaner boundary recovered about 3 μs more. Graph upload also moved into capture, so it no longer appears in the decode clock.

vLLM: optimize the small-batch decode path

Alongside IQuestRT, we specialized vLLM's attention, MoE, and sampling kernels for the small shapes repeated during single-request speculative decoding, retaining the existing execution framework.

Expose more parallelism in attention

With one or six tokens and six query heads per GPU, query and head dimensions expose limited parallelism. A dedicated Triton split-KV path distributes historical KV across thread blocks, then merges partial outputs using their log-sum-exp values. The combine kernel also applies the learned-sink correction, avoiding a separate operation while preserving important BF16 rounding boundaries. Moving those boundaries can change the numerical result even when the mathematical expression stays the same.

Fuse MoE at the intermediate write

The small-batch BF16 path combines Gate/Up GEMM, SiLU, and multiplication in one kernel, writing Down's input directly to reduce intermediate traffic and launches.

The dynamic per-token FP8 path fuses SiLU, multiplication, row-wise absmax, and scale calculation, writing quantized activations directly. Each path produces the representation its downstream computation needs; the existing block-quantized FP8 fusion remains a separate specialization.

Parallelize the Top-K vocabulary scan

The first Top-K pass divides logits into tiles of 8,192, collects candidates in parallel, and performs an ordered merge. Tile maxima let later scans skip regions that cannot contribute candidates. Collection and merge add launches but reduce total sampling latency—a cost repeated at every recursive draft depth.

Scope and measured result

Hardware, tensor shape, precision, and sampling checks determine where each specialization applies. Benefits depend on which paths the workload uses.

On a different eight-GPU system, with MTP depth five and one concurrent request, replaying 400 Claude Code trace requests from ten controlled engineering sessions exceeded 322 tokens/s, an 8% improvement over that evaluation's starting point. The opening IQuestRT figure is measured on the current GPU configuration over 400 Claude Code traces.

IQuestRT applies the same principles across the complete decode round: expose enough parallel work, avoid unnecessary intermediate transfers, and account for costs repeated during drafting. Its dedicated execution path also brings communication, scheduling, and state handling into that optimization scope.

How IQuestRT became faster

Development started with fixed graphs, fused tails, and PDL scheduling. IQuest-Q1 integration brought native block-FP8 weights, recursive drafting, and default sampling into one decode path. Replay then guided a second phase: folded routing, narrower live-row work, adaptive draft depth, and a request loop controlled entirely by the GPU.

The sequence starts from a baseline manually implemented by engineers at 6.449 ms per step. Subsequent work kept more expert tiles in flight, replaced blocking tails, and staged inputs before their consumers waited. Later work folded routing into attention, let the GPU choose draft depth, shortened the communication tails, and fused draft entry.

How it was built

AI agents contributed implementations, experiments, debugging, and code review. Engineers defined workloads, numerical constraints, and acceptance criteria, and decided which directions to pursue and which changes to merge. The execution paths above emerged through repeated hypotheses, code changes, measurements, and checks.

Turn objectives into testable questions

We optimize real-request decode speed while tracking round latency and output tokens per round. If a change makes a round faster but also changes acceptance, both effects need their own explanation.

Tasks begin with concrete questions. How much waiting does a kernel boundary expose? Can the next layer's weights arrive earlier? If padding work is skipped, are all KV entries needed by the next step still valid? An agent reads the implementation and prior experiments, proposes a candidate and its validation, and works on an independent branch. Measurements run from committed code. Results identify the source revision, environment, command, and raw output.

Engineers use this evidence to assess the objective, the strength of the result, and whether further work is justified. For small gains, a third A/A arm rebuilds the candidate in the same way as the control: otherwise a build-dependent offset can masquerade as an optimization. Each experiment leaves a record that the next task can build on.

Let real requests correct the performance model

An early attribution experiment found that real decoding was substantially slower than a synthetic step with the same KV read length. The GPU-side difference was about 0.36 ms per step, mostly in MoE computation.

Routing explained the gap. Six live verification rows select 48 expert slots, but several rows can reuse an expert. Real requests visited an average of 37.7 distinct experts per layer, compared with 11.0 for synthetic inputs. Real routing read more distinct weights and reused the cache less.

That result changed how we evaluated MoE candidates. Synthetic runs remained useful for quick checks and context-specific measurements, but changes affecting expert access had to return to real requests. A gain on repeated-token inputs needed confirmation during actual generation.

To shorten the feedback cycle, we built prefill snapshot replay. Each fixed request executes real prefill once, saving KV and hidden states. Experiments restore that state, then run normal draft prefill, sampling, graph capture, and decode. The restored path must reproduce generation from real prefill before it can compare candidates. Control and candidate alternate in the same process, reducing repeated prefill cost and process-to-process timing drift. The tool also checks that their captured graphs actually differ, so an ineffective substitution cannot masquerade as an optimization.

Use paired experiments to define the prefetch boundary

Prefetch’s value changed as the surrounding chain became faster. We measured each verification width separately: four- and five-row steps benefited, while six- to eight-row steps lost time. Restricting the fork to the useful widths recovered 20.5–23.2 μs per step in two paired trace replays.

Prediction quality needed a separate experiment. An oracle issued exactly the experts the Gate would read, at the predictor’s existing cost, bounding the value of better guesses. Refitting the predictor’s weight matrix raised held-out expert coverage from 83.5% to 87.6% and recovered 4.0 μs per step without changing the kernel. A bias-based predictor also improved coverage, but its additional kernel work cost more than it saved.

These experiments distinguish cached bytes from scheduling effects. Prefetch’s resident CTAs and their shared-memory configuration also affect which kernels can run alongside them. More hits alone do not establish a faster decode step.

Keep the experiments that did not ship

Fuse attention combine into output projection. Output projection still needs the complete attention result; fusion does not remove that cross-CTA dependency. The tested correct synchronization schemes cost more than the eliminated kernel boundary. PDL already overlaps weight preparation, reducing the time left for fusion to recover. Combine remains separate.

Rewrite MoE Gate with native FP8 MMA. A tcgen05 prototype used TMA for weights and Tensor Memory for accumulation. Across 16 real routing patterns in isolated paired tests, the prototype combined with the existing Down was approximately 3.02 μs per layer slower. Removing its MMA work still did not give the remaining data handling an advantage. Gate failed the screen, so the Down prototype contingent on its success was not pursued.

Use FP4 expert weights. The screened encoding reduced weight storage and matched the shipped FP8 path closely on task scores, but the corresponding Gate and Down kernels did not recover that saving in execution. Unpacking, staging, and the compute schedule outweighed the smaller reads. FP4 remains an explored direction, not the runtime’s weight format.

Compute guessed experts before routing. Starting Gate weight loads from the prefetch predictor looked promising, but selecting the guesses inside Gate cost more than it hid. The paired replay was slower on every request; the shipped Gate continues to consume the final routing decision.

Push partials instead of polling the switch. Multicast stores and peer-to-peer reads did not produce a usable reduction sooner than the existing switch read, and their local sums did not reproduce its bits. The NaN-polled reduction stays.

Rejected candidates leave specific reasons behind. The host-loop direction was reopened once the runtime supported device graph launches and retained after a paired replay showed a gain. A previous limit is useful evidence, not a permanent rule.

Let correctness checks change the design

The sliding-window ring expansion came from checking a state boundary. Verification writes eight KV rows before reading. If the physical ring has exactly 4,096 slots, later writes can evict history still needed by earlier queries.

We reproduced the problem using independent position labels and visible-position sets, then expanded the ring to 4,160 slots, adding one 64-slot tile. Across 200 paired real decodes, the fix increased step latency by approximately 0.28%. We retained it because every verification row must see its complete window.

The fix also exposed a validation problem. If a control implementation contains a defect, using its error bound as a tolerance can hide real mistakes. Qualification must distinguish the positions that the defect can affect and require deliberately broken negative controls to fail.

Further investigation found numerical variation in the reference itself. Appending tokens can change the number of rows in each expert GEMM, changing rounding at earlier positions and amplifying differences through later routing. We added stability checks that append one and eight tokens to distinguish candidate errors from reference variation.

Keep the evidence with the decision

After implementing a change, the agent checks requirements, tests, and the final diff. Engineers decide whether it merges. Bit-exact scheduling changes must preserve the relevant numerical values, output tokens, and acceptance sequences. Changes to floating-point computation are checked against the reference, with acceptance changes reported separately. Both forms of validation need conditions that can fail.

Performance and numerical evidence applies to a particular source revision and environment. Later changes that affect the measured path require revalidation. Each adopted change records why it was retained; each closed direction records what would make it worth revisiting. Those records give the next agent a concrete starting point.