Skip to content

Commit 2a3e1ab

Browse files
steffenlarsensteffenlarsen
authored
[SYCL] Change return type of abs_diff (#11914)
As of KhronosGroup/SYCL-Docs#458 the return type of abs_diff has been changed to be the same as the input arguments, which in turn corresponds to how the extended OpenCL instruction set defines s_abs_diff. This commit changes the preview builtins to use the correct type similar to how `sycl::abs` does. --------- Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
1 parent 1dc1372 commit 2a3e1ab

File tree

2 files changed

+45
-34
lines changed

2 files changed

+45
-34
lines changed

‎sycl/source/builtins_generator.py

Lines changed: 18 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -324,7 +324,6 @@ def __str__(self):
324324
# To help resolve template arguments, these are given the index of their parent
325325
# argument.
326326
elementtype0 = [ElementType(0)]
327-
unsignedtype0 = [ConversionTraitType("make_unsigned_t", 0)]
328327
samesizesignedint0 = [ConversionTraitType("same_size_signed_int_t", 0)]
329328
samesizeunsignedint0 = [ConversionTraitType("same_size_unsigned_int_t", 0)]
330329
intelements0 = [ConversionTraitType("int_elements_t", 0)]
@@ -400,7 +399,6 @@ def __str__(self):
400399
"intnptr0" : intnptr0,
401400
"vint32nptr0" : vint32ptr0,
402401
"elementtype0" : elementtype0,
403-
"unsignedtype0" : unsignedtype0,
404402
"samesizesignedint0" : samesizesignedint0,
405403
"samesizeunsignedint0" : samesizeunsignedint0,
406404
"intelements0" : intelements0,
@@ -691,15 +689,19 @@ def get_scalar_vec_invoke_body(self, invoke_name, return_type, arg_types, arg_na
691689
invoke_args = ', '.join(get_invoke_args(arg_types, arg_names))
692690
return f' return detail::RelConverter<{return_type}>::apply(__sycl_std::__invoke_{self.invoke_prefix}{invoke_name}<detail::internal_rel_ret_t<{return_type}>>({invoke_args}));'
693691

694-
def custom_signed_abs_scalar_invoke(return_type, _, arg_names):
695-
"""Generates the custom body for signed scalar `abs`."""
696-
args = ' ,'.join(arg_names)
697-
return f'return static_cast<{return_type}>(__sycl_std::__invoke_s_abs<detail::make_unsigned_t<{return_type}>>({args}));'
692+
def get_custom_unsigned_to_signed_scalar_invoke(invoke_name):
693+
"""
694+
Creates a function for generating the custom body for invocations returning
695+
an unsigned scalar value, which will in turn be converted to a signed value.
696+
"""
697+
return (lambda return_type, _, arg_names: f'return static_cast<{return_type}>(__sycl_std::__invoke_{invoke_name}<detail::make_unsigned_t<{return_type}>>({" ,".join(arg_names)}));')
698698

699-
def custom_signed_abs_vec_invoke(return_type, arg_types, arg_names):
700-
"""Generates the custom body for signed vector `abs`."""
701-
args = ' ,'.join(get_invoke_args(arg_types, arg_names))
702-
return f'return __sycl_std::__invoke_s_abs<detail::make_unsigned_t<{return_type}>>({args}).template convert<detail::get_elem_type_t<{return_type}>>();'
699+
def get_custom_unsigned_to_signed_vec_invoke(invoke_name):
700+
"""
701+
Creates a function for generating the custom body for invocations returning
702+
an unsigned scalar value, which will in turn be converted to a signed value.
703+
"""
704+
return (lambda return_type, arg_types, arg_names: f'return __sycl_std::__invoke_{invoke_name}<detail::make_unsigned_t<{return_type}>>({" ,".join(get_invoke_args(arg_types, arg_names))}).template convert<detail::get_elem_type_t<{return_type}>>();')
703705

704706
def get_custom_any_all_vec_invoke(invoke_name):
705707
"""
@@ -873,8 +875,10 @@ def custom_nan_invoke(return_type, arg_types, arg_names):
873875
"tgamma": [Def("genfloat", ["genfloat"])],
874876
"trunc": [Def("genfloat", ["genfloat"])],
875877
# Integer functions
876-
"abs_diff": [Def("unsignedtype0", ["igeninteger", "igeninteger"], invoke_prefix="s_", marray_use_loop=True, template_scalar_args=True),
877-
Def("unsignedtype0", ["ugeninteger", "ugeninteger"], invoke_prefix="u_", marray_use_loop=True, template_scalar_args=True)],
878+
"abs_diff": [Def("sigeninteger", ["sigeninteger", "sigeninteger"], custom_invoke=get_custom_unsigned_to_signed_scalar_invoke("s_abs_diff"), template_scalar_args=True),
879+
Def("vigeninteger", ["vigeninteger", "vigeninteger"], custom_invoke=get_custom_unsigned_to_signed_vec_invoke("s_abs_diff")),
880+
Def("migeninteger", ["migeninteger", "migeninteger"], marray_use_loop=True),
881+
Def("ugeninteger", ["ugeninteger", "ugeninteger"], invoke_prefix="u_", marray_use_loop=True, template_scalar_args=True)],
878882
"add_sat": [Def("igeninteger", ["igeninteger", "igeninteger"], invoke_prefix="s_", marray_use_loop=True, template_scalar_args=True),
879883
Def("ugeninteger", ["ugeninteger", "ugeninteger"], invoke_prefix="u_", marray_use_loop=True, template_scalar_args=True)],
880884
"hadd": [Def("igeninteger", ["igeninteger", "igeninteger"], invoke_prefix="s_", marray_use_loop=True, template_scalar_args=True),
@@ -968,8 +972,8 @@ def custom_nan_invoke(return_type, arg_types, arg_names):
968972
Def("mdoublen", ["double", "double", "mdoublen"]),
969973
Def("mhalfn", ["half", "half", "mhalfn"])],
970974
"sign": [Def("genfloat", ["genfloat"], template_scalar_args=True)],
971-
"abs": [Def("sigeninteger", ["sigeninteger"], custom_invoke=custom_signed_abs_scalar_invoke, template_scalar_args=True),
972-
Def("vigeninteger", ["vigeninteger"], custom_invoke=custom_signed_abs_vec_invoke),
975+
"abs": [Def("sigeninteger", ["sigeninteger"], custom_invoke=get_custom_unsigned_to_signed_scalar_invoke("s_abs"), template_scalar_args=True),
976+
Def("vigeninteger", ["vigeninteger"], custom_invoke=get_custom_unsigned_to_signed_vec_invoke("s_abs")),
973977
Def("migeninteger", ["migeninteger"], marray_use_loop=True),
974978
Def("ugeninteger", ["ugeninteger"], invoke_prefix="u_", marray_use_loop=True, template_scalar_args=True)],
975979
# Geometric functions

‎sycl/test/regression/abs_diff_host.cpp

Lines changed: 27 additions & 20 deletions
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,7 @@
1-
// RUN: %clangxx -fsycl %s -o %t.out
2-
// RUN: %t.out
3-
// RUN: %if preview-breaking-changes-supported %{ %clangxx -fsycl -fpreview-breaking-changes %s -o %t2.out %}
4-
// RUN: %if preview-breaking-changes-supported %{ %t2.out %}
1+
// RUN: %clangxx -fsycl %s -o %t.oRT
2+
// RUN: %t.oRT
3+
// RUN: %if preview-breaking-changes-supported %{ %clangxx -fsycl -fpreview-breaking-changes %s -o %t2.oRT %}
4+
// RUN: %if preview-breaking-changes-supported %{ %t2.oRT %}
55

66
// Test checks that sycl::abs_diff correctly handles signed operations that
77
// might overflow.
@@ -12,26 +12,33 @@
1212
#include <type_traits>
1313

1414
template <typename T> void check() {
15-
using UT = std::make_unsigned_t<T>;
15+
#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
16+
using RT = T;
17+
#else
18+
using RT = std::make_unsigned_t<T>;
19+
#endif
20+
1621
constexpr T MaxVal = std::numeric_limits<T>::max();
1722
constexpr T MinVal = std::numeric_limits<T>::min();
18-
constexpr UT UMaxVal = MaxVal;
19-
// Avoid unrepresentable negation.
20-
constexpr UT UAbsMinVal = UMaxVal - (MaxVal + MinVal);
23+
constexpr RT RMaxVal = MaxVal;
2124

2225
// Sanity checks to make sure simple distances work.
23-
assert(sycl::abs_diff(MaxVal, T(10)) == UT(MaxVal - 10));
24-
assert(sycl::abs_diff(T(10), MaxVal) == UT(MaxVal - 10));
25-
assert(sycl::abs_diff(MinVal, T(-10)) == UT(MinVal - 10));
26-
assert(sycl::abs_diff(T(-10), MinVal) == UT(MinVal - 10));
27-
28-
// Checks potential overflows.
29-
assert(sycl::abs_diff(MaxVal, T(-10)) == (UMaxVal + 10));
30-
assert(sycl::abs_diff(T(-10), MaxVal) == (UMaxVal + 10));
31-
assert(sycl::abs_diff(MinVal, T(10)) == (UAbsMinVal + 10));
32-
assert(sycl::abs_diff(T(10), MinVal) == (UAbsMinVal + 10));
33-
assert(sycl::abs_diff(MaxVal, MinVal) == (UMaxVal + UAbsMinVal));
34-
assert(sycl::abs_diff(MinVal, MaxVal) == (UMaxVal + UAbsMinVal));
26+
assert(sycl::abs_diff(MaxVal, T(10)) == RT(MaxVal - 10));
27+
assert(sycl::abs_diff(T(10), MaxVal) == RT(MaxVal - 10));
28+
assert(sycl::abs_diff(MinVal, T(-10)) == RT(MinVal - 10));
29+
assert(sycl::abs_diff(T(-10), MinVal) == RT(MinVal - 10));
30+
31+
#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
32+
// Avoid unrepresentable negation.
33+
constexpr RT RAbsMinVal = RMaxVal - (MaxVal + MinVal);
34+
// Checks potential overflows. This is UB in SYCL 2020 Rev. 8.
35+
assert(sycl::abs_diff(MaxVal, T(-10)) == (RMaxVal + 10));
36+
assert(sycl::abs_diff(T(-10), MaxVal) == (RMaxVal + 10));
37+
assert(sycl::abs_diff(MinVal, T(10)) == (RAbsMinVal + 10));
38+
assert(sycl::abs_diff(T(10), MinVal) == (RAbsMinVal + 10));
39+
assert(sycl::abs_diff(MaxVal, MinVal) == (RMaxVal + RAbsMinVal));
40+
assert(sycl::abs_diff(MinVal, MaxVal) == (RMaxVal + RAbsMinVal));
41+
#endif
3542
}
3643

3744
int main() {

0 commit comments

Comments
 (0)