blob: 34ef38f48cf2e145033ab4e8ad0bee082214178b [file]
// SPDX-FileCopyrightText: The Eigen Authors
// SPDX-License-Identifier: MPL-2.0
// EIGEN_SUFFIXES;1;2
// Core use inside a SYCL kernel does not require EIGEN_USE_SYCL.
#if defined(EIGEN_TEST_PART_2)
#include <Eigen/Core>
#endif
#include <CL/sycl.hpp>
#include <Eigen/Core>
#include <contrib/Eigen/SpecialFunctions>
template <typename Scalar>
bool sycl_first_vector_addition() {
using Vector = Eigen::Matrix<Scalar, 4, 1>;
Vector a, b, result, expected;
a << 1, 2, 3, 4;
b << 5, 6, 7, 8;
expected << 6, 8, 10, 12;
cl::sycl::queue queue;
{
cl::sycl::buffer<Scalar, 1> input_a(a.data(), cl::sycl::range<1>(4));
cl::sycl::buffer<Scalar, 1> input_b(b.data(), cl::sycl::range<1>(4));
cl::sycl::buffer<Scalar, 1> output(result.data(), cl::sycl::range<1>(4));
queue.submit([&](cl::sycl::handler& handler) {
auto x = input_a.template get_access<cl::sycl::access::mode::read>(handler);
auto y = input_b.template get_access<cl::sycl::access::mode::read>(handler);
auto z = output.template get_access<cl::sycl::access::mode::discard_write>(handler);
handler.single_task([=]() {
Eigen::Map<const Vector> a_map(&x[0]), b_map(&y[0]);
Eigen::Map<Vector> result_map(&z[0]);
result_map = a_map + b_map;
});
});
queue.wait_and_throw();
}
return result == expected;
}
template <typename Scalar>
bool sycl_scalar_math() {
using Vector = Eigen::Matrix<Scalar, 4, 1>;
Vector input, result_sin, result_lgamma;
input << 1, 2, 3, 4;
cl::sycl::queue queue;
{
cl::sycl::buffer<Scalar, 1> input_buffer(input.data(), cl::sycl::range<1>(4));
cl::sycl::buffer<Scalar, 1> output_sin(result_sin.data(), cl::sycl::range<1>(4));
cl::sycl::buffer<Scalar, 1> output_lgamma(result_lgamma.data(), cl::sycl::range<1>(4));
queue.submit([&](cl::sycl::handler& handler) {
auto x = input_buffer.template get_access<cl::sycl::access::mode::read>(handler);
auto y = output_sin.template get_access<cl::sycl::access::mode::discard_write>(handler);
auto z = output_lgamma.template get_access<cl::sycl::access::mode::discard_write>(handler);
handler.single_task([=]() {
Eigen::Map<const Vector> input_map(&x[0]);
Eigen::Map<Vector> sin_map(&y[0]);
sin_map = input_map.array().sin().matrix();
for (Eigen::Index i = 0; i < 4; ++i) z[i] = Eigen::numext::lgamma(input_map[i]);
});
});
queue.wait_and_throw();
}
const Scalar tolerance = Scalar(32) * Eigen::NumTraits<Scalar>::epsilon();
for (Eigen::Index i = 0; i < 4; ++i) {
const Scalar expected_sin = std::sin(input[i]);
const Scalar expected_lgamma = std::lgamma(input[i]);
if (!(std::abs(result_sin[i] - expected_sin) <= tolerance * (Scalar(1) + std::abs(expected_sin)))) return false;
if (!(std::abs(result_lgamma[i] - expected_lgamma) <= tolerance * (Scalar(1) + std::abs(expected_lgamma))))
return false;
}
return true;
}
template <typename Scalar>
bool sycl_low_precision_fma() {
float input[] = {1.5f, 2.f, 0.5f}, result = 0;
cl::sycl::queue queue;
{
cl::sycl::buffer<float, 1> inputs(input, cl::sycl::range<1>(3));
cl::sycl::buffer<float, 1> output(&result, cl::sycl::range<1>(1));
queue.submit([&](cl::sycl::handler& handler) {
auto x = inputs.template get_access<cl::sycl::access::mode::read>(handler);
auto y = output.template get_access<cl::sycl::access::mode::discard_write>(handler);
handler.single_task([=]() { y[0] = float(Eigen::numext::fma(Scalar(x[0]), Scalar(x[1]), Scalar(x[2]))); });
});
queue.wait_and_throw();
}
const float expected =
float(Scalar(std::fma(float(Scalar(input[0])), float(Scalar(input[1])), float(Scalar(input[2])))));
return result == expected;
}
// The shared test driver explicitly enables CUDA when __CUDACC__ is defined;
// this test must leave backend selection to Eigen, as a user translation unit does.
int main() {
if (!sycl_first_vector_addition<float>() || !sycl_scalar_math<float>() || !sycl_low_precision_fma<Eigen::half>() ||
!sycl_low_precision_fma<Eigen::bfloat16>())
return 1;
cl::sycl::queue queue;
if (queue.get_device().has(cl::sycl::aspect::fp64) &&
(!sycl_first_vector_addition<double>() || !sycl_scalar_math<double>()))
return 1;
return 0;
}