From 003eea7777e34a91e6d4d451356f6f8bc9c8d465 Mon Sep 17 00:00:00 2001 From: Denis Bakhvalov Date: Thu, 29 Apr 2021 16:56:28 -0700 Subject: [PATCH 1/3] [ESIMD] Implement ESIMD sin,cos,exp,log functions using scalar versions --- sycl/include/CL/sycl/builtins_esimd.hpp | 95 +++++++++++++++++++ .../experimental/esimd/detail/math_intrin.hpp | 56 ++++++++++- sycl/test/esimd/lane_id.cpp | 34 +++++++ sycl/test/esimd/math_impl.cpp | 41 ++++++++ 4 files changed, 224 insertions(+), 2 deletions(-) create mode 100644 sycl/include/CL/sycl/builtins_esimd.hpp create mode 100644 sycl/test/esimd/lane_id.cpp create mode 100644 sycl/test/esimd/math_impl.cpp diff --git a/sycl/include/CL/sycl/builtins_esimd.hpp b/sycl/include/CL/sycl/builtins_esimd.hpp new file mode 100644 index 0000000000000..3d4fb7c13022c --- /dev/null +++ b/sycl/include/CL/sycl/builtins_esimd.hpp @@ -0,0 +1,95 @@ +//==----------- builtins_esimd.hpp - SYCL ESIMD built-in functions ---------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#pragma once + +#include +#include +#include +#include +#include +#include + +// TODO Decide whether to mark functions with this attribute. +#define __NOEXC /*noexcept*/ + +__SYCL_INLINE_NAMESPACE(cl) { +namespace sycl { + +#define __ESIMD_NS ext::intel::experimental::esimd + +// cos +template +ESIMD_NODEBUG ESIMD_INLINE __ESIMD_NS::simd +cos(__ESIMD_NS::simd x) __NOEXC { +#ifdef __SYCL_DEVICE_ONLY__ + return __ESIMD_NS::detail::ocl_cos(x.data()); +#else + return __esimd_cos(x.data()); +#endif // __SYCL_DEVICE_ONLY__ +} + +ESIMD_NODEBUG ESIMD_INLINE __ESIMD_NS::simd +cos(__ESIMD_NS::simd x) __NOEXC { + return cos(x[0]); +} + +// sin +template +ESIMD_NODEBUG ESIMD_INLINE __ESIMD_NS::simd +sin(__ESIMD_NS::simd x) __NOEXC { +#ifdef __SYCL_DEVICE_ONLY__ + return __ESIMD_NS::detail::ocl_sin(x.data()); +#else + return __esimd_sin(x.data()); +#endif // __SYCL_DEVICE_ONLY__ +} + +ESIMD_NODEBUG ESIMD_INLINE __ESIMD_NS::simd +sin(__ESIMD_NS::simd x) __NOEXC { + return sin(x[0]); +} + +// exp +template +ESIMD_NODEBUG ESIMD_INLINE __ESIMD_NS::simd +exp(__ESIMD_NS::simd x) __NOEXC { +#ifdef __SYCL_DEVICE_ONLY__ + return __ESIMD_NS::detail::ocl_exp(x.data()); +#else + return __esimd_exp(x.data()); +#endif // __SYCL_DEVICE_ONLY__ +} + +ESIMD_NODEBUG ESIMD_INLINE __ESIMD_NS::simd +exp(__ESIMD_NS::simd x) __NOEXC { + return exp(x[0]); +} + +// log +template +ESIMD_NODEBUG ESIMD_INLINE __ESIMD_NS::simd +log(__ESIMD_NS::simd x) __NOEXC { +#ifdef __SYCL_DEVICE_ONLY__ + return __ESIMD_NS::detail::ocl_log(x.data()); +#else + return __esimd_log(x.data()); +#endif // __SYCL_DEVICE_ONLY__ +} + +ESIMD_NODEBUG ESIMD_INLINE __ESIMD_NS::simd +log(__ESIMD_NS::simd x) __NOEXC { + return log(x[0]); +} + +#undef __ESIMD_NS + +} // namespace sycl +} // __SYCL_INLINE_NAMESPACE(cl) + +#undef __NOEXC diff --git a/sycl/include/sycl/ext/intel/experimental/esimd/detail/math_intrin.hpp b/sycl/include/sycl/ext/intel/experimental/esimd/detail/math_intrin.hpp index 3113dc29b6173..619f9f598f198 100644 --- a/sycl/include/sycl/ext/intel/experimental/esimd/detail/math_intrin.hpp +++ b/sycl/include/sycl/ext/intel/experimental/esimd/detail/math_intrin.hpp @@ -11,6 +11,7 @@ #pragma once +#include #include #include #include @@ -316,7 +317,58 @@ SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SEIEED::vector_type_t __esimd_dp4(__SEIEED::vector_type_t v1, __SEIEED::vector_type_t v2); -#ifndef __SYCL_DEVICE_ONLY__ +#ifdef __SYCL_DEVICE_ONLY__ + +// lane-id for reusing scalar math functions. +// Depending upon the SIMT mode(8/16/32), the return value is +// in the range of 0-7, 0-15, or 0-31. +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION int __esimd_lane_id(); + +// Wrapper for designating a scalar region of code that will be +// vectorized by the backend compiler. +#define __ESIMD_SIMT_BEGIN(N, lane) \ + [&]() SYCL_ESIMD_FUNCTION ESIMD_NOINLINE \ + [[intel::sycl_esimd_vectorize(N)]] { \ + int lane = __esimd_lane_id(); +#define __ESIMD_SIMT_END \ + } \ + (); + +#define ESIMD_MATH_INTRINSIC_IMPL(type, func) \ + template \ + SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SEIEED::vector_type_t \ + ocl_##func(__SEIEED::vector_type_t src0) { \ + __SEIEED::vector_type_t retv; \ + __ESIMD_SIMT_BEGIN(SZ, lane) \ + retv[lane] = sycl::func(src0[lane]); \ + __ESIMD_SIMT_END \ + return retv; \ + } + +__SYCL_INLINE_NAMESPACE(cl) { +namespace sycl { +namespace ext { +namespace intel { +namespace experimental { +namespace esimd { +namespace detail { +ESIMD_MATH_INTRINSIC_IMPL(float, sin) +ESIMD_MATH_INTRINSIC_IMPL(float, cos) +ESIMD_MATH_INTRINSIC_IMPL(float, exp) +ESIMD_MATH_INTRINSIC_IMPL(float, log) +} // namespace detail +} // namespace esimd +} // namespace experimental +} // namespace intel +} // namespace ext +} // namespace sycl +} // __SYCL_INLINE_NAMESPACE(cl) + +#undef __ESIMD_SIMT_BEGIN +#undef __ESIMD_SIMT_END +#undef ESIMD_MATH_INTRINSIC_IMPL + +#else // __SYCL_DEVICE_ONLY__ template inline T extract(const uint32_t &width, const uint32_t &offset, uint32_t src, @@ -1277,6 +1329,6 @@ __esimd_reduced_smin(__SEIEED::vector_type_t src1, #undef __SEIEEED -#endif // #ifndef __SYCL_DEVICE_ONLY__ +#endif // #ifdef __SYCL_DEVICE_ONLY__ #undef __SEIEED diff --git a/sycl/test/esimd/lane_id.cpp b/sycl/test/esimd/lane_id.cpp new file mode 100644 index 0000000000000..aa9919ca8ff7f --- /dev/null +++ b/sycl/test/esimd/lane_id.cpp @@ -0,0 +1,34 @@ +// RUN: %clangxx -fsycl -c -fsycl-device-only -Xclang -emit-llvm %s -o - | FileCheck %s + +// This test checks the codegen for the basic usage of __ESIMD_SIMT_BEGIN - +// __ESIMD_SIMT_END construct. + +#include +#include + +using namespace cl::sycl; +using namespace sycl::ext::intel::experimental::esimd; + +// Wrapper for designating a scalar region of code that will be +// vectorized by the backend compiler. +#define SIMT_BEGIN(N, lane) \ + [&]() SYCL_ESIMD_FUNCTION ESIMD_NOINLINE \ + [[intel::sycl_esimd_vectorize(N)]] { \ + int lane = __esimd_lane_id(); +#define SIMT_END \ + } \ + (); + +// CHECK-LABEL: define dso_local spir_func void @_Z3fooi +//CHECK: call spir_func void @"_ZZ3fooiENK3$_0clEv"( +SYCL_ESIMD_FUNCTION SYCL_EXTERNAL simd foo(int x) { + simd v = 0; + SIMT_BEGIN(16, lane) + //CHECK: define internal spir_func void @"_ZZ3fooiENK3$_0clEv"({{.*}}) {{.*}} #[[ATTR:[0-9]+]] + //CHECK: %{{[0-9a-zA-Z_.]+}} = tail call spir_func i32 @_Z15__esimd_lane_idv() + v.select<1, 0>(lane) = x++; + SIMT_END + return v; +} + +//CHECK: attributes #[[ATTR]] = { {{.*}} "CMGenxSIMT"="16" {{.*}}} diff --git a/sycl/test/esimd/math_impl.cpp b/sycl/test/esimd/math_impl.cpp new file mode 100644 index 0000000000000..e4c1bad1f4b13 --- /dev/null +++ b/sycl/test/esimd/math_impl.cpp @@ -0,0 +1,41 @@ +// RUN: %clangxx -fsycl -c -fsycl-device-only -Xclang -emit-llvm %s -o - | FileCheck %s + +// This test checks the codegen for the following ESIMD APIs: +// sin, cos, exp, log. + +#include +#include +#include + +using namespace cl::sycl; +using namespace sycl::ext::intel::experimental::esimd; + +// Math sin,cos,log,exp functions are translated into scalar __spirv_ocl_ calls +SYCL_ESIMD_FUNCTION SYCL_EXTERNAL simd sycl_math(simd x) { + simd v = 0; + //CHECK: call spir_func float @_Z15__spirv_ocl_cosf + v = sycl::cos(x); + //CHECK: call spir_func float @_Z15__spirv_ocl_sinf + v = sycl::sin(v); + //CHECK: call spir_func float @_Z15__spirv_ocl_logf + v = sycl::log(v); + //CHECK: call spir_func float @_Z15__spirv_ocl_expf + v = sycl::exp(v); + return v; +} + +// Math sin,cos,log,exp functions from esimd namespace are translated +// into vector __esimd_ calls, which later translate into GenX intrinsics. +SYCL_ESIMD_FUNCTION SYCL_EXTERNAL simd +esimd_math(simd x) { + simd v = 0; + //CHECK: call spir_func <16 x float> @_Z11__esimd_cos + v = esimd_cos(x); + //CHECK: call spir_func <16 x float> @_Z11__esimd_sin + v = esimd_sin(v); + //CHECK: call spir_func <16 x float> @_Z11__esimd_log + v = esimd_log(v); + //CHECK: call spir_func <16 x float> @_Z11__esimd_exp + v = esimd_exp(v); + return v; +} From 1c224dddff8e341222143533569013da942a3feb Mon Sep 17 00:00:00 2001 From: Denis Bakhvalov Date: Fri, 11 Jun 2021 16:36:50 -0700 Subject: [PATCH 2/3] Fixed tests --- sycl/include/CL/sycl/builtins_esimd.hpp | 2 +- sycl/test/esimd/lane_id.cpp | 4 ++-- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/sycl/include/CL/sycl/builtins_esimd.hpp b/sycl/include/CL/sycl/builtins_esimd.hpp index 3d4fb7c13022c..cd6a31f1d3b1c 100644 --- a/sycl/include/CL/sycl/builtins_esimd.hpp +++ b/sycl/include/CL/sycl/builtins_esimd.hpp @@ -8,12 +8,12 @@ #pragma once -#include #include #include #include #include #include +#include // TODO Decide whether to mark functions with this attribute. #define __NOEXC /*noexcept*/ diff --git a/sycl/test/esimd/lane_id.cpp b/sycl/test/esimd/lane_id.cpp index aa9919ca8ff7f..df6031ac1810f 100644 --- a/sycl/test/esimd/lane_id.cpp +++ b/sycl/test/esimd/lane_id.cpp @@ -20,11 +20,11 @@ using namespace sycl::ext::intel::experimental::esimd; (); // CHECK-LABEL: define dso_local spir_func void @_Z3fooi -//CHECK: call spir_func void @"_ZZ3fooiENK3$_0clEv"( +//CHECK: call spir_func void @_ZZ3fooiENKUlvE_clEv( SYCL_ESIMD_FUNCTION SYCL_EXTERNAL simd foo(int x) { simd v = 0; SIMT_BEGIN(16, lane) - //CHECK: define internal spir_func void @"_ZZ3fooiENK3$_0clEv"({{.*}}) {{.*}} #[[ATTR:[0-9]+]] + //CHECK: define internal spir_func void @_ZZ3fooiENKUlvE_clEv({{.*}}) {{.*}} #[[ATTR:[0-9]+]] //CHECK: %{{[0-9a-zA-Z_.]+}} = tail call spir_func i32 @_Z15__esimd_lane_idv() v.select<1, 0>(lane) = x++; SIMT_END From 5df760245adb79efa8fe8eb751610c4da10c96cd Mon Sep 17 00:00:00 2001 From: Denis Bakhvalov Date: Fri, 11 Jun 2021 16:58:31 -0700 Subject: [PATCH 3/3] Fixed review comments --- sycl/include/CL/sycl/builtins_esimd.hpp | 24 ------------------------ 1 file changed, 24 deletions(-) diff --git a/sycl/include/CL/sycl/builtins_esimd.hpp b/sycl/include/CL/sycl/builtins_esimd.hpp index cd6a31f1d3b1c..5b64428c685b3 100644 --- a/sycl/include/CL/sycl/builtins_esimd.hpp +++ b/sycl/include/CL/sycl/builtins_esimd.hpp @@ -21,8 +21,6 @@ __SYCL_INLINE_NAMESPACE(cl) { namespace sycl { -#define __ESIMD_NS ext::intel::experimental::esimd - // cos template ESIMD_NODEBUG ESIMD_INLINE __ESIMD_NS::simd @@ -34,11 +32,6 @@ cos(__ESIMD_NS::simd x) __NOEXC { #endif // __SYCL_DEVICE_ONLY__ } -ESIMD_NODEBUG ESIMD_INLINE __ESIMD_NS::simd -cos(__ESIMD_NS::simd x) __NOEXC { - return cos(x[0]); -} - // sin template ESIMD_NODEBUG ESIMD_INLINE __ESIMD_NS::simd @@ -50,11 +43,6 @@ sin(__ESIMD_NS::simd x) __NOEXC { #endif // __SYCL_DEVICE_ONLY__ } -ESIMD_NODEBUG ESIMD_INLINE __ESIMD_NS::simd -sin(__ESIMD_NS::simd x) __NOEXC { - return sin(x[0]); -} - // exp template ESIMD_NODEBUG ESIMD_INLINE __ESIMD_NS::simd @@ -66,11 +54,6 @@ exp(__ESIMD_NS::simd x) __NOEXC { #endif // __SYCL_DEVICE_ONLY__ } -ESIMD_NODEBUG ESIMD_INLINE __ESIMD_NS::simd -exp(__ESIMD_NS::simd x) __NOEXC { - return exp(x[0]); -} - // log template ESIMD_NODEBUG ESIMD_INLINE __ESIMD_NS::simd @@ -82,13 +65,6 @@ log(__ESIMD_NS::simd x) __NOEXC { #endif // __SYCL_DEVICE_ONLY__ } -ESIMD_NODEBUG ESIMD_INLINE __ESIMD_NS::simd -log(__ESIMD_NS::simd x) __NOEXC { - return log(x[0]); -} - -#undef __ESIMD_NS - } // namespace sycl } // __SYCL_INLINE_NAMESPACE(cl)