Skip to content

Commit

Permalink
Remove validation that requires introspection (#11938)
Browse files Browse the repository at this point in the history
This PR removes optional validation for some APIs. Performing these validations requires data introspection, which we do not want. This PR resolves #5505.

Authors:
  - Vyas Ramasubramani (https://github.com/vyasr)
  - David Wendt (https://github.com/davidwendt)

Approvers:
  - Mark Harris (https://github.com/harrism)
  - GALI PREM SAGAR (https://github.com/galipremsagar)
  - Matthew Roeschke (https://github.com/mroeschke)
  - David Wendt (https://github.com/davidwendt)
  - Nghia Truong (https://github.com/ttnghia)
  - Jason Lowe (https://github.com/jlowe)

URL: #11938
  • Loading branch information
vyasr authored Oct 20, 2022
1 parent 416d4d5 commit ff41841
Show file tree
Hide file tree
Showing 24 changed files with 99 additions and 267 deletions.
2 changes: 1 addition & 1 deletion cpp/benchmarks/lists/copying/scatter_lists.cu
Original file line number Diff line number Diff line change
Expand Up @@ -108,7 +108,7 @@ void BM_lists_scatter(::benchmark::State& state)

for (auto _ : state) {
cuda_event_timer raii(state, true); // flush_l2_cache = true, stream = 0
scatter(table_view{{*source}}, *scatter_map, table_view{{*target}}, false, mr);
scatter(table_view{{*source}}, *scatter_map, table_view{{*target}}, mr);
}

state.SetBytesProcessed(static_cast<int64_t>(state.iterations()) * state.range(0) * 2 *
Expand Down
18 changes: 5 additions & 13 deletions cpp/include/cudf/copying.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -140,30 +140,26 @@ std::unique_ptr<column> reverse(
* If the same index appears more than once in the scatter map, the result is
* undefined.
*
* If any values in `scatter_map` are outside of the interval [-n, n) where `n`
* is the number of rows in the `target` table, behavior is undefined.
*
* A negative value `i` in the `scatter_map` is interpreted as `i+n`, where `n`
* is the number of rows in the `target` table.
*
* @throws cudf::logic_error if `check_bounds == true` and an index exists in
* `scatter_map` outside the range `[-n, n)`, where `n` is the number of rows in
* the target table. If `check_bounds == false`, the behavior is undefined.
*
* @param source The input columns containing values to be scattered into the
* target columns
* @param scatter_map A non-nullable column of integral indices that maps the
* rows in the source table to rows in the target table. The size must be equal
* to or less than the number of elements in the source columns.
* @param target The set of columns into which values from the source_table
* are to be scattered
* @param check_bounds Optionally perform bounds checking on the values of
* `scatter_map` and throw an error if any of its values are out of bounds.
* @param mr Device memory resource used to allocate the returned table's device memory
* @return Result of scattering values from source to target
*/
std::unique_ptr<table> scatter(
table_view const& source,
column_view const& scatter_map,
table_view const& target,
bool check_bounds = false,
rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());

/**
Expand All @@ -184,26 +180,22 @@ std::unique_ptr<table> scatter(
* If the same index appears more than once in the scatter map, the result is
* undefined.
*
* @throws cudf::logic_error if `check_bounds == true` and an index exists in
* `scatter_map` outside the range `[-n, n)`, where `n` is the number of rows in
* the target table. If `check_bounds == false`, the behavior is undefined.
* If any values in `scatter_map` are outside of the interval [-n, n) where `n`
* is the number of rows in the `target` table, behavior is undefined.
*
* @param source The input scalars containing values to be scattered into the
* target columns
* @param indices A non-nullable column of integral indices that indicate
* the rows in the target table to be replaced by source.
* @param target The set of columns into which values from the source_table
* are to be scattered
* @param check_bounds Optionally perform bounds checking on the values of
* `scatter_map` and throw an error if any of its values are out of bounds.
* @param mr Device memory resource used to allocate the returned table's device memory
* @return Result of scattering values from source to target
*/
std::unique_ptr<table> scatter(
std::vector<std::reference_wrapper<const scalar>> const& source,
column_view const& indices,
table_view const& target,
bool check_bounds = false,
rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());

/**
Expand Down
11 changes: 0 additions & 11 deletions cpp/include/cudf/detail/scatter.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -390,24 +390,13 @@ std::unique_ptr<table> scatter(
MapIterator scatter_map_begin,
MapIterator scatter_map_end,
table_view const& target,
bool check_bounds = false,
rmm::cuda_stream_view stream = cudf::default_stream_value,
rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource())
{
CUDF_FUNC_RANGE();

using MapType = typename thrust::iterator_traits<MapIterator>::value_type;

if (check_bounds) {
auto const begin = -target.num_rows();
auto const end = target.num_rows();
auto bounds = bounds_checker<MapType>{begin, end};
CUDF_EXPECTS(
std::distance(scatter_map_begin, scatter_map_end) ==
thrust::count_if(rmm::exec_policy(stream), scatter_map_begin, scatter_map_end, bounds),
"Scatter map index out of bounds");
}

CUDF_EXPECTS(std::distance(scatter_map_begin, scatter_map_end) <= source.num_rows(),
"scatter map size should be <= to number of rows in source");

Expand Down
18 changes: 4 additions & 14 deletions cpp/include/cudf/detail/scatter.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -45,10 +45,8 @@ namespace detail {
*
* If the same index appears more than once in the scatter map, the result is
* undefined.
*
* @throws cudf::logic_error if `check_bounds == true` and an index exists in
* `scatter_map` outside the range `[-n, n)`, where `n` is the number of rows in
* the target table. If `check_bounds == false`, the behavior is undefined.
* If any values in `scatter_map` are outside of the interval [-n, n) where `n`
* is the number of rows in the `target` table, behavior is undefined.
*
* @param source The input columns containing values to be scattered into the
* target columns
Expand All @@ -57,8 +55,6 @@ namespace detail {
* to or less than the number of elements in the source columns.
* @param target The set of columns into which values from the source_table
* are to be scattered
* @param check_bounds Optionally perform bounds checking on the values of
* `scatter_map` and throw an error if any of its values are out of bounds.
* @param stream CUDA stream used for device memory operations and kernel launches.
* @param mr Device memory resource used to allocate the returned table's device memory
* @return Result of scattering values from source to target
Expand All @@ -67,7 +63,6 @@ std::unique_ptr<table> scatter(
table_view const& source,
column_view const& scatter_map,
table_view const& target,
bool check_bounds = false,
rmm::cuda_stream_view stream = cudf::default_stream_value,
rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());

Expand All @@ -81,7 +76,6 @@ std::unique_ptr<table> scatter(
table_view const& source,
device_span<size_type const> const scatter_map,
table_view const& target,
bool check_bounds = false,
rmm::cuda_stream_view stream = cudf::default_stream_value,
rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());

Expand All @@ -101,18 +95,15 @@ std::unique_ptr<table> scatter(
* If the same index appears more than once in the scatter map, the result is
* undefined.
*
* @throws cudf::logic_error if `check_bounds == true` and an index exists in
* `scatter_map` outside the range `[-n, n)`, where `n` is the number of rows in
* the target table. If `check_bounds == false`, the behavior is undefined.
* If any values in `indices` are outside of the interval [-n, n) where `n`
* is the number of rows in the `target` table, behavior is undefined.
*
* @param source The input scalars containing values to be scattered into the
* target columns
* @param indices A non-nullable column of integral indices that indicate
* the rows in the target table to be replaced by source.
* @param target The set of columns into which values from the source_table
* are to be scattered
* @param check_bounds Optionally perform bounds checking on the values of
* `scatter_map` and throw an error if any of its values are out of bounds.
* @param stream CUDA stream used for device memory operations and kernel launches.
* @param mr Device memory resource used to allocate the returned table's device memory
* @return Result of scattering values from source to target
Expand All @@ -121,7 +112,6 @@ std::unique_ptr<table> scatter(
std::vector<std::reference_wrapper<const scalar>> const& source,
column_view const& indices,
table_view const& target,
bool check_bounds = false,
rmm::cuda_stream_view stream = cudf::default_stream_value,
rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());

Expand Down
8 changes: 2 additions & 6 deletions cpp/include/cudf/filling.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -103,26 +103,22 @@ std::unique_ptr<column> fill(
* ```
* @p count should not have null values; should not contain negative values;
* and the sum of count elements should not overflow the size_type's limit.
* It is undefined behavior if @p count has negative values or the sum overflows
* and @p check_count is set to false.
* The behavior of this function is undefined if @p count has negative values
* or the sum overflows.
*
* @throws cudf::logic_error if the data type of @p count is not size_type.
* @throws cudf::logic_error if @p input_table and @p count have different
* number of rows.
* @throws cudf::logic_error if @p count has null values.
* @throws cudf::logic_error if @p check_count is set to true and @p count
* has negative values or the sum of @p count elements overflows.
*
* @param input_table Input table
* @param count Non-nullable column of an integral type
* @param check_count Whether to check count (negative values and overflow)
* @param mr Device memory resource used to allocate the returned table's device memory
* @return The result table containing the repetitions
*/
std::unique_ptr<table> repeat(
table_view const& input_table,
column_view const& count,
bool check_count = false,
rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());

/**
Expand Down
9 changes: 2 additions & 7 deletions cpp/src/copying/copy.cu
Original file line number Diff line number Diff line change
Expand Up @@ -180,7 +180,6 @@ std::unique_ptr<column> scatter_gather_based_if_else(cudf::column_view const& lh
table_view{std::vector<column_view>{scatter_src_lhs->get_column(0).view()}},
gather_map,
table_view{std::vector<column_view>{rhs}},
false,
stream,
mr);

Expand Down Expand Up @@ -208,12 +207,8 @@ std::unique_ptr<column> scatter_gather_based_if_else(cudf::scalar const& lhs,
static_cast<cudf::size_type>(scatter_map_size),
scatter_map.begin()};

auto result = cudf::detail::scatter(scatter_source,
scatter_map_column_view,
table_view{std::vector<column_view>{rhs}},
false,
stream,
mr);
auto result = cudf::detail::scatter(
scatter_source, scatter_map_column_view, table_view{std::vector<column_view>{rhs}}, stream, mr);

return std::move(result->release()[0]);
}
Expand Down
32 changes: 6 additions & 26 deletions cpp/src/copying/scatter.cu
Original file line number Diff line number Diff line change
Expand Up @@ -285,7 +285,6 @@ struct column_scalar_scatterer_impl<struct_view, MapIterator> {
std::unique_ptr<table> scatter(table_view const& source,
column_view const& scatter_map,
table_view const& target,
bool check_bounds,
rmm::cuda_stream_view stream,
rmm::mr::device_memory_resource* mr)
{
Expand All @@ -307,13 +306,12 @@ std::unique_ptr<table> scatter(table_view const& source,
// create index type normalizing iterator for the scatter_map
auto map_begin = indexalator_factory::make_input_iterator(scatter_map);
auto map_end = map_begin + scatter_map.size();
return detail::scatter(source, map_begin, map_end, target, check_bounds, stream, mr);
return detail::scatter(source, map_begin, map_end, target, stream, mr);
}

std::unique_ptr<table> scatter(table_view const& source,
device_span<size_type const> const scatter_map,
table_view const& target,
bool check_bounds,
rmm::cuda_stream_view stream,
rmm::mr::device_memory_resource* mr)
{
Expand All @@ -322,13 +320,12 @@ std::unique_ptr<table> scatter(table_view const& source,
auto map_col = column_view(data_type{type_to_id<size_type>()},
static_cast<size_type>(scatter_map.size()),
scatter_map.data());
return scatter(source, map_col, target, check_bounds, stream, mr);
return scatter(source, map_col, target, stream, mr);
}

std::unique_ptr<table> scatter(std::vector<std::reference_wrapper<const scalar>> const& source,
column_view const& indices,
table_view const& target,
bool check_bounds,
rmm::cuda_stream_view stream,
rmm::mr::device_memory_resource* mr)
{
Expand All @@ -340,20 +337,9 @@ std::unique_ptr<table> scatter(std::vector<std::reference_wrapper<const scalar>>

// Create normalizing iterator for indices column
auto map_begin = indexalator_factory::make_input_iterator(indices);
auto map_end = map_begin + indices.size();

// Optionally check map index values are within the number of target rows.
auto const n_rows = target.num_rows();
if (check_bounds) {
CUDF_EXPECTS(
indices.size() == thrust::count_if(rmm::exec_policy(stream),
map_begin,
map_end,
[n_rows] __device__(size_type index) {
return ((index >= -n_rows) && (index < n_rows));
}),
"Scatter map index out of bounds");
}

// Transform negative indices to index + target size
auto scatter_rows = indices.size();
Expand Down Expand Up @@ -404,12 +390,8 @@ std::unique_ptr<column> boolean_mask_scatter(column_view const& input,
// The scatter map is actually a table with only one column, which is scatter map.
auto scatter_map =
detail::apply_boolean_mask(table_view{{indices->view()}}, boolean_mask, stream);
auto output_table = detail::scatter(table_view{{input}},
scatter_map->get_column(0).view(),
table_view{{target}},
false,
stream,
mr);
auto output_table = detail::scatter(
table_view{{input}}, scatter_map->get_column(0).view(), table_view{{target}}, stream, mr);

// There is only one column in output_table
return std::make_unique<column>(std::move(output_table->get_column(0)));
Expand Down Expand Up @@ -505,21 +487,19 @@ std::unique_ptr<table> boolean_mask_scatter(
std::unique_ptr<table> scatter(table_view const& source,
column_view const& scatter_map,
table_view const& target,
bool check_bounds,
rmm::mr::device_memory_resource* mr)
{
CUDF_FUNC_RANGE();
return detail::scatter(source, scatter_map, target, check_bounds, cudf::default_stream_value, mr);
return detail::scatter(source, scatter_map, target, cudf::default_stream_value, mr);
}

std::unique_ptr<table> scatter(std::vector<std::reference_wrapper<const scalar>> const& source,
column_view const& indices,
table_view const& target,
bool check_bounds,
rmm::mr::device_memory_resource* mr)
{
CUDF_FUNC_RANGE();
return detail::scatter(source, indices, target, check_bounds, cudf::default_stream_value, mr);
return detail::scatter(source, indices, target, cudf::default_stream_value, mr);
}

std::unique_ptr<table> boolean_mask_scatter(table_view const& input,
Expand Down
11 changes: 1 addition & 10 deletions cpp/src/filling/repeat.cu
Original file line number Diff line number Diff line change
Expand Up @@ -103,7 +103,6 @@ namespace cudf {
namespace detail {
std::unique_ptr<table> repeat(table_view const& input_table,
column_view const& count,
bool check_count,
rmm::cuda_stream_view stream,
rmm::mr::device_memory_resource* mr)
{
Expand All @@ -112,19 +111,12 @@ std::unique_ptr<table> repeat(table_view const& input_table,

if (input_table.num_rows() == 0) { return cudf::empty_like(input_table); }

if (check_count) { cudf::type_dispatcher(count.type(), count_checker{count}, stream); }

auto count_iter = cudf::detail::indexalator_factory::make_input_iterator(count);

rmm::device_uvector<cudf::size_type> offsets(count.size(), stream);
thrust::inclusive_scan(
rmm::exec_policy(stream), count_iter, count_iter + count.size(), offsets.begin());

if (check_count) {
CUDF_EXPECTS(thrust::is_sorted(rmm::exec_policy(stream), offsets.begin(), offsets.end()),
"count has negative values or the resulting table has too many rows.");
}

size_type output_size{offsets.back_element(stream)};
rmm::device_uvector<size_type> indices(output_size, stream);
thrust::upper_bound(rmm::exec_policy(stream),
Expand Down Expand Up @@ -162,11 +154,10 @@ std::unique_ptr<table> repeat(table_view const& input_table,

std::unique_ptr<table> repeat(table_view const& input_table,
column_view const& count,
bool check_count,
rmm::mr::device_memory_resource* mr)
{
CUDF_FUNC_RANGE();
return detail::repeat(input_table, count, check_count, cudf::default_stream_value, mr);
return detail::repeat(input_table, count, cudf::default_stream_value, mr);
}

std::unique_ptr<table> repeat(table_view const& input_table,
Expand Down
6 changes: 3 additions & 3 deletions cpp/src/groupby/sort/scan.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -178,9 +178,9 @@ void scan_result_functor::operator()<aggregation::RANK>(aggregation const& agg)
stream,
mr);
}
result = std::move(cudf::detail::scatter(
table_view{{*result}}, *gather_map, table_view{{*result}}, false, stream, mr)
->release()[0]);
result = std::move(
cudf::detail::scatter(table_view{{*result}}, *gather_map, table_view{{*result}}, stream, mr)
->release()[0]);
if (rank_agg._null_handling == null_policy::EXCLUDE) {
result->set_null_mask(cudf::detail::copy_bitmask(get_grouped_values(), stream, mr));
}
Expand Down
1 change: 0 additions & 1 deletion cpp/src/groupby/sort/sort_helper.cu
Original file line number Diff line number Diff line change
Expand Up @@ -244,7 +244,6 @@ column_view sort_groupby_helper::unsorted_keys_labels(rmm::cuda_stream_view stre
cudf::detail::scatter(table_view({group_labels_view}),
scatter_map,
table_view({temp_labels->view()}),
false,
stream,
rmm::mr::get_current_device_resource());

Expand Down
4 changes: 2 additions & 2 deletions cpp/src/partitioning/partitioning.cu
Original file line number Diff line number Diff line change
Expand Up @@ -610,7 +610,7 @@ std::pair<std::unique_ptr<table>, std::vector<size_type>> hash_partition_table(

// Use the resulting scatter map to materialize the output
auto output = detail::scatter(
input, row_partition_numbers.begin(), row_partition_numbers.end(), input, false, stream, mr);
input, row_partition_numbers.begin(), row_partition_numbers.end(), input, stream, mr);

stream.synchronize(); // Async D2H copy must finish before returning host vec
return std::pair(std::move(output), std::move(partition_offsets));
Expand Down Expand Up @@ -698,7 +698,7 @@ struct dispatch_map_type {

// Scatter the rows into their partitions
auto scattered =
cudf::detail::scatter(t, scatter_map.begin(), scatter_map.end(), t, false, stream, mr);
cudf::detail::scatter(t, scatter_map.begin(), scatter_map.end(), t, stream, mr);

return std::pair(std::move(scattered), std::move(partition_offsets));
}
Expand Down
Loading

0 comments on commit ff41841

Please sign in to comment.