Repository navigation
[RCCL] [JIRA ID : ROCM-25479] LL128: generate separate kernels for UBR - #9347
Merged
Merged
Conversation
…er kernels Registered user buffers must be accessed with system-scope cache-bypassing loads/stores for cross-GPU coherence, but that path forces the compiler into heavy AGPR usage and spilling. Non-registered buffers only need the cheaper plain/non-temporal path. Selecting between them at runtime kept both code paths (and the AGPR-spilling one) live in every LL128 kernel, hurting occupancy. Make user-buffer registration a compile-time dimension for LL128 so each launch runs a single, clean code path: - device.h / generate.py: add a UserRegMode kernel dimension (1=registered, 2=non-registered) and emit two LL128 variants per collective (ring + tree for AllReduce, ring for AllGather/Broadcast/Reduce/ReduceScatter). Kernel key packing, device/host tables and the total function count are updated accordingly. - enqueue.cc: pick the registered vs non-registered LL128 function id at launch time from work->regUsed / netRegUsed. - primitives.h / prims_ll128.h: thread UserRegMode through Primitives; the registered specialization keeps the system-scope cache-bypass path, the non-registered specialization compiles only the plain/non-temporal path (no dead cache-bypass code, no AGPR spill). - common.h / add_unroll.sh: forward UserRegMode through RunWorkBatch/RunWorkColl and the runRing/runTreeSplit dispatch helpers. The add_unroll.sh regexes use `template\s*<` so they tolerate clang-format's `template <` spacing. - test: add a unit test verifying the LL128 registered/non-registered kernel split. LL (ProtoLL) is intentionally left at the baseline: it does not benefit from the split, so its user-buffer and FIFO accesses are unchanged. Co-authored-by: Cursor <cursoragent@cursor.com>
❌ PR Check — Action Required
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🚫 Please fix the failed policies before requesting reviews. The following policy checks failed:
The |
1 task
pvallem
reviewed
Aug 1, 2026
pvallem
reviewed
Aug 1, 2026
pvallem
reviewed
Aug 1, 2026
The ll128-reg-split change made device-function symbol names conditional:
fn_sym() appends the user-buffer registration field (_1/_2) only for the
LL128 reg-variant collectives and omits it entirely when reg=="0". The
forward declarations, function-pointer table, and DEFINE_ncclDevFunc were
all updated to use fn_sym(), but the pure-RDC binary-search dispatcher
(emitted under !USE_INDIRECT_FUNCTION_CALL && !RCCL_DEVICE_LINKER) still
built the name from paste("_", "ncclDevFunc", *fn), which unpacks all
fields including reg. For every non-reg-variant kernel it emitted calls to
ncclDevFunc_..._0() -- symbols that are never declared or defined -- so the
-fgpu-rdc (--no-device-linker) link failed. The default device-linker build
was unaffected because it dispatches through the function-pointer table.
Make the pure-RDC dispatcher use fn_sym() like the other three sites so the
called symbols match the declared ones. Add a regression test asserting that
every symbol called by a pure-RDC Caller leaf is forward-declared (exercised
with a reg-variant collective so the reg-suffix logic is covered).
Co-authored-by: Cursor <cursoragent@cursor.com>
clang-format the over-length reg-variant device-func selection call to wrap at the column limit. Co-authored-by: Cursor <cursoragent@cursor.com>
Align the LL128 reg-variant test filename with the Tests.cpp naming convention used by the rest of the suite. Update the CMakeLists reference. Co-authored-by: Cursor <cursoragent@cursor.com>
Co-authored-by: Cursor <cursoragent@cursor.com> # Conflicts: # projects/rccl/src/device/generate.py # projects/rccl/src/include/device.h
pvallem
reviewed
Aug 5, 2026
pvallem
reviewed
Aug 5, 2026
pvallem
reviewed
Aug 5, 2026
Contributor
|
please post coherence soak results: force NCCL_PROTO=LL128, small sizes, high rank counts, many iterations, exact-integer datatypes, on gfx942 + gfx950 if possible |
pvallem
reviewed
Aug 5, 2026
pvallem
reviewed
Aug 5, 2026
pvallem
reviewed
Aug 5, 2026
pvallem
reviewed
Aug 5, 2026
Resolve pvallem's review feedback on the LL128 registered/non-registered kernel split (PR ROCm#9347): - enqueue.cc / device.h: factor the registration->UserRegMode mapping into a shared ncclDevFuncLL128RegMode() helper and add an Enqueue/LL128 test that pins it -- a registered (or net-registered) buffer must resolve to the reg=1 kernel and a non-registered one to reg=2, so inverting the 1:2 mapping now fails instead of passing on the "ids are distinct" checks alone. - enqueue.cc: on id<0 in ncclTasksRegAndEnqueue, WARN and return ncclInvalidUsage (mirroring the ncclPrepareTasks path) instead of silently keeping the reg=2 placeholder and losing cache-bypass for a registered buffer. - prims_ll128.h: emit store128Plain as an explicit 128-bit vector store (global_store_dwordx4) instead of two scalar u64 stores, so the data-visible-no-later-than-flag ordering is enforced rather than left to backend coalescing; expand the comment to note it is correctness-load-bearing. - prims_ll128.h: add braces to userBypass()/loadUser128()/storeUser128(). - LL128RegVariantKernelTests.cpp: assert ReduceScatter has neither a reg=1 nor a reg=2 LL128 kernel, confirming the split is scoped. - test_generate_device_table.py: match the reg field specifically (reg=1 and reg=2 pair) so the check can't pass on a reg=0 kernel's unroll suffix. - device.h: wrap the ncclDevFuncId() signature under the 120-column limit. Co-authored-by: Cursor <cursoragent@cursor.com>
Contributor
Author
gfx942 results. gfx950 pending. |
pvallem
approved these changes
Aug 6, 2026
Contributor
Author
Coherence soak — forced NCCL_PROTO=LL128 (PR #9347, ll128-reg-split @ cfd6bb2)Setup
Results — 16-node (16×8 = 128 ranks): PASS
Results — 2-node (16 ranks) cross-node IB pre-check: PASS
Verdict: No coherence regressions — 0 #wrong across all sizes, datatypes, register modes and cycles at both 16 and 128 ranks with LL128 forced on gfx950. Logs: |
alexander-sannikov
added a commit
that referenced
this pull request
Oct 30, 2026
Backported PRs: - #10670 fix(rccl): LL protocol explicit 128-bit vector loads/stores - #10339 fix(rccl): default cheap post-send fence off for gfx950 - #9336 LL: cache-bypass user buffers, plain/non-temporal comm FIFO Partially backported PRs: - #9347 [RCCL] [JIRA ID : ROCM-25479] LL128: generate separate kernels for UBR Aligned with develop: - RCCL_GFX9_CHEAP_FENCE_OFF default Co-authored-by: Cursor <cursoragent@cursor.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Motivation
LL128: generate separate registered / non-registered user-buffer kernels
Technical Details
Registered user buffers must be accessed with system-scope cache-bypassing
loads/stores for cross-GPU coherence, but that path forces the compiler into
heavy AGPR usage and spilling. Non-registered buffers only need the cheaper
plain/non-temporal path. Selecting between them at runtime kept both code paths
(and the AGPR-spilling one) live in every LL128 kernel, hurting occupancy.
Make user-buffer registration a compile-time dimension for LL128 so each launch
runs a single, clean code path:
2=non-registered) and emit two LL128 variants per collective (ring + tree for
AllReduce, ring for AllGather/Broadcast/Reduce/ReduceScatter). Kernel key
packing, device/host tables and the total function count are updated
accordingly.
time from work->regUsed / netRegUsed.
registered specialization keeps the system-scope cache-bypass path, the
non-registered specialization compiles only the plain/non-temporal path (no
dead cache-bypass code, no AGPR spill).
and the runRing/runTreeSplit dispatch helpers.
Result on gfx950: the common non-registered LL128 kernels drop from AGPR=13-32
to AGPR=1 and shed ~30 VGPR (e.g. AllReduce ring f32 197->165, AllGather i8
245->181), improving occupancy, while the registered specialization retains the
required system-scope path in isolation.
LL (ProtoLL) is intentionally left at the f2af4e8 baseline: it does not benefit
from the split and the non-temporal FIFO / runtime-split experiments regressed
it, so its user-buffer and FIFO accesses are unchanged.
Issue Tracking
ROCM-25479
Test Plan
Test Result
Submission Checklist