Refactor cub::UnitWord and cub::Uninitialized (#8119)

Author: bernhardmgruber

cub/cub/util_type.cuh

@@ -386,31 +386,55 @@ using AlignBytes CCCL_DEPRECATED_BECAUSE("Use alignof(T) directly") = detail::Al
 template <typename T>
 struct UnitWord
 {
-  static constexpr auto ALIGN_BYTES = alignof(T);
+  CCCL_DEPRECATED static constexpr auto ALIGN_BYTES = alignof(T);
 
   template <typename Unit>
-  struct IsMultiple
+  struct CCCL_DEPRECATED IsMultiple
   {
     static constexpr auto UNIT_ALIGN_BYTES = alignof(Unit);
     static constexpr bool IS_MULTIPLE =
-      (sizeof(T) % sizeof(Unit) == 0) && (int(ALIGN_BYTES) % int(UNIT_ALIGN_BYTES) == 0);
+      (sizeof(T) % sizeof(Unit) == 0) && (int(alignof(T)) % int(UNIT_ALIGN_BYTES) == 0);
   };
 
+  // MSVC < 19.39 gets an internal compile error on the __is_multiple variable template below, so use a struct instead
+#  if _CCCL_COMPILER(MSVC, <, 19, 39)
+  template <typename Unit>
+  using __is_multiple =
+    ::cuda::std::bool_constant<(sizeof(T) % sizeof(Unit) == 0) && (alignof(T) % alignof(Unit) == 0)>;
+
   /// Largest shuffle word such that sizeof(T) % sizeof(W) == 0 and alignof(W) <= alignof(T)
   using ShuffleWord =
-    ::cuda::std::_If<IsMultiple<int>::IS_MULTIPLE,
+    ::cuda::std::_If<__is_multiple<int>::value,
                      unsigned int,
-                     ::cuda::std::_If<IsMultiple<short>::IS_MULTIPLE, unsigned short, unsigned char>>;
+                     ::cuda::std::_If<__is_multiple<short>::value, unsigned short, unsigned char>>;
 
   /// Largest volatile word such that sizeof(T) % sizeof(W) == 0 and alignof(W) <= alignof(T)
-  using VolatileWord = ::cuda::std::_If<IsMultiple<long long>::IS_MULTIPLE, unsigned long long, ShuffleWord>;
+  using VolatileWord = ::cuda::std::_If<__is_multiple<long long>::value, unsigned long long, ShuffleWord>;
 
   /// Largest memory-access word such that sizeof(T) % sizeof(W) == 0 and alignof(W) <= alignof(T)
-  using DeviceWord = ::cuda::std::_If<IsMultiple<longlong2>::IS_MULTIPLE, ulonglong2, VolatileWord>;
+  using DeviceWord = ::cuda::std::_If<__is_multiple<longlong2>::value, ulonglong2, VolatileWord>;
 
   /// Biggest texture reference word such that sizeof(T) % sizeof(W) == 0 and alignof(W) <= alignof(T)
-  using TextureWord = ::cuda::std::
-    _If<IsMultiple<int4>::IS_MULTIPLE, uint4, ::cuda::std::_If<IsMultiple<int2>::IS_MULTIPLE, uint2, ShuffleWord>>;
+  using TextureWord =
+    ::cuda::std::_If<__is_multiple<int4>::value, uint4, ::cuda::std::_If<__is_multiple<int2>::value, uint2, ShuffleWord>>;
+#  else // _CCCL_COMPILER(MSVC, <, 19, 39)
+  template <typename Unit>
+  static constexpr bool __is_multiple = (sizeof(T) % sizeof(Unit) == 0) && (alignof(T) % alignof(Unit) == 0);
+
+  /// Largest shuffle word evenly dividing T and not increasing the alignment
+  using ShuffleWord = ::cuda::std::
+    _If<__is_multiple<int>, unsigned int, ::cuda::std::_If<__is_multiple<short>, unsigned short, unsigned char>>;
+
+  /// Largest volatile word evenly dividing T and not increasing the alignment
+  using VolatileWord = ::cuda::std::_If<__is_multiple<long long>, unsigned long long, ShuffleWord>;
+
+  /// Largest memory-access word evenly dividing T and not increasing the alignment
+  using DeviceWord = ::cuda::std::_If<__is_multiple<longlong2>, ulonglong2, VolatileWord>;
+
+  /// Biggest texture reference word evenly dividing T and not increasing the alignment
+  using TextureWord =
+    ::cuda::std::_If<__is_multiple<int4>, uint4, ::cuda::std::_If<__is_multiple<int2>, uint2, ShuffleWord>>;
+#  endif // _CCCL_COMPILER(MSVC, <, 19, 39)
 };
 
 // float2 specialization workaround (for SM10-SM13)
@@ -644,18 +668,18 @@ CUB_DEFINE_VECTOR_TYPE(bool,               uchar)
  * Wrapper types
  ******************************************************************************/
 
-/**
- * \brief A storage-backing wrapper that allows types with non-trivial constructors to be aliased in unions
- */
+//! \brief A storage-backing wrapper that allows types with non-trivial constructors to be aliased in unions. Has the
+//! same size as T.
 template <typename T>
 struct Uninitialized
 {
-  /// Largest memory-access word such that sizeof(T) % sizeof(W) == 0 and alignof(W) <= alignof(T)
+  /// Largest memory-access word evenly dividing T and not increasing the alignment
   using DeviceWord = typename UnitWord<T>::DeviceWord;
 
   static constexpr ::cuda::std::size_t DATA_SIZE = sizeof(T);
   static constexpr ::cuda::std::size_t WORD_SIZE = sizeof(DeviceWord);
   static constexpr ::cuda::std::size_t WORDS     = DATA_SIZE / WORD_SIZE;
+  static_assert(DATA_SIZE % WORDS == 0);
 
   /// Backing storage
   DeviceWord storage[WORDS];