Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

cudf::ast Small Refactorings #7352

Merged
merged 7 commits into from
Feb 18, 2021
Merged

cudf::ast Small Refactorings #7352

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
7 commits
Select commit Hold shift + click to select a range
4aeb8c8
Initial refactorings
codereport Feb 9, 2021
76d7b8d
More refactorings, clang-format fix
codereport Feb 9, 2021
ff056ba
Conditional expression
codereport Feb 9, 2021
396964f
Collapsing namespaces
codereport Feb 10, 2021
1bd0909
Clean up
codereport Feb 10, 2021
d084f88
Copyright
codereport Feb 17, 2021
b52f687
Add back shmem_
codereport Feb 18, 2021
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
41 changes: 20 additions & 21 deletions cpp/include/cudf/ast/detail/linearizer.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2020, NVIDIA CORPORATION.
* Copyright (c) 2020-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.
Expand All @@ -23,7 +23,6 @@
#include <cudf/utilities/error.hpp>

namespace cudf {

namespace ast {

// Forward declaration
Expand Down Expand Up @@ -64,8 +63,8 @@ struct alignas(8) device_data_reference {

const device_data_reference_type reference_type; // Source of data
const cudf::data_type data_type; // Type of data
const cudf::size_type
data_index; // The column index of a table, index of a literal, or index of an intermediate
const cudf::size_type data_index; // The column index of a table, index of a
// literal, or index of an intermediate
const table_reference table_source;

inline bool operator==(const device_data_reference& rhs) const
Expand Down Expand Up @@ -114,7 +113,7 @@ class linearizer {
* @param table The table used for evaluating the abstract syntax tree.
*/
linearizer(detail::node const& expr, cudf::table_view table)
: table(table), node_count(0), intermediate_counter()
: _table(table), _node_count(0), _intermediate_counter()
{
expr.accept(*this);
}
Expand All @@ -124,50 +123,50 @@ class linearizer {
*
* @return cudf::data_type
*/
cudf::data_type get_root_data_type() const;
cudf::data_type root_data_type() const;

/**
* @brief Get the maximum number of intermediates stored by the abstract syntax tree.
*
* @return cudf::size_type
*/
cudf::size_type get_intermediate_count() const { return intermediate_counter.get_max_used(); }
cudf::size_type intermediate_count() const { return _intermediate_counter.get_max_used(); }

/**
* @brief Get the device data references.
*
* @return std::vector<detail::device_data_reference>
*/
std::vector<detail::device_data_reference> const& get_data_references() const
std::vector<detail::device_data_reference> const& data_references() const
{
return data_references;
return _data_references;
}

/**
* @brief Get the operators.
*
* @return std::vector<ast_operator>
*/
std::vector<ast_operator> const& get_operators() const { return operators; }
std::vector<ast_operator> const& operators() const { return _operators; }

/**
* @brief Get the operator source indices.
*
* @return std::vector<cudf::size_type>
*/
std::vector<cudf::size_type> const& get_operator_source_indices() const
std::vector<cudf::size_type> const& operator_source_indices() const
{
return operator_source_indices;
return _operator_source_indices;
}

/**
* @brief Get the literal device views.
*
* @return std::vector<cudf::detail::fixed_width_scalar_device_view_base>
*/
std::vector<cudf::detail::fixed_width_scalar_device_view_base> const& get_literals() const
std::vector<cudf::detail::fixed_width_scalar_device_view_base> const& literals() const
{
return literals;
return _literals;
}

/**
Expand Down Expand Up @@ -225,13 +224,13 @@ class linearizer {
cudf::size_type add_data_reference(detail::device_data_reference data_ref);

// State information about the "linearized" GPU execution plan
cudf::table_view table;
cudf::size_type node_count;
intermediate_counter intermediate_counter;
std::vector<detail::device_data_reference> data_references;
std::vector<ast_operator> operators;
std::vector<cudf::size_type> operator_source_indices;
std::vector<cudf::detail::fixed_width_scalar_device_view_base> literals;
cudf::table_view _table;
cudf::size_type _node_count;
intermediate_counter _intermediate_counter;
std::vector<detail::device_data_reference> _data_references;
std::vector<ast_operator> _operators;
std::vector<cudf::size_type> _operator_source_indices;
std::vector<cudf::detail::fixed_width_scalar_device_view_base> _literals;
};

} // namespace detail
Expand Down
3 changes: 1 addition & 2 deletions cpp/include/cudf/ast/linearizer.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2020, NVIDIA CORPORATION.
* Copyright (c) 2020-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.
Expand All @@ -25,7 +25,6 @@
#include <cudf/utilities/error.hpp>

namespace cudf {

namespace ast {

/**
Expand Down
109 changes: 46 additions & 63 deletions cpp/src/ast/linearizer.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2020, NVIDIA CORPORATION.
* Copyright (c) 2020-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.
Expand All @@ -23,6 +23,8 @@
#include <cudf/utilities/error.hpp>
#include <cudf/utilities/traits.hpp>

#include <thrust/iterator/transform_iterator.h>

#include <algorithm>
#include <functional>
#include <iterator>
Expand Down Expand Up @@ -64,10 +66,10 @@ cudf::size_type linearizer::intermediate_counter::take()

void linearizer::intermediate_counter::give(cudf::size_type value)
{
// TODO: add comment
auto const lower_bound = std::lower_bound(used_values.cbegin(), used_values.cend(), value);
if ((*lower_bound == value) && (lower_bound != used_values.cend())) {
if ((lower_bound != used_values.cend()) && (*lower_bound == value))
used_values.erase(lower_bound);
}
}

/**
Expand All @@ -83,47 +85,33 @@ void linearizer::intermediate_counter::give(cudf::size_type value)
*/
cudf::size_type linearizer::intermediate_counter::find_first_missing() const
{
if ((used_values.empty()) || (used_values.front() != 0)) {
// Handle cases where the container is empty or first value is non-zero.
return 0;
} else {
// Search for the first non-contiguous pair of elements.
auto found = std::adjacent_find(used_values.cbegin(),
used_values.cend(),
[](auto const& a, auto const& b) { return a != b - 1; });
if (found != used_values.cend()) {
// A missing value was found and is returned.
return *found + 1;
} else {
// No missing elements. Return the next element in the sequence.
return used_values.size();
}
}
if (used_values.empty() || (used_values.front() != 0)) { return 0; }
// Search for the first non-contiguous pair of elements.
auto diff_not_one = [](auto a, auto b) { return a != b - 1; };
auto it = std::adjacent_find(used_values.cbegin(), used_values.cend(), diff_not_one);
return it != used_values.cend()
? *it + 1 // A missing value was found and is returned.
: used_values.size(); // No missing elements. Return the next element in the sequence.
}

cudf::size_type linearizer::visit(literal const& expr)
{
// Increment the node index
node_count++;
// Resolve node type
auto const data_type = expr.get_data_type();
// Construct a scalar device view
auto device_view = expr.get_value();
// Push literal
auto const literal_index = cudf::size_type(literals.size());
literals.push_back(device_view);
// Push data reference
_node_count++; // Increment the node index
auto const data_type = expr.get_data_type(); // Resolve node type
auto device_view = expr.get_value(); // Construct a scalar device view
auto const literal_index = cudf::size_type(_literals.size()); // Push literal
_literals.push_back(device_view);
auto const source = detail::device_data_reference(
detail::device_data_reference_type::LITERAL, data_type, literal_index);
detail::device_data_reference_type::LITERAL, data_type, literal_index); // Push data reference
return add_data_reference(source);
}

cudf::size_type linearizer::visit(column_reference const& expr)
{
// Increment the node index
node_count++;
_node_count++;
// Resolve node type
auto const data_type = expr.get_data_type(table);
auto const data_type = expr.get_data_type(_table);
// Push data reference
auto const source = detail::device_data_reference(detail::device_data_reference_type::COLUMN,
data_type,
Expand All @@ -135,38 +123,36 @@ cudf::size_type linearizer::visit(column_reference const& expr)
cudf::size_type linearizer::visit(expression const& expr)
{
// Increment the node index
auto const node_index = node_count++;
auto const node_index = _node_count++;
// Visit children (operands) of this node
auto const operand_data_reference_indices = visit_operands(expr.get_operands());
auto const operand_data_ref_indices = visit_operands(expr.get_operands());
// Resolve operand types
auto operand_types = std::vector<cudf::data_type>(operand_data_reference_indices.size());
std::transform(operand_data_reference_indices.cbegin(),
operand_data_reference_indices.cend(),
operand_types.begin(),
[this](auto const& data_reference_index) -> cudf::data_type {
return get_data_references()[data_reference_index].data_type;
});
auto data_ref = [this](auto const& index) { return data_references()[index].data_type; };
auto begin = thrust::make_transform_iterator(operand_data_ref_indices.cbegin(), data_ref);
auto end = begin + operand_data_ref_indices.size();
auto const operand_types = std::vector<cudf::data_type>(begin, end);

// Validate types of operand data references match
if (std::adjacent_find(operand_types.cbegin(), operand_types.cend(), std::not_equal_to<>()) !=
operand_types.cend()) {
CUDF_FAIL("An AST expression was provided non-matching operand types.");
}

// Give back intermediate storage locations that are consumed by this operation
std::for_each(
operand_data_reference_indices.cbegin(),
operand_data_reference_indices.cend(),
operand_data_ref_indices.cbegin(),
operand_data_ref_indices.cend(),
[this](auto const& data_reference_index) {
auto const operand_source = get_data_references()[data_reference_index];
auto const operand_source = data_references()[data_reference_index];
if (operand_source.reference_type == detail::device_data_reference_type::INTERMEDIATE) {
auto const intermediate_index = operand_source.data_index;
intermediate_counter.give(intermediate_index);
_intermediate_counter.give(intermediate_index);
}
});
// Resolve node type
auto const op = expr.get_operator();
auto const data_type = cudf::ast::detail::ast_operator_return_type(op, operand_types);
// Push operator
operators.push_back(op);
_operators.push_back(op);
// Push data reference
auto const output = [&]() {
if (node_index == 0) {
Expand All @@ -183,25 +169,22 @@ cudf::size_type linearizer::visit(expression const& expr)
CUDF_FAIL("The output data type is too large to be stored in an intermediate.");
}
return detail::device_data_reference(
detail::device_data_reference_type::INTERMEDIATE, data_type, intermediate_counter.take());
detail::device_data_reference_type::INTERMEDIATE, data_type, _intermediate_counter.take());
}
}();
auto const index = add_data_reference(output);
// Insert source indices from all operands (sources) and this operator (destination)
operator_source_indices.insert(operator_source_indices.end(),
operand_data_reference_indices.cbegin(),
operand_data_reference_indices.cend());
operator_source_indices.push_back(index);
_operator_source_indices.insert(_operator_source_indices.end(),
operand_data_ref_indices.cbegin(),
operand_data_ref_indices.cend());
_operator_source_indices.push_back(index);
return index;
}

cudf::data_type linearizer::get_root_data_type() const
cudf::data_type linearizer::root_data_type() const
{
if (get_data_references().empty()) {
return cudf::data_type(cudf::type_id::EMPTY);
} else {
return get_data_references().back().data_type;
}
return data_references().empty() ? cudf::data_type(cudf::type_id::EMPTY)
: data_references().back().data_type;
}

std::vector<cudf::size_type> linearizer::visit_operands(
Expand All @@ -219,12 +202,12 @@ cudf::size_type linearizer::add_data_reference(detail::device_data_reference dat
{
// If an equivalent data reference already exists, return its index. Otherwise add this data
// reference and return the new index.
auto const it = std::find(data_references.cbegin(), data_references.cend(), data_ref);
if (it != data_references.cend()) {
return std::distance(data_references.cbegin(), it);
auto const it = std::find(_data_references.cbegin(), _data_references.cend(), data_ref);
if (it != _data_references.cend()) {
return std::distance(_data_references.cbegin(), it);
} else {
data_references.push_back(data_ref);
return data_references.size() - 1;
_data_references.push_back(data_ref);
return _data_references.size() - 1;
}
}

Expand Down
24 changes: 11 additions & 13 deletions cpp/src/ast/transform.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -40,9 +40,7 @@
#include <type_traits>

namespace cudf {

namespace ast {

namespace detail {

/**
Expand Down Expand Up @@ -94,12 +92,12 @@ std::unique_ptr<column> compute_column(table_view const table,
{
// Linearize the AST
auto const expr_linearizer = linearizer(expr, table);
auto const data_references = expr_linearizer.get_data_references();
auto const literals = expr_linearizer.get_literals();
auto const operators = expr_linearizer.get_operators();
auto const data_references = expr_linearizer.data_references();
auto const literals = expr_linearizer.literals();
auto const operators = expr_linearizer.operators();
auto const num_operators = cudf::size_type(operators.size());
auto const operator_source_indices = expr_linearizer.get_operator_source_indices();
auto const expr_data_type = expr_linearizer.get_root_data_type();
auto const operator_source_indices = expr_linearizer.operator_source_indices();
auto const expr_data_type = expr_linearizer.root_data_type();

// Create ast_plan and device buffer
auto plan = ast_plan();
Expand Down Expand Up @@ -138,19 +136,19 @@ std::unique_ptr<column> compute_column(table_view const table,
cudf::mutable_column_device_view::create(output_column->mutable_view(), stream);

// Configure kernel parameters
auto const num_intermediates = expr_linearizer.get_intermediate_count();
auto const num_intermediates = expr_linearizer.intermediate_count();
auto const shmem_size_per_thread = static_cast<int>(sizeof(std::int64_t) * num_intermediates);
int device_id;
CUDA_TRY(cudaGetDevice(&device_id));
int shmem_per_block_limit;
int shmem_limit_per_block;
CUDA_TRY(
cudaDeviceGetAttribute(&shmem_per_block_limit, cudaDevAttrMaxSharedMemoryPerBlock, device_id));
cudaDeviceGetAttribute(&shmem_limit_per_block, cudaDevAttrMaxSharedMemoryPerBlock, device_id));
auto constexpr MAX_BLOCK_SIZE = 128;
auto const block_size =
(shmem_size_per_thread > 0)
? std::min(MAX_BLOCK_SIZE, shmem_per_block_limit / shmem_size_per_thread)
shmem_size_per_thread != 0
? std::min(MAX_BLOCK_SIZE, shmem_limit_per_block / shmem_size_per_thread)
: MAX_BLOCK_SIZE;
cudf::detail::grid_1d config(table_num_rows, block_size);
auto const config = cudf::detail::grid_1d{table_num_rows, block_size};
auto const shmem_size_per_block = shmem_size_per_thread * config.num_threads_per_block;

// Execute the kernel
Expand Down