diff --git a/cpp/include/cudf/column/column_child_offsets.hpp b/cpp/include/cudf/column/column_child_offsets.hpp new file mode 100644 index 000000000000..89baa563e17e --- /dev/null +++ b/cpp/include/cudf/column/column_child_offsets.hpp @@ -0,0 +1,25 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ +#pragma once + +/** + * @file column_child_offsets.hpp + * @brief Constants for child column indices within compound column types + */ + +#include +#include + +namespace CUDF_EXPORT cudf { + +static constexpr size_type offsets_column_index = 0; ///< Child index of the offsets column + +static constexpr size_type dictionary_indices_column_index = + 0; ///< Child index of the dictionary indices column + +static constexpr size_type dictionary_keys_column_index = + 1; ///< Child index of the dictionary key column + +} // namespace CUDF_EXPORT cudf diff --git a/cpp/include/cudf/column/column_device_view_base.cuh b/cpp/include/cudf/column/column_device_view_base.cuh index 3c1dce082b50..161aa969f4e0 100644 --- a/cpp/include/cudf/column/column_device_view_base.cuh +++ b/cpp/include/cudf/column/column_device_view_base.cuh @@ -1,9 +1,10 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once +#include #include #include #include @@ -61,6 +62,36 @@ struct nullate { }; }; +/** + * @brief A type tag to specify that a column should be treated as a dictionary column + * + * @tparam IndexType The type of the dictionary indices + * @tparam KeyType The type of the dictionary keys + */ +template + requires(is_index_type() && is_relationally_comparable()) +struct dictionary_element { + using index_type = IndexType; ///< The type of the dictionary indices + using key_type = KeyType; ///< The type of the dictionary keys + + key_type key{}; ///< The dictionary key for this element +}; + +/** + * @brief A type trait to determine if a type is a dictionary encoded type. + * @tparam T The type to check + */ +template +inline constexpr bool is_dictionary_encoded = false; + +/** + * @brief A type trait to determine if a type is a dictionary encoded type. + * @tparam IndexType The type of the dictionary indices + * @tparam KeyType The type of the dictionary keys + */ +template +inline constexpr bool is_dictionary_encoded> = true; + namespace detail { /** * @brief An immutable, non-owning view of device data as a column of elements @@ -68,8 +99,8 @@ namespace detail { */ class alignas(16) column_device_view_base { public: - // TODO: merge this offsets column index with `strings_column_view::offsets_column_index` - static constexpr size_type offsets_column_index{0}; ///< Child index of the offsets column + static constexpr size_type offsets_column_index = + cudf::offsets_column_index; ///< Child index of the offsets column column_device_view_base() = delete; ~column_device_view_base() = default; @@ -464,6 +495,33 @@ class alignas(16) column_device_view_core : public detail::column_device_view_ba return T{scaled_integer{data()[element_index], scale}}; } + /** + * @brief Returns a decoded copy of the 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. + * + * This function does not participate in overload resolution if `is_dictionary_encoded` is + * false. + * + * @tparam T The element type, e.g., `dictionary_element` for a dictionary column + * with `int32_t` indices and `float` keys + * @param element_index Position of the desired element + * @return The element at the specified index + */ + template )> + [[nodiscard]] __device__ decltype(auto) element(size_type element_index) const noexcept + { + auto const& indices = child(dictionary_indices_column_index); + auto const& keys = child(dictionary_keys_column_index); + auto const index = indices.template element( + element_index + offset()); // account for this view's _offset + return keys.template element(index); + } + /** * @brief Returns a nullable element at the specified index. If the element is null, returns * `nullopt`. @@ -538,7 +596,8 @@ class alignas(16) column_device_view_core : public detail::column_device_view_ba /** * @brief A mutable, non-owning view of device data as a column of elements - * that is trivially copyable and usable in CUDA device code and offline-compiled code (i.e. NVRTC). + * that is trivially copyable and usable in CUDA device code and offline-compiled code (i.e. + * NVRTC). * * @ingroup column_classes */ diff --git a/cpp/include/cudf/dictionary/dictionary_column_view.hpp b/cpp/include/cudf/dictionary/dictionary_column_view.hpp index 2c4659a508c8..df48faee484b 100644 --- a/cpp/include/cudf/dictionary/dictionary_column_view.hpp +++ b/cpp/include/cudf/dictionary/dictionary_column_view.hpp @@ -1,9 +1,10 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2024, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once +#include #include /** @@ -52,9 +53,9 @@ class dictionary_column_view : private column_view { dictionary_column_view& operator=(dictionary_column_view&&) = default; /// Index of the indices column of the dictionary column - static constexpr size_type indices_column_index{0}; + static constexpr size_type indices_column_index = cudf::dictionary_indices_column_index; /// Index of the keys column of the dictionary column - static constexpr size_type keys_column_index{1}; + static constexpr size_type keys_column_index = cudf::dictionary_keys_column_index; using column_view::has_nulls; using column_view::is_empty; diff --git a/cpp/include/cudf/transform.hpp b/cpp/include/cudf/transform.hpp index fc8e78b3fe05..25b385a5db23 100644 --- a/cpp/include/cudf/transform.hpp +++ b/cpp/include/cudf/transform.hpp @@ -95,9 +95,11 @@ struct transform_output { * * * @throws std::invalid_argument if any of the input columns have different sizes (except scalars) - * @throws std::invalid_argument if `output_type` or any of the input types are not supported. - * CUDA-supported types are fixed-width and string types, while PTX-supported types are integral and - * floating-point types. + * @throws std::invalid_argument if any of the output or input types are not supported. + * CUDA-supported input types are fixed-width, string, and their dictionary types. PTX-supported + * input types are integrals, floats, and their dictionary types. CUDA-supported output types are + * fixed-width, string, and their dictionary types. PTX-supported output types are integrals, + * floats, and their dictionary types. * @throws std::invalid_argument if the inputs only have a scalar with no column inputs and * `row_size` is not provided. This is because the row size cannot be inferred from the inputs in * this case. @@ -145,6 +147,11 @@ std::unique_ptr transform_extended( * @throws std::invalid_argument if the inputs only have a scalar with no column inputs and * `row_size` is not provided. This is because the row size cannot be inferred from the inputs in * this case. + * @throws std::invalid_argument if any of the output or input types are not supported. + * CUDA-supported input types are fixed-width, string, and their dictionary types. PTX-supported + * input types are integrals, floats, and their dictionary types. CUDA-supported output types are + * fixed-width, string, and their dictionary types. PTX-supported output types are integrals, + * floats, and their dictionary types. * @throws std::invalid_argument if string offsets are provided for non-string output columns, or * if the number of string offsets does not match the number of output columns. * diff --git a/cpp/src/jit/column_accessor.cuh b/cpp/src/jit/column_accessor.cuh index ef1536cb409d..9851bbe4b868 100644 --- a/cpp/src/jit/column_accessor.cuh +++ b/cpp/src/jit/column_accessor.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -45,7 +45,7 @@ struct column_accessor { return reinterpret_cast(cols[index]); } - static __device__ element_type element(auto const* __restrict__ cols, size_type row) + static __device__ auto element(auto const* __restrict__ cols, size_type row) { return column(cols).template element(map_index(row)); } @@ -60,8 +60,7 @@ struct column_accessor { return column(cols).is_valid(map_index(row)); } - static __device__ optional_element_type nullable_element(auto const* __restrict__ cols, - size_type row) + static __device__ auto nullable_element(auto const* __restrict__ cols, size_type row) { return column(cols).template nullable_element(map_index(row)); } diff --git a/cpp/src/jit/helpers.cpp b/cpp/src/jit/helpers.cpp index f7d78c08360f..a7620284b05c 100644 --- a/cpp/src/jit/helpers.cpp +++ b/cpp/src/jit/helpers.cpp @@ -1,16 +1,20 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #include "helpers.hpp" +#include #include +#include #include #include #include +#include + namespace cudf { namespace jit { @@ -77,18 +81,6 @@ std::map build_ptx_params(std::span ou return params; } -std::vector input_type_names( - std::span const> views) -{ - std::vector names; - - std::transform(views.begin(), views.end(), std::back_inserter(names), [&](auto const& view) { - return std::visit([](auto& a) { return type_to_name(a.type()); }, view); - }); - - return names; -} - kernel get_udf_kernel(std::string const& source_file, std::string const& kernel_name, std::string const& cuda_source) diff --git a/cpp/src/jit/helpers.hpp b/cpp/src/jit/helpers.hpp index 8a7685a63ebd..9d06d7883c33 100644 --- a/cpp/src/jit/helpers.hpp +++ b/cpp/src/jit/helpers.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -72,9 +72,6 @@ column_views_to_device(std::span views, return std::make_tuple(std::move(handles), std::move(device_array)); } -std::vector input_type_names( - std::span const> views); - kernel get_udf_kernel(std::string const& source_file, std::string const& kernel_name, std::string const& cuda_source); diff --git a/cpp/src/transform/transform.cu b/cpp/src/transform/transform.cu index 80b0509cfddf..117dc62922ab 100644 --- a/cpp/src/transform/transform.cu +++ b/cpp/src/transform/transform.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -36,6 +36,10 @@ namespace cudf { namespace { +column_view as_column_view(scalar_column_view const& scalar) { return scalar.as_column_view(); } + +column_view as_column_view(column_view const& column) { return column; } + struct mutable_fixed_width_column_view { mutable_column_view _view; @@ -200,13 +204,53 @@ void launch(cudf::kernel const& kernel, kernel.launch({cfg.min_grid_size}, {cfg.block_size}, 0, stream, args); } -std::string reflect_input_element(column_view const& c) { return type_to_name(c.type()); } +namespace { + +std::string get_element_type_name(column_view const& view); + +struct element_type_name_fn { + template + std::string operator()(column_view const& view) const + requires(is_fixed_width() || std::same_as) + { + return type_to_name(view.type()); + } -std::string reflect_input_element(scalar_column_view const& c) { return type_to_name(c.type()); } + template + std::string operator()(column_view const& view) const + requires(std::same_as) + { + return std::format("cudf::dictionary_element<{}, {}>", + get_element_type_name(view.child(cudf::dictionary_indices_column_index)), + get_element_type_name(view.child(cudf::dictionary_keys_column_index))); + } + + template + std::string operator()(column_view const& view) const + requires(!is_fixed_width() && !std::same_as && + !std::same_as) + { + CUDF_FAIL("Unsupported type for JIT compilation: " + type_to_name(view.type())); + } +}; + +std::string get_element_type_name(column_view const& view) +{ + return cudf::type_dispatcher(view.type(), element_type_name_fn{}, view); +} + +} // namespace + +std::string reflect_input_element(column_view const& c) { return get_element_type_name(c); } + +std::string reflect_input_element(scalar_column_view const& c) +{ + return get_element_type_name(c.as_column_view()); +} std::string reflect_output_element(fixed_width_column const& c) { - return type_to_name(c._col->type()); + return get_element_type_name(c._col->view()); } std::string reflect_output_element(string_views_column const&) { return "cudf::string_view"; } @@ -216,6 +260,33 @@ std::string reflect_output_element(mutable_strings_column const&) return "cuda::std::span"; } +std::string reflect_input_value_type(column_view const& c) +{ + return is_dictionary(c.type()) + ? reflect_input_value_type(c.child(cudf::dictionary_keys_column_index)) + : reflect_input_element(c); +} + +std::string reflect_input_value_type(scalar_column_view const& c) +{ + return reflect_input_value_type(c.as_column_view()); +} + +std::string reflect_output_value_type(fixed_width_column const& c) +{ + return reflect_output_element(c); +} + +std::string reflect_output_value_type(string_views_column const& c) +{ + return reflect_output_element(c); +} + +std::string reflect_output_value_type(mutable_strings_column const& c) +{ + return reflect_output_element(c); +} + std::string reflect_input_column(column_view const&) { return "cudf::column_device_view_core"; } std::string reflect_input_column(scalar_column_view const&) @@ -283,11 +354,12 @@ auto reflect(udf_source_type source_type, if (source_type == udf_source_type::PTX) { for (auto& in : inputs) { - ptx_in_types.push_back(std::visit([](auto& c) { return reflect_input_element(c); }, in)); + ptx_in_types.push_back(std::visit([](auto& c) { return reflect_input_value_type(c); }, in)); } for (auto& out : outputs) { - ptx_out_types.push_back(std::visit([](auto& c) { return reflect_output_element(c); }, out)); + ptx_out_types.push_back( + std::visit([](auto& c) { return reflect_output_value_type(c); }, out)); } } @@ -511,17 +583,20 @@ void perform_checks(udf_source_type source_type, std::span const> string_offsets) { if (source_type == udf_source_type::PTX) { - CUDF_EXPECTS(std::none_of(inputs.begin(), - inputs.end(), - [](auto& in) { - return std::visit( - [](auto& c) { - return !is_integral(c.type()) && !is_floating_point(c.type()); - }, - in); - }), - "Transforms with PTX UDFs only support integer, floating-point, and boolean", - std::invalid_argument); + static constexpr auto is_input_value_supported = [](auto const& c) { + return is_integral(c.type()) || is_floating_point(c.type()); + }; + static constexpr auto is_supported_input_type = [](auto const& c) { + auto col = std::visit([](auto& c) { return as_column_view(c); }, c); + return is_input_value_supported(col) || + (is_dictionary(col.type()) && + is_input_value_supported(col.child(dictionary_keys_column_index))); + }; + CUDF_EXPECTS( + std::none_of( + inputs.begin(), inputs.end(), [](auto const& in) { return !is_supported_input_type(in); }), + "Transforms with PTX UDFs only support integer, floating-point, and boolean", + std::invalid_argument); CUDF_EXPECTS(std::none_of(outputs.begin(), outputs.end(), [](auto& out) { @@ -542,14 +617,20 @@ void perform_checks(udf_source_type source_type, "Transforms only support output of fixed-width or string types", std::invalid_argument); - CUDF_EXPECTS(std::none_of(inputs.begin(), - inputs.end(), - [&](auto& in) { - auto type = std::visit([](auto& c) { return c.type(); }, in); - return !is_fixed_width(type) && type.id() != type_id::STRING; - }), - "Transforms only support input of fixed-width or string types", - std::invalid_argument); + static constexpr auto is_input_value_supported = [](auto const& c) { + return is_fixed_width(c.type()) || c.type().id() == type_id::STRING || is_dictionary(c.type()); + }; + static constexpr auto is_supported_input_type = [&](auto const& c) { + auto col = std::visit([](auto const& c) { return as_column_view(c); }, c); + return is_input_value_supported(col) || + (is_dictionary(col.type()) && + is_input_value_supported(col.child(dictionary_keys_column_index))); + }; + CUDF_EXPECTS( + std::none_of( + inputs.begin(), inputs.end(), [&](auto const& in) { return !is_supported_input_type(in); }), + "Transforms only support input of fixed-width, string, or dictionary types", + std::invalid_argument); if (!in_row_size.has_value()) { CUDF_EXPECTS( diff --git a/cpp/tests/transform/integration/unary_transform_test.cpp b/cpp/tests/transform/integration/unary_transform_test.cpp index d7c476e35b97..64738f54eb85 100644 --- a/cpp/tests/transform/integration/unary_transform_test.cpp +++ b/cpp/tests/transform/integration/unary_transform_test.cpp @@ -1,7 +1,7 @@ /* * SPDX-FileCopyrightText: Copyright 2018-2019 BlazingDB, Inc. * SPDX-FileCopyrightText: Copyright 2018 Christian Noboa Mardini - * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ /* @@ -31,6 +31,7 @@ #include #include +#include #include namespace transformation { @@ -413,6 +414,196 @@ __device__ inline void f(cudf::timestamp_us* output, cudf::timestamp_us input) test_udf(cuda, op, data_init, 500, cudf::udf_source_type::CUDA); } +TEST_F(UnaryOperationIntegrationTest, TransformDictionaryString) +{ + std::string const cuda = + R"***( +__device__ inline void decode(cudf::string_view * output, cudf::string_view input){ + *output = input; +})***"; + + // non-nullable + { + auto a = + cudf::test::strings_column_wrapper{ + "eee", "aaa", "ddd", "bbb", "ccc", "ccc", "ccc", "eee", "aaa"} + .release(); + + auto a_encoded = cudf::dictionary::encode(a->view()); + + cudf::transform_input inputs[] = {*a_encoded}; + + auto out = cudf::transform_extended( + inputs, cuda, cudf::data_type{cudf::type_id::STRING}, cudf::udf_source_type::CUDA); + + CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), a->view()); + } + + // nullable + { + auto a = + cudf::test::strings_column_wrapper{ + {"eee", "aaa", "ddd", "bbb", "ccc", "ccc", "ccc", "eee", "aaa"}, + {true, true, true, false, true, true, true, true, true}} + .release(); + + auto a_encoded = cudf::dictionary::encode(a->view()); + + cudf::transform_input inputs[] = {*a_encoded}; + + auto out = cudf::transform_extended( + inputs, cuda, cudf::data_type{cudf::type_id::STRING}, cudf::udf_source_type::CUDA); + + CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), a->view()); + } +} + +TEST_F(UnaryOperationIntegrationTest, TransformDictionaryFloat) +{ + std::string const cuda = + R"***( +__device__ inline void decode(float * output, float input){ + *output = input; +})***"; + + // Generated from NUMBA, using: + // + // ```py + // + // from numba import cuda, float32 + // from numba.cuda import compile_ptx_for_current_device + // + // # Define a CUDA device function + // + // @cuda.jit(device=True) + // def op(a): + // return a + // + // # Define argument types for the function + // arg_types = (float32, ) + // + // # Compile the device function as relocatable + // ptx, _ = cuda.compile_ptx_for_current_device(op, arg_types, device=True) + // + // + // # Print the PTX code + // print("Relocatable PTX Code:") + // print(ptx) + // + // + // ``` + // + std::string const ptx = R"***( +// +// Generated by NVIDIA NVVM Compiler +// +// Compiler Build ID: UNKNOWN +// Cuda compilation tools, release 13.2, V13.2.51 +// Based on NVVM 7.0.1 +// + +.version 9.2 +.target sm_86 +.address_size 64 + + // .globl _ZN8__main__3opxB2v1B96cw51cXTLSUwv1sCUt9Ww0FEw09RRQPKiLTj0gIGIFp_2b2oLQFEYYkHSQB1OQAk0Bynm21OizQ1K0UoIGvDpQE8oxrNQE_3dEf +.common .global .align 8 .u64 _ZN08NumbaEnv8__main__3opxB2v1B96cw51cXTLSUwv1sCUt9Ww0FEw09RRQPKiLTj0gIGIFp_2b2oLQFEYYkHSQB1OQAk0Bynm21OizQ1K0UoIGvDpQE8oxrNQE_3dEf; + +.visible .func (.param .b32 func_retval0) _ZN8__main__3opxB2v1B96cw51cXTLSUwv1sCUt9Ww0FEw09RRQPKiLTj0gIGIFp_2b2oLQFEYYkHSQB1OQAk0Bynm21OizQ1K0UoIGvDpQE8oxrNQE_3dEf( + .param .b64 _ZN8__main__3opxB2v1B96cw51cXTLSUwv1sCUt9Ww0FEw09RRQPKiLTj0gIGIFp_2b2oLQFEYYkHSQB1OQAk0Bynm21OizQ1K0UoIGvDpQE8oxrNQE_3dEf_param_0, + .param .b32 _ZN8__main__3opxB2v1B96cw51cXTLSUwv1sCUt9Ww0FEw09RRQPKiLTj0gIGIFp_2b2oLQFEYYkHSQB1OQAk0Bynm21OizQ1K0UoIGvDpQE8oxrNQE_3dEf_param_1 +) +{ + .reg .f32 %f<2>; + .reg .b32 %r<2>; + .reg .b64 %rd<2>; + + + ld.param.u64 %rd1, [_ZN8__main__3opxB2v1B96cw51cXTLSUwv1sCUt9Ww0FEw09RRQPKiLTj0gIGIFp_2b2oLQFEYYkHSQB1OQAk0Bynm21OizQ1K0UoIGvDpQE8oxrNQE_3dEf_param_0]; + ld.param.f32 %f1, [_ZN8__main__3opxB2v1B96cw51cXTLSUwv1sCUt9Ww0FEw09RRQPKiLTj0gIGIFp_2b2oLQFEYYkHSQB1OQAk0Bynm21OizQ1K0UoIGvDpQE8oxrNQE_3dEf_param_1]; + st.f32 [%rd1], %f1; + mov.u32 %r1, 0; + st.param.b32 [func_retval0+0], %r1; + ret; + +} +)***"; + + // non-nullable + { + auto a = cudf::test::fixed_width_column_wrapper( + {1.0F, 2.0F, 3.0F, 4.0F, 5.0F, 5.0F, 5.0F, 1.0F, 2.0F}) + .release(); + auto a_encoded = cudf::dictionary::encode(a->view()); + cudf::transform_input inputs[] = {*a_encoded}; + + auto out = cudf::transform_extended( + inputs, cuda, cudf::data_type{cudf::type_id::FLOAT32}, cudf::udf_source_type::CUDA); + + CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), a->view()); + + auto out_ptx = cudf::transform_extended( + inputs, ptx, cudf::data_type{cudf::type_id::FLOAT32}, cudf::udf_source_type::PTX); + + CUDF_TEST_EXPECT_COLUMNS_EQUAL(out_ptx->view(), a->view()); + } + + // nullable + { + auto a = cudf::test::fixed_width_column_wrapper( + {{1.0F, 2.0F, 3.0F, 4.0F, 5.0F, 5.0F, 5.0F, 1.0F, 2.0F}, + {true, true, true, true, true, false, true, true, true}}) + .release(); + + auto a_encoded = cudf::dictionary::encode(a->view()); + cudf::transform_input inputs[] = {*a_encoded}; + + auto out = cudf::transform_extended( + inputs, cuda, cudf::data_type{cudf::type_id::FLOAT32}, cudf::udf_source_type::CUDA); + + CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), a->view()); + + auto out_ptx = cudf::transform_extended( + inputs, ptx, cudf::data_type{cudf::type_id::FLOAT32}, cudf::udf_source_type::PTX); + + CUDF_TEST_EXPECT_COLUMNS_EQUAL(out_ptx->view(), a->view()); + } + + // empty column + { + auto a_empty = cudf::test::fixed_width_column_wrapper({}).release(); + auto a_encoded = cudf::dictionary::encode(a_empty->view()); + cudf::transform_input inputs[] = {*a_encoded}; + + auto out = cudf::transform_extended( + inputs, cuda, cudf::data_type{cudf::type_id::FLOAT32}, cudf::udf_source_type::CUDA); + + EXPECT_EQ(out->size(), 0); + } + + // sliced + { + auto a = cudf::test::fixed_width_column_wrapper( + {1.0F, 2.0F, 3.0F, 4.0F, 5.0F, 5.0F, 5.0F, 1.0F, 2.0F}) + .release(); + + auto a_encoded = cudf::dictionary::encode(a->view()); + auto sliced_input = cudf::slice(a_encoded->view(), {2, 7}).front(); + auto sliced_expect = cudf::slice(a->view(), {2, 7}).front(); + + cudf::transform_input inputs[] = {sliced_input}; + + auto out = cudf::transform_extended( + inputs, cuda, cudf::data_type{cudf::type_id::FLOAT32}, cudf::udf_source_type::CUDA); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), sliced_expect); + + auto out_ptx = cudf::transform_extended( + inputs, ptx, cudf::data_type{cudf::type_id::FLOAT32}, cudf::udf_source_type::PTX); + + CUDF_TEST_EXPECT_COLUMNS_EQUAL(out_ptx->view(), sliced_expect); + } +} + struct TernaryOperationTest : public cudf::test::BaseFixture {}; TEST_F(TernaryOperationTest, TransformWithScalar) diff --git a/docs/cudf/source/libcudf/api_docs/dictionary_classes.rst b/docs/cudf/source/libcudf/api_docs/dictionary_classes.rst index 00dec78c5f54..586c476afcb4 100644 --- a/docs/cudf/source/libcudf/api_docs/dictionary_classes.rst +++ b/docs/cudf/source/libcudf/api_docs/dictionary_classes.rst @@ -1,5 +1,9 @@ Dictionary Classes ================== +.. doxygenvariable:: cudf::dictionary_indices_column_index + +.. doxygenvariable:: cudf::dictionary_keys_column_index + .. doxygengroup:: dictionary_classes :members: