@@ -113,23 +113,24 @@ inline __SYCL_ALWAYS_INLINE
113
113
sycl::marray<T, N>>
114
114
tanh (sycl::marray<T, N> x) __NOEXC {
115
115
sycl::marray<T, N> res;
116
- #if defined(__SYCL_DEVICE_ONLY__) && defined(__NVPTX__)
117
- #define FUNC_VEC native::tanh
118
- #define FUNC FUNC_VEC
119
- #else
120
- #define FUNC_VEC __sycl_std::__invoke_tanh<sycl::vec<T, 2 >>
121
- #define FUNC __sycl_std::__invoke_tanh<T>
122
- #endif // defined(__SYCL_DEVICE_ONLY__) && defined(__NVPTX__)
123
116
124
117
for (size_t i = 0 ; i < N / 2 ; i++) {
125
- auto partial_res = FUNC_VEC (sycl::detail::to_vec2 (x, i * 2 ));
118
+ #if defined(__SYCL_DEVICE_ONLY__) && defined(__NVPTX__)
119
+ auto partial_res = native::tanh (sycl::detail::to_vec2 (x, i * 2 ));
120
+ #else
121
+ auto partial_res = __sycl_std::__invoke_tanh<sycl::vec<T, 2 >>(
122
+ sycl::detail::to_vec2 (x, i * 2 ));
123
+ #endif
126
124
std::memcpy (&res[i * 2 ], &partial_res, sizeof (vec<T, 2 >));
127
125
}
128
126
if (N % 2 ) {
129
- res[N - 1 ] = FUNC (x[N - 1 ]);
127
+ #if defined(__SYCL_DEVICE_ONLY__) && defined(__NVPTX__)
128
+ res[N - 1 ] = native::tanh (x[N - 1 ]);
129
+ #else
130
+ res[N - 1 ] = __sycl_std::__invoke_tanh<T>(x[N - 1 ]);
131
+ #endif
130
132
}
131
- #undef FUNC_VEC
132
- #undef FUNC
133
+
133
134
return res;
134
135
}
135
136
@@ -155,23 +156,23 @@ template <size_t N>
155
156
inline __SYCL_ALWAYS_INLINE sycl::marray<half, N>
156
157
exp2 (sycl::marray<half, N> x) __NOEXC {
157
158
sycl::marray<half, N> res;
158
- #if defined(__SYCL_DEVICE_ONLY__) && defined(__NVPTX__)
159
- #define FUNC_VEC native::exp2
160
- #define FUNC FUNC_VEC
161
- #else
162
- #define FUNC_VEC __sycl_std::__invoke_exp2<sycl::vec<half, 2 >>
163
- #define FUNC __sycl_std::__invoke_exp2<half>
164
- #endif // defined(__SYCL_DEVICE_ONLY__) && defined(__NVPTX__)
165
159
166
160
for (size_t i = 0 ; i < N / 2 ; i++) {
167
- auto partial_res = FUNC_VEC (sycl::detail::to_vec2 (x, i * 2 ));
161
+ #if defined(__SYCL_DEVICE_ONLY__) && defined(__NVPTX__)
162
+ auto partial_res = native::exp2 (sycl::detail::to_vec2 (x, i * 2 ));
163
+ #else
164
+ auto partial_res = __sycl_std::__invoke_exp2<sycl::vec<half, 2 >>(
165
+ sycl::detail::to_vec2 (x, i * 2 ));
166
+ #endif
168
167
std::memcpy (&res[i * 2 ], &partial_res, sizeof (vec<half, 2 >));
169
168
}
170
169
if (N % 2 ) {
171
- res[N - 1 ] = FUNC (x[N - 1 ]);
170
+ #if defined(__SYCL_DEVICE_ONLY__) && defined(__NVPTX__)
171
+ res[N - 1 ] = native::exp2 (x[N - 1 ]);
172
+ #else
173
+ res[N - 1 ] = __sycl_std::__invoke_exp2<half>(x[N - 1 ]);
174
+ #endif
172
175
}
173
- #undef FUNC_VEC
174
- #undef FUNC
175
176
return res;
176
177
}
177
178
0 commit comments