Skip to content

Commit

Permalink
cudf::ast Small Refactorings (#7352)
Browse files Browse the repository at this point in the history 
While reading/debugging through AST code in order to understand it, I did/am doing a few small refactorings.

Authors:
  - Conor Hoekstra (@codereport)

Approvers:
  - Ram (Ramakrishna Prabhu) (@rgsl888prabhu)
  - Jake Hemstad (@jrhemstad)

URL: #7352
  • Loading branch information
@codereport
codereport authored Feb 18, 2021
1 parent 0bc7e15 commit 593dc1c
Show file tree
Hide file tree
Showing 4 changed files with 78 additions and 99 deletions.
41 changes: 20 additions & 21 deletions cpp/include/cudf/ast/detail/linearizer.hpp
View file
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
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
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
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

0 comments on commit 593dc1c

Please sign in to comment.