[SYCL] Align nan builtin with the updated specification (#12182)

aelovikov-intel · web-flow · commit e95bde0e23eb · 2024-01-11T08:35:18.000-08:00
Specification has been changed in KhronosGroup/SYCL-Docs#519.
diff --git a/sycl/include/sycl/builtins_legacy_marray_vec.hpp b/sycl/include/sycl/builtins_legacy_marray_vec.hpp
@@ -412,14 +412,28 @@ __SYCL_MARRAY_MATH_FUNCTION_REMQUO_OVERLOAD(remquo, x[j], y[j])
 #undef __SYCL_MARRAY_MATH_FUNCTION_W_GENPTR_ARG_OVERLOAD_IMPL
 
 template <typename T, size_t N>
-std::enable_if_t<detail::is_nan_type_v<T>, marray<detail::nan_return_t<T>, N>>
+std::enable_if_t<detail::is_nan_type_v<T> &&
+                     detail::is_non_deprecated_nan_type_v<T>,
+                 marray<detail::nan_return_t<T>, N>>
 nan(marray<T, N> nancode) {
   marray<detail::nan_return_t<T>, N> res;
   for (int j = 0; j < N; j++) {
     res[j] = nan(nancode[j]);
   }
   return res;
 }
+template <typename T, size_t N>
+__SYCL_DEPRECATED(
+    "This is a deprecated argument type for SYCL nan built-in function.")
+std::enable_if_t<detail::is_nan_type_v<T> &&
+                     !detail::is_non_deprecated_nan_type_v<T>,
+                 marray<detail::nan_return_t<T>, N>> nan(marray<T, N> nancode) {
+  marray<detail::nan_return_t<T>, N> res;
+  for (int j = 0; j < N; j++) {
+    res[j] = nan(nancode[j]);
+  }
+  return res;
+}
 
 /* --------------- 4.13.5 Common functions. ---------------------------------*/
 // vgenfloath clamp (vgenfloath x, half minval, half maxval)
diff --git a/sycl/include/sycl/builtins_legacy_scalar.hpp b/sycl/include/sycl/builtins_legacy_scalar.hpp
@@ -322,8 +322,30 @@ modf(T x, T2 iptr) {
   return __sycl_std::__invoke_modf<T>(x, iptr);
 }
 
-template <typename T, typename = std::enable_if_t<detail::is_nan_type_v<T>, T>>
-detail::nan_return_t<T> nan(T nancode) {
+namespace detail {
+// SYCL 2020, revision 9 modifies accepted overloads to scalar/vec/marray of
+// uint16_t/uint32_t/uint64_t.
+template <typename T>
+inline constexpr bool is_non_deprecated_nan_type_v =
+    std::is_same_v<get_elem_type_t<T>, uint16_t> ||
+    std::is_same_v<get_elem_type_t<T>, uint32_t> ||
+    std::is_same_v<get_elem_type_t<T>, uint64_t>;
+} // namespace detail
+
+template <typename T>
+std::enable_if_t<detail::is_nan_type_v<T> &&
+                     detail::is_non_deprecated_nan_type_v<T>,
+                 detail::nan_return_t<T>>
+nan(T nancode) {
+  return __sycl_std::__invoke_nan<detail::nan_return_t<T>>(
+      detail::convert_data_type<T, detail::nan_argument_base_t<T>>()(nancode));
+}
+template <typename T>
+__SYCL_DEPRECATED(
+    "This is a deprecated argument type for SYCL nan built-in function.")
+std::enable_if_t<detail::is_nan_type_v<T> &&
+                     !detail::is_non_deprecated_nan_type_v<T>,
+                 detail::nan_return_t<T>> nan(T nancode) {
   return __sycl_std::__invoke_nan<detail::nan_return_t<T>>(
       detail::convert_data_type<T, detail::nan_argument_base_t<T>>()(nancode));
 }
diff --git a/sycl/include/sycl/builtins_utils_scalar.hpp b/sycl/include/sycl/builtins_utils_scalar.hpp
@@ -206,7 +206,6 @@ template <typename T> struct nan_return_unswizzled {
 
 template <typename T>
 using nan_return_unswizzled_t = typename nan_return_unswizzled<T>::type;
-
 } // namespace detail
 } // namespace _V1
 } // namespace sycl
diff --git a/sycl/source/builtins_generator.py b/sycl/source/builtins_generator.py
@@ -317,12 +317,18 @@ def __str__(self):
     Vec(["uint64_t", "unsigned long long"])
 ]  # Fundamental integer types non-standard. Deprecated.
 
+# vuint<N>n name is taken, use "_fixed" suffix.
+vuint16n_fixed = [Vec(["uint16_t"])]
+vuint32n_fixed = [Vec(["uint32_t"])]
+vuint64n_fixed = [Vec(["uint64_t"])]
+
 mint8n = [Marray(["int8_t"])]
 mint16n = [Marray(["int16_t"])]
 mint32n = [Marray(["int32_t"])]
 muint8n = [Marray(["uint8_t"])]
 muint16n = [Marray(["uint16_t"])]
 muint32n = [Marray(["uint32_t"])]
+muint64n = [Marray(["uint64_t"])]
 mintn = [Marray(["int"])]
 mushortn = [Marray(["unsigned short"])]
 muintn = [Marray(["unsigned int"])]
@@ -594,12 +600,16 @@ def __str__(self):
     "vuint16n": vuint16n,
     "vuint32n": vuint32n,
     "vuint64n": vuint64n,
+    "vuint16n_fixed": vuint16n_fixed,
+    "vuint32n_fixed": vuint32n_fixed,
+    "vuint64n_fixed": vuint64n_fixed,
     "mint8n": mint8n,
     "mint16n": mint16n,
     "mint32n": mint32n,
     "muint8n": muint8n,
     "muint16n": muint16n,
     "muint32n": muint32n,
+    "muint64n": muint64n,
     "mintn": mintn,
     "mushortn": mushortn,
     "muintn": muintn,
@@ -1228,16 +1238,15 @@ def custom_nan_invoke(return_type, arg_types, arg_names):
         Def("half", ["half", "halfptr"]),
     ],
     "nan": [
-        Def("nanreturn0", ["vuint32n"], custom_invoke=custom_nan_invoke),
-        Def("nanreturn0", ["muintn"], marray_use_loop=True),
-        Def("nanreturn0", ["unsigned int"], custom_invoke=custom_nan_invoke),
-        Def("nanreturn0", ["vuint64n"], custom_invoke=custom_nan_invoke),
-        Def("nanreturn0", ["mulongn"], marray_use_loop=True),
-        Def("nanreturn0", ["unsigned long"], custom_invoke=custom_nan_invoke),
-        Def("nanreturn0", ["unsigned long long"], custom_invoke=custom_nan_invoke),
-        Def("nanreturn0", ["vuint16n"], custom_invoke=custom_nan_invoke),
-        Def("nanreturn0", ["mushortn"], marray_use_loop=True),
-        Def("nanreturn0", ["unsigned short"], custom_invoke=custom_nan_invoke),
+        Def("nanreturn0", ["uint16_t"], custom_invoke=custom_nan_invoke),
+        Def("nanreturn0", ["vuint16n_fixed"], custom_invoke=custom_nan_invoke),
+        Def("nanreturn0", ["muint16n"], marray_use_loop=True),
+        Def("nanreturn0", ["uint32_t"], custom_invoke=custom_nan_invoke),
+        Def("nanreturn0", ["vuint32n_fixed"], custom_invoke=custom_nan_invoke),
+        Def("nanreturn0", ["muint32n"], marray_use_loop=True),
+        Def("nanreturn0", ["uint64_t"], custom_invoke=custom_nan_invoke),
+        Def("nanreturn0", ["vuint64n_fixed"], custom_invoke=custom_nan_invoke),
+        Def("nanreturn0", ["muint64n"], marray_use_loop=True),
     ],
     "nextafter": [Def("genfloat", ["genfloat", "genfloat"])],
     "pow": [Def("genfloat", ["genfloat", "genfloat"])],
diff --git a/sycl/test-e2e/Basic/built-ins/marray_math.cpp b/sycl/test-e2e/Basic/built-ins/marray_math.cpp
@@ -136,8 +136,13 @@ int main() {
   sycl::marray<int, 3> ma4{1, 1, 1};
   sycl::marray<float, 3> ma5{180, -180, -180};
   sycl::marray<float, 3> ma6{1.4f, 4.2f, 5.3f};
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
   sycl::marray<unsigned int, 3> ma7{1, 2, 3};
   sycl::marray<unsigned long int, 3> ma8{1, 2, 3};
+#else
+  sycl::marray<uint32_t, 3> ma7{1, 2, 3};
+  sycl::marray<uint64_t, 3> ma8{1, 2, 3};
+#endif
 
   TEST(sycl::fabs, float, 3, EXPECTED(float, 180, 180, 180), 0, ma5);
   TEST(sycl::ilogb, int, 3, EXPECTED(int, 7, 7, 7), 0, ma3);
diff --git a/sycl/test-e2e/DeviceLib/built-ins/nan.cpp b/sycl/test-e2e/DeviceLib/built-ins/nan.cpp
@@ -48,6 +48,7 @@ template <typename T, typename R> void check_vec_nan(s::queue &Queue) {
 }
 
 int main() {
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
   test_nan_call<unsigned short, s::half>();
   test_nan_call<unsigned int, float>();
   test_nan_call<unsigned long, double>();
@@ -56,6 +57,14 @@ int main() {
   test_nan_call<s::uint2, s::float2>();
   test_nan_call<s::ulong2, s::double2>();
   test_nan_call<s::vec<unsigned long long, 2>, s::double2>();
+#else
+  test_nan_call<uint16_t, s::half>();
+  test_nan_call<uint32_t, float>();
+  test_nan_call<uint64_t, double>();
+  test_nan_call<s::vec<uint16_t, 2>, s::half2>();
+  test_nan_call<s::vec<uint32_t, 2>, s::float2>();
+  test_nan_call<s::vec<uint64_t, 2>, s::double2>();
+#endif
 
   s::queue Queue([](sycl::exception_list ExceptionList) {
     for (std::exception_ptr ExceptionPtr : ExceptionList) {
@@ -70,17 +79,33 @@ int main() {
   });
 
   if (Queue.get_device().has(sycl::aspect::fp16)) {
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
     check_scalar_nan<unsigned short, s::half>(Queue);
     check_vec_nan<uint16_t, s::half>(Queue);
+#else
+    check_scalar_nan<uint16_t, s::half>(Queue);
+    check_vec_nan<uint16_t, s::half>(Queue);
+#endif
   }
 
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
   check_scalar_nan<unsigned int, float>(Queue);
   check_vec_nan<uint32_t, float>(Queue);
+#else
+  check_scalar_nan<uint32_t, float>(Queue);
+  check_vec_nan<uint32_t, float>(Queue);
+#endif
+
   if (Queue.get_device().has(sycl::aspect::fp64)) {
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
     check_scalar_nan<unsigned long, double>(Queue);
     check_scalar_nan<unsigned long long, double>(Queue);
     check_vec_nan<uint64_t, double>(Queue);
     check_vec_nan<unsigned long long, double>(Queue);
+#else
+    check_scalar_nan<uint64_t, double>(Queue);
+    check_vec_nan<uint64_t, double>(Queue);
+#endif
   }
   return 0;
 }
diff --git a/sycl/test/basic_tests/builtins_implicitly_convertible_args.cpp b/sycl/test/basic_tests/builtins_implicitly_convertible_args.cpp
@@ -43,9 +43,9 @@ template <typename T> struct ImplicitlyConvertibleType {
   CHECK_INNER(NUM_ARGS, FUNC_NAME, double)
 
 #define UGENINT_NAN_CHECK(NUM_ARGS, FUNC_NAME)                                 \
-  CHECK_INNER(NUM_ARGS, FUNC_NAME, unsigned int)                               \
-  CHECK_INNER(NUM_ARGS, FUNC_NAME, unsigned short)                             \
-  CHECK_INNER(NUM_ARGS, FUNC_NAME, unsigned long)
+  CHECK_INNER(NUM_ARGS, FUNC_NAME, uint32_t)                                   \
+  CHECK_INNER(NUM_ARGS, FUNC_NAME, uint16_t)                                   \
+  CHECK_INNER(NUM_ARGS, FUNC_NAME, uint64_t)
 
 void check() {
   GENFLOAT_CHECK(ONE_ARG, acos)
diff --git a/sycl/test/basic_tests/builtins_templates.cpp b/sycl/test/basic_tests/builtins_templates.cpp
@@ -45,9 +45,9 @@
   CHECK_INNER(NUM_ARGS, FUNC_NAME, sycl::vec<uint32_t, 4>)                     \
   CHECK_INNER(NUM_ARGS, FUNC_NAME, sycl::vec<uint16_t, 4>)                     \
   CHECK_INNER(NUM_ARGS, FUNC_NAME, sycl::vec<uint64_t, 4>)                     \
-  CHECK_INNER(NUM_ARGS, FUNC_NAME, sycl::marray<unsigned int, 4>)              \
-  CHECK_INNER(NUM_ARGS, FUNC_NAME, sycl::marray<unsigned short, 4>)            \
-  CHECK_INNER(NUM_ARGS, FUNC_NAME, sycl::marray<unsigned long, 4>)
+  CHECK_INNER(NUM_ARGS, FUNC_NAME, sycl::marray<uint32_t, 4>)                  \
+  CHECK_INNER(NUM_ARGS, FUNC_NAME, sycl::marray<uint16_t, 4>)                  \
+  CHECK_INNER(NUM_ARGS, FUNC_NAME, sycl::marray<uint64_t, 4>)
 
 #define SGENINT_CHECK(NUM_ARGS, FUNC_NAME)                                     \
   CHECK_INNER(NUM_ARGS, FUNC_NAME, signed char)                                \
diff --git a/sycl/test/warnings/sycl_2020_deprecations.cpp b/sycl/test/warnings/sycl_2020_deprecations.cpp
@@ -454,5 +454,14 @@ int main() {
   std::ignore =
       sycl::fast_distance(sycl::vec<double, 2>{0.0}, sycl::vec<double, 2>{1.0});
 
+  // clang-format off
+  // SYCL 2020, revision 9 uses fixed-width integer type, one of these has to be
+  // deprecated.
+  // expected-warning-re@+1 {{'nan<{{.*}}>' is deprecated: This is a deprecated argument type for SYCL nan built-in function.}}
+  std::ignore = (sycl::nan((unsigned long){0}), sycl::nan((unsigned long long){0}));
+  // expected-warning-re@+1 {{'nan<sycl::vec<{{.*}}, 2>>' is deprecated: This is a deprecated argument type for SYCL nan built-in function.}}
+  std::ignore = (sycl::nan(sycl::vec<unsigned long, 2>{0}), sycl::nan(sycl::vec<unsigned long long, 2>{0}));
+  // clang-format on
+
   return 0;
 }
