Skip to content

[SYCL] Add marray support to rest math built-in functions #8912

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 13 commits into from
Apr 6, 2023
Merged

[SYCL] Add marray support to rest math built-in functions #8912

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
13 commits
Select commit Hold shift + click to select a range
e7f7929
[SYCL] Add marray support to rest math built-in functions
dm-vodopyanov Mar 31, 2023
100a872
Apply CR comment, fixed other bugs
dm-vodopyanov Apr 4, 2023
59fb71a
Debugging test failures
dm-vodopyanov Apr 4, 2023
8bac796
Debugging test failures
dm-vodopyanov Apr 4, 2023
f8c049d
Debugging test failures
dm-vodopyanov Apr 4, 2023
9fdbb6e
Merge branch 'intel:sycl' into marray_for_rest_math_funcs
dm-vodopyanov Apr 4, 2023
d745d4b
Debugging test failures
dm-vodopyanov Apr 4, 2023
ae65305
Debugging test failures
dm-vodopyanov Apr 4, 2023
a084a79
Merge branch 'intel:sycl' into marray_for_rest_math_funcs
dm-vodopyanov Apr 5, 2023
f58d684
Fix marray math builtins API test
dm-vodopyanov Apr 5, 2023
6acfc04
Update marray_math.cpp
dm-vodopyanov Apr 5, 2023
052e0cc
Fix marray_math.cpp
dm-vodopyanov Apr 6, 2023
2814d7e
Update marray_math.cpp
dm-vodopyanov Apr 6, 2023
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
87 changes: 87 additions & 0 deletions sycl/include/sycl/builtins.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -12,6 +12,7 @@
#include <sycl/detail/builtins.hpp>
#include <sycl/detail/common.hpp>
#include <sycl/detail/generic_type_traits.hpp>
#include <sycl/pointers.hpp>
#include <sycl/types.hpp>

// TODO Decide whether to mark functions with this attribute.
Expand Down Expand Up @@ -775,6 +776,92 @@ detail::enable_if_t<detail::is_svgenfloat<T>::value, T> trunc(T x) __NOEXC {
return __sycl_std::__invoke_trunc<T>(x);
}

// other marray math functions

// TODO: can be optimized in the way marray math functions above are optimized
// (usage of vec<T, 2>)
#define __SYCL_MARRAY_MATH_FUNCTION_W_GENPTR_ARG_OVERLOAD_IMPL(NAME, ARGPTR, \
...) \
marray<T, N> res; \
for (int j = 0; j < N; j++) { \
multi_ptr<detail::marray_element_t<detail::remove_pointer_t<T2>>, \
T2::address_space> \
ptr = nullptr; \
res[j] = NAME(__VA_ARGS__, ptr); \
(*ARGPTR)[j] = *ptr; \
} \
return res;

#define __SYCL_MARRAY_MATH_FUNCTION_BINOP_2ND_ARG_GENFLOATPTR_OVERLOAD( \
NAME, ARG1, ARG2, ...) \
template <typename T, typename T2, size_t N> \
std::enable_if_t<detail::is_svgenfloat<T>::value && \
detail::is_genfloatptr_marray<T2>::value, \
marray<T, N>> \
NAME(marray<T, N> ARG1, T2 ARG2) __NOEXC { \
__SYCL_MARRAY_MATH_FUNCTION_W_GENPTR_ARG_OVERLOAD_IMPL(NAME, ARG2, \
__VA_ARGS__) \
}

__SYCL_MARRAY_MATH_FUNCTION_BINOP_2ND_ARG_GENFLOATPTR_OVERLOAD(fract, x, iptr,
x[j])
__SYCL_MARRAY_MATH_FUNCTION_BINOP_2ND_ARG_GENFLOATPTR_OVERLOAD(modf, x, iptr,
x[j])
__SYCL_MARRAY_MATH_FUNCTION_BINOP_2ND_ARG_GENFLOATPTR_OVERLOAD(sincos, x,
cosval, x[j])

#undef __SYCL_MARRAY_MATH_FUNCTION_BINOP_2ND_GENFLOATPTR_OVERLOAD

#define __SYCL_MARRAY_MATH_FUNCTION_BINOP_2ND_ARG_GENINTPTR_OVERLOAD( \
NAME, ARG1, ARG2, ...) \
template <typename T, typename T2, size_t N> \
std::enable_if_t<detail::is_svgenfloat<T>::value && \
detail::is_genintptr_marray<T2>::value, \
marray<T, N>> \
NAME(marray<T, N> ARG1, T2 ARG2) __NOEXC { \
__SYCL_MARRAY_MATH_FUNCTION_W_GENPTR_ARG_OVERLOAD_IMPL(NAME, ARG2, \
__VA_ARGS__) \
}

__SYCL_MARRAY_MATH_FUNCTION_BINOP_2ND_ARG_GENINTPTR_OVERLOAD(frexp, x, exp,
x[j])
__SYCL_MARRAY_MATH_FUNCTION_BINOP_2ND_ARG_GENINTPTR_OVERLOAD(lgamma_r, x, signp,
x[j])

#undef __SYCL_MARRAY_MATH_FUNCTION_BINOP_2ND_GENINTPTR_OVERLOAD

#define __SYCL_MARRAY_MATH_FUNCTION_REMQUO_OVERLOAD(NAME, ...) \
template <typename T, typename T2, size_t N> \
std::enable_if_t<detail::is_svgenfloat<T>::value && \
detail::is_genintptr_marray<T2>::value, \
marray<T, N>> \
NAME(marray<T, N> x, marray<T, N> y, T2 quo) __NOEXC { \
__SYCL_MARRAY_MATH_FUNCTION_W_GENPTR_ARG_OVERLOAD_IMPL(NAME, quo, \
__VA_ARGS__) \
}

__SYCL_MARRAY_MATH_FUNCTION_REMQUO_OVERLOAD(remquo, x[j], y[j])

#undef __SYCL_MARRAY_MATH_FUNCTION_REMQUO_OVERLOAD

#undef __SYCL_MARRAY_MATH_FUNCTION_W_GENPTR_ARG_OVERLOAD_IMPL

#define __SYCL_MARRAY_MATH_FUNCTION_NAN_OVERLOAD(NAME) \
template <typename T, size_t N> \
std::enable_if_t<detail::is_nan_type<T>::value, \
marray<detail::nan_return_t<T>, N>> \
NAME(marray<T, N> nancode) __NOEXC { \
marray<T, N> res; \
for (int j = 0; j < N; j++) { \
res[j] = NAME(nancode[j]); \
} \
return res; \
}

__SYCL_MARRAY_MATH_FUNCTION_NAN_OVERLOAD(nan)

#undef __SYCL_MARRAY_MATH_FUNCTION_NAN_OVERLOAD

/* --------------- 4.13.5 Common functions. ---------------------------------*/
// svgenfloat clamp (svgenfloat x, svgenfloat minval, svgenfloat maxval)
template <typename T>
Expand Down
12 changes: 12 additions & 0 deletions sycl/include/sycl/detail/generic_type_traits.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -232,11 +232,23 @@ using is_genintptr = bool_constant<
is_pointer<T>::value && is_genint<remove_pointer_t<T>>::value &&
is_address_space_compliant<T, gvl::nonconst_address_space_list>::value>;

template <typename T>
using is_genintptr_marray = bool_constant<
is_pointer<T>::value &&
is_genint<marray_element_t<remove_pointer_t<T>>>::value &&
is_address_space_compliant<T, gvl::nonconst_address_space_list>::value>;

template <typename T>
using is_genfloatptr = bool_constant<
is_pointer<T>::value && is_genfloat<remove_pointer_t<T>>::value &&
is_address_space_compliant<T, gvl::nonconst_address_space_list>::value>;

template <typename T>
using is_genfloatptr_marray = bool_constant<
is_pointer<T>::value &&
is_genfloat<marray_element_t<remove_pointer_t<T>>>::value &&
is_address_space_compliant<T, gvl::nonconst_address_space_list>::value>;

template <typename T>
using is_genptr = bool_constant<
is_pointer<T>::value && is_gentype<remove_pointer_t<T>>::value &&
Expand Down
122 changes: 122 additions & 0 deletions sycl/test-e2e/DeviceLib/built-ins/marray_common.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,122 @@
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
// RUN: %CPU_RUN_PLACEHOLDER %t.out
// RUN: %GPU_RUN_PLACEHOLDER %t.out
// RUN: %ACC_RUN_PLACEHOLDER %t.out

#ifdef _WIN32
#define _USE_MATH_DEFINES // To use math constants
#include <cmath>
#endif

#include <sycl/sycl.hpp>

#define TEST(FUNC, MARRAY_ELEM_TYPE, DIM, EXPECTED, DELTA, ...) \
{ \
{ \
MARRAY_ELEM_TYPE result[DIM]; \
{ \
sycl::buffer<MARRAY_ELEM_TYPE> b(result, sycl::range{DIM}); \
deviceQueue.submit([&](sycl::handler &cgh) { \
sycl::accessor res_access{b, cgh}; \
cgh.single_task([=]() { \
sycl::marray<MARRAY_ELEM_TYPE, DIM> res = FUNC(__VA_ARGS__); \
for (int i = 0; i < DIM; i++) \
res_access[i] = res[i]; \
}); \
}); \
} \
for (int i = 0; i < DIM; i++) \
assert(abs(result[i] - EXPECTED[i]) <= DELTA); \
} \
}

#define EXPECTED(TYPE, ...) ((TYPE[]){__VA_ARGS__})

int main() {
sycl::queue deviceQueue;
sycl::device dev = deviceQueue.get_device();

sycl::marray<float, 2> ma1{1.0f, 2.0f};
sycl::marray<float, 2> ma2{1.0f, 2.0f};
sycl::marray<float, 2> ma3{3.0f, 2.0f};
sycl::marray<double, 2> ma4{1.0, 2.0};
sycl::marray<float, 3> ma5{M_PI, M_PI, M_PI};
sycl::marray<double, 3> ma6{M_PI, M_PI, M_PI};
sycl::marray<sycl::half, 3> ma7{M_PI, M_PI, M_PI};
sycl::marray<float, 2> ma8{0.3f, 0.6f};
sycl::marray<double, 2> ma9{5.0, 8.0};
sycl::marray<float, 3> ma10{180, 180, 180};
sycl::marray<double, 3> ma11{180, 180, 180};
sycl::marray<sycl::half, 3> ma12{180, 180, 180};
sycl::marray<sycl::half, 3> ma13{181, 179, 181};
sycl::marray<float, 2> ma14{+0.0f, -0.6f};
sycl::marray<double, 2> ma15{-0.0, 0.6f};

// sycl::clamp
TEST(sycl::clamp, float, 2, EXPECTED(float, 1.0f, 2.0f), 0, ma1, ma2, ma3);
TEST(sycl::clamp, float, 2, EXPECTED(float, 2.0f, 2.0f), 0, ma1, 3.0f, 2.0f);
if (dev.has(sycl::aspect::fp64))
TEST(sycl::clamp, double, 2, EXPECTED(double, 2.0, 2.0), 0, ma4, 3.0, 2.0);
// sycl::degrees
TEST(sycl::degrees, float, 3, EXPECTED(float, 180, 180, 180), 0, ma5);
if (dev.has(sycl::aspect::fp64))
TEST(sycl::degrees, double, 3, EXPECTED(double, 180, 180, 180), 0, ma6);
if (dev.has(sycl::aspect::fp16))
TEST(sycl::degrees, sycl::half, 3, EXPECTED(sycl::half, 180, 180, 180), 0.2,
ma7);
// sycl::max
TEST(sycl::max, float, 2, EXPECTED(float, 3.0f, 2.0f), 0, ma1, ma3);
TEST(sycl::max, float, 2, EXPECTED(float, 1.5f, 2.0f), 0, ma1, 1.5f);
if (dev.has(sycl::aspect::fp64))
TEST(sycl::max, double, 2, EXPECTED(double, 1.5, 2.0), 0, ma4, 1.5);
// sycl::min
TEST(sycl::min, float, 2, EXPECTED(float, 1.0f, 2.0f), 0, ma1, ma3);
TEST(sycl::min, float, 2, EXPECTED(float, 1.0f, 1.5f), 0, ma1, 1.5f);
if (dev.has(sycl::aspect::fp64))
TEST(sycl::min, double, 2, EXPECTED(double, 1.0, 1.5), 0, ma4, 1.5);
// sycl::mix
TEST(sycl::mix, float, 2, EXPECTED(float, 1.6f, 2.0f), 0, ma1, ma3, ma8);
TEST(sycl::mix, float, 2, EXPECTED(float, 1.4f, 2.0f), 0, ma1, ma3, 0.2);
if (dev.has(sycl::aspect::fp64))
TEST(sycl::mix, double, 2, EXPECTED(double, 3.0, 5.0), 0, ma4, ma9, 0.5);
// sycl::radians
TEST(sycl::radians, float, 3, EXPECTED(float, M_PI, M_PI, M_PI), 0, ma10);
if (dev.has(sycl::aspect::fp64))
TEST(sycl::radians, double, 3, EXPECTED(double, M_PI, M_PI, M_PI), 0, ma11);
if (dev.has(sycl::aspect::fp16))
TEST(sycl::radians, sycl::half, 3, EXPECTED(sycl::half, M_PI, M_PI, M_PI),
0.002, ma12);
// sycl::step
TEST(sycl::step, float, 2, EXPECTED(float, 1.0f, 1.0f), 0, ma1, ma3);
if (dev.has(sycl::aspect::fp64))
TEST(sycl::step, double, 2, EXPECTED(double, 1.0, 1.0), 0, ma4, ma9);
if (dev.has(sycl::aspect::fp16))
TEST(sycl::step, sycl::half, 3, EXPECTED(sycl::half, 1.0, 0.0, 1.0), 0,
ma12, ma13);
TEST(sycl::step, float, 2, EXPECTED(float, 1.0f, 0.0f), 0, 2.5f, ma3);
if (dev.has(sycl::aspect::fp64))
TEST(sycl::step, double, 2, EXPECTED(double, 0.0f, 1.0f), 0, 6.0f, ma9);
// sycl::smoothstep
TEST(sycl::smoothstep, float, 2, EXPECTED(float, 1.0f, 0.0f), 0, ma1, ma2,
ma3);
if (dev.has(sycl::aspect::fp64))
TEST(sycl::smoothstep, double, 2, EXPECTED(double, 0.00147576, 0.00446826),
0.00000001, ma4, ma11, ma9);
if (dev.has(sycl::aspect::fp16))
TEST(sycl::smoothstep, sycl::half, 3, EXPECTED(sycl::half, 1.0, 0.0, 1.0),
0, ma12, ma12, ma13);
TEST(sycl::smoothstep, float, 2, EXPECTED(float, 0.0553936f, 0.0f), 0.0000001,
2.5f, 6.0f, ma3);
if (dev.has(sycl::aspect::fp64))
TEST(sycl::smoothstep, double, 2, EXPECTED(double, 0.0f, 1.0f), 0, 6.0f,
8.0f, ma9);
// sign
TEST(sycl::sign, float, 2, EXPECTED(float, +0.0f, -1.0f), 0, ma14);
if (dev.has(sycl::aspect::fp64))
TEST(sycl::sign, double, 2, EXPECTED(double, -0.0, 1.0), 0, ma15);
if (dev.has(sycl::aspect::fp16))
TEST(sycl::sign, sycl::half, 3, EXPECTED(sycl::half, 1.0, 1.0, 1.0), 0,
ma12);

return 0;
}
104 changes: 104 additions & 0 deletions sycl/test-e2e/DeviceLib/built-ins/marray_math.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,104 @@
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
// RUN: %CPU_RUN_PLACEHOLDER %t.out
// RUN: %GPU_RUN_PLACEHOLDER %t.out
// RUN: %ACC_RUN_PLACEHOLDER %t.out

#include <sycl/sycl.hpp>

#define TEST(FUNC, MARRAY_ELEM_TYPE, DIM, EXPECTED, DELTA, ...) \
{ \
{ \
MARRAY_ELEM_TYPE result[DIM]; \
{ \
sycl::buffer<MARRAY_ELEM_TYPE> b(result, sycl::range{DIM}); \
deviceQueue.submit([&](sycl::handler &cgh) { \
sycl::accessor res_access{b, cgh}; \
cgh.single_task([=]() { \
sycl::marray<MARRAY_ELEM_TYPE, DIM> res = FUNC(__VA_ARGS__); \
for (int i = 0; i < DIM; i++) \
res_access[i] = res[i]; \
}); \
}); \
} \
for (int i = 0; i < DIM; i++) \
assert(abs(result[i] - EXPECTED[i]) <= DELTA); \
} \
}

#define TEST2(FUNC, MARRAY_ELEM_TYPE, PTR_TYPE, DIM, EXPECTED_1, EXPECTED_2, \
DELTA, ...) \
{ \
{ \
MARRAY_ELEM_TYPE result[DIM]; \
sycl::marray<PTR_TYPE, DIM> result_ptr; \
{ \
sycl::buffer<MARRAY_ELEM_TYPE> b(result, sycl::range{DIM}); \
sycl::buffer<sycl::marray<PTR_TYPE, DIM>, 1> b_ptr(&result_ptr, \
sycl::range<1>(1)); \
deviceQueue.submit([&](sycl::handler &cgh) { \
sycl::accessor res_access{b, cgh}; \
sycl::accessor res_ptr_access{b_ptr, cgh}; \
cgh.single_task([=]() { \
sycl::global_ptr<sycl::marray<PTR_TYPE, DIM>> ptr(res_ptr_access); \
sycl::marray<MARRAY_ELEM_TYPE, DIM> res = FUNC(__VA_ARGS__, ptr); \
for (int i = 0; i < DIM; i++) \
res_access[i] = res[i]; \
}); \
}); \
} \
for (int i = 0; i < DIM; i++) { \
assert(std::abs(result[i] - EXPECTED_1[i]) <= DELTA); \
assert(std::abs(result_ptr[i] - EXPECTED_2[i]) <= DELTA); \
} \
} \
}

#define EXPECTED(TYPE, ...) ((TYPE[]){__VA_ARGS__})

int main() {
sycl::queue deviceQueue;

sycl::marray<float, 2> ma1{1.0f, 2.0f};
sycl::marray<float, 2> ma2{3.0f, 2.0f};
sycl::marray<float, 3> ma3{180, 180, 180};
sycl::marray<int, 3> ma4{1, 1, 1};
sycl::marray<float, 3> ma5{180, -180, -180};
sycl::marray<float, 3> ma6{1.4f, 4.2f, 5.3f};
sycl::marray<unsigned int, 3> ma7{1, 2, 3};
sycl::marray<unsigned long int, 3> ma8{1, 2, 3};

TEST(sycl::fabs, float, 3, EXPECTED(float, 180, 180, 180), 0, ma5);
TEST(sycl::ilogb, int, 3, EXPECTED(int, 7, 7, 7), 0, ma3);
TEST(sycl::fmax, float, 2, EXPECTED(float, 3.0f, 2.0f), 0, ma1, ma2);
TEST(sycl::fmax, float, 2, EXPECTED(float, 5.0f, 5.0f), 0, ma1, 5.0f);
TEST(sycl::fmin, float, 2, EXPECTED(float, 1.0f, 2.0f), 0, ma1, ma2);
TEST(sycl::fmin, float, 2, EXPECTED(float, 1.0f, 2.0f), 0, ma1, 5.0f);
TEST(sycl::ldexp, float, 3, EXPECTED(float, 360, 360, 360), 0, ma3, ma4);
TEST(sycl::ldexp, float, 3, EXPECTED(float, 5760, 5760, 5760), 0, ma3, 5);
TEST(sycl::pown, float, 3, EXPECTED(float, 180, 180, 180), 0.1, ma3, ma4);
TEST(sycl::pown, float, 3, EXPECTED(float, 180, 180, 180), 0.1, ma3, 1);
TEST(sycl::rootn, float, 3, EXPECTED(float, 180, 180, 180), 0.1, ma3, ma4);
TEST(sycl::rootn, float, 3, EXPECTED(float, 2.82523, 2.82523, 2.82523),
0.00001, ma3, 5);
TEST2(sycl::fract, float, float, 3, EXPECTED(float, 0.4f, 0.2f, 0.3f),
EXPECTED(float, 1, 4, 5), 0.0001, ma6);
TEST2(sycl::modf, float, float, 3, EXPECTED(float, 0.4f, 0.2f, 0.3f),
EXPECTED(float, 1, 4, 5), 0.0001, ma6);
TEST2(sycl::sincos, float, float, 3,
EXPECTED(float, 0.98545f, -0.871576f, -0.832267f),
EXPECTED(float, 0.169967, -0.490261, 0.554375), 0.0001, ma6);
TEST2(sycl::frexp, float, int, 3, EXPECTED(float, 0.7f, 0.525f, 0.6625f),
EXPECTED(int, 1, 3, 3), 0.0001, ma6);
TEST2(sycl::lgamma_r, float, int, 3,
EXPECTED(float, -0.119613f, 2.04856f, 3.63964f), EXPECTED(int, 1, 1, 1),
0.0001, ma6);
TEST2(sycl::remquo, float, int, 3, EXPECTED(float, 1.4f, 4.2f, 5.3f),
EXPECTED(int, 0, 0, 0), 0.0001, ma6, ma3);
TEST(sycl::nan, float, 3, EXPECTED(float, 0, 0, 0), 0.1, ma7);
TEST(sycl::nan, double, 3, EXPECTED(double, 0, 0, 0), 0.1, ma8);

TEST(sycl::half_precision::exp10, float, 2, EXPECTED(float, 10, 100), 0.1,
ma1);

return 0;
}