[NFCI][SYCL] Introduce convertFromOpenCLTypeFor<T>(x) helper (intel#1…

…2684)

Continuation of refactoring series initiated in
intel#12674.

sycl/include/sycl/detail/builtins.hpp

@@ -9,7 +9,7 @@
 #pragma once
 
 #include <sycl/detail/defines_elementary.hpp> // for __SYCL_ALWAYS_INLINE
-#include <sycl/detail/generic_type_traits.hpp> // for convertDataToType
+#include <sycl/detail/generic_type_traits.hpp> // for to/from OpenCLType converts
 #include <utility>
 
 // TODO Decide whether to mark functions with this attribute.
@@ -18,27 +18,31 @@
 #ifdef __SYCL_DEVICE_ONLY__
 #define __FUNC_PREFIX_OCL __spirv_ocl_
 #define __FUNC_PREFIX_CORE __spirv_
-#define __SYCL_EXTERN_IT1(Ret, prefix, call, arg1)
-#define __SYCL_EXTERN_IT2(Ret, prefix, call, arg1, arg2)
-#define __SYCL_EXTERN_IT2_SAME(Ret, prefix, call, arg)
-#define __SYCL_EXTERN_IT3(Ret, prefix, call, Arg1, Arg2, Arg3)
+#define __SYCL_EXTERN_IT1(R, prefix, call, arg1)
+#define __SYCL_EXTERN_IT2(R, prefix, call, arg1, arg2)
+#define __SYCL_EXTERN_IT2_SAME(R, prefix, call, arg)
+#define __SYCL_EXTERN_IT3(R, prefix, call, Arg1, Arg2, Arg3)
 #else
 #define __FUNC_PREFIX_OCL sycl_host_
 #define __FUNC_PREFIX_CORE sycl_host_
-#define __SYCL_EXTERN_IT1(Ret, prefix, call, arg) \
+#define __SYCL_EXTERN_IT1(R, prefix, call, arg)  \
  using Arg = decltype(arg); \
+ using Ret = sycl::detail::ConvertToOpenCLType_t<R>; \
  extern Ret __SYCL_PPCAT(prefix, call)(Arg)
-#define __SYCL_EXTERN_IT2_SAME(Ret, prefix, call, arg) \
+#define __SYCL_EXTERN_IT2_SAME(R, prefix, call, arg)  \
  using Arg = decltype(arg); \
+ using Ret = sycl::detail::ConvertToOpenCLType_t<R>; \
  extern Ret __SYCL_PPCAT(prefix, call)(Arg, Arg)
-#define __SYCL_EXTERN_IT2(Ret, prefix, call, arg1, arg2) \
+#define __SYCL_EXTERN_IT2(R, prefix, call, arg1, arg2)  \
  using Arg1 = decltype(arg1); \
  using Arg2 = decltype(arg2); \
+ using Ret = sycl::detail::ConvertToOpenCLType_t<R>; \
  extern Ret __SYCL_PPCAT(prefix, call)(Arg1, Arg2)
-#define __SYCL_EXTERN_IT3(Ret, prefix, call, arg1, arg2, arg3) \
+#define __SYCL_EXTERN_IT3(R, prefix, call, arg1, arg2, arg3)  \
  using Arg1 = decltype(arg1); \
  using Arg2 = decltype(arg2); \
  using Arg3 = decltype(arg3); \
+ using Ret = sycl::detail::ConvertToOpenCLType_t<R>; \
  extern Ret __SYCL_PPCAT(prefix, call)(Arg1, Arg2, Arg3)
 #endif
 
@@ -49,43 +53,39 @@
  template <typename R, typename T1> \
  inline __SYCL_ALWAYS_INLINE R __invoke_##call(T1 t1) __NOEXC { \
  auto arg1 = sycl::detail::convertToOpenCLType(std::move(t1)); \
- using Ret = sycl::detail::ConvertToOpenCLType_t<R>; \
- __SYCL_EXTERN_IT1(Ret, prefix, call, arg1); \
- Ret ret = __SYCL_PPCAT(prefix, call)(std::move(arg1)); \
- return sycl::detail::convertDataToType<Ret, R>(std::move(ret)); \
+ __SYCL_EXTERN_IT1(R, prefix, call, arg1); \
+ return sycl::detail::convertFromOpenCLTypeFor<R>( \
+ __SYCL_PPCAT(prefix, call)(std::move(arg1))); \
  }
 
 #define __SYCL_MAKE_CALL_ARG2(call, prefix) \
  template <typename R, typename T1, typename T2> \
  inline __SYCL_ALWAYS_INLINE R __invoke_##call(T1 t1, T2 t2) __NOEXC { \
  auto arg1 = sycl::detail::convertToOpenCLType(std::move(t1)); \
  auto arg2 = sycl::detail::convertToOpenCLType(std::move(t2)); \
- using Ret = sycl::detail::ConvertToOpenCLType_t<R>; \
- __SYCL_EXTERN_IT2(Ret, prefix, call, arg1, arg2); \
- Ret ret = __SYCL_PPCAT(prefix, call)(std::move(arg1), std::move(arg2)); \
- return sycl::detail::convertDataToType<Ret, R>(std::move(ret)); \
+ __SYCL_EXTERN_IT2(R, prefix, call, arg1, arg2); \
+ return sycl::detail::convertFromOpenCLTypeFor<R>( \
+ __SYCL_PPCAT(prefix, call)(std::move(arg1), std::move(arg2))); \
  }
 
 #define __SYCL_MAKE_CALL_ARG2_SAME(call, prefix) \
  template <typename R, typename T> \
  inline __SYCL_ALWAYS_INLINE R __invoke_##call(T t1, T t2) __NOEXC { \
  auto arg1 = sycl::detail::convertToOpenCLType(std::move(t1)); \
  auto arg2 = sycl::detail::convertToOpenCLType(std::move(t2)); \
- using Ret = sycl::detail::ConvertToOpenCLType_t<R>; \
- __SYCL_EXTERN_IT2_SAME(Ret, prefix, call, arg1); \
- Ret ret = __SYCL_PPCAT(prefix, call)(std::move(arg1), std::move(arg2)); \
- return sycl::detail::convertDataToType<Ret, R>(std::move(ret)); \
+ __SYCL_EXTERN_IT2_SAME(R, prefix, call, arg1); \
+ return sycl::detail::convertFromOpenCLTypeFor<R>( \
+ __SYCL_PPCAT(prefix, call)(std::move(arg1), std::move(arg2))); \
  }
 
 #define __SYCL_MAKE_CALL_ARG2_SAME_RESULT(call, prefix) \
  template <typename T> \
  inline __SYCL_ALWAYS_INLINE T __invoke_##call(T v1, T v2) __NOEXC { \
  auto arg1 = sycl::detail::convertToOpenCLType(std::move(v1)); \
  auto arg2 = sycl::detail::convertToOpenCLType(std::move(v2)); \
- using Type = decltype(arg1); \
- __SYCL_EXTERN_IT2_SAME(Type, prefix, call, arg1); \
- Type ret = __SYCL_PPCAT(prefix, call)(std::move(arg1), std::move(arg2)); \
- return sycl::detail::convertDataToType<Type, T>(std::move(ret)); \
+ __SYCL_EXTERN_IT2_SAME(T, prefix, call, arg1); \
+ return sycl::detail::convertFromOpenCLTypeFor<T>( \
+ __SYCL_PPCAT(prefix, call)(std::move(arg1), std::move(arg2))); \
  }
 
 #define __SYCL_MAKE_CALL_ARG3(call, prefix) \
@@ -94,11 +94,9 @@
  auto arg1 = sycl::detail::convertToOpenCLType(std::move(t1)); \
  auto arg2 = sycl::detail::convertToOpenCLType(std::move(t2)); \
  auto arg3 = sycl::detail::convertToOpenCLType(std::move(t3)); \
- using Ret = sycl::detail::ConvertToOpenCLType_t<R>; \
- __SYCL_EXTERN_IT3(Ret, prefix, call, arg1, arg2, arg3); \
- Ret ret = __SYCL_PPCAT(prefix, call)(std::move(arg1), std::move(arg2), \
- std::move(arg3)); \
- return sycl::detail::convertDataToType<Ret, R>(std::move(ret)); \
+ __SYCL_EXTERN_IT3(R, prefix, call, arg1, arg2, arg3); \
+ return sycl::detail::convertFromOpenCLTypeFor<R>(__SYCL_PPCAT( \
+ prefix, call)(std::move(arg1), std::move(arg2), std::move(arg3))); \
  }
 
 #ifndef __SYCL_DEVICE_ONLY__

sycl/include/sycl/detail/generic_type_traits.hpp

@@ -704,6 +704,18 @@ template <typename T> auto convertToOpenCLType(T &&x) {
  return convertDataToType<T, OpenCLType>(std::forward<T>(x));
 }
 
+template <typename To, typename From> auto convertFromOpenCLTypeFor(From &&x) {
+ if constexpr (std::is_same_v<To, bool> &&
+ std::is_same_v<std::remove_reference_t<From>, bool>) {
+ // FIXME: Something seems to be wrong elsewhere...
+ return x;
+ } else {
+ static_assert(std::is_same_v<std::remove_reference_t<From>,
+ ConvertToOpenCLType_t<To>>);
+ return convertDataToType<From, To>(std::forward<From>(x));
+ }
+}
+
 // Used for all, any and select relational built-in functions
 template <typename T> inline constexpr T msbMask(T) {
  using UT = make_unsigned_t<T>;

sycl/include/sycl/detail/image_ocl_types.hpp

@@ -47,9 +47,8 @@ template <typename ImageType> struct sampled_opencl_image_type;
 #define __SYCL_INVOKE_SPIRV_CALL_ARG1(call) \
  template <typename R, typename T1> inline R __invoke_##call(T1 ParT1) { \
  using Ret = sycl::detail::ConvertToOpenCLType_t<R>; \
- T1 Arg1 = ParT1; \
- Ret RetVar = __spirv_##call<Ret, T1>(Arg1); \
- return sycl::detail::convertDataToType<Ret, R>(RetVar); \
+ return sycl::detail::convertFromOpenCLTypeFor<R>( \
+ __spirv_##call<Ret, T1>(ParT1)); \
  }
 
 // The macro defines the function __invoke_ImageXXXX,
@@ -74,12 +73,10 @@ static RetType __invoke__ImageRead(ImageT Img, CoordT Coords) {
 
  // Convert from sycl types to builtin types to get correct function mangling.
  using TempRetT = sycl::detail::ConvertToOpenCLType_t<RetType>;
-
  auto TmpCoords = sycl::detail::convertToOpenCLType(Coords);
 
- TempRetT Ret =
- __spirv_ImageRead<TempRetT, ImageT, decltype(TmpCoords)>(Img, TmpCoords);
- return sycl::detail::convertDataToType<TempRetT, RetType>(Ret);
+ return sycl::detail::convertFromOpenCLTypeFor<RetType>(
+ __spirv_ImageRead<TempRetT, ImageT, decltype(TmpCoords)>(Img, TmpCoords));
 }
 
 template <typename RetType, typename SmpImageT, typename CoordT>
@@ -98,10 +95,9 @@ static RetType __invoke__ImageReadLod(SmpImageT SmpImg, CoordT Coords,
  // Sampled Image must be an object whose type is OpTypeSampledImage
  // Image Operands encodes what operands follow. Either Lod
  // or Grad image operands must be present
- TempRetT Ret =
+ return sycl::detail::convertFromOpenCLTypeFor<RetType>(
  __spirv_ImageSampleExplicitLod<SmpImageT, TempRetT, decltype(TmpCoords)>(
- SmpImg, TmpCoords, ImageOperands::Lod, Level);
- return sycl::detail::convertDataToType<TempRetT, RetType>(Ret);
+ SmpImg, TmpCoords, ImageOperands::Lod, Level));
 }
 
 template <typename RetType, typename SmpImageT, typename CoordT>
@@ -122,10 +118,9 @@ static RetType __invoke__ImageReadGrad(SmpImageT SmpImg, CoordT Coords,
  // Sampled Image must be an object whose type is OpTypeSampledImage
  // Image Operands encodes what operands follow. Either Lod
  // or Grad image operands must be present
- TempRetT Ret =
+ return sycl::detail::convertFromOpenCLTypeFor<RetType>(
  __spirv_ImageSampleExplicitLod<SmpImageT, TempRetT, decltype(TmpCoords)>(
- SmpImg, TmpCoords, ImageOperands::Grad, TmpGraddX, TmpGraddY);
- return sycl::detail::convertDataToType<TempRetT, RetType>(Ret);
+ SmpImg, TmpCoords, ImageOperands::Grad, TmpGraddX, TmpGraddY));
 }
 
 template <typename RetType, typename ImageT, typename CoordT>
@@ -150,11 +145,10 @@ static RetType __invoke__ImageReadSampler(ImageT Img, CoordT Coords,
  enum ImageOperands { Lod = 0x2 };
 
  // Lod value is zero as mipmap is not supported.
- TempRetT Ret =
+ return sycl::detail::convertFromOpenCLTypeFor<RetType>(
  __spirv_ImageSampleExplicitLod<SampledT, TempRetT, decltype(TmpCoords)>(
  __spirv_SampledImage<ImageT, SampledT>(Img, Smpl), TmpCoords,
- ImageOperands::Lod, 0.0f);
- return sycl::detail::convertDataToType<TempRetT, RetType>(Ret);
+ ImageOperands::Lod, 0.0f));
 }
 
 namespace sycl {

sycl/include/sycl/ext/oneapi/experimental/builtins.hpp

@@ -95,8 +95,7 @@ inline __SYCL_ALWAYS_INLINE std::enable_if_t<
  sycl::detail::is_svgenfloatf_v<T> || sycl::detail::is_svgenfloath_v<T>, T>
 tanh(T x) __NOEXC {
 #if defined(__NVPTX__)
- using _ocl_T = sycl::detail::ConvertToOpenCLType_t<T>;
- return sycl::detail::convertDataToType<_ocl_T, T>(
+ return sycl::detail::convertFromOpenCLTypeFor<T>(
  __clc_native_tanh(sycl::detail::convertToOpenCLType(x)));
 #else
  return __sycl_std::__invoke_tanh<T>(x);
@@ -145,8 +144,7 @@ inline __SYCL_ALWAYS_INLINE
  std::enable_if_t<sycl::detail::is_svgenfloath_v<T>, T>
  exp2(T x) __NOEXC {
 #if defined(__NVPTX__)
- using _ocl_T = sycl::detail::ConvertToOpenCLType_t<T>;
- return sycl::detail::convertDataToType<_ocl_T, T>(
+ return sycl::detail::convertFromOpenCLTypeFor<T>(
  __clc_native_exp2(sycl::detail::convertToOpenCLType(x)));
 #else
  return __sycl_std::__invoke_exp2<T>(x);