Replace stream.value() with stream for stream_view args

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<column> copy_if_else(
   constexpr int block_size = 256;
   cudf::detail::grid_1d grid{num_els, block_size, 1};
 
-  std::unique_ptr<column> out =
-    make_fixed_width_column(output_type,
-                            size,
-                            nullable ? mask_state::UNINITIALIZED : mask_state::UNALLOCATED,
-                            stream.value(),
-                            mr);
+  std::unique_ptr<column> 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);
 

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<string_view, MapIterator> {
                                      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);

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<column> scatter(
                                  std::move(child_column),
                                  cudf::UNKNOWN_NULL_COUNT,
                                  std::move(null_mask),
-                                 stream.value(),
+                                 stream,
                                  mr);
 }
 

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<cudf::column> copy_if_else(
                              std::move(chars_column),
                              null_count,
                              std::move(null_mask),
-                             stream.value(),
+                             stream,
                              mr);
 }
 

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<column> scatter(
   if (target.has_nulls()) null_mask = cudf::detail::copy_bitmask(target.parent(), stream, mr);
 
   // create string vectors
-  rmm::device_vector<string_view> target_vector =
-    create_string_vector_from_column(target, stream.value());
+  rmm::device_vector<string_view> 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());
 

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<column> compute_column(table_view const table,
     reinterpret_cast<const cudf::size_type*>(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();

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 <[email protected]>
@@ -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

cpp/src/column/column.cu

@@ -191,7 +191,7 @@ struct create_column_from_view {
   std::unique_ptr<column> 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 <typename ColumnType,
@@ -244,7 +244,7 @@ struct create_column_from_view {
   std::unique_ptr<column> 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 <typename ColumnType,
@@ -272,7 +272,7 @@ struct create_column_from_view {
                                std::move(children),
                                view.null_count(),
                                cudf::detail::copy_bitmask(view.null_mask(), begin, end, stream, mr),
-                               stream.value(),
+                               stream,
                                mr);
   }
 };

cpp/src/column/column_factories.cpp

@@ -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.
@@ -171,7 +171,7 @@ struct column_from_scalar_dispatch {
     auto output_column =
       make_fixed_width_column(value.type(), size, mask_state::UNALLOCATED, stream, mr);
     auto view = output_column->mutable_view();
-    detail::fill_in_place(view, 0, size, value, stream.value());
+    detail::fill_in_place(view, 0, size, value, stream);
     return output_column;
   }
 };

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<table> 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),

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<index_vector>(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<size_type>();
   decltype(_group_offsets->begin()) result_end;
 

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<column> dispatch_to_cudf_column::operator()<bool>(
   auto out_col = mask_to_bools(static_cast<bitmask_type*>(data.data()),
                                array.offset(),
                                array.offset() + array.length(),
-                               stream.value(),
+                               stream,
                                mr);
 
   auto const has_nulls = skip_mask ? false : array.null_bitmap_data() != nullptr;

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<column> merge_offsets(std::vector<lists_column_view> 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<size_type>(
-                                                          c.offsets(), c.offset(), stream.value())
-                                                      : 0);
+      int const local_shift =
+        shift -
+        (c.offset() > 0 ? cudf::detail::get_value<size_type>(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<column> merge_offsets(std::vector<lists_column_view> const& colu
         d_merged_offsets.begin<size_type>() + 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<column> 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);
 

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<StorageResult>()};
-    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; }
 

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<T>{dev_result.data(), minimum->data(), maximum->data()},
-                         stream.value());
+                         stream);
     return {std::unique_ptr<scalar>(minimum), std::unique_ptr<scalar>(maximum)};
   }
 

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<numeric_scalar<ResultType>>(ResultType{}, true, stream, mr);
     auto d_output = cudf::get_scalar_device_view(*result);
     cudf::detail::device_single_thread(assign_scalar_fn<InputType, ResultType>{d_input, d_output},
-                                       stream.value());
+                                       stream);
     return result;
   }