Skip to content

Commit

Permalink
[SYCL] Change return type of abs_diff (#11914)
Browse files Browse the repository at this point in the history 
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 <[email protected]>
  • Loading branch information
@steffenlarsen
steffenlarsen authored Nov 21, 2023
1 parent 1dc1372 commit 2a3e1ab
Show file tree
Hide file tree
Showing 2 changed files with 45 additions and 34 deletions.
32 changes: 18 additions & 14 deletions sycl/source/builtins_generator.py
View file
Original file line number Diff line number Diff line change
Expand Up @@ -324,7 +324,6 @@ def __str__(self):
# To help resolve template arguments, these are given the index of their parent
# argument.
elementtype0 = [ElementType(0)]
unsignedtype0 = [ConversionTraitType("make_unsigned_t", 0)]
samesizesignedint0 = [ConversionTraitType("same_size_signed_int_t", 0)]
samesizeunsignedint0 = [ConversionTraitType("same_size_unsigned_int_t", 0)]
intelements0 = [ConversionTraitType("int_elements_t", 0)]
Expand Down Expand Up @@ -400,7 +399,6 @@ def __str__(self):
"intnptr0" : intnptr0,
"vint32nptr0" : vint32ptr0,
"elementtype0" : elementtype0,
"unsignedtype0" : unsignedtype0,
"samesizesignedint0" : samesizesignedint0,
"samesizeunsignedint0" : samesizeunsignedint0,
"intelements0" : intelements0,
Expand Down Expand Up @@ -691,15 +689,19 @@ def get_scalar_vec_invoke_body(self, invoke_name, return_type, arg_types, arg_na
invoke_args = ', '.join(get_invoke_args(arg_types, arg_names))
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}));'

def custom_signed_abs_scalar_invoke(return_type, _, arg_names):
"""Generates the custom body for signed scalar `abs`."""
args = ' ,'.join(arg_names)
return f'return static_cast<{return_type}>(__sycl_std::__invoke_s_abs<detail::make_unsigned_t<{return_type}>>({args}));'
def get_custom_unsigned_to_signed_scalar_invoke(invoke_name):
"""
Creates a function for generating the custom body for invocations returning
an unsigned scalar value, which will in turn be converted to a signed value.
"""
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)}));')

def custom_signed_abs_vec_invoke(return_type, arg_types, arg_names):
"""Generates the custom body for signed vector `abs`."""
args = ' ,'.join(get_invoke_args(arg_types, arg_names))
return f'return __sycl_std::__invoke_s_abs<detail::make_unsigned_t<{return_type}>>({args}).template convert<detail::get_elem_type_t<{return_type}>>();'
def get_custom_unsigned_to_signed_vec_invoke(invoke_name):
"""
Creates a function for generating the custom body for invocations returning
an unsigned scalar value, which will in turn be converted to a signed value.
"""
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}>>();')

def get_custom_any_all_vec_invoke(invoke_name):
"""
Expand Down Expand Up @@ -873,8 +875,10 @@ def custom_nan_invoke(return_type, arg_types, arg_names):
"tgamma": [Def("genfloat", ["genfloat"])],
"trunc": [Def("genfloat", ["genfloat"])],
# Integer functions
"abs_diff": [Def("unsignedtype0", ["igeninteger", "igeninteger"], invoke_prefix="s_", marray_use_loop=True, template_scalar_args=True),
Def("unsignedtype0", ["ugeninteger", "ugeninteger"], invoke_prefix="u_", marray_use_loop=True, template_scalar_args=True)],
"abs_diff": [Def("sigeninteger", ["sigeninteger", "sigeninteger"], custom_invoke=get_custom_unsigned_to_signed_scalar_invoke("s_abs_diff"), template_scalar_args=True),
Def("vigeninteger", ["vigeninteger", "vigeninteger"], custom_invoke=get_custom_unsigned_to_signed_vec_invoke("s_abs_diff")),
Def("migeninteger", ["migeninteger", "migeninteger"], marray_use_loop=True),
Def("ugeninteger", ["ugeninteger", "ugeninteger"], invoke_prefix="u_", marray_use_loop=True, template_scalar_args=True)],
"add_sat": [Def("igeninteger", ["igeninteger", "igeninteger"], invoke_prefix="s_", marray_use_loop=True, template_scalar_args=True),
Def("ugeninteger", ["ugeninteger", "ugeninteger"], invoke_prefix="u_", marray_use_loop=True, template_scalar_args=True)],
"hadd": [Def("igeninteger", ["igeninteger", "igeninteger"], invoke_prefix="s_", marray_use_loop=True, template_scalar_args=True),
Expand Down Expand Up @@ -968,8 +972,8 @@ def custom_nan_invoke(return_type, arg_types, arg_names):
Def("mdoublen", ["double", "double", "mdoublen"]),
Def("mhalfn", ["half", "half", "mhalfn"])],
"sign": [Def("genfloat", ["genfloat"], template_scalar_args=True)],
"abs": [Def("sigeninteger", ["sigeninteger"], custom_invoke=custom_signed_abs_scalar_invoke, template_scalar_args=True),
Def("vigeninteger", ["vigeninteger"], custom_invoke=custom_signed_abs_vec_invoke),
"abs": [Def("sigeninteger", ["sigeninteger"], custom_invoke=get_custom_unsigned_to_signed_scalar_invoke("s_abs"), template_scalar_args=True),
Def("vigeninteger", ["vigeninteger"], custom_invoke=get_custom_unsigned_to_signed_vec_invoke("s_abs")),
Def("migeninteger", ["migeninteger"], marray_use_loop=True),
Def("ugeninteger", ["ugeninteger"], invoke_prefix="u_", marray_use_loop=True, template_scalar_args=True)],
# Geometric functions
Expand Down
47 changes: 27 additions & 20 deletions sycl/test/regression/abs_diff_host.cpp
View file
Original file line number Diff line number Diff line change
@@ -1,7 +1,7 @@
// RUN: %clangxx -fsycl %s -o %t.out
// RUN: %t.out
// RUN: %if preview-breaking-changes-supported %{ %clangxx -fsycl -fpreview-breaking-changes %s -o %t2.out %}
// RUN: %if preview-breaking-changes-supported %{ %t2.out %}
// RUN: %clangxx -fsycl %s -o %t.oRT
// RUN: %t.oRT
// RUN: %if preview-breaking-changes-supported %{ %clangxx -fsycl -fpreview-breaking-changes %s -o %t2.oRT %}
// RUN: %if preview-breaking-changes-supported %{ %t2.oRT %}

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

template <typename T> void check() {
using UT = std::make_unsigned_t<T>;
#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
using RT = T;
#else
using RT = std::make_unsigned_t<T>;
#endif

constexpr T MaxVal = std::numeric_limits<T>::max();
constexpr T MinVal = std::numeric_limits<T>::min();
constexpr UT UMaxVal = MaxVal;
// Avoid unrepresentable negation.
constexpr UT UAbsMinVal = UMaxVal - (MaxVal + MinVal);
constexpr RT RMaxVal = MaxVal;

// Sanity checks to make sure simple distances work.
assert(sycl::abs_diff(MaxVal, T(10)) == UT(MaxVal - 10));
assert(sycl::abs_diff(T(10), MaxVal) == UT(MaxVal - 10));
assert(sycl::abs_diff(MinVal, T(-10)) == UT(MinVal - 10));
assert(sycl::abs_diff(T(-10), MinVal) == UT(MinVal - 10));

// Checks potential overflows.
assert(sycl::abs_diff(MaxVal, T(-10)) == (UMaxVal + 10));
assert(sycl::abs_diff(T(-10), MaxVal) == (UMaxVal + 10));
assert(sycl::abs_diff(MinVal, T(10)) == (UAbsMinVal + 10));
assert(sycl::abs_diff(T(10), MinVal) == (UAbsMinVal + 10));
assert(sycl::abs_diff(MaxVal, MinVal) == (UMaxVal + UAbsMinVal));
assert(sycl::abs_diff(MinVal, MaxVal) == (UMaxVal + UAbsMinVal));
assert(sycl::abs_diff(MaxVal, T(10)) == RT(MaxVal - 10));
assert(sycl::abs_diff(T(10), MaxVal) == RT(MaxVal - 10));
assert(sycl::abs_diff(MinVal, T(-10)) == RT(MinVal - 10));
assert(sycl::abs_diff(T(-10), MinVal) == RT(MinVal - 10));

#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
// Avoid unrepresentable negation.
constexpr RT RAbsMinVal = RMaxVal - (MaxVal + MinVal);
// Checks potential overflows. This is UB in SYCL 2020 Rev. 8.
assert(sycl::abs_diff(MaxVal, T(-10)) == (RMaxVal + 10));
assert(sycl::abs_diff(T(-10), MaxVal) == (RMaxVal + 10));
assert(sycl::abs_diff(MinVal, T(10)) == (RAbsMinVal + 10));
assert(sycl::abs_diff(T(10), MinVal) == (RAbsMinVal + 10));
assert(sycl::abs_diff(MaxVal, MinVal) == (RMaxVal + RAbsMinVal));
assert(sycl::abs_diff(MinVal, MaxVal) == (RMaxVal + RAbsMinVal));
#endif
}

int main() {
Expand Down

0 comments on commit 2a3e1ab

Please sign in to comment.