Skip to content

Commit

Permalink
Move device_data_references and a few traits to new file.
Browse files Browse the repository at this point in the history
  • Loading branch information
vyasr committed Sep 9, 2021
1 parent 758a2f0 commit 7f3c775
Show file tree
Hide file tree
Showing 3 changed files with 117 additions and 81 deletions.
115 changes: 115 additions & 0 deletions cpp/include/cudf/ast/detail/device_data_references.hpp
Original file line number Diff line number Diff line change
@@ -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 <cudf/ast/expressions.hpp>
#include <cudf/scalar/scalar_device_view.cuh>
#include <cudf/types.hpp>

#include <thrust/optional.h>

//#include <functional>
//#include <numeric>
//
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 <typename T, bool has_nulls>
struct possibly_null_value;

template <typename T>
struct possibly_null_value<T, true> {
using type = thrust::optional<T>;
};

template <typename T>
struct possibly_null_value<T, false> {
using type = T;
};

template <typename T, bool has_nulls>
using possibly_null_value_t = typename possibly_null_value<T, has_nulls>::type;

// Type used for intermediate storage in expression evaluation.
template <bool has_nulls>
using IntermediateDataType = possibly_null_value_t<std::int64_t, has_nulls>;

/**
* @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<device_data_reference const> data_references;
device_span<cudf::detail::fixed_width_scalar_device_view_base const> literals;
device_span<ast_operator const> operators;
device_span<size_type const> operator_source_indices;
cudf::size_type num_intermediates;
int shmem_per_thread;
};

} // namespace detail

} // namespace ast

} // namespace cudf
2 changes: 1 addition & 1 deletion cpp/include/cudf/ast/detail/expression_evaluator.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -15,7 +15,7 @@
*/
#pragma once

#include <cudf/ast/detail/expression_parser.hpp>
#include <cudf/ast/detail/device_data_references.hpp>
#include <cudf/ast/detail/operators.hpp>
#include <cudf/ast/expressions.hpp>
#include <cudf/column/column_device_view.cuh>
Expand Down
81 changes: 1 addition & 80 deletions cpp/include/cudf/ast/detail/expression_parser.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -15,6 +15,7 @@
*/
#pragma once

#include <cudf/ast/detail/device_data_references.hpp>
#include <cudf/ast/expressions.hpp>
#include <cudf/scalar/scalar_device_view.cuh>
#include <cudf/table/table_view.hpp>
Expand All @@ -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 <typename T, bool has_nulls>
struct possibly_null_value;

template <typename T>
struct possibly_null_value<T, true> {
using type = thrust::optional<T>;
};

template <typename T>
struct possibly_null_value<T, false> {
using type = T;
};

template <typename T, bool has_nulls>
using possibly_null_value_t = typename possibly_null_value<T, has_nulls>::type;

// Type used for intermediate storage in expression evaluation.
template <bool has_nulls>
using IntermediateDataType = possibly_null_value_t<std::int64_t, has_nulls>;

/**
* @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<detail::device_data_reference const> data_references;
device_span<cudf::detail::fixed_width_scalar_device_view_base const> literals;
device_span<ast_operator const> operators;
device_span<cudf::size_type const> 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.
Expand Down

0 comments on commit 7f3c775

Please sign in to comment.