Skip to content

Commit 066adc4

Browse files
alexeyvoronov-intelromanovvlad
alexeyvoronov-intel
authored and
romanovvlad
committed
[SYCL] Enable built-in functions overloaded for half type.
Signed-off-by: Alexey Voronov <alexey.voronov@intel.com>
1 parent 821dfce commit 066adc4

File tree

3 files changed

+385
-1072
lines changed

3 files changed

+385
-1072
lines changed

‎sycl/include/CL/sycl/builtins.hpp

Lines changed: 45 additions & 60 deletions
Original file line numberDiff line numberDiff line change
@@ -17,9 +17,6 @@
1717
// TODO Decide whether to mark functions with this attribute.
1818
#define __NOEXC /*noexcept*/
1919

20-
// TODO Remove when half type will supported by SYCL Runtime
21-
#define __HALF_NO_ENABLED
22-
2320
namespace cl {
2421
namespace sycl {
2522
#ifdef __SYCL_DEVICE_ONLY__
@@ -260,15 +257,15 @@ fmod(T x, T y) __NOEXC {
260257
// genfloat fract (genfloat x, genfloatptr iptr)
261258
template <typename T, typename T2>
262259
typename std::enable_if<
263-
detail::is_genfloat<T>::value &&detail::is_genfloatptr<T2>::value, T>::type
260+
detail::is_genfloat<T>::value && detail::is_genfloatptr<T2>::value, T>::type
264261
fract(T x, T2 iptr) __NOEXC {
265262
return __sycl_std::__invoke_fract<T>(x, iptr);
266263
}
267264

268265
// genfloat frexp (genfloat x, genintptr exp)
269266
template <typename T, typename T2>
270267
typename std::enable_if<
271-
detail::is_genfloat<T>::value &&detail::is_genintptr<T2>::value, T>::type
268+
detail::is_genfloat<T>::value && detail::is_genintptr<T2>::value, T>::type
272269
frexp(T x, T2 exp) __NOEXC {
273270
return __sycl_std::__invoke_frexp<T>(x, exp);
274271
}
@@ -308,7 +305,7 @@ ldexp(T x, int k) __NOEXC {
308305
// vgenfloat ldexp (vgenfloat x, genint k)
309306
template <typename T, typename T2>
310307
typename std::enable_if<
311-
detail::is_vgenfloat<T>::value &&detail::is_intn<T2>::value, T>::type
308+
detail::is_vgenfloat<T>::value && detail::is_intn<T2>::value, T>::type
312309
ldexp(T x, T2 k) __NOEXC {
313310
return __sycl_std::__invoke_ldexp<T>(x, k);
314311
}
@@ -323,7 +320,7 @@ lgamma(T x) __NOEXC {
323320
// genfloat lgamma_r (genfloat x, genintptr signp)
324321
template <typename T, typename T2>
325322
typename std::enable_if<
326-
detail::is_genfloat<T>::value &&detail::is_genintptr<T2>::value, T>::type
323+
detail::is_genfloat<T>::value && detail::is_genintptr<T2>::value, T>::type
327324
lgamma_r(T x, T2 signp) __NOEXC {
328325
return __sycl_std::__invoke_lgamma_r<T>(x, signp);
329326
}
@@ -387,7 +384,7 @@ minmag(T x, T y) __NOEXC {
387384
// genfloat modf (genfloat x, genfloatptr iptr)
388385
template <typename T, typename T2>
389386
typename std::enable_if<
390-
detail::is_genfloat<T>::value &&detail::is_genfloatptr<T2>::value, T>::type
387+
detail::is_genfloat<T>::value && detail::is_genfloatptr<T2>::value, T>::type
391388
modf(T x, T2 iptr) __NOEXC {
392389
return __sycl_std::__invoke_modf<T>(x, iptr);
393390
}
@@ -420,7 +417,7 @@ pow(T x, T y) __NOEXC {
420417
// genfloat pown (genfloat x, genint y)
421418
template <typename T, typename T2>
422419
typename std::enable_if<
423-
detail::is_genfloat<T>::value &&detail::is_genint<T2>::value, T>::type
420+
detail::is_genfloat<T>::value && detail::is_genint<T2>::value, T>::type
424421
pown(T x, T2 y) __NOEXC {
425422
return __sycl_std::__invoke_pown<T>(x, y);
426423
}
@@ -442,7 +439,7 @@ remainder(T x, T y) __NOEXC {
442439
// genfloat remquo (genfloat x, genfloat y, genintptr quo)
443440
template <typename T, typename T2>
444441
typename std::enable_if<
445-
detail::is_genfloat<T>::value &&detail::is_genintptr<T2>::value, T>::type
442+
detail::is_genfloat<T>::value && detail::is_genintptr<T2>::value, T>::type
446443
remquo(T x, T y, T2 quo) __NOEXC {
447444
return __sycl_std::__invoke_remquo<T>(x, y, quo);
448445
}
@@ -457,7 +454,7 @@ rint(T x) __NOEXC {
457454
// genfloat rootn (genfloat x, genint y)
458455
template <typename T, typename T2>
459456
typename std::enable_if<
460-
detail::is_genfloat<T>::value &&detail::is_genint<T2>::value, T>::type
457+
detail::is_genfloat<T>::value && detail::is_genint<T2>::value, T>::type
461458
rootn(T x, T2 y) __NOEXC {
462459
return __sycl_std::__invoke_rootn<T>(x, y);
463460
}
@@ -486,7 +483,7 @@ sin(T x) __NOEXC {
486483
// genfloat sincos (genfloat x, genfloatptr cosval)
487484
template <typename T, typename T2>
488485
typename std::enable_if<
489-
detail::is_genfloat<T>::value &&detail::is_genfloatptr<T2>::value, T>::type
486+
detail::is_genfloat<T>::value && detail::is_genfloatptr<T2>::value, T>::type
490487
sincos(T x, T2 cosval) __NOEXC {
491488
return __sycl_std::__invoke_sincos<T>(x, cosval);
492489
}
@@ -1042,15 +1039,13 @@ dot(T p0, T p1) __NOEXC {
10421039
return __sycl_std::__invoke_OpDot<cl::sycl::cl_double>(p0, p1);
10431040
}
10441041

1045-
#ifndef __HALF_NO_ENABLED
10461042
// half dot (vgengeohalf p0, vgengeohalf p1)
10471043
template <typename T>
10481044
typename std::enable_if<detail::is_vgengeohalf<T>::value,
10491045
cl::sycl::cl_half>::type
10501046
dot(T p0, T p1) __NOEXC {
10511047
return __sycl_std::__invoke_OpDot<cl::sycl::cl_half>(p0, p1);
10521048
}
1053-
#endif
10541049

10551050
// float distance (gengeofloat p0, gengeofloat p1)
10561051
template <typename T, typename = typename std::enable_if<
@@ -1066,14 +1061,12 @@ cl::sycl::cl_double distance(T p0, T p1) __NOEXC {
10661061
return __sycl_std::__invoke_distance<cl::sycl::cl_double>(p0, p1);
10671062
}
10681063

1069-
#ifndef __HALF_NO_ENABLED
10701064
// half distance (gengeohalf p0, gengeohalf p1)
10711065
template <typename T, typename = typename std::enable_if<
10721066
detail::is_gengeohalf<T>::value, T>::type>
10731067
cl::sycl::cl_half distance(T p0, T p1) __NOEXC {
10741068
return __sycl_std::__invoke_distance<cl::sycl::cl_half>(p0, p1);
10751069
}
1076-
#endif
10771070

10781071
// float length (gengeofloat p)
10791072
template <typename T, typename = typename std::enable_if<
@@ -1089,14 +1082,12 @@ cl::sycl::cl_double length(T p) __NOEXC {
10891082
return __sycl_std::__invoke_length<cl::sycl::cl_double>(p);
10901083
}
10911084

1092-
#ifndef __HALF_NO_ENABLED
10931085
// half length (gengeohalf p)
10941086
template <typename T, typename = typename std::enable_if<
10951087
detail::is_gengeohalf<T>::value, T>::type>
10961088
cl::sycl::cl_half length(T p) __NOEXC {
10971089
return __sycl_std::__invoke_length<cl::sycl::cl_half>(p);
10981090
}
1099-
#endif
11001091

11011092
// gengeofloat normalize (gengeofloat p)
11021093
template <typename T>
@@ -1112,14 +1103,12 @@ normalize(T p) __NOEXC {
11121103
return __sycl_std::__invoke_normalize<T>(p);
11131104
}
11141105

1115-
#ifndef __HALF_NO_ENABLED
11161106
// gengeohalf normalize (gengeohalf p)
11171107
template <typename T>
11181108
typename std::enable_if<detail::is_gengeohalf<T>::value, T>::type
11191109
normalize(T p) __NOEXC {
11201110
return __sycl_std::__invoke_normalize<T>(p);
11211111
}
1122-
#endif
11231112

11241113
// float fast_distance (gengeofloat p0, gengeofloat p1)
11251114
template <typename T, typename = typename std::enable_if<
@@ -1173,7 +1162,7 @@ template <typename T, typename = typename std::enable_if<
11731162
detail::is_genfloat<T>::value, T>::type>
11741163
detail::common_rel_ret_t<T> isequal(T x, T y) __NOEXC {
11751164
return detail::RelConverter<T>::apply(
1176-
__sycl_std::__invoke_OpFOrdEqual<detail::rel_ret_t<T> >(x, y));
1165+
__sycl_std::__invoke_OpFOrdEqual<detail::rel_ret_t<T>>(x, y));
11771166
}
11781167

11791168
// int isnotequal (half x, half y)
@@ -1185,7 +1174,7 @@ template <typename T, typename = typename std::enable_if<
11851174
detail::is_genfloat<T>::value, T>::type>
11861175
detail::common_rel_ret_t<T> isnotequal(T x, T y) __NOEXC {
11871176
return detail::RelConverter<T>::apply(
1188-
__sycl_std::__invoke_OpFUnordNotEqual<detail::rel_ret_t<T> >(x, y));
1177+
__sycl_std::__invoke_OpFUnordNotEqual<detail::rel_ret_t<T>>(x, y));
11891178
}
11901179

11911180
// int isgreater (half x, half y)
@@ -1197,7 +1186,7 @@ template <typename T, typename = typename std::enable_if<
11971186
detail::is_genfloat<T>::value, T>::type>
11981187
detail::common_rel_ret_t<T> isgreater(T x, T y) __NOEXC {
11991188
return detail::RelConverter<T>::apply(
1200-
__sycl_std::__invoke_OpFOrdGreaterThan<detail::rel_ret_t<T> >(x, y));
1189+
__sycl_std::__invoke_OpFOrdGreaterThan<detail::rel_ret_t<T>>(x, y));
12011190
}
12021191

12031192
// int isgreaterequal (half x, half y)
@@ -1209,7 +1198,7 @@ template <typename T, typename = typename std::enable_if<
12091198
detail::is_genfloat<T>::value, T>::type>
12101199
detail::common_rel_ret_t<T> isgreaterequal(T x, T y) __NOEXC {
12111200
return detail::RelConverter<T>::apply(
1212-
__sycl_std::__invoke_OpFOrdGreaterThanEqual<detail::rel_ret_t<T> >(x, y));
1201+
__sycl_std::__invoke_OpFOrdGreaterThanEqual<detail::rel_ret_t<T>>(x, y));
12131202
}
12141203

12151204
// int isless (half x, half y)
@@ -1221,7 +1210,7 @@ template <typename T, typename = typename std::enable_if<
12211210
detail::is_genfloat<T>::value, T>::type>
12221211
detail::common_rel_ret_t<T> isless(T x, T y) __NOEXC {
12231212
return detail::RelConverter<T>::apply(
1224-
__sycl_std::__invoke_OpFOrdLessThan<detail::rel_ret_t<T> >(x, y));
1213+
__sycl_std::__invoke_OpFOrdLessThan<detail::rel_ret_t<T>>(x, y));
12251214
}
12261215

12271216
// int islessequal (half x, half y)
@@ -1233,7 +1222,7 @@ template <typename T, typename = typename std::enable_if<
12331222
detail::is_genfloat<T>::value, T>::type>
12341223
detail::common_rel_ret_t<T> islessequal(T x, T y) __NOEXC {
12351224
return detail::RelConverter<T>::apply(
1236-
__sycl_std::__invoke_OpFOrdLessThanEqual<detail::rel_ret_t<T> >(x, y));
1225+
__sycl_std::__invoke_OpFOrdLessThanEqual<detail::rel_ret_t<T>>(x, y));
12371226
}
12381227

12391228
// int islessgreater (half x, half y)
@@ -1245,7 +1234,7 @@ template <typename T, typename = typename std::enable_if<
12451234
detail::is_genfloat<T>::value, T>::type>
12461235
detail::common_rel_ret_t<T> islessgreater(T x, T y) __NOEXC {
12471236
return detail::RelConverter<T>::apply(
1248-
__sycl_std::__invoke_OpLessOrGreater<detail::rel_ret_t<T> >(x, y));
1237+
__sycl_std::__invoke_OpLessOrGreater<detail::rel_ret_t<T>>(x, y));
12491238
}
12501239

12511240
// int isfinite (half x)
@@ -1257,7 +1246,7 @@ template <typename T, typename = typename std::enable_if<
12571246
detail::is_genfloat<T>::value, T>::type>
12581247
detail::common_rel_ret_t<T> isfinite(T x) __NOEXC {
12591248
return detail::RelConverter<T>::apply(
1260-
__sycl_std::__invoke_OpIsFinite<detail::rel_ret_t<T> >(x));
1249+
__sycl_std::__invoke_OpIsFinite<detail::rel_ret_t<T>>(x));
12611250
}
12621251

12631252
// int isinf (half x)
@@ -1269,7 +1258,7 @@ template <typename T, typename = typename std::enable_if<
12691258
detail::is_genfloat<T>::value, T>::type>
12701259
detail::common_rel_ret_t<T> isinf(T x) __NOEXC {
12711260
return detail::RelConverter<T>::apply(
1272-
__sycl_std::__invoke_OpIsInf<detail::rel_ret_t<T> >(x));
1261+
__sycl_std::__invoke_OpIsInf<detail::rel_ret_t<T>>(x));
12731262
}
12741263

12751264
// int isnan (half x)
@@ -1281,7 +1270,7 @@ template <typename T, typename = typename std::enable_if<
12811270
detail::is_genfloat<T>::value, T>::type>
12821271
detail::common_rel_ret_t<T> isnan(T x) __NOEXC {
12831272
return detail::RelConverter<T>::apply(
1284-
__sycl_std::__invoke_OpIsNan<detail::rel_ret_t<T> >(x));
1273+
__sycl_std::__invoke_OpIsNan<detail::rel_ret_t<T>>(x));
12851274
}
12861275

12871276
// int isnormal (half x)
@@ -1293,7 +1282,7 @@ template <typename T, typename = typename std::enable_if<
12931282
detail::is_genfloat<T>::value, T>::type>
12941283
detail::common_rel_ret_t<T> isnormal(T x) __NOEXC {
12951284
return detail::RelConverter<T>::apply(
1296-
__sycl_std::__invoke_OpIsNormal<detail::rel_ret_t<T> >(x));
1285+
__sycl_std::__invoke_OpIsNormal<detail::rel_ret_t<T>>(x));
12971286
}
12981287

12991288
// int isordered (half x)
@@ -1305,7 +1294,7 @@ template <typename T, typename = typename std::enable_if<
13051294
detail::is_genfloat<T>::value, T>::type>
13061295
detail::common_rel_ret_t<T> isordered(T x, T y) __NOEXC {
13071296
return detail::RelConverter<T>::apply(
1308-
__sycl_std::__invoke_OpOrdered<detail::rel_ret_t<T> >(x, y));
1297+
__sycl_std::__invoke_OpOrdered<detail::rel_ret_t<T>>(x, y));
13091298
}
13101299

13111300
// int isunordered (half x, half y)
@@ -1317,7 +1306,7 @@ template <typename T, typename = typename std::enable_if<
13171306
detail::is_genfloat<T>::value, T>::type>
13181307
detail::common_rel_ret_t<T> isunordered(T x, T y) __NOEXC {
13191308
return detail::RelConverter<T>::apply(
1320-
__sycl_std::__invoke_OpUnordered<detail::rel_ret_t<T> >(x, y));
1309+
__sycl_std::__invoke_OpUnordered<detail::rel_ret_t<T>>(x, y));
13211310
}
13221311

13231312
// int signbit (half x)
@@ -1329,7 +1318,7 @@ template <typename T, typename = typename std::enable_if<
13291318
detail::is_genfloat<T>::value, T>::type>
13301319
detail::common_rel_ret_t<T> signbit(T x) __NOEXC {
13311320
return detail::RelConverter<T>::apply(
1332-
__sycl_std::__invoke_OpSignBitSet<detail::rel_ret_t<T> >(x));
1321+
__sycl_std::__invoke_OpSignBitSet<detail::rel_ret_t<T>>(x));
13331322
}
13341323

13351324
// int any (sigeninteger x)
@@ -1346,7 +1335,7 @@ typename std::enable_if<detail::is_vigeninteger<T>::value,
13461335
cl::sycl::cl_int>::type
13471336
any(T x) __NOEXC {
13481337
return detail::rel_sign_bit_test_ret_t<T>(
1349-
__sycl_std::__invoke_OpAny<detail::rel_sign_bit_test_ret_t<T> >(
1338+
__sycl_std::__invoke_OpAny<detail::rel_sign_bit_test_ret_t<T>>(
13501339
detail::rel_sign_bit_test_arg_t<T>(x)));
13511340
}
13521341

@@ -1364,7 +1353,7 @@ typename std::enable_if<detail::is_vigeninteger<T>::value,
13641353
cl::sycl::cl_int>::type
13651354
all(T x) __NOEXC {
13661355
return detail::rel_sign_bit_test_ret_t<T>(
1367-
__sycl_std::__invoke_OpAll<detail::rel_sign_bit_test_ret_t<T> >(
1356+
__sycl_std::__invoke_OpAll<detail::rel_sign_bit_test_ret_t<T>>(
13681357
detail::rel_sign_bit_test_arg_t<T>(x)));
13691358
}
13701359

@@ -1377,75 +1366,73 @@ bitselect(T a, T b, T c) __NOEXC {
13771366

13781367
// geninteger select (geninteger a, geninteger b, igeninteger c)
13791368
template <typename T, typename T2>
1380-
typename std::enable_if<
1381-
detail::is_geninteger<T>::value &&detail::is_igeninteger<T2>::value,
1382-
T>::type
1369+
typename std::enable_if<detail::is_geninteger<T>::value &&
1370+
detail::is_igeninteger<T2>::value,
1371+
T>::type
13831372
select(T a, T b, T2 c) __NOEXC {
13841373
return __sycl_std::__invoke_OpSelect<T>(detail::select_arg_c_t<T2>(c), b, a);
13851374
}
13861375

13871376
// geninteger select (geninteger a, geninteger b, ugeninteger c)
13881377
template <typename T, typename T2>
1389-
typename std::enable_if<
1390-
detail::is_geninteger<T>::value &&detail::is_ugeninteger<T2>::value,
1391-
T>::type
1378+
typename std::enable_if<detail::is_geninteger<T>::value &&
1379+
detail::is_ugeninteger<T2>::value,
1380+
T>::type
13921381
select(T a, T b, T2 c) __NOEXC {
13931382
return __sycl_std::__invoke_OpSelect<T>(detail::select_arg_c_t<T2>(c), b, a);
13941383
}
13951384

13961385
// genfloatf select (genfloatf a, genfloatf b, genint c)
13971386
template <typename T, typename T2>
13981387
typename std::enable_if<
1399-
detail::is_genfloatf<T>::value &&detail::is_genint<T2>::value, T>::type
1388+
detail::is_genfloatf<T>::value && detail::is_genint<T2>::value, T>::type
14001389
select(T a, T b, T2 c) __NOEXC {
14011390
return __sycl_std::__invoke_OpSelect<T>(detail::select_arg_c_t<T2>(c), b, a);
14021391
}
14031392

14041393
// genfloatf select (genfloatf a, genfloatf b, ugenint c)
14051394
template <typename T, typename T2>
14061395
typename std::enable_if<
1407-
detail::is_genfloatf<T>::value &&detail::is_ugenint<T2>::value, T>::type
1396+
detail::is_genfloatf<T>::value && detail::is_ugenint<T2>::value, T>::type
14081397
select(T a, T b, T2 c) __NOEXC {
14091398
return __sycl_std::__invoke_OpSelect<T>(detail::select_arg_c_t<T2>(c), b, a);
14101399
}
14111400

14121401
// genfloatd select (genfloatd a, genfloatd b, igeninteger64 c)
14131402
template <typename T, typename T2>
1414-
typename std::enable_if<
1415-
detail::is_genfloatd<T>::value &&detail::is_igeninteger64bit<T2>::value,
1416-
T>::type
1403+
typename std::enable_if<detail::is_genfloatd<T>::value &&
1404+
detail::is_igeninteger64bit<T2>::value,
1405+
T>::type
14171406
select(T a, T b, T2 c) __NOEXC {
14181407
return __sycl_std::__invoke_OpSelect<T>(detail::select_arg_c_t<T2>(c), b, a);
14191408
}
14201409

14211410
// genfloatd select (genfloatd a, genfloatd b, ugeninteger64 c)
14221411
template <typename T, typename T2>
1423-
typename std::enable_if<
1424-
detail::is_genfloatd<T>::value &&detail::is_ugeninteger64bit<T2>::value,
1425-
T>::type
1412+
typename std::enable_if<detail::is_genfloatd<T>::value &&
1413+
detail::is_ugeninteger64bit<T2>::value,
1414+
T>::type
14261415
select(T a, T b, T2 c) __NOEXC {
14271416
return __sycl_std::__invoke_OpSelect<T>(detail::select_arg_c_t<T2>(c), b, a);
14281417
}
14291418

1430-
#ifndef __HALF_NO_ENABLED
14311419
// genfloath select (genfloath a, genfloath b, igeninteger16 c)
14321420
template <typename T, typename T2>
1433-
typename std::enable_if<
1434-
detail::is_genfloath<T>::value &&detail::is_igeninteger16bit<T2>::value,
1435-
T>::type
1421+
typename std::enable_if<detail::is_genfloath<T>::value &&
1422+
detail::is_igeninteger16bit<T2>::value,
1423+
T>::type
14361424
select(T a, T b, T2 c) __NOEXC {
14371425
return __sycl_std::__invoke_OpSelect<T>(detail::select_arg_c_t<T2>(c), b, a);
14381426
}
14391427

14401428
// genfloath select (genfloath a, genfloath b, ugeninteger16 c)
14411429
template <typename T, typename T2>
1442-
typename std::enable_if<
1443-
detail::is_genfloath<T>::value &&detail::is_ugeninteger16bit<T2>::value,
1444-
T>::type
1430+
typename std::enable_if<detail::is_genfloath<T>::value &&
1431+
detail::is_ugeninteger16bit<T2>::value,
1432+
T>::type
14451433
select(T a, T b, T2 c) __NOEXC {
14461434
return __sycl_std::__invoke_OpSelect<T>(detail::select_arg_c_t<T2>(c), b, a);
14471435
}
1448-
#endif
14491436

14501437
namespace native {
14511438
/* ----------------- 4.13.3 Math functions. ---------------------------------*/
@@ -1652,6 +1639,4 @@ tan(T x) __NOEXC {
16521639
} // namespace sycl
16531640
} // namespace cl
16541641

1655-
#undef __HALF_NO_ENABLED
16561642
#undef __NOEXC
1657-
#undef __DEVICE_SIDE

0 commit comments

Comments
 (0)