fix(pcie): harden replay and CUDA IPC lifecycle - #105
Conversation
…60730' into HEAD # Conflicts: # sparkinfer/comm/pcie/pcie_dcp_a2a.cu
Rental-machine validation — PASSTested exact PR head: Tested lineage base: Host: Vast.ai Commands and resultsAll GPU commands used
Paired exact vLLM head Collect-only found 64 nodes; execution was 64 passed in 102.17 s. The first combined rejection command omitted the synthetic Observed invariants
Independent reviewers approved the collective symmetry, allocation safety, Raw logs are retained on the active rental under
Durable source: |
📝 WalkthroughWalkthroughThe PR adds PCIe control-node benchmarks, collective setup validation, retryable IPC lifecycle management, device-controlled double-buffer staging, semantic channel routing, CUDA graph checks, and distributed correctness tests. ChangesPCIe benchmark and reporting
Collective setup and resource lifecycle
Device-controlled staging and channel routing
Estimated code review effort: 5 (Critical) | ~120 minutes 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 |
|
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: |
|
@coderabbitai review |
✅ Action performedReview finished.
|
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
Actionable comments posted: 10
🧹 Nitpick comments (21)
tests/comm/test_pcie_dcp_a2a.py (1)
383-391: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valueRename the unused unpacked variable.
Ruff reports RUF059 for
requestedon Line 386. Use a dummy name to keep the lint clean.♻️ Proposed fix
def gather(local_state, group): if not local_state or isinstance(local_state[0], str): return [local_state, ()] - requested, existing = local_state + _requested, existing = local_state return [(("other",), existing), local_state]🤖 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_dcp_a2a.py` around lines 383 - 391, In the test helper function gather, rename the unused requested variable in the local_state unpacking to the project’s dummy-variable convention, while preserving the existing existing value and returned structures.Source: Linters/SAST tools
tests/comm/test_pcie_control_node_benchmark_contract.py (4)
45-71: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winAdd a case for the misaligned-vector guard.
The test covers the aligned path well. The anti-phase vectors prove that the per-iteration maximum differs from the old
max(mean(rank))aggregation, and line 65 records the old value as a contrast.
_distributed_critical_pathalso raises when rank vectors have different lengths. That guard protects against a corrupted or partially gathered record. Add a case with unequalsamples_uslengths and assertValueError.💚 Proposed additional case
assert result["graph"]["samples_us"] == [6.0, 6.0, 6.0, 6.0] assert result["graph"]["mean_us"] == 6.0 assert result["graph"]["cold_us"] == 8.0 + + misaligned = [ + { + "eager": {"cold_us": 11.0, "samples_us": [10.0, 1.0]}, + "graph": {"cold_us": 7.0, "samples_us": [6.0, 2.0]}, + }, + { + "eager": {"cold_us": 12.0, "samples_us": [1.0]}, + "graph": {"cold_us": 8.0, "samples_us": [2.0]}, + }, + ] + with pytest.raises(ValueError, match="empty or misaligned"): + worker["_distributed_critical_path"](misaligned)🤖 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_control_node_benchmark_contract.py` around lines 45 - 71, Add a test case for _distributed_critical_path using per-rank records whose samples_us vectors have different lengths, and assert that calling it raises ValueError. Keep the existing aligned-vector assertions unchanged and reuse the same eager/graph record structure as needed.
19-42: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winDocument why this test extracts functions instead of importing the module.
benchmarks/benchmark_pcie_oneshot_control_node.pyimportstorchandcuda.bindings.runtimeat module level. A normal import would make this CPU contract test require torch, cuda-python, and a GPU-capable environment. The AST extraction avoids that, and the injected namespace at lines 31-40 supplies exactly the globals the selected functions need.State that reason in a docstring. Without it, a maintainer can replace this helper with a plain import, and the contract test then stops running in CPU-only CI. Also note that any new global reference inside an extracted function fails later with a bare
NameErrorat call time.♻️ Proposed docstring
def _load_functions(relative_path: str, names: set[str]) -> dict[str, object]: + """Extract selected top-level functions without importing the module. + + Both benchmark files import torch and cuda.bindings at module scope. This + contract test must run without a GPU, so it compiles only the requested + functions and supplies their globals explicitly. Add any new global that a + selected function references to ``namespace`` below, or the call fails with + a bare NameError. + """ + path = _ROOT / relative_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 `@tests/comm/test_pcie_control_node_benchmark_contract.py` around lines 19 - 42, Add a docstring to _load_functions explaining that AST extraction avoids importing benchmark_pcie_oneshot_control_node.py and its torch/CUDA dependencies, allowing the contract test to run in CPU-only CI. Document that the injected namespace supplies the selected functions’ globals and that adding an unlisted global reference causes a runtime NameError when the extracted function is called.
100-108: 🎯 Functional Correctness | 🔵 Trivial | ⚡ Quick winAssert
delta_percentso the ratio direction is pinned.The test asserts
delta_usfor one eager metric and one graph metric. It does not assertdelta_percent, which is the field a reader uses to state the speedup or regression. A sign inversion or an inverted ratio in_paired_deltasat line 544 ofbenchmarks/run_pcie_oneshot_control_node_ab.pywould pass this test.For the given inputs, eager
delta_percentis20.0and graphdelta_percentis-10.0. Assert both. Negative means the candidate is faster.As per coding guidelines: "Benchmark the real target path before making performance claims, recording the command, commit, worktree, GPU mode, correctness state, raw timings, and ratio direction."
💚 Proposed assertions
assert deltas[0]["metrics"]["eager"]["mean_us"]["delta_us"] == 2.0 assert deltas[0]["metrics"]["graph"]["p95_us"]["delta_us"] == -2.0 + # A negative delta means the candidate is faster than the baseline. + assert deltas[0]["metrics"]["eager"]["mean_us"]["delta_percent"] == 20.0 + assert deltas[0]["metrics"]["graph"]["p95_us"]["delta_percent"] == -10.0 + assert deltas[0]["metrics"]["eager"]["mean_us"]["baseline_us"] == 10.0 + assert deltas[0]["metrics"]["eager"]["mean_us"]["candidate_us"] == 12.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 `@tests/comm/test_pcie_control_node_benchmark_contract.py` around lines 100 - 108, Extend the test around controller["_paired_deltas"] to assert delta_percent for both metrics: eager mean_us must be 20.0 and graph p95_us must be -10.0. Keep the existing delta_us assertions and ensure the expectations pin the ratio direction, with negative values indicating a faster candidate.Source: Coding guidelines
111-121: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valueThe test verifies the
TimeoutError, not group termination.The test name states that the command process group is terminated. The assertion only matches the exception message. A
_runthat calledos.kill(process.pid, ...)instead ofos.killpgwould still pass, and orphanedtorch.distributed.runranks are the failure this test exists to prevent.To verify the group, launch a child that spawns a grandchild and prints its PID, then assert the grandchild is gone after
_runraises. If that is too costly, rename the test to describe what it checks.🤖 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_control_node_benchmark_contract.py` around lines 111 - 121, The test test_controller_timeout_terminates_the_command_process_group currently verifies only the TimeoutError message, not process-group termination. Strengthen it by launching a child that spawns a grandchild and reports its PID, then after controller["_run"] raises assert that the grandchild no longer exists; otherwise rename the test to reflect that it checks only the timeout error.benchmarks/run_pcie_oneshot_control_node_ab.py (3)
199-217: 🩺 Stability & Availability | 🔵 Trivial | ⚡ Quick winReap the child process group on every exit path.
The timeout path at lines 201-213 terminates the group correctly. Other exception paths do not. If
communicateat line 200 raisesKeyboardInterrupt, the spawnedtorch.distributed.rungroup survives. Orphaned ranks keep their GPU allocations and block the next run in the AB/BA sequence.Wrap the body so any exception terminates the group. Add a timeout to the final
communicateat line 209 so a held pipe cannot block the controller.♻️ Proposed guard
try: stdout, _ = process.communicate(timeout=timeout) except subprocess.TimeoutExpired: with suppress(ProcessLookupError): os.killpg(process.pid, signal.SIGTERM) try: stdout, _ = process.communicate(timeout=10) except subprocess.TimeoutExpired: with suppress(ProcessLookupError): os.killpg(process.pid, signal.SIGKILL) - stdout, _ = process.communicate() + with suppress(subprocess.TimeoutExpired): + stdout, _ = process.communicate(timeout=10) raise TimeoutError( f"command exceeded {timeout:.1f}s and its process group was terminated: " f"{' '.join(args)}\n{stdout}" ) from None + except BaseException: + with suppress(ProcessLookupError): + os.killpg(process.pid, signal.SIGKILL) + with suppress(subprocess.TimeoutExpired): + process.communicate(timeout=10) + raiseNote that the
except subprocess.TimeoutExpiredbranch reassignsstdoutbefore theraise, so keep it ahead of the broad handler.🤖 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 `@benchmarks/run_pcie_oneshot_control_node_ab.py` around lines 199 - 217, Update the subprocess execution flow around process.communicate so every exception path, including KeyboardInterrupt, terminates and reaps the process group. Keep the existing subprocess.TimeoutExpired handler ahead of the broad exception handler, and add a bounded timeout to its final communicate call so pipe closure cannot block indefinitely; preserve the current timeout escalation and error reporting.
748-764: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winDetect stack drift after each run instead of after the whole suite.
The fingerprint checks run once, after all
2 * pairsruns finish. If the driver, software stack, or environment changes during run 2 of 8, the controller still executes the remaining six runs and then discards the suite. Each run costs a full JIT build plus 1000 timed iterations per mode.Compare each new record against the first record inside the run loop, directly after
_validate_run. Keep the final aggregate check as it is.🤖 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 `@benchmarks/run_pcie_oneshot_control_node_ab.py` around lines 748 - 764, Add per-run drift validation inside the run loop immediately after _validate_run, comparing the new record’s software, _stable_hardware, and _stable_environment fingerprints against the first record. Raise the same ValueError messages as the existing aggregate checks on any mismatch, while preserving the final aggregate fingerprint validation unchanged.
766-791: 🚀 Performance & Scalability | 🔵 TrivialConsider gating the suite on throttle evidence.
NVIDIA_SMI_FIELDSalready recordspstate,clocks.current.sm,clocks.max.sm, andclocks_event_reasons.activefor every sample. Nothing rejects a suite in which a timing phase ran while a throttle reason was active. A throttled candidate run paired against an unthrottled baseline run produces a paired delta that reflects clock state, not the algorithm.Add a derived field to
paired_summarythat reports the throttle reasons observed during theeager_timingandgraph_timingphases of each record. Then a reader can discard affected pairs without re-running the suite.🤖 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 `@benchmarks/run_pcie_oneshot_control_node_ab.py` around lines 766 - 791, Update the paired_summary construction via _paired_summary(paired_deltas) to include derived throttle reasons observed in each record’s eager_timing and graph_timing phases, using the existing NVIDIA_SMI_FIELDS data such as pstate, current/max SM clocks, and active clock-event reasons. Preserve the existing summary fields and expose enough per-record or per-pair detail for readers to discard pairs with throttle evidence.benchmarks/benchmark_pcie_oneshot_control_node.py (2)
1151-1156: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valueExpose sampler liveness instead of reading
sampler._thread.Line 1152 reads the private
_threadattribute of_NvidiaSmiSamplerfrommain(). Add a public method on the class and call it here. This also letsstop()become idempotent, so a secondstop()in thefinallypath cannot mask the original exception.♻️ Proposed accessor
def stop(self) -> None: + if self._stop.is_set(): + return self._stop.set() self._thread.join(timeout=15) if self._thread.is_alive(): raise RuntimeError("nvidia-smi sampler thread did not stop within 15s") self._phase = "final" self._sample() + + def is_running(self) -> bool: + return self._thread.is_alive()finally: - if sampler is not None and sampler._thread.is_alive(): + if sampler is not None and sampler.is_running(): sampler.stop()🤖 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 `@benchmarks/benchmark_pcie_oneshot_control_node.py` around lines 1151 - 1156, Replace main()’s direct sampler._thread access with a public liveness method on _NvidiaSmiSampler, and use that method in the finally cleanup before stopping the sampler. Make stop() idempotent so repeated calls safely return without masking the original exception.
932-940: 🩺 Stability & Availability | 🔵 Trivial | ⚡ Quick winBind the NCCL device before initialization.
dist.init_process_group("nccl", ...)runs beforetorch.cuda.set_device(local_rank). Passdevice_idtoinit_process_groupso NCCL binds the communicator during initialization instead of waiting for the first collective.♻️ Proposed reorder
+ local_rank = int(os.environ["LOCAL_RANK"]) + torch.cuda.set_device(local_rank) + device = torch.device(f"cuda:{local_rank}") dist.init_process_group( "nccl", timeout=timedelta(seconds=args.distributed_timeout_seconds), + device_id=device, ) rank = dist.get_rank() world_size = dist.get_world_size() - local_rank = int(os.environ["LOCAL_RANK"]) - �torch.cuda.set_device(local_rank) - device = torch.device(f"cuda:{local_rank}")🤖 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 `@benchmarks/benchmark_pcie_oneshot_control_node.py` around lines 932 - 940, Update the initialization flow around dist.init_process_group to determine the local CUDA device before initializing NCCL, then pass that device through the init_process_group device_id parameter. Preserve the existing local_rank and device setup while ensuring the communicator binds to the selected GPU during initialization.sparkinfer/comm/pcie/pcie_twoshot.cu (1)
297-297: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winValue-initialize
staging_.
staging_is not brace-initialized, and the constructor fills only indices[0, world_size_). Entries[world_size_, kMaxRanks)stay indeterminate, and the whole struct is copied by value into every kernel launch. The kernels read onlypeer < ngpus, so no wrong pointer is dereferenced today; the safety depends entirely on that invariant. The sibling runtimes already value-initialize the equivalent member (DoubleStaging staging_{}insparkinfer/comm/pcie/pcie_dcp_a2a.cu,DoubleRankData dbuf_data_ = {}insparkinfer/comm/pcie/pcie_oneshot.cu). Match them.♻️ Proposed change
- DoubleRankPtrs staging_; + DoubleRankPtrs staging_{};🤖 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_twoshot.cu` at line 297, Value-initialize the DoubleRankPtrs member staging_ at its declaration, matching the initialization pattern used by the sibling PCIe runtimes, while preserving the constructor’s existing population of indices below world_size_.tests/comm/test_pcie_staging_control.py (4)
137-143: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winThe second assertion is a tautology.
worker_launchesonly keeps lines that contain<<<blocks, threads, 0, stream>>>. Every element therefore contains, 0, stream>>>by construction, so line 143 can never fail. The stated intent is that no peer-waiting worker launch requests dynamic shared memory. Collect launches by the grid and block arguments, then assert the shared-memory argument separately.♻️ Proposed change
- worker_launches = [ - line - for line in source.splitlines() - if "<<<blocks, threads, 0, stream>>>" in line - ] - assert worker_launches - assert all(", 0, stream>>>" in line for line in worker_launches) + worker_launches = [ + line for line in source.splitlines() if "<<<blocks, threads," in line + ] + assert worker_launches + assert all( + "<<<blocks, threads, 0, stream>>>" in line for line in worker_launches + )🤖 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_staging_control.py` around lines 137 - 143, Update the worker-launch collection and assertions in the PCIe staging control test so filtering uses the grid and block arguments without requiring the shared-memory value. Assert separately that each collected launch uses zero for the dynamic shared-memory argument, preserving the intended validation that peer-waiting worker launches do not request shared memory.
22-67: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valueThese tests assert a local Python model, not the shipped code.
test_slot_control_alternates_across_replays_and_wraps,test_same_stream_capture_replay_handles_variable_worker_grids,test_multistream_separate_channels_match_across_rank_interleavings, andtest_control_node_graph_topology_contractonly exercise_advanceand a dict literal defined in this file. They cannot fail if the CUDA sources diverge from the modeled invariant. The source-text tests below carry the real contract.Keep them if the intent is executable documentation of the slot invariant, and state that intent in each docstring. Otherwise bind them to the sources, for example by asserting that the modeled expression matches the literal in
advance_staging_slot_kernelastest_control_kernel_is_graph_replay_dynamicalready does.Also applies to: 250-273
🤖 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_staging_control.py` around lines 22 - 67, The listed tests only validate local Python models and do not verify the shipped CUDA implementation. Either add docstrings to test_slot_control_alternates_across_replays_and_wraps, test_same_stream_capture_replay_handles_variable_worker_grids, test_multistream_separate_channels_match_across_rank_interleavings, and test_control_node_graph_topology_contract explicitly documenting their executable-documentation purpose, or bind their assertions to the corresponding source expressions in advance_staging_slot_kernel, consistent with test_control_kernel_is_graph_replay_dynamic.
235-247: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winExtend the immediate-error-check test to
pcie_dcp_a2a.cu.The parametrization covers
pcie_oneshot.cuandpcie_twoshot.cuonly.sparkinfer/comm/pcie/pcie_dcp_a2a.cualso launches the control kernel twice and does placeCHECK_CUDA_SUCCESS(cudaGetLastError());immediately after each launch, but no test guards that. The DCP control launch passesself_signal_instead ofself_sg_, so the literal must be parametrized with the launch string.♻️ Proposed change
`@pytest.mark.parametrize`( - ("source_name", "expected_checks"), - (("pcie_oneshot.cu", 4), ("pcie_twoshot.cu", 4)), + ("source_name", "signal_name", "expected_checks"), + ( + ("pcie_oneshot.cu", "self_sg_", 4), + ("pcie_twoshot.cu", "self_sg_", 4), + ("pcie_dcp_a2a.cu", "self_signal_", 4), + ), ) def test_control_and_worker_launches_have_immediate_error_checks( - source_name: str, expected_checks: int + source_name: str, signal_name: str, expected_checks: int ) -> None: source = (PCIE / source_name).read_text(encoding="utf-8") assert source.count("CHECK_CUDA_SUCCESS(cudaGetLastError());") >= expected_checks - control_launch = "advance_staging_slot_kernel<<<1, 1, 0, stream>>>(self_sg_);" + control_launch = ( + f"advance_staging_slot_kernel<<<1, 1, 0, stream>>>({signal_name});" + ) for suffix in source.split(control_launch)[1:]: assert suffix.lstrip().startswith("CHECK_CUDA_SUCCESS(cudaGetLastError());")🤖 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_staging_control.py` around lines 235 - 247, Extend test_control_and_worker_launches_have_immediate_error_checks to include pcie_dcp_a2a.cu, parameterizing the control launch literal because its advance_staging_slot_kernel invocation uses self_signal_ rather than self_sg_. Keep the existing expected check count and immediate CHECK_CUDA_SUCCESS validation for all covered sources.
156-156: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valueWhitespace-exact matching makes these assertions fragile.
Lines 156 and 164 embed the exact newline and six-space indentation of the CUDA sources. A reformat of
sparkinfer/comm/pcie/pcie_oneshot.cubreaks the test without any behavior change. The intent is that the control launch is inside the staged branch and precedes the worker dispatch. The followingindex()comparisons already prove the ordering. Normalize whitespace before matching so the guard survives formatting.♻️ Proposed change
+def _collapse(text: str) -> str: + return " ".join(text.split()) + + def test_oneshot_control_is_staged_only_and_precedes_each_worker() -> None:Then compare collapsed text:
- assert f"if (stage_input) {{\n {control_launch}" in regular + assert _collapse(f"if (stage_input) {{ {control_launch}") in _collapse(regular)Also applies to: 164-165
🤖 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_staging_control.py` at line 156, Update the assertions around the staged control and worker dispatch checks in the test to normalize or collapse whitespace before matching, rather than embedding exact CUDA newlines and indentation. Preserve the existing checks that the control launch is within the stage_input branch and that ordering is validated by the index comparisons.sparkinfer/comm/pcie/ipc_handle_registry.h (1)
41-56: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valueConsider erasing successfully closed entries instead of clearing them.
close_all_noexceptsets each closed handle toHandle{}and keeps the map entry.findtherefore still returns a hit for that key, andopen_ipc_handleinsparkinfer/comm/pcie/pcie_oneshot.cureturns that cleared value without reopening. That path is currently unreachable becausedisposedeletes the runtime, butclose_ipc_handles_strictcan throw and keep the object alive. Erasing closed entries removes the hazard and keeps the retry semantics identical, because only failed entries must persist.♻️ Proposed change
Error close_all_noexcept() noexcept { Error first_error = Success; - for (auto& entry : handles_) { - Handle& handle = entry.second; - if (handle == Handle{}) { - continue; - } - const Error error = close_(handle); - if (error == Success) { - handle = Handle{}; - } else if (first_error == Success) { - first_error = error; - } - } + for (auto it = handles_.begin(); it != handles_.end();) { + const Error error = close_(it->second); + if (error == Success) { + it = handles_.erase(it); + continue; + } + if (first_error == Success) { + first_error = error; + } + ++it; + } return first_error; }🤖 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/ipc_handle_registry.h` around lines 41 - 56, Update close_all_noexcept() to erase each map entry after close_ succeeds instead of assigning Handle{} while retaining the entry; preserve failed entries and first_error tracking so retries continue to operate only on handles that were not closed.sparkinfer/comm/pcie/pcie_dcp_a2a.cu (1)
161-169: 🚀 Performance & Scalability | 🔵 Trivial | ⚡ Quick winCache the test-delay environment reads and gate the extra block barrier.
Two costs land on every production launch here.
First,
env_intcallsstd::getenvandstd::atoion eachrun()andall_gather_heads()invocation. The neighboring overrides (dcp_threads_override,dcp_block_limit_override) already cache their value in a function-local static. These collectives are microsecond-scale, so repeated environment lookups add avoidable host work in the launch path.Second,
test_post_barrier_delayends with an unconditional__syncthreads().start_barrieralready ends with__syncthreads(), so whendelay_cycles == 0the worker kernels execute a redundant block-wide barrier. Move the barrier inside thedelay_cycles != 0branch.This comment reports avoidable work, not a measured regression. Record a benchmark before you claim a latency delta.
♻️ Proposed change
DINLINE void test_post_barrier_delay(int rank, int delayed_rank, uint64_t delay_cycles) { - if (delay_cycles != 0 && rank == delayed_rank && threadIdx.x == 0) { - const uint64_t start = clock64(); - while (clock64() - start < delay_cycles) { - } - } - __syncthreads(); + if (delay_cycles == 0) return; + if (rank == delayed_rank && threadIdx.x == 0) { + const uint64_t start = clock64(); + while (clock64() - start < delay_cycles) { + } + } + __syncthreads(); } + +static int dcp_test_delay_rank() { + static const int value = env_int("SPARKINFER_PCIE_DCP_TEST_DELAY_RANK", -1); + return value; +} + +static uint64_t dcp_test_delay_cycles() { + static const uint64_t value = static_cast<uint64_t>(std::max( + 0, env_int("SPARKINFER_PCIE_DCP_TEST_POST_BARRIER_DELAY_CYCLES", 0))); + return value; +}Then use the cached accessors at both launch sites:
- const int delayed_rank = - env_int("SPARKINFER_PCIE_DCP_TEST_DELAY_RANK", -1); - const uint64_t delay_cycles = static_cast<uint64_t>(std::max( - 0, env_int("SPARKINFER_PCIE_DCP_TEST_POST_BARRIER_DELAY_CYCLES", 0))); + const int delayed_rank = dcp_test_delay_rank(); + const uint64_t delay_cycles = dcp_test_delay_cycles();Note:
tests/comm/test_pcie_staging_control.pyasserts the literal string"test_post_barrier_delay(rank, delayed_rank, delay_cycles);", which this change preserves.Also applies to: 468-471, 533-536
🤖 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_dcp_a2a.cu` around lines 161 - 169, Cache the test-delay environment value in a function-local static accessor, matching dcp_threads_override and dcp_block_limit_override, and use that accessor at both run() and all_gather_heads() launch sites instead of repeated env_int reads. Update test_post_barrier_delay so __syncthreads() executes only when delay_cycles is nonzero, while preserving the existing call expression string and recording a benchmark before reporting any latency change.Source: Coding guidelines
tests/comm/test_pcie_oneshot.py (1)
586-591: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winGuard the retained-export registry against cross-test leakage. Many new tests in both files read and assert the module-global
_RETAINED_FAILED_IPC_EXPORTS, and several end withassert _RETAINED_FAILED_IPC_EXPORTS == {}. The root cause is shared: this dict is process-global state that no fixture resets. A single test that raises before callingretry()leaves a ticket behind, and every later== {}assertion in either file then fails with a message that points at the wrong test. Ordering and-p no:randomlybehavior also become significant.
tests/comm/test_pcie_oneshot.py#L586-L591: add an autouse fixture in this module that clears_RETAINED_FAILED_IPC_EXPORTSbefore and after each test, so the emptiness assertions describe only the test under execution.tests/comm/test_pcie_twoshot.py#L119-L123: add the same autouse fixture in this module, or move a shared fixture intotests/comm/conftest.pyand use it from both files.♻️ Proposed fixture
`@pytest.fixture`(autouse=True) def _clear_retained_ipc_setup(): _RETAINED_FAILED_IPC_EXPORTS.clear() yield _RETAINED_FAILED_IPC_EXPORTS.clear()🤖 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.py` around lines 586 - 591, Add an autouse fixture that clears _RETAINED_FAILED_IPC_EXPORTS before and after every test, either independently in tests/comm/test_pcie_oneshot.py at lines 586-591 and tests/comm/test_pcie_twoshot.py at lines 119-123, or as a shared fixture in tests/comm/conftest.py used by both modules. Ensure cleanup runs even when a test fails, preserving each test’s registry assertions as isolated.sparkinfer/comm/pcie/pcie_oneshot.cu (1)
636-709: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valueHoist the device and SM-count query out of the per-mode helper.
fused_resident_grid_capacitycallscudaGetDeviceandcudaDeviceGetAttributeon every invocation. The constructor invokes it nine times (three dtypes by three modes), so it repeats the same two device queries nine times. Query the device and SM count once in the constructor and pass the SM count to the helper. This is construction-time only, so the gain is code clarity rather than measurable latency.🤖 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 636 - 709, Update PCIeAllreduce construction and fused_resident_capacities usage so cudaGetDevice and cudaDeviceGetAttribute are performed once in the constructor, then pass the cached SM count into fused_resident_capacities for each dtype and mode instead of repeating device queries inside the per-mode helper.tests/comm/test_pcie_twoshot.py (1)
126-163: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winImport the alignment constant instead of hardcoding 256.
_local_staging_wordsrecomputes the slab layout with the literal256in three places. The production layout insparkinfer/comm/pcie/pcie_twoshot.pyusesIPC_SLAB_ALIGNMENTfor the same three alignments. If that constant changes, this helper reads the wrong device offsets, and_assert_alternating_slotsthen compares unrelated words.tests/comm/test_pcie_oneshot.pyalready importsIPC_SLAB_ALIGNMENTfromsparkinfer.comm.pcie.pcie_oneshot, so the constant is available.♻️ Proposed change
from sparkinfer.comm.pcie.pcie_oneshot import ( + IPC_SLAB_ALIGNMENT, PCIeOneshotAllReduce, _OwnedSharedBuffer, _RETAINED_FAILED_IPC_EXPORTS, )pack_stride = _align_up(max_rows_per_rank * (pool.row_elems // 16), 16) payload_bytes = pool.world_size * pack_stride * 16 - scale_offset = _align_up(payload_bytes, 256) + scale_offset = _align_up(payload_bytes, IPC_SLAB_ALIGNMENT) scale_stride = _align_up(max_rows_per_rank, 64) slot_bytes = _align_up( scale_offset + pool.world_size * scale_stride * 4, - 256, + IPC_SLAB_ALIGNMENT, ) - signal_bytes = _align_up(int(pool._ext.meta_size()), 256) + signal_bytes = _align_up(int(pool._ext.meta_size()), IPC_SLAB_ALIGNMENT)🤖 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 126 - 163, Update _local_staging_words to import and use IPC_SLAB_ALIGNMENT from the PCIe communication module for all three slab-alignment calculations, replacing the hardcoded 256 values while preserving the existing offset computation.tests/comm/test_pcie_oneshot_torture.py (1)
46-60: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winExtract the duplicated
_local_eager_wordsprobe into one shared test helper. Both GPU tests define the same eager staging-slot probe, including theoffsetparameter and thec_uint64word size. The shared root cause is the absence of a common helper module for PCIe staging inspection. Independent copies will diverge when the eager slot layout changes.
tests/comm/test_pcie_oneshot_torture.py#L46-L60: move this helper into a shared module undertests/comm/and import it here.tests/comm/test_pcie_oneshot_fused_rmsnorm_gpu.py#L80-L94: delete the local copy and import the shared helper.🤖 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_torture.py` around lines 46 - 60, Extract the duplicated _local_eager_words probe into a shared helper module under tests/comm, preserving its offset parameter, c_uint64 word handling, and synchronization behavior. In tests/comm/test_pcie_oneshot_torture.py lines 46-60, move the helper and import it; in tests/comm/test_pcie_oneshot_fused_rmsnorm_gpu.py lines 80-94, remove the local copy and import the shared helper.
🤖 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 `@benchmarks/benchmark_pcie_oneshot_control_node.py`:
- Around line 979-986: Update the graph capture context around pool.capture in
the benchmark to pass a stable, identical channel_id on every rank. Use that
shared identifier in the capture call while preserving the existing stream and
CUDA graph capture behavior.
- Around line 496-504: Update the fallback handling around
cudaGraphNodeGetParams so its result is validated before accessing
node_params.kernel. Preserve the existing error path by raising the CUDA failure
when either parameter query fails, and only assign params from the fallback
union after a successful call.
In `@sparkinfer/comm/pcie/pcie_twoshot.cu`:
- Around line 352-358: Validate the threads parameter at the start of both
reduce_scatter and all_gather, before calculating shard_packs or performing any
division: require it to be within [world_size_, 1024] and divisible by 32,
matching the PCIE DCP A2A guard. Preserve the existing block_limit capacity
rejection after this validation and before the shard calculation.
In `@tests/comm/test_pcie_dcp_a2a_gpu.py`:
- Around line 702-704: Update
test_pcie_dcp_a2a_rejects_reduced_sm_slice_before_ipc_allocation to assert that
SPARKINFER_PCIE_TEST_VISIBLE_SM_COUNT is configured to a positive value below
DCP_A2A_REQUIRED_SMS before spawning workers; fail with a clear configuration
message instead of allowing the test to reach DID NOT RAISE or
unexpected_cuda_rt.
In `@tests/comm/test_pcie_ipc_lifecycle.py`:
- Around line 383-386: Update the match pattern in the pytest.raises assertion
for the coordinated PCIe close failure to use a raw string literal, preserving
the existing regex behavior and text.
In `@tests/comm/test_pcie_oneshot_fused_rmsnorm_gpu.py`:
- Around line 128-139: Update the cudaGraphGetEdges usage in the graph-edge
query test: allocate a scalar cudaSizeT count buffer and pass it to the sizing
call, read num_edges from that buffer, then allocate the required
source/destination output arrays and pass them with num_edges to the fetch call.
Preserve the existing success and returned-edge assertions.
In `@tests/comm/test_pcie_oneshot_opposite_order_gpu.py`:
- Around line 245-255: Update the cleanup in the test’s finally block so
dist.destroy_process_group() is guaranteed to run even when pool.close() raises:
nest the process-group destruction in an inner finally around pool.close(),
preserving the existing teardown order.
In `@tests/comm/test_pcie_oneshot.py`:
- Around line 1450-1454: Update the gather helper to stop binding the unused
requested tuple element; preserve existing while using an underscore-prefixed
variable or tuple indexing to satisfy Ruff RUF059.
- Line 713: Update the match pattern in the pytest.raises call to use a raw
string literal while preserving the existing regex text and expected
RuntimeError behavior.
In `@tests/comm/test_pcie_preallocation_setup.py`:
- Around line 477-485: Update the test cleanup after the assertions around the
DCP runtime to clear the runtime pointer before destruction, matching the
one-shot test’s cleanup behavior near the earlier setup. Ensure _ptr is reset
from 456 to the non-quarantining value so PCIeDCPA2A.__del__ does not retain the
runtime in _ABANDONED_PCIE_RUNTIME_QUARANTINE.
---
Nitpick comments:
In `@benchmarks/benchmark_pcie_oneshot_control_node.py`:
- Around line 1151-1156: Replace main()’s direct sampler._thread access with a
public liveness method on _NvidiaSmiSampler, and use that method in the finally
cleanup before stopping the sampler. Make stop() idempotent so repeated calls
safely return without masking the original exception.
- Around line 932-940: Update the initialization flow around
dist.init_process_group to determine the local CUDA device before initializing
NCCL, then pass that device through the init_process_group device_id parameter.
Preserve the existing local_rank and device setup while ensuring the
communicator binds to the selected GPU during initialization.
In `@benchmarks/run_pcie_oneshot_control_node_ab.py`:
- Around line 199-217: Update the subprocess execution flow around
process.communicate so every exception path, including KeyboardInterrupt,
terminates and reaps the process group. Keep the existing
subprocess.TimeoutExpired handler ahead of the broad exception handler, and add
a bounded timeout to its final communicate call so pipe closure cannot block
indefinitely; preserve the current timeout escalation and error reporting.
- Around line 748-764: Add per-run drift validation inside the run loop
immediately after _validate_run, comparing the new record’s software,
_stable_hardware, and _stable_environment fingerprints against the first record.
Raise the same ValueError messages as the existing aggregate checks on any
mismatch, while preserving the final aggregate fingerprint validation unchanged.
- Around line 766-791: Update the paired_summary construction via
_paired_summary(paired_deltas) to include derived throttle reasons observed in
each record’s eager_timing and graph_timing phases, using the existing
NVIDIA_SMI_FIELDS data such as pstate, current/max SM clocks, and active
clock-event reasons. Preserve the existing summary fields and expose enough
per-record or per-pair detail for readers to discard pairs with throttle
evidence.
In `@sparkinfer/comm/pcie/ipc_handle_registry.h`:
- Around line 41-56: Update close_all_noexcept() to erase each map entry after
close_ succeeds instead of assigning Handle{} while retaining the entry;
preserve failed entries and first_error tracking so retries continue to operate
only on handles that were not closed.
In `@sparkinfer/comm/pcie/pcie_dcp_a2a.cu`:
- Around line 161-169: Cache the test-delay environment value in a
function-local static accessor, matching dcp_threads_override and
dcp_block_limit_override, and use that accessor at both run() and
all_gather_heads() launch sites instead of repeated env_int reads. Update
test_post_barrier_delay so __syncthreads() executes only when delay_cycles is
nonzero, while preserving the existing call expression string and recording a
benchmark before reporting any latency change.
In `@sparkinfer/comm/pcie/pcie_oneshot.cu`:
- Around line 636-709: Update PCIeAllreduce construction and
fused_resident_capacities usage so cudaGetDevice and cudaDeviceGetAttribute are
performed once in the constructor, then pass the cached SM count into
fused_resident_capacities for each dtype and mode instead of repeating device
queries inside the per-mode helper.
In `@sparkinfer/comm/pcie/pcie_twoshot.cu`:
- Line 297: Value-initialize the DoubleRankPtrs member staging_ at its
declaration, matching the initialization pattern used by the sibling PCIe
runtimes, while preserving the constructor’s existing population of indices
below world_size_.
In `@tests/comm/test_pcie_control_node_benchmark_contract.py`:
- Around line 45-71: Add a test case for _distributed_critical_path using
per-rank records whose samples_us vectors have different lengths, and assert
that calling it raises ValueError. Keep the existing aligned-vector assertions
unchanged and reuse the same eager/graph record structure as needed.
- Around line 19-42: Add a docstring to _load_functions explaining that AST
extraction avoids importing benchmark_pcie_oneshot_control_node.py and its
torch/CUDA dependencies, allowing the contract test to run in CPU-only CI.
Document that the injected namespace supplies the selected functions’ globals
and that adding an unlisted global reference causes a runtime NameError when the
extracted function is called.
- Around line 100-108: Extend the test around controller["_paired_deltas"] to
assert delta_percent for both metrics: eager mean_us must be 20.0 and graph
p95_us must be -10.0. Keep the existing delta_us assertions and ensure the
expectations pin the ratio direction, with negative values indicating a faster
candidate.
- Around line 111-121: The test
test_controller_timeout_terminates_the_command_process_group currently verifies
only the TimeoutError message, not process-group termination. Strengthen it by
launching a child that spawns a grandchild and reports its PID, then after
controller["_run"] raises assert that the grandchild no longer exists; otherwise
rename the test to reflect that it checks only the timeout error.
In `@tests/comm/test_pcie_dcp_a2a.py`:
- Around line 383-391: In the test helper function gather, rename the unused
requested variable in the local_state unpacking to the project’s dummy-variable
convention, while preserving the existing existing value and returned
structures.
In `@tests/comm/test_pcie_oneshot_torture.py`:
- Around line 46-60: Extract the duplicated _local_eager_words probe into a
shared helper module under tests/comm, preserving its offset parameter, c_uint64
word handling, and synchronization behavior. In
tests/comm/test_pcie_oneshot_torture.py lines 46-60, move the helper and import
it; in tests/comm/test_pcie_oneshot_fused_rmsnorm_gpu.py lines 80-94, remove the
local copy and import the shared helper.
In `@tests/comm/test_pcie_oneshot.py`:
- Around line 586-591: Add an autouse fixture that clears
_RETAINED_FAILED_IPC_EXPORTS before and after every test, either independently
in tests/comm/test_pcie_oneshot.py at lines 586-591 and
tests/comm/test_pcie_twoshot.py at lines 119-123, or as a shared fixture in
tests/comm/conftest.py used by both modules. Ensure cleanup runs even when a
test fails, preserving each test’s registry assertions as isolated.
In `@tests/comm/test_pcie_staging_control.py`:
- Around line 137-143: Update the worker-launch collection and assertions in the
PCIe staging control test so filtering uses the grid and block arguments without
requiring the shared-memory value. Assert separately that each collected launch
uses zero for the dynamic shared-memory argument, preserving the intended
validation that peer-waiting worker launches do not request shared memory.
- Around line 22-67: The listed tests only validate local Python models and do
not verify the shipped CUDA implementation. Either add docstrings to
test_slot_control_alternates_across_replays_and_wraps,
test_same_stream_capture_replay_handles_variable_worker_grids,
test_multistream_separate_channels_match_across_rank_interleavings, and
test_control_node_graph_topology_contract explicitly documenting their
executable-documentation purpose, or bind their assertions to the corresponding
source expressions in advance_staging_slot_kernel, consistent with
test_control_kernel_is_graph_replay_dynamic.
- Around line 235-247: Extend
test_control_and_worker_launches_have_immediate_error_checks to include
pcie_dcp_a2a.cu, parameterizing the control launch literal because its
advance_staging_slot_kernel invocation uses self_signal_ rather than self_sg_.
Keep the existing expected check count and immediate CHECK_CUDA_SUCCESS
validation for all covered sources.
- Line 156: Update the assertions around the staged control and worker dispatch
checks in the test to normalize or collapse whitespace before matching, rather
than embedding exact CUDA newlines and indentation. Preserve the existing checks
that the control launch is within the stage_input branch and that ordering is
validated by the index comparisons.
In `@tests/comm/test_pcie_twoshot.py`:
- Around line 126-163: Update _local_staging_words to import and use
IPC_SLAB_ALIGNMENT from the PCIe communication module for all three
slab-alignment calculations, replacing the hardcoded 256 values while preserving
the existing offset computation.
🪄 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: 9ec37b57-c717-4ab8-967b-54a958e5b957
📒 Files selected for processing (23)
benchmarks/benchmark_pcie_oneshot_control_node.pybenchmarks/run_pcie_oneshot_control_node_ab.pysparkinfer/comm/pcie/ipc_handle_registry.hsparkinfer/comm/pcie/pcie_dcp_a2a.cusparkinfer/comm/pcie/pcie_dcp_a2a.pysparkinfer/comm/pcie/pcie_oneshot.cusparkinfer/comm/pcie/pcie_oneshot.pysparkinfer/comm/pcie/pcie_twoshot.cusparkinfer/comm/pcie/pcie_twoshot.pytests/comm/test_pcie.pytests/comm/test_pcie_collective_validation_gpu.pytests/comm/test_pcie_control_node_benchmark_contract.pytests/comm/test_pcie_dcp_a2a.pytests/comm/test_pcie_dcp_a2a_gpu.pytests/comm/test_pcie_ipc_lifecycle.pytests/comm/test_pcie_native_teardown.pytests/comm/test_pcie_oneshot.pytests/comm/test_pcie_oneshot_fused_rmsnorm_gpu.pytests/comm/test_pcie_oneshot_opposite_order_gpu.pytests/comm/test_pcie_oneshot_torture.pytests/comm/test_pcie_preallocation_setup.pytests/comm/test_pcie_staging_control.pytests/comm/test_pcie_twoshot.py
Harden launch validation, benchmark semantic channels, graph inspection, and test teardown after CodeRabbit review.\n\nThe exact GPU-qualified provenance remains bc62980; this commit is a post-qualification review layer and has not replaced that evidence head.
There was a problem hiding this comment.
🧹 Nitpick comments (1)
tests/comm/test_pcie_control_node_benchmark_contract.py (1)
19-28: 🎯 Functional Correctness | 🔵 Trivial | ⚡ Quick winBind the channel contract to executable structure.
These assertions search the entire source file. They do not prove that the active benchmark path uses these channel IDs, that
prepare_channelsruns beforefor_streamandcapture, or that the calls use the correct arguments. A comment or dead helper can satisfy the test while runtime wiring is incorrect. Use the AST helper to inspect the target function, or execute it with fakepoolandstreamobjects and assert the recorded call order and channel IDs.🤖 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_control_node_benchmark_contract.py` around lines 19 - 28, Strengthen test_benchmark_prepares_distinct_stable_eager_and_graph_channels by validating the executable benchmark function rather than searching raw source text. Use the existing AST helper to inspect the target function and confirm prepare_channels precedes for_stream and capture with the expected channel IDs, or execute it with fake pool and stream objects and assert the recorded call order and arguments.
🤖 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.
Nitpick comments:
In `@tests/comm/test_pcie_control_node_benchmark_contract.py`:
- Around line 19-28: Strengthen
test_benchmark_prepares_distinct_stable_eager_and_graph_channels by validating
the executable benchmark function rather than searching raw source text. Use the
existing AST helper to inspect the target function and confirm prepare_channels
precedes for_stream and capture with the expected channel IDs, or execute it
with fake pool and stream objects and assert the recorded call order and
arguments.
ℹ️ Review info
⚙️ Run configuration
Configuration used: Organization UI
Review profile: CHILL
Plan: Pro Plus
Run ID: 025e73c9-36ac-472b-9e2f-7280663057d0
📒 Files selected for processing (10)
benchmarks/benchmark_pcie_oneshot_control_node.pysparkinfer/comm/pcie/pcie_twoshot.cutests/comm/test_pcie_control_node_benchmark_contract.pytests/comm/test_pcie_dcp_a2a_gpu.pytests/comm/test_pcie_ipc_lifecycle.pytests/comm/test_pcie_oneshot.pytests/comm/test_pcie_oneshot_fused_rmsnorm_gpu.pytests/comm/test_pcie_oneshot_opposite_order_gpu.pytests/comm/test_pcie_preallocation_setup.pytests/comm/test_pcie_staging_control.py
🚧 Files skipped from review as they are similar to previous changes (9)
- tests/comm/test_pcie_oneshot_fused_rmsnorm_gpu.py
- tests/comm/test_pcie_staging_control.py
- benchmarks/benchmark_pcie_oneshot_control_node.py
- tests/comm/test_pcie_oneshot_opposite_order_gpu.py
- tests/comm/test_pcie_oneshot.py
- tests/comm/test_pcie_dcp_a2a_gpu.py
- sparkinfer/comm/pcie/pcie_twoshot.cu
- tests/comm/test_pcie_preallocation_setup.py
- tests/comm/test_pcie_ipc_lifecycle.py
Review and field closeoutAll ten CodeRabbit threads are now resolved at current review head Post-review validation is 128 passed / 4 expected GPU-gated skips targeted and 243 passed / 18 skipped across the full locally collectable communication suite; changed-file Ruff, explicit RUF043/RUF059, formatting, compilation, and diff checks pass. The provenance boundary is deliberate:
The exact GPU pair passed startup/capture, 13/13 API checks twice, 20/20 cold 65K/126K needles, two error-free bounded C1/C2/C4/C8 matrices, GSM8K 96/100, GPQA 44/50, and complete GPU release within 19 seconds. This is correctness/lifecycle evidence, not an independent memory or throughput claim. Immutable recipe: https://github.com/malaiwah/glm52-exl3-vast/blob/d821001f97dd15f0ae6ace626cc1f334a2069553/docs/field-review-results/2026-07-30-vast-46335896/COUNTER-VALIDATION.md The PR body now carries the exact split between the proposed merge head and the GPU evidence boundary. Ready for human maintainer review and independent counter-validation together with vLLM #216. |
1 participant
Status and exact provenance
This PR supersedes #101 (
0959fe8) and #103 (8fd90a4). Do not stack those rejected public heads with this successor.Summary
Harden SparkInfer PCIe one-shot, two-shot, and DCP A2A for CUDA-graph replay and CUDA-IPC lifetime. Replay staging advances on device execution; graph channels have explicit semantic identity; setup and teardown are collective, ordered, and retryable; descriptor prewarm follows the active semantic scope before CUDA reports physical capture.
Focused numerical graph tests had passed on #101/#103, but strict review found contracts that a short smoke could not discharge: host-baked replay parity, variable-grid rendezvous hazards, partial IPC setup, wrong rank ordering, teardown that discarded retryable ownership, and GC that could unmap pointers still in use. The first successor (
31fa6a4) repaired those issues but failed full GLM startup because descriptor warmup collided with the static eager channel.bc629805routes warmup through the semantic scope owner stream while keeping unrelated streams fail-closed.Final contracts
Validation
Exact GPU-qualified head
bc629805Vast.ai instance 46335896; 4x RTX PRO 6000 Blackwell 96 GB; driver 610.43.02; CUDA runtime 13.2; PyTorch 2.12.0+cu132; fresh extension caches:
Paired vLLM
f99e1e7: 2/2 abnormal-finalizer tests and 16/16 containing distributed tests passed; Ruff, format, byte-compilation, and hooks passed.Post-qualification review head
caa643cNine valid CodeRabbit findings were fixed: CUDA fallback-result ordering, distinct prepared benchmark semantic channels, two-shot launch guards, reduced-SM fixture validation, guaranteed process-group cleanup, test quarantine cleanup, two Ruff regex findings, and an unused binding. The related fused-RMS fallback was hardened too.
One finding was rejected with evidence:
cuda-bindings==13.3.1explicitly supports the one-argumentcudaGraphGetEdges(graph)sizing call, and NVIDIA v13.3.1 uses the same pattern.bc629805GPU evidence remains authoritative.Full GLM-5.2 attribution gate
Exact GPU run:
field-review-final-bc62980-f99e1e7-mtp3-c8-v2; willfalco 3.25 bpw; TP4/DCP4/MTP3; dynamic NVFP4 MLA KV; max length 131,072; sequences 8; batch 2,048; graph/Trellis 32; GMU 0.92; text-only and offload disabled.This closes the 131K text-only/offload-disabled repair-attribution profile. It does not by itself qualify 512K, vision, or live LMCache.
Claim boundary and evidence
This is correctness and lifecycle work. It claims no independent VRAM, PP, or TG gain. The approximately 1.1 GiB/rank full-model recovery belongs to vLLM #210, not this PR.
Review
caa643cas the proposed merge head while usingbc629805+f99e1e7as the exact GPU evidence boundary. After acceptance, close #95 and #97 through this successor.Fixes #95.
Fixes #97.
Supersedes #101 and #103.
Changes, tests, and evidence were prepared with Codex assistance. Human maintainer review and independent counter-validation are requested.