From 6b09253ac77a6b3f163a00380e39daa0884edf8d Mon Sep 17 00:00:00 2001 From: Robert Maynard Date: Wed, 19 May 2021 08:31:58 -0400 Subject: [PATCH] Correct unused parameters in the copying algorithms (#8232) 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: https://github.com/rapidsai/cudf/pull/8232 --- cpp/include/cudf/detail/copy_if.cuh | 4 +-- cpp/src/copying/contiguous_split.cu | 21 +++++-------- cpp/src/copying/copy.cu | 47 ++++++++--------------------- cpp/src/copying/copy_range.cu | 29 +++++------------- cpp/src/copying/get_element.cu | 9 ++---- cpp/src/copying/scatter.cu | 16 +++------- cpp/src/copying/segmented_shift.cu | 32 -------------------- cpp/src/copying/shift.cu | 3 +- 8 files changed, 37 insertions(+), 124 deletions(-) diff --git a/cpp/include/cudf/detail/copy_if.cuh b/cpp/include/cudf/detail/copy_if.cuh index fbf68a20364..2051daec00b 100644 --- a/cpp/include/cudf/detail/copy_if.cuh +++ b/cpp/include/cudf/detail/copy_if.cuh @@ -278,9 +278,9 @@ struct scatter_gather_functor { std::unique_ptr 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()) { diff --git a/cpp/src/copying/contiguous_split.cu b/cpp/src/copying/contiguous_split.cu index ce01090f7eb..82dc2ad7066 100644 --- a/cpp/src/copying/contiguous_split.cu +++ b/cpp/src/copying/contiguous_split.cu @@ -255,7 +255,6 @@ __device__ void copy_buffer(uint8_t* __restrict__ dst, */ template __global__ void copy_partition(int num_src_bufs, - int num_partitions, uint8_t** src_bufs, uint8_t** dst_bufs, dst_buf_info* buf_info) @@ -447,6 +446,13 @@ struct buf_info_functor { return {current + 1, offset_stack_pos + offset_depth}; } + template + std::enable_if_t::value, std::pair> + operator()(Args&&...) + { + CUDF_FAIL("Unsupported type"); + } + private: std::pair add_null_buffer(column_view const& col, src_buf_info* current, @@ -599,17 +605,6 @@ std::pair buf_info_functor::operator() -std::pair buf_info_functor::operator()( - column_view const& col, - src_buf_info* current, - int offset_stack_pos, - int parent_offset_index, - int offset_depth) -{ - CUDF_FAIL("Unsupported type"); -} - template std::pair setup_source_buf_info(InputIter begin, InputIter end, @@ -1023,7 +1018,7 @@ std::vector contiguous_split(cudf::table_view const& input, { constexpr size_type block_size = 256; copy_partition<<>>( - 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) diff --git a/cpp/src/copying/copy.cu b/cpp/src/copying/copy.cu index 3e962c76cf7..f12c9dcf006 100644 --- a/cpp/src/copying/copy.cu +++ b/cpp/src/copying/copy.cu @@ -241,43 +241,14 @@ std::unique_ptr 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 { - template - std::unique_ptr 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 { - template - std::unique_ptr 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 @@ -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 or std::is_same_v) { + (void)left_nullable; + (void)right_nullable; + return scatter_gather_based_if_else(lhs, rhs, size, filter, stream, mr); + } + copy_if_else_functor_impl copier{}; return copier(lhs, rhs, size, left_nullable, right_nullable, filter, stream, mr); } diff --git a/cpp/src/copying/copy_range.cu b/cpp/src/copying/copy_range.cu index f4ce9ea27ac..39a947d2ab9 100644 --- a/cpp/src/copying/copy_range.cu +++ b/cpp/src/copying/copy_range.cu @@ -90,17 +90,6 @@ struct out_of_place_copy_range_dispatch { cudf::column_view const& source; cudf::column_view const& target; - template ())> - std::unique_ptr 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 ())> std::unique_ptr operator()( cudf::size_type source_begin, @@ -122,6 +111,13 @@ struct out_of_place_copy_range_dispatch { return p_ret; } + + template + std::enable_if_t(), std::unique_ptr> + operator()(Args...) + { + CUDF_FAIL("Unsupported type for out of place copy."); + } }; template <> @@ -212,17 +208,6 @@ std::unique_ptr out_of_place_copy_range_dispatch::operator() -std::unique_ptr out_of_place_copy_range_dispatch::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) -{ - CUDF_FAIL("list_view type not supported"); -} - } // namespace namespace cudf { diff --git a/cpp/src/copying/get_element.cu b/cpp/src/copying/get_element.cu index fa5902efc0e..dc0334bd37b 100644 --- a/cpp/src/copying/get_element.cu +++ b/cpp/src/copying/get_element.cu @@ -171,12 +171,9 @@ struct get_element_functor { mr); } - template ::value> *p = nullptr> - std::unique_ptr 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 + std::enable_if_t::value, std::unique_ptr> operator()( + Args &&...) { CUDF_FAIL("get_element_functor not supported for struct_view"); } diff --git a/cpp/src/copying/scatter.cu b/cpp/src/copying/scatter.cu index cedac96cee6..3fccc2122cf 100644 --- a/cpp/src/copying/scatter.cu +++ b/cpp/src/copying/scatter.cu @@ -142,12 +142,8 @@ struct column_scalar_scatterer_impl { template struct column_scalar_scatterer_impl { - std::unique_ptr operator()(std::reference_wrapper 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 + std::unique_ptr operator()(Args&&...) const { CUDF_FAIL("scatter scalar to list_view not implemented"); } @@ -155,12 +151,8 @@ struct column_scalar_scatterer_impl { template struct column_scalar_scatterer_impl { - std::unique_ptr operator()(std::reference_wrapper 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 + std::unique_ptr operator()(Args&&...) const { CUDF_FAIL("scatter scalar to struct_view not implemented"); } diff --git a/cpp/src/copying/segmented_shift.cu b/cpp/src/copying/segmented_shift.cu index a1990f6234e..6fc785a61c6 100644 --- a/cpp/src/copying/segmented_shift.cu +++ b/cpp/src/copying/segmented_shift.cu @@ -205,38 +205,6 @@ struct segmented_shift_functor { } }; -/** - * @brief Segmented shift specialization for `list_view`. - */ -template <> -struct segmented_shift_functor { - std::unique_ptr operator()(column_view const& segmented_values, - device_span 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 { - std::unique_ptr operator()(column_view const& segmented_values, - device_span 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. diff --git a/cpp/src/copying/shift.cu b/cpp/src/copying/shift.cu index cf85bf51e80..ebeaf0e3b20 100644 --- a/cpp/src/copying/shift.cu +++ b/cpp/src/copying/shift.cu @@ -46,8 +46,7 @@ inline bool __device__ out_of_bounds(size_type size, size_type idx) struct shift_functor { template - std::enable_if_t(), std::unique_ptr> operator()( - Args&&... args) + std::enable_if_t(), std::unique_ptr> operator()(Args&&...) { CUDF_FAIL("shift does not support non-fixed-width types."); }