feat: support CCL ReduceScatter

[GordonYang1]

Summary

This PR adds ReduceScatter support to the shared CCL backend abstraction, including native bindings through the existing NCCL and MCCL provider layers for the public infinicclReduceScatter() API. It also keeps inter-only communicators on the existing OpenMPI staging path during mixed-backend bootstrap flows, hardens the shared OpenMPI/MPICH ReduceScatter implementation with typed-reduction validation, overflow checks, and automatic host-buffer cleanup, defines zero-count behavior, makes the existing MPI example validate rank-specific destination blocks and propagate failures, and includes CCL-only plus OpenMPI-assisted ReduceScatter example programs covering out-of-place and canonical in-place flows with correctness and communication-bandwidth reporting.

Changes

• Public API and Dispatch

  • Enable the existing infinicclReduceScatter() API for configured CCL backends through generated bridge dispatch without changing its public signature.
  • Use a matching native CCL intra communicator when available, and otherwise delegate to the existing OpenMPI provider when the communicator has only an OpenMPI inter communicator.
  • Return success for a zero-count ReduceScatter after validating the communicator, data type, and reduction operation, without requiring non-null buffers or entering a backend.

• Common CCL Implementation

  • Add a shared provider-oriented CCL ReduceScatter implementation following the existing AllReduce structure.
  • Validate the native communicator backend, device, and handle before dispatch.
  • Map InfiniCCL data types and reduction operations through the configured provider and return NotSupported when the provider cannot represent either value.

• Existing CCL Provider Bindings

  • Extend the existing NCCL and MCCL API wrappers with thin bindings to ncclReduceScatter() and mcclReduceScatter().
  • Register ReduceScatter with the existing NCCL and MCCL provider layers.

• MPI Correctness and Safety

  • Keep arithmetic reduction on typed MPI_Reduce_scatter_block() calls and reject data types that map to MPI_BYTE, avoiding invalid byte-wise reductions for Float16 and BFloat16.
  • Add world-size, size_t multiplication, buffer-size, and MPI int count-range checks for staged transfers.
  • Manage host staging buffers with automatic cleanup across success and error paths.
  • Validate the rank-specific destination block in the existing MPI example and return a non-zero process status on failure.

• Examples, Validation, and Metrics

  • Add a thread-per-GPU single-node CCL ReduceScatter example using one shared unique ID and rank-based communicator initialization.
  • Add an OpenMPI-assisted CCL ReduceScatter example that first uses the public ReduceScatter API through the OpenMPI inter communicator to distribute the native unique ID, then initializes the native CCL communicator for GPU reduction and scatter.
  • Validate both out-of-place and canonical in-place layouts with destination-block-dependent expected values on every rank.
  • Restore the local input block before every in-place iteration and propagate cluster-wide validation failures through a native ReduceScatter status operation.
  • Report per-rank receive size, total reduced data size, average time, algorithm bandwidth as world_size * receive_bytes / time, and bus bandwidth as algorithm bandwidth multiplied by (world_size - 1) / world_size.
  • Parse numeric options without exceptions and guard example buffer-size calculations against overflow.

Platform and Backend Affected

Platform

• CPU
• NVIDIA GPU
• Iluvatar GPU
• MetaX GPU
• Moore Threads GPU
• Cambricon MLU
• HYGON DCU

Backend

• OpenMPI
• MPICH
• NCCL
• MCCL

Performance Impact

• No performance impact
• Performance improved
• Performance regression possible

This adds GPU-native NCCL and MCCL paths for ReduceScatter, avoiding the existing MPI host-staging path when a supported CCL backend and native communicator are available. The shared MPI path remains host-staged and retains typed reduction semantics with additional validation and cleanup. Other collective operations are intended to remain unchanged. The validation-log timings below are execution references rather than a quantified cross-backend performance comparison.

Known Issues & Future Work

• The CCL collective backend on this branch now covers AllReduce and ReduceScatter; other CCL collective operations remain future work.
• ReduceScatter inherits the backend/device combinations, data types, and reduction operations supported by the existing CCL providers; this PR does not add a new provider or device integration.
• The server examples validate Float32 Sum payloads on the default stream. Additional payload data types, reduction operations, non-default streams, and runtime coverage on Iluvatar and Moore Threads remain future work.
• The shared OpenMPI/MPICH fallback remains host-staged with per-call allocations, rejects Float16 and BFloat16 arithmetic reductions until a typed host conversion path is available, and rejects per-rank counts above the MPI int range; chunked transfers remain future work.

Test Results

The implementation commit is 752bc4e3714ac3e7b4605cd1af0c9dd60df1d5fe, a single commit directly based on ef4045a2d99837c75c2acd90aae57dacb83172d2. The attached evidence contains exactly 15 canonical logs from the current commit: all 15 targets passed, with Correct: YES 20 times and Correct: NO 0 times. The five additional positive results come from the second out-of-place/in-place cases in the three native ReduceScatter logs and the two additional expected validation cases in the broadcast log.

All four ReduceScatter paths passed with a 1,048,576-element Float32 receive block per rank (4.00 MiB), 2 warm-up iterations, and 20 profiled iterations. The native paths validate both out-of-place and canonical in-place layouts; the heterogeneous MPI path validates its out-of-place layout.

Path and layout	Test topology	Receive per rank / total reduced	Time	Alg BW	Bus BW
Pure NCCL, out-of-place	8 NVIDIA A100 GPUs	4.00 / 32.00 MiB	0.240 ms	139.90 GB/s	122.41 GB/s
Pure NCCL, in-place	8 NVIDIA A100 GPUs	4.00 / 32.00 MiB	0.523 ms	64.11 GB/s	56.10 GB/s
OpenMPI + NCCL hybrid, out-of-place	8 NVIDIA A100 GPUs	4.00 / 32.00 MiB	0.250 ms	134.40 GB/s	117.60 GB/s
OpenMPI + NCCL hybrid, in-place	8 NVIDIA A100 GPUs	4.00 / 32.00 MiB	0.345 ms	97.38 GB/s	85.21 GB/s
Heterogeneous OpenMPI, out-of-place	8 NVIDIA A100 + 8 MetaX C550 GPUs	4.00 / 64.00 MiB	142.754 ms	0.47 GB/s	0.44 GB/s
Pure MCCL, out-of-place	8 MetaX C550 GPUs	4.00 / 32.00 MiB	0.219 ms	153.38 GB/s	134.21 GB/s
Pure MCCL, in-place	8 MetaX C550 GPUs	4.00 / 32.00 MiB	1.385 ms	24.22 GB/s	21.19 GB/s

Algorithm bandwidth uses the total logical reduced input size, world_size * receive_bytes, divided by the measured time. Bus bandwidth applies the standard ReduceScatter correction factor (world_size - 1) / world_size. The values were independently recomputed while accounting for the displayed time being rounded to three decimal places. They are included as execution evidence, not as a cross-platform benchmark comparison.

• The pure NCCL and OpenMPI+NCCL configurations each passed both selected targets, 2/2 and 2/2, on all 8 A100 GPUs. Each native ReduceScatter log passed both out-of-place and in-place validation.
• The heterogeneous OpenMPI configuration passed all 9 MPI targets on 16 ranks. Ranks 0-7 mapped to NVIDIA devices 0-7, and ranks 8-15 mapped to MetaX devices 0-7.
• The pure MCCL configuration passed both selected targets on all 8 MetaX C550 GPUs, with ranks 0-7 mapped to devices 0-7. Its ReduceScatter log passed both out-of-place and in-place validation.
• All five formal runner invocations returned zero. There were no harness retries, MetaX vendor queue retries, timeouts, missing canonical logs, or recorded finalization errors.

Test Involved Platform

• CPU
• NVIDIA GPU
• Iluvatar GPU
• MetaX GPU
• Moore Threads GPU
• Cambricon MLU
• HYGON DCU

Test Involved Backend

• OpenMPI
• MPICH
• NCCL
• MCCL

Pure CCL (NCCL) on single-node NVIDIA:
ccl_all_reduce.log
ccl_reduce_scatter.log

CCL + MPI on single-node NVIDIA:
ccl_mpi_hybrid_all_reduce.log
ccl_mpi_hybrid_reduce_scatter.log

MPI on Heterogeneous Cluster:
mpi_all_gather.log
mpi_all_reduce.log
mpi_all_to_all.log
mpi_broadcast.log
mpi_gather.log
mpi_reduce.log
mpi_reduce_scatter.log
mpi_scatter.log
mpi_send_recv.log

Pure CCL (MCCL) on single-node MetaX:
ccl_all_reduce.log
ccl_reduce_scatter.log

---

Checklist

Every contributor must verify every item below before requesting
review. Tick each box only after the check has actually been performed —
do not tick speculatively. If an item truly does not apply, replace the
checkbox with N/A and briefly explain why in an inline comment.

Title, Branch, and Commits

• PR title follows Conventional Commits (e.g. feat: …, fix(nccl): …).
• Branch name follows <type>/xxx-yyyy-zzzz where <type> matches the PR title's Conventional Commits type and words are joined with hyphens (see CONTRIBUTING.md §Branches).
• Each commit message follows Conventional Commits.
• Small PR is a single squashable commit; or, for a large PR, every commit is meaningful, well-formed, and independently reviewable (see CONTRIBUTING.md §Pull Requests).
• No stray merge commits from master — the branch is rebased cleanly on top of the current master.
• No fixup! / squash! / wip commits remain.

Scope and Design

• Changes are minimal — no unrelated modifications were introduced (CONTRIBUTING.md §Code/General).
• No dead code, commented-out blocks, debug prints, printf/std::cout/print(...) left behind, or TODO without an owner and issue link.
• No unrelated formatting churn that would obscure the diff.
• Public API changes (if any) are intentional, documented, and reflected in affected callers/tests.

General Code Hygiene

• The code is self-explanatory; comments were added only where the intent or rationale is non-obvious (CONTRIBUTING.md §Code/General).
• Every modified or added file ends with a single trailing newline (CONTRIBUTING.md §Code/General).
• No trailing whitespace, inconsistent indentation, or mixed formatting styles remain.
• Identifiers referenced in comments or error messages are wrapped in Markdown backticks (e.g. the `AllReduce` implementation) (CONTRIBUTING.md §Code/General).
• All comments and error messages are in English (CONTRIBUTING.md §Code/General).
• Comments and error messages are complete sentences — capitalized first letter, terminal punctuation — unless the language/framework convention says otherwise (CONTRIBUTING.md §Code/General; §Python).

C++ Specific (if C++ files changed)

• Code follows the Google C++ Style Guide strictly.
• clang-format (version 16, per .github/workflows/clang-format.yml) has been run against all modified applicable files; the diff is clean.
• No exceptions are thrown. Error paths use assert with messages that include at least __FILE__, __LINE__, and __func__ (CONTRIBUTING.md §C++).
• Error and warning message wording follows the LLVM Coding Standards (CONTRIBUTING.md §C++).
• N/A- Constructor initializer list order matches member declaration order (CONTRIBUTING.md §C++).
• Exactly one blank line between classes, between classes and functions, and between functions (CONTRIBUTING.md §C++).
• Exactly one blank line between members (functions and variables) within a class (CONTRIBUTING.md §C++).
• Exactly one blank line before and after the contents of a namespace (CONTRIBUTING.md §C++).

Python Specific (if Python files changed)

• N/A- Code is PEP 8 compliant; ruff check passes cleanly on CI (see `.github/workflows/ruff.yml).
• N/A- ruff format --check passes cleanly — if not, run ruff format and commit the result.
• N/A- Comments are complete English sentences, starting with a capital letter and ending with punctuation; Markdown backticks are used for code references (CONTRIBUTING.md §Python).
• N/A- Framework-specific conventions (e.g. lowercase pytest.skip messages without terminal period) are honored where applicable (CONTRIBUTING.md §Python).
• N/A- No blank line between the function signature and the body when there is no docstring or comment (CONTRIBUTING.md §Python).
• N/A- A blank line is present before and after if, for, and similar control-flow statements (CONTRIBUTING.md §Python).
• N/A- A blank line appears before each return, except when it directly follows a control-flow statement (CONTRIBUTING.md §Python).
• N/A- Docstrings (if any) follow PEP 257 (CONTRIBUTING.md §Python).
• N/A- Type hints are added / kept consistent with the surrounding code.

Testing

• All applicable example programs have been built and tested successfully on at least one supported heterogeneous cluster setup.

Build, CI, and Tooling

• N/A- New backends or devices have been added to auto-detection in CMakeLists.txt under if(AUTO_DETECT_DEVICES) or to if(AUTO_DETECT_BACKENDS) if applicable.
• Both CI workflows (clang-format.yml, ruff.yml) are green locally (or expected to be green on CI).

Documentation

• N/A- README.md, CONTRIBUTING.md, or inline docs updated when behavior, build flags, or developer workflow changed.
• N/A- Any user-visible breaking change is called out explicitly under "Summary" and in the commit/PR title with a ! or BREAKING CHANGE: footer.

Security and Safety

• No secrets, access tokens, internal URLs, customer data, or personal hardware identifiers have been committed.
• N/A- Third-party code is license-compatible and attributed.
• No unsafe pointer arithmetic, uninitialized reads, or missing bounds checks were introduced.