Make permute deterministic given a seed

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

@@ -248,15 +248,22 @@ void make_regression_caller(raft::resources const& handle,
 
     // 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);
+      perms_samples.data(), tmp_out.data(), out, n_cols, n_rows, true, stream, 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);
+      perms_features.data(),
+      out,
+      tmp_out.data(),
+      n_rows,
+      n_cols,
+      false,
+      stream,
+      seed + 1);  // different derived seed for feature shuffle keeps sample and feature independent
 
     // 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,22 @@ void permute(IntType* perms,
              IntType D,
              IntType N,
              bool rowMajor,
-             cudaStream_t stream)
+             cudaStream_t stream,
+             uint64_t seed)
 {
   auto nblks = raft::ceildiv(N, (IntType)TPB);
 
-  // always keep 'a' to be coprime to N
-  IdxType a = rand() % N;
+  // gen seeded once, deterministically, from seed param
+  std::mt19937_64 gen(seed);
+  // 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

@@ -90,7 +90,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,
+             uint64_t seed = 0ULL)
 {
   static_assert(std::is_integral_v<IntType>,
                 "permute: The type of each element "
@@ -126,7 +127,8 @@ void permute(raft::resources const& handle,
                                                             D,
                                                             N,
                                                             is_row_major,
-                                                            resource::get_cuda_stream(handle));
+                                                            resource::get_cuda_stream(handle),
+                                                            seed);
   }
 }
 
@@ -190,9 +192,10 @@ void permute(IntType* perms,
              IntType D,
              IntType N,
              bool rowMajor,
-             cudaStream_t stream)
+             cudaStream_t stream,
+             uint64_t seed = 0ULL)  // default seed
 {
-  detail::permute<Type, IntType, IdxType, TPB>(perms, out, in, D, N, rowMajor, stream);
+  detail::permute<Type, IntType, IdxType, TPB>(perms, out, in, D, N, rowMajor, stream, seed);
 }
 
 };  // 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