From a6c839bb63f1af3315823f0bcd8e9c332080566b Mon Sep 17 00:00:00 2001 From: Robert Maynard Date: Tue, 11 May 2021 14:41:56 -0400 Subject: [PATCH 1/7] Remove unused parameter from copy_partition kernel --- cpp/src/copying/contiguous_split.cu | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/cpp/src/copying/contiguous_split.cu b/cpp/src/copying/contiguous_split.cu index 1a840a6f39a..5d73a30dc73 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) @@ -1023,7 +1022,7 @@ std::vector contiguous_split(cudf::table_view const& input, { constexpr size_type block_size = 512; 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) From 40e7b46f98e451a64c4549c82af34464e52fcf38 Mon Sep 17 00:00:00 2001 From: Robert Maynard Date: Thu, 13 May 2021 07:53:25 -0400 Subject: [PATCH 2/7] Correct 'declared but never referenced' warnings in src/copying Found by CUDA 11.3 --- cpp/src/copying/contiguous_split.cu | 17 +++++------ cpp/src/copying/copy.cu | 46 ++++++++--------------------- cpp/src/copying/copy_range.cu | 32 +++++++------------- cpp/src/copying/get_element.cu | 9 ++---- cpp/src/copying/scatter.cu | 17 ++++------- cpp/src/copying/segmented_shift.cu | 42 +++++++------------------- cpp/src/copying/shift.cu | 3 +- 7 files changed, 48 insertions(+), 118 deletions(-) diff --git a/cpp/src/copying/contiguous_split.cu b/cpp/src/copying/contiguous_split.cu index 5d73a30dc73..b6d11ee02ab 100644 --- a/cpp/src/copying/contiguous_split.cu +++ b/cpp/src/copying/contiguous_split.cu @@ -446,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, @@ -598,16 +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, diff --git a/cpp/src/copying/copy.cu b/cpp/src/copying/copy.cu index 3e962c76cf7..ab276311dac 100644 --- a/cpp/src/copying/copy.cu +++ b/cpp/src/copying/copy.cu @@ -241,42 +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 @@ -293,6 +265,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..dc80bccd0d7 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,16 @@ struct out_of_place_copy_range_dispatch { return p_ret; } + + template + std::enable_if_t(), std::unique_ptr> + operator()(Args...) + { + if constexpr (std::is_same_v) { + CUDF_FAIL("list_view type not supported"); + } + CUDF_FAIL("Unsupported type for out of place copy."); + } }; template <> @@ -212,17 +211,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..524cad1a187 100644 --- a/cpp/src/copying/scatter.cu +++ b/cpp/src/copying/scatter.cu @@ -142,25 +142,18 @@ 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"); } }; + 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..095e1b80588 100644 --- a/cpp/src/copying/segmented_shift.cu +++ b/cpp/src/copying/segmented_shift.cu @@ -121,6 +121,16 @@ struct segmented_shift_functor { template std::unique_ptr operator()(Args&&...) { + if constexpr( std::is_same_v ) + { + CUDF_FAIL("segmented_shift does not support list_view yet"); + } + + if constexpr( std::is_same_v ) + { + CUDF_FAIL("segmented_shift does not support struct_view yet"); + } + CUDF_FAIL("Unsupported type for segmented_shift."); } }; @@ -205,38 +215,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."); } From af97322a7b1bc042bc9442fd48600c531b68cfa8 Mon Sep 17 00:00:00 2001 From: Robert Maynard Date: Thu, 13 May 2021 07:55:28 -0400 Subject: [PATCH 3/7] Correct 'declared but never referenced' warnings in copy_if.cuh Found by CUDA 11.3 --- cpp/include/cudf/detail/copy_if.cuh | 4 ++-- 1 file changed, 2 insertions(+), 2 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()) { From 82f5be57e91d80d568259d88c06b78458e04d715 Mon Sep 17 00:00:00 2001 From: Robert Maynard Date: Thu, 13 May 2021 09:37:36 -0400 Subject: [PATCH 4/7] correct copy and paste bug found in review --- cpp/src/copying/copy_range.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/src/copying/copy_range.cu b/cpp/src/copying/copy_range.cu index dc80bccd0d7..63da83b3dfe 100644 --- a/cpp/src/copying/copy_range.cu +++ b/cpp/src/copying/copy_range.cu @@ -116,7 +116,7 @@ struct out_of_place_copy_range_dispatch { std::enable_if_t(), std::unique_ptr> operator()(Args...) { - if constexpr (std::is_same_v) { + if constexpr (std::is_same_v) { CUDF_FAIL("list_view type not supported"); } CUDF_FAIL("Unsupported type for out of place copy."); From e7219c18d9f985347d0fad9f5e3bf24cce2210c0 Mon Sep 17 00:00:00 2001 From: Robert Maynard Date: Thu, 13 May 2021 16:51:18 -0400 Subject: [PATCH 5/7] revise failure messages based on review feedback --- cpp/src/copying/copy_range.cu | 3 --- 1 file changed, 3 deletions(-) diff --git a/cpp/src/copying/copy_range.cu b/cpp/src/copying/copy_range.cu index 63da83b3dfe..39a947d2ab9 100644 --- a/cpp/src/copying/copy_range.cu +++ b/cpp/src/copying/copy_range.cu @@ -116,9 +116,6 @@ struct out_of_place_copy_range_dispatch { std::enable_if_t(), std::unique_ptr> operator()(Args...) { - if constexpr (std::is_same_v) { - CUDF_FAIL("list_view type not supported"); - } CUDF_FAIL("Unsupported type for out of place copy."); } }; From 13cb7b56e20c65653cdc1800c3de74eba0c2b33d Mon Sep 17 00:00:00 2001 From: Robert Maynard Date: Fri, 14 May 2021 07:53:38 -0400 Subject: [PATCH 6/7] Correct style issues found by CI --- cpp/src/copying/contiguous_split.cu | 1 - cpp/src/copying/copy.cu | 13 ++++++------- cpp/src/copying/scatter.cu | 1 - cpp/src/copying/segmented_shift.cu | 6 ++---- 4 files changed, 8 insertions(+), 13 deletions(-) diff --git a/cpp/src/copying/contiguous_split.cu b/cpp/src/copying/contiguous_split.cu index b6d11ee02ab..fe2e12adff5 100644 --- a/cpp/src/copying/contiguous_split.cu +++ b/cpp/src/copying/contiguous_split.cu @@ -605,7 +605,6 @@ std::pair buf_info_functor::operator() std::pair setup_source_buf_info(InputIter begin, InputIter end, diff --git a/cpp/src/copying/copy.cu b/cpp/src/copying/copy.cu index ab276311dac..f12c9dcf006 100644 --- a/cpp/src/copying/copy.cu +++ b/cpp/src/copying/copy.cu @@ -241,15 +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; + (void)lhs; + (void)rhs; + (void)size; + (void)is_left; + (void)stream; + (void)mr; } - /** * @brief Functor called by the `type_dispatcher` to invoke copy_if_else on combinations * of column_view and scalar diff --git a/cpp/src/copying/scatter.cu b/cpp/src/copying/scatter.cu index 524cad1a187..3fccc2122cf 100644 --- a/cpp/src/copying/scatter.cu +++ b/cpp/src/copying/scatter.cu @@ -149,7 +149,6 @@ struct column_scalar_scatterer_impl { } }; - template struct column_scalar_scatterer_impl { template diff --git a/cpp/src/copying/segmented_shift.cu b/cpp/src/copying/segmented_shift.cu index 095e1b80588..dacc722396e 100644 --- a/cpp/src/copying/segmented_shift.cu +++ b/cpp/src/copying/segmented_shift.cu @@ -121,13 +121,11 @@ struct segmented_shift_functor { template std::unique_ptr operator()(Args&&...) { - if constexpr( std::is_same_v ) - { + if constexpr (std::is_same_v) { CUDF_FAIL("segmented_shift does not support list_view yet"); } - if constexpr( std::is_same_v ) - { + if constexpr (std::is_same_v) { CUDF_FAIL("segmented_shift does not support struct_view yet"); } From f1adb67e5ed5b84a8227df818557e3673df9befd Mon Sep 17 00:00:00 2001 From: Robert Maynard Date: Fri, 14 May 2021 08:42:58 -0400 Subject: [PATCH 7/7] Generalized CUDF_FAIL messages --- cpp/src/copying/segmented_shift.cu | 8 -------- 1 file changed, 8 deletions(-) diff --git a/cpp/src/copying/segmented_shift.cu b/cpp/src/copying/segmented_shift.cu index dacc722396e..6fc785a61c6 100644 --- a/cpp/src/copying/segmented_shift.cu +++ b/cpp/src/copying/segmented_shift.cu @@ -121,14 +121,6 @@ struct segmented_shift_functor { template std::unique_ptr operator()(Args&&...) { - if constexpr (std::is_same_v) { - CUDF_FAIL("segmented_shift does not support list_view yet"); - } - - if constexpr (std::is_same_v) { - CUDF_FAIL("segmented_shift does not support struct_view yet"); - } - CUDF_FAIL("Unsupported type for segmented_shift."); } };