[WIP] Reduce compilation time

CMakeLists.txt

@@ -69,7 +69,7 @@ add_library(cuco::cuco ALIAS cuco)
 target_include_directories(cuco INTERFACE
     INTERFACE   $<BUILD_INTERFACE:${CUCO_SOURCE_DIR}/include>
                 $<INSTALL_INTERFACE:include>)
-target_link_libraries(cuco INTERFACE libcudacxx CUDA::toolkit $<BUILD_INTERFACE:cuco::Thrust>)
+target_link_libraries(cuco INTERFACE libcudacxx::libcudacxx CUDA::toolkit $<BUILD_INTERFACE:cuco::Thrust>)
 target_compile_features(cuco INTERFACE cxx_std_17 cuda_std_17)
 
 ###################################################################################################
@@ -108,7 +108,7 @@ endif(BUILD_EXAMPLES)
 ###################################################################################################
 # - Install targets -------------------------------------------------------------------------------
 
-install(TARGETS cuco libcudacxx EXPORT cuco-exports)
+install(TARGETS cuco EXPORT cuco-exports)
 install(DIRECTORY include/cuco/ DESTINATION include/cuco)
 install(FILES ${CUCO_BINARY_DIR}/include/cuco/version_config.hpp DESTINATION include/cuco)
 
@@ -131,7 +131,7 @@ structures tailored for efficient use with GPUs.
 rapids_export(
     INSTALL cuco
     EXPORT_SET cuco-exports
-    GLOBAL_TARGETS cuco libcudacxx
+    GLOBAL_TARGETS cuco
     NAMESPACE cuco::
     DOCUMENTATION doc_string)
 
@@ -146,7 +146,7 @@ endif()
 rapids_export(
     BUILD cuco
     EXPORT_SET cuco-exports
-    GLOBAL_TARGETS cuco libcudacxx
+    GLOBAL_TARGETS cuco
     NAMESPACE cuco::
     DOCUMENTATION doc_string
     FINAL_CODE_BLOCK code_string)

benchmarks/hash_table/dynamic_map_bench.cu

@@ -15,120 +15,118 @@
  */
 
 #include <benchmark/benchmark.h>
-#include <synchronization.hpp>
 #include <cuco/dynamic_map.cuh>
 #include <iostream>
 #include <random>
+#include <synchronization.hpp>
 
-enum class dist_type {
-  UNIQUE,
-  UNIFORM,
-  GAUSSIAN
-};
+enum class dist_type { UNIQUE, UNIFORM, GAUSSIAN };
 
-template<dist_type Dist, typename Key, typename OutputIt>
-static void generate_keys(OutputIt output_begin, OutputIt output_end) {
+template <dist_type Dist, typename Key, typename OutputIt>
+static void generate_keys(OutputIt output_begin, OutputIt output_end)
+{
   auto num_keys = std::distance(output_begin, output_end);
-  
+
   std::random_device rd;
   std::mt19937 gen{rd()};
 
-  switch(Dist) {
+  switch (Dist) {
     case dist_type::UNIQUE:
-      for(auto i = 0; i < num_keys; ++i) {
+      for (auto i = 0; i < num_keys; ++i) {
         output_begin[i] = i;
       }
       break;
     case dist_type::UNIFORM:
-      for(auto i = 0; i < num_keys; ++i) {
+      for (auto i = 0; i < num_keys; ++i) {
         output_begin[i] = std::abs(static_cast<Key>(gen()));
       }
       break;
     case dist_type::GAUSSIAN:
       std::normal_distribution<> dg{1e9, 1e7};
-      for(auto i = 0; i < num_keys; ++i) {
+      for (auto i = 0; i < num_keys; ++i) {
         output_begin[i] = std::abs(static_cast<Key>(dg(gen)));
       }
       break;
   }
 }
 
-static void gen_final_size(benchmark::internal::Benchmark* b) {
-  for(auto size = 10'000'000; size <= 150'000'000; size += 20'000'000) {
+static void gen_final_size(benchmark::internal::Benchmark* b)
+{
+  for (auto size = 10'000'000; size <= 150'000'000; size += 20'000'000) {
     b->Args({size});
   }
 }
 
 template <typename Key, typename Value, dist_type Dist>
-static void BM_dynamic_insert(::benchmark::State& state) {
+static void BM_dynamic_insert(::benchmark::State& state)
+{
   using map_type = cuco::dynamic_map<Key, Value>;
-  
-  std::size_t num_keys = state.range(0);
-  std::size_t initial_size = 1<<27;
-  
-  std::vector<Key> h_keys( num_keys );
-  std::vector<cuco::pair_type<Key, Value>> h_pairs ( num_keys );
-  
+
+  std::size_t num_keys     = state.range(0);
+  std::size_t initial_size = 1 << 27;
+
+  std::vector<Key> h_keys(num_keys);
+  std::vector<cuco::pair_type<Key, Value>> h_pairs(num_keys);
+
   generate_keys<Dist, Key>(h_keys.begin(), h_keys.end());
 
-  for(auto i = 0; i < num_keys; ++i) {
-    Key key = h_keys[i];
-    Value val = h_keys[i];
-    h_pairs[i].first = key;
+  for (auto i = 0; i < num_keys; ++i) {
+    Key key           = h_keys[i];
+    Value val         = h_keys[i];
+    h_pairs[i].first  = key;
     h_pairs[i].second = val;
   }
 
-  thrust::device_vector<cuco::pair_type<Key, Value>> d_pairs( h_pairs );
+  thrust::device_vector<cuco::pair_type<Key, Value>> d_pairs(h_pairs);
 
   std::size_t batch_size = 1E6;
-  for(auto _ : state) {
+  for (auto _ : state) {
     map_type map{initial_size, -1, -1};
     {
-      cuda_event_timer raii{state}; 
-      for(auto i = 0; i < num_keys; i += batch_size) {
+      cuda_event_timer raii{state};
+      for (auto i = 0; i < num_keys; i += batch_size) {
         map.insert(d_pairs.begin() + i, d_pairs.begin() + i + batch_size);
       }
     }
   }
 
-  state.SetBytesProcessed((sizeof(Key) + sizeof(Value)) *
-                          int64_t(state.iterations()) *
+  state.SetBytesProcessed((sizeof(Key) + sizeof(Value)) * int64_t(state.iterations()) *
                           int64_t(state.range(0)));
 }
 
 template <typename Key, typename Value, dist_type Dist>
-static void BM_dynamic_search_all(::benchmark::State& state) {
+static void BM_dynamic_search_all(::benchmark::State& state)
+{
   using map_type = cuco::dynamic_map<Key, Value>;
-
-  std::size_t num_keys = state.range(0);
-  std::size_t initial_size = 1<<27;
 
-  std::vector<Key> h_keys( num_keys );
-  std::vector<cuco::pair_type<Key, Value>> h_pairs ( num_keys );
+  std::size_t num_keys     = state.range(0);
+  std::size_t initial_size = 1 << 27;
+
+  std::vector<Key> h_keys(num_keys);
+  std::vector<cuco::pair_type<Key, Value>> h_pairs(num_keys);
 
   generate_keys<Dist, Key>(h_keys.begin(), h_keys.end());
-  
-  for(auto i = 0; i < num_keys; ++i) {
-    Key key = h_keys[i];
-    Value val = h_keys[i];
-    h_pairs[i].first = key;
+
+  for (auto i = 0; i < num_keys; ++i) {
+    Key key           = h_keys[i];
+    Value val         = h_keys[i];
+    h_pairs[i].first  = key;
     h_pairs[i].second = val;
   }
 
-  thrust::device_vector<Key> d_keys( h_keys );
-  thrust::device_vector<cuco::pair_type<Key, Value>> d_pairs( h_pairs );
-  thrust::device_vector<Value> d_results( num_keys );
+  thrust::device_vector<Key> d_keys(h_keys);
+  thrust::device_vector<cuco::pair_type<Key, Value>> d_pairs(h_pairs);
+  thrust::device_vector<Value> d_results(num_keys);
 
   map_type map{initial_size, -1, -1};
   map.insert(d_pairs.begin(), d_pairs.end());
 
-  for(auto _ : state) {
+  for (auto _ : state) {
     cuda_event_timer raii{state};
     map.find(d_keys.begin(), d_keys.end(), d_results.begin());
   }
 
-  state.SetBytesProcessed((sizeof(Key) + sizeof(Value)) *
-                          int64_t(state.iterations()) *
+  state.SetBytesProcessed((sizeof(Key) + sizeof(Value)) * int64_t(state.iterations()) *
                           int64_t(state.range(0)));
 }
 
@@ -161,7 +159,7 @@ BENCHMARK_TEMPLATE(BM_dynamic_search_all, int32_t, int32_t, dist_type::GAUSSIAN)
   ->Unit(benchmark::kMillisecond)
   ->Apply(gen_final_size)
   ->UseManualTime();
-  
+
 BENCHMARK_TEMPLATE(BM_dynamic_insert, int64_t, int64_t, dist_type::UNIQUE)
   ->Unit(benchmark::kMillisecond)
   ->Apply(gen_final_size)

benchmarks/hash_table/static_map_bench.cu

@@ -15,40 +15,37 @@
  */
 
 #include <benchmark/benchmark.h>
-#include "cuco/static_map.cuh"
-#include <thrust/for_each.h>
 #include <thrust/device_vector.h>
-#include <iostream>
+#include <thrust/for_each.h>
 #include <fstream>
+#include <iostream>
 #include <random>
+#include "cuco/static_map.cuh"
 
-enum class dist_type {
-  UNIQUE,
-  UNIFORM,
-  GAUSSIAN
-};
+enum class dist_type { UNIQUE, UNIFORM, GAUSSIAN };
 
-template<dist_type Dist, typename Key, typename OutputIt>
-static void generate_keys(OutputIt output_begin, OutputIt output_end) {
+template <dist_type Dist, typename Key, typename OutputIt>
+static void generate_keys(OutputIt output_begin, OutputIt output_end)
+{
   auto num_keys = std::distance(output_begin, output_end);
-  
+
   std::random_device rd;
   std::mt19937 gen{rd()};
 
-  switch(Dist) {
+  switch (Dist) {
     case dist_type::UNIQUE:
-      for(auto i = 0; i < num_keys; ++i) {
+      for (auto i = 0; i < num_keys; ++i) {
         output_begin[i] = i;
       }
       break;
     case dist_type::UNIFORM:
-      for(auto i = 0; i < num_keys; ++i) {
+      for (auto i = 0; i < num_keys; ++i) {
         output_begin[i] = std::abs(static_cast<Key>(gen()));
       }
       break;
     case dist_type::GAUSSIAN:
       std::normal_distribution<> dg{1e9, 1e7};
-      for(auto i = 0; i < num_keys; ++i) {
+      for (auto i = 0; i < num_keys; ++i) {
         output_begin[i] = std::abs(static_cast<Key>(dg(gen)));
       }
       break;
@@ -59,39 +56,39 @@ static void generate_keys(OutputIt output_begin, OutputIt output_end) {
  * @brief Generates input sizes and hash table occupancies
  *
  */
-static void generate_size_and_occupancy(benchmark::internal::Benchmark* b) {
+static void generate_size_and_occupancy(benchmark::internal::Benchmark* b)
+{
   for (auto size = 100'000'000; size <= 100'000'000; size *= 10) {
     for (auto occupancy = 10; occupancy <= 90; occupancy += 10) {
       b->Args({size, occupancy});
     }
   }
 }
 
-
-
 template <typename Key, typename Value, dist_type Dist>
-static void BM_static_map_insert(::benchmark::State& state) {
+static void BM_static_map_insert(::benchmark::State& state)
+{
   using map_type = cuco::static_map<Key, Value>;
-  
+
   std::size_t num_keys = state.range(0);
-  float occupancy = state.range(1) / float{100};
-  std::size_t size = num_keys / occupancy;
+  float occupancy      = state.range(1) / float{100};
+  std::size_t size     = num_keys / occupancy;
+
+  std::vector<Key> h_keys(num_keys);
+  std::vector<cuco::pair_type<Key, Value>> h_pairs(num_keys);
 
-  std::vector<Key> h_keys( num_keys );
-  std::vector<cuco::pair_type<Key, Value>> h_pairs( num_keys );
-
   generate_keys<Dist, Key>(h_keys.begin(), h_keys.end());
-  
-  for(auto i = 0; i < num_keys; ++i) {
-    Key key = h_keys[i];
-    Value val = h_keys[i];
-    h_pairs[i].first = key;
+
+  for (auto i = 0; i < num_keys; ++i) {
+    Key key           = h_keys[i];
+    Value val         = h_keys[i];
+    h_pairs[i].first  = key;
     h_pairs[i].second = val;
   }
 
-  thrust::device_vector<cuco::pair_type<Key, Value>> d_pairs( h_pairs );
+  thrust::device_vector<cuco::pair_type<Key, Value>> d_pairs(h_pairs);
 
-  for(auto _ : state) {
+  for (auto _ : state) {
     state.ResumeTiming();
     state.PauseTiming();
     map_type map{size, -1, -1};
@@ -102,54 +99,50 @@ static void BM_static_map_insert(::benchmark::State& state) {
     state.PauseTiming();
   }
 
-  state.SetBytesProcessed((sizeof(Key) + sizeof(Value)) *
-                          int64_t(state.iterations()) *
+  state.SetBytesProcessed((sizeof(Key) + sizeof(Value)) * int64_t(state.iterations()) *
                           int64_t(state.range(0)));
 }
 
-
-
 template <typename Key, typename Value, dist_type Dist>
-static void BM_static_map_search_all(::benchmark::State& state) {
+static void BM_static_map_search_all(::benchmark::State& state)
+{
   using map_type = cuco::static_map<Key, Value>;
-  
+
   std::size_t num_keys = state.range(0);
-  float occupancy = state.range(1) / float{100};
-  std::size_t size = num_keys / occupancy;
+  float occupancy      = state.range(1) / float{100};
+  std::size_t size     = num_keys / occupancy;
 
   map_type map{size, -1, -1};
   auto view = map.get_device_mutable_view();
 
-  std::vector<Key> h_keys( num_keys );
-  std::vector<Value> h_values( num_keys );
-  std::vector<cuco::pair_type<Key, Value>> h_pairs ( num_keys );
-  std::vector<Value> h_results (num_keys);
+  std::vector<Key> h_keys(num_keys);
+  std::vector<Value> h_values(num_keys);
+  std::vector<cuco::pair_type<Key, Value>> h_pairs(num_keys);
+  std::vector<Value> h_results(num_keys);
 
   generate_keys<Dist, Key>(h_keys.begin(), h_keys.end());
-  
-  for(auto i = 0; i < num_keys; ++i) {
-    Key key = h_keys[i];
-    Value val = h_keys[i];
-    h_pairs[i].first = key;
+
+  for (auto i = 0; i < num_keys; ++i) {
+    Key key           = h_keys[i];
+    Value val         = h_keys[i];
+    h_pairs[i].first  = key;
     h_pairs[i].second = val;
   }
 
-  thrust::device_vector<Key> d_keys( h_keys ); 
-  thrust::device_vector<Value> d_results( num_keys);
-  thrust::device_vector<cuco::pair_type<Key, Value>> d_pairs( h_pairs );
+  thrust::device_vector<Key> d_keys(h_keys);
+  thrust::device_vector<Value> d_results(num_keys);
+  thrust::device_vector<cuco::pair_type<Key, Value>> d_pairs(h_pairs);
 
   map.insert(d_pairs.begin(), d_pairs.end());
-  
-  for(auto _ : state) {
+
+  for (auto _ : state) {
     map.find(d_keys.begin(), d_keys.end(), d_results.begin());
   }
 
   state.SetBytesProcessed((sizeof(Key) + sizeof(Value)) * int64_t(state.iterations()) *
                           int64_t(state.range(0)));
 }
 
-
-
 BENCHMARK_TEMPLATE(BM_static_map_insert, int32_t, int32_t, dist_type::UNIQUE)
   ->Unit(benchmark::kMillisecond)
   ->Apply(generate_size_and_occupancy);