Yet more fixes to make cuCollections compile with LLVM.

Most of the changes are just helping the compiler to parse the template syntax.

Author: vitor1001

include/cuco/detail/bloom_filter/bloom_filter.inl

@@ -163,7 +163,7 @@ template <class Key, class Extent, cuda::thread_scope Scope, class Policy, class
 
 template <class Key, class Extent, cuda::thread_scope Scope, class Policy, class Allocator>
 [[nodiscard]] __host__ constexpr
-  typename bloom_filter<Key, Extent, Scope, Policy, Allocator>::ref_type<>
+  typename bloom_filter<Key, Extent, Scope, Policy, Allocator>::template ref_type<>
   bloom_filter<Key, Extent, Scope, Policy, Allocator>::ref() const noexcept
 {
   return ref_;

include/cuco/detail/hyperloglog/hyperloglog.inl

@@ -119,7 +119,7 @@ constexpr std::size_t hyperloglog<T, Scope, Hash, Allocator>::estimate(
 }
 
 template <class T, cuda::thread_scope Scope, class Hash, class Allocator>
-constexpr typename hyperloglog<T, Scope, Hash, Allocator>::ref_type<>
+constexpr typename hyperloglog<T, Scope, Hash, Allocator>::template ref_type<>
 hyperloglog<T, Scope, Hash, Allocator>::ref() const noexcept
 {
   return {this->sketch(), this->hash_function()};

include/cuco/detail/open_addressing/open_addressing_impl.cuh

@@ -1299,7 +1299,7 @@ class open_addressing_impl {
    */
   [[nodiscard]] constexpr key_type const& extract_key(value_type const& slot) const noexcept
   {
-    if constexpr (this->has_payload) {
+    if constexpr (has_payload) {
       return slot.first;
     } else {
       return slot;

include/cuco/detail/open_addressing/open_addressing_ref_impl.cuh

@@ -379,15 +379,15 @@ 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>(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::YES>(key, this->extract_key(slot_content));
+          this->predicate_.template operator()<is_insert::YES>(key, this->extract_key(slot_content));
 
         if constexpr (not allows_duplicates) {
           // If the key is already in the container, return false
@@ -431,7 +431,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) {
@@ -440,7 +440,7 @@ class open_addressing_ref_impl {
       auto const [state, intra_bucket_index] = [&]() {
         for (auto i = 0; i < bucket_size; ++i) {
           switch (
-            this->predicate_.operator()<is_insert::YES>(key, this->extract_key(bucket_slots[i]))) {
+            this->predicate_.template operator()<is_insert::YES>(key, this->extract_key(bucket_slots[i]))) {
             case detail::equal_result::AVAILABLE:
               return bucket_probing_results{detail::equal_result::AVAILABLE, i};
             case detail::equal_result::EQUAL: {
@@ -789,7 +789,7 @@ class open_addressing_ref_impl {
   [[nodiscard]] __device__ bool contains(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) {
@@ -798,7 +798,7 @@ class open_addressing_ref_impl {
 
       for (auto i = 0; i < bucket_size; ++i) {
         switch (
-          this->predicate_.operator()<is_insert::NO>(key, this->extract_key(bucket_slots[i]))) {
+          this->predicate_.template operator()<is_insert::NO>(key, this->extract_key(bucket_slots[i]))) {
           case detail::equal_result::UNEQUAL: continue;
           case detail::equal_result::EMPTY: return false;
           case detail::equal_result::EQUAL: return true;
@@ -1534,7 +1534,7 @@ class open_addressing_ref_impl {
   template <typename Value>
   [[nodiscard]] __host__ __device__ constexpr auto extract_key(Value const& value) const noexcept
   {
-    if constexpr (this->has_payload) {
+    if constexpr (has_payload) {
       return thrust::raw_reference_cast(value).first;
     } else {
       return thrust::raw_reference_cast(value);
@@ -1570,7 +1570,7 @@ class open_addressing_ref_impl {
   template <typename T>
   [[nodiscard]] __device__ constexpr value_type native_value(T const& value) const noexcept
   {
-    if constexpr (this->has_payload) {
+    if constexpr (has_payload) {
       return {static_cast<key_type>(this->extract_key(value)), this->extract_payload(value)};
     } else {
       return static_cast<value_type>(value);
@@ -1590,7 +1590,7 @@ class open_addressing_ref_impl {
   template <typename T>
   [[nodiscard]] __device__ constexpr auto heterogeneous_value(T const& value) const noexcept
   {
-    if constexpr (this->has_payload and not cuda::std::is_same_v<T, value_type>) {
+    if constexpr (has_payload and not cuda::std::is_same_v<T, value_type>) {
       using mapped_type = decltype(this->empty_value_sentinel());
       if constexpr (cuco::detail::is_cuda_std_pair_like<T>::value) {
         return cuco::pair{cuda::std::get<0>(value),
@@ -1612,7 +1612,7 @@ class open_addressing_ref_impl {
    */
   [[nodiscard]] __device__ constexpr value_type const erased_slot_sentinel() const noexcept
   {
-    if constexpr (this->has_payload) {
+    if constexpr (has_payload) {
       return cuco::pair{this->erased_key_sentinel(), this->empty_value_sentinel()};
     } else {
       return this->erased_key_sentinel();

include/cuco/detail/static_map.inl

@@ -402,7 +402,7 @@ template <typename Hash, typename KeyEqual>
 __device__ bool static_map<Key, Value, Scope, Allocator>::device_mutable_view::insert(
   value_type const& insert_pair, Hash hash, KeyEqual key_equal) noexcept
 {
-  auto current_slot{initial_slot(insert_pair.first, hash)};
+  auto current_slot{this->initial_slot(insert_pair.first, hash)};
 
   while (true) {
     key_type const existing_key = current_slot->first.load(cuda::std::memory_order_relaxed);
@@ -439,7 +439,7 @@ __device__ bool static_map<Key, Value, Scope, Allocator>::device_mutable_view::i
 
     // if we couldn't insert the key, but it wasn't a duplicate, then there must
     // have been some other key there, so we keep looking for a slot
-    current_slot = next_slot(current_slot);
+    current_slot = this->next_slot(current_slot);
   }
 }
 
@@ -538,7 +538,7 @@ template <typename CG, typename Hash, typename KeyEqual>
 __device__ bool static_map<Key, Value, Scope, Allocator>::device_mutable_view::insert(
   CG const& g, value_type const& insert_pair, Hash hash, KeyEqual key_equal) noexcept
 {
-  auto current_slot = initial_slot(g, insert_pair.first, hash);
+  auto current_slot = this->initial_slot(g, insert_pair.first, hash);
 
   while (true) {
     key_type const existing_key = current_slot->first.load(cuda::std::memory_order_relaxed);
@@ -592,7 +592,7 @@ __device__ bool static_map<Key, Value, Scope, Allocator>::device_mutable_view::i
     // if there are no empty slots in the current bucket,
     // we move onto the next bucket
     else {
-      current_slot = next_slot(g, current_slot);
+      current_slot = this->next_slot(g, current_slot);
     }
   }
 }
@@ -602,7 +602,7 @@ template <typename Hash, typename KeyEqual>
 __device__ bool static_map<Key, Value, Scope, Allocator>::device_mutable_view::erase(
   key_type const& k, Hash hash, KeyEqual key_equal) noexcept
 {
-  auto current_slot{initial_slot(k, hash)};
+  auto current_slot{this->initial_slot(k, hash)};
 
   value_type const insert_pair =
     make_pair<Key, Value>(this->get_erased_key_sentinel(), this->get_empty_value_sentinel());
@@ -641,7 +641,7 @@ __device__ bool static_map<Key, Value, Scope, Allocator>::device_mutable_view::e
       }
     }
 
-    current_slot = next_slot(current_slot);
+    current_slot = this->next_slot(current_slot);
   }
 }
 
@@ -650,7 +650,7 @@ template <typename CG, typename Hash, typename KeyEqual>
 __device__ bool static_map<Key, Value, Scope, Allocator>::device_mutable_view::erase(
   CG const& g, key_type const& k, Hash hash, KeyEqual key_equal) noexcept
 {
-  auto current_slot = initial_slot(g, k, hash);
+  auto current_slot = this->initial_slot(g, k, hash);
   value_type const insert_pair =
     make_pair<Key, Value>(this->get_erased_key_sentinel(), this->get_empty_value_sentinel());
 
@@ -699,7 +699,7 @@ __device__ bool static_map<Key, Value, Scope, Allocator>::device_mutable_view::e
     // empty slot found, but key not found, must not be in the map
     if (g.ballot(slot_is_empty)) { return false; }
 
-    current_slot = next_slot(g, current_slot);
+    current_slot = this->next_slot(g, current_slot);
   }
 }
 
@@ -710,7 +710,7 @@ static_map<Key, Value, Scope, Allocator>::device_view::find(Key const& k,
                                                             Hash hash,
                                                             KeyEqual key_equal) noexcept
 {
-  auto current_slot = initial_slot(k, hash);
+  auto current_slot = this->initial_slot(k, hash);
 
   while (true) {
     auto const existing_key = current_slot->first.load(cuda::std::memory_order_relaxed);
@@ -722,7 +722,7 @@ static_map<Key, Value, Scope, Allocator>::device_view::find(Key const& k,
     // Key exists, return iterator to location
     if (key_equal(existing_key, k)) { return current_slot; }
 
-    current_slot = next_slot(current_slot);
+    current_slot = this->next_slot(current_slot);
   }
 }
 
@@ -733,7 +733,7 @@ static_map<Key, Value, Scope, Allocator>::device_view::find(Key const& k,
                                                             Hash hash,
                                                             KeyEqual key_equal) const noexcept
 {
-  auto current_slot = initial_slot(k, hash);
+  auto current_slot = this->initial_slot(k, hash);
 
   while (true) {
     auto const existing_key = current_slot->first.load(cuda::std::memory_order_relaxed);
@@ -745,7 +745,7 @@ static_map<Key, Value, Scope, Allocator>::device_view::find(Key const& k,
     // Key exists, return iterator to location
     if (key_equal(existing_key, k)) { return current_slot; }
 
-    current_slot = next_slot(current_slot);
+    current_slot = this->next_slot(current_slot);
   }
 }
 
@@ -757,7 +757,7 @@ static_map<Key, Value, Scope, Allocator>::device_view::find(CG g,
                                                             Hash hash,
                                                             KeyEqual key_equal) noexcept
 {
-  auto current_slot = initial_slot(g, k, hash);
+  auto current_slot = this->initial_slot(g, k, hash);
 
   while (true) {
     auto const existing_key = current_slot->first.load(cuda::std::memory_order_relaxed);
@@ -783,7 +783,7 @@ static_map<Key, Value, Scope, Allocator>::device_view::find(CG g,
 
     // otherwise, all slots in the current bucket are full with other keys, so we move onto the
     // next bucket
-    current_slot = next_slot(g, current_slot);
+    current_slot = this->next_slot(g, current_slot);
   }
 }
 
@@ -795,7 +795,7 @@ static_map<Key, Value, Scope, Allocator>::device_view::find(CG g,
                                                             Hash hash,
                                                             KeyEqual key_equal) const noexcept
 {
-  auto current_slot = initial_slot(g, k, hash);
+  auto current_slot = this->initial_slot(g, k, hash);
 
   while (true) {
     auto const existing_key = current_slot->first.load(cuda::std::memory_order_relaxed);
@@ -823,7 +823,7 @@ static_map<Key, Value, Scope, Allocator>::device_view::find(CG g,
     // otherwise, all slots in the current bucket are full with other keys,
     // so we move onto the next bucket in the current submap
 
-    current_slot = next_slot(g, current_slot);
+    current_slot = this->next_slot(g, current_slot);
   }
 }
 
@@ -832,7 +832,7 @@ template <typename ProbeKey, typename Hash, typename KeyEqual>
 __device__ bool static_map<Key, Value, Scope, Allocator>::device_view::contains(
   ProbeKey const& k, Hash hash, KeyEqual key_equal) const noexcept
 {
-  auto current_slot = initial_slot(k, hash);
+  auto current_slot = this->initial_slot(k, hash);
 
   while (true) {
     auto const existing_key = current_slot->first.load(cuda::std::memory_order_relaxed);
@@ -841,7 +841,7 @@ __device__ bool static_map<Key, Value, Scope, Allocator>::device_view::contains(
 
     if (key_equal(existing_key, k)) { return true; }
 
-    current_slot = next_slot(current_slot);
+    current_slot = this->next_slot(current_slot);
   }
 }
 
@@ -853,7 +853,7 @@ static_map<Key, Value, Scope, Allocator>::device_view::contains(CG const& g,
                                                                 Hash hash,
                                                                 KeyEqual key_equal) const noexcept
 {
-  auto current_slot = initial_slot(g, k, hash);
+  auto current_slot = this->initial_slot(g, k, hash);
 
   while (true) {
     key_type const existing_key = current_slot->first.load(cuda::std::memory_order_relaxed);
@@ -872,7 +872,7 @@ static_map<Key, Value, Scope, Allocator>::device_view::contains(CG const& g,
 
     // otherwise, all slots in the current bucket are full with other keys, so we move onto the
     // next bucket
-    current_slot = next_slot(g, current_slot);
+    current_slot = this->next_slot(g, current_slot);
   }
 }
 }  // namespace cuco::legacy

include/cuco/detail/static_map/helpers.cuh

@@ -75,7 +75,7 @@ void dispatch_insert_or_apply(
                                              Allocator,
                                              cuco::storage<1>>;
 
-    using shared_map_ref_type = typename shared_map_type::ref_type<>;
+    using shared_map_ref_type = typename shared_map_type::template ref_type<>;
     auto constexpr bucket_extent =
       cuco::make_valid_extent<typename shared_map_ref_type::probing_scheme_type,
                               typename shared_map_ref_type::storage_ref_type>(extent_type{});

include/cuco/detail/static_map/static_map_ref.inl

@@ -471,9 +471,9 @@ class operator_impl<
   {
     auto& ref_ = static_cast<ref_type&>(*this);
     if (ref_.erased_key_sentinel() != ref_.empty_key_sentinel()) {
-      return ref_.impl_.insert<true>(group, value);
+      return ref_.impl_.template insert<true>(group, value);
     } else {
-      return ref_.impl_.insert<false>(group, value);
+      return ref_.impl_.template insert<false>(group, value);
     }
   }
 };

include/cuco/detail/trie/dynamic_bitset/dynamic_bitset.cuh

@@ -85,7 +85,8 @@ class dynamic_bitset {
   using size_type = std::size_t;  ///< size type to specify bit index
   using word_type = uint64_t;     ///< word type
   /// Type of the allocator to (de)allocate words
-  using allocator_type = typename std::allocator_traits<Allocator>::rebind_alloc<word_type>;
+  using allocator_type =
+    typename std::allocator_traits<Allocator>::template rebind_alloc<word_type>;
 
   /// Number of bits per block. Note this is a tradeoff between space efficiency and perf.
   static constexpr size_type words_per_block = 4;
@@ -326,9 +327,11 @@ class dynamic_bitset {
 
  private:
   /// Type of the allocator to (de)allocate ranks
-  using rank_allocator_type = typename std::allocator_traits<Allocator>::rebind_alloc<rank_type>;
+  using rank_allocator_type =
+    typename std::allocator_traits<Allocator>::template rebind_alloc<rank_type>;
   /// Type of the allocator to (de)allocate indices
-  using size_allocator_type = typename std::allocator_traits<Allocator>::rebind_alloc<size_type>;
+  using size_allocator_type =
+    typename std::allocator_traits<Allocator>::template rebind_alloc<size_type>;
 
   allocator_type allocator_;  ///< Words allocator
   size_type n_bits_;          ///< Number of bits dynamic_bitset currently holds

include/cuco/detail/trie/dynamic_bitset/dynamic_bitset.inl

@@ -147,8 +147,9 @@ constexpr void dynamic_bitset<Allocator>::build_ranks_and_selects(
     thrust::raw_pointer_cast(words_.data()), bit_counts_begin, num_words, flip_bits);
 
   std::size_t temp_storage_bytes = 0;
-  using temp_allocator_type = typename std::allocator_traits<allocator_type>::rebind_alloc<char>;
-  auto temp_allocator       = temp_allocator_type{this->allocator_};
+  using temp_allocator_type =
+    typename std::allocator_traits<allocator_type>::template rebind_alloc<char>;
+  auto temp_allocator = temp_allocator_type{this->allocator_};
 
   CUCO_CUDA_TRY(cub::DeviceScan::ExclusiveSum(nullptr,
                                               temp_storage_bytes,

include/cuco/static_multimap.cuh

@@ -1031,7 +1031,7 @@ class static_multimap {
   using pair_atomic_type =
     cuco::pair<atomic_key_type,
                atomic_mapped_type>;  ///< Pair type of atomic key and atomic mapped value
-  using allocator_type = typename std::allocator_traits<Allocator>::rebind_alloc<
+  using allocator_type = typename std::allocator_traits<Allocator>::template rebind_alloc<
     pair_atomic_type>;  ///< Type of the allocator to (de)allocate slots
   using probe_sequence_type =
     cuco::legacy::detail::probe_sequence<ProbeSequence, Key, Value, Scope>;  ///< Probe scheme type