Skip to content

perf(rccl): skip generic-kernel profiler scans and progress-counter syncs when disabled - #13402

Merged
rahulvaidya20 merged 2 commits into
developfrom
users/rahulvaidya20/rccl-generic-kernel-devflags
Nov 5, 2026
Merged

rahulvaidya20 merged 2 commits into
developfrom
users/rahulvaidya20/rccl-generic-kernel-devflags

Conversation

@rahulvaidya20

@rahulvaidya20 rahulvaidya20 commented Oct 13, 2026 •

Copy link
Copy Markdown
Contributor

Motivation

The NCCL 2.32.3 sync added small-message latency to the generic kernel: Gather and Scatter are 0.3–0.35 µs slower per call on one node, and most other collectives 0.1–0.2 µs. The cost is per-batch work done by thread 0 while the rest of the block waits. Two parts of it run even when nothing reads the result: the profiler stamp loops and the progress-counter synchronization. This PR skips them when they are disabled.

Technical Details

  • projects/rccl/src/include/device.h: add RCCL_DEV_FLAG_PROFILER and RCCL_DEV_FLAG_PROGRESS, and a one-byte devFlags field in ncclDevKernelArgs. The field uses existing padding before workMask, so the struct size and layout are unchanged.
  • projects/rccl/src/enqueue/enqueue.cc (finishPlan): set PROFILER when a profiler plugin is loaded and PROGRESS when the progress-counter block is allocated, once per launch.
  • projects/rccl/src/device/common.h:
    • profiler(): skip the per-work-item stamp loops when PROFILER is clear. Without a plugin every item's profilerEnabled bit is already false, so the loops wrote nothing; the signature is unchanged.
    • ncclKernelMain: read PROGRESS once from the kernel arguments and use it for the two progress-counter __syncthreads() guards, which used to test the shared-memory pointer.
  • Behaviour is unchanged. With a profiler loaded, the kernel writes the same stamps; the progress-counter code still checks its pointer, and the flag is set exactly when that pointer is non-null.

Issue Tracking

JIRA ID: AICOMRCCL-2916
JIRA ID: AICOMRCCL-2745

Test Plan

  • rccl-tests, 1K–128M, -n 100 -w 20 -N 5, with the correctness check on. Each comparison runs the pre-sync build (2.31.2), develop and this fix back to back on the same nodes, in rotated order over several rounds; numbers are mean time per call.
  • 1 node: MI355X (gfx950), ROCm 10.1, 1K–64K, 3 rounds.
  • 2 and 4 nodes: MI350X (gfx950, RoCE), ROCm 10.1, 2 rounds, plus 4 extra rounds for 4-node Gather.
  • 2, 4 and 8 nodes on MI300X (gfx942, InfiniBand) and 2 and 4 nodes on MI355X, ROCm 7.0.2, 10.0 and 10.1
  • With a profiler plugin loaded: same profiler event counts and progress-counter totals as develop.

Test Result

  • 1 node, 1K–64K, vs pre-sync: develop is +0.10 µs on average (Gather +0.35, Scatter +0.32, SendRecv +0.14 µs); with this change it is −0.13 µs (Gather −0.12, Scatter −0.14, SendRecv −0.11 µs).
  • 2 and 4 nodes, 1K–64K mean: develop +0.22 / −0.05 µs, this change −0.01 / −0.29 µs. Remaining cross-node AllToAll, AllToAllv and Gather costs on gfx950 (+0.3–0.45 µs) are also in develop and are not from this code; they are tracked separately.
  • MI300X 2/4/8 nodes: median −0.07 / −0.02 / −0.03 µs vs pre-sync.
  • No wrong results in any run.

Submission Checklist

@rahulvaidya20
rahulvaidya20 requested review from a team and a balanced review from Copilot October 13, 2026 20:17
@therock-pr-bot

therock-pr-bot Bot commented Oct 13, 2026 •

Copy link
Copy Markdown

✅ All Policy Checks Passed

Check Status Details
📝 PR Description ✅ Pass —
⛔ Forbidden Files ✅ Pass —
🧪 Unit Test ⚠️ Warning Error: Source/code files changed without an accompanying unit test.
Expected: add at least one test file named like test_<name>.py / test_<name>.cpp (or <name>_test.*).
Current: code file(s) changed: projects/rccl/src/device/common.h, projects/rccl/src/enqueue/enqueue.cc, projects/rccl/src/include/device.h, projects/rccl/test/host/enqueue-test.cc; no test file found
🚫 Draft PR 🔜 To Be Enabled —
🚩 Feature Flag 🔜 To Be Enabled —
📊 Code Coverage 🔜 To Be Enabled —

🎉 All policy checks passed!

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🟢 Approval recommended

The flags preserve existing feature-enabled behavior while safely removing disabled-path overhead.

0 open findings

What changed in this PR

Optimizes RCCL generic kernels by bypassing profiler scans and progress synchronization when those features are disabled.

Changes:

  • Adds kernel flags for profiler and progress-counter availability.
  • Initializes flags per kernel plan.
  • Uses flags to skip unnecessary device-side loops and barriers.
File Description
projects/​rccl/​src/​include/​device.h Defines flags within existing argument-structure padding.
projects/​rccl/​src/​enqueue/​enqueue.cc Initializes flags from profiler and counter state.
projects/​rccl/​src/​device/​common.h Gates profiler scans and progress barriers.

🧠 Review effort: Balanced


💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

@rahulvaidya20
rahulvaidya20 force-pushed the users/rahulvaidya20/rccl-generic-kernel-devflags branch from f79b6db to 74673ad Compare October 13, 2026 20:26

@amd-argus amd-argus Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

ARGUS

Bottom line: Worth settling before approval: devFlags at enqueue.cc:356 has no assertion, and the zero-filled args buffer makes a dropped store indistinguishable from both features being off.

Comment @argus review to re-run, or @argus re-review after pushing.

Comment thread projects/rccl/src/enqueue/enqueue.cc Resolved
Comment thread projects/rccl/src/device/common.h Resolved
Comment thread projects/rccl/src/include/device.h Resolved
@rahulvaidya20

Copy link
Copy Markdown
Contributor Author

Addressed the Argus finding on enqueue.cc:356 (no assertion on devFlags) in d42ec3a.

  • Added 4 EnqueueMicrotest cases in test/host/enqueue-test.cc over g_profilerPluginLoaded x deviceCountersBlock, using the FinishComm fixture that calls finishPlan.
  • Each case checks the exact kernelArgs->devFlags value: 0, RCCL_DEV_FLAG_PROFILER, RCCL_DEV_FLAG_PROGRESS, or both.
  • Verified that deleting the devFlags store in finishPlan fails the three cases that set a bit. The all-off case passes either way, by design, since kernelArgs is zero-filled.
  • rccl-UnitTestsMicroEnqueue: 208/208 pass; the new cases also pass in rccl-UnitTestsMicroEnqueue-devlinker.

No change to src/.

@amd-argus amd-argus Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

ARGUS

Bottom line: Looks approval ready: the devFlags assertion gap is closed by the new cases. The remaining points are optional, the largest being whether the profiler bit should key on plan-&gt;hasProfilerOps rather than the process-global check at enqueue.cc:356.

Comment @argus review to re-run, or @argus re-review after pushing.

Comment thread projects/rccl/src/device/common.h Resolved
Comment thread projects/rccl/src/device/common.h Resolved
@rahulvaidya20
rahulvaidya20 force-pushed the users/rahulvaidya20/rccl-generic-kernel-devflags branch from d42ec3a to baac8c4 Compare October 20, 2026 20:30

@amd-argus amd-argus Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

ARGUS

Bottom line: LGTM.

Comment @argus review to re-run, or @argus re-review after pushing.

@rahulvaidya20
rahulvaidya20 force-pushed the users/rahulvaidya20/rccl-generic-kernel-devflags branch from baac8c4 to 5cc5763 Compare November 5, 2026 00:00
@github-actions

github-actions Bot commented Nov 5, 2026

Copy link
Copy Markdown
Contributor

pr-rocm-systems-13402-oci-rccl-pr-smoke

View this job on CoCo

Builds

Stack Component Version Commit Status Build time
pr-13402_rccl-perf_rocm-10.1-ga gcc-13.2 13.2.0 13.2.0 cached
pr-13402_rccl-perf_rocm-10.1-ga rocm-10.1-ga 10.1.0 10.1.0 cached
pr-13402_rccl-perf_rocm-10.1-ga openmpi-5.0.8 5.0.8 5.0.8 cached
pr-13402_rccl-perf_rocm-10.1-ga rccl-pr-13402 5cc5763 5cc5763 cached
pr-13402_rccl-perf_rocm-10.1-ga rccl-tests-develop-perf develop 8de4115 cached
pr-13402_rccl-perf_rocm-10.2-latest gcc-13.2 13.2.0 13.2.0 cached
pr-13402_rccl-perf_rocm-10.2-latest rocm-10.2-latest 10.2.0rc1 10.2.0rc1 cached
pr-13402_rccl-perf_rocm-10.2-latest openmpi-5.0.8 5.0.8 5.0.8 cached
pr-13402_rccl-perf_rocm-10.2-latest rccl-pr-13402 5cc5763 5cc5763 cached
pr-13402_rccl-perf_rocm-10.2-latest rccl-tests-develop-perf develop 8de4115 cached
pr-13402_rccl-smoke_rocm-10.1-ga gcc-13.2 13.2.0 13.2.0 cached
pr-13402_rccl-smoke_rocm-10.1-ga rocm-10.1-ga 10.1.0 10.1.0 cached
pr-13402_rccl-smoke_rocm-10.1-ga openmpi-5.0.8 5.0.8 5.0.8 cached
pr-13402_rccl-smoke_rocm-10.1-ga rccl-pr-13402-smoke 5cc5763 5cc5763 cached
pr-13402_rccl-smoke_rocm-10.1-ga rccl_gtest-pr-13402 5cc5763 5cc5763 cached
rocm-10.1-ga_rccl-develop-perf gcc-13.2 13.2.0 13.2.0 cached
rocm-10.1-ga_rccl-develop-perf rocm-10.1-ga 10.1.0 10.1.0 cached
rocm-10.1-ga_rccl-develop-perf openmpi-5.0.8 5.0.8 5.0.8 cached
rocm-10.1-ga_rccl-develop-perf rccl-develop-f030609e6984 f030609 f030609 cached
rocm-10.1-ga_rccl-develop-perf rccl-tests-develop-perf develop 8de4115 cached

Failures

No failing tests in this run.

15 tests skipped.

Performance regressions

pr-13402_rccl-perf_rocm-10.1-ga vs rocm-10.1-ga_rccl-develop-perf

Readings: flat 2070.

No confirmed regressions or improvements.

@github-actions

github-actions Bot commented Nov 5, 2026

Copy link
Copy Markdown
Contributor

pr-rocm-systems-13402-ruby64-rccl-pr-smoke

View this job on CoCo

Builds

Stack Component Version Commit Status Build time
pr-13402_rccl-perf_rocm-10.1-ga gcc-11.5 11.5.0 11.5.0 cached
pr-13402_rccl-perf_rocm-10.1-ga rocm-10.1-ga 10.1.0 10.1.0 cached
pr-13402_rccl-perf_rocm-10.1-ga openmpi-4.0.2 4.0.2rc3 4.0.2rc3 cached
pr-13402_rccl-perf_rocm-10.1-ga rccl-pr-13402 5cc5763 5cc5763 cached
pr-13402_rccl-perf_rocm-10.1-ga rcclt-develop-perf develop 8de4115 cached
pr-13402_rccl-perf_rocm-10.2-latest gcc-11.5 11.5.0 11.5.0 cached
pr-13402_rccl-perf_rocm-10.2-latest rocm-10.2-latest 10.2.0rc1 10.2.0rc1 cached
pr-13402_rccl-perf_rocm-10.2-latest openmpi-4.0.2 4.0.2rc3 4.0.2rc3 cached
pr-13402_rccl-perf_rocm-10.2-latest rccl-pr-13402 5cc5763 5cc5763 cached
pr-13402_rccl-perf_rocm-10.2-latest rcclt-develop-perf develop 8de4115 cached
pr-13402_rccl-smoke_rocm-10.1-ga gcc-11.5 11.5.0 11.5.0 cached
pr-13402_rccl-smoke_rocm-10.1-ga rocm-10.1-ga 10.1.0 10.1.0 cached
pr-13402_rccl-smoke_rocm-10.1-ga openmpi-4.0.2 4.0.2rc3 4.0.2rc3 cached
pr-13402_rccl-smoke_rocm-10.1-ga rccl-pr-13402-smoke 5cc5763 5cc5763 cached
pr-13402_rccl-smoke_rocm-10.1-ga rccl_gtest-pr-13402 5cc5763 5cc5763 cached
rocm-10.1-ga_rccl-develop-perf gcc-11.5 11.5.0 11.5.0 cached
rocm-10.1-ga_rccl-develop-perf rocm-10.1-ga 10.1.0 10.1.0 cached
rocm-10.1-ga_rccl-develop-perf openmpi-4.0.2 4.0.2rc3 4.0.2rc3 cached
rocm-10.1-ga_rccl-develop-perf rccl-develop-f030609e6984 f030609 f030609 cached
rocm-10.1-ga_rccl-develop-perf rcclt-develop-perf develop 8de4115 cached

Failures

No failing tests in this run.

15 tests skipped.

Performance regressions

pr-13402_rccl-perf_rocm-10.1-ga vs rocm-10.1-ga_rccl-develop-perf

Readings: flat 2070.

Fails the run: 0 regressed point(s) past 8%, each confirmed by an adjacent size.

No confirmed regressions or improvements.

@rahulvaidya20
rahulvaidya20 merged commit 2ce5ce3 into develop Nov 5, 2026
64 of 67 checks passed
@rahulvaidya20
rahulvaidya20 deleted the users/rahulvaidya20/rccl-generic-kernel-devflags branch November 5, 2026 16:27
@systems-assistant

Copy link
Copy Markdown
Contributor

TheRock Submodule Bump Activity

Newest first

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

4 participants