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 <[email protected]>

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);
 }
 
 /**

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);
 }
 
 /**