cuda::is_bitwise_comparable

docs/libcudacxx/extended_api/type_traits.rst

@@ -8,6 +8,7 @@ Type traits
    :maxdepth: 1
 
    type_traits/is_floating_point
+   type_traits/is_bitwise_comparable
    type_traits/vector_types
 
 .. list-table::
@@ -28,3 +29,8 @@ Type traits
      - Type traits for CUDA vector types
      - CCCL 3.3.0
      - CUDA 13.3
+
+   * - :ref:`cuda::is_bitwise_comparable <libcudacxx-extended-api-type_traits-is_bitwise_comparable>`
+     - User-specializable bitwise comparability check
+     - CCCL 3.4.0
+     - CUDA 13.4

docs/libcudacxx/extended_api/type_traits/is_bitwise_comparable.rst

@@ -0,0 +1,81 @@
+.. _libcudacxx-extended-api-type_traits-is_bitwise_comparable:
+
+``cuda::is_bitwise_comparable``
+===============================
+
+Defined in the ``<cuda/type_traits>`` header.
+
+.. code:: cuda
+
+   namespace cuda {
+
+   template <typename T>
+   constexpr bool is_bitwise_comparable_v = /* see below */;
+
+   template <typename T>
+   using is_bitwise_comparable = cuda::std::bool_constant<is_bitwise_comparable_v<T>>;
+
+   } // namespace cuda
+
+``cuda::is_bitwise_comparable_v<T>`` is a user-specializable variable template that forwards to ``cuda::std::has_unique_object_representations`` but excludes extended floating-point scalar and vector types.
+
+The trait also propagates through composite types:
+
+- C-style arrays: ``T[N]`` and ``T[]`` are bitwise comparable when ``T`` is.
+- ``cuda::std::array<T, N>``: bitwise comparable when ``T`` is.
+- ``cuda::std::pair<T1, T2>``: bitwise comparable when both ``T1`` and ``T2`` are and the object has no padding.
+- ``cuda::std::tuple<Ts...>``: bitwise comparable when all ``Ts...`` are and the object has no padding.
+
+``const``, ``volatile``, and ``const volatile`` qualifications are handled transparently.
+
+.. note::
+
+  The following conditions generally prevent a type from being bitwise comparable:
+
+  - The type contains floating-point types. ``NaN`` and ``+/-0`` values are not bitwise comparable.
+  - The type has internal padding, for example a structure with ``char`` and ``int`` members.
+  - The type has special semantics for the comparison, namely a user-defined ``operator==``.
+
+.. warning::
+
+   The type trait cannot determine if a structure (``struct`` or ``class``) contains *extended* floating-point types, and thus it recognizes the type as *bitwise comparable*. The user must manually specialize the type trait for such types.
+
+Custom Specialization
+---------------------
+
+Users may specialize ``cuda::is_bitwise_comparable_v`` for their own types to indicate that two object representations can be compared bitwise, even when the compiler cannot determine this automatically.
+The specialization must be provided for the unqualified type; cv-qualified forms are handled automatically.
+
+.. code:: cuda
+
+  struct MyType {
+    double value;
+  };
+
+  template <>
+  constexpr bool cuda::is_bitwise_comparable_v<MyType> = true;
+
+  static_assert(cuda::is_bitwise_comparable_v<MyType>);
+  static_assert(cuda::is_bitwise_comparable_v<const MyType>);
+
+
+.. warning::
+
+  Users are responsible for ensuring that the type is actually bitwise comparable when specializing this variable template. Otherwise, the behavior of functionalities that rely on this trait is undefined. For the reasons described above, bitwise comparison is especially problematic for floating-point types and internal padding.
+
+Examples
+--------
+
+.. code:: cuda
+
+   #include <cuda/type_traits>
+
+   // Integer types have unique object representations
+   static_assert(cuda::is_bitwise_comparable_v<int>);
+   static_assert(cuda::is_bitwise_comparable_v<unsigned>);
+   static_assert(cuda::is_bitwise_comparable_v<char>);
+   static_assert(cuda::is_bitwise_comparable_v<int[4]>);
+
+   // Floating-point types do not
+   static_assert(!cuda::is_bitwise_comparable_v<float>);
+   static_assert(!cuda::is_bitwise_comparable_v<double>);

libcudacxx/include/cuda/__type_traits/is_bitwise_comparable.h

@@ -0,0 +1,74 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __CUDA__TYPE_TRAITS_IS_BITWISE_COMPARABLE_H
+#define __CUDA__TYPE_TRAITS_IS_BITWISE_COMPARABLE_H
+
+#include <cuda/std/detail/__config>
+
+#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
+#  pragma GCC system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
+#  pragma clang system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
+#  pragma system_header
+#endif // no system header
+
+#include <cuda/__type_traits/is_vector_type.h>
+#include <cuda/std/__cstddef/types.h>
+#include <cuda/std/__fwd/array.h>
+#include <cuda/std/__fwd/pair.h>
+#include <cuda/std/__fwd/tuple.h>
+#include <cuda/std/__type_traits/has_unique_object_representation.h>
+#include <cuda/std/__type_traits/integral_constant.h>
+#include <cuda/std/__type_traits/is_extended_floating_point.h>
+#include <cuda/std/__type_traits/remove_cv.h>
+
+#include <cuda/std/__cccl/prologue.h>
+
+_CCCL_BEGIN_NAMESPACE_CUDA
+
+//! Users are allowed to specialize this variable template for their own types
+template <typename _Tp>
+constexpr bool is_bitwise_comparable_v =
+  ::cuda::std::has_unique_object_representations_v<::cuda::std::remove_cv_t<_Tp>>
+  && !::cuda::std::__is_extended_floating_point_v<::cuda::std::remove_cv_t<_Tp>>
+#if _CCCL_HAS_CTK()
+  && !::cuda::is_extended_fp_vector_type_v<::cuda::std::remove_cv_t<_Tp>>
+#endif // _CCCL_HAS_CTK()
+  ;
+
+template <typename _Tp>
+constexpr bool is_bitwise_comparable_v<_Tp[]> = is_bitwise_comparable_v<_Tp>;
+
+template <typename _Tp, ::cuda::std::size_t _Size>
+constexpr bool is_bitwise_comparable_v<_Tp[_Size]> = is_bitwise_comparable_v<_Tp>;
+
+template <typename _Tp, ::cuda::std::size_t _Size>
+constexpr bool is_bitwise_comparable_v<::cuda::std::array<_Tp, _Size>> = is_bitwise_comparable_v<_Tp>;
+
+template <typename _T1, typename _T2>
+constexpr bool is_bitwise_comparable_v<::cuda::std::pair<_T1, _T2>> =
+  (sizeof(::cuda::std::pair<_T1, _T2>) == sizeof(_T1) + sizeof(_T2))
+  && is_bitwise_comparable_v<_T1> && is_bitwise_comparable_v<_T2>;
+
+template <typename... _Ts>
+constexpr bool is_bitwise_comparable_v<::cuda::std::tuple<_Ts...>> =
+  (sizeof...(_Ts) == 0 || sizeof(::cuda::std::tuple<_Ts...>) == (sizeof(_Ts) + ... + 0))
+  && (is_bitwise_comparable_v<_Ts> && ...);
+
+// defined as alias so users cannot specialize it (they should specialize the variable template instead)
+template <typename _Tp>
+using is_bitwise_comparable = ::cuda::std::bool_constant<is_bitwise_comparable_v<_Tp>>;
+
+_CCCL_END_NAMESPACE_CUDA
+
+#include <cuda/std/__cccl/epilogue.h>
+
+#endif // __CUDA__TYPE_TRAITS_IS_BITWISE_COMPARABLE_H

libcudacxx/include/cuda/type_traits

@@ -21,6 +21,7 @@
 #  pragma system_header
 #endif // no system header
 
+#include <cuda/__type_traits/is_bitwise_comparable.h>
 #include <cuda/__type_traits/is_floating_point.h>
 #include <cuda/__type_traits/is_vector_type.h>
 #include <cuda/__type_traits/vector_type.h>

libcudacxx/test/libcudacxx/cuda/type_traits/is_bitwise_comparable.pass.cpp

@@ -0,0 +1,131 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#include <cuda/std/array>
+#include <cuda/std/tuple>
+#include <cuda/std/utility>
+#include <cuda/type_traits>
+
+#include "test_macros.h"
+
+template <class T>
+__host__ __device__ void test_is_bitwise_comparable()
+{
+  static_assert(cuda::is_bitwise_comparable<T>::value);
+  static_assert(cuda::is_bitwise_comparable_v<T>);
+  static_assert(cuda::is_bitwise_comparable_v<const T>);
+  static_assert(cuda::is_bitwise_comparable_v<volatile T>);
+  static_assert(cuda::is_bitwise_comparable_v<const volatile T>);
+}
+
+template <class T>
+__host__ __device__ void test_is_not_bitwise_comparable()
+{
+  static_assert(!cuda::is_bitwise_comparable<T>::value);
+  static_assert(!cuda::is_bitwise_comparable_v<T>);
+  static_assert(!cuda::is_bitwise_comparable_v<const T>);
+  static_assert(!cuda::is_bitwise_comparable_v<volatile T>);
+  static_assert(!cuda::is_bitwise_comparable_v<const volatile T>);
+}
+
+struct WithPadding
+{
+  int x;
+  char y;
+};
+
+struct UserSpecialization
+{
+  double value;
+};
+
+template <>
+constexpr bool cuda::is_bitwise_comparable_v<UserSpecialization> = true;
+
+__host__ __device__ void test()
+{
+  // types with unique object representations
+  test_is_bitwise_comparable<int>();
+  test_is_bitwise_comparable<unsigned>();
+  test_is_bitwise_comparable<char>();
+  test_is_bitwise_comparable<unsigned char>();
+  test_is_bitwise_comparable<short>();
+  test_is_bitwise_comparable<long long>();
+
+  // arrays
+  static_assert(cuda::is_bitwise_comparable_v<int[4]>);
+  static_assert(cuda::is_bitwise_comparable_v<const int[4]>);
+  static_assert(cuda::is_bitwise_comparable_v<unsigned char[8]>);
+
+  // padding-free cuda::std::array, pair, tuple of bitwise comparable types
+  static_assert(cuda::is_bitwise_comparable_v<cuda::std::array<int, 4>>);
+  static_assert(cuda::is_bitwise_comparable_v<cuda::std::pair<int, unsigned>>);
+  static_assert(cuda::is_bitwise_comparable_v<cuda::std::tuple<int, unsigned>>);
+  static_assert(cuda::is_bitwise_comparable_v<cuda::std::tuple<>>);
+
+  // composites with padding are not bitwise comparable
+  static_assert(!cuda::is_bitwise_comparable_v<cuda::std::pair<int, char>>);
+  static_assert(!cuda::is_bitwise_comparable_v<cuda::std::tuple<int, unsigned, char>>);
+
+  // types without unique object representations
+  test_is_not_bitwise_comparable<float>();
+  test_is_not_bitwise_comparable<double>();
+  test_is_not_bitwise_comparable<WithPadding>();
+
+  // extended floating-point scalar types
+#if _CCCL_HAS_NVFP16()
+  test_is_not_bitwise_comparable<__half>();
+#endif // _CCCL_HAS_NVFP16()
+#if _CCCL_HAS_NVBF16()
+  test_is_not_bitwise_comparable<__nv_bfloat16>();
+#endif // _CCCL_HAS_NVBF16()
+#if _CCCL_HAS_NVFP8_E4M3()
+  test_is_not_bitwise_comparable<__nv_fp8_e4m3>();
+#endif // _CCCL_HAS_NVFP8_E4M3()
+
+  // extended floating-point vector types
+#if _CCCL_HAS_NVFP16()
+  test_is_not_bitwise_comparable<__half2>();
+#endif // _CCCL_HAS_NVFP16()
+#if _CCCL_HAS_NVBF16()
+  test_is_not_bitwise_comparable<__nv_bfloat162>();
+#endif // _CCCL_HAS_NVBF16()
+#if _CCCL_HAS_NVFP8()
+  test_is_not_bitwise_comparable<__nv_fp8x2_e4m3>();
+#endif // _CCCL_HAS_NVFP8()
+
+  // compositions of extended floating-point types
+#if _CCCL_HAS_NVFP16()
+  static_assert(!cuda::is_bitwise_comparable_v<__half[4]>);
+  static_assert(!cuda::is_bitwise_comparable_v<const __half[4]>);
+  static_assert(!cuda::is_bitwise_comparable_v<cuda::std::array<__half, 4>>);
+  static_assert(!cuda::is_bitwise_comparable_v<cuda::std::pair<__half, int>>);
+  static_assert(!cuda::is_bitwise_comparable_v<cuda::std::tuple<__half, float>>);
+#endif // _CCCL_HAS_NVFP16()
+#if _CCCL_HAS_NVBF16()
+  static_assert(!cuda::is_bitwise_comparable_v<cuda::std::array<__nv_bfloat16, 2>>);
+  static_assert(!cuda::is_bitwise_comparable_v<cuda::std::pair<__nv_bfloat16, int>>);
+#endif // _CCCL_HAS_NVBF16()
+
+  // nested compositions
+#if _CCCL_HAS_NVFP16()
+  static_assert(!cuda::is_bitwise_comparable_v<cuda::std::array<cuda::std::pair<__half, int>, 2>>);
+  static_assert(!cuda::is_bitwise_comparable_v<cuda::std::tuple<cuda::std::array<__half, 4>, int>>);
+#endif // _CCCL_HAS_NVFP16()
+
+  // user specialization of the variable template
+  static_assert(cuda::is_bitwise_comparable_v<UserSpecialization>);
+  static_assert(cuda::is_bitwise_comparable<UserSpecialization>::value);
+}
+
+int main(int, char**)
+{
+  test();
+  return 0;
+}