diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 8e237b08b128..85d622b49405 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -482,11 +482,6 @@ rtcx_embed_includes( cudf/cpp/src/join/jit INCLUDE_DIRECTORIES cudf/cpp/src ) -rtcx_embed_includes( - cudf_cuda_embed SOURCE_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/src/rolling DEST_DIRECTORY - cudf/cpp/src/rolling INCLUDE_DIRECTORIES cudf/cpp/src -) - rtcx_embed_includes( cudf_cuda_embed SOURCE_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/src/transform/jit DEST_DIRECTORY cudf/cpp/src/transform/jit INCLUDE_DIRECTORIES cudf/cpp/src diff --git a/cpp/include/cudf/aggregation.hpp b/cpp/include/cudf/aggregation.hpp index 727b62b9789a..31570b1c95ec 100644 --- a/cpp/include/cudf/aggregation.hpp +++ b/cpp/include/cudf/aggregation.hpp @@ -103,8 +103,6 @@ class aggregation { COLLECT_SET, ///< collect values into a list without duplicate entries LEAD, ///< window function, accesses row at specified offset following current row LAG, ///< window function, accesses row at specified offset preceding current row - PTX, ///< PTX based UDF aggregation - CUDA, ///< CUDA based UDF aggregation HOST_UDF, ///< host based UDF aggregation MERGE_LISTS, ///< merge multiple lists values into one list MERGE_SETS, ///< merge multiple lists values into one list then drop duplicate entries @@ -541,20 +539,6 @@ std::unique_ptr make_lag_aggregation(size_type offset); template std::unique_ptr make_lead_aggregation(size_type offset); -/** - * @brief Factory to create an aggregation base on UDF for PTX or CUDA - * - * @param[in] type The source type of the UDF aggregation - * @param[in] user_defined_aggregator A string containing the aggregator code - * @param[in] output_type expected output type - * - * @return An aggregation containing a user-defined aggregator string - */ -template -std::unique_ptr make_udf_aggregation(udf_source_type type, - std::string const& user_defined_aggregator, - data_type output_type); - // Forward declaration of `host_udf_base` for the factory function of `HOST_UDF` aggregation. class host_udf_base; diff --git a/cpp/include/cudf/detail/aggregation/aggregation.hpp b/cpp/include/cudf/detail/aggregation/aggregation.hpp index ae5078642720..01201eef9aed 100644 --- a/cpp/include/cudf/detail/aggregation/aggregation.hpp +++ b/cpp/include/cudf/detail/aggregation/aggregation.hpp @@ -581,52 +581,6 @@ class lead_lag_aggregation final [[nodiscard]] size_t hash_impl() const { return std::hash()(row_offset); } }; -/** - * @brief Derived class for specifying a custom aggregation - * specified in udf - */ -class udf_aggregation final : public clonable::derived_from { - public: - udf_aggregation(aggregation::Kind type, - std::string user_defined_aggregator, - data_type output_type) - : aggregation{type}, - _source{std::move(user_defined_aggregator)}, - _operator_name{(type == aggregation::PTX) ? "rolling_udf_ptx" : "rolling_udf_cuda"}, - _function_name{"GENERIC_ROLLING_OP"}, - _output_type{output_type} - { - CUDF_EXPECTS(type == aggregation::PTX or type == aggregation::CUDA, - "udf_aggregation can accept only PTX, CUDA"); - } - - [[nodiscard]] bool is_equal(aggregation const& _other) const override - { - if (!this->aggregation::is_equal(_other)) { return false; } - auto const& other = dynamic_cast(_other); - return (_source == other._source and _operator_name == other._operator_name and - _function_name == other._function_name and _output_type == other._output_type); - } - - [[nodiscard]] size_t do_hash() const override - { - return this->aggregation::do_hash() ^ hash_impl(); - } - - std::string const _source; - std::string const _operator_name; - std::string const _function_name; - data_type _output_type; - - protected: - [[nodiscard]] size_t hash_impl() const - { - return std::hash{}(_source) ^ std::hash{}(_operator_name) ^ - std::hash{}(_function_name) ^ - std::hash{}(static_cast(_output_type.id())); - } -}; - /** * @brief Derived class for specifying host-based UDF aggregation. */ diff --git a/cpp/src/aggregation/aggregation.cpp b/cpp/src/aggregation/aggregation.cpp index e25459f70e15..ea8e1fbf66a1 100644 --- a/cpp/src/aggregation/aggregation.cpp +++ b/cpp/src/aggregation/aggregation.cpp @@ -428,23 +428,6 @@ template CUDF_EXPORT std::unique_ptr make_lead_aggregation make_lead_aggregation(size_type offset); -/// Factory to create a UDF aggregation -template -std::unique_ptr make_udf_aggregation(udf_source_type type, - std::string const& user_defined_aggregator, - data_type output_type) -{ - auto* a = - new detail::udf_aggregation{type == udf_source_type::PTX ? aggregation::PTX : aggregation::CUDA, - user_defined_aggregator, - output_type}; - return std::unique_ptr(a); -} -template CUDF_EXPORT std::unique_ptr make_udf_aggregation( - udf_source_type type, std::string const& user_defined_aggregator, data_type output_type); -template CUDF_EXPORT std::unique_ptr make_udf_aggregation( - udf_source_type type, std::string const& user_defined_aggregator, data_type output_type); - /// Factory to create a MERGE_LISTS aggregation template std::unique_ptr make_merge_lists_aggregation() diff --git a/cpp/src/rolling/detail/rolling.cuh b/cpp/src/rolling/detail/rolling.cuh index 2a82ed5656c7..b42e558f05ce 100644 --- a/cpp/src/rolling/detail/rolling.cuh +++ b/cpp/src/rolling/detail/rolling.cuh @@ -46,20 +46,8 @@ namespace detail { static std::unique_ptr empty_output_for_rolling_aggregation(column_view const& input, rolling_aggregation const& agg) { - // TODO: - // Ideally, for UDF aggregations, the returned column would match - // the agg's return type. It currently returns empty_like(input), because: - // 1. This preserves prior behavior for empty input columns. - // 2. There is insufficient information to construct nested return columns. - // `cudf::make_udf_aggregation()` expresses the return type as a `data_type` - // which cannot express recursively nested types (e.g. `STRUCT>`.) - // 3. In any case, UDFs that return nested types are not currently supported. - // Constructing a more accurate return type for UDFs will be taken up - // at a later date. - return agg.kind == aggregation::CUDA || agg.kind == aggregation::PTX - ? empty_like(input) - : cudf::detail::dispatch_type_and_aggregation( - input.type(), agg.kind, agg_specific_empty_output{}, input, agg); + return cudf::detail::dispatch_type_and_aggregation( + input.type(), agg.kind, agg_specific_empty_output{}, input, agg); } /** diff --git a/cpp/src/rolling/detail/rolling_fixed_window.cu b/cpp/src/rolling/detail/rolling_fixed_window.cu index df50f1f30a82..6cb07148c836 100644 --- a/cpp/src/rolling/detail/rolling_fixed_window.cu +++ b/cpp/src/rolling/detail/rolling_fixed_window.cu @@ -4,7 +4,6 @@ */ #include "rolling.cuh" -#include "rolling_udf.cuh" #include "rolling_utils.cuh" #include @@ -37,24 +36,13 @@ std::unique_ptr rolling_window(column_view const& input, CUDF_EXPECTS(-(preceding_window - 1) <= following_window, "Preceding window bounds must precede the following window bounds."); - if (agg.kind == aggregation::CUDA || agg.kind == aggregation::PTX) { - // TODO: In future, might need to clamp preceding/following to column boundaries. - return cudf::detail::rolling_window_udf(input, - cudf::detail::fixed_window_wrapper(preceding_window), - cudf::detail::fixed_window_wrapper(following_window), - min_periods, - agg, - stream, - mr); - } else { - namespace utils = cudf::detail::rolling; - auto groups = utils::ungrouped{input.size()}; - auto preceding = - utils::make_clamped_window_iterator(preceding_window, groups); - auto following = - utils::make_clamped_window_iterator(following_window, groups); - return cudf::detail::rolling_window( - input, default_outputs, preceding, following, min_periods, agg, stream, mr); - } + namespace utils = cudf::detail::rolling; + auto groups = utils::ungrouped{input.size()}; + auto preceding = + utils::make_clamped_window_iterator(preceding_window, groups); + auto following = + utils::make_clamped_window_iterator(following_window, groups); + return cudf::detail::rolling_window( + input, default_outputs, preceding, following, min_periods, agg, stream, mr); } } // namespace cudf::detail diff --git a/cpp/src/rolling/detail/rolling_jit.cuh b/cpp/src/rolling/detail/rolling_jit.cuh deleted file mode 100644 index 4c32695134dc..000000000000 --- a/cpp/src/rolling/detail/rolling_jit.cuh +++ /dev/null @@ -1,109 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ - -#pragma once - -#include - -#include - -namespace cudf { -namespace detail { - -struct window_wrapper_base { - cudf::size_type const* group_offsets = nullptr; - cudf::size_type const* group_labels = nullptr; - cudf::size_type window = 0; -}; - -struct fixed_window_wrapper : public window_wrapper_base { - __device__ __host__ fixed_window_wrapper(cudf::size_type window) - : window_wrapper_base{nullptr, nullptr, window} - { - } - - __device__ __host__ fixed_window_wrapper(window_wrapper_base const& base) - : window_wrapper_base(base) - { - } - - __device__ __host__ cudf::size_type operator[](cudf::size_type) const { return window; } -}; - -struct variable_window_wrapper : public window_wrapper_base { - __device__ __host__ variable_window_wrapper(cudf::size_type const* group_offsets) - : window_wrapper_base{group_offsets, nullptr, 0} - { - } - - __device__ __host__ variable_window_wrapper(window_wrapper_base const& base) - : window_wrapper_base(base) - { - } - - __device__ __host__ cudf::size_type operator[](cudf::size_type idx) const - { - return group_offsets[idx]; - } -}; - -struct preceding_window_wrapper : public window_wrapper_base { - __device__ __host__ preceding_window_wrapper(cudf::size_type const* group_offsets, - cudf::size_type const* group_labels, - cudf::size_type window) - : window_wrapper_base{group_offsets, group_labels, window} - { - } - - __device__ __host__ preceding_window_wrapper(window_wrapper_base const& base) - : window_wrapper_base(base) - { - } - - __device__ cudf::size_type operator[](cudf::size_type idx) const - { - auto group_label = group_labels[idx]; - auto group_start = group_offsets[group_label]; - return cuda::std::min(window, idx - group_start + 1); // Preceding includes current row. - } -}; - -struct following_window_wrapper : public window_wrapper_base { - __device__ __host__ following_window_wrapper(cudf::size_type const* group_offsets, - cudf::size_type const* group_labels, - cudf::size_type window) - : window_wrapper_base{group_offsets, group_labels, window} - { - } - - __device__ __host__ following_window_wrapper(window_wrapper_base const& base) - : window_wrapper_base(base) - { - } - - __device__ cudf::size_type operator[](cudf::size_type idx) const - { - auto group_label = group_labels[idx]; - auto group_end = - group_offsets[group_label + - 1]; // Cannot fall off the end, since offsets is capped with `input.size()`. - return cuda::std::min(window, (group_end - 1) - idx); - } -}; - -static_assert(sizeof(fixed_window_wrapper) == sizeof(variable_window_wrapper)); -static_assert(alignof(fixed_window_wrapper) == alignof(variable_window_wrapper)); - -static_assert(sizeof(variable_window_wrapper) == sizeof(fixed_window_wrapper)); -static_assert(alignof(variable_window_wrapper) == alignof(fixed_window_wrapper)); - -static_assert(sizeof(fixed_window_wrapper) == sizeof(preceding_window_wrapper)); -static_assert(alignof(fixed_window_wrapper) == alignof(preceding_window_wrapper)); - -static_assert(sizeof(fixed_window_wrapper) == sizeof(following_window_wrapper)); -static_assert(alignof(fixed_window_wrapper) == alignof(following_window_wrapper)); - -} // namespace detail -} // namespace cudf diff --git a/cpp/src/rolling/detail/rolling_udf.cuh b/cpp/src/rolling/detail/rolling_udf.cuh deleted file mode 100644 index 66b1bcc09ace..000000000000 --- a/cpp/src/rolling/detail/rolling_udf.cuh +++ /dev/null @@ -1,145 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. - * SPDX-License-Identifier: Apache-2.0 - */ - -#pragma once - -#include "jit/cache.hpp" -#include "jit/helpers.hpp" -#include "jit/parser.hpp" -#include "jit/util.hpp" -#include "rolling.hpp" -#include "rolling_jit.cuh" - -#include -#include -#include -#include -#include -#include - -#include - -#include - -namespace cudf { -namespace detail { - -template -std::string reflect_window_wrapper() -{ - if constexpr (std::is_same_v) { - return "cudf::detail::fixed_window_wrapper"; - } else if constexpr (std::is_same_v) { - return "cudf::detail::variable_window_wrapper"; - } else if constexpr (std::is_same_v) { - return "cudf::detail::preceding_window_wrapper"; - } else { - static_assert(std::is_same_v, "Unsupported window wrapper type"); - return "cudf::detail::following_window_wrapper"; - } -} - -inline std::unique_ptr rolling_window_udf_impl( - column_view const& input, - std::string const& preceding_window_str, - cudf::detail::window_wrapper_base const& preceding_window, - std::string const& following_window_str, - cudf::detail::window_wrapper_base const& following_window, - size_type min_periods, - rolling_aggregation const& agg, - cuda::stream_ref stream, - rmm::device_async_resource_ref mr) -{ - static_assert(warp_size == cudf::detail::size_in_bits(), - "bitmask_type size does not match CUDA warp size"); - - if (input.has_nulls()) { - CUDF_FAIL("Currently the UDF version of rolling window does NOT support inputs with nulls."); - } - - min_periods = std::max(min_periods, 0); - - auto& udf_agg = dynamic_cast(agg); - - std::string hash = "prog_rolling." + std::to_string(std::hash{}(udf_agg._source)); - - std::string cuda_source; - switch (udf_agg.kind) { - case aggregation::Kind::PTX: - cuda_source += - cudf::jit::parse_single_function_ptx(udf_agg._source, - udf_agg._function_name, - {{0, cudf::type_to_name(udf_agg._output_type) + " *"}, - {5, "void const *"}}); // args 0 and 5 are pointers - break; - case aggregation::Kind::CUDA: - cuda_source += cudf::jit::parse_single_function_cuda(udf_agg._source, udf_agg._function_name); - break; - default: CUDF_FAIL("Unsupported UDF type."); - } - - std::unique_ptr output = make_numeric_column( - udf_agg._output_type, input.size(), cudf::mask_state::UNINITIALIZED, stream, mr); - - auto output_view = output->mutable_view(); - cudf::detail::device_scalar device_valid_count{ - 0, stream, cudf::get_current_device_resource_ref()}; - - std::string kernel_reflection = - rtcx::reflect_template("cudf::rolling::jit::rolling_window_kernel", - cudf::type_to_name(input.type()), // list of template arguments - cudf::type_to_name(output->type()), - udf_agg._operator_name, - preceding_window_str, - following_window_str); - - auto kernel = - cudf::jit::get_udf_kernel("cudf/cpp/src/rolling/jit/kernel.cu", kernel_reflection, cuda_source); - auto cfg = kernel.max_occupancy_config(0, 0); - kernel.launch_with({cfg.min_grid_size}, - {cfg.block_size}, - 0, - stream, - input.size(), - cudf::jit::get_data_ptr(input), - input.null_mask(), - cudf::jit::get_data_ptr(output->mutable_view()), - output_view.null_mask(), - device_valid_count.data(), - preceding_window, - following_window, - min_periods); - - output->set_null_count(output->size() - device_valid_count.value(stream)); - - // check the stream for debugging - CUDF_CHECK_CUDA(stream.get()); - - return output; -} - -// Applies a user-defined rolling window function to the values in a column. -template -std::unique_ptr rolling_window_udf(column_view const& input, - PrecedingWindowIterator preceding_window, - FollowingWindowIterator following_window, - size_type min_periods, - rolling_aggregation const& agg, - cuda::stream_ref stream, - rmm::device_async_resource_ref mr) -{ - return rolling_window_udf_impl(input, - reflect_window_wrapper(), - preceding_window, - reflect_window_wrapper(), - following_window, - min_periods, - agg, - stream, - mr); -} - -} // namespace detail -} // namespace cudf diff --git a/cpp/src/rolling/detail/rolling_variable_window.cu b/cpp/src/rolling/detail/rolling_variable_window.cu index 3766b52cc8b7..3ff19ed899a4 100644 --- a/cpp/src/rolling/detail/rolling_variable_window.cu +++ b/cpp/src/rolling/detail/rolling_variable_window.cu @@ -4,7 +4,6 @@ */ #include "rolling.cuh" -#include "rolling_udf.cuh" #include #include @@ -35,28 +34,16 @@ std::unique_ptr rolling_window(column_view const& input, CUDF_EXPECTS(preceding_window.size() == input.size() && following_window.size() == input.size(), "preceding_window/following_window size must match input size"); - if (agg.kind == aggregation::CUDA || agg.kind == aggregation::PTX) { - // TODO: In future, might need to clamp preceding/following to column boundaries. - return cudf::detail::rolling_window_udf( - input, - cudf::detail::variable_window_wrapper{preceding_window.begin()}, - cudf::detail::variable_window_wrapper{following_window.begin()}, - min_periods, - agg, - stream, - mr); - } else { - auto defaults_col = - cudf::is_dictionary(input.type()) ? dictionary_column_view(input).indices() : input; - return cudf::detail::rolling_window(input, - empty_like(defaults_col)->view(), - preceding_window.begin(), - following_window.begin(), - min_periods, - agg, - stream, - mr); - } + auto defaults_col = + cudf::is_dictionary(input.type()) ? dictionary_column_view(input).indices() : input; + return cudf::detail::rolling_window(input, + empty_like(defaults_col)->view(), + preceding_window.begin(), + following_window.begin(), + min_periods, + agg, + stream, + mr); } } // namespace cudf::detail diff --git a/cpp/src/rolling/grouped_rolling.cu b/cpp/src/rolling/grouped_rolling.cu index 30b3cd44ccfb..942304e39bcb 100644 --- a/cpp/src/rolling/grouped_rolling.cu +++ b/cpp/src/rolling/grouped_rolling.cu @@ -6,7 +6,6 @@ #include "detail/optimized_unbounded_window.hpp" #include "detail/range_window_bounds.hpp" #include "detail/rolling.cuh" -#include "detail/rolling_udf.cuh" #include "detail/rolling_utils.cuh" #include @@ -94,33 +93,14 @@ std::unique_ptr grouped_rolling_window(table_view const& group_keys, // groups.) // 3. [0, 500, 1000] indicates two equal-sized groups: [0,500), and [500,1000). - if (aggr.kind == aggregation::CUDA || aggr.kind == aggregation::PTX) { - cudf::detail::preceding_window_wrapper grouped_preceding_window{ - group_offsets.data(), group_labels.data(), preceding_window}; - - cudf::detail::following_window_wrapper grouped_following_window{ - group_offsets.data(), group_labels.data(), following_window}; - - return cudf::detail::rolling_window_udf( - input, - cudf::detail::preceding_window_wrapper{ - group_offsets.data(), group_labels.data(), preceding_window}, - cudf::detail::following_window_wrapper{ - group_offsets.data(), group_labels.data(), following_window}, - min_periods, - aggr, - stream, - mr); - } else { - namespace utils = cudf::detail::rolling; - auto groups = utils::grouped{group_labels.data(), group_offsets.data()}; - auto preceding = - utils::make_clamped_window_iterator(preceding_window, groups); - auto following = - utils::make_clamped_window_iterator(following_window, groups); - return cudf::detail::rolling_window( - input, default_outputs, preceding, following, min_periods, aggr, stream, mr); - } + namespace utils = cudf::detail::rolling; + auto groups = utils::grouped{group_labels.data(), group_offsets.data()}; + auto preceding = + utils::make_clamped_window_iterator(preceding_window, groups); + auto following = + utils::make_clamped_window_iterator(following_window, groups); + return cudf::detail::rolling_window( + input, default_outputs, preceding, following, min_periods, aggr, stream, mr); } } // namespace detail diff --git a/cpp/src/rolling/jit/kernel.cu b/cpp/src/rolling/jit/kernel.cu deleted file mode 100644 index e255d7b3a967..000000000000 --- a/cpp/src/rolling/jit/kernel.cu +++ /dev/null @@ -1,136 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. - * SPDX-License-Identifier: Apache-2.0 - */ - -#include -#include -#include - -#include - -#pragma nv_hdrstop // The above headers are used by the kernel below and need to be included before - // it. Each UDF will have a different operation_udf.cuh generated for it, so we - // need to put this pragma before including it to avoid PCH mismatch. - -#include -#include - -struct rolling_udf_ptx { - template - static OutType operate(InType const* in_col, cudf::size_type start, cudf::size_type count) - { - OutType ret; - GENERIC_ROLLING_OP(&ret, 0, 0, 0, 0, &in_col[start], count, sizeof(InType)); - return ret; - } -}; - -struct rolling_udf_cuda { - template - static OutType operate(InType const* in_col, cudf::size_type start, cudf::size_type count) - { - OutType ret; - GENERIC_ROLLING_OP(&ret, in_col, start, count); - return ret; - } -}; - -namespace cudf { -namespace rolling { -namespace jit { - -template -__device__ void rolling_window_kernel(cudf::size_type nrows, - void const* __restrict__ p_in_col, - cudf::bitmask_type const* __restrict__ const in_col_valid, - void* __restrict__ p_out_col, - cudf::bitmask_type* __restrict__ out_col_valid, - cudf::size_type* __restrict__ output_valid_count, - detail::window_wrapper_base b_preceding_window_begin, - detail::window_wrapper_base b_following_window_begin, - cudf::size_type min_periods) -{ - auto i = cudf::detail::grid_1d::global_thread_id(); - auto const stride = cudf::detail::grid_1d::grid_stride(); - PrecedingWindowType const preceding_window_begin = b_preceding_window_begin; - FollowingWindowType const following_window_begin = b_following_window_begin; - auto const* const in_col = static_cast(p_in_col); - auto* const out_col = static_cast(p_out_col); - - cudf::size_type warp_valid_count{0}; - - auto active_threads = __ballot_sync(0xffff'ffffu, i < nrows); - while (i < nrows) { - int64_t const preceding_window = preceding_window_begin[i]; - int64_t const following_window = following_window_begin[i]; - - // compute bounds - auto const start = static_cast( - min(static_cast(nrows), max(int64_t{0}, i - preceding_window + 1))); - auto const end = static_cast( - min(static_cast(nrows), max(int64_t{0}, i + following_window + 1))); - auto const start_index = min(start, end); - auto const end_index = max(start, end); - - // aggregate - // TODO: We should explore using shared memory to avoid redundant loads. - // This might require separating the kernel into a special version - // for dynamic and static sizes. - cudf::size_type count = end_index - start_index; - OutType val = agg_op::template operate(in_col, start_index, count); - - // check if we have enough input samples - bool const output_is_valid = (count >= min_periods); - - // set the mask - unsigned int const result_mask = __ballot_sync(active_threads, output_is_valid); - - // store the output value, one per thread - if (output_is_valid) { out_col[i] = val; } - - // only one thread writes the mask - if (0 == cudf::intra_word_index(i)) { - out_col_valid[cudf::word_index(i)] = result_mask; - warp_valid_count += __popc(result_mask); - } - - // process next element - i += stride; - active_threads = __ballot_sync(active_threads, i < nrows); - } - - // TODO: likely faster to do a single_lane_block_reduce and a single - // atomic per block but that requires jit-compiling single_lane_block_reduce... - if (0 == cudf::intra_word_index(threadIdx.x)) { atomicAdd(output_valid_count, warp_valid_count); } -} - -} // namespace jit -} // namespace rolling -} // namespace cudf - -extern "C" __global__ void cudf_kernel_entry( - cudf::size_type nrows, - void const* __restrict__ in_col, - cudf::bitmask_type const* __restrict__ in_col_valid, - void* __restrict__ out_col, - cudf::bitmask_type* __restrict__ out_col_valid, - cudf::size_type* __restrict__ output_valid_count, - cudf::detail::window_wrapper_base preceding_window_begin, - cudf::detail::window_wrapper_base following_window_begin, - cudf::size_type min_periods) -{ - CUDF_KERNEL_INSTANCE(nrows, - in_col, - in_col_valid, - out_col, - out_col_valid, - output_valid_count, - preceding_window_begin, - following_window_begin, - min_periods); -} diff --git a/cpp/tests/rolling/empty_input_test.cpp b/cpp/tests/rolling/empty_input_test.cpp index b6a73d0803c3..73c32c9d3504 100644 --- a/cpp/tests/rolling/empty_input_test.cpp +++ b/cpp/tests/rolling/empty_input_test.cpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -39,12 +39,6 @@ auto row_number() { return cudf::make_row_number_aggregation(); } -auto udf() -{ - return cudf::make_udf_aggregation( - cudf::udf_source_type::CUDA, "", cudf::data_type{cudf::type_id::INT32}); -} - // Constants for rolling_window. auto constexpr min_periods = 1; auto constexpr preceding = 2; @@ -160,7 +154,6 @@ TYPED_TEST(TypedRollingEmptyInputTest, EmptyFixedWidthInputs) aggs.emplace_back(max()); aggs.emplace_back(lead()); aggs.emplace_back(lag()); - aggs.emplace_back(udf()); rolling_output_type_matches(empty_input, aggs, cudf::type_to_id()); } @@ -233,7 +226,6 @@ TEST_F(RollingEmptyInputTest, Strings) aggs.emplace_back(max()); aggs.emplace_back(lead()); aggs.emplace_back(lag()); - aggs.emplace_back(udf()); rolling_output_type_matches(empty_input, aggs, cudf::type_id::STRING); } @@ -279,7 +271,6 @@ TEST_F(RollingEmptyInputTest, Dictionaries) aggs.emplace_back(max()); aggs.emplace_back(lead()); aggs.emplace_back(lag()); - aggs.emplace_back(udf()); rolling_output_type_matches(empty_input, aggs, cudf::type_id::DICTIONARY32); } @@ -327,7 +318,6 @@ TYPED_TEST(TypedRollingEmptyInputTest, Lists) aggs.emplace_back(max()); aggs.emplace_back(lead()); aggs.emplace_back(lag()); - aggs.emplace_back(udf()); rolling_output_type_matches(empty_input, aggs, cudf::type_id::LIST, cudf::type_to_id()); } @@ -375,7 +365,6 @@ TYPED_TEST(TypedRollingEmptyInputTest, Structs) aggs.emplace_back(max()); aggs.emplace_back(lead()); aggs.emplace_back(lag()); - aggs.emplace_back(udf()); rolling_output_type_matches(empty_input, aggs, cudf::type_id::STRUCT, cudf::type_to_id()); } diff --git a/cpp/tests/rolling/grouped_rolling_test.cpp b/cpp/tests/rolling/grouped_rolling_test.cpp index b568a909ff0e..8ee6391ec7ec 100644 --- a/cpp/tests/rolling/grouped_rolling_test.cpp +++ b/cpp/tests/rolling/grouped_rolling_test.cpp @@ -22,85 +22,6 @@ #include -std::string const cuda_func{ - R"***( - template - __device__ void CUDA_GENERIC_AGGREGATOR(OutType *ret, InType *in_col, cudf::size_type start, - cudf::size_type count) { - OutType val = 0; - for (cudf::size_type i = 0; i < count; i++) { - val += in_col[start + i]; - } - *ret = val; - } - )***"}; - -std::string const ptx_func{ - R"***( - // - // Generated by NVIDIA NVVM Compiler - // - // Compiler Build ID: CL-24817639 - // Cuda compilation tools, release 10.0, V10.0.130 - // Based on LLVM 3.4svn - // - - .version 6.3 - .target sm_70 - .address_size 64 - - // .globl _ZN8__main__7add$241E5ArrayIiLi1E1A7mutable7alignedE - .common .global .align 8 .u64 _ZN08NumbaEnv8__main__7add$241E5ArrayIiLi1E1A7mutable7alignedE; - - .visible .func (.param .b32 func_retval0) _ZN8__main__7add$241E5ArrayIiLi1E1A7mutable7alignedE( - .param .b64 _ZN8__main__7add$241E5ArrayIiLi1E1A7mutable7alignedE_paam_0, - .param .b64 _ZN8__main__7add$241E5ArrayIiLi1E1A7mutable7alignedE_paam_1, - .param .b64 _ZN8__main__7add$241E5ArrayIiLi1E1A7mutable7alignedE_paam_2, - .param .b64 _ZN8__main__7add$241E5ArrayIiLi1E1A7mutable7alignedE_paam_3, - .param .b64 _ZN8__main__7add$241E5ArrayIiLi1E1A7mutable7alignedE_paam_4, - .param .b64 _ZN8__main__7add$241E5ArrayIiLi1E1A7mutable7alignedE_paam_5, - .param .b64 _ZN8__main__7add$241E5ArrayIiLi1E1A7mutable7alignedE_paam_6, - .param .b64 _ZN8__main__7add$241E5ArrayIiLi1E1A7mutable7alignedE_paam_7 - ) - { - .reg .pred %p<3>; - .reg .b32 %r<6>; - .reg .b64 %rd<18>; - - - ld.param.u64 %rd6, [_ZN8__main__7add$241E5ArrayIiLi1E1A7mutable7alignedE_paam_0]; - ld.param.u64 %rd7, [_ZN8__main__7add$241E5ArrayIiLi1E1A7mutable7alignedE_paam_5]; - ld.param.u64 %rd8, [_ZN8__main__7add$241E5ArrayIiLi1E1A7mutable7alignedE_paam_6]; - ld.param.u64 %rd9, [_ZN8__main__7add$241E5ArrayIiLi1E1A7mutable7alignedE_paam_7]; - mov.u64 %rd15, 0; - mov.u64 %rd16, %rd15; - - BB0_1: - mov.u64 %rd2, %rd16; - mov.u32 %r5, 0; - setp.ge.s64 %p1, %rd15, %rd8; - mov.u64 %rd17, %rd15; - @%p1 bra BB0_3; - - mul.lo.s64 %rd12, %rd15, %rd9; - add.s64 %rd13, %rd12, %rd7; - ld.u32 %r5, [%rd13]; - add.s64 %rd17, %rd15, 1; - - BB0_3: - cvt.s64.s32 %rd14, %r5; - add.s64 %rd16, %rd14, %rd2; - setp.lt.s64 %p2, %rd15, %rd8; - mov.u64 %rd15, %rd17; - @%p2 bra BB0_1; - - st.u64 [%rd6], %rd2; - mov.u32 %r4, 0; - st.param.b32 [func_retval0+0], %r4; - ret; - } - )***"}; - template class GroupedRollingTest : public cudf::test::BaseFixture { protected: @@ -189,29 +110,6 @@ class GroupedRollingTest : public cudf::test::BaseFixture { following_window, min_periods, *cudf::make_row_number_aggregation()); - - // >>> test UDFs <<< - if (input.type() == cudf::data_type{cudf::type_id::INT32} && !input.has_nulls()) { - auto cuda_udf_agg = cudf::make_udf_aggregation( - cudf::udf_source_type::CUDA, cuda_func, cudf::data_type{cudf::type_id::INT64}); - run_test_col(keys, - input, - expected_grouping, - preceding_window, - following_window, - min_periods, - *cuda_udf_agg); - - auto ptx_udf_agg = cudf::make_udf_aggregation( - cudf::udf_source_type::PTX, ptx_func, cudf::data_type{cudf::type_id::INT64}); - run_test_col(keys, - input, - expected_grouping, - preceding_window, - following_window, - min_periods, - *ptx_udf_agg); - } } private: @@ -407,19 +305,6 @@ class GroupedRollingTest : public cudf::test::BaseFixture { std::conditional_t(), T, double>, true>( input, group_offsets, preceding_window, following_window, min_periods); - // >>> UDFs <<< - case cudf::aggregation::CUDA: - return create_reference_output, int64_t, T>, - false>( - input, group_offsets, preceding_window, following_window, min_periods); - case cudf::aggregation::PTX: - return create_reference_output, int64_t, T>, - false>( - input, group_offsets, preceding_window, following_window, min_periods); default: return cudf::test::fixed_width_column_wrapper({}).release(); } } diff --git a/cpp/tests/rolling/rolling_test.cpp b/cpp/tests/rolling/rolling_test.cpp index a2f178499388..df7dd539aaf5 100644 --- a/cpp/tests/rolling/rolling_test.cpp +++ b/cpp/tests/rolling/rolling_test.cpp @@ -1204,22 +1204,6 @@ TEST_F(RollingTestStrings, StringsUnsupportedOperators) EXPECT_THROW( cudf::rolling_window(input, 2, 2, 0, *cudf::make_mean_aggregation()), cudf::logic_error); - EXPECT_THROW( - cudf::rolling_window(input, - 2, - 2, - 0, - *cudf::make_udf_aggregation( - cudf::udf_source_type::PTX, std::string{}, cudf::data_type{})), - cudf::logic_error); - EXPECT_THROW( - cudf::rolling_window(input, - 2, - 2, - 0, - *cudf::make_udf_aggregation( - cudf::udf_source_type::CUDA, std::string{}, cudf::data_type{})), - cudf::logic_error); } /*TEST_F(RollingTestStrings, SimpleStatic) @@ -1235,172 +1219,6 @@ TEST_F(RollingTestStrings, StringsUnsupportedOperators) EXPECT_NO_THROW(this->run_test_col(input, window, window, 0, rolling_operator::COUNT_ALL)); }*/ -struct RollingTestUdf : public cudf::test::BaseFixture { - std::string const cuda_func{ - R"***( - template - __device__ void CUDA_GENERIC_AGGREGATOR(OutType *ret, InType *in_col, cudf::size_type start, - cudf::size_type count) { - OutType val = 0; - for (cudf::size_type i = 0; i < count; i++) { - val += in_col[start + i]; - } - *ret = val; - } - )***"}; - - std::string const ptx_func{ - R"***( - // - // Generated by NVIDIA NVVM Compiler - // - // Compiler Build ID: CL-24817639 - // Cuda compilation tools, release 10.0, V10.0.130 - // Based on LLVM 3.4svn - // - - .version 6.3 - .target sm_70 - .address_size 64 - - // .globl _ZN8__main__7add$241E5ArrayIiLi1E1A7mutable7alignedE - .common .global .align 8 .u64 _ZN08NumbaEnv8__main__7add$241E5ArrayIiLi1E1A7mutable7alignedE; - - .visible .func (.param .b32 func_retval0) _ZN8__main__7add$241E5ArrayIiLi1E1A7mutable7alignedE( - .param .b64 _ZN8__main__7add$241E5ArrayIiLi1E1A7mutable7alignedE_paam_0, - .param .b64 _ZN8__main__7add$241E5ArrayIiLi1E1A7mutable7alignedE_paam_1, - .param .b64 _ZN8__main__7add$241E5ArrayIiLi1E1A7mutable7alignedE_paam_2, - .param .b64 _ZN8__main__7add$241E5ArrayIiLi1E1A7mutable7alignedE_paam_3, - .param .b64 _ZN8__main__7add$241E5ArrayIiLi1E1A7mutable7alignedE_paam_4, - .param .b64 _ZN8__main__7add$241E5ArrayIiLi1E1A7mutable7alignedE_paam_5, - .param .b64 _ZN8__main__7add$241E5ArrayIiLi1E1A7mutable7alignedE_paam_6, - .param .b64 _ZN8__main__7add$241E5ArrayIiLi1E1A7mutable7alignedE_paam_7 - ) - { - .reg .pred %p<3>; - .reg .b32 %r<6>; - .reg .b64 %rd<18>; - - - ld.param.u64 %rd6, [_ZN8__main__7add$241E5ArrayIiLi1E1A7mutable7alignedE_paam_0]; - ld.param.u64 %rd7, [_ZN8__main__7add$241E5ArrayIiLi1E1A7mutable7alignedE_paam_5]; - ld.param.u64 %rd8, [_ZN8__main__7add$241E5ArrayIiLi1E1A7mutable7alignedE_paam_6]; - ld.param.u64 %rd9, [_ZN8__main__7add$241E5ArrayIiLi1E1A7mutable7alignedE_paam_7]; - mov.u64 %rd15, 0; - mov.u64 %rd16, %rd15; - - BB0_1: - mov.u64 %rd2, %rd16; - mov.u32 %r5, 0; - setp.ge.s64 %p1, %rd15, %rd8; - mov.u64 %rd17, %rd15; - @%p1 bra BB0_3; - - mul.lo.s64 %rd12, %rd15, %rd9; - add.s64 %rd13, %rd12, %rd7; - ld.u32 %r5, [%rd13]; - add.s64 %rd17, %rd15, 1; - - BB0_3: - cvt.s64.s32 %rd14, %r5; - add.s64 %rd16, %rd14, %rd2; - setp.lt.s64 %p2, %rd15, %rd8; - mov.u64 %rd15, %rd17; - @%p2 bra BB0_1; - - st.u64 [%rd6], %rd2; - mov.u32 %r4, 0; - st.param.b32 [func_retval0+0], %r4; - ret; - } - )***"}; -}; - -TEST_F(RollingTestUdf, StaticWindow) -{ - cudf::size_type size = 1000; - - cudf::test::fixed_width_column_wrapper input(cuda::counting_iterator{0}, - cuda::counting_iterator{size}, - cuda::make_constant_iterator(true)); - - std::unique_ptr output; - - auto start = cudf::detail::make_counting_transform_iterator(0, [size](cudf::size_type row) { - return std::accumulate(cuda::counting_iterator{std::max(0, row - 2 + 1)}, - cuda::counting_iterator{std::min(size, row + 2 + 1)}, - 0); - }); - - auto valid = cudf::detail::make_counting_transform_iterator( - 0, [size](cudf::size_type row) { return (row != 0 && row != size - 2 && row != size - 1); }); - - cudf::test::fixed_width_column_wrapper expected{start, start + size, valid}; - - // Test CUDA UDF - auto cuda_udf_agg = cudf::make_udf_aggregation( - cudf::udf_source_type::CUDA, this->cuda_func, cudf::data_type{cudf::type_id::INT64}); - - output = cudf::rolling_window(input, 2, 2, 4, *cuda_udf_agg); - - CUDF_TEST_EXPECT_COLUMNS_EQUAL(*output, expected); - - // Test NUMBA UDF - auto ptx_udf_agg = cudf::make_udf_aggregation( - cudf::udf_source_type::PTX, this->ptx_func, cudf::data_type{cudf::type_id::INT64}); - - output = cudf::rolling_window(input, 2, 2, 4, *ptx_udf_agg); - - CUDF_TEST_EXPECT_COLUMNS_EQUAL(*output, expected); -} - -TEST_F(RollingTestUdf, DynamicWindow) -{ - cudf::size_type size = 1000; - - cudf::test::fixed_width_column_wrapper input(cuda::counting_iterator{0}, - cuda::counting_iterator{size}, - cuda::make_constant_iterator(true)); - - auto prec = cudf::detail::make_counting_transform_iterator( - 0, [] __device__(cudf::size_type row) { return row % 2 + 2; }); - - auto follow = cudf::detail::make_counting_transform_iterator( - 0, [] __device__(cudf::size_type row) { return row % 2; }); - - cudf::test::fixed_width_column_wrapper preceding(prec, prec + size); - cudf::test::fixed_width_column_wrapper following(follow, follow + size); - std::unique_ptr output; - - auto start = - cudf::detail::make_counting_transform_iterator(0, [size] __device__(cudf::size_type row) { - return std::accumulate(cuda::counting_iterator{std::max(0, row - (row % 2 + 2) + 1)}, - cuda::counting_iterator{std::min(size, row + (row % 2) + 1)}, - 0); - }); - - auto valid = cudf::detail::make_counting_transform_iterator( - 0, [] __device__(cudf::size_type row) { return row != 0; }); - - cudf::test::fixed_width_column_wrapper expected{start, start + size, valid}; - - // Test CUDA UDF - auto cuda_udf_agg = cudf::make_udf_aggregation( - cudf::udf_source_type::CUDA, this->cuda_func, cudf::data_type{cudf::type_id::INT64}); - - output = cudf::rolling_window(input, preceding, following, 2, *cuda_udf_agg); - - CUDF_TEST_EXPECT_COLUMNS_EQUAL(*output, expected); - - // Test PTX UDF - auto ptx_udf_agg = cudf::make_udf_aggregation( - cudf::udf_source_type::PTX, this->ptx_func, cudf::data_type{cudf::type_id::INT64}); - - output = cudf::rolling_window(input, preceding, following, 2, *ptx_udf_agg); - - CUDF_TEST_EXPECT_COLUMNS_EQUAL(*output, expected); -} - template struct FixedPointTests : public cudf::test::BaseFixture {}; diff --git a/python/pylibcudf/pylibcudf/aggregation.pxd b/python/pylibcudf/pylibcudf/aggregation.pxd index 66909b62e322..c8d43c2895a2 100644 --- a/python/pylibcudf/pylibcudf/aggregation.pxd +++ b/python/pylibcudf/pylibcudf/aggregation.pxd @@ -1,4 +1,4 @@ -# SPDX-FileCopyrightText: Copyright (c) 2024-2025, NVIDIA CORPORATION. +# SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. # SPDX-License-Identifier: Apache-2.0 from libcpp cimport bool @@ -93,7 +93,6 @@ cpdef Aggregation collect_list(null_policy null_handling = *) cpdef Aggregation collect_set(null_handling = *, nulls_equal = *, nans_equal = *) -cpdef Aggregation udf(str operation, DataType output_type) cpdef Aggregation correlation(correlation_type type, size_type min_periods) diff --git a/python/pylibcudf/pylibcudf/aggregation.pyi b/python/pylibcudf/pylibcudf/aggregation.pyi index 6735478e52f6..d953500f89d3 100644 --- a/python/pylibcudf/pylibcudf/aggregation.pyi +++ b/python/pylibcudf/pylibcudf/aggregation.pyi @@ -1,4 +1,4 @@ -# SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION. +# SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. # SPDX-License-Identifier: Apache-2.0 from enum import IntEnum @@ -99,7 +99,6 @@ def collect_set( nulls_equal: NullEquality = NullEquality.EQUAL, nans_equal: NanEquality = NanEquality.ALL_EQUAL, ) -> Aggregation: ... -def udf(operation: str, output_type: DataType) -> Aggregation: ... def correlation(type: CorrelationType, min_periods: int) -> Aggregation: ... def covariance(min_periods: int, ddof: int) -> Aggregation: ... def rank( diff --git a/python/pylibcudf/pylibcudf/aggregation.pyx b/python/pylibcudf/pylibcudf/aggregation.pyx index b7e6b5559354..4b761bc77590 100644 --- a/python/pylibcudf/pylibcudf/aggregation.pyx +++ b/python/pylibcudf/pylibcudf/aggregation.pyx @@ -48,7 +48,6 @@ from pylibcudf.libcudf.aggregation cimport ( make_sum_aggregation, make_sum_of_squares_aggregation, make_tdigest_aggregation, - make_udf_aggregation, make_variance_aggregation, rank_method, rank_percentage, @@ -77,8 +76,6 @@ from pylibcudf.libcudf.aggregation import \ rank_method as RankMethod # no-cython-lint from pylibcudf.libcudf.aggregation import \ rank_percentage as RankPercentage # no-cython-lint -from pylibcudf.libcudf.types import \ - udf_source_type as UdfSourceType # no-cython-lint from .types cimport DataType @@ -91,7 +88,6 @@ __all__ = [ "Kind", "RankMethod", "RankPercentage", - "UdfSourceType", "all", "any", "argmax", @@ -127,7 +123,6 @@ __all__ = [ "sum", "sum_of_squares", "tdigest", - "udf", "variance", ] @@ -570,33 +565,6 @@ cpdef Aggregation collect_set( ) ) -cpdef Aggregation udf(str operation, DataType output_type): - """Create a udf aggregation. - - For details, see :cpp:func:`make_udf_aggregation`. - - Parameters - ---------- - operation : str - The operation to perform as a string of PTX code. - output_type : DataType - The output type of the aggregation. - - Returns - ------- - Aggregation - The udf aggregation. - """ - return Aggregation.from_libcudf( - move( - make_udf_aggregation[aggregation]( - UdfSourceType.PTX, - operation.encode("utf-8"), - output_type.c_obj, - ) - ) - ) - cpdef Aggregation correlation(correlation_type type, size_type min_periods): """Create a correlation aggregation. @@ -927,4 +895,3 @@ CorrelationType.__str__ = CorrelationType.__repr__ EWMHistory.__str__ = EWMHistory.__repr__ RankMethod.__str__ = RankMethod.__repr__ RankPercentage.__str__ = RankPercentage.__repr__ -UdfSourceType.__str__ = UdfSourceType.__repr__ diff --git a/python/pylibcudf/pylibcudf/libcudf/aggregation.pxd b/python/pylibcudf/pylibcudf/libcudf/aggregation.pxd index 01f7d0995d1c..295f21a5c824 100644 --- a/python/pylibcudf/pylibcudf/libcudf/aggregation.pxd +++ b/python/pylibcudf/pylibcudf/libcudf/aggregation.pxd @@ -1,10 +1,9 @@ -# SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. +# SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. # SPDX-License-Identifier: Apache-2.0 from libc.stddef cimport size_t from libc.stdint cimport int32_t from libcpp cimport bool from libcpp.memory cimport unique_ptr -from libcpp.string cimport string from libcpp.vector cimport vector from pylibcudf.exception_handler cimport libcudf_exception_handler from pylibcudf.libcudf.types cimport ( @@ -16,7 +15,6 @@ from pylibcudf.libcudf.types cimport ( null_policy, order, size_type, - udf_source_type, ) @@ -51,8 +49,6 @@ cdef extern from "cudf/aggregation.hpp" namespace "cudf" nogil: COLLECT_SET LEAD LAG - PTX - CUDA HOST_UDF MERGE_LISTS MERGE_SETS @@ -166,11 +162,6 @@ cdef extern from "cudf/aggregation.hpp" namespace "cudf" nogil: null_policy null_handling, null_equality nulls_equal, nan_equality nans_equal ) except +libcudf_exception_handler - cdef unique_ptr[T] make_udf_aggregation[T]( - udf_source_type type, - string user_defined_aggregator, - data_type output_type) except +libcudf_exception_handler - cdef unique_ptr[T] make_ewma_aggregation[T]( double com, ewm_history adjust ) except +libcudf_exception_handler