diff --git a/sycl/include/CL/sycl/builtins.hpp b/sycl/include/CL/sycl/builtins.hpp index c82433c61c091..8dde058034eed 100644 --- a/sycl/include/CL/sycl/builtins.hpp +++ b/sycl/include/CL/sycl/builtins.hpp @@ -1176,7 +1176,7 @@ template ::value, T>> detail::common_rel_ret_t islessgreater(T x, T y) __NOEXC { return detail::RelConverter::apply( - __sycl_std::__invoke_LessOrGreater>(x, y)); + __sycl_std::__invoke_FOrdNotEqual>(x, y)); } // int isfinite (half x) diff --git a/sycl/include/CL/sycl/detail/builtins.hpp b/sycl/include/CL/sycl/detail/builtins.hpp index a1e10c32f205d..f4f0475ea905b 100644 --- a/sycl/include/CL/sycl/detail/builtins.hpp +++ b/sycl/include/CL/sycl/detail/builtins.hpp @@ -265,6 +265,7 @@ __SYCL_MAKE_CALL_ARG2_SAME(FOrdGreaterThanEqual, __FUNC_PREFIX_CORE) // isgreaterequal __SYCL_MAKE_CALL_ARG2_SAME(FOrdLessThan, __FUNC_PREFIX_CORE) // isless __SYCL_MAKE_CALL_ARG2_SAME(FOrdLessThanEqual, __FUNC_PREFIX_CORE) // islessequal +__SYCL_MAKE_CALL_ARG2_SAME(FOrdNotEqual, __FUNC_PREFIX_CORE) // islessgreater __SYCL_MAKE_CALL_ARG2_SAME(LessOrGreater, __FUNC_PREFIX_CORE) // islessgreater __SYCL_MAKE_CALL_ARG1(IsFinite, __FUNC_PREFIX_CORE) // isfinite __SYCL_MAKE_CALL_ARG1(IsInf, __FUNC_PREFIX_CORE) // isinf diff --git a/sycl/source/detail/builtins_relational.cpp b/sycl/source/detail/builtins_relational.cpp index 8ae7ca1990689..3300213d36a5b 100644 --- a/sycl/source/detail/builtins_relational.cpp +++ b/sycl/source/detail/builtins_relational.cpp @@ -47,6 +47,14 @@ template inline T __vFOrdLessThanEqual(T x, T y) { template inline T __sFOrdLessThanEqual(T x, T y) { return x <= y; } +template inline T __vFOrdNotEqual(T x, T y) { + return -((x < y) || (x > y)); +} + +template inline T __sFOrdNotEqual(T x, T y) { + return ((x < y) || (x > y)); +} + template inline T __vLessOrGreater(T x, T y) { return -((x < y) || (x > y)); } @@ -247,6 +255,23 @@ MAKE_1V_2V_FUNC(FOrdLessThanEqual, __vFOrdLessThanEqual, s::cl_long, MAKE_1V_2V_FUNC(FOrdLessThanEqual, __vFOrdLessThanEqual, s::cl_short, s::cl_half, s::cl_half) +// (FOrdNotEqual) // islessgreater +__SYCL_EXPORT s::cl_int FOrdNotEqual(s::cl_float x, s::cl_float y) __NOEXC { + return __sFOrdNotEqual(x, y); +} +__SYCL_EXPORT s::cl_int FOrdNotEqual(s::cl_double x, s::cl_double y) __NOEXC { + return __sFOrdNotEqual(x, y); +} +__SYCL_EXPORT s::cl_int FOrdNotEqual(s::cl_half x, s::cl_half y) __NOEXC { + return __sFOrdNotEqual(x, y); +} +MAKE_1V_2V_FUNC(FOrdNotEqual, __vFOrdNotEqual, s::cl_int, s::cl_float, + s::cl_float) +MAKE_1V_2V_FUNC(FOrdNotEqual, __vFOrdNotEqual, s::cl_long, s::cl_double, + s::cl_double) +MAKE_1V_2V_FUNC(FOrdNotEqual, __vFOrdNotEqual, s::cl_short, s::cl_half, + s::cl_half) + // (LessOrGreater) // islessgreater __SYCL_EXPORT s::cl_int LessOrGreater(s::cl_float x, s::cl_float y) __NOEXC { return __sLessOrGreater(x, y); diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 12a2de86c252a..4d8a377088db6 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -323,6 +323,27 @@ _ZN2cl10__host_std12FOrdLessThanENS_4sycl3vecIfLi8EEES3_ _ZN2cl10__host_std12FOrdLessThanENS_4sycl6detail9half_impl4halfES4_ _ZN2cl10__host_std12FOrdLessThanEdd _ZN2cl10__host_std12FOrdLessThanEff +_ZN2cl10__host_std12FOrdNotEqualENS_4sycl3vecINS1_6detail9half_impl4halfELi16EEES6_ +_ZN2cl10__host_std12FOrdNotEqualENS_4sycl3vecINS1_6detail9half_impl4halfELi1EEES6_ +_ZN2cl10__host_std12FOrdNotEqualENS_4sycl3vecINS1_6detail9half_impl4halfELi2EEES6_ +_ZN2cl10__host_std12FOrdNotEqualENS_4sycl3vecINS1_6detail9half_impl4halfELi3EEES6_ +_ZN2cl10__host_std12FOrdNotEqualENS_4sycl3vecINS1_6detail9half_impl4halfELi4EEES6_ +_ZN2cl10__host_std12FOrdNotEqualENS_4sycl3vecINS1_6detail9half_impl4halfELi8EEES6_ +_ZN2cl10__host_std12FOrdNotEqualENS_4sycl3vecIdLi16EEES3_ +_ZN2cl10__host_std12FOrdNotEqualENS_4sycl3vecIdLi1EEES3_ +_ZN2cl10__host_std12FOrdNotEqualENS_4sycl3vecIdLi2EEES3_ +_ZN2cl10__host_std12FOrdNotEqualENS_4sycl3vecIdLi3EEES3_ +_ZN2cl10__host_std12FOrdNotEqualENS_4sycl3vecIdLi4EEES3_ +_ZN2cl10__host_std12FOrdNotEqualENS_4sycl3vecIdLi8EEES3_ +_ZN2cl10__host_std12FOrdNotEqualENS_4sycl3vecIfLi16EEES3_ +_ZN2cl10__host_std12FOrdNotEqualENS_4sycl3vecIfLi1EEES3_ +_ZN2cl10__host_std12FOrdNotEqualENS_4sycl3vecIfLi2EEES3_ +_ZN2cl10__host_std12FOrdNotEqualENS_4sycl3vecIfLi3EEES3_ +_ZN2cl10__host_std12FOrdNotEqualENS_4sycl3vecIfLi4EEES3_ +_ZN2cl10__host_std12FOrdNotEqualENS_4sycl3vecIfLi8EEES3_ +_ZN2cl10__host_std12FOrdNotEqualENS_4sycl6detail9half_impl4halfES4_ +_ZN2cl10__host_std12FOrdNotEqualEdd +_ZN2cl10__host_std12FOrdNotEqualEff _ZN2cl10__host_std12native_exp10ENS_4sycl3vecIfLi16EEE _ZN2cl10__host_std12native_exp10ENS_4sycl3vecIfLi1EEE _ZN2cl10__host_std12native_exp10ENS_4sycl3vecIfLi2EEE diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index e33424c33e2ef..3d4a525f89651 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -744,6 +744,27 @@ ?FOrdLessThanEqual@__host_std@cl@@YAHMM@Z ?FOrdLessThanEqual@__host_std@cl@@YAHNN@Z ?FOrdLessThanEqual@__host_std@cl@@YAHVhalf@half_impl@detail@sycl@2@0@Z +?FOrdNotEqual@__host_std@cl@@YA?AV?$vec@F$00@sycl@2@V?$vec@Vhalf@half_impl@detail@sycl@cl@@$00@42@0@Z +?FOrdNotEqual@__host_std@cl@@YA?AV?$vec@F$01@sycl@2@V?$vec@Vhalf@half_impl@detail@sycl@cl@@$01@42@0@Z +?FOrdNotEqual@__host_std@cl@@YA?AV?$vec@F$02@sycl@2@V?$vec@Vhalf@half_impl@detail@sycl@cl@@$02@42@0@Z +?FOrdNotEqual@__host_std@cl@@YA?AV?$vec@F$03@sycl@2@V?$vec@Vhalf@half_impl@detail@sycl@cl@@$03@42@0@Z +?FOrdNotEqual@__host_std@cl@@YA?AV?$vec@F$07@sycl@2@V?$vec@Vhalf@half_impl@detail@sycl@cl@@$07@42@0@Z +?FOrdNotEqual@__host_std@cl@@YA?AV?$vec@F$0BA@@sycl@2@V?$vec@Vhalf@half_impl@detail@sycl@cl@@$0BA@@42@0@Z +?FOrdNotEqual@__host_std@cl@@YA?AV?$vec@H$00@sycl@2@V?$vec@M$00@42@0@Z +?FOrdNotEqual@__host_std@cl@@YA?AV?$vec@H$01@sycl@2@V?$vec@M$01@42@0@Z +?FOrdNotEqual@__host_std@cl@@YA?AV?$vec@H$02@sycl@2@V?$vec@M$02@42@0@Z +?FOrdNotEqual@__host_std@cl@@YA?AV?$vec@H$03@sycl@2@V?$vec@M$03@42@0@Z +?FOrdNotEqual@__host_std@cl@@YA?AV?$vec@H$07@sycl@2@V?$vec@M$07@42@0@Z +?FOrdNotEqual@__host_std@cl@@YA?AV?$vec@H$0BA@@sycl@2@V?$vec@M$0BA@@42@0@Z +?FOrdNotEqual@__host_std@cl@@YA?AV?$vec@_J$00@sycl@2@V?$vec@N$00@42@0@Z +?FOrdNotEqual@__host_std@cl@@YA?AV?$vec@_J$01@sycl@2@V?$vec@N$01@42@0@Z +?FOrdNotEqual@__host_std@cl@@YA?AV?$vec@_J$02@sycl@2@V?$vec@N$02@42@0@Z +?FOrdNotEqual@__host_std@cl@@YA?AV?$vec@_J$03@sycl@2@V?$vec@N$03@42@0@Z +?FOrdNotEqual@__host_std@cl@@YA?AV?$vec@_J$07@sycl@2@V?$vec@N$07@42@0@Z +?FOrdNotEqual@__host_std@cl@@YA?AV?$vec@_J$0BA@@sycl@2@V?$vec@N$0BA@@42@0@Z +?FOrdNotEqual@__host_std@cl@@YAHMM@Z +?FOrdNotEqual@__host_std@cl@@YAHNN@Z +?FOrdNotEqual@__host_std@cl@@YAHVhalf@half_impl@detail@sycl@2@0@Z ?FUnordNotEqual@__host_std@cl@@YA?AV?$vec@F$00@sycl@2@V?$vec@Vhalf@half_impl@detail@sycl@cl@@$00@42@0@Z ?FUnordNotEqual@__host_std@cl@@YA?AV?$vec@F$01@sycl@2@V?$vec@Vhalf@half_impl@detail@sycl@cl@@$01@42@0@Z ?FUnordNotEqual@__host_std@cl@@YA?AV?$vec@F$02@sycl@2@V?$vec@Vhalf@half_impl@detail@sycl@cl@@$02@42@0@Z diff --git a/sycl/test/check_device_code/islessgreater.cpp b/sycl/test/check_device_code/islessgreater.cpp new file mode 100644 index 0000000000000..1c611f14553a2 --- /dev/null +++ b/sycl/test/check_device_code/islessgreater.cpp @@ -0,0 +1,9 @@ +// RUN: %clangxx -I %sycl_include -S -emit-llvm -fsycl-device-only %s -o - -Xclang -disable-llvm-passes | FileCheck %s + +#include + +SYCL_EXTERNAL void test_islessgreater(float a, float b) { + sycl::islessgreater(a, b); +} +// CHECK-NOT: __spirv_LessOrGreater +// CHECK: {{.*}} = call spir_func zeroext i1 @_Z20__spirv_FOrdNotEqualff(float {{.*}}, float {{.*}})