From 0be84414502aabc383b72affb865951a67a3b42a Mon Sep 17 00:00:00 2001 From: Dave Baranec Date: Mon, 30 Nov 2020 17:38:49 -0600 Subject: [PATCH 1/6] Handle a corner case with nested columns containing a string column with no children. --- cpp/src/copying/contiguous_split.cu | 45 ++++++++++++++++------------- cpp/tests/copying/split_tests.cpp | 16 ++++++++++ 2 files changed, 41 insertions(+), 20 deletions(-) diff --git a/cpp/src/copying/contiguous_split.cu b/cpp/src/copying/contiguous_split.cu index 961c928825c..8e8ae62285a 100644 --- a/cpp/src/copying/contiguous_split.cu +++ b/cpp/src/copying/contiguous_split.cu @@ -419,38 +419,43 @@ std::pair buf_info_functor::operator() 0) { + strings_column_view scv(col); - // info for the offsets buffer - *current = src_buf_info(type_id::INT32, - scv.offsets().begin>(), - offset_stack_pos, - parent_offset_index, - false, - col.offset()); + auto offset_col = current; + + // info for the offsets buffer + *current = src_buf_info(type_id::INT32, + scv.offsets().begin>(), + offset_stack_pos, + parent_offset_index, + false, + col.offset()); - // prevent appending buf_info for non-exist chars buffer - if (scv.chars_size() > 0) { current++; offset_stack_pos += offset_depth; - // since we are crossing an offset boundary, our offset_depth and parent_offset_index go up. - offset_depth++; - parent_offset_index = offset_col - head; - - // info for the chars buffer - *current = src_buf_info( - type_id::INT8, nullptr, offset_stack_pos, parent_offset_index, false, col.offset()); + // prevent appending buf_info for non-exist chars buffer + if (scv.chars_size() > 0) { + // since we are crossing an offset boundary, our offset_depth and parent_offset_index go up. + offset_depth++; + parent_offset_index = offset_col - head; + + // info for the chars buffer + *current = src_buf_info( + type_id::INT8, nullptr, offset_stack_pos, parent_offset_index, false, col.offset()); + current++; + offset_stack_pos += offset_depth; + } } - return {current + 1, offset_stack_pos + offset_depth}; + return {current, offset_stack_pos}; } template <> diff --git a/cpp/tests/copying/split_tests.cpp b/cpp/tests/copying/split_tests.cpp index e965df6db81..e359105c744 100644 --- a/cpp/tests/copying/split_tests.cpp +++ b/cpp/tests/copying/split_tests.cpp @@ -16,6 +16,7 @@ #include #include +#include #include #include @@ -1276,6 +1277,21 @@ TEST_F(ContiguousSplitTableCornerCases, PreSplitTable) } } +TEST_F(ContiguousSplitTableCornerCases, NestedEmptyStrings) +{ + { + auto empty_string = cudf::strings::detail::make_empty_strings_column(); + auto offsets = cudf::test::fixed_width_column_wrapper({0, 1}); + auto list = cudf::make_lists_column( + 1, offsets.release(), std::move(empty_string), 0, rmm::device_buffer{0}); + + cudf::table_view src_table({static_cast(*list)}); + + std::vector splits({0}); + EXPECT_NO_THROW(contiguous_split(src_table, splits)); + } +} + struct ContiguousSplitNestedTypesTest : public cudf::test::BaseFixture { }; From c05b46558e5a131c051b3343d0878a4b6d2ad1f1 Mon Sep 17 00:00:00 2001 From: Dave Baranec Date: Mon, 30 Nov 2020 17:51:28 -0600 Subject: [PATCH 2/6] Changelog for 6864 --- CHANGELOG.md | 1 + 1 file changed, 1 insertion(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 44a5bbbc5cb..4db5eb365e4 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -170,6 +170,7 @@ - PR #6854 Fix the parameter order of writeParquetBufferBegin - PR #6855 Fix `.str.replace_with_backrefs` docs examples - PR #6853 Fix contiguous split of null string columns +- PR #6864 Handle contiguous_split corner case for nested string columns with no children # cuDF 0.16.0 (21 Oct 2020) From 4589ecc6e7a6039ec363f8acfb4724ef83ebe5bc Mon Sep 17 00:00:00 2001 From: Dave Baranec Date: Tue, 1 Dec 2020 17:11:19 -0600 Subject: [PATCH 3/6] Changed the logic of the buf info functor to match more directly with the logic of the other related functors to avoid mismatches. --- cpp/src/copying/contiguous_split.cu | 34 +++++++++++++++++++---------- 1 file changed, 22 insertions(+), 12 deletions(-) diff --git a/cpp/src/copying/contiguous_split.cu b/cpp/src/copying/contiguous_split.cu index 8e8ae62285a..c347289518e 100644 --- a/cpp/src/copying/contiguous_split.cu +++ b/cpp/src/copying/contiguous_split.cu @@ -390,10 +390,15 @@ struct buf_info_functor { } // info for the data buffer - *current = src_buf_info( - col.type().id(), nullptr, offset_stack_pos, parent_offset_index, false, col.offset()); + if(col.head()){ + *current = src_buf_info( + col.type().id(), nullptr, offset_stack_pos, parent_offset_index, false, col.offset()); - return {current + 1, offset_stack_pos + offset_depth}; + current++; + offset_stack_pos += offset_depth; + } + + return {current, offset_stack_pos}; } private: @@ -431,18 +436,23 @@ std::pair buf_info_functor::operator()>(), - offset_stack_pos, - parent_offset_index, - false, - col.offset()); + if(scv.offsets().head()){ + CUDF_EXPECTS(scv.offsets().nullable() == false, "Encountered nullable string offsets column"); + *current = src_buf_info(type_id::INT32, + scv.offsets().begin>(), + offset_stack_pos, + parent_offset_index, + false, + col.offset()); - current++; - offset_stack_pos += offset_depth; + current++; + offset_stack_pos += offset_depth; + } // prevent appending buf_info for non-exist chars buffer - if (scv.chars_size() > 0) { + if (scv.chars().head()){ + CUDF_EXPECTS(scv.chars().nullable() == false, "Encountered nullable string chars column"); + // since we are crossing an offset boundary, our offset_depth and parent_offset_index go up. offset_depth++; parent_offset_index = offset_col - head; From d40dcaae407548f2d22c1886ebe221f913e3a673 Mon Sep 17 00:00:00 2001 From: Dave Baranec Date: Wed, 2 Dec 2020 09:57:58 -0600 Subject: [PATCH 4/6] Formatting --- cpp/src/copying/contiguous_split.cu | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/cpp/src/copying/contiguous_split.cu b/cpp/src/copying/contiguous_split.cu index c347289518e..9750d4f7af0 100644 --- a/cpp/src/copying/contiguous_split.cu +++ b/cpp/src/copying/contiguous_split.cu @@ -390,7 +390,7 @@ struct buf_info_functor { } // info for the data buffer - if(col.head()){ + if (col.head()) { *current = src_buf_info( col.type().id(), nullptr, offset_stack_pos, parent_offset_index, false, col.offset()); @@ -436,7 +436,7 @@ std::pair buf_info_functor::operator()>(), @@ -450,7 +450,7 @@ std::pair buf_info_functor::operator() Date: Mon, 4 Jan 2021 15:14:14 -0600 Subject: [PATCH 5/6] Always add a src/dst buf record for columns to be copied, even if the data is nullptr/size is 0. Cleaner and avoids odd edge cases that can crop up with strangely formed columns (eg, ones that have size 0, but a non-null data pointer). --- cpp/src/copying/contiguous_split.cu | 155 ++++++++++++----------- cpp/tests/copying/split_tests.cpp | 183 +++++++++++++++++++++++++++- 2 files changed, 264 insertions(+), 74 deletions(-) diff --git a/cpp/src/copying/contiguous_split.cu b/cpp/src/copying/contiguous_split.cu index 2b1144a1562..ca15d7c97f5 100644 --- a/cpp/src/copying/contiguous_split.cu +++ b/cpp/src/copying/contiguous_split.cu @@ -273,7 +273,7 @@ template size_t compute_offset_stack_size(InputIter begin, InputIter end, int offset_depth = 0) { return std::accumulate(begin, end, 0, [offset_depth](auto stack_size, column_view const& col) { - auto const num_buffers = (col.head() != nullptr ? 1 : 0) + (col.nullable() ? 1 : 0); + auto const num_buffers = 1 + (col.nullable() ? 1 : 0); return stack_size + (offset_depth * num_buffers) + compute_offset_stack_size( col.child_begin(), col.child_end(), offset_depth + is_offset_type(col.type().id())); @@ -302,11 +302,11 @@ OutputIter setup_src_buf_data(InputIter begin, InputIter end, OutputIter out_buf out_buf++; } // NOTE: we're always returning the base pointer here. column-level offset is accounted - // for later. - if (col.head() != nullptr) { - *out_buf = const_cast(col.head()); - out_buf++; - } + // for later. Also, for some column types (string, list, struct) this pointer will be null + // because there is no associated data with the root column. + *out_buf = const_cast(col.head()); + out_buf++; + out_buf = setup_src_buf_data(col.child_begin(), col.child_end(), out_buf); }); return out_buf; @@ -331,8 +331,7 @@ template size_type count_src_bufs(InputIter begin, InputIter end) { auto buf_iter = thrust::make_transform_iterator(begin, [](column_view const& col) { - return (col.head() != nullptr ? 1 : 0) + (col.nullable() ? 1 : 0) + - count_src_bufs(col.child_begin(), col.child_end()); + return 1 + (col.nullable() ? 1 : 0) + count_src_bufs(col.child_begin(), col.child_end()); }); return std::accumulate(buf_iter, buf_iter + std::distance(begin, end), 0); } @@ -390,15 +389,10 @@ struct buf_info_functor { } // info for the data buffer - if (col.head()) { - *current = src_buf_info( - col.type().id(), nullptr, offset_stack_pos, parent_offset_index, false, col.offset()); - - current++; - offset_stack_pos += offset_depth; - } + *current = src_buf_info( + col.type().id(), nullptr, offset_stack_pos, parent_offset_index, false, col.offset()); - return {current, offset_stack_pos}; + return {current + 1, offset_stack_pos + offset_depth}; } private: @@ -429,40 +423,45 @@ std::pair buf_info_functor::operator() 0) { + CUDF_EXPECTS(col.num_children() == 2, "Encountered malformed string column"); strings_column_view scv(col); - auto offset_col = current; - // info for the offsets buffer - if (scv.offsets().head()) { - CUDF_EXPECTS(scv.offsets().nullable() == false, "Encountered nullable string offsets column"); - *current = src_buf_info(type_id::INT32, - scv.offsets().begin>(), - offset_stack_pos, - parent_offset_index, - false, - col.offset()); - - current++; - offset_stack_pos += offset_depth; - } - - // prevent appending buf_info for non-exist chars buffer - if (scv.chars().head()) { - CUDF_EXPECTS(scv.chars().nullable() == false, "Encountered nullable string chars column"); - - // since we are crossing an offset boundary, our offset_depth and parent_offset_index go up. - offset_depth++; - parent_offset_index = offset_col - head; - - // info for the chars buffer - *current = src_buf_info( - type_id::INT8, nullptr, offset_stack_pos, parent_offset_index, false, col.offset()); - current++; - offset_stack_pos += offset_depth; - } + auto offset_col = current; + CUDF_EXPECTS(scv.offsets().nullable() == false, "Encountered nullable string offsets column"); + *current = src_buf_info(type_id::INT32, + // note: offsets can be null in the case where the string column + // has been created with empty_like(). + scv.offsets().begin>(), + offset_stack_pos, + parent_offset_index, + false, + col.offset()); + + current++; + offset_stack_pos += offset_depth; + + // since we are crossing an offset boundary, calculate our new depth and parent offset index. + offset_depth++; + parent_offset_index = offset_col - head; + + // prevent appending buf_info for non-existent chars buffer + CUDF_EXPECTS(scv.chars().nullable() == false, "Encountered nullable string chars column"); + + // info for the chars buffer + *current = src_buf_info( + type_id::INT8, nullptr, offset_stack_pos, parent_offset_index, false, col.offset()); + current++; + offset_stack_pos += offset_depth; } return {current, offset_stack_pos}; @@ -483,10 +482,20 @@ std::pair buf_info_functor::operator()>(), offset_stack_pos, parent_offset_index, @@ -495,7 +504,7 @@ std::pair buf_info_functor::operator() buf_info_functor::operator()(nullptr), 0); }(); - uint8_t const* data_ptr; - size_type size; - std::tie(data_ptr, size) = [&]() { - if (src.head()) { - auto const ptr = base_ptr + current_info->dst_offset; - // if we have data, num_elements will always be the correct size. - // we don't want to use num_rows because if we are an offset column, num_rows - // represents the # of rows of our owning parent. num_elements always represents - // the proper size for this column - auto const size = current_info->num_elements; - ++current_info; - return std::make_pair(ptr, size); - } - // Parent columns w/o data (e.g., strings, lists) don't have an associated `dst_buf_info`, - // therefore, use the first child's info if it has at least one child. Their num_rows value - // will be correct (also see comment above) - auto const size = (src.num_children() == 0) ? 0 : current_info->num_rows; - return std::make_pair(static_cast(nullptr), size); - }(); + + // size/data pointer for the column + auto const size = current_info->num_elements; + uint8_t const* data_ptr = + size == 0 || src.head() == nullptr ? nullptr : base_ptr + current_info->dst_offset; + ++current_info; + + // children auto children = std::vector{}; children.reserve(src.num_children()); current_info = build_output_columns( src.child_begin(), src.child_end(), current_info, std::back_inserter(children), base_ptr); + return column_view{src.type(), size, data_ptr, bitmask_ptr, null_count, 0, std::move(children)}; }); @@ -837,8 +843,13 @@ std::vector contiguous_split(cudf::table_view const& in int row_end = d_indices[split_index + 1] + src_info.column_offset; while (stack_size > 0) { stack_size--; - row_start = d_src_buf_info[offset_stack[stack_size]].offsets[row_start]; - row_end = d_src_buf_info[offset_stack[stack_size]].offsets[row_end]; + auto const offsets = d_src_buf_info[offset_stack[stack_size]].offsets; + // this case can happen when you have empty string or list columns constructed with + // empty_like() + if (offsets != nullptr) { + row_start = offsets[row_start]; + row_end = offsets[row_end]; + } } // final row indices and row count @@ -850,7 +861,7 @@ std::vector contiguous_split(cudf::table_view const& in int const bit_shift = src_info.is_validity ? row_start % 32 : 0; // # of rows isn't necessarily the same as # of elements to be copied. auto const num_elements = [&]() { - if (src_info.offsets != nullptr) { + if (src_info.offsets != nullptr && num_rows > 0) { return num_rows + 1; } else if (src_info.is_validity) { return (num_rows + 31) / 32; @@ -987,4 +998,4 @@ std::vector contiguous_split(cudf::table_view const& in return cudf::detail::contiguous_split(input, splits, rmm::cuda_stream_default, mr); } -}; // namespace cudf +}; // namespace cudf \ No newline at end of file diff --git a/cpp/tests/copying/split_tests.cpp b/cpp/tests/copying/split_tests.cpp index e359105c744..3c95e5e5da1 100644 --- a/cpp/tests/copying/split_tests.cpp +++ b/cpp/tests/copying/split_tests.cpp @@ -900,6 +900,81 @@ void split_structs(bool include_validity, SplitFunc Split, CompareFunc Compare) } } +template +void split_structs_no_children(SplitFunc Split, CompareFunc Compare) +{ + using namespace cudf::test; + + // no nulls + { + auto struct_column = cudf::make_structs_column(4, {}, 0, rmm::device_buffer{}); + auto expected = cudf::make_structs_column(2, {}, 0, rmm::device_buffer{}); + + // split + std::vector splits{2}; + auto result = Split(*struct_column, splits); + + EXPECT_EQ(result.size(), 2); + Compare(*expected, result[0]); + Compare(*expected, result[1]); + } + + // all nulls + { + std::vector struct_validity{false, false, false, false}; + auto struct_column = cudf::make_structs_column( + 4, {}, 4, detail::make_null_mask(struct_validity.begin(), struct_validity.end())); + + std::vector expected_validity{false, false}; + auto expected = cudf::make_structs_column( + 2, {}, 2, detail::make_null_mask(expected_validity.begin(), expected_validity.end())); + + // split + std::vector splits{2}; + auto result = Split(*struct_column, splits); + + EXPECT_EQ(result.size(), 2); + Compare(*expected, result[0]); + Compare(*expected, result[1]); + } + + // no nulls, empty output column + { + auto struct_column = cudf::make_structs_column(4, {}, 0, rmm::device_buffer{}); + auto expected0 = cudf::make_structs_column(4, {}, 0, rmm::device_buffer{}); + auto expected1 = cudf::make_structs_column(0, {}, 0, rmm::device_buffer{}); + + // split + std::vector splits{4}; + auto result = Split(*struct_column, splits); + + EXPECT_EQ(result.size(), 2); + Compare(*expected0, result[0]); + Compare(*expected1, result[1]); + } + + // all nulls, empty output column + { + std::vector struct_validity{false, false, false, false}; + auto struct_column = cudf::make_structs_column( + 4, {}, 4, detail::make_null_mask(struct_validity.begin(), struct_validity.end())); + + std::vector expected_validity0{false, false, false, false}; + auto expected0 = cudf::make_structs_column( + 4, {}, 4, detail::make_null_mask(expected_validity0.begin(), expected_validity0.end())); + + auto expected1 = cudf::make_structs_column(0, {}, 0, rmm::device_buffer{}); + + // split + std::vector splits{4}; + auto result = Split(*struct_column, splits); + + EXPECT_EQ(result.size(), 2); + Compare(*expected0, result[0]); + Compare(*expected1, result[1]); + } +} + template void split_nested_struct_of_list(SplitFunc Split, CompareFunc Compare) { @@ -1013,6 +1088,17 @@ TEST_F(SplitNestedTypesTest, StructsWithNulls) }); } +TEST_F(SplitNestedTypesTest, StructsNoChildren) +{ + split_structs_no_children( + [](cudf::column_view const& t, std::vector const& splits) { + return cudf::split(t, splits); + }, + [](cudf::column_view const& expected, cudf::column_view const& result) { + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(expected, result); + }); +} + TEST_F(SplitNestedTypesTest, StructsOfList) { split_nested_struct_of_list( @@ -1277,11 +1363,13 @@ TEST_F(ContiguousSplitTableCornerCases, PreSplitTable) } } -TEST_F(ContiguousSplitTableCornerCases, NestedEmptyStrings) +TEST_F(ContiguousSplitTableCornerCases, NestedEmpty) { + // this produces an empty strings column with no children, + // nested inside a list { auto empty_string = cudf::strings::detail::make_empty_strings_column(); - auto offsets = cudf::test::fixed_width_column_wrapper({0, 1}); + auto offsets = cudf::test::fixed_width_column_wrapper({0, 0}); auto list = cudf::make_lists_column( 1, offsets.release(), std::move(empty_string), 0, rmm::device_buffer{0}); @@ -1289,9 +1377,88 @@ TEST_F(ContiguousSplitTableCornerCases, NestedEmptyStrings) std::vector splits({0}); EXPECT_NO_THROW(contiguous_split(src_table, splits)); + + std::vector splits2({1}); + EXPECT_NO_THROW(contiguous_split(src_table, splits2)); + } + + // this produces an empty strings column with children that have no data, + // nested inside a list + { + cudf::test::strings_column_wrapper str{"abc"}; + auto empty_string = cudf::empty_like(str); + auto offsets = cudf::test::fixed_width_column_wrapper({0, 0}); + auto list = cudf::make_lists_column( + 1, offsets.release(), std::move(empty_string), 0, rmm::device_buffer{0}); + + cudf::table_view src_table({static_cast(*list)}); + + std::vector splits({0}); + EXPECT_NO_THROW(contiguous_split(src_table, splits)); + + std::vector splits2({1}); + EXPECT_NO_THROW(contiguous_split(src_table, splits2)); + } + + // this produces an empty lists column with children that have no data, + // nested inside a list + { + cudf::test::lists_column_wrapper listw{{1.0f, 2.0f}, {3.0f, 4.0f}}; + auto empty_list = cudf::empty_like(listw); + auto offsets = cudf::test::fixed_width_column_wrapper({0, 0}); + auto list = cudf::make_lists_column( + 1, offsets.release(), std::move(empty_list), 0, rmm::device_buffer{0}); + + cudf::table_view src_table({static_cast(*list)}); + + std::vector splits({0}); + EXPECT_NO_THROW(contiguous_split(src_table, splits)); + + std::vector splits2({1}); + EXPECT_NO_THROW(contiguous_split(src_table, splits2)); + } + + // this produces an empty lists column with children that have no data, + // nested inside a list + { + cudf::test::lists_column_wrapper listw{{1.0f, 2.0f}, {3.0f, 4.0f}}; + auto empty_list = cudf::empty_like(listw); + auto offsets = cudf::test::fixed_width_column_wrapper({0, 0}); + auto list = cudf::make_lists_column( + 1, offsets.release(), std::move(empty_list), 0, rmm::device_buffer{0}); + + cudf::table_view src_table({static_cast(*list)}); + + std::vector splits({0}); + EXPECT_NO_THROW(contiguous_split(src_table, splits)); + + std::vector splits2({1}); + EXPECT_NO_THROW(contiguous_split(src_table, splits2)); + } + + // this produces an empty struct column with children that have no data, + // nested inside a list + { + cudf::test::fixed_width_column_wrapper ints{0, 1, 2, 3, 4}; + cudf::test::fixed_width_column_wrapper floats{4, 3, 2, 1, 0}; + auto struct_column = cudf::test::structs_column_wrapper({ints, floats}); + auto empty_struct = cudf::empty_like(struct_column); + auto offsets = cudf::test::fixed_width_column_wrapper({0, 0}); + auto list = cudf::make_lists_column( + 1, offsets.release(), std::move(empty_struct), 0, rmm::device_buffer{0}); + + cudf::table_view src_table({static_cast(*list)}); + + std::vector splits({0}); + EXPECT_NO_THROW(contiguous_split(src_table, splits)); + + std::vector splits2({1}); + EXPECT_NO_THROW(contiguous_split(src_table, splits2)); } } +TEST_F(ContiguousSplitTableCornerCases, MalformedColumns) {} + struct ContiguousSplitNestedTypesTest : public cudf::test::BaseFixture { }; @@ -1345,6 +1512,18 @@ TEST_F(ContiguousSplitNestedTypesTest, StructsWithNulls) }); } +TEST_F(ContiguousSplitNestedTypesTest, StructsNoChildren) +{ + split_structs_no_children( + [](cudf::column_view const& c, std::vector const& splits) { + cudf::table_view t({c}); + return cudf::contiguous_split(t, splits); + }, + [](cudf::column_view const& expected, cudf::contiguous_split_result const& result) { + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(expected, result.table.column(0)); + }); +} + TEST_F(ContiguousSplitNestedTypesTest, StructsOfList) { split_nested_struct_of_list( From 849f4ff37a0e7063f0bcb94b024b176540226f55 Mon Sep 17 00:00:00 2001 From: Dave Baranec Date: Tue, 12 Jan 2021 09:56:09 -0600 Subject: [PATCH 6/6] Minor reformatting. --- CHANGELOG.md | 1 - cpp/src/copying/contiguous_split.cu | 2 +- 2 files changed, 1 insertion(+), 2 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index d1573c49e09..a3c84ba1b72 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -228,7 +228,6 @@ - PR #6853 Fix contiguous split of null string columns - PR #6861 Fix compile error in type_dispatch_benchmark.cu - PR #6864 Handle contiguous_split corner case for nested string columns with no children - - PR #6869 Avoid dependency resolution failure in latest version of pip by explicitly specifying versions for dask and distributed - PR #6806 Force install of local conda artifacts - PR #6887 Fix typo and `0-d` numpy array handling in binary operation diff --git a/cpp/src/copying/contiguous_split.cu b/cpp/src/copying/contiguous_split.cu index 089de23a060..10e2dfbdaeb 100644 --- a/cpp/src/copying/contiguous_split.cu +++ b/cpp/src/copying/contiguous_split.cu @@ -995,4 +995,4 @@ std::vector contiguous_split(cudf::table_view const& in return cudf::detail::contiguous_split(input, splits, rmm::cuda_stream_default, mr); } -}; // namespace cudf \ No newline at end of file +}; // namespace cudf