Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Enable binaryop build without relying on relaxed constexpr #17598

Draft
wants to merge 8 commits into
base: branch-25.02
Choose a base branch
from
Draft

Enable binaryop build without relying on relaxed constexpr #17598

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
8 commits
Select commit Hold shift + click to select a range
1f0428f
Remove relaxed constexpr
PointKernel Dec 16, 2024
aaa8f5d
Enable binaryop build without relying on relaxed constexpr
PointKernel Dec 16, 2024
f556d1b
Make fixed_point a CUDA header
PointKernel Dec 16, 2024
9e653b5
Merge remote-tracking branch 'upstream/branch-25.02' into remove-cons…
PointKernel Dec 16, 2024
55a453d
Revert CMake changes
PointKernel Dec 16, 2024
03a484b
Merge remote-tracking branch 'upstream/branch-25.02' into remove-cons…
PointKernel Dec 16, 2024
adf58d2
Merge remote-tracking branch 'upstream/branch-25.02' into remove-cons…
PointKernel Dec 20, 2024
775de86
Revert fixed point back to cpp header + fix an nvrtc issue
PointKernel Dec 20, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
42 changes: 30 additions & 12 deletions cpp/include/cudf/fixed_point/fixed_point.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -60,7 +60,7 @@ enum class Radix : int32_t { BASE_2 = 2, BASE_10 = 10 };
* @return `true` if the type is supported by `fixed_point` implementation
*/
template <typename T>
constexpr inline auto is_supported_representation_type()
CUDF_HOST_DEVICE constexpr inline auto is_supported_representation_type()
{
return cuda::std::is_same_v<T, int32_t> || //
cuda::std::is_same_v<T, int64_t> || //
Expand All @@ -72,6 +72,24 @@ constexpr inline auto is_supported_representation_type()
// Helper functions for `fixed_point` type
namespace detail {

/**
* @brief Returns the smaller of the given scales
*
* @param a The left-hand side value to compare
* @param b The right-hand side value to compare
* @return The smaller of the given scales
*/
CUDF_HOST_DEVICE constexpr inline scale_type min(scale_type const& a, scale_type const& b)
{
// TODO This is a temporary workaround because <cuda/std/functional> is not self-contained when
// built with NVRTC 11.8. Replace this with cuda::std::min once the underlying issue is resolved.
#ifdef __CUDA_ARCH__
return scale_type{min(static_cast<int>(a), static_cast<int>(b))};
#else
return std::min(a, b);
#endif
}

/**
* @brief A function for integer exponentiation by squaring.
*
Expand Down Expand Up @@ -267,12 +285,12 @@ class fixed_point {
* @return The `fixed_point` number in base 10 (aka human readable format)
*/
template <typename U, typename cuda::std::enable_if_t<cuda::std::is_integral_v<U>>* = nullptr>
explicit constexpr operator U() const
CUDF_HOST_DEVICE explicit constexpr operator U() const
{
// Cast to the larger of the two types (of U and Rep) before converting to Rep because in
// certain cases casting to U before shifting will result in integer overflow (i.e. if U =
// int32_t, Rep = int64_t and _value > 2 billion)
auto const value = std::common_type_t<U, Rep>(_value);
auto const value = cuda::std::common_type_t<U, Rep>(_value);
return static_cast<U>(detail::shift<Rep, Rad>(value, scale_type{-_scale}));
}

Expand Down Expand Up @@ -669,7 +687,7 @@ template <typename Rep1, Radix Rad1>
CUDF_HOST_DEVICE inline fixed_point<Rep1, Rad1> operator+(fixed_point<Rep1, Rad1> const& lhs,
fixed_point<Rep1, Rad1> const& rhs)
{
auto const scale = std::min(lhs._scale, rhs._scale);
auto const scale = detail::min(lhs._scale, rhs._scale);
auto const sum = lhs.rescaled(scale)._value + rhs.rescaled(scale)._value;

#if defined(__CUDACC_DEBUG__)
Expand All @@ -687,7 +705,7 @@ template <typename Rep1, Radix Rad1>
CUDF_HOST_DEVICE inline fixed_point<Rep1, Rad1> operator-(fixed_point<Rep1, Rad1> const& lhs,
fixed_point<Rep1, Rad1> const& rhs)
{
auto const scale = std::min(lhs._scale, rhs._scale);
auto const scale = detail::min(lhs._scale, rhs._scale);
auto const diff = lhs.rescaled(scale)._value - rhs.rescaled(scale)._value;

#if defined(__CUDACC_DEBUG__)
Expand Down Expand Up @@ -735,7 +753,7 @@ template <typename Rep1, Radix Rad1>
CUDF_HOST_DEVICE inline bool operator==(fixed_point<Rep1, Rad1> const& lhs,
fixed_point<Rep1, Rad1> const& rhs)
{
auto const scale = std::min(lhs._scale, rhs._scale);
auto const scale = detail::min(lhs._scale, rhs._scale);
return lhs.rescaled(scale)._value == rhs.rescaled(scale)._value;
}

Expand All @@ -744,7 +762,7 @@ template <typename Rep1, Radix Rad1>
CUDF_HOST_DEVICE inline bool operator!=(fixed_point<Rep1, Rad1> const& lhs,
fixed_point<Rep1, Rad1> const& rhs)
{
auto const scale = std::min(lhs._scale, rhs._scale);
auto const scale = detail::min(lhs._scale, rhs._scale);
return lhs.rescaled(scale)._value != rhs.rescaled(scale)._value;
}

Expand All @@ -753,7 +771,7 @@ template <typename Rep1, Radix Rad1>
CUDF_HOST_DEVICE inline bool operator<=(fixed_point<Rep1, Rad1> const& lhs,
fixed_point<Rep1, Rad1> const& rhs)
{
auto const scale = std::min(lhs._scale, rhs._scale);
auto const scale = detail::min(lhs._scale, rhs._scale);
return lhs.rescaled(scale)._value <= rhs.rescaled(scale)._value;
}

Expand All @@ -762,7 +780,7 @@ template <typename Rep1, Radix Rad1>
CUDF_HOST_DEVICE inline bool operator>=(fixed_point<Rep1, Rad1> const& lhs,
fixed_point<Rep1, Rad1> const& rhs)
{
auto const scale = std::min(lhs._scale, rhs._scale);
auto const scale = detail::min(lhs._scale, rhs._scale);
return lhs.rescaled(scale)._value >= rhs.rescaled(scale)._value;
}

Expand All @@ -771,7 +789,7 @@ template <typename Rep1, Radix Rad1>
CUDF_HOST_DEVICE inline bool operator<(fixed_point<Rep1, Rad1> const& lhs,
fixed_point<Rep1, Rad1> const& rhs)
{
auto const scale = std::min(lhs._scale, rhs._scale);
auto const scale = detail::min(lhs._scale, rhs._scale);
return lhs.rescaled(scale)._value < rhs.rescaled(scale)._value;
}

Expand All @@ -780,7 +798,7 @@ template <typename Rep1, Radix Rad1>
CUDF_HOST_DEVICE inline bool operator>(fixed_point<Rep1, Rad1> const& lhs,
fixed_point<Rep1, Rad1> const& rhs)
{
auto const scale = std::min(lhs._scale, rhs._scale);
auto const scale = detail::min(lhs._scale, rhs._scale);
return lhs.rescaled(scale)._value > rhs.rescaled(scale)._value;
}

Expand All @@ -789,7 +807,7 @@ template <typename Rep1, Radix Rad1>
CUDF_HOST_DEVICE inline fixed_point<Rep1, Rad1> operator%(fixed_point<Rep1, Rad1> const& lhs,
fixed_point<Rep1, Rad1> const& rhs)
{
auto const scale = std::min(lhs._scale, rhs._scale);
auto const scale = detail::min(lhs._scale, rhs._scale);
auto const remainder = lhs.rescaled(scale)._value % rhs.rescaled(scale)._value;
return fixed_point<Rep1, Rad1>{scaled_integer<Rep1>{remainder, scale}};
}
Expand Down
38 changes: 20 additions & 18 deletions cpp/src/binaryop/compiled/binary_ops.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -27,6 +27,8 @@
#include <rmm/cuda_stream_view.hpp>
#include <rmm/exec_policy.hpp>

#include <cuda/std/type_traits>

namespace cudf {
namespace binops {
namespace compiled {
Expand All @@ -51,7 +53,7 @@ struct type_casted_accessor {
{
if constexpr (column_device_view::has_element_accessor<Element>()) {
auto const element = col.element<Element>(is_scalar ? 0 : i);
if constexpr (std::is_convertible_v<Element, CastType>) {
if constexpr (cuda::std::is_convertible_v<Element, CastType>) {
return static_cast<CastType>(element);
} else if constexpr (is_fixed_point<Element>() && cuda::std::is_floating_point_v<CastType>) {
return convert_fixed_to_floating<CastType>(element);
Expand All @@ -75,7 +77,7 @@ struct typed_casted_writer {
FromType val) const
{
if constexpr (mutable_column_device_view::has_element_accessor<Element>() and
std::is_constructible_v<Element, FromType>) {
cuda::std::is_constructible_v<Element, FromType>) {
col.element<Element>(i) = static_cast<Element>(val);
} else if constexpr (is_fixed_point<Element>()) {
auto const scale = numeric::scale_type{col.type().scale()};
Expand Down Expand Up @@ -109,18 +111,18 @@ struct ops_wrapper {
template <typename TypeCommon>
__device__ void operator()(size_type i)
{
if constexpr (std::is_invocable_v<BinaryOperator, TypeCommon, TypeCommon>) {
if constexpr (cuda::std::is_invocable_v<BinaryOperator, TypeCommon, TypeCommon>) {
TypeCommon x =
type_dispatcher(lhs.type(), type_casted_accessor<TypeCommon>{}, i, lhs, is_lhs_scalar);
TypeCommon y =
type_dispatcher(rhs.type(), type_casted_accessor<TypeCommon>{}, i, rhs, is_rhs_scalar);
auto result = [&]() {
if constexpr (std::is_same_v<BinaryOperator, ops::NullEquals> or
std::is_same_v<BinaryOperator, ops::NullNotEquals> or
std::is_same_v<BinaryOperator, ops::NullLogicalAnd> or
std::is_same_v<BinaryOperator, ops::NullLogicalOr> or
std::is_same_v<BinaryOperator, ops::NullMax> or
std::is_same_v<BinaryOperator, ops::NullMin>) {
if constexpr (cuda::std::is_same_v<BinaryOperator, ops::NullEquals> or
cuda::std::is_same_v<BinaryOperator, ops::NullNotEquals> or
cuda::std::is_same_v<BinaryOperator, ops::NullLogicalAnd> or
cuda::std::is_same_v<BinaryOperator, ops::NullLogicalOr> or
cuda::std::is_same_v<BinaryOperator, ops::NullMax> or
cuda::std::is_same_v<BinaryOperator, ops::NullMin>) {
bool output_valid = false;
auto result = BinaryOperator{}.template operator()<TypeCommon, TypeCommon>(
x,
Expand All @@ -134,7 +136,7 @@ struct ops_wrapper {
return BinaryOperator{}.template operator()<TypeCommon, TypeCommon>(x, y);
}
// To suppress nvcc warning
return std::invoke_result_t<BinaryOperator, TypeCommon, TypeCommon>{};
return cuda::std::invoke_result_t<BinaryOperator, TypeCommon, TypeCommon>{};
}();
if constexpr (is_bool_result<BinaryOperator, TypeCommon, TypeCommon>())
out.element<decltype(result)>(i) = result;
Expand All @@ -161,16 +163,16 @@ struct ops2_wrapper {
__device__ void operator()(size_type i)
{
if constexpr (!has_common_type_v<TypeLhs, TypeRhs> and
std::is_invocable_v<BinaryOperator, TypeLhs, TypeRhs>) {
cuda::std::is_invocable_v<BinaryOperator, TypeLhs, TypeRhs>) {
TypeLhs x = lhs.element<TypeLhs>(is_lhs_scalar ? 0 : i);
TypeRhs y = rhs.element<TypeRhs>(is_rhs_scalar ? 0 : i);
auto result = [&]() {
if constexpr (std::is_same_v<BinaryOperator, ops::NullEquals> or
std::is_same_v<BinaryOperator, ops::NullNotEquals> or
std::is_same_v<BinaryOperator, ops::NullLogicalAnd> or
std::is_same_v<BinaryOperator, ops::NullLogicalOr> or
std::is_same_v<BinaryOperator, ops::NullMax> or
std::is_same_v<BinaryOperator, ops::NullMin>) {
if constexpr (cuda::std::is_same_v<BinaryOperator, ops::NullEquals> or
cuda::std::is_same_v<BinaryOperator, ops::NullNotEquals> or
cuda::std::is_same_v<BinaryOperator, ops::NullLogicalAnd> or
cuda::std::is_same_v<BinaryOperator, ops::NullLogicalOr> or
cuda::std::is_same_v<BinaryOperator, ops::NullMax> or
cuda::std::is_same_v<BinaryOperator, ops::NullMin>) {
bool output_valid = false;
auto result = BinaryOperator{}.template operator()<TypeLhs, TypeRhs>(
x,
Expand All @@ -184,7 +186,7 @@ struct ops2_wrapper {
return BinaryOperator{}.template operator()<TypeLhs, TypeRhs>(x, y);
}
// To suppress nvcc warning
return std::invoke_result_t<BinaryOperator, TypeLhs, TypeRhs>{};
return cuda::std::invoke_result_t<BinaryOperator, TypeLhs, TypeRhs>{};
}();
if constexpr (is_bool_result<BinaryOperator, TypeLhs, TypeRhs>())
out.element<decltype(result)>(i) = result;
Expand Down
Loading