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

Add support for unary negation operator #17560

Open
wants to merge 20 commits into
base: branch-25.02
Choose a base branch
from
Open

Add support for unary negation operator #17560

Show file tree
Hide file tree
Changes from 8 commits
Commits
Show all changes
20 commits
Select commit Hold shift + click to select a range
1ffae0b
Add support for unary negate operator in libcudf
Matt711 Dec 10, 2024
6c5290f
plumb through cudf python
Matt711 Dec 10, 2024
819b0c8
reject unsigned types
Matt711 Dec 10, 2024
3068215
clean up
Matt711 Dec 10, 2024
6d3e1cc
fix test
Matt711 Dec 10, 2024
688d60a
specialized template for DeviceNegate
Matt711 Dec 10, 2024
9d5365d
Merge branch 'branch-25.02' into fea/negate-unary-op
Matt711 Dec 11, 2024
8134e18
Update math_ops.cu
Matt711 Dec 11, 2024
4798838
extend comment to duration types
Matt711 Dec 11, 2024
4cc1033
Merge branch 'branch-25.02' into fea/negate-unary-op
Matt711 Dec 11, 2024
9910117
remove extra paren
Matt711 Dec 11, 2024
0b219fd
clean up
Matt711 Dec 11, 2024
cef040c
add support for decimal types, clean up tests
Matt711 Dec 12, 2024
67118a0
Merge branch 'branch-25.02' into fea/negate-unary-op
Matt711 Dec 12, 2024
e8779dd
partially address review
Matt711 Dec 12, 2024
b4b4f39
add missing types
Matt711 Dec 12, 2024
b19eebe
Update cpp/tests/unary/math_ops_test.cpp
Matt711 Dec 13, 2024
fe6a01a
Update cpp/tests/unary/math_ops_test.cpp
Matt711 Dec 13, 2024
a6ed9ad
Merge branch 'branch-25.02' into fea/negate-unary-op
Matt711 Dec 13, 2024
48f4a21
Merge branch 'branch-25.02' into fea/negate-unary-op
Matt711 Dec 13, 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
1 change: 1 addition & 0 deletions cpp/include/cudf/unary.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -139,6 +139,7 @@ enum class unary_operator : int32_t {
RINT, ///< Rounds the floating-point argument arg to an integer value
BIT_INVERT, ///< Bitwise Not (~)
NOT, ///< Logical Not (!)
NEGATE, ///< Unary negation (-), only for signed numeric types.
Matt711 marked this conversation as resolved.
Show resolved Hide resolved
};

/**
Expand Down
38 changes: 34 additions & 4 deletions cpp/src/unary/math_ops.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -234,6 +234,21 @@ struct DeviceNot {
}
};

// negation

struct DeviceNegate {
template <typename T, CUDF_ENABLE_IF(std::is_signed_v<T> || cudf::is_duration_t<T>::value)>
T __device__ operator()(T data)
{
return -data;
Matt711 marked this conversation as resolved.
Show resolved Hide resolved
}
template <typename T, CUDF_ENABLE_IF(!std::is_signed_v<T> && !cudf::is_duration_t<T>::value)>
T __device__ operator()(T data)
{
return data;
}
Comment on lines +245 to +249
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I am confused by this implementation. I thought we didn't want to provide an implementation for unsigned data types?

Copy link
Contributor

@bdice bdice Dec 12, 2024 •
edited by Matt711
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think that needs to happen at the dispatch layer, not SFINAE on the device operator itself. For example, we're not constraining types in other operators:

struct DeviceSqrt {
  template <typename T>
  __device__ T operator()(T data)
  {
    return std::sqrt(data);
  }
};

We only use SFINAE in this file when the implementation differs by type -- not to allow/disallow types.

Can we remove the SFINAE and just have this?

struct DeviceNegate {
  template <typename T>
  T __device__ operator()(T data)
  {
    return -data;
  }
};

Also related to https://github.com/rapidsai/cudf/pull/17560/files#r1878099060

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Let me post the compile error I got when I removed it.

};

// fixed_point ops

/*
Expand Down Expand Up @@ -361,9 +376,19 @@ std::unique_ptr<cudf::column> transform_fn(cudf::dictionary_column_view const& i
output->view(), dictionary::detail::get_indices_type_for_size(output->size()), stream, mr);
}

template <typename UFN, typename T>
struct is_supported_type : std::is_arithmetic<T> {};

template <typename T>
struct is_supported_type<DeviceNegate, T>
: std::disjunction<std::is_arithmetic<T>, cudf::is_duration_t<T>> {};
Matt711 marked this conversation as resolved.
Show resolved Hide resolved

template <typename UFN, typename T>
constexpr bool is_supported_type_v = is_supported_type<UFN, T>::value;

template <typename UFN>
struct MathOpDispatcher {
template <typename T, std::enable_if_t<std::is_arithmetic_v<T>>* = nullptr>
template <typename T, std::enable_if_t<is_supported_type_v<UFN, T>>* = nullptr>
std::unique_ptr<cudf::column> operator()(cudf::column_view const& input,
rmm::cuda_stream_view stream,
rmm::device_async_resource_ref mr)
Expand All @@ -377,7 +402,7 @@ struct MathOpDispatcher {
}

struct dictionary_dispatch {
template <typename T, std::enable_if_t<std::is_arithmetic_v<T>>* = nullptr>
template <typename T, std::enable_if_t<is_supported_type_v<UFN, T>>* = nullptr>
std::unique_ptr<cudf::column> operator()(cudf::dictionary_column_view const& input,
rmm::cuda_stream_view stream,
rmm::device_async_resource_ref mr)
Expand All @@ -394,7 +419,7 @@ struct MathOpDispatcher {

template <
typename T,
std::enable_if_t<!std::is_arithmetic_v<T> and std::is_same_v<T, dictionary32>>* = nullptr>
std::enable_if_t<!is_supported_type_v<UFN, T> and std::is_same_v<T, dictionary32>>* = nullptr>
std::unique_ptr<cudf::column> operator()(cudf::column_view const& input,
rmm::cuda_stream_view stream,
rmm::device_async_resource_ref mr)
Expand All @@ -406,7 +431,7 @@ struct MathOpDispatcher {
}

template <typename T, typename... Args>
std::enable_if_t<!std::is_arithmetic_v<T> and !std::is_same_v<T, dictionary32>,
std::enable_if_t<!is_supported_type_v<UFN, T> and !std::is_same_v<T, dictionary32>,
std::unique_ptr<cudf::column>>
operator()(Args&&...)
{
Expand Down Expand Up @@ -639,6 +664,11 @@ std::unique_ptr<cudf::column> unary_operation(cudf::column_view const& input,
case cudf::unary_operator::NOT:
return cudf::type_dispatcher(
input.type(), detail::LogicalOpDispatcher<detail::DeviceNot>{}, input, stream, mr);
case cudf::unary_operator::NEGATE:
CUDF_EXPECTS(cudf::is_duration(input.type()) || cudf::is_signed(input.type())),
"NEGATE operator requires signed numeric types or duration types.");
return cudf::type_dispatcher(
input.type(), detail::MathOpDispatcher<detail::DeviceNegate>{}, input, stream, mr);
Matt711 marked this conversation as resolved.
Show resolved Hide resolved
default: CUDF_FAIL("Undefined unary operation");
}
}
Expand Down
28 changes: 27 additions & 1 deletion cpp/tests/unary/math_ops_test.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -25,6 +25,23 @@

#include <vector>

using SignedNumericTypesNotBool =
cudf::test::Types<int8_t, int16_t, int32_t, int64_t, float, double>;

template <typename T>
struct UnaryMathOpsSignedTest : public cudf::test::BaseFixture {};

TYPED_TEST_SUITE(UnaryMathOpsSignedTest, SignedNumericTypesNotBool);

TYPED_TEST(UnaryMathOpsSignedTest, SimpleNEGATE)
{
cudf::test::fixed_width_column_wrapper<TypeParam> input{{1, 2, 3}};
auto const v = cudf::test::make_type_param_vector<TypeParam>({-1, -2, -3});
cudf::test::fixed_width_column_wrapper<TypeParam> expected(v.begin(), v.end());
auto output = cudf::unary_operation(input, cudf::unary_operator::NEGATE);
CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, output->view());
}

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

We need another test that unsigned types, timestamps, strings, and other complex types fail.

template <typename T>
struct UnaryLogicalOpsTest : public cudf::test::BaseFixture {};

Expand Down Expand Up @@ -234,6 +251,15 @@ using floating_point_type_list = ::testing::Types<float, double>;

TYPED_TEST_SUITE(UnaryMathFloatOpsTest, floating_point_type_list);

TYPED_TEST(UnaryMathFloatOpsTest, SimpleNEGATE)
{
cudf::test::fixed_width_column_wrapper<TypeParam> input{{1.0, 2.0}};
auto const v = cudf::test::make_type_param_vector<TypeParam>({-1.0, -2.0});
cudf::test::fixed_width_column_wrapper<TypeParam> expected(v.begin(), v.end());
auto output = cudf::unary_operation(input, cudf::unary_operator::NEGATE);
CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, output->view());
}

TYPED_TEST(UnaryMathFloatOpsTest, SimpleSIN)
{
cudf::test::fixed_width_column_wrapper<TypeParam> input{{0.0}};
Expand Down Expand Up @@ -274,7 +300,7 @@ TYPED_TEST(UnaryMathFloatOpsTest, SimpleTANH)
CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, output->view());
}

TYPED_TEST(UnaryMathFloatOpsTest, SimpleiASINH)
TYPED_TEST(UnaryMathFloatOpsTest, SimpleASINH)
{
cudf::test::fixed_width_column_wrapper<TypeParam> input{{0.0}};
cudf::test::fixed_width_column_wrapper<TypeParam> expected{{0.0}};
Expand Down
2 changes: 1 addition & 1 deletion python/cudf/cudf/core/frame.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1609,7 +1609,7 @@ def __neg__(self):
(
col.unary_operator("not")
if col.dtype.kind == "b"
else -1 * col
else col.unary_operator("negate")
for col in self._columns
)
)
Expand Down
1 change: 1 addition & 0 deletions python/cudf_polars/cudf_polars/dsl/expressions/unary.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -117,6 +117,7 @@ class UnaryFunction(Expr):
"abs": plc.unary.UnaryOperator.ABS,
"bit_invert": plc.unary.UnaryOperator.BIT_INVERT,
"not": plc.unary.UnaryOperator.NOT,
"negate": plc.unary.UnaryOperator.NEGATE,
}
_supported_misc_fns = frozenset(
{
Expand Down
9 changes: 9 additions & 0 deletions python/cudf_polars/tests/expressions/test_numeric_unaryops.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2,6 +2,8 @@
# SPDX-License-Identifier: Apache-2.0
from __future__ import annotations

from datetime import timedelta

import numpy as np
import pytest

Expand Down Expand Up @@ -58,6 +60,7 @@ def ldf(with_nulls, dtype):
{
"a": pl.Series(values, dtype=dtype),
"b": pl.Series([i - 4 for i in range(len(values))], dtype=pl.Float32),
"c": pl.Series([timedelta(hours=i) for i in range(len(values))]),
}
)

Expand Down Expand Up @@ -89,3 +92,9 @@ def test_log(ldf, natural):
q = ldf.select(expr)

assert_gpu_result_equal(q, check_exact=False)


@pytest.mark.parametrize("col", ["a", "b", "c"])
def test_negate(ldf, col):
q = ldf.select(-pl.col(col))
assert_gpu_result_equal(q)
1 change: 1 addition & 0 deletions python/pylibcudf/pylibcudf/libcudf/unary.pxd
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -33,6 +33,7 @@ cdef extern from "cudf/unary.hpp" namespace "cudf" nogil:
RINT
BIT_INVERT
NOT
NEGATE
Matt711 marked this conversation as resolved.
Show resolved Hide resolved

cdef extern unique_ptr[column] unary_operation(
column_view input,
Expand Down
1 change: 1 addition & 0 deletions python/pylibcudf/pylibcudf/unary.pyi
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -28,6 +28,7 @@ class UnaryOperator(IntEnum):
RINT = ...
BIT_INVERT = ...
NOT = ...
NEGATE = ...

def unary_operation(input: Column, op: UnaryOperator) -> Column: ...
def is_null(input: Column) -> Column: ...
Expand Down
Loading