GitHub-Proxy/eigen
1
0
Fork 0
mirror of https://gitlab.com/libeigen/eigen.git synced 2025-08-11 03:09:01 +08:00
Code Issues Packages Projects Releases Wiki Activity

[SYCL] Fix some SYCL tests

Browse Source
This commit is contained in:
Romain Biessy 2022-08-16 17:37:54 +00:00 committed by Rasmus Munk Larsen
parent 27367017bd
commit 2f7cce2dd5
4 changed files with 7 additions and 100 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

82
Eigen/src/Core/arch/SYCL/PacketMath.h
View File

@ -477,25 +477,19 @@ pabs<cl::sycl::cl_double2>(const cl::sycl::cl_double2& a) {
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>());
return (a <= b).template as<Packet>();
}
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>());
return (a < b).template as<Packet>();
}
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>());
return (a == b).template as<Packet>();
}
#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 <typename T> struct convert_to_integer;
template <> struct convert_to_integer<float> {
using type = std::int32_t;
using packet_type = cl::sycl::cl_int4;
};
template <> struct convert_to_integer<double> {
using type = std::int64_t;
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();

3
unsupported/test/cxx11_tensor_builtins_sycl.cpp
View File

@ -86,7 +86,8 @@ void test_unary_builtins_for_scalar(const Eigen::SyclDevice& sycl_device,
{
/* Assignment(out, Operator(out)) */
Tensor<DataType, 3, DataLayout, int64_t> 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<DataType, 3, DataLayout, int64_t> reference(out);
DataType *gpu_data_out = static_cast<DataType *>(
sycl_device.allocate(out.size() * sizeof(DataType)));

4
unsupported/test/cxx11_tensor_math_sycl.cpp
View File

@ -94,12 +94,8 @@ template<typename DataType, typename dev_Selector> void sycl_computing_test_per_
auto sycl_device = Eigen::SyclDevice(&queueInterface);
test_tanh_sycl<DataType, RowMajor, int64_t>(sycl_device);
test_tanh_sycl<DataType, ColMajor, int64_t>(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<DataType, RowMajor, int64_t>(sycl_device);
test_sigmoid_sycl<DataType, ColMajor, int64_t>(sycl_device);
#endif
}
EIGEN_DECLARE_TEST(cxx11_tensor_math_sycl) {

18
unsupported/test/cxx11_tensor_random_sycl.cpp
View File

@ -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<sizeDim0; i++)
for(IndexType j=1; j<sizeDim1; j++)
{
VERIFY_IS_NOT_EQUAL(out(i,j), out(i-1,j));
VERIFY_IS_NOT_EQUAL(out(i,j), out(i,j-1));
VERIFY_IS_NOT_EQUAL(out(i,j), out(i-1,j-1)); }
// For now we just check thes code doesn't crash.
// For now we just check the code doesn't crash.
// TODO: come up with a valid test of randomness
sycl_device.deallocate(d_out);
}
@ -66,16 +60,8 @@ void test_sycl_random_normal(const Eigen::SyclDevice& sycl_device)
Eigen::internal::NormalRandomGenerator<DataType> 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<sizeDim0; i++)
for(IndexType j=1; j<sizeDim1; j++)
{
VERIFY_IS_NOT_EQUAL(out(i,j), out(i-1,j));
VERIFY_IS_NOT_EQUAL(out(i,j), out(i,j-1));
VERIFY_IS_NOT_EQUAL(out(i,j), out(i-1,j-1));
}
// For now we just check thes code doesn't crash.
// For now we just check the code doesn't crash.
// TODO: come up with a valid test of randomness
sycl_device.deallocate(d_out);
}