From aeeeab03fa1a72348a95033db0dd076526577847 Mon Sep 17 00:00:00 2001 From: Huy Nguyen Date: Sat, 25 Jul 2026 07:19:28 +0000 Subject: [PATCH 1/4] FastIntDiv: fallback to normal division when values exceed 32-bit range during run time. --- cpp/include/raft/util/fast_int_div.cuh | 8 ++-- cpp/tests/util/fast_int_div.cu | 63 ++++++++++++++------------ 2 files changed, 38 insertions(+), 33 deletions(-) diff --git a/cpp/include/raft/util/fast_int_div.cuh b/cpp/include/raft/util/fast_int_div.cuh index d7f385a5d4..bb310a01f5 100644 --- a/cpp/include/raft/util/fast_int_div.cuh +++ b/cpp/include/raft/util/fast_int_div.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -19,9 +19,6 @@ namespace util { * @brief Perform fast integer division and modulo using a known divisor * From Hacker's Delight, Second Edition, Chapter 10 * - * @note 32b signed integer is supported. - * @note 64b signed integers is supported for an input data up to 2^31 - * because gpu-non-native int128 is avoided for performance. * @todo Extend support for signed divisors */ template @@ -102,6 +99,7 @@ struct FastIntDiv { template HDI IntT operator/(IntT n, const FastIntDiv& divisor) { + if ((int64_t(divisor.d) >> 32) != 0 || (int64_t(n) >> 32) != 0) return n / divisor.d; if (divisor.d == 1) return n; IntT ret = (int64_t(divisor.m) * int64_t(n)) >> divisor.p; if (n < 0) ++ret; @@ -118,6 +116,8 @@ HDI IntT operator/(IntT n, const FastIntDiv& divisor) template HDI IntT operator%(IntT n, const FastIntDiv& divisor) { + // TODO (huy) measure overhead + if ((int64_t(divisor.d) >> 32) != 0 || (int64_t(n) >> 32) != 0) return n % divisor.d; IntT quotient = n / divisor; IntT remainder = n - quotient * divisor.d; return remainder; diff --git a/cpp/tests/util/fast_int_div.cu b/cpp/tests/util/fast_int_div.cu index eb391c1dbb..4f9af1af7e 100644 --- a/cpp/tests/util/fast_int_div.cu +++ b/cpp/tests/util/fast_int_div.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -13,40 +13,45 @@ namespace raft::util { -constexpr int64_t kInt32Max = std::numeric_limits::max(); - -template -class FastIntDivTest : public ::testing::Test { - protected: - void CompareWithNativeDivision() - { - std::vector magnitudes{0, 1, 2, 3, 7, 13, 255, 12345, (1 << 20), kInt32Max}; - std::vector divisors{1, 2, 4, 7, 16, 31, 63, 128, 1000, (1 << 15), kInt32Max}; - - for (IntT d : divisors) { - FastIntDiv fid(d); - for (IntT mag : magnitudes) { - for (IntT n : {mag, -mag}) { - ASSERT_EQ(n / fid, n / d) << "operator/ mismatch for numerator=" << n << " divisor=" << d; - ASSERT_EQ(n % fid, n % d) << "operator% mismatch for numerator=" << n << " divisor=" << d; - } +constexpr auto kInt32Max = std::numeric_limits::max(); +constexpr auto kInt64Max = std::numeric_limits::max(); + +TEST(FastIntDivTest, UnsupportedNegativeDenumerator) +{ + ASSERT_THROW(FastIntDiv{-42}, raft::exception); + ASSERT_THROW(FastIntDiv{0}, raft::exception); +} + +TEST(FastIntDivTest, CompareWithNativeDivisionInt32) +{ + std::vector magnitudes{0, 1, 2, 3, 7, 13, 255, 12345, (1 << 20), kInt32Max}; + std::vector divisors{1, 2, 4, 7, 16, 31, 63, 128, 1000, (1 << 15), kInt32Max}; + + for (int32_t d : divisors) { + FastIntDiv fid(d); + for (int32_t mag : magnitudes) { + for (int32_t n : {mag, -mag}) { + ASSERT_EQ(n / fid, n / d) << "operator/ mismatch for numerator=" << n << " divisor=" << d; + ASSERT_EQ(n % fid, n % d) << "operator% mismatch for numerator=" << n << " divisor=" << d; } } } -}; - -using FastIntDivTypes = ::testing::Types; -TYPED_TEST_CASE(FastIntDivTest, FastIntDivTypes); - -TYPED_TEST(FastIntDivTest, CompareWithNativeDivision) { this->CompareWithNativeDivision(); } +} -TEST(FastIntDiv, Int64NumeratorPastInt32Boundary) +TEST(FastIntDivTest, CompareWithNativeDivisionInt64) { - for (int64_t d : {129, 772, 1000}) { + std::vector magnitudes{ + 0, 1013, 10007, int64_t(kInt32Max) + 3, (int64_t(1) << 33), (int64_t(3) << 40), kInt64Max}; + std::vector divisors{ + 1, 129, 772, 1000, (int64_t(1) << 31), int64_t(kInt32Max) + 1, int64_t(3) << 40, kInt64Max}; + + for (int64_t d : divisors) { FastIntDiv fid(d); - for (int64_t n : {1LL << 31, 2147704000LL, 3LL << 30}) { // in [2^31, 2^32) - ASSERT_EQ(n / fid, n / d) << "operator/ mismatch for numerator=" << n << " divisor=" << d; - ASSERT_EQ(n % fid, n % d) << "operator% mismatch for numerator=" << n << " divisor=" << d; + for (int64_t mag : magnitudes) { + for (int64_t n : {mag, -mag}) { + ASSERT_EQ(n / fid, n / d) << "operator/ mismatch for numerator=" << n << " divisor=" << d; + ASSERT_EQ(n % fid, n % d) << "operator% mismatch for numerator=" << n << " divisor=" << d; + } } } } From 22845bbeb9c807f6af42e716e6874047989e4194 Mon Sep 17 00:00:00 2001 From: Huy Nguyen Date: Mon, 27 Jul 2026 10:14:43 +0000 Subject: [PATCH 2/4] Support cross-int divisions --- cpp/include/raft/util/fast_int_div.cuh | 44 ++++++++----- cpp/tests/util/fast_int_div.cu | 89 +++++++++++++++++++++----- 2 files changed, 101 insertions(+), 32 deletions(-) diff --git a/cpp/include/raft/util/fast_int_div.cuh b/cpp/include/raft/util/fast_int_div.cuh index bb310a01f5..917e1f2fc5 100644 --- a/cpp/include/raft/util/fast_int_div.cuh +++ b/cpp/include/raft/util/fast_int_div.cuh @@ -15,6 +15,9 @@ namespace raft { namespace util { +constexpr auto kInt32Min = std::numeric_limits::min(); +constexpr auto kInt32Max = std::numeric_limits::max(); + /** * @brief Perform fast integer division and modulo using a known divisor * From Hacker's Delight, Second Edition, Chapter 10 @@ -48,12 +51,16 @@ struct FastIntDiv { * @brief host and device ctor's * @param other source object to be copied from */ - HDI FastIntDiv(const FastIntDiv& other) : d(other.d), m(other.m), p(other.p) {} + HDI FastIntDiv(const FastIntDiv& other) + : d(other.d), m(other.m), p(other.p), fallback(other.fallback) + { + } HDI FastIntDiv& operator=(const FastIntDiv& other) { - d = other.d; - m = other.m; - p = other.p; + d = other.d; + m = other.m; + p = other.p; + fallback = other.fallback; return *this; } /** @} */ @@ -64,6 +71,8 @@ struct FastIntDiv { UIntT m; /** the term 'p' as found in the reference chapter */ int p; + /** Flag for falling back to canonical division on unsupported divisor's ranges */ + bool fallback = false; private: void computeScalars() @@ -76,6 +85,9 @@ struct FastIntDiv { ASSERT(false, "FastIntDiv: division by negative numbers not supported!"); } else if (d == 0) { ASSERT(false, "FastIntDiv: got division by zero!"); + } else if (int64_t(d) > kInt32Max) { + fallback = true; + return; } int64_t nc = ((1LL << 31) / d) * d - 1; p = 31; @@ -96,12 +108,15 @@ struct FastIntDiv { * @param divisor the denominator * @return the quotient */ -template -HDI IntT operator/(IntT n, const FastIntDiv& divisor) +template +HDI std::common_type_t operator/(NumIntT n, const FastIntDiv& divisor) { - if ((int64_t(divisor.d) >> 32) != 0 || (int64_t(n) >> 32) != 0) return n / divisor.d; - if (divisor.d == 1) return n; - IntT ret = (int64_t(divisor.m) * int64_t(n)) >> divisor.p; + using CommonIntT = std::common_type_t; + if (divisor.d == 1) return CommonIntT(n); + if (divisor.fallback || n < kInt32Min || n > kInt32Max) { + return CommonIntT(n) / CommonIntT(divisor.d); + } + CommonIntT ret = (int64_t(divisor.m) * int64_t(n)) >> divisor.p; if (n < 0) ++ret; return ret; } @@ -113,13 +128,12 @@ HDI IntT operator/(IntT n, const FastIntDiv& divisor) * @param divisor the denominator * @return the remainder */ -template -HDI IntT operator%(IntT n, const FastIntDiv& divisor) +template +HDI std::common_type_t operator%(NumIntT n, const FastIntDiv& divisor) { - // TODO (huy) measure overhead - if ((int64_t(divisor.d) >> 32) != 0 || (int64_t(n) >> 32) != 0) return n % divisor.d; - IntT quotient = n / divisor; - IntT remainder = n - quotient * divisor.d; + using CommonIntT = std::common_type_t; + CommonIntT quotient = n / divisor; + CommonIntT remainder = CommonIntT(n) - quotient * CommonIntT(divisor.d); return remainder; } diff --git a/cpp/tests/util/fast_int_div.cu b/cpp/tests/util/fast_int_div.cu index 4f9af1af7e..141b330944 100644 --- a/cpp/tests/util/fast_int_div.cu +++ b/cpp/tests/util/fast_int_div.cu @@ -13,42 +13,97 @@ namespace raft::util { -constexpr auto kInt32Max = std::numeric_limits::max(); +constexpr auto kInt64Min = std::numeric_limits::min(); constexpr auto kInt64Max = std::numeric_limits::max(); -TEST(FastIntDivTest, UnsupportedNegativeDenumerator) +TEST(FastIntDivTest, UnsupportedDivisors) { ASSERT_THROW(FastIntDiv{-42}, raft::exception); ASSERT_THROW(FastIntDiv{0}, raft::exception); } -TEST(FastIntDivTest, CompareWithNativeDivisionInt32) +TEST(FastIntDivTest, Int32Division) { - std::vector magnitudes{0, 1, 2, 3, 7, 13, 255, 12345, (1 << 20), kInt32Max}; + std::vector numerators{kInt32Min, + -(1 << 20), + -12345, + -255, + -13, + -7, + -3, + -2, + 0, + 1, + 2, + 3, + 7, + 13, + 255, + 12345, + (1 << 20), + kInt32Max}; std::vector divisors{1, 2, 4, 7, 16, 31, 63, 128, 1000, (1 << 15), kInt32Max}; - for (int32_t d : divisors) { + for (auto d : divisors) { FastIntDiv fid(d); - for (int32_t mag : magnitudes) { - for (int32_t n : {mag, -mag}) { - ASSERT_EQ(n / fid, n / d) << "operator/ mismatch for numerator=" << n << " divisor=" << d; - ASSERT_EQ(n % fid, n % d) << "operator% mismatch for numerator=" << n << " divisor=" << d; - } + for (auto n : numerators) { + ASSERT_EQ(n / fid, n / d) << "operator/ mismatch for numerator=" << n << " divisor=" << d; + ASSERT_EQ(n % fid, n % d) << "operator% mismatch for numerator=" << n << " divisor=" << d; } } } -TEST(FastIntDivTest, CompareWithNativeDivisionInt64) +TEST(FastIntDivTest, Int64Divisors) { - std::vector magnitudes{ - 0, 1013, 10007, int64_t(kInt32Max) + 3, (int64_t(1) << 33), (int64_t(3) << 40), kInt64Max}; + std::vector numerators{kInt64Min, + int64_t(kInt32Min) - 12345678, + int64_t(kInt32Min), + -10007, + -1013, + 0, + 1013, + 10007, + int64_t(kInt32Max) + 3, + (int64_t(1) << 33), + (int64_t(3) << 40), + kInt64Max}; + std::vector divisors{ - 1, 129, 772, 1000, (int64_t(1) << 31), int64_t(kInt32Max) + 1, int64_t(3) << 40, kInt64Max}; + 1, 31, 129, 772, 1000, (int64_t(1) << 31), int64_t(kInt32Max) + 5, int64_t(3) << 40, kInt64Max}; - for (int64_t d : divisors) { + for (auto d : divisors) { FastIntDiv fid(d); - for (int64_t mag : magnitudes) { - for (int64_t n : {mag, -mag}) { + for (auto n : numerators) { + ASSERT_EQ(n / fid, n / d) << "operator/ mismatch for numerator=" << n << " divisor=" << d; + ASSERT_EQ(n % fid, n % d) << "operator% mismatch for numerator=" << n << " divisor=" << d; + } + } +} + +TEST(FastIntDivTest, CrossIntTypesDivision) +{ + // i32 numerators, i64 divisors + { + std::vector numerators{kInt32Min, -123, 0, 321, kInt32Max}; + std::vector divisors{1, 4321, kInt64Max}; + + for (auto d : divisors) { + FastIntDiv fid(d); + for (auto n : numerators) { + ASSERT_EQ(n / fid, n / d) << "operator/ mismatch for numerator=" << n << " divisor=" << d; + ASSERT_EQ(n % fid, n % d) << "operator% mismatch for numerator=" << n << " divisor=" << d; + } + } + } + + // i64 numerators, i32 divisors + { + std::vector numerators{kInt64Min, -1234, 0, 4321, kInt64Max}; + std::vector divisors{1, 1234, kInt32Max}; + + for (auto d : divisors) { + FastIntDiv fid(d); + for (auto n : numerators) { ASSERT_EQ(n / fid, n / d) << "operator/ mismatch for numerator=" << n << " divisor=" << d; ASSERT_EQ(n % fid, n % d) << "operator% mismatch for numerator=" << n << " divisor=" << d; } From b974d8b7868d4d7cd1befda33ed67fd6ffb0cbed Mon Sep 17 00:00:00 2001 From: Huy Nguyen Date: Mon, 27 Jul 2026 14:20:10 +0000 Subject: [PATCH 3/4] Add a bench to compare runtime between FastIntDiv and native division --- cpp/bench/prims/CMakeLists.txt | 4 +- cpp/bench/prims/util/fast_int_div.cu | 115 +++++++++++++++++++++++++ cpp/include/raft/util/fast_int_div.cuh | 4 +- 3 files changed, 119 insertions(+), 4 deletions(-) create mode 100644 cpp/bench/prims/util/fast_int_div.cu diff --git a/cpp/bench/prims/CMakeLists.txt b/cpp/bench/prims/CMakeLists.txt index 7dc510639e..d8efd5ea36 100644 --- a/cpp/bench/prims/CMakeLists.txt +++ b/cpp/bench/prims/CMakeLists.txt @@ -1,6 +1,6 @@ # ============================================================================= # cmake-format: off -# SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION. +# SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. # SPDX-License-Identifier: Apache-2.0 # cmake-format: on # ============================================================================= @@ -69,7 +69,7 @@ endfunction() if(BUILD_PRIMS_BENCH) ConfigureBench(NAME CORE_BENCH PATH core/bitset.cu core/copy.cu core/memory_tracking.cu main.cpp) - ConfigureBench(NAME UTIL_BENCH PATH util/popc.cu main.cpp) + ConfigureBench(NAME UTIL_BENCH PATH util/fast_int_div.cu util/popc.cu main.cpp) ConfigureBench( NAME diff --git a/cpp/bench/prims/util/fast_int_div.cu b/cpp/bench/prims/util/fast_int_div.cu new file mode 100644 index 0000000000..af3368c856 --- /dev/null +++ b/cpp/bench/prims/util/fast_int_div.cu @@ -0,0 +1,115 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include + +#include + +#include +#include + +#include +#include +#include + +namespace raft::bench::util { + +constexpr size_t kNumNumerators = 1000 * 1000; +constexpr size_t kNumDivisors = 100; + +/** + * A single kernel serves both variants: `DivisorT` is either `raft::util::FastIntDiv` + * or a plain `IntT` (native division), and both support `operator/` and `operator%` against + * an `IntT` numerator. + */ +template +RAFT_KERNEL divmod_kernel(const IntT* numerators, + const int64_t n_numerators, + const DivisorT* divisors, + const int64_t n_divisors, + IntT* out) +{ + int64_t tid = int64_t(blockIdx.x) * int64_t(blockDim.x) + int64_t(threadIdx.x); + int64_t stride = int64_t(gridDim.x) * int64_t(blockDim.x); + IntT acc = 0; // to prevent compiler from optimizing away the division ops + for (int64_t i = tid; i < n_numerators; i += stride) { + IntT n = numerators[i]; + for (int64_t j = 0; j < n_divisors; ++j) { + acc ^= n / divisors[j]; + acc ^= n % divisors[j]; + } + } + out[tid] = acc; +} + +template +struct fast_int_div_bench : public fixture { + using divisor_t = std::conditional_t, IntT>; + + explicit fast_int_div_bench() + : d_numerators(kNumNumerators, stream), + d_divisors(size_t(kNumDivisors) * sizeof(divisor_t), stream), + out_d(size_t(kBlocks) * size_t(kThreads), stream) + { + std::mt19937_64 rng(42); + std::uniform_int_distribution numerator_dist(raft::util::kInt32Min, + raft::util::kInt32Max); + // non-zero, non-neg divisors + std::uniform_int_distribution divisor_dist(1, raft::util::kInt32Max); + + std::vector h_numerators(kNumNumerators); + for (auto& n : h_numerators) { + n = IntT(numerator_dist(rng)); + } + + std::vector h_divisors; + h_divisors.reserve(kNumDivisors); + for (size_t i = 0; i < kNumDivisors; ++i) { + h_divisors.push_back(divisor_t(IntT(divisor_dist(rng)))); + } + + RAFT_CUDA_TRY(cudaMemcpyAsync(d_numerators.data(), + h_numerators.data(), + h_numerators.size() * sizeof(IntT), + cudaMemcpyHostToDevice, + stream)); + RAFT_CUDA_TRY(cudaMemcpyAsync(d_divisors.data(), + h_divisors.data(), + h_divisors.size() * sizeof(divisor_t), + cudaMemcpyHostToDevice, + stream)); + stream.synchronize(); + } + + void run_benchmark(::benchmark::State& state) override + { + const auto* divisors = static_cast(d_divisors.data()); + loop_on_state(state, [this, divisors]() { + divmod_kernel<<>>( + d_numerators.data(), kNumNumerators, divisors, kNumDivisors, out_d.data()); + RAFT_CUDA_TRY(cudaPeekAtLastError()); + }); + } + + private: + static constexpr int kThreads = 256; + static constexpr int kBlocks = 1024; + + rmm::device_uvector d_numerators; + rmm::device_buffer d_divisors; + rmm::device_uvector out_d; +}; + +using fast_int_div_i32 = fast_int_div_bench; +using native_div_i32 = fast_int_div_bench; +using fast_int_div_i64 = fast_int_div_bench; +using native_div_i64 = fast_int_div_bench; + +RAFT_BENCH_REGISTER(fast_int_div_i32, "FastIntDiv/int32"); +RAFT_BENCH_REGISTER(native_div_i32, "NativeIntDiv/int32"); +RAFT_BENCH_REGISTER(fast_int_div_i64, "FastIntDiv/int64"); +RAFT_BENCH_REGISTER(native_div_i64, "NativeIntDiv/int64"); + +} // namespace raft::bench::util diff --git a/cpp/include/raft/util/fast_int_div.cuh b/cpp/include/raft/util/fast_int_div.cuh index 917e1f2fc5..6cf2af468c 100644 --- a/cpp/include/raft/util/fast_int_div.cuh +++ b/cpp/include/raft/util/fast_int_div.cuh @@ -10,6 +10,7 @@ #include +#include #include namespace raft { @@ -117,8 +118,7 @@ HDI std::common_type_t operator/(NumIntT n, const FastIntDiv> divisor.p; - if (n < 0) ++ret; - return ret; + return ret + CommonIntT(n < 0); } /** From fc9681c1821fc4a1559b8970a8e814f17569f84a Mon Sep 17 00:00:00 2001 From: Huy Nguyen Date: Tue, 28 Jul 2026 09:34:52 +0000 Subject: [PATCH 4/4] Address AI reviews. Increase arith intensity --- cpp/bench/prims/util/fast_int_div.cu | 21 +++++++++++---------- cpp/include/raft/util/fast_int_div.cuh | 1 + cpp/tests/util/fast_int_div.cu | 4 +++- 3 files changed, 15 insertions(+), 11 deletions(-) diff --git a/cpp/bench/prims/util/fast_int_div.cu b/cpp/bench/prims/util/fast_int_div.cu index af3368c856..96f57dfa6d 100644 --- a/cpp/bench/prims/util/fast_int_div.cu +++ b/cpp/bench/prims/util/fast_int_div.cu @@ -34,11 +34,12 @@ RAFT_KERNEL divmod_kernel(const IntT* numerators, int64_t tid = int64_t(blockIdx.x) * int64_t(blockDim.x) + int64_t(threadIdx.x); int64_t stride = int64_t(gridDim.x) * int64_t(blockDim.x); IntT acc = 0; // to prevent compiler from optimizing away the division ops - for (int64_t i = tid; i < n_numerators; i += stride) { - IntT n = numerators[i]; - for (int64_t j = 0; j < n_divisors; ++j) { - acc ^= n / divisors[j]; - acc ^= n % divisors[j]; + for (int64_t j = 0; j < n_divisors; ++j) { + DivisorT divisor = divisors[j]; + for (int64_t i = tid; i < n_numerators; i += stride) { + IntT n = numerators[i]; + acc ^= n / divisor; + acc ^= n % divisor; } } out[tid] = acc; @@ -54,20 +55,20 @@ struct fast_int_div_bench : public fixture { out_d(size_t(kBlocks) * size_t(kThreads), stream) { std::mt19937_64 rng(42); - std::uniform_int_distribution numerator_dist(raft::util::kInt32Min, - raft::util::kInt32Max); + std::uniform_int_distribution numerator_dist(std::numeric_limits::min(), + std::numeric_limits::max()); // non-zero, non-neg divisors - std::uniform_int_distribution divisor_dist(1, raft::util::kInt32Max); + std::uniform_int_distribution divisor_dist(1, std::numeric_limits::max()); std::vector h_numerators(kNumNumerators); for (auto& n : h_numerators) { - n = IntT(numerator_dist(rng)); + n = numerator_dist(rng); } std::vector h_divisors; h_divisors.reserve(kNumDivisors); for (size_t i = 0; i < kNumDivisors; ++i) { - h_divisors.push_back(divisor_t(IntT(divisor_dist(rng)))); + h_divisors.push_back(divisor_t(divisor_dist(rng))); } RAFT_CUDA_TRY(cudaMemcpyAsync(d_numerators.data(), diff --git a/cpp/include/raft/util/fast_int_div.cuh b/cpp/include/raft/util/fast_int_div.cuh index 6cf2af468c..5e51d912e5 100644 --- a/cpp/include/raft/util/fast_int_div.cuh +++ b/cpp/include/raft/util/fast_int_div.cuh @@ -135,6 +135,7 @@ HDI std::common_type_t operator%(NumIntT n, const FastIntDiv