Skip to content

Commit c274f40

Browse files
dm-vodopyanovdm-vodopyanov
authored
[SYCL] Add marray support to integer built-in functions (#8861)
This patch adds support of sycl::marray to integer built-in functions (SYCL 2020 spec, Table 178).
1 parent b9b8513 commit c274f40

File tree

4 files changed

+459
-15
lines changed

4 files changed

+459
-15
lines changed

‎sycl/include/sycl/builtins.hpp

+270-11
Original file line numberDiff line numberDiff line change
@@ -933,15 +933,16 @@ __SYCL_MARRAY_COMMON_FUNCTION_UNOP_OVERLOAD(sign, T x, x[i])
933933
// errors.
934934
__SYCL_MARRAY_COMMON_FUNCTION_BINOP_OVERLOAD((min), T x, T y, x[i], y[i])
935935
__SYCL_MARRAY_COMMON_FUNCTION_BINOP_OVERLOAD((min), T x,
936-
detail::marray_element_type<T> y,
936+
detail::marray_element_t<T> y,
937937
x[i], y)
938938
__SYCL_MARRAY_COMMON_FUNCTION_BINOP_OVERLOAD((max), T x, T y, x[i], y[i])
939939
__SYCL_MARRAY_COMMON_FUNCTION_BINOP_OVERLOAD((max), T x,
940-
detail::marray_element_type<T> y,
940+
detail::marray_element_t<T> y,
941941
x[i], y)
942942
__SYCL_MARRAY_COMMON_FUNCTION_BINOP_OVERLOAD(step, T edge, T x, edge[i], x[i])
943-
__SYCL_MARRAY_COMMON_FUNCTION_BINOP_OVERLOAD(
944-
step, detail::marray_element_type<T> edge, T x, edge, x[i])
943+
__SYCL_MARRAY_COMMON_FUNCTION_BINOP_OVERLOAD(step,
944+
detail::marray_element_t<T> edge,
945+
T x, edge, x[i])
945946

946947
#undef __SYCL_MARRAY_COMMON_FUNCTION_BINOP_OVERLOAD
947948

@@ -956,18 +957,18 @@ __SYCL_MARRAY_COMMON_FUNCTION_BINOP_OVERLOAD(
956957
__SYCL_MARRAY_COMMON_FUNCTION_TEROP_OVERLOAD(clamp, T x, T minval, T maxval,
957958
x[i], minval[i], maxval[i])
958959
__SYCL_MARRAY_COMMON_FUNCTION_TEROP_OVERLOAD(
959-
clamp, T x, detail::marray_element_type<T> minval,
960-
detail::marray_element_type<T> maxval, x[i], minval, maxval)
960+
clamp, T x, detail::marray_element_t<T> minval,
961+
detail::marray_element_t<T> maxval, x[i], minval, maxval)
961962
__SYCL_MARRAY_COMMON_FUNCTION_TEROP_OVERLOAD(mix, T x, T y, T a, x[i], y[i],
962963
a[i])
963964
__SYCL_MARRAY_COMMON_FUNCTION_TEROP_OVERLOAD(mix, T x, T y,
964-
detail::marray_element_type<T> a,
965+
detail::marray_element_t<T> a,
965966
x[i], y[i], a)
966967
__SYCL_MARRAY_COMMON_FUNCTION_TEROP_OVERLOAD(smoothstep, T edge0, T edge1, T x,
967968
edge0[i], edge1[i], x[i])
968969
__SYCL_MARRAY_COMMON_FUNCTION_TEROP_OVERLOAD(
969-
smoothstep, detail::marray_element_type<T> edge0,
970-
detail::marray_element_type<T> edge1, T x, edge0, edge1, x[i])
970+
smoothstep, detail::marray_element_t<T> edge0,
971+
detail::marray_element_t<T> edge1, T x, edge0, edge1, x[i])
971972

972973
#undef __SYCL_MARRAY_COMMON_FUNCTION_TEROP_OVERLOAD
973974
#undef __SYCL_MARRAY_COMMON_FUNCTION_OVERLOAD_IMPL
@@ -1310,6 +1311,264 @@ mul24(T x, T y) __NOEXC {
13101311
return __sycl_std::__invoke_u_mul24<T>(x, y);
13111312
}
13121313

1314+
// marray integer functions
1315+
1316+
// TODO: can be optimized in the way math functions are optimized (usage of
1317+
// vec<T, 2>)
1318+
#define __SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, ...) \
1319+
marray<T, N> res; \
1320+
for (int j = 0; j < N; j++) { \
1321+
res[j] = NAME(__VA_ARGS__); \
1322+
} \
1323+
return res;
1324+
1325+
// Keep NAME for readability
1326+
#define __SYCL_MARRAY_INTEGER_FUNCTION_ABS_U_OVERLOAD(NAME, ARG, ...) \
1327+
template <typename T, size_t N> \
1328+
std::enable_if_t<detail::is_ugeninteger<T>::value, marray<T, N>> NAME( \
1329+
marray<T, N> ARG) __NOEXC { \
1330+
__SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
1331+
}
1332+
1333+
#define __SYCL_MARRAY_INTEGER_FUNCTION_ABS_I_OVERLOAD(NAME, ARG, ...) \
1334+
template <typename T, size_t N> \
1335+
std::enable_if_t<detail::is_igeninteger<T>::value, \
1336+
marray<detail::make_unsigned_t<T>, N>> \
1337+
NAME(marray<T, N> ARG) __NOEXC { \
1338+
__SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
1339+
}
1340+
1341+
__SYCL_MARRAY_INTEGER_FUNCTION_ABS_U_OVERLOAD(abs, x, x[j])
1342+
__SYCL_MARRAY_INTEGER_FUNCTION_ABS_I_OVERLOAD(abs, x, x[j])
1343+
1344+
#undef __SYCL_MARRAY_INTEGER_FUNCTION_ABS_I_OVERLOAD
1345+
#undef __SYCL_MARRAY_INTEGER_FUNCTION_ABS_U_OVERLOAD
1346+
1347+
#define __SYCL_MARRAY_INTEGER_FUNCTION_UNOP_OVERLOAD(NAME, ARG, ...) \
1348+
template <typename T, size_t N> \
1349+
std::enable_if_t<detail::is_geninteger<T>::value, marray<T, N>> NAME( \
1350+
marray<T, N> ARG) __NOEXC { \
1351+
__SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
1352+
}
1353+
1354+
__SYCL_MARRAY_INTEGER_FUNCTION_UNOP_OVERLOAD(clz, x, x[j])
1355+
__SYCL_MARRAY_INTEGER_FUNCTION_UNOP_OVERLOAD(ctz, x, x[j])
1356+
__SYCL_MARRAY_INTEGER_FUNCTION_UNOP_OVERLOAD(popcount, x, x[j])
1357+
1358+
#undef __SYCL_MARRAY_INTEGER_FUNCTION_UNOP_OVERLOAD
1359+
1360+
#define __SYCL_MARRAY_INTEGER_FUNCTION_BINOP_U_OVERLOAD(NAME, ARG1, ARG2, ...) \
1361+
template <typename T, size_t N> \
1362+
std::enable_if_t<detail::is_ugeninteger<T>::value, marray<T, N>> NAME( \
1363+
marray<T, N> ARG1, marray<T, N> ARG2) __NOEXC { \
1364+
__SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
1365+
}
1366+
1367+
#define __SYCL_MARRAY_INTEGER_FUNCTION_BINOP_I_RET_U_OVERLOAD(NAME, ARG1, \
1368+
ARG2, ...) \
1369+
template <typename T, size_t N> \
1370+
std::enable_if_t<detail::is_igeninteger<T>::value, \
1371+
marray<detail::make_unsigned_t<T>, N>> \
1372+
NAME(marray<T, N> ARG1, marray<T, N> ARG2) __NOEXC { \
1373+
__SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
1374+
}
1375+
1376+
#define __SYCL_MARRAY_INTEGER_FUNCTION_BINOP_I_OVERLOAD(NAME, ARG1, ARG2, ...) \
1377+
template <typename T, size_t N> \
1378+
std::enable_if_t<detail::is_igeninteger<T>::value, marray<T, N>> NAME( \
1379+
marray<T, N> ARG1, marray<T, N> ARG2) __NOEXC { \
1380+
__SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
1381+
}
1382+
1383+
#define __SYCL_MARRAY_INTEGER_FUNCTION_BINOP_U_2ND_ARG_SCALAR_OVERLOAD( \
1384+
NAME, ARG1, ARG2, ...) \
1385+
template <typename T, size_t N> \
1386+
std::enable_if_t<detail::is_ugeninteger<T>::value, marray<T, N>> NAME( \
1387+
marray<T, N> ARG1, T ARG2) __NOEXC { \
1388+
__SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
1389+
}
1390+
1391+
#define __SYCL_MARRAY_INTEGER_FUNCTION_BINOP_I_2ND_ARG_SCALAR_OVERLOAD( \
1392+
NAME, ARG1, ARG2, ...) \
1393+
template <typename T, size_t N> \
1394+
std::enable_if_t<detail::is_igeninteger<T>::value, marray<T, N>> NAME( \
1395+
marray<T, N> ARG1, T ARG2) __NOEXC { \
1396+
__SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
1397+
}
1398+
1399+
__SYCL_MARRAY_INTEGER_FUNCTION_BINOP_U_OVERLOAD(abs_diff, x, y, x[j], y[j])
1400+
__SYCL_MARRAY_INTEGER_FUNCTION_BINOP_I_RET_U_OVERLOAD(abs_diff, x, y, x[j],
1401+
y[j])
1402+
__SYCL_MARRAY_INTEGER_FUNCTION_BINOP_U_OVERLOAD(add_sat, x, y, x[j], y[j])
1403+
__SYCL_MARRAY_INTEGER_FUNCTION_BINOP_I_OVERLOAD(add_sat, x, y, x[j], y[j])
1404+
__SYCL_MARRAY_INTEGER_FUNCTION_BINOP_U_OVERLOAD(hadd, x, y, x[j], y[j])
1405+
__SYCL_MARRAY_INTEGER_FUNCTION_BINOP_I_OVERLOAD(hadd, x, y, x[j], y[j])
1406+
__SYCL_MARRAY_INTEGER_FUNCTION_BINOP_U_OVERLOAD(rhadd, x, y, x[j], y[j])
1407+
__SYCL_MARRAY_INTEGER_FUNCTION_BINOP_I_OVERLOAD(rhadd, x, y, x[j], y[j])
1408+
__SYCL_MARRAY_INTEGER_FUNCTION_BINOP_U_OVERLOAD((max), x, y, x[j], y[j])
1409+
__SYCL_MARRAY_INTEGER_FUNCTION_BINOP_I_OVERLOAD((max), x, y, x[j], y[j])
1410+
__SYCL_MARRAY_INTEGER_FUNCTION_BINOP_U_2ND_ARG_SCALAR_OVERLOAD((max), x, y,
1411+
x[j], y)
1412+
__SYCL_MARRAY_INTEGER_FUNCTION_BINOP_I_2ND_ARG_SCALAR_OVERLOAD((max), x, y,
1413+
x[j], y)
1414+
__SYCL_MARRAY_INTEGER_FUNCTION_BINOP_U_OVERLOAD((min), x, y, x[j], y[j])
1415+
__SYCL_MARRAY_INTEGER_FUNCTION_BINOP_I_OVERLOAD((min), x, y, x[j], y[j])
1416+
__SYCL_MARRAY_INTEGER_FUNCTION_BINOP_U_2ND_ARG_SCALAR_OVERLOAD((min), x, y,
1417+
x[j], y)
1418+
__SYCL_MARRAY_INTEGER_FUNCTION_BINOP_I_2ND_ARG_SCALAR_OVERLOAD((min), x, y,
1419+
x[j], y)
1420+
__SYCL_MARRAY_INTEGER_FUNCTION_BINOP_U_OVERLOAD(mul_hi, x, y, x[j], y[j])
1421+
__SYCL_MARRAY_INTEGER_FUNCTION_BINOP_I_OVERLOAD(mul_hi, x, y, x[j], y[j])
1422+
__SYCL_MARRAY_INTEGER_FUNCTION_BINOP_U_OVERLOAD(rotate, v, i, v[j], i[j])
1423+
__SYCL_MARRAY_INTEGER_FUNCTION_BINOP_I_OVERLOAD(rotate, v, i, v[j], i[j])
1424+
__SYCL_MARRAY_INTEGER_FUNCTION_BINOP_U_OVERLOAD(sub_sat, x, y, x[j], y[j])
1425+
__SYCL_MARRAY_INTEGER_FUNCTION_BINOP_I_OVERLOAD(sub_sat, x, y, x[j], y[j])
1426+
1427+
#undef __SYCL_MARRAY_INTEGER_FUNCTION_BINOP_U_2ND_ARG_SCALAR_OVERLOAD
1428+
#undef __SYCL_MARRAY_INTEGER_FUNCTION_BINOP_I_2ND_ARG_SCALAR_OVERLOAD
1429+
#undef __SYCL_MARRAY_INTEGER_FUNCTION_BINOP_I_OVERLOAD
1430+
#undef __SYCL_MARRAY_INTEGER_FUNCTION_BINOP_I_RET_U_OVERLOAD
1431+
#undef __SYCL_MARRAY_INTEGER_FUNCTION_BINOP_U_OVERLOAD
1432+
1433+
#define __SYCL_MARRAY_INTEGER_FUNCTION_TEROP_U_OVERLOAD(NAME, ARG1, ARG2, \
1434+
ARG3, ...) \
1435+
template <typename T, size_t N> \
1436+
std::enable_if_t<detail::is_ugeninteger<T>::value, marray<T, N>> NAME( \
1437+
marray<T, N> ARG1, marray<T, N> ARG2, marray<T, N> ARG3) __NOEXC { \
1438+
__SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
1439+
}
1440+
1441+
#define __SYCL_MARRAY_INTEGER_FUNCTION_TEROP_I_OVERLOAD(NAME, ARG1, ARG2, \
1442+
ARG3, ...) \
1443+
template <typename T, size_t N> \
1444+
std::enable_if_t<detail::is_igeninteger<T>::value, marray<T, N>> NAME( \
1445+
marray<T, N> ARG1, marray<T, N> ARG2, marray<T, N> ARG3) __NOEXC { \
1446+
__SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
1447+
}
1448+
1449+
#define __SYCL_MARRAY_INTEGER_FUNCTION_TEROP_U_2ND_3RD_ARGS_SCALAR_OVERLOAD( \
1450+
NAME, ARG1, ARG2, ARG3, ...) \
1451+
template <typename T, size_t N> \
1452+
std::enable_if_t<detail::is_ugeninteger<T>::value, marray<T, N>> NAME( \
1453+
marray<T, N> ARG1, T ARG2, T ARG3) __NOEXC { \
1454+
__SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
1455+
}
1456+
1457+
#define __SYCL_MARRAY_INTEGER_FUNCTION_TEROP_I_2ND_3RD_ARGS_SCALAR_OVERLOAD( \
1458+
NAME, ARG1, ARG2, ARG3, ...) \
1459+
template <typename T, size_t N> \
1460+
std::enable_if_t<detail::is_igeninteger<T>::value, marray<T, N>> NAME( \
1461+
marray<T, N> ARG1, T ARG2, T ARG3) __NOEXC { \
1462+
__SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
1463+
}
1464+
1465+
__SYCL_MARRAY_INTEGER_FUNCTION_TEROP_U_OVERLOAD(clamp, x, minval, maxval, x[j],
1466+
minval[j], maxval[j])
1467+
__SYCL_MARRAY_INTEGER_FUNCTION_TEROP_I_OVERLOAD(clamp, x, minval, maxval, x[j],
1468+
minval[j], maxval[j])
1469+
__SYCL_MARRAY_INTEGER_FUNCTION_TEROP_U_2ND_3RD_ARGS_SCALAR_OVERLOAD(
1470+
clamp, x, minval, maxval, x[j], minval, maxval)
1471+
__SYCL_MARRAY_INTEGER_FUNCTION_TEROP_I_2ND_3RD_ARGS_SCALAR_OVERLOAD(
1472+
clamp, x, minval, maxval, x[j], minval, maxval)
1473+
__SYCL_MARRAY_INTEGER_FUNCTION_TEROP_U_OVERLOAD(mad_hi, a, b, c, a[j], b[j],
1474+
c[j])
1475+
__SYCL_MARRAY_INTEGER_FUNCTION_TEROP_I_OVERLOAD(mad_hi, a, b, c, a[j], b[j],
1476+
c[j])
1477+
__SYCL_MARRAY_INTEGER_FUNCTION_TEROP_U_OVERLOAD(mad_sat, a, b, c, a[j], b[j],
1478+
c[j])
1479+
__SYCL_MARRAY_INTEGER_FUNCTION_TEROP_I_OVERLOAD(mad_sat, a, b, c, a[j], b[j],
1480+
c[j])
1481+
1482+
#undef __SYCL_MARRAY_INTEGER_FUNCTION_TEROP_U_2ND_3RD_ARGS_SCALAR_OVERLOAD
1483+
#undef __SYCL_MARRAY_INTEGER_FUNCTION_TEROP_I_2ND_3RD_ARGS_SCALAR_OVERLOAD
1484+
#undef __SYCL_MARRAY_INTEGER_FUNCTION_TEROP_I_OVERLOAD
1485+
#undef __SYCL_MARRAY_INTEGER_FUNCTION_TEROP_U_OVERLOAD
1486+
1487+
// Keep NAME for readability
1488+
#define __SYCL_MARRAY_INTEGER_FUNCTION_MAD24_U_OVERLOAD(NAME, ARG1, ARG2, \
1489+
ARG3, ...) \
1490+
template <typename T, size_t N> \
1491+
std::enable_if_t<detail::is_ugeninteger32bit<T>::value, marray<T, N>> NAME( \
1492+
marray<T, N> ARG1, marray<T, N> ARG2, marray<T, N> ARG3) __NOEXC { \
1493+
__SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
1494+
}
1495+
1496+
#define __SYCL_MARRAY_INTEGER_FUNCTION_MAD24_I_OVERLOAD(NAME, ARG1, ARG2, \
1497+
ARG3, ...) \
1498+
template <typename T, size_t N> \
1499+
std::enable_if_t<detail::is_igeninteger32bit<T>::value, marray<T, N>> NAME( \
1500+
marray<T, N> ARG1, marray<T, N> ARG2, marray<T, N> ARG3) __NOEXC { \
1501+
__SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
1502+
}
1503+
1504+
__SYCL_MARRAY_INTEGER_FUNCTION_MAD24_U_OVERLOAD(mad24, x, y, z, x[j], y[j],
1505+
z[j])
1506+
__SYCL_MARRAY_INTEGER_FUNCTION_MAD24_I_OVERLOAD(mad24, x, y, z, x[j], y[j],
1507+
z[j])
1508+
1509+
#undef __SYCL_MARRAY_INTEGER_FUNCTION_MAD24_I_OVERLOAD
1510+
#undef __SYCL_MARRAY_INTEGER_FUNCTION_MAD24_U_OVERLOAD
1511+
1512+
// Keep NAME for readability
1513+
#define __SYCL_MARRAY_INTEGER_FUNCTION_MUL24_U_OVERLOAD(NAME, ARG1, ARG2, ...) \
1514+
template <typename T, size_t N> \
1515+
std::enable_if_t<detail::is_ugeninteger32bit<T>::value, marray<T, N>> NAME( \
1516+
marray<T, N> ARG1, marray<T, N> ARG2) __NOEXC { \
1517+
__SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
1518+
}
1519+
1520+
#define __SYCL_MARRAY_INTEGER_FUNCTION_MUL24_I_OVERLOAD(NAME, ARG1, ARG2, ...) \
1521+
template <typename T, size_t N> \
1522+
std::enable_if_t<detail::is_igeninteger32bit<T>::value, marray<T, N>> NAME( \
1523+
marray<T, N> ARG1, marray<T, N> ARG2) __NOEXC { \
1524+
__SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
1525+
}
1526+
1527+
__SYCL_MARRAY_INTEGER_FUNCTION_MUL24_U_OVERLOAD(mul24, x, y, x[j], y[j])
1528+
__SYCL_MARRAY_INTEGER_FUNCTION_MUL24_I_OVERLOAD(mul24, x, y, x[j], y[j])
1529+
1530+
#undef __SYCL_MARRAY_INTEGER_FUNCTION_MUL24_I_OVERLOAD
1531+
#undef __SYCL_MARRAY_INTEGER_FUNCTION_MUL24_U_OVERLOAD
1532+
#undef __SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL
1533+
1534+
// TODO: can be optimized in the way math functions are optimized (usage of
1535+
// vec<T, 2>)
1536+
#define __SYCL_MARRAY_INTEGER_FUNCTION_UPSAMPLE_OVERLOAD_IMPL(NAME) \
1537+
detail::make_larger_t<marray<T, N>> res; \
1538+
for (int j = 0; j < N; j++) { \
1539+
res[j] = NAME(hi[j], lo[j]); \
1540+
} \
1541+
return res;
1542+
1543+
// Keep NAME for readability
1544+
#define __SYCL_MARRAY_INTEGER_FUNCTION_UPSAMPLE_UU_OVERLOAD(NAME, KBIT) \
1545+
template <typename T, size_t N> \
1546+
std::enable_if_t<detail::is_ugeninteger##KBIT<T>::value, \
1547+
detail::make_larger_t<marray<T, N>>> \
1548+
NAME(marray<T, N> hi, marray<T, N> lo) __NOEXC { \
1549+
__SYCL_MARRAY_INTEGER_FUNCTION_UPSAMPLE_OVERLOAD_IMPL(NAME) \
1550+
}
1551+
1552+
#define __SYCL_MARRAY_INTEGER_FUNCTION_UPSAMPLE_IU_OVERLOAD(NAME, KBIT) \
1553+
template <typename T, typename T2, size_t N> \
1554+
std::enable_if_t<detail::is_igeninteger##KBIT<T>::value && \
1555+
detail::is_ugeninteger##KBIT<T2>::value, \
1556+
detail::make_larger_t<marray<T, N>>> \
1557+
NAME(marray<T, N> hi, marray<T2, N> lo) __NOEXC { \
1558+
__SYCL_MARRAY_INTEGER_FUNCTION_UPSAMPLE_OVERLOAD_IMPL(NAME) \
1559+
}
1560+
1561+
__SYCL_MARRAY_INTEGER_FUNCTION_UPSAMPLE_UU_OVERLOAD(upsample, 8bit)
1562+
__SYCL_MARRAY_INTEGER_FUNCTION_UPSAMPLE_IU_OVERLOAD(upsample, 8bit)
1563+
__SYCL_MARRAY_INTEGER_FUNCTION_UPSAMPLE_UU_OVERLOAD(upsample, 16bit)
1564+
__SYCL_MARRAY_INTEGER_FUNCTION_UPSAMPLE_IU_OVERLOAD(upsample, 16bit)
1565+
__SYCL_MARRAY_INTEGER_FUNCTION_UPSAMPLE_UU_OVERLOAD(upsample, 32bit)
1566+
__SYCL_MARRAY_INTEGER_FUNCTION_UPSAMPLE_IU_OVERLOAD(upsample, 32bit)
1567+
1568+
#undef __SYCL_MARRAY_INTEGER_FUNCTION_UPSAMPLE_IU_OVERLOAD
1569+
#undef __SYCL_MARRAY_INTEGER_FUNCTION_UPSAMPLE_UU_OVERLOAD
1570+
#undef __SYCL_MARRAY_INTEGER_FUNCTION_UPSAMPLE_OVERLOAD_IMPL
1571+
13131572
/* --------------- 4.13.6 Geometric Functions. ------------------------------*/
13141573
// float3 cross (float3 p0, float3 p1)
13151574
// float4 cross (float4 p0, float4 p1)
@@ -1654,9 +1913,9 @@ detail::enable_if_t<detail::is_sgentype<T>::value, T> select(T a, T b,
16541913
// mgentype select (mgentype a, mgentype b, marray<bool, { N }> c)
16551914
template <typename T,
16561915
typename = std::enable_if_t<detail::is_mgenfloat<T>::value>>
1657-
sycl::marray<detail::marray_element_type<T>, T::size()>
1916+
sycl::marray<detail::marray_element_t<T>, T::size()>
16581917
select(T a, T b, sycl::marray<bool, T::size()> c) __NOEXC {
1659-
sycl::marray<detail::marray_element_type<T>, T::size()> res;
1918+
sycl::marray<detail::marray_element_t<T>, T::size()> res;
16601919
for (int i = 0; i < a.size(); i++) {
16611920
res[i] = select(a[i], b[i], c[i]);
16621921
}

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

+2-4
Original file line numberDiff line numberDiff line change
@@ -59,12 +59,10 @@ using is_vgenfloat = is_contained<T, gtl::vector_floating_list>;
5959
template <typename T>
6060
using is_svgenfloat = is_contained<T, gtl::scalar_vector_floating_list>;
6161

62-
template <typename T> using marray_element_type = typename T::value_type;
63-
6462
template <typename T>
6563
using is_mgenfloat = bool_constant<
66-
std::is_same<T, sycl::marray<marray_element_type<T>, T::size()>>::value &&
67-
is_svgenfloat<marray_element_type<T>>::value>;
64+
std::is_same<T, sycl::marray<marray_element_t<T>, T::size()>>::value &&
65+
is_svgenfloat<marray_element_t<T>>::value>;
6866

6967
template <typename T>
7068
using is_gengeofloat = is_contained<T, gtl::geo_float_list>;

0 commit comments

Comments
 (0)