Skip to content

Commit

Permalink
Correct unused parameters in the copying algorithms (#8232)
Browse files Browse the repository at this point in the history
Starting in CUDA 11.3, nvcc will start to unconditionally warn about unused parameters on functions/methods that are in anonymous namespaces.

Authors:
  - Robert Maynard (https://github.com/robertmaynard)

Approvers:
  - David Wendt (https://github.com/davidwendt)
  - Mark Harris (https://github.com/harrism)

URL: #8232
  • Loading branch information
robertmaynard authored May 19, 2021
1 parent 7e3249b commit 6b09253
Show file tree
Hide file tree
Showing 8 changed files with 37 additions and 124 deletions.
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 = 256;
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;
(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

0 comments on commit 6b09253

Please sign in to comment.