Skip to content

Commit 60ee031

Browse files
garimaguvladimirlaz
garimagu
authored and
vladimirlaz
committed
[SYCL] Support for relational operators
Support for relational operators over the new SYCL math_builtins implementations. Tests have been added for scalar floats and vector float4: scalar_relational.cpp and vector_relational.cpp. Signed-off-by: Vladimir Lazarev <[email protected]> Signed-off-by: Garima Gupta <[email protected]>
1 parent 0baa447 commit 60ee031

File tree

5 files changed

+1519
-99
lines changed

5 files changed

+1519
-99
lines changed

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

Lines changed: 58 additions & 57 deletions
Original file line numberDiff line numberDiff line change
@@ -1220,170 +1220,182 @@ fast_normalize(T p) __NOEXC {
12201220
return __sycl_std::__invoke_fast_normalize<T>(p);
12211221
}
12221222

1223-
/* --------------- 4.13.7 Relational functions. -----------------------------*/
1224-
// igeninteger16bit isequal (genfloath x, genfloath y)
1223+
/* --------------- 4.13.7 Relational functions. Device version --------------*/
1224+
//
1225+
// int isequal (half x, half y)
1226+
// shortn isequal (halfn x, halfn y)
12251227
// igeninteger32bit isequal (genfloatf x, genfloatf y)
1226-
// igeninteger64bit isequal (genfloatd x, genfloatd y)
1228+
// int isequal (double x,double y);
1229+
// longn isequal (doublen x, doublen y)
12271230
template <typename T, typename = typename std::enable_if<
12281231
detail::is_genfloat<T>::value, T>::type>
12291232
typename detail::float_point_to_sign_integral<T>::type isequal(T x,
12301233
T y) __NOEXC {
1231-
__NO_SUPPORT_HOST_VERSION(isequal, T)
12321234
return __sycl_std::__invoke_OpFOrdEqual<
12331235
typename detail::float_point_to_sign_integral<T>::type>(x, y);
12341236
}
12351237

1236-
// igeninteger16bit isnotequal (genfloath x, genfloath y)
1238+
// int isnotequal (half x, half y)
1239+
// shortn isnotequal (halfn x, halfn y)
12371240
// igeninteger32bit isnotequal (genfloatf x, genfloatf y)
1238-
// igeninteger64bit isnotequal (genfloatd x, genfloatd y)
1241+
// int isnotequal (double x, double y)
1242+
// longn isnotequal (doublen x, doublen y)
12391243
template <typename T, typename = typename std::enable_if<
12401244
detail::is_genfloat<T>::value, T>::type>
12411245
typename detail::float_point_to_sign_integral<T>::type isnotequal(T x,
12421246
T y) __NOEXC {
1243-
__NO_SUPPORT_HOST_VERSION(isnotequal, T);
12441247
return __sycl_std::__invoke_OpFUnordNotEqual<
12451248
typename detail::float_point_to_sign_integral<T>::type>(x, y);
12461249
}
12471250

1248-
// igeninteger16bit isgreater (genfloath x, genfloath y)
1251+
// int isgreater (half x, half y)
1252+
// shortn isgreater (halfn x, halfn y)
12491253
// igeninteger32bit isgreater (genfloatf x, genfloatf y)
1250-
// igeninteger64bit isgreater (genfloatd x, genfloatd y)
1254+
// int isgreater (double x, double y)
1255+
// longn isgreater (doublen x, doublen y)
12511256
template <typename T, typename = typename std::enable_if<
12521257
detail::is_genfloat<T>::value, T>::type>
12531258
typename detail::float_point_to_sign_integral<T>::type isgreater(T x,
12541259
T y) __NOEXC {
1255-
__NO_SUPPORT_HOST_VERSION(isgreater, T);
12561260
return __sycl_std::__invoke_OpFOrdGreaterThan<
12571261
typename detail::float_point_to_sign_integral<T>::type>(x, y);
12581262
}
12591263

1260-
// igeninteger16bit isgreaterequal (genfloath x, genfloath y)
1264+
// int isgreaterequal (half x, half y)
1265+
// shortn isgreaterequal (halfn x, halfn y)
12611266
// igeninteger32bit isgreaterequal (genfloatf x, genfloatf y)
1262-
// igeninteger64bit isgreaterequal (genfloatd x, genfloatd y)
1267+
// int isgreaterequal (double x, double y)
1268+
// longn isgreaterequal (doublen x, doublen y)
12631269
template <typename T, typename = typename std::enable_if<
12641270
detail::is_genfloat<T>::value, T>::type>
12651271
typename detail::float_point_to_sign_integral<T>::type
12661272
isgreaterequal(T x, T y) __NOEXC {
1267-
__NO_SUPPORT_HOST_VERSION(isgreaterequal, T);
12681273
return __sycl_std::__invoke_OpFOrdGreaterThanEqual<
12691274
typename detail::float_point_to_sign_integral<T>::type>(x, y);
12701275
}
12711276

1272-
// igeninteger16bit isless (genfloath x, genfloath y)
1277+
// int isless (half x, half y)
1278+
// shortn isless (halfn x, halfn y)
12731279
// igeninteger32bit isless (genfloatf x, genfloatf y)
1274-
// igeninteger64bit isless (genfloatd x, genfloatd y)
1280+
// int isless (long x, long y)
1281+
// longn isless (doublen x, doublen y)
12751282
template <typename T, typename = typename std::enable_if<
12761283
detail::is_genfloat<T>::value, T>::type>
12771284
typename detail::float_point_to_sign_integral<T>::type isless(T x,
12781285
T y) __NOEXC {
1279-
__NO_SUPPORT_HOST_VERSION(isless, T);
12801286
return __sycl_std::__invoke_OpFOrdLessThan<
12811287
typename detail::float_point_to_sign_integral<T>::type>(x, y);
12821288
}
1283-
1284-
// igeninteger16bit islessequal (genfloath x, genfloath y)
1289+
// int islessequal (half x, half y)
1290+
// shortn islessequal (halfn x, halfn y)
12851291
// igeninteger32bit islessequal (genfloatf x, genfloatf y)
1286-
// igeninteger64bit islessequal (genfloatd x, genfloatd y)
1292+
// int islessequal (double x, double y)
1293+
// longn islessequal (doublen x, doublen y)
12871294
template <typename T, typename = typename std::enable_if<
12881295
detail::is_genfloat<T>::value, T>::type>
12891296
typename detail::float_point_to_sign_integral<T>::type
12901297
islessequal(T x, T y) __NOEXC {
1291-
__NO_SUPPORT_HOST_VERSION(islessequal, T);
12921298
return __sycl_std::__invoke_OpFOrdLessThanEqual<
12931299
typename detail::float_point_to_sign_integral<T>::type>(x, y);
12941300
}
12951301

1296-
// igeninteger16bit islessgreater (genfloath x, genfloath y)
1302+
// int islessgreater (half x, half y)
1303+
// shortn islessgreater (halfn x, halfn y)
12971304
// igeninteger32bit islessgreater (genfloatf x, genfloatf y)
1298-
// igeninteger64bit islessgreater (genfloatd x, genfloatd y)
1305+
// int islessgreater (double x, double y)
1306+
// longn islessgreater (doublen x, doublen y)
12991307
template <typename T, typename = typename std::enable_if<
13001308
detail::is_genfloat<T>::value, T>::type>
13011309
typename detail::float_point_to_sign_integral<T>::type
13021310
islessgreater(T x, T y) __NOEXC {
1303-
__NO_SUPPORT_HOST_VERSION(islessgreater, T);
13041311
return __sycl_std::__invoke_OpLessOrGreater<
13051312
typename detail::float_point_to_sign_integral<T>::type>(x, y);
13061313
}
1307-
1308-
// igeninteger16bit isfinite (genfloath x)
1314+
// int isfinite (half x)
1315+
// shortn isfinite (halfn x)
13091316
// igeninteger32bit isfinite (genfloatf x)
1310-
// igeninteger64bit isfinite (genfloatd x)
1317+
// int isfinite (double x)
1318+
// longn isfinite (doublen x)
13111319
template <typename T, typename = typename std::enable_if<
13121320
detail::is_genfloat<T>::value, T>::type>
13131321
typename detail::float_point_to_sign_integral<T>::type isfinite(T x) __NOEXC {
1314-
__NO_SUPPORT_HOST_VERSION(isfinite, T)
13151322
__NO_SUPPORT_DEVICE_VERSION(isfinite, T)
13161323
return __sycl_std::__invoke_OpIsFinite<
13171324
typename detail::float_point_to_sign_integral<T>::type>(x);
13181325
}
13191326

1320-
// igeninteger16bit isinf (genfloath x)
1327+
// int isinf (half x)
1328+
// shortn isinf (halfn x)
13211329
// igeninteger32bit isinf (genfloatf x)
1322-
// igeninteger64bit isinf (genfloatd x)
1330+
// int isinf (double x)
1331+
// longn isinf (doublen x)
13231332
template <typename T, typename = typename std::enable_if<
13241333
detail::is_genfloat<T>::value, T>::type>
13251334
typename detail::float_point_to_sign_integral<T>::type isinf(T x) __NOEXC {
1326-
__NO_SUPPORT_HOST_VERSION(isinf, T)
13271335
__NO_SUPPORT_DEVICE_VERSION(isinf, T)
13281336
return __sycl_std::__invoke_OpIsInf<
13291337
typename detail::float_point_to_sign_integral<T>::type>(x);
13301338
}
13311339

1332-
// igeninteger16bit isnan (genfloath x)
1340+
// int isnan (half x)
1341+
// shortn isnan (halfn x)
13331342
// igeninteger32bit isnan (genfloatf x)
1334-
// igeninteger64bit isnan (genfloatd x)
1343+
// int isnan (double x)
1344+
// longn isnan (doublen x)
13351345
template <typename T, typename = typename std::enable_if<
13361346
detail::is_genfloat<T>::value, T>::type>
13371347
typename detail::float_point_to_sign_integral<T>::type isnan(T x) __NOEXC {
1338-
__NO_SUPPORT_HOST_VERSION(isnan, T)
13391348
__NO_SUPPORT_DEVICE_VERSION(isnan, T)
13401349
return __sycl_std::__invoke_OpIsNan<
13411350
typename detail::float_point_to_sign_integral<T>::type>(x);
13421351
}
1343-
1344-
// igeninteger16bit isnormal (genfloath x)
1352+
// int isnormal (half x)
1353+
// shortn isnormal (halfn x)
13451354
// igeninteger32bit isnormal (genfloatf x)
1346-
// igeninteger64bit isnormal (genfloatd x)
1355+
// int isnormal (double x)
1356+
// longn isnormal (doublen x)
13471357
template <typename T, typename = typename std::enable_if<
13481358
detail::is_genfloat<T>::value, T>::type>
13491359
typename detail::float_point_to_sign_integral<T>::type isnormal(T x) __NOEXC {
1350-
__NO_SUPPORT_HOST_VERSION(isnormal, T)
13511360
__NO_SUPPORT_DEVICE_VERSION(isnormal, T)
13521361
return __sycl_std::__invoke_OpIsNormal<
13531362
typename detail::float_point_to_sign_integral<T>::type>(x);
13541363
}
13551364

1356-
// igeninteger16bit isordered (genfloath x, genfloath y)
1365+
// int isordered (half x)
1366+
// shortn isordered (halfn x, halfn y)
13571367
// igeninteger32bit isordered (genfloatf x, genfloatf y)
1358-
// igeninteger64bit isordered (genfloatd x, genfloatd y)
1368+
// int isordered (double x, double y)
1369+
// longn isordered (doublen x, doublen y)
13591370
template <typename T, typename = typename std::enable_if<
13601371
detail::is_genfloat<T>::value, T>::type>
13611372
typename detail::float_point_to_sign_integral<T>::type isordered(T x,
13621373
T y) __NOEXC {
1363-
__NO_SUPPORT_HOST_VERSION(isordered, T)
13641374
return __sycl_std::__invoke_OpOrdered<
13651375
typename detail::float_point_to_sign_integral<T>::type>(x, y);
13661376
}
13671377

1368-
// igeninteger16bit isunordered (genfloath x, genfloath y)
1378+
// int isunordered (half x, half y)
1379+
// shortn isunordered (halfn x, halfn y)
13691380
// igeninteger32bit isunordered (genfloatf x, genfloatf y)
1370-
// igeninteger64bit isunordered (genfloatd x, genfloatd y)
1381+
// int isunordered (double x, double y)
1382+
// longn isunordered (doublen x, doublen y)
13711383
template <typename T, typename = typename std::enable_if<
13721384
detail::is_genfloat<T>::value, T>::type>
13731385
typename detail::float_point_to_sign_integral<T>::type
13741386
isunordered(T x, T y) __NOEXC {
1375-
__NO_SUPPORT_HOST_VERSION(isunordered, T)
13761387
return __sycl_std::__invoke_OpUnordered<
13771388
typename detail::float_point_to_sign_integral<T>::type>(x, y);
13781389
}
13791390

1380-
// igeninteger16bit signbit (genfloath x)
1391+
// int signbit (half x)
1392+
// shortn signbit (halfn x)
13811393
// igeninteger32bit signbit (genfloatf x)
1382-
// igeninteger64bit signbit (genfloatd x)
1394+
// int signbit (double)
1395+
// longn signbit (doublen x)
13831396
template <typename T, typename = typename std::enable_if<
13841397
detail::is_genfloat<T>::value, T>::type>
13851398
typename detail::float_point_to_sign_integral<T>::type signbit(T x) __NOEXC {
1386-
__NO_SUPPORT_HOST_VERSION(signbit, T)
13871399
__NO_SUPPORT_DEVICE_VERSION(signbit, T)
13881400
return __sycl_std::__invoke_OpSignBitSet<
13891401
typename detail::float_point_to_sign_integral<T>::type>(x);
@@ -1393,7 +1405,6 @@ typename detail::float_point_to_sign_integral<T>::type signbit(T x) __NOEXC {
13931405
template <typename T, typename = typename std::enable_if<
13941406
detail::is_igeninteger<T>::value, T>::type>
13951407
cl::sycl::cl_int any(T x) __NOEXC {
1396-
__NO_SUPPORT_HOST_VERSION(signbit, T)
13971408
__NO_SUPPORT_DEVICE_VERSION(signbit, T)
13981409
return __sycl_std::__invoke_OpAny<cl::sycl::cl_int>(x);
13991410
}
@@ -1402,7 +1413,6 @@ cl::sycl::cl_int any(T x) __NOEXC {
14021413
template <typename T, typename = typename std::enable_if<
14031414
detail::is_igeninteger<T>::value, T>::type>
14041415
cl::sycl::cl_int all(T x) __NOEXC {
1405-
__NO_SUPPORT_HOST_VERSION(all, T)
14061416
__NO_SUPPORT_DEVICE_VERSION(all, T)
14071417
return __sycl_std::__invoke_OpAll<cl::sycl::cl_int>(x);
14081418
}
@@ -1411,7 +1421,6 @@ cl::sycl::cl_int all(T x) __NOEXC {
14111421
template <typename T>
14121422
typename std::enable_if<detail::is_gentype<T>::value, T>::type
14131423
bitselect(T a, T b, T c) __NOEXC {
1414-
__NO_SUPPORT_HOST_VERSION(bitselect, T)
14151424
return __sycl_std::__invoke_bitselect<T>(a, b, c);
14161425
}
14171426

@@ -1421,7 +1430,6 @@ typename std::enable_if<detail::is_geninteger<T>::value &&
14211430
detail::is_igeninteger<T2>::value,
14221431
T>::type
14231432
select(T a, T b, T2 c) __NOEXC {
1424-
__NO_SUPPORT_HOST_VERSION(select, T)
14251433
__NO_SUPPORT_DEVICE_VERSION(select, T)
14261434
return __sycl_std::__invoke_OpSelect<T>(a, b, c);
14271435
}
@@ -1432,7 +1440,6 @@ typename std::enable_if<detail::is_geninteger<T>::value &&
14321440
detail::is_ugeninteger<T2>::value,
14331441
T>::type
14341442
select(T a, T b, T2 c) __NOEXC {
1435-
__NO_SUPPORT_HOST_VERSION(select, T)
14361443
__NO_SUPPORT_DEVICE_VERSION(select, T)
14371444
return __sycl_std::__invoke_OpSelect<T>(a, b, c);
14381445
}
@@ -1442,7 +1449,6 @@ template <typename T, typename T2>
14421449
typename std::enable_if<
14431450
detail::is_genfloatf<T>::value && detail::is_genint<T2>::value, T>::type
14441451
select(T a, T b, T2 c) __NOEXC {
1445-
__NO_SUPPORT_HOST_VERSION(select, T)
14461452
__NO_SUPPORT_DEVICE_VERSION(select, T)
14471453
return __sycl_std::__invoke_OpSelect<T>(a, b, c);
14481454
}
@@ -1452,7 +1458,6 @@ template <typename T, typename T2>
14521458
typename std::enable_if<
14531459
detail::is_genfloatf<T>::value && detail::is_ugenint<T2>::value, T>::type
14541460
select(T a, T b, T2 c) __NOEXC {
1455-
__NO_SUPPORT_HOST_VERSION(select, T)
14561461
__NO_SUPPORT_DEVICE_VERSION(select, T)
14571462
return __sycl_std::__invoke_OpSelect<T>(a, b, c);
14581463
}
@@ -1463,7 +1468,6 @@ typename std::enable_if<detail::is_genfloatd<T>::value &&
14631468
detail::is_igeninteger64bit<T2>::value,
14641469
T>::type
14651470
select(T a, T b, T2 c) __NOEXC {
1466-
__NO_SUPPORT_HOST_VERSION(select, T)
14671471
__NO_SUPPORT_DEVICE_VERSION(select, T)
14681472
return __sycl_std::__invoke_OpSelect<T>(a, b, c);
14691473
}
@@ -1474,7 +1478,6 @@ typename std::enable_if<detail::is_genfloatd<T>::value &&
14741478
detail::is_ugeninteger64bit<T2>::value,
14751479
T>::type
14761480
select(T a, T b, T2 c) __NOEXC {
1477-
__NO_SUPPORT_HOST_VERSION(select, T)
14781481
__NO_SUPPORT_DEVICE_VERSION(select, T)
14791482
return __sycl_std::__invoke_OpSelect<T>(a, b, c);
14801483
}
@@ -1486,7 +1489,6 @@ typename std::enable_if<detail::is_genfloath<T>::value &&
14861489
detail::is_igeninteger64bit<T2>::value,
14871490
T>::type
14881491
select(T a, T b, T2 c) __NOEXC {
1489-
__NO_SUPPORT_HOST_VERSION(select, T)
14901492
__NO_SUPPORT_DEVICE_VERSION(select, T)
14911493
return __sycl_std::__invoke_OpSelect<T>(a, b, c);
14921494
}
@@ -1497,7 +1499,6 @@ typename std::enable_if<detail::is_genfloath<T>::value &&
14971499
detail::is_ugeninteger64bit<T2>::value,
14981500
T>::type
14991501
select(T a, T b, T2 c) __NOEXC {
1500-
__NO_SUPPORT_HOST_VERSION(select, T)
15011502
__NO_SUPPORT_DEVICE_VERSION(select, T)
15021503
return __sycl_std::__invoke_OpSelect<T>(a, b, c);
15031504
}

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

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -587,7 +587,7 @@ template <> struct float_point_to_sign_integral<cl_float16> {
587587

588588
#ifndef __HALF_NO_ENABLED
589589
template <> struct float_point_to_sign_integral<cl_half> {
590-
using type = cl_short;
590+
using type = cl_int;
591591
};
592592
template <> struct float_point_to_sign_integral<cl_half2> {
593593
using type = cl_short2;
@@ -607,7 +607,7 @@ template <> struct float_point_to_sign_integral<cl_half16> {
607607
#endif
608608

609609
template <> struct float_point_to_sign_integral<cl_double> {
610-
using type = cl_long;
610+
using type = cl_int;
611611
};
612612
template <> struct float_point_to_sign_integral<cl_double2> {
613613
using type = cl_long2;
@@ -711,4 +711,4 @@ template <> struct make_upper<cl::sycl::cl_uint> {
711711
} // namespace sycl
712712
} // namespace cl
713713

714-
#undef __HALF_NO_ENABLED
714+
#undef __HALF_NO_ENABLED

0 commit comments

Comments
 (0)