Skip to content

Commit 003eea7

Browse files
DenisBakhvalovDenisBakhvalov
committed
[ESIMD] Implement ESIMD sin,cos,exp,log functions using scalar versions
1 parent 4836390 commit 003eea7

File tree

4 files changed

+224
-2
lines changed

4 files changed

+224
-2
lines changed

‎sycl/include/CL/sycl/builtins_esimd.hpp‎

Lines changed: 95 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,95 @@
1+
//==----------- builtins_esimd.hpp - SYCL ESIMD built-in functions ---------==//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
#pragma once
10+
11+
#include <CL/sycl/INTEL/esimd/detail/esimd_math_intrin.hpp>
12+
#include <CL/sycl/detail/boolean.hpp>
13+
#include <CL/sycl/detail/builtins.hpp>
14+
#include <CL/sycl/detail/common.hpp>
15+
#include <CL/sycl/detail/generic_type_traits.hpp>
16+
#include <CL/sycl/types.hpp>
17+
18+
// TODO Decide whether to mark functions with this attribute.
19+
#define __NOEXC /*noexcept*/
20+
21+
__SYCL_INLINE_NAMESPACE(cl) {
22+
namespace sycl {
23+
24+
#define __ESIMD_NS ext::intel::experimental::esimd
25+
26+
// cos
27+
template <int SZ>
28+
ESIMD_NODEBUG ESIMD_INLINE __ESIMD_NS::simd<float, SZ>
29+
cos(__ESIMD_NS::simd<float, SZ> x) __NOEXC {
30+
#ifdef __SYCL_DEVICE_ONLY__
31+
return __ESIMD_NS::detail::ocl_cos<SZ>(x.data());
32+
#else
33+
return __esimd_cos<SZ>(x.data());
34+
#endif // __SYCL_DEVICE_ONLY__
35+
}
36+
37+
ESIMD_NODEBUG ESIMD_INLINE __ESIMD_NS::simd<float, 1>
38+
cos(__ESIMD_NS::simd<float, 1> x) __NOEXC {
39+
return cos(x[0]);
40+
}
41+
42+
// sin
43+
template <int SZ>
44+
ESIMD_NODEBUG ESIMD_INLINE __ESIMD_NS::simd<float, SZ>
45+
sin(__ESIMD_NS::simd<float, SZ> x) __NOEXC {
46+
#ifdef __SYCL_DEVICE_ONLY__
47+
return __ESIMD_NS::detail::ocl_sin<SZ>(x.data());
48+
#else
49+
return __esimd_sin<SZ>(x.data());
50+
#endif // __SYCL_DEVICE_ONLY__
51+
}
52+
53+
ESIMD_NODEBUG ESIMD_INLINE __ESIMD_NS::simd<float, 1>
54+
sin(__ESIMD_NS::simd<float, 1> x) __NOEXC {
55+
return sin(x[0]);
56+
}
57+
58+
// exp
59+
template <int SZ>
60+
ESIMD_NODEBUG ESIMD_INLINE __ESIMD_NS::simd<float, SZ>
61+
exp(__ESIMD_NS::simd<float, SZ> x) __NOEXC {
62+
#ifdef __SYCL_DEVICE_ONLY__
63+
return __ESIMD_NS::detail::ocl_exp<SZ>(x.data());
64+
#else
65+
return __esimd_exp<SZ>(x.data());
66+
#endif // __SYCL_DEVICE_ONLY__
67+
}
68+
69+
ESIMD_NODEBUG ESIMD_INLINE __ESIMD_NS::simd<float, 1>
70+
exp(__ESIMD_NS::simd<float, 1> x) __NOEXC {
71+
return exp(x[0]);
72+
}
73+
74+
// log
75+
template <int SZ>
76+
ESIMD_NODEBUG ESIMD_INLINE __ESIMD_NS::simd<float, SZ>
77+
log(__ESIMD_NS::simd<float, SZ> x) __NOEXC {
78+
#ifdef __SYCL_DEVICE_ONLY__
79+
return __ESIMD_NS::detail::ocl_log<SZ>(x.data());
80+
#else
81+
return __esimd_log<SZ>(x.data());
82+
#endif // __SYCL_DEVICE_ONLY__
83+
}
84+
85+
ESIMD_NODEBUG ESIMD_INLINE __ESIMD_NS::simd<float, 1>
86+
log(__ESIMD_NS::simd<float, 1> x) __NOEXC {
87+
return log(x[0]);
88+
}
89+
90+
#undef __ESIMD_NS
91+
92+
} // namespace sycl
93+
} // __SYCL_INLINE_NAMESPACE(cl)
94+
95+
#undef __NOEXC

‎sycl/include/sycl/ext/intel/experimental/esimd/detail/math_intrin.hpp‎

Lines changed: 54 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -11,6 +11,7 @@
1111

1212
#pragma once
1313

14+
#include <CL/sycl/builtins.hpp>
1415
#include <sycl/ext/intel/experimental/esimd/common.hpp>
1516
#include <sycl/ext/intel/experimental/esimd/detail/host_util.hpp>
1617
#include <sycl/ext/intel/experimental/esimd/detail/types.hpp>
@@ -316,7 +317,58 @@ SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SEIEED::vector_type_t<Ty, N>
316317
__esimd_dp4(__SEIEED::vector_type_t<Ty, N> v1,
317318
__SEIEED::vector_type_t<Ty, N> v2);
318319

319-
#ifndef __SYCL_DEVICE_ONLY__
320+
#ifdef __SYCL_DEVICE_ONLY__
321+
322+
// lane-id for reusing scalar math functions.
323+
// Depending upon the SIMT mode(8/16/32), the return value is
324+
// in the range of 0-7, 0-15, or 0-31.
325+
SYCL_EXTERNAL SYCL_ESIMD_FUNCTION int __esimd_lane_id();
326+
327+
// Wrapper for designating a scalar region of code that will be
328+
// vectorized by the backend compiler.
329+
#define __ESIMD_SIMT_BEGIN(N, lane) \
330+
[&]() SYCL_ESIMD_FUNCTION ESIMD_NOINLINE \
331+
[[intel::sycl_esimd_vectorize(N)]] { \
332+
int lane = __esimd_lane_id();
333+
#define __ESIMD_SIMT_END \
334+
} \
335+
();
336+
337+
#define ESIMD_MATH_INTRINSIC_IMPL(type, func) \
338+
template <int SZ> \
339+
SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SEIEED::vector_type_t<type, SZ> \
340+
ocl_##func(__SEIEED::vector_type_t<type, SZ> src0) { \
341+
__SEIEED::vector_type_t<type, SZ> retv; \
342+
__ESIMD_SIMT_BEGIN(SZ, lane) \
343+
retv[lane] = sycl::func(src0[lane]); \
344+
__ESIMD_SIMT_END \
345+
return retv; \
346+
}
347+
348+
__SYCL_INLINE_NAMESPACE(cl) {
349+
namespace sycl {
350+
namespace ext {
351+
namespace intel {
352+
namespace experimental {
353+
namespace esimd {
354+
namespace detail {
355+
ESIMD_MATH_INTRINSIC_IMPL(float, sin)
356+
ESIMD_MATH_INTRINSIC_IMPL(float, cos)
357+
ESIMD_MATH_INTRINSIC_IMPL(float, exp)
358+
ESIMD_MATH_INTRINSIC_IMPL(float, log)
359+
} // namespace detail
360+
} // namespace esimd
361+
} // namespace experimental
362+
} // namespace intel
363+
} // namespace ext
364+
} // namespace sycl
365+
} // __SYCL_INLINE_NAMESPACE(cl)
366+
367+
#undef __ESIMD_SIMT_BEGIN
368+
#undef __ESIMD_SIMT_END
369+
#undef ESIMD_MATH_INTRINSIC_IMPL
370+
371+
#else // __SYCL_DEVICE_ONLY__
320372

321373
template <typename T>
322374
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<Ty, N> src1,
12771329

12781330
#undef __SEIEEED
12791331

1280-
#endif // #ifndef __SYCL_DEVICE_ONLY__
1332+
#endif // #ifdef __SYCL_DEVICE_ONLY__
12811333

12821334
#undef __SEIEED

‎sycl/test/esimd/lane_id.cpp‎

Lines changed: 34 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,34 @@
1+
// RUN: %clangxx -fsycl -c -fsycl-device-only -Xclang -emit-llvm %s -o - | FileCheck %s
2+
3+
// This test checks the codegen for the basic usage of __ESIMD_SIMT_BEGIN -
4+
// __ESIMD_SIMT_END construct.
5+
6+
#include <CL/sycl.hpp>
7+
#include <CL/sycl/INTEL/esimd.hpp>
8+
9+
using namespace cl::sycl;
10+
using namespace sycl::ext::intel::experimental::esimd;
11+
12+
// Wrapper for designating a scalar region of code that will be
13+
// vectorized by the backend compiler.
14+
#define SIMT_BEGIN(N, lane) \
15+
[&]() SYCL_ESIMD_FUNCTION ESIMD_NOINLINE \
16+
[[intel::sycl_esimd_vectorize(N)]] { \
17+
int lane = __esimd_lane_id();
18+
#define SIMT_END \
19+
} \
20+
();
21+
22+
// CHECK-LABEL: define dso_local spir_func void @_Z3fooi
23+
//CHECK: call spir_func void @"_ZZ3fooiENK3$_0clEv"(
24+
SYCL_ESIMD_FUNCTION SYCL_EXTERNAL simd<int, 16> foo(int x) {
25+
simd<int, 16> v = 0;
26+
SIMT_BEGIN(16, lane)
27+
//CHECK: define internal spir_func void @"_ZZ3fooiENK3$_0clEv"({{.*}}) {{.*}} #[[ATTR:[0-9]+]]
28+
//CHECK: %{{[0-9a-zA-Z_.]+}} = tail call spir_func i32 @_Z15__esimd_lane_idv()
29+
v.select<1, 0>(lane) = x++;
30+
SIMT_END
31+
return v;
32+
}
33+
34+
//CHECK: attributes #[[ATTR]] = { {{.*}} "CMGenxSIMT"="16" {{.*}}}

‎sycl/test/esimd/math_impl.cpp‎

Lines changed: 41 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,41 @@
1+
// RUN: %clangxx -fsycl -c -fsycl-device-only -Xclang -emit-llvm %s -o - | FileCheck %s
2+
3+
// This test checks the codegen for the following ESIMD APIs:
4+
// sin, cos, exp, log.
5+
6+
#include <CL/sycl.hpp>
7+
#include <CL/sycl/INTEL/esimd.hpp>
8+
#include <CL/sycl/builtins_esimd.hpp>
9+
10+
using namespace cl::sycl;
11+
using namespace sycl::ext::intel::experimental::esimd;
12+
13+
// Math sin,cos,log,exp functions are translated into scalar __spirv_ocl_ calls
14+
SYCL_ESIMD_FUNCTION SYCL_EXTERNAL simd<float, 16> sycl_math(simd<float, 16> x) {
15+
simd<float, 16> v = 0;
16+
//CHECK: call spir_func float @_Z15__spirv_ocl_cosf
17+
v = sycl::cos(x);
18+
//CHECK: call spir_func float @_Z15__spirv_ocl_sinf
19+
v = sycl::sin(v);
20+
//CHECK: call spir_func float @_Z15__spirv_ocl_logf
21+
v = sycl::log(v);
22+
//CHECK: call spir_func float @_Z15__spirv_ocl_expf
23+
v = sycl::exp(v);
24+
return v;
25+
}
26+
27+
// Math sin,cos,log,exp functions from esimd namespace are translated
28+
// into vector __esimd_ calls, which later translate into GenX intrinsics.
29+
SYCL_ESIMD_FUNCTION SYCL_EXTERNAL simd<float, 16>
30+
esimd_math(simd<float, 16> x) {
31+
simd<float, 16> v = 0;
32+
//CHECK: call spir_func <16 x float> @_Z11__esimd_cos
33+
v = esimd_cos(x);
34+
//CHECK: call spir_func <16 x float> @_Z11__esimd_sin
35+
v = esimd_sin(v);
36+
//CHECK: call spir_func <16 x float> @_Z11__esimd_log
37+
v = esimd_log(v);
38+
//CHECK: call spir_func <16 x float> @_Z11__esimd_exp
39+
v = esimd_exp(v);
40+
return v;
41+
}

0 commit comments

Comments
 (0)