Skip to content

Commit f1ecadd

Browse files
PointKernelPointKernel
committed
Fix
1 parent 589341d commit f1ecadd

10 files changed

Lines changed: 27 additions & 27 deletions

File tree

‎cpp/src/copying/copy.cu‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -157,7 +157,7 @@ std::unique_ptr<column> scatter_gather_based_if_else(cudf::column_view const& lh
157157
rmm::device_async_resource_ref mr)
158158
{
159159
auto gather_map = rmm::device_uvector<size_type>{static_cast<std::size_t>(size), stream};
160-
auto const gather_map_end = cudf::detail::copy_if(cuda::counting_iterator{size_type{0}},
160+
auto const gather_map_end = cudf::detail::copy_if(cuda::counting_iterator<size_type>{0},
161161
cuda::counting_iterator{size_type{size}},
162162
gather_map.begin(),
163163
is_left,
@@ -191,7 +191,7 @@ std::unique_ptr<column> scatter_gather_based_if_else(cudf::scalar const& lhs,
191191
rmm::device_async_resource_ref mr)
192192
{
193193
auto scatter_map = rmm::device_uvector<size_type>{static_cast<std::size_t>(size), stream};
194-
auto const scatter_map_end = cudf::detail::copy_if(cuda::counting_iterator{size_type{0}},
194+
auto const scatter_map_end = cudf::detail::copy_if(cuda::counting_iterator<size_type>{0},
195195
cuda::counting_iterator{size_type{size}},
196196
scatter_map.begin(),
197197
is_left,

‎cpp/src/groupby/hash/compute_global_memory_aggs.cuh‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -105,7 +105,7 @@ std::pair<std::unique_ptr<table>, rmm::device_uvector<size_type>> compute_aggs_d
105105
auto d_results_ptr = mutable_table_device_view::create(*agg_results, stream);
106106

107107
thrust::for_each_n(rmm::exec_policy_nosync(stream),
108-
cuda::counting_iterator{int64_t{0}},
108+
cuda::counting_iterator<int64_t>{0},
109109
num_rows * static_cast<int64_t>(h_agg_kinds.size()),
110110
compute_single_pass_aggs_dense_output_fn{
111111
target_indices.begin(), d_agg_kinds.data(), *d_values, *d_results_ptr});

‎cpp/src/io/orc/reader_impl_decode.cu‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -193,7 +193,7 @@ rmm::device_buffer decompress_stripe_data(
193193
// Check if any block has been failed to decompress.
194194
// Not using `thrust::any` or `thrust::count_if` to defer stream sync.
195195
thrust::for_each(rmm::exec_policy_nosync(stream),
196-
cuda::counting_iterator{std::size_t{0}},
196+
cuda::counting_iterator<std::size_t>{0},
197197
cuda::counting_iterator{inflate_res.size()},
198198
[results = inflate_res.begin(),
199199
any_block_failure = any_block_failure.device_ptr()] __device__(auto const idx) {

‎cpp/src/io/parquet/reader_impl_chunking_utils.cu‎

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -390,7 +390,7 @@ std::tuple<rmm::device_uvector<page_span>, size_t, size_t> compute_next_subpass(
390390

391391
// for each column, collect the set of pages that spans start_row / end_row
392392
rmm::device_uvector<page_span> page_bounds(num_columns, stream);
393-
auto iter = cuda::counting_iterator{size_t{0}};
393+
auto iter = cuda::counting_iterator<size_t>{0};
394394
auto page_row_index =
395395
cudf::detail::make_counting_transform_iterator(0, get_page_end_row_index{c_info});
396396
thrust::transform(
@@ -743,7 +743,7 @@ rmm::device_uvector<size_t> compute_decompression_scratch_sizes(
743743

744744
// Collect pages with matching codecs
745745
rmm::device_uvector<device_span<uint8_t const>> temp_spans(pages.size(), stream);
746-
auto iter = cuda::counting_iterator{size_t{0}};
746+
auto iter = cuda::counting_iterator<size_t>{0};
747747
thrust::for_each(
748748
rmm::exec_policy_nosync(stream),
749749
iter,
@@ -793,7 +793,7 @@ rmm::device_uvector<size_t> compute_decompression_scratch_sizes(
793793

794794
// Apply the adjustment ratio to each page's temporary cost
795795
thrust::for_each(rmm::exec_policy_nosync(stream),
796-
cuda::counting_iterator{size_t{0}},
796+
cuda::counting_iterator<size_t>{0},
797797
cuda::counting_iterator{pages.size()},
798798
[pages = pages.begin(),
799799
chunks = chunks.begin(),
@@ -821,7 +821,7 @@ void include_scratch_size(device_span<size_t const> temp_cost,
821821
device_span<cumulative_page_info> c_info,
822822
rmm::cuda_stream_view stream)
823823
{
824-
auto iter = cuda::counting_iterator{size_t{0}};
824+
auto iter = cuda::counting_iterator<size_t>{0};
825825
thrust::for_each(rmm::exec_policy_nosync(stream),
826826
iter,
827827
iter + c_info.size(),

‎cpp/src/io/parquet/reader_impl_preprocess.cu‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -108,7 +108,7 @@ void reader_impl::allocate_nesting_info()
108108

109109
// generate the number of nesting info structs needed per-page, by column
110110
std::vector<int> per_page_nesting_info_size(num_columns);
111-
auto iter = cuda::counting_iterator{size_type{0}};
111+
auto iter = cuda::counting_iterator<size_type>{0};
112112
std::transform(iter, iter + num_columns, per_page_nesting_info_size.begin(), [&](size_type i) {
113113
// Schema index of the current input column
114114
auto const schema_idx = _input_columns[i].schema_idx;
@@ -128,7 +128,7 @@ void reader_impl::allocate_nesting_info()
128128

129129
// compute total # of page_nesting infos needed and allocate space. doing this in one
130130
// buffer to keep it to a single gpu allocation
131-
auto counting_iter = cuda::counting_iterator{size_t{0}};
131+
auto counting_iter = cuda::counting_iterator<size_t>{0};
132132
size_t const total_page_nesting_infos =
133133
std::accumulate(counting_iter, counting_iter + num_columns, 0, [&](int total, size_t index) {
134134
return total + (per_page_nesting_info_size[index] * subpass.column_page_count[index]);

‎cpp/src/jit/parser.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -346,7 +346,7 @@ std::string ptx_parser::parse_function_header(std::string const& src)
346346

347347
std::vector<std::string> param_decls;
348348

349-
std::transform(cuda::counting_iterator{size_t{0}},
349+
std::transform(cuda::counting_iterator<size_t>{0},
350350
cuda::counting_iterator{ptx_params.size()},
351351
std::back_inserter(param_decls),
352352
[&](size_t param_index) {

‎cpp/src/join/mark_join.cu‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -292,7 +292,7 @@ std::unique_ptr<rmm::device_uvector<cudf::size_type>> mark_join::mark_probe_and_
292292

293293
auto materialize_probe_rows = [&](auto const& key_fn) {
294294
rmm::device_uvector<probe_key_type> probe_rows(probe.num_rows(), stream);
295-
cub::DeviceTransform::Transform(cuda::counting_iterator{size_type{0}},
295+
cub::DeviceTransform::Transform(cuda::counting_iterator<size_type>{0},
296296
probe_rows.begin(),
297297
probe.num_rows(),
298298
cuda::proclaim_return_type<probe_key_type>(key_fn),
@@ -397,7 +397,7 @@ std::unique_ptr<rmm::device_uvector<cudf::size_type>> mark_join::mark_probe_and_
397397
auto const bitmask_buffer_and_ptr = build_row_bitmask(_build, stream);
398398
auto const row_bitmask_ptr = bitmask_buffer_and_ptr.second;
399399
thrust::copy_if(rmm::exec_policy_nosync(stream),
400-
cuda::counting_iterator{size_type{0}},
400+
cuda::counting_iterator<size_type>{0},
401401
cuda::counting_iterator{_build.num_rows()},
402402
result.begin() + unmatched_valid,
403403
row_is_null{row_bitmask_ptr});
@@ -438,7 +438,7 @@ mark_join::mark_join(cudf::table_view const& build,
438438
<<<grid_size, cuco::detail::default_block_size(), 0, stream.value()>>>(
439439
build_iter,
440440
_build.num_rows(),
441-
cuda::counting_iterator{size_type{0}},
441+
cuda::counting_iterator<size_type>{0},
442442
row_is_valid{row_bitmask_ptr},
443443
insert_ref);
444444
} else {

‎cpp/src/lists/combine/concatenate_rows.cu‎

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -72,13 +72,13 @@ generate_regrouped_offsets_and_null_mask(table_device_view const& input,
7272
data_type{type_to_id<size_type>()}, input.num_rows() + 1, mask_state::UNALLOCATED, stream, mr);
7373

7474
auto keys = thrust::make_transform_iterator(
75-
cuda::counting_iterator{size_t{0}},
75+
cuda::counting_iterator<size_t>{0},
7676
cuda::proclaim_return_type<size_type>([num_columns = input.num_columns()] __device__(
7777
size_t i) -> size_type { return i / num_columns; }));
7878

7979
// generate sizes for the regrouped rows
8080
auto values = thrust::make_transform_iterator(
81-
cuda::counting_iterator{size_t{0}},
81+
cuda::counting_iterator<size_t>{0},
8282
cuda::proclaim_return_type<size_type>([input,
8383
row_null_counts = row_null_counts.data(),
8484
null_policy] __device__(size_t i) -> size_type {
@@ -152,12 +152,12 @@ rmm::device_uvector<size_type> generate_null_counts(table_device_view const& inp
152152
rmm::device_uvector<size_type> null_counts(input.num_rows(), stream);
153153

154154
auto keys = thrust::make_transform_iterator(
155-
cuda::counting_iterator{size_t{0}},
155+
cuda::counting_iterator<size_t>{0},
156156
cuda::proclaim_return_type<size_type>([num_columns = input.num_columns()] __device__(
157157
size_t i) -> size_type { return i / num_columns; }));
158158

159159
auto null_values = thrust::make_transform_iterator(
160-
cuda::counting_iterator{size_t{0}},
160+
cuda::counting_iterator<size_t>{0},
161161
cuda::proclaim_return_type<size_type>([input] __device__(size_t i) -> size_type {
162162
auto const col_index = i % input.num_columns();
163163
auto const row_index = i / input.num_columns();
@@ -226,7 +226,7 @@ std::unique_ptr<column> concatenate_rows(table_view const& input,
226226
// be nullified.
227227
if (build_null_mask) {
228228
auto [null_mask, null_count] = [&]() {
229-
auto iter = cuda::counting_iterator{size_t{0}};
229+
auto iter = cuda::counting_iterator<size_t>{0};
230230

231231
// IGNORE. Output row is nullified if all input rows are null.
232232
if (null_policy == concatenate_null_policy::IGNORE) {
@@ -264,7 +264,7 @@ std::unique_ptr<column> concatenate_rows(table_view const& input,
264264
// we had concatenated all the rows together instead of concatenating within the rows. To fix
265265
// this we can simply swap in a new set of offsets that re-groups them. bmo
266266
auto iter = thrust::make_transform_iterator(
267-
cuda::counting_iterator{size_t{0}},
267+
cuda::counting_iterator<size_t>{0},
268268
cuda::proclaim_return_type<size_type>(
269269
[num_columns = input.num_columns(),
270270
num_rows = input.num_rows()] __device__(size_t i) -> size_type {

‎cpp/src/rolling/detail/lead_lag_nested.cuh‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -131,7 +131,7 @@ std::unique_ptr<column> compute_lead_lag_for_nested(aggregation::Kind op,
131131
auto const null_index = input.size();
132132
if (op == aggregation::LEAD) {
133133
thrust::transform(rmm::exec_policy_nosync(stream),
134-
cuda::counting_iterator{size_type{0}},
134+
cuda::counting_iterator<size_type>{0},
135135
cuda::counting_iterator{size_type{input.size()}},
136136
gather_map.begin<size_type>(),
137137
cuda::proclaim_return_type<size_type>(
@@ -140,7 +140,7 @@ std::unique_ptr<column> compute_lead_lag_for_nested(aggregation::Kind op,
140140
}));
141141
} else {
142142
thrust::transform(rmm::exec_policy_nosync(stream),
143-
cuda::counting_iterator{size_type{0}},
143+
cuda::counting_iterator<size_type>{0},
144144
cuda::counting_iterator{size_type{input.size()}},
145145
gather_map.begin<size_type>(),
146146
cuda::proclaim_return_type<size_type>(
@@ -163,7 +163,7 @@ std::unique_ptr<column> compute_lead_lag_for_nested(aggregation::Kind op,
163163

164164
// Find all indices at which LEAD/LAG computed nulls previously.
165165
auto scatter_map_end =
166-
cudf::detail::copy_if(cuda::counting_iterator{size_type{0}},
166+
cudf::detail::copy_if(cuda::counting_iterator<size_type>{0},
167167
cuda::counting_iterator{size_type{input.size()}},
168168
scatter_map.begin(),
169169
is_null_index_predicate(input.size(), gather_map.begin<size_type>()),

‎cpp/tests/utilities/debug_utilities.cu‎

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -178,7 +178,7 @@ struct column_view_printer {
178178
out.resize(col.size());
179179

180180
if (col.nullable()) {
181-
std::transform(cuda::counting_iterator{size_type{0}},
181+
std::transform(cuda::counting_iterator<size_type>{0},
182182
cuda::counting_iterator{col.size()},
183183
out.begin(),
184184
[&h_data](auto idx) {
@@ -227,7 +227,7 @@ struct column_view_printer {
227227
{
228228
auto const h_data = cudf::test::to_host<Element>(col);
229229
if (col.nullable()) {
230-
std::transform(cuda::counting_iterator{size_type{0}},
230+
std::transform(cuda::counting_iterator<size_type>{0},
231231
cuda::counting_iterator{col.size()},
232232
std::back_inserter(out),
233233
[&h_data](auto idx) {
@@ -272,7 +272,7 @@ struct column_view_printer {
272272
};
273273

274274
out.resize(col.size());
275-
std::transform(cuda::counting_iterator{size_type{0}},
275+
std::transform(cuda::counting_iterator<size_type>{0},
276276
cuda::counting_iterator{col.size()},
277277
out.begin(),
278278
[&](auto idx) {
@@ -313,7 +313,7 @@ struct column_view_printer {
313313
out.resize(col.size());
314314

315315
if (col.nullable()) {
316-
std::transform(cuda::counting_iterator{size_type{0}},
316+
std::transform(cuda::counting_iterator<size_type>{0},
317317
cuda::counting_iterator{col.size()},
318318
out.begin(),
319319
[&h_data](auto idx) {

0 commit comments

Comments
 (0)