Three Layers of a Burst-Admission Stall

Every steady-state profile of our production Qwen3.6-35B deployment told the same comfortable story: after ramp, the H100 is 97% busy and kernel share is throughput. But every profile also carried the same asterisk. In the first two or three seconds after a burst of agent requests arrives, the GPU timeline shows mega-gaps: windows of 0.3 to 1.8 seconds where the accelerator does nothing at all, while the scheduler process shows almost no traced work either. The gaps only happen during ramp, so they never show up in a steady-state number. They show up somewhere worse: in the time-to-first-token of every user who arrives at a cold or freshly scaled replica.

This is the story of peeling that stall. It took three rounds of capture-fix-recapture, and the three layers turned out to be three different kinds of systems bug: a JIT compiler running synchronously on the scheduler thread, a debug-grade consistency check spinning in the idle loop, and a tokenizer. None of them were kernels. All of them were invisible to a normal profile.

TL;DR

• An "untraced wait" in a profiler gap is a stack you have not captured yet. Arming the torch profiler with Python stacks at the moment of first admission attributed all three layers in three runs.
• Layer 1 (1.8 s): the TileLang JIT was parsing and compiling a linear-attention kernel variant on the scheduler thread, in the middle of serving, because the boot warmup never exercised the multi-sequence mixed-chunk geometry that burst admission produces. Fix: warm the variant space with tiny dummies at boot (kernel shape dims are dynamic, so a 256-token dummy warms the same kernel a 16k forward uses).
• Layer 2 (1.1 s): with the JIT gone, the gap was the scheduler running an O(tree) radix-cache sanity check plus a pool-leak walk on every idle-loop iteration: 3,200 iterations while the incoming prompts tokenized. Fix: rate-limit the guards to once per second.
• Layer 3 (the rest): tokenization itself. One tokenizer worker encodes a 40-request burst of 32k-token prompts serially: 102.5 ms per prompt, 4.1 seconds per burst, GPU idle throughout. We swapped the encode path to gigatoken: 1.5 ms per prompt, the whole burst in 60 ms, and byte-exact token ids on our vocab.
• Net effect at the serving gate: 32k cold prefill −25%, 64k −15%, concurrent throughput +9.3%, mean TTFT −9%, with temp-0 outputs exactly matching the baseline. Ramp-window GPU occupancy went from 55% to 84% before the tokenizer fix even landed.
• The gates earned their keep twice: the tokenizer swap silently broke structured outputs (a type-checking consumer rejected our wrapper) and crashed the scheduler once (an array type where a list was assumed). Both were caught by boot-time checks and one-boot probes, not by users.

The shape of the problem

Our serving workload is bursty by construction: coding agents fire batches of 32k to 64k-token prompts at once, and Modal scales replicas from zero, so "a burst hitting a fresh server" is not an edge case, it is the normal arrival pattern. A concurrent profile window over one such burst looked like this: the first forward step lands, then the GPU goes dark for 762 ms, briefly wakes, and goes dark again for 248 ms, all within the first two seconds. After that the timeline is dense and stays dense.

The frustrating part was that the scheduler showed almost nothing during the gaps: under 2 ms of traced CPU ops across a 762 ms hole. Whatever was stalling us was either outside the process or outside the profiler's default view. We had been filing this as the "untraced wait, admission-stall class" for two profiling rounds. The correct response to an untraced wait is not to speculate about it, it is to make it traced: re-arm the capture with Python stacks (with_stack=True), and arm it at the moment the first request enters the server rather than at steady state. Stacked multi-request traces are big, which is why profilers default to leaving stacks off, but a ramp window is short and one capture is all it takes.

Layer 1: a compiler on the scheduler thread

The first stacked capture put a name on the 1.81 s gap immediately. The scheduler thread was inside one extend step, annotated step[EXTEND bs=2 toks=16348], and under it: roughly 23,000 events inside tilelang/language/eager/ast.py. That is the TileLang eager front-end parsing a kernel's Python AST and JIT-compiling it. Our linear-attention prefill kernel (the FlashQLA Gated DeltaNet chunk kernel) is TileLang-generated, and it was being built, from source, on the scheduler thread, in the middle of serving the first burst.

Why did the boot warmup not catch this? Because of what the warmup looks like versus what a burst looks like. The warmup compiles every token-bucket shape with single-sequence batches. Burst admission produces something else: mixed chunks, where two requests share one 16,384-token prefill chunk (here: the tail of one request's chunk plus the head of the next). The TileLang kernel specializes on static keys that include the sequence-metadata dtype and layout flags, and the mixed-chunk geometry hit a variant the single-sequence warmup had never built. One variant, one compile, 1.8 seconds, every fresh replica.

The fix leans on a property of the kernel's specialization: batch size, token count, and chunk count are all dynamic dimensions, so a 256-token dummy builds the exact same binary a 16k-token forward uses. At the first real extend call (which happens during warmup, not serving), we now run a handful of tiny dummy forwards spanning the variant space: multi-sequence varlen with a partial chunk, both int32 and int64 sequence-offset dtypes. On a cold cache that costs about two minutes at boot; at serving time it costs nothing, which is the entire point. The follow-up capture confirmed it: zero TileLang activity during ramp.

Layer 2: the idle loop that was never idle

With the compiler gone, the next capture showed a 1.13 s gap with a completely different signature: the scheduler was executing its idle loop, 3,200 iterations of it, and each iteration ran a memory-pool leak check plus a radix-tree sanity check. Both guards walk data structures whose size grows with the cache (the radix tree's two LRU lists, every allocator pool), so each iteration cost ~0.35 ms, and the loop ran while the incoming burst was still being tokenized in another process. About 620 ms of the gap was the scheduler diligently re-verifying, hundreds of times per second, that a tree nobody had touched was still consistent.

Consistency guards in a scheduler idle loop are genuinely useful: they catch leaks and cache corruption close to the moment they happen. The bug is only their frequency: they are invariant checks, not per-iteration work, and their cost scales with cache occupancy, which on a long-context agent workload is always high. We rate-limited the pair to once per second (configurable, with the old every-iteration behavior available behind the knob). That converts the idle loop back into what it should be during an admission wave: a tight poll on the request queue.

A theme worth naming: neither layer 1 nor layer 2 would appear in any kernel table. One was a compiler, one was a Python loop, and both parked themselves exactly where a GPU-centric profile has a blind spot: the host side of a gap where no kernels run. The stacked capture is what turns "the GPU is idle and we do not know why" into a file and line number.

Layer 3: the tokenizer was the ramp

The third capture (both fixes applied) showed ramp GPU occupancy up from 55% to 84%, with one 476 ms gap left, and this time the scheduler's deepest frame was just the event loop waiting for requests. Requests exist; they are 32k tokens each; something upstream is slow. That something is tokenization: the server runs one tokenizer worker, and a HuggingFace fast tokenizer encodes a 32k-token prompt in about 102 ms. Forty prompts arriving at once therefore serialize into 4.1 seconds of tokenization, drip-feeding the scheduler while the GPU waits between the early, half-empty prefill batches.

Per-request encode latency also sits directly on the time-to-first-token path, so this is not just a ramp problem. We benchmarked the encode paths on our actual Qwen3.6 tokenizer with 32k-token prompts:

We integrated gigatoken as an opt-in tokenizer backend with a deliberately narrow blast radius: it owns exactly one surface, plain-text encode, and everything else (incremental decode, chat templates, the special-token API) stays on HuggingFace. That split is not caution for its own sake, it is what the measurements said. Encode was 99.98% of tokenizer cost. Incremental decode costs 0.59 µs per token in a separate process. Chat-template rendering is 3.2 ms of Jinja, which no tokenizer library helps with. And gigatoken's own decoder measurably diverges from HuggingFace on our vocabulary, while the incremental detokenizer is exactly where byte-level BPE quirks (partial UTF-8, prefix spaces, special-token skipping) are load-bearing for user-visible text. Encode is the only surface that is both expensive and provably exact, so it is the only surface we moved.

What the gates caught

A tokenizer swap sounds like the safest change in this article. It caused both of the round's regressions, and the reason is instructive: a drop-in wrapper is only a drop-in for consumers that duck-type. Two consumers did not.

• The scheduler assumed a list. gigatoken returns array types, and a request-state hot path concatenates origin_input_ids + output_ids. For Python lists that is concatenation; for an array it became a broadcast add and aborted the scheduler on the first request. One .tolist() in the wrapper, plus the lesson that the type contract of tokenizer output is part of its API.
• The grammar engine type-checks the tokenizer class. XGrammar inspects the tokenizer object to build its vocabulary tables, rejected our wrapper class, and the server "gracefully" fell back to no grammar backend at all: structured outputs and constrained tool calls silently disabled, in a deployment whose agents rely on tool_choice=required. The fix hands grammar backends the underlying HuggingFace tokenizer (they only consume the vocabulary surface, which the encode path never touches). The catch here was a boot-log line; the confirmation is now a permanent gate that boots the production config and asserts a JSON-schema-constrained generation actually conforms.

Both failures share a moral with the previous article's expert-0 bug: the dangerous failure mode of infrastructure changes is not the crash, it is the silent downgrade. "Falling back to grammar_backend='none'" is one log line in a wall of boot output. The only reliable defense is end-to-end gates that assert the capability, not the configuration.

The serving gate

The final same-host A/B compares the full shipping configuration against itself with only the tokenizer backend changed. Because the encode is token-for-token identical, greedy outputs must match exactly, and they do, across every probe. That equality is itself the strongest correctness gate we have: any wrapper bug that survived the unit tests would flip a token somewhere in eight 32k-to-64k-token generations.

One honest caveat: a residual ~0.7 s ramp gap remains in the very first admission wave, and its stack points at the request-delivery pipeline (HTTP ingest and inter-process transfer of tokenized requests), not at any compute we have named. It is the next layer of the same onion, and it is on the list.

What we took away

• An untraced wait is a capture request. If the GPU is idle and the profile shows nothing, the answer is a stacked capture armed at the right moment, not a theory.
• Warmup must rehearse the arrival pattern, not just the shapes. Our warmup covered every token bucket and still missed the variant that only multi-request admission produces. The production traffic shape is part of the kernel-variant space.
• Anything O(state) in a scheduler loop needs a rate limit. Guards that were cheap on an empty server become the bottleneck exactly when the server is full, which is the only time you care.
• Tokenization is part of the serving critical path. It does not appear in GPU profiles, it scales with prompt length, and at agent prompt sizes it was costing us a quarter of cold-prefill latency. Measure it like a kernel.
• Swap one surface at a time, and gate the capability. The encode-only tokenizer split made exactness provable; the grammar probe made the silent fallback loud. Both regressions cost one boot each instead of an incident.

FAQ

Why not just run more tokenizer workers instead of swapping libraries? Threaded HuggingFace encoding does help (8 ms per prompt at 16 threads), and it is a fine mitigation. But per-request latency still sits on the TTFT path at 100+ ms serial per prompt, cores spent tokenizing are cores not running the scheduler, and the 68x single-thread win makes the whole question moot: the burst costs 60 ms with change to spare.

Is a 68x tokenizer actually exact? On our vocabulary, yes, and we did not take it on faith: the gate encodes long random-word prompts plus special-token, unicode, emoji and mixed-language edge cases and asserts token-id equality against HuggingFace, then the serving A/B asserts byte-identical greedy outputs end-to-end. Encode exactness is checkable; we checked it. The same harness showed gigatoken's decode is not identical on our vocab, which is why decode stays on HuggingFace.

Why was the TileLang compile on the scheduler thread at all? JIT kernel frameworks compile on first use by design; the framework assumes first use happens during warmup. The bug is in that assumption, not in the compiler: serving traffic can always produce a shape class the warmup did not. If a kernel library exposes dynamic dims, prewarming the static-key variant space is cheap insurance.

Do these fixes matter at steady state? Mostly no, and that is the point: they matter at ramp, which is where scale-from-zero serving actually lives. The exceptions are the tokenizer (its per-request latency helps every request's TTFT) and a related find from the same captures: a per-chunk host stall in position-id construction (a 3×16,384 nested Python list built per prefill chunk) that cost 2-3% of steady-state GPU time and is now vectorized.

What stack does this apply to? Single-H100 serving of Qwen3.6-35B-A3B on our SGLang-based runtime with TileLang-generated linear-attention kernels. The specific bugs are ours; the classes (JIT-at-serving, O(state) idle guards, serialized tokenization, type-checking consumers behind duck-typed wrappers) apply to any LLM serving stack with a scheduler loop and a burst arrival pattern.

Free tools from Netra

All of these run in your browser. No signup, and your data is not uploaded to us.

• LLM VRAM Calculator: size the GPU for any model, precision and context length.
• Fine-Tuning Memory Calculator: the training-time memory budget.
• Token Counter: measure prompts under the GPT, Claude and Llama tokenizers.
• Free OCR: turn scans and PDFs into LLM-ready text.
• Web to Markdown: strip a page to clean Markdown before it reaches a prompt.
• Limbus: local-first image segmentation.
• sam3.c: SAM3 inference in pure C.

References

• Kernel Surgery on a Production Qwen Deployment: the previous entry in this series, covering the kernel-side half of the same profiling pass.
• gigatoken: the Rust SIMD BPE tokenizer behind layer 3.
• TileLang: the kernel DSL whose JIT model motivates the variant prewarm.
• Radix Islands: How Eager SWA Eviction Doubled Our Serving Concurrency: background on the radix cache whose sanity check starred in layer 2.
• Up to 4x faster than vLLM: the serving benchmark methodology, including same-host serial A/Bs.