From 4fdc73b3408fe5f2959e02c56b38cfe8f0ca3364 Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?Daniel=20J=C3=BCnger?=
 <2955913+sleeepyjack@users.noreply.github.com>
Date: Tue, 13 Feb 2024 00:05:31 +0100
Subject: [PATCH] Fix ODR violation where end() might be defined by multiple
 mixin operators (#440)

This PR fixes an ODR violation that occurs when mixing multiple
operators that define `end()` in their body, e.g., `find()` and
`insert_and_find().`

Solution: Move definition of `end()` to the `_ref` class instead of
defining it in the operators.
---
 .../cuco/detail/static_map/static_map_ref.inl | 91 ++++++++-----------
 .../cuco/detail/static_set/static_set_ref.inl | 86 +++++++-----------
 include/cuco/static_map_ref.cuh               | 14 +++
 include/cuco/static_set_ref.cuh               | 14 +++
 4 files changed, 101 insertions(+), 104 deletions(-)

diff --git a/include/cuco/detail/static_map/static_map_ref.inl b/include/cuco/detail/static_map/static_map_ref.inl
index fa39a0134..53ae0793a 100644
--- a/include/cuco/detail/static_map/static_map_ref.inl
+++ b/include/cuco/detail/static_map/static_map_ref.inl
@@ -122,6 +122,45 @@ static_map_ref<Key, T, Scope, KeyEqual, ProbingScheme, StorageRef, Operators...>
   return this->impl_.key_eq();
 }
 
+template <typename Key,
+          typename T,
+          cuda::thread_scope Scope,
+          typename KeyEqual,
+          typename ProbingScheme,
+          typename StorageRef,
+          typename... Operators>
+__host__ __device__ constexpr static_map_ref<Key,
+                                             T,
+                                             Scope,
+                                             KeyEqual,
+                                             ProbingScheme,
+                                             StorageRef,
+                                             Operators...>::const_iterator
+static_map_ref<Key, T, Scope, KeyEqual, ProbingScheme, StorageRef, Operators...>::end()
+  const noexcept
+{
+  return this->impl_.end();
+}
+
+template <typename Key,
+          typename T,
+          cuda::thread_scope Scope,
+          typename KeyEqual,
+          typename ProbingScheme,
+          typename StorageRef,
+          typename... Operators>
+__host__ __device__ constexpr static_map_ref<Key,
+                                             T,
+                                             Scope,
+                                             KeyEqual,
+                                             ProbingScheme,
+                                             StorageRef,
+                                             Operators...>::iterator
+static_map_ref<Key, T, Scope, KeyEqual, ProbingScheme, StorageRef, Operators...>::end() noexcept
+{
+  return this->impl_.end();
+}
+
 template <typename Key,
           typename T,
           cuda::thread_scope Scope,
@@ -517,32 +556,6 @@ class operator_impl<
   static constexpr auto window_size = base_type::window_size;
 
  public:
-  /**
-   * @brief Returns a const_iterator to one past the last slot.
-   *
-   * @note This API is available only when `find_tag` or `insert_and_find_tag` is present.
-   *
-   * @return A const_iterator to one past the last slot
-   */
-  [[nodiscard]] __host__ __device__ constexpr const_iterator end() const noexcept
-  {
-    auto const& ref_ = static_cast<ref_type const&>(*this);
-    return ref_.impl_.end();
-  }
-
-  /**
-   * @brief Returns an iterator to one past the last slot.
-   *
-   * @note This API is available only when `find_tag` or `insert_and_find_tag` is present.
-   *
-   * @return An iterator to one past the last slot
-   */
-  [[nodiscard]] __host__ __device__ constexpr iterator end() noexcept
-  {
-    auto const& ref_ = static_cast<ref_type const&>(*this);
-    return ref_.impl_.end();
-  }
-
   /**
    * @brief Inserts the given element into the map.
    *
@@ -724,32 +737,6 @@ class operator_impl<
   static constexpr auto window_size = base_type::window_size;
 
  public:
-  /**
-   * @brief Returns a const_iterator to one past the last slot.
-   *
-   * @note This API is available only when `find_tag` or `insert_and_find_tag` is present.
-   *
-   * @return A const_iterator to one past the last slot
-   */
-  [[nodiscard]] __host__ __device__ constexpr const_iterator end() const noexcept
-  {
-    auto const& ref_ = static_cast<ref_type const&>(*this);
-    return ref_.impl_.end();
-  }
-
-  /**
-   * @brief Returns an iterator to one past the last slot.
-   *
-   * @note This API is available only when `find_tag` or `insert_and_find_tag` is present.
-   *
-   * @return An iterator to one past the last slot
-   */
-  [[nodiscard]] __host__ __device__ constexpr iterator end() noexcept
-  {
-    auto const& ref_ = static_cast<ref_type const&>(*this);
-    return ref_.impl_.end();
-  }
-
   /**
    * @brief Finds an element in the map with key equivalent to the probe key.
    *
diff --git a/include/cuco/detail/static_set/static_set_ref.inl b/include/cuco/detail/static_set/static_set_ref.inl
index 3f0a9ed71..4464916c3 100644
--- a/include/cuco/detail/static_set/static_set_ref.inl
+++ b/include/cuco/detail/static_set/static_set_ref.inl
@@ -105,6 +105,40 @@ static_set_ref<Key, Scope, KeyEqual, ProbingScheme, StorageRef, Operators...>::k
   return this->impl_.key_eq();
 }
 
+template <typename Key,
+          cuda::thread_scope Scope,
+          typename KeyEqual,
+          typename ProbingScheme,
+          typename StorageRef,
+          typename... Operators>
+__host__ __device__ constexpr static_set_ref<Key,
+                                             Scope,
+                                             KeyEqual,
+                                             ProbingScheme,
+                                             StorageRef,
+                                             Operators...>::const_iterator
+static_set_ref<Key, Scope, KeyEqual, ProbingScheme, StorageRef, Operators...>::end() const noexcept
+{
+  return this->impl_.end();
+}
+
+template <typename Key,
+          cuda::thread_scope Scope,
+          typename KeyEqual,
+          typename ProbingScheme,
+          typename StorageRef,
+          typename... Operators>
+__host__ __device__ constexpr static_set_ref<Key,
+                                             Scope,
+                                             KeyEqual,
+                                             ProbingScheme,
+                                             StorageRef,
+                                             Operators...>::iterator
+static_set_ref<Key, Scope, KeyEqual, ProbingScheme, StorageRef, Operators...>::end() noexcept
+{
+  return this->impl_.end();
+}
+
 template <typename Key,
           cuda::thread_scope Scope,
           typename KeyEqual,
@@ -286,32 +320,6 @@ class operator_impl<op::insert_and_find_tag,
   static constexpr auto window_size = base_type::window_size;
 
  public:
-  /**
-   * @brief Returns a const_iterator to one past the last slot.
-   *
-   * @note This API is available only when `find_tag` or `insert_and_find_tag` is present.
-   *
-   * @return A const_iterator to one past the last slot
-   */
-  [[nodiscard]] __host__ __device__ constexpr const_iterator end() const noexcept
-  {
-    auto const& ref_ = static_cast<ref_type const&>(*this);
-    return ref_.impl_.end();
-  }
-
-  /**
-   * @brief Returns an iterator to one past the last slot.
-   *
-   * @note This API is available only when `find_tag` or `insert_and_find_tag` is present.
-   *
-   * @return An iterator to one past the last slot
-   */
-  [[nodiscard]] __host__ __device__ constexpr iterator end() noexcept
-  {
-    auto const& ref_ = static_cast<ref_type const&>(*this);
-    return ref_.impl_.end();
-  }
-
   /**
    * @brief Inserts the given element into the set.
    *
@@ -486,32 +494,6 @@ class operator_impl<op::find_tag,
   static constexpr auto window_size = base_type::window_size;
 
  public:
-  /**
-   * @brief Returns a const_iterator to one past the last slot.
-   *
-   * @note This API is available only when `find_tag` or `insert_and_find_tag` is present.
-   *
-   * @return A const_iterator to one past the last slot
-   */
-  [[nodiscard]] __host__ __device__ constexpr const_iterator end() const noexcept
-  {
-    auto const& ref_ = static_cast<ref_type const&>(*this);
-    return ref_.impl_.end();
-  }
-
-  /**
-   * @brief Returns an iterator to one past the last slot.
-   *
-   * @note This API is available only when `find_tag` or `insert_and_find_tag` is present.
-   *
-   * @return An iterator to one past the last slot
-   */
-  [[nodiscard]] __host__ __device__ constexpr iterator end() noexcept
-  {
-    auto const& ref_ = static_cast<ref_type const&>(*this);
-    return ref_.impl_.end();
-  }
-
   /**
    * @brief Finds an element in the set with key equivalent to the probe key.
    *
diff --git a/include/cuco/static_map_ref.cuh b/include/cuco/static_map_ref.cuh
index 16499ed38..48b53ecf5 100644
--- a/include/cuco/static_map_ref.cuh
+++ b/include/cuco/static_map_ref.cuh
@@ -185,6 +185,20 @@ class static_map_ref
    */
   [[nodiscard]] __host__ __device__ constexpr key_equal key_eq() const noexcept;
 
+  /**
+   * @brief Returns a const_iterator to one past the last slot.
+   *
+   * @return A const_iterator to one past the last slot
+   */
+  [[nodiscard]] __host__ __device__ constexpr const_iterator end() const noexcept;
+
+  /**
+   * @brief Returns an iterator to one past the last slot.
+   *
+   * @return An iterator to one past the last slot
+   */
+  [[nodiscard]] __host__ __device__ constexpr iterator end() noexcept;
+
   /**
    * @brief Creates a reference with new operators from the current object.
    *
diff --git a/include/cuco/static_set_ref.cuh b/include/cuco/static_set_ref.cuh
index cfb304082..39a1e5f3b 100644
--- a/include/cuco/static_set_ref.cuh
+++ b/include/cuco/static_set_ref.cuh
@@ -163,6 +163,20 @@ class static_set_ref
    */
   [[nodiscard]] __host__ __device__ constexpr key_equal key_eq() const noexcept;
 
+  /**
+   * @brief Returns a const_iterator to one past the last slot.
+   *
+   * @return A const_iterator to one past the last slot
+   */
+  [[nodiscard]] __host__ __device__ constexpr const_iterator end() const noexcept;
+
+  /**
+   * @brief Returns an iterator to one past the last slot.
+   *
+   * @return An iterator to one past the last slot
+   */
+  [[nodiscard]] __host__ __device__ constexpr iterator end() noexcept;
+
   /**
    * @brief Creates a reference with new operators from the current object.
    *