Skip to content

Commit a4bbd09

Browse files
authored
Expose stream-ordering in column view APIs (rapidsai#17434)
Adds stream parameter to ``` cudf::detail::column_view_base::null_count(begin, end) cudf::detail::column_view_base::has_nulls(begin, end) ``` Note: Since stream-ordered prefetching is [back-logged](rapidsai#17434 (comment)), we defer modifying the `get_data` member functions to accept a stream parameter for now. Reference: 1. rapidsai#13744 2. rapidsai#16251 (comment) Authors: - Shruti Shivakumar (https://github.com/shrshi) Approvers: - Nghia Truong (https://github.com/ttnghia) - Bradley Dice (https://github.com/bdice) URL: rapidsai#17434
1 parent 7f2b2ba commit a4bbd09

File tree

6 files changed

+129
-15
lines changed

6 files changed

+129
-15
lines changed

cpp/include/cudf/column/column_view.hpp

Lines changed: 11 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
/*
2-
* Copyright (c) 2019-2024, NVIDIA CORPORATION.
2+
* Copyright (c) 2019-2025, NVIDIA CORPORATION.
33
*
44
* Licensed under the Apache License, Version 2.0 (the "License");
55
* you may not use this file except in compliance with the License.
@@ -176,9 +176,13 @@ class column_view_base {
176176
*
177177
* @param[in] begin The starting index of the range (inclusive).
178178
* @param[in] end The index of the last element in the range (exclusive).
179+
* @param[in] stream CUDA stream used for device memory operations and kernel launches
179180
* @return The count of null elements in the given range
180181
*/
181-
[[nodiscard]] size_type null_count(size_type begin, size_type end) const;
182+
[[nodiscard]] size_type null_count(
183+
size_type begin,
184+
size_type end,
185+
rmm::cuda_stream_view stream = cudf::get_default_stream()) const;
182186

183187
/**
184188
* @brief Indicates if the column contains null elements,
@@ -198,12 +202,15 @@ class column_view_base {
198202
*
199203
* @param begin The starting index of the range (inclusive).
200204
* @param end The index of the last element in the range (exclusive).
205+
* @param stream CUDA stream used for device memory operations and kernel launches
201206
* @return true One or more elements are null in the range [begin, end)
202207
* @return false All elements are valid in the range [begin, end)
203208
*/
204-
[[nodiscard]] bool has_nulls(size_type begin, size_type end) const
209+
[[nodiscard]] bool has_nulls(size_type begin,
210+
size_type end,
211+
rmm::cuda_stream_view stream = cudf::get_default_stream()) const
205212
{
206-
return null_count(begin, end) > 0;
213+
return null_count(begin, end, stream) > 0;
207214
}
208215

209216
/**

cpp/src/column/column_view.cpp

Lines changed: 5 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
/*
2-
* Copyright (c) 2019-2024, NVIDIA CORPORATION.
2+
* Copyright (c) 2019-2025, NVIDIA CORPORATION.
33
*
44
* Licensed under the Apache License, Version 2.0 (the "License");
55
* 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,
141141
}
142142
}
143143

144-
size_type column_view_base::null_count(size_type begin, size_type end) const
144+
size_type column_view_base::null_count(size_type begin,
145+
size_type end,
146+
rmm::cuda_stream_view stream) const
145147
{
146148
CUDF_EXPECTS((begin >= 0) && (end <= size()) && (begin <= end), "Range is out of bounds.");
147149
return (null_count() == 0)
148150
? 0
149-
: cudf::detail::null_count(
150-
null_mask(), offset() + begin, offset() + end, cudf::get_default_stream());
151+
: cudf::detail::null_count(null_mask(), offset() + begin, offset() + end, stream);
151152
}
152153

153154
bool is_shallow_equivalent(column_view const& lhs, column_view const& rhs)

cpp/src/copying/copy_range.cu

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
/*
2-
* Copyright (c) 2019-2024, NVIDIA CORPORATION.
2+
* Copyright (c) 2019-2025, NVIDIA CORPORATION.
33
*
44
* Licensed under the Apache License, Version 2.0 (the "License");
55
* you may not use this file except in compliance with the License.
@@ -103,7 +103,7 @@ struct out_of_place_copy_range_dispatch {
103103
rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref())
104104
{
105105
auto p_ret = std::make_unique<cudf::column>(target, stream, mr);
106-
if ((!p_ret->nullable()) && source.has_nulls(source_begin, source_end)) {
106+
if ((!p_ret->nullable()) && source.has_nulls(source_begin, source_end, stream)) {
107107
p_ret->set_null_mask(
108108
cudf::detail::create_null_mask(p_ret->size(), cudf::mask_state::ALL_VALID, stream, mr), 0);
109109
}

cpp/src/rolling/grouped_rolling.cu

Lines changed: 7 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
/*
2-
* Copyright (c) 2020-2024, NVIDIA CORPORATION.
2+
* Copyright (c) 2020-2025, NVIDIA CORPORATION.
33
*
44
* Licensed under the Apache License, Version 2.0 (the "License");
55
* you may not use this file except in compliance with the License.
@@ -419,7 +419,7 @@ template <typename T, CUDF_ENABLE_IF(cudf::column_device_view::has_element_acces
419419
/// at the beginning of the column or at the end.
420420
/// If no null values are founds, null_begin and null_end are 0.
421421
std::tuple<size_type, size_type> get_null_bounds_for_orderby_column(
422-
column_view const& orderby_column)
422+
column_view const& orderby_column, rmm::cuda_stream_view stream)
423423
{
424424
auto const num_rows = orderby_column.size();
425425
auto const num_nulls = orderby_column.null_count();
@@ -429,7 +429,7 @@ std::tuple<size_type, size_type> get_null_bounds_for_orderby_column(
429429
return std::make_tuple(0, num_nulls);
430430
}
431431

432-
auto const first_row_is_null = orderby_column.null_count(0, 1) == 1;
432+
auto const first_row_is_null = orderby_column.null_count(0, 1, stream) == 1;
433433

434434
return first_row_is_null ? std::make_tuple(0, num_nulls)
435435
: std::make_tuple(num_rows - num_nulls, num_rows);
@@ -451,7 +451,8 @@ std::unique_ptr<column> range_window_ASC(column_view const& input,
451451
rmm::cuda_stream_view stream,
452452
rmm::device_async_resource_ref mr)
453453
{
454-
auto [h_nulls_begin_idx, h_nulls_end_idx] = get_null_bounds_for_orderby_column(orderby_column);
454+
auto [h_nulls_begin_idx, h_nulls_end_idx] =
455+
get_null_bounds_for_orderby_column(orderby_column, stream);
455456
auto const p_orderby_device_view = cudf::column_device_view::create(orderby_column, stream);
456457

457458
auto const preceding_calculator = cuda::proclaim_return_type<size_type>(
@@ -740,7 +741,8 @@ std::unique_ptr<column> range_window_DESC(column_view const& input,
740741
rmm::cuda_stream_view stream,
741742
rmm::device_async_resource_ref mr)
742743
{
743-
auto [h_nulls_begin_idx, h_nulls_end_idx] = get_null_bounds_for_orderby_column(orderby_column);
744+
auto [h_nulls_begin_idx, h_nulls_end_idx] =
745+
get_null_bounds_for_orderby_column(orderby_column, stream);
744746
auto const p_orderby_device_view = cudf::column_device_view::create(orderby_column, stream);
745747

746748
auto const preceding_calculator = cuda::proclaim_return_type<size_type>(

cpp/tests/CMakeLists.txt

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -698,6 +698,7 @@ if(CUDF_BUILD_STREAMS_TEST_UTIL)
698698
endif()
699699

700700
ConfigureTest(STREAM_BINARYOP_TEST streams/binaryop_test.cpp STREAM_MODE testing)
701+
ConfigureTest(STREAM_COLUMN_VIEW_TEST streams/column_view_test.cpp STREAM_MODE testing)
701702
ConfigureTest(STREAM_CONCATENATE_TEST streams/concatenate_test.cpp STREAM_MODE testing)
702703
ConfigureTest(STREAM_COPYING_TEST streams/copying_test.cpp STREAM_MODE testing)
703704
ConfigureTest(STREAM_CSVIO_TEST streams/io/csv_test.cpp STREAM_MODE testing)
Lines changed: 103 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,103 @@
1+
/*
2+
* Copyright (c) 2025, NVIDIA CORPORATION.
3+
*
4+
* Licensed under the Apache License, Version 2.0 (the "License");
5+
* you may not use this file except in compliance with the License.
6+
* You may obtain a copy of the License at
7+
*
8+
* http://www.apache.org/licenses/LICENSE-2.0
9+
*
10+
* Unless required by applicable law or agreed to in writing, software
11+
* distributed under the License is distributed on an "AS IS" BASIS,
12+
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13+
* See the License for the specific language governing permissions and
14+
* limitations under the License.
15+
*/
16+
17+
#include <cudf_test/base_fixture.hpp>
18+
#include <cudf_test/column_wrapper.hpp>
19+
#include <cudf_test/default_stream.hpp>
20+
#include <cudf_test/type_lists.hpp>
21+
22+
#include <cudf/column/column_view.hpp>
23+
#include <cudf/null_mask.hpp>
24+
#include <cudf/transform.hpp>
25+
26+
#include <random>
27+
#include <vector>
28+
29+
template <typename T>
30+
struct TypedColumnTest : public cudf::test::BaseFixture {
31+
cudf::data_type type() { return cudf::data_type{cudf::type_to_id<T>()}; }
32+
33+
TypedColumnTest(rmm::cuda_stream_view stream = cudf::test::get_default_stream())
34+
: data{_num_elements * cudf::size_of(type()), stream},
35+
mask{cudf::bitmask_allocation_size_bytes(_num_elements), stream}
36+
{
37+
std::vector<char> h_data(std::max(data.size(), mask.size()));
38+
std::iota(h_data.begin(), h_data.end(), 0);
39+
CUDF_CUDA_TRY(
40+
cudaMemcpyAsync(data.data(), h_data.data(), data.size(), cudaMemcpyDefault, stream.value()));
41+
CUDF_CUDA_TRY(
42+
cudaMemcpyAsync(mask.data(), h_data.data(), mask.size(), cudaMemcpyDefault, stream.value()));
43+
}
44+
45+
cudf::size_type num_elements() { return _num_elements; }
46+
47+
std::random_device r;
48+
std::default_random_engine generator{r()};
49+
std::uniform_int_distribution<cudf::size_type> distribution{200, 1000};
50+
cudf::size_type _num_elements{distribution(generator)};
51+
rmm::device_buffer data{};
52+
rmm::device_buffer mask{};
53+
rmm::device_buffer all_valid_mask{create_null_mask(
54+
num_elements(), cudf::mask_state::ALL_VALID, cudf::test::get_default_stream())};
55+
rmm::device_buffer all_null_mask{
56+
create_null_mask(num_elements(), cudf::mask_state::ALL_NULL, cudf::test::get_default_stream())};
57+
};
58+
59+
TYPED_TEST_SUITE(TypedColumnTest, cudf::test::Types<int32_t>);
60+
61+
/**
62+
* @brief Verifies equality of the properties and data of a `column`'s views.
63+
*
64+
* @param col The `column` to verify
65+
*/
66+
void verify_column_views(cudf::column& col)
67+
{
68+
cudf::column_view view = col;
69+
cudf::mutable_column_view mutable_view = col;
70+
EXPECT_EQ(col.type(), view.type());
71+
EXPECT_EQ(col.type(), mutable_view.type());
72+
EXPECT_EQ(col.size(), view.size());
73+
EXPECT_EQ(col.size(), mutable_view.size());
74+
EXPECT_EQ(col.null_count(), view.null_count());
75+
EXPECT_EQ(col.null_count(), mutable_view.null_count());
76+
EXPECT_EQ(view.null_count(0, col.size(), cudf::test::get_default_stream()),
77+
mutable_view.null_count(0, col.size(), cudf::test::get_default_stream()));
78+
EXPECT_EQ(view.has_nulls(0, col.size(), cudf::test::get_default_stream()),
79+
mutable_view.has_nulls(0, col.size(), cudf::test::get_default_stream()));
80+
EXPECT_EQ(col.null_count(), mutable_view.null_count());
81+
EXPECT_EQ(col.nullable(), view.nullable());
82+
EXPECT_EQ(col.nullable(), mutable_view.nullable());
83+
EXPECT_EQ(col.num_children(), view.num_children());
84+
EXPECT_EQ(col.num_children(), mutable_view.num_children());
85+
EXPECT_EQ(view.head(), mutable_view.head());
86+
EXPECT_EQ(view.data<char>(), mutable_view.data<char>());
87+
EXPECT_EQ(view.offset(), mutable_view.offset());
88+
}
89+
90+
TYPED_TEST(TypedColumnTest, CopyConstructorWithMask)
91+
{
92+
cudf::column original{
93+
this->type(), this->num_elements(), std::move(this->data), std::move(this->all_valid_mask), 0};
94+
cudf::column copy{original, cudf::test::get_default_stream()};
95+
verify_column_views(copy);
96+
CUDF_TEST_EXPECT_COLUMNS_EQUAL(original, copy);
97+
98+
// Verify deep copy
99+
cudf::column_view original_view = original;
100+
cudf::column_view copy_view = copy;
101+
EXPECT_NE(original_view.head(), copy_view.head());
102+
EXPECT_NE(original_view.null_mask(), copy_view.null_mask());
103+
}

0 commit comments

Comments
 (0)