Skip to content

[None][perf] Experimental: compile and tune Triton, DeepGEMM and CuTe DSL JIT variants before serving - #19843

Draft
chenfeiz0326 wants to merge 99 commits into
NVIDIA:mainfrom
chenfeiz0326:chenfeiz/jit-prefetch-triton-mamba
Draft

chenfeiz0326 wants to merge 99 commits into
NVIDIA:mainfrom
chenfeiz0326:chenfeiz/jit-prefetch-triton-mamba

Conversation

@chenfeiz0326

@chenfeiz0326 chenfeiz0326 commented Oct 3, 2026 •

Copy link
Copy Markdown
Collaborator

@coderabbitai summary

Draft, experimental, not for merge. Off by default: TLLM_JIT_AHEAD=1 turns it on. No LLM-args change. Report with per-model data (NVIDIA internal).

Design

Startup and serving before and after this PR, with TRT-LLM's two executors: before, warmup pass 1 is mostly JIT compiles and tunes on the executor thread with the GPU idle, and a serving stall follows; after, CPU-only helpers compile ahead from model construction, tune ahead runs at the end of pass 1, and serving finds every kernel built and tuned. How compile ahead works, per rank: (1) one provider per kernel family finds its layers through model.modules(), (2) enumerates every shape class within the engine limits, (3) replays the real launcher on meta tensors for the exact compile key, (4) skips what is cached and queues the rest by priority, (5) CPU-only helpers build it, (6) the launch loads the build or waits for it, never compiling twice. Helpers per rank, sized to the node's CPUs: up to 16 each for Triton, DeepGEMM and fmha_sm100, 2 to 16 for CuTe DSL, one FlashInfer nvcc helper per module. Queue order: what a launch or the next warmup forward needs, then warmup shapes, graph sizes and tune-ahead candidates, then the background set at nice 19. Tune ahead benchmarks the untuned Triton autotune keys on the GPU at the end of each warmup.
  • Kernel selection is untouched. Helpers only compile; Triton's autotuner, DeepGEMM's heuristic and TRT-LLM's autotuner still choose. The serving loop is unchanged: no per-iteration step.
  • Isolated from the runtime code. The model loader, the engine's warmup and its serving forwards only call in at fixed points; jit_prefetch_engine.py (EngineJitAhead) runs each startup phase, and each kernel family's module decides, from the model and the engine limits, which kernels to build or tune (start_model_builds, start_weight_load_builds, start_warmup_builds, register_providers).
  • Warmup is not extended. It skips the passes that only compiled, whose kernels the helpers now build or that built nothing new: the DeepGEMM swap-AB sweep and the DSA metadata sweep (DeepGEMM helpers; on SM100 the metadata kernel has no batch-size template, so the sweep rebuilt one kernel per next_n), the multi-sequence one of [None][perf] Close Mamba hybrid warmup gap in autotuner warmup #16177's two Mamba passes (Triton helpers; with tune ahead, the other one, whose autotune moves to the end of warmup), the gen-FMHA NVRTC grid on disaggregated context workers (they run no generation request), and repeated MXFP8 graph tuning (each graph shape is tuned once and reuses its tactics). A pass whose helpers are not running still runs; TLLM_WARMUP_JIT_COVERAGE=1 runs them all.
  • Dead or idle helpers cost nothing. One that exits releases its builds, so a waiting launch builds in place, and a warmup pass skipped for its kernels runs again. With every helper of a kind killed as it starts, a server is left with what main's warmup leaves: on Nemotron Nano the same 11 kernels compiled while serving as with the switch off (none with live helpers), on DeepSeek-V3.2 none. Idle helpers stop once serving runs.
  • With the switch off, two changes remain, both against MiniMax-M3 recompiling in pass 2: its MSA KV-pool views are marked dynamic (no decoder recompile), and its MSA decode kernels do not specialize on the block table's stride (no decode-kernel recompiles).
  • Switches: TLLM_JIT_AHEAD=1; TLLM_JIT_TUNE_AHEAD=0|1 overrides tune ahead (=1 alone: in-process, no helpers); TLLM_JIT_STATS=1 logs every compile or autotune left on the executor, and tests/scripts/jit_prefetch_coverage.py turns that log into a coverage report. TLLM_JIT_PREFETCH is a deprecated alias. Knobs: TLLM_JIT_PREFETCH_{,DG_,CD_,FMHA_}WORKERS, TLLM_JIT_PREFETCH_WAIT_TIMEOUT_S, TLLM_JIT_TUNE_AHEAD_THREADS, TLLM_JIT_CUTE_DSL_DIR.

Performance Comparison

GB300, every server from empty JIT caches. Before = switch off, after = TLLM_JIT_AHEAD=1.

Server-ready time from empty JIT caches, before (switch off) and after (TLLM_JIT_AHEAD=1), for 12 model configs: 24 to 68 percent shorter. On this head: Nemotron Nano 12B v2, 239 to 110 s; Nemotron 3 Super 120B NVFP4 TP4, 440 to 213 s. No executor JIT after warmup in any after arm. Timing diagram of one cold start per arm on one node, Nemotron 3 Super 120B NVFP4 TP4, rank 0: with the switch off the server is ready after 440 s, its executor thread building 190 JIT variants (246 s, the GPU idle) during warmup; with TLLM_JIT_AHEAD=1 it is ready after 206 s: CPU-only helpers build the FlashInfer modules, the CuTe DSL NVFP4 GEMM tactics and the Mamba2 SSD candidates while the weights load and warmup runs, and the executor builds nothing. Warmup pass 1 by phase on the slowest rank, before and after, for 10 configs: the JIT builds each phase ran on the executor move to helpers, and pass 1 is 56 to 92 percent shorter. MiniMax-M3 ctx TP2 under nsys, GPU busy on rank 0: with the switch off, warmup is a queue of nvcc builds on the executor with the GPU idle, and after ready a variant warmup missed stalls serving for 144 s with the GPU idle; with TLLM_JIT_AHEAD=1 the helpers build them during startup and serving has no stall.
  • MiniMax-M3 ctx TP2 on fd2095d, both arms at once on one node: the off arm built 2 kernels while serving (two ~99 s stalls) and served 54.0K input tok/s over the 420 s window; with TLLM_JIT_AHEAD=1, none and 110.7K tok/s.
  • DeepSeek-V4.1-Flash TP4 on a1a4692 vs 8427b4b, same node, both on: the CuTe DSL MoE grouped GEMM tactics (12 per rank) and all but one of the 13 top-k and CSA2 variants now build in helpers, so pass 1 compiles 14 s less on each executor: pass 1 144.5 → 130.1 s, ready 380.8 → 370.6 s. On 474dc80 FlashInfer's CuTe DSL MXFP8 GEMMs (4 per rank, 2.4 s) also build in helpers; pass 1 moves within its run-to-run noise.
  • DeepSeek-V3.2 TP4 on 27abdf9 vs df55628 (TRTLLM MoE on FlashInfer, MTP3), same node: FlashInfer's TRTLLM-Gen fused MoE launcher (15 min of nvcc per rank) and RoPE now build in helpers from model construction; ready 1332 → 1177 s and, in reverse order, 1263 → 1217 s. Serving is unchanged (699.4 vs 696.6 tok/s, TPOT 3.48 vs 3.49 ms) and no longer compiles CuTe DSL's occupancy probe (3–4 per server).
  • Nothing new compiles while serving: after ready, no server measured here compiled or autotuned a kernel on its executor. That includes Kimi K3's 16 disaggregated context ranks, which skip the gen-FMHA grid: each built only its warmup forward's own FMHA kernel (6–7 s, before ready) and none in 46 min of serving at concurrency 360.
  • Kernels are unchanged: per-step GPU time agrees within ±0.4 % on matched steps. Right after a cold start the PR serves faster, as nothing compiles: Super's first 160 requests (be3a227, ISL/OSL 1024, concurrency 32), 6,952 / 6,948 → 7,137 / 7,214 tok/s, TPOT 10.19 / 10.19 → 9.95 / 9.85 ms; the off arms compiled 15 kernels across ranks during them.
  • Cost: CPU only. Helpers run at nice 19 (FlashInfer modules a layer needs at process priority). Weight loading runs slower while they build, 13–26 s on MiniMax-M3 and 22–36 s on Nemotron 3 Super; the ready times include it.
  • Still built on the executor in a cold start, and why not in a helper: trtllm-gen FMHA's NVRTC kernels (19–45 s of attention_jit), which its prebuilt library compiles and loads in one call into an in-process module cache, with no compile-only or load-from-file entry; Kimi K3's low-M split-K GEMM (~36 s), which FlashInfer builds from tensors it allocates on the GPU (a helper prototype matched bitwise but stalled in the 4-node server); MiniMax-M3's MSA sparse-attention kernels (~3.5 s per rank), compiled inside the vendored fmha_sm100 package's plan/run; each sparse layer's first top-k variant (0.6 s), compiled while its siblings are queued; and about 3 s per rank of small CuTe DSL glue kernels on DeepSeek-V4.1 (CSA2's quantize-scatter, scale-finalize and compaction kernels, the QK norm), each of which needs its own builder and plan, as CSA2's candidate publication has.

🤖 Generated with Claude Code

…orward pass

Off by default; TLLM_JIT_PREFETCH=1 enables it, TLLM_JIT_STATS=1 enables
only the measurement.

When a batch is scheduled, a per-model variant provider replays the Triton
launches the batch will issue on meta tensors. Triton's own binder computes
each launch's cache key and checks the kernel's in-memory cache, which is how
we know whether a variant is already compiled. Misses are serialized with
Triton's specialization format and compiled by CPU-only helper subprocesses
into the on-disk Triton cache, so the real launch only loads the cubin. Only
the Mamba2 SSD prefill kernels have a provider so far.

Every compile and autotune benchmark that still runs on the executor thread
after warmup is logged with its duration.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
chenfeiz0326 and others added 6 commits October 3, 2026 08:30
The planner built the token-major conv output as empty(conv_dim, T).t(),
so x/B/C had transposed strides and Triton specialized them differently
(stride == 1 became a constexpr on the wrong axis). The mixer allocates
empty(T, conv_dim).t() and slices the .t() of that, which is
contiguous [T, conv_dim]. Also widen the planner's dedup key to the chunk
counts that reach kernel signatures (nchunks, chunk-index length), and
add a GPU test that every real launch's Triton cache key is planned.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
…rs in warmup

A new process plans against an empty in-memory cache, so warm runs used to
send every variant to a helper even when its cubin was already in the on-disk
cache, and the executor waited ~2 s per helper start-up for the answer.

- Before submitting, check Triton's on-disk cache with the same key
  triton.compile uses; a hit is skipped (the launch loads it in ~1 ms).
- Helpers announce readiness after importing Triton and hashing the install;
  the model engine waits for that at the end of warmup.
- Compute triton_key() at prefetcher start-up so the first disk check does
  not pay the one-time install hash.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
The autotuner's own benchmark must decide which config runs; reusing a config
tuned for a different key can select a slower kernel.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
…elpers

The first prefill needs nearly every variant ~2.5 s after it is scheduled,
before two helpers can finish them. Two changes, neither adds warmup work or
touches autotuner config selection:

- On the first plan, a background thread enumerates one batch per kernel
  signature class reachable within max_num_tokens / max_batch_size and
  queues those variants at background priority. Variants the batch being
  scheduled needs always go to the front of the queue.
- Default helpers per rank: the CPUs available to the rank, capped at 16
  (was 2). TLLM_JIT_PREFETCH_WORKERS still overrides it.

Also stop calling Autotuner.prune_configs, which reads autotuner.nargs that
the real Autotuner.run owns, from another thread; plan every config instead
(none of the Mamba kernels prune). Add a test that the enumeration alone
covers real launches.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
Approach B, alongside the per-module providers (approach A). With
TLLM_JIT_RECORD_DIR set, every rank appends Triton's own specialization JSON
of each variant it compiles or loads (jit_post_compile_hook) to
jit_record.rank<r>.jsonl. On the next start, with TLLM_JIT_PREFETCH=1, the
record is handed to the helpers before the first request, so a process whose
compile cache is empty (new container, new node) starts compiling everything
the previous one needed before it needs it. The record header pins the
Triton version and GPU target; a mismatching record is discarded. Covers any
@triton.jit / @triton.autotune kernel, with no per-module code.

- TLLM_JIT_PREFETCH_PROVIDERS=none runs record/replay alone.
- Queue priorities: executor-blocked > scheduled batch > replay > provider
  enumeration. A variant the executor is about to wait on is moved to the
  front, so it never waits behind queued background work.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
@chenfeiz0326
chenfeiz0326 requested a review from hyukn October 8, 2026 02:46
chenfeiz0326 and others added 21 commits October 8, 2026 21:12
Extends the DeepGEMM patch (renamed deepgemm_trtllm.patch; it still carries
the configurable barrier timeout) so a kernel variant can be compiled into
DG_JIT_CACHE_DIR without being launched or loaded:

- Compiler::build_only(name, code): build()'s cache key, compile and atomic
  rename, minus the cubin load. build() now shares the compile step.
- sm100_fp8_fp4_gemm_1d1d_code(): config selection + code generation with no
  tensors, used by both the real op and the compile-only path, so both emit
  byte-identical code and hit the same cache entry.
- set_compile_target(arch_major, arch_minor, num_sms) and
  DG_JIT_COMPILE_ONLY=1: heuristics and the NVCC arch flag take the target
  from here, and no cuBLASLt handle or CUDA workspace is created, so a process
  with no visible GPU can compile for the target.
- compile_only_fp8_fp4_gemm_nt(m, n, k, dtypes, recipe, ...), bound with the
  GIL released.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
…elpers)

Adds a deep_gemm kind to the JIT prefetcher for the SM100 FP8 block-scale
Linear path (fp8_swap_ab_gemm -> deep_gemm.fp8_gemm_nt):

- FP8LinearDeepGemmProvider plans (M, N, K) per scheduled batch (A) and
  enumerates one M per 16-row window in the background.
- CPU-only helpers (jit_prefetch_dg_helper.py) call the patched DeepGEMM's
  compile_only_fp8_fp4_gemm_nt with this rank's arch and SM count, writing
  the cubin into DG_JIT_CACHE_DIR; nothing is launched on the GPU.
- Requests are recorded with kind=deep_gemm and replayed (B); a launch
  observer records shapes no provider planned and counts executor compiles.
- TLLM_DEEP_GEMM_JIT_WARMUP_STEP16=0 selects the step-128 grid (experiment
  switch for the NVIDIA#18813 warmup grid).

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
…ation order, stats in all arms

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
…EMM path too

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
…en prefetch is on

NVIDIA#18813 added two warmup changes whose only purpose is to make warmup
compile JIT variants: the 16-row DeepGEMM grid for the FP8 swap-AB /
block-scaling runners, and the aligned/ragged split of the Mamba-hybrid
warmup batches (KDA). With TLLM_JIT_PREFETCH=1 the helpers compile those
variants, so both now default to off; TLLM_WARMUP_JIT_COVERAGE=0/1
overrides either way. Without prefetch nothing changes.

Replaces the TLLM_DEEP_GEMM_JIT_WARMUP_STEP16 experiment switch.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
…lt kernels

DeepGEMM patch: compile_only_m_grouped_fp8_fp4_gemm_nt_masked, the
compile-only twin of the SM100 masked grouped GEMM that DeepGemmFusedMoE
launches. On SM100 that kernel has one layout candidate and M is not a
compiled dimension, so it depends on (num_groups, N, K) only.

- DeepGemmMoEProvider queues one request per MoE weight shape (gate+up,
  down) at batch priority as soon as prefetch starts.
- The masked grouped GEMM reports executor compiles to the same counter.
- The record now holds one line per kernel this process produced: a helper
  build, or a compile the executor had to do itself. Planned requests and
  launches are no longer recorded (most resolve to a kernel already on
  disk).

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
…lper, load in executor)

cute.compile never consults CuTe DSL's on-disk cache, so the helper exports
instead: it compiles the variant from fake tensors with the real call's
layout and an explicit --gpu-arch, with no visible GPU, and writes
<dir>/<key>.o with export_to_c. At the point the KDA op would call
cute.compile, a loader installed by the prefetcher loads that file with
cute.runtime.load_module (about 1 ms instead of about 6 s).

- jit_prefetch_cute_dsl.py: specs keyed on the K123 compile cache key,
  KdaPrefillProvider (heads, bias and safe gate from the model's KDA layers;
  the aligned or ragged variant from a batch's chunk lengths), build / load.
- jit_prefetch_cute_dsl_helper.py: CPU-only helper; preloads
  libcute_dsl_runtime.so so export resolves its CUDA entry points.
- The loader waits for an in-flight helper compile of the same variant,
  counts and records any variant the executor still compiles.
- Test: both variants, exported by a GPU-less helper and loaded by the real
  kda_prefill, give bitwise the same output and state as an in-process
  compile, and are loaded exactly once.
- DeepGEMM MoE test: pass packed UE8M0 scales, as DeepGemmFusedMoE does.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
…KDA (Kimi K3) layers

LinearAttnProvider replays, on meta tensors, the batch-dependent Triton
launches of GatedDeltaNet and KimiKDALinearAttention:

- the multi-row gated RMSNorm of both output gates (LAUNCH_WITH_PDL flips
  at ceil(tokens * heads / 4) == num_SMs; M's integer class),
- GDN's prefill / mixed-batch input prep (extract-transpose and the fused
  post-conv kernel: HAS_DECODE and the token counts' integer classes).

The executor hook now also passes generation tokens and plans eager
generation batches; batches that replay a captured CUDA graph are skipped.

Test: for prefill-only and mixed batches across the integer classes and
both sides of the PDL threshold, every real launch's Triton cache key is
among the planned ones (GDN with and without FP8 output, KDA).

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
…k V3.2/V4, GLM-5)

_fused_dsa_decode_metadata_kernel specializes only on
BLOCK_S = next_power_of_2(num_seqs) (num_seqs / num_tokens are
do_not_specialize). Captured decode sizes compile theirs at warmup; eager
decode batches compiled the rest on the executor.

fused_dsa_decode_metadata records its batch-independent launch parameters
(buffer strides, next_n, cache layout) on the first call, which warmup
always makes when the fused path is on. DsaDecodeMetadataProvider then
plans one variant per power of two up to max_batch_size, once, in the
background.

Test: after one call fills the template, every real launch at decode batch
sizes from 1 to 256 hits a planned cache key.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
…2/V4, GLM-5)

DeepGEMM patch: compile-only entries for the three SM100 kernels the DSA
indexer launches -- prefill MQA logits, paged decode MQA logits, and the
paged-logits scheduling metadata. Their generated code depends only on
config scalars and dtypes plus, for the decode pair, next_n (1 for decode
and MTP draft steps, 1 + max_draft for verify).

The indexer records each real call's compiled-in values; at prefetch start
DsaIndexerProvider queues every variant they imply, including the next_n
not run yet.

Test (GPU): the helper builds each variant into an empty cache, and the
real op then adds no cache entry (prefill; paged at next_n 1 and 4).

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
…with a LaunchArgs placeholder

LaunchArgs has no default constructor, so value-initializing the runtime
Args failed to compile.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
…utor thread

Wraps cutlass.cute.compile (and its compile[...] option form) so the JIT
stats cover CuTe DSL as well as Triton and DeepGEMM: one event per compile
plus cd_executor_compiles / cd_executor_compile_s in the summary.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
…kernel for GDN/KDA

DeepGEMM patch: compile_only_fp8_fp4_mega_moe. The MegaMoE kernel is built
with the layer's fixed shape plus the block config that
get_block_config_for_mega_moe picks from the per-rank token count (six
bands of expected tokens per expert). The ring sizing is factored into
get_mega_moe_ring_tokens so the compile-only entry and the real op agree.
Each real call records its layer config; MegaMoEProvider queues one token
count per band.

LinearAttnProvider also plans _cu_seqlens_triton_kernel, which
Mamba2Metadata.prepare launches for multi-sequence context batches with a
cached prefix (Qwen3.5-3.8, Kimi K3, GLM-5-next); the Kimi K3 diagnostic run
compiled it on every context rank.

Not yet run on a GPU: the MegaMoE test and the new chunk-index cases.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
LinearAttnProvider imported GatedDeltaNet, which gdn_mixer does not define
(the class is Qwen3NextGatedDeltaNet), so building the provider raised
ImportError in warmup whenever prefetch was on. The provider tests built it
with __new__ and never ran the constructor. Add a test that builds the
provider from real Qwen3NextGatedDeltaNet and KimiKDALinearAttention layers.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
… once per shape

Helpers (and the compilers they fork) run at nice 19, so compiling only takes
CPU time the executor and its threads leave idle.

The DeepGEMM compile counter flagged every launch slower than 0.2 s on the
executor thread. On a Kimi K3 context server near its memory limit the same
(m, n, k) crossed it 8-13 times per rank: allocator stalls, not compiles,
since DeepGEMM keeps a kernel in memory once loaded. Count only the first
launch of a shape.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
…read GEMMs and MLA prefix chunks

Three kernels the Kimi K3 context server still compiled on the executor
thread with prefetch on:

* _extract_transpose_prefill_kernel: KDA prefill repacks the [q | k | v]
  columns of the fused qkvg projection with the same kernel GDN uses. The
  KDA planner now replays it (prefill tokens only), and the KDA test checks
  both kernels a prefill launches.
* DeepGEMM FP8 GEMMs of _Fp8BlockScaleWeightReadLinear (Kimi K3 attention
  and shared-expert projections): they call trtllm::fp8_swap_ab_gemm but are
  not Linear modules, so FP8LinearDeepGemmProvider never saw their shapes.
* MLA kv_b_proj over a cached prefix: each loop runs it over up to
  chunk_size x chunked_prefill_buffer_batch_size cached tokens (32768 here),
  past max_num_tokens, so the M enumeration now reaches that bound when
  chunked prefill and block reuse are on.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
From the Kimi K3 disaggregated run (context server, 16 ranks, 4 per node),
where the executor still waited 98 s on helpers after warmup:

* CuTe DSL helpers set CUTE_DSL_ARCH to the rank's arch. Without it CuTe DSL
  queries the driver for the arch, which with no visible GPU costs about
  27 s per K123 compile and serializes across the node's helpers: one
  compile alone takes 32 s, eight at once 222 s each. With it, 4.7 s.
* The ranks of a node share the CuTe DSL output directory; a per-variant
  file lock makes one helper compile while the others wait and then load it.
* A Triton variant no helper has started yet is compiled on the executor
  (about 0.1 s here) instead of waited for; the queued request becomes a
  disk-cache hit. A helper's first compile also cost 1.4 s more than later
  ones (backend init); helpers now compile a trivial kernel before they
  report ready, inside warmup.
* Background enumeration starts when the providers are registered, at the
  end of warmup, rather than at the first batch.
* The DeepGEMM swap-AB timer covers only the DeepGEMM call, not the FP8
  quantize and output allocation before it; a caching-allocator stall there
  read as a compile.

Writes to the record file skip a closed handle, and the prefetcher tests
restore Triton's process-wide hooks, so one test's prefetcher does not see
another test's compiles.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
…s during warmup

A Triton helper's first request cost 1.44 s against 0.1 s for later ones.
The trivial warm-up compile added in the previous commit did not help: the
cost is importing the kernel's module (1.3 s for the linear-attention
modules, through their own imports such as torch), not the compiler.

Providers now list the modules defining their kernels (kernel_modules), and
register() sends each Triton helper a {"preload": [...]} request, which the
helper imports without replying. The first request after a preload takes
0.2 s. The warm-up compile is removed. Helper stdin writes go through one
lock per helper, since register() and the feed thread both write.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
… patch to b64107f

main moved DeepGEMM to b64107f, where the JIT runtime lives in the DeepJIT
submodule (third-party/deep_jit), and moved warmup out of model_engine.py
into DecoderRunner. Conflicts: model_engine.py keeps main's warmup
delegation and calls _init_jit_prefetch() after it; fetch_content.json keeps
main's pin and adds the patch; the fused MoE DeepGEMM call keeps main's
argument list.

deepgemm_trtllm.patch is rewritten for b64107f (it applies with
`patch -p1` from the DeepGEMM source root, including the submodule):

* DeepJIT: DJ_COMPILE_TARGET=major,minor,num_sms,smem_optin,l2 makes
  cuda::Device use those properties instead of querying a device, so a
  process with no visible GPU compiles for the rank's target. This replaces
  set_compile_target / DG_JIT_COMPILE_ONLY of the old patch.
* compile_only_source(tag, source): the cache key jit->compile would use;
  compiles and commits the entry on a miss, no load or launch.
* Compile-only twins, each sharing the real launcher's source function so
  both produce the same cache entry: fp8_fp4_gemm_nt and
  m_grouped_fp8_fp4_gemm_nt_masked (SM100 1D1D), the DSA indexer's
  mqa_logits (now with clean_logits / has_schedule_meta), paged_mqa_logits
  and paged_mqa_logits_metadata, and fp8_fp4_mega_moe (ring sizes computed
  by the same function as the symmetric-buffer size).

The DeepGEMM helper passes the target through DJ_COMPILE_TARGET and uses
PyTorch's cuBLASLt handle (DeepGEMM's own needs a device). MegaMoE token
bands come from DeepGEMM's get_block_m_for_mega_moe instead of a copy of
its thresholds, which changed at this pin, and the layer records its MMA
type so the compiled weight dtype matches. The DSA indexer records the
kernel's head count after FP4 padding and clean_logits.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
… prefetch compiles it

The FP8 swap-AB GEMMs (trtllm::fp8_swap_ab_gemm, the default FP8 block-scale
Linear on SM100/103, and trtllm::fp8_prequantized_swap_ab_gemm) have a single
tactic. They go through the autotuner only so that its pass over every
16-row M bucket compiles each DeepGEMM layout during warmup, and they set
exclude_from_cache so that a warm tuning cache cannot skip that pass
(NVIDIA#15626, then the step-16 grid of NVIDIA#18813 and NVIDIA#19183).

With JIT prefetch on, FP8LinearDeepGemmProvider builds every one of those
layouts in CPU-only helpers, without launching anything on the GPU, so the
two ops now call their single tactic directly; DeepGEMM finds each layout
on disk. The tuning pass, its exclude_from_cache and the step-16 grid stay
for every other case: prefetch off, TLLM_JIT_PREFETCH_PROVIDERS=none (no
providers, so nothing builds the layouts ahead), a DeepGEMM without the
compile-only entry point, and other architectures (SM107). The Hopper FP8
block-scaling runner keeps its grid unconditionally; prefetch does not
cover that path.

This also drops the earlier gate inside deep_gemm_jit_warmup_buckets: with
the swap-AB ops off the autotuner, nothing that prefetch covers reads the
grid, so it returns to main's step-16 grid for the runners that still tune.

TLLM_WARMUP_JIT_COVERAGE=1 restores the tuning pass with prefetch on.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
…mup buckets

_warmup_dg_paged_mqa_logits_metadata (NVIDIA#16178) runs DeepGEMM's
get_paged_mqa_logits_metadata once per 32-aligned batch bucket up to
max_batch_size x (1 + max_draft_tokens), so each bucket's kernel is compiled
during warmup instead of on the first batch that reaches it.

At the current DeepGEMM pin the SM100 kernel is no longer templated on the
batch size: its compiled-in values are next_n, the SM count and split_kv, so
the bucket loop builds one variant and launches the kernel once per bucket.
With JIT prefetch on, DsaIndexerProvider compiles that kernel for every
next_n the indexer uses (1 and 1 + max_draft_tokens) in CPU-only helpers,
with nothing launched, so the warmup loop is skipped. It still runs with
prefetch off, without the patched DeepGEMM, and on SM90 / SM120, whose
metadata kernels prefetch does not build.

TLLM_WARMUP_JIT_COVERAGE=1 restores the loop with prefetch on.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
…he weights load

The sampling module's helper started at the beginning of warmup, and its
nvcc build takes a fixed 63-76 s per rank. Once warmup's first forwards got
faster (DeepGEMM and the low-M GEMM variants built while the weights load),
warmup reached the join before CUDA graph capture first: on the Kimi K3
high-throughput generation worker (helix16, bs 128) sampling_module_join
waited 25-28 s on every rank.

- start_model_builds now appends "sampling" to the model's FlashInfer
  modules, after the layers' own (the first forward needs those first), if
  the engine builds it in warmup: warmup_builds_sampling_module, true for a
  decoder runner (resolve_runner_type). The helper builds it during weight
  load, ~300 s on K3.
- The decoder runner's warmup takes a pending sampling build (new
  flashinfer_build_pending) as its own: it neither starts a second helper
  nor builds the module in place, and joins it before CUDA graph capture
  as before. Warmup's finally now stops any unjoined helper even without a
  KV cache to warm up, so a loader-started build never outlives warmup.

Kimi K3 high-throughput generation worker (GB300, helix16 bs 128, prefetch on
in both arms, same 4 nodes, empty JIT caches): sampling_module_join 26.6 s ->
0.0 s, final executor creation 181 s -> 152 s, server ready 536 s -> 501 s;
decode probe 6381 vs 6372 tok/s, 0 compiles on the executor after warmup.

Tests: test_sampling_module_starts_after_the_layer_modules_at_load (cpu_only,
decoder vs encoder engine), test_warmup_joins_the_sampling_build_the_model_
loader_started (cpu_only, pending or not), test_load_time_sampling_build_is_
joined_not_restarted (GPU: the helper's build is loaded without a rebuild).
test_jit_prefetch_deep_gemm_load_time.py is added to l0_cpu.

Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
…nd per-iteration planning

The PR does two different kinds of work. Compile ahead and tune ahead cover
the reachable variant space from the model and the engine limits, and finish
by the end of warmup: the CPU-only helpers, the builds started at model load
and before warmup's first forward, the background enumeration, the Mamba SSD
tune ahead, and skipping warmup's compile-only passes. Per-iteration planning
runs on the executor thread for every scheduled batch and queues that batch's
variants (Triton providers, DeepGEMM M, CuTe DSL K123) ahead of the
enumeration. One switch turned both on.

- TLLM_JIT_AHEAD=1 turns on compile ahead and tune ahead, and the
  TLLM_WARMUP_JIT_COVERAGE / TLLM_JIT_TUNE_AHEAD defaults follow it.
- TLLM_JIT_PLAN_PER_ITER turns on per-iteration planning. It needs
  TLLM_JIT_AHEAD (whose helpers do the compiling; set alone it is ignored
  with a warning) and follows it while unset, so TLLM_JIT_AHEAD=1 alone
  behaves as TLLM_JIT_PREFETCH=1 did.
- With TLLM_JIT_PLAN_PER_ITER=0 the executor thread is still bound, so the
  JIT stats and the waits on helper-started compiles keep working; under
  TLLM_JIT_STATS the report then lists what compile ahead alone misses.
  The [JIT stats] SUMMARY line reports plan_per_iter.
- TLLM_JIT_TUNE_AHEAD still overrides tune ahead either way; =1 alone is the
  in-process mode without helpers.
- TLLM_JIT_PREFETCH is read as TLLM_JIT_AHEAD while that is unset, with a
  deprecation warning.

Tests (GB300, this change as a Python overlay on a wheel of the branch merged
with main): test_jit_ahead_and_plan_per_iter_switches (cpu_only: the switch
matrix, the legacy alias, TLLM_JIT_TUNE_AHEAD overrides, coverage defaults),
test_plan_per_iter_off_plans_no_batch_but_binds_executor (cpu_only:
JitPrefetcher plans no batch, DeepGEMM or CuTe DSL request without it, binds
the executor either way) and
test_plan_jit_prefetch_binds_the_executor_and_plans_only_with_plan_per_iter
(cpu_only: the engine hook). The existing tests now set TLLM_JIT_AHEAD and
clear the legacy name. 330 pass across test_jit_prefetch.py (mamba),
test_pytorch_model_engine_warmup.py, test_distributed_warmup_oom.py,
test_jit_prefetch_deep_gemm_load_time.py, test_jit_prefetch_cute_dsl.py,
test_jit_prefetch_deep_gemm.py, test_jit_prefetch_flashinfer.py,
test_minimax_m3_fmha_jit_prefetch_plan.py and
test_kpool_jit_prefetch_modules.py; -m cpu_only collects 159 of them, as
CI's CPU stage does through l0_cpu.yml.

Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
Main split DecoderRunner by phase (NVIDIA#20134): warmup moved to
decoder/warmup.py (DecoderWarmup) and CUDA graph capture to
decoder/graphs.py (DecoderCudaGraphs). The PR's warmup changes move with
them:
- DecoderWarmup._warmup_scheduled starts the sampling-module and MXFP8
  FlashInfer helpers, calls before_warmup_forwards_callable after the
  pre-JIT prewarms, and joins the helpers before CUDA graph capture
  (_warmup_scheduled_forwards, _join_sampling_module_build).
  DecoderRunner.before_warmup_forwards_callable forwards to it.
- _tunes_flashinfer_mxfp8 / flashinfer_mxfp8_pass_selected move to
  DecoderWarmup; DecoderRunner.flashinfer_mxfp8_pass_selected keeps its
  classmethod signature for jit_prefetch.tunes_flashinfer_mxfp8_at_load,
  passing the runner class's _context_warmups_supported. The MXFP8 helper
  starts only with compile ahead on.
- The DSA paged_mqa_logits_metadata skip, the context-only gen FMHA grid
  skip (through main's trtllm_gen_jit_warmup argument) and the Mamba
  tune-ahead pass selection move to DecoderWarmup.
- The once-per-shape FlashInfer MXFP8 decode-graph tuning moves to
  DecoderCudaGraphs.
- Main restores the DSv4 indexer-Q CuTe DSL prewarm with a working import
  path (NVIDIA#20253); it is kept.

Other conflicts:
- glm_kpool/backend.py: main adds the native FP8 TRTLLM-GEN decode path;
  the BF16 branch keeps native_decode_geometry, which jit_prefetch's
  fmha_gen planner shares.
- test_distributed_warmup_oom.py: the Mamba tune-ahead test runs on main's
  DecoderWarmup fixture.

Main moved the CuTe DSL kernels out of _torch/cute_dsl_kernels (NVIDIA#19275);
jit_prefetch_cute_dsl.py imports them from _torch/kernels and
_torch/attention/kernels. The PR's runner tests now target DecoderWarmup
and DecoderCudaGraphs.

Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
… set

_executor_loop_cleanup stops the JIT prefetcher of the executor's model
engine. The cleanup unit tests run it on stubs without a model_engine
(test_py_executor.py TestExecutorLoopCleanup, test_gpu_keepalive.py
test_fill_gate_wiring), which raised AttributeError; read the attribute
defensively, as the prefetcher lookup already was.

Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
Compile ahead and tune ahead leave nothing for per-iteration planning to
do. Every model the PR measures (DeepSeek-V4.1-Flash ctx-only and
aggregated, GLM-5.3-Flash ctx, MiniMax-M3 ctx and gen, Kimi K3 ctx and
gen-only, Nemotron Nano and Super) ran with and without it on one PR
head: it never queued a variant, executor JIT after warmup was the same
in both arms, and it cost 9-60 us of host time per iteration.

Remove it: TLLM_JIT_PLAN_PER_ITER, the executor-loop hook, the engine's
batch planner, JitPrefetcher.plan / plan_deep_gemm / plan_cute_dsl, and
the providers' batch-only planning paths. The executor thread the JIT
stats count is now bound by the engine's first forward after warmup.

With no batch queuing its exact M, a DeepGEMM launch looks up the
request the enumeration queued for its 16-row M window (same layout):
it waits for that build instead of compiling the kernel itself, and is
no longer reported as an unplanned launch.

Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
@chenfeiz0326 chenfeiz0326 changed the title [None][perf] Experimental: compile Triton, DeepGEMM and CuTe DSL JIT variants ahead of the forward pass [None][perf] Experimental: compile and tune Triton, DeepGEMM and CuTe DSL JIT variants before serving Oct 23, 2026
…ant ahead

The MSA indexer plans rows with a long cached prefix as segments split
along KV (msa_indexer.segmented_proxy_plan), and warmup loads that plan's
two fmha_sm100 variants before its first forward
(prewarm_split_proxy_variants). The provider listed only unsplit variants,
so with prefetch on, the executor still compiled the KV-split one itself
in warmup (1_0_1_2_1_1_0, 128 s under nsys on a cold M3 ctx TP2 server).

List the KV-split variant for every index site, and put both of the
plan's variants first in the build order, as warmup needs them first.
planned_variant_keys now holds the provider's list; a GPU test checks it
against the variants real segmented proxy plans select.

Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
Three conflicts, each two additions to one statement:
- decoder/graphs.py: main runs a warmup-only generation graph forward
  through forward_discarded_warmup; it now runs inside the PR's
  once-per-shape FlashInfer MXFP8 tuning context.
- decoder/warmup.py: main wraps the attention warmup forward in
  discard_nccl_window_tensor_outputs; it keeps the PR's
  trtllm_gen_jit_warmup=gen_fmha_grid (context-only workers skip the
  generation FMHA grid).
- model_loader.py: main's validate_lm_head_output_dtype runs first, then
  the PR's start_model_builds.

Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
…warmup's Mamba2 prefills planned, helper exits and idle helpers handled

Warmup time:
- FlashInfer's TRTLLM-GEN launcher fmha_gen (~50 s of nvcc) is built from
  model construction for every model whose regular attention layers can run
  on FlashInfer's TRTLLM-GEN FMHA library (their FmhaManager lists it), not
  only for GLM-5.3-Flash and MiniMax-M3. Each FlashInfer module gets a
  helper of its own. The layers' modules build at this process's priority
  and are joined before warmup's first forward (flashinfer_layer_join):
  a forward reaching fmha_gen while the helper still fetched its kernel
  header fetched it again, and FlashInfer's ninja then rebuilt the module
  in place after waiting for the helper.
- Before each warmup forward with context requests runs, the providers that
  plan warmup batches (the Mamba2 SSD one) queue that batch's variants not
  compiled yet ahead of all other work, every candidate config of an untuned
  autotune key included, so the autotuner warmup's prefill no longer
  compiles ~70 SSD configs one by one on the executor. The SSD provider
  registers before the first forward and queues tune ahead's batch with
  cached prefixes there.
- FlashInfer's CuTe DSL RMSNorm variants, and with them the CuTe DSL
  helpers, start at model construction, not at the end of warmup, where
  wait_ready waited ~10 s for the helpers' start-up.

Robustness and resources:
- A helper that exits releases the build it was working on and, as the last
  of its kind, its queue: a launch waiting for it builds the variant in place
  at once instead of after the wait timeout. wait_ready has one deadline for
  all helpers and stops waiting for helpers that exited. A wait that reaches
  TLLM_JIT_PREFETCH_WAIT_TIMEOUT_S (parsed once; a bad value is ignored)
  gives up on that build. The pool's threads end with their helper.
- Once serving has started, each kind's helpers stop when no enumeration is
  still queueing and their queue has drained; idle, they only held host
  memory (18 helpers, 14.6 GB of RSS per Nemotron Nano rank).
- Helpers get the JIT cache locations as absolute paths (they run in /tmp).
- Warmup's swap-AB DeepGEMM pass is skipped only while DeepGEMM helpers run.
- MLA's kv_b_proj rows over a cached prefix are planned whenever chunked
  prefill is on, with or without block reuse.
- DeepGEMM compiles on paths no launch hook observes (DSA indexer, MegaMoE)
  are reported while serving, at most once a second, not only at executor
  shutdown; the JIT stats events list is bounded.

Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
…s while the weights load

Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
…ecialize on the block-table stride

The MSA decode producers specialized on stride_bt_b, the block table's row
stride. It follows max_seq_len, which differs between the KV-cache
estimation executor and the final one (MiniMax-M3 ctx TP2: 6192 and 8193
blocks per sequence), so its divisibility by 16 differed and warmup pass 2
compiled every decode variant again, into identical binaries: 10 kernels
per rank, 6.4 s. The producers now do not specialize on it, and the MSA
decode planner's replays, which run the real launcher, follow.

Measured on MiniMax-M3 ctx TP2 (GB300, empty JIT caches, same node,
TLLM_JIT_AHEAD=1): decode compiles on the executor in pass 2, 10 -> 0 per
rank; pass 2, 98.8 -> 89.0 s; serving unchanged.

Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
…tive_state_update module while the weights load

A Mamba2 layer whose shape FlashInfer supports runs its single-token decode
step through FlashInfer's selective_state_update. FlashInfer built that op's
JIT module at its first call, in warmup pass 1's first generation forward,
with the GPU idle: 24.8 s on Nemotron 3 Super 120B NVFP4. The module follows
from the layer alone (the SSM cache dtype, the activation dtype, float32 dt
and A, int32 state indices, head dim, state size, and the Philox rounds of
stochastic rounding), so jit_prefetch_flashinfer.flashinfer_modules names it
at model construction, a FlashInfer helper builds it from that name (its
JitSpec name spells out the generator's arguments), and warmup joins it with
the other layer modules before its first forward. A model with speculative
decoding verifies drafts through another module or kernel and is left out.

Measured on Nemotron 3 Super 120B NVFP4 TP4 (GB300, empty JIT caches,
TLLM_JIT_AHEAD=1, ABBA on one node): the helper builds the module in
31.7-32.4 s while the weights load; the executor's JIT builds in warmup,
1 (24.7-25.7 s) -> none; warmup pass 1, 75.5 / 62.8 -> 55.6 / 43.0 s;
server ready, 245.2 / 228.6 -> 206.0 / 203.8 s.

Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
…ss as warmup records it

The MSA decode provider planned from the launches warmup recorded, but was
registered only after the first warmup, so every decode variant pass 1
reached, one per batch size in its forwards and CUDA graph captures,
compiled on the executor first: 14 kernels per rank, 6.7 s. The provider is
now registered before warmup's first forward, and each new launch class
warmup records is planned at once (JitPrefetcher.plan_recorded, on a thread
of its own, ahead of other work), so warmup's later launches of that class
find their variants built or being built. The prefetcher keeps its
providers by name, so both warmups use the first one and plan only the
classes it has not planned.

Measured on MiniMax-M3 ctx TP2 (GB300, empty JIT caches, TLLM_JIT_AHEAD=1,
both servers at once on one node, so they share its CPUs): MSA decode
compiles on the executor in pass 1, 14 -> 1 per rank (6.7 -> 0.7 s; the
helpers built the other 22 variants); warmup pass 1, 294.6 -> 283.5 s;
server ready, 555.9 -> 535.9 s.

Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
…on variants before warmup reaches them

The radix top-k and CSA2 candidate-publication providers started only at
the end of the first warmup, so every top-k variant pass 1 reached compiled
on the executor first. On DeepSeek-V4.1-Flash: 7 FilteredTopKKernelVarlen-
Decode, 4 FilteredTopKKernelVarlenPrefill and 2 CandidatePublishKernel
compiles per rank (6.5 s), and one more decode variant in pass 2.

- CSA2's provider follows from the model and now starts before warmup's
  first forward.
- A sparse backend builds its indexer, and the indexer's TopK, in its first
  forward, and sets the TopK's radix implementations after constructing it,
  so before warmup there is no TopK to plan from. TopK.on_first_forward
  reports each TopK at its first forward; its variants are enumerated on a
  thread of their own and queued ahead of other work.
- The enumeration spans every score-width bucket from one column to the one
  above max_seq_len's (a CUDA graph's decode scores are a little wider than
  max_seq_len), CSA2's capacities one power of two further for the same
  reason, and lists the variants in the order warmup reaches them (prefill,
  then decode, narrow to wide) instead of their specs' text order, which put
  the 4096-column prefill variant behind dozens of decode ones.
- enable_cute_dsl_topk is additive (each call queues only new variants);
  enable_csa2_publish acts once.

Measured on DeepSeek-V4.1-Flash TP4 / EP4 with attention DP and DSpark
(GB300, empty JIT caches, TLLM_JIT_AHEAD=1, same node, together with the
next commit's MoE grouped GEMM build-ahead): top-k compiles on the executor
in pass 1, 11 (6.1 s) -> 1 (0.6 s: the first forward's own variant,
compiled while its siblings are queued), and none in pass 2 (was 1); CSA2
candidate publication, 2 -> 0.

Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
…M tactics while the weights load

CuteDslFusedMoE's NVFP4 forward has the autotuner race the routing tile
sizes (128 and 256 rows on SM100/SM103), and each tile's forward tunes FC1
(gather + grouped GEMM + activation) and FC2 (grouped GEMM + finalize) over
every tactic their runners list, compiling each at its first benchmark on
the executor thread: on DeepSeek-V4.1-Flash, 4 FC1 and 8 FC2 kernels per
rank (8.3 s) in warmup pass 1.

Both kernels take every extent as a runtime scalar and every operand as a
pointer, so a kernel follows from the tile, the tactic and the layer's
constants: the routed top-k, the local expert count, the activation with
its SwiGLU clamp or SiTU caps, and whether finalize is fused.
MoeGemmProvider reads those off each CuteDslFusedMoE right after
construction (its weights exist, on meta), asks the runners' own
get_valid_tactics on meta tensors of the weight shapes, and the CuTe DSL
helpers compile and export every kernel while the weights load; both
runners consult moe_gemm_loader before cute.compile. Layers on the
locality-domain kernels, and DeepEP's count-native metadata (top-k 1 over
per-expert counts, chosen at the first forward), are not planned: those
kernels compile in place as before.

Measured on DeepSeek-V4.1-Flash TP4 / EP4 with attention DP and DSpark
(GB300, empty JIT caches, TLLM_JIT_AHEAD=1, same node, with the previous
commit): MoE grouped GEMM compiles on the executor in pass 1, 12 (8.3 s) ->
0; with the top-k change, warmup pass 1 144.5 -> 130.1 s, server ready
380.8 -> 370.6 s.

Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
…that warmup's MXFP8 pass races

The MXFP8 Linears' warmup pass (DeepSeek-V4.1's Engram WKV and
MiniMax-M3's MXFP8 Linears, with decode CUDA graphs) races FlashInfer's
CUTLASS and CuTe DSL MXFP8 GEMMs per token bucket. FlashInfer compiles each
CuTe DSL kernel at its first benchmark, on the executor thread, and keeps
it in memory only: 4 kernels per rank in pass 1, 2.2 s on
DeepSeek-V4.1-Flash and 2.5 s on MiniMax-M3.

The race runs FlashInfer's default tactic, which follows from the bucket's
row count alone (a low-M swap-AB tile for up to 32 rows, the default tile
above), and the kernel is shape-generic (TVM-FFI, symbolic extents), so a
handful per output dtype covers every layer. FlashInferMxfp8GemmProvider
runs FlashInfer's own runner on meta operands for every power-of-two
bucket up to max_num_tokens, with the compile capturing its key; the CuTe
DSL helpers compile each through FlashInfer's own compile closure, and a
wrapper of gemm_base._compile_block_scaled_gemm loads them into
FlashInfer's kernel cache (mm_fp4's calls pass through). The provider runs
when start_model_builds predicts the pass (tunes_flashinfer_mxfp8_at_load)
and there is no pipeline parallelism, as warmup decides.

Measured on DeepSeek-V4.1-Flash TP4 / EP4 with attention DP and DSpark
(GB300, empty JIT caches, TLLM_JIT_AHEAD=1, same node, on a1a4692): the
helpers build the 4 kernels while the weights load, and FlashInfer MXFP8
GEMM compiles on the executor in pass 1 go 4 (2.4 s) -> 0; pass 1 is
132.5 -> 133.0 s, within run-to-run noise (its CUDA graph capture phase,
where they compiled, 50.6 -> 49.8 s).

Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
…multi-sequence pass when no helper builds their kernels

With compile ahead on, warmup skips the passes that only compile, because
the helpers build their kernels. The DeepGEMM swap-AB tuning pass already
runs again whenever no DeepGEMM helper is running. The DSA indexer's
paged-MQA-logits metadata sweep and the Mamba hybrid warmup's
multi-sequence pass did not: if their helpers never started, or all exited
during warmup, their kernels compiled at the first serving step that
needed them.

Both now run as without compile ahead unless those helpers are running:
the DSA sweep unless DeepGEMM helpers are, which the indexer's variants are
queued to (jit_prefetch_deep_gemm.dsa_indexer_built_ahead), the Mamba pass
unless Triton helpers are on every rank, as its forwards are collective
(jit_prefetch.prefetch_helpers_running, agreed across the group).

Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
…etch test a unique basename

tests/unittest/_torch/modules/mamba/test_jit_prefetch_flashinfer.py shared
its basename with tests/unittest/_torch/test_jit_prefetch_flashinfer.py.
These directories are not packages, so a pytest session that collects both
fails with "import file mismatch". Rename it to test_mamba_ssu_jit_prefetch.py.

Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
…d RoPE modules while the weights load

A TRTLLM-Gen MoE layer on FlashInfer's op backend runs FlashInfer's fused
MoE ops, whose launcher module FlashInfer builds at the first call, and
RotaryEmbedding's query/key path builds FlashInfer's RoPE module at its
first call. On DeepSeek-V3.2 (TP4, TRTLLM MoE backend, empty JIT caches)
the executor built both in warmup pass 1, after the ~2 min weight load:
fused_moe_trtllm_sm100 in 885-901 s per rank and rope in 27 s.

Both follow from the constructed model, so start_model_builds now names
them (moe.fused_moe.jit_prefetch_flashinfer: a MoE layer whose op backend
is FlashInfer's; attention.jit_prefetch_flashinfer_rope: a RotaryEmbedding)
and FlashInfer helpers build them while the weights load; warmup joins them
before its first forward, as it does the other layer modules. The helper's
module table also takes generator keyword arguments, for the fused MoE
module's Rubin variant.

Warmup's join of the layers' modules waited at most
TLLM_JIT_PREFETCH_WAIT_TIMEOUT_S (300 s) and then stopped the helper, so the
forward rebuilt the module from scratch; with the fused MoE module that
made the server 215 s slower than building it in place. Stopping a layer
module's build can only restart it, so the join now waits up to an hour
unless that variable is set.

Measured on DeepSeek-V3.2 TP4 (GB300, empty JIT caches, TLLM_JIT_AHEAD=1,
two same-node pairs, with the next commit): the executor builds no
FlashInfer module in warmup (it built fused_moe_trtllm_sm100 in 905-918 s
per rank and rope in 27 s); server ready 1332 -> 1177 s and, in reverse
order, 1263 -> 1217 s. Serving is unchanged: 699.4 vs 696.6 output tok/s
and TPOT 3.48 vs 3.49 ms at concurrency 4, with three drafts accepted per
step.

Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
… kernel once per device

CuTe DSL's HardwareInfo.get_max_active_clusters compiles an empty kernel at
every call, into a fresh temporary dump directory, so CuTe DSL's compile
cache never hits: 60-100 ms on the executor each time a runner asks for
its cluster occupancy, which runners do whenever they compile or load a
kernel. On DeepSeek-V3.2 (TP4, empty JIT caches) three of four ranks
compiled it once while serving, about 14 s after ready; MiniMax-M3's
warmup compiled it 4 times per rank.

The kernel is the same for every query on a device: only the driver
occupancy queries that follow it depend on the cluster size. With compile
ahead on, start_model_builds now has HardwareInfo keep that kernel per
device (jit_prefetch_cute_dsl.install_occupancy_kernel_cache), so only a
process's first query compiles it.

Measured on DeepSeek-V3.2 (the runs of the previous commit): executor JIT
after ready, 3-4 occupancy-probe compiles per server -> none.

Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
out.append(n)
n *= 2
if out and out[-1] < self.max_batch_size:
out.append(self.max_batch_size)

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I am not sure if only enumerating the batch size is good enough to trigger all dsa kernels needed.

hd, H = t["head_dim"], t["num_heads"]
q_dtype, q_aligned = t["q"]
w_row, w_head, w_aligned, w_dtype = t["weights"]
for n in (r for r in _SCORE_ROWS if r <= self.max_num_tokens):

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

same here.. how do we make sure this logic can trigger all needed kernels?

…efetch_engine.py and the provider modules

The runtime files now only call into JIT compile ahead and tune ahead at
fixed points; what starts in each phase, and which kernels each family
builds or tunes, lives in new or existing jit_prefetch files.

* jit_prefetch_engine.py (new): EngineJitAhead, one per PyTorchModelEngine,
  runs the startup phases in the order they ran before: model constructed
  and weights allocated (from ModelLoader), the engine's warmup (decode
  recording, the runner's before-forwards and per-batch hooks, the end of
  warmup), and the serving-forward hook for the JIT stats. start_model_builds
  and start_weight_load_builds move here from jit_prefetch.py.
* Each family decides its kernels from the engine's model and limits next
  to its providers: jit_prefetch_deep_gemm.py (engine limits, the FP8 Linear
  and MoE providers, the DSA indexer and MegaMoE kernels warmup ran, the
  swap-AB tuning bypass), jit_prefetch_cute_dsl.py (model-construction and
  weight-load tactics, radix top-k / CSA2, end-of-warmup providers), the
  Mamba, MiniMax-M3, DSA and GLM kpool provider modules (register_providers,
  start_warmup_builds), modules/jit_prefetch_mxfp8.py (new: the autotuner
  warmup's FlashInfer MXFP8 pass predictor and tuning each decode graph shape
  once) and models/jit_prefetch_kimi_linear.py (new: Kimi K3's weight-read
  GEMM shapes).
* model_engine.py keeps a cached EngineJitAhead, one `with` around the
  runner's warmup and a two-line serving-forward hook; model_loader.py two
  hook calls; DecoderWarmup's two passes are one method again (helpers are
  cancelled in run()); graphs.py and runner.py keep only their hooks.

No behavior change: every moved function is AST-identical to its previous
body up to the renamed receiver, and each phase calls the families in the
order the engine methods did.

Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
(``jit_prefetch_flashinfer_helper``); warmup joins it before its first forward.
"""

from __future__ import annotations

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

FYI, @rosong11 is working on the removal of FI from TRTLLMGen.

@xxi-nv xxi-nv left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

On fused_moe_deepgemm.py#L740-L748:

Question on the design here: why does the masked grouped GEMM only report to done_observer, without a pre-launch wait like the other JIT paths?

Every other path waits for a helper that is already building the variant it needs:

Here the only hook is done_observer, and it fires after the launch, so it can only count a compile. The comment at jit_prefetch_deep_gemm.py#L437-L441 states this, but not why. As far as I can tell, if a warmup forward reaches a (num_groups, N, K) that a DeepGEMM helper is still compiling, the executor compiles the same kernel again. The result stays correct, since DeepJIT writes its cache atomically, but the second compile is wasted host time on the executor thread.

Is skipping the wait intentional? For example:

  • the window is too small to matter, because DeepGemmMoEProvider queues these requests as soon as prefetch starts; or
  • _queued_deep_gemm_spec only resolves FP8 Linear specs (spec_for), and mapping -num_groups to moe_spec_for was judged not worth it.

If it is intentional, could you add a short comment stating that a duplicate in-place compile is accepted here, and why? If not, a note_launch(-num_groups, n, k) before the launch, plus a _queued_deep_gemm_spec that resolves MoE specs, would bring this path in line with the others.

This branch has not been deployed

No deployments
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants