Skip to content

test(rccl): host microtests for DDA IPC AllReduce dispatch, gates and flat/tree choice #13240

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
pvallem merged 8 commits into from
Oct 19, 2026
Merged

test(rccl): host microtests for DDA IPC AllReduce dispatch, gates and flat/tree choice #13240

Show file tree
Hide file tree
Changes from 7 commits
Commits
Show all changes
8 commits
Select commit Hold shift + click to select a range
c072ac0
test(rccl): host microtests for DDA IPC all-reduce eligibility, dispa…
pvallem Oct 8, 2026
2a2e370
test(rccl): cover rcclDdaEnabled unset launch order, gfx950/gfx1250 r…
pvallem Oct 8, 2026
8c924ec
test(rccl): keep DDA IPC tree-selection checks valid with in-kernel s…
pvallem Oct 10, 2026
4f8c04b
test(rccl): fake the HIP profile shadow registration for rccl-UnitTes…
pvallem Oct 11, 2026
3839b78
test(rccl): pin DDA IPC peer pointers, stream, scratch-copy error and…
pvallem Oct 12, 2026
9112f95
Merge remote-tracking branch 'origin/develop' into users/pvallem/upsh…
pvallem Oct 17, 2026
6bd5dbc
test(rccl): pin the exact DDA IPC tree kernel and the launch grid/block
pvallem Oct 17, 2026
2f2d789
Merge branch 'develop' into users/pvallem/upshift-dda-allreduce-host-…
pvallem Oct 17, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
24 changes: 24 additions & 0 deletions projects/rccl/test/host/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -403,6 +403,15 @@ function(rccl_define_micro_source_lists)
../common/main_altrsmi.cpp
PARENT_SCOPE)

# rccl-UnitTestsMicroDdaIpc -- src/algorithms/dda/all_reduce/dda_all_reduce_ipc.cu. Own binary:
# rccl-UnitTestsMicro links fakes/dda_fakes.cc, which defines the same ncclAllReduceDdaIpc* symbols.
set(TEST_MICRO_DDA_IPC_SOURCE_FILES
dda-allreduce-ipc-test.cc
fakes/hip_fakes.cc
fakes/nccl_fakes.cc
../common/main_altrsmi.cpp
PARENT_SCOPE)

# rccl-UnitTestsMicroRomeModels
set(TEST_MICRO_ROME_MODELS_SOURCE_FILES
rome-models-test.cc
Expand Down Expand Up @@ -988,6 +997,21 @@ rccl_add_micro_binary(rccl-UnitTestsMicroGinHost
GIN_HOST_CC_PATH="${RCCL_HOST_HIPIFY_DIR}/src/gin/gin_host.cc"
CATEGORIES_YAML "${CMAKE_CURRENT_SOURCE_DIR}/../test_categories_micro_gin_host.yaml")

# ---------------------------------------------------------------------------
# rccl-UnitTestsMicroDdaIpc -- src/algorithms/dda/all_reduce/dda_all_reduce_ipc.cu
# ---------------------------------------------------------------------------
rccl_add_micro_binary(rccl-UnitTestsMicroDdaIpc
SOURCES ${TEST_MICRO_DDA_IPC_SOURCE_FILES}
EXTRA_DEFS
${_MICRO_MODE_DEFS}
DDA_ALL_REDUCE_IPC_CC_PATH="${RCCL_HOST_HIPIFY_DIR}/src/algorithms/dda/all_reduce/dda_all_reduce_ipc.cu.cpp"
CATEGORIES_YAML "${CMAKE_CURRENT_SOURCE_DIR}/../test_categories_micro_dda_ipc.yaml")
# The TU defines __global__ kernels; an embedded (empty) GPU binary keeps clang from referencing an
# external __hip_fatbin_<cuid>. Registration itself is faked in fakes/hip_fakes.cc.
file(WRITE "${CMAKE_CURRENT_BINARY_DIR}/empty_gpubinary" "")
target_compile_options(rccl-UnitTestsMicroDdaIpc PRIVATE
"SHELL:-Xclang -fcuda-include-gpubinary -Xclang ${CMAKE_CURRENT_BINARY_DIR}/empty_gpubinary")

# sym_kernels_host.cc only exists when sym kernels are generated (GENERATE_SYM_KERNELS
# in-build / present on disk standalone).
if(_MICRO_BUILD_SYM_KERNELS)
Expand Down
11 changes: 9 additions & 2 deletions projects/rccl/test/host/MICROTEST_README.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -322,6 +322,12 @@ export colliding non-`static` symbols; otherwise a unit needs its own binary:
`NCCL_GIN_PROXY_NTHREADS` progress-thread assignment. Its own binary, not
sharing `rccl-UnitTestsMicro`: `gin-plugin-init-test.cc` already defines
`ncclParamGinEnable` there. See `test_categories_micro_gin_host.yaml`.
- **`rccl-UnitTestsMicroDdaIpc`**: `algorithms/dda/all_reduce/dda_all_reduce_ipc.cu`
(`DDA_ALL_REDUCE_IPC_CC_PATH`, from `dda-allreduce-ipc-test.cc`); suites
`DdaAllReduceIpc*`. Eligibility gates, datatype dispatch and flat/tree kernel
selection, observed through the `g_hipExtLaunchKernel` seam. Its own binary:
`rccl-UnitTestsMicro` links `fakes/dda_fakes.cc`, which defines the same symbols.
See `test_categories_micro_dda_ipc.yaml`.

- **`rccl-UnitTestsMicroRmaProxy`** — `src/rma/rma_proxy.cc` (via
`RMA_PROXY_CC_PATH`, from `rma-proxy-test.cc`); suite
Expand Down Expand Up @@ -856,8 +862,8 @@ above (`./install.sh -t`, wired via `add_subdirectory(host)`), the same file
can be configured **directly** to build every host binary — `rccl-HostUnitTests`,
`rccl-UnitTestsMicro`, `rccl-UnitTestsMicroWarpSpeed`,
`rccl-UnitTestsMicroInit[-uncached|-faultinj]`, `rccl-UnitTestsMicroEnqueue[-devlinker]`,
`rccl-UnitTestsMicroGinHost`, `rccl-UnitTestsMicroCe`, `rccl-UnitTestsMicroDiagnostics`,
`rccl-UnitTestsMicroSymKernels`, `rccl-UnitTestsMicroTaskPrep` and
`rccl-UnitTestsMicroGinHost`, `rccl-UnitTestsMicroDdaIpc`, `rccl-UnitTestsMicroCe`,
`rccl-UnitTestsMicroDiagnostics`, `rccl-UnitTestsMicroSymKernels`, `rccl-UnitTestsMicroTaskPrep` and
`rccl-UnitTestsMicroRmaProxy` — **without configuring/building all of
librccl**. It compiles just the tests + fakes + the hipified unit-under-test
sources.
Expand Down Expand Up @@ -899,6 +905,7 @@ cmake --build build -j"$(nproc)"
./build/rccl-UnitTestsMicroConnect # src/graph/connect.cc Preset/Postset tests
./build/rccl-UnitTestsMicroDiagnostics # src/diagnostics/{p2p,ib_write_bw}.cc tests
./build/rccl-UnitTestsMicroGinHost # src/gin/gin_host.cc GIN_PROXY_NTHREADS
./build/rccl-UnitTestsMicroDdaIpc # src/algorithms/dda/all_reduce/dda_all_reduce_ipc.cu
./build/rccl-UnitTestsMicroDiagnostics # src/diagnostics/p2p.cc tests
./build/rccl-UnitTestsMicroCe # src/ce_coll.cc tests
./build/rccl-UnitTestsMicroRmaProxy # src/rma/rma_proxy.cc tests
Expand Down
Loading
Loading