Fix issue with newer thrust.

Implicit conversions from `nvbench::cuda_stream` -> `cudaStream_t` no longer work when creating Thrust execution policies, as the `cudaStream_t` overload of `thrust::device.on(...)` was removed and replaced with `cuda::stream_ref`.

Explicitly pulling the `cudaStream_t` out of the `nvbench::cuda_stream` and passing it to `on` so that the implicit constructor of `stream_ref` can handle it seems to be the only way to fix this while preserving support back to CTK 12.0 (`stream_ref` was added in CTK 12.3).

Author: alliepiper

examples/exec_tag_sync.cu

@@ -53,7 +53,7 @@ void sequence_bench(nvbench::state &state)
 
   // nvbench::exec_tag::sync indicates that this will implicitly sync:
   state.exec(nvbench::exec_tag::sync, [&data](nvbench::launch &launch) {
-    thrust::sequence(thrust::device.on(launch.get_stream()), data.begin(), data.end());
+    thrust::sequence(thrust::device.on(launch.get_stream().get_stream()), data.begin(), data.end());
   });
 }
 NVBENCH_BENCH(sequence_bench);

examples/exec_tag_timer.cu

@@ -57,7 +57,7 @@ void mod2_inplace(nvbench::state &state)
                (void)num_values; // clang thinks this is unused...
 
                // Reset working data:
-               thrust::copy(thrust::device.on(launch.get_stream()),
+               thrust::copy(thrust::device.on(launch.get_stream().get_stream()),
                             input.cbegin(),
                             input.cend(),
                             data.begin());