From 1968b9c85d7bad09dafdd858d62b960449241ba2 Mon Sep 17 00:00:00 2001
From: Michael Schellenberger Costa <miscco@nvidia.com>
Date: Fri, 23 Aug 2024 08:16:35 +0200
Subject: [PATCH] Work around gcc bug in pair_alignment

There seems to be a gcc bug where gcc versions prior to gcc13 get confused by passing in the output of cuda::std::bit_ceil into a constant expression.

This only manifests when trying to use the constant expression in an alignas direction.

However, we can convince gcc that it is indeed a constant expression by storing it in a constexpr variable before passing it along.

---------

Co-authored-by: Wesley Maxey <wmaxey@nvidia.com>
---
 include/cuco/detail/bitwise_compare.cuh | 6 ++++--
 include/cuco/detail/pair/helpers.cuh    | 5 ++++-
 2 files changed, 8 insertions(+), 3 deletions(-)

diff --git a/include/cuco/detail/bitwise_compare.cuh b/include/cuco/detail/bitwise_compare.cuh
index a8a5a69d1..d828bbf86 100644
--- a/include/cuco/detail/bitwise_compare.cuh
+++ b/include/cuco/detail/bitwise_compare.cuh
@@ -18,6 +18,7 @@
 
 #include <cuco/utility/traits.hpp>
 
+#include <cuda/functional>
 #include <cuda/std/bit>
 
 #include <cstdint>
@@ -67,9 +68,10 @@ struct bitwise_compare_impl<8> {
  * size of type, or 16, whichever is smaller.
  */
 template <typename T>
-constexpr std::size_t alignment()
+__host__ __device__ constexpr std::size_t alignment()
 {
-  return std::min(std::size_t{16}, cuda::std::bit_ceil(sizeof(T)));
+  constexpr std::size_t alignment = cuda::std::bit_ceil(sizeof(T));
+  return cuda::std::min(std::size_t{16}, alignment);
 }
 
 /**
diff --git a/include/cuco/detail/pair/helpers.cuh b/include/cuco/detail/pair/helpers.cuh
index 3cc343687..aa040ab94 100644
--- a/include/cuco/detail/pair/helpers.cuh
+++ b/include/cuco/detail/pair/helpers.cuh
@@ -19,6 +19,8 @@
 #include <cuda/std/bit>
 #include <cuda/std/type_traits>
 
+#include <cstdint>
+
 namespace cuco::detail {
 
 /**
@@ -29,7 +31,8 @@ namespace cuco::detail {
 template <typename First, typename Second>
 __host__ __device__ constexpr std::size_t pair_alignment()
 {
-  return cuda::std::min(std::size_t{16}, cuda::std::bit_ceil(sizeof(First) + sizeof(Second)));
+  constexpr std::size_t alignment = cuda::std::bit_ceil(sizeof(First) + sizeof(Second));
+  return cuda::std::min(std::size_t{16}, alignment);
 }
 
 /**