DPC++ Runtime
Runtime libraries for oneAPI DPC++
builtins_integer.cpp
Go to the documentation of this file.
1 //==---------- builtins_integer.cpp - SYCL built-in integer 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 // This file defines the host versions of functions defined
10 // in SYCL SPEC section - 4.13.4 Integer functions.
11 
12 #include "builtins_helper.hpp"
13 #include <CL/sycl/detail/export.hpp>
14 
15 #include <algorithm>
16 #include <type_traits>
17 
18 namespace s = cl::sycl;
19 namespace d = s::detail;
20 
21 __SYCL_INLINE_NAMESPACE(cl) {
22 namespace __host_std {
23 namespace {
24 
25 template <typename T> inline T __abs_diff(T x, T y) {
26  static_assert(std::is_integral<T>::value,
27  "Only integral types are supported");
28  return (x > y) ? (x - y) : (y - x);
29 }
30 
31 template <typename T> inline T __u_add_sat(T x, T y) {
32  return (x < (d::max_v<T>() - y) ? x + y : d::max_v<T>());
33 }
34 
35 template <typename T> inline T __s_add_sat(T x, T y) {
36  if (x > 0 && y > 0)
37  return (x < (d::max_v<T>() - y) ? (x + y) : d::max_v<T>());
38  if (x < 0 && y < 0)
39  return (x > (d::min_v<T>() - y) ? (x + y) : d::min_v<T>());
40  return x + y;
41 }
42 
43 template <typename T> inline T __hadd(T x, T y) {
44  const T one = 1;
45  return (x >> one) + (y >> one) + ((y & x) & one);
46 }
47 
48 template <typename T> inline T __rhadd(T x, T y) {
49  const T one = 1;
50  return (x >> one) + (y >> one) + ((y | x) & one);
51 }
52 
53 template <typename T> inline T __clamp(T x, T minval, T maxval) {
54  return std::min(std::max(x, minval), maxval);
55 }
56 
57 template <typename T> inline constexpr T __clz_impl(T x, T m, T n = 0) {
58  return (x & m) ? n : __clz_impl(x, T(m >> 1), ++n);
59 }
60 
61 template <typename T> inline constexpr T __clz(T x) {
62  using UT = typename std::make_unsigned<T>::type;
63  return (x == T(0)) ? sizeof(T) * 8 : __clz_impl<UT>(x, d::msbMask<UT>(x));
64 }
65 
66 template <typename T> inline constexpr T __ctz_impl(T x, T m, T n = 0) {
67  return (x & m) ? n : __ctz_impl(x, T(m << 1), ++n);
68 }
69 
70 template <typename T> inline constexpr T __ctz(T x) {
71  using UT = typename std::make_unsigned<T>::type;
72  return (x == T(0)) ? sizeof(T) * 8 : __ctz_impl<UT>(x, 1);
73 }
74 
75 template <typename T> T __mul_hi(T a, T b) {
76  using UPT = typename d::make_larger<T>::type;
77  UPT a_s = a;
78  UPT b_s = b;
79  UPT mul = a_s * b_s;
80  return (mul >> (sizeof(T) * 8));
81 }
82 
83 // A helper function for mul_hi built-in for long
84 template <typename T> inline T __get_high_half(T a0b0, T a0b1, T a1b0, T a1b1) {
85  constexpr int halfsize = (sizeof(T) * 8) / 2;
86  // To get the upper 64 bits:
87  // 64 bits from a1b1, upper 32 bits from [a1b0 + (a0b1 + a0b0>>32 (carry bit
88  // in 33rd bit))] with carry bit on 64th bit - use of hadd. Add the a1b1 to
89  // the above 32 bit result.
90  return a1b1 + (__hadd(a1b0, (a0b1 + (a0b0 >> halfsize))) >> (halfsize - 1));
91 }
92 
93 // A helper function for mul_hi built-in for long
94 template <typename T>
95 inline void __get_half_products(T a, T b, T &a0b0, T &a0b1, T &a1b0, T &a1b1) {
96  constexpr s::cl_int halfsize = (sizeof(T) * 8) / 2;
97  T a1 = a >> halfsize;
98  T a0 = (a << halfsize) >> halfsize;
99  T b1 = b >> halfsize;
100  T b0 = (b << halfsize) >> halfsize;
101 
102  // a1b1 - for bits - [64-128)
103  // a1b0 a0b1 for bits - [32-96)
104  // a0b0 for bits - [0-64)
105  a1b1 = a1 * b1;
106  a0b1 = a0 * b1;
107  a1b0 = a1 * b0;
108  a0b0 = a0 * b0;
109 }
110 
111 // T is minimum of 64 bits- long or longlong
112 template <typename T> inline T __u_long_mul_hi(T a, T b) {
113  T a0b0, a0b1, a1b0, a1b1;
114  __get_half_products(a, b, a0b0, a0b1, a1b0, a1b1);
115  T result = __get_high_half(a0b0, a0b1, a1b0, a1b1);
116  return result;
117 }
118 
119 template <typename T> inline T __s_long_mul_hi(T a, T b) {
120  using UT = typename std::make_unsigned<T>::type;
121  UT absA = std::abs(a);
122  UT absB = std::abs(b);
123 
124  UT a0b0, a0b1, a1b0, a1b1;
125  __get_half_products(absA, absB, a0b0, a0b1, a1b0, a1b1);
126  T result = __get_high_half(a0b0, a0b1, a1b0, a1b1);
127 
128  bool isResultNegative = (a < 0) != (b < 0);
129  if (isResultNegative) {
130  result = ~result;
131 
132  // Find the low half to see if we need to carry
133  constexpr int halfsize = (sizeof(T) * 8) / 2;
134  UT low = a0b0 + ((a0b1 + a1b0) << halfsize);
135  if (low == 0)
136  ++result;
137  }
138 
139  return result;
140 }
141 
142 template <typename T> inline T __mad_hi(T a, T b, T c) {
143  return __mul_hi(a, b) + c;
144 }
145 
146 template <typename T> inline T __u_long_mad_hi(T a, T b, T c) {
147  return __u_long_mul_hi(a, b) + c;
148 }
149 
150 template <typename T> inline T __s_long_mad_hi(T a, T b, T c) {
151  return __s_long_mul_hi(a, b) + c;
152 }
153 
154 template <typename T> inline T __s_mad_sat(T a, T b, T c) {
155  using UPT = typename d::make_larger<T>::type;
156  UPT mul = UPT(a) * UPT(b);
157  UPT res = mul + UPT(c);
158  const UPT max = d::max_v<T>();
159  const UPT min = d::min_v<T>();
160  res = std::min(std::max(res, min), max);
161  return T(res);
162 }
163 
164 template <typename T> inline T __s_long_mad_sat(T a, T b, T c) {
165  bool neg_prod = (a < 0) ^ (b < 0);
166  T mulhi = __s_long_mul_hi(a, b);
167 
168  // check mul_hi. If it is any value != 0.
169  // if prod is +ve, any value in mulhi means we need to saturate.
170  // if prod is -ve, any value in mulhi besides -1 means we need to saturate.
171  if (!neg_prod && mulhi != 0)
172  return d::max_v<T>();
173  if (neg_prod && mulhi != -1)
174  return d::min_v<T>(); // essentially some other negative value.
175  return __s_add_sat(T(a * b), c);
176 }
177 
178 template <typename T> inline T __u_mad_sat(T a, T b, T c) {
179  using UPT = typename d::make_larger<T>::type;
180  UPT mul = UPT(a) * UPT(b);
181  const UPT min = d::min_v<T>();
182  const UPT max = d::max_v<T>();
183  mul = std::min(std::max(mul, min), max);
184  return __u_add_sat(T(mul), c);
185 }
186 
187 template <typename T> inline T __u_long_mad_sat(T a, T b, T c) {
188  T mulhi = __u_long_mul_hi(a, b);
189  // check mul_hi. If it is any value != 0.
190  if (mulhi != 0)
191  return d::max_v<T>();
192  return __u_add_sat(T(a * b), c);
193 }
194 
195 template <typename T> inline T __rotate(T x, T n) {
196  using UT = typename std::make_unsigned<T>::type;
197  // Shrink the shift width so that it's in the range [0, num_bits(T)). Cast
198  // everything to unsigned to avoid type conversion issues.
199  constexpr UT size = sizeof(x) * 8;
200  UT xu = UT(x);
201  UT nu = UT(n) & (size - 1);
202  return (xu << nu) | (xu >> (size - nu));
203 }
204 
205 template <typename T> inline T __u_sub_sat(T x, T y) {
206  return (y < (x - d::min_v<T>())) ? (x - y) : d::min_v<T>();
207 }
208 
209 template <typename T> inline T __s_sub_sat(T x, T y) {
210  using UT = typename std::make_unsigned<T>::type;
211  T result = UT(x) - UT(y);
212  // Saturate result if (+) - (-) = (-) or (-) - (+) = (+).
213  if (((x < 0) ^ (y < 0)) && ((x < 0) ^ (result < 0)))
214  result = result < 0 ? d::max_v<T>() : d::min_v<T>();
215  return result;
216 }
217 
218 template <typename T1, typename T2>
219 typename d::make_larger<T1>::type inline __upsample(T1 hi, T2 lo) {
220  using UT = typename d::make_larger<T1>::type;
221  return (UT(hi) << (sizeof(T1) * 8)) | lo;
222 }
223 
224 template <typename T> inline constexpr T __popcount_impl(T x, size_t n = 0) {
225  return (x == T(0)) ? n : __popcount_impl(x >> 1, ((x & T(1)) ? ++n : n));
226 }
227 
228 template <typename T> inline constexpr T __popcount(T x) {
229  using UT = typename d::make_unsigned<T>::type;
230  return __popcount_impl(UT(x));
231 }
232 
233 template <typename T> inline T __mad24(T x, T y, T z) { return (x * y) + z; }
234 
235 template <typename T> inline T __mul24(T x, T y) { return (x * y); }
236 
237 } // namespace
238 
239 // --------------- 4.13.4 Integer functions. Host implementations --------------
240 // u_abs
241 __SYCL_EXPORT s::cl_uchar u_abs(s::cl_uchar x) __NOEXC { return x; }
242 __SYCL_EXPORT s::cl_ushort u_abs(s::cl_ushort x) __NOEXC { return x; }
243 __SYCL_EXPORT s::cl_uint u_abs(s::cl_uint x) __NOEXC { return x; }
244 __SYCL_EXPORT s::cl_ulong u_abs(s::cl_ulong x) __NOEXC { return x; }
245 MAKE_1V(u_abs, s::cl_uchar, s::cl_uchar)
246 MAKE_1V(u_abs, s::cl_ushort, s::cl_ushort)
247 MAKE_1V(u_abs, s::cl_uint, s::cl_uint)
248 MAKE_1V(u_abs, s::cl_ulong, s::cl_ulong)
249 
250 // s_abs
251 __SYCL_EXPORT s::cl_uchar s_abs(s::cl_char x) __NOEXC { return std::abs(x); }
252 __SYCL_EXPORT s::cl_ushort s_abs(s::cl_short x) __NOEXC { return std::abs(x); }
253 __SYCL_EXPORT s::cl_uint s_abs(s::cl_int x) __NOEXC { return std::abs(x); }
254 __SYCL_EXPORT s::cl_ulong s_abs(s::cl_long x) __NOEXC { return std::abs(x); }
255 MAKE_1V(s_abs, s::cl_uchar, s::cl_char)
256 MAKE_1V(s_abs, s::cl_ushort, s::cl_short)
257 MAKE_1V(s_abs, s::cl_uint, s::cl_int)
258 MAKE_1V(s_abs, s::cl_ulong, s::cl_long)
259 
260 // u_abs_diff
261 __SYCL_EXPORT s::cl_uchar u_abs_diff(s::cl_uchar x, s::cl_uchar y) __NOEXC {
262  return __abs_diff(x, y);
263 }
264 __SYCL_EXPORT s::cl_ushort u_abs_diff(s::cl_ushort x, s::cl_ushort y) __NOEXC {
265  return __abs_diff(x, y);
266 }
267 __SYCL_EXPORT s::cl_uint u_abs_diff(s::cl_uint x, s::cl_uint y) __NOEXC {
268  return __abs_diff(x, y);
269 }
270 __SYCL_EXPORT s::cl_ulong u_abs_diff(s::cl_ulong x, s::cl_ulong y) __NOEXC {
271  return __abs_diff(x, y);
272 }
273 
274 MAKE_1V_2V(u_abs_diff, s::cl_uchar, s::cl_uchar, s::cl_uchar)
275 MAKE_1V_2V(u_abs_diff, s::cl_ushort, s::cl_ushort, s::cl_ushort)
276 MAKE_1V_2V(u_abs_diff, s::cl_uint, s::cl_uint, s::cl_uint)
277 MAKE_1V_2V(u_abs_diff, s::cl_ulong, s::cl_ulong, s::cl_ulong)
278 
279 // s_abs_diff
280 __SYCL_EXPORT s::cl_uchar s_abs_diff(s::cl_char x, s::cl_char y) __NOEXC {
281  return __abs_diff(x, y);
282 }
283 __SYCL_EXPORT s::cl_ushort s_abs_diff(s::cl_short x, s::cl_short y) __NOEXC {
284  return __abs_diff(x, y);
285 }
286 __SYCL_EXPORT s::cl_uint s_abs_diff(s::cl_int x, s::cl_int y) __NOEXC {
287  return __abs_diff(x, y);
288 }
289 __SYCL_EXPORT s::cl_ulong s_abs_diff(s::cl_long x, s::cl_long y) __NOEXC {
290  return __abs_diff(x, y);
291 }
292 MAKE_1V_2V(s_abs_diff, s::cl_uchar, s::cl_char, s::cl_char)
293 MAKE_1V_2V(s_abs_diff, s::cl_ushort, s::cl_short, s::cl_short)
294 MAKE_1V_2V(s_abs_diff, s::cl_uint, s::cl_int, s::cl_int)
295 MAKE_1V_2V(s_abs_diff, s::cl_ulong, s::cl_long, s::cl_long)
296 
297 // u_add_sat
298 __SYCL_EXPORT s::cl_uchar u_add_sat(s::cl_uchar x, s::cl_uchar y) __NOEXC {
299  return __u_add_sat(x, y);
300 }
301 __SYCL_EXPORT s::cl_ushort u_add_sat(s::cl_ushort x, s::cl_ushort y) __NOEXC {
302  return __u_add_sat(x, y);
303 }
304 __SYCL_EXPORT s::cl_uint u_add_sat(s::cl_uint x, s::cl_uint y) __NOEXC {
305  return __u_add_sat(x, y);
306 }
307 __SYCL_EXPORT s::cl_ulong u_add_sat(s::cl_ulong x, s::cl_ulong y) __NOEXC {
308  return __u_add_sat(x, y);
309 }
310 MAKE_1V_2V(u_add_sat, s::cl_uchar, s::cl_uchar, s::cl_uchar)
311 MAKE_1V_2V(u_add_sat, s::cl_ushort, s::cl_ushort, s::cl_ushort)
312 MAKE_1V_2V(u_add_sat, s::cl_uint, s::cl_uint, s::cl_uint)
313 MAKE_1V_2V(u_add_sat, s::cl_ulong, s::cl_ulong, s::cl_ulong)
314 
315 // s_add_sat
316 __SYCL_EXPORT s::cl_char s_add_sat(s::cl_char x, s::cl_char y) __NOEXC {
317  return __s_add_sat(x, y);
318 }
319 __SYCL_EXPORT s::cl_short s_add_sat(s::cl_short x, s::cl_short y) __NOEXC {
320  return __s_add_sat(x, y);
321 }
322 __SYCL_EXPORT s::cl_int s_add_sat(s::cl_int x, s::cl_int y) __NOEXC {
323  return __s_add_sat(x, y);
324 }
325 __SYCL_EXPORT s::cl_long s_add_sat(s::cl_long x, s::cl_long y) __NOEXC {
326  return __s_add_sat(x, y);
327 }
328 MAKE_1V_2V(s_add_sat, s::cl_char, s::cl_char, s::cl_char)
329 MAKE_1V_2V(s_add_sat, s::cl_short, s::cl_short, s::cl_short)
330 MAKE_1V_2V(s_add_sat, s::cl_int, s::cl_int, s::cl_int)
331 MAKE_1V_2V(s_add_sat, s::cl_long, s::cl_long, s::cl_long)
332 
333 // u_hadd
334 __SYCL_EXPORT s::cl_uchar u_hadd(s::cl_uchar x, s::cl_uchar y) __NOEXC {
335  return __hadd(x, y);
336 }
337 __SYCL_EXPORT s::cl_ushort u_hadd(s::cl_ushort x, s::cl_ushort y) __NOEXC {
338  return __hadd(x, y);
339 }
340 __SYCL_EXPORT s::cl_uint u_hadd(s::cl_uint x, s::cl_uint y) __NOEXC {
341  return __hadd(x, y);
342 }
343 __SYCL_EXPORT s::cl_ulong u_hadd(s::cl_ulong x, s::cl_ulong y) __NOEXC {
344  return __hadd(x, y);
345 }
346 MAKE_1V_2V(u_hadd, s::cl_uchar, s::cl_uchar, s::cl_uchar)
347 MAKE_1V_2V(u_hadd, s::cl_ushort, s::cl_ushort, s::cl_ushort)
348 MAKE_1V_2V(u_hadd, s::cl_uint, s::cl_uint, s::cl_uint)
349 MAKE_1V_2V(u_hadd, s::cl_ulong, s::cl_ulong, s::cl_ulong)
350 
351 // s_hadd
352 __SYCL_EXPORT s::cl_char s_hadd(s::cl_char x, s::cl_char y) __NOEXC {
353  return __hadd(x, y);
354 }
355 __SYCL_EXPORT s::cl_short s_hadd(s::cl_short x, s::cl_short y) __NOEXC {
356  return __hadd(x, y);
357 }
358 __SYCL_EXPORT s::cl_int s_hadd(s::cl_int x, s::cl_int y) __NOEXC {
359  return __hadd(x, y);
360 }
361 __SYCL_EXPORT s::cl_long s_hadd(s::cl_long x, s::cl_long y) __NOEXC {
362  return __hadd(x, y);
363 }
364 MAKE_1V_2V(s_hadd, s::cl_char, s::cl_char, s::cl_char)
365 MAKE_1V_2V(s_hadd, s::cl_short, s::cl_short, s::cl_short)
366 MAKE_1V_2V(s_hadd, s::cl_int, s::cl_int, s::cl_int)
367 MAKE_1V_2V(s_hadd, s::cl_long, s::cl_long, s::cl_long)
368 
369 // u_rhadd
370 __SYCL_EXPORT s::cl_uchar u_rhadd(s::cl_uchar x, s::cl_uchar y) __NOEXC {
371  return __rhadd(x, y);
372 }
373 __SYCL_EXPORT s::cl_ushort u_rhadd(s::cl_ushort x, s::cl_ushort y) __NOEXC {
374  return __rhadd(x, y);
375 }
376 __SYCL_EXPORT s::cl_uint u_rhadd(s::cl_uint x, s::cl_uint y) __NOEXC {
377  return __rhadd(x, y);
378 }
379 __SYCL_EXPORT s::cl_ulong u_rhadd(s::cl_ulong x, s::cl_ulong y) __NOEXC {
380  return __rhadd(x, y);
381 }
382 MAKE_1V_2V(u_rhadd, s::cl_uchar, s::cl_uchar, s::cl_uchar)
383 MAKE_1V_2V(u_rhadd, s::cl_ushort, s::cl_ushort, s::cl_ushort)
384 MAKE_1V_2V(u_rhadd, s::cl_uint, s::cl_uint, s::cl_uint)
385 MAKE_1V_2V(u_rhadd, s::cl_ulong, s::cl_ulong, s::cl_ulong)
386 
387 // s_rhadd
388 __SYCL_EXPORT s::cl_char s_rhadd(s::cl_char x, s::cl_char y) __NOEXC {
389  return __rhadd(x, y);
390 }
391 __SYCL_EXPORT s::cl_short s_rhadd(s::cl_short x, s::cl_short y) __NOEXC {
392  return __rhadd(x, y);
393 }
394 __SYCL_EXPORT s::cl_int s_rhadd(s::cl_int x, s::cl_int y) __NOEXC {
395  return __rhadd(x, y);
396 }
397 __SYCL_EXPORT s::cl_long s_rhadd(s::cl_long x, s::cl_long y) __NOEXC {
398  return __rhadd(x, y);
399 }
400 MAKE_1V_2V(s_rhadd, s::cl_char, s::cl_char, s::cl_char)
401 MAKE_1V_2V(s_rhadd, s::cl_short, s::cl_short, s::cl_short)
402 MAKE_1V_2V(s_rhadd, s::cl_int, s::cl_int, s::cl_int)
403 MAKE_1V_2V(s_rhadd, s::cl_long, s::cl_long, s::cl_long)
404 
405 // u_clamp
406 __SYCL_EXPORT s::cl_uchar u_clamp(s::cl_uchar x, s::cl_uchar minval,
407  s::cl_uchar maxval) __NOEXC {
408  return __clamp(x, minval, maxval);
409 }
410 __SYCL_EXPORT s::cl_ushort u_clamp(s::cl_ushort x, s::cl_ushort minval,
411  s::cl_ushort maxval) __NOEXC {
412  return __clamp(x, minval, maxval);
413 }
414 __SYCL_EXPORT s::cl_uint u_clamp(s::cl_uint x, s::cl_uint minval,
415  s::cl_uint maxval) __NOEXC {
416  return __clamp(x, minval, maxval);
417 }
418 __SYCL_EXPORT s::cl_ulong u_clamp(s::cl_ulong x, s::cl_ulong minval,
419  s::cl_ulong maxval) __NOEXC {
420  return __clamp(x, minval, maxval);
421 }
422 MAKE_1V_2V_3V(u_clamp, s::cl_uchar, s::cl_uchar, s::cl_uchar, s::cl_uchar)
423 MAKE_1V_2V_3V(u_clamp, s::cl_ushort, s::cl_ushort, s::cl_ushort, s::cl_ushort)
424 MAKE_1V_2V_3V(u_clamp, s::cl_uint, s::cl_uint, s::cl_uint, s::cl_uint)
425 MAKE_1V_2V_3V(u_clamp, s::cl_ulong, s::cl_ulong, s::cl_ulong, s::cl_ulong)
426 MAKE_1V_2S_3S(u_clamp, s::cl_uchar, s::cl_uchar, s::cl_uchar, s::cl_uchar)
427 MAKE_1V_2S_3S(u_clamp, s::cl_ushort, s::cl_ushort, s::cl_ushort, s::cl_ushort)
428 MAKE_1V_2S_3S(u_clamp, s::cl_uint, s::cl_uint, s::cl_uint, s::cl_uint)
429 MAKE_1V_2S_3S(u_clamp, s::cl_ulong, s::cl_ulong, s::cl_ulong, s::cl_ulong)
430 
431 // s_clamp
432 __SYCL_EXPORT s::cl_char s_clamp(s::cl_char x, s::cl_char minval,
433  s::cl_char maxval) __NOEXC {
434  return __clamp(x, minval, maxval);
435 }
436 __SYCL_EXPORT s::cl_short s_clamp(s::cl_short x, s::cl_short minval,
437  s::cl_short maxval) __NOEXC {
438  return __clamp(x, minval, maxval);
439 }
440 __SYCL_EXPORT s::cl_int s_clamp(s::cl_int x, s::cl_int minval,
441  s::cl_int maxval) __NOEXC {
442  return __clamp(x, minval, maxval);
443 }
444 __SYCL_EXPORT s::cl_long s_clamp(s::cl_long x, s::cl_long minval,
445  s::cl_long maxval) __NOEXC {
446  return __clamp(x, minval, maxval);
447 }
448 MAKE_1V_2V_3V(s_clamp, s::cl_char, s::cl_char, s::cl_char, s::cl_char)
449 MAKE_1V_2V_3V(s_clamp, s::cl_short, s::cl_short, s::cl_short, s::cl_short)
450 MAKE_1V_2V_3V(s_clamp, s::cl_int, s::cl_int, s::cl_int, s::cl_int)
451 MAKE_1V_2V_3V(s_clamp, s::cl_long, s::cl_long, s::cl_long, s::cl_long)
452 MAKE_1V_2S_3S(s_clamp, s::cl_char, s::cl_char, s::cl_char, s::cl_char)
453 MAKE_1V_2S_3S(s_clamp, s::cl_short, s::cl_short, s::cl_short, s::cl_short)
454 MAKE_1V_2S_3S(s_clamp, s::cl_int, s::cl_int, s::cl_int, s::cl_int)
455 MAKE_1V_2S_3S(s_clamp, s::cl_long, s::cl_long, s::cl_long, s::cl_long)
456 
457 // clz
458 __SYCL_EXPORT s::cl_uchar clz(s::cl_uchar x) __NOEXC { return __clz(x); }
459 __SYCL_EXPORT s::cl_char clz(s::cl_char x) __NOEXC { return __clz(x); }
460 __SYCL_EXPORT s::cl_ushort clz(s::cl_ushort x) __NOEXC { return __clz(x); }
461 __SYCL_EXPORT s::cl_short clz(s::cl_short x) __NOEXC { return __clz(x); }
462 __SYCL_EXPORT s::cl_uint clz(s::cl_uint x) __NOEXC { return __clz(x); }
463 __SYCL_EXPORT s::cl_int clz(s::cl_int x) __NOEXC { return __clz(x); }
464 __SYCL_EXPORT s::cl_ulong clz(s::cl_ulong x) __NOEXC { return __clz(x); }
465 __SYCL_EXPORT s::cl_long clz(s::cl_long x) __NOEXC { return __clz(x); }
466 MAKE_1V(clz, s::cl_uchar, s::cl_uchar)
467 MAKE_1V(clz, s::cl_char, s::cl_char)
468 MAKE_1V(clz, s::cl_ushort, s::cl_ushort)
469 MAKE_1V(clz, s::cl_short, s::cl_short)
470 MAKE_1V(clz, s::cl_uint, s::cl_uint)
471 MAKE_1V(clz, s::cl_int, s::cl_int)
472 MAKE_1V(clz, s::cl_ulong, s::cl_ulong)
473 MAKE_1V(clz, s::cl_long, s::cl_long)
474 
475 // ctz
476 __SYCL_EXPORT s::cl_uchar ctz(s::cl_uchar x) __NOEXC { return __ctz(x); }
477 __SYCL_EXPORT s::cl_char ctz(s::cl_char x) __NOEXC { return __ctz(x); }
478 __SYCL_EXPORT s::cl_ushort ctz(s::cl_ushort x) __NOEXC { return __ctz(x); }
479 __SYCL_EXPORT s::cl_short ctz(s::cl_short x) __NOEXC { return __ctz(x); }
480 __SYCL_EXPORT s::cl_uint ctz(s::cl_uint x) __NOEXC { return __ctz(x); }
481 __SYCL_EXPORT s::cl_int ctz(s::cl_int x) __NOEXC { return __ctz(x); }
482 __SYCL_EXPORT s::cl_ulong ctz(s::cl_ulong x) __NOEXC { return __ctz(x); }
483 __SYCL_EXPORT s::cl_long ctz(s::cl_long x) __NOEXC { return __ctz(x); }
484 MAKE_1V(ctz, s::cl_uchar, s::cl_uchar)
485 MAKE_1V(ctz, s::cl_char, s::cl_char)
486 MAKE_1V(ctz, s::cl_ushort, s::cl_ushort)
487 MAKE_1V(ctz, s::cl_short, s::cl_short)
488 MAKE_1V(ctz, s::cl_uint, s::cl_uint)
489 MAKE_1V(ctz, s::cl_int, s::cl_int)
490 MAKE_1V(ctz, s::cl_ulong, s::cl_ulong)
491 MAKE_1V(ctz, s::cl_long, s::cl_long)
492 
493 // s_mul_hi
494 __SYCL_EXPORT s::cl_char s_mul_hi(s::cl_char a, s::cl_char b) {
495  return __mul_hi(a, b);
496 }
497 __SYCL_EXPORT s::cl_short s_mul_hi(s::cl_short a, s::cl_short b) {
498  return __mul_hi(a, b);
499 }
500 __SYCL_EXPORT s::cl_int s_mul_hi(s::cl_int a, s::cl_int b) {
501  return __mul_hi(a, b);
502 }
503 __SYCL_EXPORT s::cl_long s_mul_hi(s::cl_long x, s::cl_long y) __NOEXC {
504  return __s_long_mul_hi(x, y);
505 }
506 MAKE_1V_2V(s_mul_hi, s::cl_char, s::cl_char, s::cl_char)
507 MAKE_1V_2V(s_mul_hi, s::cl_short, s::cl_short, s::cl_short)
508 MAKE_1V_2V(s_mul_hi, s::cl_int, s::cl_int, s::cl_int)
509 MAKE_1V_2V(s_mul_hi, s::cl_long, s::cl_long, s::cl_long)
510 
511 // u_mul_hi
512 __SYCL_EXPORT s::cl_uchar u_mul_hi(s::cl_uchar a, s::cl_uchar b) {
513  return __mul_hi(a, b);
514 }
515 __SYCL_EXPORT s::cl_ushort u_mul_hi(s::cl_ushort a, s::cl_ushort b) {
516  return __mul_hi(a, b);
517 }
518 __SYCL_EXPORT s::cl_uint u_mul_hi(s::cl_uint a, s::cl_uint b) {
519  return __mul_hi(a, b);
520 }
521 __SYCL_EXPORT s::cl_ulong u_mul_hi(s::cl_ulong x, s::cl_ulong y) __NOEXC {
522  return __u_long_mul_hi(x, y);
523 }
524 MAKE_1V_2V(u_mul_hi, s::cl_uchar, s::cl_uchar, s::cl_uchar)
525 MAKE_1V_2V(u_mul_hi, s::cl_ushort, s::cl_ushort, s::cl_ushort)
526 MAKE_1V_2V(u_mul_hi, s::cl_uint, s::cl_uint, s::cl_uint)
527 MAKE_1V_2V(u_mul_hi, s::cl_ulong, s::cl_ulong, s::cl_ulong)
528 
529 // s_mad_hi
530 __SYCL_EXPORT s::cl_char s_mad_hi(s::cl_char x, s::cl_char minval,
531  s::cl_char maxval) __NOEXC {
532  return __mad_hi(x, minval, maxval);
533 }
534 __SYCL_EXPORT s::cl_short s_mad_hi(s::cl_short x, s::cl_short minval,
535  s::cl_short maxval) __NOEXC {
536  return __mad_hi(x, minval, maxval);
537 }
538 __SYCL_EXPORT s::cl_int s_mad_hi(s::cl_int x, s::cl_int minval,
539  s::cl_int maxval) __NOEXC {
540  return __mad_hi(x, minval, maxval);
541 }
542 __SYCL_EXPORT s::cl_long s_mad_hi(s::cl_long x, s::cl_long minval,
543  s::cl_long maxval) __NOEXC {
544  return __s_long_mad_hi(x, minval, maxval);
545 }
546 MAKE_1V_2V_3V(s_mad_hi, s::cl_char, s::cl_char, s::cl_char, s::cl_char)
547 MAKE_1V_2V_3V(s_mad_hi, s::cl_short, s::cl_short, s::cl_short, s::cl_short)
548 MAKE_1V_2V_3V(s_mad_hi, s::cl_int, s::cl_int, s::cl_int, s::cl_int)
549 MAKE_1V_2V_3V(s_mad_hi, s::cl_long, s::cl_long, s::cl_long, s::cl_long)
550 
551 // u_mad_hi
552 __SYCL_EXPORT s::cl_uchar u_mad_hi(s::cl_uchar x, s::cl_uchar minval,
553  s::cl_uchar maxval) __NOEXC {
554  return __mad_hi(x, minval, maxval);
555 }
556 __SYCL_EXPORT s::cl_ushort u_mad_hi(s::cl_ushort x, s::cl_ushort minval,
557  s::cl_ushort maxval) __NOEXC {
558  return __mad_hi(x, minval, maxval);
559 }
560 __SYCL_EXPORT s::cl_uint u_mad_hi(s::cl_uint x, s::cl_uint minval,
561  s::cl_uint maxval) __NOEXC {
562  return __mad_hi(x, minval, maxval);
563 }
564 __SYCL_EXPORT s::cl_ulong u_mad_hi(s::cl_ulong x, s::cl_ulong minval,
565  s::cl_ulong maxval) __NOEXC {
566  return __u_long_mad_hi(x, minval, maxval);
567 }
568 MAKE_1V_2V_3V(u_mad_hi, s::cl_uchar, s::cl_uchar, s::cl_uchar, s::cl_uchar)
569 MAKE_1V_2V_3V(u_mad_hi, s::cl_ushort, s::cl_ushort, s::cl_ushort, s::cl_ushort)
570 MAKE_1V_2V_3V(u_mad_hi, s::cl_uint, s::cl_uint, s::cl_uint, s::cl_uint)
571 MAKE_1V_2V_3V(u_mad_hi, s::cl_ulong, s::cl_ulong, s::cl_ulong, s::cl_ulong)
572 
573 // s_mad_sat
574 __SYCL_EXPORT s::cl_char s_mad_sat(s::cl_char a, s::cl_char b,
575  s::cl_char c) __NOEXC {
576  return __s_mad_sat(a, b, c);
577 }
578 __SYCL_EXPORT s::cl_short s_mad_sat(s::cl_short a, s::cl_short b,
579  s::cl_short c) __NOEXC {
580  return __s_mad_sat(a, b, c);
581 }
582 __SYCL_EXPORT s::cl_int s_mad_sat(s::cl_int a, s::cl_int b,
583  s::cl_int c) __NOEXC {
584  return __s_mad_sat(a, b, c);
585 }
586 __SYCL_EXPORT s::cl_long s_mad_sat(s::cl_long a, s::cl_long b,
587  s::cl_long c) __NOEXC {
588  return __s_long_mad_sat(a, b, c);
589 }
590 MAKE_1V_2V_3V(s_mad_sat, s::cl_char, s::cl_char, s::cl_char, s::cl_char)
591 MAKE_1V_2V_3V(s_mad_sat, s::cl_short, s::cl_short, s::cl_short, s::cl_short)
592 MAKE_1V_2V_3V(s_mad_sat, s::cl_int, s::cl_int, s::cl_int, s::cl_int)
593 MAKE_1V_2V_3V(s_mad_sat, s::cl_long, s::cl_long, s::cl_long, s::cl_long)
594 
595 // u_mad_sat
596 __SYCL_EXPORT s::cl_uchar u_mad_sat(s::cl_uchar a, s::cl_uchar b,
597  s::cl_uchar c) __NOEXC {
598  return __u_mad_sat(a, b, c);
599 }
600 __SYCL_EXPORT s::cl_ushort u_mad_sat(s::cl_ushort a, s::cl_ushort b,
601  s::cl_ushort c) __NOEXC {
602  return __u_mad_sat(a, b, c);
603 }
604 __SYCL_EXPORT s::cl_uint u_mad_sat(s::cl_uint a, s::cl_uint b,
605  s::cl_uint c) __NOEXC {
606  return __u_mad_sat(a, b, c);
607 }
608 __SYCL_EXPORT s::cl_ulong u_mad_sat(s::cl_ulong a, s::cl_ulong b,
609  s::cl_ulong c) __NOEXC {
610  return __u_long_mad_sat(a, b, c);
611 }
612 MAKE_1V_2V_3V(u_mad_sat, s::cl_uchar, s::cl_uchar, s::cl_uchar, s::cl_uchar)
613 MAKE_1V_2V_3V(u_mad_sat, s::cl_ushort, s::cl_ushort, s::cl_ushort, s::cl_ushort)
614 MAKE_1V_2V_3V(u_mad_sat, s::cl_uint, s::cl_uint, s::cl_uint, s::cl_uint)
615 MAKE_1V_2V_3V(u_mad_sat, s::cl_ulong, s::cl_ulong, s::cl_ulong, s::cl_ulong)
616 
617 // s_max
618 __SYCL_EXPORT s::cl_char s_max(s::cl_char x, s::cl_char y) __NOEXC {
619  return std::max(x, y);
620 }
621 __SYCL_EXPORT s::cl_short s_max(s::cl_short x, s::cl_short y) __NOEXC {
622  return std::max(x, y);
623 }
624 __SYCL_EXPORT s::cl_int s_max(s::cl_int x, s::cl_int y) __NOEXC {
625  return std::max(x, y);
626 }
627 __SYCL_EXPORT s::cl_long s_max(s::cl_long x, s::cl_long y) __NOEXC {
628  return std::max(x, y);
629 }
630 MAKE_1V_2V(s_max, s::cl_char, s::cl_char, s::cl_char)
631 MAKE_1V_2V(s_max, s::cl_short, s::cl_short, s::cl_short)
632 MAKE_1V_2V(s_max, s::cl_int, s::cl_int, s::cl_int)
633 MAKE_1V_2V(s_max, s::cl_long, s::cl_long, s::cl_long)
634 MAKE_1V_2S(s_max, s::cl_char, s::cl_char, s::cl_char)
635 MAKE_1V_2S(s_max, s::cl_short, s::cl_short, s::cl_short)
636 MAKE_1V_2S(s_max, s::cl_int, s::cl_int, s::cl_int)
637 MAKE_1V_2S(s_max, s::cl_long, s::cl_long, s::cl_long)
638 
639 // u_max
640 __SYCL_EXPORT s::cl_uchar u_max(s::cl_uchar x, s::cl_uchar y) __NOEXC {
641  return std::max(x, y);
642 }
643 __SYCL_EXPORT s::cl_ushort u_max(s::cl_ushort x, s::cl_ushort y) __NOEXC {
644  return std::max(x, y);
645 }
646 __SYCL_EXPORT s::cl_uint u_max(s::cl_uint x, s::cl_uint y) __NOEXC {
647  return std::max(x, y);
648 }
649 __SYCL_EXPORT s::cl_ulong u_max(s::cl_ulong x, s::cl_ulong y) __NOEXC {
650  return std::max(x, y);
651 }
652 MAKE_1V_2V(u_max, s::cl_uchar, s::cl_uchar, s::cl_uchar)
653 MAKE_1V_2V(u_max, s::cl_ushort, s::cl_ushort, s::cl_ushort)
654 MAKE_1V_2V(u_max, s::cl_uint, s::cl_uint, s::cl_uint)
655 MAKE_1V_2V(u_max, s::cl_ulong, s::cl_ulong, s::cl_ulong)
656 MAKE_1V_2S(u_max, s::cl_uchar, s::cl_uchar, s::cl_uchar)
657 MAKE_1V_2S(u_max, s::cl_ushort, s::cl_ushort, s::cl_ushort)
658 MAKE_1V_2S(u_max, s::cl_uint, s::cl_uint, s::cl_uint)
659 MAKE_1V_2S(u_max, s::cl_ulong, s::cl_ulong, s::cl_ulong)
660 
661 // s_min
662 __SYCL_EXPORT s::cl_char s_min(s::cl_char x, s::cl_char y) __NOEXC {
663  return std::min(x, y);
664 }
665 __SYCL_EXPORT s::cl_short s_min(s::cl_short x, s::cl_short y) __NOEXC {
666  return std::min(x, y);
667 }
668 __SYCL_EXPORT s::cl_int s_min(s::cl_int x, s::cl_int y) __NOEXC {
669  return std::min(x, y);
670 }
671 __SYCL_EXPORT s::cl_long s_min(s::cl_long x, s::cl_long y) __NOEXC {
672  return std::min(x, y);
673 }
674 MAKE_1V_2V(s_min, s::cl_char, s::cl_char, s::cl_char)
675 MAKE_1V_2V(s_min, s::cl_short, s::cl_short, s::cl_short)
676 MAKE_1V_2V(s_min, s::cl_int, s::cl_int, s::cl_int)
677 MAKE_1V_2V(s_min, s::cl_long, s::cl_long, s::cl_long)
678 MAKE_1V_2S(s_min, s::cl_char, s::cl_char, s::cl_char)
679 MAKE_1V_2S(s_min, s::cl_short, s::cl_short, s::cl_short)
680 MAKE_1V_2S(s_min, s::cl_int, s::cl_int, s::cl_int)
681 MAKE_1V_2S(s_min, s::cl_long, s::cl_long, s::cl_long)
682 
683 // u_min
684 __SYCL_EXPORT s::cl_uchar u_min(s::cl_uchar x, s::cl_uchar y) __NOEXC {
685  return std::min(x, y);
686 }
687 __SYCL_EXPORT s::cl_ushort u_min(s::cl_ushort x, s::cl_ushort y) __NOEXC {
688  return std::min(x, y);
689 }
690 __SYCL_EXPORT s::cl_uint u_min(s::cl_uint x, s::cl_uint y) __NOEXC {
691  return std::min(x, y);
692 }
693 __SYCL_EXPORT s::cl_ulong u_min(s::cl_ulong x, s::cl_ulong y) __NOEXC {
694  return std::min(x, y);
695 }
696 MAKE_1V_2V(u_min, s::cl_uchar, s::cl_uchar, s::cl_uchar)
697 MAKE_1V_2V(u_min, s::cl_ushort, s::cl_ushort, s::cl_ushort)
698 MAKE_1V_2V(u_min, s::cl_uint, s::cl_uint, s::cl_uint)
699 MAKE_1V_2V(u_min, s::cl_ulong, s::cl_ulong, s::cl_ulong)
700 MAKE_1V_2S(u_min, s::cl_uchar, s::cl_uchar, s::cl_uchar)
701 MAKE_1V_2S(u_min, s::cl_ushort, s::cl_ushort, s::cl_ushort)
702 MAKE_1V_2S(u_min, s::cl_uint, s::cl_uint, s::cl_uint)
703 MAKE_1V_2S(u_min, s::cl_ulong, s::cl_ulong, s::cl_ulong)
704 
705 // rotate
706 __SYCL_EXPORT s::cl_uchar rotate(s::cl_uchar x, s::cl_uchar y) __NOEXC {
707  return __rotate(x, y);
708 }
709 __SYCL_EXPORT s::cl_ushort rotate(s::cl_ushort x, s::cl_ushort y) __NOEXC {
710  return __rotate(x, y);
711 }
712 __SYCL_EXPORT s::cl_uint rotate(s::cl_uint x, s::cl_uint y) __NOEXC {
713  return __rotate(x, y);
714 }
715 __SYCL_EXPORT s::cl_ulong rotate(s::cl_ulong x, s::cl_ulong y) __NOEXC {
716  return __rotate(x, y);
717 }
718 __SYCL_EXPORT s::cl_char rotate(s::cl_char x, s::cl_char y) __NOEXC {
719  return __rotate(x, y);
720 }
721 __SYCL_EXPORT s::cl_short rotate(s::cl_short x, s::cl_short y) __NOEXC {
722  return __rotate(x, y);
723 }
724 __SYCL_EXPORT s::cl_int rotate(s::cl_int x, s::cl_int y) __NOEXC {
725  return __rotate(x, y);
726 }
727 __SYCL_EXPORT s::cl_long rotate(s::cl_long x, s::cl_long y) __NOEXC {
728  return __rotate(x, y);
729 }
730 MAKE_1V_2V(rotate, s::cl_uchar, s::cl_uchar, s::cl_uchar)
731 MAKE_1V_2V(rotate, s::cl_ushort, s::cl_ushort, s::cl_ushort)
732 MAKE_1V_2V(rotate, s::cl_uint, s::cl_uint, s::cl_uint)
733 MAKE_1V_2V(rotate, s::cl_ulong, s::cl_ulong, s::cl_ulong)
734 MAKE_1V_2V(rotate, s::cl_char, s::cl_char, s::cl_char)
735 MAKE_1V_2V(rotate, s::cl_short, s::cl_short, s::cl_short)
736 MAKE_1V_2V(rotate, s::cl_int, s::cl_int, s::cl_int)
737 MAKE_1V_2V(rotate, s::cl_long, s::cl_long, s::cl_long)
738 
739 // u_sub_sat
740 __SYCL_EXPORT s::cl_uchar u_sub_sat(s::cl_uchar x, s::cl_uchar y) __NOEXC {
741  return __u_sub_sat(x, y);
742 }
743 __SYCL_EXPORT s::cl_ushort u_sub_sat(s::cl_ushort x, s::cl_ushort y) __NOEXC {
744  return __u_sub_sat(x, y);
745 }
746 __SYCL_EXPORT s::cl_uint u_sub_sat(s::cl_uint x, s::cl_uint y) __NOEXC {
747  return __u_sub_sat(x, y);
748 }
749 __SYCL_EXPORT s::cl_ulong u_sub_sat(s::cl_ulong x, s::cl_ulong y) __NOEXC {
750  return __u_sub_sat(x, y);
751 }
752 MAKE_1V_2V(u_sub_sat, s::cl_uchar, s::cl_uchar, s::cl_uchar)
753 MAKE_1V_2V(u_sub_sat, s::cl_ushort, s::cl_ushort, s::cl_ushort)
754 MAKE_1V_2V(u_sub_sat, s::cl_uint, s::cl_uint, s::cl_uint)
755 MAKE_1V_2V(u_sub_sat, s::cl_ulong, s::cl_ulong, s::cl_ulong)
756 
757 // s_sub_sat
758 __SYCL_EXPORT s::cl_char s_sub_sat(s::cl_char x, s::cl_char y) __NOEXC {
759  return __s_sub_sat(x, y);
760 }
761 __SYCL_EXPORT s::cl_short s_sub_sat(s::cl_short x, s::cl_short y) __NOEXC {
762  return __s_sub_sat(x, y);
763 }
764 __SYCL_EXPORT s::cl_int s_sub_sat(s::cl_int x, s::cl_int y) __NOEXC {
765  return __s_sub_sat(x, y);
766 }
767 __SYCL_EXPORT s::cl_long s_sub_sat(s::cl_long x, s::cl_long y) __NOEXC {
768  return __s_sub_sat(x, y);
769 }
770 MAKE_1V_2V(s_sub_sat, s::cl_char, s::cl_char, s::cl_char)
771 MAKE_1V_2V(s_sub_sat, s::cl_short, s::cl_short, s::cl_short)
772 MAKE_1V_2V(s_sub_sat, s::cl_int, s::cl_int, s::cl_int)
773 MAKE_1V_2V(s_sub_sat, s::cl_long, s::cl_long, s::cl_long)
774 
775 // u_upsample
776 __SYCL_EXPORT s::cl_ushort u_upsample(s::cl_uchar x, s::cl_uchar y) __NOEXC {
777  return __upsample(x, y);
778 }
779 __SYCL_EXPORT s::cl_uint u_upsample(s::cl_ushort x, s::cl_ushort y) __NOEXC {
780  return __upsample(x, y);
781 }
782 __SYCL_EXPORT s::cl_ulong u_upsample(s::cl_uint x, s::cl_uint y) __NOEXC {
783  return __upsample(x, y);
784 }
785 MAKE_1V_2V(u_upsample, s::cl_ushort, s::cl_uchar, s::cl_uchar)
786 MAKE_1V_2V(u_upsample, s::cl_uint, s::cl_ushort, s::cl_ushort)
787 MAKE_1V_2V(u_upsample, s::cl_ulong, s::cl_uint, s::cl_uint)
788 
789 __SYCL_EXPORT s::cl_short s_upsample(s::cl_char x, s::cl_uchar y) __NOEXC {
790  return __upsample(x, y);
791 }
792 __SYCL_EXPORT s::cl_int s_upsample(s::cl_short x, s::cl_ushort y) __NOEXC {
793  return __upsample(x, y);
794 }
795 __SYCL_EXPORT s::cl_long s_upsample(s::cl_int x, s::cl_uint y) __NOEXC {
796  return __upsample(x, y);
797 }
798 MAKE_1V_2V(s_upsample, s::cl_short, s::cl_char, s::cl_uchar)
799 MAKE_1V_2V(s_upsample, s::cl_int, s::cl_short, s::cl_ushort)
800 MAKE_1V_2V(s_upsample, s::cl_long, s::cl_int, s::cl_uint)
801 
802 // popcount
803 __SYCL_EXPORT s::cl_uchar popcount(s::cl_uchar x) __NOEXC {
804  return __popcount(x);
805 }
806 __SYCL_EXPORT s::cl_ushort popcount(s::cl_ushort x) __NOEXC {
807  return __popcount(x);
808 }
809 __SYCL_EXPORT s::cl_uint popcount(s::cl_uint x) __NOEXC {
810  return __popcount(x);
811 }
812 __SYCL_EXPORT s::cl_ulong popcount(s::cl_ulong x) __NOEXC {
813  return __popcount(x);
814 }
815 MAKE_1V(popcount, s::cl_uchar, s::cl_uchar)
816 MAKE_1V(popcount, s::cl_ushort, s::cl_ushort)
817 MAKE_1V(popcount, s::cl_uint, s::cl_uint)
818 MAKE_1V(popcount, s::cl_ulong, s::cl_ulong)
819 
820 __SYCL_EXPORT s::cl_char popcount(s::cl_char x) __NOEXC {
821  return __popcount(x);
822 }
823 __SYCL_EXPORT s::cl_short popcount(s::cl_short x) __NOEXC {
824  return __popcount(x);
825 }
826 __SYCL_EXPORT s::cl_int popcount(s::cl_int x) __NOEXC { return __popcount(x); }
827 __SYCL_EXPORT s::cl_long popcount(s::cl_long x) __NOEXC {
828  return __popcount(x);
829 }
830 MAKE_1V(popcount, s::cl_char, s::cl_char)
831 MAKE_1V(popcount, s::cl_short, s::cl_short)
832 MAKE_1V(popcount, s::cl_int, s::cl_int)
833 MAKE_1V(popcount, s::cl_long, s::cl_long)
834 
835 // u_mad24
836 __SYCL_EXPORT s::cl_uint u_mad24(s::cl_uint x, s::cl_uint y,
837  s::cl_uint z) __NOEXC {
838  return __mad24(x, y, z);
839 }
840 MAKE_1V_2V_3V(u_mad24, s::cl_uint, s::cl_uint, s::cl_uint, s::cl_uint)
841 
842 // s_mad24
843 __SYCL_EXPORT s::cl_int s_mad24(s::cl_int x, s::cl_int y, s::cl_int z) __NOEXC {
844  return __mad24(x, y, z);
845 }
846 MAKE_1V_2V_3V(s_mad24, s::cl_int, s::cl_int, s::cl_int, s::cl_int)
847 
848 // u_mul24
849 __SYCL_EXPORT s::cl_uint u_mul24(s::cl_uint x, s::cl_uint y) __NOEXC {
850  return __mul24(x, y);
851 }
852 MAKE_1V_2V(u_mul24, s::cl_uint, s::cl_uint, s::cl_uint)
853 
854 // s_mul24
855 __SYCL_EXPORT s::cl_int s_mul24(s::cl_int x, s::cl_int y) __NOEXC {
856  return __mul24(x, y);
857 }
858 MAKE_1V_2V(s_mul24, s::cl_int, s::cl_int, s::cl_int)
859 
860 } // namespace __host_std
861 } // __SYCL_INLINE_NAMESPACE(cl)
MAKE_1V_2S
#define MAKE_1V_2S(Fun, Ret, Arg1, Arg2)
Definition: builtins_helper.hpp:164
cl::__host_std::u_min
s::cl_ulong u_min(s::cl_ulong x, s::cl_ulong y) __NOEXC
Definition: builtins_integer.cpp:693
cl::sycl::cl_uchar
std::uint8_t cl_uchar
Definition: aliases.hpp:80
MAKE_1V_2V_3V
#define MAKE_1V_2V_3V(Fun, Ret, Arg1, Arg2, Arg3)
Definition: builtins_helper.hpp:140
cl::sycl::cl_long
std::int64_t cl_long
Definition: aliases.hpp:85
T
cl::__host_std::s_upsample
s::cl_long s_upsample(s::cl_int x, s::cl_uint y) __NOEXC
Definition: builtins_integer.cpp:795
cl::__host_std::ctz
s::cl_long ctz(s::cl_long x) __NOEXC
Definition: builtins_integer.cpp:483
cl::__host_std::u_mad24
s::cl_uint u_mad24(s::cl_uint x, s::cl_uint y, s::cl_uint z) __NOEXC
Definition: builtins_integer.cpp:836
cl::sycl::detail::min_v
static constexpr T min_v()
Definition: generic_type_traits.hpp:599
cl::sycl::ext::oneapi::experimental::matrix::matrix_use::b
@ b
cl::__host_std::clz
s::cl_long clz(s::cl_long x) __NOEXC
Definition: builtins_integer.cpp:465
cl::__host_std::s_hadd
s::cl_long s_hadd(s::cl_long x, s::cl_long y) __NOEXC
Definition: builtins_integer.cpp:361
cl::sycl
Definition: access.hpp:14
cl::__host_std::u_rhadd
s::cl_ulong u_rhadd(s::cl_ulong x, s::cl_ulong y) __NOEXC
Definition: builtins_integer.cpp:379
cl::__host_std::u_mul24
s::cl_uint u_mul24(s::cl_uint x, s::cl_uint y) __NOEXC
Definition: builtins_integer.cpp:849
max
simd< _Tp, _Abi > max(const simd< _Tp, _Abi > &, const simd< _Tp, _Abi > &) noexcept
MAKE_1V_2V
#define MAKE_1V_2V(Fun, Ret, Arg1, Arg2)
Definition: builtins_helper.hpp:129
cl::__host_std::s_max
s::cl_long s_max(s::cl_long x, s::cl_long y) __NOEXC
Definition: builtins_integer.cpp:627
cl::__host_std::u_hadd
s::cl_ulong u_hadd(s::cl_ulong x, s::cl_ulong y) __NOEXC
Definition: builtins_integer.cpp:343
cl::__host_std::s_mad_sat
s::cl_long s_mad_sat(s::cl_long a, s::cl_long b, s::cl_long c) __NOEXC
Definition: builtins_integer.cpp:586
cl::__host_std::s_add_sat
s::cl_long s_add_sat(s::cl_long x, s::cl_long y) __NOEXC
Definition: builtins_integer.cpp:325
cl::__host_std::s_sub_sat
s::cl_long s_sub_sat(s::cl_long x, s::cl_long y) __NOEXC
Definition: builtins_integer.cpp:767
export.hpp
cl::__host_std::s_mul24
s::cl_int s_mul24(s::cl_int x, s::cl_int y) __NOEXC
Definition: builtins_integer.cpp:855
MAKE_1V_2S_3S
#define MAKE_1V_2S_3S(Fun, Ret, Arg1, Arg2, Arg3)
Definition: builtins_helper.hpp:172
cl::sycl::cl_ushort
std::uint16_t cl_ushort
Definition: aliases.hpp:82
cl::sycl::abs
detail::enable_if_t< detail::is_genfloat< T >::value, T > abs(T x) __NOEXC
Definition: builtins.hpp:530
cl::__host_std::u_abs
s::cl_ulong u_abs(s::cl_ulong x) __NOEXC
Definition: builtins_integer.cpp:244
cl::__host_std::rotate
s::cl_long rotate(s::cl_long x, s::cl_long y) __NOEXC
Definition: builtins_integer.cpp:727
__NOEXC
#define __NOEXC
Definition: builtins.hpp:18
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
cl::__host_std::s_rhadd
s::cl_long s_rhadd(s::cl_long x, s::cl_long y) __NOEXC
Definition: builtins_integer.cpp:397
cl::sycl::cl_char
std::int8_t cl_char
Definition: aliases.hpp:79
cl::__host_std::s_mad24
s::cl_int s_mad24(s::cl_int x, s::cl_int y, s::cl_int z) __NOEXC
Definition: builtins_integer.cpp:843
builtins_helper.hpp
cl::__host_std::u_clamp
s::cl_ulong u_clamp(s::cl_ulong x, s::cl_ulong minval, s::cl_ulong maxval) __NOEXC
Definition: builtins_integer.cpp:418
cl::sycl::cl_ulong
std::uint64_t cl_ulong
Definition: aliases.hpp:86
cl::__host_std::u_mul_hi
s::cl_ulong u_mul_hi(s::cl_ulong x, s::cl_ulong y) __NOEXC
Definition: builtins_integer.cpp:521
cl::sycl::image_channel_order::a
@ a
cl::__host_std::u_max
s::cl_ulong u_max(s::cl_ulong x, s::cl_ulong y) __NOEXC
Definition: builtins_integer.cpp:649
cl::sycl::cl_int
std::int32_t cl_int
Definition: aliases.hpp:83
cl::__host_std::s_clamp
s::cl_long s_clamp(s::cl_long x, s::cl_long minval, s::cl_long maxval) __NOEXC
Definition: builtins_integer.cpp:444
cl::__host_std::s_mad_hi
s::cl_long s_mad_hi(s::cl_long x, s::cl_long minval, s::cl_long maxval) __NOEXC
Definition: builtins_integer.cpp:542
cl::__host_std::u_mad_hi
s::cl_ulong u_mad_hi(s::cl_ulong x, s::cl_ulong minval, s::cl_ulong maxval) __NOEXC
Definition: builtins_integer.cpp:564
cl::__host_std::u_upsample
s::cl_ulong u_upsample(s::cl_uint x, s::cl_uint y) __NOEXC
Definition: builtins_integer.cpp:782
cl::__host_std::s_abs_diff
s::cl_ulong s_abs_diff(s::cl_long x, s::cl_long y) __NOEXC
Definition: builtins_integer.cpp:289
cl::__host_std::u_abs_diff
s::cl_ulong u_abs_diff(s::cl_ulong x, s::cl_ulong y) __NOEXC
Definition: builtins_integer.cpp:270
cl::sycl::cl_uint
std::uint32_t cl_uint
Definition: aliases.hpp:84
cl::__host_std::popcount
s::cl_long popcount(s::cl_long x) __NOEXC
Definition: builtins_integer.cpp:827
cl::__host_std::u_sub_sat
s::cl_ulong u_sub_sat(s::cl_ulong x, s::cl_ulong y) __NOEXC
Definition: builtins_integer.cpp:749
MAKE_1V
#define MAKE_1V(Fun, Ret, Arg1)
Definition: builtins_helper.hpp:119
cl::sycl::cl_short
std::int16_t cl_short
Definition: aliases.hpp:81
cl::__host_std::s_abs
s::cl_ulong s_abs(s::cl_long x) __NOEXC
Definition: builtins_integer.cpp:254
cl::__host_std::u_add_sat
s::cl_ulong u_add_sat(s::cl_ulong x, s::cl_ulong y) __NOEXC
Definition: builtins_integer.cpp:307
cl::__host_std::u_mad_sat
s::cl_ulong u_mad_sat(s::cl_ulong a, s::cl_ulong b, s::cl_ulong c) __NOEXC
Definition: builtins_integer.cpp:608
cl::__host_std::s_min
s::cl_long s_min(s::cl_long x, s::cl_long y) __NOEXC
Definition: builtins_integer.cpp:671
cl::__host_std::s_mul_hi
s::cl_long s_mul_hi(s::cl_long x, s::cl_long y) __NOEXC
Definition: builtins_integer.cpp:503
min
simd< _Tp, _Abi > min(const simd< _Tp, _Abi > &, const simd< _Tp, _Abi > &) noexcept
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:12