feat(deepep-efa): TensorRT-LLM NcclEP MoE all-to-all over EFA (NCCL-GIN CPU-proxy) - #1240
feat(deepep-efa): TensorRT-LLM NcclEP MoE all-to-all over EFA (NCCL-GIN CPU-proxy)#1240dmvevents wants to merge 11 commits into
Conversation
…IN CPU-proxy) New test case 3.test_cases/pytorch/tensorrt-llm/nccl-ep-efa/: TRT-LLM's own expert-parallel backend (nccl.ep via nccl4py, selected by TRTLLM_FORCE_COMM_METHOD=NCCL_EP) over AWS EFA through aws-ofi-nccl's GIN CPU-proxy. NGC-from-scratch image (release:1.3.0rc24 — GA v1.2.1 has no NcclEP backend), pinned NCCL 2.30.4 + nccl4py + gdrcopy + aws-ofi-nccl@9c44d34+PR#1351, opt-in default-OFF layer for TensorRT-LLM PR#17715 (algorithm/layout selectability). Recipe: verify-image gate, 16-rank cross-node factory/dispatch probe, single-node EP8 serve (mpirun cannot bootstrap across VPC-CNI pods — documented), concurrency benchmark, 2-node K8s manifest. Mirrors the vllm/deepep-v2-efa + sglang/dsr1-deepep-efa sibling shape; does NOT vendor setup_deepep_efa.sh (no NVSHMEM DeepEP build on this path). Signed-off-by: Anton Alexander <dmvevents@gmail.com>
…e index, de-link unmerged sibling - env_vars.example: default via :- so APPLY_HT_FLAT_PATCH=1 build-push.sh no longer clobbers to 0 - add tensorrt-llm/README.md engine index (Available test cases table, parity with sglang/vllm/nvidia-dynamo) - README.md: de-link the not-yet-merged vLLM DeepEP-V2 sibling (awslabs#1230) to avoid a 404-on-merge - Dockerfile: add org.opencontainers.image.source LABEL (review-optional trivial) Signed-off-by: Anton Alexander <dmvevents@gmail.com>
KeitaW
left a comment
There was a problem hiding this comment.
Review Batch 1/6 — Transport Substrate & Upstream Pins
Everything here is pinned to an immutable SHA with a stated reason, which is the standard this repo asks for and rarely gets. The problem is not pinning discipline — it is what two of the pins point at. The GIN transport is built from an aws-ofi-nccl commit that exists on no branch, tag or PR upstream, plus a cherry-pick from a pull request that was closed unmerged. Neither can receive an upstream fix, neither has an upgrade path, and neither is reproducible from a released artifact. The released tag v1.21.1 supports the CPU-proxy GIN mode this sample uses, so I think both can simply go away.
To be explicit about scope: NCCL_GIN_TYPE=2 (CPU-proxy) is a perfectly reasonable choice for this backend and nothing below asks you to change it. The ask is that the software stack under it be released software.
| # shipped. --no-deps so pip does not drag torch/TRT-LLM's pinned dependency graph backwards; | ||
| # we are deliberately overriding exactly one pin. nccl4py ships the `nccl.ep` python package | ||
| # + libnccl_ep.so (0.1.0 — the version whose HT-kernel ABI details the README documents). | ||
| ARG NVIDIA_NCCL_CU13=2.30.4 |
There was a problem hiding this comment.
NVIDIA_NCCL_CU13=2.30.4 — the one pin that may be genuinely forced; please measure it and say so (should fix)
This is the pin I am least sure should move, and I want to be careful about it rather than lump it in with the two above.
The floor is not what holds it here: nccl_ep_utils.py at v1.3.0rc24 has _MIN_NCCL_RUNTIME_VERSION = "2.30.4" compared with runtime < Version(...), so 2.31.2 satisfies TRT-LLM. nccl4py does not pin it either — its metadata declares a bare nvidia-nccl-cu13 under the cu13 extra with no version bound, and the wheel bundles no libnccl, linking libnccl.so.2 by soname. And nvidia-nccl-cu13==2.31.2 is published.
What might hold it is ABI: libnccl_ep.so 0.1.0's device kernels take ncclDevComm* and ncclWindow_vidmem* by pointer (visible in the exported symbol names), and NCCL's device-API structs did change between 2.30 and 2.31 — ncclDevCommRequirements gained fields, and the ncclGinType_t enum grew GPI = 4 / EFA_GDA = 5. A prebuilt third-party binary compiled against the 2.30 device API running on a 2.31 runtime is exactly the case that can fail quietly rather than loudly.
So the ask is a measurement, not an edit: try 2.31.2, and if it works, take it — if it does not, put that in the pin table ("held at 2.30.4 because libnccl_ep 0.1.0 is built against the 2.30 device API; symptom: …"). Either outcome converts an unexplained old pin into a justified one, which is the part that matters. Note this is also the pin that decides the GIN backend menu — NCCL_GIN_TYPE_EFA_GDA = 5 does not exist in 2.30.4's nccl_device/core.h — but since CPU-proxy is the intended mode, that is context rather than a reason to move.
There was a problem hiding this comment.
Partially addressed — I took the "justify" fork of your ask, not yet the measurement. 558cd52 + e8494dd rewrite the Dockerfile comment and pin-table row to record exactly why it is held: 2.31.2 clears TRT-LLM's _MIN_NCCL_RUNTIME_VERSION floor, but libnccl_ep 0.1.0 is a prebuilt binary against the 2.30 device API, and the 2.30→2.31 device-struct changes you list are exactly the fails-quietly class — so the row frames 2.30.4 as the measured-matching floor, not an upper bound, with the bump condition stated. The live 2.31.2 trial needs a cluster window (a build-only check cannot catch a quiet device-ABI failure); it is planned alongside the EFA 1.50.0 re-measure, and the row updates with whichever result it produces.
There was a problem hiding this comment.
Held at 2.30.4 with the reason now recorded in 558cd52 (the Dockerfile comment): libnccl_ep 0.1.0 is built against the NCCL 2.30 device-side API the GIN CPU-proxy calls into, so moving to 2.31.x is a device-ABI change. I have not measured 2.31.2 on this substrate, so I am not moving the pin blind — it is documented as the measured-matching floor, to be bumped together with a libnccl_ep rebuilt on the newer device API and re-run through verify-image.sh + run-kernel-test.sh.
There was a problem hiding this comment.
Measured on cgk 2× p5en.48xlarge (H200), 16 ranks cross-node. Splitting the answer along the two ABI layers your comment separates, because only one of them is settled:
Library / binding ABI on 2.31.2 — clean, no quiet break (measured). Built the image --build-arg NVIDIA_NCCL_CU13=2.31.2 as a single-variable A/B against a 2.30.4 control. On all 16 ranks is_nccl_ep_installed() passes (_get_nccl_ep_unavailable_reason → None) and CommunicationFactory.create_strategy returns NcclEP — no silent symbol/binding break at the nccl.ep Python layer, and nvidia-nccl-cu13==2.31.2 links libnccl.so.2 by soname exactly as you noted. So the floor and the import/selection path are genuinely fine on 2.31.2.
Device kernel ABI on 2.31.2 — NOT yet certified (this is precisely your ncclDevComm* / ncclGinType_t concern, and I don't want to overstate it). This is the layer that can fail quietly, and I can't yet claim it passes or fails. Two honest obstacles on this specific substrate:
- cgk's compute hosts run gdrdrv 2.4 (userspace libgdrapi 2.5.2), and aws-ofi-nccl v1.21.1's GIN CPU-proxy
forced_pcie_copy()gates onmin(userspace, kernel) >= 2.5→ GIN init hard-fails version-independently on both arms (the sample documents gdrdrv ≥ 2.5 as a host prerequisite; this pair doesn't meet it). - With a throwaway instrument patch (
forced_pcie_copy() -> true, applied identically to both arms so the only differential stays the NCCL runtime version) GIN did initialize and both arms advanced toNcclEPselection — but the run then hit a latent bug in the probe's own diagnostic line (NcclEpContext._ep_algorithm, absent on 1.3.0rc24 — upstream hardcodes LOW_LATENCY and doesn't store the algo on the context) one statement before the firstdispatch(). So the device dispatch/combine round-trip was never exercised on either arm. That probe bug is nowgetattr-guarded (commitee0b062con this branch); certifying the device kernels needs a rebuild + re-run past that fix on a gdrdrv ≥ 2.5 host (or with the instrument patch).
Pin decision, recorded as you asked. Held at 2.30.4 — the version the 2026-08-07 correctness E2E (real trtllm-serve HTTP-200-correct + 16-rank cross-node dispatch/combine, IMA=0, efa-direct on every rank) actually ran on — not as an upper bound but as the measured-matching floor. The Layer-4 Dockerfile comment now states this explicitly: library ABI clean on 2.31.2 (measured), device round-trip not yet measured on either arm, and the exact re-run needed to bump it.
I'll update this thread with the device-level verdict once the rebuilt image runs on a gdrdrv ≥ 2.5 host — leaving it unresolved until there's a PROBE-PASS / PROBE-MISMATCH on the device path, since that's the half your question actually turns on.
There was a problem hiding this comment.
Following up with the header-level answer to the part your comment actually turns on — the device-side ncclDevComm* / ncclWindow_vidmem* ABI. I diffed the device-API headers between v2.30.4-1 and v2.31.2-1 (nccl_device/core.h, nccl_device/impl/impl_comm__types.h, .../impl_core__types.h), which lets me name the specific reason to hold, rather than leaving it as "built against the 2.30 API."
ncclGinType_t — append-only; PROXY=2 is byte-identical (agrees with your read). 2.30.4 has {NONE=0, PROXY=2, GDAKI=3}; 2.31.2 appends {GPI=4, EFA_GDA=5, MAX_TYPES=6} and leaves NCCL_GIN_TYPE_PROXY=2 unchanged. So, as you said, EFA_GDA=5 not existing at 2.30.4 is context, not a reason to move — the CPU-proxy value is stable across the bump.
ncclDevComm — NOT append-only; this is the crux, and it's concrete. The struct grows 23→35 field-lines and, more importantly, the first divergence is mid-struct at field #10:
- 2.30.4:
… resourceWindow_inlined; lsaMultimem; lsaBarrier; railGinBarrier; ginConnectionCount; ginNetDeviceTypes[]; … - 2.31.2:
… resourceWindow_inlined;hybridDenseGinBarrier;lsaMultimem; lsaBarrier; railGinBarrier; ginConnectionCount;backendIndex;ginNetDeviceTypes[]; …
hybridDenseGinBarrier (a ncclGinBarrierHandle_t) is inserted before lsaMultimem, and backendIndex (uint8_t) is inserted inside the GIN block, before ginNetDeviceTypes. There's a third shift on the same struct: ginIsRailed (one bool at 2.30.4) is replaced by ginConnectionStride + ginContextStride (2× int) + ginStrongLegacySignals. And the inlined field #9 itself (resourceWindow_inlined, embedded by value) shrank — 2.30.4 defines ncclResourceWindow_vidmem with explicit reserved padding and the comment "Same size as ncclWindow_vidmem for backward compatibility"; 2.31.2 drops that padding to a bare {lsaFlatBase, stride4G, mcOffset4K}. Any one of these shifts the byte offsets of the GIN fields the CPU-proxy device code reads (railGinBarrier, ginHandles[], ginSignalShadows, ginContextCount, …); together they guarantee it. A prebuilt libnccl_ep.so 0.1.0 compiled against the 2.30 layout, handed a ncclDevComm allocated by a 2.31.2 runtime, reads those fields at the wrong offsets — the fails-quietly case you flagged, now with a named field rather than a hand-wave.
ncclWindow_vidmem — near-compatible. For completeness on the other pointer arg you named: this one is far tamer — ginWins[]→ginWinsDefaultBackend[] is a rename of the same ncclGinWindow_t[NCCL_GIN_MAX_CONNECTIONS] type, and int cftFlatRank is appended at the tail. So ncclWindow_vidmem alone would be layout-compatible; ncclDevComm is the one that isn't.
The honest bound — why this justifies the pin but the empirical run still certifies it. 2.31.2 did add version-negotiation machinery that 2.30.4 lacks entirely: ncclDevCommRequirements gained bool useRuntimeVersion (+ devCommRuntimeVersionSize), alongside the existing magic/version header on ncclDevComm. Its initializer defaults useRuntimeVersion=false — documented as the "device code is not the runtime version" (i.e. AOT/prebuilt) case, which is exactly libnccl_ep's case. So NCCL 2.31 is aware of version skew and might lay out a back-compat devComm for an older-compiled kernel — but whether it does so correctly for a foreign prebuilt AOT binary is runtime-internal, not visible in the headers. So the header diff makes the risk structurally real and specific (it justifies holding the pin, which was your ask — "convert an unexplained old pin into a justified one"), and it also confirms why a green smoke test alone couldn't prove safety here. The device dispatch/combine round-trip on a gdrdrv ≥ 2.5 host — the run I still owe this thread — stays the empirical certifier, and I'll post the PROBE-PASS/PROBE-MISMATCH when that host is available.
Net: held at 2.30.4, now with the concrete reason recorded — ncclDevComm is not append-only across 2.30.4→2.31.2 (hybridDenseGinBarrier inserted at field 10 before lsaMultimem, plus backendIndex mid-GIN-block), shifting the offsets of the GIN fields the CPU-proxy path dereferences. I'll fold that one-liner into the Layer-4 pin comment. Thanks for pushing on this one specifically — you were right that it was the pin that deserved a real answer.
KeitaW
left a comment
There was a problem hiding this comment.
Review Batch 2/6 — Runtime Knobs That Do Nothing
Three of the variables the README presents as the transport contract are not read by any component in the stack. The first is the one I would fix before anything else, because benchmarks/README.md promotes it to a provenance knob that published numbers are supposed to carry.
KeitaW
left a comment
There was a problem hiding this comment.
Review Batch 3/6 — Deployment Manifest & Operational Correctness
The manifest is unusually well commented — the QoS rationale, the anti-affinity reasoning and the scale seam are all things I would otherwise have asked for. Two things in it, though, mean the documented workflow cannot be run as written.
KeitaW
left a comment
There was a problem hiding this comment.
Review Batch 4/6 — Probe & Launcher Robustness
The probe is the strongest artifact in this PR — driving the real CommunicationFactory.create_strategy() entrypoint rather than constructing NcclEP directly is the right design, and I verified the whole call shape against v1.3.0rc24 (details in the positives batch). Three robustness gaps in it and in the launchers.
KeitaW
left a comment
There was a problem hiding this comment.
Review Batch 5/6 — Documentation Consistency
No .dockerignore beside a build context that holds setup/env_vars (nit)
setup/build-push.sh runs docker build ... . from the sample root, and .gitignore does not apply to Docker build contexts — so the gitignored setup/env_vars is uploaded to the daemon and to any remote builder. Nothing bakes it into a layer (the Dockerfile uses targeted COPYs, never COPY . .), and the contents here are a registry URI and a region rather than a credential, so the exposure is small. A two-line .dockerignore covering setup/env_vars and benchmarks/raw/ closes it and keeps the context small.
Size (informational)
15 files / +1,293 lines is at the upper end of what one pass reviews well, and this repo's own PR-size history is consistent that review cost climbs faster than diff size. A natural split here would have been image + verify-image.sh first, then recipe + manifest — worth keeping in mind for the vLLM/SGLang siblings rather than a request to re-cut this one.
KeitaW
left a comment
There was a problem hiding this comment.
Review Batch 6/6 — Evaluation, Positives & Sources
Things That Look Great
I checked the PR's own claims rather than taking them, and they hold up unusually well:
- The base-tag rationale is exactly right. GA
v1.2.1really does lack the backend:tensorrt_llm/_torch/modules/fused_moe/nccl_ep_utils.py404s at that tag, and.../fused_moe/communication/containsdeep_ep.py,nvlink_one_sided.pyand six others but nonccl_ep.py. Atv1.3.0rc24nccl_ep.pyis present.v1.2.1is also genuinely the latest non-rc release. The "GA-over-prerelease exception, documented" framing is warranted. - The probe drives the real entrypoint, and the call shape is correct.
_get_nccl_ep_unavailable_reason(act_dtype, quant_config, num_slots, hidden_size, max_num_tokens, moe_max_num_tokens, top_k)— your seven positional args match the signature exactly.Mapping.dp_sizeistp_size if enable_attention_dp else 1, sotp_size=WORLD, enable_attention_dp=Trueclearscreate_strategy'sdp_size == 1early return, andmoe_tp_size=1clears theAllGatherReduceScatterbranch — trap 2 is real and your Mapping is built to satisfy it._get_context()does hardcodeLayout.RANK_MAJOR, so the LL/RANK_MAJOR claim and the patched-image guard are both grounded. - Trap 1 and trap 3 are real.
TLLM_LOG_LEVELis whatlogger.pyreads; and_MIN_NCCL_RUNTIME_VERSION = "2.30.4"is compared against the loaded runtime, withNcclEP.__init__raisingRuntimeError("nccl-ep is not installed.")when it fails — so the NGC base's older NCCL really is dead-on-arrival as shipped. gdrcopyc91ad9fis thev2.5.2tag object —GET /repos/NVIDIA/gdrcopy/git/ref/tags/v2.5.2resolves to exactly that SHA. Commit-not-tag pinning with a correct equivalence is the standard this repo asks for and rarely gets.- The patch layer's design is right even though its SHA list is currently broken (see the merge-commit finding in batch 1): it defaults OFF so the baseline image has no dependence on an unmerged PR,
git apply --checkfails the build rather than shipping a half-applied tree, andverify-image.shcross-checks the marker file against what is actually in site-packages so a mislabelled image cannot ship quietly. The marker/content consistency check in particular is a genuinely good idea and I would keep it. ignore_eosis a declared field on TRT-LLM's chat-completion request, and the request model isConfigDict(extra="forbid")— so it is genuinely honoured rather than silently dropped, which is the part that is easy to get wrong. Much of the surrounding discipline is better than most benchmark folders in this repo: successes-only percentiles,200-without-usagecounted as a failure rather than a zero-token success, the exit code as sole pass/fail authority with no second grep gate, and the explicit "no alternative-backend baseline was measured" caveat. Two of the stated controls do not hold as implemented (the per-level index reset and the missing warmup, both in batch 4) — the reasoning behind them is sound and it is the scoping that slipped.- The EFA scale seam is correct.
aws ec2 describe-instance-typesgivesp5en.48xlarge16 EFA interfaces / 192 vCPU / 2048 GiB andp5.48xlarge32 — matching the manifest'svpc.amazonaws.com/efa: "16", thecpu: "90"/memory: 1024Giheadroom, and thep5=32comment. NCCL_GIN_TYPE=2is right, and stable.NCCL_GIN_TYPE_PROXY = 2innccl_device/core.hat bothv2.30.4-1andv2.31.2-1.- MIT-0 header on all 15 files — including
.gitignore, the YAML manifest andenv_vars.example, with the shebang correctly first on the executables. This is the most-missed convention in this repo and you got every file. - Nothing floats. Every dependency is a SHA or an exact version,
IMAGE_TAGis immutable with a comment explaining whylatestis wrong underimagePullPolicy: IfNotPresent, no registry is hardcoded, and the "Known limitations" section volunteers that this exact assembly is build-staged rather than re-measured and that no performance numbers are claimed. That last part is the hardest thing to write and the most useful thing in the folder. - The
#1230cross-reference is still accurate — that PR is open, not merged.
Sources
TensorRT-LLM (v1.3.0rc24 unless noted)
tensorrt_llm/_torch/modules/fused_moe/communication/communication_factory.py—create_strategygating,_create_forced_method,_get_nccl_ep_unavailable_reasontensorrt_llm/_torch/modules/fused_moe/communication/nccl_ep.py—NcclEP.__init__,_get_context(Layout.RANK_MAJOR),dispatch/combinesignaturestensorrt_llm/_torch/modules/fused_moe/nccl_ep_utils.py—_MIN_NCCL_RUNTIME_VERSION,_MIN_NCCL_EP_INT32_TOPK_VERSIONtensorrt_llm/mapping.py—dp_sizepropertytensorrt_llm/serve/openai_protocol.py—ignore_eos,ConfigDict(extra="forbid")v1.2.1tree:.../fused_moe/communication/(nonccl_ep.py),nccl_ep_utils.py(404) — verified live, 2026-08-25- NVIDIA/TensorRT-LLM#17715 — open, 4 commits, head
e14a6f64— verified live, 2026-08-25
aws-ofi-nccl
- aws/aws-ofi-nccl#1351 —
state: closed, merged: false, closed 2026-08-21, commitsc2e773d+63698ea— verified live, 2026-08-25 9c44d34...v1.21.1compare →diverged, ahead 145 / behind 5;9c44d34not an ancestor ofmaster, contained by no branch or tag, associated with no PR — verified live, 2026-08-25src/rdma/gin/nccl_ofi_gin_api.cppatv1.21.1exportsncclGinPlugin_v11+ncclGinPlugin_v13(CPU-proxy);ncclGinPlugin_v14comes fromnccl_ofi_gin_gdaki.cpp— verified live, 2026-08-25v1.21.1vendors the GIN plugin headers at3rd-party/nccl/cuda/include/nccl/{gin_v13.h,gin_v14.h}— verified live, 2026-08-25AC_ARG_WITHset inconfigure.ac+m4/at9c44d34andv1.21.1—cuda, gdrcopy, hwloc, libfabric, lttng, mpi, nvtx, rocm, valgrind(+nvccat v1.21.1); nonccl-headers, nonccl.--enable-cudart-dynamic/--enable-platform-awsboth present at both refs — verified live, 2026-08-25src/nccl_ofi_gdrcopy.cppatv1.21.1vs9c44d34— both attemptgdr_pin_buffer_v2withGDR_PIN_FLAG_FORCE_PCIEand fall back toflags = 0; the only delta is9c44d34's extra EINVAL/ENOTTY → v1 branch — verified live, 2026-08-25include/nccl_ofi_param.hat9c44d34vsv1.21.1—GIN_GDAKI(bool, stub) vsGDAKI_EFA_HW_COUNTER(auto/on/off); noGIN_MAX_REQUESTSat either — verified live, 2026-08-25src/nccl_ofi_gdrcopy.cpp:132at9c44d34—FALLBACK_V1_FOR_GDRDRV_24, absent atv1.21.1/master— verified live, 2026-08-25
NCCL
src/include/nccl_device/core.hatv2.30.4-1—NCCL_GIN_TYPE_PROXY = 2,GDAKI = 3; atv2.31.2-1— addsGPI = 4,EFA_GDA = 5— verified live, 2026-08-25src/gin/gin_host.ccatv2.30.4-1—NCCL_GIN_ENABLE,NCCL_GIN_NCONNECTIONS; noGIN_MAX_REQUESTS— verified live, 2026-08-25
Packaging
- PyPI
nccl4py0.3.1 wheel →nccl/ep/lib/libnccl_ep.so,NCCL EP version 0.1.0+cuda13.3; 0.4.1 wheel → nonccl/ep— verified live, 2026-08-25 readelf -d libnccl_ep.so→NEEDED libnccl.so.2by soname, no bundled libnccl; nccl4py 0.3.1METADATAdeclaresnvidia-nccl-cu13; extra == "cu13"with no version bound — verified live, 2026-08-25- PyPI
nvidia-nccl-cu13— 2.30.4 and 2.31.2 both published — verified live, 2026-08-25 https://efa-installer.amazonaws.com/aws-efa-installer-{1.48.0,1.49.0,1.50.0}.tar.gz→ 200; 1.51.0 → 403 — verified live, 2026-08-25NVIDIA/gdrcopytagv2.5.2→c91ad9f178e5fb729fc5b6dc62a77c3bb364d6c9— verified live, 2026-08-25
AWS
aws ec2 describe-instance-types --instance-types p5en.48xlarge p5.48xlarge→ 16 / 32 EFA interfaces — verified live, 2026-08-25
Repo precedent
micro-benchmarks/expert-parallelism/deepep-v2-benchmark/{README.md,deepep.Dockerfile,setup_deepep_gin.sh}(merged as #1234 = this PR's base commit9baf5aa7)micro-benchmarks/expert-parallelism/deepep-benchmark/deepep.Dockerfile:1593.test_cases/pytorch/sglang/dsr1-deepep-efa/recipe/{serve.sh:190,verify-image.sh:26}3.test_cases/megatron/megatron-bridge/run-ab-rawpods.sh:90-99
…EFA 1.50.0; justify NCCL 2.30.4 pin - aws-ofi-nccl: retired dev commit 9c44d34 + closed PR#1351 -> released tag v1.21.1, whose gdr_pin_buffer_v2 makes forced-PCIe-with-fallback the released default (no cherry-pick, no OFI_NCCL_GDRCOPY_FORCED_PCIE_COPY override needed). - EFA installer 1.48.0 -> 1.50.0 to match the sibling deepep-v2-benchmark build; labeled a build-pin (the 2026-08-07 correctness E2E ran on 1.48.0, not yet cluster-re-measured). - NCCL 2.30.4: record the ABI reason it is held rather than bumped to 2.31.2 (libnccl_ep 0.1.0 is built against the NCCL 2.30 device API the GIN proxy calls; 2.31.x is an unmeasured device-ABI change) — measured-matching floor, not an upper bound. - nccl4py 0.3.1 pin comment: it ships nccl.ep + libnccl_ep 0.1.0; 0.4.1 drops nccl.ep. - HT_FLAT opt-in patch: pin PR#17715's three single-parent commits; ENV NCCL_GIN_PLUGIN / NVIDIA_GDRCOPY=enabled baked so the manifest and image agree. Signed-off-by: Anton Alexander <dmvevents@gmail.com>
…eral-storage reservation, non-privileged gdrdrv variant - RUN_SERVE env gates whether ordinal 0 becomes the serve; the startup/readiness probes key off it so a RUN_SERVE=0 idle ordinal-0 does not restart-loop against a :8000 that never comes up (probe peers gate on /tmp/probe-peer.ready instead). - Cross-node probe rendezvous uses the stable headless-Service DNS name (trtllm-nccl-ep-0.trtllm-nccl-ep.trtllm-nccl-ep.svc.cluster.local) that publishNotReadyAddresses exists to serve — survives pod restart, no raw pod IP. - ephemeral-storage 200Gi added to limits AND requests so the scheduler reserves room for the work emptyDir (sizeLimit bounds but does not reserve; eviction risk mid-run). - Documented the non-privileged variant: commented /dev/gdrdrv hostPath mount + CharDevice volume, and kept IPC_LOCK (redundant under privileged, load-bearing when unprivileged). - No livenessProbe by deliberate choice (a wedged serve stays inspectable) — documented. Signed-off-by: Anton Alexander <dmvevents@gmail.com>
…e-safe probe, SIGPIPE-safe verify, benchmark warmup+monotonic index - verify-image.sh: drain libnccl lookup with awk NR==1 instead of head -1 (head closes the pipe -> grep takes SIGPIPE-141 under pipefail, worst exactly in the 2+-entry case the check exists to catch). - probe_nccl_ep.py: all_reduce(MIN) on the per-rank selectability verdict so ranks fail together; assert E % WORLD before the collective; CPU-snapshot oracle; module-scope RANK/WORLD/LOCAL defaults; Split docstring. - serve.sh / run-kernel-test.sh: fail-loud NCCL_LIB resolution; ENABLE_CONFIGURABLE_MOE set in run-kernel-test.sh so the correctness gate exercises the served config; SERVE_REVISION. - benchmark_probe.py: discarded warmup before level 0 (lazy CUDA/tokenizer init no longer inflates the first level); run-global monotonic prompt index across levels (per-level range(n) let a KV cache serve later levels' prefills free); short-generation check (completion_tokens != max_tokens is a failure, not a deflated-denominator success); reject --requests-per-level-mult < 1 (0 false-passes the exit-code gate); BENCH_CONCURRENCIES env-read so a higher sweep needs no image rebuild. Signed-off-by: Anton Alexander <dmvevents@gmail.com>
…svc DNS commands, transport-vs-selection env split, second int32 exit
- Pins table + diagram + References: aws-ofi-nccl released tag v1.21.1 (dev commit 9c44d34 /
closed PR#1351 dropped), EFA installer 1.50.0 (build-pin note), nccl4py 0.3.1 REQUIRED
(libnccl_ep is 0.1.0). benchmarks/README: drop the retired OFI_NCCL_GIN_MAX_REQUESTS row.
- Prerequisites: add the host gdrdrv kernel-module line (lsmod | grep gdrdrv) — privileged
lets the container open /dev/gdrdrv but cannot conjure the node if the module never loaded.
- Documented commands use the stable .svc DNS name / kubectl exec loopback, no raw pod IP.
- Runtime-requirements: split into the transport contract (manifest env + every launcher)
vs the launcher-owned NcclEP-selection knobs (deliberately NOT in the manifest env — a
TRTLLM_FORCE_COMM_METHOD on the idle probe-peer pod would mislead).
- Trap 4: document the second, package-side exit (_MIN_NCCL_EP_INT32_TOPK_VERSION=0.2 — a
libnccl_ep>=0.2 retires the int64 boundary with no patch; unreachable today, a PyPI watch.
- Build block: source setup/env_vars + reference ${REGISTRY}/${IMAGE_NAME}:${IMAGE_TAG}.
Signed-off-by: Anton Alexander <dmvevents@gmail.com>
|
Thanks for the review — all 31 threads are addressed across four commits, with a per-thread reply above mapping each finding to its fix:
Two honest partials, flagged in their threads: the NCCL 2.31.2 trial and the EFA 1.50.0 re-measure both need a cluster window — the pin comments carry those caveats until the runs happen. |
…erence (reorg awslabs#1119) Upstream awslabs#1119 (merged) reorganized the repo away from the numbered `3.test_cases/pytorch/<framework>/` tree to a framework-centric `examples/inference/<engine>/<name>/` layout. This PR was opened against the now-deleted path, so its files were orphaned there after merging upstream/main and the PR went un-mergeable. Move the TensorRT-LLM NcclEP-over-EFA sample and its framework-level README to their reorganized home: 3.test_cases/pytorch/tensorrt-llm/ -> examples/inference/tensorrt-llm/ - Pure `git mv` — all 15 files are 100% renames, zero content change. - `tensorrt-llm/` is a net-new engine directory. Per AGENTS.md "extend before create", the closest sibling is `examples/inference/sglang/dsr1-deepep-efa` (also a DeepEP-family MoE-over-EFA serving case), but TRT-LLM's execution model differs materially: it drives wide-EP all-to-all through TensorRT-LLM's `NcclEP` backend (`nccl.ep` / `libnccl_ep` via NCCL-GIN CPU-proxy), not the `deep_ep` package the sglang/vllm samples use, and the framework index already lists TRT-LLM as planned. A parallel engine dir is the house-correct placement, not an extension of sglang. - Add the `tensorrt-llm` row to the `examples/inference/README.md` engine index. - Relative links preserved: old and new paths are the same depth, and sibling frameworks moved to parallel homes, so `../../sglang/dsr1-deepep-efa` and `../../../../micro-benchmarks/expert-parallelism` still resolve. Merged (not rebased) upstream/main to preserve the exact commit SHAs cited in the resolved review threads. Signed-off-by: Anton Alexander <dmvevents@gmail.com>
…2 device-ABI trial result Three fixes surfaced running the KeitaW awslabs#1240 device-ABI trial on cgk p5en: 1. Dockerfile: drop `--disable-build-ngc` from the EFA installer invocation. That flag existed in aws-efa-installer 1.48.0 but was REMOVED in 1.49.0+; under the pinned 1.50.0 installer getopt now rejects the unknown long-opt and efa_installer.sh prints usage + exits 1, failing the layer. `--disable-ngc` alone is the correct flag (matches the sibling dsv3-uccl-nixl sample). 2. Dockerfile: stop prepending the EFA installer's bundled OpenMPI (/opt/amazon/openmpi) to PATH/LD. Its OpenMPI 4.1.7 libopen-pal.so.40 does not export opal_libevent2022_event_assign; on the NGC TRT-LLM base it shadows the HPC-X libopen-pal that tensorrt_llm's MPI_Init resolves through, giving an undefined-symbol abort on `import tensorrt_llm`. Neither recipe uses EFA's mpirun (serve.sh is single-node; the probe uses torchrun), so EFA's OpenMPI is unnecessary and actively harmful on LD. 3. run-kernel-test.sh: the line-17 `${1:?usage: ... {leader|worker} ...}` message contained a literal `}` which closes the parameter expansion early, so ROLE captured the junk tail and the case below never matched -> FATAL on every invocation. Split the role check into a brace-free `${ROLE:?...}` + case. Also sharpens the NCCL pin-table comment to record the measured 2.31.2 trial: the nccl.ep LIBRARY ABI loads clean on 2.31.2 (is_nccl_ep_installed passes, factory selects NcclEP on all 16 ranks), but the DEVICE dispatch/combine round-trip was unmeasurable on cgk p5en for a NCCL-version-independent reason — the GIN CPU-proxy hard-requires GDRCopy >= 2.5 (aws-ofi-nccl v1.21.1 forced_pcie_copy -> min(runtime,driver) >= 2.5) and the host gdrdrv is 2.4, so GIN init hard-fails on BOTH 2.30.4 and 2.31.2. Pin stays 2.30.4 (the version the 2026-08-07 correctness E2E passed on); a 2.31.2 device-path cert is owed on a host with gdrdrv >= 2.5. Signed-off-by: Anton Alexander <dmvevents@gmail.com>
…bsent on 1.3.0rc24) recipe/probe_nccl_ep.py read ctx._ep_algorithm.name in three diagnostic spots. On TRT-LLM 1.3.0rc24 NcclEpContext exposes only `layout` — no _ep_algorithm (upstream hardcodes LOW_LATENCY, does not store the algorithm on the context). This was latent: on a stock image the aws-ofi-nccl GIN CPU-proxy init fails the GDRCopy>=2.5 gate and create_group() raises before the probe reaches the diagnostic line, so it was never hit. On a substrate where GIN does init (gdrdrv>=2.5), all ranks pass selection then AttributeError one line before the first strategy.dispatch(), so the device round-trip is never exercised. Fix: getattr(ctx,'_ep_algorithm',None), fall back to the env-requested algo / the observable ctx.layout for the HIGH_THROUGHPUT check. Prefer the real attr if a future base exposes it. Also refresh the Layer-4 pin rationale comment: the 2.31.2 device path is not "unmeasurable on this substrate" — it is reachable (measured this trial), the remaining step is a rebuild + re-run past this fix. Signed-off-by: Anton Alexander <dmvevents@gmail.com>
… 2.30.4 pin The Layer-4 pin comment said only that libnccl_ep 0.1.0 is 'built against the NCCL 2.30 device-side API' — true but a hand-wave, and exactly the gap a reviewer flagged. Replace it with the named-field reason from a raw-header diff of nccl_device/impl/impl_comm__types.h at tags v2.30.4-1 vs v2.31.2-1: - ncclDevComm (the device struct libnccl_ep dereferences by pointer) is NOT append-only across the bump. 2.31.2 inserts hybridDenseGinBarrier at field 10 (before lsaMultimem) and backendIndex mid-GIN-block, and shrinks the by-value resourceWindow_inlined member (drops its reserved padding) — each shifts the byte offsets of the GIN fields the CPU-proxy kernels read. - ncclGinType_t IS append-only (PROXY=2 unchanged, +EFA_GDA=5), so the enum is not the issue — the devComm layout is. This is the silent-corruption class, so a build-only check cannot catch it, which is why the pin is held rather than bumped on a green library-ABI smoke. Comment-only; no build behavior changes. Signed-off-by: Anton Alexander <dmvevents@gmail.com>
What this adds
A new self-contained test case,
3.test_cases/pytorch/tensorrt-llm/nccl-ep-efa/, runningTensorRT-LLM's own MoE expert-parallel backend (
NcclEP, thenccl.ep/nccl4py path — not thedeep_epPython package) over AWS EFA via aws-ofi-nccl's GIN CPU-proxy. It fills theTensorRT-LLM gap in the expert-parallelism example matrix alongside the merged
sglang/dsr1-deepep-efa(NVSHMEM host-proxy) and the vLLMdeepep-v2-efa(#1230, same NCCL-GINsubstrate, different EP kernel package) test cases, and mirrors their folder shape:
nvcr.io/nvidia/tensorrt-llm/release:1.3.0rc24), publicsources only, every pin an immutable SHA with a WHY comment, fail-loud draining asserts,
in-tree COPY'd setup script (aws-ofi-nccl GIN @
9c44d34+ PR#1351 immutable head).An opt-in, default-OFF layer (
APPLY_HT_FLAT_PATCH=1) bakes [#17714][feat] NcclEP: make algorithm/layout selectable so HIGH_THROUGHPUT+FLAT works on EFA NVIDIA/TensorRT-LLM#17715(algorithm/layout selectability); the baseline image has zero dependence on the unmerged PR.
verify-image.sh(pre-model gate: efa-direct, single-NCCL-wins, GIN symbol,factory imports, patch-marker consistency),
run-kernel-test.sh+probe_nccl_ep.py(16-rankcross-node probe of the REAL
CommunicationFactory.create_strategy()entrypoint +dispatch/combine numerics vs an oracle),
serve.sh(single-node EP8trtllm-serve),benchmark.sh+ stdlib-only probe (successes-only percentiles, unique prompt prefixes,ignore_eos, exit-code authority).publishNotReadyAddresses, Guaranteed QoS with the CPU-proxy rationale, privileged/gdrdrvjustification, verbatim transport-env contract).
TLLM_LOG_LEVEL-not-TRTLLM_LOG_LEVELsilent-ignore, the attention-DP factory gate, thebaked-NCCL-too-old trap, the int64/int32 harness-green≠serve-green boundary, the
GA-lacks-the-backend trap, and more), honest Known limitations.
Why the pins look the way they do
1.3.0rc24, not GA: GA v1.2.1 does not contain the NcclEP backend at all(
nccl_ep_utils.pyabsent at the tag) — the rc pin is required, not a preference, and theREADME documents re-pinning when a GA carrying the backend ships.
setup_deepep_efa.shis NOT vendored — this path builds no NVSHMEM DeepEP, so the folderis intentionally outside the
deepep-vendor-sync.ymlgate (its own setup script says so).What is measured vs. what is not (honest status)
Measured (2026-08-07, 2× p5en.48xlarge/H200): with the same component versions
(NCCL 2.30.4 + nccl_ep 0.1.0 + EFA 1.48.0 + gdrcopy 2.5.2 + aws-ofi-nccl@9c44d34+PR#1351) and the
rc24 NcclEP modules grafted onto an rc9 container: a real
trtllm-serve(Qwen3-30B-A3B, EP8)reached startup-complete with
NCCL EP group createdon all 8 ranks and answered/v1/chat/completionscorrectly (incl. an arithmetic-correctness check), and the 16-rankcross-node factory/dispatch/combine probe passed 16/16 with zero illegal-memory-access and the
efa-directbanner on every rank — under BOTH the upstream LOW_LATENCY/RANK_MAJOR default AND(patched image) HIGH_THROUGHPUT/FLAT.
Not measured, stated plainly: this exact image assembly (rc24-native, no graft) is
build-staged but not yet cluster-re-run — the recipe's gates exist to re-verify it; no
performance numbers are published (benchmarks/README is methodology + provenance only, with the
no-alternative-baseline caveat); the served completion is single-node EP8 (a node-spanning
trtllm-serveneeds mpirun, whose OOB cannot bootstrap across VPC-CNI /32 pods — documented inthe README and manifest, with the torchrun probe as the cross-node proof).
References
dispatch-return fix (the opt-in patch layer; retired when it merges)
OFI_NCCL_GDRCOPY_FORCED_PCIE_COPY(pinned at immutable head)3.test_cases/pytorch/sglang/dsr1-deepep-efa/(merged),3.test_cases/pytorch/vllm/deepep-v2-efa/(feat(deepep-efa): vLLM DeepEP-V2 MoE all-to-all over EFA (eager + non-eager) #1230)