WIP Adding async scan algorithms, iterate on async testing.

testing/CMakeLists.txt

@@ -151,6 +151,7 @@ foreach(thrust_target IN LISTS THRUST_TARGETS)
 endforeach()
 
 # Add specialized tests:
+add_subdirectory(async)
 add_subdirectory(cpp)
 add_subdirectory(cuda)
 add_subdirectory(omp)

testing/async/CMakeLists.txt

@@ -0,0 +1,82 @@
+# The async tests perform a large amount of codegen, making them expensive to
+# build and test. To keep compilation and runtimes manageable, the tests are
+# broken up into many files per algorithm to enable parallelism during
+# compilation and testing. The structure of these test directories are:
+#
+# thrust/testing/async/<algorithm_name>/<unit_test>.cu
+#
+# These generate executables and CTest tests named
+# ${config_prefix}.test.async.<algorithm_name>.<unit_test>.
+
+# The async tests only support CUDA enabled configs. Create a list of valid
+# thrust targets:
+set(cuda_configs)
+foreach(thrust_target IN LISTS THRUST_TARGETS)
+  thrust_get_target_property(config_device ${thrust_target} DEVICE)
+  if (config_device STREQUAL CUDA)
+    list(APPEND cuda_configs ${thrust_target})
+  endif()
+endforeach()
+
+list(LENGTH cuda_configs num_cuda_configs)
+if (num_cuda_configs EQUAL 0)
+  return() # No valid configs found, nothing to do.
+endif()
+
+# Process a single algorithm directory, adding all .cu/cpp files as tests for
+# each valid backend. algo_name is the name of the subdir (<algorithm_name>
+# above) and is used for naming the executable/targets.
+function(thrust_add_async_test_dir algo_name)
+  file(GLOB test_srcs
+    RELATIVE "${CMAKE_CURRENT_LIST_DIR}"
+    CONFIGURE_DEPENDS
+    "${algo_name}/*.cu"
+    "${algo_name}/*.cpp"
+  )
+
+  # Per-algorithm, all-config metatarget: thrust.all.async.[algo].all
+  set(algo_meta_target thrust.all.async.${algo_name}.all)
+  add_custom_target(${algo_meta_target})
+
+  foreach(thrust_target IN LISTS cuda_configs)
+    # Per-algorithm, per-config metatarget: thrust.[config].async.[algo].all
+    thrust_get_target_property(config_prefix ${thrust_target} PREFIX)
+    set(algo_config_meta_target ${config_prefix}.async.${algo_name}.all)
+    add_custom_target(${algo_config_meta_target})
+    add_dependencies(${algo_meta_target} ${algo_config_meta_target})
+
+    foreach(test_src IN LISTS test_srcs)
+      get_filename_component(test_name "${test_src}" NAME_WLE)
+      string(PREPEND test_name async.${algo_name}.)
+      message("Adding test ${config_prefix}.${test_name} at ${test_src}.")
+
+      thrust_add_test(test_target ${test_name} "${test_src}" ${thrust_target})
+      if(THRUST_ENABLE_TESTS_WITH_RDC)
+        thrust_enable_rdc_for_cuda_target(${test_target})
+      endif()
+
+      add_dependencies(${algo_config_meta_target} ${test_target})
+    endforeach()
+  endforeach()
+endfunction()
+
+# Grab all algorithm subdirectories:
+set(test_dirs)
+file(GLOB contents
+  CONFIGURE_DEPENDS
+  "${CMAKE_CURRENT_LIST_DIR}/*"
+)
+message("Contents: ${contents}")
+foreach(test_dir IN LISTS contents)
+  if(IS_DIRECTORY "${test_dir}")
+    list(APPEND test_dirs "${test_dir}")
+    message("Found test_dir: ${test_dir}")
+  endif()
+endforeach()
+
+# Process all test dirs:
+foreach(test_dir IN LISTS test_dirs)
+  get_filename_component(algo_name "${test_dir}" NAME_WLE)
+  message("Processing algorithm ${algo_name} from ${test_dir}.")
+  thrust_add_async_test_dir(${algo_name})
+endforeach()

testing/async/exclusive_scan/basic.cu

@@ -0,0 +1,42 @@
+#pragma once
+
+#include <thrust/detail/config.h>
+#include <thrust/detail/cpp14_required.h>
+
+#if THRUST_CPP_DIALECT >= 2014
+
+#include <async/test_policy_overloads.h>
+
+#include <async/exclusive_scan/mixin.h>
+
+template <typename input_value_type,
+          typename output_value_type   = input_value_type,
+          typename initial_value_type  = input_value_type,
+          typename alternate_binary_op = thrust::maximum<>>
+struct basic_invoker
+    : testing::async::mixin::input::device_vector<input_value_type>
+    , testing::async::mixin::output::device_vector<output_value_type>
+    , testing::async::exclusive_scan::mixin::postfix_args::
+        all_overloads<initial_value_type, alternate_binary_op>
+    , testing::async::exclusive_scan::mixin::invoke_reference::
+        host_synchronous<input_value_type, output_value_type>
+    , testing::async::exclusive_scan::mixin::invoke_async::basic
+    , testing::async::mixin::compare_outputs::assert_maybe_fuzzy_equal_quiet
+{
+  static std::string description()
+  {
+    return "basic invocation with device vectors";
+  }
+};
+
+template <typename T>
+struct TestBasic
+{
+  void operator()(std::size_t num_values) const
+  {
+    testing::async::test_policy_overloads<basic_invoker<T>>::run(num_values);
+  }
+};
+DECLARE_GENERIC_SIZED_UNITTEST_WITH_TYPES(TestBasic, NumericTypes);
+
+#endif // C++14

testing/async/exclusive_scan/counting_iterator.cu

@@ -0,0 +1,49 @@
+#pragma once
+
+#include <thrust/detail/config.h>
+#include <thrust/detail/cpp14_required.h>
+
+#if THRUST_CPP_DIALECT >= 2014
+
+#include <async/test_policy_overloads.h>
+
+#include <async/exclusive_scan/mixin.h>
+
+#include <algorithm>
+#include <limits>
+
+template <typename input_value_type,
+          typename output_value_type   = input_value_type,
+          typename initial_value_type  = input_value_type,
+          typename alternate_binary_op = thrust::maximum<>>
+struct invoker
+    : testing::async::mixin::input::counting_iterator<input_value_type>
+    , testing::async::mixin::output::device_vector<output_value_type>
+    , testing::async::exclusive_scan::mixin::postfix_args::
+        all_overloads<initial_value_type, alternate_binary_op>
+    , testing::async::exclusive_scan::mixin::invoke_reference::
+        host_synchronous<input_value_type, output_value_type>
+    , testing::async::exclusive_scan::mixin::invoke_async::basic
+    , testing::async::mixin::compare_outputs::assert_maybe_fuzzy_equal_quiet
+{
+  static std::string description()
+  {
+    return "fancy input iterator (counting_iterator)";
+  }
+};
+
+template <typename T>
+struct TestCountingIterator
+{
+  void operator()(std::size_t num_values) const
+  {
+    num_values = unittest::truncate_to_max_representable<T>(num_values);
+    testing::async::test_policy_overloads<invoker<T>>::run(num_values);
+  }
+};
+// Use built-in types only, counting_iterator doesn't seem to be compatible with
+// the custom_numeric.
+DECLARE_GENERIC_SIZED_UNITTEST_WITH_TYPES(TestCountingIterator,
+                                          BuiltinNumericTypes);
+
+#endif // C++14

testing/async/exclusive_scan/discard_output.cu

@@ -0,0 +1,41 @@
+#pragma once
+
+#include <thrust/detail/config.h>
+#include <thrust/detail/cpp14_required.h>
+
+#if THRUST_CPP_DIALECT >= 2014
+
+#include <async/test_policy_overloads.h>
+
+#include <async/exclusive_scan/mixin.h>
+
+// Compilation test with discard iterators. No runtime validation is actually
+// performed, other than testing whether the algorithm completes without
+// exception.
+
+template <typename input_value_type,
+          typename initial_value_type  = input_value_type,
+          typename alternate_binary_op = thrust::maximum<>>
+struct discard_invoker
+    : testing::async::mixin::input::device_vector<input_value_type>
+    , testing::async::mixin::output::discard_iterator
+    , testing::async::exclusive_scan::mixin::postfix_args::
+        all_overloads<initial_value_type, alternate_binary_op>
+    , testing::async::mixin::invoke_reference::noop
+    , testing::async::exclusive_scan::mixin::invoke_async::basic
+    , testing::async::mixin::compare_outputs::noop
+{
+  static std::string description() { return "discard output"; }
+};
+
+template <typename T>
+struct TestDiscard
+{
+  void operator()(std::size_t num_values) const
+  {
+    testing::async::test_policy_overloads<discard_invoker<T>>::run(num_values);
+  }
+};
+DECLARE_GENERIC_SIZED_UNITTEST_WITH_TYPES(TestDiscard, NumericTypes);
+
+#endif // C++14

testing/async/exclusive_scan/in_place.cu

@@ -0,0 +1,42 @@
+#pragma once
+
+#include <thrust/detail/config.h>
+#include <thrust/detail/cpp14_required.h>
+
+#if THRUST_CPP_DIALECT >= 2014
+
+#include <async/test_policy_overloads.h>
+
+#include <async/exclusive_scan/mixin.h>
+
+template <typename input_value_type,
+          typename initial_value_type  = input_value_type,
+          typename alternate_binary_op = thrust::maximum<>>
+struct basic_inplace_invoker
+    : testing::async::mixin::input::device_vector<input_value_type>
+    , testing::async::mixin::output::device_vector_reuse_input<input_value_type>
+    , testing::async::exclusive_scan::mixin::postfix_args::
+        all_overloads<initial_value_type, alternate_binary_op>
+    , testing::async::exclusive_scan::mixin::invoke_reference::host_synchronous<
+        input_value_type>
+    , testing::async::exclusive_scan::mixin::invoke_async::basic
+    , testing::async::mixin::compare_outputs::assert_maybe_fuzzy_equal_quiet
+{
+  static std::string description()
+  {
+    return "basic in-place invocation with device vectors";
+  }
+};
+
+template <typename T>
+struct TestBasicInPlace
+{
+  void operator()(std::size_t num_values) const
+  {
+    using invoker = basic_inplace_invoker<T>;
+    testing::async::test_policy_overloads<invoker>::run(num_values);
+  }
+};
+DECLARE_GENERIC_SIZED_UNITTEST_WITH_TYPES(TestBasicInPlace, NumericTypes);
+
+#endif // C++14

testing/async/exclusive_scan/mixed_types.cu

@@ -0,0 +1,103 @@
+#pragma once
+
+#include <thrust/detail/config.h>
+#include <thrust/detail/cpp14_required.h>
+
+#if THRUST_CPP_DIALECT >= 2014
+
+#include <async/test_policy_overloads.h>
+
+#include <async/exclusive_scan/mixin.h>
+
+// Test using mixed int/float types for:
+// - input_value_type       | (int, float)
+// - output_value_type      | (int, float)
+// - initial_value_type     | (int, float, <none>)
+// - thrust::plus<T> T-type | (int, float, void>)
+//
+// The initial_value_type and thrust::plus<T> types are covered by the
+// mixin::postfix_args::scan_mixed_types_overloads component.
+//
+// The testing/scan.cu TestMixedTypes test spells out the expected behavior,
+// which is defined by https://wg21.link/P0571.
+
+template <typename value_type>
+struct mixed_type_input_generator
+{
+  using input_type = thrust::device_vector<value_type>;
+
+  static input_type generate_input(std::size_t num_values)
+  {
+    input_type input(num_values);
+    thrust::sequence(input.begin(),
+                     input.end(),
+                     // fractional values are chosen deliberately to test
+                     // casting orders and accumulator types:
+                     static_cast<value_type>(1.5),
+                     static_cast<value_type>(1));
+    return input;
+  }
+};
+
+// A fractional value is used to ensure that a different result is obtained when
+// using float vs. int.
+template <typename value_type>
+struct mixed_types_postfix_args
+{
+  using postfix_args_type = std::tuple<         // Overloads to test:
+    std::tuple<>,                               // - no extra args
+    std::tuple<value_type>,                     // - initial_value
+    std::tuple<value_type, thrust::plus<>>,     // - initial_value, plus<>
+    std::tuple<value_type, thrust::plus<int>>,  // - initial_value, plus<int>
+    std::tuple<value_type, thrust::plus<float>> // - initial_value, plus<float>
+    >;
+
+  static postfix_args_type generate_postfix_args()
+  {
+    return {{},
+            {static_cast<value_type>(5.5)},
+            {static_cast<value_type>(5.5), thrust::plus<>{}},
+            {static_cast<value_type>(5.5), thrust::plus<int>{}},
+            {static_cast<value_type>(5.5), thrust::plus<float>{}}};
+  }
+};
+
+template <typename input_value_type,
+          typename output_value_type,
+          typename initial_value_type>
+struct invoker
+    : mixed_type_input_generator<input_value_type>
+    , testing::async::mixin::output::device_vector<output_value_type>
+    , mixed_types_postfix_args<initial_value_type>
+    , testing::async::exclusive_scan::mixin::invoke_reference::
+        host_synchronous<input_value_type, output_value_type>
+    , testing::async::exclusive_scan::mixin::invoke_async::basic
+    , testing::async::mixin::compare_outputs::assert_maybe_fuzzy_equal_quiet
+{
+  static std::string description()
+  {
+    return "mixed input/output/initial type tests";
+  }
+};
+
+void TestScanMixedTypes()
+{
+  // 10 values are enough to check the behavior we want to test while staying
+  // small enough to reason about.
+  constexpr std::size_t num_values = 10;
+
+  // invoker template params are input_value_type, output_vt, initial_vt:
+  using testing::async::test_policy_overloads;
+  test_policy_overloads<invoker<int, int, int>>::run(num_values);
+  test_policy_overloads<invoker<int, int, float>>::run(num_values);
+  test_policy_overloads<invoker<int, float, int>>::run(num_values);
+  test_policy_overloads<invoker<int, float, float>>::run(num_values);
+  test_policy_overloads<invoker<float, int, int>>::run(num_values);
+  test_policy_overloads<invoker<float, int, float>>::run(num_values);
+  test_policy_overloads<invoker<float, float, int>>::run(num_values);
+  // We all float down here
+  test_policy_overloads<invoker<float, float, float>>::run(num_values);
+}
+DECLARE_UNITTEST(TestScanMixedTypes);
+
+#endif // C++14