From 4bca172b58f33ba5fc11b448daf254427e66f042 Mon Sep 17 00:00:00 2001
From: Yunsong Wang <yunsongw@nvidia.com>
Date: Wed, 2 Mar 2022 14:11:31 -0500
Subject: [PATCH 1/8] Reorder header groups in cuco headers

---
 include/cuco/detail/dynamic_map_kernels.cuh   | 10 +++++++--
 include/cuco/detail/error.hpp                 |  3 ++-
 include/cuco/detail/pair.cuh                  |  1 +
 include/cuco/detail/probe_sequence_impl.cuh   |  6 +++---
 include/cuco/detail/static_map_kernels.cuh    |  8 ++++++-
 .../cuco/detail/static_multimap/kernels.cuh   |  9 +++++---
 .../static_multimap/static_multimap.inl       |  6 +++---
 include/cuco/dynamic_map.cuh                  | 13 +++++++-----
 include/cuco/static_map.cuh                   | 21 ++++++++-----------
 include/cuco/static_multimap.cuh              | 20 +++++++++---------
 10 files changed, 57 insertions(+), 40 deletions(-)

diff --git a/include/cuco/detail/dynamic_map_kernels.cuh b/include/cuco/detail/dynamic_map_kernels.cuh
index c1e21e863..a18aee736 100644
--- a/include/cuco/detail/dynamic_map_kernels.cuh
+++ b/include/cuco/detail/dynamic_map_kernels.cuh
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,6 +14,12 @@
  * limitations under the License.
  */
 
+#include <cub/cub.cuh>
+
+#include <cuda/std/atomic>
+
+#include <cooperative_groups.h>
+
 namespace cuco {
 namespace detail {
 namespace cg = cooperative_groups;
@@ -457,4 +463,4 @@ __global__ void contains(InputIt first,
   }
 }
 }  // namespace detail
-}  // namespace cuco
\ No newline at end of file
+}  // namespace cuco
diff --git a/include/cuco/detail/error.hpp b/include/cuco/detail/error.hpp
index f5f331222..ee29bfb5f 100644
--- a/include/cuco/detail/error.hpp
+++ b/include/cuco/detail/error.hpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -17,6 +17,7 @@
 #pragma once
 
 #include <cuda_runtime_api.h>
+
 #include <stdexcept>
 #include <string>
 
diff --git a/include/cuco/detail/pair.cuh b/include/cuco/detail/pair.cuh
index 792b6c2ca..1b013b771 100644
--- a/include/cuco/detail/pair.cuh
+++ b/include/cuco/detail/pair.cuh
@@ -19,6 +19,7 @@
 #include <thrust/device_reference.h>
 #include <thrust/pair.h>
 #include <thrust/tuple.h>
+
 #include <tuple>
 #include <type_traits>
 
diff --git a/include/cuco/detail/probe_sequence_impl.cuh b/include/cuco/detail/probe_sequence_impl.cuh
index 549d59df3..747df89c7 100644
--- a/include/cuco/detail/probe_sequence_impl.cuh
+++ b/include/cuco/detail/probe_sequence_impl.cuh
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -16,11 +16,11 @@
 
 #pragma once
 
-#include <cuda/std/atomic>
-
 #include <cuco/detail/hash_functions.cuh>
 #include <cuco/detail/pair.cuh>
 
+#include <cuda/std/atomic>
+
 namespace cuco {
 namespace detail {
 
diff --git a/include/cuco/detail/static_map_kernels.cuh b/include/cuco/detail/static_map_kernels.cuh
index 642373135..2133bfa24 100644
--- a/include/cuco/detail/static_map_kernels.cuh
+++ b/include/cuco/detail/static_map_kernels.cuh
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,6 +14,12 @@
  * limitations under the License.
  */
 
+#include <cub/cub.cuh>
+
+#include <cuda/std/atomic>
+
+#include <cooperative_groups.h>
+
 namespace cuco {
 namespace detail {
 namespace cg = cooperative_groups;
diff --git a/include/cuco/detail/static_multimap/kernels.cuh b/include/cuco/detail/static_multimap/kernels.cuh
index 83aa2c640..8bb5c486e 100644
--- a/include/cuco/detail/static_multimap/kernels.cuh
+++ b/include/cuco/detail/static_multimap/kernels.cuh
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,10 +14,13 @@
  * limitations under the License.
  */
 
-#include <cooperative_groups/memcpy_async.h>
+#include <cuco/detail/pair.cuh>
+
 #include <thrust/type_traits/is_contiguous_iterator.h>
 
-#include <cuco/detail/pair.cuh>
+#include <cub/cub.cuh>
+
+#include <cooperative_groups/memcpy_async.h>
 
 namespace cuco {
 namespace detail {
diff --git a/include/cuco/detail/static_multimap/static_multimap.inl b/include/cuco/detail/static_multimap/static_multimap.inl
index f1f325e8c..d8d6b8683 100644
--- a/include/cuco/detail/static_multimap/static_multimap.inl
+++ b/include/cuco/detail/static_multimap/static_multimap.inl
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,12 +14,12 @@
  * limitations under the License.
  */
 
+#include <cuco/detail/utils.hpp>
+
 #include <thrust/count.h>
 #include <thrust/iterator/transform_iterator.h>
 #include <thrust/tuple.h>
 
-#include <cuco/detail/utils.hpp>
-
 namespace {
 /**
  * @brief Device functor used to determine if a slot is filled.
diff --git a/include/cuco/dynamic_map.cuh b/include/cuco/dynamic_map.cuh
index d762de41e..9867aa7f3 100644
--- a/include/cuco/dynamic_map.cuh
+++ b/include/cuco/dynamic_map.cuh
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -16,13 +16,16 @@
 
 #pragma once
 
-#include <cooperative_groups.h>
-#include <cub/cub.cuh>
 #include <cuco/detail/dynamic_map_kernels.cuh>
 #include <cuco/detail/error.hpp>
 #include <cuco/static_map.cuh>
-#include <cuda/std/atomic>
+
 #include <thrust/device_vector.h>
+#include <thrust/functional.h>
+
+#include <cuda/std/atomic>
+
+#include <memory>
 
 namespace cuco {
 
@@ -258,4 +261,4 @@ class dynamic_map {
 };
 }  // namespace cuco
 
-#include <cuco/detail/dynamic_map.inl>
\ No newline at end of file
+#include <cuco/detail/dynamic_map.inl>
diff --git a/include/cuco/static_map.cuh b/include/cuco/static_map.cuh
index 199dcd838..24b7fbdea 100644
--- a/include/cuco/static_map.cuh
+++ b/include/cuco/static_map.cuh
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2021, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -16,16 +16,16 @@
 
 #pragma once
 
-#include <cooperative_groups.h>
-#include <cub/cub.cuh>
-#include <cuda/std/atomic>
-#include <memory>
-#include <thrust/distance.h>
-#include <thrust/functional.h>
-
 #include <cuco/allocator.hpp>
+#include <cuco/detail/error.hpp>
+#include <cuco/detail/hash_functions.cuh>
+#include <cuco/detail/pair.cuh>
+#include <cuco/detail/static_map_kernels.cuh>
 #include <cuco/traits.hpp>
 
+#include <thrust/functional.h>
+
+#include <cuda/std/atomic>
 #if defined(CUDART_VERSION) && (CUDART_VERSION >= 11000) && defined(__CUDA_ARCH__) && \
   (__CUDA_ARCH__ >= 700)
 #define CUCO_HAS_CUDA_BARRIER
@@ -35,10 +35,7 @@
 #include <cuda/barrier>
 #endif
 
-#include <cuco/detail/error.hpp>
-#include <cuco/detail/hash_functions.cuh>
-#include <cuco/detail/pair.cuh>
-#include <cuco/detail/static_map_kernels.cuh>
+#include <memory>
 
 namespace cuco {
 
diff --git a/include/cuco/static_multimap.cuh b/include/cuco/static_multimap.cuh
index fea5b0565..2953b5dcf 100644
--- a/include/cuco/static_multimap.cuh
+++ b/include/cuco/static_multimap.cuh
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -16,16 +16,16 @@
 
 #pragma once
 
-#include <cooperative_groups.h>
-#include <cub/cub.cuh>
-#include <memory>
-#include <thrust/distance.h>
-#include <thrust/functional.h>
-
 #include <cuco/allocator.hpp>
+#include <cuco/detail/error.hpp>
+#include <cuco/detail/prime.hpp>
+#include <cuco/detail/static_multimap/kernels.cuh>
 #include <cuco/probe_sequences.cuh>
 #include <cuco/traits.hpp>
 
+#include <thrust/functional.h>
+
+#include <cuda/std/atomic>
 #if defined(CUDART_VERSION) && (CUDART_VERSION >= 11000) && defined(__CUDA_ARCH__) && \
   (__CUDA_ARCH__ >= 700)
 #define CUCO_HAS_CUDA_BARRIER
@@ -40,9 +40,9 @@
 #include <cuda/barrier>
 #endif
 
-#include <cuco/detail/error.hpp>
-#include <cuco/detail/prime.hpp>
-#include <cuco/detail/static_multimap/kernels.cuh>
+#include <cooperative_groups.h>
+
+#include <memory>
 
 namespace cuco {
 

From 5418d2545a5ed1b6cf96f186bad906f5ddb59543 Mon Sep 17 00:00:00 2001
From: Yunsong Wang <yunsongw@nvidia.com>
Date: Wed, 2 Mar 2022 15:10:22 -0500
Subject: [PATCH 2/8] Reorder header groups in benchmarks, examples and tests

---
 benchmarks/hash_table/dynamic_map_bench.cu          |  9 ++++++---
 benchmarks/hash_table/static_map_bench.cu           | 13 ++++++++-----
 .../hash_table/static_multimap/count_bench.cu       | 10 +++++-----
 .../hash_table/static_multimap/insert_bench.cu      | 10 +++++-----
 .../static_multimap/optimal_retrieve_bench.cu       |  7 +++----
 .../static_multimap/pair_retrieve_bench.cu          | 10 +++++-----
 .../hash_table/static_multimap/query_bench.cu       | 10 +++++-----
 .../hash_table/static_multimap/retrieve_bench.cu    |  7 +++----
 benchmarks/synchronization.hpp                      |  5 +++--
 examples/static_map/custom_type_example.cu          |  6 +++---
 examples/static_map/static_map_example.cu           |  6 +++---
 examples/static_multimap/static_multimap_example.cu |  6 +++---
 include/cuco/detail/static_multimap/kernels.cuh     |  2 ++
 tests/dynamic_map/unique_sequence_test.cu           |  9 +++++----
 tests/static_map/custom_type_test.cu                | 11 +++++++----
 tests/static_map/key_sentinel_test.cu               |  9 +++++----
 tests/static_map/shared_memory_test.cu              | 11 ++++++-----
 tests/static_map/stream_test.cu                     |  9 +++++----
 tests/static_map/unique_sequence_test.cu            |  9 +++++----
 tests/static_multimap/custom_pair_retrieve_test.cu  | 12 ++++++++----
 tests/static_multimap/custom_type_test.cu           | 12 ++++++++----
 tests/static_multimap/insert_if_test.cu             |  9 +++++----
 tests/static_multimap/multiplicity_test.cu          |  9 +++++----
 tests/static_multimap/non_match_test.cu             |  9 +++++----
 tests/static_multimap/pair_function_test.cu         | 11 ++++++-----
 tests/utils.hpp                                     |  4 +++-
 26 files changed, 127 insertions(+), 98 deletions(-)

diff --git a/benchmarks/hash_table/dynamic_map_bench.cu b/benchmarks/hash_table/dynamic_map_bench.cu
index 995e53903..2ce5b1a48 100644
--- a/benchmarks/hash_table/dynamic_map_bench.cu
+++ b/benchmarks/hash_table/dynamic_map_bench.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,11 +14,14 @@
  * limitations under the License.
  */
 
-#include <benchmark/benchmark.h>
+#include <synchronization.hpp>
+
 #include <cuco/dynamic_map.cuh>
+
+#include <benchmark/benchmark.h>
+
 #include <iostream>
 #include <random>
-#include <synchronization.hpp>
 
 enum class dist_type { UNIQUE, UNIFORM, GAUSSIAN };
 
diff --git a/benchmarks/hash_table/static_map_bench.cu b/benchmarks/hash_table/static_map_bench.cu
index 59c325d5d..96def361b 100644
--- a/benchmarks/hash_table/static_map_bench.cu
+++ b/benchmarks/hash_table/static_map_bench.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,13 +14,16 @@
  * limitations under the License.
  */
 
-#include "cuco/static_map.cuh"
+#include <cuco/static_map.cuh>
+
+#include <thrust/device_vector.h>
+#include <thrust/for_each.h>
+
 #include <benchmark/benchmark.h>
+
 #include <fstream>
 #include <iostream>
 #include <random>
-#include <thrust/device_vector.h>
-#include <thrust/for_each.h>
 
 enum class dist_type { UNIQUE, UNIFORM, GAUSSIAN };
 
@@ -189,4 +192,4 @@ BENCHMARK_TEMPLATE(BM_static_map_insert, int64_t, int64_t, dist_type::GAUSSIAN)
 
 BENCHMARK_TEMPLATE(BM_static_map_search_all, int64_t, int64_t, dist_type::GAUSSIAN)
   ->Unit(benchmark::kMillisecond)
-  ->Apply(generate_size_and_occupancy);
\ No newline at end of file
+  ->Apply(generate_size_and_occupancy);
diff --git a/benchmarks/hash_table/static_multimap/count_bench.cu b/benchmarks/hash_table/static_multimap/count_bench.cu
index 518198f27..9cdb3bcea 100644
--- a/benchmarks/hash_table/static_multimap/count_bench.cu
+++ b/benchmarks/hash_table/static_multimap/count_bench.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,13 +14,13 @@
  * limitations under the License.
  */
 
-#include <random>
+#include <key_generator.hpp>
+
+#include <cuco/static_multimap.cuh>
 
 #include <nvbench/nvbench.cuh>
-#include <thrust/device_vector.h>
 
-#include <cuco/static_multimap.cuh>
-#include <key_generator.hpp>
+#include <thrust/device_vector.h>
 
 /**
  * @brief A benchmark evaluating multi-value `count` performance:
diff --git a/benchmarks/hash_table/static_multimap/insert_bench.cu b/benchmarks/hash_table/static_multimap/insert_bench.cu
index 3798a69fe..2bc5b8e7f 100644
--- a/benchmarks/hash_table/static_multimap/insert_bench.cu
+++ b/benchmarks/hash_table/static_multimap/insert_bench.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,13 +14,13 @@
  * limitations under the License.
  */
 
-#include <random>
+#include <key_generator.hpp>
+
+#include <cuco/static_multimap.cuh>
 
 #include <nvbench/nvbench.cuh>
-#include <thrust/device_vector.h>
 
-#include <cuco/static_multimap.cuh>
-#include <key_generator.hpp>
+#include <thrust/device_vector.h>
 
 /**
  * @brief A benchmark evaluating multi-value `insert` performance:
diff --git a/benchmarks/hash_table/static_multimap/optimal_retrieve_bench.cu b/benchmarks/hash_table/static_multimap/optimal_retrieve_bench.cu
index 265988444..28a79fe36 100644
--- a/benchmarks/hash_table/static_multimap/optimal_retrieve_bench.cu
+++ b/benchmarks/hash_table/static_multimap/optimal_retrieve_bench.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,13 +14,12 @@
  * limitations under the License.
  */
 
+#include <cuco/static_multimap.cuh>
+
 #include <nvbench/nvbench.cuh>
 
-#include <random>
 #include <thrust/device_vector.h>
 
-#include "cuco/static_multimap.cuh"
-
 /**
  * @brief Generates input keys by a given number of repetitions per key.
  *
diff --git a/benchmarks/hash_table/static_multimap/pair_retrieve_bench.cu b/benchmarks/hash_table/static_multimap/pair_retrieve_bench.cu
index 5b89e22be..b73a14039 100644
--- a/benchmarks/hash_table/static_multimap/pair_retrieve_bench.cu
+++ b/benchmarks/hash_table/static_multimap/pair_retrieve_bench.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,15 +14,15 @@
  * limitations under the License.
  */
 
-#include <random>
+#include <key_generator.hpp>
+
+#include <cuco/static_multimap.cuh>
 
 #include <nvbench/nvbench.cuh>
+
 #include <thrust/device_vector.h>
 #include <thrust/iterator/discard_iterator.h>
 
-#include <cuco/static_multimap.cuh>
-#include <key_generator.hpp>
-
 namespace {
 // Custom pair equal
 template <typename Key, typename Value>
diff --git a/benchmarks/hash_table/static_multimap/query_bench.cu b/benchmarks/hash_table/static_multimap/query_bench.cu
index de76efb1b..36daf33ea 100644
--- a/benchmarks/hash_table/static_multimap/query_bench.cu
+++ b/benchmarks/hash_table/static_multimap/query_bench.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,13 +14,13 @@
  * limitations under the License.
  */
 
-#include <random>
+#include <key_generator.hpp>
+
+#include <cuco/static_multimap.cuh>
 
 #include <nvbench/nvbench.cuh>
-#include <thrust/device_vector.h>
 
-#include <cuco/static_multimap.cuh>
-#include <key_generator.hpp>
+#include <thrust/device_vector.h>
 
 /**
  * @brief A benchmark evaluating multi-value query (`count` + `retrieve`) performance:
diff --git a/benchmarks/hash_table/static_multimap/retrieve_bench.cu b/benchmarks/hash_table/static_multimap/retrieve_bench.cu
index ff1787f43..8c13aee39 100644
--- a/benchmarks/hash_table/static_multimap/retrieve_bench.cu
+++ b/benchmarks/hash_table/static_multimap/retrieve_bench.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,14 +14,13 @@
  * limitations under the License.
  */
 
-#include <cuco/static_multimap.cuh>
 #include <key_generator.hpp>
 
-#include <thrust/device_vector.h>
+#include <cuco/static_multimap.cuh>
 
 #include <nvbench/nvbench.cuh>
 
-#include <random>
+#include <thrust/device_vector.h>
 
 /**
  * @brief A benchmark evaluating multi-value `retrieve` performance:
diff --git a/benchmarks/synchronization.hpp b/benchmarks/synchronization.hpp
index d9d6019ca..a3be1ae5a 100644
--- a/benchmarks/synchronization.hpp
+++ b/benchmarks/synchronization.hpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -18,6 +18,7 @@
 
 // Google Benchmark library
 #include <benchmark/benchmark.h>
+
 #include <cuda_runtime_api.h>
 
 #define BENCH_CUDA_TRY(call)                                                         \
@@ -129,4 +130,4 @@ class cuda_event_timer {
   cudaEvent_t stop_;
   cudaStream_t stream_;
   benchmark::State* p_state;
-};
\ No newline at end of file
+};
diff --git a/examples/static_map/custom_type_example.cu b/examples/static_map/custom_type_example.cu
index 443323e36..a84a395dc 100644
--- a/examples/static_map/custom_type_example.cu
+++ b/examples/static_map/custom_type_example.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,12 +14,12 @@
  * limitations under the License.
  */
 
+#include <cuco/static_map.cuh>
+
 #include <thrust/device_vector.h>
 #include <thrust/logical.h>
 #include <thrust/transform.h>
 
-#include <cuco/static_map.cuh>
-
 // User-defined key type
 #ifdef CUCO_NO_INDEPENDENT_THREADS
 struct custom_key_type {
diff --git a/examples/static_map/static_map_example.cu b/examples/static_map/static_map_example.cu
index 12d12578d..4f365bc50 100644
--- a/examples/static_map/static_map_example.cu
+++ b/examples/static_map/static_map_example.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,13 +14,13 @@
  * limitations under the License.
  */
 
-#include <limits>
+#include <cuco/static_map.cuh>
 
 #include <thrust/device_vector.h>
 #include <thrust/sequence.h>
 #include <thrust/transform.h>
 
-#include <cuco/static_map.cuh>
+#include <limits>
 
 int main(void)
 {
diff --git a/examples/static_multimap/static_multimap_example.cu b/examples/static_multimap/static_multimap_example.cu
index cda8f2073..d229c8455 100644
--- a/examples/static_multimap/static_multimap_example.cu
+++ b/examples/static_multimap/static_multimap_example.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,13 +14,13 @@
  * limitations under the License.
  */
 
-#include <limits>
+#include <cuco/static_multimap.cuh>
 
 #include <thrust/device_vector.h>
 #include <thrust/sequence.h>
 #include <thrust/transform.h>
 
-#include <cuco/static_multimap.cuh>
+#include <limits>
 
 int main(void)
 {
diff --git a/include/cuco/detail/static_multimap/kernels.cuh b/include/cuco/detail/static_multimap/kernels.cuh
index 8bb5c486e..e3990a373 100644
--- a/include/cuco/detail/static_multimap/kernels.cuh
+++ b/include/cuco/detail/static_multimap/kernels.cuh
@@ -20,6 +20,8 @@
 
 #include <cub/cub.cuh>
 
+#include <cuda/std/atomic>
+
 #include <cooperative_groups/memcpy_async.h>
 
 namespace cuco {
diff --git a/tests/dynamic_map/unique_sequence_test.cu b/tests/dynamic_map/unique_sequence_test.cu
index 8d5af41e3..7aa891278 100644
--- a/tests/dynamic_map/unique_sequence_test.cu
+++ b/tests/dynamic_map/unique_sequence_test.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2021, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,12 +14,13 @@
  * limitations under the License.
  */
 
-#include <catch2/catch.hpp>
-#include <thrust/device_vector.h>
+#include <utils.hpp>
 
 #include <cuco/dynamic_map.cuh>
 
-#include <utils.hpp>
+#include <thrust/device_vector.h>
+
+#include <catch2/catch.hpp>
 
 TEMPLATE_TEST_CASE_SIG("Unique sequence of keys",
                        "",
diff --git a/tests/static_map/custom_type_test.cu b/tests/static_map/custom_type_test.cu
index 1e75bf826..2789e34b9 100644
--- a/tests/static_map/custom_type_test.cu
+++ b/tests/static_map/custom_type_test.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2021, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,12 +14,15 @@
  * limitations under the License.
  */
 
-#include <catch2/catch.hpp>
-#include <thrust/device_vector.h>
+#include <utils.hpp>
 
 #include <cuco/static_map.cuh>
 
-#include <utils.hpp>
+#include <thrust/device_vector.h>
+
+#include <catch2/catch.hpp>
+
+#include <tuple>
 
 // User-defined key type
 template <typename T>
diff --git a/tests/static_map/key_sentinel_test.cu b/tests/static_map/key_sentinel_test.cu
index a96a60af4..c1d0e4296 100644
--- a/tests/static_map/key_sentinel_test.cu
+++ b/tests/static_map/key_sentinel_test.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2021, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,12 +14,13 @@
  * limitations under the License.
  */
 
-#include <catch2/catch.hpp>
-#include <thrust/device_vector.h>
+#include <utils.hpp>
 
 #include <cuco/static_map.cuh>
 
-#include <utils.hpp>
+#include <thrust/device_vector.h>
+
+#include <catch2/catch.hpp>
 
 #define SIZE 10
 __device__ int A[SIZE];
diff --git a/tests/static_map/shared_memory_test.cu b/tests/static_map/shared_memory_test.cu
index 00bbbcd08..89b8800c6 100644
--- a/tests/static_map/shared_memory_test.cu
+++ b/tests/static_map/shared_memory_test.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2021, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,14 +14,15 @@
  * limitations under the License.
  */
 
-#include <limits>
+#include <utils.hpp>
+
+#include <cuco/static_map.cuh>
 
-#include <catch2/catch.hpp>
 #include <thrust/device_vector.h>
 
-#include <cuco/static_map.cuh>
+#include <catch2/catch.hpp>
 
-#include <utils.hpp>
+#include <limits>
 
 template <typename MapType, int CAPACITY>
 __global__ void shared_memory_test_kernel(
diff --git a/tests/static_map/stream_test.cu b/tests/static_map/stream_test.cu
index aadee823a..bb7dac85f 100644
--- a/tests/static_map/stream_test.cu
+++ b/tests/static_map/stream_test.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,12 +14,13 @@
  * limitations under the License.
  */
 
-#include <catch2/catch.hpp>
-#include <thrust/device_vector.h>
+#include <utils.hpp>
 
 #include <cuco/static_map.cuh>
 
-#include <utils.hpp>
+#include <thrust/device_vector.h>
+
+#include <catch2/catch.hpp>
 
 TEMPLATE_TEST_CASE_SIG("Unique sequence of keys on given stream",
                        "",
diff --git a/tests/static_map/unique_sequence_test.cu b/tests/static_map/unique_sequence_test.cu
index da4b51a3f..ca8cc82ea 100644
--- a/tests/static_map/unique_sequence_test.cu
+++ b/tests/static_map/unique_sequence_test.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2021, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,12 +14,13 @@
  * limitations under the License.
  */
 
-#include <catch2/catch.hpp>
-#include <thrust/device_vector.h>
+#include <utils.hpp>
 
 #include <cuco/static_map.cuh>
 
-#include <utils.hpp>
+#include <thrust/device_vector.h>
+
+#include <catch2/catch.hpp>
 
 TEMPLATE_TEST_CASE_SIG("Unique sequence of keys",
                        "",
diff --git a/tests/static_multimap/custom_pair_retrieve_test.cu b/tests/static_multimap/custom_pair_retrieve_test.cu
index 322d5dc64..9c2915da5 100644
--- a/tests/static_multimap/custom_pair_retrieve_test.cu
+++ b/tests/static_multimap/custom_pair_retrieve_test.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,13 +14,17 @@
  * limitations under the License.
  */
 
-#include <catch2/catch.hpp>
+#include <utils.hpp>
+
+#include <cuco/static_multimap.cuh>
+
 #include <thrust/device_vector.h>
+#include <thrust/iterator/counting_iterator.h>
 #include <thrust/iterator/discard_iterator.h>
 
-#include <cuco/static_multimap.cuh>
+#include <catch2/catch.hpp>
 
-#include <utils.hpp>
+#include <cooperative_groups.h>
 
 // Custom pair equal
 template <typename Key, typename Value>
diff --git a/tests/static_multimap/custom_type_test.cu b/tests/static_multimap/custom_type_test.cu
index fd69cae8a..c734d3cb9 100644
--- a/tests/static_multimap/custom_type_test.cu
+++ b/tests/static_multimap/custom_type_test.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,12 +14,16 @@
  * limitations under the License.
  */
 
-#include <catch2/catch.hpp>
-#include <thrust/device_vector.h>
+#include <utils.hpp>
 
 #include <cuco/static_multimap.cuh>
 
-#include <utils.hpp>
+#include <thrust/device_vector.h>
+#include <thrust/iterator/counting_iterator.h>
+
+#include <catch2/catch.hpp>
+
+#include <tuple>
 
 // User-defined key type
 struct key_pair {
diff --git a/tests/static_multimap/insert_if_test.cu b/tests/static_multimap/insert_if_test.cu
index b3b3598dc..6465c5396 100644
--- a/tests/static_multimap/insert_if_test.cu
+++ b/tests/static_multimap/insert_if_test.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,12 +14,13 @@
  * limitations under the License.
  */
 
-#include <catch2/catch.hpp>
-#include <thrust/device_vector.h>
+#include <utils.hpp>
 
 #include <cuco/static_multimap.cuh>
 
-#include <utils.hpp>
+#include <thrust/device_vector.h>
+
+#include <catch2/catch.hpp>
 
 template <typename Key, typename Map, typename PairIt, typename KeyIt>
 __inline__ void test_insert_if(Map& map, PairIt pair_begin, KeyIt key_begin, std::size_t size)
diff --git a/tests/static_multimap/multiplicity_test.cu b/tests/static_multimap/multiplicity_test.cu
index 2bc04baec..de89bb1ea 100644
--- a/tests/static_multimap/multiplicity_test.cu
+++ b/tests/static_multimap/multiplicity_test.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,12 +14,13 @@
  * limitations under the License.
  */
 
-#include <catch2/catch.hpp>
-#include <thrust/device_vector.h>
+#include <utils.hpp>
 
 #include <cuco/static_multimap.cuh>
 
-#include <utils.hpp>
+#include <thrust/device_vector.h>
+
+#include <catch2/catch.hpp>
 
 template <typename Map>
 __inline__ void test_multiplicity_two(Map& map, std::size_t num_items)
diff --git a/tests/static_multimap/non_match_test.cu b/tests/static_multimap/non_match_test.cu
index e91ca70ba..9d0582fed 100644
--- a/tests/static_multimap/non_match_test.cu
+++ b/tests/static_multimap/non_match_test.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,12 +14,13 @@
  * limitations under the License.
  */
 
-#include <catch2/catch.hpp>
-#include <thrust/device_vector.h>
+#include <utils.hpp>
 
 #include <cuco/static_multimap.cuh>
 
-#include <utils.hpp>
+#include <thrust/device_vector.h>
+
+#include <catch2/catch.hpp>
 
 template <typename Key, typename Value, typename Map, typename PairIt, typename KeyIt>
 __inline__ void test_non_matches(Map& map, PairIt pair_begin, KeyIt key_begin, std::size_t num_keys)
diff --git a/tests/static_multimap/pair_function_test.cu b/tests/static_multimap/pair_function_test.cu
index 29073374b..66e3b1d27 100644
--- a/tests/static_multimap/pair_function_test.cu
+++ b/tests/static_multimap/pair_function_test.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,13 +14,14 @@
  * limitations under the License.
  */
 
-#include <catch2/catch.hpp>
-#include <thrust/device_vector.h>
-#include <thrust/iterator/discard_iterator.h>
+#include <utils.hpp>
 
 #include <cuco/static_multimap.cuh>
 
-#include <utils.hpp>
+#include <thrust/device_vector.h>
+#include <thrust/iterator/discard_iterator.h>
+
+#include <catch2/catch.hpp>
 
 // Custom pair equal
 template <typename Key, typename Value>
diff --git a/tests/utils.hpp b/tests/utils.hpp
index 37c7ac6f2..4b1ece781 100644
--- a/tests/utils.hpp
+++ b/tests/utils.hpp
@@ -16,9 +16,11 @@
 
 #pragma once
 
+#include <utils.cuh>
+
 #include <thrust/functional.h>
 
-#include <utils.cuh>
+#include <cooperative_groups.h>
 
 namespace cuco {
 namespace test {

From cc7dc149fda08c6025cdc0b36ab31b97114baa64 Mon Sep 17 00:00:00 2001
From: Yunsong Wang <yunsongw@nvidia.com>
Date: Wed, 2 Mar 2022 15:12:38 -0500
Subject: [PATCH 3/8] Update copyright year

---
 include/cuco/detail/pair.cuh | 2 +-
 tests/utils.hpp              | 2 +-
 2 files changed, 2 insertions(+), 2 deletions(-)

diff --git a/include/cuco/detail/pair.cuh b/include/cuco/detail/pair.cuh
index 1b013b771..7c3f55c25 100644
--- a/include/cuco/detail/pair.cuh
+++ b/include/cuco/detail/pair.cuh
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2021, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
diff --git a/tests/utils.hpp b/tests/utils.hpp
index 4b1ece781..dd2f6545f 100644
--- a/tests/utils.hpp
+++ b/tests/utils.hpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2021, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.

From 029edcb16a8d0b1f8250783048d1a9e107b5f524 Mon Sep 17 00:00:00 2001
From: Yunsong Wang <yunsongw@nvidia.com>
Date: Wed, 11 May 2022 17:26:57 -0400
Subject: [PATCH 4/8] Uncomment

---
 benchmarks/hash_table/static_map_bench.cu | 58 +++++++++++------------
 1 file changed, 28 insertions(+), 30 deletions(-)

diff --git a/benchmarks/hash_table/static_map_bench.cu b/benchmarks/hash_table/static_map_bench.cu
index eddc5bf07..a6c30c40c 100644
--- a/benchmarks/hash_table/static_map_bench.cu
+++ b/benchmarks/hash_table/static_map_bench.cu
@@ -202,63 +202,61 @@ static void BM_static_map_erase_all(::benchmark::State& state)
   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)
   ->UseManualTime();
-  */
 
 BENCHMARK_TEMPLATE(BM_static_map_search_all, int32_t, int32_t, dist_type::UNIQUE)
   ->Unit(benchmark::kMillisecond)
   ->Apply(generate_size_and_occupancy);
-/*
 
 BENCHMARK_TEMPLATE(BM_static_map_insert, int32_t, int32_t, dist_type::UNIFORM)
-->Unit(benchmark::kMillisecond)
-->Apply(generate_size_and_occupancy)
-->UseManualTime();
+  ->Unit(benchmark::kMillisecond)
+  ->Apply(generate_size_and_occupancy)
+  ->UseManualTime();
 
 BENCHMARK_TEMPLATE(BM_static_map_search_all, int32_t, int32_t, dist_type::UNIFORM)
-->Unit(benchmark::kMillisecond)
-->Apply(generate_size_and_occupancy);
+  ->Unit(benchmark::kMillisecond)
+  ->Apply(generate_size_and_occupancy);
 
 BENCHMARK_TEMPLATE(BM_static_map_insert, int32_t, int32_t, dist_type::GAUSSIAN)
-->Unit(benchmark::kMillisecond)
-->Apply(generate_size_and_occupancy)
-->UseManualTime();
+  ->Unit(benchmark::kMillisecond)
+  ->Apply(generate_size_and_occupancy)
+  ->UseManualTime();
 
 BENCHMARK_TEMPLATE(BM_static_map_search_all, int32_t, int32_t, dist_type::GAUSSIAN)
-->Unit(benchmark::kMillisecond)
-->Apply(generate_size_and_occupancy);
+  ->Unit(benchmark::kMillisecond)
+  ->Apply(generate_size_and_occupancy);
 
 BENCHMARK_TEMPLATE(BM_static_map_insert, int64_t, int64_t, dist_type::UNIQUE)
-->Unit(benchmark::kMillisecond)
-->Apply(generate_size_and_occupancy)
-->UseManualTime();
+  ->Unit(benchmark::kMillisecond)
+  ->Apply(generate_size_and_occupancy)
+  ->UseManualTime();
 
 BENCHMARK_TEMPLATE(BM_static_map_search_all, int64_t, int64_t, dist_type::UNIQUE)
-->Unit(benchmark::kMillisecond)
-->Apply(generate_size_and_occupancy);
+  ->Unit(benchmark::kMillisecond)
+  ->Apply(generate_size_and_occupancy);
 
 BENCHMARK_TEMPLATE(BM_static_map_insert, int64_t, int64_t, dist_type::UNIFORM)
-->Unit(benchmark::kMillisecond)
-->Apply(generate_size_and_occupancy)
-->UseManualTime();
+  ->Unit(benchmark::kMillisecond)
+  ->Apply(generate_size_and_occupancy)
+  ->UseManualTime();
 
 BENCHMARK_TEMPLATE(BM_static_map_search_all, int64_t, int64_t, dist_type::UNIFORM)
-->Unit(benchmark::kMillisecond)
-->Apply(generate_size_and_occupancy);
+  ->Unit(benchmark::kMillisecond)
+  ->Apply(generate_size_and_occupancy);
 
 BENCHMARK_TEMPLATE(BM_static_map_insert, int64_t, int64_t, dist_type::GAUSSIAN)
-->Unit(benchmark::kMillisecond)
-->Apply(generate_size_and_occupancy)
-->UseManualTime();
+  ->Unit(benchmark::kMillisecond)
+  ->Apply(generate_size_and_occupancy)
+  ->UseManualTime();
 
 BENCHMARK_TEMPLATE(BM_static_map_search_all, int64_t, int64_t, dist_type::GAUSSIAN)
-->Unit(benchmark::kMillisecond)
-->Apply(generate_size_and_occupancy);
+  ->Unit(benchmark::kMillisecond)
+  ->Apply(generate_size_and_occupancy);
 
 BENCHMARK_TEMPLATE(BM_static_map_erase_all, int32_t, int32_t, dist_type::UNIQUE)
-->Unit(benchmark::kMillisecond)
-->Apply(generate_size_and_occupancy);*/
+  ->Unit(benchmark::kMillisecond)
+  ->Apply(generate_size_and_occupancy);

From 3f1fee52b60c09a5d0696fcaca13a7451ef61285 Mon Sep 17 00:00:00 2001
From: Yunsong Wang <yunsongw@nvidia.com>
Date: Wed, 11 May 2022 17:32:59 -0400
Subject: [PATCH 5/8] Include more headers

---
 include/cuco/static_multimap.cuh | 3 +++
 1 file changed, 3 insertions(+)

diff --git a/include/cuco/static_multimap.cuh b/include/cuco/static_multimap.cuh
index 2953b5dcf..5f7ef35e6 100644
--- a/include/cuco/static_multimap.cuh
+++ b/include/cuco/static_multimap.cuh
@@ -42,7 +42,10 @@
 
 #include <cooperative_groups.h>
 
+#include <cstddef>
 #include <memory>
+#include <type_traits>
+#include <utility>
 
 namespace cuco {
 

From 9e90d87a486302a7c31252282ee79eb031c93fbd Mon Sep 17 00:00:00 2001
From: Yunsong Wang <yunsongw@nvidia.com>
Date: Wed, 11 May 2022 17:36:13 -0400
Subject: [PATCH 6/8] Make sentinel ctor constexpr

---
 include/cuco/sentinel.hpp | 8 ++++----
 1 file changed, 4 insertions(+), 4 deletions(-)

diff --git a/include/cuco/sentinel.hpp b/include/cuco/sentinel.hpp
index 53814c438..2dc28b1b0 100644
--- a/include/cuco/sentinel.hpp
+++ b/include/cuco/sentinel.hpp
@@ -21,21 +21,21 @@ namespace sentinel {
 
 template <typename T>
 struct empty_key {
-  __host__ __device__ empty_key(T v) : value{v} {}
+  constexpr empty_key(T v) : value{v} {}
   T value;
 };
 
 template <typename T>
 struct empty_value {
-  __host__ __device__ empty_value(T v) : value{v} {}
+  constexpr empty_value(T v) : value{v} {}
   T value;
 };
 
 template <typename T>
 struct erased_key {
-  __host__ __device__ erased_key(T v) : value{v} {}
+  constexpr erased_key(T v) : value{v} {}
   T value;
 };
 
 }  // namespace sentinel
-}  // namespace cuco
\ No newline at end of file
+}  // namespace cuco

From 48688ced72332204dac34a614ea20040845de7e8 Mon Sep 17 00:00:00 2001
From: Yunsong Wang <yunsongw@nvidia.com>
Date: Thu, 12 May 2022 18:09:27 -0400
Subject: [PATCH 7/8] Include more headers in dynamic_map

---
 include/cuco/dynamic_map.cuh | 3 +++
 1 file changed, 3 insertions(+)

diff --git a/include/cuco/dynamic_map.cuh b/include/cuco/dynamic_map.cuh
index 9867aa7f3..3357ece11 100644
--- a/include/cuco/dynamic_map.cuh
+++ b/include/cuco/dynamic_map.cuh
@@ -25,7 +25,10 @@
 
 #include <cuda/std/atomic>
 
+#include <cstddef>
 #include <memory>
+#include <type_traits>
+#include <vector>
 
 namespace cuco {
 

From ca5d052251f85a7785d2aa51a463854979f02c61 Mon Sep 17 00:00:00 2001
From: Yunsong Wang <yunsongw@nvidia.com>
Date: Wed, 18 May 2022 12:38:31 -0400
Subject: [PATCH 8/8] Cleanups: use proper headers

---
 include/cuco/detail/dynamic_map_kernels.cuh             | 2 +-
 include/cuco/detail/pair.cuh                            | 1 +
 include/cuco/detail/static_map_kernels.cuh              | 2 +-
 include/cuco/detail/static_multimap/kernels.cuh         | 2 +-
 include/cuco/detail/static_multimap/static_multimap.inl | 2 ++
 5 files changed, 6 insertions(+), 3 deletions(-)

diff --git a/include/cuco/detail/dynamic_map_kernels.cuh b/include/cuco/detail/dynamic_map_kernels.cuh
index a18aee736..4bb514618 100644
--- a/include/cuco/detail/dynamic_map_kernels.cuh
+++ b/include/cuco/detail/dynamic_map_kernels.cuh
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include <cub/cub.cuh>
+#include <cub/block/block_reduce.cuh>
 
 #include <cuda/std/atomic>
 
diff --git a/include/cuco/detail/pair.cuh b/include/cuco/detail/pair.cuh
index 7c3f55c25..fe7e7d7e8 100644
--- a/include/cuco/detail/pair.cuh
+++ b/include/cuco/detail/pair.cuh
@@ -20,6 +20,7 @@
 #include <thrust/pair.h>
 #include <thrust/tuple.h>
 
+#include <algorithm>
 #include <tuple>
 #include <type_traits>
 
diff --git a/include/cuco/detail/static_map_kernels.cuh b/include/cuco/detail/static_map_kernels.cuh
index 091a08a91..4f4daa0d5 100644
--- a/include/cuco/detail/static_map_kernels.cuh
+++ b/include/cuco/detail/static_map_kernels.cuh
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include <cub/cub.cuh>
+#include <cub/block/block_reduce.cuh>
 
 #include <cuda/std/atomic>
 
diff --git a/include/cuco/detail/static_multimap/kernels.cuh b/include/cuco/detail/static_multimap/kernels.cuh
index e3990a373..3d0cbb2f0 100644
--- a/include/cuco/detail/static_multimap/kernels.cuh
+++ b/include/cuco/detail/static_multimap/kernels.cuh
@@ -18,7 +18,7 @@
 
 #include <thrust/type_traits/is_contiguous_iterator.h>
 
-#include <cub/cub.cuh>
+#include <cub/block/block_reduce.cuh>
 
 #include <cuda/std/atomic>
 
diff --git a/include/cuco/detail/static_multimap/static_multimap.inl b/include/cuco/detail/static_multimap/static_multimap.inl
index 3fea8478a..175852acf 100644
--- a/include/cuco/detail/static_multimap/static_multimap.inl
+++ b/include/cuco/detail/static_multimap/static_multimap.inl
@@ -20,6 +20,8 @@
 #include <thrust/iterator/transform_iterator.h>
 #include <thrust/tuple.h>
 
+#include <iterator>
+
 namespace {
 /**
  * @brief Device functor used to determine if a slot is filled.