From 88be4c337af4a73a1df2f5e303f4a2c71c1f1e9a Mon Sep 17 00:00:00 2001 From: Karthikeyan Natarajan Date: Wed, 28 Jul 2021 09:59:34 -0700 Subject: [PATCH 1/8] fix SCALAR_TEST memory errors --- cpp/src/scalar/scalar.cpp | 9 +++++---- cpp/src/structs/utilities.cpp | 3 +-- 2 files changed, 6 insertions(+), 6 deletions(-) diff --git a/cpp/src/scalar/scalar.cpp b/cpp/src/scalar/scalar.cpp index 045bfbe0327..3d208928f8d 100644 --- a/cpp/src/scalar/scalar.cpp +++ b/cpp/src/scalar/scalar.cpp @@ -17,7 +17,7 @@ #include #include -#include +#include #include #include #include @@ -567,12 +567,13 @@ void struct_scalar::superimpose_nulls(rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { // push validity mask down - std::vector host_validity({0}); - auto validity = cudf::detail::make_device_uvector_sync(host_validity, stream, mr); + std::vector host_validity( + cudf::bitmask_allocation_size_bytes(1) / sizeof(bitmask_type), 0); + auto validity = cudf::detail::create_null_mask(1, mask_state::ALL_NULL, stream); auto iter = thrust::make_counting_iterator(0); std::for_each(iter, iter + _data.num_columns(), [&](size_type i) { cudf::structs::detail::superimpose_parent_nulls( - validity.data(), 1, _data.get_column(i), stream, mr); + static_cast(validity.data()), 1, _data.get_column(i), stream, mr); }); } diff --git a/cpp/src/structs/utilities.cpp b/cpp/src/structs/utilities.cpp index 80bea2ab55e..aa32c555324 100644 --- a/cpp/src/structs/utilities.cpp +++ b/cpp/src/structs/utilities.cpp @@ -187,8 +187,7 @@ void superimpose_parent_nulls(bitmask_type const* parent_null_mask, { if (!child.nullable()) { // Child currently has no null mask. Copy parent's null mask. - child.set_null_mask(rmm::device_buffer{ - parent_null_mask, cudf::bitmask_allocation_size_bytes(child.size()), stream, mr}); + child.set_null_mask(cudf::detail::copy_bitmask(parent_null_mask, 0, child.size(), stream, mr)); child.set_null_count(parent_null_count); } else { // Child should have a null mask. From 2b4987b23f0dd218c46b6f9124fe76d827b18fa1 Mon Sep 17 00:00:00 2001 From: Karthikeyan Natarajan Date: Wed, 28 Jul 2021 10:00:09 -0700 Subject: [PATCH 2/8] fix MERGE_TEST memory errors --- cpp/src/dictionary/detail/merge.cu | 8 ++++++-- 1 file changed, 6 insertions(+), 2 deletions(-) diff --git a/cpp/src/dictionary/detail/merge.cu b/cpp/src/dictionary/detail/merge.cu index 2ff0a3e0a2a..e972403cad3 100644 --- a/cpp/src/dictionary/detail/merge.cu +++ b/cpp/src/dictionary/detail/merge.cu @@ -17,6 +17,7 @@ #include #include #include +#include #include #include #include @@ -62,8 +63,11 @@ std::unique_ptr merge(dictionary_column_view const& lcol, return make_dictionary_column( std::make_unique(lcol.keys(), stream, mr), std::move(indices_column), - rmm::device_buffer{ - lcol.has_nulls() || rcol.has_nulls() ? static_cast(merged_size) : 0, stream, mr}, + cudf::detail::create_null_mask( + lcol.has_nulls() || rcol.has_nulls() ? static_cast(merged_size) : 0, + mask_state::UNINITIALIZED, + stream, + mr), lcol.null_count() + rcol.null_count()); } From 7de2dc811f956461cfe68f1d518d46a841b099c5 Mon Sep 17 00:00:00 2001 From: Karthikeyan Natarajan Date: Wed, 28 Jul 2021 10:00:36 -0700 Subject: [PATCH 3/8] fix part of COPYING_TEST memory errors --- cpp/tests/copying/concatenate_tests.cu | 10 ++++------ 1 file changed, 4 insertions(+), 6 deletions(-) diff --git a/cpp/tests/copying/concatenate_tests.cu b/cpp/tests/copying/concatenate_tests.cu index 7d3b7beb2cb..c513716859e 100644 --- a/cpp/tests/copying/concatenate_tests.cu +++ b/cpp/tests/copying/concatenate_tests.cu @@ -48,8 +48,6 @@ using Table = cudf::table; template struct TypedColumnTest : public cudf::test::BaseFixture { - static std::size_t data_size() { return 1000; } - static std::size_t mask_size() { return 100; } cudf::data_type type() { return cudf::data_type{cudf::type_to_id()}; } TypedColumnTest(rmm::cuda_stream_view stream = rmm::cuda_stream_default) @@ -58,14 +56,14 @@ struct TypedColumnTest : public cudf::test::BaseFixture { { auto typed_data = static_cast(data.data()); auto typed_mask = static_cast(mask.data()); - std::vector h_data(data_size()); + std::vector h_data(data.size()); std::iota(h_data.begin(), h_data.end(), char{0}); - std::vector h_mask(mask_size()); + std::vector h_mask(mask.size()); std::iota(h_mask.begin(), h_mask.end(), char{0}); CUDA_TRY(cudaMemcpyAsync( - typed_data, h_data.data(), data_size(), cudaMemcpyHostToDevice, stream.value())); + typed_data, h_data.data(), data.size(), cudaMemcpyHostToDevice, stream.value())); CUDA_TRY(cudaMemcpyAsync( - typed_mask, h_mask.data(), mask_size(), cudaMemcpyHostToDevice, stream.value())); + typed_mask, h_mask.data(), mask.size(), cudaMemcpyHostToDevice, stream.value())); stream.synchronize(); } From 7840f7568456ea4d2c12d062b23101d54e9fded0 Mon Sep 17 00:00:00 2001 From: Karthikeyan Natarajan Date: Sun, 1 Aug 2021 21:44:57 +0530 Subject: [PATCH 4/8] binary ops ref fix --- cpp/src/binaryop/compiled/binary_ops.cu | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/cpp/src/binaryop/compiled/binary_ops.cu b/cpp/src/binaryop/compiled/binary_ops.cu index 0576606a58c..2b38224864a 100644 --- a/cpp/src/binaryop/compiled/binary_ops.cu +++ b/cpp/src/binaryop/compiled/binary_ops.cu @@ -45,7 +45,7 @@ struct scalar_as_column_device_view { rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { - auto h_scalar_type_view = static_cast&>(const_cast(s)); + auto& h_scalar_type_view = static_cast&>(const_cast(s)); auto col_v = column_view(s.type(), 1, h_scalar_type_view.data(), (bitmask_type const*)s.validity_data()); return std::pair{column_device_view::create(col_v, stream), std::unique_ptr(nullptr)}; @@ -63,8 +63,8 @@ scalar_as_column_device_view::operator()(scalar const& s, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { - using T = cudf::string_view; - auto h_scalar_type_view = static_cast&>(const_cast(s)); + using T = cudf::string_view; + auto& h_scalar_type_view = static_cast&>(const_cast(s)); // build offsets column from the string size auto offsets_transformer_itr = From 09eebd37ac21f5bb82a41095be1f9720f48813b7 Mon Sep 17 00:00:00 2001 From: Karthikeyan Natarajan Date: Sun, 1 Aug 2021 21:45:30 +0530 Subject: [PATCH 5/8] zero size list disable scatter_if --- cpp/tests/utilities/column_utilities.cu | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/cpp/tests/utilities/column_utilities.cu b/cpp/tests/utilities/column_utilities.cu index 88e9e3d1384..0d38e0006e1 100644 --- a/cpp/tests/utilities/column_utilities.cu +++ b/cpp/tests/utilities/column_utilities.cu @@ -118,9 +118,12 @@ std::unique_ptr generate_child_row_indices(lists_column_view const& c, 0, [row_indices = row_indices.begin(), validity = c.null_mask(), + offsets = c.offsets().begin(), offset = c.offset()] __device__(int index) { auto const true_index = row_indices[index] + offset; - return !validity || cudf::bit_is_set(validity, true_index) ? 1 : 0; + return !validity || cudf::bit_is_set(validity, true_index) + ? (offsets[true_index + 1] - offsets[true_index]) != 0 + : 0; }); auto output_row_iter = cudf::detail::make_counting_transform_iterator( 0, From 3099fef1e1b086723e1183bd254365a4118da655 Mon Sep 17 00:00:00 2001 From: Karthikeyan Natarajan Date: Sun, 1 Aug 2021 21:47:00 +0530 Subject: [PATCH 6/8] char column size fix in unit test --- cpp/tests/copying/concatenate_tests.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/tests/copying/concatenate_tests.cu b/cpp/tests/copying/concatenate_tests.cu index c513716859e..5496fa9d72e 100644 --- a/cpp/tests/copying/concatenate_tests.cu +++ b/cpp/tests/copying/concatenate_tests.cu @@ -482,7 +482,7 @@ TEST_F(OverflowTest, Presliced) auto offset_gen = cudf::detail::make_counting_transform_iterator( 0, [string_size](size_type index) { return index * string_size; }); cudf::test::fixed_width_column_wrapper offsets(offset_gen, offset_gen + num_rows + 1); - auto many_chars = cudf::make_fixed_width_column(data_type{type_id::INT8}, num_rows); + auto many_chars = cudf::make_fixed_width_column(data_type{type_id::INT8}, total_chars_size); auto col = cudf::make_strings_column( num_rows, offsets.release(), std::move(many_chars), 0, rmm::device_buffer{}); From ff918854d18bcdc165c2b5d5ee4cbb2af7414972 Mon Sep 17 00:00:00 2001 From: Karthikeyan Natarajan Date: Tue, 3 Aug 2021 13:06:34 +0530 Subject: [PATCH 7/8] addressed review comments --- cpp/tests/copying/concatenate_tests.cu | 2 +- cpp/tests/utilities/column_utilities.cu | 10 +++++----- 2 files changed, 6 insertions(+), 6 deletions(-) diff --git a/cpp/tests/copying/concatenate_tests.cu b/cpp/tests/copying/concatenate_tests.cu index 5496fa9d72e..f78ab14d6b0 100644 --- a/cpp/tests/copying/concatenate_tests.cu +++ b/cpp/tests/copying/concatenate_tests.cu @@ -513,7 +513,7 @@ TEST_F(OverflowTest, Presliced) offsets->view().begin(), offsets->view().end(), offsets->mutable_view().begin()); - auto many_chars = cudf::make_fixed_width_column(data_type{type_id::INT8}, num_rows); + auto many_chars = cudf::make_fixed_width_column(data_type{type_id::INT8}, total_chars_size); auto col = cudf::make_strings_column( num_rows, std::move(offsets), std::move(many_chars), 0, rmm::device_buffer{}); diff --git a/cpp/tests/utilities/column_utilities.cu b/cpp/tests/utilities/column_utilities.cu index 0d38e0006e1..41f9c516fb2 100644 --- a/cpp/tests/utilities/column_utilities.cu +++ b/cpp/tests/utilities/column_utilities.cu @@ -114,7 +114,7 @@ std::unique_ptr generate_child_row_indices(lists_column_view const& c, // // result = [6, 1, 11, 1, 1] // - auto validity_iter = cudf::detail::make_counting_transform_iterator( + auto is_output_row = cudf::detail::make_counting_transform_iterator( 0, [row_indices = row_indices.begin(), validity = c.null_mask(), @@ -139,7 +139,7 @@ std::unique_ptr generate_child_row_indices(lists_column_view const& c, output_row_iter, output_row_iter + row_indices.size(), output_row_start->view().begin(), - validity_iter, + is_output_row, result->mutable_view().begin()); // generate keys for each output row @@ -153,10 +153,10 @@ std::unique_ptr generate_child_row_indices(lists_column_view const& c, keys->mutable_view().end(), [] __device__() { return 0; }); thrust::scatter_if(rmm::exec_policy(), - validity_iter, - validity_iter + row_indices.size(), + is_output_row, + is_output_row + row_indices.size(), output_row_start->view().begin(), - validity_iter, + is_output_row, keys->mutable_view().begin()); thrust::inclusive_scan(rmm::exec_policy(), keys->view().begin(), From 45d213b4a08475660df06f6cfd02149bb468cd0a Mon Sep 17 00:00:00 2001 From: Karthikeyan Natarajan Date: Fri, 6 Aug 2021 12:49:33 +0530 Subject: [PATCH 8/8] reuse row_size_iter --- cpp/tests/utilities/column_utilities.cu | 25 ++++++++----------------- 1 file changed, 8 insertions(+), 17 deletions(-) diff --git a/cpp/tests/utilities/column_utilities.cu b/cpp/tests/utilities/column_utilities.cu index 41f9c516fb2..f3002bc4b1a 100644 --- a/cpp/tests/utilities/column_utilities.cu +++ b/cpp/tests/utilities/column_utilities.cu @@ -114,17 +114,6 @@ std::unique_ptr generate_child_row_indices(lists_column_view const& c, // // result = [6, 1, 11, 1, 1] // - auto is_output_row = cudf::detail::make_counting_transform_iterator( - 0, - [row_indices = row_indices.begin(), - validity = c.null_mask(), - offsets = c.offsets().begin(), - offset = c.offset()] __device__(int index) { - auto const true_index = row_indices[index] + offset; - return !validity || cudf::bit_is_set(validity, true_index) - ? (offsets[true_index + 1] - offsets[true_index]) != 0 - : 0; - }); auto output_row_iter = cudf::detail::make_counting_transform_iterator( 0, [row_indices = row_indices.begin(), @@ -139,8 +128,9 @@ std::unique_ptr generate_child_row_indices(lists_column_view const& c, output_row_iter, output_row_iter + row_indices.size(), output_row_start->view().begin(), - is_output_row, - result->mutable_view().begin()); + row_size_iter, + result->mutable_view().begin(), + [] __device__(auto row_size) { return row_size != 0; }); // generate keys for each output row // @@ -153,11 +143,12 @@ std::unique_ptr generate_child_row_indices(lists_column_view const& c, keys->mutable_view().end(), [] __device__() { return 0; }); thrust::scatter_if(rmm::exec_policy(), - is_output_row, - is_output_row + row_indices.size(), + row_size_iter, + row_size_iter + row_indices.size(), output_row_start->view().begin(), - is_output_row, - keys->mutable_view().begin()); + row_size_iter, + keys->mutable_view().begin(), + [] __device__(auto row_size) { return row_size != 0; }); thrust::inclusive_scan(rmm::exec_policy(), keys->view().begin(), keys->view().end(),