From a2ea0524e2d598b947b113d4863727fe320c67cf Mon Sep 17 00:00:00 2001 From: Shane Ding Date: Mon, 26 Jul 2021 20:31:19 +0000 Subject: [PATCH 01/14] added initial function --- cpp/src/datetime/datetime_ops.cu | 10 ++++++++++ 1 file changed, 10 insertions(+) diff --git a/cpp/src/datetime/datetime_ops.cu b/cpp/src/datetime/datetime_ops.cu index 41f3e7dcfee..7eec1a2028f 100644 --- a/cpp/src/datetime/datetime_ops.cu +++ b/cpp/src/datetime/datetime_ops.cu @@ -77,6 +77,16 @@ struct extract_component_operator { } }; +template +struct ceil_timestamp { + template + CUDA_DEVICE_CALLABLE Timestamp operator()(Timestamp const ts) const + { + using namespace cuda::std::chrono; + return Timestamp{ts}; + } +}; + // Number of days until month indexed by leap year and month (0-based index) static __device__ int16_t const days_until_month[2][13] = { {0, 31, 59, 90, 120, 151, 181, 212, 243, 273, 304, 334, 365}, // For non leap years From a1ddeb33953913842572921946b8eb682e706621 Mon Sep 17 00:00:00 2001 From: Shane Ding Date: Tue, 3 Aug 2021 22:10:40 +0000 Subject: [PATCH 02/14] initial version --- cpp/include/cudf/datetime.hpp | 4 ++ cpp/src/datetime/datetime_ops.cu | 70 +++++++++++++++++++++++- cpp/tests/datetime/datetime_ops_test.cpp | 31 +++++++++++ 3 files changed, 102 insertions(+), 3 deletions(-) diff --git a/cpp/include/cudf/datetime.hpp b/cpp/include/cudf/datetime.hpp index 980c824fdf2..d9ea910c258 100644 --- a/cpp/include/cudf/datetime.hpp +++ b/cpp/include/cudf/datetime.hpp @@ -207,5 +207,9 @@ std::unique_ptr is_leap_year( rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** @} */ // end of group + +std::unique_ptr ceil_day( + cudf::column_view const& column, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); } // namespace datetime } // namespace cudf diff --git a/cpp/src/datetime/datetime_ops.cu b/cpp/src/datetime/datetime_ops.cu index 7eec1a2028f..216bb9a5c9a 100644 --- a/cpp/src/datetime/datetime_ops.cu +++ b/cpp/src/datetime/datetime_ops.cu @@ -26,8 +26,11 @@ #include #include +#include #include #include +#include "cudf/utilities/error.hpp" +#include "cudf/utilities/type_dispatcher.hpp" namespace cudf { namespace datetime { @@ -77,13 +80,19 @@ struct extract_component_operator { } }; -template +template struct ceil_timestamp { - template CUDA_DEVICE_CALLABLE Timestamp operator()(Timestamp const ts) const { using namespace cuda::std::chrono; - return Timestamp{ts}; + switch (Component) { + case datetime_component::DAY: + return time_point_cast( + cuda::std::chrono::ceil(ts)); + default: + return time_point_cast( + cuda::std::chrono::ceil(ts)); + } } }; @@ -148,6 +157,47 @@ struct is_leap_year_op { } }; +// Specific function for applying ceil/floor date ops +template +struct dispatch_ceil { + template + std::enable_if_t(), std::unique_ptr> operator()( + // cudf::data_type type, + cudf::column_view const& column, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) const + { + CUDF_EXPECTS(is_timestamp(column.type()), "Column type should be timestamp"); + auto size = column.size(); + auto output_col_type = data_type{cudf::type_to_id()}; + + // Return an empty column if source column is empty + if (size == 0) return make_empty_column(output_col_type); + + auto output = make_fixed_width_column(output_col_type, + size, + cudf::detail::copy_bitmask(column, stream, mr), + column.null_count(), + stream, + mr); + + thrust::transform(rmm::exec_policy(stream), + column.begin(), + column.end(), + output->mutable_view().begin(), + ceil_timestamp{}); + + return output; + } + + template + std::enable_if_t(), std::unique_ptr> operator()( + Args&&...) + { + CUDF_FAIL("Must be cudf::timestamp"); + } +}; + // Apply the functor for every element/row in the input column to create the output column template struct launch_functor { @@ -300,6 +350,14 @@ std::unique_ptr add_calendrical_months(column_view const& timestamp_colu return output; } +std::unique_ptr ceil_days(column_view const& column, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + return cudf::type_dispatcher( + column.type(), dispatch_ceil{}, column, stream, mr); +} + std::unique_ptr extract_year(column_view const& column, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) @@ -388,6 +446,12 @@ std::unique_ptr is_leap_year(column_view const& column, } // namespace detail +std::unique_ptr ceil_day(column_view const& column, rmm::mr::device_memory_resource* mr) +{ + CUDF_FUNC_RANGE(); + return detail::ceil_days(column, rmm::cuda_stream_default, mr); +} + std::unique_ptr extract_year(column_view const& column, rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/tests/datetime/datetime_ops_test.cpp b/cpp/tests/datetime/datetime_ops_test.cpp index c200dc6c935..2a98c79540a 100644 --- a/cpp/tests/datetime/datetime_ops_test.cpp +++ b/cpp/tests/datetime/datetime_ops_test.cpp @@ -347,6 +347,37 @@ TEST_F(BasicDatetimeOpsTest, TestLastDayOfMonthWithDate) verbosity); } +TEST_F(BasicDatetimeOpsTest, TestCeilDay) +{ + using namespace cudf::test; + using namespace cudf::datetime; + using namespace cuda::std::chrono; + + auto timestamps_s = + cudf::test::fixed_width_column_wrapper{ + { + // null + 915148800L, // 1999-01-01 00:00:00 GMT - non leap year + -11663029161L, // 1600-5-31 05:40:39 GMT - leap year + 707904541L, // 1992-06-07 08:09:01 GMT - leap year + 2181048447L, // 1900-11-20 09:12:33 GMT - non leap year + }, + {true, true, true, true}}; + + CUDF_TEST_EXPECT_COLUMNS_EQUAL( + *ceil_day(timestamps_s), + fixed_width_column_wrapper{ + { + // null + 915148800L, // 1999-01-01 00:00:00 GMT - non leap year + -11662963200L, // 1600-5-31 05:40:39 GMT - leap year + 707961600L, // 1992-06-07 08:09:01 GMT - leap year + 2181081600L, // 1900-11-20 09:12:33 GMT - non leap year + }, + {true, true, true, true}}, + verbosity); +} + TEST_F(BasicDatetimeOpsTest, TestDayOfYearWithDate) { using namespace cudf::test; From ce0c3008192b9343d875abd6b66fd1a65a0de4c9 Mon Sep 17 00:00:00 2001 From: Shane Ding Date: Thu, 5 Aug 2021 23:12:47 +0000 Subject: [PATCH 03/14] added ceil for all resolutions --- cpp/include/cudf/datetime.hpp | 77 ++++++++++++ cpp/include/cudf/wrappers/durations.hpp | 8 ++ cpp/src/datetime/datetime_ops.cu | 148 ++++++++++++++++++++++-- 3 files changed, 222 insertions(+), 11 deletions(-) diff --git a/cpp/include/cudf/datetime.hpp b/cpp/include/cudf/datetime.hpp index d9ea910c258..cf9161e4c58 100644 --- a/cpp/include/cudf/datetime.hpp +++ b/cpp/include/cudf/datetime.hpp @@ -208,8 +208,85 @@ std::unique_ptr is_leap_year( /** @} */ // end of group +/** + * @brief Round up to the nearest day + * + * @param[in] cudf::column_view of the input datetime values + * + * @returns cudf::column of the same datetime resolution as the input column + * @throw cudf::logic_error if input column datatype is not TIMESTAMP + */ std::unique_ptr ceil_day( cudf::column_view const& column, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + +/** + * @brief Round up to the nearest hour + * + * @param[in] cudf::column_view of the input datetime values + * + * @returns cudf::column of the same datetime resolution as the input column + * @throw cudf::logic_error if input column datatype is not TIMESTAMP + */ +std::unique_ptr ceil_hour( + cudf::column_view const& column, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + +/** + * @brief Round up to the nearest minute + * + * @param[in] cudf::column_view of the input datetime values + * + * @returns cudf::column of the same datetime resolution as the input column + * @throw cudf::logic_error if input column datatype is not TIMESTAMP + */ +std::unique_ptr ceil_minute( + cudf::column_view const& column, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + +/** + * @brief Round up to the nearest second + * + * @param[in] cudf::column_view of the input datetime values + * + * @returns cudf::column of the same datetime resolution as the input column + * @throw cudf::logic_error if input column datatype is not TIMESTAMP + */ +std::unique_ptr ceil_second( + cudf::column_view const& column, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + +/** + * @brief Round up to the nearest millisecond + * + * @param[in] cudf::column_view of the input datetime values + * + * @returns cudf::column of the same datetime resolution as the input column + * @throw cudf::logic_error if input column datatype is not TIMESTAMP + */ +std::unique_ptr ceil_millisecond(column_view const& column, + rmm::mr::device_memory_resource* mr); + +/** + * @brief Round up to the nearest microsecond + * + * @param[in] cudf::column_view of the input datetime values + * + * @returns cudf::column of the same datetime resolution as the input column + * @throw cudf::logic_error if input column datatype is not TIMESTAMP + */ +std::unique_ptr ceil_microsecond(column_view const& column, + rmm::mr::device_memory_resource* mr); + +/** + * @brief Round up to the nearest nanosecond + * + * @param[in] cudf::column_view of the input datetime values + * + * @returns cudf::column of the same datetime resolution as the input column + * @throw cudf::logic_error if input column datatype is not TIMESTAMP + */ +std::unique_ptr ceil_nanosecond(column_view const& column, + rmm::mr::device_memory_resource* mr); } // namespace datetime } // namespace cudf diff --git a/cpp/include/cudf/wrappers/durations.hpp b/cpp/include/cudf/wrappers/durations.hpp index 07bcc1976a8..8bc8b7a7e6e 100644 --- a/cpp/include/cudf/wrappers/durations.hpp +++ b/cpp/include/cudf/wrappers/durations.hpp @@ -33,6 +33,14 @@ namespace cudf { * @brief Type alias representing an int32_t duration of days. */ using duration_D = cuda::std::chrono::duration; +/** + * @brief Type alias representing an int32_t duration of hours. + */ +using duration_h = cuda::std::chrono::duration; +/** + * @brief Type alias representing an int32_t duration of minutes. + */ +using duration_m = cuda::std::chrono::duration; /** * @brief Type alias representing an int64_t duration of seconds. */ diff --git a/cpp/src/datetime/datetime_ops.cu b/cpp/src/datetime/datetime_ops.cu index 216bb9a5c9a..5c35078cccd 100644 --- a/cpp/src/datetime/datetime_ops.cu +++ b/cpp/src/datetime/datetime_ops.cu @@ -25,6 +25,7 @@ #include #include #include +#include #include #include @@ -44,6 +45,9 @@ enum class datetime_component { HOUR, MINUTE, SECOND, + MILLISECOND, + MICROSECOND, + NANOSECOND }; template @@ -80,18 +84,29 @@ struct extract_component_operator { } }; -template +template struct ceil_timestamp { + template CUDA_DEVICE_CALLABLE Timestamp operator()(Timestamp const ts) const { using namespace cuda::std::chrono; + // want to use this with D, H, T (minute), S, L (millisecond), U switch (Component) { case datetime_component::DAY: - return time_point_cast( - cuda::std::chrono::ceil(ts)); - default: - return time_point_cast( - cuda::std::chrono::ceil(ts)); + return time_point_cast(ceil(ts)); + case datetime_component::HOUR: + return time_point_cast(ceil(ts)); + case datetime_component::MINUTE: + return time_point_cast(ceil(ts)); + case datetime_component::SECOND: + return time_point_cast(ceil(ts)); + case datetime_component::MILLISECOND: + return time_point_cast(ceil(ts)); + case datetime_component::MICROSECOND: + return time_point_cast(ceil(ts)); + case datetime_component::NANOSECOND: + return time_point_cast(ceil(ts)); + default: return time_point_cast(ceil(ts)); } } }; @@ -158,7 +173,7 @@ struct is_leap_year_op { }; // Specific function for applying ceil/floor date ops -template +template struct dispatch_ceil { template std::enable_if_t(), std::unique_ptr> operator()( @@ -185,7 +200,7 @@ struct dispatch_ceil { column.begin(), column.end(), output->mutable_view().begin(), - ceil_timestamp{}); + TransformFunctor{}); return output; } @@ -350,12 +365,84 @@ std::unique_ptr add_calendrical_months(column_view const& timestamp_colu return output; } -std::unique_ptr ceil_days(column_view const& column, +std::unique_ptr ceil_day(column_view const& column, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + return cudf::type_dispatcher(column.type(), + dispatch_ceil>{}, + column, + stream, + mr); +} + +std::unique_ptr ceil_hour(column_view const& column, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) +{ + return cudf::type_dispatcher(column.type(), + dispatch_ceil>{}, + column, + stream, + mr); +} + +std::unique_ptr ceil_minute(column_view const& column, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + return cudf::type_dispatcher(column.type(), + dispatch_ceil>{}, + column, + stream, + mr); +} + +std::unique_ptr ceil_second(column_view const& column, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + return cudf::type_dispatcher(column.type(), + dispatch_ceil>{}, + column, + stream, + mr); +} + +std::unique_ptr ceil_millisecond(column_view const& column, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + return cudf::type_dispatcher( + column.type(), + dispatch_ceil>{}, + column, + stream, + mr); +} + +std::unique_ptr ceil_microsecond(column_view const& column, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) { return cudf::type_dispatcher( - column.type(), dispatch_ceil{}, column, stream, mr); + column.type(), + dispatch_ceil>{}, + column, + stream, + mr); +} + +std::unique_ptr ceil_nanosecond(column_view const& column, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + return cudf::type_dispatcher( + column.type(), + dispatch_ceil>{}, + column, + stream, + mr); } std::unique_ptr extract_year(column_view const& column, @@ -449,7 +536,46 @@ std::unique_ptr is_leap_year(column_view const& column, std::unique_ptr ceil_day(column_view const& column, rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); - return detail::ceil_days(column, rmm::cuda_stream_default, mr); + return detail::ceil_day(column, rmm::cuda_stream_default, mr); +} + +std::unique_ptr ceil_hour(column_view const& column, rmm::mr::device_memory_resource* mr) +{ + CUDF_FUNC_RANGE(); + return detail::ceil_hour(column, rmm::cuda_stream_default, mr); +} + +std::unique_ptr ceil_minute(column_view const& column, rmm::mr::device_memory_resource* mr) +{ + CUDF_FUNC_RANGE(); + return detail::ceil_minute(column, rmm::cuda_stream_default, mr); +} + +std::unique_ptr ceil_second(column_view const& column, rmm::mr::device_memory_resource* mr) +{ + CUDF_FUNC_RANGE(); + return detail::ceil_second(column, rmm::cuda_stream_default, mr); +} + +std::unique_ptr ceil_millisecond(column_view const& column, + rmm::mr::device_memory_resource* mr) +{ + CUDF_FUNC_RANGE(); + return detail::ceil_millisecond(column, rmm::cuda_stream_default, mr); +} + +std::unique_ptr ceil_microsecond(column_view const& column, + rmm::mr::device_memory_resource* mr) +{ + CUDF_FUNC_RANGE(); + return detail::ceil_microsecond(column, rmm::cuda_stream_default, mr); +} + +std::unique_ptr ceil_nanosecond(column_view const& column, + rmm::mr::device_memory_resource* mr) +{ + CUDF_FUNC_RANGE(); + return detail::ceil_nanosecond(column, rmm::cuda_stream_default, mr); } std::unique_ptr extract_year(column_view const& column, rmm::mr::device_memory_resource* mr) From 1c8e1d3f67002d47fb3302857ca0341d81780e92 Mon Sep 17 00:00:00 2001 From: Shane Ding Date: Sat, 7 Aug 2021 20:10:09 +0000 Subject: [PATCH 04/14] updated tests --- cpp/tests/datetime/datetime_ops_test.cpp | 42 ++++++++++++++++++++++++ 1 file changed, 42 insertions(+) diff --git a/cpp/tests/datetime/datetime_ops_test.cpp b/cpp/tests/datetime/datetime_ops_test.cpp index 2a98c79540a..513929ab8c3 100644 --- a/cpp/tests/datetime/datetime_ops_test.cpp +++ b/cpp/tests/datetime/datetime_ops_test.cpp @@ -347,6 +347,48 @@ TEST_F(BasicDatetimeOpsTest, TestLastDayOfMonthWithDate) verbosity); } +TYPED_TEST(TypedDatetimeOpsTest, TestCeilDatetime) +{ + // using T = TypeParam; + using namespace cudf::test; + using namespace cudf::datetime; + using namespace cuda::std::chrono; + + auto start = milliseconds(-2500000000000); // Sat, 11 Oct 1890 19:33:20 GMT + auto stop_ = milliseconds(2500000000000); // Mon, 22 Mar 2049 04:26:40 GMT + auto timestamps = generate_timestamps( + this->size(), time_point_ms(start), time_point_ms(stop_)); + + std::vector expected_day = {time_point_ms(milliseconds(-2499984000000L)), + time_point_ms(milliseconds(-1999987200000L))}; + + auto expected_ceil_day = fixed_width_column_wrapper{ + -2499984000000L, // 1890-10-12 + -1999987200000L, // 1906-8-17 + -1499990400000L, // 1922-6-21 + -999993600000L, // 1938-4-25 + -499996800000L, // 1954-2-27 + 0L, // 1970-1-1 (exactly midnight) + 500083200000L, // 1985-11-6 + 1000080000000L, // 2001-9-10 + 1500076800000L, // 2017-7-15 + 2000073600000L // 2033-5-19 + }; + // auto expected_weekdays = fixed_width_column_wrapper{6, 4, 2, 7, 5, 4, 2, 7, 5, 3}; + // auto expected_hours = fixed_width_column_wrapper{19, 20, 21, 22, 23, 0, 0, 1, 2, + // 3}; auto expected_minutes = fixed_width_column_wrapper{33, 26, 20, 13, 6, 0, 53, 46, + // 40, 33}; auto expected_seconds = fixed_width_column_wrapper{20, 40, 0, 20, 40, 0, 20, + // 40, 0, 20}; + + // CUDF_TEST_EXPECT_COLUMNS_EQUAL(*extract_year(timestamps), expected_years); + // CUDF_TEST_EXPECT_COLUMNS_EQUAL(*extract_month(timestamps), expected_months); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*ceil_day(timestamps), expected_ceil_day); + // CUDF_TEST_EXPECT_COLUMNS_EQUAL(*extract_weekday(timestamps), expected_weekdays); + // CUDF_TEST_EXPECT_COLUMNS_EQUAL(*extract_hour(timestamps), expected_hours); + // CUDF_TEST_EXPECT_COLUMNS_EQUAL(*extract_minute(timestamps), expected_minutes); + // CUDF_TEST_EXPECT_COLUMNS_EQUAL(*extract_second(timestamps), expected_seconds); +} + TEST_F(BasicDatetimeOpsTest, TestCeilDay) { using namespace cudf::test; From 46d5b3e2e0defd4fc1a7f18ceee1b856c6ff711f Mon Sep 17 00:00:00 2001 From: Shane Ding Date: Wed, 11 Aug 2021 17:46:29 +0000 Subject: [PATCH 05/14] added testcases --- cpp/include/cudf/datetime.hpp | 15 ++-- cpp/tests/datetime/datetime_ops_test.cpp | 99 ++++++++++++++++-------- 2 files changed, 76 insertions(+), 38 deletions(-) diff --git a/cpp/include/cudf/datetime.hpp b/cpp/include/cudf/datetime.hpp index cf9161e4c58..d51bd06b77d 100644 --- a/cpp/include/cudf/datetime.hpp +++ b/cpp/include/cudf/datetime.hpp @@ -264,8 +264,9 @@ std::unique_ptr ceil_second( * @returns cudf::column of the same datetime resolution as the input column * @throw cudf::logic_error if input column datatype is not TIMESTAMP */ -std::unique_ptr ceil_millisecond(column_view const& column, - rmm::mr::device_memory_resource* mr); +std::unique_ptr ceil_millisecond( + column_view const& column, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** * @brief Round up to the nearest microsecond @@ -275,8 +276,9 @@ std::unique_ptr ceil_millisecond(column_view const& column, * @returns cudf::column of the same datetime resolution as the input column * @throw cudf::logic_error if input column datatype is not TIMESTAMP */ -std::unique_ptr ceil_microsecond(column_view const& column, - rmm::mr::device_memory_resource* mr); +std::unique_ptr ceil_microsecond( + column_view const& column, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** * @brief Round up to the nearest nanosecond @@ -286,7 +288,8 @@ std::unique_ptr ceil_microsecond(column_view const& column, * @returns cudf::column of the same datetime resolution as the input column * @throw cudf::logic_error if input column datatype is not TIMESTAMP */ -std::unique_ptr ceil_nanosecond(column_view const& column, - rmm::mr::device_memory_resource* mr); +std::unique_ptr ceil_nanosecond( + column_view const& column, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); } // namespace datetime } // namespace cudf diff --git a/cpp/tests/datetime/datetime_ops_test.cpp b/cpp/tests/datetime/datetime_ops_test.cpp index 513929ab8c3..f99532b1a3b 100644 --- a/cpp/tests/datetime/datetime_ops_test.cpp +++ b/cpp/tests/datetime/datetime_ops_test.cpp @@ -354,39 +354,74 @@ TYPED_TEST(TypedDatetimeOpsTest, TestCeilDatetime) using namespace cudf::datetime; using namespace cuda::std::chrono; - auto start = milliseconds(-2500000000000); // Sat, 11 Oct 1890 19:33:20 GMT - auto stop_ = milliseconds(2500000000000); // Mon, 22 Mar 2049 04:26:40 GMT - auto timestamps = generate_timestamps( - this->size(), time_point_ms(start), time_point_ms(stop_)); - - std::vector expected_day = {time_point_ms(milliseconds(-2499984000000L)), - time_point_ms(milliseconds(-1999987200000L))}; - - auto expected_ceil_day = fixed_width_column_wrapper{ - -2499984000000L, // 1890-10-12 - -1999987200000L, // 1906-8-17 - -1499990400000L, // 1922-6-21 - -999993600000L, // 1938-4-25 - -499996800000L, // 1954-2-27 - 0L, // 1970-1-1 (exactly midnight) - 500083200000L, // 1985-11-6 - 1000080000000L, // 2001-9-10 - 1500076800000L, // 2017-7-15 - 2000073600000L // 2033-5-19 + std::vector timestamps{ + 118800987L, // 1970-01-02 09:00:00 + 129390186056L, // 2011-01-01 17:11:00 + 318402007832L, // 1980-02-03 05:00:00 + 604996200000L, // 1989-03-04 06:30:00 + 1270413572078L, // 2010-04-04 20:39:32 + 1588734621000L, // 2020-05-06 03:10:21 + 2550814152105L, // 2050-10-31 07:29:12 + 4102518778005L, // 2100-01-01 20:32:58 + 702696234000L, // 1992-04-08 01:23:54 + 6516816203020L, // 2176-07-05 02:43:23 + 26472091292001L, // 2808-11-12 22:41:32 + 4133857172000L, // 2100-12-30 01:39:32 + 1560248892060L, // 2019-06-19 12:54:52 + 4115217620600L, // 2100-05-28 20:00:00 + -265880250000L, // 1961-07-29 16:22:30 }; - // auto expected_weekdays = fixed_width_column_wrapper{6, 4, 2, 7, 5, 4, 2, 7, 5, 3}; - // auto expected_hours = fixed_width_column_wrapper{19, 20, 21, 22, 23, 0, 0, 1, 2, - // 3}; auto expected_minutes = fixed_width_column_wrapper{33, 26, 20, 13, 6, 0, 53, 46, - // 40, 33}; auto expected_seconds = fixed_width_column_wrapper{20, 40, 0, 20, 40, 0, 20, - // 40, 0, 20}; - - // CUDF_TEST_EXPECT_COLUMNS_EQUAL(*extract_year(timestamps), expected_years); - // CUDF_TEST_EXPECT_COLUMNS_EQUAL(*extract_month(timestamps), expected_months); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(*ceil_day(timestamps), expected_ceil_day); - // CUDF_TEST_EXPECT_COLUMNS_EQUAL(*extract_weekday(timestamps), expected_weekdays); - // CUDF_TEST_EXPECT_COLUMNS_EQUAL(*extract_hour(timestamps), expected_hours); - // CUDF_TEST_EXPECT_COLUMNS_EQUAL(*extract_minute(timestamps), expected_minutes); - // CUDF_TEST_EXPECT_COLUMNS_EQUAL(*extract_second(timestamps), expected_seconds); + + auto input = fixed_width_column_wrapper( + timestamps.begin(), timestamps.end()); + + std::vector ms_timestamp(timestamps.size()); + std::transform(timestamps.begin(), timestamps.end(), ms_timestamp.begin(), [](auto i) { + return time_point_ms(milliseconds(i)); + }); + + std::vector ceiled_day(timestamps.size()); + std::transform(ms_timestamp.begin(), ms_timestamp.end(), ceiled_day.begin(), [](auto i) { + return ceil(i); + }); + + std::vector ceiled_hour(timestamps.size()); + std::transform(ms_timestamp.begin(), ms_timestamp.end(), ceiled_hour.begin(), [](auto i) { + return ceil(i); + }); + + std::vector ceiled_minute(timestamps.size()); + std::transform(ms_timestamp.begin(), ms_timestamp.end(), ceiled_minute.begin(), [](auto i) { + return ceil(i); + }); + + std::vector ceiled_second(timestamps.size()); + std::transform(ms_timestamp.begin(), ms_timestamp.end(), ceiled_second.begin(), [](auto i) { + return ceil(i); + }); + + std::vector ceiled_millisecond(timestamps.size()); + std::transform(ms_timestamp.begin(), ms_timestamp.end(), ceiled_millisecond.begin(), [](auto i) { + return ceil(i); + }); + + auto expected_day = + fixed_width_column_wrapper(ceiled_day.begin(), ceiled_day.end()); + auto expected_hour = + fixed_width_column_wrapper(ceiled_hour.begin(), ceiled_hour.end()); + auto expected_minute = + fixed_width_column_wrapper(ceiled_minute.begin(), ceiled_minute.end()); + auto expected_second = + fixed_width_column_wrapper(ceiled_second.begin(), ceiled_second.end()); + auto expected_millisecond = fixed_width_column_wrapper( + ceiled_millisecond.begin(), ceiled_millisecond.end()); + + // auto res = *ceil_millisecond(input); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*ceil_day(input), expected_day); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*ceil_hour(input), expected_hour); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*ceil_minute(input), expected_minute); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*ceil_second(input), expected_second); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*ceil_millisecond(input), expected_millisecond); } TEST_F(BasicDatetimeOpsTest, TestCeilDay) From 659cbf6135eb9cc7b5ffccd44ce17b9ae8b80a2a Mon Sep 17 00:00:00 2001 From: Shane Ding Date: Wed, 11 Aug 2021 17:59:13 +0000 Subject: [PATCH 06/14] added testcases --- cpp/tests/datetime/datetime_ops_test.cpp | 33 ------------------------ 1 file changed, 33 deletions(-) diff --git a/cpp/tests/datetime/datetime_ops_test.cpp b/cpp/tests/datetime/datetime_ops_test.cpp index f99532b1a3b..5b5e57f0832 100644 --- a/cpp/tests/datetime/datetime_ops_test.cpp +++ b/cpp/tests/datetime/datetime_ops_test.cpp @@ -349,7 +349,6 @@ TEST_F(BasicDatetimeOpsTest, TestLastDayOfMonthWithDate) TYPED_TEST(TypedDatetimeOpsTest, TestCeilDatetime) { - // using T = TypeParam; using namespace cudf::test; using namespace cudf::datetime; using namespace cuda::std::chrono; @@ -416,7 +415,6 @@ TYPED_TEST(TypedDatetimeOpsTest, TestCeilDatetime) auto expected_millisecond = fixed_width_column_wrapper( ceiled_millisecond.begin(), ceiled_millisecond.end()); - // auto res = *ceil_millisecond(input); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*ceil_day(input), expected_day); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*ceil_hour(input), expected_hour); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*ceil_minute(input), expected_minute); @@ -424,37 +422,6 @@ TYPED_TEST(TypedDatetimeOpsTest, TestCeilDatetime) CUDF_TEST_EXPECT_COLUMNS_EQUAL(*ceil_millisecond(input), expected_millisecond); } -TEST_F(BasicDatetimeOpsTest, TestCeilDay) -{ - using namespace cudf::test; - using namespace cudf::datetime; - using namespace cuda::std::chrono; - - auto timestamps_s = - cudf::test::fixed_width_column_wrapper{ - { - // null - 915148800L, // 1999-01-01 00:00:00 GMT - non leap year - -11663029161L, // 1600-5-31 05:40:39 GMT - leap year - 707904541L, // 1992-06-07 08:09:01 GMT - leap year - 2181048447L, // 1900-11-20 09:12:33 GMT - non leap year - }, - {true, true, true, true}}; - - CUDF_TEST_EXPECT_COLUMNS_EQUAL( - *ceil_day(timestamps_s), - fixed_width_column_wrapper{ - { - // null - 915148800L, // 1999-01-01 00:00:00 GMT - non leap year - -11662963200L, // 1600-5-31 05:40:39 GMT - leap year - 707961600L, // 1992-06-07 08:09:01 GMT - leap year - 2181081600L, // 1900-11-20 09:12:33 GMT - non leap year - }, - {true, true, true, true}}, - verbosity); -} - TEST_F(BasicDatetimeOpsTest, TestDayOfYearWithDate) { using namespace cudf::test; From 20f7a0a27a042000786f2a0da0de19cf0e64fa79 Mon Sep 17 00:00:00 2001 From: Shane Ding Date: Thu, 12 Aug 2021 18:02:14 +0000 Subject: [PATCH 07/14] refactored test cases --- cpp/tests/datetime/datetime_ops_test.cpp | 87 +++++++++--------------- 1 file changed, 34 insertions(+), 53 deletions(-) diff --git a/cpp/tests/datetime/datetime_ops_test.cpp b/cpp/tests/datetime/datetime_ops_test.cpp index 5b5e57f0832..6f620796ee8 100644 --- a/cpp/tests/datetime/datetime_ops_test.cpp +++ b/cpp/tests/datetime/datetime_ops_test.cpp @@ -349,76 +349,57 @@ TEST_F(BasicDatetimeOpsTest, TestLastDayOfMonthWithDate) TYPED_TEST(TypedDatetimeOpsTest, TestCeilDatetime) { + using T = TypeParam; using namespace cudf::test; using namespace cudf::datetime; using namespace cuda::std::chrono; - std::vector timestamps{ - 118800987L, // 1970-01-02 09:00:00 - 129390186056L, // 2011-01-01 17:11:00 - 318402007832L, // 1980-02-03 05:00:00 - 604996200000L, // 1989-03-04 06:30:00 - 1270413572078L, // 2010-04-04 20:39:32 - 1588734621000L, // 2020-05-06 03:10:21 - 2550814152105L, // 2050-10-31 07:29:12 - 4102518778005L, // 2100-01-01 20:32:58 - 702696234000L, // 1992-04-08 01:23:54 - 6516816203020L, // 2176-07-05 02:43:23 - 26472091292001L, // 2808-11-12 22:41:32 - 4133857172000L, // 2100-12-30 01:39:32 - 1560248892060L, // 2019-06-19 12:54:52 - 4115217620600L, // 2100-05-28 20:00:00 - -265880250000L, // 1961-07-29 16:22:30 - }; + auto start = milliseconds(-2500000000000); // Sat, 11 Oct 1890 19:33:20 GMT + auto stop_ = milliseconds(2500000000000); // Mon, 22 Mar 2049 04:26:40 GMT - auto input = fixed_width_column_wrapper( - timestamps.begin(), timestamps.end()); + auto input = generate_timestamps(this->size(), time_point_ms(start), time_point_ms(stop_)); - std::vector ms_timestamp(timestamps.size()); - std::transform(timestamps.begin(), timestamps.end(), ms_timestamp.begin(), [](auto i) { - return time_point_ms(milliseconds(i)); - }); + auto host_val = to_host(input); + thrust::host_vector timestamps = host_val.first; - std::vector ceiled_day(timestamps.size()); - std::transform(ms_timestamp.begin(), ms_timestamp.end(), ceiled_day.begin(), [](auto i) { - return ceil(i); + thrust::host_vector ceiled_day(timestamps.size()); + thrust::transform(timestamps.begin(), timestamps.end(), ceiled_day.begin(), [](auto i) { + return time_point_cast(ceil(i)); }); + auto expected_day = + fixed_width_column_wrapper(ceiled_day.begin(), ceiled_day.end()); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*ceil_day(input), expected_day); - std::vector ceiled_hour(timestamps.size()); - std::transform(ms_timestamp.begin(), ms_timestamp.end(), ceiled_hour.begin(), [](auto i) { - return ceil(i); + thrust::host_vector ceiled_hour(timestamps.size()); + thrust::transform(timestamps.begin(), timestamps.end(), ceiled_hour.begin(), [](auto i) { + return time_point_cast(ceil(i)); }); + auto expected_hour = fixed_width_column_wrapper(ceiled_hour.begin(), + ceiled_hour.end()); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*ceil_hour(input), expected_hour); - std::vector ceiled_minute(timestamps.size()); - std::transform(ms_timestamp.begin(), ms_timestamp.end(), ceiled_minute.begin(), [](auto i) { - return ceil(i); + std::vector ceiled_minute(timestamps.size()); + std::transform(timestamps.begin(), timestamps.end(), ceiled_minute.begin(), [](auto i) { + return time_point_cast(ceil(i)); }); + auto expected_minute = fixed_width_column_wrapper( + ceiled_minute.begin(), ceiled_minute.end()); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*ceil_minute(input), expected_minute); - std::vector ceiled_second(timestamps.size()); - std::transform(ms_timestamp.begin(), ms_timestamp.end(), ceiled_second.begin(), [](auto i) { - return ceil(i); + std::vector ceiled_second(timestamps.size()); + std::transform(timestamps.begin(), timestamps.end(), ceiled_second.begin(), [](auto i) { + return time_point_cast(ceil(i)); }); + auto expected_second = fixed_width_column_wrapper( + ceiled_second.begin(), ceiled_second.end()); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*ceil_second(input), expected_second); - std::vector ceiled_millisecond(timestamps.size()); - std::transform(ms_timestamp.begin(), ms_timestamp.end(), ceiled_millisecond.begin(), [](auto i) { - return ceil(i); + std::vector ceiled_millisecond(timestamps.size()); + std::transform(timestamps.begin(), timestamps.end(), ceiled_millisecond.begin(), [](auto i) { + return time_point_cast(ceil(i)); }); - - auto expected_day = - fixed_width_column_wrapper(ceiled_day.begin(), ceiled_day.end()); - auto expected_hour = - fixed_width_column_wrapper(ceiled_hour.begin(), ceiled_hour.end()); - auto expected_minute = - fixed_width_column_wrapper(ceiled_minute.begin(), ceiled_minute.end()); - auto expected_second = - fixed_width_column_wrapper(ceiled_second.begin(), ceiled_second.end()); - auto expected_millisecond = fixed_width_column_wrapper( + auto expected_millisecond = fixed_width_column_wrapper( ceiled_millisecond.begin(), ceiled_millisecond.end()); - - CUDF_TEST_EXPECT_COLUMNS_EQUAL(*ceil_day(input), expected_day); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(*ceil_hour(input), expected_hour); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(*ceil_minute(input), expected_minute); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(*ceil_second(input), expected_second); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*ceil_millisecond(input), expected_millisecond); } From bee2578f597cb7623e37b84c4d6aa34c8fb1e64f Mon Sep 17 00:00:00 2001 From: Shane Ding Date: Tue, 17 Aug 2021 15:00:20 +0000 Subject: [PATCH 08/14] refactored code --- cpp/include/cudf/datetime.hpp | 5 +- cpp/src/datetime/datetime_ops.cu | 169 +++++++++++++++---------------- 2 files changed, 82 insertions(+), 92 deletions(-) diff --git a/cpp/include/cudf/datetime.hpp b/cpp/include/cudf/datetime.hpp index d51bd06b77d..76c9dfd8a39 100644 --- a/cpp/include/cudf/datetime.hpp +++ b/cpp/include/cudf/datetime.hpp @@ -209,7 +209,7 @@ std::unique_ptr is_leap_year( /** @} */ // end of group /** - * @brief Round up to the nearest day + * @brief Round up to the nearest day * * @param[in] cudf::column_view of the input datetime values * @@ -221,7 +221,7 @@ std::unique_ptr ceil_day( rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** - * @brief Round up to the nearest hour + * @brief Round up to the nearest hour * * @param[in] cudf::column_view of the input datetime values * @@ -291,5 +291,6 @@ std::unique_ptr ceil_microsecond( std::unique_ptr ceil_nanosecond( column_view const& column, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + } // namespace datetime } // namespace cudf diff --git a/cpp/src/datetime/datetime_ops.cu b/cpp/src/datetime/datetime_ops.cu index 5c35078cccd..12e5198b758 100644 --- a/cpp/src/datetime/datetime_ops.cu +++ b/cpp/src/datetime/datetime_ops.cu @@ -24,14 +24,13 @@ #include #include #include +#include #include +#include #include -#include #include #include -#include "cudf/utilities/error.hpp" -#include "cudf/utilities/type_dispatcher.hpp" namespace cudf { namespace datetime { @@ -84,14 +83,14 @@ struct extract_component_operator { } }; -template +template struct ceil_timestamp { template CUDA_DEVICE_CALLABLE Timestamp operator()(Timestamp const ts) const { using namespace cuda::std::chrono; // want to use this with D, H, T (minute), S, L (millisecond), U - switch (Component) { + switch (COMPONENT) { case datetime_component::DAY: return time_point_cast(ceil(ts)); case datetime_component::HOUR: @@ -177,7 +176,6 @@ template struct dispatch_ceil { template std::enable_if_t(), std::unique_ptr> operator()( - // cudf::data_type type, cudf::column_view const& column, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) const @@ -365,84 +363,68 @@ std::unique_ptr add_calendrical_months(column_view const& timestamp_colu return output; } -std::unique_ptr ceil_day(column_view const& column, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) -{ - return cudf::type_dispatcher(column.type(), - dispatch_ceil>{}, - column, - stream, - mr); -} - -std::unique_ptr ceil_hour(column_view const& column, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) -{ - return cudf::type_dispatcher(column.type(), - dispatch_ceil>{}, - column, - stream, - mr); -} - -std::unique_ptr ceil_minute(column_view const& column, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) -{ - return cudf::type_dispatcher(column.type(), - dispatch_ceil>{}, - column, - stream, - mr); -} - -std::unique_ptr ceil_second(column_view const& column, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) -{ - return cudf::type_dispatcher(column.type(), - dispatch_ceil>{}, - column, - stream, - mr); -} - -std::unique_ptr ceil_millisecond(column_view const& column, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) -{ - return cudf::type_dispatcher( - column.type(), - dispatch_ceil>{}, - column, - stream, - mr); -} - -std::unique_ptr ceil_microsecond(column_view const& column, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) -{ - return cudf::type_dispatcher( - column.type(), - dispatch_ceil>{}, - column, - stream, - mr); -} - -std::unique_ptr ceil_nanosecond(column_view const& column, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) +std::unique_ptr ceil_general(column_view const& column, + datetime_component Component, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) { - return cudf::type_dispatcher( - column.type(), - dispatch_ceil>{}, - column, - stream, - mr); + switch (Component) { + case datetime_component::DAY: + return cudf::type_dispatcher(column.type(), + dispatch_ceil>{}, + column, + stream, + mr); + case datetime_component::HOUR: + return cudf::type_dispatcher( + column.type(), + dispatch_ceil>{}, + column, + stream, + mr); + case datetime_component::MINUTE: + return cudf::type_dispatcher( + column.type(), + dispatch_ceil>{}, + column, + stream, + mr); + case datetime_component::SECOND: + return cudf::type_dispatcher( + column.type(), + dispatch_ceil>{}, + column, + stream, + mr); + case datetime_component::MILLISECOND: + return cudf::type_dispatcher( + column.type(), + dispatch_ceil>{}, + column, + stream, + mr); + case datetime_component::MICROSECOND: + return cudf::type_dispatcher( + column.type(), + dispatch_ceil>{}, + column, + stream, + mr); + case datetime_component::NANOSECOND: + return cudf::type_dispatcher( + column.type(), + dispatch_ceil>{}, + column, + stream, + mr); + default: + return cudf::type_dispatcher( + column.type(), + dispatch_ceil>{}, + column, + stream, + mr); + } } std::unique_ptr extract_year(column_view const& column, @@ -536,46 +518,53 @@ std::unique_ptr is_leap_year(column_view const& column, std::unique_ptr ceil_day(column_view const& column, rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); - return detail::ceil_day(column, rmm::cuda_stream_default, mr); + return detail::ceil_general( + column, detail::datetime_component::DAY, rmm::cuda_stream_default, mr); } std::unique_ptr ceil_hour(column_view const& column, rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); - return detail::ceil_hour(column, rmm::cuda_stream_default, mr); + return detail::ceil_general( + column, detail::datetime_component::HOUR, rmm::cuda_stream_default, mr); } std::unique_ptr ceil_minute(column_view const& column, rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); - return detail::ceil_minute(column, rmm::cuda_stream_default, mr); + return detail::ceil_general( + column, detail::datetime_component::MINUTE, rmm::cuda_stream_default, mr); } std::unique_ptr ceil_second(column_view const& column, rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); - return detail::ceil_second(column, rmm::cuda_stream_default, mr); + return detail::ceil_general( + column, detail::datetime_component::SECOND, rmm::cuda_stream_default, mr); } std::unique_ptr ceil_millisecond(column_view const& column, rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); - return detail::ceil_millisecond(column, rmm::cuda_stream_default, mr); + return detail::ceil_general( + column, detail::datetime_component::MILLISECOND, rmm::cuda_stream_default, mr); } std::unique_ptr ceil_microsecond(column_view const& column, rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); - return detail::ceil_microsecond(column, rmm::cuda_stream_default, mr); + return detail::ceil_general( + column, detail::datetime_component::MICROSECOND, rmm::cuda_stream_default, mr); } std::unique_ptr ceil_nanosecond(column_view const& column, rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); - return detail::ceil_nanosecond(column, rmm::cuda_stream_default, mr); + return detail::ceil_general( + column, detail::datetime_component::NANOSECOND, rmm::cuda_stream_default, mr); } std::unique_ptr extract_year(column_view const& column, rmm::mr::device_memory_resource* mr) From 4af1707a5f4e67b715d5b2c29d6c54a93212c4ea Mon Sep 17 00:00:00 2001 From: Shane Ding Date: Tue, 17 Aug 2021 15:01:27 +0000 Subject: [PATCH 09/14] refactored code --- cpp/src/datetime/datetime_ops.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/src/datetime/datetime_ops.cu b/cpp/src/datetime/datetime_ops.cu index 12e5198b758..b2c8108aa28 100644 --- a/cpp/src/datetime/datetime_ops.cu +++ b/cpp/src/datetime/datetime_ops.cu @@ -105,7 +105,7 @@ struct ceil_timestamp { return time_point_cast(ceil(ts)); case datetime_component::NANOSECOND: return time_point_cast(ceil(ts)); - default: return time_point_cast(ceil(ts)); + default: return time_point_cast(ceil(ts)); } } }; From 9ba60cc1d937315a7f1f455cfdefe43fe1658f39 Mon Sep 17 00:00:00 2001 From: Shane Ding Date: Tue, 17 Aug 2021 15:12:25 +0000 Subject: [PATCH 10/14] updated doc --- cpp/include/cudf/datetime.hpp | 38 +++++++++++++++++------------------ 1 file changed, 19 insertions(+), 19 deletions(-) diff --git a/cpp/include/cudf/datetime.hpp b/cpp/include/cudf/datetime.hpp index 76c9dfd8a39..c46aa20162c 100644 --- a/cpp/include/cudf/datetime.hpp +++ b/cpp/include/cudf/datetime.hpp @@ -211,10 +211,10 @@ std::unique_ptr is_leap_year( /** * @brief Round up to the nearest day * - * @param[in] cudf::column_view of the input datetime values + * @param cudf::column_view of the input datetime values * - * @returns cudf::column of the same datetime resolution as the input column * @throw cudf::logic_error if input column datatype is not TIMESTAMP + * @return cudf::column of the same datetime resolution as the input column */ std::unique_ptr ceil_day( cudf::column_view const& column, @@ -223,70 +223,70 @@ std::unique_ptr ceil_day( /** * @brief Round up to the nearest hour * - * @param[in] cudf::column_view of the input datetime values + * @param cudf::column_view of the input datetime values * - * @returns cudf::column of the same datetime resolution as the input column * @throw cudf::logic_error if input column datatype is not TIMESTAMP + * @return cudf::column of the same datetime resolution as the input column */ std::unique_ptr ceil_hour( cudf::column_view const& column, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** - * @brief Round up to the nearest minute + * @brief Round up to the nearest minute * - * @param[in] cudf::column_view of the input datetime values + * @param cudf::column_view of the input datetime values * - * @returns cudf::column of the same datetime resolution as the input column * @throw cudf::logic_error if input column datatype is not TIMESTAMP + * @return cudf::column of the same datetime resolution as the input column */ std::unique_ptr ceil_minute( cudf::column_view const& column, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** - * @brief Round up to the nearest second + * @brief Round up to the nearest second * - * @param[in] cudf::column_view of the input datetime values + * @param cudf::column_view of the input datetime values * - * @returns cudf::column of the same datetime resolution as the input column * @throw cudf::logic_error if input column datatype is not TIMESTAMP + * @return cudf::column of the same datetime resolution as the input column */ std::unique_ptr ceil_second( cudf::column_view const& column, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** - * @brief Round up to the nearest millisecond + * @brief Round up to the nearest millisecond * - * @param[in] cudf::column_view of the input datetime values + * @param cudf::column_view of the input datetime values * - * @returns cudf::column of the same datetime resolution as the input column * @throw cudf::logic_error if input column datatype is not TIMESTAMP + * @return cudf::column of the same datetime resolution as the input column */ std::unique_ptr ceil_millisecond( column_view const& column, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** - * @brief Round up to the nearest microsecond + * @brief Round up to the nearest microsecond * - * @param[in] cudf::column_view of the input datetime values + * @param cudf::column_view of the input datetime values * - * @returns cudf::column of the same datetime resolution as the input column * @throw cudf::logic_error if input column datatype is not TIMESTAMP + * @return cudf::column of the same datetime resolution as the input column */ std::unique_ptr ceil_microsecond( column_view const& column, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** - * @brief Round up to the nearest nanosecond + * @brief Round up to the nearest nanosecond * - * @param[in] cudf::column_view of the input datetime values + * @param cudf::column_view of the input datetime values * - * @returns cudf::column of the same datetime resolution as the input column * @throw cudf::logic_error if input column datatype is not TIMESTAMP + * @return cudf::column of the same datetime resolution as the input column */ std::unique_ptr ceil_nanosecond( column_view const& column, From 3e8e88457dbe14530fc994c94772dd1de6b7a93f Mon Sep 17 00:00:00 2001 From: Shane Ding Date: Wed, 18 Aug 2021 17:23:32 +0000 Subject: [PATCH 11/14] updated format --- cpp/src/datetime/datetime_ops.cu | 13 +++---------- 1 file changed, 3 insertions(+), 10 deletions(-) diff --git a/cpp/src/datetime/datetime_ops.cu b/cpp/src/datetime/datetime_ops.cu index b2c8108aa28..c5fcb7d1af0 100644 --- a/cpp/src/datetime/datetime_ops.cu +++ b/cpp/src/datetime/datetime_ops.cu @@ -180,7 +180,6 @@ struct dispatch_ceil { rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) const { - CUDF_EXPECTS(is_timestamp(column.type()), "Column type should be timestamp"); auto size = column.size(); auto output_col_type = data_type{cudf::type_to_id()}; @@ -364,11 +363,11 @@ std::unique_ptr add_calendrical_months(column_view const& timestamp_colu } std::unique_ptr ceil_general(column_view const& column, - datetime_component Component, + datetime_component component, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { - switch (Component) { + switch (component) { case datetime_component::DAY: return cudf::type_dispatcher(column.type(), dispatch_ceil>{}, @@ -417,13 +416,7 @@ std::unique_ptr ceil_general(column_view const& column, column, stream, mr); - default: - return cudf::type_dispatcher( - column.type(), - dispatch_ceil>{}, - column, - stream, - mr); + default: CUDF_FAIL("Unexpected resolution"); } } From ddf0a62c8a8a48fb19241d942512cd9806562363 Mon Sep 17 00:00:00 2001 From: Shane Ding Date: Fri, 20 Aug 2021 15:53:52 +0000 Subject: [PATCH 12/14] changed assert --- cpp/src/datetime/datetime_ops.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/src/datetime/datetime_ops.cu b/cpp/src/datetime/datetime_ops.cu index c5fcb7d1af0..041162c7034 100644 --- a/cpp/src/datetime/datetime_ops.cu +++ b/cpp/src/datetime/datetime_ops.cu @@ -105,7 +105,7 @@ struct ceil_timestamp { return time_point_cast(ceil(ts)); case datetime_component::NANOSECOND: return time_point_cast(ceil(ts)); - default: return time_point_cast(ceil(ts)); + default: cudf_assert(false && "Unexpected resolution"); ; } } }; From 3c5c9f96642bf2d1be562636374423066e6c9759 Mon Sep 17 00:00:00 2001 From: Shane Ding Date: Wed, 25 Aug 2021 06:45:47 +0000 Subject: [PATCH 13/14] refactored --- cpp/src/datetime/datetime_ops.cu | 85 +++++++------------------------- 1 file changed, 18 insertions(+), 67 deletions(-) diff --git a/cpp/src/datetime/datetime_ops.cu b/cpp/src/datetime/datetime_ops.cu index 041162c7034..b1b504681a0 100644 --- a/cpp/src/datetime/datetime_ops.cu +++ b/cpp/src/datetime/datetime_ops.cu @@ -105,7 +105,7 @@ struct ceil_timestamp { return time_point_cast(ceil(ts)); case datetime_component::NANOSECOND: return time_point_cast(ceil(ts)); - default: cudf_assert(false && "Unexpected resolution"); ; + default: cudf_assert(false && "Unexpected resolution"); } } }; @@ -362,62 +362,13 @@ std::unique_ptr add_calendrical_months(column_view const& timestamp_colu return output; } +template std::unique_ptr ceil_general(column_view const& column, - datetime_component component, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { - switch (component) { - case datetime_component::DAY: - return cudf::type_dispatcher(column.type(), - dispatch_ceil>{}, - column, - stream, - mr); - case datetime_component::HOUR: - return cudf::type_dispatcher( - column.type(), - dispatch_ceil>{}, - column, - stream, - mr); - case datetime_component::MINUTE: - return cudf::type_dispatcher( - column.type(), - dispatch_ceil>{}, - column, - stream, - mr); - case datetime_component::SECOND: - return cudf::type_dispatcher( - column.type(), - dispatch_ceil>{}, - column, - stream, - mr); - case datetime_component::MILLISECOND: - return cudf::type_dispatcher( - column.type(), - dispatch_ceil>{}, - column, - stream, - mr); - case datetime_component::MICROSECOND: - return cudf::type_dispatcher( - column.type(), - dispatch_ceil>{}, - column, - stream, - mr); - case datetime_component::NANOSECOND: - return cudf::type_dispatcher( - column.type(), - dispatch_ceil>{}, - column, - stream, - mr); - default: CUDF_FAIL("Unexpected resolution"); - } + return cudf::type_dispatcher( + column.type(), dispatch_ceil>{}, column, stream, mr); } std::unique_ptr extract_year(column_view const& column, @@ -511,53 +462,53 @@ std::unique_ptr is_leap_year(column_view const& column, std::unique_ptr ceil_day(column_view const& column, rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); - return detail::ceil_general( - column, detail::datetime_component::DAY, rmm::cuda_stream_default, mr); + return detail::ceil_general( + column, rmm::cuda_stream_default, mr); } std::unique_ptr ceil_hour(column_view const& column, rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); - return detail::ceil_general( - column, detail::datetime_component::HOUR, rmm::cuda_stream_default, mr); + return detail::ceil_general( + column, rmm::cuda_stream_default, mr); } std::unique_ptr ceil_minute(column_view const& column, rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); - return detail::ceil_general( - column, detail::datetime_component::MINUTE, rmm::cuda_stream_default, mr); + return detail::ceil_general( + column, rmm::cuda_stream_default, mr); } std::unique_ptr ceil_second(column_view const& column, rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); - return detail::ceil_general( - column, detail::datetime_component::SECOND, rmm::cuda_stream_default, mr); + return detail::ceil_general( + column, rmm::cuda_stream_default, mr); } std::unique_ptr ceil_millisecond(column_view const& column, rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); - return detail::ceil_general( - column, detail::datetime_component::MILLISECOND, rmm::cuda_stream_default, mr); + return detail::ceil_general( + column, rmm::cuda_stream_default, mr); } std::unique_ptr ceil_microsecond(column_view const& column, rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); - return detail::ceil_general( - column, detail::datetime_component::MICROSECOND, rmm::cuda_stream_default, mr); + return detail::ceil_general( + column, rmm::cuda_stream_default, mr); } std::unique_ptr ceil_nanosecond(column_view const& column, rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); - return detail::ceil_general( - column, detail::datetime_component::NANOSECOND, rmm::cuda_stream_default, mr); + return detail::ceil_general( + column, rmm::cuda_stream_default, mr); } std::unique_ptr extract_year(column_view const& column, rmm::mr::device_memory_resource* mr) From 12e23d79faa9b6fe31294541f1d7203419bacd12 Mon Sep 17 00:00:00 2001 From: Shane Ding Date: Wed, 25 Aug 2021 17:48:43 +0000 Subject: [PATCH 14/14] returned default value ffor ceil --- cpp/src/datetime/datetime_ops.cu | 2 ++ 1 file changed, 2 insertions(+) diff --git a/cpp/src/datetime/datetime_ops.cu b/cpp/src/datetime/datetime_ops.cu index b1b504681a0..914f9484056 100644 --- a/cpp/src/datetime/datetime_ops.cu +++ b/cpp/src/datetime/datetime_ops.cu @@ -107,6 +107,8 @@ struct ceil_timestamp { return time_point_cast(ceil(ts)); default: cudf_assert(false && "Unexpected resolution"); } + + return {}; } };