From e2b3ce7e7f428115aa9060e3c8aa9b02a885a981 Mon Sep 17 00:00:00 2001 From: says1117 Date: Mon, 6 Jul 2026 16:48:35 -0400 Subject: [PATCH 1/4] Make permute deterministic given a seed permute() drew its affine permutation coefficients from the unseeded global rand(), so make_regression produced different output on repeated calls with the same random_state (cuML issue #7871). Derive the coefficients from a local std::mt19937_64 seeded per call, thread the seed through make_regression's row and feature shuffles, and add a determinism test. Public wrappers default the seed so existing callers are unaffected. --- .../raft/random/detail/make_regression.cuh | 11 ++- cpp/include/raft/random/detail/permute.cuh | 16 +++- cpp/include/raft/random/permute.cuh | 11 ++- cpp/tests/random/permute.cu | 73 ++++++++++++++++++- 4 files changed, 100 insertions(+), 11 deletions(-) diff --git a/cpp/include/raft/random/detail/make_regression.cuh b/cpp/include/raft/random/detail/make_regression.cuh index 773eae7b39..0a75490cc1 100644 --- a/cpp/include/raft/random/detail/make_regression.cuh +++ b/cpp/include/raft/random/detail/make_regression.cuh @@ -248,7 +248,7 @@ void make_regression_caller(raft::resources const& handle, // Shuffle the samples from out to tmp_out raft::random::permute( - 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(n_rows, Nthreads); _gather2d_kernel<<>>( values, _values, perms_samples.data(), n_rows, n_targets); @@ -256,7 +256,14 @@ void make_regression_caller(raft::resources const& handle, // Shuffle the features from tmp_out to out raft::random::permute( - 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) { diff --git a/cpp/include/raft/random/detail/permute.cuh b/cpp/include/raft/random/detail/permute.cuh index d991d6fc50..da4a228562 100644 --- a/cpp/include/raft/random/detail/permute.cuh +++ b/cpp/include/raft/random/detail/permute.cuh @@ -13,6 +13,7 @@ #include #include +#include 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 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 in, std::optional> permsOut, - std::optional> out) + std::optional> out, + uint64_t seed = 0ULL) { static_assert(std::is_integral_v, "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(perms, out, in, D, N, rowMajor, stream); + detail::permute(perms, out, in, D, N, rowMajor, stream, seed); } }; // namespace random diff --git a/cpp/tests/random/permute.cu b/cpp/tests/random/permute.cu index 44b5b9c66c..3b1cef59d9 100644 --- a/cpp/tests/random/permute.cu +++ b/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 +class PermSeedTest : public ::testing::TestWithParam> { + 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>::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( + perms1.data(), out.data(), in.data(), D, N, params.rowMajor, stream, params.seed); + permute( + perms2.data(), out.data(), in.data(), D, N, params.rowMajor, stream, params.seed); + resource::sync_stream(handle); + } + + protected: + raft::resources handle; + PermInputs params; + rmm::device_uvector in, out; + rmm::device_uvector perms1, perms2; +}; + +using PermSeedTestF = PermSeedTest; +TEST_P(PermSeedTestF, SameSeedIsDeterministic) +{ + auto stream = resource::get_cuda_stream(handle); + int N = params.N; + std::vector 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 perms3(N, stream); + permute( + perms3.data(), out.data(), in.data(), params.D, N, params.rowMajor, stream, params.seed + 1); + resource::sync_stream(handle); + std::vector 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 From 5021723aef2bbf47d72ff3b0f1dff109a2749d6d Mon Sep 17 00:00:00 2001 From: says1117 Date: Thu, 9 Jul 2026 05:19:33 -0400 Subject: [PATCH 2/4] Address review: single generator by ref, optional seed with rand() fallback. Corrected function parameters from seed to gen. --- .../raft/random/detail/make_regression.cuh | 25 +++++++++-------- cpp/include/raft/random/detail/permute.cuh | 4 +-- cpp/include/raft/random/permute.cuh | 28 +++++++++++-------- 3 files changed, 32 insertions(+), 25 deletions(-) diff --git a/cpp/include/raft/random/detail/make_regression.cuh b/cpp/include/raft/random/detail/make_regression.cuh index 0a75490cc1..72bee7465b 100644 --- a/cpp/include/raft/random/detail/make_regression.cuh +++ b/cpp/include/raft/random/detail/make_regression.cuh @@ -240,6 +240,10 @@ 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 tmp_out(n_rows * n_cols, stream); rmm::device_uvector perms_samples(n_rows, stream); rmm::device_uvector perms_features(n_cols, stream); @@ -247,23 +251,22 @@ void make_regression_caller(raft::resources const& handle, constexpr IdxT Nthreads = 256; // Shuffle the samples from out to tmp_out - raft::random::permute( - perms_samples.data(), tmp_out.data(), out, n_cols, n_rows, true, stream, seed); + raft::random::detail::permute(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(n_rows, Nthreads); _gather2d_kernel<<>>( values, _values, perms_samples.data(), n_rows, n_targets); RAFT_CUDA_TRY(cudaPeekAtLastError()); // Shuffle the features from tmp_out to out - raft::random::permute( - 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 + raft::random::detail::permute( + perms_features.data(), out, tmp_out.data(), n_rows, n_cols, false, stream, gen); // Shuffle the coefficients accordingly if (coef != nullptr) { diff --git a/cpp/include/raft/random/detail/permute.cuh b/cpp/include/raft/random/detail/permute.cuh index da4a228562..517ad9806f 100644 --- a/cpp/include/raft/random/detail/permute.cuh +++ b/cpp/include/raft/random/detail/permute.cuh @@ -122,12 +122,10 @@ void permute(IntType* perms, IntType N, bool rowMajor, cudaStream_t stream, - uint64_t seed) + std::mt19937_64& gen) // switched from uint64_t seed { auto nblks = raft::ceildiv(N, (IntType)TPB); - // 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 dist(0, N - 1); diff --git a/cpp/include/raft/random/permute.cuh b/cpp/include/raft/random/permute.cuh index 443551caad..dcd46096ab 100644 --- a/cpp/include/raft/random/permute.cuh +++ b/cpp/include/raft/random/permute.cuh @@ -16,6 +16,7 @@ #include #include +#include #include namespace raft { @@ -91,7 +92,7 @@ void permute(raft::resources const& handle, raft::device_matrix_view in, std::optional> permsOut, std::optional> out, - uint64_t seed = 0ULL) + std::optional seed = std::nullopt) { static_assert(std::is_integral_v, "permute: The type of each element " @@ -121,14 +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(permsOut_ptr, - out_ptr, - in.data_handle(), - D, - N, - is_row_major, - resource::get_cuda_stream(handle), - seed); + std::mt19937_64 gen(seed.has_value() ? *seed + : rand()); // use the given seed, else a random one + detail::permute( + permsOut_ptr, + out_ptr, + in.data_handle(), + D, + N, + is_row_major, + resource::get_cuda_stream(handle), + gen); // switched from seed for deterministic swap } } @@ -193,9 +197,11 @@ void permute(IntType* perms, IntType N, bool rowMajor, cudaStream_t stream, - uint64_t seed = 0ULL) // default seed + std::optional seed = std::nullopt) { - detail::permute(perms, out, in, D, N, rowMajor, stream, seed); + std::mt19937_64 gen(seed.has_value() ? *seed : rand()); + // generator build and passed + detail::permute(perms, out, in, D, N, rowMajor, stream, gen); } }; // namespace random From c4796632c83cff929dac1d2ea84cee0477f7737e Mon Sep 17 00:00:00 2001 From: says1117 Date: Thu, 9 Jul 2026 05:59:19 -0400 Subject: [PATCH 3/4] Address CodeRabbit findings: forward seed through convenience overload, add Doxygen, use test_data_type alias --- cpp/include/raft/random/permute.cuh | 15 +++++++++++++-- cpp/tests/random/permute.cu | 10 +++++++--- 2 files changed, 20 insertions(+), 5 deletions(-) diff --git a/cpp/include/raft/random/permute.cuh b/cpp/include/raft/random/permute.cuh index dcd46096ab..b8edf32f08 100644 --- a/cpp/include/raft/random/permute.cuh +++ b/cpp/include/raft/random/permute.cuh @@ -75,6 +75,9 @@ using perms_out_view_t = typename perms_out_view in, PermsOutType&& permsOut, - OutType&& out) + OutType&& out, + std::optional seed = std::nullopt) { // If PermsOutType is std::optional> // for some T, then that type T need not be related to any of the @@ -165,7 +173,7 @@ void permute(raft::resources const& handle, std::optional permsOut_arg = std::forward(permsOut); std::optional out_arg = std::forward(out); - permute(handle, in, permsOut_arg, out_arg); + permute(handle, in, permsOut_arg, out_arg, seed); } /** @} */ @@ -188,6 +196,9 @@ void permute(raft::resources const& handle, * @param[in] rowMajor true if the matrices are row major, * false if they are column major * @param[in] stream CUDA stream on which to run + * @param[in] seed If provided, seeds the permutation so that it is + * reproducible; if not provided, falls back to a non-deterministic + * `rand()`-derived seed (matching this function's legacy behavior). */ template void permute(IntType* perms, diff --git a/cpp/tests/random/permute.cu b/cpp/tests/random/permute.cu index 3b1cef59d9..f4111d8ee3 100644 --- a/cpp/tests/random/permute.cu +++ b/cpp/tests/random/permute.cu @@ -344,6 +344,9 @@ INSTANTIATE_TEST_CASE_P(PermMdspanTests, PermMdspanTestD, ::testing::ValuesIn(in // seed must yield identical permutation indices across calls. template class PermSeedTest : public ::testing::TestWithParam> { + public: + using test_data_type = T; + protected: PermSeedTest() : in(0, resource::get_cuda_stream(handle)), @@ -385,8 +388,9 @@ class PermSeedTest : public ::testing::TestWithParam> { using PermSeedTestF = PermSeedTest; TEST_P(PermSeedTestF, SameSeedIsDeterministic) { - auto stream = resource::get_cuda_stream(handle); - int N = params.N; + using test_data_type = PermSeedTestF::test_data_type; + auto stream = resource::get_cuda_stream(handle); + int N = params.N; std::vector h1(N), h2(N); raft::update_host(h1.data(), perms1.data(), N, stream); raft::update_host(h2.data(), perms2.data(), N, stream); @@ -398,7 +402,7 @@ TEST_P(PermSeedTestF, SameSeedIsDeterministic) // collision is astronomically unlikely. if (N >= 1024) { rmm::device_uvector perms3(N, stream); - permute( + permute( perms3.data(), out.data(), in.data(), params.D, N, params.rowMajor, stream, params.seed + 1); resource::sync_stream(handle); std::vector h3(N); From 9f84b8be0d8f6beb62228d81649ac5c2853c1666 Mon Sep 17 00:00:00 2001 From: says1117 Date: Fri, 10 Jul 2026 05:39:58 -0400 Subject: [PATCH 4/4] Address review: thread-safe seed fallback, dedupe generator construction --- .../raft/random/detail/make_regression.cuh | 18 +++---- cpp/include/raft/random/detail/permute.cuh | 4 +- cpp/include/raft/random/permute.cuh | 47 ++++++++++--------- cpp/tests/random/permute.cu | 32 ++++++++++++- 4 files changed, 65 insertions(+), 36 deletions(-) diff --git a/cpp/include/raft/random/detail/make_regression.cuh b/cpp/include/raft/random/detail/make_regression.cuh index 72bee7465b..5e8ff54722 100644 --- a/cpp/include/raft/random/detail/make_regression.cuh +++ b/cpp/include/raft/random/detail/make_regression.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -24,6 +24,7 @@ #include #include +#include namespace raft { namespace random { @@ -240,8 +241,9 @@ 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 + // One generator shared by both shuffles below: each consumes distinct draws + // from its sequence, so the sample and feature permutations stay independent + // while remaining reproducible for a given seed. std::mt19937_64 gen(seed); rmm::device_uvector tmp_out(n_rows * n_cols, stream); @@ -251,14 +253,8 @@ void make_regression_caller(raft::resources const& handle, constexpr IdxT Nthreads = 256; // Shuffle the samples from out to tmp_out - raft::random::detail::permute(perms_samples.data(), - tmp_out.data(), - out, - n_cols, - n_rows, - true, - stream, - gen); // now passes generator rather than seed + raft::random::detail::permute( + perms_samples.data(), tmp_out.data(), out, n_cols, n_rows, true, stream, gen); IdxT nblks_rows = raft::ceildiv(n_rows, Nthreads); _gather2d_kernel<<>>( values, _values, perms_samples.data(), n_rows, n_targets); diff --git a/cpp/include/raft/random/detail/permute.cuh b/cpp/include/raft/random/detail/permute.cuh index 517ad9806f..1874a5a768 100644 --- a/cpp/include/raft/random/detail/permute.cuh +++ b/cpp/include/raft/random/detail/permute.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -122,7 +122,7 @@ void permute(IntType* perms, IntType N, bool rowMajor, cudaStream_t stream, - std::mt19937_64& gen) // switched from uint64_t seed + std::mt19937_64& gen) { auto nblks = raft::ceildiv(N, (IntType)TPB); diff --git a/cpp/include/raft/random/permute.cuh b/cpp/include/raft/random/permute.cuh index b8edf32f08..3d24ee596b 100644 --- a/cpp/include/raft/random/permute.cuh +++ b/cpp/include/raft/random/permute.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -24,6 +24,12 @@ namespace random { namespace permute_impl { +inline std::mt19937_64 make_permute_generator(std::optional seed) +{ + // use the given seed, else a random one + return std::mt19937_64(seed.has_value() ? *seed : std::random_device{}()); +} + template struct perms_out_view {}; @@ -76,8 +82,8 @@ using perms_out_view_t = typename perms_out_view( - permsOut_ptr, - out_ptr, - in.data_handle(), - D, - N, - is_row_major, - resource::get_cuda_stream(handle), - gen); // switched from seed for deterministic swap + const IdxType N = in.extent(0); + const IdxType D = in.extent(1); + std::mt19937_64 gen = permute_impl::make_permute_generator(seed); + detail::permute(permsOut_ptr, + out_ptr, + in.data_handle(), + D, + N, + is_row_major, + resource::get_cuda_stream(handle), + gen); } } @@ -144,8 +148,8 @@ void permute(raft::resources const& handle, * for either or both of `permsOut` and `out`. * * @param[in] seed If provided, seeds the permutation so that it is - * reproducible; if not provided, falls back to a non-deterministic - * `rand()`-derived seed (matching this function's legacy behavior). + * reproducible; if not provided, falls back to a non-deterministic, + * thread-safe `std::random_device`-derived seed. */ template void permute(IntType* perms, @@ -210,8 +214,7 @@ void permute(IntType* perms, cudaStream_t stream, std::optional seed = std::nullopt) { - std::mt19937_64 gen(seed.has_value() ? *seed : rand()); - // generator build and passed + std::mt19937_64 gen = permute_impl::make_permute_generator(seed); detail::permute(perms, out, in, D, N, rowMajor, stream, gen); } diff --git a/cpp/tests/random/permute.cu b/cpp/tests/random/permute.cu index f4111d8ee3..98b57a4d6b 100644 --- a/cpp/tests/random/permute.cu +++ b/cpp/tests/random/permute.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2018-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2018-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -411,7 +411,37 @@ TEST_P(PermSeedTestF, SameSeedIsDeterministic) ASSERT_NE(h1, h3); } } + INSTANTIATE_TEST_CASE_P(PermSeedTests, PermSeedTestF, ::testing::ValuesIn(inputsf)); +using PermSeedTestD = PermSeedTest; +TEST_P(PermSeedTestD, SameSeedIsDeterministic) +{ + using test_data_type = PermSeedTestD::test_data_type; + auto stream = resource::get_cuda_stream(handle); + int N = params.N; + std::vector 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 perms3(N, stream); + permute( + perms3.data(), out.data(), in.data(), params.D, N, params.rowMajor, stream, params.seed + 1); + resource::sync_stream(handle); + std::vector 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, PermSeedTestD, ::testing::ValuesIn(inputsd)); + } // end namespace random } // end namespace raft