diff --git a/cpp/include/cudf/unary.hpp b/cpp/include/cudf/unary.hpp
index 046e9745a71..5cb7fb4378f 100644
--- a/cpp/include/cudf/unary.hpp
+++ b/cpp/include/cudf/unary.hpp
@@ -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 and duration types.
 };
 
 /**
diff --git a/cpp/include/cudf/utilities/traits.hpp b/cpp/include/cudf/utilities/traits.hpp
index dae1cd38832..160ada6831c 100644
--- a/cpp/include/cudf/utilities/traits.hpp
+++ b/cpp/include/cudf/utilities/traits.hpp
@@ -94,6 +94,8 @@ constexpr inline bool has_common_type_v = detail::has_common_type_impl<void, Ts.
 /// Checks if a type is a timestamp type.
 template <typename T>
 using is_timestamp_t = cuda::std::disjunction<std::is_same<cudf::timestamp_D, T>,
+                                              std::is_same<cudf::timestamp_h, T>,
+                                              std::is_same<cudf::timestamp_m, T>,
                                               std::is_same<cudf::timestamp_s, T>,
                                               std::is_same<cudf::timestamp_ms, T>,
                                               std::is_same<cudf::timestamp_us, T>,
@@ -102,6 +104,8 @@ using is_timestamp_t = cuda::std::disjunction<std::is_same<cudf::timestamp_D, T>
 /// Checks if a type is a duration type.
 template <typename T>
 using is_duration_t = cuda::std::disjunction<std::is_same<cudf::duration_D, T>,
+                                             std::is_same<cudf::duration_h, T>,
+                                             std::is_same<cudf::duration_m, T>,
                                              std::is_same<cudf::duration_s, T>,
                                              std::is_same<cudf::duration_ms, T>,
                                              std::is_same<cudf::duration_us, T>,
diff --git a/cpp/src/unary/math_ops.cu b/cpp/src/unary/math_ops.cu
index 1d506c59cd9..d597b4d4bca 100644
--- a/cpp/src/unary/math_ops.cu
+++ b/cpp/src/unary/math_ops.cu
@@ -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;
+  }
+  template <typename T, CUDF_ENABLE_IF(!std::is_signed_v<T> && !cudf::is_duration_t<T>::value)>
+  T __device__ operator()(T data)
+  {
+    return data;
+  }
+};
+
 // fixed_point ops
 
 /*
@@ -278,6 +293,12 @@ struct fixed_point_abs {
   __device__ T operator()(T data) { return numeric::detail::abs(data); }
 };
 
+template <typename T>
+struct fixed_point_negate {
+  T n;
+  __device__ T operator()(T data) { return -data; }
+};
+
 template <typename T, template <typename> typename FixedPointFunctor>
 std::unique_ptr<column> unary_op_with(column_view const& input,
                                       rmm::cuda_stream_view stream,
@@ -361,9 +382,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>> {};
+
+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)
@@ -377,7 +408,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)
@@ -394,7 +425,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)
@@ -406,7 +437,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&&...)
   {
@@ -550,9 +581,10 @@ struct FixedPointOpDispatcher {
   {
     // clang-format off
     switch (op) {
-      case cudf::unary_operator::CEIL:  return unary_op_with<T, fixed_point_ceil>(input, stream, mr);
-      case cudf::unary_operator::FLOOR: return unary_op_with<T, fixed_point_floor>(input, stream, mr);
-      case cudf::unary_operator::ABS:   return unary_op_with<T, fixed_point_abs>(input, stream, mr);
+      case cudf::unary_operator::CEIL:   return unary_op_with<T, fixed_point_ceil>(input, stream, mr);
+      case cudf::unary_operator::FLOOR:  return unary_op_with<T, fixed_point_floor>(input, stream, mr);
+      case cudf::unary_operator::ABS:    return unary_op_with<T, fixed_point_abs>(input, stream, mr);
+      case cudf::unary_operator::NEGATE: return unary_op_with<T, fixed_point_negate>(input, stream, mr);
       default: CUDF_FAIL("Unsupported fixed_point unary operation");
     }
     // clang-format on
@@ -639,6 +671,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_signed(input.type()) || cudf::is_duration(input.type()),
+                   "NEGATE operator requires signed numeric types or duration types.");
+      return cudf::type_dispatcher(
+        input.type(), detail::MathOpDispatcher<detail::DeviceNegate>{}, input, stream, mr);
     default: CUDF_FAIL("Undefined unary operation");
   }
 }
diff --git a/cpp/tests/unary/math_ops_test.cpp b/cpp/tests/unary/math_ops_test.cpp
index 663a919f3f4..cad7e98096e 100644
--- a/cpp/tests/unary/math_ops_test.cpp
+++ b/cpp/tests/unary/math_ops_test.cpp
@@ -25,6 +25,73 @@
 
 #include <vector>
 
+using TypesToNegate = cudf::test::Types<int8_t,
+                                        int16_t,
+                                        int32_t,
+                                        int64_t,
+                                        float,
+                                        double,
+                                        cudf::duration_D,
+                                        cudf::duration_h,
+                                        cudf::duration_m,
+                                        cudf::duration_s,
+                                        cudf::duration_ms,
+                                        cudf::duration_us,
+                                        cudf::duration_ns>;
+
+template <typename T>
+struct UnaryNegateTests : public cudf::test::BaseFixture {};
+
+TYPED_TEST_SUITE(UnaryNegateTests, TypesToNegate);
+
+TYPED_TEST(UnaryNegateTests, SimpleNEGATE)
+{
+  using T = TypeParam;
+  cudf::test::fixed_width_column_wrapper<T> input{{0, 1, 2, 3}};
+  auto const v = cudf::test::make_type_param_vector<T>({0, -1, -2, -3});
+  cudf::test::fixed_width_column_wrapper<T> expected(v.begin(), v.end());
+  auto output = cudf::unary_operation(input, cudf::unary_operator::NEGATE);
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, output->view());
+}
+
+using TypesNotToNegate = cudf::test::Types<uint8_t,
+                                           uint16_t,
+                                           uint32_t,
+                                           uint64_t,
+                                           cudf::timestamp_D,
+                                           cudf::timestamp_h,
+                                           // cudf::timestamp_m,
+                                           // cudf::timestamp_s,
+                                           cudf::timestamp_ms,
+                                           cudf::timestamp_us,
+                                           cudf::timestamp_ns>;
+
+template <typename T>
+struct UnaryNegateErrorTests : public cudf::test::BaseFixture {};
+
+TYPED_TEST_SUITE(UnaryNegateErrorTests, TypesNotToNegate);
+
+TYPED_TEST(UnaryNegateErrorTests, UnsupportedTypesFail)
+{
+  using T = TypeParam;
+  cudf::test::fixed_width_column_wrapper<T> input({1, 2, 3, 4});
+  EXPECT_THROW(cudf::unary_operation(input, cudf::unary_operator::NEGATE), cudf::logic_error);
+}
+
+struct UnaryNegateComplexTypesErrorTests : public cudf::test::BaseFixture {};
+
+TEST_F(UnaryNegateComplexTypesErrorTests, NegateStringColumnFail)
+{
+  cudf::test::strings_column_wrapper input({"foo", "bar"});
+  EXPECT_THROW(cudf::unary_operation(input, cudf::unary_operator::NEGATE), cudf::logic_error);
+}
+
+TEST_F(UnaryNegateComplexTypesErrorTests, NegateListsColumnFail)
+{
+  cudf::test::lists_column_wrapper<int32_t> input{{1, 2}, {3, 4}};
+  EXPECT_THROW(cudf::unary_operation(input, cudf::unary_operator::NEGATE), cudf::logic_error);
+}
+
 template <typename T>
 struct UnaryLogicalOpsTest : public cudf::test::BaseFixture {};
 
@@ -274,7 +341,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}};
diff --git a/cpp/tests/unary/unary_ops_test.cpp b/cpp/tests/unary/unary_ops_test.cpp
index 3c616461c74..55a51219f9b 100644
--- a/cpp/tests/unary/unary_ops_test.cpp
+++ b/cpp/tests/unary/unary_ops_test.cpp
@@ -266,6 +266,20 @@ struct FixedPointUnaryTests : public cudf::test::BaseFixture {};
 
 TYPED_TEST_SUITE(FixedPointUnaryTests, cudf::test::FixedPointTypes);
 
+TYPED_TEST(FixedPointUnaryTests, FixedPointUnaryNegate)
+{
+  using namespace numeric;
+  using decimalXX  = TypeParam;
+  using RepType    = cudf::device_storage_type_t<decimalXX>;
+  using fp_wrapper = cudf::test::fixed_point_column_wrapper<RepType>;
+
+  auto const input    = fp_wrapper{{0, -1234, -3456, -6789, 1234, 3456, 6789}, scale_type{-3}};
+  auto const expected = fp_wrapper{{0, 1234, 3456, 6789, -1234, -3456, -6789}, scale_type{-3}};
+  auto const result   = cudf::unary_operation(input, cudf::unary_operator::NEGATE);
+
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view());
+}
+
 TYPED_TEST(FixedPointUnaryTests, FixedPointUnaryAbs)
 {
   using namespace numeric;
diff --git a/python/cudf/cudf/core/column/decimal.py b/python/cudf/cudf/core/column/decimal.py
index 9e6a73f1a9c..9956f785994 100644
--- a/python/cudf/cudf/core/column/decimal.py
+++ b/python/cudf/cudf/core/column/decimal.py
@@ -226,6 +226,17 @@ def as_numerical_column(
     ) -> "cudf.core.column.NumericalColumn":
         return unary.cast(self, dtype)  # type: ignore[return-value]
 
+    def unary_operator(self, unaryop: str) -> ColumnBase:
+        # TODO: Support Callable unary operations
+        unaryop = unaryop.upper()
+        try:
+            unaryop = plc.unary.UnaryOperator[unaryop]
+            return unary.unary_operation(self, unaryop)
+        except Exception:
+            raise TypeError(
+                f"Operation {unaryop} not supported for dtype {self.dtype}."
+            )
+
 
 class Decimal32Column(DecimalBaseColumn):
     def __init__(
diff --git a/python/cudf/cudf/core/frame.py b/python/cudf/cudf/core/frame.py
index 00199cca828..bee829f3057 100644
--- a/python/cudf/cudf/core/frame.py
+++ b/python/cudf/cudf/core/frame.py
@@ -1652,7 +1652,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
                 )
             )
diff --git a/python/cudf/cudf/tests/test_unaops.py b/python/cudf/cudf/tests/test_unaops.py
index b714beb0069..2bd7c7084e8 100644
--- a/python/cudf/cudf/tests/test_unaops.py
+++ b/python/cudf/cudf/tests/test_unaops.py
@@ -3,6 +3,7 @@
 import itertools
 import operator
 import re
+from decimal import Decimal
 
 import numpy as np
 import pandas as pd
@@ -135,3 +136,9 @@ def test_series_bool_neg():
     sr = Series([True, False, True, None, False, None, True, True])
     psr = sr.to_pandas(nullable=True)
     assert_eq((-sr).to_pandas(nullable=True), -psr, check_dtype=True)
+
+
+def test_series_decimal_neg():
+    sr = Series([Decimal("0.0"), Decimal("1.23"), Decimal("4.567")])
+    psr = sr.to_pandas()
+    assert_eq((-sr).to_pandas(), -psr, check_dtype=True)
diff --git a/python/cudf_polars/cudf_polars/dsl/expressions/unary.py b/python/cudf_polars/cudf_polars/dsl/expressions/unary.py
index 10caaff6811..f364e562e8c 100644
--- a/python/cudf_polars/cudf_polars/dsl/expressions/unary.py
+++ b/python/cudf_polars/cudf_polars/dsl/expressions/unary.py
@@ -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(
         {
diff --git a/python/cudf_polars/tests/expressions/test_numeric_unaryops.py b/python/cudf_polars/tests/expressions/test_numeric_unaryops.py
index ac3aecf88e6..6b89b5b349e 100644
--- a/python/cudf_polars/tests/expressions/test_numeric_unaryops.py
+++ b/python/cudf_polars/tests/expressions/test_numeric_unaryops.py
@@ -2,6 +2,8 @@
 # SPDX-License-Identifier: Apache-2.0
 from __future__ import annotations
 
+from datetime import timedelta
+
 import numpy as np
 import pytest
 
@@ -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))]),
         }
     )
 
@@ -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)
diff --git a/python/pylibcudf/pylibcudf/libcudf/unary.pxd b/python/pylibcudf/pylibcudf/libcudf/unary.pxd
index 4666012623e..711113af98d 100644
--- a/python/pylibcudf/pylibcudf/libcudf/unary.pxd
+++ b/python/pylibcudf/pylibcudf/libcudf/unary.pxd
@@ -33,6 +33,7 @@ cdef extern from "cudf/unary.hpp" namespace "cudf" nogil:
         RINT
         BIT_INVERT
         NOT
+        NEGATE
 
     cdef extern unique_ptr[column] unary_operation(
         column_view input,
diff --git a/python/pylibcudf/pylibcudf/unary.pyi b/python/pylibcudf/pylibcudf/unary.pyi
index 7aa23b618f4..4d06a51c03a 100644
--- a/python/pylibcudf/pylibcudf/unary.pyi
+++ b/python/pylibcudf/pylibcudf/unary.pyi
@@ -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: ...