Refactor c2h gen to ensure teardown before main (#7067)

Fixes: #7063

Author: bernhardmgruber

c2h/generators.cu

@@ -9,6 +9,8 @@
 #include <thrust/iterator/transform_iterator.h>
 #include <thrust/tabulate.h>
 
+#include <cuda/std/optional>
+
 #include <c2h/bfloat16.cuh>
 #include <c2h/custom_type.h>
 #include <c2h/detail/generators.cuh>
@@ -44,29 +46,77 @@ struct i_to_rnd_t
 };
 #endif // !C2H_HAS_CURAND
 
-void generator_t::generate()
+class generator_t
 {
+public:
+  generator_t()
+  {
+#if C2H_HAS_CURAND
+    curandCreateGenerator(&m_gen, CURAND_RNG_PSEUDO_DEFAULT);
+#endif
+  }
+
+  ~generator_t()
+  {
+#if C2H_HAS_CURAND
+    curandDestroyGenerator(m_gen);
+#endif
+  }
+
+  float* prepare_random_generator(seed_t seed, std::size_t num_items)
+  {
+    m_distribution.resize(num_items);
+
 #if C2H_HAS_CURAND
-  curandGenerateUniform(m_gen, thrust::raw_pointer_cast(m_distribution.data()), m_distribution.size());
+    curandSetPseudoRandomGeneratorSeed(m_gen, seed.get());
 #else
-  thrust::tabulate(device_policy, m_distribution.begin(), m_distribution.end(), i_to_rnd_t{m_re});
-  m_re.discard(m_distribution.size());
+    m_gen.seed(seed.get());
 #endif
-}
 
-float* generator_t::prepare_random_generator(seed_t seed, std::size_t num_items)
-{
-  m_distribution.resize(num_items);
+    generate();
 
+    return thrust::raw_pointer_cast(m_distribution.data());
+  }
+
+  // re-fills the currently held distribution vector with new random values
+  void generate()
+  {
+#if C2H_HAS_CURAND
+    curandGenerateUniform(m_gen, thrust::raw_pointer_cast(m_distribution.data()), m_distribution.size());
+#else
+    thrust::tabulate(device_policy, m_distribution.begin(), m_distribution.end(), i_to_rnd_t{m_gen});
+    m_gen.discard(m_distribution.size());
+#endif
+  }
+
+private:
 #if C2H_HAS_CURAND
-  curandSetPseudoRandomGeneratorSeed(m_gen, seed.get());
+  curandGenerator_t
 #else
-  m_re.seed(seed.get());
+  thrust::default_random_engine
 #endif
+    m_gen;
+  c2h::device_vector<float> m_distribution;
+};
 
-  generate();
+// global generator state
+cuda::std::optional<generator_t> generator;
 
-  return thrust::raw_pointer_cast(m_distribution.data());
+void init_generator()
+{
+  _CCCL_VERIFY(!generator.has_value(), "");
+  generator.emplace();
+}
+
+float* prepare_random_data(seed_t seed, std::size_t num_items)
+{
+  return generator.value().prepare_random_generator(seed, num_items);
+}
+
+void cleanup_generator()
+{
+  _CCCL_VERIFY(generator.has_value(), "");
+  generator.reset();
 }
 
 struct random_to_custom_t
@@ -94,7 +144,7 @@ void gen_custom_type_state(
   std::size_t element_size)
 {
   // FIXME(bgruber): implement min/max handling for custom_type_state_t
-  float* d_in = generator.prepare_random_generator(seed, elements * 2);
+  float* d_in = prepare_random_data(seed, elements * 2);
   thrust::for_each(device_policy,
                    thrust::counting_iterator<std::size_t>{0},
                    thrust::counting_iterator<std::size_t>{elements},

c2h/generators_gen_values.cu

@@ -16,7 +16,7 @@ namespace c2h::detail
 template <typename T>
 void gen_values_between(seed_t seed, ::cuda::std::span<T> data, T min, T max)
 {
-  const auto* dist = generator.prepare_random_generator(seed, data.size());
+  const auto* dist = prepare_random_data(seed, data.size());
   thrust::transform(device_policy, dist, dist + data.size(), data.begin(), random_to_item_t<T>(min, max));
 }
 

c2h/generators_vector.cu

@@ -53,7 +53,7 @@ struct random_to_vec_item_t
     template <>                                                                                                   \
     void gen_values_between(seed_t seed, ::cuda::std::span<T> data, T min, T max)                                 \
     {                                                                                                             \
-      const auto* dist = generator.prepare_random_generator(seed, data.size());                                   \
+      const auto* dist = prepare_random_data(seed, data.size());                                                  \
       auto op          = random_to_vec_item_t<T, ::cuda::std::tuple_size_v<T>>{min, max, dist, data.data()};      \
       thrust::for_each(                                                                                           \
         device_policy, thrust::counting_iterator<size_t>{0}, thrust::counting_iterator<size_t>{data.size()}, op); \

c2h/include/c2h/catch2_main.h

@@ -5,7 +5,7 @@
 
 #include <thrust/detail/config/device_system.h>
 
-#include <iostream>
+#include <c2h/detail/generators.cuh>
 
 //! @file
 //! This file includes a custom Catch2 main function. When CMake is configured to build each test as a separate
@@ -43,6 +43,9 @@ int main(int argc, char* argv[])
 
   set_device(device_id);
 #  endif // THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
-  return session.run();
+  c2h::detail::init_generator();
+  const auto ret = session.run();
+  c2h::detail::cleanup_generator();
+  return ret;
 }
 #endif // C2H_CONFIG_MAIN

c2h/include/c2h/detail/generators.cuh

@@ -4,50 +4,17 @@
 #include <cuda/std/complex>
 
 #include <c2h/generators.h>
-#include <c2h/vector.h>
-
-#if C2H_HAS_CURAND
-#  include <curand.h>
-#else
-#  include <thrust/random.h>
-#endif
 
 namespace c2h::detail
 {
-class generator_t
-{
-public:
-  generator_t()
-  {
-#if C2H_HAS_CURAND
-    curandCreateGenerator(&m_gen, CURAND_RNG_PSEUDO_DEFAULT);
-#endif
-  }
-
-  ~generator_t()
-  {
-#if C2H_HAS_CURAND
-    curandDestroyGenerator(m_gen);
-#endif
-  }
-
-  // sets the seed and resizes the distribution vector, fills it by calling generate(), and returns a pointer the start
-  // of the data
-  float* prepare_random_generator(seed_t seed, std::size_t num_items);
+// called once from main to set up the generator state
+void init_generator();
 
-  // re-fills the currently held distribution vector with new random values
-  void generate();
-
-private:
-#if C2H_HAS_CURAND
-  curandGenerator_t m_gen;
-#else
-  thrust::default_random_engine m_re;
-#endif
-  c2h::device_vector<float> m_distribution;
-};
+// sets the seed and resizes the distribution vector, fills it, and returns a pointer the start of the data
+float* prepare_random_data(seed_t seed, std::size_t num_items);
 
-inline generator_t generator;
+// called once before main returns to clean up the generator state
+void cleanup_generator();
 
 template <typename T, bool = ::cuda::is_floating_point_v<T>>
 struct random_to_item_t