Fix wave-size assumption in EvaluateSplitsKernel (RDNA + mixed-arch HIP builds)

[hrylx]

Fix wave-size assumption in EvaluateSplitsKernel (RDNA + mixed-arch HIP builds)

amd_xgboost==3.1.1 segfaults on device="cuda" on every RDNA-class GPU.
Reproduced on RX 9070 XT (gfx1201) under ROCm 7.2.2 in rocm/dev-ubuntu-24.04.
There are two separate problems.

1. The wheel's fatbin is missing RDNA code objects

$ objcopy -O binary --only-section=.hip_fatbin libxgboost.so /tmp/fat.bin
$ strings /tmp/fat.bin | grep amdhsa | sort -u
hipv4-amdgcn-amd-amdhsa--gfx942
amdgcn-amd-amdhsa--gfx942

Build-pipeline issue. Adding gfx1100/gfx1200/gfx1201 to the wheel's
--offload-arch set fixes it. This PR's source change is what makes that
meaningful, so they need to land together.

2. The source assumes wave64 wherever HIP is enabled

src/tree/gpu_hist/evaluate_splits.cu:

#if defined(XGBOOST_USE_HIP)
uint32_t constexpr kBlockThreads = 64;  // AMD wavefront size
#else
uint32_t constexpr kBlockThreads = 32;  // NVIDIA warp size
#endif

That's only true for CDNA (gfx9). RDNA (gfx10/11/12) is wave32, so a 64-thread
block on RDNA hardware contains two wavefronts. EvaluateSplitAgent then breaks:
cub::WarpReduce reduces within one wavefront, __shfl_sync(kFullMask, x, 0)
broadcasts within one wavefront, and the two wavefronts race into the same
__shared__ best_split. No memory access is OOB, so the kernel does not crash;
it just picks the wrong split. Trees are silently wrong. On
make_classification(10k, 50) accuracy drops from 0.96 to 0.85. The
static_assert(kBlockSize == WARP_SIZE) that would catch this is
#if defined(XGBOOST_USE_CUDA)-guarded.

Prior history

Branch release_2.0.3-rocm_navi has eight fix navi wave32 commits by @hliuca
between 2025-03 and 2025-04 (tip ee5234f2c). They diagnose the same problem and
dispatch via a runtime hipDeviceGetAttribute query. git merge-base ee5234f2c release/3.1.1 is empty: those commits never made it forward, so 3.1.1 ships with
the original wave64-only assumption.

Fix

Same approach as the navi patches: runtime-query the wave size and dispatch to
the matching EvaluateSplitsKernel<32> or EvaluateSplitsKernel<64>. The
duplicated launch block from the navi version is folded into a single
integral_constant-driven lambda, the lazy-init global is gone, and the
printf + exit(-1) is replaced with LOG(FATAL). One source file, no CMake
changes.

Both kernel templates are emitted, so a single binary covers wave32 and wave64.
Multi-arch builds work as expected:

$ cmake -DCMAKE_HIP_ARCHITECTURES="gfx942;gfx1201" ...
$ strings /tmp/fat.bin | grep amdhsa | sort -u
amdgcn-amd-amdhsa--gfx1201
amdgcn-amd-amdhsa--gfx942
hipv4-amdgcn-amd-amdhsa--gfx1201
hipv4-amdgcn-amd-amdhsa--gfx942

That should let one amd_xgboost wheel cover both families.

Validation (RX 9070 XT, ROCm 7.2.2)

Single-arch (gfx1201) and mixed-arch (gfx942 + gfx1201) builds give the same
results on RDNA hardware. 6/7 of binary, multiclass, regression, deep-tree, and
500-round configurations match CPU train-loss within 5%. The seventh, a wide
500-feature problem, has 13.7% train-loss diff on a single seed; over 8 seeds
with held-out evaluation the test AUC differs from CPU by 0.04%, which is
FP-reduction-order variance.

CDNA (gfx9) is unchanged. The runtime dispatch picks kBlockThreads = 64 on
gfx9 hardware, identical to the original code path.

CI

The CDNA matrix cannot catch this class of bug because block_size == wavefront
always holds on gfx9. Adding any gfx10/11/12 target would have caught both this
issue and the earlier 2.0 to 3.1 forward-port loss.

[edrich]

Thanks for the PR. I experienced the same behavior on my RX 9070 XT.