Several fixes to make cuCo compile with LLVM: (#733)

- add `template` keyword to help parsing.
- do not mark as `constexpr` functions that cannot be evaluated at
compile-time.

---------

Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>
Co-authored-by: Yunsong Wang <[email protected]>

Author: 3 people

include/cuco/detail/bloom_filter/bloom_filter_impl.cuh

@@ -363,44 +363,44 @@ class bloom_filter_impl {
   // const;
 
   template <class InputIt, class OutputIt>
-  __host__ constexpr void contains(InputIt first,
-                                   InputIt last,
-                                   OutputIt output_begin,
-                                   cuda::stream_ref stream) const
+  __host__ void contains(InputIt first,
+                         InputIt last,
+                         OutputIt output_begin,
+                         cuda::stream_ref stream) const
   {
     this->contains_async(first, last, output_begin, stream);
     stream.wait();
   }
 
   template <class InputIt, class OutputIt>
-  __host__ constexpr void contains_async(InputIt first,
-                                         InputIt last,
-                                         OutputIt output_begin,
-                                         cuda::stream_ref stream) const noexcept
+  __host__ void contains_async(InputIt first,
+                               InputIt last,
+                               OutputIt output_begin,
+                               cuda::stream_ref stream) const noexcept
   {
     auto const always_true = thrust::constant_iterator<bool>{true};
     this->contains_if_async(first, last, always_true, cuda::std::identity{}, output_begin, stream);
   }
 
   template <class InputIt, class StencilIt, class Predicate, class OutputIt>
-  __host__ constexpr void contains_if(InputIt first,
-                                      InputIt last,
-                                      StencilIt stencil,
-                                      Predicate pred,
-                                      OutputIt output_begin,
-                                      cuda::stream_ref stream) const
+  __host__ void contains_if(InputIt first,
+                            InputIt last,
+                            StencilIt stencil,
+                            Predicate pred,
+                            OutputIt output_begin,
+                            cuda::stream_ref stream) const
   {
     this->contains_if_async(first, last, stencil, pred, output_begin, stream);
     stream.wait();
   }
 
   template <class InputIt, class StencilIt, class Predicate, class OutputIt>
-  __host__ constexpr void contains_if_async(InputIt first,
-                                            InputIt last,
-                                            StencilIt stencil,
-                                            Predicate pred,
-                                            OutputIt output_begin,
-                                            cuda::stream_ref stream) const noexcept
+  __host__ void contains_if_async(InputIt first,
+                                  InputIt last,
+                                  StencilIt stencil,
+                                  Predicate pred,
+                                  OutputIt output_begin,
+                                  cuda::stream_ref stream) const noexcept
   {
     auto const num_keys = cuco::detail::distance(first, last);
     if (num_keys == 0) { return; }

include/cuco/detail/hyperloglog/kernels.cuh

@@ -42,7 +42,7 @@ CUCO_KERNEL void add_shmem_vectorized(typename RefType::value_type const* first,
 {
   using value_type     = typename RefType::value_type;
   using vector_type    = cuda::std::array<value_type, VectorSize>;
-  using local_ref_type = typename RefType::with_scope<cuda::thread_scope_block>;
+  using local_ref_type = typename RefType::template with_scope<cuda::thread_scope_block>;
 
   // Base address of dynamic shared memory is guaranteed to be aligned to at least 16 bytes which is
   // sufficient for this purpose
@@ -91,7 +91,7 @@ CUCO_KERNEL void add_shmem_vectorized(typename RefType::value_type const* first,
 template <class InputIt, class RefType>
 CUCO_KERNEL void add_shmem(InputIt first, cuco::detail::index_type n, RefType ref)
 {
-  using local_ref_type = typename RefType::with_scope<cuda::thread_scope_block>;
+  using local_ref_type = typename RefType::template with_scope<cuda::thread_scope_block>;
 
   // TODO assert alignment
   extern __shared__ cuda::std::byte local_sketch[];

include/cuco/detail/open_addressing/kernels.cuh

@@ -691,19 +691,19 @@ CUCO_KERNEL void retrieve(InputProbeIt input_probe,
 
   if (block_begin_offset < block_end_offset) {
     if constexpr (IsOuter) {
-      ref.retrieve_outer<BlockSize>(block,
-                                    input_probe + block_begin_offset,
-                                    input_probe + block_end_offset,
-                                    output_probe,
-                                    output_match,
-                                    atomic_counter);
+      ref.template retrieve_outer<BlockSize>(block,
+                                             input_probe + block_begin_offset,
+                                             input_probe + block_end_offset,
+                                             output_probe,
+                                             output_match,
+                                             atomic_counter);
     } else {
-      ref.retrieve<BlockSize>(block,
-                              input_probe + block_begin_offset,
-                              input_probe + block_end_offset,
-                              output_probe,
-                              output_match,
-                              atomic_counter);
+      ref.template retrieve<BlockSize>(block,
+                                       input_probe + block_begin_offset,
+                                       input_probe + block_end_offset,
+                                       output_probe,
+                                       output_match,
+                                       atomic_counter);
     }
   }
 }

include/cuco/detail/open_addressing/open_addressing_ref_impl.cuh

@@ -522,17 +522,18 @@ class open_addressing_ref_impl {
       "insert_and_find is not supported for pair types larger than 8 bytes on pre-Volta GPUs.");
 #endif
 
-    auto const val      = this->heterogeneous_value(value);
-    auto const key      = this->extract_key(val);
-    auto probing_iter   = probing_scheme_.make_iterator<bucket_size>(key, storage_ref_.extent());
+    auto const val = this->heterogeneous_value(value);
+    auto const key = this->extract_key(val);
+    auto probing_iter =
+      probing_scheme_.template make_iterator<bucket_size>(key, storage_ref_.extent());
     auto const init_idx = *probing_iter;
 
     while (true) {
       auto const bucket_slots = storage_ref_[*probing_iter];
 
       for (auto i = 0; i < bucket_size; ++i) {
-        auto const eq_res =
-          this->predicate_.operator()<is_insert::YES>(key, this->extract_key(bucket_slots[i]));
+        auto const eq_res = this->predicate_.template operator()<is_insert::YES>(
+          key, this->extract_key(bucket_slots[i]));
         auto* slot_ptr = this->get_slot_ptr(*probing_iter, i);
 
         // If the key is already in the container, return false
@@ -598,7 +599,7 @@ class open_addressing_ref_impl {
     auto const val = this->heterogeneous_value(value);
     auto const key = this->extract_key(val);
     auto probing_iter =
-      probing_scheme_.make_iterator<bucket_size>(group, key, storage_ref_.extent());
+      probing_scheme_.template make_iterator<bucket_size>(group, key, storage_ref_.extent());
     auto const init_idx = *probing_iter;
 
     while (true) {
@@ -607,8 +608,8 @@ class open_addressing_ref_impl {
       auto const [state, intra_bucket_index] = [&]() {
         auto res = detail::equal_result::UNEQUAL;
         for (auto i = 0; i < bucket_size; ++i) {
-          res =
-            this->predicate_.operator()<is_insert::YES>(key, this->extract_key(bucket_slots[i]));
+          res = this->predicate_.template operator()<is_insert::YES>(
+            key, this->extract_key(bucket_slots[i]));
           if (res != detail::equal_result::UNEQUAL) { return bucket_probing_results{res, i}; }
         }
         // returns dummy index `-1` for UNEQUAL
@@ -685,15 +686,16 @@ class open_addressing_ref_impl {
   {
     static_assert(cg_size == 1, "Non-CG operation is incompatible with the current probing scheme");
 
-    auto probing_iter   = probing_scheme_.make_iterator<bucket_size>(key, storage_ref_.extent());
+    auto probing_iter =
+      probing_scheme_.template make_iterator<bucket_size>(key, storage_ref_.extent());
     auto const init_idx = *probing_iter;
 
     while (true) {
       auto const bucket_slots = storage_ref_[*probing_iter];
 
       for (auto& slot_content : bucket_slots) {
         auto const eq_res =
-          this->predicate_.operator()<is_insert::NO>(key, this->extract_key(slot_content));
+          this->predicate_.template operator()<is_insert::NO>(key, this->extract_key(slot_content));
 
         // Key doesn't exist, return false
         if (eq_res == detail::equal_result::EMPTY) { return false; }
@@ -729,7 +731,7 @@ class open_addressing_ref_impl {
                         ProbeKey const& key) noexcept
   {
     auto probing_iter =
-      probing_scheme_.make_iterator<bucket_size>(group, key, storage_ref_.extent());
+      probing_scheme_.template make_iterator<bucket_size>(group, key, storage_ref_.extent());
     auto const init_idx = *probing_iter;
 
     while (true) {
@@ -738,7 +740,8 @@ class open_addressing_ref_impl {
       auto const [state, intra_bucket_index] = [&]() {
         auto res = detail::equal_result::UNEQUAL;
         for (auto i = 0; i < bucket_size; ++i) {
-          res = this->predicate_.operator()<is_insert::NO>(key, this->extract_key(bucket_slots[i]));
+          res = this->predicate_.template operator()<is_insert::NO>(
+            key, this->extract_key(bucket_slots[i]));
           if (res != detail::equal_result::UNEQUAL) { return bucket_probing_results{res, i}; }
         }
         // returns dummy index `-1` for UNEQUAL
@@ -824,7 +827,7 @@ class open_addressing_ref_impl {
     cooperative_groups::thread_block_tile<cg_size> const& group, ProbeKey const& key) const noexcept
   {
     auto probing_iter =
-      probing_scheme_.make_iterator<bucket_size>(group, key, storage_ref_.extent());
+      probing_scheme_.template make_iterator<bucket_size>(group, key, storage_ref_.extent());
     auto const init_idx = *probing_iter;
 
     while (true) {
@@ -833,7 +836,8 @@ class open_addressing_ref_impl {
       auto const state = [&]() {
         auto res = detail::equal_result::UNEQUAL;
         for (auto i = 0; i < bucket_size; ++i) {
-          res = this->predicate_.operator()<is_insert::NO>(key, this->extract_key(bucket_slots[i]));
+          res = this->predicate_.template operator()<is_insert::NO>(
+            key, this->extract_key(bucket_slots[i]));
           if (res != detail::equal_result::UNEQUAL) { return res; }
         }
         return res;
@@ -863,16 +867,17 @@ class open_addressing_ref_impl {
   [[nodiscard]] __device__ iterator find(ProbeKey const& key) const noexcept
   {
     static_assert(cg_size == 1, "Non-CG operation is incompatible with the current probing scheme");
-    auto probing_iter   = probing_scheme_.make_iterator<bucket_size>(key, storage_ref_.extent());
+    auto probing_iter =
+      probing_scheme_.template make_iterator<bucket_size>(key, storage_ref_.extent());
     auto const init_idx = *probing_iter;
 
     while (true) {
       // TODO atomic_ref::load if insert operator is present
       auto const bucket_slots = storage_ref_[*probing_iter];
 
       for (auto i = 0; i < bucket_size; ++i) {
-        switch (
-          this->predicate_.operator()<is_insert::NO>(key, this->extract_key(bucket_slots[i]))) {
+        switch (this->predicate_.template operator()<is_insert::NO>(
+          key, this->extract_key(bucket_slots[i]))) {
           case detail::equal_result::EMPTY: {
             return this->end();
           }
@@ -905,7 +910,7 @@ class open_addressing_ref_impl {
     cooperative_groups::thread_block_tile<cg_size> const& group, ProbeKey const& key) const noexcept
   {
     auto probing_iter =
-      probing_scheme_.make_iterator<bucket_size>(group, key, storage_ref_.extent());
+      probing_scheme_.template make_iterator<bucket_size>(group, key, storage_ref_.extent());
     auto const init_idx = *probing_iter;
 
     while (true) {
@@ -914,7 +919,8 @@ class open_addressing_ref_impl {
       auto const [state, intra_bucket_index] = [&]() {
         auto res = detail::equal_result::UNEQUAL;
         for (auto i = 0; i < bucket_size; ++i) {
-          res = this->predicate_.operator()<is_insert::NO>(key, this->extract_key(bucket_slots[i]));
+          res = this->predicate_.template operator()<is_insert::NO>(
+            key, this->extract_key(bucket_slots[i]));
           if (res != detail::equal_result::UNEQUAL) { return bucket_probing_results{res, i}; }
         }
         // returns dummy index `-1` for UNEQUAL
@@ -954,7 +960,8 @@ class open_addressing_ref_impl {
     if constexpr (not allows_duplicates) {
       return static_cast<size_type>(this->contains(key));
     } else {
-      auto probing_iter   = probing_scheme_.make_iterator<bucket_size>(key, storage_ref_.extent());
+      auto probing_iter =
+        probing_scheme_.template make_iterator<bucket_size>(key, storage_ref_.extent());
       auto const init_idx = *probing_iter;
       size_type count     = 0;
 
@@ -999,7 +1006,7 @@ class open_addressing_ref_impl {
     cooperative_groups::thread_block_tile<cg_size> const& group, ProbeKey const& key) const noexcept
   {
     auto probing_iter =
-      probing_scheme_.make_iterator<bucket_size>(group, key, storage_ref_.extent());
+      probing_scheme_.template make_iterator<bucket_size>(group, key, storage_ref_.extent());
     auto const init_idx = *probing_iter;
     size_type count     = 0;
 
@@ -1222,7 +1229,7 @@ class open_addressing_ref_impl {
         // perform probing
         // make sure the flushing_tile is converged at this point to get a coalesced load
         auto const probe_key = *(input_probe + idx);
-        auto probing_iter    = probing_scheme_.make_iterator<bucket_size>(
+        auto probing_iter    = probing_scheme_.template make_iterator<bucket_size>(
           probing_tile, probe_key, storage_ref_.extent());
         auto const init_idx = *probing_iter;
 
@@ -1242,7 +1249,7 @@ class open_addressing_ref_impl {
               equals[i] = false;
               if (running) {
                 // inspect slot content
-                switch (this->predicate_.operator()<is_insert::NO>(
+                switch (this->predicate_.template operator()<is_insert::NO>(
                   probe_key, this->extract_key(bucket_slots[i]))) {
                   case detail::equal_result::EMPTY: {
                     running = false;
@@ -1354,16 +1361,17 @@ class open_addressing_ref_impl {
   __device__ void for_each(ProbeKey const& key, CallbackOp&& callback_op) const noexcept
   {
     static_assert(cg_size == 1, "Non-CG operation is incompatible with the current probing scheme");
-    auto probing_iter   = probing_scheme_.make_iterator<bucket_size>(key, storage_ref_.extent());
+    auto probing_iter =
+      probing_scheme_.template make_iterator<bucket_size>(key, storage_ref_.extent());
     auto const init_idx = *probing_iter;
 
     while (true) {
       // TODO atomic_ref::load if insert operator is present
       auto const bucket_slots = this->storage_ref_[*probing_iter];
 
       for (int32_t i = 0; i < bucket_size; ++i) {
-        switch (
-          this->predicate_.operator()<is_insert::NO>(key, this->extract_key(bucket_slots[i]))) {
+        switch (this->predicate_.template operator()<is_insert::NO>(
+          key, this->extract_key(bucket_slots[i]))) {
           case detail::equal_result::EMPTY: {
             return;
           }
@@ -1404,7 +1412,7 @@ class open_addressing_ref_impl {
                            CallbackOp&& callback_op) const noexcept
   {
     auto probing_iter =
-      probing_scheme_.make_iterator<bucket_size>(group, key, storage_ref_.extent());
+      probing_scheme_.template make_iterator<bucket_size>(group, key, storage_ref_.extent());
     auto const init_idx = *probing_iter;
     bool empty          = false;
 
@@ -1413,8 +1421,8 @@ class open_addressing_ref_impl {
       auto const bucket_slots = this->storage_ref_[*probing_iter];
 
       for (int32_t i = 0; i < bucket_size and !empty; ++i) {
-        switch (
-          this->predicate_.operator()<is_insert::NO>(key, this->extract_key(bucket_slots[i]))) {
+        switch (this->predicate_.template operator()<is_insert::NO>(
+          key, this->extract_key(bucket_slots[i]))) {
           case detail::equal_result::EMPTY: {
             empty = true;
             continue;
@@ -1469,7 +1477,7 @@ class open_addressing_ref_impl {
                            SyncOp&& sync_op) const noexcept
   {
     auto probing_iter =
-      probing_scheme_.make_iterator<bucket_size>(group, key, storage_ref_.extent());
+      probing_scheme_.template make_iterator<bucket_size>(group, key, storage_ref_.extent());
     auto const init_idx = *probing_iter;
     bool empty          = false;
 
@@ -1478,8 +1486,8 @@ class open_addressing_ref_impl {
       auto const bucket_slots = this->storage_ref_[*probing_iter];
 
       for (int32_t i = 0; i < bucket_size and !empty; ++i) {
-        switch (
-          this->predicate_.operator()<is_insert::NO>(key, this->extract_key(bucket_slots[i]))) {
+        switch (this->predicate_.template operator()<is_insert::NO>(
+          key, this->extract_key(bucket_slots[i]))) {
           case detail::equal_result::EMPTY: {
             empty = true;
             continue;

include/cuco/detail/static_map/kernels.cuh

@@ -191,7 +191,8 @@ CUCO_KERNEL __launch_bounds__(BlockSize) void insert_or_apply_shmem(
 
   // Shared map initialization
   __shared__ typename SharedMapRefType::value_type slots[bucket_extent.value()];
-  auto storage           = SharedMapRefType::storage_ref_type(bucket_extent, slots);
+  using storage_ref_type = typename SharedMapRefType::storage_ref_type;
+  auto storage           = storage_ref_type(bucket_extent, slots);
   auto const num_buckets = storage.num_buckets();
 
   using atomic_type = cuda::atomic<int32_t, cuda::thread_scope_block>;

include/cuco/detail/static_map/static_map.inl

@@ -703,7 +703,7 @@ std::pair<KeyOut, ValueOut>
 static_map<Key, T, Extent, Scope, KeyEqual, ProbingScheme, Allocator, Storage>::retrieve_all(
   KeyOut keys_out, ValueOut values_out, cuda::stream_ref stream) const
 {
-  auto const zipped_out_begin = thrust::make_zip_iterator(cuda::std::tuple{keys_out, values_out});
+  auto const zipped_out_begin = thrust::make_zip_iterator(keys_out, values_out);
   auto const zipped_out_end   = impl_->retrieve_all(zipped_out_begin, stream);
   auto const num              = std::distance(zipped_out_begin, zipped_out_end);