Skip to content

[SYCL][Test] Devicelib test #1256

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 21 commits into from
Apr 15, 2020
Merged

[SYCL][Test] Devicelib test #1256

Show file tree
Hide file tree
Changes from 9 commits
Commits
Show all changes
21 commits
Select commit Hold shift + click to select a range
c3a5cd5
[SYCL] add test for all math function.
haonanya1 Mar 6, 2020
1816430
[SYCL] in llvm/sycl/source/detail/devicelib/CMakeLists.txt, libsycl-c…
haonanya1 Mar 6, 2020
5540691
[SYCL] Put all math function calls into a single kernel.
haonanya1 Mar 9, 2020
79dce3e
[SYCL] Remove using namespace std in math_utils.hpp
haonanya1 Mar 9, 2020
3836b41
[SYCL] modify std_complex_math_fp64_test.cpp and std_complex_math_tes…
haonanya1 Mar 9, 2020
e28efb4
[SYCL] [Test] Put all math function calls into a single kernel
haonanya1 Mar 11, 2020
1251e56
[SYCL] [Test] Deal with nan and inf input in is_about_FP function in …
haonanya1 Mar 11, 2020
9d62247
[SYCL] Format code
haonanya1 Mar 11, 2020
b670056
[SYCL] Code clean
haonanya1 Mar 13, 2020
fd16ef2
[SYCL] [Test] Improve readability for math_utils.hpp
haonanya1 Mar 16, 2020
764518e
Update sycl/test/devicelib/math_utils.hpp
haonanya Mar 18, 2020
10c938b
[SYCL] Add comment for is_about_FP function in math_utils.hpp
haonanya1 Mar 18, 2020
e8bfe40
[SYCL] [Test] Rename is_about_FP function and remove ret to use early…
haonanya1 Mar 19, 2020
014d3b2
[SYCL] [Test] Rename complex compare function.
haonanya1 Mar 19, 2020
6ef6fac
[SYCL] [Test] There is a rename for float point compare in math_utils…
haonanya1 Mar 19, 2020
0e84acf
Merge branch 'sycl' into devicelibTest
haonanya1 Mar 20, 2020
d6cd3fc
Merge branch 'sycl' into devicelibTest
haonanya1 Mar 23, 2020
88fc254
Fix use else after return issue
haonanya1 Apr 7, 2020
048e5d0
[SYCL][Test] Enable cmath lit test.
haonanya1 Apr 9, 2020
e8ac45e
[SYCL][Test] Enable Host device for math_override.cpp.
haonanya1 Apr 9, 2020
a54c5f1
[SYCL] [Test] Fix format issue.
haonanya1 Apr 9, 2020
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion sycl/source/detail/devicelib/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -64,7 +64,7 @@ add_custom_command(OUTPUT ${devicelib-obj-cmath}
${CMAKE_CURRENT_SOURCE_DIR}/cmath_wrapper.cpp
-o ${devicelib-obj-cmath}
MAIN_DEPENDENCY cmath_wrapper.cpp
DEPENDS device_complex.h clang
DEPENDS device_math.h clang
VERBATIM)

set(devicelib-obj-cmath-fp64 ${binary_dir}/libsycl-cmath-fp64.o)
Expand Down
113 changes: 113 additions & 0 deletions sycl/test/devicelib/cmath_fp64_test.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,113 @@
// UNSUPPORTED: windows
// RUN: %clangxx -fsycl -c %s -o %t.o
// RUN: %clangxx -fsycl %t.o %sycl_libs_dir/libsycl-cmath-fp64.o -o %t.out
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Please add a RUN line for execution here and in all other tests.
Otherwise this just tests for compilation and no runtime validation is performed.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

If these are going to run the compiled program, they should be moved to end-to-end testing.

Copy link
Contributor

@jinge90 jinge90 Apr 9, 2020 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Hi, @andykaylor and @asavonic
Could we enable the run command for cmath tests on CPU device currently?
For GPU, only OCL RT can work. For CPU, all cmath and complex tests should work in theory but OCL CPU RT has an issue which will lead to crash on some platforms, the fix has not been merged into latest OCL CPU RT, so we may encounter some unexpected crash if enable complex tests.
Thank you very much.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@andykaylor

If these are going to run the compiled program, they should be moved to end-to-end testing.

Most of the tests in sycl/test directory have RUN lines to test on available devices.
This test in particular has a lot of runtime asserts.

Could we enable the run command for cmath tests on CPU device currently?

I think they should be enabled for all devices you expect this code to work. I assume it should work for CPU and GPU, right? FPGA is probably not going work, because AOT compilation is not support for device libraries yet.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

If a test doesn't work on some platform due to a bug, it should be disabled temporarily.

#include <CL/sycl.hpp>
#include <cmath>
#include <iostream>
#include "math_utils.hpp"

namespace s = cl::sycl;
constexpr s::access::mode sycl_read = s::access::mode::read;
constexpr s::access::mode sycl_write = s::access::mode::write;

#define TEST_NUM 38

double ref[TEST_NUM] = {
1, 0, 0, 0, 0, 0, 0, 1, 1, 0.5,
0, 2, 0, 0, 1, 0, 2, 0, 0, 0,
0, 0, 1, 0, 1, 2, 0, 1, 2, 5,
0, 0, 0, 0, 0.5, 0.5, NAN, NAN,};

double refIptr = 1;

template <class T>
void device_cmath_test(s::queue &deviceQueue) {
s::range<1> numOfItems{TEST_NUM};
T result[TEST_NUM] = {-1};

// Variable exponent is an integer value to store the exponent in frexp function
int exponent = -1;

// Variable iptr stores the integral part of float point in modf function
T iptr = -1;

// Variable quo stores the sign and some bits of x/y in remquo function
int quo = -1;
{
s::buffer<T, 1> buffer1(result, numOfItems);
s::buffer<int, 1> buffer2(&exponent, s::range<1>{1});
s::buffer<T, 1> buffer3(&iptr, s::range<1>{1});
s::buffer<int, 1> buffer4(&quo, s::range<1>{1});
deviceQueue.submit([&](cl::sycl::handler &cgh) {
auto res_access = buffer1.template get_access<sycl_write>(cgh);
auto exp_access = buffer2.template get_access<sycl_write>(cgh);
auto iptr_access = buffer3.template get_access<sycl_write>(cgh);
auto quo_access = buffer4.template get_access<sycl_write>(cgh);
cgh.single_task<class DeviceMathTest>([=]() {
int i = 0;
res_access[i++] = std::cos(0.0);
res_access[i++] = std::sin(0.0);
res_access[i++] = std::log(1.0);
res_access[i++] = std::acos(1.0);
res_access[i++] = std::asin(0.0);
res_access[i++] = std::atan(0.0);
res_access[i++] = std::atan2(0.0, 1.0);
res_access[i++] = std::cosh(0.0);
res_access[i++] = std::exp(0.0);
res_access[i++] = std::fmod(1.5, 1.0);
res_access[i++] = std::frexp(0.0, &exp_access[0]);
res_access[i++] = std::ldexp(1.0, 1);
res_access[i++] = std::log10(1.0);
res_access[i++] = std::modf(1.0, &iptr_access[0]);
res_access[i++] = std::pow(1.0, 1.0);
res_access[i++] = std::sinh(0.0);
res_access[i++] = std::sqrt(4.0);
res_access[i++] = std::tan(0.0);
res_access[i++] = std::tanh(0.0);
res_access[i++] = std::acosh(1.0);
res_access[i++] = std::asinh(0.0);
res_access[i++] = std::atanh(0.0);
res_access[i++] = std::cbrt(1.0);
res_access[i++] = std::erf(0.0);
res_access[i++] = std::erfc(0.0);
res_access[i++] = std::exp2(1.0);
res_access[i++] = std::expm1(0.0);
res_access[i++] = std::fdim(1.0, 0.0);
res_access[i++] = std::fma(1.0, 1.0, 1.0);
res_access[i++] = std::hypot(3.0, 4.0);
res_access[i++] = std::ilogb(1.0);
res_access[i++] = std::log1p(0.0);
res_access[i++] = std::log2(1.0);
res_access[i++] = std::logb(1.0);
res_access[i++] = std::remainder(0.5, 1.0);
res_access[i++] = std::remquo(0.5, 1.0, &quo_access[0]);
T a = NAN;
res_access[i++] = std::tgamma(a);
res_access[i++] = std::lgamma(a);
});
});
}

// Compare result with reference
for (int i = 0; i < TEST_NUM; ++i) {
assert(is_about_FP(result[i], ref[i]));
}

// Test modf integral part
assert(is_about_FP(iptr, refIptr));

// Test frexp exponent
assert(exponent == 0);

// Test remquo sign
assert(quo == 0);
}

int main() {
s::queue deviceQueue;
if (deviceQueue.get_device().has_extension("cl_khr_fp64")) {
device_cmath_test<double>(deviceQueue);
std::cout << "Pass" << std::endl;
}
return 0;
}
129 changes: 78 additions & 51 deletions sycl/test/devicelib/cmath_test.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2,78 +2,105 @@
// RUN: %clangxx -fsycl -c %s -o %t.o
// RUN: %clangxx -fsycl %t.o %sycl_libs_dir/libsycl-cmath.o -o %t.out
#include <CL/sycl.hpp>
#include <iostream>
#include <cmath>
#include <iostream>
#include "math_utils.hpp"

namespace s = cl::sycl;
constexpr s::access::mode sycl_read = s::access::mode::read;
constexpr s::access::mode sycl_write = s::access::mode::write;

template <class T>
class DeviceCos;
#define TEST_NUM 38

template <class T>
void device_cos_test(s::queue &deviceQueue) {
s::range<1> numOfItems{1};
T result = -1;
{
s::buffer<T, 1> buffer1(&result, numOfItems);
deviceQueue.submit([&](cl::sycl::handler &cgh) {
auto res_access = buffer1.template get_access<sycl_write>(cgh);
cgh.single_task<class DeviceCos<T> >([=]() {
res_access[0] = std::cos(0);
});
});
}
float ref[TEST_NUM] = {
1, 0, 0, 0, 0, 0, 0, 1, 1, 0.5,
0, 2, 0, 0, 1, 0, 2, 0, 0, 0,
0, 0, 1, 0, 1, 2, 0, 1, 2, 5,
0, 0, 0, 0, 0.5, 0.5, NAN, NAN,};

assert(result == 1);
}
float refIptr = 1;

template <class T>
class DeviceSin;

template <class T>
void device_sin_test(s::queue &deviceQueue) {
s::range<1> numOfItems{1};
T result = -1;
{
s::buffer<T, 1> buffer1(&result, numOfItems);
deviceQueue.submit([&](cl::sycl::handler &cgh) {
auto res_access = buffer1.template get_access<sycl_write>(cgh);
cgh.single_task<class DeviceSin<T> >([=]() {
res_access[0] = std::sin(0);
});
});
}
void device_cmath_test(s::queue &deviceQueue) {
s::range<1> numOfItems{TEST_NUM};
T result[TEST_NUM] = {-1};

assert(result == 0);
}
// Variable exponent is an integer value to store the exponent in frexp function
int exponent = -1;

template <class T>
class DeviceLog;
// Variable iptr stores the integral part of float point in modf function
T iptr = -1;

template <class T>
void device_log_test(s::queue &deviceQueue) {
s::range<1> numOfItems{1};
T result = -1;
// Variable quo stores the sign and some bits of x/y in remquo function
int quo = -1;
{
s::buffer<T, 1> buffer1(&result, numOfItems);
s::buffer<T, 1> buffer1(result, numOfItems);
s::buffer<int, 1> buffer2(&exponent, s::range<1>{1});
s::buffer<T, 1> buffer3(&iptr, s::range<1>{1});
s::buffer<int, 1> buffer4(&quo, s::range<1>{1});
deviceQueue.submit([&](cl::sycl::handler &cgh) {
auto res_access = buffer1.template get_access<sycl_write>(cgh);
cgh.single_task<class DeviceLog<T> >([=]() {
res_access[0] = std::log(1);
auto exp_access = buffer2.template get_access<sycl_write>(cgh);
auto iptr_access = buffer3.template get_access<sycl_write>(cgh);
auto quo_access = buffer4.template get_access<sycl_write>(cgh);
cgh.single_task<class DeviceMathTest>([=]() {
int i = 0;
res_access[i++] = std::cos(0.0f);
res_access[i++] = std::sin(0.0f);
res_access[i++] = std::log(1.0f);
res_access[i++] = std::acos(1.0f);
res_access[i++] = std::asin(0.0f);
res_access[i++] = std::atan(0.0f);
res_access[i++] = std::atan2(0.0f, 1.0f);
res_access[i++] = std::cosh(0.0f);
res_access[i++] = std::exp(0.0f);
res_access[i++] = std::fmod(1.5f, 1.0f);
res_access[i++] = std::frexp(0.0f, &exp_access[0]);
res_access[i++] = std::ldexp(1.0f, 1);
res_access[i++] = std::log10(1.0f);
res_access[i++] = std::modf(1.0f, &iptr_access[0]);
res_access[i++] = std::pow(1.0f, 1.0f);
res_access[i++] = std::sinh(0.0f);
res_access[i++] = std::sqrt(4.0f);
res_access[i++] = std::tan(0.0f);
res_access[i++] = std::tanh(0.0f);
res_access[i++] = std::acosh(1.0f);
res_access[i++] = std::asinh(0.0f);
res_access[i++] = std::atanh(0.0f);
res_access[i++] = std::cbrt(1.0f);
res_access[i++] = std::erf(0.0f);
res_access[i++] = std::erfc(0.0f);
res_access[i++] = std::exp2(1.0f);
res_access[i++] = std::expm1(0.0f);
res_access[i++] = std::fdim(1.0f, 0.0f);
res_access[i++] = std::fma(1.0f, 1.0f, 1.0f);
res_access[i++] = std::hypot(3.0f, 4.0f);
res_access[i++] = std::ilogb(1.0f);
res_access[i++] = std::log1p(0.0f);
res_access[i++] = std::log2(1.0f);
res_access[i++] = std::logb(1.0f);
res_access[i++] = std::remainder(0.5f, 1.0f);
res_access[i++] = std::remquo(0.5f, 1.0f, &quo_access[0]);
T a = NAN;
res_access[i++] = std::tgamma(a);
res_access[i++] = std::lgamma(a);
});
});
}

assert(result == 0);
}
// Compare result with reference
for (int i = 0; i < TEST_NUM; ++i) {
assert(is_about_FP(result[i], ref[i]));
}

template <class T>
void device_cmath_test(s::queue &deviceQueue) {
device_cos_test<T>(deviceQueue);
device_sin_test<T>(deviceQueue);
device_log_test<T>(deviceQueue);
// Test modf integral part
assert(is_about_FP(iptr, refIptr));

// Test frexp exponent
assert(exponent == 0);

// Test remquo sign
assert(quo == 0);
}

int main() {
Expand Down
86 changes: 0 additions & 86 deletions sycl/test/devicelib/cmath_test_fp64.cpp
View file
Open in desktop

This file was deleted.

Loading