add tests for legacy and per-thread default streams

testing/CMakeLists.txt

@@ -111,6 +111,14 @@ function(thrust_add_test target_name_var test_name test_src thrust_target)
   if (("OMP" IN_LIST config_systems) OR ("TBB" IN_LIST config_systems))
     set_tests_properties(${test_target} PROPERTIES RUN_SERIAL ON)
   endif()
+
+  # Check for per-test script. Script will be included in the current scope
+  # to allow custom property modifications.
+  get_filename_component(test_cmake_script "${test_src}" NAME_WLE)
+  set(test_cmake_script "${CMAKE_CURRENT_LIST_DIR}/${test_cmake_script}.cmake")
+  if (EXISTS "${test_cmake_script}")
+    include("${test_cmake_script}")
+  endif()
 endfunction()
 
 file(GLOB test_srcs

testing/cuda/stream_legacy.cu

@@ -0,0 +1,21 @@
+#include <unittest/unittest.h>
+#include <thrust/execution_policy.h>
+#include <thrust/system/cuda/detail/util.h>
+
+#include <thread>
+
+void verify_stream()
+{
+  auto exec = thrust::device;
+  auto stream = thrust::cuda_cub::stream(exec);
+  ASSERT_EQUAL(stream, cudaStreamLegacy);
+}
+
+void TestLegacyDefaultStream()
+{
+  verify_stream();
+
+  std::thread t(verify_stream);
+  t.join();
+}
+DECLARE_UNITTEST(TestLegacyDefaultStream);

testing/cuda/stream_per_thread.cmake

@@ -0,0 +1,6 @@
+# This test should always use per-thread streams on NVCC.
+# NVC++ does not have an equivalent option.
+set_target_properties(${test_target} PROPERTIES
+  COMPILE_OPTIONS
+    $<$<AND:$<COMPILE_LANGUAGE:CUDA>,$<CUDA_COMPILER_ID:NVIDIA>>:--default-stream=per-thread>
+)

testing/cuda/stream_per_thread.cu

@@ -0,0 +1,21 @@
+#include <unittest/unittest.h>
+#include <thrust/execution_policy.h>
+#include <thrust/system/cuda/detail/util.h>
+
+#include <thread>
+
+void verify_stream()
+{
+  auto exec = thrust::device;
+  auto stream = thrust::cuda_cub::stream(exec);
+  ASSERT_EQUAL(stream, cudaStreamPerThread);
+}
+
+void TestPerThreadDefaultStream()
+{
+  verify_stream();
+
+  std::thread t(verify_stream);
+  t.join();
+}
+DECLARE_UNITTEST(TestPerThreadDefaultStream);

testing/cuda/stream_per_thread.mk

@@ -0,0 +1 @@
+CUDACC_FLAGS += --default-stream per-thread