[SYCL] Add sycl cpp header for imf libdevice API

[jinge90]

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);
...
``
Thanks very much.
Signed-off-by: jinge90 [email protected]

[xtian-github]

LGTM

[jinge90]

Hi, @pvchupin
could you take a look at this PR?

Thanks very much.

[zettai-reido]

LGTM.

[akolesov-nv]

Looks good to me. Is this file meant to be edited manually for every new function or any script exist? Which team should own (and update) the file? Thanks!

[jinge90]

Hi, @intel/llvm-reviewers-runtime and @cperkinsintel
Could you help review this patch? We need it ASAP.

Thanks very much.

[steffenlarsen]

Since you went with SFINAE for saturate, maybe to keep it consistent it would make sense to do it for copysign too? That is, it would be something like:

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));
}

That should at least make the misuse of it give a similar error to that of saturate.

[jinge90]

Hi, @steffenlarsen
Your suggestion does make sense, already updated patch.

Thanks very much.

[zettai-reido]

New changes are fine.

[xtian-github]

LGTM

[steffenlarsen]

LGTM!

[steffenlarsen]

@jinge90 - Are there tests for this?

[jinge90]

Hi, @pvchupin
Could you help merge this PR?
Thanks very much.

[jinge90]

@jinge90 - Are there tests for this?

We will extend imf libdevice imf tests: https://github.com/intel/llvm-test-suite/blob/intel/SYCL/DeviceLib/imf_fp32_test.cpp
https://github.com/intel/llvm-test-suite/blob/intel/SYCL/DeviceLib/imf_fp64_test.cpp

Currently, we directly call "_imf*" functions in SYCL kernel in those tests, we add function call to corresponding cpp wrappers.

Thanks very much.