Skip to content

Commit 895bb1e

Browse files
author
iclsrc
committed
Merge from 'sycl' to 'sycl-web' (2 commits)
2 parents 7903415 + 313c4f0 commit 895bb1e

File tree

14 files changed

+135
-31
lines changed

14 files changed

+135
-31
lines changed

‎sycl/include/sycl/detail/builtins/helper_macros.hpp

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -197,7 +197,8 @@
197197
[](NUM_ARGS##_AUTO_ARG) { return (NS::NAME)(NUM_ARGS##_ARG); }, \
198198
NUM_ARGS##_ARG); \
199199
} else { \
200-
return __VA_ARGS__(NUM_ARGS##_CONVERTED_ARG); \
200+
return bit_cast<detail::ENABLER<NUM_ARGS##_TEMPLATE_TYPE>>( \
201+
__VA_ARGS__(NUM_ARGS##_CONVERTED_ARG)); \
201202
} \
202203
}
203204

‎sycl/include/sycl/detail/builtins/math_functions.inc

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -254,7 +254,9 @@ auto builtin_delegate_ptr_impl(FuncTy F, PtrTy p, Ts... xs) {
254254
detail::NON_SCALAR_ENABLER<SYCL_CONCAT(LESS_ONE(NUM_ARGS), _TEMPLATE_TYPE), \
255255
PtrTy> \
256256
NAME(SYCL_CONCAT(LESS_ONE(NUM_ARGS), _TEMPLATE_TYPE_ARG), PtrTy p) { \
257-
return detail::NAME##_impl(SYCL_CONCAT(LESS_ONE(NUM_ARGS), _ARG), p); \
257+
return bit_cast<detail::NON_SCALAR_ENABLER< \
258+
SYCL_CONCAT(LESS_ONE(NUM_ARGS), _TEMPLATE_TYPE), PtrTy>>( \
259+
detail::NAME##_impl(SYCL_CONCAT(LESS_ONE(NUM_ARGS), _ARG), p)); \
258260
}
259261

260262
#if __SYCL_DEVICE_ONLY__

‎sycl/include/sycl/detail/builtins/relational_functions.inc

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -63,7 +63,7 @@ auto builtin_device_rel_impl(FuncTy F, const Ts &...xs) {
6363
// the relation builtin (vector of int16_t/int32_t/int64_t depending on the
6464
// arguments' element type).
6565
auto ret = F(builtins::convert_arg(xs)...);
66-
vec<signed char, num_elements<T>::value> tmp{ret};
66+
auto tmp = bit_cast<vec<signed char, num_elements<T>::value>>(ret);
6767
using res_elem_type = fixed_width_signed<sizeof(get_elem_type_t<T>)>;
6868
static_assert(is_scalar_arithmetic_v<res_elem_type>);
6969
return tmp.template convert<res_elem_type>();

‎sycl/include/sycl/detail/spirv.hpp

Lines changed: 10 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -946,11 +946,12 @@ EnableIfNativeShuffle<T> Shuffle(GroupT g, T x, id<1> local_id) {
946946
return result;
947947
} else if constexpr (ext::oneapi::experimental::is_user_constructed_group_v<
948948
GroupT>) {
949-
return __spirv_GroupNonUniformShuffle(group_scope<GroupT>::value,
950-
convertToOpenCLType(x), LocalId);
949+
return convertFromOpenCLTypeFor<T>(__spirv_GroupNonUniformShuffle(
950+
group_scope<GroupT>::value, convertToOpenCLType(x), LocalId));
951951
} else {
952952
// Subgroup.
953-
return __spirv_SubgroupShuffleINTEL(convertToOpenCLType(x), LocalId);
953+
return convertFromOpenCLTypeFor<T>(
954+
__spirv_SubgroupShuffleINTEL(convertToOpenCLType(x), LocalId));
954955
}
955956
#else
956957
if constexpr (ext::oneapi::experimental::is_user_constructed_group_v<
@@ -987,8 +988,8 @@ EnableIfNativeShuffle<T> ShuffleXor(GroupT g, T x, id<1> mask) {
987988
convertToOpenCLType(x), TargetId);
988989
} else {
989990
// Subgroup.
990-
return __spirv_SubgroupShuffleXorINTEL(convertToOpenCLType(x),
991-
static_cast<uint32_t>(mask.get(0)));
991+
return convertFromOpenCLTypeFor<T>(__spirv_SubgroupShuffleXorINTEL(
992+
convertToOpenCLType(x), static_cast<uint32_t>(mask.get(0))));
992993
}
993994
#else
994995
if constexpr (ext::oneapi::experimental::is_user_constructed_group_v<
@@ -1035,8 +1036,8 @@ EnableIfNativeShuffle<T> ShuffleDown(GroupT g, T x, uint32_t delta) {
10351036
convertToOpenCLType(x), TargetId);
10361037
} else {
10371038
// Subgroup.
1038-
return __spirv_SubgroupShuffleDownINTEL(convertToOpenCLType(x),
1039-
convertToOpenCLType(x), delta);
1039+
return convertFromOpenCLTypeFor<T>(__spirv_SubgroupShuffleDownINTEL(
1040+
convertToOpenCLType(x), convertToOpenCLType(x), delta));
10401041
}
10411042
#else
10421043
if constexpr (ext::oneapi::experimental::is_user_constructed_group_v<
@@ -1079,8 +1080,8 @@ EnableIfNativeShuffle<T> ShuffleUp(GroupT g, T x, uint32_t delta) {
10791080
convertToOpenCLType(x), TargetId);
10801081
} else {
10811082
// Subgroup.
1082-
return __spirv_SubgroupShuffleUpINTEL(convertToOpenCLType(x),
1083-
convertToOpenCLType(x), delta);
1083+
return convertFromOpenCLTypeFor<T>(__spirv_SubgroupShuffleUpINTEL(
1084+
convertToOpenCLType(x), convertToOpenCLType(x), delta));
10841085
}
10851086
#else
10861087
if constexpr (ext::oneapi::experimental::is_user_constructed_group_v<

‎sycl/include/sycl/detail/vector_convert.hpp

Lines changed: 30 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -882,6 +882,36 @@ using ConvertBoolAndByteT =
882882
template <typename DataT, int NumElements>
883883
template <typename convertT, rounding_mode roundingMode>
884884
vec<convertT, NumElements> vec<DataT, NumElements>::convert() const {
885+
#if !__SYCL_USE_LIBSYCL8_VEC_IMPL
886+
auto getValue = [this](int Index) {
887+
using RetType = typename std::conditional_t<
888+
detail::is_byte_v<DataT>, int8_t,
889+
#ifdef __SYCL_DEVICE_ONLY__
890+
typename detail::map_type<
891+
DataT,
892+
#if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
893+
std::byte, /*->*/ std::uint8_t, //
894+
#endif
895+
bool, /*->*/ std::uint8_t, //
896+
sycl::half, /*->*/ sycl::detail::half_impl::StorageT, //
897+
sycl::ext::oneapi::bfloat16,
898+
/*->*/ sycl::ext::oneapi::bfloat16::Bfloat16StorageT, //
899+
char, /*->*/ detail::ConvertToOpenCLType_t<char>, //
900+
DataT, /*->*/ DataT //
901+
>::type
902+
#else
903+
DataT
904+
#endif
905+
>;
906+
907+
#ifdef __SYCL_DEVICE_ONLY__
908+
if constexpr (std::is_same_v<DataT, sycl::ext::oneapi::bfloat16>)
909+
return sycl::bit_cast<RetType>(this->m_Data[Index]);
910+
else
911+
#endif
912+
return static_cast<RetType>(this->m_Data[Index]);
913+
};
914+
#endif
885915
using T = detail::ConvertBoolAndByteT<DataT>;
886916
using R = detail::ConvertBoolAndByteT<convertT>;
887917
using bfloat16 = sycl::ext::oneapi::bfloat16;

‎sycl/include/sycl/handler.hpp

Lines changed: 16 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -384,16 +384,6 @@ template <int Dims> bool range_size_fits_in_size_t(const range<Dims> &r) {
384384
return true;
385385
}
386386

387-
template <typename KernelNameType>
388-
std::vector<kernel_param_desc_t> getKernelParamDescs() {
389-
std::vector<kernel_param_desc_t> Result;
390-
int NumParams = getKernelNumParams<KernelNameType>();
391-
Result.reserve(NumParams);
392-
for (int I = 0; I < NumParams; ++I) {
393-
Result.push_back(getKernelParamDesc<KernelNameType>(I));
394-
}
395-
return Result;
396-
}
397387
} // namespace detail
398388

399389
/// Command group handler class.
@@ -485,20 +475,30 @@ class __SYCL_EXPORT handler {
485475
"a single kernel or explicit memory operation.");
486476
}
487477

488-
/// Extracts and prepares kernel arguments from the lambda using information
489-
/// from the built-ins or integration header.
478+
#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
479+
// TODO: Those functions are not used anymore, remove it in the next
480+
// ABI-breaking window.
490481
void extractArgsAndReqsFromLambda(
491482
char *LambdaPtr,
492483
const std::vector<detail::kernel_param_desc_t> &ParamDescs, bool IsESIMD);
493-
// TODO Unused, remove during ABI breaking window
494484
void
495485
extractArgsAndReqsFromLambda(char *LambdaPtr, size_t KernelArgsNum,
496486
const detail::kernel_param_desc_t *KernelArgs,
497487
bool IsESIMD);
488+
#endif
489+
/// Extracts and prepares kernel arguments from the lambda using information
490+
/// from the built-ins or integration header.
491+
void extractArgsAndReqsFromLambda(
492+
char *LambdaPtr, detail::kernel_param_desc_t (*ParamDescGetter)(int),
493+
size_t NumKernelParams, bool IsESIMD);
498494

499495
/// Extracts and prepares kernel arguments set via set_arg(s).
500496
void extractArgsAndReqs();
501497

498+
#if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
499+
// TODO: processArg need not to be public
500+
__SYCL_DLL_LOCAL
501+
#endif
502502
void processArg(void *Ptr, const detail::kernel_param_kind_t &Kind,
503503
const int Size, const size_t Index, size_t &IndexShift,
504504
bool IsKernelCreatedFromSource, bool IsESIMD);
@@ -770,9 +770,11 @@ class __SYCL_EXPORT handler {
770770
// header, so don't perform things that require it.
771771
if constexpr (KernelHasName) {
772772
// TODO support ESIMD in no-integration-header case too.
773+
773774
clearArgs();
774775
extractArgsAndReqsFromLambda(MHostKernel->getPtr(),
775-
detail::getKernelParamDescs<KernelName>(),
776+
&(detail::getKernelParamDesc<KernelName>),
777+
detail::getKernelNumParams<KernelName>(),
776778
detail::isKernelESIMD<KernelName>());
777779
MKernelName = detail::getKernelName<KernelName>();
778780
} else {

‎sycl/include/sycl/vector.hpp

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -423,9 +423,11 @@ class __SYCL_EBO Swizzle
423423
using element_type = DataT;
424424
using value_type = DataT;
425425

426+
#if __SYCL_USE_LIBSYCL8_VEC_IMPL
426427
#ifdef __SYCL_DEVICE_ONLY__
427428
using vector_t = typename vec<DataT, NumElements>::vector_t;
428429
#endif // __SYCL_DEVICE_ONLY__
430+
#endif
429431

430432
Swizzle() = delete;
431433
Swizzle(const Swizzle &) = delete;
@@ -497,6 +499,7 @@ class __SYCL_EBO vec :
497499

498500
using Base = detail::vec_base<DataT, NumElements>;
499501

502+
#if __SYCL_USE_LIBSYCL8_VEC_IMPL
500503
#ifdef __SYCL_DEVICE_ONLY__
501504
using element_type_for_vector_t = typename detail::map_type<
502505
DataT,
@@ -541,6 +544,7 @@ class __SYCL_EBO vec :
541544

542545
private:
543546
#endif // __SYCL_DEVICE_ONLY__
547+
#endif
544548

545549
#if __SYCL_USE_LIBSYCL8_VEC_IMPL
546550
template <int... Indexes>
@@ -618,6 +622,7 @@ class __SYCL_EBO vec :
618622
static constexpr size_t get_size() { return byte_size(); }
619623
static constexpr size_t byte_size() noexcept { return sizeof(Base); }
620624

625+
#if __SYCL_USE_LIBSYCL8_VEC_IMPL
621626
private:
622627
// getValue should be able to operate on different underlying
623628
// types: enum cl_float#N , builtin vector float#N, builtin type float.
@@ -640,6 +645,8 @@ class __SYCL_EBO vec :
640645
}
641646

642647
public:
648+
#endif
649+
643650
// Out-of-class definition is in `sycl/detail/vector_convert.hpp`
644651
template <typename convertT,
645652
rounding_mode roundingMode = rounding_mode::automatic>

‎sycl/source/handler.cpp

Lines changed: 38 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1115,6 +1115,43 @@ void handler::extractArgsAndReqs() {
11151115
}
11161116
}
11171117

1118+
void handler::extractArgsAndReqsFromLambda(
1119+
char *LambdaPtr, detail::kernel_param_desc_t (*ParamDescGetter)(int),
1120+
size_t NumKernelParams, bool IsESIMD) {
1121+
size_t IndexShift = 0;
1122+
impl->MArgs.reserve(MaxNumAdditionalArgs * NumKernelParams);
1123+
1124+
for (size_t I = 0; I < NumKernelParams; ++I) {
1125+
detail::kernel_param_desc_t ParamDesc = ParamDescGetter(I);
1126+
void *Ptr = LambdaPtr + ParamDesc.offset;
1127+
const detail::kernel_param_kind_t &Kind = ParamDesc.kind;
1128+
const int &Size = ParamDesc.info;
1129+
if (Kind == detail::kernel_param_kind_t::kind_accessor) {
1130+
// For args kind of accessor Size is information about accessor.
1131+
// The first 11 bits of Size encodes the accessor target.
1132+
const access::target AccTarget =
1133+
static_cast<access::target>(Size & AccessTargetMask);
1134+
if ((AccTarget == access::target::device ||
1135+
AccTarget == access::target::constant_buffer) ||
1136+
(AccTarget == access::target::image ||
1137+
AccTarget == access::target::image_array)) {
1138+
detail::AccessorBaseHost *AccBase =
1139+
static_cast<detail::AccessorBaseHost *>(Ptr);
1140+
Ptr = detail::getSyclObjImpl(*AccBase).get();
1141+
} else if (AccTarget == access::target::local) {
1142+
detail::LocalAccessorBaseHost *LocalAccBase =
1143+
static_cast<detail::LocalAccessorBaseHost *>(Ptr);
1144+
Ptr = detail::getSyclObjImpl(*LocalAccBase).get();
1145+
}
1146+
}
1147+
processArg(Ptr, Kind, Size, I, IndexShift,
1148+
/*IsKernelCreatedFromSource=*/false, IsESIMD);
1149+
}
1150+
}
1151+
1152+
#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
1153+
// TODO: Those functions are not used anymore, remove it in the next
1154+
// ABI-breaking window.
11181155
void handler::extractArgsAndReqsFromLambda(
11191156
char *LambdaPtr, const std::vector<detail::kernel_param_desc_t> &ParamDescs,
11201157
bool IsESIMD) {
@@ -1149,14 +1186,14 @@ void handler::extractArgsAndReqsFromLambda(
11491186
}
11501187
}
11511188

1152-
// TODO Unused, remove during ABI breaking window
11531189
void handler::extractArgsAndReqsFromLambda(
11541190
char *LambdaPtr, size_t KernelArgsNum,
11551191
const detail::kernel_param_desc_t *KernelArgs, bool IsESIMD) {
11561192
std::vector<detail::kernel_param_desc_t> ParamDescs(
11571193
KernelArgs, KernelArgs + KernelArgsNum);
11581194
extractArgsAndReqsFromLambda(LambdaPtr, ParamDescs, IsESIMD);
11591195
}
1196+
#endif // __INTEL_PREVIEW_BREAKING_CHANGES
11601197

11611198
// Calling methods of kernel_impl requires knowledge of class layout.
11621199
// As this is impossible in header, there's a function that calls necessary

‎sycl/test-e2e/DeviceLib/built-ins/printf.cpp

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -59,10 +59,11 @@ int main() {
5959
sycl::vec<int, 4> v4{5, 6, 7, 8};
6060
#if defined(__SYCL_DEVICE_ONLY__) && (defined(__SPIR__) || defined(__SPIRV__))
6161
// On SPIRV devices, vectors can be printed via native OpenCL types:
62-
using ocl_int4 = sycl::vec<int, 4>::vector_t;
62+
using ocl_int4 = int __attribute__((ext_vector_type(4)));
6363
{
6464
static const CONSTANT char format[] = "%v4hld\n";
65-
ext::oneapi::experimental::printf(format, (ocl_int4)v4);
65+
ext::oneapi::experimental::printf(format,
66+
sycl::bit_cast<ocl_int4>(v4));
6667
}
6768

6869
// However, you are still able to print them by-element:

‎sycl/test/abi/sycl_symbols_linux.dump

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -3567,6 +3567,7 @@ _ZN4sycl3_V17handler27addLifetimeSharedPtrStorageESt10shared_ptrIKvE
35673567
_ZN4sycl3_V17handler27computeFallbackKernelBoundsEmm
35683568
_ZN4sycl3_V17handler28extractArgsAndReqsFromLambdaEPcRKSt6vectorINS0_6detail19kernel_param_desc_tESaIS5_EEb
35693569
_ZN4sycl3_V17handler28extractArgsAndReqsFromLambdaEPcmPKNS0_6detail19kernel_param_desc_tEb
3570+
_ZN4sycl3_V17handler28extractArgsAndReqsFromLambdaEPcPFNS0_6detail19kernel_param_desc_tEiEmb
35703571
_ZN4sycl3_V17handler28memcpyToHostOnlyDeviceGlobalEPKvS3_mbmm
35713572
_ZN4sycl3_V17handler28setArgsToAssociatedAccessorsEv
35723573
_ZN4sycl3_V17handler28setStateExplicitKernelBundleEv

0 commit comments

Comments
 (0)