Repository navigation
feat(rocshmem): rework the tile-tests - #11770
Conversation
✅ All Policy Checks Passed
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
There was a problem hiding this comment.
🟡 Changes recommended
Unresolved critical and moderate correctness, timing, sizing, and resource issues remain.
Get a fresh assessment by requesting another Copilot review.
Pull request overview
This PR reworks rocSHMEM tile functional tests with multi-wave collectives, configurable RMA sweeps, additional layouts, and expanded test coverage.
Changes:
- Adds wave/WG row-major and column-major RMA variants.
- Updates broadcast, allgather, and reduce tests for multiple wave contexts.
- Expands CTest and driver coverage, including 1 MiB variants.
File summaries
| File | Description and final review comments |
|---|---|
projects/rocshmem/tests/functional_tests/tile_rma_tester.hpp |
Adds configurable tile dimensions. |
projects/rocshmem/tests/functional_tests/tile_rma_tester.cpp |
Implements layouts, size sweeps, verification, and accounting. Critical · 3 votes: Row-major wave slots use an incorrect stride, causing overlapping source/destination slices. Critical · 2 votes: Wave verification uses the service-thread count and can skip verification entirely. Moderate · 2 votes: Arbitrary layouts cannot exercise the advertised 1 MiB payload. Moderate · 1 vote: Requests below one column silently become 16 KiB tests. Moderate · 1 vote: End timestamps can precede completion of all quiet operations. |
projects/rocshmem/tests/functional_tests/tile_reduce_tester.hpp |
Updates the wave reduction interface. |
projects/rocshmem/tests/functional_tests/tile_reduce_tester.cpp |
Adds multi-wave contexts and teams. Moderate · 1 vote: Timing buffers and message counts are not populated, producing invalid performance results. |
projects/rocshmem/tests/functional_tests/tile_broadcast_tester.hpp |
Updates the wave broadcast interface. |
projects/rocshmem/tests/functional_tests/tile_broadcast_tester.cpp |
Adds multi-wave broadcast handling. Moderate · 1 vote: Timing buffers and message counts are not populated, producing invalid performance results. |
projects/rocshmem/tests/functional_tests/tile_allgather_tester.hpp |
Updates the wave allgather interface. |
projects/rocshmem/tests/functional_tests/tile_allgather_tester.cpp |
Adds multi-wave allgather handling. Moderate · 1 vote: Timing buffers and message counts are not populated, producing invalid performance results. |
projects/rocshmem/tests/functional_tests/tester.hpp |
Registers new RMA test types. |
projects/rocshmem/tests/functional_tests/tester.cpp |
Creates testers for the new variants. |
projects/rocshmem/tests/functional_tests/CTestFunctionalTests.cmake |
Registers expanded test configurations. Moderate · 3 votes: WG tests multiply context allocation by NUM_WF despite using one context per WG, potentially exhausting backend resources. |
projects/rocshmem/scripts/functional_tests/driver.sh |
Adds driver mappings and test invocations. Moderate · 1 vote: New WG 1 MiB cases can allocate approximately 1–2 GiB per buffer and exhaust the 6 GiB heap. |
Review details
Suppressed comments (6)
projects/rocshmem/scripts/functional_tests/driver.sh:1045
- The new WG 1 MiB row-major cases pass
-z $((WAVE_SIZE * 16));TileRMATesterallocatesargs.num_threadsslots, so on a 64-lane GPU each row-major slot is 2 MiB and each of the symmetric source/destination buffers is about 2 GiB even though a WG collective uses one tile. The corresponding col-major/contiguous cases are also about 1 GiB per buffer, which can exhaust the 6 GiB heap and makes the new driver sweep unnecessarily large. Allocate by collective count or use a wave-count path that does not inflate the host allocation count.
ExecTest "tile_put_wg_rowmajor" 2 1 $((WAVE_SIZE * 16))
ExecTest "tile_put_wg_rowmajor" 2 1 $((WAVE_SIZE * 16)) 1048576
projects/rocshmem/tests/functional_tests/tile_allgather_tester.cpp:348
- This new wave/WG launch path still does not populate
start_time/end_timeornum_timed_msgs.Tester::print()subsequently reads the uninitialized timer buffers and computes zero message rate/bandwidth, so these tests cannot produce the comparable performance numbers described by the PR. Populate the tester timing/count contract or use the HIP event timing for these kernels.
hipLaunchKernelGGL(TileAllgatherWaveTest, gridSize, blockSize,
wave_shared, stream,
teams, source, dest, tile_extent_0, tile_extent_1,
args.myid, n_pes, _shmem_context,
wf_size, num_warps, error_flag);
projects/rocshmem/tests/functional_tests/tile_broadcast_tester.cpp:329
- This new wave/WG launch path still does not populate
start_time/end_timeornum_timed_msgs.Tester::print()subsequently reads the uninitialized timer buffers and computes zero message rate/bandwidth, so these tests cannot produce the comparable performance numbers described by the PR. Populate the tester timing/count contract or use the HIP event timing for these kernels.
hipLaunchKernelGGL(TileBroadcastWaveTest, gridSize, blockSize,
wave_shared, stream,
teams, source, dest, tile_extent_0, tile_extent_1,
args.myid, pe_root, _shmem_context,
wf_size, num_warps, error_flag);
projects/rocshmem/tests/functional_tests/tile_reduce_tester.cpp:526
- This new wave/WG launch path still does not populate
start_time/end_timeornum_timed_msgs.Tester::print()subsequently reads the uninitialized timer buffers and computes zero message rate/bandwidth, so these tests cannot produce the comparable performance numbers described by the PR. Populate the tester timing/count contract or use the HIP event timing for these kernels.
hipLaunchKernelGGL(TileReduceWaveTest, gridSize, blockSize,
wave_shared, stream,
teams, source, sum_dest, max_dest, min_dest,
short_source, short_sum_dest, short_max_dest,
short_min_dest, int_source, int_sum_dest,
projects/rocshmem/tests/functional_tests/tile_rma_tester.cpp:359
- When
-sis smaller than one 64-row column (256 bytes),derivedis zero and this leavestile_extent_1at the default 64. The tester then changes the requested cap into a 16 KiB test instead of honoring it or using the smallest valid 256-byte tile. Clamp the extent to one column (or reject the invalid request) rather than silently falling back to 64.
if (args.max_msg_size_set) {
int derived = static_cast<int>(args.max_msg_size / (tile_extent_0 * sizeof(float)));
if (derived >= 1) {
tile_extent_1 = derived;
}
projects/rocshmem/tests/functional_tests/tile_rma_tester.cpp:709
- These new wave cases report one message per wave, but
TileRMATestrecordsend_time[wg_id]from every thread immediately after only thread 0 callsrocshmem_ctx_quiet; there is no post-quiet barrier. A nonzero thread can therefore publish an end time before the quiet has completed all wave operations, overstating bandwidth for these newly reported multi-wave counts. Record the end timestamp after a WG-wide barrier followingquiet, from one thread.
case TilePutWaveRowMajorTestType:
case TilePutWaveColumnMajorTestType:
case TileGetWaveRowMajorTestType:
case TileGetWaveColumnMajorTestType:
tiles_per_loop = gridSize.x * num_warps;
- Files reviewed: 12/12 changed files
- Comments generated: 4
- Review effort level: Lite
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
3d022f8 to
d474016
Compare
ddebonis-amd
left a comment
There was a problem hiding this comment.
Missing tile reduce tests for types (long / int / short) and ops (min / max). Only testing float / sum.
d474016 to
d2172ed
Compare
…ams and contexts ## Motivation Apply the same per-wave context pattern introduced in dd5b69b for the wave-level collective testers to the tile_broadcast wave and wg tests. ## Technical Details tile_broadcast_tester.cpp / .hpp: - TileBroadcastWaveTest: replaced single __shared__ rocshmem_ctx_t ctx + single team with extern __shared__ rocshmem_ctx_t ctx_array[] sized to num_waves_per_wg; each wave creates its own context from a dedicated team via rocshmem_wg_team_create_ctx and operates on an independent source/dest buffer slice indexed by flat_wf_id (wg_id * num_waves_per_wg + wf_id); all waves now execute concurrently - TileBroadcastTest (WG-level): migrated from __shared__ ctx to extern __shared__ ctx_array[0] for consistency with the pattern - Constructor sizes the teams array to num_wgs * num_warps (one team per wave, the maximum needed across all test types) - preLaunchKernel/postLaunchKernel: create and destroy the correct number of teams per test type (wave: num_wgs * num_warps, WG: num_wgs, thread: 1) - launchKernel: passes num_warps * sizeof(rocshmem_ctx_t) as dynamic shared memory for the wave test; sizeof(rocshmem_ctx_t) for the WG test driver.sh: - tile_broadcast_wave runs with --num-wf 4 (4 waves per WG) and ROCSHMEM_MAX_NUM_CONTEXTS = NUM_WGS * 4, matching the per-wave context requirement; the 4-PE variant now also uses 4 WGs - export/unset pattern mirrors the existing wave collective tests CTestFunctionalTests.cmake: - _add_single_rocshmem_test: compute ROCSHMEM_MAX_NUM_CONTEXTS as WORKGROUPS * NUM_WF when NUM_WF is set, instead of just WORKGROUPS; this correctly handles all current and future wave-level tests - tile_broadcast_wave entries updated to NUM_WF 4, with the 4-PE variant using WORKGROUPS 4 (yielding MAX_NUM_CONTEXTS=16) Co-Authored-By: Claude <noreply@anthropic.com>
…ams and contexts Apply the same per-wave context pattern from dd5b69b and the tile_broadcast rework to the tile_allgather wave and wg tests. tile_allgather_tester.cpp / .hpp: - TileAllgatherWaveTest: replaced single __shared__ rocshmem_ctx_t ctx + single team with extern __shared__ rocshmem_ctx_t ctx_array[] sized to num_waves_per_wg; each wave creates its own context from a dedicated team and operates on independent buffer slices indexed by flat_wf_id: source at flat_wf_id * tile_size, dest at flat_wf_id * tile_size * n_pes - TileAllgatherTest (WG-level): migrated from __shared__ ctx to extern __shared__ ctx_array[0] - Constructor sizes teams array to num_wgs * num_warps; allocates source and dest with num_teams slots to cover all wave slots - preLaunchKernel/postLaunchKernel: create/destroy teams per test type (wave: num_wgs * num_warps, WG: num_wgs, thread: 1) - launchKernel: passes dynamic shared memory sized to num_warps * sizeof(rocshmem_ctx_t) for wave test; sizeof(rocshmem_ctx_t) for WG test - resetBuffers: initializes source as myid * 100 + slot * 1000 + i so each wave slot carries a distinct value; verification checks the same driver.sh: - tile_allgather_wave runs with --num-wf 4 and ROCSHMEM_MAX_NUM_CONTEXTS = NUM_WGS * 4; 4-PE variant uses 4 WGs CTestFunctionalTests.cmake: - tile_allgather_wave entries updated to NUM_WF 4, 4-PE variant uses WORKGROUPS 4 (MAX_NUM_CONTEXTS computed as 4 and 16 respectively by the existing WORKGROUPS * NUM_WF formula) Co-Authored-By: Claude <noreply@anthropic.com>
… and contexts Apply the same per-wave context pattern from dd5b69b and the previous tile_broadcast/allgather reworks to the tile_reduce wave and wg tests. tile_reduce_tester.cpp / .hpp: - TileReduceWaveTest: replaced single __shared__ rocshmem_ctx_t ctx + single team with extern __shared__ rocshmem_ctx_t ctx_array[] sized to num_waves_per_wg; each wave creates its own context and operates on an independent buffer slice (offset = flat_wf_id * matrix_size) for all four data types (float, short, int, long); verification runs per-wave on the root PE using t_id % wf_size == 0 - TileReduceTest (WG-level): migrated from __shared__ ctx to extern __shared__ ctx_array[0] - Constructor: sizes num_teams to num_wgs * num_warps for wave tests, num_wgs for WG tests, 1 for thread tests, so one allocation covers all - preLaunchKernel/postLaunchKernel: simplified to a single loop over num_teams (already set correctly per test type in the constructor) - launchKernel: passes num_warps * sizeof(rocshmem_ctx_t) dynamic shared memory for wave test; sizeof(rocshmem_ctx_t) for WG test; wave test now launches gridSize blocks instead of dim3(1) driver.sh: - tile_reduce_wave runs with --num-wf 4 and ROCSHMEM_MAX_NUM_CONTEXTS = NUM_WGS * 4; 4-PE variant uses 4 WGs CTestFunctionalTests.cmake: - tile_reduce_wave entries updated to NUM_WF 4, 4-PE variant uses WORKGROUPS 4; tile_reduce_wg entries updated to NUM_WF 1 WORKGROUPS 4 Co-Authored-By: Claude <noreply@anthropic.com>
…gurable sweeps ## Motivation The tile RMA tests (tile_put/get_*) previously used hardcoded 64×64 tile dimensions and ignored the -s / MAX_MSG_SIZE argument entirely, making it impossible to sweep data sizes from the driver or CTest. Reported bandwidths were meaningless for sub-tile message sizes since the kernel always transferred the full tile regardless. ## Technical Details tile_rma_tester.hpp: - Added tile_extent_0 and tile_extent_1 as protected member fields tile_rma_tester.cpp: - Constructor derives tile_extent_1 from args.max_msg_size when set (max_msg_size = tile_extent_0 * tile_extent_1 * sizeof(float)) - tile_extent_0 is fixed at 64 rows; tile_extent_1 scales with message size - Arbitrary-stride layout clamps tile_extent_1 to 85 (floor(257/3)) to keep column accesses within the fixed row_stride=257 - Sets this->args.min_msg_size = one column (tile_extent_0 * sizeof(float)) and max_msg_size = full tile bytes, so the framework sweep runs from one column up to the full tile — every data point is a valid tile shape - launchKernel: derives ke_tile_extent_1 from the current iteration's size so each sweep point transfers the correctly-sized tile and reports accurate bandwidth - resetBuffers: re-initializes the source buffer using the current iteration's tile dimensions (ke_tile_extent_1-based offsets) so that source layout matches exactly what the kernel expects for each sweep point - verifyResults: re-derives tile dimensions from size using the same logic as launchKernel so only the actually-transferred elements are checked - Buffer allocation in the constructor uses the max tile_extent_1 to cover all sweep points within a single allocation driver.sh: - Added a 1MB-max variant alongside each existing tile RMA test; sweeps from 256B (1 column) to 1048576B (4096 columns) CTestFunctionalTests.cmake: - Added MAX_MSG_SIZE 1048576 variants for all tile put/get test groups Co-Authored-By: Claude <noreply@anthropic.com>
…t and tile_get The tile RMA wave-level tests previously covered only the contiguous layout. Row-major and col-major variants exist at the thread level but were missing for the wave-collective operations, preventing performance comparisons across memory layouts at the wave level. tester.hpp: - Added 4 new enum entries appended after the last existing entry (SdmaPutNbi=161): TilePutWaveRowMajor=162, TilePutWaveColumnMajor=163, TileGetWaveRowMajor=164, TileGetWaveColumnMajor=165 tile_rma_tester.cpp: - Offset calculation: new types added to the wave-collective branch (offset = matrix_size * flat_wave_id) - Kernel branches: 4 new if-constexpr blocks using rocshmem_ctx_tile_put_wave / rocshmem_ctx_tile_get_wave with row-major (row_stride=2*t1) and col-major (row_stride=1, col_stride=t0) tensor layouts - Constructor: new types routed to correct buffer_elements_per_thread cases, source/dest assignment (put→local/remote, get→remote/local) - resetBuffers: new types added to the appropriate stride/slot-size cases - launchKernel: 4 new LAUNCH_TILE_RMA_TEST dispatch cases; new types added to the wave tiles_per_loop branch (gridSize.x * num_warps) - verifyResults: new types added to check_id, row/col_stride, num_tiles_transferred, and buffer_elements_per_thread switches tester.cpp: - 4 new TileRMATester instantiation cases in Tester::create() driver.sh: - Added TEST_NUMBERS entries (162-165) - Added ExecTest calls with default and 1MB size variants alongside the existing wave contiguous tests CTestFunctionalTests.cmake: - Added set(TEST_tile_*) entries for the 4 new test numbers - Added add_rocshmem_functional_test entries (default + MAX_MSG_SIZE 1048576) in the wave put and get test groups Co-Authored-By: Claude <noreply@anthropic.com>
…wg and tile_get_wg Apply the same pattern from the wave variant commit (ddde4d522e) to WG-collective tile RMA operations. 4 new enum entries appended after TileGetWaveColumnMajor (165): TilePutWGRowMajor=166, TilePutWGColumnMajor=167, TileGetWGRowMajor=168, TileGetWGColumnMajor=169 tile_rma_tester.cpp: offset uses wg_id; kernel branches use rocshmem_ctx_tile_put_wg / rocshmem_ctx_tile_get_wg with row-major (row_stride=2*t1) and col-major (row_stride=1, col_stride=t0) layouts; num_timed_msgs uses gridSize.x (one tile per WG per loop); all switch statements updated consistently across offset, buffer sizing, source/dest assignment, resetBuffers, launchKernel, and verifyResults. tester.cpp: 4 new TileRMATester instantiation cases. driver.sh: TEST_NUMBERS entries 166-169; ExecTest calls with default and 1MB size variants using $((WAVE_SIZE * 16)) threads (matching existing tile_put_wg_contiguous pattern). CTestFunctionalTests.cmake: set() entries and add_rocshmem_functional_test entries (NUM_WF 16, default + MAX_MSG_SIZE 1048576) in the WG put/get groups. Co-Authored-By: Claude <noreply@anthropic.com>
tile_memcpy_contig<Kind, TileScope::Wg> was calling memcpy_lane through thread 0 only, leaving the other 1023 threads idle for the contiguous fast path. Replace with memcpy_wg, which distributes the transfer across all threads in the workgroup — matching the pattern already used by the Wave scope (memcpy_wave) and the RowContig/ColContig strided paths. Observed impact on tile_put_wg_contiguous / tile_get_wg_contiguous at 1MB with 1 WG / 1024 threads: Before: ~6000 µs / 0.16 GB/s (serialized through thread 0) After: ~44 µs / 22 GB/s (all 1024 threads cooperate) Co-Authored-By: Claude <noreply@anthropic.com>
The 1MB sweep (min_msg_size=256B, max_msg_size=1MB) already covers all data points of the default no-size run (max_msg_size=16384B), making the latter redundant. Remove the no-size ExecTest calls from driver.sh and the corresponding add_rocshmem_functional_test entries from CTestFunctionalTests.cmake for all tile put/get tests. Co-Authored-By: Claude <noreply@anthropic.com>
…testers ## tile_broadcast and tile_allgather testers - Constructor: derive tile_extent_1 from max_msg_size; set min_msg_size = tile_extent_0 * sizeof(float) (one column) and max_msg_size = full tile so the framework sweeps from one column to the configured maximum - Constructor: fix num_timers for wave variant (num_wgs * num_warps) - Kernels: add loop/skip timing with wall_clock64; start_time/end_time arrays now instrumented for all three variants (thread, wave, WG) - launchKernel: derive ke_tile_extent_1 from current iteration's size and pass it to the kernel; set num_msgs/num_timed_msgs correctly - resetBuffers: reinitialize source using ke_tile_extent_1-sized slot strides so source layout matches what the kernel expects each sweep step - verifyResults: print PASSED only at the last sweep iteration (size == max_msg_size); exit on first failure covers all prior iterations ## tile_reduce tester — full template refactor Replaced the monolithic TileReduceTester (4 types, 3 ops, all timed together) with TileReduceTester<T, Op> templated on element type and reduction operation, matching the pattern used by TeamReductionTester: - TILE_REDUCE_DEF_GEN macro specialises tile_reduce_lane/wave/wg<T,Op> for float/short/int/long × SUM/MAX/MIN dispatching to the typed tile reduce API - Three kernel templates (TileReduceThreadTest, TileReduceWaveTest, TileReduceWGTest) each time exactly one type+op in a loop - Host class: same constructor sweep-pin and ke_tile_extent_1 pattern as broadcast/allgather; verifyResults checks result on root PE host-side - tester.cpp: instantiates TileReduceTester<float, ROCSHMEM_SUM> for all three test types with init (s=1, r=0) and verify (v == n_pes) lambdas - CMakeLists.txt: tile_reduce_tester.cpp removed from direct compilation (it is included by the header as a template implementation) ## driver.sh and CTestFunctionalTests.cmake - All tile collective tests (broadcast, allgather, reduce) now run up to 1MB message size - Wave collective variants use 1 WG with NUM_WF=4 for both 2-PE and 4-PE cases (was 4 WGs for 4-PE case); ROCSHMEM_MAX_NUM_CONTEXTS = 1 * 4 = 4 - Stale old-style reduce size entries (262144B, 2048B) replaced uniformly with 1048576B across all three reduce variants Co-Authored-By: Claude <noreply@anthropic.com>
Both the IPC and GDA backends previously silently ignored ndim values
outside {1, 2}: tile_memcpy_rma and tile_make_view have no else branch
for ndim > 2, resulting in a no-op (no data transferred, no error).
The reduce inner impls rejected ndim <= 0 but not ndim > 2.
Add explicit ndim < 1 || ndim > 2 guards at the entry point of all six
RMA functions in each backend (tile_put, tile_put_wave, tile_put_wg,
tile_get, tile_get_wave, tile_get_wg), returning ROCSHMEM_ERROR
immediately. Collectives (allgather, broadcast) propagate naturally since
they delegate to the guarded RMA functions internally.
Extend the existing ndim <= 0 check in both backends' inner reduce
implementations to ndim < 1 || ndim > 2, covering the previously
unhandled ndim > 2 case.
Co-Authored-By: Claude <noreply@anthropic.com>
…nsts Replace bare integer literals used as tile dimensions with named constexpr members in each tester class: - TileRMATester: DEFAULT_TILE_ROWS=64, DEFAULT_TILE_COLS=64, ARBITRARY_MAX_COLS=85 (floor(257/3), the max col count for the arbitrary-stride layout where row_stride=257 and col_stride=3) - TileBroadcastTester: DEFAULT_TILE_ROWS=8, DEFAULT_TILE_COLS=8 - TileAllgatherTester: DEFAULT_TILE_ROWS=8, DEFAULT_TILE_COLS=8 - TileReduceTester: DEFAULT_TILE_ROWS=8, DEFAULT_TILE_COLS=8 Co-Authored-By: Claude <noreply@anthropic.com>
…ss tests The template refactor reduced each tile_reduce test type to a single TileReduceTester<float, ROCSHMEM_SUM> instance, dropping the short/int/long types and MAX/MIN operations that the previous monolithic tester covered. Restore full coverage by pushing 12 instances per test type (float/short/int/long × SUM/MAX/MIN) using a TILE_REDUCE_PUSH macro. Verification uses simple constant source values (s=1, r=0 for sum/max; r=2 for min) so expected results are type-safe and unambiguous: SUM → v == n_pes, MAX → v == 1, MIN → v == 1 Co-Authored-By: Claude <noreply@anthropic.com>
Running 12 type+op combinations per test invocation makes the 1MB 4-PE reduce tests significantly slower. Cap them at 65536B to keep the suite runtime reasonable while still sweeping through all message sizes up to 64KB (256B columns → 16KB tile at the maximum). Co-Authored-By: Claude <noreply@anthropic.com>
d2172ed to
4b61dba
Compare
TheRock Submodule Bump ActivityNewest first
|
Motivation
Rework tile tests to cover a wider range of possible configurations, and produce performance numbers that are more easily comparable to the non-tile variants of the tests.
Technical Details
tile_broadcast / tile_allgather / tile_reduce wave & WG tests: enable using multiple waves in a wg and have each wave use a separate context. Matches the changes made to the wave-collectives last months. driver.sh and CTest updated to use --num-wf 4 with ROCSHMEM_MAX_NUM_CONTEXTS = NUM_WGS × NUM_WF.
Tile RMA size-configurable sweeps: Wired up max_msg_size in the tile RMA tester so tile dimensions scale with -s; the message sweep now runs from one column to the full tile (one valid tile shape per data point) with accurate volume and bandwidth reporting. Added 1 MB variants to driver.sh and CTest for all tile put/get tests.
New wave-level row-major and col-major tile RMA tests
New wg-level row-major and col-major tile RMA tests
Update:
Issue Tracking
JIRA ID: AIROCSHMEM-517
Submission Checklist