diff --git a/Eigen/src/Core/arch/SYCL/PacketMath.h b/Eigen/src/Core/arch/SYCL/PacketMath.h index 92accc88f..5bc323589 100644 --- a/Eigen/src/Core/arch/SYCL/PacketMath.h +++ b/Eigen/src/Core/arch/SYCL/PacketMath.h @@ -477,25 +477,19 @@ pabs(const cl::sycl::cl_double2& a) { template EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Packet sycl_pcmp_le(const Packet &a, const Packet &b) { - return ((a <= b) - .template convert::type, - cl::sycl::rounding_mode::automatic>()); + return (a <= b).template as(); } template EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Packet sycl_pcmp_lt(const Packet &a, const Packet &b) { - return ((a < b) - .template convert::type, - cl::sycl::rounding_mode::automatic>()); + return (a < b).template as(); } template EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Packet sycl_pcmp_eq(const Packet &a, const Packet &b) { - return ((a == b) - .template convert::type, - cl::sycl::rounding_mode::automatic>()); + return (a == b).template as(); } #define SYCL_PCMP(OP, TYPE) \ @@ -513,76 +507,6 @@ SYCL_PCMP(lt, cl::sycl::cl_double2) SYCL_PCMP(eq, cl::sycl::cl_double2) #undef SYCL_PCMP -template struct convert_to_integer; - -template <> struct convert_to_integer { - using type = std::int32_t; - using packet_type = cl::sycl::cl_int4; -}; -template <> struct convert_to_integer { - using type = std::int64_t; - using packet_type = cl::sycl::cl_long2; -}; - -template -EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE typename convert_to_integer< - typename unpacket_traits::type>::packet_type -vector_as_int(const PacketIn &p) { - return ( - p.template convert::type>::type, - cl::sycl::rounding_mode::automatic>()); -} - -template -EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packetOut -convert_vector(const PacketIn &p) { - return (p.template convert::type, - cl::sycl::rounding_mode::automatic>()); -} - -#define SYCL_PAND(TYPE) \ - template <> \ - EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TYPE pand(const TYPE &a, \ - const TYPE &b) { \ - return convert_vector(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(const TYPE &a, \ - const TYPE &b) { \ - return convert_vector(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(const TYPE &a, \ - const TYPE &b) { \ - return convert_vector(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(const TYPE &a, \ - const TYPE &b) { \ - return convert_vector(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& kernel) { float tmp = kernel.packet[0].y(); diff --git a/unsupported/test/cxx11_tensor_builtins_sycl.cpp b/unsupported/test/cxx11_tensor_builtins_sycl.cpp index df142fe8a..27a82540f 100644 --- a/unsupported/test/cxx11_tensor_builtins_sycl.cpp +++ b/unsupported/test/cxx11_tensor_builtins_sycl.cpp @@ -86,7 +86,8 @@ void test_unary_builtins_for_scalar(const Eigen::SyclDevice& sycl_device, { /* Assignment(out, Operator(out)) */ Tensor out(tensor_range); - out = out.random() + DataType(0.01); + // Offset with 1 to avoid tiny output (< 1e-6) as they can easily fail. + out = out.random() + DataType(1); Tensor reference(out); DataType *gpu_data_out = static_cast( sycl_device.allocate(out.size() * sizeof(DataType))); diff --git a/unsupported/test/cxx11_tensor_math_sycl.cpp b/unsupported/test/cxx11_tensor_math_sycl.cpp index 03cf0a80b..029653e27 100644 --- a/unsupported/test/cxx11_tensor_math_sycl.cpp +++ b/unsupported/test/cxx11_tensor_math_sycl.cpp @@ -94,12 +94,8 @@ template void sycl_computing_test_per_ auto sycl_device = Eigen::SyclDevice(&queueInterface); test_tanh_sycl(sycl_device); test_tanh_sycl(sycl_device); - // Sigmoid broke for SYCL as of 0b5873 because of missing functions in PacketMath.h. - // Disable the test for now. -#if 0 test_sigmoid_sycl(sycl_device); test_sigmoid_sycl(sycl_device); -#endif } EIGEN_DECLARE_TEST(cxx11_tensor_math_sycl) { diff --git a/unsupported/test/cxx11_tensor_random_sycl.cpp b/unsupported/test/cxx11_tensor_random_sycl.cpp index 6c83894a3..14a7c489c 100644 --- a/unsupported/test/cxx11_tensor_random_sycl.cpp +++ b/unsupported/test/cxx11_tensor_random_sycl.cpp @@ -37,14 +37,8 @@ static void test_sycl_random_uniform(const Eigen::SyclDevice& sycl_device) gpu_out.device(sycl_device)=gpu_out.random(); sycl_device.memcpyDeviceToHost(out.data(), d_out,out_bytes); - for(IndexType i=1; i gen(true); gpu_out.device(sycl_device)=gpu_out.random(gen); sycl_device.memcpyDeviceToHost(out.data(), d_out,out_bytes); - for(IndexType i=1; i