Repository navigation
test(rccl): guard cache bits on LL flag-poll loads (AICOMRCCL-2742) - #12836
dhanavikram wants to merge 6 commits into
Conversation
✅ All Policy Checks Passed
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
d65e88b to
ecfe982
Compare
1b74dbb to
ddccf85
Compare
ddccf85 to
af51844
Compare
There was a problem hiding this comment.
ARGUS
Findings
2 more finding(s), not posted inline
projects/rccl/test/ll-cache-bits/tests/test_ll_isa_detector.py:113(Low) No detector test covers the function-boundary break in find_polls (ll_isa.py:82); add one with the compares in the next function.projects/rccl/test/ll-cache-bits/tests/ll_isa.py:48(Nit) CANDIDATE_RE restates PROTO_LL_RE's pattern; build it from PROTO_LL_RE.pattern.
Bottom line: Worth settling before approval: the unpinned pytest install that runs after Configure AWS Credentials (therock-rccl-ci-linux.yml:187), and targets() dropping a second feature variant of one arch (ll_isa.py:146).
Comment @argus review to re-run, or @argus re-review after pushing.
pr-rocm-systems-12836-oci-rccl-pr-buildBuilds
FailuresNo failing tests in this run. 189 tests skipped. pr-rocm-systems-12836-oci-rccl-pr-smokeBuilds
Failures
15 tests skipped. |
pr-rocm-systems-12836-ruby64-rccl-pr-smokeBuilds
FailuresNo failing tests in this run. 15 tests skipped. |
af51844 to
1545b00
Compare
The LL receiver re-reads a FIFO line until its flags change. If the compiler drops the cache bits on that load, the poll can spin on a stale cached copy and hang, with no build error. Add a pytest guard that disassembles a built librccl.so, finds every LL flag-poll load by its shape, and checks it carries the expected bits (nt on gfx942 and gfx950). Other arches are reported as skipped.
Run the full guard in the TheRock RCCL build job on the unsplit librccl.so in dist/rocm, for gfx94X and gfx950. Run the detector unit tests in the host-test guards phase. RCCL_LL_GUARD_REQUIRED=1 turns skips into failures so the CI step cannot pass without running. Also look for the LLVM tools under lib/llvm/bin, and skip cleanly when the library has no device code.
- Run the CI guard after the artifacts push with continue-on-error, and fail a separate gate job instead, so a guard failure no longer skips the nightly test suites. Note in the README that the job is nightly. - Expect scope:SCOPE_SYS on gfx1250 poll loads, and accept v_cmp_eq as well as v_cmp_ne flag compares. - Move the no-poll ProtoLL check into ll_isa and cover it in the detector tests. - Run the suite locally without a venv in the test dir.
…COMRCCL-2742) The flag compares can follow the poll load by ~35 instructions when the inlined abort check sits in between. With a 16-instruction window the detector missed 468 of 2829 polls per arch in a Release gfx942/gfx950 build (all in runTreeSplit), and every poll in a Debug build, where the guard then failed. Scan 48 lines instead.
…-2742) - Pin pytest and run the CI guard before the AWS credentials step; continue-on-error still lets the artifacts push run. - Check every feature variant of an arch (e.g. gfx950:xnack- and xnack+), not just the last one listed. - Point the guard's docstring at the README instead of repeating it. - Test that compares in the next function don't count and that targets() keeps every variant; build CANDIDATE_RE from PROTO_LL_RE.
…2A (AICOMRCCL-2742) With the device linker, .hip_fatbin holds one offload bundle per fat object and clang-offload-bundler reads only the first, so the DDA LL kernels and the symmetric (LLA2A) kernels were never disassembled. - Split .hip_fatbin at each 4 KiB-aligned bundle header and check every code object. - Add the DDA LL and symmetric LL kernels to the candidate functions, and require a poll in each of them as in ProtoLL functions. - Follow flag copies (v_mov) and flag-to-flag compares, and stop when a flag register is overwritten instead of at the next 128-bit load, so the DDA and batched LLA2A polls are found. - Accept nt or sc0 sc1 on gfx942/gfx950: DDA and LLA2A load at system scope. On a Release gfx942+gfx950 build this finds 3512 polls per arch (2829 nt, 683 sc0 sc1), up from 2829, with every LL function covered.
c737118 to
0218a21
Compare
pr-rocm-systems-12836-oci-rccl-pr-buildBuilds
FailuresNo failing tests in this run. 189 tests skipped. |
pr-rocm-systems-12836-ruby64-rccl-pr-buildBuilds
FailuresNo failing tests in this run. 189 tests skipped. |
pr-rocm-systems-12836-oci-rccl-pr-buildBuilds
FailuresNo failing tests in this run. 189 tests skipped. pr-rocm-systems-12836-oci-rccl-pr-smokeBuilds
Failures
15 tests skipped. |
pr-rocm-systems-12836-ruby64-rccl-pr-smokeBuilds
FailuresNo failing tests in this run. 15 tests skipped. |
Motivation
The LL receiver keeps re-reading one FIFO line until its flags change. That load must skip stale cached copies, or the poll never sees the new flag and the collective hangs. The compiler picks the cache bits, so a compiler or code change can drop them silently.
Technical Details
New pytest guard in
test/ll-cache-bits:librccl.soand disassembles itnton gfx942 and gfx950Needs a built RCCL and ROCm LLVM tools. No GPU.
Where it runs in CI:
therock-rccl-ci-linux.yml, nightly and manual dispatch, gfx94X and gfx950): a step after "Push artifacts" runs the full suite onTheRock/build/dist/rocm/lib/librccl.so, which still has its device code (packaging only splits the artifact copy). The step usescontinue-on-error; a separateLL cache-bit gatejob turns a failure red, so the nightly test suites still run.run_host_tests.sh guards): runs only the detector unit tests (-m detector). That job doesn't buildlibrccl.so.Other changes:
RCCL_LL_GUARD_REQUIRED=1(set in CI) turns skips into failures, so the step can't pass without running.$ROCM_PATH/lib/llvm/bin(TheRock dist layout).runTreeSplit, which a 16-line window missed)..hip_fatbin(the device linker adds one per fat object), so the DDA LL and symmetric LLA2A polls are covered too. gfx942/gfx950 acceptnt(LL primitives) orsc0 sc1(DDA, LLA2A).scope:SCOPE_SYS(RCCL_LL_FIFO_SYS_SCOPE). The detector acceptsv_cmp_eqas well asv_cmp_neflag compares.Issue Tracking
JIRA ID: AICOMRCCL-2742
Test Plan
loadLLLinechanged to a plain load, to check the guard catches it.run_host_tests.sh guardsafter rebasing ondevelop.RCCL_LL_GUARD_REQUIRED=1fails when the tools or the lib are missing, or the lib has no device code (a packaged TheRocklibrccl.so).__builtin_amdgcn_global_load_b128(p, "")poll loop, not a full gfx1250librccl.sobuild.Test Result
nt, 683 DDA/LLA2Asc0 sc1)guards: pass (11 detector tests)Submission Checklist