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.

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

Disable column_view data accessors for unsupported types #7725

Merged
merged 43 commits into from Mar 30, 2021
Merged
Show file tree
Hide file tree
Changes from 35 commits
Commits
Show all changes
43 commits
Select commit Hold shift + click to select a range
d73219e
Disable column_view data accessors for non fixed-width types.
jrhemstad Mar 25, 2021
2f8e0dc
Fix scatter code path specialization.
jrhemstad Mar 25, 2021
759bf6d
Clarify that the static_assert will be removed in favor of enable_if.
jrhemstad Mar 25, 2021
3df7685
Use dispatch_storage_type in list scatter.
jrhemstad Mar 25, 2021
c2604cb
Add is_rep_layout_compatible trait.
jrhemstad Mar 25, 2021
5e6311c
Provide error path for dlpack.
jrhemstad Mar 25, 2021
5f818ad
Error path for from_arrow.
jrhemstad Mar 25, 2021
55ff95b
Error path for to_arrow.
jrhemstad Mar 25, 2021
24d437f
Error path for JIT.
jrhemstad Mar 25, 2021
608459f
Add error path for merge.
jrhemstad Mar 25, 2021
1453153
Modify specialization of gather code paths to add error path.
jrhemstad Mar 25, 2021
f99318e
Format.
jrhemstad Mar 25, 2021
a4788a0
Use dispatch_storage_type in get_data_ptr.
jrhemstad Mar 26, 2021
95d55ea
Error path for copy_range.
jrhemstad Mar 26, 2021
98072b0
Use is_rep_layout_compatible instead of is_fixed_width
jrhemstad Mar 26, 2021
0238112
Error path for fill.
jrhemstad Mar 26, 2021
978af0b
Error path for copy test.
jrhemstad Mar 26, 2021
9516e09
Update specializations and error path for copy_if_else.
jrhemstad Mar 26, 2021
fb27cf1
Disable column_device_view accessors for invalid types.
jrhemstad Mar 26, 2021
400dbaf
Add error path for type dispatcher benchmark.
jrhemstad Mar 26, 2021
dcd75e2
Error path in row_hasher for types that don't have an element accessor.
jrhemstad Mar 26, 2021
1134362
format.
jrhemstad Mar 26, 2021
86437a6
Error path for AST.
jrhemstad Mar 26, 2021
7b7beae
Add CUDF_ENABLE_IF macro.
jrhemstad Mar 27, 2021
daf2e72
Move explicit specializations of element<T> to SFINAE.
jrhemstad Mar 27, 2021
160a660
Add has_element_accessor to query if a type has
jrhemstad Mar 27, 2021
fd24294
Use has_element_accessor to provide error path.
jrhemstad Mar 27, 2021
7d61ace
SFINAE column_view accessors instead of static_assert.
jrhemstad Mar 27, 2021
1d475e7
Docs.
jrhemstad Mar 27, 2021
7c279ac
SFINAE mutable_column_view accessors.
jrhemstad Mar 27, 2021
dc70ba6
Consistent column_view docs.
jrhemstad Mar 29, 2021
9a9b0bc
Moved has_element_accessor_impl to base class.
jrhemstad Mar 29, 2021
3d08a57
SFINAE remaining device_view accessors.
jrhemstad Mar 29, 2021
e896a26
format.
jrhemstad Mar 29, 2021
3a6405d
Update cpp/benchmarks/type_dispatcher/type_dispatcher_benchmark.cu
jrhemstad Mar 29, 2021
5982248
Use the CUDF_ENABLE_IF macro everywhere.
jrhemstad Mar 29, 2021
d6ce9f0
Merge branch 'disable_column_view_begin' of github.com:jrhemstad/cudf…
jrhemstad Mar 29, 2021
7e9abd4
Format.
jrhemstad Mar 29, 2021
6a56c02
Remove extra spaces.
jrhemstad Mar 29, 2021
d20e9ae
Apply suggestions from code review
jrhemstad Mar 29, 2021
a057e71
Update cpp/src/copying/copy.cu
jrhemstad Mar 29, 2021
85c8c11
Merge branch 'disable_column_view_begin' of github.com:jrhemstad/cudf…
jrhemstad Mar 29, 2021
f92d462
Format.
jrhemstad Mar 29, 2021
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Jump to
Jump to file
Failed to load files.
Diff view
Diff view
19 changes: 17 additions & 2 deletions cpp/benchmarks/type_dispatcher/type_dispatcher_benchmark.cu
Expand Up @@ -31,6 +31,7 @@
#include <type_traits>
#include "../fixture/benchmark_fixture.hpp"
#include "../synchronization/synchronization.hpp"
#include <cudf/utilities/traits.hpp>

using namespace cudf;

Expand Down Expand Up @@ -87,14 +88,22 @@ __global__ void host_dispatching_kernel(mutable_column_device_view source_column

template <FunctorType functor_type>
struct ColumnHandle {
template <typename ColumnType>
template <typename ColumnType,
std::enable_if_t<cudf::is_rep_layout_compatible<ColumnType>()>* = nullptr>
void operator()(mutable_column_device_view source_column, int work_per_thread)
{
cudf::detail::grid_1d grid_config{source_column.size(), block_size};
int grid_size = grid_config.num_blocks;
// Launch the kernel.
host_dispatching_kernel<functor_type, ColumnType><<<grid_size, block_size>>>(source_column);
}

template <typename ColumnType,
std::enable_if_t<not cudf::is_rep_layout_compatible<ColumnType>()>* = nullptr>
void operator()(mutable_column_device_view source_column, int work_per_thread)
{
CUDF_FAIL("Invalid type to benchmark.");
}
};

// The following is for DEVICE_DISPATCHING:
Expand All @@ -104,12 +113,18 @@ struct ColumnHandle {
// n_rows * n_cols.
template <FunctorType functor_type>
struct RowHandle {
template <typename T>
template <typename T, std::enable_if_t<cudf::is_rep_layout_compatible<T>()>* = nullptr>
__device__ void operator()(mutable_column_device_view source, cudf::size_type index)
{
using F = Functor<T, functor_type>;
source.data<T>()[index] = F::f(source.data<T>()[index]);
}

template <typename T, std::enable_if_t<not cudf::is_rep_layout_compatible<T>()>* = nullptr>
__device__ void operator()(mutable_column_device_view source, cudf::size_type index)
{
cudf_assert(false && "Unsupported type.");
}
};

// This is for DEVICE_DISPATCHING
Expand Down
26 changes: 23 additions & 3 deletions cpp/include/cudf/ast/detail/transform.cuh
Expand Up @@ -20,10 +20,12 @@
#include <cudf/ast/operators.hpp>
#include <cudf/column/column_device_view.cuh>
#include <cudf/column/column_factories.hpp>
#include <cudf/detail/utilities/assert.cuh>
#include <cudf/scalar/scalar_device_view.cuh>
#include <cudf/table/table_device_view.cuh>
#include <cudf/table/table_view.hpp>
#include <cudf/types.hpp>
#include <cudf/utilities/traits.hpp>

#include <rmm/cuda_stream_view.hpp>

Expand Down Expand Up @@ -55,10 +57,19 @@ struct row_output {
* @param row_index Row index of data column.
* @param result Value to assign to output.
*/
template <typename Element>
template <typename Element, std::enable_if_t<is_rep_layout_compatible<Element>()>* = nullptr>
__device__ void resolve_output(detail::device_data_reference device_data_reference,
cudf::size_type row_index,
Element result) const;
// Definition below after row_evaluator is a complete type

template <typename Element, std::enable_if_t<not is_rep_layout_compatible<Element>()>* = nullptr>
__device__ void resolve_output(detail::device_data_reference device_data_reference,
cudf::size_type row_index,
Element result) const
{
cudf_assert(false && "Invalid type in resolve_output.");
}

private:
row_evaluator const& evaluator;
Expand Down Expand Up @@ -167,7 +178,7 @@ struct row_evaluator {
* @param row_index Row index of data column.
* @return Element
*/
template <typename Element>
template <typename Element, CUDF_ENABLE_IF(column_device_view::has_element_accessor<Element>())>
__device__ Element resolve_input(detail::device_data_reference device_data_reference,
cudf::size_type row_index) const
{
Expand All @@ -187,6 +198,15 @@ struct row_evaluator {
}
}

template <typename Element,
CUDF_ENABLE_IF(not column_device_view::has_element_accessor<Element>())>
__device__ Element resolve_input(detail::device_data_reference device_data_reference,
cudf::size_type row_index) const
{
cudf_assert(false && "Unsupported type in resolve_input.");
return {};
}

/**
* @brief Callable to perform a unary operation.
*
Expand Down Expand Up @@ -249,7 +269,7 @@ struct row_evaluator {
mutable_column_device_view* output_column;
};

template <typename Element>
template <typename Element, std::enable_if_t<is_rep_layout_compatible<Element>()>*>
__device__ void row_output::resolve_output(detail::device_data_reference device_data_reference,
cudf::size_type row_index,
Element result) const
Expand Down