Skip to content

Commit 3d5e41f

Browse files
[SYCL] Fix integer vec conversions (#11821)
This is a bugfix for regressions caused by #11770 The following cases are fixed/updated: - Signed to unsigned (and vice versa) conversions are made to always go through per-element path, because native SPIR-V instruction behavies differently from what SYCL CTS expects. Strictly speaking, this could be a CTS issue, because `vec::convert` is underspecified (see KhronosGroup/SYCL-Docs#492), but until the spec is clarified we will stick with existing behavior; - Conversions to bool are made to always go through per-element path, because under the hood bool is represented as an integer and there is no guarantee that regular integer conversion will produce the right bit patterns for `true` and `false` which are not defined by any specification; Added an on-device test to check that `vec::convert` for integer types behaves correctly and consistently on both host and device.
1 parent 396eb5a commit 3d5e41f

File tree

3 files changed

+199
-73
lines changed

3 files changed

+199
-73
lines changed

sycl/include/sycl/detail/vector_convert.hpp

Lines changed: 14 additions & 60 deletions
Original file line numberDiff line numberDiff line change
@@ -85,12 +85,9 @@ using is_uint_to_uint =
8585
std::bool_constant<is_sugeninteger_v<T> && is_sugeninteger_v<R>>;
8686

8787
template <typename T, typename R>
88-
using is_sint_to_uint =
89-
std::bool_constant<is_sigeninteger_v<T> && is_sugeninteger_v<R>>;
90-
91-
template <typename T, typename R>
92-
using is_uint_to_sint =
93-
std::bool_constant<is_sugeninteger_v<T> && is_sigeninteger_v<R>>;
88+
using is_sint_to_from_uint = std::bool_constant<
89+
(detail::is_sigeninteger_v<T> && detail::is_sugeninteger_v<R>) ||
90+
(detail::is_sugeninteger_v<T> && detail::is_sigeninteger_v<R>)>;
9491

9592
template <typename T, typename R>
9693
using is_sint_to_float =
@@ -146,18 +143,6 @@ To ConvertUToF(From Value) {
146143
return static_cast<To>(Value);
147144
}
148145

149-
template <typename From, typename To, int VecSize,
150-
typename Enable = std::enable_if_t<VecSize == 1>>
151-
To SatConvertSToU(From Value) {
152-
return static_cast<To>(Value);
153-
}
154-
155-
template <typename From, typename To, int VecSize,
156-
typename Enable = std::enable_if_t<VecSize == 1>>
157-
To SatConvertUToS(From Value) {
158-
return static_cast<To>(Value);
159-
}
160-
161146
template <typename From, typename To, int VecSize,
162147
typename Enable = std::enable_if_t<VecSize == 1>,
163148
sycl::rounding_mode RM>
@@ -319,44 +304,6 @@ __SYCL_INT_INT_CONVERT(U, ulong)
319304
#undef __SYCL_VECTOR_INT_INT_CONVERT
320305
#undef __SYCL_INT_INT_CONVERT
321306

322-
// signed to unsigned, unsigned to signed conversions
323-
#define __SYCL_SCALAR_SINT_UINT_CONVERT(Op, DestType) \
324-
template <typename From, typename To, int VecSize, typename Enable> \
325-
enable_if_to_int_scalar_t<sycl::opencl::cl_##DestType, Enable, VecSize, To> \
326-
SatConvert##Op(From value) { \
327-
return __spirv_SatConvert##Op##_R##DestType(value); \
328-
}
329-
330-
#define __SYCL_VECTOR_SINT_UINT_CONVERT(Op, N, DestType) \
331-
template <typename From, typename To, int VecSize, typename Enable> \
332-
enable_if_to_int_vector_t<sycl::opencl::cl_##DestType, Enable, N, VecSize, \
333-
To> \
334-
SatConvert##Op(From value) { \
335-
return __spirv_SatConvert##Op##_R##DestType##N(value); \
336-
}
337-
338-
#define __SYCL_SINT_UINT_CONVERT(Op, DestType) \
339-
__SYCL_SCALAR_SINT_UINT_CONVERT(Op, DestType) \
340-
__SYCL_VECTOR_SINT_UINT_CONVERT(Op, 2, DestType) \
341-
__SYCL_VECTOR_SINT_UINT_CONVERT(Op, 3, DestType) \
342-
__SYCL_VECTOR_SINT_UINT_CONVERT(Op, 4, DestType) \
343-
__SYCL_VECTOR_SINT_UINT_CONVERT(Op, 8, DestType) \
344-
__SYCL_VECTOR_SINT_UINT_CONVERT(Op, 16, DestType)
345-
346-
__SYCL_SINT_UINT_CONVERT(UToS, char)
347-
__SYCL_SINT_UINT_CONVERT(UToS, short)
348-
__SYCL_SINT_UINT_CONVERT(UToS, int)
349-
__SYCL_SINT_UINT_CONVERT(UToS, long)
350-
351-
__SYCL_SINT_UINT_CONVERT(SToU, uchar)
352-
__SYCL_SINT_UINT_CONVERT(SToU, ushort)
353-
__SYCL_SINT_UINT_CONVERT(SToU, uint)
354-
__SYCL_SINT_UINT_CONVERT(SToU, ulong)
355-
356-
#undef __SYCL_SCALAR_SINT_UINT_CONVERT
357-
#undef __SYCL_VECTOR_SINT_UINT_CONVERT
358-
#undef __SYCL_SINT_UINT_CONVERT
359-
360307
// float to signed, float to unsigned conversion
361308
#define __SYCL_SCALAR_FLOAT_INT_CONVERT(Op, DestType, RoundingMode, \
362309
RoundingModeCondition) \
@@ -591,11 +538,18 @@ NativeToT convertImpl(NativeFromT Value) {
591538
else if constexpr (is_float_to_uint<FromT, ToT>::value)
592539
return ConvertFToU<NativeFromT, NativeToT, VecSize, ElemTy, RoundingMode>(
593540
Value);
594-
else if constexpr (is_sint_to_uint<FromT, ToT>::value)
595-
return SatConvertSToU<NativeFromT, NativeToT, VecSize, ElemTy>(Value);
596541
else {
597-
static_assert(is_uint_to_sint<FromT, ToT>::value);
598-
return SatConvertUToS<NativeFromT, NativeToT, VecSize, ElemTy>(Value);
542+
static_assert(is_sint_to_from_uint<FromT, ToT>::value,
543+
"Unexpected conversion type");
544+
static_assert(VecSize == 1, "Conversion between signed and unsigned data "
545+
"types is only available for scalars");
546+
// vec::convert is underspecified and therefore it is not entirely clear
547+
// what to do here. 'static_cast' implementation matches SYCL CTS and it
548+
// matches our old implementation. Unfortunately, OpSetConvertUToS and
549+
// OpSatConvertSToU behave differently and we can't use them here until the
550+
// behavior of conversions is well-defined by the SYCL 2020 specificiation.
551+
// See https://github.com/KhronosGroup/SYCL-Docs/issues/492
552+
return static_cast<NativeToT>(Value);
599553
}
600554
}
601555

sycl/include/sycl/types.hpp

Lines changed: 27 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -752,10 +752,9 @@ template <typename Type, int NumElements> class vec {
752752
if constexpr (!std::is_same_v<DataT, convertT>) {
753753
// Dummy conversion for cases like vec<signed char> -> vec<char>
754754
vec<convertT, NumElements> Result;
755-
for (size_t I = 0; I < NumElements; ++I) {
756-
Result.setValue(I, vec_data<convertT>::get(static_cast<convertT>(
757-
vec_data<DataT>::get(getValue(I)))));
758-
}
755+
for (size_t I = 0; I < NumElements; ++I)
756+
Result.setValue(I, static_cast<convertT>(getValue(I)));
757+
759758
return Result;
760759
} else {
761760
// No conversion necessary
@@ -779,15 +778,33 @@ template <typename Type, int NumElements> class vec {
779778
using OpenCLT = detail::ConvertToOpenCLType_t<T>;
780779
using OpenCLR = detail::ConvertToOpenCLType_t<R>;
781780
vec<convertT, NumElements> Result;
781+
782782
#if defined(__INTEL_PREVIEW_BREAKING_CHANGES) && defined(__SYCL_DEVICE_ONLY__)
783783
using OpenCLVecT = OpenCLT __attribute__((ext_vector_type(NumElements)));
784784
using OpenCLVecR = OpenCLR __attribute__((ext_vector_type(NumElements)));
785-
if constexpr (NativeVec && vec<convertT, NumElements>::NativeVec &&
786-
std::is_convertible_v<decltype(m_Data), OpenCLVecT> &&
787-
std::is_convertible_v<decltype(Result.m_Data), OpenCLR>) {
788-
// If both vectors are representable as native vectors and these native
789-
// vectors can be converted to valid OpenCL representations, then we can
790-
// use a single vector-wide operation to do a conversion:
785+
// Whole vector conversion can only be done, if:
786+
constexpr bool canUseNativeVectorConvert =
787+
#ifdef __NVPTX__
788+
// - we are not on CUDA, see intel/llvm#11840
789+
false &&
790+
#endif
791+
// - both vectors are represented using native vector types;
792+
NativeVec && vec<convertT, NumElements>::NativeVec &&
793+
// - vec storage has an equivalent OpenCL native vector it is implicitly
794+
// convertible to. There are some corner cases where it is not the
795+
// case with char, long and long long types.
796+
std::is_convertible_v<decltype(m_Data), OpenCLVecT> &&
797+
std::is_convertible_v<decltype(Result.m_Data), OpenCLVecR> &&
798+
// - it is not a signed to unsigned (or vice versa) conversion
799+
// see comments within 'convertImpl' for more details;
800+
!detail::is_sint_to_from_uint<T, R>::value &&
801+
// - destination type is not bool. bool is stored as integer under the
802+
// hood and therefore conversion to bool looks like conversion between
803+
// two integer types. Since bit pattern for true and false is not
804+
// defined, there is no guarantee that integer conversion yields
805+
// right results here;
806+
!std::is_same_v<convertT, bool>;
807+
if constexpr (canUseNativeVectorConvert) {
791808
Result.m_Data = detail::convertImpl<T, R, roundingMode, NumElements,
792809
OpenCLVecT, OpenCLVecR>(m_Data);
793810
} else
@@ -803,9 +820,6 @@ template <typename Type, int NumElements> class vec {
803820
}
804821
}
805822

806-
if constexpr (std::is_same_v<convertT, bool>) {
807-
Result.ConvertToDataT();
808-
}
809823
return Result;
810824
}
811825

Lines changed: 158 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,158 @@
1+
// Basic acceptance test which checks vec::convert implementation on both
2+
// host and device. Coverage is limited to vec<T, 4> only, rest of vector sizes
3+
// are covered by SYCL-CTS.
4+
//
5+
// Macro is passed to silence warnings about sycl::byte
6+
//
7+
// XFAIL: cuda
8+
// FIXME: un-xfail the test once intel/llvm#11840 is resolved
9+
//
10+
// RUN: %{build} -o %t.out -DSYCL2020_DISABLE_DEPRECATION_WARNINGS
11+
// RUN: %{run} %t.out
12+
//
13+
// RUN: %if preview-breaking-changes-supported %{ %{build} -fpreview-breaking-changes -DSYCL2020_DISABLE_DEPRECATION_WARNINGS %s -o %t2.out %}
14+
// RUN: %if preview-breaking-changes-supported %{ %{run} %t2.out %}
15+
16+
#include <sycl/sycl.hpp>
17+
18+
#include <cstdint>
19+
#include <iostream>
20+
#include <string>
21+
#include <type_traits>
22+
23+
template <typename T> std::string to_string() { return "unknown type"; }
24+
template <> std::string to_string<std::byte>() { return "std::byte"; }
25+
template <> std::string to_string<char>() { return "char"; }
26+
template <> std::string to_string<signed char>() { return "signed char"; }
27+
template <> std::string to_string<short>() { return "short"; }
28+
template <> std::string to_string<int>() { return "int"; }
29+
template <> std::string to_string<long>() { return "long"; }
30+
template <> std::string to_string<long long>() { return "long long"; }
31+
template <> std::string to_string<unsigned char>() { return "unsigned char"; }
32+
template <> std::string to_string<unsigned short>() { return "unsigned short"; }
33+
template <> std::string to_string<unsigned int>() { return "unsigned int"; }
34+
template <> std::string to_string<unsigned long>() { return "unsigned long"; }
35+
template <> std::string to_string<unsigned long long>() {
36+
return "unsigned long long";
37+
}
38+
template <> std::string to_string<bool>() { return "bool"; }
39+
40+
template <typename T>
41+
bool check_vectors_equal(sycl::vec<T, 4> a, sycl::vec<T, 4> b,
42+
const std::string &fail_message) {
43+
bool result =
44+
a.x() == b.x() && a.y() == b.y() && a.z() == b.z() && a.w() == b.w();
45+
if (!result) {
46+
std::cout << fail_message << std::endl;
47+
std::cout << "\t{" << static_cast<int>(a.x()) << ", "
48+
<< static_cast<int>(a.y()) << ", " << static_cast<int>(a.z())
49+
<< ", " << static_cast<int>(a.w()) << "} vs {"
50+
<< static_cast<int>(b.x()) << ", " << static_cast<int>(b.y())
51+
<< ", " << static_cast<int>(b.z()) << ", "
52+
<< static_cast<int>(b.w()) << "}" << std::endl;
53+
}
54+
55+
return result;
56+
}
57+
58+
template <typename From, typename To> bool check_convert() {
59+
sycl::vec<From, 4> input;
60+
if constexpr (std::is_signed_v<From>) {
61+
input = sycl::vec<From, 4>{static_cast<From>(37), static_cast<From>(0),
62+
static_cast<From>(-11), static_cast<From>(13)};
63+
} else {
64+
input = sycl::vec<From, 4>{static_cast<From>(37), static_cast<From>(0),
65+
static_cast<From>(11), static_cast<From>(13)};
66+
}
67+
68+
sycl::vec<To, 4> hostResult = input.template convert<To>();
69+
70+
sycl::buffer<sycl::vec<To, 4>> buf(sycl::range{1});
71+
sycl::queue q;
72+
q.submit([&](sycl::handler &cgh) {
73+
sycl::accessor acc(buf, cgh);
74+
cgh.single_task([=]() { acc[0] = input.template convert<To>(); });
75+
}).wait();
76+
77+
auto acc = buf.get_host_access();
78+
auto deviceResult = acc[0];
79+
80+
std::string test =
81+
"(vec<" + to_string<From>() + ", 4>::convert<" + to_string<To>() + ">)";
82+
83+
// Host and device results must match.
84+
bool host_and_device_match = check_vectors_equal(
85+
hostResult, deviceResult, "host and device results do not match " + test);
86+
// And they should match with a reference, which is for integer conversions
87+
// can be computed with a simple static_cast.
88+
// Strictly speaking, integer conversions are underspecified in the SYCL 2020
89+
// spec, but `static_cast` implementation matches SYCL-CTS, so we will leave
90+
// it here for now as well.
91+
// See https://github.com/KhronosGroup/SYCL-Docs/issues/492
92+
sycl::vec<To, 4> reference{
93+
static_cast<To>(input.x()), static_cast<To>(input.y()),
94+
static_cast<To>(input.z()), static_cast<To>(input.w())};
95+
bool device_matches_reference = check_vectors_equal(
96+
deviceResult, reference, "device results don't match reference " + test);
97+
bool host_matches_reference = check_vectors_equal(
98+
hostResult, reference, "host resutls don't match reference " + test);
99+
100+
return host_and_device_match && device_matches_reference &&
101+
host_matches_reference;
102+
}
103+
104+
template <class T>
105+
constexpr auto has_unsigned_v =
106+
std::is_integral_v<T> && !std::is_same_v<T, bool> &&
107+
!std::is_same_v<T, sycl::byte> && !std::is_same_v<T, std::byte>;
108+
109+
template <typename From, typename To> bool check_signed_unsigned_convert_to() {
110+
bool pass = true;
111+
pass &= check_convert<From, To>();
112+
if constexpr (has_unsigned_v<To>)
113+
pass &= check_convert<From, std::make_unsigned_t<To>>();
114+
if constexpr (has_unsigned_v<From>)
115+
pass &= check_convert<std::make_unsigned_t<From>, To>();
116+
if constexpr (has_unsigned_v<To> && has_unsigned_v<From>)
117+
pass &=
118+
check_convert<std::make_unsigned_t<From>, std::make_unsigned_t<To>>();
119+
return pass;
120+
}
121+
122+
template <typename From> bool check_convert_from() {
123+
bool pass = true;
124+
pass &= check_signed_unsigned_convert_to<From, sycl::byte>();
125+
pass &= check_signed_unsigned_convert_to<From, std::byte>();
126+
pass &= check_signed_unsigned_convert_to<From, std::int8_t>();
127+
pass &= check_signed_unsigned_convert_to<From, std::int16_t>();
128+
pass &= check_signed_unsigned_convert_to<From, std::int32_t>();
129+
pass &= check_signed_unsigned_convert_to<From, std::int64_t>();
130+
pass &= check_signed_unsigned_convert_to<From, bool>();
131+
pass &= check_signed_unsigned_convert_to<From, char>();
132+
pass &= check_signed_unsigned_convert_to<From, signed char>();
133+
pass &= check_signed_unsigned_convert_to<From, short>();
134+
pass &= check_signed_unsigned_convert_to<From, int>();
135+
pass &= check_signed_unsigned_convert_to<From, long>();
136+
pass &= check_signed_unsigned_convert_to<From, long long>();
137+
138+
return pass;
139+
}
140+
141+
int main() {
142+
bool pass = true;
143+
pass &= check_convert_from<sycl::byte>();
144+
pass &= check_convert_from<std::byte>();
145+
pass &= check_convert_from<std::int8_t>();
146+
pass &= check_convert_from<std::int16_t>();
147+
pass &= check_convert_from<std::int32_t>();
148+
pass &= check_convert_from<std::int64_t>();
149+
pass &= check_convert_from<char>();
150+
pass &= check_convert_from<signed char>();
151+
pass &= check_convert_from<short>();
152+
pass &= check_convert_from<int>();
153+
pass &= check_convert_from<long>();
154+
pass &= check_convert_from<long long>();
155+
pass &= check_convert_from<bool>();
156+
157+
return static_cast<int>(!pass);
158+
}

0 commit comments

Comments
 (0)