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

Correct unused parameters in the copying algorithms #8232

Merged
Show file tree
Hide file tree
Changes from all commits
Commits
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
4 changes: 2 additions & 2 deletions cpp/include/cudf/detail/copy_if.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -278,9 +278,9 @@ struct scatter_gather_functor {
std::unique_ptr<cudf::column> operator()(
cudf::column_view const& input,
cudf::size_type const& output_size,
cudf::size_type const* block_offsets,
cudf::size_type const*,
Filter filter,
cudf::size_type per_thread,
cudf::size_type,
rmm::cuda_stream_view stream,
rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource())
{
Expand Down
21 changes: 8 additions & 13 deletions cpp/src/copying/contiguous_split.cu
Original file line number Diff line number Diff line change
Expand Up @@ -255,7 +255,6 @@ __device__ void copy_buffer(uint8_t* __restrict__ dst,
*/
template <int block_size>
__global__ void copy_partition(int num_src_bufs,
int num_partitions,
uint8_t** src_bufs,
uint8_t** dst_bufs,
dst_buf_info* buf_info)
Expand Down Expand Up @@ -447,6 +446,13 @@ struct buf_info_functor {
return {current + 1, offset_stack_pos + offset_depth};
}

template <typename T, typename... Args>
std::enable_if_t<std::is_same<T, cudf::dictionary32>::value, std::pair<src_buf_info*, size_type>>
operator()(Args&&...)
{
CUDF_FAIL("Unsupported type");
}

private:
std::pair<src_buf_info*, size_type> add_null_buffer(column_view const& col,
src_buf_info* current,
Expand Down Expand Up @@ -599,17 +605,6 @@ std::pair<src_buf_info*, size_type> buf_info_functor::operator()<cudf::struct_vi
offset_depth);
}

template <>
std::pair<src_buf_info*, size_type> buf_info_functor::operator()<cudf::dictionary32>(
column_view const& col,
src_buf_info* current,
int offset_stack_pos,
int parent_offset_index,
int offset_depth)
{
CUDF_FAIL("Unsupported type");
}

template <typename InputIter>
std::pair<src_buf_info*, size_type> setup_source_buf_info(InputIter begin,
InputIter end,
Expand Down Expand Up @@ -1023,7 +1018,7 @@ std::vector<packed_table> contiguous_split(cudf::table_view const& input,
{
constexpr size_type block_size = 512;
copy_partition<block_size><<<num_bufs, block_size, 0, stream.value()>>>(
num_src_bufs, num_partitions, d_src_bufs, d_dst_bufs, d_dst_buf_info);
num_src_bufs, d_src_bufs, d_dst_bufs, d_dst_buf_info);
}

// DtoH dst info (to retrieve null counts)
Expand Down
47 changes: 12 additions & 35 deletions cpp/src/copying/copy.cu
Original file line number Diff line number Diff line change
Expand Up @@ -241,43 +241,14 @@ std::unique_ptr<column> scatter_gather_based_if_else(Left const& lhs,
// 1. Struct scalars are not yet available.
// 2. List scalars do not yet support explosion to a full column.
CUDF_FAIL("Scalars of nested types are not currently supported!");
(void)lhs;
robertmaynard marked this conversation as resolved.
Show resolved Hide resolved
(void)rhs;
(void)size;
(void)is_left;
(void)stream;
(void)mr;
}

/**
* @brief Specialization of copy_if_else_functor for list_views.
*/
template <>
struct copy_if_else_functor_impl<list_view> {
template <typename Left, typename Right, typename Filter>
std::unique_ptr<column> operator()(Left const& lhs,
Right const& rhs,
size_type size,
bool left_nullable,
bool right_nullable,
Filter filter,
rmm::cuda_stream_view stream,
rmm::mr::device_memory_resource* mr)
{
return scatter_gather_based_if_else(lhs, rhs, size, filter, stream, mr);
}
};

template <>
struct copy_if_else_functor_impl<struct_view> {
template <typename Left, typename Right, typename Filter>
std::unique_ptr<column> operator()(Left const& lhs,
Right const& rhs,
size_type size,
bool left_nullable,
bool right_nullable,
Filter filter,
rmm::cuda_stream_view stream,
rmm::mr::device_memory_resource* mr)
{
return scatter_gather_based_if_else(lhs, rhs, size, filter, stream, mr);
}
};

/**
* @brief Functor called by the `type_dispatcher` to invoke copy_if_else on combinations
* of column_view and scalar
Expand All @@ -293,6 +264,12 @@ struct copy_if_else_functor {
rmm::cuda_stream_view stream,
rmm::mr::device_memory_resource* mr)
{
if constexpr (std::is_same_v<T, cudf::list_view> or std::is_same_v<T, cudf::struct_view>) {
(void)left_nullable;
(void)right_nullable;
return scatter_gather_based_if_else(lhs, rhs, size, filter, stream, mr);
}

copy_if_else_functor_impl<T> copier{};
return copier(lhs, rhs, size, left_nullable, right_nullable, filter, stream, mr);
}
Expand Down
29 changes: 7 additions & 22 deletions cpp/src/copying/copy_range.cu
Original file line number Diff line number Diff line change
Expand Up @@ -90,17 +90,6 @@ struct out_of_place_copy_range_dispatch {
cudf::column_view const& source;
cudf::column_view const& target;

template <typename T, CUDF_ENABLE_IF(not cudf::is_rep_layout_compatible<T>())>
std::unique_ptr<cudf::column> operator()(
cudf::size_type source_begin,
cudf::size_type source_end,
cudf::size_type target_begin,
rmm::cuda_stream_view stream,
rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource())
{
CUDF_FAIL("Unsupported type for out of place copy.");
}

template <typename T, CUDF_ENABLE_IF(cudf::is_rep_layout_compatible<T>())>
std::unique_ptr<cudf::column> operator()(
cudf::size_type source_begin,
Expand All @@ -122,6 +111,13 @@ struct out_of_place_copy_range_dispatch {

return p_ret;
}

template <typename T, typename... Args>
std::enable_if_t<not cudf::is_rep_layout_compatible<T>(), std::unique_ptr<cudf::column>>
operator()(Args...)
{
CUDF_FAIL("Unsupported type for out of place copy.");
}
};

template <>
Expand Down Expand Up @@ -212,17 +208,6 @@ std::unique_ptr<cudf::column> out_of_place_copy_range_dispatch::operator()<cudf:
null_count);
}

template <>
std::unique_ptr<cudf::column> out_of_place_copy_range_dispatch::operator()<cudf::list_view>(
cudf::size_type source_begin,
cudf::size_type source_end,
cudf::size_type target_begin,
rmm::cuda_stream_view stream,
rmm::mr::device_memory_resource* mr)
{
CUDF_FAIL("list_view type not supported");
}

} // namespace

namespace cudf {
Expand Down
9 changes: 3 additions & 6 deletions cpp/src/copying/get_element.cu
Original file line number Diff line number Diff line change
Expand Up @@ -171,12 +171,9 @@ struct get_element_functor {
mr);
}

template <typename T, std::enable_if_t<std::is_same<T, struct_view>::value> *p = nullptr>
std::unique_ptr<scalar> operator()(
column_view const &input,
size_type index,
rmm::cuda_stream_view stream,
rmm::mr::device_memory_resource *mr = rmm::mr::get_current_device_resource())
template <typename T, typename... Args>
std::enable_if_t<std::is_same<T, struct_view>::value, std::unique_ptr<scalar>> operator()(
Args &&...)
{
CUDF_FAIL("get_element_functor not supported for struct_view");
}
Expand Down
16 changes: 4 additions & 12 deletions cpp/src/copying/scatter.cu
Original file line number Diff line number Diff line change
Expand Up @@ -142,25 +142,17 @@ struct column_scalar_scatterer_impl<string_view, MapIterator> {

template <typename MapIterator>
struct column_scalar_scatterer_impl<list_view, MapIterator> {
std::unique_ptr<column> operator()(std::reference_wrapper<const scalar> const& source,
MapIterator scatter_iter,
size_type scatter_rows,
column_view const& target,
rmm::cuda_stream_view stream,
rmm::mr::device_memory_resource* mr) const
template <typename... Args>
std::unique_ptr<column> operator()(Args&&...) const
{
CUDF_FAIL("scatter scalar to list_view not implemented");
}
};

template <typename MapIterator>
struct column_scalar_scatterer_impl<struct_view, MapIterator> {
std::unique_ptr<column> operator()(std::reference_wrapper<const scalar> const& source,
MapIterator scatter_iter,
size_type scatter_rows,
column_view const& target,
rmm::cuda_stream_view stream,
rmm::mr::device_memory_resource* mr) const
template <typename... Args>
std::unique_ptr<column> operator()(Args&&...) const
{
CUDF_FAIL("scatter scalar to struct_view not implemented");
}
Expand Down
32 changes: 0 additions & 32 deletions cpp/src/copying/segmented_shift.cu
Original file line number Diff line number Diff line change
Expand Up @@ -205,38 +205,6 @@ struct segmented_shift_functor<string_view> {
}
};

/**
* @brief Segmented shift specialization for `list_view`.
*/
template <>
struct segmented_shift_functor<list_view> {
std::unique_ptr<column> operator()(column_view const& segmented_values,
device_span<size_type const> segment_offsets,
size_type offset,
scalar const& fill_value,
rmm::cuda_stream_view stream,
rmm::mr::device_memory_resource* mr)
{
CUDF_FAIL("segmented_shift does not support list_view yet");
}
};

/**
* @brief Segmented shift specialization for `struct_view`.
*/
template <>
struct segmented_shift_functor<struct_view> {
std::unique_ptr<column> operator()(column_view const& segmented_values,
device_span<size_type const> segment_offsets,
size_type offset,
scalar const& fill_value,
rmm::cuda_stream_view stream,
rmm::mr::device_memory_resource* mr)
{
CUDF_FAIL("segmented_shift does not support struct_view yet");
}
};

/**
* @brief Functor to instantiate the specializations for segmented shift and
* forward arguments.
Expand Down
3 changes: 1 addition & 2 deletions cpp/src/copying/shift.cu
Original file line number Diff line number Diff line change
Expand Up @@ -46,8 +46,7 @@ inline bool __device__ out_of_bounds(size_type size, size_type idx)

struct shift_functor {
template <typename T, typename... Args>
std::enable_if_t<not cudf::is_fixed_width<T>(), std::unique_ptr<column>> operator()(
Args&&... args)
std::enable_if_t<not cudf::is_fixed_width<T>(), std::unique_ptr<column>> operator()(Args&&...)
{
CUDF_FAIL("shift does not support non-fixed-width types.");
}
Expand Down