From ac9d01bef0a26bc842f34f631275b3d00108b620 Mon Sep 17 00:00:00 2001 From: KornevNikita Date: Mon, 20 Mar 2023 05:55:05 -0700 Subject: [PATCH 1/3] [SYCL] Implement the rest of geometric built-ins --- sycl/include/sycl/builtins.hpp | 87 +++++++++++++++++++ .../sycl/detail/generic_type_lists.hpp | 12 +++ .../sycl/detail/generic_type_traits.hpp | 9 ++ 3 files changed, 108 insertions(+) diff --git a/sycl/include/sycl/builtins.hpp b/sycl/include/sycl/builtins.hpp index fe6eeaaec2e4e..6ca8c0995fe5e 100644 --- a/sycl/include/sycl/builtins.hpp +++ b/sycl/include/sycl/builtins.hpp @@ -25,6 +25,17 @@ namespace detail { template vec to_vec2(marray x, size_t start) { return {x[start], x[start + 1]}; } +template vec to_vec(marray x) { + vec vec; + std::memcpy(&vec, &x, sizeof(x)); + return vec; +} +template marray to_marray(vec x) { + marray marray; + for (int i = 0; i < N; i++) + marray[i] = x[i]; + return marray; +} } // namespace detail #ifdef __SYCL_DEVICE_ONLY__ @@ -1140,6 +1151,16 @@ detail::enable_if_t::value, T> cross(T p0, return __sycl_std::__invoke_cross(p0, p1); } +template +detail::enable_if_t::value, T> +cross(T p0, T p1) __NOEXC { + vec, T::size()> result_v; + result_v = __sycl_std::__invoke_cross< + vec, T::size()>>(detail::to_vec(p0), + detail::to_vec(p1)); + return detail::to_marray(result_v); +} + // float dot (float p0, float p1) // double dot (double p0, double p1) // half dot (half p0, half p1) @@ -1169,6 +1190,16 @@ detail::enable_if_t::value, half> dot(T p0, return __sycl_std::__invoke_Dot(p0, p1); } +// float dot (mfloatn p0, mfloatn p1) (n = 2, 3, 4) +// double dot (mdoublen p0, mdoublen p1) (n = 2, 3, 4) +template +detail::enable_if_t::value, + detail::marray_element_type> +dot(T p0, T p1) __NOEXC { + return __sycl_std::__invoke_Dot>( + detail::to_vec(p0), detail::to_vec(p1)); +} + // float distance (gengeofloat p0, gengeofloat p1) template ::value, T>> @@ -1190,6 +1221,16 @@ half distance(T p0, T p1) __NOEXC { return __sycl_std::__invoke_distance(p0, p1); } +// float distance (mfloatn p0, mfloatn p1) +// double distance (mfloatn p0, mfloatn p1) +template ::value, + detail::marray_element_type>> +detail::marray_element_type distance(T p0, T p1) __NOEXC { + return __sycl_std::__invoke_distance>( + detail::to_vec(p0), detail::to_vec(p1)); +} + // float length (gengeofloat p) template ::value, T>> @@ -1211,6 +1252,15 @@ half length(T p) __NOEXC { return __sycl_std::__invoke_length(p); } +// float length (mfloatn p) (n = 2, 3, 4) +// double length (mdoublen p) (n = 2, 3, 4) +template ::value, T>> +detail::marray_element_type length(T p) __NOEXC { + return __sycl_std::__invoke_length>( + detail::to_vec(p)); +} + // gengeofloat normalize (gengeofloat p) template detail::enable_if_t::value, T> @@ -1231,6 +1281,17 @@ detail::enable_if_t::value, T> normalize(T p) __NOEXC { return __sycl_std::__invoke_normalize(p); } +// mfloatn normalize (mfloatn p) (n = 2, 3, 4) +// mdoublen normalize (mdoublen p) (n = 2, 3, 4) +template +detail::enable_if_t::value, T> +normalize(T p) __NOEXC { + vec, T::size()> result_v; + result_v = __sycl_std::__invoke_normalize< + vec, T::size()>>(detail::to_vec(p)); + return detail::to_marray(result_v); +} + // float fast_distance (gengeofloat p0, gengeofloat p1) template ::value, T>> @@ -1245,6 +1306,14 @@ double fast_distance(T p0, T p1) __NOEXC { return __sycl_std::__invoke_fast_distance(p0, p1); } +// float fast_distance (mfloatn p0, mfloatn p1) (n = 2, 3, 4) +template ::value, T>> +detail::marray_element_type fast_distance(T p0, T p1) __NOEXC { + return __sycl_std::__invoke_fast_distance(detail::to_vec(p0), + detail::to_vec(p1)); +} + // float fast_length (gengeofloat p) template ::value, T>> @@ -1259,6 +1328,14 @@ double fast_length(T p) __NOEXC { return __sycl_std::__invoke_fast_length(p); } +// float fast_length (mfloatn p0, mfloatn p1) (n = 2, 3, 4) +template ::value, T>> +detail::marray_element_type fast_length(T p) __NOEXC { + return __sycl_std::__invoke_fast_length>( + detail::to_vec(p)); +} + // gengeofloat fast_normalize (gengeofloat p) template detail::enable_if_t::value, T> @@ -1273,6 +1350,16 @@ fast_normalize(T p) __NOEXC { return __sycl_std::__invoke_fast_normalize(p); } +// mfloatn fast_normalize (mfloatn p) (n = 2, 3, 4) +template +detail::enable_if_t::value, T> +fast_normalize(T p) __NOEXC { + vec, T::size()> result_v; + result_v = __sycl_std::__invoke_fast_normalize< + vec, T::size()>>(detail::to_vec(p)); + return detail::to_marray(result_v); +} + /* SYCL 1.2.1 ---- 4.13.7 Relational functions. -----------------------------*/ /* SYCL 2020 ---- 4.17.9 Relational functions. -----------------------------*/ diff --git a/sycl/include/sycl/detail/generic_type_lists.hpp b/sycl/include/sycl/detail/generic_type_lists.hpp index fa75c28aff2aa..d9fb70e5c86a9 100644 --- a/sycl/include/sycl/detail/generic_type_lists.hpp +++ b/sycl/include/sycl/detail/generic_type_lists.hpp @@ -108,6 +108,12 @@ using vector_geo_float_list = using vector_geo_double_list = type_list, vec, vec, vec>; +using marray_geo_float_list = + type_list, marray, marray>; + +using marray_geo_double_list = + type_list, marray, marray>; + using geo_half_list = type_list; using geo_float_list = type_list; @@ -121,6 +127,9 @@ using scalar_geo_list = type_list; +using marray_geo_list = + type_list; + using geo_list = type_list; // cross floating point types @@ -133,6 +142,9 @@ using cross_double_list = type_list, vec>; using cross_floating_list = type_list; +using cross_marray_list = type_list, marray, + marray, marray>; + using scalar_default_char_list = type_list; using vector_default_char_list = diff --git a/sycl/include/sycl/detail/generic_type_traits.hpp b/sycl/include/sycl/detail/generic_type_traits.hpp index 09e1bb4850139..0448af57fd05e 100644 --- a/sycl/include/sycl/detail/generic_type_traits.hpp +++ b/sycl/include/sycl/detail/generic_type_traits.hpp @@ -70,6 +70,12 @@ using is_gengeofloat = is_contained; template using is_gengeodouble = is_contained; +template +using is_gengeomarrayfloat = is_contained; + +template +using is_gengeomarray = is_contained; + template using is_gengeohalf = is_contained; template @@ -97,6 +103,9 @@ using is_gencrosshalf = is_contained; template using is_gencross = is_contained; +template +using is_gencrossmarray = is_contained; + template using is_charn = is_contained; From eb1fcd3645ac4d801ee20d63abe1cc8a35bd1d70 Mon Sep 17 00:00:00 2001 From: Dmitry Vodopyanov Date: Mon, 10 Apr 2023 18:27:55 -0400 Subject: [PATCH 2/3] Refactoring + adding tests from llvm-test-suite --- sycl/include/sycl/builtins.hpp | 134 ++++++++--------- .../Basic/built-ins/marray_geometric.cpp | 137 ++++++++++++++++++ 2 files changed, 198 insertions(+), 73 deletions(-) create mode 100644 sycl/test-e2e/Basic/built-ins/marray_geometric.cpp diff --git a/sycl/include/sycl/builtins.hpp b/sycl/include/sycl/builtins.hpp index 362f8a223c3b7..68ea4b3527222 100644 --- a/sycl/include/sycl/builtins.hpp +++ b/sycl/include/sycl/builtins.hpp @@ -27,7 +27,8 @@ template vec to_vec2(marray x, size_t start) { } template vec to_vec(marray x) { vec vec; - std::memcpy(&vec, &x, sizeof(x)); + for (int i = 0; i < N; i++) + vec[i] = x[i]; return vec; } template marray to_marray(vec x) { @@ -1593,16 +1594,6 @@ detail::enable_if_t::value, T> cross(T p0, return __sycl_std::__invoke_cross(p0, p1); } -template -detail::enable_if_t::value, T> -cross(T p0, T p1) __NOEXC { - vec, T::size()> result_v; - result_v = __sycl_std::__invoke_cross< - vec, T::size()>>(detail::to_vec(p0), - detail::to_vec(p1)); - return detail::to_marray(result_v); -} - // float dot (float p0, float p1) // double dot (double p0, double p1) // half dot (half p0, half p1) @@ -1632,16 +1623,6 @@ detail::enable_if_t::value, half> dot(T p0, return __sycl_std::__invoke_Dot(p0, p1); } -// float dot (mfloatn p0, mfloatn p1) (n = 2, 3, 4) -// double dot (mdoublen p0, mdoublen p1) (n = 2, 3, 4) -template -detail::enable_if_t::value, - detail::marray_element_type> -dot(T p0, T p1) __NOEXC { - return __sycl_std::__invoke_Dot>( - detail::to_vec(p0), detail::to_vec(p1)); -} - // float distance (gengeofloat p0, gengeofloat p1) template ::value, T>> @@ -1663,16 +1644,6 @@ half distance(T p0, T p1) __NOEXC { return __sycl_std::__invoke_distance(p0, p1); } -// float distance (mfloatn p0, mfloatn p1) -// double distance (mfloatn p0, mfloatn p1) -template ::value, - detail::marray_element_type>> -detail::marray_element_type distance(T p0, T p1) __NOEXC { - return __sycl_std::__invoke_distance>( - detail::to_vec(p0), detail::to_vec(p1)); -} - // float length (gengeofloat p) template ::value, T>> @@ -1694,15 +1665,6 @@ half length(T p) __NOEXC { return __sycl_std::__invoke_length(p); } -// float length (mfloatn p) (n = 2, 3, 4) -// double length (mdoublen p) (n = 2, 3, 4) -template ::value, T>> -detail::marray_element_type length(T p) __NOEXC { - return __sycl_std::__invoke_length>( - detail::to_vec(p)); -} - // gengeofloat normalize (gengeofloat p) template detail::enable_if_t::value, T> @@ -1723,17 +1685,6 @@ detail::enable_if_t::value, T> normalize(T p) __NOEXC { return __sycl_std::__invoke_normalize(p); } -// mfloatn normalize (mfloatn p) (n = 2, 3, 4) -// mdoublen normalize (mdoublen p) (n = 2, 3, 4) -template -detail::enable_if_t::value, T> -normalize(T p) __NOEXC { - vec, T::size()> result_v; - result_v = __sycl_std::__invoke_normalize< - vec, T::size()>>(detail::to_vec(p)); - return detail::to_marray(result_v); -} - // float fast_distance (gengeofloat p0, gengeofloat p1) template ::value, T>> @@ -1748,14 +1699,6 @@ double fast_distance(T p0, T p1) __NOEXC { return __sycl_std::__invoke_fast_distance(p0, p1); } -// float fast_distance (mfloatn p0, mfloatn p1) (n = 2, 3, 4) -template ::value, T>> -detail::marray_element_type fast_distance(T p0, T p1) __NOEXC { - return __sycl_std::__invoke_fast_distance(detail::to_vec(p0), - detail::to_vec(p1)); -} - // float fast_length (gengeofloat p) template ::value, T>> @@ -1770,14 +1713,6 @@ double fast_length(T p) __NOEXC { return __sycl_std::__invoke_fast_length(p); } -// float fast_length (mfloatn p0, mfloatn p1) (n = 2, 3, 4) -template ::value, T>> -detail::marray_element_type fast_length(T p) __NOEXC { - return __sycl_std::__invoke_fast_length>( - detail::to_vec(p)); -} - // gengeofloat fast_normalize (gengeofloat p) template detail::enable_if_t::value, T> @@ -1792,14 +1727,67 @@ fast_normalize(T p) __NOEXC { return __sycl_std::__invoke_fast_normalize(p); } -// mfloatn fast_normalize (mfloatn p) (n = 2, 3, 4) +// marray geometric functions + +#define __SYCL_MARRAY_GEOMETRIC_FUNCTION_OVERLOAD_IMPL(NAME, ...) \ + vec, T::size()> result_v; \ + result_v = NAME(__VA_ARGS__); \ + return detail::to_marray(result_v); + template -detail::enable_if_t::value, T> +std::enable_if_t::value, T> cross(T p0, + T p1) __NOEXC { + __SYCL_MARRAY_GEOMETRIC_FUNCTION_OVERLOAD_IMPL(cross, detail::to_vec(p0), + detail::to_vec(p1)) +} + +template +std::enable_if_t::value, T> normalize(T p) __NOEXC { + __SYCL_MARRAY_GEOMETRIC_FUNCTION_OVERLOAD_IMPL(normalize, detail::to_vec(p)) +} + +template +std::enable_if_t::value, T> fast_normalize(T p) __NOEXC { - vec, T::size()> result_v; - result_v = __sycl_std::__invoke_fast_normalize< - vec, T::size()>>(detail::to_vec(p)); - return detail::to_marray(result_v); + __SYCL_MARRAY_GEOMETRIC_FUNCTION_OVERLOAD_IMPL(normalize, detail::to_vec(p)) +} + +#undef __SYCL_MARRAY_GEOMETRIC_FUNCTION_OVERLOAD_IMPL + +#define __SYCL_MARRAY_GEOMETRIC_FUNCTION_IS_GENGEOMARRAY_BINOP_OVERLOAD(NAME) \ + template \ + std::enable_if_t::value, \ + detail::marray_element_t> \ + NAME(T p0, T p1) __NOEXC { \ + return NAME(detail::to_vec(p0), detail::to_vec(p1)); \ + } + +// clang-format off +__SYCL_MARRAY_GEOMETRIC_FUNCTION_IS_GENGEOMARRAY_BINOP_OVERLOAD(dot) +__SYCL_MARRAY_GEOMETRIC_FUNCTION_IS_GENGEOMARRAY_BINOP_OVERLOAD(distance) +// clang-format on + +#undef __SYCL_MARRAY_GEOMETRIC_FUNCTION_IS_GENGEOMARRAY_BINOP_OVERLOAD + +template +std::enable_if_t::value, detail::marray_element_t> +length(T p) __NOEXC { + return __sycl_std::__invoke_length>( + detail::to_vec(p)); +} + +template +std::enable_if_t::value, + detail::marray_element_t> +fast_distance(T p0, T p1) __NOEXC { + return fast_distance(detail::to_vec(p0), detail::to_vec(p1)); +} + +template +std::enable_if_t::value, + detail::marray_element_t> +fast_length(T p) __NOEXC { + return fast_length(detail::to_vec(p)); } /* SYCL 1.2.1 ---- 4.13.7 Relational functions. -----------------------------*/ diff --git a/sycl/test-e2e/Basic/built-ins/marray_geometric.cpp b/sycl/test-e2e/Basic/built-ins/marray_geometric.cpp new file mode 100644 index 0000000000000..133363d96ee83 --- /dev/null +++ b/sycl/test-e2e/Basic/built-ins/marray_geometric.cpp @@ -0,0 +1,137 @@ +// 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 + +#define TEST(FUNC, MARRAY_ELEM_TYPE, DIM, EXPECTED, DELTA, ...) \ + { \ + { \ + MARRAY_ELEM_TYPE result[DIM]; \ + { \ + sycl::buffer b(result, sycl::range{DIM}); \ + Queue.submit([&](sycl::handler &cgh) { \ + sycl::accessor res_access{b, cgh}; \ + cgh.single_task([=]() { \ + sycl::marray 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, TYPE, EXPECTED, DELTA, ...) \ + { \ + { \ + TYPE result; \ + { \ + sycl::buffer b(&result, 1); \ + Queue.submit([&](sycl::handler &cgh) { \ + sycl::accessor res_access{b, cgh}; \ + cgh.single_task([=]() { res_access[0] = FUNC(__VA_ARGS__); }); \ + }); \ + } \ + assert(abs(result - EXPECTED) <= DELTA); \ + } \ + } + +#define EXPECTED(TYPE, ...) ((TYPE[]){__VA_ARGS__}) + +int main() { + sycl::device Dev; + sycl::queue Queue(Dev); + // clang-format off + sycl::marray MFloatD2 = {1.f, 2.f}; + sycl::marray MFloatD2_2 = {3.f, 5.f}; + sycl::marray MFloatD3 = {1.f, 2.f, 3.f}; + sycl::marray MFloatD3_2 = {1.f, 5.f, 7.f}; + sycl::marray MFloatD4 = {1.f, 2.f, 3.f, 4.f}; + sycl::marray MFloatD4_2 = {1.f, 5.f, 7.f, 4.f}; + + sycl::marray MDoubleD2 = {1.0, 2.0}; + sycl::marray MDoubleD2_2 = {3.0, 5.0}; + sycl::marray MDoubleD3 = {1.0, 2.0, 3.0}; + sycl::marray MDoubleD3_2 = {1.0, 5.0, 7.0}; + sycl::marray MDoubleD4 = {1.0, 2.0, 3.0, 4.0}; + sycl::marray MDoubleD4_2 = {1.0, 5.0, 7.0, 4.0}; + // clang-format on + + TEST(sycl::cross, float, 3, EXPECTED(float, -1.f, -4.f, 3.f), 0, MFloatD3, + MFloatD3_2); + TEST(sycl::cross, float, 4, EXPECTED(float, -1.f, -4.f, 3.f, 0.f), 0, + MFloatD4, MFloatD4_2); + if (Dev.has(sycl::aspect::fp64)) { + TEST(sycl::cross, double, 3, EXPECTED(double, -1.f, -4.f, 3.f), 0, + MDoubleD3, MDoubleD3_2); + TEST(sycl::cross, double, 4, EXPECTED(double, -1.f, -4.f, 3.f, 0.f), 0, + MDoubleD4, MDoubleD4_2); + } + + TEST2(sycl::dot, float, 13.f, 0, MFloatD2, MFloatD2_2); + TEST2(sycl::dot, float, 32.f, 0, MFloatD3, MFloatD3_2); + TEST2(sycl::dot, float, 48.f, 0, MFloatD4, MFloatD4_2); + if (Dev.has(sycl::aspect::fp64)) { + TEST2(sycl::dot, double, 13, 0, MDoubleD2, MDoubleD2_2); + TEST2(sycl::dot, double, 32, 0, MDoubleD3, MDoubleD3_2); + TEST2(sycl::dot, double, 48, 0, MDoubleD4, MDoubleD4_2); + } + + TEST2(sycl::length, float, 2.236068f, 1e-6, MFloatD2); + TEST2(sycl::length, float, 3.741657f, 1e-6, MFloatD3); + TEST2(sycl::length, float, 5.477225f, 1e-6, MFloatD4); + if (Dev.has(sycl::aspect::fp64)) { + TEST2(sycl::length, double, 2.236068, 1e-6, MDoubleD2); + TEST2(sycl::length, double, 3.741657, 1e-6, MDoubleD3); + TEST2(sycl::length, double, 5.477225, 1e-6, MDoubleD4); + } + + TEST2(sycl::distance, float, 3.605551f, 1e-6, MFloatD2, MFloatD2_2); + TEST2(sycl::distance, float, 5.f, 0, MFloatD3, MFloatD3_2); + TEST2(sycl::distance, float, 5.f, 0, MFloatD4, MFloatD4_2); + if (Dev.has(sycl::aspect::fp64)) { + TEST2(sycl::distance, double, 3.605551, 1e-6, MDoubleD2, MDoubleD2_2); + TEST2(sycl::distance, double, 5.0, 0, MDoubleD3, MDoubleD3_2); + TEST2(sycl::distance, double, 5.0, 0, MDoubleD4, MDoubleD4_2); + } + + TEST(sycl::normalize, float, 2, EXPECTED(float, 0.447213f, 0.894427f), 1e-6, + MFloatD2); + TEST(sycl::normalize, float, 3, + EXPECTED(float, 0.267261f, 0.534522f, 0.801784f), 1e-6, MFloatD3); + TEST(sycl::normalize, float, 4, + EXPECTED(float, 0.182574f, 0.365148f, 0.547723f, 0.730297f), 1e-6, + MFloatD4); + if (Dev.has(sycl::aspect::fp64)) { + TEST(sycl::normalize, double, 2, EXPECTED(double, 0.447213, 0.894427), 1e-6, + MDoubleD2); + TEST(sycl::normalize, double, 3, + EXPECTED(double, 0.267261, 0.534522, 0.801784), 1e-6, MDoubleD3); + TEST(sycl::normalize, double, 4, + EXPECTED(double, 0.182574, 0.365148, 0.547723, 0.730297), 1e-6, + MDoubleD4); + } + + TEST2(sycl::fast_distance, float, 3.605551f, 1e-6, MFloatD2, MFloatD2_2); + TEST2(sycl::fast_distance, float, 5.f, 0, MFloatD3, MFloatD3_2); + TEST2(sycl::fast_distance, float, 5.f, 0, MFloatD4, MFloatD4_2); + + TEST2(sycl::fast_length, float, 2.236068f, 1e-6, MFloatD2); + TEST2(sycl::fast_length, float, 3.741657f, 1e-6, MFloatD3); + TEST2(sycl::fast_length, float, 5.477225f, 1e-6, MFloatD4); + + TEST(sycl::fast_normalize, float, 2, EXPECTED(float, 0.447213f, 0.894427f), + 1e-3, MFloatD2); + TEST(sycl::fast_normalize, float, 3, + EXPECTED(float, 0.267261f, 0.534522f, 0.801784f), 1e-3, MFloatD3); + TEST(sycl::fast_normalize, float, 4, + EXPECTED(float, 0.182574f, 0.365148f, 0.547723f, 0.730297f), 1e-3, + MFloatD4); + + return 0; +} From 07d4b2fe6e1b5ce41d01e3e568ac9e48e7f3db12 Mon Sep 17 00:00:00 2001 From: Dmitry Vodoypanov Date: Tue, 11 Apr 2023 09:26:23 -0400 Subject: [PATCH 3/3] Apply CR comments --- sycl/include/sycl/builtins.hpp | 7 ++++--- 1 file changed, 4 insertions(+), 3 deletions(-) diff --git a/sycl/include/sycl/builtins.hpp b/sycl/include/sycl/builtins.hpp index 68ea4b3527222..7f8f85e733a22 100644 --- a/sycl/include/sycl/builtins.hpp +++ b/sycl/include/sycl/builtins.hpp @@ -27,13 +27,13 @@ template vec to_vec2(marray x, size_t start) { } template vec to_vec(marray x) { vec vec; - for (int i = 0; i < N; i++) + for (size_t i = 0; i < N; i++) vec[i] = x[i]; return vec; } template marray to_marray(vec x) { marray marray; - for (int i = 0; i < N; i++) + for (size_t i = 0; i < N; i++) marray[i] = x[i]; return marray; } @@ -1749,7 +1749,8 @@ std::enable_if_t::value, T> normalize(T p) __NOEXC { template std::enable_if_t::value, T> fast_normalize(T p) __NOEXC { - __SYCL_MARRAY_GEOMETRIC_FUNCTION_OVERLOAD_IMPL(normalize, detail::to_vec(p)) + __SYCL_MARRAY_GEOMETRIC_FUNCTION_OVERLOAD_IMPL(fast_normalize, + detail::to_vec(p)) } #undef __SYCL_MARRAY_GEOMETRIC_FUNCTION_OVERLOAD_IMPL