[None][feat] Support Inkling-NVFP4 model - #17062
Conversation
Add the Inkling NVFP4 model to the _torch stack: modeling_inkling.py, a Triton score_mod attention backend (SWA + global + relative-position bias) over KVCacheManagerV2, HF NVFP4 weight mapper/configs, trtllm-gen blockScaleMoe runner with sink-renorm routing, reasoning-parser/effort rendering, lm_eval post-processing, and the inkling_* unittest suite. Progress (working snapshot): - Component validation passes in isolation on the TP=4 NVFP4 / trtllm-gen MoE stack: weight load & accounting, attention source-activation replay, MoE replay, and full-model source-logit replay. - Baseline (cuda_graph=off, overlap=off) accuracy vs the SGLang reference: GSM8K 0.916 vs 0.972 (-5.6pt), full MMLU 82.22 vs 85.66 (-3.44pt). The gap is dominated by runaway / non-terminating generation on hard prompts: 76% of GSM8K errors are 7k-8.6k-token spirals where the model reaches the answer but never emits EOS, while SGLang commits. Per-layer localization is exhausted; the residual is a diffuse fp4 / kernel-family divergence (bf16 Triton attention + trtllm-gen fp4 MoE vs SGLang flashinfer), not a single fixable layer bug. - Enabled (cuda_graph=on) is blocked by a decode collapse (B2) localized to the global-attention block under CUDA-graph capture/replay at TP=4: h_attn goes non-finite at the first global-attention layer once decode crosses a KV-page boundary; a reduced-model TP=2 harness reproduces it. Next: decode-side termination fix for the baseline runaway (finish_reason-based detection + EOS/stop handling), and resolve B2 in the global-attention-under- graph path. Excludes build artifacts (libtensorrt_llm.so), locks, and the ext/ submodule. Signed-off-by: kleinc <kleinc@nvidia.com>
…(6*448)) Root cause of the baseline NVFP4 accuracy gap vs SGLang. The Inkling checkpoint ships routed-expert activation calibration as a RAW `.input_amax`, but the fused-MoE loader / trtllm `fp4_quantize` expect the ModelOpt per-tensor `input_scale = amax / (E2M1_MAX * E4M3_MAX) = amax / (6*448)`. The mapper renamed `.input_amax` -> `input_scale` WITHOUT the conversion, making the activation global scale 2688x too small: every routed expert's fp4 output was ~0.62 rel_rms from the bf16 ground truth (7.6x SGLang's 0.082). Mirrors sglang inkling.py:1222 / inkling_common/dense_mlp.py:497. Positional bisection confirmed the weight/block-scale/gate-up-interleave layout was already element-wise correct (24/24 L3 experts, 18.8M elems, 0 diverged) -- the defect was ONLY the activation input scale, shared by both TRT MoE backends. Fix: `inkling_weight_mapper._map_expert` divides `.input_amax` by (6*448). Result (baseline: cuda_graph=off, overlap=off, TP=4, CUTLASS MoE, fp4 fix active): - L3 routed-expert vs bf16 truth: rel_rms 0.624 -> 0.081 (== SGLang 0.082) - GSM8K paired 5x100: TRT 0.916 -> 0.968 vs SGLang 0.974; mean_delta -0.056 -> -0.006 (Gate 2 PASS, within +/-0.02); runaway canary passes; collapse 0 - MMLU paired 6x100: TRT ~0.822 -> 0.862 vs SGLang 0.872; gap -3.44pt -> -1.0pt (Gate 3 accuracy PASS); B-bias error-is-B 65.9% -> 40%, TRT-B 33.2% -> 29.8% Residual (not this fix): TRT (CUTLASS) vs SGLang (flashinfer) fp4 kernel-family near-tie noise -- accuracy-neutral, not bit-reproducible at batch>1. The strict within-2pp-of-gold MMLU B-bias criterion is model-inherent (SGLang itself over-picks B) and is left to human adjudication, not a TRT defect. Also adds env-gated per-layer/per-module dump instrumentation (modeling_inkling.py dump_sink; inkling_perlayer_localize_test.py) and the trtllm-gen MoE backend path used to localize the bug. Signed-off-by: kleinc <kleinc@nvidia.com>
…ring graph capture
Root cause of the Inkling TP=4 enabled-runtime (cuda_graph=on + overlap) decode collapse
("B2"): the AutoTuner-selected all-reduce (`tunable_allreduce`, AUTO strategy) is not
CUDA-graph-capture-safe. Frozen into a decode graph, the tuned tactic produces a non-finite
result on replay, so decode goes NaN from the first global-attention layer and collapses to
a token-0 repeat ("Paris!!!!"). Eager is finite at identical metadata; the fault is baked
into the captured graph.
Localized by single-variable determinism isolation (autotuner ON -> collapse, autotuner OFF
-> clean), NOT by op-fingerprints -- any in-graph probe shifts the graph memory pool / tactic
selection and suppresses the bug (a Heisenbug).
Fix (`tensorrt_llm/_torch/distributed/ops.py`, AllReduce.forward): during CUDA-graph capture
(`torch.cuda.is_current_stream_capturing()`), skip `tunable_allreduce` and fall back to the
static, graph-safe `all_reduce_op(AUTO)`. Warm-up / eager (not capturing) still autotune, so
steady-state performance is unchanged. This is a mainline TRT-LLM robustness fix (any TP
model that captures an AUTO all-reduce benefits), not Inkling-specific.
Confirmed: enabled generation_parity is BIT-IDENTICAL to the baseline (cuda_graph=off) --
tf_mismatch/neartie/confident=16/11/5, freerun_collapse=0, identical logit_checksum; enabled
5x100 GSM8K (0.966 vs SGLang 0.974) and MMLU (0.875 vs 0.872) within +/-0.02, zero errors,
no regression vs the accepted CUTLASS baseline.
Also includes env-gated per-layer/per-op B2 localization instrumentation
(`modeling_inkling.py` dump_sink / INKLING_FP*; `inkling_fp_localize_test.py`), zero-cost
when the INKLING_FP* env is unset.
Signed-off-by: kleinc <kleinc@nvidia.com>
…MLP tower, image fusion (WIP) WIP -- Stage-1 progress snapshot on the Inkling NVFP4 multimodal tower, on top of the accepted text tower. Not finished: the vision path is verified clean but the MMMU Accounting gap is still open (decode-side), and audio / MTP are still deferred. Committed to record current progress, not as a complete feature. Text decode/attention/MoE paths are untouched. Model / config: - configs/inkling.py: add image_token_id (200054, the in-vocab chat-template <|unused_200054|>) and audio_token_id. The SGLang-internal -101 sentinel is rejected by TensorRT-LLM's executor token-id validation; the two ids are interchangeable for parity since both are overwritten by vision embeddings. - modeling_inkling_vision.py (new): hMLP vision tower InklingVisionModel and InklingInputProcessor, which expands the <image> placeholder to one token per vision patch and attaches vision_patches_bthwc features. - modeling_inkling.py: InklingForConditionalGeneration registers the input processor, builds the vision tower as a replicated bf16 submodule, and fuses per-patch embeddings into the text stream via fuse_input_embeds with explicit text/mm indices (OOV-safe). Fixes the "image not visible" hallucination: the fused stream must NOT be re-normed, since SGLang scatters raw vision rows in after embed_norm; the extra RMSNorm corrupted the image rows. Tests / diagnostics (tests/unittest/_torch/modeling/, 25 new files): MMMU harness alignment against the SGLang scorer, input-processor and vision-tower unit checks, image e2e / fusion / logit-replay / generation-parity drivers, and the localization probes used to isolate the vision-vs-decode split (vision verified bitwise-clean; the residual Accounting gap is decode-side). These require GPU + the NVFP4 checkpoint (TP=4) and were not run for this commit. Signed-off-by: kleinc <kleinc@nvidia.com>
…s_token SamplingParams._setup() looked for the end-of-generation token in only two places: tokenizer.eos_token_id, then generation_config.eos_token_id. When a checkpoint provides neither, end_id stayed None, nothing could terminate a request, and every generation silently ran to max_tokens. That combination is reachable. Multi-part chat formats have no single terminator -- a message end, an end-of-sampling marker and a document separator are distinct tokens -- so such checkpoints register their control tokens under extra_special_tokens / additional_special_tokens, neither of which populates tokenizer.eos_token_id, and declare the real stop token as eos_token_id in config.json. Those checkpoints also tend to ship no generation_config.json. _setup() already received hf_model_config but never consulted it. Observed on a checkpoint of that shape: responses ran to the token limit while emitting the configured eos_token_id up to 441 times in a single response. Fall back to hf_model_config.eos_token_id when the first two sources yield nothing. A list value sets end_id from its first entry and appends the rest to stop_token_ids, mirroring the existing generation_config path. Priority is otherwise unchanged: an explicit SamplingParams(end_id=...) still wins and tokenizer.eos_token_id still takes precedence over config.json, so models that already resolved an end_id see no behavioural change. Warn when all three sources come up empty. Generating to the cap on every request with no diagnostic is the part that makes this expensive to find. Add unit tests for tokenizer priority, the config fallback, the list form, an explicit end_id not being overridden, the all-empty case, and a missing hf_model_config. Signed-off-by: KleinBlueC <102355066+KleinBlueC@users.noreply.github.com>
Adds the audio and video modalities alongside the existing vision path.
Audio:
* InklingAudioPreprocessor -- dMel feature extraction (mel basis, hz<->mel,
per-frame bin quantization) producing int32 [N, 80] bins, one audio token
per frame.
* InklingAudioModel -- codebook encoder (bin m occupies codebook rows
[m*V, (m+1)*V), summed over bins) plus optional final norm, bf16, loaded
strictly from the real checkpoint's audio tensors.
* Placeholder expansion and fail-loud count checks in the input processor.
Video:
* sample_video_frames / sample_video_as_images / DecodedVideo -- frame
sampling ported to match SGLang's sample_video_frames semantics.
* The <image>-per-frame path: every sampled frame becomes its own image
span through the existing vision tower.
Audio and video land together because both wire into the same
InklingInputProcessor.assemble dispatch; splitting them would leave an
intermediate commit whose processor references helpers that do not exist yet.
Tests (all GPU-verified on TP=4):
* inkling_audio_tower_test.py -- 10 passed, incl. real-weight CUDA forward
(AUDIO_TOWER_CUDA_OK, out=(n_frames, 6144) bf16 finite) and a
reference-math allclose(atol=1e-5) check of the codebook sum.
* inkling_video_utils_test.py -- 12 passed, incl. a port of SGLang's
test_video_utils.py::test_sample_video_frames_lengths (same 4 cases and
the same expected frame indices) and a real-weight multi-frame CUDA
forward (VIDEO_TOWER_CUDA_OK, out=(total_patches, 6144)).
* inkling_audio_e2e_test.py / inkling_video_e2e_test.py -- strict TP=4
end-to-end smokes: every prompt must be finite, non-empty and
non-collapsed. Green both baseline (5/5) and with cuda_graph+overlap
enabled (5/5).
Scope note: these prove the modalities run and are shape/dtype/finiteness
correct; they are not a cross-stack numerical parity check against SGLang.
Signed-off-by: KleinBlueC <102355066+KleinBlueC@users.noreply.github.com>
`logits_rows` is a view of the same slice that is written back, so the assignment self-overlaps and torch rejects it with "... refer to a single memory location. Please clone()...". The processors edit in place, so the write-back is redundant anyway; cloning the source makes it overlap-safe. Only reached for requests carrying a py_logits_post_processor, so normal requests are unaffected. Signed-off-by: KleinBlueC <102355066+KleinBlueC@users.noreply.github.com>
… parity runs evaluate/interface.py gains summarize_generation_stats/log_generation_stats, and lm_eval emits a greppable GEN_STATS marker per batch. Without it a runaway / no-EOS regression hides behind a parseable answer buried in a wall of repeated text -- the failure mode that cost several bring-up iterations. The MMMU harness/runner changes carry the sharded union runs used for the TRT-vs-SGLang comparison (shard plan, incremental atomic per-item writes so a wall-killed shard keeps everything it scored, cap/max_seq plumbing), plus unit coverage of the answer parser. Signed-off-by: KleinBlueC <102355066+KleinBlueC@users.noreply.github.com>
Adds streaming-vs-batch equivalence over arbitrary split points, tool-call and repetition segmentation, and end-tokens split across deltas. Signed-off-by: KleinBlueC <102355066+KleinBlueC@users.noreply.github.com>
Standalone probes used to localize the TRT-vs-SGLang vision divergence: prefill logits, transformers-reference decode, termination behaviour, and a per-layer activation dump. Diagnostics only -- not part of any test suite. Signed-off-by: KleinBlueC <102355066+KleinBlueC@users.noreply.github.com>
dfc574c to
f5fcaf1
Compare
Two in-progress pieces, neither reached by a runtime path yet. MTP static tier: InklingMTPConfig carries the checkpoint's num_nextn_predict_layers / chain_hidden_post_norm / local_layer_ids plus the per-depth banded-attention geometry injected from the text tower, so draft depths in local_layer_ids get SWA head geometry and the rest stay global. The weight mapper gains inkling_expected_mtp_keys() and accounts model.mtp.* as consumed rather than deferred when an mtp_config is supplied; the default mtp_config=None leaves the text-tower accounting byte-identical. Unit coverage for config parse, weight accounting, the BF16/unquantized requirement, and per-depth banding against the checkpoint shapes. MMMU harness: INKLING_MMMU_TEXT_ONLY reruns the same items as pure text (image placeholder stripped, no image attached) so the shared decoder is exercised at the identical bs / cap / overlap regime without the vision path. Default-off — the vision scoring path is unchanged when unset. Signed-off-by: KleinBlueC <102355066+KleinBlueC@users.noreply.github.com>
…del code The cuda-graph decode collapse is a defect in symmetric all-reduce: a captured NCCL_SYMMETRIC reduce whose send buffer is unregistered while its recv buffer is a registered NCCL window corrupts the run at a 12288 B message. Inkling hits that size exactly -- hidden 6144, bf16, one decode token -- so the first global-attention layer goes non-finite and decode collapses to a repeated token 0. Revert the shared-code mitigation in distributed/ops.py: that file is now byte-identical to its pre-Inkling state, so no other model's all-reduce path changes on our account. Mitigate in modeling_inkling.py instead. The all-reduces that trigger this are built by generic modules -- attention o_proj, MoE down_proj -- so the strategy cannot be passed at construction without editing shared code; rebuilding each AllReduce after super().__init__() keeps the mitigation model-local. Each rebuilt instance carries the module's own mapping and dtype over, so strategy is the only delta. Pinning ONESHOT also drops the window requirement, since AllReduce only takes an NCCL window under NCCL_SYMMETRIC/NCCL/AUTO -- two of the five trigger conditions go away, not just one. Active by default; INKLING_ALLREDUCE_STRATEGY=AUTO restores stock behaviour, and the defect with it, for A/B runs. Measured on job 5728192, alternating arms in one job against one binary: default 0/3 collapse, AUTO 3/3, 331 modules swapped. Cost: symmetric is disabled on every Inkling all-reduce, eager included, and roughly a third of captured decode all-reduces pick it today. The performance impact is unmeasured and should be measured before this is treated as final. This is containment, not a root-cause fix -- the defect remains for any other model that meets all five conditions. Signed-off-by: KleinBlueC <102355066+KleinBlueC@users.noreply.github.com>
f5fcaf1 to
35a4a46
Compare
Removes the bring-up's debug scaffolding and completes the deliverables the Inkling NVFP4 change was missing. No model behaviour changes. Deleted 41 debug/localization test files (divergence probes, dump/isolate helpers, teacher-forcing and graph-capture localizers) that existed to find the bring-up's defects and have no role now that it works. Nothing imports them: every deleted basename was grepped across the tree, zero references remain. The two environment variables the model still reads are the ones worth keeping — INKLING_ALLREDUCE_STRATEGY (the escape hatch for the ONESHOT all-reduce mitigation) and INKLING_MOE_BACKEND (the trtllm-gen MoE kernel select) — and both are documented; the 19 debug-only INKLING_* knobs are gone. Completed the deliverables: - docs/source/models/supported-models.md — architecture row, multimodal feature-matrix row, and footnote [^14] covering modality coverage, the unsupported set (MTP, LoRA, function calling, constrained decoding, EPD, mm-hash caching), and the all-reduce mitigation plus its escape hatch. - TestInkling_NVFP4::test_nvfp4 added to the shared multimodal accuracy file rather than a standalone per-model test, with a sourced MMMU reference. - Registered that id in test-db/l0_b200.yml and qa/llm_function_core.txt. - Dropped the orphaned inkling_vision_tower_artifact.json; the surviving regression test consumes the generated artifact instead. Static checks: git diff --check clean, every modified Python file compiles, no stale references to the deleted files. Runtime coverage as of the last completed suite: unit tiers green (vision 43, audio 10, video 12, text 31, collect 95 with zero residual debug files) and GSM8K cg0ov0 parity at delta=0.0. Signed-off-by: KleinBlueC <102355066+KleinBlueC@users.noreply.github.com>
d1649bb to
8a32f22
Compare
The MTP / next-N draft work never reached a runtime path: nothing builds the draft layers and nothing consumes them. Only the config plumbing and the load-accounting existed, so remove them rather than ship dead code. - configs/inkling.py: drop InklingMTPConfig, InklingConfig.has_mtp and _as_mtp_config. mtp_config goes back to a plain retained blob, so the checkpoint still round-trips. - inkling_weight_mapper.py: drop inkling_expected_mtp_keys and the consumed_mtp bucket; inkling_account_checkpoint loses its mtp_config parameter. - test_modeling_inkling.py: drop the four Stage-9 MTP tests. - supported-models.md: the footnote no longer claims the draft weights are weight-accounted. Checkpoint accounting is unaffected: "model.mtp." stays in INKLING_DEFERRED_PREFIXES, so the draft weights are classified as deferred exactly like the audio and vision blocks, and `unaccounted` stays empty. Verified: every touched file compiles, and no reference to any removed symbol remains anywhere in the tree. Signed-off-by: KleinBlueC <102355066+KleinBlueC@users.noreply.github.com>
…ences MMMU already had a reference entry; the two text benchmarks the bring-up actually measured did not, so the numbers lived only in run logs. Both references are the cached SGLang NVFP4 measurement under the same harness, matching how the MMMU entry was sourced: - GSM8K 95.53 (full set, 0.9553). TRT-Inkling was validated against it with the paired 5x100 protocol: flexible-parse mean 0.968, and all four cuda-graph x overlap-scheduler corners scored 0.98-0.99 on the 100 paired items where SGLang scored 0.99. - MMLU 85.66 (full Hendrycks, 14042 samples, weighted_accuracy 85.6573). The TRT side was measured with the 5-seed text-regression protocol on the harness' 114-item subset (83.33 / 84.21 / 85.09 / 85.96 / 87.72, mean ~85.3). The comment says so explicitly: it tracks the reference, but no full-set TRT-LLM MMLU run has been recorded yet. Signed-off-by: KleinBlueC <102355066+KleinBlueC@users.noreply.github.com>
…that matters
The bring-up left 22 Inkling files under tests/unittest/_torch/modeling. Most
were per-defect regressions or SGLang-alignment scaffolding whose outcome is
now covered end-to-end by the MMMU/GSM8K accuracy tests. Keep one test per
thing that nothing else covers, drop the rest.
Kept:
- test_modeling_inkling.py config parse, layer classification, weight
accounting over the real checkpoint
- inkling_vision_tower_test.py vision tower
- inkling_audio_tower_test.py audio tower (no accuracy benchmark covers it)
- inkling_video_utils_test.py video utils (no accuracy benchmark covers it)
- inkling_input_processor_test.py multimodal input processor
- inkling_moe_backend_select_test.py the INKLING_MOE_BACKEND kernel select
inkling_mmmu_real_align_test.py was doing double duty: the vision-tower and
input-processor tests import its MMMU item fetch/cache and its importlib
loader. Split that half out as inkling_mmmu_fixtures.py (a fixture module, no
tests) and drop the SGLang-alignment machinery with the rest.
Deleted (16): the mmmu align/harness/run/parser set, image_prompts, the three
per-modality e2e smokes, image_fusion, image_norm_fix, attn_decode_meta,
gate_up_deinterleave, kv_manager_v2, generation_parity and
source_logit_replay.
Every kept file compiles and no reference to a deleted module remains in the
tree.
Signed-off-by: KleinBlueC <102355066+KleinBlueC@users.noreply.github.com>
The multimodal file already carried TestInkling_NVFP4::test_nvfp4 for MMMU, so the vision path had an integration test but the shared text decoder -- what GSM8K and MMLU actually exercise -- had none. Adds TestInkling_NVFP4 to test_llm_api_pytorch.py running both text benchmarks against the references recorded earlier. It mirrors the multimodal class: NVFP4 assert, 16384-token budget for the long chain of thought, and extract_inkling_content as the post-processor so the <|content_thinking|> channel is dropped and only the visible answer is scored. Registered the new id in both lists that already carry the multimodal one: test-db/l0_b200.yml and qa/llm_function_core.txt. Signed-off-by: KleinBlueC <102355066+KleinBlueC@users.noreply.github.com>
…keep CUTLASS The Inkling routed experts now run only the default CUTLASS MoE backend. The trtllm-gen blockScaleMoe path was an opt-in experiment behind INKLING_MOE_BACKEND=TRTLLM, never the default, and it needed a dedicated routing enum plumbed all the way into the CUDA runner to work. Restored to their pre-bring-up state (the additions were Inkling-only, so these files are now byte-identical to 44f0521): - cpp/.../trtllmGenKernels/blockScaleMoe/runner.h the InklingSinkRenorm = 9 enum value and its name case - cpp/.../trtllmGenKernels/blockScaleMoe/runner.cu the precomputed-routing dispatch branch - _torch/modules/fused_moe/routing.py the matching Python enum value and its autotuner-dummy mapping Removed from the model: - _inkling_trtllm_moe_backend / _moe_config_with_trtllm_backend and the frozen-config copy they needed to retarget moe_backend - InklingMoeRoutingMethod._trtllm_backend, so routing_method_type is always Unspecified and requires_separated_routing goes back to the default - the per-layer INKLING_MOE_SELECT backend-introspection log Also deletes inkling_moe_backend_select_test.py, which existed only to cover that knob, and the stale comment that pointed at it as the fix for the fused combine's cross-row non-determinism. INKLING_ALLREDUCE_STRATEGY is now the only environment variable the model reads. Everything compiles and git diff --check is clean. Signed-off-by: KleinBlueC <102355066+KleinBlueC@users.noreply.github.com>
The MMLU comment added in 5f28628 claimed no full-set TRT-LLM run had been recorded. That is wrong. One exists: 2026-07-20, all 14042 Hendrycks samples, weighted_accuracy 82.2176 -- 3.44 points below the 85.66 reference, and it failed the bring-up's 2-point gate. The gap was later traced to an fp4 expert-GEMM family difference (CUTLASS/trtllm-gen vs SGLang's flashinfer) plus non-terminating generation on hard prompts. What is true is narrower: the full set has not been re-measured since those fixes. Post-fix MMLU evidence is subset-only -- a 570-item stratified canary at 84.21 (-1.45, inside the gate) and a 5-seed 114-item regression averaging ~85.3. The comment now says exactly that, including that meeting 85.66 at full scale is unverified. GSM8K's comment was not wrong but was too vague about coverage. TRT-LLM has never been measured on the full 1319-item set: the evidence is the paired 5x100 protocol (flexible mean 0.968), the four cuda-graph x overlap corners at 0.98-0.99 on 100 paired items, and one early full-set attempt that completed only its first 120-item chunk (0.925 vs 0.975). Spelled out. Both files still record the SGLang reference as the accuracy value, so TestInkling_NVFP4::test_nvfp4 grades against a bar the model has not been shown to clear at full scale. Signed-off-by: KleinBlueC <102355066+KleinBlueC@users.noreply.github.com>
…timodal
The module already carried the image, audio and video paths, so the "vision"
name no longer described it. Rename it and reorganize the contents into four
labelled sections -- vision tower, audio tower, video tower, and the shared
multimodal input processor -- so each modality is easy to locate.
Behavior-preserving. Along the way:
* factor the duplicated tower weight-loading into ``_load_tower_weights``
and the duplicated RMSNorm into a single shared ``InklingRMSNorm``
(was ``InklingVisionRMSNorm``, used by both towers);
* factor the input processor's repeated media-list coercion and its
placeholder/feature-row count checks into small helpers, dropping a
tautological per-item check (``num_tokens`` is built from ``num_patches``);
* trim the docstrings to what the code needs, dropping the development-time
stage/goal references and reference-implementation file paths.
Signed-off-by: KleinBlueC <102355066+KleinBlueC@users.noreply.github.com>
Bring-up leftovers that do not belong in the production model:
* cuda_graph_runner: drop the "[cuda-graph] CAPTURED/REPLAYED" evidence
logging and its once-only gate flag. This was pure instrumentation added
to prove the runtime really captured and replayed a graph, and it sat in
shared (non-Inkling) code.
* modeling_inkling: drop the INKLING_ALLREDUCE_STRATEGY environment knob.
The ONESHOT pin is a correctness mitigation, not a tuning option, so it
is applied unconditionally instead of via an A/B toggle (and the log line
that reported it is gone).
* modeling_inkling: drop the explicit-window short-conv decode path
(InklingShortConv.forward_decode, the decoder layer's third branch, and
the attention's conv_states/return_conv_state arguments). Only the
bring-up replay harness ever passed those; the runtime always drives the
short convs through the per-request state pool.
* modeling_inkling: drop the attention's decode_seq_lens / decode_page_table
/ skip_kv_write arguments. The runtime publishes decode metadata into the
layer's stable GPU buffers before capture, so the pre-supplied static
tensors had no caller left; the eager fallback that builds them from the
host block table stays.
* modeling_inkling: drop InklingConvStateCache.reset (unused) and merge
InklingConvRuntime.from_metadata into build (the split existed only so the
replay harness could publish slots itself).
Also rewrite the comments and docstrings across the Inkling _torch files to
describe the code as it stands: no development stage/goal numbers, job ids,
absolute paths into local reference checkouts, or references to the deleted
replay harness. The stale "audio / vision / MTP are deferred" module docstring
now reflects that only MTP is unimplemented.
Signed-off-by: KleinBlueC <102355066+KleinBlueC@users.noreply.github.com>
The Inkling unit tests were five files carrying a bring-up harness rather than a test suite: they loaded the SGLang reference from an absolute path in a local checkout, downloaded MMMU rows over the network into a gitignored cache, wrote JSON artifacts, hardcoded a personal /lustre checkpoint path, and shipped __main__ runners that printed comparison tables. None of that can run in CI. Fold everything multimodal into test_modeling_inkling_multimodal.py -- one file covering the vision, audio and video paths plus the shared input processor -- with small synthetic configs and inputs: no checkpoint, no GPU, no network, no SGLang import, no artifacts. 27 tests, well under a second. Slim test_modeling_inkling.py the same way. The config and layer-classification tests now build their config explicitly instead of requiring the checkpoint, so they actually run; the weight-accounting and tensor-shape tests keep the checkpoint (index JSON only, no weights) and resolve it through the standard llm_models_root() with an INKLING_CHECKPOINT override, so they skip cleanly instead of always skipping on a path that only existed on one machine. Removed: inkling_vision_tower_test.py, inkling_input_processor_test.py, inkling_audio_tower_test.py, inkling_video_utils_test.py (folded in) inkling_mmmu_fixtures.py (MMMU downloader + SGLang loader; no longer used) Also drop the .gitignore entries for the deleted caches and artifacts. Signed-off-by: KleinBlueC <102355066+KleinBlueC@users.noreply.github.com>
…nkling PR Run the repo's pre-commit formatters (isort, yapf, ruff, ruff-format) over the files this PR touches, and fix the eleven ruff-legacy D205/D209 docstring regressions it flags in tensorrt_llm/evaluate/interface.py and tests/unittest/llmapi/test_reasoning_parser.py. Formatting and docstring wording only; no behavior change. Signed-off-by: KleinBlueC <102355066+KleinBlueC@users.noreply.github.com>
Two bring-up leftovers outside _torch:
* evaluate: drop summarize_generation_stats / log_generation_stats and both
call sites (Evaluator.evaluate and LmEvalWrapper). This logged a greppable
GEN_STATS line with the finish_reason mix and generated-token distribution
so a runaway/no-EOS regression would be visible while bringing the model up.
It is observability scaffolding, not part of the eval contract: it never
influenced what was scored, nothing consumed the marker, and no test
covered it.
* docs: drop "The text decoder is also usable standalone (text-only) via the
InklingForCausalLM architecture" from the Inkling footnote. It is not true.
InklingForCausalLM carries no @register_auto_model, so it is not in
MODEL_CLASS_MAPPING and no checkpoint can select it; the config registry has
no inkling_text entry either. The class is the base that
InklingForConditionalGeneration derives from, and the inkling_text handling
in config_utils/_util exists for that nested sub-config, not for standalone
loading. The published checkpoint declares InklingForConditionalGeneration
and the text-only accuracy test loads it through that same architecture, so
registering the class would advertise a path with no checkpoint to exercise
it and no test coverage.
Signed-off-by: KleinBlueC <102355066+KleinBlueC@users.noreply.github.com>
8658713 to
c15285c
Compare
📝 WalkthroughWalkthroughChangesInkling NVFP4 text and multimodal support is added to TensorRT-LLM. The changes include model configuration, checkpoint mapping, Triton attention, convolution state management, vision/audio processing, reasoning parsing, evaluation integration, serving updates, documentation, and accuracy coverage. Inkling support
Estimated code review effort: 5 (Critical) | ~120 minutes Sequence Diagram(s)sequenceDiagram
participant Request
participant InklingInputProcessor
participant InklingForConditionalGeneration
participant InklingAttention
participant inkling_triton
Request->>InklingInputProcessor: submit text and optional media
InklingInputProcessor->>InklingForConditionalGeneration: assemble tokens and modality features
InklingForConditionalGeneration->>InklingAttention: run hybrid decoder layers
InklingAttention->>inkling_triton: execute prefill or paged decode
inkling_triton-->>InklingAttention: return attention results
InklingForConditionalGeneration-->>Request: return generated output
Possibly related PRs
Suggested reviewers: 🚥 Pre-merge checks | ✅ 4 | ❌ 1❌ Failed checks (1 warning)
✅ Passed checks (4 passed)
✨ Finishing Touches 💡 1⚔️ Resolve merge conflicts 💡
🧪 Generate unit tests (beta)
Comment |
There was a problem hiding this comment.
Actionable comments posted: 15
🧹 Nitpick comments (21)
tests/integration/defs/accuracy/test_llm_api_pytorch.py (1)
7815-7815: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winUse PascalCase test class names and update direct selectors.
Rename both classes to
TestInklingNVFP4. Update the listed selectors to the new class name.
tests/integration/defs/accuracy/test_llm_api_pytorch.py#L7815-L7815: renameTestInkling_NVFP4.tests/integration/defs/accuracy/test_llm_api_pytorch_multimodal.py#L683-L683: renameTestInkling_NVFP4.tests/integration/test_lists/qa/llm_function_core.txt#L812-L813: update the two direct class selectors.tests/integration/test_lists/test-db/l0_b200.yml#L355-L356: update the two direct class selectors.As per coding guidelines, use PascalCase for classes.
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@tests/integration/defs/accuracy/test_llm_api_pytorch.py` at line 7815, Rename both TestInkling_NVFP4 classes to TestInklingNVFP4 in tests/integration/defs/accuracy/test_llm_api_pytorch.py:7815-7815 and tests/integration/defs/accuracy/test_llm_api_pytorch_multimodal.py:683-683. Update both direct class selectors to TestInklingNVFP4 in tests/integration/test_lists/qa/llm_function_core.txt:812-813 and tests/integration/test_lists/test-db/l0_b200.yml:355-356.Source: Coding guidelines
tensorrt_llm/_torch/configs/inkling.py (2)
137-148: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valueAnnotate
_local_idswith the element type and consider caching it.
_local_idsrebuilds aseton everyis_local_layercall.num_kv_heads_per_layer()calls it once per layer. The cost is small, but a precomputed set is simpler and the annotation should beset[int]per the typing guideline.♻️ Proposed refactor
- `@property` - def _local_ids(self) -> set: - return set(self.local_layer_ids) + `@property` + def _local_ids(self) -> set[int]: + cached = getattr(self, "_local_ids_cache", None) + if cached is None or len(cached) != len(self.local_layer_ids): + cached = set(self.local_layer_ids) + self._local_ids_cache = cached + return cachedAs per coding guidelines: "prefer built-in generic types and
|".🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@tensorrt_llm/_torch/configs/inkling.py` around lines 137 - 148, The _local_ids property should use the built-in generic annotation set[int] and avoid rebuilding the set for every is_local_layer call. Precompute or cache the converted local_layer_ids set at the configuration level, then have is_local_layer reuse that cached set while preserving its current membership behavior.Source: Coding guidelines
195-248: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valueAdd type annotations to the constructor parameters and
_as_config.
text_config,audio_config,vision_config,mtp_config, and_as_config(value)carry no annotations. The repository guideline requires annotating every function. Usedict[str, Any] | PretrainedConfig | None.♻️ Proposed refactor
def __init__( self, - text_config=None, - audio_config=None, - vision_config=None, - mtp_config=None, + text_config: "dict[str, Any] | InklingTextConfig | None" = None, + audio_config: "dict[str, Any] | PretrainedConfig | None" = None, + vision_config: "dict[str, Any] | PretrainedConfig | None" = None, + mtp_config: "dict[str, Any] | PretrainedConfig | None" = None, eos_token_id: int = 200006,`@staticmethod` - def _as_config(value): + def _as_config( + value: "dict[str, Any] | PretrainedConfig | None", + ) -> "PretrainedConfig | None":As per coding guidelines: "Annotate every function, use
Nonefor non-returning functions".🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@tensorrt_llm/_torch/configs/inkling.py` around lines 195 - 248, Add type annotations to InklingConfig.__init__ parameters text_config, audio_config, vision_config, and mtp_config using dict[str, Any] | PretrainedConfig | None, and annotate _as_config(value) with the same input type and an appropriate return type. Add the required Any import, and annotate __init__ as returning None while preserving the existing configuration handling.Source: Coding guidelines
tests/unittest/_torch/modeling/test_modeling_inkling.py (2)
52-68: 🚀 Performance & Scalability | 🔵 Trivial | 💤 Low valueCache the parsed safetensors index.
_safetensors_shapecalls_index(ckpt)for every tensor lookup, so the index JSON is re-read and re-parsed once per assertion. Use the module-scoped fixture pattern orfunctools.lru_cacheto read it once.♻️ Proposed refactor
+from functools import lru_cache + +@lru_cache(maxsize=None) def _index(ckpt: str) -> dict: with open(os.path.join(ckpt, "model.safetensors.index.json")) as f: return json.load(f)["weight_map"]🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@tests/unittest/_torch/modeling/test_modeling_inkling.py` around lines 52 - 68, Cache the parsed weight map used by _safetensors_shape so model.safetensors.index.json is read and parsed only once per checkpoint. Apply a module-scoped fixture or functools.lru_cache to _index, preserving _safetensors_shape’s existing lookup behavior.
122-141: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winCoverage gap: the gate/up de-interleave is untested.
The tests cover accounting (
inkling_account_checkpoint,inkling_nvfp4_expert_layers) and checkpoint shapes, but not the load-path remapping ininkling_weight_mapper.py._split_interleaved_gate_upis correctness-critical: a contiguouschunk(2)instead of the strided even/odd split pairs the wrong SwiGLU channels and produces silently wrong output, with no shape error to catch it. That function is pure CPU and needs no checkpoint.Coverage assessment: sufficient for config/accounting, insufficient for the mapper load path. Add a small CPU test in this file, for example:
💚 Suggested test
def test_split_interleaved_gate_up_takes_even_odd_rows(): from tensorrt_llm._torch.models.checkpoints.hf.inkling_weight_mapper import ( _split_interleaved_gate_up, ) t = torch.arange(8).reshape(4, 2) # rows [g0, u0, g1, u1] gate, up = _split_interleaved_gate_up(t, dim=0) assert torch.equal(gate, t[0::2]) assert torch.equal(up, t[1::2]) with pytest.raises(ValueError): _split_interleaved_gate_up(torch.zeros(3, 2), dim=0)Do you want me to open an issue to track this?
As per path instructions: "suggest concrete list file names and whether coverage is sufficient, insufficient, or needs follow-up outside the PR".
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@tests/unittest/_torch/modeling/test_modeling_inkling.py` around lines 122 - 141, Add a CPU unit test in the test file for _split_interleaved_gate_up, using an even-row tensor to verify gate receives even-indexed rows and up receives odd-indexed rows, and assert that an odd-sized split dimension raises ValueError. Import the mapper helper and pytest as needed; no checkpoint setup is required.Source: Path instructions
tensorrt_llm/_torch/models/modeling_inkling_multimodal.py (1)
1021-1028: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valueAnnotate the rejecting getters as non-returning.
get_num_tokens_per_image,get_num_tokens_per_audio, andget_num_tokens_per_videoare annotated-> intbut always raise. Usetyping.NoReturnso type checkers do not report a missing return.♻️ Proposed refactor
- def get_num_tokens_per_image(self, *, image=None, **kwargs) -> int: + def get_num_tokens_per_image(self, *, image=None, **kwargs) -> NoReturn: self._reject_mm_hash_cache("image") - def get_num_tokens_per_audio(self, *, audio=None, **kwargs) -> int: + def get_num_tokens_per_audio(self, *, audio=None, **kwargs) -> NoReturn: self._reject_mm_hash_cache("audio") - def get_num_tokens_per_video(self, *, video=None, **kwargs) -> int: + def get_num_tokens_per_video(self, *, video=None, **kwargs) -> NoReturn: self._reject_mm_hash_cache("video")As per coding guidelines: "Annotate every function ... use
Literal,overload,TypeVar, orProtocolwhen appropriate."🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@tensorrt_llm/_torch/models/modeling_inkling_multimodal.py` around lines 1021 - 1028, Update the return annotations of get_num_tokens_per_image, get_num_tokens_per_audio, and get_num_tokens_per_video from int to typing.NoReturn, ensuring NoReturn is imported or referenced through the existing typing conventions.Source: Coding guidelines
tests/unittest/_torch/modeling/test_modeling_inkling_multimodal.py (3)
173-181: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valueDerive the expected bin geometry from the module constants.
Lines 177 and 179 hardcode
80and16. Those areDEFAULT_AUDIO_N_MELSandDEFAULT_AUDIO_NUM_DMEL_BINS. If a default changes, the test fails with an opaque number instead of pointing at the constant.♻️ Proposed refactor
- assert bins.ndim == 2 and bins.shape[1] == 80 + assert bins.ndim == 2 and bins.shape[1] == DEFAULT_AUDIO_N_MELS assert bins.shape[0] >= 8, f"expected ~12 frames for 0.6s, got {bins.shape[0]}" - assert 0 <= int(bins.min()) and int(bins.max()) < 16 + assert 0 <= int(bins.min()) and int(bins.max()) < DEFAULT_AUDIO_NUM_DMEL_BINSThe same applies to line 193 in
test_dmel_multi_clip_concat_and_empty_clip.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@tests/unittest/_torch/modeling/test_modeling_inkling_multimodal.py` around lines 173 - 181, Update test_dmel_preprocess_shape_and_range and test_dmel_multi_clip_concat_and_empty_clip to derive the expected mel-bin width and value range from DEFAULT_AUDIO_N_MELS and DEFAULT_AUDIO_NUM_DMEL_BINS instead of hardcoded 80 and 16. Keep the existing assertions and behavior unchanged while referencing the module constants directly.Source: Path instructions
124-167: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winCoverage gap: image input coercion is untested.
Every vision test feeds a numpy array.
_to_pil_rgbalso accepts raw bytes, a PIL image, afile://URI, and a local path, and it explicitly rejectshttp://,https://, anddata:inputs. That rejection is a deliberate guard against the preprocessor reaching for the network; nothing pins it today.Coverage assessment: sufficient for the tower geometry and forward, insufficient for input coercion. Add a short CPU test:
💚 Suggested test
def test_to_pil_rgb_rejects_remote_sources_and_accepts_bytes(): from tensorrt_llm._torch.models.modeling_inkling_multimodal import _to_pil_rgb for url in ("http://x/y.png", "https://x/y.png", "data:image/png;base64,AA=="): with pytest.raises(ValueError, match="resolve"): _to_pil_rgb(url) with pytest.raises(TypeError): _to_pil_rgb(object()) assert _to_pil_rgb(_image(8, 8)).mode == "RGB"Do you want me to open an issue to track this?
As per path instructions: "suggest concrete list file names and whether coverage is sufficient, insufficient, or needs follow-up outside the PR".
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@tests/unittest/_torch/modeling/test_modeling_inkling_multimodal.py` around lines 124 - 167, Add a CPU-focused test near the existing vision input tests covering _to_pil_rgb: assert http://, https://, and data: inputs raise ValueError matching “resolve”, unsupported objects raise TypeError, and a generated image is accepted with RGB mode. This closes the input-coercion coverage gap while preserving the existing tower tests.Source: Path instructions
290-306: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low value
test_registrationreaches into private registry internals.Line 304 reads
INPUT_PROCESSOR_REGISTRY._input_processors_cls_by_model_typeand line 306 readsInklingInputProcessor._registered_model_type. Both are private. A rename intensorrt_llm/inputs/registry.pybreaks this test without any behavior change in Inkling. If the registry exposes a public lookup, prefer it.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@tests/unittest/_torch/modeling/test_modeling_inkling_multimodal.py` around lines 290 - 306, Update test_registration to validate processor registration through the public lookup API exposed by INPUT_PROCESSOR_REGISTRY instead of accessing _input_processors_cls_by_model_type. Also avoid asserting InklingInputProcessor._registered_model_type directly; verify the expected "inkling_mm_model" association through the public registry interface while preserving the existing registration expectations.Source: Path instructions
tensorrt_llm/_torch/models/modeling_inkling.py (4)
265-274: 🚀 Performance & Scalability | 🔵 Trivial | 💤 Low valueAvoid the per-forward pageable allocation in
write_state_indices.
torch.tensor(slots, dtype=torch.int32)allocates a pageable host tensor on every forward, then copies it into the pinned staging buffer. Write the slots into the pinned buffer directly.♻️ Proposed refactor
n = len(slots) - self.state_indices_cpu[:n].copy_(torch.tensor(slots, dtype=torch.int32)) + for i, s in enumerate(slots): + self.state_indices_cpu[i] = s self.state_indices[:n].copy_(self.state_indices_cpu[:n], non_blocking=True)🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@tensorrt_llm/_torch/models/modeling_inkling.py` around lines 265 - 274, Update write_state_indices to populate the pinned state_indices_cpu staging buffer directly from slots, removing the per-forward torch.tensor allocation. Preserve the existing n-based slicing, non-blocking device copy into state_indices, and CUDA graph pointer validation.
1476-1504: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valueThe two decoder-layer paths diverge on
**kwargs.The pool path at line 1492 forwards
**kwargstoself.attn. The stateless path at line 1479 does not. Any kwarg a future caller adds reaches the attention module on one path only. Forward**kwargson both paths, or on neither.♻️ Proposed fix
- hidden_states = self.attn(position_ids, hidden_states, attn_metadata) + hidden_states = self.attn(position_ids, hidden_states, attn_metadata, **kwargs)🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@tensorrt_llm/_torch/models/modeling_inkling.py` around lines 1476 - 1504, Make `**kwargs` handling consistent between the `conv_rt is None` and runtime state-pool branches in the decoder layer: forward the same keyword arguments to `self.attn` in the stateless path as in the pool path, or remove forwarding from both paths. Preserve all existing attention and residual behavior.
1412-1427: 🚀 Performance & Scalability | 🔵 Trivial | ⚡ Quick winOne root cause:
InklingGate.routing_methodreturns a new object per access, so the joint renorm cannot be shared. Because the property constructs a freshInklingMoeRoutingMethodon every read, the instancecreate_moeholds cannot publish itsshared_gammasback toInklingMoE.forward, which is why the renorm runs a second time per layer.
tensorrt_llm/_torch/models/modeling_inkling.py#L1412-L1427: drop the secondinkling_joint_renormcall and readshared_gammasfrom the stored routing-method instance.tensorrt_llm/_torch/models/modeling_inkling.py#L1292-L1301: build theInklingMoeRoutingMethodonce inInklingGate.__init__and return the stored instance from the property, so it can carry the per-forwardshared_gammas.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@tensorrt_llm/_torch/models/modeling_inkling.py` around lines 1412 - 1427, Update tensorrt_llm/_torch/models/modeling_inkling.py lines 1412-1427 in InklingMoE.forward to remove the second inkling_joint_renorm call and obtain shared_gammas from the stored InklingMoeRoutingMethod instance. Update lines 1292-1301 in InklingGate.__init__ and its routing_method property to construct InklingMoeRoutingMethod once and return that instance, allowing shared_gammas to persist across the forward call.
969-991: 🚀 Performance & Scalability | 🔵 Trivial | 🏗️ Heavy lift
_build_rel_logitsmaterializes a large fp32 tensor per layer at prefill.The result is
[T, local_heads, rel_extent]in fp32, plus a.contiguous()copy. For a global layer withT = 8192, 16 local heads andrel_extent = 1024, that is about 512 MiB live per layer during prefill, and theeinsumoutput plus the contiguous copy can both be resident. Long-context prefill can therefore hit an out-of-memory condition that scales with the chunk size rather than with the model size.Two mitigations to consider: emit the tensor in bf16 if the Triton kernel accepts it, and make
einsumproduce a contiguous result directly so the extra copy disappears.♻️ Proposed change to drop the extra copy
- rel = torch.einsum( - "thd,de->the", r.float(), self.rel_logits_proj.float() - ) # [T, H, rel_extent] + # matmul on the flattened [T*H, d_rel] view returns a contiguous result, + # so the trailing .contiguous() never copies. + rel = (r.float().reshape(-1, self.d_rel) @ self.rel_logits_proj.float()).view( + -1, self.local_num_heads, self.rel_extent + )Please confirm the peak prefill memory against the configured
max_num_tokens.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@tensorrt_llm/_torch/models/modeling_inkling.py` around lines 969 - 991, Update _build_rel_logits to avoid materializing both the einsum output and a separate .contiguous() copy; make the einsum result contiguous directly where supported, and emit bf16 when the consuming Triton kernels accept it while preserving the required computation and scaling behavior. Verify peak prefill memory using the configured max_num_tokens rather than only the example chunk size.tensorrt_llm/_torch/attention_backend/inkling_triton.py (1)
53-58: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valueUpdate the
_NEGcomment._NEGis a module-leveltl.constexpr, and Triton 3.6.0 supports this usage. Remove the claim that the value is inlined because@triton.jitcannot read module-level globals.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@tensorrt_llm/_torch/attention_backend/inkling_triton.py` around lines 53 - 58, Update the comment above the module-level `_NEG` constant to remove the inaccurate claim that Triton `@jit` functions cannot read module-level globals or that the value is inlined for that reason. Preserve the existing explanation of its finite, sufficiently negative value and its role in avoiding NaNs during masked softmax bookkeeping.tests/unittest/llmapi/test_sampling_params.py (1)
329-333: 🎯 Functional Correctness | 🔵 Trivial | ⚡ Quick winAdd cases for a tuple EOS and an empty EOS list.
SamplingParams._setupbranches onisinstance(config_eos, (list, tuple))and then guards withif config_eos:. The tests cover only the non-empty list. The tuple path and the empty-sequence path are untested, and the empty-sequence guard is what prevents anIndexErroronconfig_eos[0].Coverage for the scalar, list, explicit-override, and missing-config paths is sufficient.
💚 Proposed tests
def test_setup_model_config_eos_tuple_is_handled(): params = SamplingParams() params._setup(_FakeTokenizer(eos_token_id=None), _FakeConfig(eos_token_id=(7, 8)), None) assert params.end_id == 7 assert params.stop_token_ids == [8] def test_setup_model_config_empty_eos_list_leaves_end_id_none(): params = SamplingParams() params._setup(_FakeTokenizer(eos_token_id=None), _FakeConfig(eos_token_id=[]), None) assert params.end_id is None🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@tests/unittest/llmapi/test_sampling_params.py` around lines 329 - 333, Add tests alongside test_setup_model_config_eos_list_sets_end_id_and_stop_tokens for the tuple EOS path and empty EOS sequence path: verify tuple (7, 8) sets end_id to 7 and stop_token_ids to [8], and verify an empty list leaves end_id as None without raising an error.tensorrt_llm/evaluate/lm_eval.py (1)
259-260: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valueForward
post_process_fnthroughsuper().__init__()for consistency.
keep_special_tokensis forwarded at Line 292, butpost_process_fnis set later at Line 301 instead. The base constructor already acceptspost_process_fn. The result is the same today, but the split assignment is easy to break during future edits.♻️ Proposed consistency change
preserve_caller_max_tokens=preserve_caller_max_tokens, + post_process_fn=post_process_fn, keep_special_tokens=keep_special_tokens, )Then remove the duplicate assignment at Line 301.
Also applies to: 279-280, 292-292
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@tensorrt_llm/evaluate/lm_eval.py` around lines 259 - 260, Update the constructor around the visible post_process_fn and keep_special_tokens parameters to pass post_process_fn through super().__init__() alongside keep_special_tokens, then remove the later duplicate post_process_fn assignment. Preserve the existing callback behavior while keeping all base-constructor initialization together.tensorrt_llm/serve/openai_server.py (1)
1514-1519: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winDo not read the factory's private
_parsersmap from the server.
ReasoningParserFactoryexposes onlycreate_reasoning_parser()andkeys()as public API. This code reaches into_parsersand depends on the entry being a tuple whose first element is the parser class. The adjacent tool-parser block uses the publicToolParserFactory.parsers.If the factory later changes the stored entry shape,
reasoning_entry[0]no longer yields the class.getattr(..., 'needs_raw_special_tokens', False)then returnsFalse,skip_special_tokensstaysTrue, and Inkling reasoning splitting silently stops working. The failure produces no error.Add a public lookup on the factory and call it here.
♻️ Proposed public accessor
In
tensorrt_llm/llmapi/reasoning_parser.py:`@classmethod` def needs_raw_special_tokens(cls, reasoning_parser: str) -> bool: """Return True if the named parser requires undecoded special tokens.""" entry = cls._parsers.get(reasoning_parser.lower()) if entry is None: return False return getattr(entry[0], "needs_raw_special_tokens", False)In this file:
if self.generator.args.reasoning_parser: - reasoning_entry = ReasoningParserFactory._parsers.get( - self.generator.args.reasoning_parser.lower()) - if reasoning_entry and getattr( - reasoning_entry[0], 'needs_raw_special_tokens', False): + if ReasoningParserFactory.needs_raw_special_tokens( + self.generator.args.reasoning_parser): sampling_params.skip_special_tokens = False🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@tensorrt_llm/serve/openai_server.py` around lines 1514 - 1519, Add the public ReasoningParserFactory.needs_raw_special_tokens() accessor in reasoning_parser.py, encapsulating the private _parsers lookup and parser-entry inspection. Update the server logic around reasoning_parser to call this accessor instead of reading _parsers directly, while preserving the existing skip_special_tokens behavior.tests/unittest/llmapi/test_reasoning_parser.py (1)
944-968: 🎯 Functional Correctness | 🔵 Trivial | ⚡ Quick winAdd a test that reuses one parser across two streaming sessions.
Every streaming test builds a fresh parser or calls
finish()once. No test drivesparse_delta→finish→parse_deltaon the same instance. That path leavesself._kindstale, as noted onInklingReasoningParser.finishintensorrt_llm/llmapi/reasoning_parser.py.Current coverage is otherwise sufficient for the batch and single-stream paths.
💚 Proposed test
def test_inkling_reasoning_parser_instance_reuse_after_finish(): """A second stream on the same parser must not inherit the previous block kind.""" parser = ReasoningParserFactory.create_reasoning_parser("inkling") parser.parse_delta(f"{INK_CH}first stream reasoning") parser.finish() r = parser.parse_delta("plain second stream answer") assert r.content == "plain second stream answer" assert r.reasoning_content == ""🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@tests/unittest/llmapi/test_reasoning_parser.py` around lines 944 - 968, Add a streaming reuse test alongside the existing Inkling parser tests that calls parse_delta, then finish, then parse_delta again on the same parser instance. Verify the second stream’s plain text is returned as content with empty reasoning_content, confirming finish resets the parser state such as _kind instead of carrying over the first stream’s reasoning block.tests/unittest/others/test_lm_eval.py (1)
856-867: 🎯 Functional Correctness | 🔵 Trivial | ⚡ Quick winAdd a case for the invalid
--post_process_fnvalue.The resolver test covers the three valid keys. It does not cover the
elsebranch in_resolve_post_process_fn, which raisesclick.BadParameter. That branch is the only validation for the GSM8K option, which is declared astype=strrather thanclick.Choice. It should be pinned.Coverage for the three supported processors is sufficient.
💚 Proposed test
def test_lm_eval_post_process_resolver_rejects_unknown_name(): with pytest.raises(click.BadParameter): _resolve_post_process_fn("not_a_processor") def test_lm_eval_post_process_resolver_passes_through_callables(): fn, keep_special_tokens = _resolve_post_process_fn(extract_inkling_content) assert fn is extract_inkling_content assert keep_special_tokens is FalseThe second case pins the non-string passthrough at the top of the resolver, which
LmEvalEvaluatorrelies on when a caller supplies a callable directly.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@tests/unittest/others/test_lm_eval.py` around lines 856 - 867, Extend the resolver tests around _resolve_post_process_fn with a pytest.raises(click.BadParameter) case for an unknown string such as "not_a_processor", covering the invalid-name branch. Also add a callable passthrough case using extract_inkling_content and assert it is returned with keep_special_tokens set to False.tensorrt_llm/llmapi/reasoning_parser.py (2)
719-746: 🎯 Functional Correctness | 🔵 Trivial | 🏗️ Heavy liftThe Inkling control-token format is implemented twice, and the two implementations disagree.
InklingReasoningParser._consumeandextract_inkling_contenteach define their own copy of the control-token alphabet and their own block state machine for the same wire format. Two divergences already exist: tool-invocation blocks route to visible content in the serving parser but are dropped in the offline extractor, and text before the first control token is kept as content in the serving parser but dropped in the offline extractor. The offline path is documented as "the offline analog" of the serving path, so these must agree. Neither divergence affects GSM8K, MMLU, or MMMU, because those generations contain no tool blocks and no leading unmarked text.
tensorrt_llm/llmapi/reasoning_parser.py#L719-L746: extract the token alphabet and the block-classification state machine into one shared helper, then have_consumecall it and map the resulting blocks ontocontentandreasoning. Keep the tool-block-to-content routing as the single agreed rule.tensorrt_llm/evaluate/post_processing.py#L164-L208: delete the duplicated_INK_*constants and_INK_CONTROL_RE, import the shared alphabet, and makeextract_inkling_contentreuse the shared state machine so tool blocks and leading unmarked text are classified identically to the serving path.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@tensorrt_llm/llmapi/reasoning_parser.py` around lines 719 - 746, The duplicated Inkling control-token alphabet and block state machine must be unified. In tensorrt_llm/llmapi/reasoning_parser.py lines 719-746, extract the shared alphabet and classification logic into a helper, then have InklingReasoningParser._consume map its blocks to content or reasoning while preserving tool blocks as visible content. In tensorrt_llm/evaluate/post_processing.py lines 164-208, remove the local _INK_* constants and _INK_CONTROL_RE and make extract_inkling_content reuse that helper so leading unmarked text and tool blocks follow the same classification.
757-771: 🚀 Performance & Scalability | 🔵 Trivial | 💤 Low valueConsider a precomputed prefix set for the streaming holdback.
_partial_control_lengthruns on every streaming delta. It performs up to_INKLING_MAX_CONTROL_LEN - 1iterations, and each iteration scans all 18 control tokens. That is roughly 500startswithcalls per delta per request.A module-level set of all strict prefixes turns the check into a hash lookup per length. Build it once next to
_INKLING_CONTROL_TOKENS.♻️ Proposed optimization
_INKLING_MAX_CONTROL_LEN = max(len(t) for t in _INKLING_CONTROL_TOKENS) +# Every strict prefix of every control token, for the streaming holdback. +_INKLING_CONTROL_PREFIXES = frozenset( + t[:i] for t in _INKLING_CONTROL_TOKENS for i in range(1, len(t)))max_len = min(len(text), _INKLING_MAX_CONTROL_LEN - 1) for length in range(max_len, 0, -1): - suffix = text[-length:] - if any( - len(suffix) < len(tok) and tok.startswith(suffix) - for tok in _INKLING_CONTROL_TOKENS): + if text[-length:] in _INKLING_CONTROL_PREFIXES: return length return 0🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@tensorrt_llm/llmapi/reasoning_parser.py` around lines 757 - 771, Optimize _partial_control_length by defining a module-level set containing every strict prefix of the tokens in _INKLING_CONTROL_TOKENS, initialized alongside that token collection. Replace the per-length any(...startswith...) scan with membership checks against the precomputed prefix set while preserving the longest-suffix and zero-result behavior.
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
Inline comments:
In `@docs/source/models/supported-models.md`:
- Line 97: Update the Inkling footnote’s accuracy-gate description to state that
the text decoder is gated by both GSM8K and MMLU, while retaining MMMU as the
vision-path gate and leaving the remaining support details unchanged.
In `@tensorrt_llm/_torch/attention_backend/inkling_triton.py`:
- Around line 337-339: Update the affected attention paths, including
inkling_decode_attention, to avoid unpacking the unused total_tokens value and
validate that num_heads is evenly divisible by num_kv_heads before computing
kv_group_num. Raise the established appropriate error for invalid head
configurations, while preserving the existing mapping for divisible inputs.
In `@tensorrt_llm/_torch/configs/inkling.py`:
- Around line 121-124: Update the initializer around `moe_intermediate_size` to
read the explicitly supplied `moe_intermediate_size` kwarg before assigning its
default, following the existing `max_position_embeddings` handling; preserve the
checkpoint value when present and fall back to `intermediate_size` only when it
is absent.
In `@tensorrt_llm/_torch/model_config.py`:
- Around line 387-400: Update the MIXED_PRECISION branch’s kv_cache_quant_algo
re-derivation to call the existing _algo_or_none helper instead of constructing
QuantAlgo directly. Preserve the helper’s normalization of None, "none", "null",
and empty strings while retaining normal enum conversion for valid algorithm
values.
In `@tensorrt_llm/_torch/models/checkpoints/hf/inkling_weight_mapper.py`:
- Around line 260-268: Update the mapper initialization or configuration
handling used by ModelLoader.load and
InklingForConditionalGeneration.load_weights so preprocess_weights and
_map_expert resolve the text sub-model configuration via
_text_sub_model_config(model_config). Ensure vocab_size and related text-model
fields come from that resolved config, including on the supplied-mapper path.
In `@tensorrt_llm/_torch/models/modeling_inkling_multimodal.py`:
- Around line 1046-1073: In __init__, after resolving image_token_id and
audio_token_id via _resolve_image_token_id and _resolve_audio_token_id, validate
that they are distinct and reject any collision with a clear configuration error
before placeholder processing can run. Keep _check_placeholder_count and the
multimodal expansion loop unchanged.
- Around line 396-408: Update plan_out_scales and its dependency configuration
so the linear_sum_assignment branch does not rely on an undeclared runtime scipy
import: either declare scipy in the runtime requirements consumed by setup.py,
or replace the assignment logic with an equivalent dependency-free
implementation. Preserve the existing scale-index selection and endpoint
behavior.
In `@tensorrt_llm/_torch/models/modeling_inkling.py`:
- Around line 783-795: Refactor the decode preparation around refresh and
prepare_inkling_attn_decode so seq_lens is computed and its host/device buffer
is allocated and copied once per generation step, then shared across all layers
instead of rebuilt for each layer. Replace the per-row page-table construction
loop and torch.tensor allocations with a single vectorized fill while preserving
padding and invalid-block handling. Benchmark decode-step latency before and
after the change.
- Around line 763-797: Bind the published ready state in refresh to a
per-forward token containing the current cache metadata, including
num_cached_tokens_per_seq, rather than only request IDs and batch size. Update
PyTorchModelEngine and InklingForCausalLM.forward so each forward invalidates or
validates that token before reading seq_lens or page_table, and ensure the eager
fallback cannot consume stale buffers when refresh is skipped.
In `@tensorrt_llm/evaluate/lm_eval.py`:
- Around line 805-812: Update the MMLU command in tensorrt_llm/evaluate/mmlu.py
to declare and propagate --post_process_fn, applying the selected processor to
outputs before compute_score(). Also change the GSM8K option declaration near
LmEvalEvaluator.command_harness to use click.Choice(["strip_thinking_mmmu",
"inkling", "inkling_mmmu"]) instead of type=str.
In `@tensorrt_llm/llmapi/reasoning_parser.py`:
- Around line 787-796: Update the finish() method to reset self._kind alongside
self._buffer before handling remaining content, matching parse() so reused
parser instances start each stream without stale block state.
In `@tests/integration/defs/accuracy/test_llm_api_pytorch_multimodal.py`:
- Around line 710-713: Update the validation description near the vision-path
reference to state that Inkling is validated against the reference using
aggregate accuracy, not item-for-item matching; preserve the existing
TensorRT-LLM/SGLang baseline context.
In `@tests/integration/defs/accuracy/test_llm_api_pytorch.py`:
- Line 7843: The test_nvfp4 methods in
tests/integration/defs/accuracy/test_llm_api_pytorch.py:7843-7843 and
tests/integration/defs/accuracy/test_llm_api_pytorch_multimodal.py:707-707 both
require explicit return-type annotations. Update each test_nvfp4 definition to
include -> None.
- Around line 7836-7840: Annotate both shared EXTRA_EVALUATOR_KWARGS mappings
with ClassVar[...] in
tests/integration/defs/accuracy/test_llm_api_pytorch.py:7836-7840 and
tests/integration/defs/accuracy/test_llm_api_pytorch_multimodal.py:700-704,
preserving their existing dictionary contents and evaluate() behavior.
In `@tests/unittest/llmapi/test_sampling_params.py`:
- Around line 342-345: Update
test_setup_without_any_eos_warns_and_leaves_end_id_none to match its actual
coverage: either configure capture for tensorrt_llm.logger and assert the
expected warning through caplog, or remove the unused caplog fixture and rename
the test to describe only the end_id-is-None behavior.
---
Nitpick comments:
In `@tensorrt_llm/_torch/attention_backend/inkling_triton.py`:
- Around line 53-58: Update the comment above the module-level `_NEG` constant
to remove the inaccurate claim that Triton `@jit` functions cannot read
module-level globals or that the value is inlined for that reason. Preserve the
existing explanation of its finite, sufficiently negative value and its role in
avoiding NaNs during masked softmax bookkeeping.
In `@tensorrt_llm/_torch/configs/inkling.py`:
- Around line 137-148: The _local_ids property should use the built-in generic
annotation set[int] and avoid rebuilding the set for every is_local_layer call.
Precompute or cache the converted local_layer_ids set at the configuration
level, then have is_local_layer reuse that cached set while preserving its
current membership behavior.
- Around line 195-248: Add type annotations to InklingConfig.__init__ parameters
text_config, audio_config, vision_config, and mtp_config using dict[str, Any] |
PretrainedConfig | None, and annotate _as_config(value) with the same input type
and an appropriate return type. Add the required Any import, and annotate
__init__ as returning None while preserving the existing configuration handling.
In `@tensorrt_llm/_torch/models/modeling_inkling_multimodal.py`:
- Around line 1021-1028: Update the return annotations of
get_num_tokens_per_image, get_num_tokens_per_audio, and get_num_tokens_per_video
from int to typing.NoReturn, ensuring NoReturn is imported or referenced through
the existing typing conventions.
In `@tensorrt_llm/_torch/models/modeling_inkling.py`:
- Around line 265-274: Update write_state_indices to populate the pinned
state_indices_cpu staging buffer directly from slots, removing the per-forward
torch.tensor allocation. Preserve the existing n-based slicing, non-blocking
device copy into state_indices, and CUDA graph pointer validation.
- Around line 1476-1504: Make `**kwargs` handling consistent between the
`conv_rt is None` and runtime state-pool branches in the decoder layer: forward
the same keyword arguments to `self.attn` in the stateless path as in the pool
path, or remove forwarding from both paths. Preserve all existing attention and
residual behavior.
- Around line 1412-1427: Update tensorrt_llm/_torch/models/modeling_inkling.py
lines 1412-1427 in InklingMoE.forward to remove the second inkling_joint_renorm
call and obtain shared_gammas from the stored InklingMoeRoutingMethod instance.
Update lines 1292-1301 in InklingGate.__init__ and its routing_method property
to construct InklingMoeRoutingMethod once and return that instance, allowing
shared_gammas to persist across the forward call.
- Around line 969-991: Update _build_rel_logits to avoid materializing both the
einsum output and a separate .contiguous() copy; make the einsum result
contiguous directly where supported, and emit bf16 when the consuming Triton
kernels accept it while preserving the required computation and scaling
behavior. Verify peak prefill memory using the configured max_num_tokens rather
than only the example chunk size.
In `@tensorrt_llm/evaluate/lm_eval.py`:
- Around line 259-260: Update the constructor around the visible post_process_fn
and keep_special_tokens parameters to pass post_process_fn through
super().__init__() alongside keep_special_tokens, then remove the later
duplicate post_process_fn assignment. Preserve the existing callback behavior
while keeping all base-constructor initialization together.
In `@tensorrt_llm/llmapi/reasoning_parser.py`:
- Around line 719-746: The duplicated Inkling control-token alphabet and block
state machine must be unified. In tensorrt_llm/llmapi/reasoning_parser.py lines
719-746, extract the shared alphabet and classification logic into a helper,
then have InklingReasoningParser._consume map its blocks to content or reasoning
while preserving tool blocks as visible content. In
tensorrt_llm/evaluate/post_processing.py lines 164-208, remove the local _INK_*
constants and _INK_CONTROL_RE and make extract_inkling_content reuse that helper
so leading unmarked text and tool blocks follow the same classification.
- Around line 757-771: Optimize _partial_control_length by defining a
module-level set containing every strict prefix of the tokens in
_INKLING_CONTROL_TOKENS, initialized alongside that token collection. Replace
the per-length any(...startswith...) scan with membership checks against the
precomputed prefix set while preserving the longest-suffix and zero-result
behavior.
In `@tensorrt_llm/serve/openai_server.py`:
- Around line 1514-1519: Add the public
ReasoningParserFactory.needs_raw_special_tokens() accessor in
reasoning_parser.py, encapsulating the private _parsers lookup and parser-entry
inspection. Update the server logic around reasoning_parser to call this
accessor instead of reading _parsers directly, while preserving the existing
skip_special_tokens behavior.
In `@tests/integration/defs/accuracy/test_llm_api_pytorch.py`:
- Line 7815: Rename both TestInkling_NVFP4 classes to TestInklingNVFP4 in
tests/integration/defs/accuracy/test_llm_api_pytorch.py:7815-7815 and
tests/integration/defs/accuracy/test_llm_api_pytorch_multimodal.py:683-683.
Update both direct class selectors to TestInklingNVFP4 in
tests/integration/test_lists/qa/llm_function_core.txt:812-813 and
tests/integration/test_lists/test-db/l0_b200.yml:355-356.
In `@tests/unittest/_torch/modeling/test_modeling_inkling_multimodal.py`:
- Around line 173-181: Update test_dmel_preprocess_shape_and_range and
test_dmel_multi_clip_concat_and_empty_clip to derive the expected mel-bin width
and value range from DEFAULT_AUDIO_N_MELS and DEFAULT_AUDIO_NUM_DMEL_BINS
instead of hardcoded 80 and 16. Keep the existing assertions and behavior
unchanged while referencing the module constants directly.
- Around line 124-167: Add a CPU-focused test near the existing vision input
tests covering _to_pil_rgb: assert http://, https://, and data: inputs raise
ValueError matching “resolve”, unsupported objects raise TypeError, and a
generated image is accepted with RGB mode. This closes the input-coercion
coverage gap while preserving the existing tower tests.
- Around line 290-306: Update test_registration to validate processor
registration through the public lookup API exposed by INPUT_PROCESSOR_REGISTRY
instead of accessing _input_processors_cls_by_model_type. Also avoid asserting
InklingInputProcessor._registered_model_type directly; verify the expected
"inkling_mm_model" association through the public registry interface while
preserving the existing registration expectations.
In `@tests/unittest/_torch/modeling/test_modeling_inkling.py`:
- Around line 52-68: Cache the parsed weight map used by _safetensors_shape so
model.safetensors.index.json is read and parsed only once per checkpoint. Apply
a module-scoped fixture or functools.lru_cache to _index, preserving
_safetensors_shape’s existing lookup behavior.
- Around line 122-141: Add a CPU unit test in the test file for
_split_interleaved_gate_up, using an even-row tensor to verify gate receives
even-indexed rows and up receives odd-indexed rows, and assert that an odd-sized
split dimension raises ValueError. Import the mapper helper and pytest as
needed; no checkpoint setup is required.
In `@tests/unittest/llmapi/test_reasoning_parser.py`:
- Around line 944-968: Add a streaming reuse test alongside the existing Inkling
parser tests that calls parse_delta, then finish, then parse_delta again on the
same parser instance. Verify the second stream’s plain text is returned as
content with empty reasoning_content, confirming finish resets the parser state
such as _kind instead of carrying over the first stream’s reasoning block.
In `@tests/unittest/llmapi/test_sampling_params.py`:
- Around line 329-333: Add tests alongside
test_setup_model_config_eos_list_sets_end_id_and_stop_tokens for the tuple EOS
path and empty EOS sequence path: verify tuple (7, 8) sets end_id to 7 and
stop_token_ids to [8], and verify an empty list leaves end_id as None without
raising an error.
In `@tests/unittest/others/test_lm_eval.py`:
- Around line 856-867: Extend the resolver tests around _resolve_post_process_fn
with a pytest.raises(click.BadParameter) case for an unknown string such as
"not_a_processor", covering the invalid-name branch. Also add a callable
passthrough case using extract_inkling_content and assert it is returned with
keep_special_tokens set to False.
🪄 Autofix (Beta)
Fix all unresolved CodeRabbit comments on this PR:
- Push a commit to this branch (recommended)
- Create a new PR with the fixes
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 43b4b5e1-4a1e-4209-adbf-ee6275b1e8ef
📒 Files selected for processing (33)
docs/source/models/supported-models.mdtensorrt_llm/_torch/attention_backend/inkling_triton.pytensorrt_llm/_torch/configs/__init__.pytensorrt_llm/_torch/configs/inkling.pytensorrt_llm/_torch/model_config.pytensorrt_llm/_torch/models/__init__.pytensorrt_llm/_torch/models/checkpoints/__init__.pytensorrt_llm/_torch/models/checkpoints/hf/inkling_weight_mapper.pytensorrt_llm/_torch/models/modeling_inkling.pytensorrt_llm/_torch/models/modeling_inkling_multimodal.pytensorrt_llm/_torch/pyexecutor/_util.pytensorrt_llm/_torch/pyexecutor/config_utils.pytensorrt_llm/_torch/pyexecutor/model_engine.pytensorrt_llm/_torch/pyexecutor/py_executor_creator.pytensorrt_llm/_torch/pyexecutor/resource_manager.pytensorrt_llm/evaluate/lm_eval.pytensorrt_llm/evaluate/post_processing.pytensorrt_llm/llmapi/llm_utils.pytensorrt_llm/llmapi/reasoning_parser.pytensorrt_llm/sampling_params.pytensorrt_llm/serve/openai_server.pytests/integration/defs/accuracy/references/gsm8k.yamltests/integration/defs/accuracy/references/mmlu.yamltests/integration/defs/accuracy/references/mmmu.yamltests/integration/defs/accuracy/test_llm_api_pytorch.pytests/integration/defs/accuracy/test_llm_api_pytorch_multimodal.pytests/integration/test_lists/qa/llm_function_core.txttests/integration/test_lists/test-db/l0_b200.ymltests/unittest/_torch/modeling/test_modeling_inkling.pytests/unittest/_torch/modeling/test_modeling_inkling_multimodal.pytests/unittest/llmapi/test_reasoning_parser.pytests/unittest/llmapi/test_sampling_params.pytests/unittest/others/test_lm_eval.py
Brings the branch up to 93e7c91 (402 upstream commits since 44f0521). Six files conflicted; all resolutions keep both sides: * supported-models.md upstream took footnote [^14] for MiniCPM-V, so the Inkling note moves to [^15] (3 table rows renumbered) * _util.py union of the import blocks; the new is_hybrid_linear V2 guard sits alongside the Inkling one * model_engine.py upstream's steady-gen fast path and the Inkling runtime publish are independent additions; upstream also fixed the logits self-assignment defect by dropping the redundant write-back, which supersedes the clone() this branch carried * py_executor_creator union of the config_utils imports * lm_eval.py six additive-kwarg conflicts, union throughout * test_lm_eval.py both sides appended an independent test section One resolution is not textual. Upstream's new _apply_steady_gen_fast_prepare returns early from _prepare_tp_inputs, bypassing the tail where Inkling publishes its short-conv pool rows and per-layer attention decode metadata. git reports no conflict there -- the two sides touch different lines -- but an unchanged generation-only batch is exactly Inkling's steady-state decode, so a plain merge would have left the captured graph reading capture-time page tables. The fast path now publishes before returning; it is a no-op for models without a registered CONV_STATE_MANAGER. Also initializes the new 3rdparty/MSA submodule (upstream NVIDIA#16291), which setup.py now requires. Signed-off-by: KleinBlueC <102355066+KleinBlueC@users.noreply.github.com>
…d it Two review findings on InklingDecodeMeta, both on the decode critical path. **Stale publication.** `ready` was a plain flag that only `refresh` cleared, so it described "some forward published these buffers", not "this forward did". A forward that reached the decode kernel without the engine having prepared -- an eager fallback, or a prepare path that returns early -- read the previous step's seq_lens and page table. The same requests advance num_cached_tokens_per_seq every step, so a page table one step out of date silently drops a newly allocated page. Request ids and the batch size cannot detect that. `prepare_inkling_attn_decode` now bumps a shared per-forward counter before doing anything else, and `ready` is a property comparing the layer's recorded epoch against it. A layer that is not refreshed this forward reports False by construction, so the attention takes its host fallback instead of reading stale buffers. This is the same contract the upstream merge broke once already (a new fast path bypassed the publish site); making the invariant self-checking means the next such refactor fails loudly rather than silently. **Per-layer work on every token.** `refresh` ran for all 66 layers per decode step, and each call allocated two pinned host tensors and built the page table with one `torch.tensor` per row -- 132 cudaHostAllocs and 66*num_gen tensor constructions per token. The total-KV lengths do not depend on the layer, so they are now built once per step and shared; the per-layer pinned page-table staging is allocated once and reused, and filled through its numpy view with a vectorized row assignment. Behavior is unchanged when the engine publishes every forward, which is the runtime path. Signed-off-by: KleinBlueC <102355066+KleinBlueC@users.noreply.github.com>
Three review findings on the way a checkpoint is read. The weight mapper resolved its geometry from whatever config it was handed. ``ModelLoader.load`` initializes it with the TOP-LEVEL ``InklingConfig``, which keeps the decoder geometry under ``text_config`` and has no ``vocab_size`` or ``n_routed_experts`` of its own, so ``preprocess_weights`` raised ``AttributeError`` on the unpadded-vocab lookup and ``_map_expert`` silently fell back to ``tensor.shape[0]`` for the expert count. The mapper now resolves ``text_config`` itself, which covers both entry points: the loader-supplied mapper and the one ``InklingForConditionalGeneration.load_weights`` builds. ``InklingTextConfig`` unconditionally set ``moe_intermediate_size`` from ``intermediate_size``, discarding the value a ``config.json`` had already stored through ``**kwargs``. It now reads the kwarg back first, matching how ``max_position_embeddings`` is handled two lines above. ``load_modelopt_quant_config`` gained ``_algo_or_none`` to map both JSON null and the strings "none"/"null" onto ``None``, but the MIXED_PRECISION branch still built ``QuantAlgo`` directly, so a mixed-precision checkpoint spelling ``kv_cache_quant_algo: "none"`` would still raise. That branch uses the helper now. Signed-off-by: KleinBlueC <102355066+KleinBlueC@users.noreply.github.com>
Each of these turned a bad input into wrong output rather than an error. The Triton kernels map a query head to its KV head as ``cur_head // kv_group_num``. Both entry points computed ``kv_group_num = num_heads // num_kv_heads`` without checking divisibility, so a non-divisible pair produced a wrong mapping inside the kernel with no diagnostic. Both now assert. The prefill entry also unpacked ``total_tokens`` without using it (ruff RUF059); renamed to ``_total_tokens``. ``InklingInputProcessor.assemble`` dispatches on the token value, and both placeholder ids are config-overridable. If a config gave them the same value every occurrence was consumed as an image, the audio branch never ran, and the placeholder-count check reported the mismatch against the wrong modality. The constructor now rejects the collision. ``InklingReasoningParser.finish`` cleared ``_buffer`` but left ``_kind`` at the last block kind, so reusing an instance for a second stream would route that stream's first segment by the previous stream's channel. ``parse`` already resets both. Not reachable today (the factory builds one parser per request), but the asymmetry is the kind that survives a refactor. Signed-off-by: KleinBlueC <102355066+KleinBlueC@users.noreply.github.com>
The supported-models footnote named GSM8K and MMMU as the accuracy gates, but TestInkling_NVFP4 also evaluates MMLU on the text decoder. Name both. The MMMU test docstring claimed Inkling was validated against SGLang "item-for-item". It is not: 74 of the 858 pool items score differently, in both directions, and the harness gates the aggregate score. Say what is actually validated, and why per-item identity is the wrong metric here -- two different NVFP4 implementations resolve near-ties differently deep in a long-CoT decode. Signed-off-by: KleinBlueC <102355066+KleinBlueC@users.noreply.github.com>
…ion tower `plan_out_scales` imported `scipy.optimize.linear_sum_assignment` to place the hMLP scales, but neither requirements.txt nor pyproject.toml declares scipy, and setup.py derives install_requires from requirements.txt. With the checkpoint geometry (patch 40 -> four prime factors, temporal 2 -> one, so six candidate scales against four layers) that branch runs on every InklingVisionModel construction, so a clean install would ImportError on the first vision request. The cost matrix is |a_i - b_j| with both sequences strictly increasing, which makes it Monge: an optimal assignment is guaranteed to be non-crossing, so an O(rows*cols) DP is exact. Verified against scipy over 447 geometries -- 443 identical, and the four that differ have the same total cost, where scipy returned a crossing permutation. Non-crossing is what this caller needs anyway: it consumes the result as an ordered scale progression, and a crossing assignment yields a non-monotonic plan whose fold does not divide. Adds a test over the assignment branch asserting the plan stays strictly growing and integrally foldable, alongside the existing test that pins the checkpoint's exact progression. Also renames test_setup_without_any_eos_warns_and_leaves_end_id_none: it took a caplog fixture it never inspected, so the name promised an assertion the body did not make. Signed-off-by: KleinBlueC <102355066+KleinBlueC@users.noreply.github.com>
d7b7931 to
3d079c0
Compare
|
/bot run --disable-fail-fast |
|
PR_Github #63406 [ run ] triggered by Bot. Commit: |
|
PR_Github #63406 [ run ] completed with state
|
|
/bot run --disable-fail-fast |
|
PR_Github #63448 [ run ] triggered by Bot. Commit: |
|
PR_Github #63448 [ run ] completed with state
|
|
/bot run --disable-fail-fast |
|
PR_Github #63478 [ run ] triggered by Bot. Commit: |
|
PR_Github #63478 [ run ] completed with state
|
Summary by CodeRabbit
New Features
Bug Fixes
Documentation
Description
Adds PyTorch-backend support for Inkling (
thinkingmachines/Inkling-NVFP4), aRoPE-free hybrid-attention MoE reasoning model with vision and audio towers.
Registered as
InklingForConditionalGeneration.Architecture
pre-softmax, per (query token, head, relative distance).
interleaved with 11 global full-causal layers (8 KV heads).
on the residual stream, carrying per-request state across decode steps.
bias, then a log-sigmoid renorm spanning the routed and two shared-expert logits.
logits_mup_width_multiplierbefore the head; logits sliced from 201024 down to 200058.
fusing one row per patch / per frame into the text embedding stream. Video is
multi-frame images; there is no separate video encoder.
NVFP4 covers the text decoder and routed experts (layers 3–65); layer-2 experts,
attention, shared experts, and both towers stay BF16 per the checkpoint's
hf_quant_config.json.Implementation
modeling_inkling.pyattention_backend/inkling_triton.pyscore_mod. No existing fused backend exposes such a hook (context FMHA is disabled forkRELATIVE, trtllm-gen rejects a relative bias, FlashInfer has no additive per-token bias).rel_logitsis static-shape, so the decode kernel is CUDA-graph capturablemodeling_inkling_multimodal.py<image>placeholder per patch and one<audio>per dMel frame. Pure numpy + torch preprocessingconfigs/inkling.pycheckpoints/hf/inkling_weight_mapper.pypyexecutor/*CONV_STATE_MANAGER. Pool rows and attention decode metadata are published into stable CUDA buffers eagerly from input prep, so the captured decode forward does no host→device copy and each replay reads the current batch (same stable-pointer pattern asMamba2Metadata)KV cache: the per-layer KV-head split (local 16 / global 8) structurally requires
KVCacheManagerV2— V1's unified pool would coerce it to one value and mis-size theper-layer KV bytes, a correctness bug. The model defaults to V2 and the incompatible
path raises rather than silently downgrading.
Serving / eval plumbing (small, each needed end to end):
--reasoning_parser inklingfor Inkling's typed-content blocks, with streaming.needs_raw_special_tokensso a delimiter-based reasoning parser actually sees itsmarkers (previously only the tool-parser path preserved them).
end_idfalls back to the config'seos_token_id— checkpoints whose terminatorlives only in
config.jsonotherwise never stop and run tomax_tokens.hf_quant_config.jsonmay spell "no quantization" as the string"none", whichpreviously reached
QuantAlgo("none")and raised.vocab otherwise fails KV-cache estimation with "Token ID out of range".
--post_process_fn inkling/inkling_mmmufor trtllm-eval.Accuracy
Measured on the complete datasets, TRT-LLM and SGLang side by side at TP=4 with
CUDA graph and the overlap scheduler on, batch 8, driven by the same client so
prompt rendering and scoring are shared code:
Not supported in this release
MTP / speculative decoding (the checkpoint ships next-N draft weights; nothing
builds or loads them), LoRA, function calling, constrained/guided decoding, EPD
disaggregated serving, and multimodal-hash prefix caching (refused loudly per
modality, so multimodal requests still run — just uncached).
One workaround worth flagging for review: every Inkling all-reduce is rebuilt with
ONESHOTafter construction. Under CUDA-graph capture a symmetric all-reducecorrupts the run when its send buffer is unregistered while its recv buffer is a
registered NCCL window at a 12288 B message — which Inkling hits exactly (hidden
6144, bf16, one decode token), sending the first global-attention layer non-finite.
The all-reduces involved are built by generic modules (attention
o_proj, MoEdown_proj), so pinning the strategy afterwards keeps the mitigation model-local.Test Coverage
Accuracy (integration) — added to
l0_b200andllm_function_core:accuracy/test_llm_api_pytorch.py::TestInkling_NVFP4::test_nvfp4— GSM8K + MMLUon the text decoder.
accuracy/test_llm_api_pytorch_multimodal.py::TestInkling_NVFP4::test_nvfp4—MMMU on the vision path.
Unit — CPU-only, no checkpoint / GPU / network needed, all well under a second:
unittest/_torch/modeling/test_modeling_inkling.py— config parsing,registration, per-layer classification; plus checkpoint-gated weight accounting
and tensor-shape checks that read only the safetensors index and skip cleanly
when the checkpoint is absent.
unittest/_torch/modeling/test_modeling_inkling_multimodal.py— the three mediapaths and the input processor on synthetic configs: hMLP scale plan and module
tree, the fold's value preservation, dMel preprocessing, the audio codebook
forward against a reference, frame sampling, and the fail-loud placeholder
contract.
unittest/llmapi/test_reasoning_parser.py— full-parse and streaming equivalencefor the Inkling parser, including control tokens split across delta boundaries.
unittest/llmapi/test_sampling_params.py— theend_idconfig fallback.unittest/others/test_lm_eval.py— the offline post-processing hook.PR Checklist
Please review the following before submitting your PR:
PR description clearly explains what and why. If using CodeRabbit's summary, please make sure it makes sense.
PR Follows TRT-LLM CODING GUIDELINES to the best of your knowledge.
Test cases are provided for new code paths (see test instructions)
If PR introduces API changes, an appropriate PR label is added - either
api-compatibleorapi-breaking. Forapi-breaking, includeBREAKINGin the PR title.Any new dependencies have been scanned for license and vulnerabilities
CODEOWNERS updated if ownership changes
Documentation updated as needed
Update tava architecture diagram if there is a significant design change in PR.
The reviewers assigned automatically/manually are appropriate for the PR.
Please check this after reviewing the above items as appropriate for this PR.
GitHub Bot Help
To see a list of available CI bot commands, please comment
/bot help.