Skip to content

Commit 46ed826

Browse files
committed
[SYCL] Fix integer vec conversions
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. The test uncovered that we have a problem with `std::byte` and `vec::convert`, but it seems like we have never tested a case like this and it may not be even a regression. In any case it will be fixed in a separate PR.
1 parent aa0171d commit 46ed826

File tree

3 files changed

+176
-63
lines changed

3 files changed

+176
-63
lines changed

sycl/include/sycl/detail/vector_convert.hpp

Lines changed: 13 additions & 60 deletions
Original file line numberDiff line numberDiff line change
@@ -80,12 +80,9 @@ using is_uint_to_uint =
8080
std::bool_constant<is_sugeninteger_v<T> && is_sugeninteger_v<R>>;
8181

8282
template <typename T, typename R>
83-
using is_sint_to_uint =
84-
std::bool_constant<is_sigeninteger_v<T> && is_sugeninteger_v<R>>;
85-
86-
template <typename T, typename R>
87-
using is_uint_to_sint =
88-
std::bool_constant<is_sugeninteger_v<T> && is_sigeninteger_v<R>>;
83+
using is_sint_to_from_uint = std::bool_constant<
84+
(detail::is_sigeninteger_v<T> && detail::is_sugeninteger_v<R>) ||
85+
(detail::is_sugeninteger_v<T> && detail::is_sigeninteger_v<R>)>;
8986

9087
template <typename T, typename R>
9188
using is_sint_to_float =
@@ -141,18 +138,6 @@ To ConvertUToF(From Value) {
141138
return static_cast<To>(Value);
142139
}
143140

144-
template <typename From, typename To, int VecSize,
145-
typename Enable = std::enable_if_t<VecSize == 1>>
146-
To SatConvertSToU(From Value) {
147-
return static_cast<To>(Value);
148-
}
149-
150-
template <typename From, typename To, int VecSize,
151-
typename Enable = std::enable_if_t<VecSize == 1>>
152-
To SatConvertUToS(From Value) {
153-
return static_cast<To>(Value);
154-
}
155-
156141
template <typename From, typename To, int VecSize,
157142
typename Enable = std::enable_if_t<VecSize == 1>,
158143
sycl::rounding_mode RM>
@@ -314,44 +299,6 @@ __SYCL_INT_INT_CONVERT(U, ulong)
314299
#undef __SYCL_VECTOR_INT_INT_CONVERT
315300
#undef __SYCL_INT_INT_CONVERT
316301

317-
// signed to unsigned, unsigned to signed conversions
318-
#define __SYCL_SCALAR_SINT_UINT_CONVERT(Op, DestType) \
319-
template <typename From, typename To, int VecSize, typename Enable> \
320-
enable_if_to_int_scalar_t<sycl::opencl::cl_##DestType, Enable, VecSize, To> \
321-
SatConvert##Op(From value) { \
322-
return __spirv_SatConvert##Op##_R##DestType(value); \
323-
}
324-
325-
#define __SYCL_VECTOR_SINT_UINT_CONVERT(Op, N, DestType) \
326-
template <typename From, typename To, int VecSize, typename Enable> \
327-
enable_if_to_int_vector_t<sycl::opencl::cl_##DestType, Enable, N, VecSize, \
328-
To> \
329-
SatConvert##Op(From value) { \
330-
return __spirv_SatConvert##Op##_R##DestType##N(value); \
331-
}
332-
333-
#define __SYCL_SINT_UINT_CONVERT(Op, DestType) \
334-
__SYCL_SCALAR_SINT_UINT_CONVERT(Op, DestType) \
335-
__SYCL_VECTOR_SINT_UINT_CONVERT(Op, 2, DestType) \
336-
__SYCL_VECTOR_SINT_UINT_CONVERT(Op, 3, DestType) \
337-
__SYCL_VECTOR_SINT_UINT_CONVERT(Op, 4, DestType) \
338-
__SYCL_VECTOR_SINT_UINT_CONVERT(Op, 8, DestType) \
339-
__SYCL_VECTOR_SINT_UINT_CONVERT(Op, 16, DestType)
340-
341-
__SYCL_SINT_UINT_CONVERT(UToS, char)
342-
__SYCL_SINT_UINT_CONVERT(UToS, short)
343-
__SYCL_SINT_UINT_CONVERT(UToS, int)
344-
__SYCL_SINT_UINT_CONVERT(UToS, long)
345-
346-
__SYCL_SINT_UINT_CONVERT(SToU, uchar)
347-
__SYCL_SINT_UINT_CONVERT(SToU, ushort)
348-
__SYCL_SINT_UINT_CONVERT(SToU, uint)
349-
__SYCL_SINT_UINT_CONVERT(SToU, ulong)
350-
351-
#undef __SYCL_SCALAR_SINT_UINT_CONVERT
352-
#undef __SYCL_VECTOR_SINT_UINT_CONVERT
353-
#undef __SYCL_SINT_UINT_CONVERT
354-
355302
// float to signed, float to unsigned conversion
356303
#define __SYCL_SCALAR_FLOAT_INT_CONVERT(Op, DestType, RoundingMode, \
357304
RoundingModeCondition) \
@@ -586,11 +533,17 @@ NativeToT convertImpl(NativeFromT Value) {
586533
else if constexpr (is_float_to_uint<FromT, ToT>::value)
587534
return ConvertFToU<NativeFromT, NativeToT, VecSize, ElemTy, RoundingMode>(
588535
Value);
589-
else if constexpr (is_sint_to_uint<FromT, ToT>::value)
590-
return SatConvertSToU<NativeFromT, NativeToT, VecSize, ElemTy>(Value);
591536
else {
592-
static_assert(is_uint_to_sint<FromT, ToT>::value);
593-
return SatConvertUToS<NativeFromT, NativeToT, VecSize, ElemTy>(Value);
537+
static_assert(is_sint_to_from_uint<FromT, ToT>::value);
538+
static_assert(VecSize == 1, "Conversion between signed and unsigned data "
539+
"types is only available for scalars");
540+
// vec::convert is underspecified and therefore it is not entirely clear
541+
// what to do here. 'static_cast' implementation matches SYCL CTS and it
542+
// matches our old implementation. Unfortunately, OpSetConvertUToS and
543+
// OpSatConvertSToU behave differently and we can't use them here until the
544+
// behavior of conversions is well-defined by the SYCL 2020 speficiation.
545+
// See https://github.com/KhronosGroup/SYCL-Docs/issues/492
546+
return static_cast<NativeToT>(Value);
594547
}
595548
}
596549

sycl/include/sycl/types.hpp

Lines changed: 15 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -676,10 +676,22 @@ template <typename Type, int NumElements> class vec {
676676
using R = vec_data_t<convertT>;
677677
vec<convertT, NumElements> Result;
678678

679-
if constexpr (NativeVec && vec<convertT, NumElements>::NativeVec) {
679+
// Whole vector conversion can only be done, if:
680+
constexpr bool canUseNativeVectorConvert =
681+
// - both vectors are represented using native vector types;
682+
NativeVec && vec<convertT, NumElements>::NativeVec &&
683+
// - it is not a signed to unsigned (or vice versa) conversion
684+
// see comments within 'convertImpl' for more details;
685+
!detail::is_sint_to_from_uint<T, R>::value &&
686+
// - destination type is not bool. bool is stored as integer under the
687+
// hood and therefore conversion to bool looks like conversion between
688+
// two integer types. Since bit pattern for true and false is not
689+
// defined, there is no guarantee that integer conversion yields
690+
// right results here;
691+
!std::is_same_v<convertT, bool>;
692+
693+
if constexpr (canUseNativeVectorConvert) {
680694
#ifdef __SYCL_DEVICE_ONLY__
681-
// If both vectors are representable as native vectors, then we can use
682-
// a single vector-wide operation to do a conversion:
683695
Result.m_Data = detail::convertImpl<
684696
T, R, roundingMode, NumElements, VectorDataType,
685697
typename vec<convertT, NumElements>::VectorDataType>(m_Data);
Lines changed: 148 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,148 @@
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+
// RUN: %{build} -o %t.out -DSYCL2020_DISABLE_DEPRECATION_WARNINGS
8+
// RUN: %{run} %t.out
9+
10+
#include <sycl/sycl.hpp>
11+
12+
#include <cstdint>
13+
#include <iostream>
14+
#include <string>
15+
#include <type_traits>
16+
17+
// Debug prints are hidden under macro to reduce amount of output in CI runs
18+
// and thus speed up tests. However, they are useful when debugging the test
19+
// locally and can be quickly turned on in there.
20+
#ifdef ENABLE_DEBUG_OUTPUT
21+
22+
template <typename T> std::string to_string() { return "unknown type"; }
23+
template <> std::string to_string<std::byte>() { return "std::byte"; }
24+
template <> std::string to_string<char>() { return "char"; }
25+
template <> std::string to_string<signed char>() { return "signed char"; }
26+
template <> std::string to_string<short>() { return "short"; }
27+
template <> std::string to_string<int>() { return "int"; }
28+
template <> std::string to_string<long>() { return "long"; }
29+
template <> std::string to_string<long long>() { return "long long"; }
30+
template <> std::string to_string<unsigned char>() { return "unsigned char"; }
31+
template <> std::string to_string<unsigned short>() { return "unsigned short"; }
32+
template <> std::string to_string<unsigned int>() { return "unsigned int"; }
33+
template <> std::string to_string<unsigned long>() { return "unsigned long"; }
34+
template <> std::string to_string<unsigned long long>() {
35+
return "unsigned long long";
36+
}
37+
template <> std::string to_string<bool>() { return "bool"; }
38+
39+
#define DEBUG_PRINT(x) std::cout << x << std::endl;
40+
41+
#else
42+
#define DEBUG_PRINT(x)
43+
#endif
44+
45+
template <typename T>
46+
void check_vectors_equal(sycl::vec<T, 4> a, sycl::vec<T, 4> b) {
47+
bool all_good =
48+
a.x() == b.x() && a.y() == b.y() && a.z() == b.z() && a.w() == b.w();
49+
if (!all_good) {
50+
DEBUG_PRINT("host and device results mismatch:");
51+
DEBUG_PRINT(
52+
"\t{" << static_cast<int>(a.x()) << ", " << static_cast<int>(a.y())
53+
<< ", " << static_cast<int>(a.z()) << ", "
54+
<< static_cast<int>(a.w()) << "} vs {" << static_cast<int>(b.x())
55+
<< ", " << static_cast<int>(b.y()) << ", "
56+
<< static_cast<int>(b.z()) << ", " << static_cast<int>(b.w())
57+
<< "}");
58+
}
59+
assert(all_good);
60+
}
61+
62+
template <typename From, typename To> void check_convert() {
63+
DEBUG_PRINT("checking vec<" << to_string<From>() << ", 4>::convert<"
64+
<< to_string<To>() << ">()");
65+
66+
sycl::vec<From, 4> input;
67+
if constexpr (std::is_signed_v<From>) {
68+
input = sycl::vec<From, 4>{static_cast<From>(37), static_cast<From>(0),
69+
static_cast<From>(-11), static_cast<From>(13)};
70+
} else {
71+
input = sycl::vec<From, 4>{static_cast<From>(37), static_cast<From>(0),
72+
static_cast<From>(11), static_cast<From>(13)};
73+
}
74+
75+
sycl::vec<To, 4> hostResult = input.template convert<To>();
76+
77+
sycl::buffer<sycl::vec<To, 4>> buf(sycl::range{1});
78+
sycl::queue q;
79+
q.submit([&](sycl::handler &cgh) {
80+
sycl::accessor acc(buf, cgh);
81+
cgh.single_task([=]() { acc[0] = input.template convert<To>(); });
82+
}).wait();
83+
84+
auto acc = buf.get_host_access();
85+
auto deviceResult = acc[0];
86+
87+
// Host and device results must match.
88+
check_vectors_equal(hostResult, deviceResult);
89+
90+
// And they should match with a reference, which is for integer conversions
91+
// can be computed with a simple static_cast.
92+
// Strictly speaking, integer conversions are underspecified in the SYCL 2020
93+
// spec, but `static_cast` implementation matches SYCL-CTS, so we will leave
94+
// it here for now as well.
95+
// See https://github.com/KhronosGroup/SYCL-Docs/issues/492
96+
assert(deviceResult.x() == static_cast<To>(input.x()));
97+
assert(deviceResult.y() == static_cast<To>(input.y()));
98+
assert(deviceResult.z() == static_cast<To>(input.z()));
99+
assert(deviceResult.w() == static_cast<To>(input.w()));
100+
}
101+
102+
template <class T>
103+
constexpr auto has_unsigned_v =
104+
std::is_integral_v<T> && !std::is_same_v<T, bool> &&
105+
!std::is_same_v<T, sycl::byte> && !std::is_same_v<T, std::byte>;
106+
107+
template <typename From, typename To> void check_signed_unsigned_convert_to() {
108+
check_convert<From, To>();
109+
if constexpr (has_unsigned_v<To>)
110+
check_convert<From, std::make_unsigned_t<To>>();
111+
if constexpr (has_unsigned_v<From>)
112+
check_convert<std::make_unsigned_t<From>, To>();
113+
if constexpr (has_unsigned_v<To> && has_unsigned_v<From>)
114+
check_convert<std::make_unsigned_t<From>, std::make_unsigned_t<To>>();
115+
}
116+
117+
template <typename From> void check_convert_from() {
118+
check_signed_unsigned_convert_to<From, sycl::byte>();
119+
// check_signed_unsigned_convert_to<From, std::byte>();
120+
check_signed_unsigned_convert_to<From, std::int8_t>();
121+
check_signed_unsigned_convert_to<From, std::int16_t>();
122+
check_signed_unsigned_convert_to<From, std::int32_t>();
123+
check_signed_unsigned_convert_to<From, std::int64_t>();
124+
check_signed_unsigned_convert_to<From, bool>();
125+
check_signed_unsigned_convert_to<From, char>();
126+
check_signed_unsigned_convert_to<From, signed char>();
127+
check_signed_unsigned_convert_to<From, short>();
128+
check_signed_unsigned_convert_to<From, int>();
129+
check_signed_unsigned_convert_to<From, long>();
130+
check_signed_unsigned_convert_to<From, long long>();
131+
}
132+
133+
int main() {
134+
check_convert_from<sycl::byte>();
135+
// FIXME: enable test cases below once compilation issues for them are fixed
136+
// check_convert_from<std::byte>();
137+
check_convert_from<std::int8_t>();
138+
check_convert_from<std::int16_t>();
139+
check_convert_from<std::int32_t>();
140+
check_convert_from<std::int64_t>();
141+
check_convert_from<char>();
142+
check_convert_from<signed char>();
143+
check_convert_from<short>();
144+
check_convert_from<int>();
145+
check_convert_from<long>();
146+
check_convert_from<long long>();
147+
check_convert_from<bool>();
148+
}

0 commit comments

Comments
 (0)