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..96f57dfa6d --- /dev/null +++ b/cpp/bench/prims/util/fast_int_div.cu @@ -0,0 +1,116 @@ +/* + * 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 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; +} + +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(std::numeric_limits::min(), + std::numeric_limits::max()); + // non-zero, non-neg divisors + std::uniform_int_distribution divisor_dist(1, std::numeric_limits::max()); + + std::vector h_numerators(kNumNumerators); + for (auto& n : h_numerators) { + 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(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 d7f385a5d4..5e51d912e5 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 */ @@ -10,18 +10,19 @@ #include +#include #include 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 * - * @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 @@ -51,12 +52,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; } /** @} */ @@ -67,6 +72,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() @@ -79,6 +86,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; @@ -99,13 +109,16 @@ 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 (divisor.d == 1) return n; - IntT ret = (int64_t(divisor.m) * int64_t(n)) >> divisor.p; - if (n < 0) ++ret; - return ret; + 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; + return ret + CommonIntT(n < 0); } /** @@ -115,12 +128,14 @@ 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) { - 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; + // return n % divisor.d; } }; // namespace util diff --git a/cpp/tests/util/fast_int_div.cu b/cpp/tests/util/fast_int_div.cu index eb391c1dbb..3a5afb6684 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,42 +13,104 @@ namespace raft::util { -constexpr int64_t kInt32Max = std::numeric_limits::max(); +constexpr auto kInt64Min = std::numeric_limits::min(); +constexpr auto kInt64Max = 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}; +TEST(FastIntDivTest, UnsupportedDivisors) +{ + ASSERT_THROW(FastIntDiv{-42}, raft::exception); + ASSERT_THROW(FastIntDiv{0}, raft::exception); +} - 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; - } - } +TEST(FastIntDivTest, Int32Division) +{ + 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 (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; } } -}; +} -using FastIntDivTypes = ::testing::Types; -TYPED_TEST_CASE(FastIntDivTest, FastIntDivTypes); +TEST(FastIntDivTest, Int64Divisors) +{ + std::vector numerators{kInt64Min, + int64_t(kInt32Min) - 12345678, + int64_t(kInt32Min), + -10007, + -1013, + 0, + 1013, + 10007, + int64_t(kInt32Max) + 1, + int64_t(kInt32Max) + 2, + (int64_t(1) << 32), + (int64_t(1) << 33), + (int64_t(3) << 40), + kInt64Max}; -TYPED_TEST(FastIntDivTest, CompareWithNativeDivision) { this->CompareWithNativeDivision(); } + std::vector divisors{ + 1, 31, 129, 772, 1000, (int64_t(1) << 31), int64_t(kInt32Max) + 5, int64_t(3) << 40, kInt64Max}; -TEST(FastIntDiv, Int64NumeratorPastInt32Boundary) -{ - for (int64_t d : {129, 772, 1000}) { + for (auto d : divisors) { FastIntDiv fid(d); - for (int64_t n : {1LL << 31, 2147704000LL, 3LL << 30}) { // in [2^31, 2^32) + 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; + } + } + } +} + } // namespace raft::util