[SYCL Backend]

 * Adding Missing operations for vector comparison in SYCL. This caused compiler error for vector comparison when compiling SYCL
 * Fixing the compiler error for placement new in TensorForcedEval.h This caused compiler error when compiling SYCL backend
 * Reducing the SYCL warning by  removing the abort function inside the kernel
 * Adding Strong inline to functions inside SYCL interop.

1 files changed, 109 insertions, 0 deletions

diff --git a/Eigen/src/Core/arch/SYCL/PacketMath.h b/Eigen/src/Core/arch/SYCL/PacketMath.h
index a9adb64ba..b11b5af9d 100644
--- a/Eigen/src/Core/arch/SYCL/PacketMath.h
+++ b/Eigen/src/Core/arch/SYCL/PacketMath.h
@@ -472,6 +472,115 @@ pabs<cl::sycl::cl_double2>(const cl::sycl::cl_double2& a) {
   return cl::sycl::cl_double2(cl::sycl::fabs(a.x()), cl::sycl::fabs(a.y()));
 }
 
+template <typename Packet>
+EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Packet sycl_pcmp_le(const Packet &a,
+                                                          const Packet &b) {
+  return ((a <= b)
+              .template convert<typename unpacket_traits<Packet>::type,
+                                cl::sycl::rounding_mode::automatic>());
+}
+
+template <typename Packet>
+EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Packet sycl_pcmp_lt(const Packet &a,
+                                                          const Packet &b) {
+  return ((a < b)
+              .template convert<typename unpacket_traits<Packet>::type,
+                                cl::sycl::rounding_mode::automatic>());
+}
+
+template <typename Packet>
+EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Packet sycl_pcmp_eq(const Packet &a,
+                                                          const Packet &b) {
+  return ((a == b)
+              .template convert<typename unpacket_traits<Packet>::type,
+                                cl::sycl::rounding_mode::automatic>());
+}
+
+#define SYCL_PCMP(OP, TYPE)                                                    \
+  template <>                                                                  \
+  EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE TYPE pcmp_##OP<TYPE>(const TYPE &a,    \
+                                                             const TYPE &b) {  \
+    return sycl_pcmp_##OP<TYPE>(a, b);                                         \
+  }
+
+SYCL_PCMP(le, cl::sycl::cl_float4)
+SYCL_PCMP(lt, cl::sycl::cl_float4)
+SYCL_PCMP(eq, cl::sycl::cl_float4)
+SYCL_PCMP(le, cl::sycl::cl_double2)
+SYCL_PCMP(lt, cl::sycl::cl_double2)
+SYCL_PCMP(eq, cl::sycl::cl_double2)
+#undef SYCL_PCMP
+
+template <typename T> struct convert_to_integer;
+
+template <> struct convert_to_integer<float> {
+  using type = int;
+  using packet_type = cl::sycl::cl_int4;
+};
+template <> struct convert_to_integer<double> {
+  using type = long;
+  using packet_type = cl::sycl::cl_long2;
+};
+
+template <typename PacketIn>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE typename convert_to_integer<
+    typename unpacket_traits<PacketIn>::type>::packet_type
+vector_as_int(const PacketIn &p) {
+  return (
+      p.template convert<typename convert_to_integer<
+                             typename unpacket_traits<PacketIn>::type>::type,
+                         cl::sycl::rounding_mode::automatic>());
+}
+
+template <typename packetOut, typename PacketIn>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packetOut
+convert_vector(const PacketIn &p) {
+  return (p.template convert<typename unpacket_traits<packetOut>::type,
+                             cl::sycl::rounding_mode::automatic>());
+}
+
+#define SYCL_PAND(TYPE)                                                        \
+  template <>                                                                  \
+  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TYPE pand<TYPE>(const TYPE &a,         \
+                                                        const TYPE &b) {       \
+    return convert_vector<TYPE>(vector_as_int(a) & vector_as_int(b));          \
+  }
+SYCL_PAND(cl::sycl::cl_float4)
+SYCL_PAND(cl::sycl::cl_double2)
+#undef SYCL_PAND
+
+#define SYCL_POR(TYPE)                                                         \
+  template <>                                                                  \
+  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TYPE por<TYPE>(const TYPE &a,          \
+                                                       const TYPE &b) {        \
+    return convert_vector<TYPE>(vector_as_int(a) | vector_as_int(b));          \
+  }
+
+SYCL_POR(cl::sycl::cl_float4)
+SYCL_POR(cl::sycl::cl_double2)
+#undef SYCL_POR
+
+#define SYCL_PXOR(TYPE)                                                        \
+  template <>                                                                  \
+  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TYPE pxor<TYPE>(const TYPE &a,         \
+                                                        const TYPE &b) {       \
+    return convert_vector<TYPE>(vector_as_int(a) ^ vector_as_int(b));          \
+  }
+
+SYCL_PXOR(cl::sycl::cl_float4)
+SYCL_PXOR(cl::sycl::cl_double2)
+#undef SYCL_PXOR
+
+#define SYCL_PANDNOT(TYPE)                                                     \
+  template <>                                                                  \
+  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TYPE pandnot<TYPE>(const TYPE &a,      \
+                                                           const TYPE &b) {    \
+    return convert_vector<TYPE>(vector_as_int(a) & (~vector_as_int(b)));       \
+  }
+SYCL_PANDNOT(cl::sycl::cl_float4)
+SYCL_PANDNOT(cl::sycl::cl_double2)
+#undef SYCL_PANDNOT
+
 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE void ptranspose(
     PacketBlock<cl::sycl::cl_float4, 4>& kernel) {
   float tmp = kernel.packet[0].y();