From d4a1e264e6fe859059ffaf77013ae9e2fc516b50 Mon Sep 17 00:00:00 2001 From: hrylx <57755200+hrylx@users.noreply.github.com> Date: Tue, 28 Apr 2026 18:20:05 +0200 Subject: [PATCH] Fix wave-size assumption in EvaluateSplitsKernel release/3.1.1 hardcodes kBlockThreads = 64 for HIP, treating it as synonymous with wave64. RDNA (gfx10/11/12) is wave32, so on RDNA hardware a 64-thread block contains two wavefronts and the warp-scoped reduce/broadcast in EvaluateSplitAgent runs twice independently. The result is silently wrong splits (no crash; the static_assert that would catch this is CUDA-only). On make_classification(10k, 50) accuracy drops from 0.96 to 0.85. Query the wave size at launch via hipDeviceGetAttribute and dispatch to EvaluateSplitsKernel<32> or <64>. Both templates are emitted, so one binary covers wave32 and wave64 including multi-arch fatbins. Same approach as the navi PRs (#6-#11) on the v2.0.x line, folded into a single launch via an integral_constant-driven lambda. Those PRs were never ported forward; git merge-base with release/3.1.1 is empty, so 3.1.1 ships with the original wave64-only assumption. --- src/tree/gpu_hist/evaluate_splits.cu | 45 +++++++++++++++++++++++----- 1 file changed, 37 insertions(+), 8 deletions(-) diff --git a/src/tree/gpu_hist/evaluate_splits.cu b/src/tree/gpu_hist/evaluate_splits.cu index 414b061f1505..dc99097edbde 100644 --- a/src/tree/gpu_hist/evaluate_splits.cu +++ b/src/tree/gpu_hist/evaluate_splits.cu @@ -388,17 +388,46 @@ void GPUHistEvaluator::LaunchEvaluateSplits( dh::TemporaryArray feature_best_splits( combined_num_features, DeviceSplitCandidate()); - // One block for each feature + // One block for each feature. + // + // The kernel uses cub::WarpReduce to compute per-feature reductions and __shfl_sync to + // broadcast the result across threads in the block. Both are warp-scoped, so the kernel + // only produces a correct global argmax when the block is exactly one wavefront wide -- + // otherwise the second wavefront in the block does its own private reduction and races + // into shared memory, silently producing the wrong split (no crash, just wrong models). + // + // Wavefront size is therefore the only valid block size: + // NVIDIA -> warp 32 + // AMD CDNA / GCN (gfx9)-> wave 64 + // AMD RDNA (gfx10/11/12)-> wave 32 + // + // On HIP we query the device at launch time and dispatch to the matching template + // instantiation, which keeps a single binary working across both AMD wavefront sizes + // (e.g. fatbins shipped with multiple --offload-arch entries). + auto launch = [&](auto block_threads_const) { + constexpr uint32_t kBlockThreads = decltype(block_threads_const)::value; + dh::LaunchKernel{static_cast(combined_num_features), kBlockThreads, 0, // NOLINT + ctx->CUDACtx()->Stream()}( + EvaluateSplitsKernel, max_active_features, d_inputs, shared_inputs, + this->SortedIdx(d_inputs.size(), shared_inputs.feature_values.size()), evaluator, + dh::ToSpan(feature_best_splits)); + }; + #if defined(XGBOOST_USE_HIP) -uint32_t constexpr kBlockThreads = 64; // AMD wavefront size + int dev = 0, warp_size = 0; + dh::safe_cuda(hipGetDevice(&dev)); + dh::safe_cuda(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, dev)); + if (warp_size == 32) { + launch(std::integral_constant{}); + } else if (warp_size == 64) { + launch(std::integral_constant{}); + } else { + LOG(FATAL) << "xgboost: unsupported AMD wavefront size " << warp_size + << " (expected 32 or 64)"; + } #else -uint32_t constexpr kBlockThreads = 32; // NVIDIA warp size + launch(std::integral_constant{}); #endif - dh::LaunchKernel{static_cast(combined_num_features), kBlockThreads, 0, // NOLINT - ctx->CUDACtx()->Stream()}( - EvaluateSplitsKernel, max_active_features, d_inputs, shared_inputs, - this->SortedIdx(d_inputs.size(), shared_inputs.feature_values.size()), evaluator, - dh::ToSpan(feature_best_splits)); // Reduce to get best candidate for left and right child over all features auto reduce_offset = dh::MakeTransformIterator(