add tests for legacy and per-thread default streams
MODIFY
testing/CMakeLists.txt
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

ADD
testing/cuda/stream_legacy.cu
+#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);

ADD
testing/cuda/stream_per_thread.cmake
+# This test should always use per-thread streams on NVCC.
+set_target_properties(${test_target} PROPERTIES
+ COMPILE_OPTIONS
+ $<$<AND:$<COMPILE_LANGUAGE:CUDA>,$<CUDA_COMPILER_ID:NVIDIA>>:--default-stream=per-thread>
+)
+
+# NVC++ does not have an equivalent option, and will always
+# use the global stream by default.
+if (CMAKE_CUDA_COMPILER_ID STREQUAL "Feta")
+ set_tests_properties(${test_target} PROPERTIES WILL_FAIL ON)
+endif()

ADD
testing/cuda/stream_per_thread.cu
+#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);

ADD
testing/cuda/stream_per_thread.mk
+CUDACC_FLAGS += --default-stream per-thread

add tests for legacy and per-thread default streams