Update TRT-LLM code by kaiyux · Pull Request #3 · NVIDIA/TensorRT-LLM

kaiyux · 2023-09-28T16:00:53Z

No description provided.

# This is the 1st commit message: add download models form www.modelscope.cn # This is the commit message NVIDIA#2: debug # This is the commit message NVIDIA#3: debug

Update TRT-LLM code

Add support for CPP inference with decoder

# This is the 1st commit message: kernel Signed-off-by: Ubuntu <dafrimi@nvidia.com> wip Signed-off-by: Ubuntu <dafrimi@nvidia.com> remove prints Signed-off-by: Ubuntu <dafrimi@nvidia.com> test pass Signed-off-by: Ubuntu <dafrimi@nvidia.com> test refactor with more use cases Signed-off-by: Ubuntu <dafrimi@nvidia.com> refacor Signed-off-by: Ubuntu <dafrimi@nvidia.com> refacor_2 Signed-off-by: Ubuntu <dafrimi@nvidia.com> add tuner wip Signed-off-by: Ubuntu <dafrimi@nvidia.com> autotuner works Signed-off-by: Ubuntu <dafrimi@nvidia.com> bfloat16 works. moer changes to the thop file Signed-off-by: Ubuntu <dafrimi@nvidia.com> is tune for autotuner is True --> gets real tactics configs Signed-off-by: Ubuntu <dafrimi@nvidia.com> wip Signed-off-by: Ubuntu <dafrimi@nvidia.com> wip Signed-off-by: Ubuntu <dafrimi@nvidia.com> zeros + quant mode is works Signed-off-by: Ubuntu <dafrimi@nvidia.com> act int8 Signed-off-by: Ubuntu <dafrimi@nvidia.com> removed fp8 for now Signed-off-by: Ubuntu <dafrimi@nvidia.com> wip Signed-off-by: Ubuntu <dafrimi@nvidia.com> w4a16 linear module Signed-off-by: Ubuntu <dafrimi@nvidia.com> wip Signed-off-by: Ubuntu <dafrimi@nvidia.com> changed cutalss for sm==89 Signed-off-by: Ubuntu <dafrimi@nvidia.com> wip Signed-off-by: Ubuntu <dafrimi@nvidia.com> test linear work Signed-off-by: Ubuntu <dafrimi@nvidia.com> add license Signed-off-by: Ubuntu <dafrimi@nvidia.com> works! Signed-off-by: Ubuntu <dafrimi@nvidia.com> refactor + linear test pass Signed-off-by: Ubuntu <dafrimi@nvidia.com> preprocess in load weights Signed-off-by: Ubuntu <dafrimi@nvidia.com> wip Signed-off-by: Ubuntu <dafrimi@nvidia.com> wip Signed-off-by: Ubuntu <dafrimi@nvidia.com> wip Signed-off-by: Ubuntu <dafrimi@nvidia.com> wip Signed-off-by: Ubuntu <dafrimi@nvidia.com> refactor + rebase Signed-off-by: Ubuntu <dafrimi@nvidia.com> wip Signed-off-by: Ubuntu <dafrimi@nvidia.com> wip Signed-off-by: Ubuntu <dafrimi@nvidia.com> Blackwell not supported Signed-off-by: Daniel Afrimi <dafrimi@nvidia.com> wip Signed-off-by: Daniel Afrimi <dafrimi@nvidia.com> skip blackwell Signed-off-by: Daniel Afrimi <dafrimi@nvidia.com> wip Signed-off-by: Daniel Afrimi <dafrimi@nvidia.com> works Signed-off-by: Ubuntu <dafrimi@nvidia.com> # This is the commit message NVIDIA#2: rebased Signed-off-by: Ubuntu <dafrimi@nvidia.com> # This is the commit message NVIDIA#3: align with my pld worked version of linear Signed-off-by: Ubuntu <dafrimi@nvidia.com> # This is the commit message NVIDIA#4: wip Signed-off-by: Ubuntu <dafrimi@nvidia.com> # This is the commit message NVIDIA#5: refactor Signed-off-by: Daniel Afrimi <danielafrimi8@gmail.com> # This is the commit message NVIDIA#6: refactor Signed-off-by: Daniel Afrimi <danielafrimi8@gmail.com> # This is the commit message NVIDIA#7: refactor Signed-off-by: Daniel Afrimi <danielafrimi8@gmail.com> # This is the commit message NVIDIA#8: refactor Signed-off-by: Daniel Afrimi <danielafrimi8@gmail.com> # This is the commit message NVIDIA#9: sys path Signed-off-by: Daniel Afrimi <danielafrimi8@gmail.com> # This is the commit message NVIDIA#10: sys path Signed-off-by: Daniel Afrimi <danielafrimi8@gmail.com>

- Revert utils.py docstring changes (NVIDIA#1) - Remove disable_deep_gemm flag, use module-level quant_config=None (NVIDIA#2) - Replace Flux2SwiGLU with shared swiglu from _torch/modules (NVIDIA#3) - Unify Flux2PosEmbed into FluxPosEmbed with parameterized axes_dim (NVIDIA#4) - Replace Flux2FeedForward with shared GatedMLP (NVIDIA#6) - Revert pipeline.py post_load_weights; add create_weights() in FLUX transformer __init__ to match WAN's __post_init__ pattern (NVIDIA#7) - Refactor _get_quant_algo_for_layer to encapsulate hasattr logic (NVIDIA#9) - Remove partial docstrings that just repeat class names (NVIDIA#10) - Simplify bias flags to direct booleans (NVIDIA#11) - Unify apply_rotary_emb: delete FLUX copy, FluxPosEmbed outputs [1,S,1,D] matching WAN format so shared function unchanged (NVIDIA#12) - Remove dynamic=True from maybe_compile (NVIDIA#13) - Add @torch.compiler.disable on FluxPosEmbed.forward - Tighten FLUX.2 PSNR test threshold from 12 to 20 dB Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> Signed-off-by: Kanghwan Jang <861393+karljang@users.noreply.github.com>

…model-kv-cache-test [None] Disable two-model eagle3 testing in CI

…A entry state Diag NVIDIA#2 (sA vs sA_pool static layout) at 5c8d830 showed bit-identical outer.shape / outer.stride / outer.cosize and same inner Swizzle type (OCI 2429536: linear1 4/4 PASSED, byte-equal streak 16 consecutive commits; fused [2-512-512-2] still CUDA_ERROR_MISALIGNED_ADDRESS 716). This clears the static layout-mismatch hypothesis and narrows the root cause to runtime mechanics — specifically: - Candidate NVIDIA#2: pipeline state carryover across FC1/FC2 phases. a_consumer_state / b_consumer_state / acc_producer_state are shared between FC1 and FC2 mainloops; each phase calls reset_count() but does NOT reset .index / .phase. If FC2 enters with a stale .phase bit, mbarrier parity checks may short-circuit, causing UMMA to consume un-produced sA_pool SMEM. - Candidate NVIDIA#3: UMMA descriptor SMEM base mismatch between sA-based tCrA (FC1, works) and sA_pool-based tCrA_l2 (FC2, fails). If the make_fragment_A encoding produces different SMEM descriptor bases for the two fragments (beyond the expected sA vs sA_pool delta), the PTX-emitted .shared::cta offset in FC2 UMMA could point outside the kernel's allocated SMEM region. This patch adds one cute.printf per FC1 MMA-entry and one per FC2 MMA-entry, printed from a single lane of the MMA warp via ``cute.arch.elect_one()``. Each line dumps: - a_consumer_state.count / .index / .phase - b_consumer_state.count / .index / .phase - acc_producer_state.count / .index / .phase - acc_stage_index / acc_stage_index_l2 (the computed stage used by the epilogue) - k_tile_cnt / k_tile_cnt_l2 - tCrA(_l2).iterator and tCrB(_l2).iterator (SMEM bases of the A/B fragments feeding UMMA) Both probes are gated on ``cutlass.const_expr(self.enable_linear2)``. When enable_linear2=False (Linear1 anchor byte-equal invariant), both blocks DCE-strip. No device side-effect except the printfs themselves (observability only). Both printfs live inside the existing ``if warp_idx == self.mma_warp_id:`` block (L4292), so only one warp per CTA emits; ``elect_one()`` further restricts to a single lane, keeping output volume bounded on [2-512-512-2] (expected ~few lines/CTA × few CTAs). One OCI iteration collects both Candidate NVIDIA#2 and NVIDIA#3 diagnostic data. Signed-off-by: xxi <xxi@nvidia.com>

…rintf operands Previous rt-diag commit (f25ac81) failed MLIR IR verification at compile time on OCI 2429682: 'cute.print' op operand NVIDIA#11 must be variadic of Type interface for cute types can be printed by `cute.print`, but got '!cute_nvgpu.smem_desc' Operand NVIDIA#11 in the FC2 printf was ``tCrA_l2.iterator`` and operand NVIDIA#12 was ``tCrB_l2.iterator``; the symmetric FC1 probe used ``tCrA.iterator`` / ``tCrB.iterator`` which are also smem_desc. The ``.iterator`` of a fragment bound to an SMEM tensor returns a ``!cute_nvgpu.smem_desc`` DSL-native type (the 64-bit-packed UMMA descriptor itself); ``cute.printf`` only accepts cute types with a printable Type interface, so this trips the MLIR verifier. The compile failure is itself informative for Candidate NVIDIA#3 (UMMA descriptor SMEM base): it confirms ``tCrA_l2.iterator`` IS the UMMA descriptor, not a raw address, and comparing FC1 vs FC2 descriptors requires a different extraction path (e.g. ``llvm.inline_asm`` that reads the descriptor word to a register, then prints via ``%lx``, or a Uint64 cast attempt). Defer until Candidate NVIDIA#2 is cleared. Fix: remove the ``.iterator`` args from both printfs; keep pipeline state fields (``a_consumer_state`` / ``b_consumer_state`` / ``acc_producer_state`` ``.count`` / ``.index`` / ``.phase``, ``acc_stage_index(_l2)``, ``k_tile_cnt(_l2)``). These are all ``cutlass.Int32`` / ``Uint32`` and compile cleanly with ``%d``. Pipeline state data alone is sufficient to accept or reject Candidate NVIDIA#2 (pipeline state carryover across FC1/FC2 phases causing stale ``.phase`` bit during FC2 consumer barrier acquire). Both probes remain gated on ``cutlass.const_expr(self.enable_linear2)`` so the Linear1 anchor (byte-equal invariant, 16 consecutive commits) DCE-strips the blocks. Signed-off-by: xxi <xxi@nvidia.com>

… PRs landed Apply nine consistency fixes against the post-NVIDIA#13671/NVIDIA#13672/NVIDIA#13673/NVIDIA#13674 state of the investigation: 1. Front-matter Status block: replace the "sig NVIDIA#6 root-caused, validation in flight" wording with the post-run8 picture (all 6 TRT-LLM PRs in review; NVIDIA#7 is an out-of-scope NIXL bug; deadline work is the TRT-LLM-side fallback for NVIDIA#7). 2. Front-matter Branches in this worktree: add the four new sig #4 / #5 / NVIDIA#6 branches. 3. Front-matter Related PRs: add NVIDIA#13674 / NVIDIA#13671 / NVIDIA#13672 / NVIDIA#13673 with chained-on-NVIDIA#13640 callout for NVIDIA#13673. 4. "Configurations that did not reproduce": #5 and NVIDIA#6 now do reproduce in single-process unit tests via the new tests added by NVIDIA#13672 and NVIDIA#13673; only #3 and NVIDIA#7 remain field-only. 5. Phase 6 close: the sig #4 regression test is no longer isolated in local/rc11-disagg-repro - it is now in the chained NVIDIA#13674 / NVIDIA#13671 pair. 6. Signature NVIDIA#6 section: drop the "(suspected)" qualifier and the "(most likely)" hedging on Where-it-lives - both are confirmed by run7 and run8. Rename the section header to describe the actual failure shape (recv buffer index leak via !isReady early return wedging assignBufferIndexForRecv) rather than the early control-path-stall hypothesis. Mirror the rename in the Signature - PR Map row. 7. File / Branch Index "New unit tests": add the new sig #5 (test_cancel_queued_gen_request_fulfills_receiver_future) and sig NVIDIA#6 (test_cancelled_after_ready_does_not_leak_recv_buffer_index, NIXL backend) tests. 8. Signature #3 status hypothesis: add a one-paragraph note that NIXL (signature NVIDIA#7) is now also a candidate cause of the half-initialized state, so a future field hit is not misattributed to a fresh TRT-LLM bug. 9. Phase 5 narrative: add a forward link explaining that the underlying terminal driver of the Phase-5 wedge was already NVIDIA#7 (NIXL), but #4 was the visible TRT-LLM-side symptom because the gen event loop was self-blocking before any of the later layers could surface. Signed-off-by: Chien-Chun Hung <2679986+chienchunhung@users.noreply.github.com> Made-with: Cursor

… section Fold two related pieces of analysis into the report as a new section between the Investigation Timeline and "Why the Existing Tests Did Not Catch This": (1) Signature taxonomy refining the naive "burst -> timeout -> cancellation -> bug" framing. Four-of-seven signatures (#1, #3, #5, NVIDIA#6) are direct cancellation-handling bugs; #4 is a structural latent blocking bug that cancellations expose; #2 is an eviction-driven bug that burst traffic exposes via memory pressure; NVIDIA#7 is a NIXL-internal contention bug that the same load shape happens to trigger but which is not strictly a cancellation bug. Includes a refined trigger chain diagram and two precise corrections (burst alone is not the trigger; "cancellation" is one of several entry points to cleanup paths). (2) Cascade map distinguishing two kinds of inter-signature tangling: - Type 1 (a fix produces a new signature): only one case, the #1 fix produces NVIDIA#6 by making the receiver-side !isReady early-return path reachable in production where a latent recv-buffer leak existed. This is why NVIDIA#6 PR (NVIDIA#13673) is explicitly chained on #1 fix PR (NVIDIA#13640). - Type 2 (a fix exposes a pre-existing signature): three cases where #4 fix exposes #5 / NVIDIA#6 and NVIDIA#6 fix exposes NVIDIA#7 because the upstream fix removes the masking effect on the downstream bug. These are not regressions of the fixes; they were latent pre-existing issues. - Subtler third relationship: #4 is structurally a defensive catcher for any upstream bug that produces a never-resolving receiver future. The #4 fix is independently valuable as defence in depth, not just a symptomatic patch. Also includes a fix-to-file mapping showing that the fixes do not overlap in code; the only structural dependency is the NVIDIA#6 -> #1 chain enforced by the PR base. Signed-off-by: Chien-Chun Hung <2679986+chienchunhung@users.noreply.github.com> Made-with: Cursor

…s fallback) to §03 venues Per user question: NSDI primary with MLSys same-cycle fallback is mechanically possible (NSDI spring notification ~Aug → MLSys deadline ~Oct/Nov gives ~6-10 weeks for revision), unlike the inverse pattern which is blocked by deadline ordering. Adds Variant B as a fully- scoped alternative to Variant A (the existing MLSys-primary plan). §5 strategic submission plan restructured: - Renamed existing plan to "Variant A — MLSys primary, NSDI next-cycle fallback (default recommendation)" with explicit note that same- cycle fallback is mechanically blocked. - New "Variant B — NSDI primary, MLSys same-cycle fallback" subsection with timeline, critical constraints (concurrent submission policy, ~6-10 week revision window, reframing cost, USENIX rebuttal cycle), and Phase 4 further-fallback to EuroSys 2029 winter. - "When to pick which" comparison table — first-shot venue, fallback availability, total shots through 2029 cycle, reframing burden, risk profile. - "What 'success' looks like" expanded to cover both variants with three outcome scenarios each. - Action item #3 updated to "pick Variant A or B" with deadlines for each. New action item NVIDIA#6 for Variant B specifically: pre-stage MLSys-framed intro/related-work during NSDI review period to preserve the tight revision window for reviewer-comment work. Variant A remains the recommended default (more conservative, doesn't burn two top venues in one cycle, more time per revision). Variant B is fully scoped for teams that prefer two same-cycle attempts. Signed-off-by: Chien-Chun Hung <2679986+chienchunhung@users.noreply.github.com>

…view Apply 5 review-driven edits to §16 to address residual concerns: Walk ordering (concern NVIDIA#8): change GMS RO and MX-receiver alias walks from per-module to top-level model.setup_aliases(). Matches the §7 mitigation contract from ai-dynamo/dynamo PR NVIDIA#7053 ("Call model.post_load_weights() (top-level only) before materialize_module_from_gms()"). transform_weights() and cache_derived_state() walks remain per-module since those bodies live on submodules. New "Why setup_aliases() is top-level-only" callout documents the asymmetry. Lifecycle of _weights_transformed (concern #4): new subsection specifying explicit set/reset/orthogonality rules. Includes a 2x2 truth table showing _weights_removed and _weights_transformed track different lifecycles and can take any combination. Reset is the orchestrator's responsibility (e.g., ModelLoader.reload() resets the flag before re-binding tensors); subclasses do not manage reset. Hard preconditions (concern #5): promote MX source-identity completeness from "open question" to "hard precondition P1." Lists transform-affecting parameters that MX identity must cover (attn_backend, quant backend list, FP8/NVFP4 fusion strategy, TP/EP layout, model revision). Specifies an in-tree backend-fingerprint fail-safe as the fallback if upstream MX cannot guarantee completeness. P2 documents that orchestrator owns _weights_transformed reset. Removes redundant open question NVIDIA#6 from the table. Cosmetic fixes (concerns #2, NVIDIA#6): "four stages" -> "three per-module stages plus orchestrator-managed per-process finalization." cache_derived_state description softened to "reserved for data-dependent state where it exists; many existing modules will have empty bodies." Scope clarifications (concerns #1, #3, NVIDIA#7): - Tiny PR scope reframed as "duck-typed helpers, not inheritance" with citations to existing model_loader.py walker pattern. Lists 4 walker helpers (_setup_aliases, _walk_transform, _walk_cache_state, _walk_full_post_load). - Migration callout: when migrating a subclass, the old post_load_weights() override MUST be removed; otherwise the new staged calls silently no-op. - Family PR #2 (Linear/Attention) gains a "quant-method callback decision" note with default = keep quant_method.post_load_weights callback name (no rename). No code changes. Drives Tiny prep PR scope and family-PR migration sequence. References: - TRT-LLM PR NVIDIA#13926 (GMS-only) - TRT-LLM PR NVIDIA#14151 (MX shim refactor) - ai-dynamo/dynamo PR NVIDIA#7053 (upstream GMS prototype) Signed-off-by: Chien-Chun Hung <chienchunh@nvidia.com> Signed-off-by: Chien-Chun Hung <2679986+chienchunhung@users.noreply.github.com>

- asyncUlyssesOp.cpp: - Tighten kNumSlots from 4 to 3 (minimum for V/Q/K within a single forward_async); document the slot-reuse hazard analysis. - Add cudaStreamIsCapturing guard at the top of getOrAllocSlot's alloc branch: first-time slot allocation (empty_strided_p2p + rendezvous + cudaMalloc) is not graph-capture-safe. - Annotate the 10s barrier timeout (hang-diagnostic budget). - utils_ltx2.py: document the sigma tensor-path zero-check skip (avoid .item() D2H sync that deadlocks under nsys + cuda_graph). - transformer_wan.py / transformer_ltx2.py: TODO on fuse_qk_norm_rope inert under _use_async_ulysses; points at PR follow-up NVIDIA#3. - test_ulysses_async.py (new): multi-rank (ws=2) end-to-end tests covering slot-ring wrap-around, CUDAGraph capture+replay, and multi-PG isolation in getOrCreateOp / set_group_info. Signed-off-by: Yiyun Lu <55233584+luyiyun1021@users.noreply.github.com>

#1 FD <-> chunk position misalignment in P2pHandleExporter Old code kept POSIX FDs in a flat mExportedFds and addressed them in removeHandles() by accumulating chunk counts across earlier pools. Pool-extension paths in exportHandles() append chunks to an existing pool but FDs to the END of mExportedFds, breaking that invariant. Removing a registered range could close another pool's FDs and leak the actual ones, also corrupting any in-flight UDS handshake. Fix: store FDs on P2pMemPool.fds in lockstep with chunks. Removal walks only the owning pool's fds - no cross-pool indexing possible. NVIDIA#2 UDS socket exposed POSIX FDs to other local users /tmp/trt_llm_p2p_fd_*.sock was created with default umask, no peer-credential check. Any same-host process that could connect could SCM_RIGHTS-receive the FDs and import the exporter's GPU memory. Fix: tight umask + explicit chmod 0600 around bind, fail closed if either fails; SO_PEERCRED gate at accept() to refuse non-same-uid clients (defense in depth). NVIDIA#3 P2pTransferContextPool never erased per-thread Contexts Long-lived processes whose callers are transient threads leaked ~64MB cubTempStorage + a worker pool per dead thread; OS thread::id reuse could also hand a fresh thread a stale Context built in a different CUDA context. Fix: pool is now held by shared_ptr (factory + private ctor enforce this); contextForCurrentThread installs a thread_local guard that, on thread exit, calls eraseForCurrentThread via weak_ptr (no-op if pool is already gone, so destruction order between agent and caller threads is safe in either direction). NVIDIA#4 Reads of P2pHandleExporter::mLocalInfo were unprotected getLocalAgentDesc did isSupported() then getLocalInfo().serialize() - a TOCTOU pair where a concurrent registerMemory could append a pool between the two reads, producing a half-written wire blob. Fix: single mInfoMutex protects mLocalInfo / mDetectedHandleType / mUdsPath end-to-end. New serializeIfSupported() does both checks under the lock and replaces the pair at the only caller. The previous mPoolsMutex (added for FDs) is the same mutex, renamed. NVIDIA#8 Document acquire-then-record contract on CudaEventPool Recycled events carry their prior recording; calling cudaEventQuery before record() returns stale state. All current callers already follow acquire->record->query, but the contract is now spelled out on the API doc so future callers don't diverge. Tests: p2pTransferAgentTest +2 regressions - SerializeIfSupportedEmptyOnFreshAgent: locked accessor returns empty on a freshly constructed agent. - TransientCallerThreadsReleaseContextOnExit: 32 transient caller threads each take a Context from the same agent; after join, pool->numContexts() must be exactly 0 (proves thread-exit guard fires). Local validation: p2pMemInfoTest 11/11 transferAgentTest 18/18 (AgentDesc + VmmDescSplitter) p2pTransferAgentTest 36/36 nixlP2pE2ETest 7/7 (+2 env-gated, validated separately) Signed-off-by: Chuang Zhu <111838961+chuangz0@users.noreply.github.com>

_msa_index_proxy_and_topk drove block selection through fmha_sm100.sparse_topk_select, which only accepts a scalar num_valid_pages and scalar force_begin_blocks/force_end_blocks. Those are applied uniformly to every (head, token) row, so the forced local block was pinned to the batch-global nvp-1 instead of each query's own last valid block, and the OOB clamp used a batch-wide bound. For any batch with mixed sequence lengths (and for prefill, where each query token has its own causal extent) this forced the wrong local block and admitted out-of-range block ids for every request shorter than the batch-longest. Replace the kernel selection with a per-query torch selection (_select_blocks_from_maxscore + _per_token_valid_blocks) that mirrors the reference _index_attention_and_select logic (init/local forcing + per-query valid-block masking + top-k) on the amax-reduced per-KV-head scores. The MSA path already disables CUDA graphs, so the host-side selection is safe. The amax-vs-union index-head reduction for group>1 (a TP artifact of replicated index heads / sharded KV heads) is left as-is and now explicitly flagged as NEEDS REVISIT (bug #3) at the reduction site, in the docstring, and in the parity test's Stage 2 / Stage 2b (which assert group==1 exact parity and a group>1 subset invariant). Signed-off-by: Balaram Buddharaju <169953907+brb-nv@users.noreply.github.com>

…E. All 8 acceptance criteria hold at runtime against... QA verdict: APPROVE (weighted_score 8.6) Change vs main: 7 files changed, 1376 insertions(+), 8 deletions(-) Build iterations run: 13 Code-quality verification: clean after 2 round(s) - round 1 (REJECT): ROUND MIS-SCOPED — patch not actually reviewed. This round's /simplify, checklist, and /code-review executed with cwd = the agent-flow framework repo (branch kleinc/code_quality_improvement_v1) and reviewed the code-quality-stage feature itself, NOT this task's PR patch. The ChatGLM3 code is on branch agent-team/chatglm3-6b-bringup-v1 in the TensorRT-LLM snapshot repo (a different directory), so `/simplify`/`/code-review` (which diff the cwd's HEAD) never touched it. Root cause: the Quality agent runs in the framework repo's cwd, not the snapshot repo; another round will misfire identically unless the Quality agent runs with the TensorRT-LLM task branch as its working directory. I re-reviewed the ACTUAL patch (git diff main in TensorRT-LLM). Real, unaddressed findings: - TESTS (duplication): `_patch_chatglm3_hf_tied_compat()` duplicated in all 4 new test files; `_find_cuda_graph_runner()` and `_assert_cuda_graph_hard_path()` duplicated across test_chatglm3_gsm8k.py and test_chatglm3_replay.py; CHATGLM3_CKPT const, skip decorators, and RuntimeCfg class repeated. Consolidate to a shared util/conftest (keep all coverage). - modeling_chatglm.py (simplification): redundant post-normalization reads of attention_bias/mlp_bias/head_dim via `getattr(...) or ...` after normalize_chatglm_config already set them; normalize_chatglm_config returns config but both callers discard it; a few verbose comments. - ALTITUDE (larger, behavior-sensitive; flagged not applied): custom imperative load_weights vs framework BaseWeightMapper; ChatGLM-specific normalization special-cased in shared config_utils.py; partial_rotary_factor/rope_fusion hardcoded. No fixes applied: the round did not validly review this patch, and edits to modeling/test code cannot be validated against the GSM8K accuracy criterion in this environment (risking a criterion regression). Recommend fixing the Quality agent's working directory (must be the snapshot repo/task branch) before the next round, then applying the test-dedup and post-normalization simplifications above. - round 2 (REJECT): Round 2 finally assessed the ACTUAL task patch (git diff main...agent-team/chatglm3-6b-bringup-v1 in the TensorRT-LLM repo), unlike round 1 whose skill steps mis-scoped to the agent-flow framework repo (cwd). Note: the automated stage still mis-scopes — /simplify and /code-review diff the cwd (agent-flow), not this task's snapshot repo — so I targeted the correct repo manually. /code-review (correctness): CLEAN — returned []. No correctness bugs survive verification. The high-risk items are validated: partial-RoPE, QKV/GQA split sizes (32 heads / 2 KV / head_dim 128), and transformer.* weight-name mapping are exercised by the GSM8K-parity acceptance gate QA approved; the fused-residual path is completed by the final self.norm(hidden_states, residual); embedding matches the Qwen sibling and load_weights handles attention-DP (tp_size = 1 if enable_attention_dp else mapping.tp_size). Speculative "other ChatGLM variant" concerns (_set_missing fill-if-None, head_dim None if kv_channels absent, missing dense/output biases) are not reachable for chatglm3-6b and are guarded by strict _take/consumed accounting. /simplify (round 2): applied one behavior-preserving cleanup (dropped the unused `return config` from normalize_chatglm_config, annotation -> None); it was subsequently reverted in the working tree, so the patch is unchanged. OUTSTANDING (why REJECT): the checklist's test item is not satisfied — heavy test-helper duplication across the 4 new test files (`_patch_chatglm3_hf_tied_compat` in all 4; `_find_cuda_graph_runner`/`_assert_cuda_graph_hard_path`/`RuntimeCfg` across gsm8k+replay; repeated CHATGLM3_CKPT/skip decorators), ~100+ duplicated lines to consolidate into a shared util; plus minor post-normalization config-read redundancy in modeling_chatglm.py. NOT applied here because tensorrt_llm is not importable in this environment and there is no GSM8K/test rig, so a test-file move cannot be validated for collection/pass (risking the acceptance criterion), and the working-tree edit I did make was reverted. RECOMMENDATION: this cannot be resolved by another automated round as currently wired — fix the Quality agent's working directory so /simplify and /code-review target the snapshot repo/task branch, then apply the test-helper de-duplication with the test rig available (or the author applies it directly). Correctness and framework-conformance are clean; the residual is a non-behavioral maintainability cleanup. - round 1 (REJECT): Code-quality round 1 for the chatglm3-6b bringup patch (modeling_chatglm.py, config_utils.py, _torch/models/__init__.py + 4 new test files). /simplify applied (safe, behavior-preserving, verified against repo convention): - modeling_chatglm.py: consume the normalized config fields instead of re-deriving booleans inline — `bias=config.attention_bias`, `dense_bias=config.mlp_bias`, GatedMLP `bias=config.mlp_bias`, and `head_dim=config.head_dim` at both sites (attention + load_weights). Confirmed no base infra reads these fields and that the model already hard-depends on normalize_chatglm_config having run; matches Llama/Phi3/Nemotron/Parakeet. - Both integration tests: removed the unreachable BFS-over-object-graph fallback in `_find_cuda_graph_runner`, keeping the deterministic `_executor.engine.model_engine.cuda_graph_runner` chain (verified in source; the runner is at a fixed location in single-process mode and in a subprocess otherwise, where BFS also fails). Codex checklist pass: the working tree carried minor docstring/comment trims in the ChatGLM3 test files (removed helper docstrings on `_patch_chatglm3_hf_tied_compat`, `_pick_metric_key`, `greedy_generate`, and an inline comment); no behavioral change. /code-review (high) found 10 items. Fixed this turn: (NVIDIA#4) `_pick_metric_key` in test_chatglm3_gsm8k.py now raises loudly instead of silently gating on an arbitrary non-exact_match metric — strict hardening, no change to any real gsm8k run. OUTSTANDING (why not auto-fixed): - CORRECTNESS, latent/out-of-target: (NVIDIA#1) load_weights only handles chatglm3-6b's bias layout — it unconditionally fetches the QKV bias and never loads dense/MLP biases, so ChatGLM2/3 variants with add_qkv_bias=False or add_bias_linear=True would KeyError or trip the strict unconsumed-weight ValueError. (NVIDIA#2) normalize_chatglm_config never maps ChatGLM `rope_ratio` to `rope_theta` (get_hf_rope_theta reads config.rope_theta), so chatglm3-6b-32k/128k silently get theta=10000 and degrade at long context. Both are no-ops for base chatglm3-6b (add_qkv_bias=True/add_bias_linear=False/rope_ratio=1) so QA passed; fixing them adds untestable variant paths (no GPU/checkpoint here) and is scope creep on a base bringup — hand to the coder to fix+test or narrow the module's "ChatGLM2/3" claim. - TEST ROBUSTNESS: (NVIDIA#3) the CUDA-graph hard-path tests require TLLM_WORKER_USE_SINGLE_PROCESS=1 but only name it in an assert string and never set it; without it, ENABLED-cfg tests hard-fail on `assert runner is not None` and BASELINE silently skips its checks. Not auto-fixed because gating/skipping would weaken the author's intended hard-path assertion — needs a CI-env or fixture decision. - ALTITUDE/REUSE (larger refactors, unsafe without test execution): (NVIDIA#5) hand-rolled load_weights + consumed-set + ignorable-suffix lists reimplement BaseWeightMapper/register_mapper (cf. Glm4WeightLoader in modeling_glm.py); (NVIDIA#9) normalize special-cased in shared load_pretrained_config + duplicated in __init__ vs a ChatGLMConfigLoader; (NVIDIA#6) `_patch_chatglm3_hf_tied_compat` (4 copies) + RuntimeCfg/_build_llm/_find_cuda_graph_runner/_assert_cuda_graph_hard_path (2 copies) duplicated across two separate, GPU-gated test trees; (NVIDIA#10) _GreedyHFLM._model_generate hand-rolls HF greedy decode. - EFFICIENCY: (NVIDIA#7) test_chatglm3_source_activation_replay reloads the 6B HF model + rebuilds TRT model for both parametrized scenarios though only the trailing cuda_graph block differs (use the _HFRef cache pattern). - SCOPE CREEP: (NVIDIA#8) the diff reflows unrelated bart/minimaxm3/qwen imports (+ a mistral import in config_utils.py); left as-is because formatter ownership is ambiguous (line 33 is 89 chars vs isort/yapf's 80) and reverting risks fighting the pre-commit hook. Rejecting: a substantive edit was made this turn and several actionable findings remain (test env-var robustness is the most important; the two latent correctness gaps and the altitude/reuse refactors should be triaged by the coder). - round 2 (APPROVE): Code-quality round 2 for the chatglm3-6b bringup patch. /simplify (round 2) applied 3 trivial, behavior-neutral cosmetic cleanups in modeling_chatglm.py: inlined the single-use `head_dim` local in ChatGLMAttention (`head_dim=config.head_dim` passed directly to super()), inlined the single-use `head_dim` temp in normalize_chatglm_config, and dropped the noise `: set` annotation on `consumed`. Reuse/efficiency/altitude agents found nothing new safely-actionable and confirmed the two "small fold" hypotheses don't hold (base `skip_modules` matches model module names, not the checkpoint keys `_IGNORABLE_*` filters; the second normalize_chatglm_config call is load-bearing for the direct-ModelConfig test path) — the only remaining altitude/reuse items are large refactors. Codex checklist pass (round 2): no net changes — independent verification confirmed the only diffs since round 1 were the 4 quality edits above/below, all behavior-neutral. /code-review (high) re-surfaced the same stable finding set (the substantive code was unchanged; no coder commits landed between rounds). I acted on the most actionable one this turn: - FIXED (NVIDIA#3, test robustness): the CUDA-graph hard-path tests required TLLM_WORKER_USE_SINGLE_PROCESS=1 (their ENABLED config hard-asserts the in-process runner) but never set it, so they would spuriously fail/silently-skip when run with a checkpoint outside a harness that exports it. Added an autouse `monkeypatch.setenv` fixture to both integration files (test_chatglm3_gsm8k.py, test_chatglm3_replay.py). Verified the env is read at runtime (utils.py:380, os.environ.get inside a function) so the fixture takes effect before _build_llm; the fix aligns each test with its own documented requirement and cannot regress the already-validated path. - Previously fixed (round 1): _pick_metric_key now fails loudly instead of gating on an arbitrary metric; consumed bias/head_dim inline derivation replaced with normalized fields; BFS runner-finder fallback removed. RESIDUAL (documented follow-ups — not safely resolvable by the automated quality loop): - Out-of-target latent correctness: (NVIDIA#1) load_weights only handles the chatglm3-6b/chatglm2-6b bias layout (add_qkv_bias=True/add_bias_linear=False) — it unconditionally fetches the QKV bias and never loads dense/MLP biases, so ChatGLM variants with add_qkv_bias=False (KeyError) or add_bias_linear=True (unloaded bias + unconsumed-weight ValueError) fail; (NVIDIA#2) rope_ratio is never mapped to rope_theta, so long-context chatglm3-6b-32k/128k (rope_ratio>1) silently use theta=10000. Both are no-ops for the bringup target (base chatglm3-6b/chatglm2-6b, rope_ratio=1); fixing needs variant checkpoints to test and expands scope beyond the bringup — for the coder/a follow-up. - Larger altitude/reuse/efficiency refactors needing test-execution access: (NVIDIA#4/NVIDIA#7) adopt BaseWeightMapper/register_mapper for load_weights + a ChatGLMConfigLoader/config-subclass to remove the shared-runtime normalize special-case and its second call; (NVIDIA#5) consolidate the 4x/2x-duplicated test helpers across the two GPU-gated test trees; (NVIDIA#6) cache the 6B HF/TRT model in the parametrized attention test; (NVIDIA#8) the _GreedyHFLM hand-rolled reference decode. APPROVE rationale: base chatglm3-6b is fully correct (QA-approved, acceptance criteria met); two thorough rounds applied all safe in-scope polish culminating in the NVIDIA#3 test-robustness fix; the quality loop has converged (round 2 reproduced round 1's findings with no coder changes); and every remaining finding is either out-of-target latent correctness for untested ChatGLM variants or a larger refactor requiring test-execution access — none of which another automated /simplify→/code-review round can safely resolve. Residuals are surfaced above for human/coder triage. QA report: DECISION: APPROVE. All 8 acceptance criteria hold at runtime against the current committed code (HEAD 2323424, iter11). Reference is the checkpoint's own trust_remote_code modeling_chatglm.py — an independent HF reference sharing zero code with the TRT-LLM path. PER-CRITERION (runtime evidence): 1. Bootstrap chain — PASS. slurm/core_gpu_tests.4147161.log: "Built target build_wheel_targets", "Successfully installed tensorrt_llm-1.3.0rc21 transformers-5.5.4" (pip install -e .[devel]), "Claude Code successfully installed!", "claude-agent-sdk 0.2.93"; all inside srun pyxis container with --container-mounts=<repo>:<repo> --container-workdir=<repo> under set -eo pipefail (tests ran only because bootstrap rc=0). 2. config_registration_and_weight_accounting — PASS (TIER1 rc=0). Real ckpt loads unedited, arch resolves to ChatGLMForCausalLM, TRTLLM backend + KVCacheManagerV2, strict state-dict accounting (raises on missing/unexpected/shape-mismatch; only rotary_pos_emb.inv_freq tolerated). 3. source_activation_replay — PASS both graph modes. layer0 cosine=1.000000 mean_abs=0.00000, layer27 cosine=0.999999; decode graph-vs-eager max_abs=0.0 cosine=1.0. Companion partial_rope_boundary_and_theta PASS (dims [0:64] rotate, [64:128] pass-through, theta=10000^(-i/32)). 4. source_logit_replay — PASS both configs. trt_tok==hf_tok (5231/30910/4802), cosine~0.99999; enabled num_captured_graphs=8, baseline 0. 5. generation_parity — PASS both configs. 5 prompts x 32 steps, token_mismatches=0, min per-step cosine~0.99998. 6. llm_api_smoke — PASS both configs (nonempty deterministic; V2 + TRTLLM asserted; enabled enabled=True graphs=8, baseline graphs=0). 7. gsm8k accuracy_canary — PASS (slurm/gsm8k_gate.4145862.log): HF=40.00, baseline=40.00 gap 0.00, enabled=40.00 gap 0.00; enabled num_captured_graphs=32. 8. gsm8k full_trtllm_eval — PASS: HF=52.77, baseline=52.99 gap 0.23, enabled=53.07 gap 0.30 (both < 2.0 tol), enabled captured 32 CUDA graphs; RESULT: PASS. Absolute ~53% matches published ChatGLM3-6B GSM8K, ruling out a shared-bug artifact. PROVENANCE: Core log 4147161 started 05:55, after the 05:33 HEAD commit -> definitively iter11 (criteria 1-6). git diff iter10->iter11 (2d58cc7..2323424) touched ONLY the two test files (additive CUDA-graph hard-path assertions); tensorrt_llm/ is byte-identical. GSM8K log 4145862 prints the iter11-only [cuda_graph_hard_path] output -> it ran the current test file against the unchanged runtime code (live bind-mount imported iter11 mid-bootstrap). Criteria 7-8 are therefore verified against current code. INDEPENDENT RERUN THIS TURN: I resubmitted both jobs. Core (4149172) matched the prior post-commit run and was cancelled as redundant. My GSM8K resubmit (4149173) FAILED in ~1.5 min inside build_wheel.py (FileNotFoundError: inherit_graph_6.md5 in copy_resolving_symlink) — a concurrent-build race from my own simultaneous submission of two bootstrap-heavy jobs on one live repo mount, NOT a ChatGLM defect. Accepted evidence is the sequential post-commit runs + full code/provenance review. RED-TEAM (all addressed): independent HF reference (no shared helpers); partial GPT-J RoPE covered analytically + behaviorally (layer0 exact); CUDA-graph silent fallback ruled out by in-process CUDAGraphRunner introspection (runner.enabled + >=1 torch.cuda.CUDAGraph; baseline 0); KV-cache/mask/scale drift ruled out by 32-step parity (0 mismatches) + GSM8K absolute-score match; transformers 5.x bridges (all_tied_weights_keys={}, max_length, num_hidden_layers) are loader-compat only and do not alter ChatGLM forward semantics. EVALUATION SCORES: - Functionality (2.0): 9 — every criterion passes with exact argmax parity, cosine >0.9999, GSM8K within 0.3 pts of HF on both configs. - Code Quality (1.0): 8 — clean, well-documented 373-line model reusing existing modules; narrow config hook; strict weight accounting; minor incidental import reformatting in __init__.py; a pre-existing (non-diff) ruff E501 at __init__.py:130. - Performance (1.5): 8 — parity-first bring-up (no perf gate in task.yaml); production TRTLLM backend + CUDA graph + overlap scheduler + fused QKV/gate-up + KVCacheManagerV2 all working. - Completeness (1.0): 9 — complete, buildable, no placeholders, full state-dict accounting, both runtime configs. - Technical Sophistication (1.0): 9 — correct unfused partial GPT-J RoPE with TRTLLM backend + CUDA-graph capture, compact MQA KV, config/transformers bridges. Weighted = (9x2.0 + 8x1.0 + 8x1.5 + 9x1.0 + 9x1.0)/6.5 = 56/6.5 = 8.6. STRENGTHS: genuine independent HF reference; real CUDA-graph hard-path proof (not just config flags); strict weight accounting; large GSM8K margin; minimal composable design (no new schemas). WEAKNESSES: multi-GPU/perf/chunked-prefill deferred (out of scope per task.yaml); harness must run the two jobs sequentially to avoid the concurrent-build race; pre-existing E501 worth a one-line fix. RECOMMENDATION: APPROVE — no code defect identified; task.yaml completion_criteria (GSM8K within 2 pts with and without CUDA graph + overlap scheduler) is met on both configs.</summary> </invoke>

Follow-up: user requires each pre-merge pick to already be actively running in L0 CI post-merge on origin/main — i.e. no QA-lane borrowing, no uncommenting of previously-disabled lines, no new config yamls, no first-time-in-L0 e2e/gen_only combinations. Prior revision's two picks that failed that bar: * gpt-oss 8k1k con512 e2e dep2 — active in QA lane only; the e2e mode of this 8k1k dep2 config was not in L0 anywhere. * kimi 8k1k con4 tep8 mtp3 gen_only GB200 — was commented out in the GB200 12-GPU 3-node post_merge block since April. Replaced with: * gpt-oss 8k1k con1024 tp4 e2e MTP0 — was active in L0 post-merge on GB200 ctx1_node1_gpu1_gen1_node1_gpu4 (8 GPUs, 2 nodes). Gen-worker parallelism shifts dep→tp for this slot. * kimi 8k1k con4 tep8 mtp3 gen_only — was active in L0 post-merge on GB300 ctx1_node1_gpu4_gen1_node2_gpu8 (12 GPUs, 3 nodes). GPU family shifts GB200→GB300 for this slot. File changes: - l0_gb200_multi_nodes_perf_sanity_ctx1_node1_gpu1_gen1_node1_gpu2.yml: reverted to origin/main state (no pre_merge block). - l0_gb200_multi_nodes_perf_sanity_ctx1_node1_gpu1_gen1_node1_gpu4.yml: new pre_merge block with gpt-oss 8k1k con1024 e2e; test line removed from post_merge to avoid double-run. - l0_gb200_multi_nodes_perf_sanity_ctx1_node1_gpu4_gen1_node2_gpu8.yml: reverted to origin/main state (kimi 8k1k stays commented out). - l0_gb300_multi_nodes_perf_sanity_ctx1_node1_gpu4_gen1_node2_gpu8.yml: new pre_merge block with kimi 8k1k con4 tep8 mtp3 gen_only; same test removed from post_merge. Jenkins: - Dropped stages: FUNCTIONAL-ONLY-CTX1-NODE1-GPU1-GEN1-NODE1-GPU2 (GB200) and FUNCTIONAL-ONLY-CTX1-NODE1-GPU4-GEN1-NODE2-GPU8 (GB200 kimi variant). - New stages: FUNCTIONAL-ONLY-CTX1-NODE1-GPU1-GEN1-NODE1-GPU4 (GB200 gpt-oss) and FUNCTIONAL-ONLY-CTX1-NODE1-GPU4-GEN1-NODE2-GPU8 (GB300 kimi variant). - Post-merge decrement: GB200-8_GPUs-2_Nodes CTX1-NODE1-GPU1-GEN1- NODE1-GPU4 7→6; GB300-12_GPUs-3_Nodes CTX1-NODE1-GPU4-GEN1-NODE2- GPU8 5→4. Others revert to origin/main testCounts. Shape matrix (all 4 pre-merge slots now at 8k1k): NVIDIA#1 DSR1 B200 8k1k con1536 dep e2e=NO gen_only MTP✓ NVIDIA#2 gpt-oss GB200 8k1k con1024 tp e2e=YES MTP✗ NVIDIA#3 glm-5 GB300 8k1k con1024 dep e2e=YES MTP✓ NVIDIA#4 kimi GB300 8k1k con4 tep e2e=NO gen_only MTP✓ Signed-off-by: Chenfei Zhang <chenfeiz@nvidia.com>

kaiyux added 4 commits September 28, 2023 09:00

Update code

6e9e318

Add .a libs

766926c

Update submodule

496456e

Update submodule

6111f52

kaiyux merged commit 279e329 into main Sep 28, 2023

Nam-ang mentioned this pull request Jan 9, 2024

When run llama2, Caught signal 11 (Segmentation fault) #752

Closed

tdeng521 mentioned this pull request Mar 7, 2024

batch size will affect llm inference results? #1250

Closed

4 tasks

hademircii mentioned this pull request Mar 12, 2024

Flan-T5 models with Tensor Parallelism #1286

Closed

4 tasks

zxs789 mentioned this pull request Jun 4, 2024

H20 Using random weights to infer llama2-13B results in a divide-by-zero error. #1717

Closed

4 tasks

wu1du2 pushed a commit to wu1du2/TensorRT-LLM that referenced this pull request May 11, 2025

Merge pull request NVIDIA#3 from NVIDIA/kaiyu/update

3aa8467

Update TRT-LLM code

kipraveen pushed a commit to kipraveen/TensorRT-LLM that referenced this pull request May 12, 2025

Merge pull request NVIDIA#3 from kipraveen/canary_cpp

0d73b5c

Add support for CPP inference with decoder

justremotephone mentioned this pull request Feb 5, 2026

[Bug]: Random "CUDA error: an illegal instruction was encountered" for gpt-oss 120B at Nvidia Spark. #11313

Closed

5 tasks

kaiyux mentioned this pull request Mar 9, 2026

[TRTLLM-10421][perf] Add fused cat+fp8_quantize CUDA kernel for DSA indexer #11899

Merged

3 tasks

thorjohnsen mentioned this pull request Mar 16, 2026

[TRTLLM-11513][fix] eliminate dangling reference and TOCTOU races in KVCacheManager #12251

Closed

3 tasks

dc3671 pushed a commit to dc3671/TensorRT-LLM that referenced this pull request Apr 17, 2026

Merge pull request NVIDIA#3 from ziyixiong-nv/dev-fxiong-disable-two-…

92565d8

…model-kv-cache-test [None] Disable two-model eagle3 testing in CI

chienchunhung mentioned this pull request Apr 28, 2026

[https://nvbugs/6104831][test] Add reproducers for KV-cache-block cascade prune assertion #13571

Closed

1 task

chienchunhung mentioned this pull request May 5, 2026

[https://nvbugs/6104831][fix] Disaggregated KV transfer: lifecycle, cancellation, and quiescence hardening #13713

Closed

1 task

xxi-nv mentioned this pull request May 8, 2026

[None][refactor] MoEScheduler split + MegaMoE EPLB / multi-chunk / CI integration #13908

Merged

2 tasks

wu6u3tw mentioned this pull request May 20, 2026

[None][feat] Wan 2.2 serve: extend video generation API #13783

Closed

1 task

longcheng-nv mentioned this pull request May 21, 2026

[None][perf] DSv4 indexer Top-K: align kFTarget=kK, uniform kSeqSmall=4096, diagnostic env knobs #14413

Merged

chang-l mentioned this pull request May 24, 2026

[TRTLLM-11457][feat] Async Ulysses pipeline (Enabled for LTX-2 + WAN) #13978

Merged

1 task

limin2021 mentioned this pull request Jun 1, 2026

[TRTLLM-35882][feat] Add cute dsl gvr top-k decode kernel #14602

Merged

1 task

sunnyqgg mentioned this pull request Jul 1, 2026

[None][test] Optimize LLM test startup overhead #15777

Merged

3 tasks

Wanli-Jiang mentioned this pull request Jul 15, 2026

[TRTLLM-14026][feat] BREAKING: Remove C++ modules for legacy TRT backend #16369

Merged

1 task

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

Update TRT-LLM code#3

Update TRT-LLM code#3
kaiyux merged 4 commits into
mainfrom
kaiyu/update

kaiyux commented Sep 28, 2023

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant

Uh oh!

Conversation

kaiyux commented Sep 28, 2023

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant