Merge branch 'main' of github.com:NVIDIA/cccl into fea/use-sccache-build-cluster

Author: trxcllnt

.github/workflows/build-rapids.yml

@@ -64,8 +64,8 @@ jobs:
       matrix:
         include:
           - { cuda: '13.0', libs: 'rmm kvikio cudf cudf_kafka' }
-          - { cuda: '13.0', libs: 'rmm ucxx raft cuvs cumlprims_mg cuml' }
-          - { cuda: '13.0', libs: 'rmm ucxx raft cugraph'                }
+          - { cuda: '13.0', libs: 'rmm ucxx raft cuvs cuml'    }
+          - { cuda: '13.0', libs: 'rmm ucxx raft cugraph'      }
     permissions:
       id-token: write
       contents: read
@@ -106,7 +106,6 @@ jobs:
           # RAPIDS_cugraph_GIT_REPO: '{"upstream": "rapidsai", "tag": "main"}'
           # RAPIDS_cugraph_gnn_GIT_REPO: '{"upstream": "rapidsai", "tag": "main"}'
           # RAPIDS_cuml_GIT_REPO: '{"upstream": "rapidsai", "tag": "main"}'
-          # RAPIDS_cumlprims_mg_GIT_REPO: '{"upstream": "rapidsai", "tag": "main"}'
           # RAPIDS_cuvs_GIT_REPO: '{"upstream": "rapidsai", "tag": "main"}'
           # RAPIDS_kvikio_GIT_REPO: '{"upstream": "rapidsai", "tag": "main"}'
           # RAPIDS_raft_GIT_REPO: '{"upstream": "rapidsai", "tag": "main"}'

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

ci/rapids/cuda13.0-conda/devcontainer.json

@@ -29,7 +29,6 @@
     "RAPIDS_cugraph_GIT_REPO": "${localEnv:RAPIDS_cugraph_GIT_REPO}",
     "RAPIDS_cugraph_gnn_GIT_REPO": "${localEnv:RAPIDS_cugraph_gnn_GIT_REPO}",
     "RAPIDS_cuml_GIT_REPO": "${localEnv:RAPIDS_cuml_GIT_REPO}",
-    "RAPIDS_cumlprims_mg_GIT_REPO": "${localEnv:RAPIDS_cumlprims_mg_GIT_REPO}",
     "RAPIDS_cuvs_GIT_REPO": "${localEnv:RAPIDS_cuvs_GIT_REPO}",
     "RAPIDS_kvikio_GIT_REPO": "${localEnv:RAPIDS_kvikio_GIT_REPO}",
     "RAPIDS_LIBS": "${localEnv:RAPIDS_LIBS}",
@@ -50,7 +49,7 @@
   "initializeCommand": [
     "/bin/bash",
     "-c",
-    "mkdir -m 0755 -p ${localWorkspaceFolder}/.{aws,cache,config,local/state} ${localWorkspaceFolder}/ci/rapids/.{conda,log/devcontainer-utils} ${localWorkspaceFolder}/ci/rapids/.repos/{rmm,kvikio,ucxx,cudf,raft,cuvs,cumlprims_mg,cuml,cugraph,cugraph-gnn}"
+    "mkdir -m 0755 -p ${localWorkspaceFolder}/.{aws,cache,config,local/state} ${localWorkspaceFolder}/ci/rapids/.{conda,log/devcontainer-utils} ${localWorkspaceFolder}/ci/rapids/.repos/{rmm,kvikio,ucxx,cudf,raft,cuvs,cuml,cugraph,cugraph-gnn}"
   ],
   "postCreateCommand": [
     "/bin/bash",
@@ -75,7 +74,6 @@
     "source=${localWorkspaceFolder}/ci/rapids/.repos/cudf,target=/home/coder/cudf,type=bind,consistency=consistent",
     "source=${localWorkspaceFolder}/ci/rapids/.repos/raft,target=/home/coder/raft,type=bind,consistency=consistent",
     "source=${localWorkspaceFolder}/ci/rapids/.repos/cuvs,target=/home/coder/cuvs,type=bind,consistency=consistent",
-    "source=${localWorkspaceFolder}/ci/rapids/.repos/cumlprims_mg,target=/home/coder/cumlprims_mg,type=bind,consistency=consistent",
     "source=${localWorkspaceFolder}/ci/rapids/.repos/cuml,target=/home/coder/cuml,type=bind,consistency=consistent",
     "source=${localWorkspaceFolder}/ci/rapids/.repos/cugraph,target=/home/coder/cugraph,type=bind,consistency=consistent",
     "source=${localWorkspaceFolder}/ci/rapids/.repos/cugraph-gnn,target=/home/coder/cugraph-gnn,type=bind,consistency=consistent",

cub/cub/block/block_adjacent_difference.cuh

@@ -213,7 +213,7 @@ public:
       , linear_tid(RowMajorTid(BlockDimX, BlockDimY, BlockDimZ))
   {}
 
-  //! @} end member group
+  //! @}
   //! @name Read left operations
   //! @{
 
@@ -638,7 +638,7 @@ public:
     }
   }
 
-  //! @} end member group
+  //! @}
   //! @name Read right operations
   //! @{
   //!

cub/cub/block/block_discontinuity.cuh

@@ -238,7 +238,7 @@ public:
       , linear_tid(RowMajorTid(BlockDimX, BlockDimY, BlockDimZ))
   {}
 
-  //! @} end member group
+  //! @}
   //! @name Head flag operations
   //! @{
 
@@ -482,7 +482,7 @@ public:
     FlagHeads(head_flags, input, preds, flag_op, tile_predecessor_item);
   }
 
-  //! @} end member group
+  //! @}
   //! @name Tail flag operations
   //! @{
 
@@ -675,7 +675,7 @@ public:
     Iterate::FlagTails(linear_tid, tail_flags, input, flag_op);
   }
 
-  //! @} end member group
+  //! @}
   //! @name Head & tail flag operations
   //! @{
 
@@ -1194,7 +1194,7 @@ public:
     Iterate::FlagTails(linear_tid, tail_flags, input, flag_op);
   }
 
-  //! @} end member group
+  //! @}
 };
 
 CUB_NAMESPACE_END

cub/cub/block/block_exchange.cuh

@@ -823,7 +823,7 @@ public:
       : temp_storage(temp_storage.Alias())
   {}
 
-  //! @} end member group
+  //! @}
   //! @name Structured exchanges
   //! @{
 
@@ -1034,7 +1034,7 @@ public:
     BlockedToWarpStriped(input_items, output_items, detail::bool_constant_v<WarpTimeSlicing>);
   }
 
-  //! @}  end member group
+  //! @}
   //! @name Scatter exchanges
   //! @{
 
@@ -1196,7 +1196,7 @@ public:
     }
   }
 
-  //! @}  end member group
+  //! @}
 
 #ifndef _CCCL_DOXYGEN_INVOKED // Do not document