From e7f0ce2e8650c31c1871a4dbd33287d4408892ef Mon Sep 17 00:00:00 2001
From: Krish Sapru <ksapru@bu.edu>
Date: Wed, 29 Apr 2026 12:28:59 -0400
Subject: [PATCH 01/10] Fix dangling allocator reference in storage deleters

---
 include/cuco/bucket_storage.cuh              | 4 ++--
 include/cuco/detail/storage/storage_base.cuh | 4 ++--
 2 files changed, 4 insertions(+), 4 deletions(-)

diff --git a/include/cuco/bucket_storage.cuh b/include/cuco/bucket_storage.cuh
index 9d99e3d7b..240f95e9b 100644
--- a/include/cuco/bucket_storage.cuh
+++ b/include/cuco/bucket_storage.cuh
@@ -252,10 +252,10 @@ class bucket_storage {
   struct aligned_deleter {
     value_type* raw_ptr_;
     std::size_t size_;
-    allocator_type& allocator_;
+    allocator_type allocator_;
     cuda::stream_ref stream_;
 
-    void operator()(value_type*) const { allocator_.deallocate(raw_ptr_, size_, stream_); }
+    void operator()(value_type*) { allocator_.deallocate(raw_ptr_, size_, stream_); }
   };
 
   extent_type extent_;        ///< Storage extent
diff --git a/include/cuco/detail/storage/storage_base.cuh b/include/cuco/detail/storage/storage_base.cuh
index 0f8e013f4..84097da62 100644
--- a/include/cuco/detail/storage/storage_base.cuh
+++ b/include/cuco/detail/storage/storage_base.cuh
@@ -39,7 +39,7 @@ struct custom_deleter {
    * @param allocator Allocator used for deallocating device storage
    * @param stream Stream to use for deallocation
    */
-  explicit constexpr custom_deleter(SizeType size, Allocator& allocator, cuda::stream_ref stream)
+  explicit constexpr custom_deleter(SizeType size, Allocator const& allocator, cuda::stream_ref stream)
     : size_{size}, allocator_{allocator}, stream_{stream}
   {
   }
@@ -52,7 +52,7 @@ struct custom_deleter {
   void operator()(pointer ptr) { allocator_.deallocate(ptr, size_, stream_); }
 
   SizeType size_;            ///< Number of values to delete
-  Allocator& allocator_;     ///< Allocator used deallocating values
+  Allocator allocator_;      ///< Allocator used deallocating values
   cuda::stream_ref stream_;  ///< Stream used for deallocation
 };
 

From 3d06afd05125a44a2121963f22340817e6bdf603 Mon Sep 17 00:00:00 2001
From: Krish Sapru <ksapru@bu.edu>
Date: Fri, 1 May 2026 16:45:53 -0400
Subject: [PATCH 02/10] test(static_set): add uint8_t and uint16_t to type axes

---
 tests/static_set/for_each_test.cu        | 27 ++++++++++++++++++------
 tests/static_set/insert_and_find_test.cu | 11 +++++++++-
 tests/static_set/retrieve_all_test.cu    | 11 +++++++++-
 tests/static_set/retrieve_if_test.cu     |  7 +++++-
 tests/static_set/retrieve_test.cu        | 11 +++++++++-
 tests/static_set/shared_memory_test.cu   | 11 ++++++++--
 tests/static_set/stream_test.cu          |  8 ++++++-
 tests/static_set/unique_sequence_test.cu | 11 +++++++++-
 8 files changed, 83 insertions(+), 14 deletions(-)

diff --git a/tests/static_set/for_each_test.cu b/tests/static_set/for_each_test.cu
index 5374f8fac..8ad84d482 100644
--- a/tests/static_set/for_each_test.cu
+++ b/tests/static_set/for_each_test.cu
@@ -16,6 +16,8 @@
 
 #include <test_utils.hpp>
 
+#include <cstdint>
+
 #include <cuco/detail/__config>
 #include <cuco/static_set.cuh>
 
@@ -29,7 +31,7 @@
 using size_type = std::size_t;
 
 template <typename Set>
-void test_for_each(Set& set, size_type num_keys)
+void test_for_each(Set& set, size_type num_keys, size_type expected_evens, size_type expected_odds)
 {
   using Key = typename Set::key_type;
 
@@ -40,7 +42,7 @@ void test_for_each(Set& set, size_type num_keys)
   // Insert keys
   auto keys_begin = cuda::make_transform_iterator(
     cuda::counting_iterator<size_type>{0}, cuda::proclaim_return_type<Key>([] __device__(auto i) {
-      // generates a sequence of 1, 2, 1, 2, ...
+      // generates a sequence of 0, 1, 2, ...
       return static_cast<Key>(i);
     }));
   set.insert(keys_begin, keys_begin + num_keys, stream);
@@ -56,7 +58,7 @@ void test_for_each(Set& set, size_type num_keys)
       if (slot % 2 == 0) { counter->fetch_add(slot, cuda::memory_order_relaxed); }
     },
     stream);
-  REQUIRE(counter_storage.load_to_host(stream) == 249'500);
+  REQUIRE(counter_storage.load_to_host(stream) == expected_evens);
 
   counter_storage.reset(stream);
 
@@ -68,13 +70,19 @@ void test_for_each(Set& set, size_type num_keys)
       if (!(slot % 2 == 0)) { counter->fetch_add(slot, cuda::memory_order_relaxed); }
     },
     stream);
-  REQUIRE(counter_storage.load_to_host(stream) == 250'000);
+  REQUIRE(counter_storage.load_to_host(stream) == expected_odds);
 }
 
 TEMPLATE_TEST_CASE_SIG(
   "static_set for_each tests",
   "",
   ((typename Key, cuco::test::probe_sequence Probe, int CGSize), Key, Probe, CGSize),
+  (uint8_t, cuco::test::probe_sequence::double_hashing, 1),
+  (uint8_t, cuco::test::probe_sequence::linear_probing, 1),
+  (uint16_t, cuco::test::probe_sequence::double_hashing, 1),
+  (uint16_t, cuco::test::probe_sequence::double_hashing, 2),
+  (uint16_t, cuco::test::probe_sequence::linear_probing, 1),
+  (uint16_t, cuco::test::probe_sequence::linear_probing, 2),
   (int32_t, cuco::test::probe_sequence::double_hashing, 1),
   (int32_t, cuco::test::probe_sequence::double_hashing, 2),
   (int64_t, cuco::test::probe_sequence::double_hashing, 1),
@@ -92,7 +100,14 @@ TEMPLATE_TEST_CASE_SIG(
 #endif
 )
 {
-  constexpr size_type num_keys{1'000};
+  // Limit key count for small types: leave room for the 0xFF/0xFFFF sentinel.
+  // Expected sums are pre-computed per type class:
+  //   uint8_t  (num_keys=100): sum of evens 0..98 = 2450, sum of odds 1..99 = 2500
+  //   uint16_t+ (num_keys=1000): sum of evens 0..998 = 249'500, sum of odds 1..999 = 250'000
+  constexpr size_type num_keys         = (sizeof(Key) == 1) ? 100 : 1'000;
+  constexpr size_type expected_evens   = (sizeof(Key) == 1) ? 2'450 : 249'500;
+  constexpr size_type expected_odds    = (sizeof(Key) == 1) ? 2'500 : 250'000;
+
   using probe = std::conditional_t<
     Probe == cuco::test::probe_sequence::linear_probing,
     cuco::linear_probing<CGSize, cuco::murmurhash3_32<Key>>,
@@ -107,5 +122,5 @@ TEMPLATE_TEST_CASE_SIG(
                                  cuco::storage<2>>;
 
   auto set = set_t{num_keys, cuco::empty_key<Key>{static_cast<Key>(-1)}};
-  test_for_each(set, num_keys);
+  test_for_each(set, num_keys, expected_evens, expected_odds);
 }
diff --git a/tests/static_set/insert_and_find_test.cu b/tests/static_set/insert_and_find_test.cu
index 7eefdfb6c..791244b39 100644
--- a/tests/static_set/insert_and_find_test.cu
+++ b/tests/static_set/insert_and_find_test.cu
@@ -16,6 +16,8 @@
 
 #include <test_utils.hpp>
 
+#include <cstdint>
+
 #include <cuco/detail/__config>
 #include <cuco/static_set.cuh>
 
@@ -56,6 +58,12 @@ TEMPLATE_TEST_CASE_SIG(
   "static_set Insert and find",
   "",
   ((typename Key, cuco::test::probe_sequence Probe, int CGSize), Key, Probe, CGSize),
+  (uint8_t, cuco::test::probe_sequence::double_hashing, 1),
+  (uint8_t, cuco::test::probe_sequence::linear_probing, 1),
+  (uint16_t, cuco::test::probe_sequence::double_hashing, 1),
+  (uint16_t, cuco::test::probe_sequence::double_hashing, 2),
+  (uint16_t, cuco::test::probe_sequence::linear_probing, 1),
+  (uint16_t, cuco::test::probe_sequence::linear_probing, 2),
   (int32_t, cuco::test::probe_sequence::double_hashing, 1),
   (int32_t, cuco::test::probe_sequence::double_hashing, 2),
   (int64_t, cuco::test::probe_sequence::double_hashing, 1),
@@ -73,7 +81,8 @@ TEMPLATE_TEST_CASE_SIG(
 #endif
 )
 {
-  constexpr std::size_t num_keys{400};
+  // Limit key count for small types: leave room for the 0xFF/0xFFFF sentinel
+  constexpr std::size_t num_keys = (sizeof(Key) == 1) ? 100 : 400;
 
   using probe = std::conditional_t<Probe == cuco::test::probe_sequence::linear_probing,
                                    cuco::linear_probing<CGSize, cuco::default_hash_function<Key>>,
diff --git a/tests/static_set/retrieve_all_test.cu b/tests/static_set/retrieve_all_test.cu
index ed2e70abb..0be1c08a7 100644
--- a/tests/static_set/retrieve_all_test.cu
+++ b/tests/static_set/retrieve_all_test.cu
@@ -16,6 +16,8 @@
 
 #include <test_utils.hpp>
 
+#include <cstdint>
+
 #include <cuco/detail/__config>
 #include <cuco/static_set.cuh>
 
@@ -61,6 +63,12 @@ TEMPLATE_TEST_CASE_SIG(
   "static_set::retrieve_all tests",
   "",
   ((typename Key, cuco::test::probe_sequence Probe, int CGSize), Key, Probe, CGSize),
+  (uint8_t, cuco::test::probe_sequence::double_hashing, 1),
+  (uint8_t, cuco::test::probe_sequence::linear_probing, 1),
+  (uint16_t, cuco::test::probe_sequence::double_hashing, 1),
+  (uint16_t, cuco::test::probe_sequence::double_hashing, 2),
+  (uint16_t, cuco::test::probe_sequence::linear_probing, 1),
+  (uint16_t, cuco::test::probe_sequence::linear_probing, 2),
   (int32_t, cuco::test::probe_sequence::double_hashing, 1),
   (int32_t, cuco::test::probe_sequence::double_hashing, 2),
   (int64_t, cuco::test::probe_sequence::double_hashing, 1),
@@ -78,7 +86,8 @@ TEMPLATE_TEST_CASE_SIG(
 #endif
 )
 {
-  constexpr std::size_t num_keys{400};
+  // Limit key count for small types: leave room for the 0xFF/0xFFFF sentinel
+  constexpr std::size_t num_keys        = (sizeof(Key) == 1) ? 100 : 400;
   constexpr double desired_load_factor = 1.;
 
   using probe = std::conditional_t<Probe == cuco::test::probe_sequence::linear_probing,
diff --git a/tests/static_set/retrieve_if_test.cu b/tests/static_set/retrieve_if_test.cu
index b82969926..d97222e05 100644
--- a/tests/static_set/retrieve_if_test.cu
+++ b/tests/static_set/retrieve_if_test.cu
@@ -16,6 +16,8 @@
 
 #include <test_utils.hpp>
 
+#include <cstdint>
+
 #include <cuco/detail/__config>
 #include <cuco/static_set.cuh>
 
@@ -110,6 +112,8 @@ __global__ void test_retrieve_if_all_true_kernel(
 TEMPLATE_TEST_CASE_SIG("static_set retrieve_if",
                        "",
                        ((typename Key), Key),
+                       (uint8_t),
+                       (uint16_t),
                        (int32_t),
                        (int64_t)
 #if defined(CUCO_HAS_128BIT_ATOMICS)
@@ -118,7 +122,8 @@ TEMPLATE_TEST_CASE_SIG("static_set retrieve_if",
 #endif
 )
 {
-  constexpr size_type num_keys{400};
+  // Limit key count for small types: keys start at 1, sentinel is 0xFF/0xFFFF
+  constexpr size_type num_keys = (sizeof(Key) == 1) ? 100 : 400;
 
   using container_type = cuco::static_set<Key>;
 
diff --git a/tests/static_set/retrieve_test.cu b/tests/static_set/retrieve_test.cu
index f5b7763dc..68cf44f25 100644
--- a/tests/static_set/retrieve_test.cu
+++ b/tests/static_set/retrieve_test.cu
@@ -16,6 +16,8 @@
 
 #include <test_utils.hpp>
 
+#include <cstdint>
+
 #include <cuco/detail/__config>
 #include <cuco/static_set.cuh>
 
@@ -75,6 +77,12 @@ TEMPLATE_TEST_CASE_SIG(
   "static_set retrieve tests",
   "",
   ((typename Key, cuco::test::probe_sequence Probe, int CGSize), Key, Probe, CGSize),
+  (uint8_t, cuco::test::probe_sequence::double_hashing, 1),
+  (uint8_t, cuco::test::probe_sequence::linear_probing, 1),
+  (uint16_t, cuco::test::probe_sequence::double_hashing, 1),
+  (uint16_t, cuco::test::probe_sequence::double_hashing, 2),
+  (uint16_t, cuco::test::probe_sequence::linear_probing, 1),
+  (uint16_t, cuco::test::probe_sequence::linear_probing, 2),
   (int32_t, cuco::test::probe_sequence::double_hashing, 1),
   (int32_t, cuco::test::probe_sequence::double_hashing, 2),
   (int64_t, cuco::test::probe_sequence::double_hashing, 1),
@@ -92,7 +100,8 @@ TEMPLATE_TEST_CASE_SIG(
 #endif
 )
 {
-  constexpr std::size_t num_keys{400};
+  // Limit key count for small types: leave room for the 0xFF/0xFFFF sentinel
+  constexpr std::size_t num_keys        = (sizeof(Key) == 1) ? 100 : 400;
   constexpr double desired_load_factor = 1.;
 
   using probe = std::conditional_t<Probe == cuco::test::probe_sequence::linear_probing,
diff --git a/tests/static_set/shared_memory_test.cu b/tests/static_set/shared_memory_test.cu
index 08f805927..329c9ddf3 100644
--- a/tests/static_set/shared_memory_test.cu
+++ b/tests/static_set/shared_memory_test.cu
@@ -16,6 +16,8 @@
 
 #include <test_utils.hpp>
 
+#include <cstdint>
+
 #include <cuco/detail/__config>
 #include <cuco/static_set.cuh>
 
@@ -67,6 +69,8 @@ __global__ void shared_memory_test_kernel(Ref* sets,
 TEMPLATE_TEST_CASE_SIG("static_set shared memory tests",
                        "",
                        ((typename Key), Key),
+                       (uint8_t),
+                       (uint16_t),
                        (int32_t),
                        (int64_t)
 #if defined(CUCO_HAS_128BIT_ATOMICS)
@@ -75,8 +79,11 @@ TEMPLATE_TEST_CASE_SIG("static_set shared memory tests",
 #endif
 )
 {
-  constexpr std::size_t number_of_sets  = 1000;
-  constexpr std::size_t elements_in_set = 500;
+  // For uint8_t: sentinel = 0xFF (255), so usable key range is 0..254.
+  // thrust::sequence over number_of_sets*elements_in_set keys must not wrap.
+  // Use a smaller set count and element count for 1-byte types.
+  constexpr std::size_t number_of_sets  = (sizeof(Key) == 1) ? 2 : 1000;
+  constexpr std::size_t elements_in_set = (sizeof(Key) == 1) ? 100 : 500;
   constexpr std::size_t set_capacity    = 2 * elements_in_set;
 
   using extent_type = cuco::extent<std::size_t, set_capacity>;
diff --git a/tests/static_set/stream_test.cu b/tests/static_set/stream_test.cu
index f9277ed8f..2413bd578 100644
--- a/tests/static_set/stream_test.cu
+++ b/tests/static_set/stream_test.cu
@@ -16,6 +16,8 @@
 
 #include <test_utils.hpp>
 
+#include <cstdint>
+
 #include <cuco/detail/__config>
 #include <cuco/static_set.cuh>
 
@@ -33,6 +35,8 @@
 TEMPLATE_TEST_CASE_SIG("static_set: operations on different stream than constructor",
                        "",
                        ((typename Key), Key),
+                       (uint8_t),
+                       (uint16_t),
                        (int32_t),
                        (int64_t)
 #if defined(CUCO_HAS_128BIT_ATOMICS)
@@ -47,7 +51,9 @@ TEMPLATE_TEST_CASE_SIG("static_set: operations on different stream than construc
   CUCO_CUDA_TRY(cudaStreamCreate(&operation_stream));
 
   {  // Scope ensures set is destroyed before streams
-    constexpr std::size_t num_keys{500'000};
+    // Scale num_keys to fit in the key type's value range (sentinel = 0xFF/0xFFFF).
+    constexpr std::size_t num_keys =
+      (sizeof(Key) == 1) ? 100 : (sizeof(Key) == 2) ? 1'000 : 500'000;
     auto set = cuco::static_set{num_keys * 2,
                                 cuco::empty_key<Key>{static_cast<Key>(-1)},
                                 {},
diff --git a/tests/static_set/unique_sequence_test.cu b/tests/static_set/unique_sequence_test.cu
index 2db91c2b6..2774c5e60 100644
--- a/tests/static_set/unique_sequence_test.cu
+++ b/tests/static_set/unique_sequence_test.cu
@@ -16,6 +16,8 @@
 
 #include <test_utils.hpp>
 
+#include <cstdint>
+
 #include <cuco/detail/__config>
 #include <cuco/static_set.cuh>
 
@@ -141,6 +143,12 @@ TEMPLATE_TEST_CASE_SIG(
   "static_set unique sequence tests",
   "",
   ((typename Key, cuco::test::probe_sequence Probe, int CGSize), Key, Probe, CGSize),
+  (uint8_t, cuco::test::probe_sequence::double_hashing, 1),
+  (uint8_t, cuco::test::probe_sequence::linear_probing, 1),
+  (uint16_t, cuco::test::probe_sequence::double_hashing, 1),
+  (uint16_t, cuco::test::probe_sequence::double_hashing, 2),
+  (uint16_t, cuco::test::probe_sequence::linear_probing, 1),
+  (uint16_t, cuco::test::probe_sequence::linear_probing, 2),
   (int32_t, cuco::test::probe_sequence::double_hashing, 1),
   (int32_t, cuco::test::probe_sequence::double_hashing, 2),
   (int64_t, cuco::test::probe_sequence::double_hashing, 1),
@@ -158,7 +166,8 @@ TEMPLATE_TEST_CASE_SIG(
 #endif
 )
 {
-  constexpr size_type num_keys{400};
+  // Limit key count for small types: leave room for the 0xFF/0xFFFF sentinel
+  constexpr size_type num_keys = (sizeof(Key) == 1) ? 100 : 400;
   using probe = std::conditional_t<Probe == cuco::test::probe_sequence::linear_probing,
                                    cuco::linear_probing<CGSize, cuco::default_hash_function<Key>>,
                                    cuco::double_hashing<CGSize, cuco::default_hash_function<Key>>>;

From e77f232c0c8dd83f1cdde15f8b81d600e6b61e2d Mon Sep 17 00:00:00 2001
From: Krish Sapru <ksapru@bu.edu>
Date: Fri, 1 May 2026 16:53:38 -0400
Subject: [PATCH 03/10] test(static_set): add additional edge case

---
 include/cuco/detail/pair/helpers.cuh     | 16 ++++++++++++++++
 tests/static_set/retrieve_all_test.cu    | 11 ++++++++---
 tests/static_set/unique_sequence_test.cu | 11 ++++++++---
 3 files changed, 32 insertions(+), 6 deletions(-)

diff --git a/include/cuco/detail/pair/helpers.cuh b/include/cuco/detail/pair/helpers.cuh
index c21cbdb86..b8f4d8f70 100644
--- a/include/cuco/detail/pair/helpers.cuh
+++ b/include/cuco/detail/pair/helpers.cuh
@@ -47,6 +47,22 @@ struct packed {
   using type = void;  ///< `void` type by default
 };
 
+/**
+ * @brief Denotes the packed type when the size of the object is 1.
+ */
+template <>
+struct packed<sizeof(uint8_t)> {
+  using type = uint8_t;  ///< Packed type as `uint8_t` if the size of the object is 1
+};
+
+/**
+ * @brief Denotes the packed type when the size of the object is 2.
+ */
+template <>
+struct packed<sizeof(uint16_t)> {
+  using type = uint16_t;  ///< Packed type as `uint16_t` if the size of the object is 2
+};
+
 /**
  * @brief Denotes the packed type when the size of the object is 8.
  */
diff --git a/tests/static_set/retrieve_all_test.cu b/tests/static_set/retrieve_all_test.cu
index 0be1c08a7..4e2e53529 100644
--- a/tests/static_set/retrieve_all_test.cu
+++ b/tests/static_set/retrieve_all_test.cu
@@ -96,10 +96,15 @@ TEMPLATE_TEST_CASE_SIG(
 
   constexpr std::size_t gold_capacity = [&]() {
     if constexpr (cuco::is_double_hashing<probe>::value) {
-      return (CGSize == 1) ? 401   // 401 x 1 x 1
-                           : 422;  // 211 x 2 x 1
+      if constexpr (num_keys == 100) {
+        return (CGSize == 1) ? 101   // 101 x 1 x 1
+                             : 106;  // 53 x 2 x 1
+      } else {
+        return (CGSize == 1) ? 401   // 401 x 1 x 1
+                             : 422;  // 211 x 2 x 1
+      }
     } else {
-      return 400;
+      return num_keys;
     }
   }();
 
diff --git a/tests/static_set/unique_sequence_test.cu b/tests/static_set/unique_sequence_test.cu
index 2774c5e60..a1b6663cf 100644
--- a/tests/static_set/unique_sequence_test.cu
+++ b/tests/static_set/unique_sequence_test.cu
@@ -174,10 +174,15 @@ TEMPLATE_TEST_CASE_SIG(
 
   constexpr size_type gold_capacity = [&]() {
     if constexpr (cuco::is_double_hashing<probe>::value) {
-      return (CGSize == 1) ? 422   // 211 x 1 x 2
-                           : 404;  // 101 x 2 x 2
+      if constexpr (num_keys == 100) {
+        return (CGSize == 1) ? 106   // 53 x 1 x 2
+                             : 106;  // 53 x 2 x 1... only CGSize=1 used for uint8_t
+      } else {
+        return (CGSize == 1) ? 422   // 211 x 1 x 2
+                             : 404;  // 101 x 2 x 2
+      }
     } else {
-      return 400;
+      return num_keys;
     }
   }();
 

From 22ad7ae29733a95cf3ddafa979b9863962c6653d Mon Sep 17 00:00:00 2001
From: Krish Sapru <ksapru@bu.edu>
Date: Sun, 3 May 2026 08:45:10 -0400
Subject: [PATCH 04/10] Fix sign conversion error

---
 tests/static_set/atomic_storage_test.cu |  1 +
 tests/static_set/shared_memory_test.cu  | 12 ++++++++----
 2 files changed, 9 insertions(+), 4 deletions(-)

diff --git a/tests/static_set/atomic_storage_test.cu b/tests/static_set/atomic_storage_test.cu
index b611214bb..746608428 100644
--- a/tests/static_set/atomic_storage_test.cu
+++ b/tests/static_set/atomic_storage_test.cu
@@ -69,6 +69,7 @@ TEST_CASE("atomic_storage_test", "")
                                                   cuda::proclaim_return_type<Key>(build_fn{}));
 
   set.insert_async(keys_begin, keys_begin + num_keys);
+  CUCO_CUDA_TRY(cudaDeviceSynchronize());
   auto const count = set.size();
 
   REQUIRE(count == num_keys);
diff --git a/tests/static_set/shared_memory_test.cu b/tests/static_set/shared_memory_test.cu
index 329c9ddf3..345c769c6 100644
--- a/tests/static_set/shared_memory_test.cu
+++ b/tests/static_set/shared_memory_test.cu
@@ -49,6 +49,8 @@ __global__ void shared_memory_test_kernel(Ref* sets,
   auto insert_ref = sets[set_id].make_copy(g, sm_buffer, cuco::thread_scope_block);
   auto find_ref   = insert_ref.rebind_operators(cuco::op::find);
 
+  __syncthreads();
+
   for (int i = g.thread_rank(); i < number_of_elements; i += g.size()) {
     auto found_it = find_ref.find(insterted_keys[offset + i]);
 
@@ -80,10 +82,11 @@ TEMPLATE_TEST_CASE_SIG("static_set shared memory tests",
 )
 {
   // For uint8_t: sentinel = 0xFF (255), so usable key range is 0..254.
+  // For uint16_t: sentinel = 0xFFFF (65535), so usable key range is 0..65534.
   // thrust::sequence over number_of_sets*elements_in_set keys must not wrap.
-  // Use a smaller set count and element count for 1-byte types.
-  constexpr std::size_t number_of_sets  = (sizeof(Key) == 1) ? 2 : 1000;
-  constexpr std::size_t elements_in_set = (sizeof(Key) == 1) ? 100 : 500;
+  // Use smaller set count and element count for smaller types.
+  constexpr std::size_t number_of_sets  = (sizeof(Key) <= 2) ? (sizeof(Key) == 1 ? 2 : 100) : 1000;
+  constexpr std::size_t elements_in_set = (sizeof(Key) <= 2) ? (sizeof(Key) == 1 ? 100 : 600) : 500;
   constexpr std::size_t set_capacity    = 2 * elements_in_set;
 
   using extent_type = cuco::extent<std::size_t, set_capacity>;
@@ -103,7 +106,8 @@ TEMPLATE_TEST_CASE_SIG("static_set shared memory tests",
   // operator yet
   std::vector<std::unique_ptr<set_type>> sets;
   for (std::size_t set_id = 0; set_id < number_of_sets; ++set_id) {
-    sets.push_back(std::make_unique<set_type>(extent_type{}, cuco::empty_key<Key>{-1}));
+    sets.push_back(
+      std::make_unique<set_type>(extent_type{}, cuco::empty_key<Key>{static_cast<Key>(-1)}));
   }
 
   thrust::device_vector<bool> d_keys_exist(number_of_sets * elements_in_set);

From 93800bfc24f7f596067bda12a6306501e9cd4b6c Mon Sep 17 00:00:00 2001
From: Krish Sapru <ksapru@bu.edu>
Date: Thu, 7 May 2026 11:33:21 -0400
Subject: [PATCH 05/10] Add comments

---
 include/cuco/detail/pair/helpers.cuh    |  8 -------
 tests/static_map/find_test.cu           | 32 +++++++++++++++++++++++++
 tests/static_set/atomic_storage_test.cu |  1 -
 3 files changed, 32 insertions(+), 9 deletions(-)

diff --git a/include/cuco/detail/pair/helpers.cuh b/include/cuco/detail/pair/helpers.cuh
index b8f4d8f70..c1a533bce 100644
--- a/include/cuco/detail/pair/helpers.cuh
+++ b/include/cuco/detail/pair/helpers.cuh
@@ -47,14 +47,6 @@ struct packed {
   using type = void;  ///< `void` type by default
 };
 
-/**
- * @brief Denotes the packed type when the size of the object is 1.
- */
-template <>
-struct packed<sizeof(uint8_t)> {
-  using type = uint8_t;  ///< Packed type as `uint8_t` if the size of the object is 1
-};
-
 /**
  * @brief Denotes the packed type when the size of the object is 2.
  */
diff --git a/tests/static_map/find_test.cu b/tests/static_map/find_test.cu
index 7e1289bb6..6f2d397aa 100644
--- a/tests/static_map/find_test.cu
+++ b/tests/static_map/find_test.cu
@@ -194,3 +194,35 @@ TEMPLATE_TEST_CASE_SIG(
 
   test_unique_sequence(map, num_keys);
 }
+
+TEMPLATE_TEST_CASE_SIG(
+  "static_map: find tests (small types)",
+  "",
+  ((typename Key, typename Value, cuco::test::probe_sequence Probe, int CGSize),
+   Key,
+   Value,
+   Probe,
+   CGSize),
+  (uint8_t, uint8_t, cuco::test::probe_sequence::double_hashing, 1),
+  (uint16_t, uint16_t, cuco::test::probe_sequence::double_hashing, 1)
+)
+{
+  constexpr size_type num_keys{100};
+
+  using extent_type = cuco::extent<size_type, num_keys>;
+  using probe       = cuco::double_hashing<CGSize, cuco::murmurhash3_32<Key>, cuco::murmurhash3_32<Key>>;
+
+  auto map = cuco::static_map<Key,
+                              Value,
+                              extent_type,
+                              cuda::thread_scope_device,
+                              cuda::std::equal_to<Key>,
+                              probe,
+                              cuco::cuda_allocator<cuda::std::byte>,
+                              cuco::storage<2>>{
+    extent_type{},
+    cuco::empty_key<Key>{static_cast<Key>(SENTINEL)},
+    cuco::empty_value<Value>{static_cast<Value>(SENTINEL)}};
+
+  test_unique_sequence(map, num_keys);
+}
diff --git a/tests/static_set/atomic_storage_test.cu b/tests/static_set/atomic_storage_test.cu
index 746608428..b611214bb 100644
--- a/tests/static_set/atomic_storage_test.cu
+++ b/tests/static_set/atomic_storage_test.cu
@@ -69,7 +69,6 @@ TEST_CASE("atomic_storage_test", "")
                                                   cuda::proclaim_return_type<Key>(build_fn{}));
 
   set.insert_async(keys_begin, keys_begin + num_keys);
-  CUCO_CUDA_TRY(cudaDeviceSynchronize());
   auto const count = set.size();
 
   REQUIRE(count == num_keys);

From e4e74cc031ca5cbd754cfc96bf62ae35f22b8b9a Mon Sep 17 00:00:00 2001
From: Krish Sapru <ksapru@bu.edu>
Date: Thu, 7 May 2026 11:39:19 -0400
Subject: [PATCH 06/10] Use signed integers (as in the rest of the tests)

---
 tests/static_map/find_test.cu            |  4 ++--
 tests/static_set/for_each_test.cu        | 18 +++++++++---------
 tests/static_set/insert_and_find_test.cu | 14 +++++++-------
 tests/static_set/retrieve_all_test.cu    | 14 +++++++-------
 tests/static_set/retrieve_if_test.cu     |  6 +++---
 tests/static_set/retrieve_test.cu        | 14 +++++++-------
 tests/static_set/shared_memory_test.cu   |  8 ++++----
 tests/static_set/stream_test.cu          |  6 +++---
 tests/static_set/unique_sequence_test.cu | 16 ++++++++--------
 9 files changed, 50 insertions(+), 50 deletions(-)

diff --git a/tests/static_map/find_test.cu b/tests/static_map/find_test.cu
index 6f2d397aa..12d595b73 100644
--- a/tests/static_map/find_test.cu
+++ b/tests/static_map/find_test.cu
@@ -203,8 +203,8 @@ TEMPLATE_TEST_CASE_SIG(
    Value,
    Probe,
    CGSize),
-  (uint8_t, uint8_t, cuco::test::probe_sequence::double_hashing, 1),
-  (uint16_t, uint16_t, cuco::test::probe_sequence::double_hashing, 1)
+  (int8_t, int8_t, cuco::test::probe_sequence::double_hashing, 1),
+  (int16_t, int16_t, cuco::test::probe_sequence::double_hashing, 1)
 )
 {
   constexpr size_type num_keys{100};
diff --git a/tests/static_set/for_each_test.cu b/tests/static_set/for_each_test.cu
index 8ad84d482..fc92b7e14 100644
--- a/tests/static_set/for_each_test.cu
+++ b/tests/static_set/for_each_test.cu
@@ -77,12 +77,12 @@ TEMPLATE_TEST_CASE_SIG(
   "static_set for_each tests",
   "",
   ((typename Key, cuco::test::probe_sequence Probe, int CGSize), Key, Probe, CGSize),
-  (uint8_t, cuco::test::probe_sequence::double_hashing, 1),
-  (uint8_t, cuco::test::probe_sequence::linear_probing, 1),
-  (uint16_t, cuco::test::probe_sequence::double_hashing, 1),
-  (uint16_t, cuco::test::probe_sequence::double_hashing, 2),
-  (uint16_t, cuco::test::probe_sequence::linear_probing, 1),
-  (uint16_t, cuco::test::probe_sequence::linear_probing, 2),
+  (int8_t, cuco::test::probe_sequence::double_hashing, 1),
+  (int8_t, cuco::test::probe_sequence::linear_probing, 1),
+  (int16_t, cuco::test::probe_sequence::double_hashing, 1),
+  (int16_t, cuco::test::probe_sequence::double_hashing, 2),
+  (int16_t, cuco::test::probe_sequence::linear_probing, 1),
+  (int16_t, cuco::test::probe_sequence::linear_probing, 2),
   (int32_t, cuco::test::probe_sequence::double_hashing, 1),
   (int32_t, cuco::test::probe_sequence::double_hashing, 2),
   (int64_t, cuco::test::probe_sequence::double_hashing, 1),
@@ -100,10 +100,10 @@ TEMPLATE_TEST_CASE_SIG(
 #endif
 )
 {
-  // Limit key count for small types: leave room for the 0xFF/0xFFFF sentinel.
+  // Limit key count for small types: leave room for the -1 sentinel.
   // Expected sums are pre-computed per type class:
-  //   uint8_t  (num_keys=100): sum of evens 0..98 = 2450, sum of odds 1..99 = 2500
-  //   uint16_t+ (num_keys=1000): sum of evens 0..998 = 249'500, sum of odds 1..999 = 250'000
+  //   int8_t  (num_keys=100): sum of evens 0..98 = 2450, sum of odds 1..99 = 2500
+  //   int16_t+ (num_keys=1000): sum of evens 0..998 = 249'500, sum of odds 1..999 = 250'000
   constexpr size_type num_keys         = (sizeof(Key) == 1) ? 100 : 1'000;
   constexpr size_type expected_evens   = (sizeof(Key) == 1) ? 2'450 : 249'500;
   constexpr size_type expected_odds    = (sizeof(Key) == 1) ? 2'500 : 250'000;
diff --git a/tests/static_set/insert_and_find_test.cu b/tests/static_set/insert_and_find_test.cu
index 791244b39..a9d4eeb93 100644
--- a/tests/static_set/insert_and_find_test.cu
+++ b/tests/static_set/insert_and_find_test.cu
@@ -58,12 +58,12 @@ TEMPLATE_TEST_CASE_SIG(
   "static_set Insert and find",
   "",
   ((typename Key, cuco::test::probe_sequence Probe, int CGSize), Key, Probe, CGSize),
-  (uint8_t, cuco::test::probe_sequence::double_hashing, 1),
-  (uint8_t, cuco::test::probe_sequence::linear_probing, 1),
-  (uint16_t, cuco::test::probe_sequence::double_hashing, 1),
-  (uint16_t, cuco::test::probe_sequence::double_hashing, 2),
-  (uint16_t, cuco::test::probe_sequence::linear_probing, 1),
-  (uint16_t, cuco::test::probe_sequence::linear_probing, 2),
+  (int8_t, cuco::test::probe_sequence::double_hashing, 1),
+  (int8_t, cuco::test::probe_sequence::linear_probing, 1),
+  (int16_t, cuco::test::probe_sequence::double_hashing, 1),
+  (int16_t, cuco::test::probe_sequence::double_hashing, 2),
+  (int16_t, cuco::test::probe_sequence::linear_probing, 1),
+  (int16_t, cuco::test::probe_sequence::linear_probing, 2),
   (int32_t, cuco::test::probe_sequence::double_hashing, 1),
   (int32_t, cuco::test::probe_sequence::double_hashing, 2),
   (int64_t, cuco::test::probe_sequence::double_hashing, 1),
@@ -81,7 +81,7 @@ TEMPLATE_TEST_CASE_SIG(
 #endif
 )
 {
-  // Limit key count for small types: leave room for the 0xFF/0xFFFF sentinel
+  // Limit key count for small types: leave room for the -1 sentinel
   constexpr std::size_t num_keys = (sizeof(Key) == 1) ? 100 : 400;
 
   using probe = std::conditional_t<Probe == cuco::test::probe_sequence::linear_probing,
diff --git a/tests/static_set/retrieve_all_test.cu b/tests/static_set/retrieve_all_test.cu
index 4e2e53529..b340a8dbf 100644
--- a/tests/static_set/retrieve_all_test.cu
+++ b/tests/static_set/retrieve_all_test.cu
@@ -63,12 +63,12 @@ TEMPLATE_TEST_CASE_SIG(
   "static_set::retrieve_all tests",
   "",
   ((typename Key, cuco::test::probe_sequence Probe, int CGSize), Key, Probe, CGSize),
-  (uint8_t, cuco::test::probe_sequence::double_hashing, 1),
-  (uint8_t, cuco::test::probe_sequence::linear_probing, 1),
-  (uint16_t, cuco::test::probe_sequence::double_hashing, 1),
-  (uint16_t, cuco::test::probe_sequence::double_hashing, 2),
-  (uint16_t, cuco::test::probe_sequence::linear_probing, 1),
-  (uint16_t, cuco::test::probe_sequence::linear_probing, 2),
+  (int8_t, cuco::test::probe_sequence::double_hashing, 1),
+  (int8_t, cuco::test::probe_sequence::linear_probing, 1),
+  (int16_t, cuco::test::probe_sequence::double_hashing, 1),
+  (int16_t, cuco::test::probe_sequence::double_hashing, 2),
+  (int16_t, cuco::test::probe_sequence::linear_probing, 1),
+  (int16_t, cuco::test::probe_sequence::linear_probing, 2),
   (int32_t, cuco::test::probe_sequence::double_hashing, 1),
   (int32_t, cuco::test::probe_sequence::double_hashing, 2),
   (int64_t, cuco::test::probe_sequence::double_hashing, 1),
@@ -86,7 +86,7 @@ TEMPLATE_TEST_CASE_SIG(
 #endif
 )
 {
-  // Limit key count for small types: leave room for the 0xFF/0xFFFF sentinel
+  // Limit key count for small types: leave room for the -1 sentinel
   constexpr std::size_t num_keys        = (sizeof(Key) == 1) ? 100 : 400;
   constexpr double desired_load_factor = 1.;
 
diff --git a/tests/static_set/retrieve_if_test.cu b/tests/static_set/retrieve_if_test.cu
index d97222e05..f242e4ad6 100644
--- a/tests/static_set/retrieve_if_test.cu
+++ b/tests/static_set/retrieve_if_test.cu
@@ -112,8 +112,8 @@ __global__ void test_retrieve_if_all_true_kernel(
 TEMPLATE_TEST_CASE_SIG("static_set retrieve_if",
                        "",
                        ((typename Key), Key),
-                       (uint8_t),
-                       (uint16_t),
+                       (int8_t),
+                       (int16_t),
                        (int32_t),
                        (int64_t)
 #if defined(CUCO_HAS_128BIT_ATOMICS)
@@ -122,7 +122,7 @@ TEMPLATE_TEST_CASE_SIG("static_set retrieve_if",
 #endif
 )
 {
-  // Limit key count for small types: keys start at 1, sentinel is 0xFF/0xFFFF
+  // Limit key count for small types: keys start at 1, sentinel is -1
   constexpr size_type num_keys = (sizeof(Key) == 1) ? 100 : 400;
 
   using container_type = cuco::static_set<Key>;
diff --git a/tests/static_set/retrieve_test.cu b/tests/static_set/retrieve_test.cu
index 68cf44f25..cf3adb920 100644
--- a/tests/static_set/retrieve_test.cu
+++ b/tests/static_set/retrieve_test.cu
@@ -77,12 +77,12 @@ TEMPLATE_TEST_CASE_SIG(
   "static_set retrieve tests",
   "",
   ((typename Key, cuco::test::probe_sequence Probe, int CGSize), Key, Probe, CGSize),
-  (uint8_t, cuco::test::probe_sequence::double_hashing, 1),
-  (uint8_t, cuco::test::probe_sequence::linear_probing, 1),
-  (uint16_t, cuco::test::probe_sequence::double_hashing, 1),
-  (uint16_t, cuco::test::probe_sequence::double_hashing, 2),
-  (uint16_t, cuco::test::probe_sequence::linear_probing, 1),
-  (uint16_t, cuco::test::probe_sequence::linear_probing, 2),
+  (int8_t, cuco::test::probe_sequence::double_hashing, 1),
+  (int8_t, cuco::test::probe_sequence::linear_probing, 1),
+  (int16_t, cuco::test::probe_sequence::double_hashing, 1),
+  (int16_t, cuco::test::probe_sequence::double_hashing, 2),
+  (int16_t, cuco::test::probe_sequence::linear_probing, 1),
+  (int16_t, cuco::test::probe_sequence::linear_probing, 2),
   (int32_t, cuco::test::probe_sequence::double_hashing, 1),
   (int32_t, cuco::test::probe_sequence::double_hashing, 2),
   (int64_t, cuco::test::probe_sequence::double_hashing, 1),
@@ -100,7 +100,7 @@ TEMPLATE_TEST_CASE_SIG(
 #endif
 )
 {
-  // Limit key count for small types: leave room for the 0xFF/0xFFFF sentinel
+  // Limit key count for small types: leave room for the -1 sentinel
   constexpr std::size_t num_keys        = (sizeof(Key) == 1) ? 100 : 400;
   constexpr double desired_load_factor = 1.;
 
diff --git a/tests/static_set/shared_memory_test.cu b/tests/static_set/shared_memory_test.cu
index 345c769c6..05fe246fe 100644
--- a/tests/static_set/shared_memory_test.cu
+++ b/tests/static_set/shared_memory_test.cu
@@ -71,8 +71,8 @@ __global__ void shared_memory_test_kernel(Ref* sets,
 TEMPLATE_TEST_CASE_SIG("static_set shared memory tests",
                        "",
                        ((typename Key), Key),
-                       (uint8_t),
-                       (uint16_t),
+                       (int8_t),
+                       (int16_t),
                        (int32_t),
                        (int64_t)
 #if defined(CUCO_HAS_128BIT_ATOMICS)
@@ -81,8 +81,8 @@ TEMPLATE_TEST_CASE_SIG("static_set shared memory tests",
 #endif
 )
 {
-  // For uint8_t: sentinel = 0xFF (255), so usable key range is 0..254.
-  // For uint16_t: sentinel = 0xFFFF (65535), so usable key range is 0..65534.
+  // For int8_t: sentinel = -1, so usable key range is -128..127 (excluding -1).
+  // For int16_t: sentinel = -1, so usable key range is -32768..32767 (excluding -1).
   // thrust::sequence over number_of_sets*elements_in_set keys must not wrap.
   // Use smaller set count and element count for smaller types.
   constexpr std::size_t number_of_sets  = (sizeof(Key) <= 2) ? (sizeof(Key) == 1 ? 2 : 100) : 1000;
diff --git a/tests/static_set/stream_test.cu b/tests/static_set/stream_test.cu
index 2413bd578..8811900a1 100644
--- a/tests/static_set/stream_test.cu
+++ b/tests/static_set/stream_test.cu
@@ -35,8 +35,8 @@
 TEMPLATE_TEST_CASE_SIG("static_set: operations on different stream than constructor",
                        "",
                        ((typename Key), Key),
-                       (uint8_t),
-                       (uint16_t),
+                       (int8_t),
+                       (int16_t),
                        (int32_t),
                        (int64_t)
 #if defined(CUCO_HAS_128BIT_ATOMICS)
@@ -51,7 +51,7 @@ TEMPLATE_TEST_CASE_SIG("static_set: operations on different stream than construc
   CUCO_CUDA_TRY(cudaStreamCreate(&operation_stream));
 
   {  // Scope ensures set is destroyed before streams
-    // Scale num_keys to fit in the key type's value range (sentinel = 0xFF/0xFFFF).
+    // Scale num_keys to fit in the key type's value range (sentinel = -1).
     constexpr std::size_t num_keys =
       (sizeof(Key) == 1) ? 100 : (sizeof(Key) == 2) ? 1'000 : 500'000;
     auto set = cuco::static_set{num_keys * 2,
diff --git a/tests/static_set/unique_sequence_test.cu b/tests/static_set/unique_sequence_test.cu
index a1b6663cf..a8f227ac7 100644
--- a/tests/static_set/unique_sequence_test.cu
+++ b/tests/static_set/unique_sequence_test.cu
@@ -143,12 +143,12 @@ TEMPLATE_TEST_CASE_SIG(
   "static_set unique sequence tests",
   "",
   ((typename Key, cuco::test::probe_sequence Probe, int CGSize), Key, Probe, CGSize),
-  (uint8_t, cuco::test::probe_sequence::double_hashing, 1),
-  (uint8_t, cuco::test::probe_sequence::linear_probing, 1),
-  (uint16_t, cuco::test::probe_sequence::double_hashing, 1),
-  (uint16_t, cuco::test::probe_sequence::double_hashing, 2),
-  (uint16_t, cuco::test::probe_sequence::linear_probing, 1),
-  (uint16_t, cuco::test::probe_sequence::linear_probing, 2),
+  (int8_t, cuco::test::probe_sequence::double_hashing, 1),
+  (int8_t, cuco::test::probe_sequence::linear_probing, 1),
+  (int16_t, cuco::test::probe_sequence::double_hashing, 1),
+  (int16_t, cuco::test::probe_sequence::double_hashing, 2),
+  (int16_t, cuco::test::probe_sequence::linear_probing, 1),
+  (int16_t, cuco::test::probe_sequence::linear_probing, 2),
   (int32_t, cuco::test::probe_sequence::double_hashing, 1),
   (int32_t, cuco::test::probe_sequence::double_hashing, 2),
   (int64_t, cuco::test::probe_sequence::double_hashing, 1),
@@ -166,7 +166,7 @@ TEMPLATE_TEST_CASE_SIG(
 #endif
 )
 {
-  // Limit key count for small types: leave room for the 0xFF/0xFFFF sentinel
+  // Limit key count for small types: leave room for the -1 sentinel
   constexpr size_type num_keys = (sizeof(Key) == 1) ? 100 : 400;
   using probe = std::conditional_t<Probe == cuco::test::probe_sequence::linear_probing,
                                    cuco::linear_probing<CGSize, cuco::default_hash_function<Key>>,
@@ -176,7 +176,7 @@ TEMPLATE_TEST_CASE_SIG(
     if constexpr (cuco::is_double_hashing<probe>::value) {
       if constexpr (num_keys == 100) {
         return (CGSize == 1) ? 106   // 53 x 1 x 2
-                             : 106;  // 53 x 2 x 1... only CGSize=1 used for uint8_t
+                             : 106;  // 53 x 2 x 1... only CGSize=1 used for int8_t
       } else {
         return (CGSize == 1) ? 422   // 211 x 1 x 2
                              : 404;  // 101 x 2 x 2

From af7bc7c45a194564f0ea3dd160dba894c75606df Mon Sep 17 00:00:00 2001
From: "pre-commit-ci[bot]"
 <66853113+pre-commit-ci[bot]@users.noreply.github.com>
Date: Thu, 7 May 2026 15:39:47 +0000
Subject: [PATCH 07/10] [pre-commit.ci] auto code formatting

---
 tests/static_map/find_test.cu            | 29 ++++++++++++------------
 tests/static_set/for_each_test.cu        | 10 ++++----
 tests/static_set/insert_and_find_test.cu |  4 ++--
 tests/static_set/retrieve_all_test.cu    |  6 ++---
 tests/static_set/retrieve_if_test.cu     |  4 ++--
 tests/static_set/retrieve_test.cu        |  6 ++---
 tests/static_set/shared_memory_test.cu   |  3 +--
 tests/static_set/stream_test.cu          | 19 ++++++++--------
 tests/static_set/unique_sequence_test.cu |  4 ++--
 9 files changed, 42 insertions(+), 43 deletions(-)

diff --git a/tests/static_map/find_test.cu b/tests/static_map/find_test.cu
index 12d595b73..91ba2f1a5 100644
--- a/tests/static_map/find_test.cu
+++ b/tests/static_map/find_test.cu
@@ -204,25 +204,24 @@ TEMPLATE_TEST_CASE_SIG(
    Probe,
    CGSize),
   (int8_t, int8_t, cuco::test::probe_sequence::double_hashing, 1),
-  (int16_t, int16_t, cuco::test::probe_sequence::double_hashing, 1)
-)
+  (int16_t, int16_t, cuco::test::probe_sequence::double_hashing, 1))
 {
   constexpr size_type num_keys{100};
 
   using extent_type = cuco::extent<size_type, num_keys>;
-  using probe       = cuco::double_hashing<CGSize, cuco::murmurhash3_32<Key>, cuco::murmurhash3_32<Key>>;
-
-  auto map = cuco::static_map<Key,
-                              Value,
-                              extent_type,
-                              cuda::thread_scope_device,
-                              cuda::std::equal_to<Key>,
-                              probe,
-                              cuco::cuda_allocator<cuda::std::byte>,
-                              cuco::storage<2>>{
-    extent_type{},
-    cuco::empty_key<Key>{static_cast<Key>(SENTINEL)},
-    cuco::empty_value<Value>{static_cast<Value>(SENTINEL)}};
+  using probe = cuco::double_hashing<CGSize, cuco::murmurhash3_32<Key>, cuco::murmurhash3_32<Key>>;
+
+  auto map =
+    cuco::static_map<Key,
+                     Value,
+                     extent_type,
+                     cuda::thread_scope_device,
+                     cuda::std::equal_to<Key>,
+                     probe,
+                     cuco::cuda_allocator<cuda::std::byte>,
+                     cuco::storage<2>>{extent_type{},
+                                       cuco::empty_key<Key>{static_cast<Key>(SENTINEL)},
+                                       cuco::empty_value<Value>{static_cast<Value>(SENTINEL)}};
 
   test_unique_sequence(map, num_keys);
 }
diff --git a/tests/static_set/for_each_test.cu b/tests/static_set/for_each_test.cu
index fc92b7e14..c9d4530bc 100644
--- a/tests/static_set/for_each_test.cu
+++ b/tests/static_set/for_each_test.cu
@@ -16,8 +16,6 @@
 
 #include <test_utils.hpp>
 
-#include <cstdint>
-
 #include <cuco/detail/__config>
 #include <cuco/static_set.cuh>
 
@@ -28,6 +26,8 @@
 
 #include <catch2/catch_template_test_macros.hpp>
 
+#include <cstdint>
+
 using size_type = std::size_t;
 
 template <typename Set>
@@ -104,9 +104,9 @@ TEMPLATE_TEST_CASE_SIG(
   // Expected sums are pre-computed per type class:
   //   int8_t  (num_keys=100): sum of evens 0..98 = 2450, sum of odds 1..99 = 2500
   //   int16_t+ (num_keys=1000): sum of evens 0..998 = 249'500, sum of odds 1..999 = 250'000
-  constexpr size_type num_keys         = (sizeof(Key) == 1) ? 100 : 1'000;
-  constexpr size_type expected_evens   = (sizeof(Key) == 1) ? 2'450 : 249'500;
-  constexpr size_type expected_odds    = (sizeof(Key) == 1) ? 2'500 : 250'000;
+  constexpr size_type num_keys       = (sizeof(Key) == 1) ? 100 : 1'000;
+  constexpr size_type expected_evens = (sizeof(Key) == 1) ? 2'450 : 249'500;
+  constexpr size_type expected_odds  = (sizeof(Key) == 1) ? 2'500 : 250'000;
 
   using probe = std::conditional_t<
     Probe == cuco::test::probe_sequence::linear_probing,
diff --git a/tests/static_set/insert_and_find_test.cu b/tests/static_set/insert_and_find_test.cu
index a9d4eeb93..3d7d2fb25 100644
--- a/tests/static_set/insert_and_find_test.cu
+++ b/tests/static_set/insert_and_find_test.cu
@@ -16,8 +16,6 @@
 
 #include <test_utils.hpp>
 
-#include <cstdint>
-
 #include <cuco/detail/__config>
 #include <cuco/static_set.cuh>
 
@@ -27,6 +25,8 @@
 
 #include <catch2/catch_template_test_macros.hpp>
 
+#include <cstdint>
+
 template <typename Set>
 void test_insert_and_find(Set& set, std::size_t num_keys)
 {
diff --git a/tests/static_set/retrieve_all_test.cu b/tests/static_set/retrieve_all_test.cu
index b340a8dbf..01003e35a 100644
--- a/tests/static_set/retrieve_all_test.cu
+++ b/tests/static_set/retrieve_all_test.cu
@@ -16,8 +16,6 @@
 
 #include <test_utils.hpp>
 
-#include <cstdint>
-
 #include <cuco/detail/__config>
 #include <cuco/static_set.cuh>
 
@@ -29,6 +27,8 @@
 
 #include <catch2/catch_template_test_macros.hpp>
 
+#include <cstdint>
+
 template <typename Set>
 void test_unique_sequence(Set& set, std::size_t num_keys)
 {
@@ -87,7 +87,7 @@ TEMPLATE_TEST_CASE_SIG(
 )
 {
   // Limit key count for small types: leave room for the -1 sentinel
-  constexpr std::size_t num_keys        = (sizeof(Key) == 1) ? 100 : 400;
+  constexpr std::size_t num_keys       = (sizeof(Key) == 1) ? 100 : 400;
   constexpr double desired_load_factor = 1.;
 
   using probe = std::conditional_t<Probe == cuco::test::probe_sequence::linear_probing,
diff --git a/tests/static_set/retrieve_if_test.cu b/tests/static_set/retrieve_if_test.cu
index f242e4ad6..41dcfca1e 100644
--- a/tests/static_set/retrieve_if_test.cu
+++ b/tests/static_set/retrieve_if_test.cu
@@ -16,8 +16,6 @@
 
 #include <test_utils.hpp>
 
-#include <cstdint>
-
 #include <cuco/detail/__config>
 #include <cuco/static_set.cuh>
 
@@ -29,6 +27,8 @@
 
 #include <catch2/catch_template_test_macros.hpp>
 
+#include <cstdint>
+
 using size_type = std::size_t;
 
 template <class Container>
diff --git a/tests/static_set/retrieve_test.cu b/tests/static_set/retrieve_test.cu
index cf3adb920..4b5021eb1 100644
--- a/tests/static_set/retrieve_test.cu
+++ b/tests/static_set/retrieve_test.cu
@@ -16,8 +16,6 @@
 
 #include <test_utils.hpp>
 
-#include <cstdint>
-
 #include <cuco/detail/__config>
 #include <cuco/static_set.cuh>
 
@@ -30,6 +28,8 @@
 
 #include <catch2/catch_template_test_macros.hpp>
 
+#include <cstdint>
+
 static constexpr int key_sentinel = -1;
 
 template <typename Set>
@@ -101,7 +101,7 @@ TEMPLATE_TEST_CASE_SIG(
 )
 {
   // Limit key count for small types: leave room for the -1 sentinel
-  constexpr std::size_t num_keys        = (sizeof(Key) == 1) ? 100 : 400;
+  constexpr std::size_t num_keys       = (sizeof(Key) == 1) ? 100 : 400;
   constexpr double desired_load_factor = 1.;
 
   using probe = std::conditional_t<Probe == cuco::test::probe_sequence::linear_probing,
diff --git a/tests/static_set/shared_memory_test.cu b/tests/static_set/shared_memory_test.cu
index 05fe246fe..95b4908d5 100644
--- a/tests/static_set/shared_memory_test.cu
+++ b/tests/static_set/shared_memory_test.cu
@@ -16,8 +16,6 @@
 
 #include <test_utils.hpp>
 
-#include <cstdint>
-
 #include <cuco/detail/__config>
 #include <cuco/static_set.cuh>
 
@@ -30,6 +28,7 @@
 
 #include <catch2/catch_template_test_macros.hpp>
 
+#include <cstdint>
 #include <limits>
 
 template <std::size_t ValidSize, typename Ref>
diff --git a/tests/static_set/stream_test.cu b/tests/static_set/stream_test.cu
index 8811900a1..a8d73bf33 100644
--- a/tests/static_set/stream_test.cu
+++ b/tests/static_set/stream_test.cu
@@ -16,8 +16,6 @@
 
 #include <test_utils.hpp>
 
-#include <cstdint>
-
 #include <cuco/detail/__config>
 #include <cuco/static_set.cuh>
 
@@ -32,6 +30,8 @@
 
 #include <catch2/catch_template_test_macros.hpp>
 
+#include <cstdint>
+
 TEMPLATE_TEST_CASE_SIG("static_set: operations on different stream than constructor",
                        "",
                        ((typename Key), Key),
@@ -52,15 +52,16 @@ TEMPLATE_TEST_CASE_SIG("static_set: operations on different stream than construc
 
   {  // Scope ensures set is destroyed before streams
     // Scale num_keys to fit in the key type's value range (sentinel = -1).
-    constexpr std::size_t num_keys =
-      (sizeof(Key) == 1) ? 100 : (sizeof(Key) == 2) ? 1'000 : 500'000;
-    auto set = cuco::static_set{num_keys * 2,
+    constexpr std::size_t num_keys = (sizeof(Key) == 1)   ? 100
+                                     : (sizeof(Key) == 2) ? 1'000
+                                                          : 500'000;
+    auto set                       = cuco::static_set{num_keys * 2,
                                 cuco::empty_key<Key>{static_cast<Key>(-1)},
-                                {},
+                                                      {},
                                 cuco::linear_probing<1, cuco::default_hash_function<Key>>{},
-                                {},
-                                {},
-                                {},
+                                                      {},
+                                                      {},
+                                                      {},
                                 constructor_stream};
 
     thrust::device_vector<Key> d_keys(num_keys);
diff --git a/tests/static_set/unique_sequence_test.cu b/tests/static_set/unique_sequence_test.cu
index a8f227ac7..262f4982f 100644
--- a/tests/static_set/unique_sequence_test.cu
+++ b/tests/static_set/unique_sequence_test.cu
@@ -16,8 +16,6 @@
 
 #include <test_utils.hpp>
 
-#include <cstdint>
-
 #include <cuco/detail/__config>
 #include <cuco/static_set.cuh>
 
@@ -31,6 +29,8 @@
 
 #include <catch2/catch_template_test_macros.hpp>
 
+#include <cstdint>
+
 using size_type = int32_t;
 
 int32_t constexpr SENTINEL = -1;

From fd402c5c497bf8ab413483f99a6ec00f3d12c33f Mon Sep 17 00:00:00 2001
From: Krish Sapru <ksapru@bu.edu>
Date: Sat, 9 May 2026 12:37:07 -0400
Subject: [PATCH 08/10] Final tests and CI failures

---
 include/cuco/static_map.cuh              | 2 ++
 include/cuco/static_set.cuh              | 2 ++
 tests/static_map/contains_test.cu        | 1 +
 tests/static_map/find_test.cu            | 1 -
 tests/static_map/insert_and_find_test.cu | 1 +
 tests/static_map/retrieve_if_test.cu     | 1 +
 tests/static_map/shared_memory_test.cu   | 8 +++++---
 tests/static_map/stream_test.cu          | 7 ++++---
 tests/static_set/for_each_test.cu        | 5 ++---
 tests/static_set/insert_and_find_test.cu | 3 +--
 tests/static_set/retrieve_all_test.cu    | 3 +--
 tests/static_set/retrieve_if_test.cu     | 1 -
 tests/static_set/retrieve_test.cu        | 3 +--
 tests/static_set/shared_memory_test.cu   | 3 +--
 tests/static_set/stream_test.cu          | 1 -
 tests/static_set/unique_sequence_test.cu | 5 ++---
 16 files changed, 24 insertions(+), 23 deletions(-)

diff --git a/include/cuco/static_map.cuh b/include/cuco/static_map.cuh
index c9e687e31..3beb51c42 100644
--- a/include/cuco/static_map.cuh
+++ b/include/cuco/static_map.cuh
@@ -106,6 +106,8 @@ class static_map {
                                                  Allocator,
                                                  Storage>;
 
+  static_assert(sizeof(Key) > 1, "cuCollections does not support single-byte keys.");
+
  public:
   static constexpr auto cg_size      = impl_type::cg_size;       ///< CG size used for probing
   static constexpr auto bucket_size  = impl_type::bucket_size;   ///< Bucket size used for probing
diff --git a/include/cuco/static_set.cuh b/include/cuco/static_set.cuh
index ad9ce7f6d..d6fb2fda6 100644
--- a/include/cuco/static_set.cuh
+++ b/include/cuco/static_set.cuh
@@ -92,6 +92,8 @@ class static_set {
   using impl_type = detail::
     open_addressing_impl<Key, Key, Extent, Scope, KeyEqual, ProbingScheme, Allocator, Storage>;
 
+  static_assert(sizeof(Key) > 1, "cuCollections does not support single-byte keys.");
+
  public:
   static constexpr auto cg_size      = impl_type::cg_size;       ///< CG size used for probing
   static constexpr auto bucket_size  = impl_type::bucket_size;   ///< Bucket size used for probing
diff --git a/tests/static_map/contains_test.cu b/tests/static_map/contains_test.cu
index 0b3604528..13ac31728 100644
--- a/tests/static_map/contains_test.cu
+++ b/tests/static_map/contains_test.cu
@@ -124,6 +124,7 @@ TEMPLATE_TEST_CASE_SIG(
    Value,
    Probe,
    CGSize),
+  (int16_t, int16_t, cuco::test::probe_sequence::double_hashing, 1),
   (int32_t, int32_t, cuco::test::probe_sequence::double_hashing, 1),
   (int32_t, int32_t, cuco::test::probe_sequence::double_hashing, 2),
   (int64_t, int64_t, cuco::test::probe_sequence::double_hashing, 1),
diff --git a/tests/static_map/find_test.cu b/tests/static_map/find_test.cu
index 91ba2f1a5..b1e3fab23 100644
--- a/tests/static_map/find_test.cu
+++ b/tests/static_map/find_test.cu
@@ -203,7 +203,6 @@ TEMPLATE_TEST_CASE_SIG(
    Value,
    Probe,
    CGSize),
-  (int8_t, int8_t, cuco::test::probe_sequence::double_hashing, 1),
   (int16_t, int16_t, cuco::test::probe_sequence::double_hashing, 1))
 {
   constexpr size_type num_keys{100};
diff --git a/tests/static_map/insert_and_find_test.cu b/tests/static_map/insert_and_find_test.cu
index 17665eb2b..df4de0595 100644
--- a/tests/static_map/insert_and_find_test.cu
+++ b/tests/static_map/insert_and_find_test.cu
@@ -37,6 +37,7 @@ TEMPLATE_TEST_CASE_SIG(
    Value,
    Probe,
    CGSize),
+  (int16_t, int16_t, cuco::test::probe_sequence::double_hashing, 1),
   (int32_t, int32_t, cuco::test::probe_sequence::double_hashing, 1),
   (int32_t, int64_t, cuco::test::probe_sequence::double_hashing, 1),
   (int32_t, int32_t, cuco::test::probe_sequence::double_hashing, 2),
diff --git a/tests/static_map/retrieve_if_test.cu b/tests/static_map/retrieve_if_test.cu
index fe4441822..44675a09a 100644
--- a/tests/static_map/retrieve_if_test.cu
+++ b/tests/static_map/retrieve_if_test.cu
@@ -110,6 +110,7 @@ __global__ void test_retrieve_if_all_true_kernel(
 TEMPLATE_TEST_CASE_SIG("static_map retrieve_if",
                        "",
                        ((typename Key, typename T), Key, T),
+                       (int16_t, int16_t),
                        (int32_t, int32_t),
                        (int64_t, int64_t)
 #if defined(CUCO_HAS_128BIT_ATOMICS)
diff --git a/tests/static_map/shared_memory_test.cu b/tests/static_map/shared_memory_test.cu
index 15731319c..3624a6a5f 100644
--- a/tests/static_map/shared_memory_test.cu
+++ b/tests/static_map/shared_memory_test.cu
@@ -70,6 +70,7 @@ __global__ void shared_memory_test_kernel(Ref* maps,
 TEMPLATE_TEST_CASE_SIG("static_map shared memory tests",
                        "",
                        ((typename Key, typename Value), Key, Value),
+                       (int16_t, int16_t),
                        (int32_t, int32_t),
                        (int32_t, int64_t),
                        (int64_t, int32_t),
@@ -80,8 +81,9 @@ TEMPLATE_TEST_CASE_SIG("static_map shared memory tests",
 #endif
 )
 {
-  constexpr std::size_t number_of_maps  = 1000;
-  constexpr std::size_t elements_in_map = 500;
+  // For small types, use smaller set count and element count
+  constexpr std::size_t number_of_maps = (sizeof(Key) + sizeof(Value) <= 2) ? 100 : 1000;
+  constexpr std::size_t elements_in_map = (sizeof(Key) + sizeof(Value) <= 2) ? 100 : 500;
   constexpr std::size_t map_capacity    = 2 * elements_in_map;
 
   using extent_type = cuco::extent<std::size_t, map_capacity>;
@@ -105,7 +107,7 @@ TEMPLATE_TEST_CASE_SIG("static_map shared memory tests",
   std::vector<std::unique_ptr<map_type>> maps;
   for (std::size_t map_id = 0; map_id < number_of_maps; ++map_id) {
     maps.push_back(std::make_unique<map_type>(
-      extent_type{}, cuco::empty_key<Key>{-1}, cuco::empty_value<Value>{-1}));
+      extent_type{}, cuco::empty_key<Key>{static_cast<Key>(-1)}, cuco::empty_value<Value>{static_cast<Value>(-1)}));
   }
 
   thrust::device_vector<bool> d_keys_exist(number_of_maps * elements_in_map);
diff --git a/tests/static_map/stream_test.cu b/tests/static_map/stream_test.cu
index 6c1701d98..445a9550d 100644
--- a/tests/static_map/stream_test.cu
+++ b/tests/static_map/stream_test.cu
@@ -32,6 +32,7 @@
 TEMPLATE_TEST_CASE_SIG("static_map: unique sequence of keys on given stream",
                        "",
                        ((typename Key, typename Value), Key, Value),
+                       (int16_t, int16_t),
                        (int32_t, int32_t),
                        (int32_t, int64_t),
                        (int64_t, int32_t),
@@ -46,10 +47,10 @@ TEMPLATE_TEST_CASE_SIG("static_map: unique sequence of keys on given stream",
   CUCO_CUDA_TRY(cudaStreamCreate(&stream));
 
   {  // Scope ensures map is destroyed before stream
-    constexpr std::size_t num_keys{500'000};
+    const std::size_t num_keys = (sizeof(Key) == 1) ? 100 : 500'000;
     auto map = cuco::static_map{num_keys * 2,
-                                cuco::empty_key<Key>{-1},
-                                cuco::empty_value<Value>{-1},
+                                cuco::empty_key<Key>{static_cast<Key>(-1)},
+                                cuco::empty_value<Value>{static_cast<Value>(-1)},
                                 {},
                                 cuco::linear_probing<1, cuco::default_hash_function<Key>>{},
                                 {},
diff --git a/tests/static_set/for_each_test.cu b/tests/static_set/for_each_test.cu
index c9d4530bc..e4c3760be 100644
--- a/tests/static_set/for_each_test.cu
+++ b/tests/static_set/for_each_test.cu
@@ -77,8 +77,7 @@ TEMPLATE_TEST_CASE_SIG(
   "static_set for_each tests",
   "",
   ((typename Key, cuco::test::probe_sequence Probe, int CGSize), Key, Probe, CGSize),
-  (int8_t, cuco::test::probe_sequence::double_hashing, 1),
-  (int8_t, cuco::test::probe_sequence::linear_probing, 1),
+  (int16_t, cuco::test::probe_sequence::double_hashing, 1),
   (int16_t, cuco::test::probe_sequence::double_hashing, 1),
   (int16_t, cuco::test::probe_sequence::double_hashing, 2),
   (int16_t, cuco::test::probe_sequence::linear_probing, 1),
@@ -102,7 +101,7 @@ TEMPLATE_TEST_CASE_SIG(
 {
   // Limit key count for small types: leave room for the -1 sentinel.
   // Expected sums are pre-computed per type class:
-  //   int8_t  (num_keys=100): sum of evens 0..98 = 2450, sum of odds 1..99 = 2500
+  //   int16_t (num_keys=100): sum of evens 0..98 = 2450, sum of odds 1..99 = 2500
   //   int16_t+ (num_keys=1000): sum of evens 0..998 = 249'500, sum of odds 1..999 = 250'000
   constexpr size_type num_keys       = (sizeof(Key) == 1) ? 100 : 1'000;
   constexpr size_type expected_evens = (sizeof(Key) == 1) ? 2'450 : 249'500;
diff --git a/tests/static_set/insert_and_find_test.cu b/tests/static_set/insert_and_find_test.cu
index 3d7d2fb25..5e4345c5e 100644
--- a/tests/static_set/insert_and_find_test.cu
+++ b/tests/static_set/insert_and_find_test.cu
@@ -58,8 +58,7 @@ TEMPLATE_TEST_CASE_SIG(
   "static_set Insert and find",
   "",
   ((typename Key, cuco::test::probe_sequence Probe, int CGSize), Key, Probe, CGSize),
-  (int8_t, cuco::test::probe_sequence::double_hashing, 1),
-  (int8_t, cuco::test::probe_sequence::linear_probing, 1),
+  (int16_t, cuco::test::probe_sequence::double_hashing, 1),
   (int16_t, cuco::test::probe_sequence::double_hashing, 1),
   (int16_t, cuco::test::probe_sequence::double_hashing, 2),
   (int16_t, cuco::test::probe_sequence::linear_probing, 1),
diff --git a/tests/static_set/retrieve_all_test.cu b/tests/static_set/retrieve_all_test.cu
index 01003e35a..c9a5dfe64 100644
--- a/tests/static_set/retrieve_all_test.cu
+++ b/tests/static_set/retrieve_all_test.cu
@@ -63,8 +63,7 @@ TEMPLATE_TEST_CASE_SIG(
   "static_set::retrieve_all tests",
   "",
   ((typename Key, cuco::test::probe_sequence Probe, int CGSize), Key, Probe, CGSize),
-  (int8_t, cuco::test::probe_sequence::double_hashing, 1),
-  (int8_t, cuco::test::probe_sequence::linear_probing, 1),
+  (int16_t, cuco::test::probe_sequence::double_hashing, 1),
   (int16_t, cuco::test::probe_sequence::double_hashing, 1),
   (int16_t, cuco::test::probe_sequence::double_hashing, 2),
   (int16_t, cuco::test::probe_sequence::linear_probing, 1),
diff --git a/tests/static_set/retrieve_if_test.cu b/tests/static_set/retrieve_if_test.cu
index 41dcfca1e..8faf181bf 100644
--- a/tests/static_set/retrieve_if_test.cu
+++ b/tests/static_set/retrieve_if_test.cu
@@ -112,7 +112,6 @@ __global__ void test_retrieve_if_all_true_kernel(
 TEMPLATE_TEST_CASE_SIG("static_set retrieve_if",
                        "",
                        ((typename Key), Key),
-                       (int8_t),
                        (int16_t),
                        (int32_t),
                        (int64_t)
diff --git a/tests/static_set/retrieve_test.cu b/tests/static_set/retrieve_test.cu
index 4b5021eb1..faa154d0a 100644
--- a/tests/static_set/retrieve_test.cu
+++ b/tests/static_set/retrieve_test.cu
@@ -77,8 +77,7 @@ TEMPLATE_TEST_CASE_SIG(
   "static_set retrieve tests",
   "",
   ((typename Key, cuco::test::probe_sequence Probe, int CGSize), Key, Probe, CGSize),
-  (int8_t, cuco::test::probe_sequence::double_hashing, 1),
-  (int8_t, cuco::test::probe_sequence::linear_probing, 1),
+  (int16_t, cuco::test::probe_sequence::double_hashing, 1),
   (int16_t, cuco::test::probe_sequence::double_hashing, 1),
   (int16_t, cuco::test::probe_sequence::double_hashing, 2),
   (int16_t, cuco::test::probe_sequence::linear_probing, 1),
diff --git a/tests/static_set/shared_memory_test.cu b/tests/static_set/shared_memory_test.cu
index 95b4908d5..655413d74 100644
--- a/tests/static_set/shared_memory_test.cu
+++ b/tests/static_set/shared_memory_test.cu
@@ -70,7 +70,6 @@ __global__ void shared_memory_test_kernel(Ref* sets,
 TEMPLATE_TEST_CASE_SIG("static_set shared memory tests",
                        "",
                        ((typename Key), Key),
-                       (int8_t),
                        (int16_t),
                        (int32_t),
                        (int64_t)
@@ -80,7 +79,7 @@ TEMPLATE_TEST_CASE_SIG("static_set shared memory tests",
 #endif
 )
 {
-  // For int8_t: sentinel = -1, so usable key range is -128..127 (excluding -1).
+  // For int16_t: sentinel = -1, so usable key range is -32768..32767 (excluding -1).
   // For int16_t: sentinel = -1, so usable key range is -32768..32767 (excluding -1).
   // thrust::sequence over number_of_sets*elements_in_set keys must not wrap.
   // Use smaller set count and element count for smaller types.
diff --git a/tests/static_set/stream_test.cu b/tests/static_set/stream_test.cu
index a8d73bf33..b5affa190 100644
--- a/tests/static_set/stream_test.cu
+++ b/tests/static_set/stream_test.cu
@@ -35,7 +35,6 @@
 TEMPLATE_TEST_CASE_SIG("static_set: operations on different stream than constructor",
                        "",
                        ((typename Key), Key),
-                       (int8_t),
                        (int16_t),
                        (int32_t),
                        (int64_t)
diff --git a/tests/static_set/unique_sequence_test.cu b/tests/static_set/unique_sequence_test.cu
index 262f4982f..a27840824 100644
--- a/tests/static_set/unique_sequence_test.cu
+++ b/tests/static_set/unique_sequence_test.cu
@@ -143,8 +143,7 @@ TEMPLATE_TEST_CASE_SIG(
   "static_set unique sequence tests",
   "",
   ((typename Key, cuco::test::probe_sequence Probe, int CGSize), Key, Probe, CGSize),
-  (int8_t, cuco::test::probe_sequence::double_hashing, 1),
-  (int8_t, cuco::test::probe_sequence::linear_probing, 1),
+  (int16_t, cuco::test::probe_sequence::double_hashing, 1),
   (int16_t, cuco::test::probe_sequence::double_hashing, 1),
   (int16_t, cuco::test::probe_sequence::double_hashing, 2),
   (int16_t, cuco::test::probe_sequence::linear_probing, 1),
@@ -176,7 +175,7 @@ TEMPLATE_TEST_CASE_SIG(
     if constexpr (cuco::is_double_hashing<probe>::value) {
       if constexpr (num_keys == 100) {
         return (CGSize == 1) ? 106   // 53 x 1 x 2
-                             : 106;  // 53 x 2 x 1... only CGSize=1 used for int8_t
+                             : 106;
       } else {
         return (CGSize == 1) ? 422   // 211 x 1 x 2
                              : 404;  // 101 x 2 x 2

From 35a9c5469c10b4e35d8864905c0f0dd411f63c92 Mon Sep 17 00:00:00 2001
From: Krish Sapru <ksapru@bu.edu>
Date: Sat, 9 May 2026 12:37:35 -0400
Subject: [PATCH 09/10] Final test retrieve_test

---
 tests/static_map/retrieve_test.cu | 1 +
 1 file changed, 1 insertion(+)

diff --git a/tests/static_map/retrieve_test.cu b/tests/static_map/retrieve_test.cu
index 8a4cd2acb..ad7380b43 100644
--- a/tests/static_map/retrieve_test.cu
+++ b/tests/static_map/retrieve_test.cu
@@ -105,6 +105,7 @@ TEMPLATE_TEST_CASE_SIG(
    Value,
    Probe,
    CGSize),
+  (int16_t, int16_t, cuco::test::probe_sequence::double_hashing, 1),
   (int32_t, int32_t, cuco::test::probe_sequence::double_hashing, 1),
   (int32_t, int64_t, cuco::test::probe_sequence::double_hashing, 1),
   (int32_t, int32_t, cuco::test::probe_sequence::double_hashing, 2),

From ce7e89fe735c8f89893686c765d9bda06fa81abb Mon Sep 17 00:00:00 2001
From: "pre-commit-ci[bot]"
 <66853113+pre-commit-ci[bot]@users.noreply.github.com>
Date: Sat, 9 May 2026 16:39:18 +0000
Subject: [PATCH 10/10] [pre-commit.ci] auto code formatting

---
 tests/static_map/shared_memory_test.cu   |  7 ++++---
 tests/static_map/stream_test.cu          | 10 +++++-----
 tests/static_set/unique_sequence_test.cu |  2 +-
 3 files changed, 10 insertions(+), 9 deletions(-)

diff --git a/tests/static_map/shared_memory_test.cu b/tests/static_map/shared_memory_test.cu
index 3624a6a5f..38624839d 100644
--- a/tests/static_map/shared_memory_test.cu
+++ b/tests/static_map/shared_memory_test.cu
@@ -82,7 +82,7 @@ TEMPLATE_TEST_CASE_SIG("static_map shared memory tests",
 )
 {
   // For small types, use smaller set count and element count
-  constexpr std::size_t number_of_maps = (sizeof(Key) + sizeof(Value) <= 2) ? 100 : 1000;
+  constexpr std::size_t number_of_maps  = (sizeof(Key) + sizeof(Value) <= 2) ? 100 : 1000;
   constexpr std::size_t elements_in_map = (sizeof(Key) + sizeof(Value) <= 2) ? 100 : 500;
   constexpr std::size_t map_capacity    = 2 * elements_in_map;
 
@@ -106,8 +106,9 @@ TEMPLATE_TEST_CASE_SIG("static_map shared memory tests",
   // operator yet
   std::vector<std::unique_ptr<map_type>> maps;
   for (std::size_t map_id = 0; map_id < number_of_maps; ++map_id) {
-    maps.push_back(std::make_unique<map_type>(
-      extent_type{}, cuco::empty_key<Key>{static_cast<Key>(-1)}, cuco::empty_value<Value>{static_cast<Value>(-1)}));
+    maps.push_back(std::make_unique<map_type>(extent_type{},
+                                              cuco::empty_key<Key>{static_cast<Key>(-1)},
+                                              cuco::empty_value<Value>{static_cast<Value>(-1)}));
   }
 
   thrust::device_vector<bool> d_keys_exist(number_of_maps * elements_in_map);
diff --git a/tests/static_map/stream_test.cu b/tests/static_map/stream_test.cu
index 445a9550d..85737a7ef 100644
--- a/tests/static_map/stream_test.cu
+++ b/tests/static_map/stream_test.cu
@@ -48,14 +48,14 @@ TEMPLATE_TEST_CASE_SIG("static_map: unique sequence of keys on given stream",
 
   {  // Scope ensures map is destroyed before stream
     const std::size_t num_keys = (sizeof(Key) == 1) ? 100 : 500'000;
-    auto map = cuco::static_map{num_keys * 2,
+    auto map                   = cuco::static_map{num_keys * 2,
                                 cuco::empty_key<Key>{static_cast<Key>(-1)},
                                 cuco::empty_value<Value>{static_cast<Value>(-1)},
-                                {},
+                                                  {},
                                 cuco::linear_probing<1, cuco::default_hash_function<Key>>{},
-                                {},
-                                {},
-                                {},
+                                                  {},
+                                                  {},
+                                                  {},
                                 stream};
 
     thrust::device_vector<Key> d_keys(num_keys);
diff --git a/tests/static_set/unique_sequence_test.cu b/tests/static_set/unique_sequence_test.cu
index a27840824..fa2da8725 100644
--- a/tests/static_set/unique_sequence_test.cu
+++ b/tests/static_set/unique_sequence_test.cu
@@ -174,7 +174,7 @@ TEMPLATE_TEST_CASE_SIG(
   constexpr size_type gold_capacity = [&]() {
     if constexpr (cuco::is_double_hashing<probe>::value) {
       if constexpr (num_keys == 100) {
-        return (CGSize == 1) ? 106   // 53 x 1 x 2
+        return (CGSize == 1) ? 106  // 53 x 1 x 2
                              : 106;
       } else {
         return (CGSize == 1) ? 422   // 211 x 1 x 2