Fix static_map::rehash and add unit test

include/cuco/detail/static_map/static_map.inl

@@ -358,7 +358,7 @@ void static_map<Key, T, Extent, Scope, KeyEqual, ProbingScheme, Allocator, Stora
   cuda_stream_ref stream)
 {
   auto const is_filled = static_map_ns::detail::slot_is_filled<Key, T>(this->empty_key_sentinel());
-  this->impl_.rehash(this, is_filled, stream);
+  this->impl_->rehash(*this, is_filled, stream);
 }
 
 template <class Key,
@@ -374,7 +374,7 @@ void static_map<Key, T, Extent, Scope, KeyEqual, ProbingScheme, Allocator, Stora
 {
   auto const is_filled = static_map_ns::detail::slot_is_filled<Key, T>(this->empty_key_sentinel());
   auto const extent    = make_window_extent<static_map>(capacity);
-  this->impl_.rehash(extent, this, is_filled, stream);
+  this->impl_->rehash(extent, *this, is_filled, stream);
 }
 
 template <class Key,
@@ -389,7 +389,7 @@ void static_map<Key, T, Extent, Scope, KeyEqual, ProbingScheme, Allocator, Stora
   cuda_stream_ref stream)
 {
   auto const is_filled = static_map_ns::detail::slot_is_filled<Key, T>(this->empty_key_sentinel());
-  this->impl_.rehash_async(this, is_filled, stream);
+  this->impl_->rehash_async(*this, is_filled, stream);
 }
 
 template <class Key,
@@ -405,7 +405,7 @@ void static_map<Key, T, Extent, Scope, KeyEqual, ProbingScheme, Allocator, Stora
 {
   auto const is_filled = static_map_ns::detail::slot_is_filled<Key, T>(this->empty_key_sentinel());
   auto const extent    = make_window_extent<static_map>(capacity);
-  this->impl_.rehash_async(extent, this, is_filled, stream);
+  this->impl_->rehash_async(extent, *this, is_filled, stream);
 }
 
 template <class Key,

tests/CMakeLists.txt

@@ -80,7 +80,8 @@ ConfigureTest(STATIC_MAP_TEST
     static_map/key_sentinel_test.cu
     static_map/shared_memory_test.cu
     static_map/stream_test.cu
-    static_map/unique_sequence_test.cu)
+    static_map/unique_sequence_test.cu
+    static_map/rehash_test.cu)
 
 ###################################################################################################
 # - dynamic_map tests -----------------------------------------------------------------------------

tests/static_map/rehash_test.cu

@@ -0,0 +1,55 @@
+/*
+ * Copyright (c) 2023, NVIDIA CORPORATION.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#include <cuco/static_map.cuh>
+
+#include <thrust/device_vector.h>
+#include <thrust/iterator/zip_iterator.h>
+#include <thrust/sequence.h>
+#include <thrust/tuple.h>
+
+#include <catch2/catch_test_macros.hpp>
+
+TEST_CASE("Rehash", "")
+{
+  using key_type    = int;
+  using mapped_type = long;
+  constexpr std::size_t num_keys{400};
+
+  cuco::experimental::static_map<key_type, mapped_type> map{
+    num_keys, cuco::empty_key<key_type>{-1}, cuco::empty_value<mapped_type>{-1}};
+
+  thrust::device_vector<key_type> d_keys(num_keys);
+  thrust::device_vector<mapped_type> d_values(num_keys);
+
+  thrust::sequence(d_keys.begin(), d_keys.end());
+  thrust::sequence(d_values.begin(), d_values.end());
+
+  auto pairs_begin =
+    thrust::make_zip_iterator(thrust::make_tuple(d_keys.begin(), d_values.begin()));
+
+  map.insert(pairs_begin, pairs_begin + num_keys);
+
+  map.rehash();
+  REQUIRE(map.size() == num_keys);
+
+  map.rehash(num_keys * 2);
+  REQUIRE(map.size() == num_keys);
+
+  // TODO erase num_erased keys
+  // map.rehash()
+  // REQUIRE(map.size() == num_keys - num_erased);
+}