Enable gfx1100 (RDNA3, wave32) builds#11
Open
jeffdaily wants to merge 2 commits into
Open
Conversation
hipRaft already carries the pieces needed to run on a 32-lane wavefront:
the device-side per-arch warp width (raft::WarpSize / hip_warp_primitives::
WAVEFRONT_SIZE resolves to 32 on RDNA, 64 on CDNA) and the host-side runtime
warp-size query (raft::host_warp_size via rocprim::host_warp_size), which the
grid_1d_* launch-geometry helpers and the wave-sensitive kernels already use.
What was missing was wiring the build for a wave32 target.
The cuco hashmap dependency selects its lane width at configure time from the
USE_WARPSIZE_32 option and hard-errors when a gfx11xx architecture is requested
without it. This change auto-enables USE_WARPSIZE_32 in cpp/CMakeLists.txt when
every requested HIP architecture is a wave32 target (gfx10xx/gfx11xx/gfx12xx),
so `build.sh --gpu-arch=gfx1100` configures and builds without manual flags.
Mixing wave32 and wave64 architectures in one build is rejected with a clear
message, since cuco's lane width is a single compile-time choice. CUDA and CDNA
(wave64) builds are unchanged: the new block is gated on a wave32-only arch list
and leaves USE_WARPSIZE_32 at its default otherwise.
No device or host source changed; the wave-size handling was already correct.
The fix is purely the build-system gate.
Test Plan
Built and validated on an AMD Radeon Pro W7800 (gfx1100, RDNA3, wave32),
ROCm 7.2.1, with one GPU pinned via HIP_VISIBLE_DEVICES.
Library (gfx1100, wave32):
```
./build.sh -v libraft --compile-lib --gpu-arch="gfx1100"
```
Compiles clean; cuco builds with -DCUCO_USE_WARPSIZE_32 and every translation
unit with --offload-arch=gfx1100.
Tests (gfx1100, wave32):
```
./build.sh -v libraft tests --gpu-arch="gfx1100"
HIP_VISIBLE_DEVICES=<idle-gpu> ./cpp/build/gtests/<group>
```
Per-group results on gfx1100 (wave32):
```
UTILS_TEST 181/181 (popc/ballot, reduction, binary-reduce, atomic-inc-warp, bitonic)
LINALG_TEST 2070/2070 (coalesced/strided reduction, reduce, map_then_reduce)
SRC_LINALG_TEST 37/37 (coalesced reduction)
NEIGHBORS_UTILS_TEST 14/14 (warp-sort merge network, warp/block select; 10 wave64-only cases skip by design)
MATRIX_TEST 835/835
MATRIX_SELECT_TEST 567/567
CORE_TEST 269/269
RANDOM_TEST 622/622
RUNTIME_RANDOM_TEST 48/48
LABEL_TEST 14/14
STATS_TEST 916/918 (2 MeanTestD failures at N=2^27, double, tolerance 1e-8: a
large-N reduction accumulation-order rounding artifact, ~1.7e-7
drift; the float N=2^27 cases pass at 1e-4 tolerance)
```
No HSA memory faults and no ballot/reduction wrong-result nondeterminism were
observed in any group; the wave32 warp-primitive, reduction, and ballot paths
are correct.
This work was authored with the assistance of Claude, an AI assistant.
The two N=2^27 double MeanTest cases assert the computed column mean is within 1e-8 of the population mean 0.1. The data is normal(0.1, 1e-4), so the sampling standard error of a mean of 2^27 draws is 1e-4/sqrt(2^27) ~ 8.6e-9; a correct result routinely deviates from 0.1 by a couple of sigma (observed 1.7e-8), which exceeds the 1e-8 tolerance. The deviation is in the sampled data, not the arithmetic: the GPU result matches a host long-double Kahan mean of the same buffer to ~1e-17. On a 32-lane wavefront the lane- indexed RNG draws a different sample than on a 64-lane wavefront, so the over-tight tolerance only became visible there. Widen the tolerance for those two cases to 5e-8 (~6 sigma) so the test reflects the sampling noise of the estimator rather than failing on the particular RNG draw, while still catching a gross reduction error. Only the two large-N double cases change; every other mean case is untouched. Test Plan ``` cmake --build . --target STATS_TEST -j16 HIP_VISIBLE_DEVICES=<idle gpu> ./gtests/STATS_TEST \ --gtest_filter='*MeanTestD.Result/20:*MeanTestD.Result/21' ``` On an AMD Radeon Pro W7800 (gfx1100, ROCm 7.2.1) both cases pass; the full STATS_TEST is then 918/918. This work was authored with the assistance of Claude, an AI assistant.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
hipRaft already carries the wave-size machinery to run on a 32-lane wavefront: the device-side per-arch warp width (
raft::WarpSize/hip_warp_primitives::WAVEFRONT_SIZEresolves to 32 on RDNA, 64 on CDNA) and the host-side runtime query (raft::host_warp_size) used by thegrid_1d_*launch helpers. The only missing piece was wiring the build for a wave32 target.The cuco hashmap dependency selects its lane width at configure time from
USE_WARPSIZE_32and hard-errors when a gfx11xx architecture is requested without it. This change auto-enablesUSE_WARPSIZE_32when every requested HIP architecture is a wave32 target (gfx10xx/gfx11xx/gfx12xx), sobuild.sh --gpu-arch=gfx1100configures without manual flags. Mixing wave32 and wave64 architectures in one build is rejected with a clear message, since cuco's lane width is a single compile-time choice. CUDA and CDNA (wave64) builds are unchanged.It also widens the two N=2^27 double
MeanTestDtolerances from 1e-8 to 5e-8. The data isnormal(0.1, 1e-4), whose sample-mean standard error over 2^27 draws is 1e-4/sqrt(2^27) ~ 8.6e-9, so a correct result deviates from the population mean 0.1 by a couple of sigma (observed 1.7e-8), exceeding the 1e-8 tolerance. The deviation is in the sampled data, not the arithmetic: the GPU mean matches a host long-double Kahan mean of the same buffer to ~1e-17. The lane-indexed RNG draws a different sample on a 32-lane wavefront than on a 64-lane one, so the over-tight tolerance only became visible there; 5e-8 (~6 sigma) reflects the estimator's sampling noise while still catching a gross reduction error.Validation
AMD Radeon Pro W7800 (gfx1100, RDNA3, wave32), ROCm 7.2.1, one GPU pinned via
HIP_VISIBLE_DEVICES:No HSA memory faults and no ballot/reduction nondeterminism in any group; the warp-primitive, reduction, ballot, and select paths are correct on wave32.
Note: building the full test suite under ROCm 7.2.1 also requires the companion fixes in #10 (that toolchain marks
hipError_tnodiscard, which breaks several test executables under-Werror); on the ROCm this branch targets the tests compile without it.This work was authored with the assistance of Claude, an AI assistant.