fix(pcie): prevent two-slot selector overflow - #101
Conversation
|
Note Reviews pausedIt looks like this branch is under active development. To avoid overwhelming you with review comments due to an influx of new commits, CodeRabbit has automatically paused this review. You can configure this behavior by changing the Use the following commands to manage reviews:
Use the checkboxes below for quick actions:
No actionable comments were generated in the recent review. 🎉 ℹ️ Recent review info⚙️ Run configurationConfiguration used: Organization UI Review profile: CHILL Plan: Pro Plus Run ID: 📒 Files selected for processing (1)
🚧 Files skipped from review as they are similar to previous changes (1)
📝 WalkthroughWalkthroughPCIe staging-slot selection now uses a constexpr host selector for DCP A2A and CUDA-side generation metadata for eager allreduce and two-shot collectives. Kernels receive both slot options. Graph and integration tests verify alternating device-selected slots. ChangesPCIe staging-slot selection
Estimated code review effort: 4 (Complex) | ~45 minutes Sequence Diagram(s)sequenceDiagram
participant PCIeCollective
participant SlotMetadata
participant PCIeKernel
PCIeCollective->>SlotMetadata: pass paired staging-slot metadata
SlotMetadata->>PCIeKernel: select active slot
PCIeKernel->>PCIeKernel: derive staging pointers and execute collective
Possibly related issues
Possibly related PRs
🚥 Pre-merge checks | ✅ 3 | ❌ 2❌ Failed checks (2 warnings)
✅ Passed checks (3 passed)
✨ Finishing Touches🧪 Generate unit tests (beta)
Thanks for using CodeRabbit! It's free for OSS, and your support helps us grow. If you like it, consider giving us a shout-out. Comment |
|
Post-push status:
Production was not rebuilt or restarted. No live-GPU result is claimed; the modified CUDA translation unit was compiled/loaded and its host-side behavior was exercised CPU-isolated. |
malaiwah
changed the title
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
Actionable comments posted: 1
🧹 Nitpick comments (3)
tests/comm/test_pcie_oneshot_fused_rmsnorm_gpu.py (1)
75-89: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valueSlot-probe and alternation helpers are copy-pasted across three test modules.
tests/comm/test_pcie_twoshot.pyintroduced reusable_local_staging_words/_assert_alternating_slots; the oneshot tests re-implement the identical probe and inline the identical alternation assertions, so any future change to the slot contract must be edited in several places.
tests/comm/test_pcie_oneshot_fused_rmsnorm_gpu.py#L75-L89: move_local_eager_wordsinto a shared test helper module and import it here.tests/comm/test_pcie_oneshot_torture.py#L46-L60: delete the duplicate definition and import the shared helper.tests/comm/test_pcie_oneshot_fused_rmsnorm_gpu.py#L214-L230: replace the inlinechanged_slotscomprehension with the shared_assert_alternating_slots(snapshots)used by the twoshot test (and do the same attests/comm/test_pcie_oneshot_torture.pyLines 130-141).🤖 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/comm/test_pcie_oneshot_fused_rmsnorm_gpu.py` around lines 75 - 89, The slot-probe and alternation logic is duplicated across the oneshot tests. Move _local_eager_words into the shared helper module used by tests/comm/test_pcie_twoshot.py and import it in tests/comm/test_pcie_oneshot_fused_rmsnorm_gpu.py (75-89) and tests/comm/test_pcie_oneshot_torture.py (46-60), deleting both local definitions; replace the inline changed_slots assertions in tests/comm/test_pcie_oneshot_fused_rmsnorm_gpu.py (214-230) and tests/comm/test_pcie_oneshot_torture.py (130-141) with the shared _assert_alternating_slots(snapshots) helper.tests/comm/test_pcie_twoshot.py (1)
25-62: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winReuse the production layout math instead of re-deriving it.
pack_stride/scale_offset/scale_stride/slot_bytes/signal_bytesare recomputed here with a literal256, duplicatingPCIeTwoShotSP.from_exchange_group, which usesIPC_SLAB_ALIGNMENTand the same_align_up. If either the alignment constant or the slab layout changes, this probe silently reads the wrong offsets and the alternation assertions become vacuous. Import_align_up/IPC_SLAB_ALIGNMENT(or better, expose the computed offsets from the runtime) rather than mirroring the arithmetic.🤖 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/comm/test_pcie_twoshot.py` around lines 25 - 62, Update _local_staging_words to reuse the production layout calculations and alignment constant from PCIeTwoShotSP.from_exchange_group, importing _align_up and IPC_SLAB_ALIGNMENT or using exposed runtime offsets instead of duplicating the arithmetic and literal 256. Keep the probe’s source-offset and word-reading behavior unchanged while ensuring it tracks runtime layout changes.sparkinfer/comm/pcie/pcie_oneshot.cu (1)
187-204: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winSelector primitives copy-pasted into both translation units, each with an avoidable SMEM hop. Both files define the same
advance_staging_slotkernel (with a redundantblockIdx.x == 0 && threadIdx.x == 0guard for a<<<1,1>>>launch) and the same thread-0-plus-__syncthreads()selection helper, even thoughactive_staging_slotis a uniform, broadcast-cached load every thread can read directly. The sharedtwo_slot_selector.hheader — whose include was dropped from both files — is the natural home for this contract.
sparkinfer/comm/pcie/pcie_oneshot.cu#L187-L204: moveadvance_staging_slotand the slot-index accessor into the shared header, drop the redundant thread guard, and letselect_rank_dataread the slot per-thread and indexoptions.slots[slot]without__shared__ RankDataor the extra barrier.sparkinfer/comm/pcie/pcie_twoshot.cu#L91-L107: delete the duplicated kernel/helper and reuse the shared header versions, keeping the by-valueRankPtrsselection but without the SMEM staging copy.🤖 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 `@sparkinfer/comm/pcie/pcie_oneshot.cu` around lines 187 - 204, Consolidate the duplicated staging-slot primitives into shared two_slot_selector.h: for sparkinfer/comm/pcie/pcie_oneshot.cu lines 187-204, move advance_staging_slot and the slot-index accessor there, remove the redundant launch guard, and make select_rank_data read active_staging_slot per thread without shared-memory staging or __syncthreads(); for sparkinfer/comm/pcie/pcie_twoshot.cu lines 91-107, delete the duplicate kernel/helper and reuse the header versions while preserving by-value RankPtrs selection without the SMEM copy.
🤖 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 `@sparkinfer/comm/pcie/pcie_oneshot.cu`:
- Around line 741-751: Remove the per-entry advance_staging_slot GPU launch from
the allreduce path in sparkinfer/comm/pcie/pcie_oneshot.cu at lines 741-751 and
fold the generation advance into pcie_allreduce_kernel. Apply the corresponding
staging-slot reuse or FP8-kernel integration in
sparkinfer/comm/pcie/pcie_oneshot.cu lines 836-843 and
sparkinfer/comm/pcie/pcie_twoshot.cu lines 346-354 and 362-367; if any separate
advance launch remains, add real-target GPU benchmarks covering the requested
allreduce sizes, fused RMSNorm configurations, and representative twoshot shard
shapes.
---
Nitpick comments:
In `@sparkinfer/comm/pcie/pcie_oneshot.cu`:
- Around line 187-204: Consolidate the duplicated staging-slot primitives into
shared two_slot_selector.h: for sparkinfer/comm/pcie/pcie_oneshot.cu lines
187-204, move advance_staging_slot and the slot-index accessor there, remove the
redundant launch guard, and make select_rank_data read active_staging_slot per
thread without shared-memory staging or __syncthreads(); for
sparkinfer/comm/pcie/pcie_twoshot.cu lines 91-107, delete the duplicate
kernel/helper and reuse the header versions while preserving by-value RankPtrs
selection without the SMEM copy.
In `@tests/comm/test_pcie_oneshot_fused_rmsnorm_gpu.py`:
- Around line 75-89: The slot-probe and alternation logic is duplicated across
the oneshot tests. Move _local_eager_words into the shared helper module used by
tests/comm/test_pcie_twoshot.py and import it in
tests/comm/test_pcie_oneshot_fused_rmsnorm_gpu.py (75-89) and
tests/comm/test_pcie_oneshot_torture.py (46-60), deleting both local
definitions; replace the inline changed_slots assertions in
tests/comm/test_pcie_oneshot_fused_rmsnorm_gpu.py (214-230) and
tests/comm/test_pcie_oneshot_torture.py (130-141) with the shared
_assert_alternating_slots(snapshots) helper.
In `@tests/comm/test_pcie_twoshot.py`:
- Around line 25-62: Update _local_staging_words to reuse the production layout
calculations and alignment constant from PCIeTwoShotSP.from_exchange_group,
importing _align_up and IPC_SLAB_ALIGNMENT or using exposed runtime offsets
instead of duplicating the arithmetic and literal 256. Keep the probe’s
source-offset and word-reading behavior unchanged while ensuring it tracks
runtime layout changes.
🪄 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: Organization UI
Review profile: CHILL
Plan: Pro Plus
Run ID: b4e9998b-822a-4328-ae92-43481f880178
📒 Files selected for processing (5)
sparkinfer/comm/pcie/pcie_oneshot.cusparkinfer/comm/pcie/pcie_twoshot.cutests/comm/test_pcie_oneshot_fused_rmsnorm_gpu.pytests/comm/test_pcie_oneshot_torture.pytests/comm/test_pcie_twoshot.py
|
Final review update at
Verification: 33 passed, 2 opt-in GPU skips; fresh-cache one-shot/two-shot CUDA compile+load; Ruff format/lint pass. Multi-GPU execution is owned by the separate field-test agent; no rental/production result is claimed here. @coderabbitai review |
|
✅ Action performedReview finished.
|
Rental-machine validation — FAIL (strict safety review supersedes earlier PASS)Public PR head: The earlier comment correctly fixed the scratch-reuse test oracle, but its PASS conclusion was too narrow. A subsequent strict synchronization/lifecycle review found P1 defects that the numerical smoke does not make safe:
The control-node repair did pass useful bounded evidence: cold/warm communication union, fresh first-use capture at 2/4 ranks, fused RMS at 2/4 ranks, 1,025 opposite-order and scratch-reuse graph replays, and two-shot correctness. A four-pair AB/BA benchmark (4 ranks, 200 warmups, 2,000 aligned samples/run, fresh build each run) found eager unchanged ( Immutable evidence and chronology: A collective setup/teardown + operation-wide DCP control-node replacement is being tested now. I am withholding any push to this PR until that candidate passes deterministic skew/variable-grid DCP2/DCP4 tests and independent review. |
Frozen replacement qualification — PASS (candidate only)This is not a PASS for the current public #101 head Frozen source: Host: Vast.ai Exact commands and resultsWith
The decisive opposite-order gate collectively pre-prepared Raw logs are retained on the active rental under
Conclusion: the frozen replacement passes the live four-GPU gates relevant to |
|
Superseding field result: the frozen semantic-channel successor is not yet The focused four-rank eager/capture/torture and atomic vLLM/SparkInfer suites The engine failed before KV allocation/health, so there is no performance or The complete call trace makes the repair direction deliberately narrow:
Re-qualification will include the same full-model boot/traffic/shutdown gate, Complete immutable server log and analysis: |
|
Closing as superseded, not merged. Public head 0959fe8 remains rejected for standalone promotion. The reviewed replacement is #105 at bc62980, paired atomically with field vLLM f99e1e7b8636ca3811ab6d23084ac6da63420dc3 and its patch-identical clean upstream PR local-inference-lab/vllm#216. The final corrected pair passed the complete GLM gate; evidence and exact reproduction contract: https://github.com/malaiwah/glm52-exl3-vast/blob/5c76a2536e7fc9a5f1cb6bf182531889f5385e65/docs/field-review-results/2026-07-30-vast-46335896/UPSTREAM-REPAIR-CAMPAIGN.md and https://github.com/malaiwah/glm52-exl3-vast/blob/5c76a2536e7fc9a5f1cb6bf182531889f5385e65/docs/field-review-results/2026-07-30-vast-46335896/COUNTER-VALIDATION.md. Do not cherry-pick #101 alone and call it equivalent. |
1 participant
Summary
Fix every reachable signed
% 2staging counter identified in #97, and make reusable one-shot/two-shot staging ownership advance on CUDA graph execution rather than host invocation.Fixes #97.
Confirmed failure modes
The original DCP, one-shot, and two-shot host counters fed a signed pre-increment value to
% 2. After signed overflow,-1could index before the intended staging table.A second capture defect applied to the reusable collectives: host slot selection runs once while a CUDA graph is captured, so every replay reused that recorded slab. Adjacent odd-operation replays could overwrite staging still consumed by a peer.
Change
TwoSlotSelectorfor the independently mergeable overflow fix. PR fix(pcie): retire capture channel aliases #103 replaces this host selector with DCP execution-owned parity.0, 1, 0, 1, ....Invariants
Verification
Live multi-GPU execution is delegated to the separate field-test agent; no rental or production-GPU result is claimed here. Production was not rebuilt or restarted.
Review follow-up
b24d875: removed replay-frozen host selection from reusable collectives and added focused staging-observation regressions.b518259: addressed the performance review by folding launch-global generation selection into each existing collective kernel; fused graph kernel count remains one.Stack
PR #103 is stacked on this branch. Merge #101 first; #103 then removes the DCP-only host selector after migrating DCP to device execution parity.
Summary by CodeRabbit
Bug Fixes
Tests