diff --git a/cpp/include/cudf/column/column_view.hpp b/cpp/include/cudf/column/column_view.hpp index 6db5c8b3c7b..43ce78d6c67 100644 --- a/cpp/include/cudf/column/column_view.hpp +++ b/cpp/include/cudf/column/column_view.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2024, NVIDIA CORPORATION. + * Copyright (c) 2019-2025, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -176,9 +176,13 @@ class column_view_base { * * @param[in] begin The starting index of the range (inclusive). * @param[in] end The index of the last element in the range (exclusive). + * @param[in] stream CUDA stream used for device memory operations and kernel launches * @return The count of null elements in the given range */ - [[nodiscard]] size_type null_count(size_type begin, size_type end) const; + [[nodiscard]] size_type null_count( + size_type begin, + size_type end, + rmm::cuda_stream_view stream = cudf::get_default_stream()) const; /** * @brief Indicates if the column contains null elements, @@ -198,12 +202,15 @@ class column_view_base { * * @param begin The starting index of the range (inclusive). * @param end The index of the last element in the range (exclusive). + * @param stream CUDA stream used for device memory operations and kernel launches * @return true One or more elements are null in the range [begin, end) * @return false All elements are valid in the range [begin, end) */ - [[nodiscard]] bool has_nulls(size_type begin, size_type end) const + [[nodiscard]] bool has_nulls(size_type begin, + size_type end, + rmm::cuda_stream_view stream = cudf::get_default_stream()) const { - return null_count(begin, end) > 0; + return null_count(begin, end, stream) > 0; } /** diff --git a/cpp/src/column/column_view.cpp b/cpp/src/column/column_view.cpp index ea940676f6a..a7718d19a94 100644 --- a/cpp/src/column/column_view.cpp +++ b/cpp/src/column/column_view.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2024, NVIDIA CORPORATION. + * Copyright (c) 2019-2025, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -141,13 +141,14 @@ column_view_base::column_view_base(data_type type, } } -size_type column_view_base::null_count(size_type begin, size_type end) const +size_type column_view_base::null_count(size_type begin, + size_type end, + rmm::cuda_stream_view stream) const { CUDF_EXPECTS((begin >= 0) && (end <= size()) && (begin <= end), "Range is out of bounds."); return (null_count() == 0) ? 0 - : cudf::detail::null_count( - null_mask(), offset() + begin, offset() + end, cudf::get_default_stream()); + : cudf::detail::null_count(null_mask(), offset() + begin, offset() + end, stream); } bool is_shallow_equivalent(column_view const& lhs, column_view const& rhs) diff --git a/cpp/src/copying/copy_range.cu b/cpp/src/copying/copy_range.cu index bffb48a8ec0..06b5ff6be5a 100644 --- a/cpp/src/copying/copy_range.cu +++ b/cpp/src/copying/copy_range.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2024, NVIDIA CORPORATION. + * Copyright (c) 2019-2025, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -103,7 +103,7 @@ struct out_of_place_copy_range_dispatch { rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()) { auto p_ret = std::make_unique(target, stream, mr); - if ((!p_ret->nullable()) && source.has_nulls(source_begin, source_end)) { + if ((!p_ret->nullable()) && source.has_nulls(source_begin, source_end, stream)) { p_ret->set_null_mask( cudf::detail::create_null_mask(p_ret->size(), cudf::mask_state::ALL_VALID, stream, mr), 0); } diff --git a/cpp/src/rolling/grouped_rolling.cu b/cpp/src/rolling/grouped_rolling.cu index 3cf292f5abb..66d46a51577 100644 --- a/cpp/src/rolling/grouped_rolling.cu +++ b/cpp/src/rolling/grouped_rolling.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2024, NVIDIA CORPORATION. + * Copyright (c) 2020-2025, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -419,7 +419,7 @@ template get_null_bounds_for_orderby_column( - column_view const& orderby_column) + column_view const& orderby_column, rmm::cuda_stream_view stream) { auto const num_rows = orderby_column.size(); auto const num_nulls = orderby_column.null_count(); @@ -429,7 +429,7 @@ std::tuple get_null_bounds_for_orderby_column( return std::make_tuple(0, num_nulls); } - auto const first_row_is_null = orderby_column.null_count(0, 1) == 1; + auto const first_row_is_null = orderby_column.null_count(0, 1, stream) == 1; return first_row_is_null ? std::make_tuple(0, num_nulls) : std::make_tuple(num_rows - num_nulls, num_rows); @@ -451,7 +451,8 @@ std::unique_ptr range_window_ASC(column_view const& input, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) { - auto [h_nulls_begin_idx, h_nulls_end_idx] = get_null_bounds_for_orderby_column(orderby_column); + auto [h_nulls_begin_idx, h_nulls_end_idx] = + get_null_bounds_for_orderby_column(orderby_column, stream); auto const p_orderby_device_view = cudf::column_device_view::create(orderby_column, stream); auto const preceding_calculator = cuda::proclaim_return_type( @@ -740,7 +741,8 @@ std::unique_ptr range_window_DESC(column_view const& input, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) { - auto [h_nulls_begin_idx, h_nulls_end_idx] = get_null_bounds_for_orderby_column(orderby_column); + auto [h_nulls_begin_idx, h_nulls_end_idx] = + get_null_bounds_for_orderby_column(orderby_column, stream); auto const p_orderby_device_view = cudf::column_device_view::create(orderby_column, stream); auto const preceding_calculator = cuda::proclaim_return_type( diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 6a89b1e48d6..4451f6b64c5 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -698,6 +698,7 @@ if(CUDF_BUILD_STREAMS_TEST_UTIL) endif() ConfigureTest(STREAM_BINARYOP_TEST streams/binaryop_test.cpp STREAM_MODE testing) +ConfigureTest(STREAM_COLUMN_VIEW_TEST streams/column_view_test.cpp STREAM_MODE testing) ConfigureTest(STREAM_CONCATENATE_TEST streams/concatenate_test.cpp STREAM_MODE testing) ConfigureTest(STREAM_COPYING_TEST streams/copying_test.cpp STREAM_MODE testing) ConfigureTest(STREAM_CSVIO_TEST streams/io/csv_test.cpp STREAM_MODE testing) diff --git a/cpp/tests/streams/column_view_test.cpp b/cpp/tests/streams/column_view_test.cpp new file mode 100644 index 00000000000..c7483223973 --- /dev/null +++ b/cpp/tests/streams/column_view_test.cpp @@ -0,0 +1,103 @@ +/* + * Copyright (c) 2025, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include +#include +#include +#include + +#include +#include +#include + +#include +#include + +template +struct TypedColumnTest : public cudf::test::BaseFixture { + cudf::data_type type() { return cudf::data_type{cudf::type_to_id()}; } + + TypedColumnTest(rmm::cuda_stream_view stream = cudf::test::get_default_stream()) + : data{_num_elements * cudf::size_of(type()), stream}, + mask{cudf::bitmask_allocation_size_bytes(_num_elements), stream} + { + std::vector h_data(std::max(data.size(), mask.size())); + std::iota(h_data.begin(), h_data.end(), 0); + CUDF_CUDA_TRY( + cudaMemcpyAsync(data.data(), h_data.data(), data.size(), cudaMemcpyDefault, stream.value())); + CUDF_CUDA_TRY( + cudaMemcpyAsync(mask.data(), h_data.data(), mask.size(), cudaMemcpyDefault, stream.value())); + } + + cudf::size_type num_elements() { return _num_elements; } + + std::random_device r; + std::default_random_engine generator{r()}; + std::uniform_int_distribution distribution{200, 1000}; + cudf::size_type _num_elements{distribution(generator)}; + rmm::device_buffer data{}; + rmm::device_buffer mask{}; + rmm::device_buffer all_valid_mask{create_null_mask( + num_elements(), cudf::mask_state::ALL_VALID, cudf::test::get_default_stream())}; + rmm::device_buffer all_null_mask{ + create_null_mask(num_elements(), cudf::mask_state::ALL_NULL, cudf::test::get_default_stream())}; +}; + +TYPED_TEST_SUITE(TypedColumnTest, cudf::test::Types); + +/** + * @brief Verifies equality of the properties and data of a `column`'s views. + * + * @param col The `column` to verify + */ +void verify_column_views(cudf::column& col) +{ + cudf::column_view view = col; + cudf::mutable_column_view mutable_view = col; + EXPECT_EQ(col.type(), view.type()); + EXPECT_EQ(col.type(), mutable_view.type()); + EXPECT_EQ(col.size(), view.size()); + EXPECT_EQ(col.size(), mutable_view.size()); + EXPECT_EQ(col.null_count(), view.null_count()); + EXPECT_EQ(col.null_count(), mutable_view.null_count()); + EXPECT_EQ(view.null_count(0, col.size(), cudf::test::get_default_stream()), + mutable_view.null_count(0, col.size(), cudf::test::get_default_stream())); + EXPECT_EQ(view.has_nulls(0, col.size(), cudf::test::get_default_stream()), + mutable_view.has_nulls(0, col.size(), cudf::test::get_default_stream())); + EXPECT_EQ(col.null_count(), mutable_view.null_count()); + EXPECT_EQ(col.nullable(), view.nullable()); + EXPECT_EQ(col.nullable(), mutable_view.nullable()); + EXPECT_EQ(col.num_children(), view.num_children()); + EXPECT_EQ(col.num_children(), mutable_view.num_children()); + EXPECT_EQ(view.head(), mutable_view.head()); + EXPECT_EQ(view.data(), mutable_view.data()); + EXPECT_EQ(view.offset(), mutable_view.offset()); +} + +TYPED_TEST(TypedColumnTest, CopyConstructorWithMask) +{ + cudf::column original{ + this->type(), this->num_elements(), std::move(this->data), std::move(this->all_valid_mask), 0}; + cudf::column copy{original, cudf::test::get_default_stream()}; + verify_column_views(copy); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(original, copy); + + // Verify deep copy + cudf::column_view original_view = original; + cudf::column_view copy_view = copy; + EXPECT_NE(original_view.head(), copy_view.head()); + EXPECT_NE(original_view.null_mask(), copy_view.null_mask()); +}