[https://nvbugs/6402058][fix] [None][fix] Remove stale prepare_inputs event sync to unhang DS-V3 EP4 MTP - #15849
Conversation
|
No actionable comments were generated in the recent review. 🎉 ℹ️ Recent review info⚙️ Run configurationConfiguration used: Path: .coderabbit.yaml Review profile: CHILL Plan: Enterprise Run ID: 📒 Files selected for processing (2)
💤 Files with no reviewable changes (1)
WalkthroughThe MLA BMM path now makes non-contiguous inputs contiguous on SM120/121 before dispatch. The pull request also removes a waiver for ChangesMLA BMM layout handling
Estimated code review effort: 2 (Simple) | ~10 minutes Possibly related PRs
Suggested labels: Suggested reviewers: 🚥 Pre-merge checks | ✅ 5✅ Passed checks (5 passed)
✨ Finishing Touches🧪 Generate unit tests (beta)
Comment |
3986aaa to
ce042d5
Compare
ce042d5 to
787c18a
Compare
|
AI suggests this should wait #16088 to be merged first. DetailFor the target DS-V3-Lite / V1 KV-cache path, yes: the original behavior is still correct because the race that Why
|
| padded_requests.context_requests, | ||
| padded_requests.generation_requests, | ||
| ) | ||
|
|
There was a problem hiding this comment.
787c18a to
4de2722
Compare
0a5fa44 to
4143fcf
Compare
…MU fault MLA._bmm_bf16_out passes the LHS as a head-major transpose of a [tokens, heads, dim] buffer, so its batch stride is the token count rather than the tile extent. On SM120/121 cuBLAS selects a TMA-based nvjet kernel for that layout, which cuTensorMapEncodeTiled cannot describe; the kernel faults with an MMU page fault that surfaces either as CUBLAS_STATUS_INTERNAL_ERROR or as a later illegal memory access at the next synchronize. Make the LHS contiguous on those architectures when it is not already, so a non-TMA kernel is selected. Only the LHS matters: bisection showed densifying the output buffer alone does not avoid the fault. This mirrors the existing SM120 accommodations in fp8_block_scaling_bmm_out and the CUTLASS TMA guard in fused_moe_cutlass.py. Also unwaive the test the fault was masking. Signed-off-by: trtllm-agent <296075020+trtllm-agent@users.noreply.github.com>
4143fcf to
a783a26
Compare
|
Note GitHub couldn't provide a complete incremental comparison for this pull request, so CodeRabbit is performing a full review instead. This review may take a little longer. |
| # cuTensorMapEncodeTiled cannot describe it -- the kernel then faults with | ||
| # an MMU page fault (surfaced as CUBLAS_STATUS_INTERNAL_ERROR or a later | ||
| # illegal memory access). Densify so a non-TMA kernel is selected. | ||
| a = a.contiguous() |
There was a problem hiding this comment.
This solution is completely different from the PR title and description, please update them.
brnguyen2
left a comment
There was a problem hiding this comment.
Two things before merge:
-
The PR title and description describe a different fix entirely (removing
_prepare_inputs_event/wait_for_input_copy()fromPyTorchModelEngine). The diff is an SM120/121 densify inMLA._bmm_bf16_outplus the waiver removal. Please rewrite both; the title also carries a doubled prefix —[https://nvbugs/6402058][fix] [None][fix]— instead of[https://nvbugs/6402058][fix] <description>. -
The waiver is global but the failure was on a 4-GPU post-merge RTX Pro 6000 stage, which pre-merge CI won't run. Trigger that stage explicitly (
/bot run --extra-stage "...") so the unwaive gets real coverage, and please note in the PR how the fault was tied to the reported timeout.
| if self.use_cute_dsl_bf16_bmm and is_sm_100f(): | ||
| torch.ops.trtllm.cute_dsl_bf16_bmm_blackwell(a, b_no_transpose, output) | ||
| else: | ||
| if get_sm_version() in (120, 121) and not a.is_contiguous(): |
There was a problem hiding this comment.
The guard densifies only the LHS, but at the call sites the output is also a non-contiguous transposed view (attn_output.transpose(0, 1), mla.py:1578 and mla.py:1700), and b_transposed is a transposed view too. If cuBLAS's kernel heuristic keys off those strides as well, a contiguous a may not be sufficient outside the shapes that were bisected. Worth stating in the comment which operand layouts were actually swept, or asserting/densifying defensively on the output as well.
Summary
PyTorchModelEngineheld a leftover_prepare_inputs_eventrecorded after_prepare_and_schedule_batchand synchronized viawait_for_input_copy(). In the DeepSeek-V3-Lite EP4 + MTP2 + attention_dp + CUDA-graph + overlap-scheduler + torch.compile configuration, this extra CUDA event synchronization interacted badly with the overlap scheduler on the 4-GPU post-merge stage, causing the executor to block indefinitely on IPC poll._prepare_inputs_eventfield, its record site in the forward path, and thewait_for_input_copy()helper, since input-copy ordering is already guaranteed by the stream/graph semantics elsewhere. The corresponding waive entry for nvbugs/6402058 was dropped fromwaives.txtso the previously hanging DeepSeek-V3-Lite EP4 MTP2 test re-enters CI to confirm the fix.Test plan
Links
Dev Engineer Review
_prepare_inputs_eventsynchronization fromPyTorchModelEngine.wait_for_input_copy().#15546for the DS-V3-Lite/V1 KV-cache path.MLA._bmm_bf16_outto make the first BMM input contiguous on SM120/121 when required.nvbugs/6402058waiver entry. The waiver format and scope are consistent with the targeted test.QA Engineer Review
TestDeepSeekV3Lite::test_bfloat16_4gpuswaiver fromtests/integration/test_lists/waives.txt.