Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

[NFCI][SYCL] Introduce convertToOpenCLType helper #12674

Merged
merged 5 commits into from
Feb 12, 2024
Merged

[NFCI][SYCL] Introduce convertToOpenCLType helper #12674

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
36facd8
[NFCI][SYCL] Introduce convertToOpenCLType helper
aelovikov-intel Feb 9, 2024
5da6934
clang-format
aelovikov-intel Feb 9, 2024
31f5060
Fix type detection
aelovikov-intel Feb 9, 2024
4708c42
clang-format
aelovikov-intel Feb 9, 2024
f1444d8
Merge remote-tracking branch 'origin/sycl' into convert-to-ocl-helper
aelovikov-intel Feb 9, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
60 changes: 30 additions & 30 deletions sycl/include/sycl/detail/builtins.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,20 +18,27 @@
#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_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)
#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(Ret, prefix, call, arg) \
using Arg = decltype(arg); \
Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Had to move those aliases here to avoid "unused type alias" warning/error for device compilation.

extern Ret __SYCL_PPCAT(prefix, call)(Arg)
#define __SYCL_EXTERN_IT2_SAME(Ret, prefix, call, Arg) \
#define __SYCL_EXTERN_IT2_SAME(Ret, prefix, call, arg) \
using Arg = decltype(arg); \
extern Ret __SYCL_PPCAT(prefix, call)(Arg, Arg)
#define __SYCL_EXTERN_IT2(Ret, prefix, call, Arg1, Arg2) \
#define __SYCL_EXTERN_IT2(Ret, prefix, call, arg1, arg2) \
using Arg1 = decltype(arg1); \
using Arg2 = decltype(arg2); \
extern Ret __SYCL_PPCAT(prefix, call)(Arg1, Arg2)
#define __SYCL_EXTERN_IT3(Ret, prefix, call, Arg1, Arg2, Arg3) \
#define __SYCL_EXTERN_IT3(Ret, prefix, call, arg1, arg2, arg3) \
using Arg1 = decltype(arg1); \
using Arg2 = decltype(arg2); \
using Arg3 = decltype(arg3); \
extern Ret __SYCL_PPCAT(prefix, call)(Arg1, Arg2, Arg3)
#endif

Expand All @@ -41,61 +48,54 @@
#define __SYCL_MAKE_CALL_ARG1(call, prefix) \
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>; \
using Arg1 = sycl::detail::ConvertToOpenCLType_t<T1>; \
__SYCL_EXTERN_IT1(Ret, prefix, call, Arg1); \
Arg1 arg1 = sycl::detail::convertDataToType<T1, Arg1>(std::move(t1)); \
__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)); \
}

#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>; \
using Arg1 = sycl::detail::ConvertToOpenCLType_t<T1>; \
using Arg2 = sycl::detail::ConvertToOpenCLType_t<T2>; \
__SYCL_EXTERN_IT2(Ret, prefix, call, Arg1, Arg2); \
Arg1 arg1 = sycl::detail::convertDataToType<T1, Arg1>(std::move(t1)); \
Arg2 arg2 = sycl::detail::convertDataToType<T2, Arg2>(std::move(t2)); \
__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)); \
}

#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>; \
using Arg = sycl::detail::ConvertToOpenCLType_t<T>; \
__SYCL_EXTERN_IT2_SAME(Ret, prefix, call, Arg); \
Arg arg1 = sycl::detail::convertDataToType<T, Arg>(std::move(t1)); \
Arg arg2 = sycl::detail::convertDataToType<T, Arg>(std::move(t2)); \
__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)); \
}

#define __SYCL_MAKE_CALL_ARG2_SAME_RESULT(call, prefix) \
template <typename T> \
inline __SYCL_ALWAYS_INLINE T __invoke_##call(T v1, T v2) __NOEXC { \
using Type = sycl::detail::ConvertToOpenCLType_t<T>; \
__SYCL_EXTERN_IT2_SAME(Type, prefix, call, Type); \
Type arg1 = sycl::detail::convertDataToType<T, Type>(std::move(v1)); \
Type arg2 = sycl::detail::convertDataToType<T, Type>(std::move(v2)); \
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)); \
}

#define __SYCL_MAKE_CALL_ARG3(call, prefix) \
template <typename R, typename T1, typename T2, typename T3> \
inline __SYCL_ALWAYS_INLINE R __invoke_##call(T1 t1, T2 t2, T3 t3) __NOEXC { \
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>; \
using Arg1 = sycl::detail::ConvertToOpenCLType_t<T1>; \
using Arg2 = sycl::detail::ConvertToOpenCLType_t<T2>; \
using Arg3 = sycl::detail::ConvertToOpenCLType_t<T3>; \
__SYCL_EXTERN_IT3(Ret, prefix, call, Arg1, Arg2, Arg3); \
Arg1 arg1 = sycl::detail::convertDataToType<T1, Arg1>(std::move(t1)); \
Arg2 arg2 = sycl::detail::convertDataToType<T2, Arg2>(std::move(t2)); \
Arg3 arg3 = sycl::detail::convertDataToType<T3, Arg3>(std::move(t3)); \
__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)); \
Expand Down
7 changes: 7 additions & 0 deletions sycl/include/sycl/detail/generic_type_traits.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -697,6 +697,13 @@ convertDataToType(FROM t) {
return ConvertNonVectorType<TO>(t);
}

// Now fuse the above into a simpler helper that's easy to use.
// TODO: That should probably be moved outside of "type_traits".
template <typename T> auto convertToOpenCLType(T &&x) {
using OpenCLType = ConvertToOpenCLType_t<std::remove_reference_t<T>>;
return convertDataToType<T, OpenCLType>(std::forward<T>(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
52 changes: 23 additions & 29 deletions sycl/include/sycl/detail/image_ocl_types.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -62,24 +62,23 @@ template <typename ImageT, typename CoordT, typename ValT>
static void __invoke__ImageWrite(ImageT Img, CoordT Coords, ValT Val) {

// Convert from sycl types to builtin types to get correct function mangling.
using TmpValT = sycl::detail::ConvertToOpenCLType_t<ValT>;
using TmpCoordT = sycl::detail::ConvertToOpenCLType_t<CoordT>;
auto TmpCoords = sycl::detail::convertToOpenCLType(Coords);
auto TmpVal = sycl::detail::convertToOpenCLType(Val);

TmpCoordT TmpCoord =
sycl::detail::convertDataToType<CoordT, TmpCoordT>(Coords);
TmpValT TmpVal = sycl::detail::convertDataToType<ValT, TmpValT>(Val);
__spirv_ImageWrite<ImageT, TmpCoordT, TmpValT>(Img, TmpCoord, TmpVal);
__spirv_ImageWrite<ImageT, decltype(TmpCoords), decltype(TmpVal)>(
Img, TmpCoords, TmpVal);
}

template <typename RetType, typename ImageT, typename CoordT>
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>;
using TempArgT = sycl::detail::ConvertToOpenCLType_t<CoordT>;

TempArgT Arg = sycl::detail::convertDataToType<CoordT, TempArgT>(Coords);
TempRetT Ret = __spirv_ImageRead<TempRetT, ImageT, TempArgT>(Img, Arg);
auto TmpCoords = sycl::detail::convertToOpenCLType(Coords);

TempRetT Ret =
__spirv_ImageRead<TempRetT, ImageT, decltype(TmpCoords)>(Img, TmpCoords);
return sycl::detail::convertDataToType<TempRetT, RetType>(Ret);
}

Expand All @@ -89,10 +88,7 @@ static RetType __invoke__ImageReadLod(SmpImageT SmpImg, CoordT Coords,

// Convert from sycl types to builtin types to get correct function mangling.
using TempRetT = sycl::detail::ConvertToOpenCLType_t<RetType>;
using TempArgT = sycl::detail::ConvertToOpenCLType_t<CoordT>;

TempArgT TmpCoords =
sycl::detail::convertDataToType<CoordT, TempArgT>(Coords);
auto TmpCoords = sycl::detail::convertToOpenCLType(Coords);

enum ImageOperands { Lod = 0x2 };

Expand All @@ -102,8 +98,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 = __spirv_ImageSampleExplicitLod<SmpImageT, TempRetT, TempArgT>(
SmpImg, TmpCoords, ImageOperands::Lod, Level);
TempRetT Ret =
__spirv_ImageSampleExplicitLod<SmpImageT, TempRetT, decltype(TmpCoords)>(
SmpImg, TmpCoords, ImageOperands::Lod, Level);
return sycl::detail::convertDataToType<TempRetT, RetType>(Ret);
}

Expand All @@ -113,12 +110,9 @@ static RetType __invoke__ImageReadGrad(SmpImageT SmpImg, CoordT Coords,

// Convert from sycl types to builtin types to get correct function mangling.
using TempRetT = sycl::detail::ConvertToOpenCLType_t<RetType>;
using TempArgT = sycl::detail::ConvertToOpenCLType_t<CoordT>;

TempArgT TmpCoords =
sycl::detail::convertDataToType<CoordT, TempArgT>(Coords);
TempArgT TmpGraddX = sycl::detail::convertDataToType<CoordT, TempArgT>(Dx);
TempArgT TmpGraddY = sycl::detail::convertDataToType<CoordT, TempArgT>(Dy);
auto TmpCoords = sycl::detail::convertToOpenCLType(Coords);
auto TmpGraddX = sycl::detail::convertToOpenCLType(Dx);
auto TmpGraddY = sycl::detail::convertToOpenCLType(Dy);

enum ImageOperands { Grad = 0x3 };

Expand All @@ -128,8 +122,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 = __spirv_ImageSampleExplicitLod<SmpImageT, TempRetT, TempArgT>(
SmpImg, TmpCoords, ImageOperands::Grad, TmpGraddX, TmpGraddY);
TempRetT Ret =
__spirv_ImageSampleExplicitLod<SmpImageT, TempRetT, decltype(TmpCoords)>(
SmpImg, TmpCoords, ImageOperands::Grad, TmpGraddX, TmpGraddY);
return sycl::detail::convertDataToType<TempRetT, RetType>(Ret);
}

Expand All @@ -139,12 +134,10 @@ static RetType __invoke__ImageReadSampler(ImageT Img, CoordT Coords,

// Convert from sycl types to builtin types to get correct function mangling.
using TempRetT = sycl::detail::ConvertToOpenCLType_t<RetType>;
using TempArgT = sycl::detail::ConvertToOpenCLType_t<CoordT>;
using SampledT =
typename sycl::detail::sampled_opencl_image_type<ImageT>::type;

TempArgT TmpCoords =
sycl::detail::convertDataToType<CoordT, TempArgT>(Coords);
auto TmpCoords = sycl::detail::convertToOpenCLType(Coords);
// According to validation rules(SPIR-V specification, section 2.16.1) result
// of __spirv_SampledImage is allowed to be an operand of image lookup
// and query instructions explicitly specified to take an operand whose
Expand All @@ -157,9 +150,10 @@ static RetType __invoke__ImageReadSampler(ImageT Img, CoordT Coords,
enum ImageOperands { Lod = 0x2 };

// Lod value is zero as mipmap is not supported.
TempRetT Ret = __spirv_ImageSampleExplicitLod<SampledT, TempRetT, TempArgT>(
__spirv_SampledImage<ImageT, SampledT>(Img, Smpl), TmpCoords,
ImageOperands::Lod, 0.0f);
TempRetT Ret =
__spirv_ImageSampleExplicitLod<SampledT, TempRetT, decltype(TmpCoords)>(
__spirv_SampledImage<ImageT, SampledT>(Img, Smpl), TmpCoords,
ImageOperands::Lod, 0.0f);
return sycl::detail::convertDataToType<TempRetT, RetType>(Ret);
}

Expand Down
Loading
Loading