Skip to content

Commit e8c32cb

Browse files
DenisBakhvalovDenisBakhvalov
authored
[ESIMD] Overloading sycl sin,cos,exp,log functions for ESIMD arguments (#3717)
* [ESIMD] Implement ESIMD sin,cos,exp,log functions using scalar versions
1 parent 65d1562 commit e8c32cb

File tree

4 files changed

+200
-2
lines changed

4 files changed

+200
-2
lines changed

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

Lines changed: 71 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,71 @@
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/detail/boolean.hpp>
12+
#include <CL/sycl/detail/builtins.hpp>
13+
#include <CL/sycl/detail/common.hpp>
14+
#include <CL/sycl/detail/generic_type_traits.hpp>
15+
#include <CL/sycl/types.hpp>
16+
#include <sycl/ext/intel/experimental/esimd/detail/math_intrin.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+
// cos
25+
template <int SZ>
26+
ESIMD_NODEBUG ESIMD_INLINE __ESIMD_NS::simd<float, SZ>
27+
cos(__ESIMD_NS::simd<float, SZ> x) __NOEXC {
28+
#ifdef __SYCL_DEVICE_ONLY__
29+
return __ESIMD_NS::detail::ocl_cos<SZ>(x.data());
30+
#else
31+
return __esimd_cos<SZ>(x.data());
32+
#endif // __SYCL_DEVICE_ONLY__
33+
}
34+
35+
// sin
36+
template <int SZ>
37+
ESIMD_NODEBUG ESIMD_INLINE __ESIMD_NS::simd<float, SZ>
38+
sin(__ESIMD_NS::simd<float, SZ> x) __NOEXC {
39+
#ifdef __SYCL_DEVICE_ONLY__
40+
return __ESIMD_NS::detail::ocl_sin<SZ>(x.data());
41+
#else
42+
return __esimd_sin<SZ>(x.data());
43+
#endif // __SYCL_DEVICE_ONLY__
44+
}
45+
46+
// exp
47+
template <int SZ>
48+
ESIMD_NODEBUG ESIMD_INLINE __ESIMD_NS::simd<float, SZ>
49+
exp(__ESIMD_NS::simd<float, SZ> x) __NOEXC {
50+
#ifdef __SYCL_DEVICE_ONLY__
51+
return __ESIMD_NS::detail::ocl_exp<SZ>(x.data());
52+
#else
53+
return __esimd_exp<SZ>(x.data());
54+
#endif // __SYCL_DEVICE_ONLY__
55+
}
56+
57+
// log
58+
template <int SZ>
59+
ESIMD_NODEBUG ESIMD_INLINE __ESIMD_NS::simd<float, SZ>
60+
log(__ESIMD_NS::simd<float, SZ> x) __NOEXC {
61+
#ifdef __SYCL_DEVICE_ONLY__
62+
return __ESIMD_NS::detail::ocl_log<SZ>(x.data());
63+
#else
64+
return __esimd_log<SZ>(x.data());
65+
#endif // __SYCL_DEVICE_ONLY__
66+
}
67+
68+
} // namespace sycl
69+
} // __SYCL_INLINE_NAMESPACE(cl)
70+
71+
#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 @_ZZ3fooiENKUlvE_clEv(
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 @_ZZ3fooiENKUlvE_clEv({{.*}}) {{.*}} #[[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)