[SYCL] Add sycl cpp header for imf libdevice API (#6487)

We added some imf(intel math) APIs in libdevice: #5999, #6327, those APIs are C interface which means we don't expect users to call them directly in SYCL kernel since SYCL programming uses cpp. This PR will provide cpp wrapper for those imf APIs which users can call directly in source, the approach is similar to STL , std math functions will call corresponding C math functions or compiler math builtins. The user scenario should be following:
cpp
#include <sycl/ext/intel/math.hpp>
...
float x, y, z;
z = sycl::ext::intel::math::copysign(x, y);
...
``

Author: jinge90

sycl/include/sycl/ext/intel/math.hpp

@@ -0,0 +1,68 @@
+//==------------- math.hpp - Intel specific math API -----------------------==//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+// The main header of Intel specific math API
+//===----------------------------------------------------------------------===//
+
+#pragma once
+#include <sycl/half_type.hpp>
+#include <type_traits>
+
+// _iml_half_internal is internal representation for fp16 type used in intel
+// math device library. The definition here should align with definition in
+// https://github.com/intel/llvm/blob/sycl/libdevice/imf_half.hpp
+#if defined(__SPIR__)
+using _iml_half_internal = _Float16;
+#else
+using _iml_half_internal = uint16_t;
+#endif
+
+extern "C" {
+float __imf_saturatef(float);
+float __imf_copysignf(float, float);
+double __imf_copysign(double, double);
+_iml_half_internal __imf_copysignf16(_iml_half_internal, _iml_half_internal);
+};
+
+namespace sycl {
+__SYCL_INLINE_VER_NAMESPACE(_V1) {
+namespace ext {
+namespace intel {
+namespace math {
+
+#if __cplusplus >= 201703L
+template <typename Tp>
+std::enable_if_t<std::is_same_v<Tp, float>, float> saturate(Tp x) {
+  return __imf_saturatef(x);
+}
+
+template <typename Tp>
+std::enable_if_t<std::is_same_v<Tp, float>, float> copysign(Tp x, Tp y) {
+  return __imf_copysignf(x, y);
+}
+
+template <typename Tp>
+std::enable_if_t<std::is_same_v<Tp, double>, double> copysign(Tp x, Tp y) {
+  return __imf_copysign(x, y);
+}
+
+template <typename Tp>
+std::enable_if_t<std::is_same_v<Tp, sycl::half>, sycl::half> copysign(Tp x,
+                                                                      Tp y) {
+  static_assert(sizeof(sycl::half) == sizeof(_iml_half_internal),
+                "sycl::half is not compatible with _iml_half_internal.");
+  _iml_half_internal xi = __builtin_bit_cast(_iml_half_internal, x);
+  _iml_half_internal yi = __builtin_bit_cast(_iml_half_internal, y);
+  return __builtin_bit_cast(sycl::half, __imf_copysignf16(xi, yi));
+}
+
+#endif
+} // namespace math
+} // namespace intel
+} // namespace ext
+} // __SYCL_INLINE_VER_NAMESPACE(_V1)
+} // namespace sycl