blob: c8668d847b613c5e8123845c35671735ad5cf2d6 [file] [log] [blame]
// This file is part of Eigen, a lightweight C++ template library
// for linear algebra.
//
// Copyright (C) 2016
// Mehdi Goli Codeplay Software Ltd.
// Ralph Potter Codeplay Software Ltd.
// Luke Iwanski Codeplay Software Ltd.
// Contact: <eigen@codeplay.com>
//
// This Source Code Form is subject to the terms of the Mozilla
// Public License v. 2.0. If a copy of the MPL was not distributed
// with this file, You can obtain one at http://mozilla.org/MPL/2.0/.
#define EIGEN_TEST_NO_LONGDOUBLE
#define EIGEN_TEST_NO_COMPLEX
#define EIGEN_DEFAULT_DENSE_INDEX_TYPE int64_t
#define EIGEN_USE_SYCL
#include "main.h"
#include <unsupported/Eigen/CXX11/Tensor>
using Eigen::array;
using Eigen::SyclDevice;
using Eigen::Tensor;
using Eigen::TensorMap;
// Functions used to compare the TensorMap implementation on the device with
// the equivalent on the host
namespace SYCL {
template <typename T> T abs(T x) {
return cl::sycl::abs(x);
}
template <> Eigen::half abs(Eigen::half x) {
return Eigen::half(cl::sycl::fabs(static_cast<cl::sycl::half>(x)));
}
template <> float abs(float x) {
return cl::sycl::fabs(x);
}
template <> double abs(double x) {
return cl::sycl::fabs(x);
}
template <typename T> T square(T x) { return x * x; }
template <typename T> T cube(T x) { return x * x * x; }
template <typename T> T inverse(T x) { return T(1) / x; }
template <typename T> T cwiseMax(T x, T y) {
return cl::sycl::max(x, y);
}
template <> Eigen::half cwiseMax(Eigen::half x, Eigen::half y) {
return Eigen::half(cl::sycl::max(static_cast<cl::sycl::half>(x), static_cast<cl::sycl::half>(y)));
}
template <typename T> T cwiseMin(T x, T y) {
return cl::sycl::min(x, y);
}
template <> Eigen::half cwiseMin(Eigen::half x, Eigen::half y) {
return Eigen::half(cl::sycl::min(static_cast<cl::sycl::half>(x), static_cast<cl::sycl::half>(y)));
}
template <typename T> T sqrt(T x) {
return cl::sycl::sqrt(x);
}
template <> Eigen::half sqrt(Eigen::half x) {
return Eigen::half(cl::sycl::sqrt(static_cast<cl::sycl::half>(x)));
}
template <typename T> T rsqrt(T x) {
return cl::sycl::rsqrt(x);
}
template <> Eigen::half rsqrt(Eigen::half x) {
return Eigen::half(cl::sycl::rsqrt(static_cast<cl::sycl::half>(x)));
}
template <typename T> T tanh(T x) {
return cl::sycl::tanh(x);
}
template <> Eigen::half tanh(Eigen::half x) {
return Eigen::half(cl::sycl::tanh(static_cast<cl::sycl::half>(x)));
}
template <typename T> T exp(T x) {
return cl::sycl::exp(x);
}
template <> Eigen::half exp(Eigen::half x) {
return Eigen::half(cl::sycl::exp(static_cast<cl::sycl::half>(x)));
}
template <typename T> T expm1(T x) {
return cl::sycl::expm1(x);
}
template <> Eigen::half expm1(Eigen::half x) {
return Eigen::half(cl::sycl::expm1(static_cast<cl::sycl::half>(x)));
}
template <typename T> T log(T x) {
return cl::sycl::log(x);
}
template <> Eigen::half log(Eigen::half x) {
return Eigen::half(cl::sycl::log(static_cast<cl::sycl::half>(x)));
}
template <typename T> T ceil(T x) {
return cl::sycl::ceil(x);
}
template <> Eigen::half ceil(Eigen::half x) {
return Eigen::half(cl::sycl::ceil(static_cast<cl::sycl::half>(x)));
}
template <typename T> T floor(T x) {
return cl::sycl::floor(x);
}
template <> Eigen::half floor(Eigen::half x) {
return Eigen::half(cl::sycl::floor(static_cast<cl::sycl::half>(x)));
}
template <typename T> T round(T x) {
return cl::sycl::round(x);
}
template <> Eigen::half round(Eigen::half x) {
return Eigen::half(cl::sycl::round(static_cast<cl::sycl::half>(x)));
}
template <typename T> T log1p(T x) {
return cl::sycl::log1p(x);
}
template <> Eigen::half log1p(Eigen::half x) {
return Eigen::half(cl::sycl::log1p(static_cast<cl::sycl::half>(x)));
}
template <typename T> T sign(T x) {
return cl::sycl::sign(x);
}
template <> Eigen::half sign(Eigen::half x) {
return Eigen::half(cl::sycl::sign(static_cast<cl::sycl::half>(x)));
}
template <typename T> T isnan(T x) {
return cl::sycl::isnan(x);
}
template <> Eigen::half isnan(Eigen::half x) {
return Eigen::half(cl::sycl::isnan(static_cast<cl::sycl::half>(x)));
}
template <typename T> T isfinite(T x) {
return cl::sycl::isfinite(x);
}
template <> Eigen::half isfinite(Eigen::half x) {
return Eigen::half(cl::sycl::isfinite(static_cast<cl::sycl::half>(x)));
}
template <typename T> T isinf(T x) {
return cl::sycl::isinf(x);
}
template <> Eigen::half isinf(Eigen::half x) {
return Eigen::half(cl::sycl::isinf(static_cast<cl::sycl::half>(x)));
}
}
#define DECLARE_UNARY_STRUCT_NON_SYCL(FUNC) \
struct op_##FUNC { \
template <typename T> \
auto operator()(const T& x) { \
return SYCL::FUNC(x); \
} \
template <typename T> \
auto operator()(const TensorMap<T>& x) { \
return x.FUNC(); \
} \
};
DECLARE_UNARY_STRUCT_NON_SYCL(abs)
DECLARE_UNARY_STRUCT_NON_SYCL(square)
DECLARE_UNARY_STRUCT_NON_SYCL(cube)
DECLARE_UNARY_STRUCT_NON_SYCL(inverse)
#define DECLARE_BINARY_STRUCT_NON_SYCL(FUNC) \
struct op_##FUNC { \
template <typename T1, typename T2> \
auto operator()(const T1& x, const T2& y){ \
return SYCL::FUNC(x, y); \
} \
template <typename T1, typename T2> \
auto operator()(const TensorMap<T1>& x, const TensorMap<T2>& y) { \
return x.FUNC(y); \
} \
};
DECLARE_BINARY_STRUCT_NON_SYCL(cwiseMax)
DECLARE_BINARY_STRUCT_NON_SYCL(cwiseMin)
struct EqualAssignment {
template <typename Lhs, typename Rhs>
void operator()(Lhs& lhs, const Rhs& rhs) { lhs = rhs; }
};
struct PlusEqualAssignment {
template <typename Lhs, typename Rhs>
void operator()(Lhs& lhs, const Rhs& rhs) { lhs += rhs; }
};
template <typename DataType, int DataLayout,
typename Assignment, typename Operator>
void test_unary_builtins_for_scalar(const Eigen::SyclDevice& sycl_device,
const array<int64_t, 3>& tensor_range) {
Operator op;
Assignment asgn;
{
/* Assignment(out, Operator(in)) */
Tensor<DataType, 3, DataLayout, int64_t> in(tensor_range);
Tensor<DataType, 3, DataLayout, int64_t> out(tensor_range);
in = in.random() + DataType(0.01);
out = out.random() + DataType(0.01);
Tensor<DataType, 3, DataLayout, int64_t> reference(out);
DataType *gpu_data = static_cast<DataType *>(
sycl_device.allocate(in.size() * sizeof(DataType)));
DataType *gpu_data_out = static_cast<DataType *>(
sycl_device.allocate(out.size() * sizeof(DataType)));
TensorMap<Tensor<DataType, 3, DataLayout, int64_t>> gpu(gpu_data, tensor_range);
TensorMap<Tensor<DataType, 3, DataLayout, int64_t>> gpu_out(gpu_data_out, tensor_range);
sycl_device.memcpyHostToDevice(gpu_data, in.data(),
(in.size()) * sizeof(DataType));
sycl_device.memcpyHostToDevice(gpu_data_out, out.data(),
(out.size()) * sizeof(DataType));
auto device_expr = gpu_out.device(sycl_device);
asgn(device_expr, op(gpu));
sycl_device.memcpyDeviceToHost(out.data(), gpu_data_out,
(out.size()) * sizeof(DataType));
for (int64_t i = 0; i < out.size(); ++i) {
DataType ver = reference(i);
asgn(ver, op(in(i)));
VERIFY_IS_APPROX(out(i), ver);
}
sycl_device.deallocate(gpu_data);
sycl_device.deallocate(gpu_data_out);
}
{
/* Assignment(out, Operator(out)) */
Tensor<DataType, 3, DataLayout, int64_t> out(tensor_range);
// 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)));
TensorMap<Tensor<DataType, 3, DataLayout, int64_t>> gpu_out(gpu_data_out, tensor_range);
sycl_device.memcpyHostToDevice(gpu_data_out, out.data(),
(out.size()) * sizeof(DataType));
auto device_expr = gpu_out.device(sycl_device);
asgn(device_expr, op(gpu_out));
sycl_device.memcpyDeviceToHost(out.data(), gpu_data_out,
(out.size()) * sizeof(DataType));
for (int64_t i = 0; i < out.size(); ++i) {
DataType ver = reference(i);
asgn(ver, op(reference(i)));
VERIFY_IS_APPROX(out(i), ver);
}
sycl_device.deallocate(gpu_data_out);
}
}
#define DECLARE_UNARY_STRUCT(FUNC) \
struct op_##FUNC { \
template <typename T> \
auto operator()(const T& x) -> decltype(SYCL::FUNC(x)) { \
return SYCL::FUNC(x); \
} \
template <typename T> \
auto operator()(const TensorMap<T>& x) -> decltype(x.FUNC()) { \
return x.FUNC(); \
} \
};
DECLARE_UNARY_STRUCT(sqrt)
DECLARE_UNARY_STRUCT(rsqrt)
DECLARE_UNARY_STRUCT(tanh)
DECLARE_UNARY_STRUCT(exp)
DECLARE_UNARY_STRUCT(expm1)
DECLARE_UNARY_STRUCT(log)
DECLARE_UNARY_STRUCT(ceil)
DECLARE_UNARY_STRUCT(floor)
DECLARE_UNARY_STRUCT(round)
DECLARE_UNARY_STRUCT(log1p)
DECLARE_UNARY_STRUCT(sign)
DECLARE_UNARY_STRUCT(isnan)
DECLARE_UNARY_STRUCT(isfinite)
DECLARE_UNARY_STRUCT(isinf)
template <typename DataType, int DataLayout, typename Assignment>
void test_unary_builtins_for_assignement(const Eigen::SyclDevice& sycl_device,
const array<int64_t, 3>& tensor_range) {
#define RUN_UNARY_TEST(FUNC) \
test_unary_builtins_for_scalar<DataType, DataLayout, Assignment, \
op_##FUNC>(sycl_device, tensor_range)
RUN_UNARY_TEST(abs);
RUN_UNARY_TEST(sqrt);
RUN_UNARY_TEST(rsqrt);
RUN_UNARY_TEST(square);
RUN_UNARY_TEST(cube);
RUN_UNARY_TEST(inverse);
RUN_UNARY_TEST(tanh);
RUN_UNARY_TEST(exp);
RUN_UNARY_TEST(expm1);
RUN_UNARY_TEST(log);
RUN_UNARY_TEST(ceil);
RUN_UNARY_TEST(floor);
RUN_UNARY_TEST(round);
RUN_UNARY_TEST(log1p);
RUN_UNARY_TEST(sign);
}
template <typename DataType, int DataLayout, typename Operator>
void test_unary_builtins_return_bool(const Eigen::SyclDevice& sycl_device,
const array<int64_t, 3>& tensor_range) {
/* out = op(in) */
Operator op;
Tensor<DataType, 3, DataLayout, int64_t> in(tensor_range);
Tensor<bool, 3, DataLayout, int64_t> out(tensor_range);
in = in.random() + DataType(0.01);
DataType *gpu_data = static_cast<DataType *>(
sycl_device.allocate(in.size() * sizeof(DataType)));
bool *gpu_data_out =
static_cast<bool *>(sycl_device.allocate(out.size() * sizeof(bool)));
TensorMap<Tensor<DataType, 3, DataLayout, int64_t>> gpu(gpu_data, tensor_range);
TensorMap<Tensor<bool, 3, DataLayout, int64_t>> gpu_out(gpu_data_out, tensor_range);
sycl_device.memcpyHostToDevice(gpu_data, in.data(),
(in.size()) * sizeof(DataType));
gpu_out.device(sycl_device) = op(gpu);
sycl_device.memcpyDeviceToHost(out.data(), gpu_data_out,
(out.size()) * sizeof(bool));
for (int64_t i = 0; i < out.size(); ++i) {
VERIFY_IS_EQUAL(out(i), op(in(i)));
}
sycl_device.deallocate(gpu_data);
sycl_device.deallocate(gpu_data_out);
}
template <typename DataType, int DataLayout>
void test_unary_builtins(const Eigen::SyclDevice& sycl_device,
const array<int64_t, 3>& tensor_range) {
test_unary_builtins_for_assignement<DataType, DataLayout,
PlusEqualAssignment>(sycl_device, tensor_range);
test_unary_builtins_for_assignement<DataType, DataLayout,
EqualAssignment>(sycl_device, tensor_range);
test_unary_builtins_return_bool<DataType, DataLayout,
op_isnan>(sycl_device, tensor_range);
test_unary_builtins_return_bool<DataType, DataLayout,
op_isfinite>(sycl_device, tensor_range);
test_unary_builtins_return_bool<DataType, DataLayout,
op_isinf>(sycl_device, tensor_range);
}
template <typename DataType>
static void test_builtin_unary_sycl(const Eigen::SyclDevice &sycl_device) {
int64_t sizeDim1 = 10;
int64_t sizeDim2 = 10;
int64_t sizeDim3 = 10;
array<int64_t, 3> tensor_range = {{sizeDim1, sizeDim2, sizeDim3}};
test_unary_builtins<DataType, RowMajor>(sycl_device, tensor_range);
test_unary_builtins<DataType, ColMajor>(sycl_device, tensor_range);
}
template <typename DataType, int DataLayout, typename Operator>
void test_binary_builtins_func(const Eigen::SyclDevice& sycl_device,
const array<int64_t, 3>& tensor_range) {
/* out = op(in_1, in_2) */
Operator op;
Tensor<DataType, 3, DataLayout, int64_t> in_1(tensor_range);
Tensor<DataType, 3, DataLayout, int64_t> in_2(tensor_range);
Tensor<DataType, 3, DataLayout, int64_t> out(tensor_range);
in_1 = in_1.random() + DataType(0.01);
in_2 = in_2.random() + DataType(0.01);
Tensor<DataType, 3, DataLayout, int64_t> reference(out);
DataType *gpu_data_1 = static_cast<DataType *>(
sycl_device.allocate(in_1.size() * sizeof(DataType)));
DataType *gpu_data_2 = static_cast<DataType *>(
sycl_device.allocate(in_2.size() * sizeof(DataType)));
DataType *gpu_data_out = static_cast<DataType *>(
sycl_device.allocate(out.size() * sizeof(DataType)));
TensorMap<Tensor<DataType, 3, DataLayout, int64_t>> gpu_1(gpu_data_1, tensor_range);
TensorMap<Tensor<DataType, 3, DataLayout, int64_t>> gpu_2(gpu_data_2, tensor_range);
TensorMap<Tensor<DataType, 3, DataLayout, int64_t>> gpu_out(gpu_data_out, tensor_range);
sycl_device.memcpyHostToDevice(gpu_data_1, in_1.data(),
(in_1.size()) * sizeof(DataType));
sycl_device.memcpyHostToDevice(gpu_data_2, in_2.data(),
(in_2.size()) * sizeof(DataType));
gpu_out.device(sycl_device) = op(gpu_1, gpu_2);
sycl_device.memcpyDeviceToHost(out.data(), gpu_data_out,
(out.size()) * sizeof(DataType));
for (int64_t i = 0; i < out.size(); ++i) {
VERIFY_IS_APPROX(out(i), op(in_1(i), in_2(i)));
}
sycl_device.deallocate(gpu_data_1);
sycl_device.deallocate(gpu_data_2);
sycl_device.deallocate(gpu_data_out);
}
template <typename DataType, int DataLayout, typename Operator>
void test_binary_builtins_fixed_arg2(const Eigen::SyclDevice& sycl_device,
const array<int64_t, 3>& tensor_range) {
/* out = op(in_1, 2) */
Operator op;
const DataType arg2(2);
Tensor<DataType, 3, DataLayout, int64_t> in_1(tensor_range);
Tensor<DataType, 3, DataLayout, int64_t> out(tensor_range);
in_1 = in_1.random();
Tensor<DataType, 3, DataLayout, int64_t> reference(out);
DataType *gpu_data_1 = static_cast<DataType *>(
sycl_device.allocate(in_1.size() * sizeof(DataType)));
DataType *gpu_data_out = static_cast<DataType *>(
sycl_device.allocate(out.size() * sizeof(DataType)));
TensorMap<Tensor<DataType, 3, DataLayout, int64_t>> gpu_1(gpu_data_1, tensor_range);
TensorMap<Tensor<DataType, 3, DataLayout, int64_t>> gpu_out(gpu_data_out, tensor_range);
sycl_device.memcpyHostToDevice(gpu_data_1, in_1.data(),
(in_1.size()) * sizeof(DataType));
gpu_out.device(sycl_device) = op(gpu_1, arg2);
sycl_device.memcpyDeviceToHost(out.data(), gpu_data_out,
(out.size()) * sizeof(DataType));
for (int64_t i = 0; i < out.size(); ++i) {
VERIFY_IS_APPROX(out(i), op(in_1(i), arg2));
}
sycl_device.deallocate(gpu_data_1);
sycl_device.deallocate(gpu_data_out);
}
#define DECLARE_BINARY_STRUCT(FUNC) \
struct op_##FUNC { \
template <typename T1, typename T2> \
auto operator()(const T1& x, const T2& y) -> decltype(cl::sycl::FUNC(x, y)) { \
return cl::sycl::FUNC(x, y); \
} \
template <typename T1, typename T2> \
auto operator()(const TensorMap<T1>& x, const TensorMap<T2>& y) -> decltype(x.FUNC(y)) { \
return x.FUNC(y); \
} \
};
#define DECLARE_BINARY_STRUCT_OP(NAME, OPERATOR) \
struct op_##NAME { \
template <typename T1, typename T2> \
auto operator()(const T1& x, const T2& y) -> decltype(x OPERATOR y) { \
return x OPERATOR y; \
} \
};
DECLARE_BINARY_STRUCT_OP(plus, +)
DECLARE_BINARY_STRUCT_OP(minus, -)
DECLARE_BINARY_STRUCT_OP(times, *)
DECLARE_BINARY_STRUCT_OP(divide, /)
DECLARE_BINARY_STRUCT_OP(modulo, %)
template <typename DataType, int DataLayout>
void test_binary_builtins(const Eigen::SyclDevice& sycl_device,
const array<int64_t, 3>& tensor_range) {
test_binary_builtins_func<DataType, DataLayout,
op_cwiseMax>(sycl_device, tensor_range);
test_binary_builtins_func<DataType, DataLayout,
op_cwiseMin>(sycl_device, tensor_range);
test_binary_builtins_func<DataType, DataLayout,
op_plus>(sycl_device, tensor_range);
test_binary_builtins_func<DataType, DataLayout,
op_minus>(sycl_device, tensor_range);
test_binary_builtins_func<DataType, DataLayout,
op_times>(sycl_device, tensor_range);
test_binary_builtins_func<DataType, DataLayout,
op_divide>(sycl_device, tensor_range);
}
template <typename DataType>
static void test_floating_builtin_binary_sycl(const Eigen::SyclDevice &sycl_device) {
int64_t sizeDim1 = 10;
int64_t sizeDim2 = 10;
int64_t sizeDim3 = 10;
array<int64_t, 3> tensor_range = {{sizeDim1, sizeDim2, sizeDim3}};
test_binary_builtins<DataType, RowMajor>(sycl_device, tensor_range);
test_binary_builtins<DataType, ColMajor>(sycl_device, tensor_range);
}
template <typename DataType>
static void test_integer_builtin_binary_sycl(const Eigen::SyclDevice &sycl_device) {
int64_t sizeDim1 = 10;
int64_t sizeDim2 = 10;
int64_t sizeDim3 = 10;
array<int64_t, 3> tensor_range = {{sizeDim1, sizeDim2, sizeDim3}};
test_binary_builtins_fixed_arg2<DataType, RowMajor,
op_modulo>(sycl_device, tensor_range);
test_binary_builtins_fixed_arg2<DataType, ColMajor,
op_modulo>(sycl_device, tensor_range);
}
EIGEN_DECLARE_TEST(cxx11_tensor_builtins_sycl) {
for (const auto& device :Eigen::get_sycl_supported_devices()) {
QueueInterface queueInterface(device);
Eigen::SyclDevice sycl_device(&queueInterface);
CALL_SUBTEST_1(test_builtin_unary_sycl<half>(sycl_device));
CALL_SUBTEST_2(test_floating_builtin_binary_sycl<half>(sycl_device));
CALL_SUBTEST_3(test_builtin_unary_sycl<float>(sycl_device));
CALL_SUBTEST_4(test_floating_builtin_binary_sycl<float>(sycl_device));
CALL_SUBTEST_5(test_integer_builtin_binary_sycl<int>(sycl_device));
}
}