Make permute deterministic given a seed

cpp/include/raft/random/detail/make_regression.cuh

@@ -240,23 +240,33 @@ void make_regression_caller(raft::resources const& handle,
   }
 
   if (shuffle) {
+    // creates shared generator. pulls 2 random numbers from
+    // internal sequence to generate a and b, remembers it gave out those numbers
+    std::mt19937_64 gen(seed);
+
     rmm::device_uvector<DataT> tmp_out(n_rows * n_cols, stream);
     rmm::device_uvector<IdxT> perms_samples(n_rows, stream);
     rmm::device_uvector<IdxT> perms_features(n_cols, stream);
 
     constexpr IdxT Nthreads = 256;
 
     // Shuffle the samples from out to tmp_out
-    raft::random::permute<DataT, IdxT, IdxT>(
-      perms_samples.data(), tmp_out.data(), out, n_cols, n_rows, true, stream);
+    raft::random::detail::permute<DataT, IdxT, IdxT>(perms_samples.data(),
+                                                     tmp_out.data(),
+                                                     out,
+                                                     n_cols,
+                                                     n_rows,
+                                                     true,
+                                                     stream,
+                                                     gen);  // now passes generator rather than seed
     IdxT nblks_rows = raft::ceildiv<IdxT>(n_rows, Nthreads);
     _gather2d_kernel<<<nblks_rows, Nthreads, 0, stream>>>(
       values, _values, perms_samples.data(), n_rows, n_targets);
     RAFT_CUDA_TRY(cudaPeekAtLastError());
 
     // Shuffle the features from tmp_out to out
-    raft::random::permute<DataT, IdxT, IdxT>(
-      perms_features.data(), out, tmp_out.data(), n_rows, n_cols, false, stream);
+    raft::random::detail::permute<DataT, IdxT, IdxT>(
+      perms_features.data(), out, tmp_out.data(), n_rows, n_cols, false, stream, gen);
 
     // Shuffle the coefficients accordingly
     if (coef != nullptr) {

cpp/include/raft/random/detail/permute.cuh

@@ -13,6 +13,7 @@
 #include <cooperative_groups.h>
 
 #include <memory>
+#include <random>
 
 namespace raft {
 namespace random {
@@ -120,15 +121,20 @@ void permute(IntType* perms,
              IntType D,
              IntType N,
              bool rowMajor,
-             cudaStream_t stream)
+             cudaStream_t stream,
+             std::mt19937_64& gen)  // switched from uint64_t seed
 {
   auto nblks = raft::ceildiv(N, (IntType)TPB);
 
-  // always keep 'a' to be coprime to N
-  IdxType a = rand() % N;
+  // maps gen's raw bit stream to a uniform integer
+  std::uniform_int_distribution<IdxType> dist(0, N - 1);
+
+  // a must be coprime to N, otherwise (a * tid) % N does not hit every row exactly once
+  IdxType a = dist(gen);
   while (raft::gcd(a, N) != 1)
     a = (a + 1) % N;
-  IdxType b = rand() % N;
+  // additive offset
+  IdxType b = dist(gen);
 
   if (rowMajor) {
     permute_impl_t<Type,

cpp/include/raft/random/permute.cuh

@@ -16,6 +16,7 @@
 #include <raft/core/resources.hpp>
 
 #include <optional>
+#include <random>
 #include <type_traits>
 
 namespace raft {
@@ -90,7 +91,8 @@ template <typename InputOutputValueType, typename IntType, typename IdxType, typ
 void permute(raft::resources const& handle,
              raft::device_matrix_view<const InputOutputValueType, IdxType, Layout> in,
              std::optional<raft::device_vector_view<IntType, IdxType>> permsOut,
-             std::optional<raft::device_matrix_view<InputOutputValueType, IdxType, Layout>> out)
+             std::optional<raft::device_matrix_view<InputOutputValueType, IdxType, Layout>> out,
+             std::optional<uint64_t> seed = std::nullopt)
 {
   static_assert(std::is_integral_v<IntType>,
                 "permute: The type of each element "
@@ -120,13 +122,17 @@ void permute(raft::resources const& handle,
   if (permsOut_ptr != nullptr || out_ptr != nullptr) {
     const IdxType N = in.extent(0);
     const IdxType D = in.extent(1);
-    detail::permute<InputOutputValueType, IntType, IdxType>(permsOut_ptr,
-                                                            out_ptr,
-                                                            in.data_handle(),
-                                                            D,
-                                                            N,
-                                                            is_row_major,
-                                                            resource::get_cuda_stream(handle));
+    std::mt19937_64 gen(seed.has_value() ? *seed
+                                         : rand());  // use the given seed, else a random one
+    detail::permute<InputOutputValueType, IntType, IdxType>(
+      permsOut_ptr,
+      out_ptr,
+      in.data_handle(),
+      D,
+      N,
+      is_row_major,
+      resource::get_cuda_stream(handle),
+      gen);  // switched from seed for deterministic swap
   }
 }
 
@@ -190,9 +196,12 @@ void permute(IntType* perms,
              IntType D,
              IntType N,
              bool rowMajor,
-             cudaStream_t stream)
+             cudaStream_t stream,
+             std::optional<uint64_t> seed = std::nullopt)
 {
-  detail::permute<Type, IntType, IdxType, TPB>(perms, out, in, D, N, rowMajor, stream);
+  std::mt19937_64 gen(seed.has_value() ? *seed : rand());
+  // generator build and passed
+  detail::permute<Type, IntType, IdxType, TPB>(perms, out, in, D, N, rowMajor, stream, gen);
 }
 
 };  // namespace random

cpp/tests/random/permute.cu

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2018-2024, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2018-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -338,5 +338,76 @@ TEST_P(PermMdspanTestD, Result)
 }
 INSTANTIATE_TEST_CASE_P(PermMdspanTests, PermMdspanTestD, ::testing::ValuesIn(inputsd));
 
+// Determinism test for the seeded `permute` overload (cuML issue #7871:
+// make_regression was not reproducible with shuffle=True because permute
+// drew its affine coefficients from the unseeded global rand()). The same
+// seed must yield identical permutation indices across calls.
+template <typename T>
+class PermSeedTest : public ::testing::TestWithParam<PermInputs<T>> {
+ protected:
+  PermSeedTest()
+    : in(0, resource::get_cuda_stream(handle)),
+      out(0, resource::get_cuda_stream(handle)),
+      perms1(0, resource::get_cuda_stream(handle)),
+      perms2(0, resource::get_cuda_stream(handle))
+  {
+  }
+
+  void SetUp() override
+  {
+    auto stream = resource::get_cuda_stream(handle);
+    params      = ::testing::TestWithParam<PermInputs<T>>::GetParam();
+    int N       = params.N;
+    int D       = params.D;
+    int len     = N * D;
+    raft::random::RngState r(params.seed);
+    in.resize(len, stream);
+    out.resize(len, stream);
+    perms1.resize(N, stream);
+    perms2.resize(N, stream);
+    uniform(handle, r, in.data(), len, T(-1.0), T(1.0));
+
+    // Same seed twice -> the two permutations must be identical.
+    permute<T, int, int>(
+      perms1.data(), out.data(), in.data(), D, N, params.rowMajor, stream, params.seed);
+    permute<T, int, int>(
+      perms2.data(), out.data(), in.data(), D, N, params.rowMajor, stream, params.seed);
+    resource::sync_stream(handle);
+  }
+
+ protected:
+  raft::resources handle;
+  PermInputs<T> params;
+  rmm::device_uvector<T> in, out;
+  rmm::device_uvector<int> perms1, perms2;
+};
+
+using PermSeedTestF = PermSeedTest<float>;
+TEST_P(PermSeedTestF, SameSeedIsDeterministic)
+{
+  auto stream = resource::get_cuda_stream(handle);
+  int N       = params.N;
+  std::vector<int> h1(N), h2(N);
+  raft::update_host(h1.data(), perms1.data(), N, stream);
+  raft::update_host(h2.data(), perms2.data(), N, stream);
+  RAFT_CUDA_TRY(cudaStreamSynchronize(stream));
+  ASSERT_EQ(h1, h2);
+
+  // Different seeds should produce different permutations. Only checked for
+  // larger N, where the space of affine permutations is large enough that a
+  // collision is astronomically unlikely.
+  if (N >= 1024) {
+    rmm::device_uvector<int> perms3(N, stream);
+    permute<float, int, int>(
+      perms3.data(), out.data(), in.data(), params.D, N, params.rowMajor, stream, params.seed + 1);
+    resource::sync_stream(handle);
+    std::vector<int> h3(N);
+    raft::update_host(h3.data(), perms3.data(), N, stream);
+    RAFT_CUDA_TRY(cudaStreamSynchronize(stream));
+    ASSERT_NE(h1, h3);
+  }
+}
+INSTANTIATE_TEST_CASE_P(PermSeedTests, PermSeedTestF, ::testing::ValuesIn(inputsf));
+
 }  // end namespace random
 }  // end namespace raft