Fix GB300 and GH200

cpp/CMakeLists.txt

@@ -1520,6 +1520,7 @@ if(CUDF_BUILD_TESTUTIL)
 endif()
 
 # * build cudf_identify_stream_usage --------------------------------------------------------------
+message(STATUS "CUDF: Build with stream testing enabled: ${CUDF_BUILD_STREAMS_TEST_UTIL}")
 if(CUDF_BUILD_STREAMS_TEST_UTIL)
   # Libraries for stream-related testing. We build the library twice, one with STREAM_MODE_TESTING
   # on and one with it set to off. Each test will then be configured to use the appropriate library

cpp/include/cudf_test/column_wrapper.hpp

@@ -850,6 +850,7 @@ class strings_column_wrapper : public detail::column_wrapper {
       cudf::detail::make_device_uvector(offsets, stream, mr.get_output_mr()),
       rmm::device_buffer{},
       0);
+    cudf::test::get_default_stream().synchronize();
     wrapped =
       cudf::make_strings_column(num_strings, std::move(d_offsets), d_chars.release(), 0, {});
   }
@@ -1890,6 +1891,7 @@ class lists_column_wrapper : public detail::column_wrapper {
     // construct the list column
     wrapped = make_lists_column(
       cols.size(), std::move(offsets), std::move(data), null_count, std::move(null_mask));
+    cudf::test::get_default_stream().synchronize();
   }
 
   /**

cpp/src/interop/from_arrow_device.cu

@@ -198,7 +198,7 @@ dispatch_tuple_t dispatch_from_arrow_device::operator()<cudf::string_view>(
     // gather strings into output column
     auto out_col = cudf::make_strings_column(d_indices, stream, mr);
     owned.emplace_back(std::move(out_col));
-    cudf::detail::sync_stream(stream);
+    cudf::detail::sync_stream(stream);  // variadic_ptrs goes out of scope
     return std::make_tuple<column_view, owned_columns_t>(owned.front()->view(), std::move(owned));
   }
 

cpp/src/interop/from_arrow_host_strings.cu

@@ -114,7 +114,7 @@ std::unique_ptr<column> from_arrow_stringview(ArrowSchemaView const* schema,
       return {data, size};
     });
 
-  cudf::detail::sync_stream(stream);
+  cudf::detail::sync_stream(stream);  // variadic_ptrs goes out of scope
   return cudf::make_strings_column(d_indices, stream, mr);
 }
 

cpp/src/io/json/column_tree_construction.cu

@@ -20,6 +20,7 @@
 #include <cuda/iterator>
 #include <cuda/std/tuple>
 #include <cuda/stream>
+#include <thrust/fill.h>
 #include <thrust/for_each.h>
 #include <thrust/scan.h>
 #include <thrust/sort.h>
@@ -221,8 +222,8 @@ std::tuple<compressed_sparse_row, column_tree_properties> reduce_to_column_tree(
         }),
         cuda::std::plus<NodeIndexT>{});
     } else {
-      auto single_node = 1;
-      row_idx.set_element_async(1, single_node, stream);
+      auto exec_policy = rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref());
+      thrust::fill(exec_policy, row_idx.begin() + 1, row_idx.end(), NodeIndexT{1});
     }
 
 #ifdef CSR_DEBUG_PRINT

cpp/src/io/orc/stripe_enc.cu

@@ -1425,7 +1425,7 @@ void decimal_sizes_to_offsets(device_2dspan<rowgroup_rows const> rg_bounds,
   auto const num_blocks = elem_sizes.size() * rg_bounds.size().first;
   decimal_sizes_to_offsets_kernel<block_size>
     <<<num_blocks, block_size, 0, stream.get()>>>(rg_bounds, d_sizes);
-  stream.sync();
+  stream.sync();  // h_sizes goes out of scope
 }
 
 }  // namespace cudf::io::orc::detail

cpp/src/io/orc/writer_impl.cu

@@ -1513,7 +1513,7 @@ encoded_footer_statistics finish_statistic_blobs(Footer const& footer,
       file_blobs[i].assign(stat_begin, stat_end);
     }
 
-    stream.sync();
+    stream.sync();  // h_stat_chunks goes out of scope
     return {{}, std::move(file_blobs)};
   }
 
@@ -1596,7 +1596,7 @@ encoded_footer_statistics finish_statistic_blobs(Footer const& footer,
     file_blobs[i].assign(stat_begin, stat_end);
   }
 
-  stream.sync();
+  stream.sync();  // h_srcs, h_dsts, h_lens go out of scope
   return {std::move(stripe_blobs), std::move(file_blobs)};
 }
 
@@ -2042,7 +2042,7 @@ orc_table_view make_orc_table_view(table_view const& table,
     },
     stream);
 
-  stream.sync();
+  stream.sync();  // type_kinds goes out of scope
   return {std::move(orc_columns),
           std::move(d_orc_columns),
           str_col_indexes,

cpp/src/io/parquet/experimental/deletion_vectors_helpers.cu

@@ -132,7 +132,7 @@ std::unique_ptr<cudf::column> compute_row_index_column(
     row_indices_iter,
     row_indices_iter);
 
-  stream.sync();
+  stream.sync();  // row_group_span_offsets goes out of scope
   return std::make_unique<cudf::column>(cudf::data_type{cudf::type_id::UINT64},
                                         num_rows,
                                         std::move(row_indices),

cpp/src/row_operator/row_operators.cu

@@ -707,13 +707,15 @@ std::shared_ptr<preprocessed_table> preprocessed_table::create(
     }();
 
   auto const has_ranked_children = !transformed_columns.empty();
-  return create(transformed_input,
-                std::move(verticalized_col_depths),
-                std::move(transformed_columns),
-                new_column_order,
-                new_null_precedence,
-                has_ranked_children,
-                stream);
+  auto result                    = create(transformed_input,
+                       std::move(verticalized_col_depths),
+                       std::move(transformed_columns),
+                       new_column_order,
+                       new_null_precedence,
+                       has_ranked_children,
+                       stream);
+  cudf::detail::sync_stream(stream);  // new_column_order and new_null_precedence go out of scope
+  return result;
 }
 
 std::pair<std::shared_ptr<preprocessed_table>, std::shared_ptr<preprocessed_table>>
@@ -777,20 +779,22 @@ preprocessed_table::create(table_view const& lhs,
   auto const has_ranked_children_lhs = !transformed_columns_lhs.empty();
   auto const has_ranked_children_rhs = !transformed_columns_rhs.empty();
 
-  return {create(transformed_lhs,
-                 std::move(verticalized_col_depths_lhs),
-                 std::move(transformed_columns_lhs),
-                 new_column_order_lhs,
-                 new_null_precedence_lhs,
-                 has_ranked_children_lhs,
-                 stream),
-          create(transformed_rhs,
-                 std::move(verticalized_col_depths_rhs),
-                 std::move(transformed_columns_rhs),
-                 new_column_order_lhs,
-                 new_null_precedence_lhs,
-                 has_ranked_children_rhs,
-                 stream)};
+  auto result = std::pair{create(transformed_lhs,
+                                 std::move(verticalized_col_depths_lhs),
+                                 std::move(transformed_columns_lhs),
+                                 new_column_order_lhs,
+                                 new_null_precedence_lhs,
+                                 has_ranked_children_lhs,
+                                 stream),
+                          create(transformed_rhs,
+                                 std::move(verticalized_col_depths_rhs),
+                                 std::move(transformed_columns_rhs),
+                                 new_column_order_lhs,
+                                 new_null_precedence_lhs,
+                                 has_ranked_children_rhs,
+                                 stream)};
+  cudf::detail::sync_stream(stream);  // new_column_order_* and new_null_prec_* go out of scope
+  return result;
 }
 
 preprocessed_table::preprocessed_table(