From 7f3c77536878d57496bed5dfaa9675e3dba01396 Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Thu, 9 Sep 2021 10:04:32 -0700 Subject: [PATCH] Move device_data_references and a few traits to new file. --- .../ast/detail/device_data_references.hpp | 115 ++++++++++++++++++ .../cudf/ast/detail/expression_evaluator.cuh | 2 +- .../cudf/ast/detail/expression_parser.hpp | 81 +----------- 3 files changed, 117 insertions(+), 81 deletions(-) create mode 100644 cpp/include/cudf/ast/detail/device_data_references.hpp diff --git a/cpp/include/cudf/ast/detail/device_data_references.hpp b/cpp/include/cudf/ast/detail/device_data_references.hpp new file mode 100644 index 00000000000..eab9613c832 --- /dev/null +++ b/cpp/include/cudf/ast/detail/device_data_references.hpp @@ -0,0 +1,115 @@ +/* + * Copyright (c) 2021, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ +#pragma once + +#include +#include +#include + +#include + +//#include +//#include +// +namespace cudf { +namespace ast { +namespace detail { + +/** + * @brief Node data reference types. + * + * This enum is device-specific. For instance, intermediate data references are generated by the + * linearization process but cannot be explicitly created by the user. + */ +enum class device_data_reference_type { + COLUMN, // A value in a table column + LITERAL, // A literal value + INTERMEDIATE // An internal temporary value +}; + +/** + * @brief A device data reference describes a source of data used by a expression. + * + * This is a POD class used to create references describing data type and locations for consumption + * by the `row_evaluator`. + */ +struct alignas(8) device_data_reference { + device_data_reference(device_data_reference_type reference_type, + cudf::data_type data_type, + cudf::size_type data_index, + table_reference table_source); + + device_data_reference(device_data_reference_type reference_type, + cudf::data_type data_type, + cudf::size_type data_index); + + device_data_reference_type const reference_type; // Source of data + cudf::data_type const data_type; // Type of data + cudf::size_type const data_index; // The column index of a table, index of a + // literal, or index of an intermediate + table_reference const table_source; + + bool operator==(device_data_reference const& rhs) const + { + return std::tie(data_index, reference_type, table_source) == + std::tie(rhs.data_index, rhs.reference_type, rhs.table_source); + } +}; + +// Type trait for wrapping nullable types in a thrust::optional. Non-nullable +// types are returned as is. +template +struct possibly_null_value; + +template +struct possibly_null_value { + using type = thrust::optional; +}; + +template +struct possibly_null_value { + using type = T; +}; + +template +using possibly_null_value_t = typename possibly_null_value::type; + +// Type used for intermediate storage in expression evaluation. +template +using IntermediateDataType = possibly_null_value_t; + +/** + * @brief A container of all device data required to evaluate an expression on tables. + * + * This struct should never be instantiated directly. It is created by the + * `expression_parser` on construction, and the resulting member is publicly accessible + * for passing to kernels for constructing an `expression_evaluator`. + * + */ +struct expression_device_view { + device_span data_references; + device_span literals; + device_span operators; + device_span operator_source_indices; + cudf::size_type num_intermediates; + int shmem_per_thread; +}; + +} // namespace detail + +} // namespace ast + +} // namespace cudf diff --git a/cpp/include/cudf/ast/detail/expression_evaluator.cuh b/cpp/include/cudf/ast/detail/expression_evaluator.cuh index 8d38fb4064c..7da43064330 100644 --- a/cpp/include/cudf/ast/detail/expression_evaluator.cuh +++ b/cpp/include/cudf/ast/detail/expression_evaluator.cuh @@ -15,7 +15,7 @@ */ #pragma once -#include +#include #include #include #include diff --git a/cpp/include/cudf/ast/detail/expression_parser.hpp b/cpp/include/cudf/ast/detail/expression_parser.hpp index 1f35b54ea61..5848d4f4c9b 100644 --- a/cpp/include/cudf/ast/detail/expression_parser.hpp +++ b/cpp/include/cudf/ast/detail/expression_parser.hpp @@ -15,6 +15,7 @@ */ #pragma once +#include #include #include #include @@ -30,86 +31,6 @@ namespace cudf { namespace ast { namespace detail { -/** - * @brief Node data reference types. - * - * This enum is device-specific. For instance, intermediate data references are generated by the - * linearization process but cannot be explicitly created by the user. - */ -enum class device_data_reference_type { - COLUMN, // A value in a table column - LITERAL, // A literal value - INTERMEDIATE // An internal temporary value -}; - -/** - * @brief A device data reference describes a source of data used by a expression. - * - * This is a POD class used to create references describing data type and locations for consumption - * by the `row_evaluator`. - */ -struct alignas(8) device_data_reference { - device_data_reference(device_data_reference_type reference_type, - cudf::data_type data_type, - cudf::size_type data_index, - table_reference table_source); - - device_data_reference(device_data_reference_type reference_type, - cudf::data_type data_type, - cudf::size_type data_index); - - device_data_reference_type const reference_type; // Source of data - cudf::data_type const data_type; // Type of data - cudf::size_type const data_index; // The column index of a table, index of a - // literal, or index of an intermediate - table_reference const table_source; - - bool operator==(device_data_reference const& rhs) const - { - return std::tie(data_index, reference_type, table_source) == - std::tie(rhs.data_index, rhs.reference_type, rhs.table_source); - } -}; - -// Type trait for wrapping nullable types in a thrust::optional. Non-nullable -// types are returned as is. -template -struct possibly_null_value; - -template -struct possibly_null_value { - using type = thrust::optional; -}; - -template -struct possibly_null_value { - using type = T; -}; - -template -using possibly_null_value_t = typename possibly_null_value::type; - -// Type used for intermediate storage in expression evaluation. -template -using IntermediateDataType = possibly_null_value_t; - -/** - * @brief A container of all device data required to evaluate an expression on tables. - * - * This struct should never be instantiated directly. It is created by the - * `expression_parser` on construction, and the resulting member is publicly accessible - * for passing to kernels for constructing an `expression_evaluator`. - * - */ -struct expression_device_view { - device_span data_references; - device_span literals; - device_span operators; - device_span operator_source_indices; - cudf::size_type num_intermediates; - int shmem_per_thread; -}; - /** * @brief The expression_parser traverses an expression and converts it into a form suitable for * execution on the device.