diff --git a/cpp/include/cudf/detail/copy_if_else.cuh b/cpp/include/cudf/detail/copy_if_else.cuh index 46180f3e1f3..3faba5ef51b 100644 --- a/cpp/include/cudf/detail/copy_if_else.cuh +++ b/cpp/include/cudf/detail/copy_if_else.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2020, NVIDIA CORPORATION. + * Copyright (c) 2019-2021, 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,12 +176,8 @@ std::unique_ptr copy_if_else( constexpr int block_size = 256; cudf::detail::grid_1d grid{num_els, block_size, 1}; - std::unique_ptr out = - make_fixed_width_column(output_type, - size, - nullable ? mask_state::UNINITIALIZED : mask_state::UNALLOCATED, - stream.value(), - mr); + std::unique_ptr out = make_fixed_width_column( + output_type, size, nullable ? mask_state::UNINITIALIZED : mask_state::UNALLOCATED, stream, mr); auto out_v = mutable_column_device_view::create(*out); diff --git a/cpp/include/cudf/detail/scatter.cuh b/cpp/include/cudf/detail/scatter.cuh index 305aa51afb9..0918f071443 100644 --- a/cpp/include/cudf/detail/scatter.cuh +++ b/cpp/include/cudf/detail/scatter.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020, NVIDIA CORPORATION. + * Copyright (c) 2020-2021, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -115,7 +115,7 @@ struct column_scatterer_impl { rmm::mr::device_memory_resource* mr) const { using strings::detail::create_string_vector_from_column; - auto const source_vector = create_string_vector_from_column(source, stream.value()); + auto const source_vector = create_string_vector_from_column(source, stream); auto const begin = source_vector.begin(); auto const end = begin + std::distance(scatter_map_begin, scatter_map_end); return strings::detail::scatter(begin, end, scatter_map_begin, target, stream, mr); diff --git a/cpp/include/cudf/lists/detail/scatter.cuh b/cpp/include/cudf/lists/detail/scatter.cuh index 53f0472fedc..496f9ccc838 100644 --- a/cpp/include/cudf/lists/detail/scatter.cuh +++ b/cpp/include/cudf/lists/detail/scatter.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020, NVIDIA CORPORATION. + * Copyright (c) 2020-2021, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -375,7 +375,7 @@ struct list_child_constructor { num_child_rows, child_null_mask.first, child_null_mask.second, - stream.value(), + stream, mr); auto copy_child_values_for_list_index = [d_scattered_lists = @@ -514,7 +514,7 @@ struct list_child_constructor { std::move(string_chars), child_null_mask.second, // Null count. std::move(child_null_mask.first), // Null mask. - stream.value(), + stream, mr); } @@ -616,7 +616,7 @@ struct list_child_constructor { std::move(child_column), child_null_mask.second, // Null count std::move(child_null_mask.first), // Null mask - stream.value(), + stream, mr); } @@ -710,7 +710,7 @@ struct list_child_constructor { std::move(child_columns), child_null_mask.second, std::move(child_null_mask.first), - stream.value(), + stream, mr); } }; @@ -814,7 +814,7 @@ std::unique_ptr scatter( std::move(child_column), cudf::UNKNOWN_NULL_COUNT, std::move(null_mask), - stream.value(), + stream, mr); } diff --git a/cpp/include/cudf/strings/detail/copy_if_else.cuh b/cpp/include/cudf/strings/detail/copy_if_else.cuh index f0d0e4a8f7e..176a548da4d 100644 --- a/cpp/include/cudf/strings/detail/copy_if_else.cuh +++ b/cpp/include/cudf/strings/detail/copy_if_else.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2020, NVIDIA CORPORATION. + * Copyright (c) 2019-2021, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -112,7 +112,7 @@ std::unique_ptr copy_if_else( std::move(chars_column), null_count, std::move(null_mask), - stream.value(), + stream, mr); } diff --git a/cpp/include/cudf/strings/detail/scatter.cuh b/cpp/include/cudf/strings/detail/scatter.cuh index 999fb22b162..7e2513a7633 100644 --- a/cpp/include/cudf/strings/detail/scatter.cuh +++ b/cpp/include/cudf/strings/detail/scatter.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2020, NVIDIA CORPORATION. + * Copyright (c) 2019-2021, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -69,8 +69,7 @@ std::unique_ptr scatter( if (target.has_nulls()) null_mask = cudf::detail::copy_bitmask(target.parent(), stream, mr); // create string vectors - rmm::device_vector target_vector = - create_string_vector_from_column(target, stream.value()); + rmm::device_vector target_vector = create_string_vector_from_column(target, stream); // do the scatter thrust::scatter(rmm::exec_policy(stream), begin, end, scatter_map, target_vector.begin()); diff --git a/cpp/src/ast/transform.cu b/cpp/src/ast/transform.cu index ffc80a926fb..b61f48f7a84 100644 --- a/cpp/src/ast/transform.cu +++ b/cpp/src/ast/transform.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020, NVIDIA CORPORATION. + * Copyright (c) 2020-2021, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -128,14 +128,14 @@ std::unique_ptr compute_column(table_view const table, reinterpret_cast(device_data_buffer_ptr + buffer_offsets[3]); // Create table device view - auto table_device = table_device_view::create(table, stream.value()); + auto table_device = table_device_view::create(table, stream); auto const table_num_rows = table.num_rows(); // Prepare output column auto output_column = cudf::make_fixed_width_column( - expr_data_type, table_num_rows, mask_state::UNALLOCATED, stream.value(), mr); + expr_data_type, table_num_rows, mask_state::UNALLOCATED, stream, mr); auto mutable_output_device = - cudf::mutable_column_device_view::create(output_column->mutable_view(), stream.value()); + cudf::mutable_column_device_view::create(output_column->mutable_view(), stream); // Configure kernel parameters auto const num_intermediates = expr_linearizer.get_intermediate_count(); diff --git a/cpp/src/binaryop/binaryop.cpp b/cpp/src/binaryop/binaryop.cpp index 2b1bc348d45..9602839f6e4 100644 --- a/cpp/src/binaryop/binaryop.cpp +++ b/cpp/src/binaryop/binaryop.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2020, NVIDIA CORPORATION. + * Copyright (c) 2019-2021, NVIDIA CORPORATION. * * Copyright 2018-2019 BlazingDB, Inc. * Copyright 2018 Christian Noboa Mardini @@ -110,7 +110,7 @@ void binary_operation(mutable_column_view& out, { if (is_null_dependent(op)) { cudf::jit::launcher( - hash, code::kernel, header_names, cudf::jit::compiler_flags, headers_code, stream.value()) + hash, code::kernel, header_names, cudf::jit::compiler_flags, headers_code, stream) .set_kernel_inst("kernel_v_s_with_validity", // name of the kernel we are // launching {cudf::jit::get_type_name(out.type()), // list of template arguments @@ -127,7 +127,7 @@ void binary_operation(mutable_column_view& out, lhs.is_valid()); } else { cudf::jit::launcher( - hash, code::kernel, header_names, cudf::jit::compiler_flags, headers_code, stream.value()) + hash, code::kernel, header_names, cudf::jit::compiler_flags, headers_code, stream) .set_kernel_inst("kernel_v_s", // name of the kernel we are // launching {cudf::jit::get_type_name(out.type()), // list of template arguments @@ -149,7 +149,7 @@ void binary_operation(mutable_column_view& out, { if (is_null_dependent(op)) { cudf::jit::launcher( - hash, code::kernel, header_names, cudf::jit::compiler_flags, headers_code, stream.value()) + hash, code::kernel, header_names, cudf::jit::compiler_flags, headers_code, stream) .set_kernel_inst("kernel_v_s_with_validity", // name of the kernel we are // launching {cudf::jit::get_type_name(out.type()), // list of template arguments @@ -166,7 +166,7 @@ void binary_operation(mutable_column_view& out, rhs.is_valid()); } else { cudf::jit::launcher( - hash, code::kernel, header_names, cudf::jit::compiler_flags, headers_code, stream.value()) + hash, code::kernel, header_names, cudf::jit::compiler_flags, headers_code, stream) .set_kernel_inst("kernel_v_s", // name of the kernel we are // launching {cudf::jit::get_type_name(out.type()), // list of template arguments @@ -188,7 +188,7 @@ void binary_operation(mutable_column_view& out, { if (is_null_dependent(op)) { cudf::jit::launcher( - hash, code::kernel, header_names, cudf::jit::compiler_flags, headers_code, stream.value()) + hash, code::kernel, header_names, cudf::jit::compiler_flags, headers_code, stream) .set_kernel_inst("kernel_v_v_with_validity", // name of the kernel we are // launching {cudf::jit::get_type_name(out.type()), // list of template arguments @@ -206,7 +206,7 @@ void binary_operation(mutable_column_view& out, rhs.offset()); } else { cudf::jit::launcher( - hash, code::kernel, header_names, cudf::jit::compiler_flags, headers_code, stream.value()) + hash, code::kernel, header_names, cudf::jit::compiler_flags, headers_code, stream) .set_kernel_inst("kernel_v_v", // name of the kernel we are // launching {cudf::jit::get_type_name(out.type()), // list of template arguments @@ -235,7 +235,7 @@ void binary_operation(mutable_column_view& out, cudf::jit::parse_single_function_ptx(ptx, "GENERIC_BINARY_OP", output_type_name) + code::kernel; cudf::jit::launcher( - ptx_hash, cuda_source, header_names, cudf::jit::compiler_flags, headers_code, stream.value()) + ptx_hash, cuda_source, header_names, cudf::jit::compiler_flags, headers_code, stream) .set_kernel_inst("kernel_v_v", // name of the kernel // we are launching {output_type_name, // list of template arguments diff --git a/cpp/src/column/column.cu b/cpp/src/column/column.cu index 796f522d949..6fb66b6001a 100644 --- a/cpp/src/column/column.cu +++ b/cpp/src/column/column.cu @@ -191,7 +191,7 @@ struct create_column_from_view { std::unique_ptr operator()() { cudf::strings_column_view sview(view); - return cudf::strings::detail::copy_slice(sview, 0, view.size(), 1, stream.value(), mr); + return cudf::strings::detail::copy_slice(sview, 0, view.size(), 1, stream, mr); } template operator()() { auto lists_view = lists_column_view(view); - return cudf::lists::detail::copy_slice(lists_view, 0, view.size(), stream.value(), mr); + return cudf::lists::detail::copy_slice(lists_view, 0, view.size(), stream, mr); } template mutable_view(); - detail::fill_in_place(view, 0, size, value, stream.value()); + detail::fill_in_place(view, 0, size, value, stream); return output_column; } }; diff --git a/cpp/src/copying/sample.cu b/cpp/src/copying/sample.cu index 2e6c1549afa..9e0f432fb1d 100644 --- a/cpp/src/copying/sample.cu +++ b/cpp/src/copying/sample.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020, NVIDIA CORPORATION. + * Copyright (c) 2020-2021, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -62,8 +62,8 @@ std::unique_ptr sample(table_view const& input, return detail::gather(input, begin, end, out_of_bounds_policy::DONT_CHECK, stream, mr); } else { - auto gather_map = make_numeric_column( - data_type{type_id::INT32}, num_rows, mask_state::UNALLOCATED, stream.value()); + auto gather_map = + make_numeric_column(data_type{type_id::INT32}, num_rows, mask_state::UNALLOCATED, stream); auto gather_map_mutable_view = gather_map->mutable_view(); // Shuffle all the row indices thrust::shuffle_copy(rmm::exec_policy(stream), diff --git a/cpp/src/groupby/sort/sort_helper.cu b/cpp/src/groupby/sort/sort_helper.cu index f8cb49fe1aa..8f3070c3497 100644 --- a/cpp/src/groupby/sort/sort_helper.cu +++ b/cpp/src/groupby/sort/sort_helper.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2020, NVIDIA CORPORATION. + * Copyright (c) 2019-2021, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -162,7 +162,7 @@ sort_groupby_helper::index_vector const& sort_groupby_helper::group_offsets( _group_offsets = std::make_unique(num_keys(stream) + 1); - auto device_input_table = table_device_view::create(_keys, stream.value()); + auto device_input_table = table_device_view::create(_keys, stream); auto sorted_order = key_sort_order().data(); decltype(_group_offsets->begin()) result_end; diff --git a/cpp/src/interop/from_arrow.cpp b/cpp/src/interop/from_arrow.cpp index 0bdb1f5eeb6..729b98d85a8 100644 --- a/cpp/src/interop/from_arrow.cpp +++ b/cpp/src/interop/from_arrow.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020, NVIDIA CORPORATION. + * Copyright (c) 2020-2021, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -183,7 +183,7 @@ std::unique_ptr dispatch_to_cudf_column::operator()( auto out_col = mask_to_bools(static_cast(data.data()), array.offset(), array.offset() + array.length(), - stream.value(), + stream, mr); auto const has_nulls = skip_mask ? false : array.null_bitmap_data() != nullptr; diff --git a/cpp/src/lists/copying/concatenate.cu b/cpp/src/lists/copying/concatenate.cu index 425b8f43065..c6ca56085c8 100644 --- a/cpp/src/lists/copying/concatenate.cu +++ b/cpp/src/lists/copying/concatenate.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020, NVIDIA CORPORATION. + * Copyright (c) 2020-2021, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -65,9 +65,9 @@ std::unique_ptr merge_offsets(std::vector const& colu std::for_each(columns.begin(), columns.end(), [&](lists_column_view const& c) { if (c.size() > 0) { // handle sliced columns - int const local_shift = shift - (c.offset() > 0 ? cudf::detail::get_value( - c.offsets(), c.offset(), stream.value()) - : 0); + int const local_shift = + shift - + (c.offset() > 0 ? cudf::detail::get_value(c.offsets(), c.offset(), stream) : 0); column_device_view offsets(c.offsets(), nullptr, nullptr); thrust::transform( rmm::exec_policy(stream), @@ -76,7 +76,7 @@ std::unique_ptr merge_offsets(std::vector const& colu d_merged_offsets.begin() + count, [local_shift] __device__(size_type offset) { return offset + local_shift; }); - shift += c.get_sliced_child(stream.value()).size(); + shift += c.get_sliced_child(stream).size(); count += c.size(); } }); @@ -110,7 +110,7 @@ std::unique_ptr concatenate( [&total_list_count, &children, stream](lists_column_view const& l) { // count total # of lists total_list_count += l.size(); - children.push_back(l.get_sliced_child(stream.value())); + children.push_back(l.get_sliced_child(stream)); }); auto data = cudf::detail::concatenate(children, stream, mr); diff --git a/cpp/src/quantiles/quantile.cu b/cpp/src/quantiles/quantile.cu index 3fda828d355..519feea3d7c 100644 --- a/cpp/src/quantiles/quantile.cu +++ b/cpp/src/quantiles/quantile.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2020, NVIDIA CORPORATION. + * Copyright (c) 2019-2021, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -63,8 +63,7 @@ struct quantile_functor { auto const type = is_fixed_point(input.type()) ? input.type() : data_type{type_to_id()}; - auto output = - make_fixed_width_column(type, q.size(), mask_state::UNALLOCATED, stream.value(), mr); + auto output = make_fixed_width_column(type, q.size(), mask_state::UNALLOCATED, stream, mr); if (output->size() == 0) { return output; } diff --git a/cpp/src/reductions/minmax.cu b/cpp/src/reductions/minmax.cu index 100c5ee7aec..58db90b600d 100644 --- a/cpp/src/reductions/minmax.cu +++ b/cpp/src/reductions/minmax.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020, NVIDIA CORPORATION. + * Copyright (c) 2020-2021, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -193,7 +193,7 @@ struct minmax_functor { auto maximum = new ScalarType(T{}, true, stream, mr); // copy dev_result to the output scalars device_single_thread(assign_min_max{dev_result.data(), minimum->data(), maximum->data()}, - stream.value()); + stream); return {std::unique_ptr(minimum), std::unique_ptr(maximum)}; } diff --git a/cpp/src/reductions/simple.cuh b/cpp/src/reductions/simple.cuh index ab780e70a17..f1c68e30dc9 100644 --- a/cpp/src/reductions/simple.cuh +++ b/cpp/src/reductions/simple.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2020, NVIDIA CORPORATION. + * Copyright (c) 2019-2021, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -204,7 +204,7 @@ struct cast_numeric_scalar_fn { auto result = std::make_unique>(ResultType{}, true, stream, mr); auto d_output = cudf::get_scalar_device_view(*result); cudf::detail::device_single_thread(assign_scalar_fn{d_input, d_output}, - stream.value()); + stream); return result; } diff --git a/cpp/src/search/search.cu b/cpp/src/search/search.cu index 5021af49baa..051d302c710 100644 --- a/cpp/src/search/search.cu +++ b/cpp/src/search/search.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2020, NVIDIA CORPORATION. + * Copyright (c) 2019-2021, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -243,7 +243,7 @@ struct multi_contains_dispatch { return result; } - auto hash_set = cudf::detail::unordered_multiset::create(needles, stream.value()); + auto hash_set = cudf::detail::unordered_multiset::create(needles, stream); auto device_hash_set = hash_set.to_device(); auto d_haystack_ptr = column_device_view::create(haystack, stream); diff --git a/cpp/src/stream_compaction/drop_duplicates.cu b/cpp/src/stream_compaction/drop_duplicates.cu index 7ccfdc5e6cb..25213502bf5 100644 --- a/cpp/src/stream_compaction/drop_duplicates.cu +++ b/cpp/src/stream_compaction/drop_duplicates.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2020, NVIDIA CORPORATION. + * Copyright (c) 2019-2021, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -143,7 +143,7 @@ column_view get_unique_ordered_indices(cudf::table_view const& keys, rmm::mr::get_current_device_resource()); // extract unique indices - auto device_input_table = cudf::table_device_view::create(keys, stream.value()); + auto device_input_table = cudf::table_device_view::create(keys, stream); if (cudf::has_nulls(keys)) { auto comp = row_equality_comparator( diff --git a/cpp/src/strings/findall.cu b/cpp/src/strings/findall.cu index 9b74e81bf9f..b9f2f7046a3 100644 --- a/cpp/src/strings/findall.cu +++ b/cpp/src/strings/findall.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2020, NVIDIA CORPORATION. + * Copyright (c) 2019-2021, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -119,12 +119,12 @@ std::unique_ptr
findall_re( rmm::cuda_stream_view stream = rmm::cuda_stream_default) { auto strings_count = strings.size(); - auto strings_column = column_device_view::create(strings.parent(), stream.value()); + auto strings_column = column_device_view::create(strings.parent(), stream); auto d_strings = *strings_column; auto d_flags = detail::get_character_flags_table(); // compile regex into device object - auto prog = reprog_device::create(pattern, d_flags, strings_count, stream.value()); + auto prog = reprog_device::create(pattern, d_flags, strings_count, stream); auto d_prog = *prog; int regex_insts = prog->insts_counts(); @@ -187,7 +187,7 @@ std::unique_ptr
findall_re( d_indices, findall_fn{d_strings, d_prog, column_index, d_find_counts}); // - results.emplace_back(make_strings_column(indices, stream.value(), mr)); + results.emplace_back(make_strings_column(indices, stream, mr)); } return std::make_unique
(std::move(results)); }