Skip to content

Commit

Permalink
Merge branch 'sycl' into przemek/sampled-image-fetch
Browse files Browse the repository at this point in the history
  • Loading branch information
@przemektmalon
przemektmalon committed Feb 13, 2024
2 parents 7998596 + c872cad commit 0c1119a
Show file tree
Hide file tree
Showing 32 changed files with 428 additions and 392 deletions.
65 changes: 35 additions & 30 deletions sycl/include/sycl/access/access.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -256,10 +256,44 @@ struct deduce_AS
};
#endif

template <typename T> struct remove_decoration_impl {
using type = T;
};

#ifdef __SYCL_DEVICE_ONLY__
template <typename T> struct remove_decoration_impl<__OPENCL_GLOBAL_AS__ T> {
using type = T;
};

#ifdef __ENABLE_USM_ADDR_SPACE__
template <typename T>
struct remove_decoration_impl<__OPENCL_GLOBAL_DEVICE_AS__ T> {
using type = T;
};

template <typename T>
struct remove_decoration_impl<__OPENCL_GLOBAL_HOST_AS__ T> {
using type = T;
};

#endif // __ENABLE_USM_ADDR_SPACE__

template <typename T> struct remove_decoration_impl<__OPENCL_PRIVATE_AS__ T> {
using type = T;
};

template <typename T> struct remove_decoration_impl<__OPENCL_LOCAL_AS__ T> {
using type = T;
};

template <typename T> struct remove_decoration_impl<__OPENCL_CONSTANT_AS__ T> {
using type = T;
};
#endif // __SYCL_DEVICE_ONLY__
} // namespace detail

template <typename T> struct remove_decoration {
using type = T;
using type = typename detail::remove_decoration_impl<T>::type;
};

// Propagate through const qualifier.
Expand Down Expand Up @@ -287,35 +321,6 @@ template <typename T> struct remove_decoration<const T &> {
using type = const typename remove_decoration<T>::type &;
};

#ifdef __SYCL_DEVICE_ONLY__
template <typename T> struct remove_decoration<__OPENCL_GLOBAL_AS__ T> {
using type = T;
};

#ifdef __ENABLE_USM_ADDR_SPACE__
template <typename T> struct remove_decoration<__OPENCL_GLOBAL_DEVICE_AS__ T> {
using type = T;
};

template <typename T> struct remove_decoration<__OPENCL_GLOBAL_HOST_AS__ T> {
using type = T;
};

#endif // __ENABLE_USM_ADDR_SPACE__

template <typename T> struct remove_decoration<__OPENCL_PRIVATE_AS__ T> {
using type = T;
};

template <typename T> struct remove_decoration<__OPENCL_LOCAL_AS__ T> {
using type = T;
};

template <typename T> struct remove_decoration<__OPENCL_CONSTANT_AS__ T> {
using type = T;
};
#endif // __SYCL_DEVICE_ONLY__

template <typename T>
using remove_decoration_t = typename remove_decoration<T>::type;

Expand Down
92 changes: 45 additions & 47 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,20 +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 @@ -41,64 +52,51 @@
#define __SYCL_MAKE_CALL_ARG1(call, prefix) \
template <typename R, typename T1> \
inline __SYCL_ALWAYS_INLINE R __invoke_##call(T1 t1) __NOEXC { \
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)); \
Ret ret = __SYCL_PPCAT(prefix, call)(std::move(arg1)); \
return sycl::detail::convertDataToType<Ret, R>(std::move(ret)); \
auto arg1 = sycl::detail::convertToOpenCLType(std::move(t1)); \
__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 { \
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)); \
Ret ret = __SYCL_PPCAT(prefix, call)(std::move(arg1), std::move(arg2)); \
return sycl::detail::convertDataToType<Ret, R>(std::move(ret)); \
auto arg1 = sycl::detail::convertToOpenCLType(std::move(t1)); \
auto arg2 = sycl::detail::convertToOpenCLType(std::move(t2)); \
__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 { \
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)); \
Ret ret = __SYCL_PPCAT(prefix, call)(std::move(arg1), std::move(arg2)); \
return sycl::detail::convertDataToType<Ret, R>(std::move(ret)); \
auto arg1 = sycl::detail::convertToOpenCLType(std::move(t1)); \
auto arg2 = sycl::detail::convertToOpenCLType(std::move(t2)); \
__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 { \
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)); \
Type ret = __SYCL_PPCAT(prefix, call)(std::move(arg1), std::move(arg2)); \
return sycl::detail::convertDataToType<Type, T>(std::move(ret)); \
auto arg1 = sycl::detail::convertToOpenCLType(std::move(v1)); \
auto arg2 = sycl::detail::convertToOpenCLType(std::move(v2)); \
__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) \
template <typename R, typename T1, typename T2, typename T3> \
inline __SYCL_ALWAYS_INLINE R __invoke_##call(T1 t1, T2 t2, T3 t3) __NOEXC { \
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)); \
Ret ret = __SYCL_PPCAT(prefix, call)(std::move(arg1), std::move(arg2), \
std::move(arg3)); \
return sycl::detail::convertDataToType<Ret, R>(std::move(ret)); \
auto arg1 = sycl::detail::convertToOpenCLType(std::move(t1)); \
auto arg2 = sycl::detail::convertToOpenCLType(std::move(t2)); \
auto arg3 = sycl::detail::convertToOpenCLType(std::move(t3)); \
__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
19 changes: 19 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 @@ -697,6 +697,25 @@ 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));
}

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
60 changes: 24 additions & 36 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 @@ -62,25 +61,22 @@ 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>;
auto TmpCoords = sycl::detail::convertToOpenCLType(Coords);

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

template <typename RetType, typename ImageT, typename CoordT>
Expand All @@ -101,10 +97,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 @@ -114,9 +107,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);
return sycl::detail::convertDataToType<TempRetT, RetType>(Ret);
return sycl::detail::convertFromOpenCLTypeFor<RetType>(
__spirv_ImageSampleExplicitLod<SmpImageT, TempRetT, decltype(TmpCoords)>(
SmpImg, TmpCoords, ImageOperands::Lod, Level));
}

template <typename RetType, typename SmpImageT, typename CoordT>
Expand All @@ -125,12 +118,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 @@ -140,9 +130,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);
return sycl::detail::convertDataToType<TempRetT, RetType>(Ret);
return sycl::detail::convertFromOpenCLTypeFor<RetType>(
__spirv_ImageSampleExplicitLod<SmpImageT, TempRetT, decltype(TmpCoords)>(
SmpImg, TmpCoords, ImageOperands::Grad, TmpGraddX, TmpGraddY));
}

template <typename RetType, typename ImageT, typename CoordT>
Expand All @@ -151,12 +141,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 @@ -169,10 +157,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);
return sycl::detail::convertDataToType<TempRetT, RetType>(Ret);
return sycl::detail::convertFromOpenCLTypeFor<RetType>(
__spirv_ImageSampleExplicitLod<SampledT, TempRetT, decltype(TmpCoords)>(
__spirv_SampledImage<ImageT, SampledT>(Img, Smpl), TmpCoords,
ImageOperands::Lod, 0.0f));
}

namespace sycl {
Expand Down
Loading

0 comments on commit 0c1119a

Please sign in to comment.