[dlamd1dai]

Supersedes dlamd1dai/rocm-systems#1 (fork PR retargeted to upstream develop).

Summary

Adds device-side (in-kernel wall_clock64) timing for the GIN-SDMA AllToAll collective, measuring pure GPU device-function latency (excludes host launch / graph capture / teardown overhead).

• gin_sdma_devtime.h — portable gin_devtime::measure() scaffold: queries the GPU wall-clock rate, allocates per-CTA start/end stamps, launches a persistent skip+loop timed kernel, reduces min(start)..max(end) over CTAs, takes the slowest local GPU, then MPI_MAX across ranks. Fails if wall-clock rate query fails (no silent fallback).
• gin_sdma_devtime_host.h — pure host helpers shared by BenchTime, gin_sdma_devtime.h, and the new GTest target: elapsed normalization, cuda-graph scaling, zero/negative deltaSec skip guard, iter-stats gating, and stamp-reduction math (reduceMinStart / reduceMaxEnd / gridBusyWindowPerIterUs).
• alltoall.cu — factors the GIN/Hybrid collective bodies into reusable __device__ functions, adds persistent GinAlltoAllTimedKernel/HybridAlltoAllTimedKernel loop variants, and implements the AlltoAllDeviceTime hook (CLI-controlled loop/skip/tier overrides; selects the timed kernel via deviceImpl).
• common.h / common.cu — adds the optional testColl::deviceTime hook and exposes deviceImpl (non-static) so per-collective hooks can pick the matching timed kernel. Integrates device-time modes into BenchTime:
  • --device_timing 1 (-B 1) — augment: normal run plus an extra #[a2a-devtime] line.
  • --device_timing 2 (-B 2) — device-time-only: reports in-kernel device latency as the metric (skips the host graph loop for the out-of-place pass).
  • Host-timed path skips rows when normalized deltaSec <= 0 (avoids getBw divide-by-zero / SIGFPE on fast collectives or coarse timers).
• CMakeLists.txt — hipify/copy gin_sdma_devtime.h with the collective tests; enable CTest and add test/ host unit-test subdirectory.
• test/test_AllToAllDevtime.py — opt-in GPU smoke pytest (GIN -D 3, -g 1): exercises -B 1 (#[a2a-devtime]), -B 2, and -B 2 -H 1; wired into run-device-api-ci.sh after the device-api benchmark matrix.
• test/gin_sdma_devtime_host_test.cpp — CPU-only GTest (rccl-tests-GinSdmaDevtimeHostUnit, 6 cases) covering BenchTime zero-time guards and devtime stamp-reduction math; no GPU/MPI required.

CLI flags (replaces env vars; per @abouteiller review)

Flag	Purpose	Default
-B / --device_timing	0=off, 1=augment, 2=device-only	0
-L / --devtime_loop	Timed-kernel loop iterations	10
-P / --devtime_skip	Timed-kernel warmup skip iterations	10
-W / --devtime_loop_mid	Loop at per-peer ≥ 8 MiB (0=use -L)	0
-Q / --devtime_loop_large	Loop at per-peer ≥ 64 MiB (0=use tier below)	0
-j / --devtime_skip_mid	Skip at mid tier (-1 → min(-P, 2))	-1
-k / --devtime_skip_large	Skip at large tier (-1 → min(-P, 1))	-1
-H / --devtime_check	Validate timed-kernel output before datacheck	0

Example (mode 1 augment, MI355 sweep):

alltoall_perf -D 3 -V 8 -g 1 -R 2 -B 1 -L 20 -P 5 -b 128 -e 4G -f 2 -w 5 -n 20

JIRA ID: AICOMRCCL-1458

AICOMRCCL-1459, AICOMRCCL-1460

Stacking

Base branch: develop. Depends on #10420 (128 MiB Anvil-SDMA put cap), merged 2026-08-24.

Rebased onto upstream develop @ 60ca803 (2026-08-24). Devtime tip: 9c17f5f. Zero diff in projects/rccl/; changes in rccl-tests only. Eight commits preserved (not squashed).

Build gate

The timed path (AlltoAllDeviceTime + the *TimedKernel templates) is compiled only under ENABLE_DEVICE_API and NCCL_VERSION_CODE >= NCCL_VERSION(2,28,7); otherwise a #else stub is compiled instead. ENABLE_DEVICE_API defaults OFF in install.sh/Makefile, so a default build does not compile this code. Build with --enable-device-api (cmake -DENABLE_DEVICE_API=ON) against RCCL >= 2.28.7 to exercise it; CI covering this collective should do the same.

Compile-verified with -DENABLE_DEVICE_API=ON:

• gfx942 (MI300X): alltoall_perf builds & links; AlltoAllDeviceTime + GinAlltoAllTimedKernel/HybridAlltoAllTimedKernel present; Allreduce<T> cross-thread reduction resolves at link time.
• gfx950 (MI355X, ROCm 7.13 docker): BUILD SANITY OK — GIN host plugins in librccl.so, Anvil SDMA factory linked, AlltoAllDeviceTime present in alltoall_perf. GIN Anvil-SDMA runtime smoke OK. Use the default device linker (do not pass --no-device-linker to RCCL install.sh on ROCm 7.13; otherwise cooperative_groups::this_cluster() duplicate symbols fail the link).

Test plan

• Compile-verified on MI355X / gfx950 (docker, ROCm 7.13): alltoall_perf builds & links (AlltoAllDeviceTime + timed kernels present), GIN Anvil-SDMA smoke OK.
• Mode 1 (-B 1 -L 20 -P 5, augment) on MI355X: normal numbers plus #[a2a-devtime] line per size (128 B–4 GB); datacheck Out of bounds values : 0 OK.
• Mode 2 (-B 2 -L 20 -P 5, device-time-only) on MI355X: reported metric is in-kernel device latency; #wrong=0 across 128 B–4 GB.
• Mode 2 + -H 1 on MI355X: timed-kernel output validated against expected — 0 wrong elements across all sizes through 4 GB; datacheck clean.
• Mode 1/2 + -H 1 validated earlier on MI300X (see below; pre-rebase stack).
• Host unit test rccl-tests-GinSdmaDevtimeHostUnit (MI355X): 6/6 GTest cases pass — BenchTime zero-time skip, elapsed normalization, cuda-graph scaling, iter-stats gating, and gin_sdma_devtime.h stamp-reduction math.

Build/run host unit test (no GPU/MPI):

cd projects/rccl-tests
mkdir build-host-unit && cd build-host-unit
cmake .. -DCMAKE_BUILD_TYPE=Release -DBUILD_RCCL_TESTS_HOST_UNIT_TESTS=ON
cmake --build . --target rccl-tests-GinSdmaDevtimeHostUnit -j$(nproc)
./test/rccl-tests-GinSdmaDevtimeHostUnit

Review updates

@pvallem (addressed on fork PR #1; review threads linked there)

• measure() split into launch-all then sync-all passes to avoid the -g N>1 device-barrier deadlock.
• Cross-rank/thread reduction routed through the harness Allreduce() helper (MPI serialized to the last thread; also combines across threads) instead of a direct MPI_Allreduce on MPI_COMM_WORLD.
• AlltoAllDeviceTime gated to deviceImpl 3/4 (GIN/Hybrid) only.
• skip clamped to >= 0 so the start stamp is always written.
• Opt-in --devtime_check validates the timed kernel's output against expected before the datacheck loop re-inits.
• Per-GPU wall-clock rate query; fail on bad rate (no khz=1 fallback).

@amd-jiali

• Added GPU smoke pytest test/test_AllToAllDevtime.py (GIN -D 3, modes -B 1/-B 2/-B 2 -H 1) and hooked it into device-api CI via run-device-api-ci.sh.

@abouteiller

• Device timing controlled via CLI flags (generic names, not ANVIL-specific env vars).
• Removed legacy NCCL_GIN_ANVIL_A2A_DEVICE_TIMING and all other devtime env vars.
• Replaced AUTOREDUCE with explicit --devtime_loop_mid / --devtime_loop_large tier knobs.

Runtime validation (MI355X / gfx950, GIN Anvil-SDMA, CLI flags)

8-GPU MI355X node (smci355-ccs-aus-m03-17), docker image rccl-gin-devtime-pr9927:latest (ROCm 7.13, -DENABLE_DEVICE_API=ON).

Config: NCCL_GIN_TYPE=6, -D 3 -V 8 -g 1 -R 2 -w 5 -n 20, sweep 128 B–4 GB (-b 128 -e 4G -f 2).

Gate	CLI	Result
Docker build + smoke	—	PASS — GIN smoke OK
Test#5 mode 1 (augment)	-B 1 -L 20 -P 5	PASS — #wrong=0, 26 #[a2a-devtime] lines; avg busbw 142.9 GB/s (host)
Test#5 mode 2 (device-only)	-B 2 -L 20 -P 5	PASS — #wrong=0; avg busbw 131.8 GB/s (device-time metric)
Test#5 mode 2 + datacheck	-B 2 -L 20 -P 5 -H 1	PASS — timed-kernel check 0 wrong; final datacheck #wrong=0
4 GB peak (mode 1)	—	host 419.7 GB/s busbw; #[a2a-devtime] 349.9 GB/s busbw @ 10.74 ms devtime

Runtime validation (MI300X / gfx942, earlier stack)

Ran on 8-GPU MI300X with alltoall_perf -D3 (NCCL_GIN_TYPE=5 on the pre-pr9927 stack; current stack uses NCCL_GIN_TYPE=6), -n 8 -g 1, 128 B–1 MiB:

• Mode 1 (-B 1, augment): normal numbers plus a #[a2a-devtime] line per size; datacheck Out of bounds values : 0 OK.
• Mode 2 (-B 2, device-time-only): reported metric is the in-kernel device latency and tracks the mode-1 #[a2a-devtime] value (e.g. 1 MiB: device ~40.9 µs vs host-graph ~28.7 µs); datacheck clean.
• Mode 2 + -H 1: the timed-kernel output is validated against expected — 0 wrong elements across all sizes; datacheck clean.

Made with Cursor

[therock-pr-bot]

✅ All Policy Checks Passed

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	✅ Pass	—
🚫 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?

[dlamd1dai]

@pvallem — both Aug 27 review items are addressed in 7386034:

1. Hybrid join at stamps only — hybridTimedGridJoin runs at i == skip and after the loop, not every iteration.
2. testStreamSynchronize in devtime — gin_devtime::measure() uses testStreamSynchronize so mode 2 honors -T / ncclCommAbort.

Replies are on the two threads. Could you take another look when you have a moment?

CI: device-api (incl. AllToAll devtime smoke 3/3), GIN, host unit tests, and rccl precheckin are green. mci/precheckin(rccl-tests) is still red (Test stage); unrelated infra reds (therock-pr-bot, Multi-Arch summary) also present.

[systems-assistant]

TheRock Submodule Bump Activity

Newest first

• 2026-08-28 — Included in TheRock via ROCm/TheRock#7710.