[SYCL][SPIRV] Implement islessgreater with FOrdNotEqual instead (#5076)

SPIR-V OpLessOrGreater is deprecated: https://www.khronos.org/registry/SPIR-V/specs/unified1/SPIRV.html#OpLessOrGreater OpFOrdNotEqual has the same semantics as OpLessOrGreater. Signed-off-by: Yilong Guo [email protected]
intel · Dec 9, 2021 · 0b8df3b · 0b8df3b
1 parent 7bcd22c
commit 0b8df3b
Show file tree

Hide file tree

Showing 6 changed files with 78 additions and 1 deletion.
diff --git a/sycl/include/CL/sycl/builtins.hpp b/sycl/include/CL/sycl/builtins.hpp
@@ -1176,7 +1176,7 @@ template <typename T,
           typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
 detail::common_rel_ret_t<T> islessgreater(T x, T y) __NOEXC {
   return detail::RelConverter<T>::apply(
-      __sycl_std::__invoke_LessOrGreater<detail::rel_ret_t<T>>(x, y));
+      __sycl_std::__invoke_FOrdNotEqual<detail::rel_ret_t<T>>(x, y));
 }
 
 // int isfinite (half x)

diff --git 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
@@ -47,6 +47,14 @@ template <typename T> inline T __vFOrdLessThanEqual(T x, T y) {
 
 template <typename T> inline T __sFOrdLessThanEqual(T x, T y) { return x <= y; }
 
+template <typename T> inline T __vFOrdNotEqual(T x, T y) {
+  return -((x < y) || (x > y));
+}
+
+template <typename T> inline T __sFOrdNotEqual(T x, T y) {
+  return ((x < y) || (x > y));
+}
+
 template <typename T> 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
@@ -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
@@ -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
@@ -0,0 +1,9 @@
+// RUN: %clangxx -I %sycl_include -S -emit-llvm -fsycl-device-only %s -o - -Xclang -disable-llvm-passes | FileCheck %s
+
+#include <CL/sycl.hpp>
+
+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 {{.*}})