diff --git a/cpp/benchmarks/type_dispatcher/type_dispatcher_benchmark.cu b/cpp/benchmarks/type_dispatcher/type_dispatcher_benchmark.cu index df3a373c576..18ef5a1168e 100644 --- a/cpp/benchmarks/type_dispatcher/type_dispatcher_benchmark.cu +++ b/cpp/benchmarks/type_dispatcher/type_dispatcher_benchmark.cu @@ -27,6 +27,7 @@ #include +#include #include #include #include "../fixture/benchmark_fixture.hpp" @@ -87,7 +88,7 @@ __global__ void host_dispatching_kernel(mutable_column_device_view source_column template struct ColumnHandle { - template + template ())> void operator()(mutable_column_device_view source_column, int work_per_thread) { cudf::detail::grid_1d grid_config{source_column.size(), block_size}; @@ -95,6 +96,12 @@ struct ColumnHandle { // Launch the kernel. host_dispatching_kernel<<>>(source_column); } + + template ())> + void operator()(mutable_column_device_view source_column, int work_per_thread) + { + CUDF_FAIL("Invalid type to benchmark."); + } }; // The following is for DEVICE_DISPATCHING: @@ -104,12 +111,18 @@ struct ColumnHandle { // n_rows * n_cols. template struct RowHandle { - template + template ())> __device__ void operator()(mutable_column_device_view source, cudf::size_type index) { using F = Functor; source.data()[index] = F::f(source.data()[index]); } + + template ())> + __device__ void operator()(mutable_column_device_view source, cudf::size_type index) + { + cudf_assert(false && "Unsupported type."); + } }; // This is for DEVICE_DISPATCHING diff --git a/cpp/include/cudf/ast/detail/transform.cuh b/cpp/include/cudf/ast/detail/transform.cuh index 2719a8b5077..da15ac07c63 100644 --- a/cpp/include/cudf/ast/detail/transform.cuh +++ b/cpp/include/cudf/ast/detail/transform.cuh @@ -20,10 +20,12 @@ #include #include #include +#include #include #include #include #include +#include #include @@ -55,10 +57,19 @@ struct row_output { * @param row_index Row index of data column. * @param result Value to assign to output. */ - template + template ())> __device__ void resolve_output(detail::device_data_reference device_data_reference, cudf::size_type row_index, Element result) const; + // Definition below after row_evaluator is a complete type + + template ())> + __device__ void resolve_output(detail::device_data_reference device_data_reference, + cudf::size_type row_index, + Element result) const + { + cudf_assert(false && "Invalid type in resolve_output."); + } private: row_evaluator const& evaluator; @@ -167,7 +178,7 @@ struct row_evaluator { * @param row_index Row index of data column. * @return Element */ - template + template ())> __device__ Element resolve_input(detail::device_data_reference device_data_reference, cudf::size_type row_index) const { @@ -187,6 +198,15 @@ struct row_evaluator { } } + template ())> + __device__ Element resolve_input(detail::device_data_reference device_data_reference, + cudf::size_type row_index) const + { + cudf_assert(false && "Unsupported type in resolve_input."); + return {}; + } + /** * @brief Callable to perform a unary operation. * @@ -249,7 +269,7 @@ struct row_evaluator { mutable_column_device_view* output_column; }; -template +template ()>*> __device__ void row_output::resolve_output(detail::device_data_reference device_data_reference, cudf::size_type row_index, Element result) const diff --git a/cpp/include/cudf/column/column_device_view.cuh b/cpp/include/cudf/column/column_device_view.cuh index 14d44b77fad..a842e51c94a 100644 --- a/cpp/include/cudf/column/column_device_view.cuh +++ b/cpp/include/cudf/column/column_device_view.cuh @@ -71,10 +71,14 @@ class alignas(16) column_device_view_base { * a column, and instead, accessing the elements should be done via *`data()`. * + * This function will only participate in overload resolution if `is_rep_layout_compatible()` + * or `std::is_same::value` are true. + * * @tparam The type to cast to * @return T const* Typed pointer to underlying data */ - template + template ::value or is_rep_layout_compatible())> __host__ __device__ T const* head() const noexcept { return static_cast(_data); @@ -89,10 +93,13 @@ class alignas(16) column_device_view_base { * For columns with children, the pointer returned is undefined * and should not be used. * + * This function does not participate in overload resolution if `is_rep_layout_compatible` is + * false. + * * @tparam T The type to cast to * @return T const* Typed pointer to underlying data, including the offset */ - template + template ())> __host__ __device__ T const* data() const noexcept { return head() + _offset; @@ -235,6 +242,18 @@ class alignas(16) column_device_view_base { : _type{type}, _size{size}, _data{data}, _null_mask{null_mask}, _offset{offset} { } + + template + struct has_element_accessor_impl : std::false_type { + }; + + template + struct has_element_accessor_impl< + C, + T, + void_t().template element(std::declval()))>> + : std::true_type { + }; }; // Forward declaration @@ -283,15 +302,145 @@ class alignas(16) column_device_view : public detail::column_device_view_base { * * This function accounts for the offset. * + * This function does not participate in overload resolution if `is_rep_layout_compatible` is + * false. Specializations of this function may exist for types `T` where + *`is_rep_layout_compatible` is false. + * * @tparam T The element type * @param element_index Position of the desired element */ - template - __device__ T const element(size_type element_index) const noexcept + template ())> + __device__ T element(size_type element_index) const noexcept { return data()[element_index]; } + /** + * @brief Returns `string_view` to the string element at the specified index. + * + * If the element at the specified index is NULL, i.e., `is_null(element_index) + * == true`, then any attempt to use the result will lead to undefined behavior. + * + * This function accounts for the offset. + * + * @param element_index Position of the desired string element + * @return string_view instance representing this element at this index + */ + template ::value)> + __device__ T element(size_type element_index) const noexcept + { + size_type index = element_index + offset(); // account for this view's _offset + const int32_t* d_offsets = + d_children[strings_column_view::offsets_column_index].data(); + const char* d_strings = d_children[strings_column_view::chars_column_index].data(); + size_type offset = d_offsets[index]; + return string_view{d_strings + offset, d_offsets[index + 1] - offset}; + } + + private: + /** + * @brief Dispatch functor for resolving the index value for a dictionary element. + * + * The basic dictionary elements are the indices which can be any index type. + */ + struct index_element_fn { + template () and std::is_unsigned::value)> + __device__ size_type operator()(column_device_view const& indices, size_type index) + { + return static_cast(indices.element(index)); + } + + template () and + std::is_unsigned::value))> + __device__ size_type operator()(Args&&... args) + { + cudf_assert(false and "dictionary indices must be an unsigned integral type"); + return 0; + } + }; + + public: + /** + * @brief Returns `dictionary32` element at the specified index for a + * dictionary column. + * + * `dictionary32` is a strongly typed wrapper around an `int32_t` value that holds the + * offset into the dictionary keys for the specified element. + * + * For example, given a dictionary column `d` with: + * ```c++ + * keys: {"foo", "bar", "baz"} + * indices: {2, 0, 2, 1, 0} + * + * d.element(0) == dictionary32{2}; + * d.element(1) == dictionary32{0}; + * ``` + * + * If the element at the specified index is NULL, i.e., `is_null(element_index) == true`, + * then any attempt to use the result will lead to undefined behavior. + * + * This function accounts for the offset. + * + * @param element_index Position of the desired element + * @return dictionary32 instance representing this element at this index + */ + template ::value)> + __device__ T element(size_type element_index) const noexcept + { + size_type index = element_index + offset(); // account for this view's _offset + auto const indices = d_children[0]; + return dictionary32{type_dispatcher(indices.type(), index_element_fn{}, indices, index)}; + } + + /** + * @brief Returns a `numeric::decimal32` element at the specified index for a `fixed_point` + * column. + * + * If the element at the specified index is NULL, i.e., `is_null(element_index) == true`, + * then any attempt to use the result will lead to undefined behavior. + * + * @param element_index Position of the desired element + * @return numeric::decimal32 representing the element at this index + */ + template ::value)> + __device__ T element(size_type element_index) const noexcept + { + using namespace numeric; + auto const scale = scale_type{_type.scale()}; + return decimal32{scaled_integer{data()[element_index], scale}}; + } + + /** + * @brief Returns a `numeric::decimal64` element at the specified index for a `fixed_point` + * column. + * + * If the element at the specified index is NULL, i.e., `is_null(element_index) == true`, + * then any attempt to use the result will lead to undefined behavior. + * + * @param element_index Position of the desired element + * @return numeric::decimal64 representing the element at this index + */ + template ::value)> + __device__ T element(size_type element_index) const noexcept + { + using namespace numeric; + auto const scale = scale_type{_type.scale()}; + return decimal64{scaled_integer{data()[element_index], scale}}; + } + + /** + * @brief For a given `T`, indicates if `column_device_view::element()` has a valid overload. + * + */ + template + static constexpr bool has_element_accessor() + { + return has_element_accessor_impl::value; + } + /** * @brief Iterator for navigating this column */ @@ -306,9 +455,12 @@ class alignas(16) column_device_view : public detail::column_device_view_base { * with columns where `has_nulls() == true` will result in undefined behavior * when accessing null elements. * + * This function does not participate in overload resolution if + * `column_device_view::has_element_accessor()` is false. + * * For columns with null elements, use `make_null_replacement_iterator`. */ - template + template ())> const_iterator begin() const { return const_iterator{count_it{0}, detail::value_accessor{*this}}; @@ -321,9 +473,12 @@ class alignas(16) column_device_view : public detail::column_device_view_base { * with columns where `has_nulls() == true` will result in undefined behavior * when accessing null elements. * + * This function does not participate in overload resolution if + * `column_device_view::has_element_accessor()` is false. + * * For columns with null elements, use `make_null_replacement_iterator`. */ - template + template ())> const_iterator end() const { return const_iterator{count_it{size()}, detail::value_accessor{*this}}; @@ -357,11 +512,16 @@ class alignas(16) column_device_view : public detail::column_device_view_base { * Else, if the element at `i` is null, then the value of `p.first` is * undefined and `p.second == false`. * + * This function does not participate in overload resolution if + * `column_device_view::has_element_accessor()` is false. + * * @throws cudf::logic_error if tparam `has_nulls == true` and * `nullable() == false` * @throws cudf::logic_error if column datatype and Element type mismatch. */ - template + template ())> const_pair_iterator pair_begin() const { return const_pair_iterator{count_it{0}, @@ -382,11 +542,16 @@ class alignas(16) column_device_view : public detail::column_device_view_base { * Else, if the element at `i` is null, then the value of `p.first` is * undefined and `p.second == false`. * + * This function does not participate in overload resolution if + * `column_device_view::has_element_accessor()` is false. + * * @throws cudf::logic_error if tparam `has_nulls == true` and * `nullable() == false` * @throws cudf::logic_error if column datatype and Element type mismatch. */ - template + template ())> const_pair_rep_iterator pair_rep_begin() const { return const_pair_rep_iterator{count_it{0}, @@ -397,11 +562,16 @@ class alignas(16) column_device_view : public detail::column_device_view_base { * @brief Return a pair iterator to the element following the last element of * the column. * + * This function does not participate in overload resolution if + * `column_device_view::has_element_accessor()` is false. + * * @throws cudf::logic_error if tparam `has_nulls == true` and * `nullable() == false` * @throws cudf::logic_error if column datatype and Element type mismatch. */ - template + template ())> const_pair_iterator pair_end() const { return const_pair_iterator{count_it{size()}, @@ -412,11 +582,16 @@ class alignas(16) column_device_view : public detail::column_device_view_base { * @brief Return a pair iterator to the element following the last element of * the column. * + * This function does not participate in overload resolution if + * `column_device_view::has_element_accessor()` is false. + * * @throws cudf::logic_error if tparam `has_nulls == true` and * `nullable() == false` * @throws cudf::logic_error if column datatype and Element type mismatch. */ - template + template ())> const_pair_rep_iterator pair_rep_end() const { return const_pair_rep_iterator{count_it{size()}, @@ -549,6 +724,9 @@ class alignas(16) mutable_column_device_view : public detail::column_device_view * @brief Returns pointer to the base device memory allocation casted to * the specified type. * + * This function will only participate in overload resolution if `is_rep_layout_compatible()` + * or `std::is_same::value` are true. + * * @note If `offset() == 0`, then `head() == data()` * * @note It should be rare to need to access the `head()` allocation of @@ -558,7 +736,8 @@ class alignas(16) mutable_column_device_view : public detail::column_device_view * @tparam The type to cast to * @return T* Typed pointer to underlying data */ - template + template ::value or is_rep_layout_compatible())> __host__ __device__ T* head() const noexcept { return const_cast(detail::column_device_view_base::head()); @@ -568,14 +747,15 @@ class alignas(16) mutable_column_device_view : public detail::column_device_view * @brief Returns the underlying data casted to the specified type, plus the * offset. * - * @note If `offset() == 0`, then `head() == data()` + * This function does not participate in overload resolution if `is_rep_layout_compatible` is + * false. * - * This pointer is undefined for columns with children. + * @note If `offset() == 0`, then `head() == data()` * * @tparam T The type to cast to * @return T* Typed pointer to underlying data, including the offset */ - template + template ())> __host__ __device__ T* data() const noexcept { return const_cast(detail::column_device_view_base::data()); @@ -586,15 +766,31 @@ class alignas(16) mutable_column_device_view : public detail::column_device_view * * This function accounts for the offset. * + * This function does not participate in overload resolution if `is_rep_layout_compatible` is + * false. Specializations of this function may exist for types `T` where + *`is_rep_layout_compatible` is false. + * + * * @tparam T The element type * @param element_index Position of the desired element */ - template - __device__ T& element(size_type element_index) noexcept + template ())> + __device__ T& element(size_type element_index) const noexcept { return data()[element_index]; } + /** + * @brief For a given `T`, indicates if `mutable_column_device_view::element()` has a valid + * overload. + * + */ + template + static constexpr bool has_element_accessor() + { + return has_element_accessor_impl::value; + } + /** * @brief Returns raw pointer to the underlying bitmask allocation. * @@ -618,11 +814,14 @@ class alignas(16) mutable_column_device_view : public detail::column_device_view * @brief Return first element (accounting for offset) after underlying data * is casted to the specified type. * + * This function does not participate in overload resolution if + * `mutable_column_device_view::has_element_accessor()` is false. + * * @tparam T The desired type * @return T* Pointer to the first element after casting */ - template - std::enable_if_t(), iterator> begin() + template ())> + iterator begin() { return iterator{count_it{0}, detail::mutable_value_accessor{*this}}; } @@ -631,11 +830,14 @@ class alignas(16) mutable_column_device_view : public detail::column_device_view * @brief Return one past the last element after underlying data is casted to * the specified type. * + * This function does not participate in overload resolution if + * `mutable_column_device_view::has_element_accessor()` is false. + * * @tparam T The desired type * @return T const* Pointer to one past the last element after casting */ - template - std::enable_if_t(), iterator> end() + template ())> + iterator end() { return iterator{count_it{size()}, detail::mutable_value_accessor{*this}}; } @@ -740,121 +942,6 @@ class alignas(16) mutable_column_device_view : public detail::column_device_view mutable_column_device_view(mutable_column_view source); }; -/** - * @brief Returns `string_view` to the string element at the specified index. - * - * If the element at the specified index is NULL, i.e., `is_null(element_index) - * == true`, then any attempt to use the result will lead to undefined behavior. - * - * This function accounts for the offset. - * - * @param element_index Position of the desired string element - * @return string_view instance representing this element at this index - */ -template <> -__device__ inline string_view const column_device_view::element( - size_type element_index) const noexcept -{ - size_type index = element_index + offset(); // account for this view's _offset - const int32_t* d_offsets = d_children[strings_column_view::offsets_column_index].data(); - const char* d_strings = d_children[strings_column_view::chars_column_index].data(); - size_type offset = d_offsets[index]; - return string_view{d_strings + offset, d_offsets[index + 1] - offset}; -} - -/** - * @brief Dispatch functor for resolving the index value for a dictionary element. - * - * The basic dictionary elements are the indices which can be any index type. - */ -struct index_element_fn { - template < - typename IndexType, - std::enable_if_t() and std::is_unsigned::value>* = nullptr> - __device__ size_type operator()(column_device_view const& input, size_type index) - { - return static_cast(input.element(index)); - } - template () and - std::is_unsigned::value)>* = nullptr> - __device__ size_type operator()(Args&&... args) - { - cudf_assert(false and "dictionary indices must be an unsigned integral type"); - return 0; - } -}; - -/** - * @brief Returns `dictionary32` element at the specified index for a - * dictionary column. - * - * `dictionary32` is a strongly typed wrapper around an `int32_t` value that holds the - * offset into the dictionary keys for the specified element. - * - * For example, given a dictionary column `d` with: - * ```c++ - * keys: {"foo", "bar", "baz"} - * indices: {2, 0, 2, 1, 0} - * - * d.element(0) == dictionary32{2}; - * d.element(1) == dictionary32{0}; - * ``` - * - * If the element at the specified index is NULL, i.e., `is_null(element_index) == true`, - * then any attempt to use the result will lead to undefined behavior. - * - * This function accounts for the offset. - * - * @param element_index Position of the desired element - * @return dictionary32 instance representing this element at this index - */ -template <> -__device__ inline dictionary32 const column_device_view::element( - size_type element_index) const noexcept -{ - size_type index = element_index + offset(); // account for this view's _offset - auto const indices = d_children[0]; - return dictionary32{type_dispatcher(indices.type(), index_element_fn{}, indices, index)}; -} - -/** - * @brief Returns a `numeric::decimal32` element at the specified index for a `fixed_point` column. - * - * If the element at the specified index is NULL, i.e., `is_null(element_index) == true`, - * then any attempt to use the result will lead to undefined behavior. - * - * @param element_index Position of the desired element - * @return numeric::decimal32 representing the element at this index - */ -template <> -__device__ inline numeric::decimal32 const column_device_view::element( - size_type element_index) const noexcept -{ - using namespace numeric; - auto const scale = scale_type{_type.scale()}; - return decimal32{scaled_integer{data()[element_index], scale}}; -} - -/** - * @brief Returns a `numeric::decimal64` element at the specified index for a `fixed_point` column. - * - * If the element at the specified index is NULL, i.e., `is_null(element_index) == true`, - * then any attempt to use the result will lead to undefined behavior. - * - * @param element_index Position of the desired element - * @return numeric::decimal64 representing the element at this index - */ -template <> -__device__ inline numeric::decimal64 const column_device_view::element( - size_type element_index) const noexcept -{ - using namespace numeric; - auto const scale = scale_type{_type.scale()}; - return decimal64{scaled_integer{data()[element_index], scale}}; -} - namespace detail { #ifdef __CUDACC__ // because set_bit in bit.hpp is wrapped with __CUDACC__ @@ -896,7 +983,6 @@ __device__ inline bitmask_type get_mask_offset_word(bitmask_type const* __restri * * @tparam T The type of elements in the column */ - template struct value_accessor { column_device_view const col; ///< column view of column in device @@ -1023,8 +1109,8 @@ struct mutable_value_accessor { }; /** - * @brief Helper function for use by column_device_view and mutable_column_device_view constructors - * to build device_views from views. + * @brief Helper function for use by column_device_view and mutable_column_device_view + * constructors to build device_views from views. * * It is used to build the array of child columns in device memory. Since child columns can * also have child columns, this uses recursion to build up the flat device buffer to contain diff --git a/cpp/include/cudf/column/column_view.hpp b/cpp/include/cudf/column/column_view.hpp index 168db61f672..82326a21d7d 100644 --- a/cpp/include/cudf/column/column_view.hpp +++ b/cpp/include/cudf/column/column_view.hpp @@ -16,6 +16,8 @@ #pragma once #include +#include + #include /** @@ -55,10 +57,14 @@ class column_view_base { *a column, and instead, accessing the elements should be done via *`data()`. * + * This function will only participate in overload resolution if `is_rep_layout_compatible()` + * or `std::is_same::value` are true. + * * @tparam The type to cast to * @return T const* Typed pointer to underlying data */ - template + template ::value or is_rep_layout_compatible())> T const* head() const noexcept { return static_cast(_data); @@ -70,12 +76,13 @@ class column_view_base { * * @note If `offset() == 0`, then `head() == data()` * - * @TODO Clarify behavior for variable-width types. + * This function does not participate in overload resolution if `is_rep_layout_compatible` is + * false. * * @tparam T The type to cast to * @return T const* Typed pointer to underlying data, including the offset */ - template + template ())> T const* data() const noexcept { return head() + _offset; @@ -85,10 +92,13 @@ class column_view_base { * @brief Return first element (accounting for offset) after underlying data * is casted to the specified type. * + * This function does not participate in overload resolution if `is_rep_layout_compatible` is + * false. + * * @tparam T The desired type * @return T const* Pointer to the first element after casting */ - template + template ())> T const* begin() const noexcept { return data(); @@ -98,10 +108,13 @@ class column_view_base { * @brief Return one past the last element after underlying data is casted to * the specified type. * + * This function does not participate in overload resolution if `is_rep_layout_compatible` is + * false. + * * @tparam T The desired type * @return T const* Pointer to one past the last element after casting */ - template + template ())> T const* end() const noexcept { return begin() + size(); @@ -438,6 +451,9 @@ class mutable_column_view : public detail::column_view_base { * @brief Returns pointer to the base device memory allocation casted to * the specified type. * + * This function will only participate in overload resolution if `is_rep_layout_compatible()` + * or `std::is_same::value` are true. + * * @note If `offset() == 0`, then `head() == data()` * * @note It should be rare to need to access the `head()` allocation of a @@ -446,7 +462,8 @@ class mutable_column_view : public detail::column_view_base { * @tparam The type to cast to * @return T* Typed pointer to underlying data */ - template + template ::value or is_rep_layout_compatible())> T* head() const noexcept { return const_cast(detail::column_view_base::head()); @@ -456,14 +473,15 @@ class mutable_column_view : public detail::column_view_base { * @brief Returns the underlying data casted to the specified type, plus the * offset. * - * @note If `offset() == 0`, then `head() == data()` + * This function does not participate in overload resolution if `is_rep_layout_compatible` is + * false. * - * @TODO Clarify behavior for variable-width types. + * @note If `offset() == 0`, then `head() == data()` * * @tparam T The type to cast to * @return T* Typed pointer to underlying data, including the offset */ - template + template ())> T* data() const noexcept { return const_cast(detail::column_view_base::data()); @@ -473,10 +491,13 @@ class mutable_column_view : public detail::column_view_base { * @brief Return first element (accounting for offset) when underlying data is * casted to the specified type. * + * This function does not participate in overload resolution if `is_rep_layout_compatible` is + * false. + * * @tparam T The desired type * @return T* Pointer to the first element after casting */ - template + template ())> T* begin() const noexcept { return const_cast(detail::column_view_base::begin()); @@ -486,10 +507,13 @@ class mutable_column_view : public detail::column_view_base { * @brief Return one past the last element after underlying data is casted to * the specified type. * + * This function does not participate in overload resolution if `is_rep_layout_compatible` is + * false. + * * @tparam T The desired type * @return T* Pointer to one past the last element after casting */ - template + template ())> T* end() const noexcept { return const_cast(detail::column_view_base::end()); diff --git a/cpp/include/cudf/detail/gather.cuh b/cpp/include/cudf/detail/gather.cuh index 73647ac2292..bf488621d52 100644 --- a/cpp/include/cudf/detail/gather.cuh +++ b/cpp/include/cudf/detail/gather.cuh @@ -139,6 +139,46 @@ void gather_helper(InputItr source_itr, } } +// Error case when no other overload or specialization is available +template +struct column_gatherer_impl { + std::unique_ptr operator()(...) { CUDF_FAIL("Unsupported type in gather."); } +}; + +/** + * @brief Function object for gathering a type-erased + * column. To be used with the cudf::type_dispatcher. + */ +struct column_gatherer { + /** + * @brief Type-dispatched function to gather from one column to another based + * on a `gather_map`. + * + * @tparam Element Dispatched type for the column being gathered + * @tparam MapIterator Iterator type for the gather map + * @param source_column View into the column to gather from + * @param gather_map_begin Beginning of iterator range of integral values representing the gather + * map + * @param gather_map_end End of iterator range of integral values representing the gather map + * @param nullify_out_of_bounds Nullify values in `gather_map` that are out of bounds + * @param stream CUDA stream used for device memory operations and kernel launches. + * @param mr Device memory resource used to allocate the returned column's device memory + */ + template + std::unique_ptr operator()(column_view const& source_column, + MapIterator gather_map_begin, + MapIterator gather_map_end, + bool nullify_out_of_bounds, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) + { + column_gatherer_impl gatherer{}; + + return gatherer( + source_column, gather_map_begin, gather_map_end, nullify_out_of_bounds, stream, mr); + } +}; + /** * @brief Function object for gathering a type-erased column. * @@ -148,8 +188,8 @@ void gather_helper(InputItr source_itr, * @tparam Element Dispatched type for the column being gathered * @tparam MapIterator Iterator type for the gather map */ -template -struct column_gatherer_impl { +template +struct column_gatherer_impl()>> { /** * @brief Type-dispatched function to gather from one column to another based * on a `gather_map`. @@ -164,6 +204,7 @@ struct column_gatherer_impl { * @param stream CUDA stream used for device memory operations and kernel launches. * @param mr Device memory resource used to allocate the returned column's device memory */ + template std::unique_ptr operator()(column_view const& source_column, MapIterator gather_map_begin, MapIterator gather_map_end, @@ -195,8 +236,8 @@ struct column_gatherer_impl { * * @tparam MapIterator Iterator type for the gather map */ -template -struct column_gatherer_impl { +template <> +struct column_gatherer_impl { /** * @brief Type-dispatched function to gather from one column to another based * on a `gather_map`. This handles string_view type column_views only. @@ -209,6 +250,7 @@ struct column_gatherer_impl { * @param stream CUDA stream used for device memory operations and kernel launches. * @param mr Device memory resource used to allocate the returned column's device memory */ + template std::unique_ptr operator()(column_view const& source_column, MapItType gather_map_begin, MapItType gather_map_end, @@ -234,8 +276,8 @@ struct column_gatherer_impl { * This functor is invoked only on the root column of a hierarchy of list * columns. Recursion is handled internally. */ -template -struct column_gatherer_impl { +template <> +struct column_gatherer_impl { /** * @brief Gather a list column from a hierarchy of list columns. * @@ -282,6 +324,7 @@ struct column_gatherer_impl { * @returns column with elements gathered based on the gather map * */ + template std::unique_ptr operator()(column_view const& column, MapItRoot gather_map_begin, MapItRoot gather_map_end, @@ -326,45 +369,11 @@ struct column_gatherer_impl { } }; -/** - * @brief Function object for gathering a type-erased - * column. To be used with the cudf::type_dispatcher. - */ -struct column_gatherer { - /** - * @brief Type-dispatched function to gather from one column to another based - * on a `gather_map`. - * - * @tparam Element Dispatched type for the column being gathered - * @tparam MapIterator Iterator type for the gather map - * @param source_column View into the column to gather from - * @param gather_map_begin Beginning of iterator range of integral values representing the gather - * map - * @param gather_map_end End of iterator range of integral values representing the gather map - * @param nullify_out_of_bounds Nullify values in `gather_map` that are out of bounds - * @param stream CUDA stream used for device memory operations and kernel launches. - * @param mr Device memory resource used to allocate the returned column's device memory - */ - template - std::unique_ptr operator()(column_view const& source_column, - MapIterator gather_map_begin, - MapIterator gather_map_end, - bool nullify_out_of_bounds, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) - { - column_gatherer_impl gatherer{}; - - return gatherer( - source_column, gather_map_begin, gather_map_end, nullify_out_of_bounds, stream, mr); - } -}; - /** * @brief Column gather specialization for dictionary column type. */ -template -struct column_gatherer_impl { +template <> +struct column_gatherer_impl { /** * @brief Type-dispatched function to gather from one column to another based * on a `gather_map`. @@ -378,6 +387,7 @@ struct column_gatherer_impl { * @param mr Device memory resource used to allocate the returned column's device memory * @return New dictionary column with gathered rows. */ + template std::unique_ptr operator()(column_view const& source_column, MapItType gather_map_begin, MapItType gather_map_end, @@ -426,6 +436,56 @@ struct column_gatherer_impl { } }; +template <> +struct column_gatherer_impl { + template + std::unique_ptr operator()(column_view const& column, + MapItRoot gather_map_begin, + MapItRoot gather_map_end, + bool nullify_out_of_bounds, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) + { + structs_column_view structs_column(column); + auto gather_map_size{std::distance(gather_map_begin, gather_map_end)}; + if (gather_map_size == 0) { return empty_like(column); } + + std::vector> output_struct_members; + std::transform(structs_column.child_begin(), + structs_column.child_end(), + std::back_inserter(output_struct_members), + [&gather_map_begin, &gather_map_end, nullify_out_of_bounds, stream, mr]( + cudf::column_view const& col) { + return cudf::type_dispatcher(col.type(), + column_gatherer{}, + col, + gather_map_begin, + gather_map_end, + nullify_out_of_bounds, + stream, + mr); + }); + + gather_bitmask( + // Table view of struct column. + cudf::table_view{ + std::vector{structs_column.child_begin(), structs_column.child_end()}}, + gather_map_begin, + output_struct_members, + nullify_out_of_bounds ? gather_bitmask_op::NULLIFY : gather_bitmask_op::DONT_CHECK, + stream, + mr); + + return cudf::make_structs_column( + gather_map_size, + std::move(output_struct_members), + 0, + rmm::device_buffer{0, stream, mr}, // Null mask will be fixed up in cudf::gather(). + stream, + mr); + } +}; + /** * @brief Function object for applying a transformation on the gathermap * that converts negative indices to positive indices @@ -538,55 +598,6 @@ void gather_bitmask(table_view const& source, } } -template -struct column_gatherer_impl { - std::unique_ptr operator()(column_view const& column, - MapItRoot gather_map_begin, - MapItRoot gather_map_end, - bool nullify_out_of_bounds, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) - { - structs_column_view structs_column(column); - auto gather_map_size{std::distance(gather_map_begin, gather_map_end)}; - if (gather_map_size == 0) { return empty_like(column); } - - std::vector> output_struct_members; - std::transform(structs_column.child_begin(), - structs_column.child_end(), - std::back_inserter(output_struct_members), - [&gather_map_begin, &gather_map_end, nullify_out_of_bounds, stream, mr]( - cudf::column_view const& col) { - return cudf::type_dispatcher(col.type(), - column_gatherer{}, - col, - gather_map_begin, - gather_map_end, - nullify_out_of_bounds, - stream, - mr); - }); - - gather_bitmask( - // Table view of struct column. - cudf::table_view{ - std::vector{structs_column.child_begin(), structs_column.child_end()}}, - gather_map_begin, - output_struct_members, - nullify_out_of_bounds ? gather_bitmask_op::NULLIFY : gather_bitmask_op::DONT_CHECK, - stream, - mr); - - return cudf::make_structs_column( - gather_map_size, - std::move(output_struct_members), - 0, - rmm::device_buffer{0, stream, mr}, // Null mask will be fixed up in cudf::gather(). - stream, - mr); - } -}; - /** * @brief Gathers the specified rows of a set of columns according to a gather map. * diff --git a/cpp/include/cudf/detail/scatter.cuh b/cpp/include/cudf/detail/scatter.cuh index 2cb1cbffc68..30764b9b89f 100644 --- a/cpp/include/cudf/detail/scatter.cuh +++ b/cpp/include/cudf/detail/scatter.cuh @@ -79,8 +79,14 @@ auto scatter_to_gather(MapIterator scatter_map_begin, return gather_map; } -template +template struct column_scatterer_impl { + std::unique_ptr operator()(...) const { CUDF_FAIL("Unsupported type for scatter."); } +}; + +template +struct column_scatterer_impl()>> { + template std::unique_ptr operator()(column_view const& source, MapIterator scatter_map_begin, MapIterator scatter_map_end, @@ -103,8 +109,9 @@ struct column_scatterer_impl { } }; -template -struct column_scatterer_impl { +template <> +struct column_scatterer_impl { + template std::unique_ptr operator()(column_view const& source, MapIterator scatter_map_begin, MapIterator scatter_map_end, @@ -119,8 +126,9 @@ struct column_scatterer_impl { } }; -template -struct column_scatterer_impl { +template <> +struct column_scatterer_impl { + template std::unique_ptr operator()(column_view const& source, MapIterator scatter_map_begin, MapIterator scatter_map_end, @@ -133,23 +141,9 @@ struct column_scatterer_impl { } }; -template -struct column_scatterer { - template - std::unique_ptr operator()(column_view const& source, - MapIterator scatter_map_begin, - MapIterator scatter_map_end, - column_view const& target, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) const - { - column_scatterer_impl scatterer{}; - return scatterer(source, scatter_map_begin, scatter_map_end, target, stream, mr); - } -}; - -template -struct column_scatterer_impl { +template <> +struct column_scatterer_impl { + template std::unique_ptr operator()(column_view const& source_in, MapIterator scatter_map_begin, MapIterator scatter_map_end, @@ -206,6 +200,20 @@ struct column_scatterer_impl { } }; +struct column_scatterer { + template + std::unique_ptr operator()(column_view const& source, + MapIterator scatter_map_begin, + MapIterator scatter_map_end, + column_view const& target, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) const + { + column_scatterer_impl scatterer{}; + return scatterer(source, scatter_map_begin, scatter_map_end, target, stream, mr); + } +}; + /** * @brief Scatters the rows of the source table into a copy of the target table * according to a scatter map. @@ -276,15 +284,13 @@ std::unique_ptr scatter( auto result = std::vector>(target.num_columns()); - auto scatter_functor = column_scatterer{}; - std::transform(source.begin(), source.end(), target.begin(), result.begin(), [=](auto const& source_col, auto const& target_col) { return type_dispatcher(source_col.type(), - scatter_functor, + column_scatterer{}, source_col, updated_scatter_map_begin, updated_scatter_map_end, diff --git a/cpp/include/cudf/lists/detail/scatter.cuh b/cpp/include/cudf/lists/detail/scatter.cuh index 5dd3db1117c..8e2ecdf49a7 100644 --- a/cpp/include/cudf/lists/detail/scatter.cuh +++ b/cpp/include/cudf/lists/detail/scatter.cuh @@ -577,15 +577,15 @@ struct list_child_constructor { auto child_offsets = cudf::strings::detail::make_offsets_child_column( begin, begin + child_list_views.size(), stream, mr); - auto child_column = - cudf::type_dispatcher(source_lists_column_view.child().child(1).type(), - list_child_constructor{}, - child_list_views, - child_offsets->view(), - cudf::lists_column_view(source_lists_column_view.child()), - cudf::lists_column_view(target_lists_column_view.child()), - stream, - mr); + auto child_column = cudf::type_dispatcher( + source_lists_column_view.child().child(1).type(), + list_child_constructor{}, + child_list_views, + child_offsets->view(), + cudf::lists_column_view(source_lists_column_view.child()), + cudf::lists_column_view(target_lists_column_view.child()), + stream, + mr); auto child_null_mask = source_lists_column_view.child().nullable() || target_lists_column_view.child().nullable() @@ -672,7 +672,7 @@ struct list_child_constructor { iter_target_member_as_list, std::back_inserter(child_columns), [&](auto source_struct_member_as_list, auto target_struct_member_as_list) { - return cudf::type_dispatcher( + return cudf::type_dispatcher( source_struct_member_as_list->child(cudf::lists_column_view::child_column_index).type(), list_child_constructor{}, list_vector, @@ -780,14 +780,14 @@ std::unique_ptr scatter( auto offsets_column = cudf::strings::detail::make_offsets_child_column( list_size_begin, list_size_begin + target.size(), stream, mr); - auto child_column = cudf::type_dispatcher(child_column_type, - list_child_constructor{}, - target_vector, - offsets_column->view(), - source_lists_column_view, - target_lists_column_view, - stream, - mr); + auto child_column = cudf::type_dispatcher(child_column_type, + list_child_constructor{}, + target_vector, + offsets_column->view(), + source_lists_column_view, + target_lists_column_view, + stream, + mr); auto null_mask = target.has_nulls() ? copy_bitmask(target, stream, mr) : rmm::device_buffer{0, stream, mr}; diff --git a/cpp/include/cudf/table/row_operators.cuh b/cpp/include/cudf/table/row_operators.cuh index 5af3c29a3d9..decd2879f54 100644 --- a/cpp/include/cudf/table/row_operators.cuh +++ b/cpp/include/cudf/table/row_operators.cuh @@ -28,6 +28,8 @@ #include #include +#include + namespace cudf { /** @@ -407,39 +409,47 @@ class row_lexicographic_comparator { template