Skip to content

Commit

Permalink
[NFCI][SYCL] Introduce convertFromOpenCLTypeFor<T>(x) helper (#12684)
Browse files Browse the repository at this point in the history 
Continuation of refactoring series initiated in
#12674.
  • Loading branch information
@aelovikov-intel
aelovikov-intel authored Feb 12, 2024
1 parent 6639e78 commit 3843e6b
Show file tree
Hide file tree
Showing 4 changed files with 52 additions and 50 deletions.
58 changes: 28 additions & 30 deletions sycl/include/sycl/detail/builtins.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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.
Expand All @@ -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

Expand All @@ -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) \
Expand All @@ -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__
Expand Down
12 changes: 12 additions & 0 deletions sycl/include/sycl/detail/generic_type_traits.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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>;
Expand Down
26 changes: 10 additions & 16 deletions sycl/include/sycl/detail/image_ocl_types.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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,
Expand All @@ -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>
Expand All @@ -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>
Expand All @@ -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>
Expand All @@ -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 {
Expand Down
6 changes: 2 additions & 4 deletions sycl/include/sycl/ext/oneapi/experimental/builtins.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down Expand Up @@ -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);
Expand Down

0 comments on commit 3843e6b

Please sign in to comment.