Ensure that all tests launch kernels on cudf's default stream

[vyasr]

Description

Currently many of our tests are only stream-safe because libcudf runs everything on the default stream. This PR updates tests to ensure that any function that launches a kernel and supports passing streams will act on cudf's default stream even when it is not CUDA's default stream.

There are other aspects required for stream-safety that are not addressed in this PR. For instance, some of our tests make use of thrust::device_vector, and its initialization is implicitly always on the default stream. I'll work on that in a separate PR since that also requires some discussion with the team on what expectations a stream-based libcudf API could like like for consumers that make use of thrust (i.e. do we start requiring device syncs for such consumers?). There are also numerous tests that fail when swapping in an alternate default stream, indicating other potential dependencies on streams. I'll work through those remaining issues separately as well to limit the scope of this PR.

Checklist

• I am familiar with the Contributing Guidelines.
• New or existing tests cover these changes.
• The documentation is up to date with these changes.

[codecov]

Codecov Report

❗ No coverage uploaded for pull request base (branch-22.10@0528b38). Click here to learn what that means.
Patch has no changes to coverable lines.

Additional details and impacted files

@@               Coverage Diff               @@
##             branch-22.10   #11726   +/-   ##
===============================================
  Coverage                ?   87.54%           
===============================================
  Files                   ?      133           
  Lines                   ?    21769           
  Branches                ?        0           
===============================================
  Hits                    ?    19057           
  Misses                  ?     2712           
  Partials                ?        0           

Help us with your feedback. Take ten seconds to tell us how you rate us. Have a feature suggestion? Share it here.

☔ View full report at Codecov.
📢 Do you have feedback about the report comment? Let us know in this issue.

[robertmaynard]

git grep thrust::device returns the following places we need to change as well:

copying/detail_gather_tests.cu:  thrust::sequence(thrust::device, gather_map.begin(), gather_map.end());
io/json_type_cast_test.cu:  thrust::transform(thrust::device,
io/json_type_cast_test.cu:  thrust::transform(thrust::device,
io/json_type_cast_test.cu:  thrust::transform(thrust::device,
iterator/optional_iterator_test_numeric.cu:  thrust::transform(thrust::device,
iterator/optional_iterator_test_numeric.cu:  auto result = thrust::reduce(thrust::device, results.begin(), results.end(), T_output{});
join/conditional_join_tests.cu:    thrust::device_vector<thrust::pair<cudf::size_type, cudf::size_type>> result_pairs(
join/conditional_join_tests.cu:    thrust::device_vector<thrust::pair<cudf::size_type, cudf::size_type>> reference_pairs(
join/conditional_join_tests.cu:    thrust::transform(thrust::device,
join/conditional_join_tests.cu:    thrust::transform(thrust::device,
join/conditional_join_tests.cu:    thrust::sort(thrust::device, result_pairs.begin(), result_pairs.end());
join/conditional_join_tests.cu:    thrust::sort(thrust::device, reference_pairs.begin(), reference_pairs.end());
join/conditional_join_tests.cu:      thrust::device, reference_pairs.begin(), reference_pairs.end(), result_pairs.begin()));
join/conditional_join_tests.cu:    thrust::sort(thrust::device, result->begin(), result->end());
join/conditional_join_tests.cu:    thrust::sort(thrust::device, reference->begin(), reference->end());
join/conditional_join_tests.cu:    EXPECT_TRUE(thrust::equal(thrust::device, result->begin(), result->end(), reference->begin()));
reductions/segmented_reduction_tests.cpp:  auto const d_offsets = thrust::device_vector<size_type>(offsets);
reductions/segmented_reduction_tests.cpp:  auto const d_offsets = thrust::device_vector<size_type>(offsets);
reductions/segmented_reduction_tests.cpp:  auto const d_offsets = thrust::device_vector<size_type>(offsets);
reductions/segmented_reduction_tests.cpp:  auto const d_offsets = thrust::device_vector<size_type>(offsets);
reductions/segmented_reduction_tests.cpp:  auto const d_offsets = thrust::device_vector<size_type>(offsets);
reductions/segmented_reduction_tests.cpp:  auto const d_offsets = thrust::device_vector<size_type>(offsets);
reductions/segmented_reduction_tests.cpp:  auto const d_offsets = thrust::device_vector<size_type>(offsets);
reductions/segmented_reduction_tests.cpp:  auto const d_offsets = thrust::device_vector<size_type>(offsets);
reductions/segmented_reduction_tests.cpp:  auto const d_offsets = thrust::device_vector<size_type>(offsets);
reductions/segmented_reduction_tests.cpp:  auto const d_offsets = thrust::device_vector<size_type>(offsets);
reductions/segmented_reduction_tests.cpp:  auto const d_offsets = thrust::device_vector<size_type>(offsets);
reductions/segmented_reduction_tests.cpp:  auto const d_offsets = thrust::device_vector<size_type>(offsets);
reductions/segmented_reduction_tests.cpp:  auto const d_offsets = thrust::device_vector<size_type>(offsets);
reductions/segmented_reduction_tests.cpp:  auto const d_offsets = thrust::device_vector<size_type>(offsets);
reductions/segmented_reduction_tests.cpp:  auto const d_offsets = thrust::device_vector<size_type>(offsets);
reductions/segmented_reduction_tests.cpp:  auto const d_offsets = thrust::device_vector<size_type>(offsets);
reductions/segmented_reduction_tests.cpp:    auto const d_offsets = thrust::device_vector<size_type>(offsets);
reductions/segmented_reduction_tests.cpp:    auto const d_offsets = thrust::device_vector<size_type>(offsets);
reductions/segmented_reduction_tests.cpp:    auto const d_offsets = thrust::device_vector<size_type>(offsets);
reductions/segmented_reduction_tests.cpp:    auto const d_offsets = thrust::device_vector<size_type>(offsets);
reductions/segmented_reduction_tests.cpp:    auto const d_offsets = thrust::device_vector<size_type>(offsets);
reductions/segmented_reduction_tests.cpp:    auto const d_offsets = thrust::device_vector<size_type>(offsets);
reductions/segmented_reduction_tests.cpp:  auto const d_offsets = thrust::device_vector<size_type>(offsets);
strings/factories_test.cu:  thrust::transform(thrust::device,
transform/row_bit_count_test.cu:    thrust::device, ints_view.begin<int32_t>(), ints_view.end<int32_t>(), thrust::identity{});
transform/row_bit_count_test.cu:  thrust::tabulate(thrust::device,
transform/row_bit_count_test.cu:  thrust::fill_n(thrust::device,
utilities/column_utilities.cu:  EXPECT_TRUE(thrust::equal(thrust::device, typed_lhs, typed_lhs + size_bytes, typed_rhs));
utilities_tests/span_tests.cu:  auto d_thrust_vector = thrust::device_vector<int>(1);

[vyasr]

A lot of those are thrust::device_vector, which is fine, but I do see some that are actually the execution policy. I'll have to investigate and see why my code didn't detect those.

That also raises another question, though. If someone does something like thrust::host_vector<T> x = my_device_vector;, thrust is going to synchronously copy that, right? Assuming that is correct, implicit device_vector<->host_vector copies in thrust are not stream safe (if cudf's default stream is switched away from the CUDA default stream), which is definitely something that thrust users rely on. That's something we need to keep in mind. Once I extend my code to also check APIs like cudaMemcpy it should start catching cases like this, but I'm not sure what we want to do about it.

[vyasr]

I was omitting to account for cudaStreamLegacy. Now I should get the rest.

[robertmaynard]

A lot of those are thrust::device_vector, which is fine, but I do see some that are actually the execution policy. I'll have to investigate and see why my code didn't detect those.

That also raises another question, though. If someone does something like thrust::host_vector<T> x = my_device_vector;, thrust is going to synchronously copy that, right? Assuming that is correct, implicit device_vector<->host_vector copies in thrust are not stream safe (if cudf's default stream is switched away from the CUDA default stream), which is definitely something that thrust users rely on. That's something we need to keep in mind. Once I extend my code to also check APIs like cudaMemcpy it should start catching cases like this, but I'm not sure what we want to do about it.

As far as I understand your assumptions are correct and those will copies will occur on the default CUDA stream.

[ttnghia]

Curious if using the default stream for tests has been documented in the dev guide.

[vyasr]

@gpucibot merge

[vyasr]

@ttnghia I don't think so, but at this point I'm not even sure that's the best place for it. Since we're thinking about how to expose streams publicly, this is going to end up being a user-facing issue that needs documentation not just something for devs.