Skip to content
This repository was archived by the owner on Mar 28, 2023. It is now read-only.

[SYCL] Add --ffast-math tests #919

Merged
merged 1 commit into from
Jun 24, 2022
Merged

[SYCL] Add --ffast-math tests #919

Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
82 changes: 82 additions & 0 deletions SYCL/DeviceLib/built-ins/fast-math-flag.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,82 @@
// RUN: %clangxx -fsycl -ffast-math -fsycl-targets=%sycl_triple %s -o %t.out
// RUN: %HOST_RUN_PLACEHOLDER %t.out
// RUN: %CPU_RUN_PLACEHOLDER %t.out
// RUN: %GPU_RUN_PLACEHOLDER %t.out
// RUN: %ACC_RUN_PLACEHOLDER %t.out

#include <CL/sycl.hpp>
#include <cassert>

#define __TEST_FFMATH_BINARY(func) \
int test_ffmath_##func() { \
sycl::float4 r[2]; \
sycl::float4 val[2] = {{1.0004f, 1e-4f, 1.4f, 14.0f}, \
{1.0004f, 1e-4f, 1.4f, 14.0f}}; \
{ \
sycl::buffer<sycl::float4, 1> output(&r[0], sycl::range<1>(2)); \
sycl::buffer<sycl::float4, 1> input(&val[0], sycl::range<1>(2)); \
sycl::queue q; \
q.submit([&](sycl::handler &cgh) { \
auto AccO = \
output.template get_access<sycl::access::mode::write>(cgh); \
auto AccI = input.template get_access<sycl::access::mode::read>(cgh); \
cgh.single_task([=]() { \
AccO[0] = sycl::func(AccI[0], AccI[1]); \
AccO[1] = sycl::native::func(AccI[0], AccI[1]); \
}); \
}); \
} \
return sycl::all(r[0] == r[1]); \
Copy link

@s-kanaev s-kanaev Mar 14, 2022

Choose a reason for hiding this comment

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

I can't get purpose of sycl::all for bool converted for integer here.
What's purpose of the test?

Copy link
Author

@pgorlani pgorlani Mar 14, 2022 •
edited
Loading

Choose a reason for hiding this comment

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

The purpose of the test is to check if all vector components are equal implying that the ::native function and the regular one return the same values.

Let me know if it is fine, otherwise I can change it.

Copy link

@s-kanaev s-kanaev Mar 14, 2022

Choose a reason for hiding this comment

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

Per SYCL 2020 spec, all(igeninter) Returns 1 if the most significant bit in all components of x is set; otherwise returns 0.
On line 29 a boolean is provided instead.

Copy link
Author

@pgorlani pgorlani Mar 14, 2022 •
edited
Loading

Choose a reason for hiding this comment

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

Thanks for pointing this out, I misunderstood the initial message.

I think that the operator== returns a signed integer vector of size 4, in which each component is equal to -1 (all bits are set) if the corresponding elements of r[0] and r[1] have the same value, 0 otherwise (no bit is set). For reference, see Table 141.

Copy link

@s-kanaev s-kanaev Mar 16, 2022

Choose a reason for hiding this comment

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

Right. operator== on host return vector of rel_t which is defined here https://github.com/intel/llvm/blob/cb5e8ae54d2e55e9edc0bad35b57fcad39e4f090/sycl/include/CL/sycl/types.hpp#L124

Copy link

@s-kanaev s-kanaev Mar 16, 2022

Choose a reason for hiding this comment

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

For char, short, int and long it's cahr short int and long respectively. For any other type it's a mere boolean.

Copy link
Author

@pgorlani pgorlani Mar 16, 2022

Choose a reason for hiding this comment

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

It depends on the size of the type. So, in the case of floats the returned vector type is made of cl_int as mandated by the SYCL specifications linked in my previous message.

}

#define __TEST_FFMATH_UNARY(func) \
int test_ffmath_##func() { \
sycl::float4 val = {1.0004f, 1e-4f, 1.4f, 14.0f}; \
sycl::float4 r[2]; \
{ \
sycl::buffer<sycl::float4, 1> output(&r[0], sycl::range<1>(2)); \
sycl::buffer<sycl::float4, 1> input(&val, sycl::range<1>(1)); \
sycl::queue q; \
q.submit([&](sycl::handler &cgh) { \
auto AccO = \
output.template get_access<sycl::access::mode::write>(cgh); \
auto AccI = input.template get_access<sycl::access::mode::read>(cgh); \
cgh.single_task([=]() { \
AccO[0] = sycl::func(AccI[0]); \
AccO[1] = sycl::native::func(AccI[0]); \
}); \
}); \
} \
return sycl::all(r[0] == r[1]); \
}

__TEST_FFMATH_UNARY(cos)
__TEST_FFMATH_UNARY(exp)
__TEST_FFMATH_UNARY(exp2)
__TEST_FFMATH_UNARY(exp10)
__TEST_FFMATH_UNARY(log)
__TEST_FFMATH_UNARY(log2)
__TEST_FFMATH_UNARY(log10)
__TEST_FFMATH_BINARY(powr)
__TEST_FFMATH_UNARY(rsqrt)
__TEST_FFMATH_UNARY(sin)
__TEST_FFMATH_UNARY(sqrt)
__TEST_FFMATH_UNARY(tan)

int main() {

assert(test_ffmath_cos());
assert(test_ffmath_exp());
assert(test_ffmath_exp2());
assert(test_ffmath_exp10());
assert(test_ffmath_log());
assert(test_ffmath_log2());
assert(test_ffmath_log10());
assert(test_ffmath_powr());
assert(test_ffmath_rsqrt());
assert(test_ffmath_sin());
assert(test_ffmath_sqrt());
assert(test_ffmath_tan());

return 0;
}