lib/Headers/__clang_cuda_device_functions.h

19 // __DEVICE__ is a helper macro with common set of attributes for the wrappers
22 #pragma push_macro("__DEVICE__")
24 #define __DEVICE__ static __attribute__((always_inline, nothrow))
26 #define __DEVICE__ static __device__ __forceinline__
29 __DEVICE__ int __all(int __a) { return __nvvm_vote_all(__a); }
30 __DEVICE__ int __any(int __a) { return __nvvm_vote_any(__a); }
31 __DEVICE__ unsigned int __ballot(int __a) { return __nvvm_vote_ballot(__a); }
32 __DEVICE__ unsigned int __brev(unsigned int __a) { return __nv_brev(__a); }
33 __DEVICE__ unsigned long long __brevll(unsigned long long __a) {
37 __DEVICE__ void __brkpt() { __asm__ __volatile__("brkpt;"); }
38 __DEVICE__ void __brkpt(int __a) { __brkpt(); }
40 __DEVICE__ void __attribute__((overloadable)) __brkpt(void) {
43 __DEVICE__ void __attribute__((overloadable)) __brkpt(int __a) { __brkpt(); }
45 __DEVICE__ unsigned int __byte_perm(unsigned int __a, unsigned int __b,
49 __DEVICE__ int __clz(int __a) { return __nv_clz(__a); }
50 __DEVICE__ int __clzll(long long __a) { return __nv_clzll(__a); }
51 __DEVICE__ float __cosf(float __a) { return __nv_fast_cosf(__a); }
52 __DEVICE__ double __dAtomicAdd(double *__p, double __v) {
55 __DEVICE__ double __dAtomicAdd_block(double *__p, double __v) {
58 __DEVICE__ double __dAtomicAdd_system(double *__p, double __v) {
61 __DEVICE__ double __dadd_rd(double __a, double __b) {
64 __DEVICE__ double __dadd_rn(double __a, double __b) {
67 __DEVICE__ double __dadd_ru(double __a, double __b) {
70 __DEVICE__ double __dadd_rz(double __a, double __b) {
73 __DEVICE__ double __ddiv_rd(double __a, double __b) {
76 __DEVICE__ double __ddiv_rn(double __a, double __b) {
79 __DEVICE__ double __ddiv_ru(double __a, double __b) {
82 __DEVICE__ double __ddiv_rz(double __a, double __b) {
85 __DEVICE__ double __dmul_rd(double __a, double __b) {
88 __DEVICE__ double __dmul_rn(double __a, double __b) {
91 __DEVICE__ double __dmul_ru(double __a, double __b) {
94 __DEVICE__ double __dmul_rz(double __a, double __b) {
97 __DEVICE__ float __double2float_rd(double __a) {
100 __DEVICE__ float __double2float_rn(double __a) {
103 __DEVICE__ float __double2float_ru(double __a) {
106 __DEVICE__ float __double2float_rz(double __a) {
109 __DEVICE__ int __double2hiint(double __a) { return __nv_double2hiint(__a); }
110 __DEVICE__ int __double2int_rd(double __a) { return __nv_double2int_rd(__a); }
111 __DEVICE__ int __double2int_rn(double __a) { return __nv_double2int_rn(__a); }
112 __DEVICE__ int __double2int_ru(double __a) { return __nv_double2int_ru(__a); }
113 __DEVICE__ int __double2int_rz(double __a) { return __nv_double2int_rz(__a); }
114 __DEVICE__ long long __double2ll_rd(double __a) {
117 __DEVICE__ long long __double2ll_rn(double __a) {
120 __DEVICE__ long long __double2ll_ru(double __a) {
123 __DEVICE__ long long __double2ll_rz(double __a) {
126 __DEVICE__ int __double2loint(double __a) { return __nv_double2loint(__a); }
127 __DEVICE__ unsigned int __double2uint_rd(double __a) {
130 __DEVICE__ unsigned int __double2uint_rn(double __a) {
133 __DEVICE__ unsigned int __double2uint_ru(double __a) {
136 __DEVICE__ unsigned int __double2uint_rz(double __a) {
139 __DEVICE__ unsigned long long __double2ull_rd(double __a) {
142 __DEVICE__ unsigned long long __double2ull_rn(double __a) {
145 __DEVICE__ unsigned long long __double2ull_ru(double __a) {
148 __DEVICE__ unsigned long long __double2ull_rz(double __a) {
151 __DEVICE__ long long __double_as_longlong(double __a) {
154 __DEVICE__ double __drcp_rd(double __a) { return __nv_drcp_rd(__a); }
155 __DEVICE__ double __drcp_rn(double __a) { return __nv_drcp_rn(__a); }
156 __DEVICE__ double __drcp_ru(double __a) { return __nv_drcp_ru(__a); }
157 __DEVICE__ double __drcp_rz(double __a) { return __nv_drcp_rz(__a); }
158 __DEVICE__ double __dsqrt_rd(double __a) { return __nv_dsqrt_rd(__a); }
159 __DEVICE__ double __dsqrt_rn(double __a) { return __nv_dsqrt_rn(__a); }
160 __DEVICE__ double __dsqrt_ru(double __a) { return __nv_dsqrt_ru(__a); }
161 __DEVICE__ double __dsqrt_rz(double __a) { return __nv_dsqrt_rz(__a); }
162 __DEVICE__ double __dsub_rd(double __a, double __b) {
165 __DEVICE__ double __dsub_rn(double __a, double __b) {
168 __DEVICE__ double __dsub_ru(double __a, double __b) {
171 __DEVICE__ double __dsub_rz(double __a, double __b) {
174 __DEVICE__ float __exp10f(float __a) { return __nv_fast_exp10f(__a); }
175 __DEVICE__ float __expf(float __a) { return __nv_fast_expf(__a); }
176 __DEVICE__ float __fAtomicAdd(float *__p, float __v) {
179 __DEVICE__ float __fAtomicAdd_block(float *__p, float __v) {
182 __DEVICE__ float __fAtomicAdd_system(float *__p, float __v) {
185 __DEVICE__ float __fAtomicExch(float *__p, float __v) {
189 __DEVICE__ float __fAtomicExch_block(float *__p, float __v) {
193 __DEVICE__ float __fAtomicExch_system(float *__p, float __v) {
197 __DEVICE__ float __fadd_rd(float __a, float __b) {
200 __DEVICE__ float __fadd_rn(float __a, float __b) {
203 __DEVICE__ float __fadd_ru(float __a, float __b) {
206 __DEVICE__ float __fadd_rz(float __a, float __b) {
209 __DEVICE__ float __fdiv_rd(float __a, float __b) {
212 __DEVICE__ float __fdiv_rn(float __a, float __b) {
215 __DEVICE__ float __fdiv_ru(float __a, float __b) {
218 __DEVICE__ float __fdiv_rz(float __a, float __b) {
221 __DEVICE__ float __fdividef(float __a, float __b) {
224 __DEVICE__ int __ffs(int __a) { return __nv_ffs(__a); }
225 __DEVICE__ int __ffsll(long long __a) { return __nv_ffsll(__a); }
226 __DEVICE__ int __finite(double __a) { return __nv_isfinited(__a); }
227 __DEVICE__ int __finitef(float __a) { return __nv_finitef(__a); }
229 __DEVICE__ int __finitel(long double __a);
231 __DEVICE__ int __float2int_rd(float __a) { return __nv_float2int_rd(__a); }
232 __DEVICE__ int __float2int_rn(float __a) { return __nv_float2int_rn(__a); }
233 __DEVICE__ int __float2int_ru(float __a) { return __nv_float2int_ru(__a); }
234 __DEVICE__ int __float2int_rz(float __a) { return __nv_float2int_rz(__a); }
235 __DEVICE__ long long __float2ll_rd(float __a) { return __nv_float2ll_rd(__a); }
236 __DEVICE__ long long __float2ll_rn(float __a) { return __nv_float2ll_rn(__a); }
237 __DEVICE__ long long __float2ll_ru(float __a) { return __nv_float2ll_ru(__a); }
238 __DEVICE__ long long __float2ll_rz(float __a) { return __nv_float2ll_rz(__a); }
239 __DEVICE__ unsigned int __float2uint_rd(float __a) {
242 __DEVICE__ unsigned int __float2uint_rn(float __a) {
245 __DEVICE__ unsigned int __float2uint_ru(float __a) {
248 __DEVICE__ unsigned int __float2uint_rz(float __a) {
251 __DEVICE__ unsigned long long __float2ull_rd(float __a) {
254 __DEVICE__ unsigned long long __float2ull_rn(float __a) {
257 __DEVICE__ unsigned long long __float2ull_ru(float __a) {
260 __DEVICE__ unsigned long long __float2ull_rz(float __a) {
263 __DEVICE__ int __float_as_int(float __a) { return __nv_float_as_int(__a); }
264 __DEVICE__ unsigned int __float_as_uint(float __a) {
267 __DEVICE__ double __fma_rd(double __a, double __b, double __c) {
270 __DEVICE__ double __fma_rn(double __a, double __b, double __c) {
273 __DEVICE__ double __fma_ru(double __a, double __b, double __c) {
276 __DEVICE__ double __fma_rz(double __a, double __b, double __c) {
279 __DEVICE__ float __fmaf_ieee_rd(float __a, float __b, float __c) {
282 __DEVICE__ float __fmaf_ieee_rn(float __a, float __b, float __c) {
285 __DEVICE__ float __fmaf_ieee_ru(float __a, float __b, float __c) {
288 __DEVICE__ float __fmaf_ieee_rz(float __a, float __b, float __c) {
291 __DEVICE__ float __fmaf_rd(float __a, float __b, float __c) {
294 __DEVICE__ float __fmaf_rn(float __a, float __b, float __c) {
297 __DEVICE__ float __fmaf_ru(float __a, float __b, float __c) {
300 __DEVICE__ float __fmaf_rz(float __a, float __b, float __c) {
303 __DEVICE__ float __fmul_rd(float __a, float __b) {
306 __DEVICE__ float __fmul_rn(float __a, float __b) {
309 __DEVICE__ float __fmul_ru(float __a, float __b) {
312 __DEVICE__ float __fmul_rz(float __a, float __b) {
315 __DEVICE__ float __frcp_rd(float __a) { return __nv_frcp_rd(__a); }
316 __DEVICE__ float __frcp_rn(float __a) { return __nv_frcp_rn(__a); }
317 __DEVICE__ float __frcp_ru(float __a) { return __nv_frcp_ru(__a); }
318 __DEVICE__ float __frcp_rz(float __a) { return __nv_frcp_rz(__a); }
319 __DEVICE__ float __frsqrt_rn(float __a) { return __nv_frsqrt_rn(__a); }
320 __DEVICE__ float __fsqrt_rd(float __a) { return __nv_fsqrt_rd(__a); }
321 __DEVICE__ float __fsqrt_rn(float __a) { return __nv_fsqrt_rn(__a); }
322 __DEVICE__ float __fsqrt_ru(float __a) { return __nv_fsqrt_ru(__a); }
323 __DEVICE__ float __fsqrt_rz(float __a) { return __nv_fsqrt_rz(__a); }
324 __DEVICE__ float __fsub_rd(float __a, float __b) {
327 __DEVICE__ float __fsub_rn(float __a, float __b) {
330 __DEVICE__ float __fsub_ru(float __a, float __b) {
333 __DEVICE__ float __fsub_rz(float __a, float __b) {
336 __DEVICE__ int __hadd(int __a, int __b) { return __nv_hadd(__a, __b); }
337 __DEVICE__ double __hiloint2double(int __a, int __b) {
340 __DEVICE__ int __iAtomicAdd(int *__p, int __v) {
343 __DEVICE__ int __iAtomicAdd_block(int *__p, int __v) {
346 __DEVICE__ int __iAtomicAdd_system(int *__p, int __v) {
349 __DEVICE__ int __iAtomicAnd(int *__p, int __v) {
352 __DEVICE__ int __iAtomicAnd_block(int *__p, int __v) {
355 __DEVICE__ int __iAtomicAnd_system(int *__p, int __v) {
358 __DEVICE__ int __iAtomicCAS(int *__p, int __cmp, int __v) {
361 __DEVICE__ int __iAtomicCAS_block(int *__p, int __cmp, int __v) {
364 __DEVICE__ int __iAtomicCAS_system(int *__p, int __cmp, int __v) {
367 __DEVICE__ int __iAtomicExch(int *__p, int __v) {
370 __DEVICE__ int __iAtomicExch_block(int *__p, int __v) {
373 __DEVICE__ int __iAtomicExch_system(int *__p, int __v) {
376 __DEVICE__ int __iAtomicMax(int *__p, int __v) {
379 __DEVICE__ int __iAtomicMax_block(int *__p, int __v) {
382 __DEVICE__ int __iAtomicMax_system(int *__p, int __v) {
385 __DEVICE__ int __iAtomicMin(int *__p, int __v) {
388 __DEVICE__ int __iAtomicMin_block(int *__p, int __v) {
391 __DEVICE__ int __iAtomicMin_system(int *__p, int __v) {
394 __DEVICE__ int __iAtomicOr(int *__p, int __v) {
397 __DEVICE__ int __iAtomicOr_block(int *__p, int __v) {
400 __DEVICE__ int __iAtomicOr_system(int *__p, int __v) {
403 __DEVICE__ int __iAtomicXor(int *__p, int __v) {
406 __DEVICE__ int __iAtomicXor_block(int *__p, int __v) {
409 __DEVICE__ int __iAtomicXor_system(int *__p, int __v) {
412 __DEVICE__ long long __illAtomicMax(long long *__p, long long __v) {
415 __DEVICE__ long long __illAtomicMax_block(long long *__p, long long __v) {
418 __DEVICE__ long long __illAtomicMax_system(long long *__p, long long __v) {
421 __DEVICE__ long long __illAtomicMin(long long *__p, long long __v) {
424 __DEVICE__ long long __illAtomicMin_block(long long *__p, long long __v) {
427 __DEVICE__ long long __illAtomicMin_system(long long *__p, long long __v) {
430 __DEVICE__ double __int2double_rn(int __a) { return __nv_int2double_rn(__a); }
431 __DEVICE__ float __int2float_rd(int __a) { return __nv_int2float_rd(__a); }
432 __DEVICE__ float __int2float_rn(int __a) { return __nv_int2float_rn(__a); }
433 __DEVICE__ float __int2float_ru(int __a) { return __nv_int2float_ru(__a); }
434 __DEVICE__ float __int2float_rz(int __a) { return __nv_int2float_rz(__a); }
435 __DEVICE__ float __int_as_float(int __a) { return __nv_int_as_float(__a); }
436 __DEVICE__ int __isfinited(double __a) { return __nv_isfinited(__a); }
437 __DEVICE__ int __isinf(double __a) { return __nv_isinfd(__a); }
438 __DEVICE__ int __isinff(float __a) { return __nv_isinff(__a); }
440 __DEVICE__ int __isinfl(long double __a);
442 __DEVICE__ int __isnan(double __a) { return __nv_isnand(__a); }
443 __DEVICE__ int __isnanf(float __a) { return __nv_isnanf(__a); }
445 __DEVICE__ int __isnanl(long double __a);
447 __DEVICE__ double __ll2double_rd(long long __a) {
450 __DEVICE__ double __ll2double_rn(long long __a) {
453 __DEVICE__ double __ll2double_ru(long long __a) {
456 __DEVICE__ double __ll2double_rz(long long __a) {
459 __DEVICE__ float __ll2float_rd(long long __a) { return __nv_ll2float_rd(__a); }
460 __DEVICE__ float __ll2float_rn(long long __a) { return __nv_ll2float_rn(__a); }
461 __DEVICE__ float __ll2float_ru(long long __a) { return __nv_ll2float_ru(__a); }
462 __DEVICE__ float __ll2float_rz(long long __a) { return __nv_ll2float_rz(__a); }
463 __DEVICE__ long long __llAtomicAnd(long long *__p, long long __v) {
466 __DEVICE__ long long __llAtomicAnd_block(long long *__p, long long __v) {
469 __DEVICE__ long long __llAtomicAnd_system(long long *__p, long long __v) {
472 __DEVICE__ long long __llAtomicOr(long long *__p, long long __v) {
475 __DEVICE__ long long __llAtomicOr_block(long long *__p, long long __v) {
478 __DEVICE__ long long __llAtomicOr_system(long long *__p, long long __v) {
481 __DEVICE__ long long __llAtomicXor(long long *__p, long long __v) {
484 __DEVICE__ long long __llAtomicXor_block(long long *__p, long long __v) {
487 __DEVICE__ long long __llAtomicXor_system(long long *__p, long long __v) {
490 __DEVICE__ float __log10f(float __a) { return __nv_fast_log10f(__a); }
491 __DEVICE__ float __log2f(float __a) { return __nv_fast_log2f(__a); }
492 __DEVICE__ float __logf(float __a) { return __nv_fast_logf(__a); }
493 __DEVICE__ double __longlong_as_double(long long __a) {
496 __DEVICE__ int __mul24(int __a, int __b) { return __nv_mul24(__a, __b); }
497 __DEVICE__ long long __mul64hi(long long __a, long long __b) {
500 __DEVICE__ int __mulhi(int __a, int __b) { return __nv_mulhi(__a, __b); }
501 __DEVICE__ unsigned int __pm0(void) { return __nvvm_read_ptx_sreg_pm0(); }
502 __DEVICE__ unsigned int __pm1(void) { return __nvvm_read_ptx_sreg_pm1(); }
503 __DEVICE__ unsigned int __pm2(void) { return __nvvm_read_ptx_sreg_pm2(); }
504 __DEVICE__ unsigned int __pm3(void) { return __nvvm_read_ptx_sreg_pm3(); }
505 __DEVICE__ int __popc(int __a) { return __nv_popc(__a); }
506 __DEVICE__ int __popcll(long long __a) { return __nv_popcll(__a); }
507 __DEVICE__ float __powf(float __a, float __b) {
513 __DEVICE__ int __rhadd(int __a, int __b) { return __nv_rhadd(__a, __b); }
514 __DEVICE__ unsigned int __sad(int __a, int __b, unsigned int __c) {
517 __DEVICE__ float __saturatef(float __a) { return __nv_saturatef(__a); }
518 __DEVICE__ int __signbitd(double __a) { return __nv_signbitd(__a); }
519 __DEVICE__ int __signbitf(float __a) { return __nv_signbitf(__a); }
520 __DEVICE__ void __sincosf(float __a, float *__s, float *__c) {
523 __DEVICE__ float __sinf(float __a) { return __nv_fast_sinf(__a); }
524 __DEVICE__ int __syncthreads_and(int __a) { return __nvvm_bar0_and(__a); }
525 __DEVICE__ int __syncthreads_count(int __a) { return __nvvm_bar0_popc(__a); }
526 __DEVICE__ int __syncthreads_or(int __a) { return __nvvm_bar0_or(__a); }
527 __DEVICE__ float __tanf(float __a) { return __nv_fast_tanf(__a); }
528 __DEVICE__ void __threadfence(void) { __nvvm_membar_gl(); }
529 __DEVICE__ void __threadfence_block(void) { __nvvm_membar_cta(); };
530 __DEVICE__ void __threadfence_system(void) { __nvvm_membar_sys(); };
531 __DEVICE__ void __trap(void) { __asm__ __volatile__("trap;"); }
532 __DEVICE__ unsigned int __uAtomicAdd(unsigned int *__p, unsigned int __v) {
535 __DEVICE__ unsigned int __uAtomicAdd_block(unsigned int *__p,
539 __DEVICE__ unsigned int __uAtomicAdd_system(unsigned int *__p,
543 __DEVICE__ unsigned int __uAtomicAnd(unsigned int *__p, unsigned int __v) {
546 __DEVICE__ unsigned int __uAtomicAnd_block(unsigned int *__p,
550 __DEVICE__ unsigned int __uAtomicAnd_system(unsigned int *__p,
554 __DEVICE__ unsigned int __uAtomicCAS(unsigned int *__p, unsigned int __cmp,
558 __DEVICE__ unsigned int
562 __DEVICE__ unsigned int
566 __DEVICE__ unsigned int __uAtomicDec(unsigned int *__p, unsigned int __v) {
569 __DEVICE__ unsigned int __uAtomicDec_block(unsigned int *__p,
573 __DEVICE__ unsigned int __uAtomicDec_system(unsigned int *__p,
577 __DEVICE__ unsigned int __uAtomicExch(unsigned int *__p, unsigned int __v) {
580 __DEVICE__ unsigned int __uAtomicExch_block(unsigned int *__p,
584 __DEVICE__ unsigned int __uAtomicExch_system(unsigned int *__p,
588 __DEVICE__ unsigned int __uAtomicInc(unsigned int *__p, unsigned int __v) {
591 __DEVICE__ unsigned int __uAtomicInc_block(unsigned int *__p,
595 __DEVICE__ unsigned int __uAtomicInc_system(unsigned int *__p,
599 __DEVICE__ unsigned int __uAtomicMax(unsigned int *__p, unsigned int __v) {
602 __DEVICE__ unsigned int __uAtomicMax_block(unsigned int *__p,
606 __DEVICE__ unsigned int __uAtomicMax_system(unsigned int *__p,
610 __DEVICE__ unsigned int __uAtomicMin(unsigned int *__p, unsigned int __v) {
613 __DEVICE__ unsigned int __uAtomicMin_block(unsigned int *__p,
617 __DEVICE__ unsigned int __uAtomicMin_system(unsigned int *__p,
621 __DEVICE__ unsigned int __uAtomicOr(unsigned int *__p, unsigned int __v) {
624 __DEVICE__ unsigned int __uAtomicOr_block(unsigned int *__p, unsigned int __v) {
627 __DEVICE__ unsigned int __uAtomicOr_system(unsigned int *__p,
631 __DEVICE__ unsigned int __uAtomicXor(unsigned int *__p, unsigned int __v) {
634 __DEVICE__ unsigned int __uAtomicXor_block(unsigned int *__p,
638 __DEVICE__ unsigned int __uAtomicXor_system(unsigned int *__p,
642 __DEVICE__ unsigned int __uhadd(unsigned int __a, unsigned int __b) {
645 __DEVICE__ double __uint2double_rn(unsigned int __a) {
648 __DEVICE__ float __uint2float_rd(unsigned int __a) {
651 __DEVICE__ float __uint2float_rn(unsigned int __a) {
654 __DEVICE__ float __uint2float_ru(unsigned int __a) {
657 __DEVICE__ float __uint2float_rz(unsigned int __a) {
660 __DEVICE__ float __uint_as_float(unsigned int __a) {
663 __DEVICE__ double __ull2double_rd(unsigned long long __a) {
666 __DEVICE__ double __ull2double_rn(unsigned long long __a) {
669 __DEVICE__ double __ull2double_ru(unsigned long long __a) {
672 __DEVICE__ double __ull2double_rz(unsigned long long __a) {
675 __DEVICE__ float __ull2float_rd(unsigned long long __a) {
678 __DEVICE__ float __ull2float_rn(unsigned long long __a) {
681 __DEVICE__ float __ull2float_ru(unsigned long long __a) {
684 __DEVICE__ float __ull2float_rz(unsigned long long __a) {
687 __DEVICE__ unsigned long long __ullAtomicAdd(unsigned long long *__p,
691 __DEVICE__ unsigned long long __ullAtomicAdd_block(unsigned long long *__p,
695 __DEVICE__ unsigned long long __ullAtomicAdd_system(unsigned long long *__p,
699 __DEVICE__ unsigned long long __ullAtomicAnd(unsigned long long *__p,
703 __DEVICE__ unsigned long long __ullAtomicAnd_block(unsigned long long *__p,
707 __DEVICE__ unsigned long long __ullAtomicAnd_system(unsigned long long *__p,
711 __DEVICE__ unsigned long long __ullAtomicCAS(unsigned long long *__p,
716 __DEVICE__ unsigned long long __ullAtomicCAS_block(unsigned long long *__p,
721 __DEVICE__ unsigned long long __ullAtomicCAS_system(unsigned long long *__p,
726 __DEVICE__ unsigned long long __ullAtomicExch(unsigned long long *__p,
730 __DEVICE__ unsigned long long __ullAtomicExch_block(unsigned long long *__p,
734 __DEVICE__ unsigned long long __ullAtomicExch_system(unsigned long long *__p,
738 __DEVICE__ unsigned long long __ullAtomicMax(unsigned long long *__p,
742 __DEVICE__ unsigned long long __ullAtomicMax_block(unsigned long long *__p,
746 __DEVICE__ unsigned long long __ullAtomicMax_system(unsigned long long *__p,
750 __DEVICE__ unsigned long long __ullAtomicMin(unsigned long long *__p,
754 __DEVICE__ unsigned long long __ullAtomicMin_block(unsigned long long *__p,
758 __DEVICE__ unsigned long long __ullAtomicMin_system(unsigned long long *__p,
762 __DEVICE__ unsigned long long __ullAtomicOr(unsigned long long *__p,
766 __DEVICE__ unsigned long long __ullAtomicOr_block(unsigned long long *__p,
770 __DEVICE__ unsigned long long __ullAtomicOr_system(unsigned long long *__p,
774 __DEVICE__ unsigned long long __ullAtomicXor(unsigned long long *__p,
778 __DEVICE__ unsigned long long __ullAtomicXor_block(unsigned long long *__p,
782 __DEVICE__ unsigned long long __ullAtomicXor_system(unsigned long long *__p,
786 __DEVICE__ unsigned int __umul24(unsigned int __a, unsigned int __b) {
789 __DEVICE__ unsigned long long __umul64hi(unsigned long long __a,
793 __DEVICE__ unsigned int __umulhi(unsigned int __a, unsigned int __b) {
796 __DEVICE__ unsigned int __urhadd(unsigned int __a, unsigned int __b) {
799 __DEVICE__ unsigned int __usad(unsigned int __a, unsigned int __b,
805 __DEVICE__ unsigned int __vabs2(unsigned int __a) { return __nv_vabs2(__a); }
806 __DEVICE__ unsigned int __vabs4(unsigned int __a) { return __nv_vabs4(__a); }
807 __DEVICE__ unsigned int __vabsdiffs2(unsigned int __a, unsigned int __b) {
810 __DEVICE__ unsigned int __vabsdiffs4(unsigned int __a, unsigned int __b) {
813 __DEVICE__ unsigned int __vabsdiffu2(unsigned int __a, unsigned int __b) {
816 __DEVICE__ unsigned int __vabsdiffu4(unsigned int __a, unsigned int __b) {
819 __DEVICE__ unsigned int __vabsss2(unsigned int __a) {
822 __DEVICE__ unsigned int __vabsss4(unsigned int __a) {
825 __DEVICE__ unsigned int __vadd2(unsigned int __a, unsigned int __b) {
828 __DEVICE__ unsigned int __vadd4(unsigned int __a, unsigned int __b) {
831 __DEVICE__ unsigned int __vaddss2(unsigned int __a, unsigned int __b) {
834 __DEVICE__ unsigned int __vaddss4(unsigned int __a, unsigned int __b) {
837 __DEVICE__ unsigned int __vaddus2(unsigned int __a, unsigned int __b) {
840 __DEVICE__ unsigned int __vaddus4(unsigned int __a, unsigned int __b) {
843 __DEVICE__ unsigned int __vavgs2(unsigned int __a, unsigned int __b) {
846 __DEVICE__ unsigned int __vavgs4(unsigned int __a, unsigned int __b) {
849 __DEVICE__ unsigned int __vavgu2(unsigned int __a, unsigned int __b) {
852 __DEVICE__ unsigned int __vavgu4(unsigned int __a, unsigned int __b) {
855 __DEVICE__ unsigned int __vcmpeq2(unsigned int __a, unsigned int __b) {
858 __DEVICE__ unsigned int __vcmpeq4(unsigned int __a, unsigned int __b) {
861 __DEVICE__ unsigned int __vcmpges2(unsigned int __a, unsigned int __b) {
864 __DEVICE__ unsigned int __vcmpges4(unsigned int __a, unsigned int __b) {
867 __DEVICE__ unsigned int __vcmpgeu2(unsigned int __a, unsigned int __b) {
870 __DEVICE__ unsigned int __vcmpgeu4(unsigned int __a, unsigned int __b) {
873 __DEVICE__ unsigned int __vcmpgts2(unsigned int __a, unsigned int __b) {
876 __DEVICE__ unsigned int __vcmpgts4(unsigned int __a, unsigned int __b) {
879 __DEVICE__ unsigned int __vcmpgtu2(unsigned int __a, unsigned int __b) {
882 __DEVICE__ unsigned int __vcmpgtu4(unsigned int __a, unsigned int __b) {
885 __DEVICE__ unsigned int __vcmples2(unsigned int __a, unsigned int __b) {
888 __DEVICE__ unsigned int __vcmples4(unsigned int __a, unsigned int __b) {
891 __DEVICE__ unsigned int __vcmpleu2(unsigned int __a, unsigned int __b) {
894 __DEVICE__ unsigned int __vcmpleu4(unsigned int __a, unsigned int __b) {
897 __DEVICE__ unsigned int __vcmplts2(unsigned int __a, unsigned int __b) {
900 __DEVICE__ unsigned int __vcmplts4(unsigned int __a, unsigned int __b) {
903 __DEVICE__ unsigned int __vcmpltu2(unsigned int __a, unsigned int __b) {
906 __DEVICE__ unsigned int __vcmpltu4(unsigned int __a, unsigned int __b) {
909 __DEVICE__ unsigned int __vcmpne2(unsigned int __a, unsigned int __b) {
912 __DEVICE__ unsigned int __vcmpne4(unsigned int __a, unsigned int __b) {
915 __DEVICE__ unsigned int __vhaddu2(unsigned int __a, unsigned int __b) {
918 __DEVICE__ unsigned int __vhaddu4(unsigned int __a, unsigned int __b) {
921 __DEVICE__ unsigned int __vmaxs2(unsigned int __a, unsigned int __b) {
924 __DEVICE__ unsigned int __vmaxs4(unsigned int __a, unsigned int __b) {
927 __DEVICE__ unsigned int __vmaxu2(unsigned int __a, unsigned int __b) {
930 __DEVICE__ unsigned int __vmaxu4(unsigned int __a, unsigned int __b) {
933 __DEVICE__ unsigned int __vmins2(unsigned int __a, unsigned int __b) {
936 __DEVICE__ unsigned int __vmins4(unsigned int __a, unsigned int __b) {
939 __DEVICE__ unsigned int __vminu2(unsigned int __a, unsigned int __b) {
942 __DEVICE__ unsigned int __vminu4(unsigned int __a, unsigned int __b) {
945 __DEVICE__ unsigned int __vneg2(unsigned int __a) { return __nv_vneg2(__a); }
946 __DEVICE__ unsigned int __vneg4(unsigned int __a) { return __nv_vneg4(__a); }
947 __DEVICE__ unsigned int __vnegss2(unsigned int __a) {
950 __DEVICE__ unsigned int __vnegss4(unsigned int __a) {
953 __DEVICE__ unsigned int __vsads2(unsigned int __a, unsigned int __b) {
956 __DEVICE__ unsigned int __vsads4(unsigned int __a, unsigned int __b) {
959 __DEVICE__ unsigned int __vsadu2(unsigned int __a, unsigned int __b) {
962 __DEVICE__ unsigned int __vsadu4(unsigned int __a, unsigned int __b) {
965 __DEVICE__ unsigned int __vseteq2(unsigned int __a, unsigned int __b) {
968 __DEVICE__ unsigned int __vseteq4(unsigned int __a, unsigned int __b) {
971 __DEVICE__ unsigned int __vsetges2(unsigned int __a, unsigned int __b) {
974 __DEVICE__ unsigned int __vsetges4(unsigned int __a, unsigned int __b) {
977 __DEVICE__ unsigned int __vsetgeu2(unsigned int __a, unsigned int __b) {
980 __DEVICE__ unsigned int __vsetgeu4(unsigned int __a, unsigned int __b) {
983 __DEVICE__ unsigned int __vsetgts2(unsigned int __a, unsigned int __b) {
986 __DEVICE__ unsigned int __vsetgts4(unsigned int __a, unsigned int __b) {
989 __DEVICE__ unsigned int __vsetgtu2(unsigned int __a, unsigned int __b) {
992 __DEVICE__ unsigned int __vsetgtu4(unsigned int __a, unsigned int __b) {
995 __DEVICE__ unsigned int __vsetles2(unsigned int __a, unsigned int __b) {
998 __DEVICE__ unsigned int __vsetles4(unsigned int __a, unsigned int __b) {
1001 __DEVICE__ unsigned int __vsetleu2(unsigned int __a, unsigned int __b) {
1004 __DEVICE__ unsigned int __vsetleu4(unsigned int __a, unsigned int __b) {
1007 __DEVICE__ unsigned int __vsetlts2(unsigned int __a, unsigned int __b) {
1010 __DEVICE__ unsigned int __vsetlts4(unsigned int __a, unsigned int __b) {
1013 __DEVICE__ unsigned int __vsetltu2(unsigned int __a, unsigned int __b) {
1016 __DEVICE__ unsigned int __vsetltu4(unsigned int __a, unsigned int __b) {
1019 __DEVICE__ unsigned int __vsetne2(unsigned int __a, unsigned int __b) {
1022 __DEVICE__ unsigned int __vsetne4(unsigned int __a, unsigned int __b) {
1025 __DEVICE__ unsigned int __vsub2(unsigned int __a, unsigned int __b) {
1028 __DEVICE__ unsigned int __vsub4(unsigned int __a, unsigned int __b) {
1031 __DEVICE__ unsigned int __vsubss2(unsigned int __a, unsigned int __b) {
1034 __DEVICE__ unsigned int __vsubss4(unsigned int __a, unsigned int __b) {
1037 __DEVICE__ unsigned int __vsubus2(unsigned int __a, unsigned int __b) {
1040 __DEVICE__ unsigned int __vsubus4(unsigned int __a, unsigned int __b) {
1051 __DEVICE__ unsigned int __bool2mask(unsigned int __a, int shift) {
1054 __DEVICE__ unsigned int __vabs2(unsigned int __a) {
1061 __DEVICE__ unsigned int __vabs4(unsigned int __a) {
1068 __DEVICE__ unsigned int __vabsdiffs2(unsigned int __a, unsigned int __b) {
1076 __DEVICE__ unsigned int __vabsdiffs4(unsigned int __a, unsigned int __b) {
1083 __DEVICE__ unsigned int __vabsdiffu2(unsigned int __a, unsigned int __b) {
1090 __DEVICE__ unsigned int __vabsdiffu4(unsigned int __a, unsigned int __b) {
1097 __DEVICE__ unsigned int __vabsss2(unsigned int __a) {
1104 __DEVICE__ unsigned int __vabsss4(unsigned int __a) {
1111 __DEVICE__ unsigned int __vadd2(unsigned int __a, unsigned int __b) {
1118 __DEVICE__ unsigned int __vadd4(unsigned int __a, unsigned int __b) {
1125 __DEVICE__ unsigned int __vaddss2(unsigned int __a, unsigned int __b) {
1132 __DEVICE__ unsigned int __vaddss4(unsigned int __a, unsigned int __b) {
1139 __DEVICE__ unsigned int __vaddus2(unsigned int __a, unsigned int __b) {
1146 __DEVICE__ unsigned int __vaddus4(unsigned int __a, unsigned int __b) {
1153 __DEVICE__ unsigned int __vavgs2(unsigned int __a, unsigned int __b) {
1160 __DEVICE__ unsigned int __vavgs4(unsigned int __a, unsigned int __b) {
1167 __DEVICE__ unsigned int __vavgu2(unsigned int __a, unsigned int __b) {
1174 __DEVICE__ unsigned int __vavgu4(unsigned int __a, unsigned int __b) {
1181 __DEVICE__ unsigned int __vseteq2(unsigned int __a, unsigned int __b) {
1188 __DEVICE__ unsigned int __vcmpeq2(unsigned int __a, unsigned int __b) {
1191 __DEVICE__ unsigned int __vseteq4(unsigned int __a, unsigned int __b) {
1198 __DEVICE__ unsigned int __vcmpeq4(unsigned int __a, unsigned int __b) {
1201 __DEVICE__ unsigned int __vsetges2(unsigned int __a, unsigned int __b) {
1208 __DEVICE__ unsigned int __vcmpges2(unsigned int __a, unsigned int __b) {
1211 __DEVICE__ unsigned int __vsetges4(unsigned int __a, unsigned int __b) {
1218 __DEVICE__ unsigned int __vcmpges4(unsigned int __a, unsigned int __b) {
1221 __DEVICE__ unsigned int __vsetgeu2(unsigned int __a, unsigned int __b) {
1228 __DEVICE__ unsigned int __vcmpgeu2(unsigned int __a, unsigned int __b) {
1231 __DEVICE__ unsigned int __vsetgeu4(unsigned int __a, unsigned int __b) {
1238 __DEVICE__ unsigned int __vcmpgeu4(unsigned int __a, unsigned int __b) {
1241 __DEVICE__ unsigned int __vsetgts2(unsigned int __a, unsigned int __b) {
1248 __DEVICE__ unsigned int __vcmpgts2(unsigned int __a, unsigned int __b) {
1251 __DEVICE__ unsigned int __vsetgts4(unsigned int __a, unsigned int __b) {
1258 __DEVICE__ unsigned int __vcmpgts4(unsigned int __a, unsigned int __b) {
1261 __DEVICE__ unsigned int __vsetgtu2(unsigned int __a, unsigned int __b) {
1268 __DEVICE__ unsigned int __vcmpgtu2(unsigned int __a, unsigned int __b) {
1271 __DEVICE__ unsigned int __vsetgtu4(unsigned int __a, unsigned int __b) {
1278 __DEVICE__ unsigned int __vcmpgtu4(unsigned int __a, unsigned int __b) {
1281 __DEVICE__ unsigned int __vsetles2(unsigned int __a, unsigned int __b) {
1288 __DEVICE__ unsigned int __vcmples2(unsigned int __a, unsigned int __b) {
1291 __DEVICE__ unsigned int __vsetles4(unsigned int __a, unsigned int __b) {
1298 __DEVICE__ unsigned int __vcmples4(unsigned int __a, unsigned int __b) {
1301 __DEVICE__ unsigned int __vsetleu2(unsigned int __a, unsigned int __b) {
1308 __DEVICE__ unsigned int __vcmpleu2(unsigned int __a, unsigned int __b) {
1311 __DEVICE__ unsigned int __vsetleu4(unsigned int __a, unsigned int __b) {
1318 __DEVICE__ unsigned int __vcmpleu4(unsigned int __a, unsigned int __b) {
1321 __DEVICE__ unsigned int __vsetlts2(unsigned int __a, unsigned int __b) {
1328 __DEVICE__ unsigned int __vcmplts2(unsigned int __a, unsigned int __b) {
1331 __DEVICE__ unsigned int __vsetlts4(unsigned int __a, unsigned int __b) {
1338 __DEVICE__ unsigned int __vcmplts4(unsigned int __a, unsigned int __b) {
1341 __DEVICE__ unsigned int __vsetltu2(unsigned int __a, unsigned int __b) {
1348 __DEVICE__ unsigned int __vcmpltu2(unsigned int __a, unsigned int __b) {
1351 __DEVICE__ unsigned int __vsetltu4(unsigned int __a, unsigned int __b) {
1358 __DEVICE__ unsigned int __vcmpltu4(unsigned int __a, unsigned int __b) {
1361 __DEVICE__ unsigned int __vsetne2(unsigned int __a, unsigned int __b) {
1368 __DEVICE__ unsigned int __vcmpne2(unsigned int __a, unsigned int __b) {
1371 __DEVICE__ unsigned int __vsetne4(unsigned int __a, unsigned int __b) {
1378 __DEVICE__ unsigned int __vcmpne4(unsigned int __a, unsigned int __b) {
1387 __DEVICE__ unsigned int __vhaddu2(unsigned int __a, unsigned int __b) {
1390 __DEVICE__ unsigned int __vhaddu4(unsigned int __a, unsigned int __b) {
1394 __DEVICE__ unsigned int __vmaxs2(unsigned int __a, unsigned int __b) {
1408 __DEVICE__ unsigned int __vmaxs4(unsigned int __a, unsigned int __b) {
1415 __DEVICE__ unsigned int __vmaxu2(unsigned int __a, unsigned int __b) {
1422 __DEVICE__ unsigned int __vmaxu4(unsigned int __a, unsigned int __b) {
1429 __DEVICE__ unsigned int __vmins2(unsigned int __a, unsigned int __b) {
1436 __DEVICE__ unsigned int __vmins4(unsigned int __a, unsigned int __b) {
1443 __DEVICE__ unsigned int __vminu2(unsigned int __a, unsigned int __b) {
1450 __DEVICE__ unsigned int __vminu4(unsigned int __a, unsigned int __b) {
1457 __DEVICE__ unsigned int __vsads2(unsigned int __a, unsigned int __b) {
1464 __DEVICE__ unsigned int __vsads4(unsigned int __a, unsigned int __b) {
1471 __DEVICE__ unsigned int __vsadu2(unsigned int __a, unsigned int __b) {
1478 __DEVICE__ unsigned int __vsadu4(unsigned int __a, unsigned int __b) {
1486 __DEVICE__ unsigned int __vsub2(unsigned int __a, unsigned int __b) {
1493 __DEVICE__ unsigned int __vneg2(unsigned int __a) { return __vsub2(0, __a); }
1495 __DEVICE__ unsigned int __vsub4(unsigned int __a, unsigned int __b) {
1502 __DEVICE__ unsigned int __vneg4(unsigned int __a) { return __vsub4(0, __a); }
1503 __DEVICE__ unsigned int __vsubss2(unsigned int __a, unsigned int __b) {
1510 __DEVICE__ unsigned int __vnegss2(unsigned int __a) {
1513 __DEVICE__ unsigned int __vsubss4(unsigned int __a, unsigned int __b) {
1520 __DEVICE__ unsigned int __vnegss4(unsigned int __a) {
1523 __DEVICE__ unsigned int __vsubus2(unsigned int __a, unsigned int __b) {
1530 __DEVICE__ unsigned int __vsubus4(unsigned int __a, unsigned int __b) {
1542 __DEVICE__ /* clock_t= */ int clock() { return __nvvm_read_ptx_sreg_clock(); }
1544 __DEVICE__ long long clock64() { return __nvvm_read_ptx_sreg_clock64(); }
1549 __DEVICE__ void *memcpy(void *__a, const void *__b, size_t __c) {
1552 __DEVICE__ void *memset(void *__a, int __b, size_t __c) {
1557 #pragma pop_macro("__DEVICE__")