Use philox for faster sampling on GPU. (#12223)

Author: trivialfis

src/common/random.cu

@@ -8,6 +8,7 @@
 #include "algorithm.cuh"     // for ArgSort
 #include "cuda_context.cuh"  // for CUDAContext
 #include "device_helpers.cuh"
+#include "random.cuh"  // for DefaultRng, UniformRealDistribution
 #include "random.h"
 #include "xgboost/base.h"                // for bst_feature_t
 #include "xgboost/context.h"             // for Context
@@ -30,10 +31,9 @@ void WeightedSamplingWithoutReplacement(Context const *ctx, common::Span<bst_fea
   constexpr auto kEps = kRtEps;  // avoid CUDA compilation error
   thrust::for_each_n(cuctx->CTP(), thrust::make_counting_iterator(0ul), array.size(),
                      [=] XGBOOST_DEVICE(std::size_t i) {
-                       thrust::default_random_engine rng;
-                       rng.seed(seed);
+                       DefaultRng rng{seed};
                        rng.discard(i);
-                       thrust::uniform_real_distribution<float> dist;
+                       UniformRealDistribution<float> dist;
 
                        auto w = std::max(weights[i], kEps);
                        auto u = dist(rng);

src/common/random.cuh

@@ -0,0 +1,27 @@
+/**
+ * Copyright 2026, XGBoost Contributors
+ */
+#pragma once
+
+#if __has_include(<cuda/std/random>)
+#define xgboost_HAS_CUDA_STD_RANDOM 1
+#include <cuda/std/random>
+#else
+#include <thrust/random.h>
+#endif
+
+namespace xgboost::common::cuda_impl {
+#if defined(xgboost_HAS_CUDA_STD_RANDOM)
+using DefaultRng = cuda::std::philox4x64;
+template <typename T>
+using UniformRealDistribution = cuda::std::uniform_real_distribution<T>;
+#else
+using DefaultRng = thrust::default_random_engine;
+template <typename T>
+using UniformRealDistribution = thrust::uniform_real_distribution<T>;
+#endif
+}  // namespace xgboost::common::cuda_impl
+
+#if defined(xgboost_HAS_CUDA_STD_RANDOM)
+#undef xgboost_HAS_CUDA_STD_RANDOM
+#endif

src/tree/gpu_hist/sampler.cu

@@ -25,8 +25,8 @@
 
 #include "../../common/cuda_context.cuh"    // for CUDAContext
 #include "../../common/device_helpers.cuh"  // for MakeTransformIterator
-#include "../../common/random.h"
-#include "../hist/sampler.h"  // for kDefaultMvsLambda
+#include "../../common/random.cuh"          // for DefaultRng, UniformRealDistribution
+#include "../hist/sampler.h"                // for kDefaultMvsLambda
 #include "../param.h"
 #include "quantiser.cuh"  // for GradientQuantiser
 #include "sampler.cuh"
@@ -38,8 +38,8 @@ class RandomWeight {
   explicit RandomWeight(std::size_t seed) : seed_(seed) {}
 
   XGBOOST_DEVICE float operator()(std::size_t i) const {
-    thrust::default_random_engine rng(seed_);
-    thrust::uniform_real_distribution<float> dist;
+    common::cuda_impl::DefaultRng rng{seed_};
+    common::cuda_impl::UniformRealDistribution<float> dist;
     rng.discard(i);
     return dist(rng);
   }

tests/cpp/tree/test_gpu_hist.cu

@@ -75,7 +75,7 @@ TEST(GpuHist, UniformSampling) {
   auto preds_h = preds.ConstHostVector();
   auto preds_sampling_h = preds_sampling.ConstHostVector();
   for (size_t i = 0; i < kRows; i++) {
-    EXPECT_NEAR(preds_h[i], preds_sampling_h[i], 1e-8);
+    EXPECT_NEAR(preds_h[i], preds_sampling_h[i], 1e-3);
   }
 }