[SYCL] Remove detail::make_larger (#15667)

We only used that for `sycl::mad_sat`/`sycl::mul_hi` host
implementations. We can just perform the computation on `[u]int64_t` and
let the compiler optimize it.

sycl/include/sycl/builtins_utils_scalar.hpp

@@ -8,14 +8,14 @@
 
 #pragma once
 
-#include <sycl/access/access.hpp>              // for address_space, decorated
-#include <sycl/aliases.hpp>                    // for half
-#include <sycl/detail/defines_elementary.hpp>  // for __SYCL_ALWAYS_INLINE
-#include <sycl/detail/generic_type_traits.hpp> // for is_svgenfloat, is_sge...
-#include <sycl/detail/type_list.hpp>           // for is_contained, type_list
-#include <sycl/detail/type_traits.hpp>         // for make_larger_t, marray...
-#include <sycl/half_type.hpp>                  // for half, intel
-#include <sycl/multi_ptr.hpp>                  // for address_space_cast
+#include <sycl/access/access.hpp>
+#include <sycl/aliases.hpp>
+#include <sycl/detail/defines_elementary.hpp>
+#include <sycl/detail/generic_type_traits.hpp>
+#include <sycl/detail/type_list.hpp>
+#include <sycl/detail/type_traits.hpp>
+#include <sycl/half_type.hpp>
+#include <sycl/multi_ptr.hpp>
 
 #include <algorithm>
 #include <cstring>

sycl/include/sycl/detail/type_list.hpp

@@ -69,10 +69,6 @@ using is_one_of_spaces =
 template <typename T1, typename T2>
 struct is_type_size_equal : std::bool_constant<(sizeof(T1) == sizeof(T2))> {};
 
-template <typename T1, typename T2>
-struct is_type_size_double_of
-    : std::bool_constant<(sizeof(T1) == (sizeof(T2) * 2))> {};
-
 // find required type
 template <typename TypeList, template <typename, typename> class Comp,
           typename T>
@@ -90,10 +86,6 @@ using find_type_t = typename find_type<TypeList, Comp, T>::type;
 template <typename TypeList, typename T>
 using find_same_size_type_t = find_type_t<TypeList, is_type_size_equal, T>;
 
-template <typename TypeList, typename T>
-using find_twice_as_large_type_t =
-    find_type_t<TypeList, is_type_size_double_of, T>;
-
 } // namespace detail
 } // namespace _V1
 } // namespace sycl

sycl/include/sycl/detail/type_traits.hpp

@@ -399,53 +399,6 @@ template <typename T, int N, typename TL> struct make_type_impl<vec<T, N>, TL> {
 template <typename T, typename TL>
 using make_type_t = typename make_type_impl<T, TL>::type;
 
-// make_larger_t
-template <typename T, typename Enable = void> struct make_larger_impl;
-template <typename T>
-struct make_larger_impl<
-    T, std::enable_if_t<is_contained<T, gtl::scalar_floating_list>::value, T>> {
-  using type = find_twice_as_large_type_t<gtl::scalar_floating_list, T>;
-};
-
-template <typename T>
-struct make_larger_impl<
-    T, std::enable_if_t<is_contained<T, gtl::scalar_signed_integer_list>::value,
-                        T>> {
-  using type = find_twice_as_large_type_t<gtl::scalar_signed_integer_list, T>;
-};
-
-template <typename T>
-struct make_larger_impl<
-    T, std::enable_if_t<
-           is_contained<T, gtl::scalar_unsigned_integer_list>::value, T>> {
-  using type = find_twice_as_large_type_t<gtl::scalar_unsigned_integer_list, T>;
-};
-
-template <typename T, int N> struct make_larger_impl<vec<T, N>, vec<T, N>> {
-  using base_type = vector_element_t<vec<T, N>>;
-  using upper_type = typename make_larger_impl<base_type, base_type>::type;
-  using new_type = vec<upper_type, N>;
-  static constexpr bool found = !std::is_same_v<upper_type, void>;
-  using type = std::conditional_t<found, new_type, void>;
-};
-
-template <typename T, size_t N>
-struct make_larger_impl<marray<T, N>, marray<T, N>> {
-  using base_type = marray_element_t<marray<T, N>>;
-  using upper_type = typename make_larger_impl<base_type, base_type>::type;
-  using new_type = marray<upper_type, N>;
-  static constexpr bool found = !std::is_same_v<upper_type, void>;
-  using type = std::conditional_t<found, new_type, void>;
-};
-
-// TODO: this type trait is not used anyweher in SYCL headers and it should
-// be moved to the library code
-template <typename T> struct make_larger {
-  using type = typename make_larger_impl<T, T>::type;
-};
-
-template <typename T> using make_larger_t = typename make_larger<T>::type;
-
 #if defined(RESTRICT_WRITE_ACCESS_TO_CONSTANT_PTR)
 template <access::address_space AS, class DataT>
 using const_if_const_AS =

sycl/source/builtins/integer_functions.cpp

@@ -150,7 +150,7 @@ BUILTIN_GENINT_SU(
                                                   // negative value.
           return sycl::add_sat(T(a * b), c);
         } else {
-          using UPT = sycl::detail::make_larger_t<T>;
+          using UPT = int64_t;
           UPT mul = UPT(a) * UPT(b);
           UPT res = mul + UPT(c);
           const UPT max = std::numeric_limits<T>::max();
@@ -166,7 +166,7 @@ BUILTIN_GENINT_SU(
             return std::numeric_limits<T>::max();
           return sycl::add_sat(T(a * b), c);
         } else {
-          using UPT = sycl::detail::make_larger_t<T>;
+          using UPT = uint64_t;
           UPT mul = UPT(a) * UPT(b);
           const UPT min = std::numeric_limits<T>::min();
           const UPT max = std::numeric_limits<T>::max();
@@ -176,6 +176,11 @@ BUILTIN_GENINT_SU(
       }
     })
 
+// Moved outside of macro invocation:
+template <typename T>
+using same_signedness_int64_t =
+    std::conditional_t<std::is_signed_v<T>, int64_t, uint64_t>;
+
 BUILTIN_GENINT_SU(TWO_ARGS, mul_hi, [](auto a, auto b) -> decltype(a) {
   using T = decltype(a);
   if constexpr (sizeof(T) == 8) {
@@ -184,7 +189,7 @@ BUILTIN_GENINT_SU(TWO_ARGS, mul_hi, [](auto a, auto b) -> decltype(a) {
     else
       return __u_long_mul_hi(a, b);
   } else {
-    using UPT = sycl::detail::make_larger_t<T>;
+    using UPT = same_signedness_int64_t<T>;
     UPT a_s = a;
     UPT b_s = b;
     UPT mul = a_s * b_s;

sycl/test/basic_tests/generic_type_traits.cpp

@@ -78,7 +78,6 @@ int main() {
   float_point_to_sign_integeral
 
   make_unsigned
-  make_larger
   */
 
   // checks for some type conversions.

sycl/test/type_traits/type_list.cpp

@@ -75,21 +75,12 @@ int main() {
 
   test_predicate<d::is_type_size_equal, int8_t, s::opencl::cl_char>();
   test_predicate<d::is_type_size_equal, int16_t, s::opencl::cl_char, false>();
-  test_predicate<d::is_type_size_double_of, int8_t, s::opencl::cl_char,
-                 false>();
-  test_predicate<d::is_type_size_double_of, int16_t, s::opencl::cl_char>();
-  test_predicate<d::is_type_size_double_of, int32_t, s::opencl::cl_char,
-                 false>();
 
   // if void is found, the required type is not found
   test_trait<d::find_same_size_type_t, d::type_list<int8_t>, s::opencl::cl_char,
              int8_t>();
   test_trait<d::find_same_size_type_t, d::type_list<int16_t>,
              s::opencl::cl_char, void>();
-  test_trait<d::find_twice_as_large_type_t, d::type_list<int8_t, int16_t>,
-             s::opencl::cl_char, int16_t>();
-  test_trait<d::find_twice_as_large_type_t, d::type_list<int8_t, int32_t>,
-             s::opencl::cl_char, void>();
 
   return 0;
 }

sycl/test/type_traits/type_traits.cpp

@@ -34,11 +34,6 @@ void test_make_type_t() {
                 "");
 }
 
-template <typename T, typename CheckedT, bool Expected = true>
-void test_make_larger_t() {
-  static_assert(is_same<d::make_larger_t<T>, CheckedT>::value == Expected, "");
-}
-
 template <typename T, typename T2, typename CheckedT, bool Expected = true>
 void test_change_base_type_t() {
   static_assert(
@@ -116,17 +111,6 @@ int main() {
   test_make_type_t<s::vec<s::opencl::cl_int, 3>, d::gtl::scalar_float_list,
                    s::vec<s::opencl::cl_float, 3>>();
 
-  test_make_larger_t<s::half, float>();
-  test_make_larger_t<s::half3, s::float3>();
-  test_make_larger_t<float, double>();
-  test_make_larger_t<s::float3, s::double3>();
-  test_make_larger_t<double, void>();
-  test_make_larger_t<s::double3, void>();
-  test_make_larger_t<int32_t, int64_t>();
-  test_make_larger_t<s::vec<int32_t, 8>, s::vec<int64_t, 8>>();
-  test_make_larger_t<int64_t, void>();
-  test_make_larger_t<s::vec<int64_t, 8>, void>();
-
   test_change_base_type_t<int, float, float>();
   test_change_base_type_t<s::int2, float, s::float2>();
   test_change_base_type_t<long, float, float>();