From a9e1adc6a55f03fb98199ab8c7f4bc82e9849a73 Mon Sep 17 00:00:00 2001 From: Louis Sugy Date: Sat, 21 Jan 2023 21:40:14 +0100 Subject: [PATCH] Improvement of the math API wrappers (#1146) Solves #1025 Provides a centralized collection of host- and device-friendly wrappers around common math operations, with generalizations when useful. Deprecates former `myXxx` wrappers. Those wrappers are mostly intended to future-proof the API as well as simplify the definition of host-device functions. Authors: - Louis Sugy (https://github.com/Nyrio) - Corey J. Nolet (https://github.com/cjnolet) Approvers: - Corey J. Nolet (https://github.com/cjnolet) URL: https://github.com/rapidsai/raft/pull/1146 --- cpp/include/raft/core/math.hpp | 320 ++++++++++++++++ cpp/include/raft/core/operators.hpp | 27 +- cpp/include/raft/distance/detail/canberra.cuh | 4 +- .../raft/distance/detail/chebyshev.cuh | 4 +- .../raft/distance/detail/correlation.cuh | 4 +- .../raft/distance/detail/euclidean.cuh | 8 +- .../raft/distance/detail/fused_l2_nn.cuh | 4 +- .../raft/distance/detail/hellinger.cuh | 4 +- .../raft/distance/detail/jensen_shannon.cuh | 8 +- .../raft/distance/detail/kl_divergence.cuh | 14 +- cpp/include/raft/distance/detail/l1.cuh | 2 +- .../raft/distance/detail/minkowski.cuh | 8 +- cpp/include/raft/linalg/detail/lstsq.cuh | 4 +- cpp/include/raft/matrix/detail/math.cuh | 10 +- .../raft/random/detail/make_regression.cuh | 6 +- cpp/include/raft/random/detail/rng_device.cuh | 24 +- .../sparse/distance/detail/l2_distance.cuh | 12 +- .../sparse/distance/detail/lp_distance.cuh | 6 +- .../spatial/knn/detail/ball_cover/common.cuh | 4 +- .../spatial/knn/detail/haversine_distance.cuh | 8 +- .../raft/spectral/detail/spectral_util.cuh | 4 +- cpp/include/raft/stats/detail/stddev.cuh | 6 +- cpp/include/raft/util/cuda_utils.cuh | 90 ++--- cpp/test/CMakeLists.txt | 2 + cpp/test/core/math_device.cu | 352 ++++++++++++++++++ cpp/test/core/math_host.cpp | 195 ++++++++++ cpp/test/distance/distance_base.cuh | 22 +- cpp/test/distance/fused_l2_nn.cu | 2 +- cpp/test/linalg/matrix_vector.cu | 4 +- cpp/test/linalg/norm.cu | 10 +- cpp/test/linalg/power.cu | 6 +- cpp/test/linalg/sqrt.cu | 4 +- cpp/test/matrix/math.cu | 4 +- cpp/test/neighbors/ann_utils.cuh | 2 +- cpp/test/random/rng.cu | 14 +- 35 files changed, 1034 insertions(+), 164 deletions(-) create mode 100644 cpp/include/raft/core/math.hpp create mode 100644 cpp/test/core/math_device.cu create mode 100644 cpp/test/core/math_host.cpp diff --git a/cpp/include/raft/core/math.hpp b/cpp/include/raft/core/math.hpp new file mode 100644 index 0000000000..c5f08b84b7 --- /dev/null +++ b/cpp/include/raft/core/math.hpp @@ -0,0 +1,320 @@ +/* + * Copyright (c) 2023, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#pragma once + +#include +#include +#include + +#include + +namespace raft { + +/** + * @defgroup Absolute Absolute value + * @{ + */ +template +RAFT_INLINE_FUNCTION auto abs(T x) + -> std::enable_if_t || std::is_same_v || + std::is_same_v || std::is_same_v || + std::is_same_v, + T> +{ +#ifdef __CUDA_ARCH__ + return ::abs(x); +#else + return std::abs(x); +#endif +} +template +constexpr RAFT_INLINE_FUNCTION auto abs(T x) + -> std::enable_if_t && !std::is_same_v && + !std::is_same_v && !std::is_same_v && + !std::is_same_v, + T> +{ + return x < T{0} ? -x : x; +} +/** @} */ + +/** + * @defgroup Trigonometry Trigonometry functions + * @{ + */ +/** Inverse cosine */ +template +RAFT_INLINE_FUNCTION auto acos(T x) +{ +#ifdef __CUDA_ARCH__ + return ::acos(x); +#else + return std::acos(x); +#endif +} + +/** Inverse sine */ +template +RAFT_INLINE_FUNCTION auto asin(T x) +{ +#ifdef __CUDA_ARCH__ + return ::asin(x); +#else + return std::asin(x); +#endif +} + +/** Inverse hyperbolic tangent */ +template +RAFT_INLINE_FUNCTION auto atanh(T x) +{ +#ifdef __CUDA_ARCH__ + return ::atanh(x); +#else + return std::atanh(x); +#endif +} + +/** Cosine */ +template +RAFT_INLINE_FUNCTION auto cos(T x) +{ +#ifdef __CUDA_ARCH__ + return ::cos(x); +#else + return std::cos(x); +#endif +} + +/** Sine */ +template +RAFT_INLINE_FUNCTION auto sin(T x) +{ +#ifdef __CUDA_ARCH__ + return ::sin(x); +#else + return std::sin(x); +#endif +} + +/** Sine and cosine */ +template +RAFT_INLINE_FUNCTION std::enable_if_t || std::is_same_v> sincos( + const T& x, T* s, T* c) +{ +#ifdef __CUDA_ARCH__ + ::sincos(x, s, c); +#else + *s = std::sin(x); + *c = std::cos(x); +#endif +} + +/** Hyperbolic tangent */ +template +RAFT_INLINE_FUNCTION auto tanh(T x) +{ +#ifdef __CUDA_ARCH__ + return ::tanh(x); +#else + return std::tanh(x); +#endif +} +/** @} */ + +/** + * @defgroup Exponential Exponential and logarithm + * @{ + */ +/** Exponential function */ +template +RAFT_INLINE_FUNCTION auto exp(T x) +{ +#ifdef __CUDA_ARCH__ + return ::exp(x); +#else + return std::exp(x); +#endif +} + +/** Natural logarithm */ +template +RAFT_INLINE_FUNCTION auto log(T x) +{ +#ifdef __CUDA_ARCH__ + return ::log(x); +#else + return std::log(x); +#endif +} +/** @} */ + +/** + * @defgroup Maximum Maximum of two or more values. + * + * The CUDA Math API has overloads for all combinations of float/double. We provide similar + * functionality while wrapping around std::max, which only supports arguments of the same type. + * However, though the CUDA Math API supports combinations of unsigned and signed integers, this is + * very error-prone so we do not support that and require the user to cast instead. (e.g the max of + * -1 and 1u is 4294967295u...) + * + * When no overload matches, we provide a generic implementation but require that both types be the + * same (and that the less-than operator be defined). + * @{ + */ +template +RAFT_INLINE_FUNCTION auto max(const T1& x, const T2& y) +{ +#ifdef __CUDA_ARCH__ + // Combinations of types supported by the CUDA Math API + if constexpr ((std::is_integral_v && std::is_integral_v && std::is_same_v) || + ((std::is_same_v || std::is_same_v)&&( + std::is_same_v || std::is_same_v))) { + return ::max(x, y); + } + // Else, check that the types are the same and provide a generic implementation + else { + static_assert( + std::is_same_v, + "No native max overload for these types. Both argument types must be the same to use " + "the generic max. Please cast appropriately."); + return (x < y) ? y : x; + } +#else + if constexpr (std::is_same_v && std::is_same_v) { + return std::max(static_cast(x), y); + } else if constexpr (std::is_same_v && std::is_same_v) { + return std::max(x, static_cast(y)); + } else { + static_assert( + std::is_same_v, + "std::max requires that both argument types be the same. Please cast appropriately."); + return std::max(x, y); + } +#endif +} + +/** Many-argument overload to avoid verbose nested calls or use with variadic arguments */ +template +RAFT_INLINE_FUNCTION auto max(const T1& x, const T2& y, Args&&... args) +{ + return raft::max(x, raft::max(y, std::forward(args)...)); +} + +/** One-argument overload for convenience when using with variadic arguments */ +template +constexpr RAFT_INLINE_FUNCTION auto max(const T& x) +{ + return x; +} +/** @} */ + +/** + * @defgroup Minimum Minimum of two or more values. + * + * The CUDA Math API has overloads for all combinations of float/double. We provide similar + * functionality while wrapping around std::min, which only supports arguments of the same type. + * However, though the CUDA Math API supports combinations of unsigned and signed integers, this is + * very error-prone so we do not support that and require the user to cast instead. (e.g the min of + * -1 and 1u is 1u...) + * + * When no overload matches, we provide a generic implementation but require that both types be the + * same (and that the less-than operator be defined). + * @{ + */ +template +RAFT_INLINE_FUNCTION auto min(const T1& x, const T2& y) +{ +#ifdef __CUDA_ARCH__ + // Combinations of types supported by the CUDA Math API + if constexpr ((std::is_integral_v && std::is_integral_v && std::is_same_v) || + ((std::is_same_v || std::is_same_v)&&( + std::is_same_v || std::is_same_v))) { + return ::min(x, y); + } + // Else, check that the types are the same and provide a generic implementation + else { + static_assert( + std::is_same_v, + "No native min overload for these types. Both argument types must be the same to use " + "the generic min. Please cast appropriately."); + return (y < x) ? y : x; + } +#else + if constexpr (std::is_same_v && std::is_same_v) { + return std::min(static_cast(x), y); + } else if constexpr (std::is_same_v && std::is_same_v) { + return std::min(x, static_cast(y)); + } else { + static_assert( + std::is_same_v, + "std::min requires that both argument types be the same. Please cast appropriately."); + return std::min(x, y); + } +#endif +} + +/** Many-argument overload to avoid verbose nested calls or use with variadic arguments */ +template +RAFT_INLINE_FUNCTION auto min(const T1& x, const T2& y, Args&&... args) +{ + return raft::min(x, raft::min(y, std::forward(args)...)); +} + +/** One-argument overload for convenience when using with variadic arguments */ +template +constexpr RAFT_INLINE_FUNCTION auto min(const T& x) +{ + return x; +} +/** @} */ + +/** + * @defgroup Power Power and root functions + * @{ + */ +/** Power */ +template +RAFT_INLINE_FUNCTION auto pow(T1 x, T2 y) +{ +#ifdef __CUDA_ARCH__ + return ::pow(x, y); +#else + return std::pow(x, y); +#endif +} + +/** Square root */ +template +RAFT_INLINE_FUNCTION auto sqrt(T x) +{ +#ifdef __CUDA_ARCH__ + return ::sqrt(x); +#else + return std::sqrt(x); +#endif +} +/** @} */ + +/** Sign */ +template +RAFT_INLINE_FUNCTION auto sgn(T val) -> int +{ + return (T(0) < val) - (val < T(0)); +} + +} // namespace raft diff --git a/cpp/include/raft/core/operators.hpp b/cpp/include/raft/core/operators.hpp index 398354df46..de27c2b271 100644 --- a/cpp/include/raft/core/operators.hpp +++ b/cpp/include/raft/core/operators.hpp @@ -23,6 +23,7 @@ #include #include +#include namespace raft { @@ -75,9 +76,9 @@ struct value_op { struct sqrt_op { template - constexpr RAFT_INLINE_FUNCTION auto operator()(const Type& in, UnusedArgs...) const + RAFT_INLINE_FUNCTION auto operator()(const Type& in, UnusedArgs...) const { - return std::sqrt(in); + return raft::sqrt(in); } }; @@ -91,9 +92,9 @@ struct nz_op { struct abs_op { template - constexpr RAFT_INLINE_FUNCTION auto operator()(const Type& in, UnusedArgs...) const + RAFT_INLINE_FUNCTION auto operator()(const Type& in, UnusedArgs...) const { - return std::abs(in); + return raft::abs(in); } }; @@ -148,27 +149,25 @@ struct div_checkzero_op { struct pow_op { template - constexpr RAFT_INLINE_FUNCTION auto operator()(const Type& a, const Type& b) const + RAFT_INLINE_FUNCTION auto operator()(const Type& a, const Type& b) const { - return std::pow(a, b); + return raft::pow(a, b); } }; struct min_op { - template - constexpr RAFT_INLINE_FUNCTION auto operator()(const Type& a, const Type& b) const + template + RAFT_INLINE_FUNCTION auto operator()(Args&&... args) const { - if (a > b) { return b; } - return a; + return raft::min(std::forward(args)...); } }; struct max_op { - template - constexpr RAFT_INLINE_FUNCTION auto operator()(const Type& a, const Type& b) const + template + RAFT_INLINE_FUNCTION auto operator()(Args&&... args) const { - if (b > a) { return b; } - return a; + return raft::max(std::forward(args)...); } }; diff --git a/cpp/include/raft/distance/detail/canberra.cuh b/cpp/include/raft/distance/detail/canberra.cuh index 43a904edba..f17a26dc4b 100644 --- a/cpp/include/raft/distance/detail/canberra.cuh +++ b/cpp/include/raft/distance/detail/canberra.cuh @@ -73,8 +73,8 @@ static void canberraImpl(const DataT* x, // Accumulation operation lambda auto core_lambda = [] __device__(AccT & acc, DataT & x, DataT & y) { - const auto diff = raft::myAbs(x - y); - const auto add = raft::myAbs(x) + raft::myAbs(y); + const auto diff = raft::abs(x - y); + const auto add = raft::abs(x) + raft::abs(y); // deal with potential for 0 in denominator by // forcing 1/0 instead acc += ((add != 0) * diff / (add + (add == 0))); diff --git a/cpp/include/raft/distance/detail/chebyshev.cuh b/cpp/include/raft/distance/detail/chebyshev.cuh index 52573bd170..43b36e7921 100644 --- a/cpp/include/raft/distance/detail/chebyshev.cuh +++ b/cpp/include/raft/distance/detail/chebyshev.cuh @@ -73,8 +73,8 @@ static void chebyshevImpl(const DataT* x, // Accumulation operation lambda auto core_lambda = [] __device__(AccT & acc, DataT & x, DataT & y) { - const auto diff = raft::myAbs(x - y); - acc = raft::myMax(acc, diff); + const auto diff = raft::abs(x - y); + acc = raft::max(acc, diff); }; // epilogue operation lambda for final value calculation diff --git a/cpp/include/raft/distance/detail/correlation.cuh b/cpp/include/raft/distance/detail/correlation.cuh index 9bdbbf112c..f7fe3678e6 100644 --- a/cpp/include/raft/distance/detail/correlation.cuh +++ b/cpp/include/raft/distance/detail/correlation.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021-2022, NVIDIA CORPORATION. + * Copyright (c) 2021-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -125,7 +125,7 @@ static void correlationImpl(const DataT* x, auto Q_denom = k * regx2n[i] - (regxn[i] * regxn[i]); auto R_denom = k * regy2n[j] - (regyn[j] * regyn[j]); - acc[i][j] = 1 - (numer / raft::mySqrt(Q_denom * R_denom)); + acc[i][j] = 1 - (numer / raft::sqrt(Q_denom * R_denom)); } } }; diff --git a/cpp/include/raft/distance/detail/euclidean.cuh b/cpp/include/raft/distance/detail/euclidean.cuh index 4184810fff..1a2db63f5c 100644 --- a/cpp/include/raft/distance/detail/euclidean.cuh +++ b/cpp/include/raft/distance/detail/euclidean.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2022, NVIDIA CORPORATION. + * Copyright (c) 2018-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -34,7 +34,7 @@ struct L2ExpandedOp { __device__ AccT operator()(DataT& aNorm, const DataT& bNorm, DataT& accVal) const noexcept { AccT outVal = aNorm + bNorm - DataT(2.0) * accVal; - return sqrt ? raft::mySqrt(outVal) : outVal; + return sqrt ? raft::sqrt(outVal) : outVal; } __device__ AccT operator()(DataT aData) const noexcept { return aData; } @@ -130,7 +130,7 @@ void euclideanExpImpl(const DataT* x, for (int i = 0; i < KPolicy::AccRowsPerTh; ++i) { #pragma unroll for (int j = 0; j < KPolicy::AccColsPerTh; ++j) { - acc[i][j] = raft::mySqrt(acc[i][j]); + acc[i][j] = raft::sqrt(acc[i][j]); } } } @@ -350,7 +350,7 @@ void euclideanUnExpImpl(const DataT* x, for (int i = 0; i < KPolicy::AccRowsPerTh; ++i) { #pragma unroll for (int j = 0; j < KPolicy::AccColsPerTh; ++j) { - acc[i][j] = raft::mySqrt(acc[i][j]); + acc[i][j] = raft::sqrt(acc[i][j]); } } } diff --git a/cpp/include/raft/distance/detail/fused_l2_nn.cuh b/cpp/include/raft/distance/detail/fused_l2_nn.cuh index c9750df8ad..447359ffe6 100644 --- a/cpp/include/raft/distance/detail/fused_l2_nn.cuh +++ b/cpp/include/raft/distance/detail/fused_l2_nn.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021-2022, NVIDIA CORPORATION. + * Copyright (c) 2021-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -175,7 +175,7 @@ __global__ __launch_bounds__(P::Nthreads, 2) void fusedL2NNkernel(OutT* min, #pragma unroll for (int j = 0; j < P::AccColsPerTh; ++j) { auto acc_ij = acc[i][j]; - acc[i][j] = acc_ij > DataT{0} ? raft::mySqrt(acc_ij) : DataT{0}; + acc[i][j] = acc_ij > DataT{0} ? raft::sqrt(acc_ij) : DataT{0}; } } } diff --git a/cpp/include/raft/distance/detail/hellinger.cuh b/cpp/include/raft/distance/detail/hellinger.cuh index 51f462ab36..13507fe84f 100644 --- a/cpp/include/raft/distance/detail/hellinger.cuh +++ b/cpp/include/raft/distance/detail/hellinger.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021-2022, NVIDIA CORPORATION. + * Copyright (c) 2021-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -105,7 +105,7 @@ static void hellingerImpl(const DataT* x, // Adjust to replace NaN in sqrt with 0 if input to sqrt is negative const auto finalVal = (1 - acc[i][j]); const auto rectifier = (!signbit(finalVal)); - acc[i][j] = raft::mySqrt(rectifier * finalVal); + acc[i][j] = raft::sqrt(rectifier * finalVal); } } }; diff --git a/cpp/include/raft/distance/detail/jensen_shannon.cuh b/cpp/include/raft/distance/detail/jensen_shannon.cuh index 92ee071cf5..f96da01b87 100644 --- a/cpp/include/raft/distance/detail/jensen_shannon.cuh +++ b/cpp/include/raft/distance/detail/jensen_shannon.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021, NVIDIA CORPORATION. + * Copyright (c) 2021-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -78,11 +78,11 @@ static void jensenShannonImpl(const DataT* x, auto core_lambda = [] __device__(AccT & acc, DataT & x, DataT & y) { const DataT m = 0.5f * (x + y); const bool m_zero = (m == 0); - const auto logM = (!m_zero) * raft::myLog(m + m_zero); + const auto logM = (!m_zero) * raft::log(m + m_zero); const bool x_zero = (x == 0); const bool y_zero = (y == 0); - acc += (-x * (logM - raft::myLog(x + x_zero))) + (-y * (logM - raft::myLog(y + y_zero))); + acc += (-x * (logM - raft::log(x + x_zero))) + (-y * (logM - raft::log(y + y_zero))); }; // epilogue operation lambda for final value calculation @@ -95,7 +95,7 @@ static void jensenShannonImpl(const DataT* x, for (int i = 0; i < KPolicy::AccRowsPerTh; ++i) { #pragma unroll for (int j = 0; j < KPolicy::AccColsPerTh; ++j) { - acc[i][j] = raft::mySqrt(0.5 * acc[i][j]); + acc[i][j] = raft::sqrt(0.5 * acc[i][j]); } } }; diff --git a/cpp/include/raft/distance/detail/kl_divergence.cuh b/cpp/include/raft/distance/detail/kl_divergence.cuh index 4c0c4b6ace..7ebeaf4de9 100644 --- a/cpp/include/raft/distance/detail/kl_divergence.cuh +++ b/cpp/include/raft/distance/detail/kl_divergence.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021, NVIDIA CORPORATION. + * Copyright (c) 2021-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -81,10 +81,10 @@ static void klDivergenceImpl(const DataT* x, auto core_lambda = [] __device__(AccT & acc, DataT & x, DataT & y) { if (isRowMajor) { const bool x_zero = (x == 0); - acc += x * (raft::myLog(x + x_zero) - y); + acc += x * (raft::log(x + x_zero) - y); } else { const bool y_zero = (y == 0); - acc += y * (raft::myLog(y + y_zero) - x); + acc += y * (raft::log(y + y_zero) - x); } }; @@ -92,23 +92,23 @@ static void klDivergenceImpl(const DataT* x, if (isRowMajor) { const bool x_zero = (x == 0); const bool y_zero = (y == 0); - acc += x * (raft::myLog(x + x_zero) - (!y_zero) * raft::myLog(y + y_zero)); + acc += x * (raft::log(x + x_zero) - (!y_zero) * raft::log(y + y_zero)); } else { const bool y_zero = (y == 0); const bool x_zero = (x == 0); - acc += y * (raft::myLog(y + y_zero) - (!x_zero) * raft::myLog(x + x_zero)); + acc += y * (raft::log(y + y_zero) - (!x_zero) * raft::log(x + x_zero)); } }; auto unaryOp_lambda = [] __device__(DataT input) { const bool x_zero = (input == 0); - return (!x_zero) * raft::myLog(input + x_zero); + return (!x_zero) * raft::log(input + x_zero); }; auto unaryOp_lambda_reverse = [] __device__(DataT input) { // reverse previous log (x) back to x using (e ^ log(x)) const bool x_zero = (input == 0); - return (!x_zero) * raft::myExp(input); + return (!x_zero) * raft::exp(input); }; // epilogue operation lambda for final value calculation diff --git a/cpp/include/raft/distance/detail/l1.cuh b/cpp/include/raft/distance/detail/l1.cuh index 87893bab7c..bf10651b60 100644 --- a/cpp/include/raft/distance/detail/l1.cuh +++ b/cpp/include/raft/distance/detail/l1.cuh @@ -71,7 +71,7 @@ static void l1Impl(const DataT* x, // Accumulation operation lambda auto core_lambda = [] __device__(AccT & acc, DataT & x, DataT & y) { - const auto diff = raft::myAbs(x - y); + const auto diff = raft::abs(x - y); acc += diff; }; diff --git a/cpp/include/raft/distance/detail/minkowski.cuh b/cpp/include/raft/distance/detail/minkowski.cuh index bda83babf1..42af8cd281 100644 --- a/cpp/include/raft/distance/detail/minkowski.cuh +++ b/cpp/include/raft/distance/detail/minkowski.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021-2022, NVIDIA CORPORATION. + * Copyright (c) 2021-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -74,8 +74,8 @@ void minkowskiUnExpImpl(const DataT* x, // Accumulation operation lambda auto core_lambda = [p] __device__(AccT & acc, DataT & x, DataT & y) { - const auto diff = raft::myAbs(x - y); - acc += raft::myPow(diff, p); + const auto diff = raft::abs(x - y); + acc += raft::pow(diff, p); }; // epilogue operation lambda for final value calculation @@ -89,7 +89,7 @@ void minkowskiUnExpImpl(const DataT* x, for (int i = 0; i < KPolicy::AccRowsPerTh; ++i) { #pragma unroll for (int j = 0; j < KPolicy::AccColsPerTh; ++j) { - acc[i][j] = raft::myPow(acc[i][j], one_over_p); + acc[i][j] = raft::pow(acc[i][j], one_over_p); } } }; diff --git a/cpp/include/raft/linalg/detail/lstsq.cuh b/cpp/include/raft/linalg/detail/lstsq.cuh index 1273956b21..f0cf300e2f 100644 --- a/cpp/include/raft/linalg/detail/lstsq.cuh +++ b/cpp/include/raft/linalg/detail/lstsq.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2022, NVIDIA CORPORATION. + * Copyright (c) 2018-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -104,7 +104,7 @@ struct DivideByNonZero { operator()(const math_t a, const math_t b) const { - return raft::myAbs(b) >= eps ? a / b : a; + return raft::abs(b) >= eps ? a / b : a; } }; diff --git a/cpp/include/raft/matrix/detail/math.cuh b/cpp/include/raft/matrix/detail/math.cuh index c559da3942..f5c33d1cf6 100644 --- a/cpp/include/raft/matrix/detail/math.cuh +++ b/cpp/include/raft/matrix/detail/math.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021-2022, NVIDIA CORPORATION. + * Copyright (c) 2021-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -87,10 +87,10 @@ void seqRoot(math_t* in, if (a < math_t(0)) { return math_t(0); } else { - return sqrt(a * scalar); + return raft::sqrt(a * scalar); } } else { - return sqrt(a * scalar); + return raft::sqrt(a * scalar); } }, stream); @@ -278,7 +278,7 @@ void matrixVectorBinaryDivSkipZero(Type* data, rowMajor, bcastAlongRows, [] __device__(Type a, Type b) { - if (raft::myAbs(b) < Type(1e-10)) + if (raft::abs(b) < Type(1e-10)) return Type(0); else return a / b; @@ -294,7 +294,7 @@ void matrixVectorBinaryDivSkipZero(Type* data, rowMajor, bcastAlongRows, [] __device__(Type a, Type b) { - if (raft::myAbs(b) < Type(1e-10)) + if (raft::abs(b) < Type(1e-10)) return a; else return a / b; diff --git a/cpp/include/raft/random/detail/make_regression.cuh b/cpp/include/raft/random/detail/make_regression.cuh index cb0949c458..057196cd74 100644 --- a/cpp/include/raft/random/detail/make_regression.cuh +++ b/cpp/include/raft/random/detail/make_regression.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2022, NVIDIA CORPORATION. + * Copyright (c) 2019-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -44,8 +44,8 @@ static __global__ void _singular_profile_kernel(DataT* out, IdxT n, DataT tail_s IdxT tid = threadIdx.x + blockIdx.x * blockDim.x; if (tid < n) { DataT sval = static_cast(tid) / rank; - DataT low_rank = ((DataT)1.0 - tail_strength) * raft::myExp(-sval * sval); - DataT tail = tail_strength * raft::myExp((DataT)-0.1 * sval); + DataT low_rank = ((DataT)1.0 - tail_strength) * raft::exp(-sval * sval); + DataT tail = tail_strength * raft::exp((DataT)-0.1 * sval); out[tid] = low_rank + tail; } } diff --git a/cpp/include/raft/random/detail/rng_device.cuh b/cpp/include/raft/random/detail/rng_device.cuh index 6c75a4fa78..7f994fb07f 100644 --- a/cpp/include/raft/random/detail/rng_device.cuh +++ b/cpp/include/raft/random/detail/rng_device.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022, NVIDIA CORPORATION. + * Copyright (c) 2022-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -143,10 +143,10 @@ DI void box_muller_transform(Type& val1, Type& val2, Type sigma1, Type mu1, Type { constexpr Type twoPi = Type(2.0) * Type(3.141592654); constexpr Type minus2 = -Type(2.0); - Type R = raft::mySqrt(minus2 * raft::myLog(val1)); + Type R = raft::sqrt(minus2 * raft::log(val1)); Type theta = twoPi * val2; Type s, c; - raft::mySinCos(theta, s, c); + raft::sincos(theta, &s, &c); val1 = R * c * sigma1 + mu1; val2 = R * s * sigma2 + mu2; } @@ -323,7 +323,7 @@ DI void custom_next( gen.next(res); } while (res == OutType(0.0)); - *val = params.mu - params.beta * raft::myLog(-raft::myLog(res)); + *val = params.mu - params.beta * raft::log(-raft::log(res)); } template @@ -340,8 +340,8 @@ DI void custom_next(GenType& gen, gen.next(res2); box_muller_transform(res1, res2, params.sigma, params.mu); - *val = raft::myExp(res1); - *(val + 1) = raft::myExp(res2); + *val = raft::exp(res1); + *(val + 1) = raft::exp(res2); } template @@ -358,7 +358,7 @@ DI void custom_next(GenType& gen, } while (res == OutType(0.0)); constexpr OutType one = (OutType)1.0; - *val = params.mu - params.scale * raft::myLog(one / res - one); + *val = params.mu - params.scale * raft::log(one / res - one); } template @@ -371,7 +371,7 @@ DI void custom_next(GenType& gen, OutType res; gen.next(res); constexpr OutType one = (OutType)1.0; - *val = -raft::myLog(one - res) / params.lambda; + *val = -raft::log(one - res) / params.lambda; } template @@ -386,7 +386,7 @@ DI void custom_next(GenType& gen, constexpr OutType one = (OutType)1.0; constexpr OutType two = (OutType)2.0; - *val = raft::mySqrt(-two * raft::myLog(one - res)) * params.sigma; + *val = raft::sqrt(-two * raft::log(one - res)) * params.sigma; } template @@ -409,9 +409,9 @@ DI void custom_next(GenType& gen, // The <= comparison here means, number of samples going in `if` branch are more by 1 than `else` // branch. However it does not matter as for 0.5 both branches evaluate to same result. if (res <= oneHalf) { - out = params.mu + params.scale * raft::myLog(two * res); + out = params.mu + params.scale * raft::log(two * res); } else { - out = params.mu - params.scale * raft::myLog(two * (one - res)); + out = params.mu - params.scale * raft::log(two * (one - res)); } *val = out; } @@ -424,7 +424,7 @@ DI void custom_next( gen.next(res); params.inIdxPtr[idx] = idx; constexpr OutType one = (OutType)1.0; - auto exp = -raft::myLog(one - res); + auto exp = -raft::log(one - res); if (params.wts != nullptr) { *val = exp / params.wts[idx]; } else { diff --git a/cpp/include/raft/sparse/distance/detail/l2_distance.cuh b/cpp/include/raft/sparse/distance/detail/l2_distance.cuh index 3c852235df..2f165b3ff2 100644 --- a/cpp/include/raft/sparse/distance/detail/l2_distance.cuh +++ b/cpp/include/raft/sparse/distance/detail/l2_distance.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021-2022, NVIDIA CORPORATION. + * Copyright (c) 2021-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -112,7 +112,7 @@ __global__ void compute_correlation_warp_kernel(value_t* __restrict__ C, value_t Q_denom = n * Q_l2 - (Q_l1 * Q_l1); value_t R_denom = n * R_l2 - (R_l1 * R_l1); - value_t val = 1 - (numer / sqrt(Q_denom * R_denom)); + value_t val = 1 - (numer / raft::sqrt(Q_denom * R_denom)); // correct for small instabilities C[(size_t)i * n_cols + j] = val * (fabs(val) >= 0.0001); @@ -292,7 +292,7 @@ class l2_sqrt_expanded_distances_t : public l2_expanded_distances_tconfig_->a_nrows * this->config_->b_nrows, [] __device__(value_t input) { int neg = input < 0 ? -1 : 1; - return sqrt(abs(input) * neg); + return raft::sqrt(abs(input) * neg); }, this->config_->handle.get_stream()); } @@ -379,7 +379,7 @@ class cosine_expanded_distances_t : public distances_t { config_->b_nrows, config_->handle.get_stream(), [] __device__ __host__(value_t dot, value_t q_norm, value_t r_norm) { - value_t norms = sqrt(q_norm) * sqrt(r_norm); + value_t norms = raft::sqrt(q_norm) * raft::sqrt(r_norm); // deal with potential for 0 in denominator by forcing 0/1 instead value_t cos = ((norms != 0) * dot) / ((norms == 0) + norms); @@ -429,7 +429,7 @@ class hellinger_expanded_distances_t : public distances_t { out_dists, *config_, coo_rows.data(), - [] __device__(value_t a, value_t b) { return sqrt(a) * sqrt(b); }, + [] __device__(value_t a, value_t b) { return raft::sqrt(a) * raft::sqrt(b); }, raft::add_op(), raft::atomic_add_op()); @@ -440,7 +440,7 @@ class hellinger_expanded_distances_t : public distances_t { [=] __device__(value_t input) { // Adjust to replace NaN in sqrt with 0 if input to sqrt is negative bool rectifier = (1 - input) > 0; - return sqrt(rectifier * (1 - input)); + return raft::sqrt(rectifier * (1 - input)); }, config_->handle.get_stream()); } diff --git a/cpp/include/raft/sparse/distance/detail/lp_distance.cuh b/cpp/include/raft/sparse/distance/detail/lp_distance.cuh index a973aebbab..f67109afbc 100644 --- a/cpp/include/raft/sparse/distance/detail/lp_distance.cuh +++ b/cpp/include/raft/sparse/distance/detail/lp_distance.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021-2022, NVIDIA CORPORATION. + * Copyright (c) 2021-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -132,7 +132,7 @@ class l2_sqrt_unexpanded_distances_t : public l2_unexpanded_distances_tconfig_->a_nrows * this->config_->b_nrows, [] __device__(value_t input) { int neg = input < 0 ? -1 : 1; - return sqrt(abs(input) * neg); + return raft::sqrt(abs(input) * neg); }, this->config_->handle.get_stream()); } @@ -274,7 +274,7 @@ class jensen_shannon_unexpanded_distances_t : public distances_t { out_dists, out_dists, config_->a_nrows * config_->b_nrows, - [=] __device__(value_t input) { return sqrt(0.5 * input); }, + [=] __device__(value_t input) { return raft::sqrt(0.5 * input); }, config_->handle.get_stream()); } diff --git a/cpp/include/raft/spatial/knn/detail/ball_cover/common.cuh b/cpp/include/raft/spatial/knn/detail/ball_cover/common.cuh index b09cf0da10..0a6718f5a5 100644 --- a/cpp/include/raft/spatial/knn/detail/ball_cover/common.cuh +++ b/cpp/include/raft/spatial/knn/detail/ball_cover/common.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021-2022, NVIDIA CORPORATION. + * Copyright (c) 2021-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -71,7 +71,7 @@ struct EuclideanFunc : public DistFunc { sum_sq += diff * diff; } - return sqrt(sum_sq); + return raft::sqrt(sum_sq); } }; diff --git a/cpp/include/raft/spatial/knn/detail/haversine_distance.cuh b/cpp/include/raft/spatial/knn/detail/haversine_distance.cuh index e073841dd3..9cecc0adf4 100644 --- a/cpp/include/raft/spatial/knn/detail/haversine_distance.cuh +++ b/cpp/include/raft/spatial/knn/detail/haversine_distance.cuh @@ -32,11 +32,11 @@ namespace detail { template DI value_t compute_haversine(value_t x1, value_t y1, value_t x2, value_t y2) { - value_t sin_0 = sin(0.5 * (x1 - y1)); - value_t sin_1 = sin(0.5 * (x2 - y2)); - value_t rdist = sin_0 * sin_0 + cos(x1) * cos(y1) * sin_1 * sin_1; + value_t sin_0 = raft::sin(0.5 * (x1 - y1)); + value_t sin_1 = raft::sin(0.5 * (x2 - y2)); + value_t rdist = sin_0 * sin_0 + raft::cos(x1) * raft::cos(y1) * sin_1 * sin_1; - return 2 * asin(sqrt(rdist)); + return 2 * raft::asin(raft::sqrt(rdist)); } /** diff --git a/cpp/include/raft/spectral/detail/spectral_util.cuh b/cpp/include/raft/spectral/detail/spectral_util.cuh index 3a0ad1f96f..5991e71ec6 100644 --- a/cpp/include/raft/spectral/detail/spectral_util.cuh +++ b/cpp/include/raft/spectral/detail/spectral_util.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2022, NVIDIA CORPORATION. + * Copyright (c) 2020-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -72,7 +72,7 @@ static __global__ void scale_obs_kernel(index_type_t m, index_type_t n, value_ty // scale by alpha alpha = __shfl_sync(warp_full_mask(), alpha, blockDim.x - 1, blockDim.x); - alpha = std::sqrt(alpha); + alpha = raft::sqrt(alpha); for (j = threadIdx.y + blockIdx.y * blockDim.y; j < n; j += blockDim.y * gridDim.y) { for (i = threadIdx.x; i < m; i += blockDim.x) { // blockDim.x=32 index = i + j * m; diff --git a/cpp/include/raft/stats/detail/stddev.cuh b/cpp/include/raft/stats/detail/stddev.cuh index ccea2ea5da..2f7e22ca8a 100644 --- a/cpp/include/raft/stats/detail/stddev.cuh +++ b/cpp/include/raft/stats/detail/stddev.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021-2022, NVIDIA CORPORATION. + * Copyright (c) 2021-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -63,7 +63,7 @@ __global__ void stddevKernelColMajor( thread_data += diff * diff; } Type acc = BlockReduce(temp_storage).Sum(thread_data); - if (threadIdx.x == 0) { std[blockIdx.x] = raft::mySqrt(acc / N); } + if (threadIdx.x == 0) { std[blockIdx.x] = raft::sqrt(acc / N); } } template @@ -126,7 +126,7 @@ void stddev(Type* std, std, mu, D, - [ratio] __device__(Type a, Type b) { return raft::mySqrt(a * ratio - b * b); }, + [ratio] __device__(Type a, Type b) { return raft::sqrt(a * ratio - b * b); }, stream); } else { stddevKernelColMajor<<>>(std, data, mu, D, N); diff --git a/cpp/include/raft/util/cuda_utils.cuh b/cpp/include/raft/util/cuda_utils.cuh index 61dd6e0ad8..5be9dc999a 100644 --- a/cpp/include/raft/util/cuda_utils.cuh +++ b/cpp/include/raft/util/cuda_utils.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2022, NVIDIA CORPORATION. + * Copyright (c) 2018-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -21,6 +21,7 @@ #include #include +#include #include #ifndef ENABLE_MEMCPY_ASYNC @@ -259,12 +260,14 @@ DI double myAtomicMax(double* address, double val) template HDI T myMax(T x, T y); template <> -HDI float myMax(float x, float y) +[[deprecated("use raft::max from raft/core/math.hpp instead")]] HDI float myMax(float x, + float y) { return fmaxf(x, y); } template <> -HDI double myMax(double x, double y) +[[deprecated("use raft::max from raft/core/math.hpp instead")]] HDI double myMax(double x, + double y) { return fmax(x, y); } @@ -277,12 +280,14 @@ HDI double myMax(double x, double y) template HDI T myMin(T x, T y); template <> -HDI float myMin(float x, float y) +[[deprecated("use raft::min from raft/core/math.hpp instead")]] HDI float myMin(float x, + float y) { return fminf(x, y); } template <> -HDI double myMin(double x, double y) +[[deprecated("use raft::min from raft/core/math.hpp instead")]] HDI double myMin(double x, + double y) { return fmin(x, y); } @@ -298,7 +303,7 @@ HDI double myMin(double x, double y) template DI T myAtomicMin(T* address, T val) { - myAtomicReduce(address, val, myMin); + myAtomicReduce(address, val, raft::min_op{}); return *address; } @@ -312,19 +317,10 @@ DI T myAtomicMin(T* address, T val) template DI T myAtomicMax(T* address, T val) { - myAtomicReduce(address, val, myMax); + myAtomicReduce(address, val, raft::max_op{}); return *address; } -/** - * Sign function - */ -template -HDI int sgn(const T val) -{ - return (T(0) < val) - (val < T(0)); -} - /** * @defgroup Exp Exponential function * @{ @@ -332,14 +328,14 @@ HDI int sgn(const T val) template HDI T myExp(T x); template <> -HDI float myExp(float x) +[[deprecated("use raft::exp from raft/core/math.hpp instead")]] HDI float myExp(float x) { return expf(x); } template <> -HDI double myExp(double x) +[[deprecated("use raft::exp from raft/core/math.hpp instead")]] HDI double myExp(double x) { - return exp(x); + return ::exp(x); } /** @} */ @@ -368,14 +364,14 @@ inline __device__ double myInf() template HDI T myLog(T x); template <> -HDI float myLog(float x) +[[deprecated("use raft::log from raft/core/math.hpp instead")]] HDI float myLog(float x) { return logf(x); } template <> -HDI double myLog(double x) +[[deprecated("use raft::log from raft/core/math.hpp instead")]] HDI double myLog(double x) { - return log(x); + return ::log(x); } /** @} */ @@ -386,14 +382,14 @@ HDI double myLog(double x) template HDI T mySqrt(T x); template <> -HDI float mySqrt(float x) +[[deprecated("use raft::sqrt from raft/core/math.hpp instead")]] HDI float mySqrt(float x) { return sqrtf(x); } template <> -HDI double mySqrt(double x) +[[deprecated("use raft::sqrt from raft/core/math.hpp instead")]] HDI double mySqrt(double x) { - return sqrt(x); + return ::sqrt(x); } /** @} */ @@ -404,14 +400,18 @@ HDI double mySqrt(double x) template DI void mySinCos(T x, T& s, T& c); template <> -DI void mySinCos(float x, float& s, float& c) +[[deprecated("use raft::sincos from raft/core/math.hpp instead")]] DI void mySinCos(float x, + float& s, + float& c) { sincosf(x, &s, &c); } template <> -DI void mySinCos(double x, double& s, double& c) +[[deprecated("use raft::sincos from raft/core/math.hpp instead")]] DI void mySinCos(double x, + double& s, + double& c) { - sincos(x, &s, &c); + ::sincos(x, &s, &c); } /** @} */ @@ -422,14 +422,14 @@ DI void mySinCos(double x, double& s, double& c) template DI T mySin(T x); template <> -DI float mySin(float x) +[[deprecated("use raft::sin from raft/core/math.hpp instead")]] DI float mySin(float x) { return sinf(x); } template <> -DI double mySin(double x) +[[deprecated("use raft::sin from raft/core/math.hpp instead")]] DI double mySin(double x) { - return sin(x); + return ::sin(x); } /** @} */ @@ -443,12 +443,12 @@ DI T myAbs(T x) return x < 0 ? -x : x; } template <> -DI float myAbs(float x) +[[deprecated("use raft::abs from raft/core/math.hpp instead")]] DI float myAbs(float x) { return fabsf(x); } template <> -DI double myAbs(double x) +[[deprecated("use raft::abs from raft/core/math.hpp instead")]] DI double myAbs(double x) { return fabs(x); } @@ -461,14 +461,16 @@ DI double myAbs(double x) template HDI T myPow(T x, T power); template <> -HDI float myPow(float x, float power) +[[deprecated("use raft::pow from raft/core/math.hpp instead")]] HDI float myPow(float x, + float power) { return powf(x, power); } template <> -HDI double myPow(double x, double power) +[[deprecated("use raft::pow from raft/core/math.hpp instead")]] HDI double myPow(double x, + double power) { - return pow(x, power); + return ::pow(x, power); } /** @} */ @@ -479,14 +481,14 @@ HDI double myPow(double x, double power) template HDI T myTanh(T x); template <> -HDI float myTanh(float x) +[[deprecated("use raft::tanh from raft/core/math.hpp instead")]] HDI float myTanh(float x) { return tanhf(x); } template <> -HDI double myTanh(double x) +[[deprecated("use raft::tanh from raft/core/math.hpp instead")]] HDI double myTanh(double x) { - return tanh(x); + return ::tanh(x); } /** @} */ @@ -497,14 +499,14 @@ HDI double myTanh(double x) template HDI T myATanh(T x); template <> -HDI float myATanh(float x) +[[deprecated("use raft::atanh from raft/core/math.hpp instead")]] HDI float myATanh(float x) { return atanhf(x); } template <> -HDI double myATanh(double x) +[[deprecated("use raft::atanh from raft/core/math.hpp instead")]] HDI double myATanh(double x) { - return atanh(x); + return ::atanh(x); } /** @} */ @@ -526,7 +528,7 @@ struct SqrtOp { [[deprecated("SqrtOp is deprecated. Use sqrt_op instead.")]] HDI Type operator()(Type in, IdxType i = 0) const { - return mySqrt(in); + return raft::sqrt(in); } }; @@ -544,7 +546,7 @@ struct L1Op { [[deprecated("L1Op is deprecated. Use abs_op instead.")]] HDI Type operator()(Type in, IdxType i = 0) const { - return myAbs(in); + return raft::abs(in); } }; diff --git a/cpp/test/CMakeLists.txt b/cpp/test/CMakeLists.txt index 8ca30a5c82..a4b3758faa 100644 --- a/cpp/test/CMakeLists.txt +++ b/cpp/test/CMakeLists.txt @@ -86,6 +86,8 @@ if(BUILD_TESTS) CORE_TEST PATH test/core/logger.cpp + test/core/math_device.cu + test/core/math_host.cpp test/core/operators_device.cu test/core/operators_host.cpp test/core/handle.cpp diff --git a/cpp/test/core/math_device.cu b/cpp/test/core/math_device.cu new file mode 100644 index 0000000000..ff4b343d9e --- /dev/null +++ b/cpp/test/core/math_device.cu @@ -0,0 +1,352 @@ +/* + * Copyright (c) 2023, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include + +#include "../test_utils.h" +#include +#include +#include + +template +__global__ void math_eval_kernel(OutT* out, OpT op, Args... args) +{ + out[0] = op(std::forward(args)...); +} + +template +auto math_eval(OpT op, Args&&... args) +{ + typedef decltype(op(args...)) OutT; + auto stream = rmm::cuda_stream_default; + rmm::device_scalar result(stream); + math_eval_kernel<<<1, 1, 0, stream>>>(result.data(), op, std::forward(args)...); + return result.value(stream); +} + +struct abs_test_op { + template + constexpr RAFT_INLINE_FUNCTION auto operator()(const Type& in) const + { + return raft::abs(in); + } +}; + +TEST(MathDevice, Abs) +{ + // Integer abs + ASSERT_TRUE( + raft::match(int8_t{123}, math_eval(abs_test_op{}, int8_t{-123}), raft::Compare())); + ASSERT_TRUE(raft::match(12345, math_eval(abs_test_op{}, -12345), raft::Compare())); + ASSERT_TRUE(raft::match(12345l, math_eval(abs_test_op{}, -12345l), raft::Compare())); + ASSERT_TRUE(raft::match(123451234512345ll, + math_eval(abs_test_op{}, -123451234512345ll), + raft::Compare())); + // Floating-point abs + ASSERT_TRUE( + raft::match(12.34f, math_eval(abs_test_op{}, -12.34f), raft::CompareApprox(0.0001f))); + ASSERT_TRUE( + raft::match(12.34, math_eval(abs_test_op{}, -12.34), raft::CompareApprox(0.000001))); +} + +struct acos_test_op { + template + constexpr RAFT_INLINE_FUNCTION auto operator()(const Type& in) const + { + return raft::acos(in); + } +}; + +TEST(MathDevice, Acos) +{ + ASSERT_TRUE(raft::match( + std::acos(0.123f), math_eval(acos_test_op{}, 0.123f), raft::CompareApprox(0.0001f))); + ASSERT_TRUE(raft::match( + std::acos(0.123), math_eval(acos_test_op{}, 0.123), raft::CompareApprox(0.000001))); +} + +struct asin_test_op { + template + constexpr RAFT_INLINE_FUNCTION auto operator()(const Type& in) const + { + return raft::asin(in); + } +}; + +TEST(MathDevice, Asin) +{ + ASSERT_TRUE(raft::match( + std::asin(0.123f), math_eval(asin_test_op{}, 0.123f), raft::CompareApprox(0.0001f))); + ASSERT_TRUE(raft::match( + std::asin(0.123), math_eval(asin_test_op{}, 0.123), raft::CompareApprox(0.000001))); +} + +struct atanh_test_op { + template + constexpr RAFT_INLINE_FUNCTION auto operator()(const Type& in) const + { + return raft::atanh(in); + } +}; + +TEST(MathDevice, Atanh) +{ + ASSERT_TRUE(raft::match( + std::atanh(0.123f), math_eval(atanh_test_op{}, 0.123f), raft::CompareApprox(0.0001f))); + ASSERT_TRUE(raft::match( + std::atanh(0.123), math_eval(atanh_test_op{}, 0.123), raft::CompareApprox(0.000001))); +} + +struct cos_test_op { + template + constexpr RAFT_INLINE_FUNCTION auto operator()(const Type& in) const + { + return raft::cos(in); + } +}; + +TEST(MathDevice, Cos) +{ + ASSERT_TRUE(raft::match( + std::cos(12.34f), math_eval(cos_test_op{}, 12.34f), raft::CompareApprox(0.0001f))); + ASSERT_TRUE(raft::match( + std::cos(12.34), math_eval(cos_test_op{}, 12.34), raft::CompareApprox(0.000001))); +} + +struct exp_test_op { + template + constexpr RAFT_INLINE_FUNCTION auto operator()(const Type& in) const + { + return raft::exp(in); + } +}; + +TEST(MathDevice, Exp) +{ + ASSERT_TRUE(raft::match( + std::exp(12.34f), math_eval(exp_test_op{}, 12.34f), raft::CompareApprox(0.0001f))); + ASSERT_TRUE(raft::match( + std::exp(12.34), math_eval(exp_test_op{}, 12.34), raft::CompareApprox(0.000001))); +} + +struct log_test_op { + template + constexpr RAFT_INLINE_FUNCTION auto operator()(const Type& in) const + { + return raft::log(in); + } +}; + +TEST(MathDevice, Log) +{ + ASSERT_TRUE(raft::match( + std::log(12.34f), math_eval(log_test_op{}, 12.34f), raft::CompareApprox(0.0001f))); + ASSERT_TRUE(raft::match( + std::log(12.34), math_eval(log_test_op{}, 12.34), raft::CompareApprox(0.000001))); +} + +struct max_test_op { + template + constexpr RAFT_INLINE_FUNCTION auto operator()(Args&&... args) const + { + return raft::max(std::forward(args)...); + } +}; + +TEST(MathDevice, Max2) +{ + ASSERT_TRUE(raft::match(1234, math_eval(max_test_op{}, -1234, 1234), raft::Compare())); + ASSERT_TRUE( + raft::match(1234u, math_eval(max_test_op{}, 1234u, 123u), raft::Compare())); + ASSERT_TRUE( + raft::match(1234ll, math_eval(max_test_op{}, -1234ll, 1234ll), raft::Compare())); + ASSERT_TRUE(raft::match( + 1234ull, math_eval(max_test_op{}, 1234ull, 123ull), raft::Compare())); + + ASSERT_TRUE( + raft::match(12.34f, math_eval(max_test_op{}, -12.34f, 12.34f), raft::Compare())); + ASSERT_TRUE(raft::match(12.34, math_eval(max_test_op{}, -12.34, 12.34), raft::Compare())); + ASSERT_TRUE(raft::match( + 12.34, math_eval(max_test_op{}, -12.34f, 12.34), raft::CompareApprox(0.000001))); + ASSERT_TRUE(raft::match( + 12.34, math_eval(max_test_op{}, -12.34, 12.34f), raft::CompareApprox(0.000001))); +} + +TEST(MathDevice, Max3) +{ + ASSERT_TRUE(raft::match(1234, math_eval(max_test_op{}, 1234, 0, -1234), raft::Compare())); + ASSERT_TRUE(raft::match(1234, math_eval(max_test_op{}, -1234, 1234, 0), raft::Compare())); + ASSERT_TRUE(raft::match(1234, math_eval(max_test_op{}, 0, -1234, 1234), raft::Compare())); + + ASSERT_TRUE(raft::match( + 12.34, math_eval(max_test_op{}, 12.34f, 0., -12.34), raft::CompareApprox(0.000001))); + ASSERT_TRUE(raft::match( + 12.34, math_eval(max_test_op{}, -12.34, 12.34f, 0.), raft::CompareApprox(0.000001))); + ASSERT_TRUE(raft::match( + 12.34, math_eval(max_test_op{}, 0., -12.34, 12.34f), raft::CompareApprox(0.000001))); +} + +struct min_test_op { + template + constexpr RAFT_INLINE_FUNCTION auto operator()(Args&&... args) const + { + return raft::min(std::forward(args)...); + } +}; + +TEST(MathDevice, Min2) +{ + ASSERT_TRUE(raft::match(-1234, math_eval(min_test_op{}, -1234, 1234), raft::Compare())); + ASSERT_TRUE( + raft::match(123u, math_eval(min_test_op{}, 1234u, 123u), raft::Compare())); + ASSERT_TRUE(raft::match( + -1234ll, math_eval(min_test_op{}, -1234ll, 1234ll), raft::Compare())); + ASSERT_TRUE(raft::match( + 123ull, math_eval(min_test_op{}, 1234ull, 123ull), raft::Compare())); + + ASSERT_TRUE( + raft::match(-12.34f, math_eval(min_test_op{}, -12.34f, 12.34f), raft::Compare())); + ASSERT_TRUE( + raft::match(-12.34, math_eval(min_test_op{}, -12.34, 12.34), raft::Compare())); + ASSERT_TRUE(raft::match( + -12.34, math_eval(min_test_op{}, -12.34f, 12.34), raft::CompareApprox(0.000001))); + ASSERT_TRUE(raft::match( + -12.34, math_eval(min_test_op{}, -12.34, 12.34f), raft::CompareApprox(0.000001))); +} + +TEST(MathDevice, Min3) +{ + ASSERT_TRUE(raft::match(-1234, math_eval(min_test_op{}, 1234, 0, -1234), raft::Compare())); + ASSERT_TRUE(raft::match(-1234, math_eval(min_test_op{}, -1234, 1234, 0), raft::Compare())); + ASSERT_TRUE(raft::match(-1234, math_eval(min_test_op{}, 0, -1234, 1234), raft::Compare())); + + ASSERT_TRUE(raft::match( + -12.34, math_eval(min_test_op{}, 12.34f, 0., -12.34), raft::CompareApprox(0.000001))); + ASSERT_TRUE(raft::match( + -12.34, math_eval(min_test_op{}, -12.34, 12.34f, 0.), raft::CompareApprox(0.000001))); + ASSERT_TRUE(raft::match( + -12.34, math_eval(min_test_op{}, 0., -12.34, 12.34f), raft::CompareApprox(0.000001))); +} + +struct pow_test_op { + template + constexpr RAFT_INLINE_FUNCTION auto operator()(const Type& x, const Type& y) const + { + return raft::pow(x, y); + } +}; + +TEST(MathDevice, Pow) +{ + ASSERT_TRUE(raft::match(std::pow(12.34f, 2.f), + math_eval(pow_test_op{}, 12.34f, 2.f), + raft::CompareApprox(0.0001f))); + ASSERT_TRUE(raft::match(std::pow(12.34, 2.), + math_eval(pow_test_op{}, 12.34, 2.), + raft::CompareApprox(0.000001))); +} + +struct sgn_test_op { + template + constexpr RAFT_INLINE_FUNCTION auto operator()(const Type& in) const + { + return raft::sgn(in); + } +}; + +TEST(MathDevice, Sgn) +{ + ASSERT_TRUE(raft::match(-1, math_eval(sgn_test_op{}, -1234), raft::Compare())); + ASSERT_TRUE(raft::match(0, math_eval(sgn_test_op{}, 0), raft::Compare())); + ASSERT_TRUE(raft::match(1, math_eval(sgn_test_op{}, 1234), raft::Compare())); + ASSERT_TRUE(raft::match(-1, math_eval(sgn_test_op{}, -12.34f), raft::Compare())); + ASSERT_TRUE(raft::match(0, math_eval(sgn_test_op{}, 0.f), raft::Compare())); + ASSERT_TRUE(raft::match(1, math_eval(sgn_test_op{}, 12.34f), raft::Compare())); +} + +struct sin_test_op { + template + constexpr RAFT_INLINE_FUNCTION auto operator()(const Type& in) const + { + return raft::sin(in); + } +}; + +TEST(MathDevice, Sin) +{ + ASSERT_TRUE(raft::match( + std::sin(12.34f), math_eval(sin_test_op{}, 12.34f), raft::CompareApprox(0.0001f))); + ASSERT_TRUE(raft::match( + std::sin(12.34), math_eval(sin_test_op{}, 12.34), raft::CompareApprox(0.000001))); +} + +struct sincos_test_op { + template + constexpr RAFT_INLINE_FUNCTION auto operator()(const Type& x, Type* s, Type* c) const + { + raft::sincos(x, s, c); + return x; // unused, just to avoid creating another helper + } +}; + +TEST(MathDevice, SinCos) +{ + auto stream = rmm::cuda_stream_default; + float xf = 12.34f; + rmm::device_scalar sf(stream); + rmm::device_scalar cf(stream); + math_eval(sincos_test_op{}, xf, sf.data(), cf.data()); + ASSERT_TRUE(raft::match(std::sin(12.34f), sf.value(stream), raft::CompareApprox(0.0001f))); + ASSERT_TRUE(raft::match(std::cos(12.34f), cf.value(stream), raft::CompareApprox(0.0001f))); + double xd = 12.34f; + rmm::device_scalar sd(stream); + rmm::device_scalar cd(stream); + math_eval(sincos_test_op{}, xd, sd.data(), cd.data()); + ASSERT_TRUE(raft::match(std::sin(12.34), sd.value(stream), raft::CompareApprox(0.0001f))); + ASSERT_TRUE(raft::match(std::cos(12.34), cd.value(stream), raft::CompareApprox(0.0001f))); +} + +struct sqrt_test_op { + template + constexpr RAFT_INLINE_FUNCTION auto operator()(const Type& in) const + { + return raft::sqrt(in); + } +}; + +TEST(MathDevice, Sqrt) +{ + ASSERT_TRUE(raft::match( + std::sqrt(12.34f), math_eval(sqrt_test_op{}, 12.34f), raft::CompareApprox(0.0001f))); + ASSERT_TRUE(raft::match( + std::sqrt(12.34), math_eval(sqrt_test_op{}, 12.34), raft::CompareApprox(0.000001))); +} + +struct tanh_test_op { + template + constexpr RAFT_INLINE_FUNCTION auto operator()(const Type& in) const + { + return raft::tanh(in); + } +}; + +TEST(MathDevice, Tanh) +{ + ASSERT_TRUE(raft::match( + std::tanh(12.34f), math_eval(tanh_test_op{}, 12.34f), raft::CompareApprox(0.0001f))); + ASSERT_TRUE(raft::match( + std::tanh(12.34), math_eval(tanh_test_op{}, 12.34), raft::CompareApprox(0.000001))); +} diff --git a/cpp/test/core/math_host.cpp b/cpp/test/core/math_host.cpp new file mode 100644 index 0000000000..5808905713 --- /dev/null +++ b/cpp/test/core/math_host.cpp @@ -0,0 +1,195 @@ +/* + * Copyright (c) 2023, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include + +#include "../test_utils.h" +#include + +TEST(MathHost, Abs) +{ + // Integer abs + ASSERT_TRUE(raft::match(int8_t{123}, raft::abs(int8_t{-123}), raft::Compare())); + ASSERT_TRUE(raft::match(12345, raft::abs(-12345), raft::Compare())); + ASSERT_TRUE(raft::match(12345l, raft::abs(-12345l), raft::Compare())); + ASSERT_TRUE( + raft::match(123451234512345ll, raft::abs(-123451234512345ll), raft::Compare())); + // Floating-point abs + ASSERT_TRUE(raft::match(12.34f, raft::abs(-12.34f), raft::CompareApprox(0.0001f))); + ASSERT_TRUE(raft::match(12.34, raft::abs(-12.34), raft::CompareApprox(0.000001))); +} + +TEST(MathHost, Acos) +{ + ASSERT_TRUE( + raft::match(std::acos(0.123f), raft::acos(0.123f), raft::CompareApprox(0.0001f))); + ASSERT_TRUE( + raft::match(std::acos(0.123), raft::acos(0.123), raft::CompareApprox(0.000001))); +} + +TEST(MathHost, Asin) +{ + ASSERT_TRUE( + raft::match(std::asin(0.123f), raft::asin(0.123f), raft::CompareApprox(0.0001f))); + ASSERT_TRUE( + raft::match(std::asin(0.123), raft::asin(0.123), raft::CompareApprox(0.000001))); +} + +TEST(MathHost, Atanh) +{ + ASSERT_TRUE( + raft::match(std::atanh(0.123f), raft::atanh(0.123f), raft::CompareApprox(0.0001f))); + ASSERT_TRUE( + raft::match(std::atanh(0.123), raft::atanh(0.123), raft::CompareApprox(0.000001))); +} + +TEST(MathHost, Cos) +{ + ASSERT_TRUE( + raft::match(std::cos(12.34f), raft::cos(12.34f), raft::CompareApprox(0.0001f))); + ASSERT_TRUE( + raft::match(std::cos(12.34), raft::cos(12.34), raft::CompareApprox(0.000001))); +} + +TEST(MathHost, Exp) +{ + ASSERT_TRUE( + raft::match(std::exp(12.34f), raft::exp(12.34f), raft::CompareApprox(0.0001f))); + ASSERT_TRUE( + raft::match(std::exp(12.34), raft::exp(12.34), raft::CompareApprox(0.000001))); +} + +TEST(MathHost, Log) +{ + ASSERT_TRUE( + raft::match(std::log(12.34f), raft::log(12.34f), raft::CompareApprox(0.0001f))); + ASSERT_TRUE( + raft::match(std::log(12.34), raft::log(12.34), raft::CompareApprox(0.000001))); +} + +TEST(MathHost, Max2) +{ + ASSERT_TRUE(raft::match(1234, raft::max(-1234, 1234), raft::Compare())); + ASSERT_TRUE(raft::match(1234u, raft::max(1234u, 123u), raft::Compare())); + ASSERT_TRUE(raft::match(1234ll, raft::max(-1234ll, 1234ll), raft::Compare())); + ASSERT_TRUE( + raft::match(1234ull, raft::max(1234ull, 123ull), raft::Compare())); + + ASSERT_TRUE(raft::match(12.34f, raft::max(-12.34f, 12.34f), raft::Compare())); + ASSERT_TRUE(raft::match(12.34, raft::max(-12.34, 12.34), raft::Compare())); + ASSERT_TRUE(raft::match(12.34, raft::max(-12.34f, 12.34), raft::CompareApprox(0.000001))); + ASSERT_TRUE(raft::match(12.34, raft::max(-12.34, 12.34f), raft::CompareApprox(0.000001))); +} + +TEST(MathHost, Max3) +{ + ASSERT_TRUE(raft::match(1234, raft::max(1234, 0, -1234), raft::Compare())); + ASSERT_TRUE(raft::match(1234, raft::max(-1234, 1234, 0), raft::Compare())); + ASSERT_TRUE(raft::match(1234, raft::max(0, -1234, 1234), raft::Compare())); + + ASSERT_TRUE( + raft::match(12.34, raft::max(12.34f, 0., -12.34), raft::CompareApprox(0.000001))); + ASSERT_TRUE( + raft::match(12.34, raft::max(-12.34, 12.34f, 0.), raft::CompareApprox(0.000001))); + ASSERT_TRUE( + raft::match(12.34, raft::max(0., -12.34, 12.34f), raft::CompareApprox(0.000001))); +} + +TEST(MathHost, Min2) +{ + ASSERT_TRUE(raft::match(-1234, raft::min(-1234, 1234), raft::Compare())); + ASSERT_TRUE(raft::match(123u, raft::min(1234u, 123u), raft::Compare())); + ASSERT_TRUE(raft::match(-1234ll, raft::min(-1234ll, 1234ll), raft::Compare())); + ASSERT_TRUE( + raft::match(123ull, raft::min(1234ull, 123ull), raft::Compare())); + + ASSERT_TRUE(raft::match(-12.34f, raft::min(-12.34f, 12.34f), raft::Compare())); + ASSERT_TRUE(raft::match(-12.34, raft::min(-12.34, 12.34), raft::Compare())); + ASSERT_TRUE( + raft::match(-12.34, raft::min(-12.34f, 12.34), raft::CompareApprox(0.000001))); + ASSERT_TRUE( + raft::match(-12.34, raft::min(-12.34, 12.34f), raft::CompareApprox(0.000001))); +} + +TEST(MathHost, Min3) +{ + ASSERT_TRUE(raft::match(-1234, raft::min(1234, 0, -1234), raft::Compare())); + ASSERT_TRUE(raft::match(-1234, raft::min(-1234, 1234, 0), raft::Compare())); + ASSERT_TRUE(raft::match(-1234, raft::min(0, -1234, 1234), raft::Compare())); + + ASSERT_TRUE( + raft::match(-12.34, raft::min(12.34f, 0., -12.34), raft::CompareApprox(0.000001))); + ASSERT_TRUE( + raft::match(-12.34, raft::min(-12.34, 12.34f, 0.), raft::CompareApprox(0.000001))); + ASSERT_TRUE( + raft::match(-12.34, raft::min(0., -12.34, 12.34f), raft::CompareApprox(0.000001))); +} + +TEST(MathHost, Pow) +{ + ASSERT_TRUE(raft::match( + std::pow(12.34f, 2.f), raft::pow(12.34f, 2.f), raft::CompareApprox(0.0001f))); + ASSERT_TRUE( + raft::match(std::pow(12.34, 2.), raft::pow(12.34, 2.), raft::CompareApprox(0.000001))); +} + +TEST(MathHost, Sgn) +{ + ASSERT_TRUE(raft::match(-1, raft::sgn(-1234), raft::Compare())); + ASSERT_TRUE(raft::match(0, raft::sgn(0), raft::Compare())); + ASSERT_TRUE(raft::match(1, raft::sgn(1234), raft::Compare())); + ASSERT_TRUE(raft::match(-1, raft::sgn(-12.34f), raft::Compare())); + ASSERT_TRUE(raft::match(0, raft::sgn(0.f), raft::Compare())); + ASSERT_TRUE(raft::match(1, raft::sgn(12.34f), raft::Compare())); +} + +TEST(MathHost, Sin) +{ + ASSERT_TRUE( + raft::match(std::sin(12.34f), raft::sin(12.34f), raft::CompareApprox(0.0001f))); + ASSERT_TRUE( + raft::match(std::sin(12.34), raft::sin(12.34), raft::CompareApprox(0.000001))); +} + +TEST(MathHost, SinCos) +{ + float xf = 12.34f; + float sf, cf; + raft::sincos(xf, &sf, &cf); + ASSERT_TRUE(raft::match(std::sin(12.34f), sf, raft::CompareApprox(0.0001f))); + ASSERT_TRUE(raft::match(std::cos(12.34f), cf, raft::CompareApprox(0.0001f))); + double xd = 12.34f; + double sd, cd; + raft::sincos(xd, &sd, &cd); + ASSERT_TRUE(raft::match(std::sin(12.34), sd, raft::CompareApprox(0.000001))); + ASSERT_TRUE(raft::match(std::cos(12.34), cd, raft::CompareApprox(0.000001))); +} + +TEST(MathHost, Sqrt) +{ + ASSERT_TRUE( + raft::match(std::sqrt(12.34f), raft::sqrt(12.34f), raft::CompareApprox(0.0001f))); + ASSERT_TRUE( + raft::match(std::sqrt(12.34), raft::sqrt(12.34), raft::CompareApprox(0.000001))); +} + +TEST(MathHost, Tanh) +{ + ASSERT_TRUE( + raft::match(std::tanh(12.34f), raft::tanh(12.34f), raft::CompareApprox(0.0001f))); + ASSERT_TRUE( + raft::match(std::tanh(12.34), raft::tanh(12.34), raft::CompareApprox(0.000001))); +} diff --git a/cpp/test/distance/distance_base.cuh b/cpp/test/distance/distance_base.cuh index cbfd97ebc6..fedbee919d 100644 --- a/cpp/test/distance/distance_base.cuh +++ b/cpp/test/distance/distance_base.cuh @@ -52,7 +52,7 @@ __global__ void naiveDistanceKernel(DataType* dist, } if (type == raft::distance::DistanceType::L2SqrtExpanded || type == raft::distance::DistanceType::L2SqrtUnexpanded) - acc = raft::mySqrt(acc); + acc = raft::sqrt(acc); int outidx = isRowMajor ? midx * n + nidx : midx + m * nidx; dist[outidx] = acc; } @@ -79,9 +79,9 @@ __global__ void naiveL1_Linf_CanberraDistanceKernel(DataType* dist, auto b = y[yidx]; auto diff = (a > b) ? (a - b) : (b - a); if (type == raft::distance::DistanceType::Linf) { - acc = raft::myMax(acc, diff); + acc = raft::max(acc, diff); } else if (type == raft::distance::DistanceType::Canberra) { - const auto add = raft::myAbs(a) + raft::myAbs(b); + const auto add = raft::abs(a) + raft::abs(b); // deal with potential for 0 in denominator by // forcing 1/0 instead acc += ((add != 0) * diff / (add + (add == 0))); @@ -119,7 +119,7 @@ __global__ void naiveCosineDistanceKernel( int outidx = isRowMajor ? midx * n + nidx : midx + m * nidx; // Use 1.0 - (cosine similarity) to calc the distance - dist[outidx] = (DataType)1.0 - acc_ab / (raft::mySqrt(acc_a) * raft::mySqrt(acc_b)); + dist[outidx] = (DataType)1.0 - acc_ab / (raft::sqrt(acc_a) * raft::sqrt(acc_b)); } template @@ -137,7 +137,7 @@ __global__ void naiveHellingerDistanceKernel( int yidx = isRowMajor ? i + nidx * k : i * n + nidx; auto a = x[xidx]; auto b = y[yidx]; - acc_ab += raft::mySqrt(a) * raft::mySqrt(b); + acc_ab += raft::sqrt(a) * raft::sqrt(b); } int outidx = isRowMajor ? midx * n + nidx : midx + m * nidx; @@ -145,7 +145,7 @@ __global__ void naiveHellingerDistanceKernel( // Adjust to replace NaN in sqrt with 0 if input to sqrt is negative acc_ab = 1 - acc_ab; auto rectifier = (!signbit(acc_ab)); - dist[outidx] = raft::mySqrt(rectifier * acc_ab); + dist[outidx] = raft::sqrt(rectifier * acc_ab); } template @@ -167,11 +167,11 @@ __global__ void naiveLpUnexpDistanceKernel(DataType* dist, int yidx = isRowMajor ? i + nidx * k : i * n + nidx; auto a = x[xidx]; auto b = y[yidx]; - auto diff = raft::myAbs(a - b); - acc += raft::myPow(diff, p); + auto diff = raft::abs(a - b); + acc += raft::pow(diff, p); } auto one_over_p = 1 / p; - acc = raft::myPow(acc, one_over_p); + acc = raft::pow(acc, one_over_p); int outidx = isRowMajor ? midx * n + nidx : midx + m * nidx; dist[outidx] = acc; } @@ -222,7 +222,7 @@ __global__ void naiveJensenShannonDistanceKernel( acc += (-a * (!p_zero * log(p + p_zero))) + (-b * (!q_zero * log(q + q_zero))); } - acc = raft::mySqrt(0.5f * acc); + acc = raft::sqrt(0.5f * acc); int outidx = isRowMajor ? midx * n + nidx : midx + m * nidx; dist[outidx] = acc; } @@ -297,7 +297,7 @@ __global__ void naiveCorrelationDistanceKernel( auto Q_denom = k * a_sq_norm - (a_norm * a_norm); auto R_denom = k * b_sq_norm - (b_norm * b_norm); - acc = 1 - (numer / raft::mySqrt(Q_denom * R_denom)); + acc = 1 - (numer / raft::sqrt(Q_denom * R_denom)); int outidx = isRowMajor ? midx * n + nidx : midx + m * nidx; dist[outidx] = acc; diff --git a/cpp/test/distance/fused_l2_nn.cu b/cpp/test/distance/fused_l2_nn.cu index e746a2382d..54de12307a 100644 --- a/cpp/test/distance/fused_l2_nn.cu +++ b/cpp/test/distance/fused_l2_nn.cu @@ -60,7 +60,7 @@ __global__ void naiveKernel(raft::KeyValuePair* min, auto diff = midx >= m || nidx >= n ? DataT(0) : x[xidx] - y[yidx]; acc += diff * diff; } - if (Sqrt) { acc = raft::mySqrt(acc); } + if (Sqrt) { acc = raft::sqrt(acc); } ReduceOpT redOp; typedef cub::WarpReduce> WarpReduce; __shared__ typename WarpReduce::TempStorage temp[NWARPS]; diff --git a/cpp/test/linalg/matrix_vector.cu b/cpp/test/linalg/matrix_vector.cu index 7018e1da96..fb1e2235f9 100644 --- a/cpp/test/linalg/matrix_vector.cu +++ b/cpp/test/linalg/matrix_vector.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022, NVIDIA CORPORATION. + * Copyright (c) 2022-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -116,7 +116,7 @@ void naive_matrix_vector_op_launch(const raft::handle_t& handle, } }; auto operation_bin_div_skip_zero = [] __device__(T mat_element, T vec_element) { - if (raft::myAbs(vec_element) < T(1e-10)) + if (raft::abs(vec_element) < T(1e-10)) return T(0); else return mat_element / vec_element; diff --git a/cpp/test/linalg/norm.cu b/cpp/test/linalg/norm.cu index 94540b9ff6..90cfbd8f89 100644 --- a/cpp/test/linalg/norm.cu +++ b/cpp/test/linalg/norm.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2022, NVIDIA CORPORATION. + * Copyright (c) 2018-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -56,10 +56,10 @@ __global__ void naiveRowNormKernel( if (type == L2Norm) { acc += data[rowStart * D + i] * data[rowStart * D + i]; } else { - acc += raft::myAbs(data[rowStart * D + i]); + acc += raft::abs(data[rowStart * D + i]); } } - dots[rowStart] = do_sqrt ? raft::mySqrt(acc) : acc; + dots[rowStart] = do_sqrt ? raft::sqrt(acc) : acc; } } @@ -131,10 +131,10 @@ __global__ void naiveColNormKernel( Type acc = 0; for (IdxT i = 0; i < N; i++) { Type v = data[colID + i * D]; - acc += type == L2Norm ? v * v : raft::myAbs(v); + acc += type == L2Norm ? v * v : raft::abs(v); } - dots[colID] = do_sqrt ? raft::mySqrt(acc) : acc; + dots[colID] = do_sqrt ? raft::sqrt(acc) : acc; } template diff --git a/cpp/test/linalg/power.cu b/cpp/test/linalg/power.cu index 54c2e2a7aa..5cb63a5697 100644 --- a/cpp/test/linalg/power.cu +++ b/cpp/test/linalg/power.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2022, NVIDIA CORPORATION. + * Copyright (c) 2018-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -27,7 +27,7 @@ template __global__ void naivePowerElemKernel(Type* out, const Type* in1, const Type* in2, int len) { int idx = threadIdx.x + blockIdx.x * blockDim.x; - if (idx < len) { out[idx] = raft::myPow(in1[idx], in2[idx]); } + if (idx < len) { out[idx] = raft::pow(in1[idx], in2[idx]); } } template @@ -43,7 +43,7 @@ template __global__ void naivePowerScalarKernel(Type* out, const Type* in1, const Type in2, int len) { int idx = threadIdx.x + blockIdx.x * blockDim.x; - if (idx < len) { out[idx] = raft::myPow(in1[idx], in2); } + if (idx < len) { out[idx] = raft::pow(in1[idx], in2); } } template diff --git a/cpp/test/linalg/sqrt.cu b/cpp/test/linalg/sqrt.cu index 9008313b58..93150ca77d 100644 --- a/cpp/test/linalg/sqrt.cu +++ b/cpp/test/linalg/sqrt.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2022, NVIDIA CORPORATION. + * Copyright (c) 2018-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -27,7 +27,7 @@ template __global__ void naiveSqrtElemKernel(Type* out, const Type* in1, int len) { int idx = threadIdx.x + blockIdx.x * blockDim.x; - if (idx < len) { out[idx] = raft::mySqrt(in1[idx]); } + if (idx < len) { out[idx] = raft::sqrt(in1[idx]); } } template diff --git a/cpp/test/matrix/math.cu b/cpp/test/matrix/math.cu index f2c1a6249c..9dcbfc8899 100644 --- a/cpp/test/matrix/math.cu +++ b/cpp/test/matrix/math.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2022, NVIDIA CORPORATION. + * Copyright (c) 2018-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -51,7 +51,7 @@ template __global__ void naiveSqrtKernel(Type* in, Type* out, int len) { int idx = threadIdx.x + blockIdx.x * blockDim.x; - if (idx < len) { out[idx] = std::sqrt(in[idx]); } + if (idx < len) { out[idx] = raft::sqrt(in[idx]); } } template diff --git a/cpp/test/neighbors/ann_utils.cuh b/cpp/test/neighbors/ann_utils.cuh index b88b6abd9e..bb2f334db4 100644 --- a/cpp/test/neighbors/ann_utils.cuh +++ b/cpp/test/neighbors/ann_utils.cuh @@ -131,7 +131,7 @@ __global__ void naive_distance_kernel(EvalT* dist, } if (type == raft::distance::DistanceType::L2SqrtExpanded || type == raft::distance::DistanceType::L2SqrtUnexpanded) - acc = raft::mySqrt(acc); + acc = raft::sqrt(acc); dist[midx * n + nidx] = acc; } } diff --git a/cpp/test/random/rng.cu b/cpp/test/random/rng.cu index bdce79b76e..0bf494b624 100644 --- a/cpp/test/random/rng.cu +++ b/cpp/test/random/rng.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2022, NVIDIA CORPORATION. + * Copyright (c) 2018-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -145,8 +145,8 @@ class RngTest : public ::testing::TestWithParam> { case RNG_LogNormal: { auto var = params.end * params.end; auto mu = params.start; - meanvar[0] = raft::myExp(mu + var * T(0.5)); - meanvar[1] = (raft::myExp(var) - T(1.0)) * raft::myExp(T(2.0) * mu + var); + meanvar[0] = raft::exp(mu + var * T(0.5)); + meanvar[1] = (raft::exp(var) - T(1.0)) * raft::exp(T(2.0) * mu + var); break; } case RNG_Uniform: @@ -169,7 +169,7 @@ class RngTest : public ::testing::TestWithParam> { meanvar[1] = meanvar[0] * meanvar[0]; break; case RNG_Rayleigh: - meanvar[0] = params.start * raft::mySqrt(T(3.1415 / 2.0)); + meanvar[0] = params.start * raft::sqrt(T(3.1415 / 2.0)); meanvar[1] = ((T(4.0) - T(3.1415)) / T(2.0)) * params.start * params.start; break; case RNG_Laplace: @@ -239,8 +239,8 @@ class RngMdspanTest : public ::testing::TestWithParam> { case RNG_LogNormal: { auto var = params.end * params.end; auto mu = params.start; - meanvar[0] = raft::myExp(mu + var * T(0.5)); - meanvar[1] = (raft::myExp(var) - T(1.0)) * raft::myExp(T(2.0) * mu + var); + meanvar[0] = raft::exp(mu + var * T(0.5)); + meanvar[1] = (raft::exp(var) - T(1.0)) * raft::exp(T(2.0) * mu + var); break; } case RNG_Uniform: @@ -263,7 +263,7 @@ class RngMdspanTest : public ::testing::TestWithParam> { meanvar[1] = meanvar[0] * meanvar[0]; break; case RNG_Rayleigh: - meanvar[0] = params.start * raft::mySqrt(T(3.1415 / 2.0)); + meanvar[0] = params.start * raft::sqrt(T(3.1415 / 2.0)); meanvar[1] = ((T(4.0) - T(3.1415)) / T(2.0)) * params.start * params.start; break; case RNG_Laplace: