Enable single-threaded device add

NVIDIA · Aug 27, 2024 · 26ab9f4 · 26ab9f4
1 parent 36b1532
commit 26ab9f4
Show file tree

Hide file tree

Showing 2 changed files with 15 additions and 3 deletions.
diff --git a/include/cuco/bloom_filter_ref.cuh b/include/cuco/bloom_filter_ref.cuh
@@ -58,9 +58,8 @@ class bloom_filter_ref {
 
   __host__ void clear_async(cuda::stream_ref stream = {});
 
-  // TODO
-  // template <class ProbeKey>
-  // __device__ void add(ProbeKey const& key);
+  template <class ProbeKey>
+  __device__ void add(ProbeKey const& key);
 
   template <class CG, class ProbeKey>
   __device__ void add(CG const& group, ProbeKey const& key);

diff --git a/include/cuco/detail/bloom_filter/bloom_filter_ref.inl b/include/cuco/detail/bloom_filter/bloom_filter_ref.inl
@@ -77,6 +77,19 @@ __host__ void bloom_filter_ref<Key, Extent, Scope, Hash, BlockWords, Word>::clea
   impl_.clear_async(stream);
 }
 
+template <class Key,
+          class Extent,
+          cuda::thread_scope Scope,
+          class Hash,
+          std::uint32_t BlockWords,
+          class Word>
+template <class ProbeKey>
+__device__ void bloom_filter_ref<Key, Extent, Scope, Hash, BlockWords, Word>::add(
+  ProbeKey const& key)
+{
+  impl_.add(key);
+}
+
 template <class Key,
           class Extent,
           cuda::thread_scope Scope,