Skip to content

Commit 9286183

Browse files
steffenlarsenaelovikov-intel
steffenlarsen
and
aelovikov-intel
authored
[SYCL] Fix templated and convertible builtins (#11977)
With SYCL 2020 revision 8 builtins were clarified to have more specific definitions. This commit fixes a handful of definitions in accordance with the changes to the specification, primarily in relation to templated scalar builtins and builtins accepting implicitly convertible types and swizzles. --------- Signed-off-by: Larsen, Steffen <[email protected]> Co-authored-by: aelovikov-intel <[email protected]>
1 parent 5a5fcad commit 9286183

File tree

5 files changed

+468
-30
lines changed

5 files changed

+468
-30
lines changed

‎sycl/include/sycl/detail/generic_type_traits.hpp

Lines changed: 6 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -352,7 +352,12 @@ template <typename T>
352352
inline constexpr bool is_nan_type_v = is_contained_v<T, gtl::nan_list>;
353353

354354
// nan_types
355-
template <typename T, typename Enable = void> struct nan_types;
355+
template <typename T, typename Enable = void> struct nan_types {
356+
// Nonsensical case for types implicitly convertible to scalar to avoid
357+
// templated overloads which are SFINAE'd out to cause compilation errors.
358+
using ret_type = void;
359+
using arg_type = int;
360+
};
356361

357362
template <typename T>
358363
struct nan_types<

‎sycl/include/sycl/detail/type_traits.hpp

Lines changed: 12 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -218,6 +218,12 @@ template <class T> struct make_signed<const T> {
218218
template <class T, int N> struct make_signed<vec<T, N>> {
219219
using type = vec<make_signed_t<T>, N>;
220220
};
221+
template <typename VecT, typename OperationLeftT, typename OperationRightT,
222+
template <typename> class OperationCurrentT, int... Indexes>
223+
struct make_signed<SwizzleOp<VecT, OperationLeftT, OperationRightT,
224+
OperationCurrentT, Indexes...>> {
225+
using type = make_signed_t<std::remove_cv_t<VecT>>;
226+
};
221227
template <class T, std::size_t N> struct make_signed<marray<T, N>> {
222228
using type = marray<make_signed_t<T>, N>;
223229
};
@@ -233,6 +239,12 @@ template <class T> struct make_unsigned<const T> {
233239
template <class T, int N> struct make_unsigned<vec<T, N>> {
234240
using type = vec<make_unsigned_t<T>, N>;
235241
};
242+
template <typename VecT, typename OperationLeftT, typename OperationRightT,
243+
template <typename> class OperationCurrentT, int... Indexes>
244+
struct make_unsigned<SwizzleOp<VecT, OperationLeftT, OperationRightT,
245+
OperationCurrentT, Indexes...>> {
246+
using type = make_unsigned_t<std::remove_cv_t<VecT>>;
247+
};
236248
template <class T, std::size_t N> struct make_unsigned<marray<T, N>> {
237249
using type = marray<make_unsigned_t<T>, N>;
238250
};

‎sycl/source/builtins_generator.py

Lines changed: 34 additions & 29 deletions
Original file line numberDiff line numberDiff line change
@@ -935,17 +935,17 @@ def custom_nan_invoke(return_type, arg_types, arg_names):
935935
Def("vfloatn", ["vfloatn", "float"], invoke_name="fmax_common", convert_args=[(1,0)]),
936936
Def("vdoublen", ["vdoublen", "double"], invoke_name="fmax_common", convert_args=[(1,0)]),
937937
Def("vhalfn", ["vhalfn", "half"], invoke_name="fmax_common", convert_args=[(1,0)]), # Non-standard. Deprecated.
938-
Def("igeninteger", ["igeninteger", "igeninteger"], invoke_name="s_max", marray_use_loop=True),
939-
Def("ugeninteger", ["ugeninteger", "ugeninteger"], invoke_name="u_max", marray_use_loop=True),
938+
Def("igeninteger", ["igeninteger", "igeninteger"], invoke_name="s_max", marray_use_loop=True, template_scalar_args=True),
939+
Def("ugeninteger", ["ugeninteger", "ugeninteger"], invoke_name="u_max", marray_use_loop=True, template_scalar_args=True),
940940
Def("vigeninteger", ["vigeninteger", "elementtype0"], invoke_name="s_max"),
941941
Def("vugeninteger", ["vugeninteger", "elementtype0"], invoke_name="u_max"),
942942
Def("mgentype", ["mgentype", "elementtype0"], marray_use_loop=True)],
943943
"(min)": [Def("genfloat", ["genfloat", "genfloat"], invoke_name="fmin_common", template_scalar_args=True),
944944
Def("vfloatn", ["vfloatn", "float"], invoke_name="fmin_common", convert_args=[(1,0)]),
945945
Def("vdoublen", ["vdoublen", "double"], invoke_name="fmin_common", convert_args=[(1,0)]),
946946
Def("vhalfn", ["vhalfn", "half"], invoke_name="fmin_common", convert_args=[(1,0)]), # Non-standard. Deprecated.
947-
Def("igeninteger", ["igeninteger", "igeninteger"], invoke_name="s_min", marray_use_loop=True),
948-
Def("ugeninteger", ["ugeninteger", "ugeninteger"], invoke_name="u_min", marray_use_loop=True),
947+
Def("igeninteger", ["igeninteger", "igeninteger"], invoke_name="s_min", marray_use_loop=True, template_scalar_args=True),
948+
Def("ugeninteger", ["ugeninteger", "ugeninteger"], invoke_name="u_min", marray_use_loop=True, template_scalar_args=True),
949949
Def("vigeninteger", ["vigeninteger", "elementtype0"], invoke_name="s_min"),
950950
Def("vugeninteger", ["vugeninteger", "elementtype0"], invoke_name="u_min"),
951951
Def("mgentype", ["mgentype", "elementtype0"], marray_use_loop=True)],
@@ -957,7 +957,7 @@ def custom_nan_invoke(return_type, arg_types, arg_names):
957957
Def("mdoublen", ["mdoublen", "mdoublen", "double"]),
958958
Def("mhalfn", ["mhalfn", "mhalfn", "half"])], # Non-standard. Deprecated.
959959
"radians": [Def("genfloat", ["genfloat"], template_scalar_args=True)],
960-
"step": [Def("genfloat", ["genfloat", "genfloat"]),
960+
"step": [Def("genfloat", ["genfloat", "genfloat"], template_scalar_args=True),
961961
Def("vfloatn", ["float", "vfloatn"], convert_args=[(0,1)]),
962962
Def("vdoublen", ["double", "vdoublen"], convert_args=[(0,1)]),
963963
Def("vhalfn", ["half", "vhalfn"], convert_args=[(0,1)]), # Non-standard. Deprecated.
@@ -989,25 +989,25 @@ def custom_nan_invoke(return_type, arg_types, arg_names):
989989
Def("float", ["mgeofloat", "mgeofloat"], invoke_name="Dot"),
990990
Def("double", ["mgeodouble", "mgeodouble"], invoke_name="Dot"),
991991
Def("half", ["mgeohalf", "mgeohalf"], invoke_name="Dot"),
992-
Def("sgenfloat", ["sgenfloat", "sgenfloat"], custom_invoke=(lambda return_types, arg_types, arg_names: ' return ' + ' * '.join(arg_names) + ';'))],
993-
"distance": [Def("float", ["gengeofloat", "gengeofloat"]),
994-
Def("double", ["gengeodouble", "gengeodouble"]),
995-
Def("half", ["gengeohalf", "gengeohalf"])],
996-
"length": [Def("float", ["gengeofloat"]),
997-
Def("double", ["gengeodouble"]),
998-
Def("half", ["gengeohalf"])],
999-
"normalize": [Def("gengeofloat", ["gengeofloat"]),
1000-
Def("gengeodouble", ["gengeodouble"]),
1001-
Def("gengeohalf", ["gengeohalf"])],
1002-
"fast_distance": [Def("float", ["gengeofloat", "gengeofloat"]),
1003-
Def("double", ["gengeodouble", "gengeodouble"]),
1004-
Def("half", ["gengeohalf", "gengeohalf"])],
1005-
"fast_length": [Def("float", ["gengeofloat"]),
1006-
Def("double", ["gengeodouble"]),
1007-
Def("half", ["gengeohalf"])],
1008-
"fast_normalize": [Def("gengeofloat", ["gengeofloat"]),
1009-
Def("gengeodouble", ["gengeodouble"]),
1010-
Def("gengeohalf", ["gengeohalf"])],
992+
Def("sgenfloat", ["sgenfloat", "sgenfloat"], template_scalar_args=True, custom_invoke=(lambda return_types, arg_types, arg_names: ' return ' + ' * '.join(arg_names) + ';'))],
993+
"distance": [Def("float", ["gengeofloat", "gengeofloat"], template_scalar_args=True),
994+
Def("double", ["gengeodouble", "gengeodouble"], template_scalar_args=True),
995+
Def("half", ["gengeohalf", "gengeohalf"], template_scalar_args=True)],
996+
"length": [Def("float", ["gengeofloat"], template_scalar_args=True),
997+
Def("double", ["gengeodouble"], template_scalar_args=True),
998+
Def("half", ["gengeohalf"], template_scalar_args=True)],
999+
"normalize": [Def("gengeofloat", ["gengeofloat"], template_scalar_args=True),
1000+
Def("gengeodouble", ["gengeodouble"], template_scalar_args=True),
1001+
Def("gengeohalf", ["gengeohalf"], template_scalar_args=True)],
1002+
"fast_distance": [Def("float", ["gengeofloat", "gengeofloat"], template_scalar_args=True),
1003+
Def("double", ["gengeodouble", "gengeodouble"], template_scalar_args=True),
1004+
Def("half", ["gengeohalf", "gengeohalf"], template_scalar_args=True)],
1005+
"fast_length": [Def("float", ["gengeofloat"], template_scalar_args=True),
1006+
Def("double", ["gengeodouble"], template_scalar_args=True),
1007+
Def("half", ["gengeohalf"], template_scalar_args=True)],
1008+
"fast_normalize": [Def("gengeofloat", ["gengeofloat"], template_scalar_args=True),
1009+
Def("gengeodouble", ["gengeodouble"], template_scalar_args=True),
1010+
Def("gengeohalf", ["gengeohalf"], template_scalar_args=True)],
10111011
# Relational functions
10121012
"isequal": [RelDef("samesizesignedint0", ["vgenfloat", "vgenfloat"], invoke_name="FOrdEqual"),
10131013
RelDef("bool", ["sgenfloat", "sgenfloat"], invoke_name="FOrdEqual"),
@@ -1052,13 +1052,13 @@ def custom_nan_invoke(return_type, arg_types, arg_names):
10521052
RelDef("bool", ["sgenfloat"], invoke_name="SignBitSet"),
10531053
RelDef("boolelements0", ["mgenfloat"])],
10541054
"any": [Def("int", ["vigeninteger"], custom_invoke=get_custom_any_all_vec_invoke("Any")),
1055-
Def("bool", ["sigeninteger"], custom_invoke=(lambda return_type, arg_types, arg_names: f' return bool(int(detail::msbIsSet({arg_names[0]})));')),
1055+
Def("bool", ["sigeninteger"], template_scalar_args=True, custom_invoke=(lambda return_type, arg_types, arg_names: f' return bool(int(detail::msbIsSet({arg_names[0]})));')),
10561056
Def("bool", ["migeninteger"], custom_invoke=get_custom_any_all_marray_invoke("any"))],
10571057
"all": [Def("int", ["vigeninteger"], custom_invoke=get_custom_any_all_vec_invoke("All")),
1058-
Def("bool", ["sigeninteger"], custom_invoke=(lambda return_type, arg_types, arg_names: f' return bool(int(detail::msbIsSet({arg_names[0]})));')),
1058+
Def("bool", ["sigeninteger"], template_scalar_args=True, custom_invoke=(lambda return_type, arg_types, arg_names: f' return bool(int(detail::msbIsSet({arg_names[0]})));')),
10591059
Def("bool", ["migeninteger"], custom_invoke=get_custom_any_all_marray_invoke("all"))],
10601060
"bitselect": [Def("vgentype", ["vgentype", "vgentype", "vgentype"]),
1061-
Def("sgentype", ["sgentype", "sgentype", "sgentype"]),
1061+
Def("sgentype", ["sgentype", "sgentype", "sgentype"], template_scalar_args=True),
10621062
Def("mgentype", ["mgentype", "mgentype", "mgentype"], marray_use_loop=True)],
10631063
"select": [Def("vint8n", ["vint8n", "vint8n", "vint8n"]),
10641064
Def("vint16n", ["vint16n", "vint16n", "vint16n"]),
@@ -1082,7 +1082,7 @@ def custom_nan_invoke(return_type, arg_types, arg_names):
10821082
Def("vfloatn", ["vfloatn", "vfloatn", "vuint32n"]),
10831083
Def("vdoublen", ["vdoublen", "vdoublen", "vuint64n"]),
10841084
Def("vhalfn", ["vhalfn", "vhalfn", "vuint16n"]),
1085-
Def("sgentype", ["sgentype", "sgentype", "bool"], custom_invoke=custom_bool_select_invoke),
1085+
Def("sgentype", ["sgentype", "sgentype", "bool"], template_scalar_args=True, custom_invoke=custom_bool_select_invoke),
10861086
Def("mgentype", ["mgentype", "mgentype", "mbooln"], marray_use_loop=True)]}
10871087
# List of all builtins definitions in the sycl::native namespace.
10881088
native_builtins = {"cos": [Def("genfloatf", ["genfloatf"], invoke_prefix="native_")],
@@ -1210,10 +1210,15 @@ def type_combinations(return_type, arg_types, template_scalars):
12101210
Generates all return and argument type combinations for a given builtin
12111211
definition.
12121212
"""
1213-
unique_types = list(dict.fromkeys(arg_types + [return_type]))
1213+
unique_types = list(dict.fromkeys(arg_types))
12141214
unique_type_lists = [builtin_types[unique_type] for unique_type in unique_types]
12151215
if template_scalars:
12161216
unique_type_lists = [convert_scalars_to_templated(type_list) for type_list in unique_type_lists]
1217+
if return_type not in unique_types:
1218+
# Add return type after scalars have been turned to template arguments if
1219+
# it is unique, to avoid undeducible return types.
1220+
unique_types.append(return_type)
1221+
unique_type_lists.append(builtin_types[return_type])
12171222
combinations = list(itertools.product(*unique_type_lists))
12181223
result = []
12191224
for combination in combinations:

‎sycl/test/basic_tests/builtins_implicitly_convertible_args.cpp

Lines changed: 161 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,161 @@
1+
// RUN: %clangxx -fsycl -fpreview-breaking-changes -fsyntax-only %s -o %t
2+
// REQUIRES: preview-breaking-changes-supported
3+
4+
// Checks that builtins without template arguments allow for implicit
5+
// conversions of arguments.
6+
7+
#include <sycl/sycl.hpp>
8+
9+
template <typename T> struct ImplicitlyConvertibleType {
10+
operator T() const { return {}; }
11+
};
12+
13+
#define ONE_ARG_DECLVAL_IMPLICITLY_CONVERTIBLE(...) \
14+
std::declval<ImplicitlyConvertibleType<__VA_ARGS__>>()
15+
16+
#define TWO_ARGS_DECLVAL_IMPLICITLY_CONVERTIBLE(...) \
17+
ONE_ARG_DECLVAL_IMPLICITLY_CONVERTIBLE(__VA_ARGS__), \
18+
ONE_ARG_DECLVAL_IMPLICITLY_CONVERTIBLE(__VA_ARGS__)
19+
20+
#define THREE_ARGS_DECLVAL_IMPLICITLY_CONVERTIBLE(...) \
21+
TWO_ARGS_DECLVAL_IMPLICITLY_CONVERTIBLE(__VA_ARGS__), \
22+
ONE_ARG_DECLVAL_IMPLICITLY_CONVERTIBLE(__VA_ARGS__)
23+
24+
#define ONE_ARG_DECLVAL(...) std::declval<__VA_ARGS__>()
25+
26+
#define TWO_ARGS_DECLVAL(...) \
27+
ONE_ARG_DECLVAL(__VA_ARGS__), ONE_ARG_DECLVAL(__VA_ARGS__)
28+
29+
#define THREE_ARGS_DECLVAL(...) \
30+
TWO_ARGS_DECLVAL(__VA_ARGS__), ONE_ARG_DECLVAL(__VA_ARGS__)
31+
32+
#define CHECK_INNER(NUM_ARGS, FUNC_NAME, ...) \
33+
static_assert(std::is_same_v< \
34+
decltype(sycl::FUNC_NAME( \
35+
NUM_ARGS##_DECLVAL_IMPLICITLY_CONVERTIBLE(__VA_ARGS__))), \
36+
decltype(sycl::FUNC_NAME(NUM_ARGS##_DECLVAL(__VA_ARGS__)))>);
37+
38+
#define FLOAT_CHECK(NUM_ARGS, FUNC_NAME) CHECK_INNER(NUM_ARGS, FUNC_NAME, float)
39+
40+
#define GENFLOAT_CHECK(NUM_ARGS, FUNC_NAME) \
41+
FLOAT_CHECK(NUM_ARGS, FUNC_NAME) \
42+
CHECK_INNER(NUM_ARGS, FUNC_NAME, sycl::half) \
43+
CHECK_INNER(NUM_ARGS, FUNC_NAME, double)
44+
45+
#define UGENINT_NAN_CHECK(NUM_ARGS, FUNC_NAME) \
46+
CHECK_INNER(NUM_ARGS, FUNC_NAME, unsigned int) \
47+
CHECK_INNER(NUM_ARGS, FUNC_NAME, unsigned short) \
48+
CHECK_INNER(NUM_ARGS, FUNC_NAME, unsigned long)
49+
50+
void check() {
51+
GENFLOAT_CHECK(ONE_ARG, acos)
52+
GENFLOAT_CHECK(ONE_ARG, acosh)
53+
GENFLOAT_CHECK(ONE_ARG, acospi)
54+
GENFLOAT_CHECK(ONE_ARG, asin)
55+
GENFLOAT_CHECK(ONE_ARG, asinh)
56+
GENFLOAT_CHECK(ONE_ARG, asinpi)
57+
GENFLOAT_CHECK(ONE_ARG, atan)
58+
GENFLOAT_CHECK(ONE_ARG, atanh)
59+
GENFLOAT_CHECK(ONE_ARG, atanpi)
60+
GENFLOAT_CHECK(TWO_ARGS, atan2)
61+
GENFLOAT_CHECK(TWO_ARGS, atan2pi)
62+
GENFLOAT_CHECK(ONE_ARG, cbrt)
63+
GENFLOAT_CHECK(ONE_ARG, ceil)
64+
GENFLOAT_CHECK(TWO_ARGS, copysign)
65+
GENFLOAT_CHECK(ONE_ARG, cos)
66+
GENFLOAT_CHECK(ONE_ARG, cosh)
67+
GENFLOAT_CHECK(ONE_ARG, cospi)
68+
GENFLOAT_CHECK(ONE_ARG, erfc)
69+
GENFLOAT_CHECK(ONE_ARG, erf)
70+
GENFLOAT_CHECK(ONE_ARG, exp)
71+
GENFLOAT_CHECK(ONE_ARG, exp2)
72+
GENFLOAT_CHECK(ONE_ARG, exp10)
73+
GENFLOAT_CHECK(ONE_ARG, expm1)
74+
GENFLOAT_CHECK(ONE_ARG, fabs)
75+
GENFLOAT_CHECK(TWO_ARGS, fdim)
76+
GENFLOAT_CHECK(ONE_ARG, floor)
77+
GENFLOAT_CHECK(THREE_ARGS, fma)
78+
GENFLOAT_CHECK(TWO_ARGS, fmax)
79+
GENFLOAT_CHECK(TWO_ARGS, fmin)
80+
GENFLOAT_CHECK(TWO_ARGS, fmod)
81+
GENFLOAT_CHECK(TWO_ARGS, hypot)
82+
GENFLOAT_CHECK(ONE_ARG, ilogb)
83+
GENFLOAT_CHECK(ONE_ARG, lgamma)
84+
GENFLOAT_CHECK(ONE_ARG, log)
85+
GENFLOAT_CHECK(ONE_ARG, log2)
86+
GENFLOAT_CHECK(ONE_ARG, log10)
87+
GENFLOAT_CHECK(ONE_ARG, log1p)
88+
GENFLOAT_CHECK(ONE_ARG, logb)
89+
GENFLOAT_CHECK(THREE_ARGS, mad)
90+
GENFLOAT_CHECK(TWO_ARGS, maxmag)
91+
GENFLOAT_CHECK(TWO_ARGS, minmag)
92+
UGENINT_NAN_CHECK(ONE_ARG, nan)
93+
GENFLOAT_CHECK(TWO_ARGS, nextafter)
94+
GENFLOAT_CHECK(TWO_ARGS, pow)
95+
GENFLOAT_CHECK(TWO_ARGS, powr)
96+
GENFLOAT_CHECK(TWO_ARGS, remainder)
97+
GENFLOAT_CHECK(ONE_ARG, rint)
98+
GENFLOAT_CHECK(ONE_ARG, round)
99+
GENFLOAT_CHECK(ONE_ARG, rsqrt)
100+
GENFLOAT_CHECK(ONE_ARG, sin)
101+
GENFLOAT_CHECK(ONE_ARG, sinh)
102+
GENFLOAT_CHECK(ONE_ARG, sinpi)
103+
GENFLOAT_CHECK(ONE_ARG, sqrt)
104+
GENFLOAT_CHECK(ONE_ARG, tan)
105+
GENFLOAT_CHECK(ONE_ARG, tanh)
106+
GENFLOAT_CHECK(ONE_ARG, tanpi)
107+
GENFLOAT_CHECK(ONE_ARG, tgamma)
108+
GENFLOAT_CHECK(ONE_ARG, trunc)
109+
110+
FLOAT_CHECK(ONE_ARG, native::cos)
111+
FLOAT_CHECK(TWO_ARGS, native::divide)
112+
FLOAT_CHECK(ONE_ARG, native::exp)
113+
FLOAT_CHECK(ONE_ARG, native::exp2)
114+
FLOAT_CHECK(ONE_ARG, native::exp10)
115+
FLOAT_CHECK(ONE_ARG, native::log)
116+
FLOAT_CHECK(ONE_ARG, native::log2)
117+
FLOAT_CHECK(ONE_ARG, native::log10)
118+
FLOAT_CHECK(TWO_ARGS, native::powr)
119+
FLOAT_CHECK(ONE_ARG, native::recip)
120+
FLOAT_CHECK(ONE_ARG, native::rsqrt)
121+
FLOAT_CHECK(ONE_ARG, native::sin)
122+
FLOAT_CHECK(ONE_ARG, native::sqrt)
123+
FLOAT_CHECK(ONE_ARG, native::tan)
124+
125+
FLOAT_CHECK(ONE_ARG, half_precision::cos)
126+
FLOAT_CHECK(TWO_ARGS, half_precision::divide)
127+
FLOAT_CHECK(ONE_ARG, half_precision::exp)
128+
FLOAT_CHECK(ONE_ARG, half_precision::exp2)
129+
FLOAT_CHECK(ONE_ARG, half_precision::exp10)
130+
FLOAT_CHECK(ONE_ARG, half_precision::log)
131+
FLOAT_CHECK(ONE_ARG, half_precision::log2)
132+
FLOAT_CHECK(ONE_ARG, half_precision::log10)
133+
FLOAT_CHECK(TWO_ARGS, half_precision::powr)
134+
FLOAT_CHECK(ONE_ARG, half_precision::recip)
135+
FLOAT_CHECK(ONE_ARG, half_precision::rsqrt)
136+
FLOAT_CHECK(ONE_ARG, half_precision::sin)
137+
FLOAT_CHECK(ONE_ARG, half_precision::sqrt)
138+
FLOAT_CHECK(ONE_ARG, half_precision::tan)
139+
140+
GENFLOAT_CHECK(TWO_ARGS, isequal)
141+
GENFLOAT_CHECK(TWO_ARGS, isnotequal)
142+
GENFLOAT_CHECK(TWO_ARGS, isgreater)
143+
GENFLOAT_CHECK(TWO_ARGS, isgreaterequal)
144+
GENFLOAT_CHECK(TWO_ARGS, isless)
145+
GENFLOAT_CHECK(TWO_ARGS, islessequal)
146+
GENFLOAT_CHECK(TWO_ARGS, islessgreater)
147+
GENFLOAT_CHECK(ONE_ARG, isfinite)
148+
GENFLOAT_CHECK(ONE_ARG, isinf)
149+
GENFLOAT_CHECK(ONE_ARG, isnan)
150+
GENFLOAT_CHECK(ONE_ARG, isnormal)
151+
GENFLOAT_CHECK(TWO_ARGS, isordered)
152+
GENFLOAT_CHECK(TWO_ARGS, isunordered)
153+
GENFLOAT_CHECK(ONE_ARG, signbit)
154+
}
155+
156+
int main() {
157+
check();
158+
159+
sycl::queue Q;
160+
Q.single_task([=]() { check(); });
161+
}

0 commit comments

Comments
 (0)