-
Notifications
You must be signed in to change notification settings - Fork 61
/
Copy path__clang_cuda_device_functions.h
1558 lines (1543 loc) · 56.7 KB
/
__clang_cuda_device_functions.h
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
282
283
284
285
286
287
288
289
290
291
292
293
294
295
296
297
298
299
300
301
302
303
304
305
306
307
308
309
310
311
312
313
314
315
316
317
318
319
320
321
322
323
324
325
326
327
328
329
330
331
332
333
334
335
336
337
338
339
340
341
342
343
344
345
346
347
348
349
350
351
352
353
354
355
356
357
358
359
360
361
362
363
364
365
366
367
368
369
370
371
372
373
374
375
376
377
378
379
380
381
382
383
384
385
386
387
388
389
390
391
392
393
394
395
396
397
398
399
400
401
402
403
404
405
406
407
408
409
410
411
412
413
414
415
416
417
418
419
420
421
422
423
424
425
426
427
428
429
430
431
432
433
434
435
436
437
438
439
440
441
442
443
444
445
446
447
448
449
450
451
452
453
454
455
456
457
458
459
460
461
462
463
464
465
466
467
468
469
470
471
472
473
474
475
476
477
478
479
480
481
482
483
484
485
486
487
488
489
490
491
492
493
494
495
496
497
498
499
500
501
502
503
504
505
506
507
508
509
510
511
512
513
514
515
516
517
518
519
520
521
522
523
524
525
526
527
528
529
530
531
532
533
534
535
536
537
538
539
540
541
542
543
544
545
546
547
548
549
550
551
552
553
554
555
556
557
558
559
560
561
562
563
564
565
566
567
568
569
570
571
572
573
574
575
576
577
578
579
580
581
582
583
584
585
586
587
588
589
590
591
592
593
594
595
596
597
598
599
600
601
602
603
604
605
606
607
608
609
610
611
612
613
614
615
616
617
618
619
620
621
622
623
624
625
626
627
628
629
630
631
632
633
634
635
636
637
638
639
640
641
642
643
644
645
646
647
648
649
650
651
652
653
654
655
656
657
658
659
660
661
662
663
664
665
666
667
668
669
670
671
672
673
674
675
676
677
678
679
680
681
682
683
684
685
686
687
688
689
690
691
692
693
694
695
696
697
698
699
700
701
702
703
704
705
706
707
708
709
710
711
712
713
714
715
716
717
718
719
720
721
722
723
724
725
726
727
728
729
730
731
732
733
734
735
736
737
738
739
740
741
742
743
744
745
746
747
748
749
750
751
752
753
754
755
756
757
758
759
760
761
762
763
764
765
766
767
768
769
770
771
772
773
774
775
776
777
778
779
780
781
782
783
784
785
786
787
788
789
790
791
792
793
794
795
796
797
798
799
800
801
802
803
804
805
806
807
808
809
810
811
812
813
814
815
816
817
818
819
820
821
822
823
824
825
826
827
828
829
830
831
832
833
834
835
836
837
838
839
840
841
842
843
844
845
846
847
848
849
850
851
852
853
854
855
856
857
858
859
860
861
862
863
864
865
866
867
868
869
870
871
872
873
874
875
876
877
878
879
880
881
882
883
884
885
886
887
888
889
890
891
892
893
894
895
896
897
898
899
900
901
902
903
904
905
906
907
908
909
910
911
912
913
914
915
916
917
918
919
920
921
922
923
924
925
926
927
928
929
930
931
932
933
934
935
936
937
938
939
940
941
942
943
944
945
946
947
948
949
950
951
952
953
954
955
956
957
958
959
960
961
962
963
964
965
966
967
968
969
970
971
972
973
974
975
976
977
978
979
980
981
982
983
984
985
986
987
988
989
990
991
992
993
994
995
996
997
998
999
1000
/*===---- __clang_cuda_device_functions.h - CUDA runtime support -----------===
*
* Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
* See https://llvm.org/LICENSE.txt for license information.
* SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
*
*===-----------------------------------------------------------------------===
*/
#ifndef __CLANG_CUDA_DEVICE_FUNCTIONS_H__
#define __CLANG_CUDA_DEVICE_FUNCTIONS_H__
#ifndef __OPENMP_NVPTX__
#if CUDA_VERSION < 9000
#error This file is intended to be used with CUDA-9+ only.
#endif
#endif
// __DEVICE__ is a helper macro with common set of attributes for the wrappers
// we implement in this file. We need static in order to avoid emitting unused
// functions and __forceinline__ helps inlining these wrappers at -O1.
#pragma push_macro("__DEVICE__")
#ifdef __OPENMP_NVPTX__
#define __DEVICE__ static __attribute__((always_inline, nothrow))
#else
#define __DEVICE__ static __device__ __forceinline__
#endif
__DEVICE__ int __all(int __a) { return __nvvm_vote_all(__a); }
__DEVICE__ int __any(int __a) { return __nvvm_vote_any(__a); }
__DEVICE__ unsigned int __ballot(int __a) { return __nvvm_vote_ballot(__a); }
__DEVICE__ unsigned int __brev(unsigned int __a) { return __nv_brev(__a); }
__DEVICE__ unsigned long long __brevll(unsigned long long __a) {
return __nv_brevll(__a);
}
#if defined(__cplusplus)
__DEVICE__ void __brkpt() { __asm__ __volatile__("brkpt;"); }
__DEVICE__ void __brkpt(int __a) { __brkpt(); }
#else
__DEVICE__ void __attribute__((overloadable)) __brkpt(void) {
__asm__ __volatile__("brkpt;");
}
__DEVICE__ void __attribute__((overloadable)) __brkpt(int __a) { __brkpt(); }
#endif
__DEVICE__ unsigned int __byte_perm(unsigned int __a, unsigned int __b,
unsigned int __c) {
return __nv_byte_perm(__a, __b, __c);
}
__DEVICE__ int __clz(int __a) { return __nv_clz(__a); }
__DEVICE__ int __clzll(long long __a) { return __nv_clzll(__a); }
__DEVICE__ float __cosf(float __a) { return __nv_fast_cosf(__a); }
__DEVICE__ double __dAtomicAdd(double *__p, double __v) {
return __nvvm_atom_add_gen_d(__p, __v);
}
__DEVICE__ double __dAtomicAdd_block(double *__p, double __v) {
return __nvvm_atom_cta_add_gen_d(__p, __v);
}
__DEVICE__ double __dAtomicAdd_system(double *__p, double __v) {
return __nvvm_atom_sys_add_gen_d(__p, __v);
}
__DEVICE__ double __dadd_rd(double __a, double __b) {
return __nv_dadd_rd(__a, __b);
}
__DEVICE__ double __dadd_rn(double __a, double __b) {
return __nv_dadd_rn(__a, __b);
}
__DEVICE__ double __dadd_ru(double __a, double __b) {
return __nv_dadd_ru(__a, __b);
}
__DEVICE__ double __dadd_rz(double __a, double __b) {
return __nv_dadd_rz(__a, __b);
}
__DEVICE__ double __ddiv_rd(double __a, double __b) {
return __nv_ddiv_rd(__a, __b);
}
__DEVICE__ double __ddiv_rn(double __a, double __b) {
return __nv_ddiv_rn(__a, __b);
}
__DEVICE__ double __ddiv_ru(double __a, double __b) {
return __nv_ddiv_ru(__a, __b);
}
__DEVICE__ double __ddiv_rz(double __a, double __b) {
return __nv_ddiv_rz(__a, __b);
}
__DEVICE__ double __dmul_rd(double __a, double __b) {
return __nv_dmul_rd(__a, __b);
}
__DEVICE__ double __dmul_rn(double __a, double __b) {
return __nv_dmul_rn(__a, __b);
}
__DEVICE__ double __dmul_ru(double __a, double __b) {
return __nv_dmul_ru(__a, __b);
}
__DEVICE__ double __dmul_rz(double __a, double __b) {
return __nv_dmul_rz(__a, __b);
}
__DEVICE__ float __double2float_rd(double __a) {
return __nv_double2float_rd(__a);
}
__DEVICE__ float __double2float_rn(double __a) {
return __nv_double2float_rn(__a);
}
__DEVICE__ float __double2float_ru(double __a) {
return __nv_double2float_ru(__a);
}
__DEVICE__ float __double2float_rz(double __a) {
return __nv_double2float_rz(__a);
}
__DEVICE__ int __double2hiint(double __a) { return __nv_double2hiint(__a); }
__DEVICE__ int __double2int_rd(double __a) { return __nv_double2int_rd(__a); }
__DEVICE__ int __double2int_rn(double __a) { return __nv_double2int_rn(__a); }
__DEVICE__ int __double2int_ru(double __a) { return __nv_double2int_ru(__a); }
__DEVICE__ int __double2int_rz(double __a) { return __nv_double2int_rz(__a); }
__DEVICE__ long long __double2ll_rd(double __a) {
return __nv_double2ll_rd(__a);
}
__DEVICE__ long long __double2ll_rn(double __a) {
return __nv_double2ll_rn(__a);
}
__DEVICE__ long long __double2ll_ru(double __a) {
return __nv_double2ll_ru(__a);
}
__DEVICE__ long long __double2ll_rz(double __a) {
return __nv_double2ll_rz(__a);
}
__DEVICE__ int __double2loint(double __a) { return __nv_double2loint(__a); }
__DEVICE__ unsigned int __double2uint_rd(double __a) {
return __nv_double2uint_rd(__a);
}
__DEVICE__ unsigned int __double2uint_rn(double __a) {
return __nv_double2uint_rn(__a);
}
__DEVICE__ unsigned int __double2uint_ru(double __a) {
return __nv_double2uint_ru(__a);
}
__DEVICE__ unsigned int __double2uint_rz(double __a) {
return __nv_double2uint_rz(__a);
}
__DEVICE__ unsigned long long __double2ull_rd(double __a) {
return __nv_double2ull_rd(__a);
}
__DEVICE__ unsigned long long __double2ull_rn(double __a) {
return __nv_double2ull_rn(__a);
}
__DEVICE__ unsigned long long __double2ull_ru(double __a) {
return __nv_double2ull_ru(__a);
}
__DEVICE__ unsigned long long __double2ull_rz(double __a) {
return __nv_double2ull_rz(__a);
}
__DEVICE__ long long __double_as_longlong(double __a) {
return __nv_double_as_longlong(__a);
}
__DEVICE__ double __drcp_rd(double __a) { return __nv_drcp_rd(__a); }
__DEVICE__ double __drcp_rn(double __a) { return __nv_drcp_rn(__a); }
__DEVICE__ double __drcp_ru(double __a) { return __nv_drcp_ru(__a); }
__DEVICE__ double __drcp_rz(double __a) { return __nv_drcp_rz(__a); }
__DEVICE__ double __dsqrt_rd(double __a) { return __nv_dsqrt_rd(__a); }
__DEVICE__ double __dsqrt_rn(double __a) { return __nv_dsqrt_rn(__a); }
__DEVICE__ double __dsqrt_ru(double __a) { return __nv_dsqrt_ru(__a); }
__DEVICE__ double __dsqrt_rz(double __a) { return __nv_dsqrt_rz(__a); }
__DEVICE__ double __dsub_rd(double __a, double __b) {
return __nv_dsub_rd(__a, __b);
}
__DEVICE__ double __dsub_rn(double __a, double __b) {
return __nv_dsub_rn(__a, __b);
}
__DEVICE__ double __dsub_ru(double __a, double __b) {
return __nv_dsub_ru(__a, __b);
}
__DEVICE__ double __dsub_rz(double __a, double __b) {
return __nv_dsub_rz(__a, __b);
}
__DEVICE__ float __exp10f(float __a) { return __nv_fast_exp10f(__a); }
__DEVICE__ float __expf(float __a) { return __nv_fast_expf(__a); }
__DEVICE__ float __fAtomicAdd(float *__p, float __v) {
return __nvvm_atom_add_gen_f(__p, __v);
}
__DEVICE__ float __fAtomicAdd_block(float *__p, float __v) {
return __nvvm_atom_cta_add_gen_f(__p, __v);
}
__DEVICE__ float __fAtomicAdd_system(float *__p, float __v) {
return __nvvm_atom_sys_add_gen_f(__p, __v);
}
__DEVICE__ float __fAtomicExch(float *__p, float __v) {
return __nv_int_as_float(
__nvvm_atom_xchg_gen_i((int *)__p, __nv_float_as_int(__v)));
}
__DEVICE__ float __fAtomicExch_block(float *__p, float __v) {
return __nv_int_as_float(
__nvvm_atom_cta_xchg_gen_i((int *)__p, __nv_float_as_int(__v)));
}
__DEVICE__ float __fAtomicExch_system(float *__p, float __v) {
return __nv_int_as_float(
__nvvm_atom_sys_xchg_gen_i((int *)__p, __nv_float_as_int(__v)));
}
__DEVICE__ float __fadd_rd(float __a, float __b) {
return __nv_fadd_rd(__a, __b);
}
__DEVICE__ float __fadd_rn(float __a, float __b) {
return __nv_fadd_rn(__a, __b);
}
__DEVICE__ float __fadd_ru(float __a, float __b) {
return __nv_fadd_ru(__a, __b);
}
__DEVICE__ float __fadd_rz(float __a, float __b) {
return __nv_fadd_rz(__a, __b);
}
__DEVICE__ float __fdiv_rd(float __a, float __b) {
return __nv_fdiv_rd(__a, __b);
}
__DEVICE__ float __fdiv_rn(float __a, float __b) {
return __nv_fdiv_rn(__a, __b);
}
__DEVICE__ float __fdiv_ru(float __a, float __b) {
return __nv_fdiv_ru(__a, __b);
}
__DEVICE__ float __fdiv_rz(float __a, float __b) {
return __nv_fdiv_rz(__a, __b);
}
__DEVICE__ float __fdividef(float __a, float __b) {
return __nv_fast_fdividef(__a, __b);
}
__DEVICE__ int __ffs(int __a) { return __nv_ffs(__a); }
__DEVICE__ int __ffsll(long long __a) { return __nv_ffsll(__a); }
__DEVICE__ int __finite(double __a) { return __nv_isfinited(__a); }
__DEVICE__ int __finitef(float __a) { return __nv_finitef(__a); }
#ifdef _MSC_VER
__DEVICE__ int __finitel(long double __a);
#endif
__DEVICE__ int __float2int_rd(float __a) { return __nv_float2int_rd(__a); }
__DEVICE__ int __float2int_rn(float __a) { return __nv_float2int_rn(__a); }
__DEVICE__ int __float2int_ru(float __a) { return __nv_float2int_ru(__a); }
__DEVICE__ int __float2int_rz(float __a) { return __nv_float2int_rz(__a); }
__DEVICE__ long long __float2ll_rd(float __a) { return __nv_float2ll_rd(__a); }
__DEVICE__ long long __float2ll_rn(float __a) { return __nv_float2ll_rn(__a); }
__DEVICE__ long long __float2ll_ru(float __a) { return __nv_float2ll_ru(__a); }
__DEVICE__ long long __float2ll_rz(float __a) { return __nv_float2ll_rz(__a); }
__DEVICE__ unsigned int __float2uint_rd(float __a) {
return __nv_float2uint_rd(__a);
}
__DEVICE__ unsigned int __float2uint_rn(float __a) {
return __nv_float2uint_rn(__a);
}
__DEVICE__ unsigned int __float2uint_ru(float __a) {
return __nv_float2uint_ru(__a);
}
__DEVICE__ unsigned int __float2uint_rz(float __a) {
return __nv_float2uint_rz(__a);
}
__DEVICE__ unsigned long long __float2ull_rd(float __a) {
return __nv_float2ull_rd(__a);
}
__DEVICE__ unsigned long long __float2ull_rn(float __a) {
return __nv_float2ull_rn(__a);
}
__DEVICE__ unsigned long long __float2ull_ru(float __a) {
return __nv_float2ull_ru(__a);
}
__DEVICE__ unsigned long long __float2ull_rz(float __a) {
return __nv_float2ull_rz(__a);
}
__DEVICE__ int __float_as_int(float __a) { return __nv_float_as_int(__a); }
__DEVICE__ unsigned int __float_as_uint(float __a) {
return __nv_float_as_uint(__a);
}
__DEVICE__ double __fma_rd(double __a, double __b, double __c) {
return __nv_fma_rd(__a, __b, __c);
}
__DEVICE__ double __fma_rn(double __a, double __b, double __c) {
return __nv_fma_rn(__a, __b, __c);
}
__DEVICE__ double __fma_ru(double __a, double __b, double __c) {
return __nv_fma_ru(__a, __b, __c);
}
__DEVICE__ double __fma_rz(double __a, double __b, double __c) {
return __nv_fma_rz(__a, __b, __c);
}
__DEVICE__ float __fmaf_ieee_rd(float __a, float __b, float __c) {
return __nv_fmaf_ieee_rd(__a, __b, __c);
}
__DEVICE__ float __fmaf_ieee_rn(float __a, float __b, float __c) {
return __nv_fmaf_ieee_rn(__a, __b, __c);
}
__DEVICE__ float __fmaf_ieee_ru(float __a, float __b, float __c) {
return __nv_fmaf_ieee_ru(__a, __b, __c);
}
__DEVICE__ float __fmaf_ieee_rz(float __a, float __b, float __c) {
return __nv_fmaf_ieee_rz(__a, __b, __c);
}
__DEVICE__ float __fmaf_rd(float __a, float __b, float __c) {
return __nv_fmaf_rd(__a, __b, __c);
}
__DEVICE__ float __fmaf_rn(float __a, float __b, float __c) {
return __nv_fmaf_rn(__a, __b, __c);
}
__DEVICE__ float __fmaf_ru(float __a, float __b, float __c) {
return __nv_fmaf_ru(__a, __b, __c);
}
__DEVICE__ float __fmaf_rz(float __a, float __b, float __c) {
return __nv_fmaf_rz(__a, __b, __c);
}
__DEVICE__ float __fmul_rd(float __a, float __b) {
return __nv_fmul_rd(__a, __b);
}
__DEVICE__ float __fmul_rn(float __a, float __b) {
return __nv_fmul_rn(__a, __b);
}
__DEVICE__ float __fmul_ru(float __a, float __b) {
return __nv_fmul_ru(__a, __b);
}
__DEVICE__ float __fmul_rz(float __a, float __b) {
return __nv_fmul_rz(__a, __b);
}
__DEVICE__ float __frcp_rd(float __a) { return __nv_frcp_rd(__a); }
__DEVICE__ float __frcp_rn(float __a) { return __nv_frcp_rn(__a); }
__DEVICE__ float __frcp_ru(float __a) { return __nv_frcp_ru(__a); }
__DEVICE__ float __frcp_rz(float __a) { return __nv_frcp_rz(__a); }
__DEVICE__ float __frsqrt_rn(float __a) { return __nv_frsqrt_rn(__a); }
__DEVICE__ float __fsqrt_rd(float __a) { return __nv_fsqrt_rd(__a); }
__DEVICE__ float __fsqrt_rn(float __a) { return __nv_fsqrt_rn(__a); }
__DEVICE__ float __fsqrt_ru(float __a) { return __nv_fsqrt_ru(__a); }
__DEVICE__ float __fsqrt_rz(float __a) { return __nv_fsqrt_rz(__a); }
__DEVICE__ float __fsub_rd(float __a, float __b) {
return __nv_fsub_rd(__a, __b);
}
__DEVICE__ float __fsub_rn(float __a, float __b) {
return __nv_fsub_rn(__a, __b);
}
__DEVICE__ float __fsub_ru(float __a, float __b) {
return __nv_fsub_ru(__a, __b);
}
__DEVICE__ float __fsub_rz(float __a, float __b) {
return __nv_fsub_rz(__a, __b);
}
__DEVICE__ int __hadd(int __a, int __b) { return __nv_hadd(__a, __b); }
__DEVICE__ double __hiloint2double(int __a, int __b) {
return __nv_hiloint2double(__a, __b);
}
__DEVICE__ int __iAtomicAdd(int *__p, int __v) {
return __nvvm_atom_add_gen_i(__p, __v);
}
__DEVICE__ int __iAtomicAdd_block(int *__p, int __v) {
return __nvvm_atom_cta_add_gen_i(__p, __v);
}
__DEVICE__ int __iAtomicAdd_system(int *__p, int __v) {
return __nvvm_atom_sys_add_gen_i(__p, __v);
}
__DEVICE__ int __iAtomicAnd(int *__p, int __v) {
return __nvvm_atom_and_gen_i(__p, __v);
}
__DEVICE__ int __iAtomicAnd_block(int *__p, int __v) {
return __nvvm_atom_cta_and_gen_i(__p, __v);
}
__DEVICE__ int __iAtomicAnd_system(int *__p, int __v) {
return __nvvm_atom_sys_and_gen_i(__p, __v);
}
__DEVICE__ int __iAtomicCAS(int *__p, int __cmp, int __v) {
return __nvvm_atom_cas_gen_i(__p, __cmp, __v);
}
__DEVICE__ int __iAtomicCAS_block(int *__p, int __cmp, int __v) {
return __nvvm_atom_cta_cas_gen_i(__p, __cmp, __v);
}
__DEVICE__ int __iAtomicCAS_system(int *__p, int __cmp, int __v) {
return __nvvm_atom_sys_cas_gen_i(__p, __cmp, __v);
}
__DEVICE__ int __iAtomicExch(int *__p, int __v) {
return __nvvm_atom_xchg_gen_i(__p, __v);
}
__DEVICE__ int __iAtomicExch_block(int *__p, int __v) {
return __nvvm_atom_cta_xchg_gen_i(__p, __v);
}
__DEVICE__ int __iAtomicExch_system(int *__p, int __v) {
return __nvvm_atom_sys_xchg_gen_i(__p, __v);
}
__DEVICE__ int __iAtomicMax(int *__p, int __v) {
return __nvvm_atom_max_gen_i(__p, __v);
}
__DEVICE__ int __iAtomicMax_block(int *__p, int __v) {
return __nvvm_atom_cta_max_gen_i(__p, __v);
}
__DEVICE__ int __iAtomicMax_system(int *__p, int __v) {
return __nvvm_atom_sys_max_gen_i(__p, __v);
}
__DEVICE__ int __iAtomicMin(int *__p, int __v) {
return __nvvm_atom_min_gen_i(__p, __v);
}
__DEVICE__ int __iAtomicMin_block(int *__p, int __v) {
return __nvvm_atom_cta_min_gen_i(__p, __v);
}
__DEVICE__ int __iAtomicMin_system(int *__p, int __v) {
return __nvvm_atom_sys_min_gen_i(__p, __v);
}
__DEVICE__ int __iAtomicOr(int *__p, int __v) {
return __nvvm_atom_or_gen_i(__p, __v);
}
__DEVICE__ int __iAtomicOr_block(int *__p, int __v) {
return __nvvm_atom_cta_or_gen_i(__p, __v);
}
__DEVICE__ int __iAtomicOr_system(int *__p, int __v) {
return __nvvm_atom_sys_or_gen_i(__p, __v);
}
__DEVICE__ int __iAtomicXor(int *__p, int __v) {
return __nvvm_atom_xor_gen_i(__p, __v);
}
__DEVICE__ int __iAtomicXor_block(int *__p, int __v) {
return __nvvm_atom_cta_xor_gen_i(__p, __v);
}
__DEVICE__ int __iAtomicXor_system(int *__p, int __v) {
return __nvvm_atom_sys_xor_gen_i(__p, __v);
}
__DEVICE__ long long __illAtomicMax(long long *__p, long long __v) {
return __nvvm_atom_max_gen_ll(__p, __v);
}
__DEVICE__ long long __illAtomicMax_block(long long *__p, long long __v) {
return __nvvm_atom_cta_max_gen_ll(__p, __v);
}
__DEVICE__ long long __illAtomicMax_system(long long *__p, long long __v) {
return __nvvm_atom_sys_max_gen_ll(__p, __v);
}
__DEVICE__ long long __illAtomicMin(long long *__p, long long __v) {
return __nvvm_atom_min_gen_ll(__p, __v);
}
__DEVICE__ long long __illAtomicMin_block(long long *__p, long long __v) {
return __nvvm_atom_cta_min_gen_ll(__p, __v);
}
__DEVICE__ long long __illAtomicMin_system(long long *__p, long long __v) {
return __nvvm_atom_sys_min_gen_ll(__p, __v);
}
__DEVICE__ double __int2double_rn(int __a) { return __nv_int2double_rn(__a); }
__DEVICE__ float __int2float_rd(int __a) { return __nv_int2float_rd(__a); }
__DEVICE__ float __int2float_rn(int __a) { return __nv_int2float_rn(__a); }
__DEVICE__ float __int2float_ru(int __a) { return __nv_int2float_ru(__a); }
__DEVICE__ float __int2float_rz(int __a) { return __nv_int2float_rz(__a); }
__DEVICE__ float __int_as_float(int __a) { return __nv_int_as_float(__a); }
__DEVICE__ int __isfinited(double __a) { return __nv_isfinited(__a); }
__DEVICE__ int __isinf(double __a) { return __nv_isinfd(__a); }
__DEVICE__ int __isinff(float __a) { return __nv_isinff(__a); }
#ifdef _MSC_VER
__DEVICE__ int __isinfl(long double __a);
#endif
__DEVICE__ int __isnan(double __a) { return __nv_isnand(__a); }
__DEVICE__ int __isnanf(float __a) { return __nv_isnanf(__a); }
#ifdef _MSC_VER
__DEVICE__ int __isnanl(long double __a);
#endif
__DEVICE__ double __ll2double_rd(long long __a) {
return __nv_ll2double_rd(__a);
}
__DEVICE__ double __ll2double_rn(long long __a) {
return __nv_ll2double_rn(__a);
}
__DEVICE__ double __ll2double_ru(long long __a) {
return __nv_ll2double_ru(__a);
}
__DEVICE__ double __ll2double_rz(long long __a) {
return __nv_ll2double_rz(__a);
}
__DEVICE__ float __ll2float_rd(long long __a) { return __nv_ll2float_rd(__a); }
__DEVICE__ float __ll2float_rn(long long __a) { return __nv_ll2float_rn(__a); }
__DEVICE__ float __ll2float_ru(long long __a) { return __nv_ll2float_ru(__a); }
__DEVICE__ float __ll2float_rz(long long __a) { return __nv_ll2float_rz(__a); }
__DEVICE__ long long __llAtomicAnd(long long *__p, long long __v) {
return __nvvm_atom_and_gen_ll(__p, __v);
}
__DEVICE__ long long __llAtomicAnd_block(long long *__p, long long __v) {
return __nvvm_atom_cta_and_gen_ll(__p, __v);
}
__DEVICE__ long long __llAtomicAnd_system(long long *__p, long long __v) {
return __nvvm_atom_sys_and_gen_ll(__p, __v);
}
__DEVICE__ long long __llAtomicOr(long long *__p, long long __v) {
return __nvvm_atom_or_gen_ll(__p, __v);
}
__DEVICE__ long long __llAtomicOr_block(long long *__p, long long __v) {
return __nvvm_atom_cta_or_gen_ll(__p, __v);
}
__DEVICE__ long long __llAtomicOr_system(long long *__p, long long __v) {
return __nvvm_atom_sys_or_gen_ll(__p, __v);
}
__DEVICE__ long long __llAtomicXor(long long *__p, long long __v) {
return __nvvm_atom_xor_gen_ll(__p, __v);
}
__DEVICE__ long long __llAtomicXor_block(long long *__p, long long __v) {
return __nvvm_atom_cta_xor_gen_ll(__p, __v);
}
__DEVICE__ long long __llAtomicXor_system(long long *__p, long long __v) {
return __nvvm_atom_sys_xor_gen_ll(__p, __v);
}
__DEVICE__ float __log10f(float __a) { return __nv_fast_log10f(__a); }
__DEVICE__ float __log2f(float __a) { return __nv_fast_log2f(__a); }
__DEVICE__ float __logf(float __a) { return __nv_fast_logf(__a); }
__DEVICE__ double __longlong_as_double(long long __a) {
return __nv_longlong_as_double(__a);
}
__DEVICE__ int __mul24(int __a, int __b) { return __nv_mul24(__a, __b); }
__DEVICE__ long long __mul64hi(long long __a, long long __b) {
return __nv_mul64hi(__a, __b);
}
__DEVICE__ int __mulhi(int __a, int __b) { return __nv_mulhi(__a, __b); }
__DEVICE__ unsigned int __pm0(void) { return __nvvm_read_ptx_sreg_pm0(); }
__DEVICE__ unsigned int __pm1(void) { return __nvvm_read_ptx_sreg_pm1(); }
__DEVICE__ unsigned int __pm2(void) { return __nvvm_read_ptx_sreg_pm2(); }
__DEVICE__ unsigned int __pm3(void) { return __nvvm_read_ptx_sreg_pm3(); }
__DEVICE__ int __popc(unsigned int __a) { return __nv_popc(__a); }
__DEVICE__ int __popcll(unsigned long long __a) { return __nv_popcll(__a); }
__DEVICE__ float __powf(float __a, float __b) {
return __nv_fast_powf(__a, __b);
}
// Parameter must have a known integer value.
#define __prof_trigger(__a) __asm__ __volatile__("pmevent \t%0;" ::"i"(__a))
__DEVICE__ int __rhadd(int __a, int __b) { return __nv_rhadd(__a, __b); }
__DEVICE__ unsigned int __sad(int __a, int __b, unsigned int __c) {
return __nv_sad(__a, __b, __c);
}
__DEVICE__ float __saturatef(float __a) { return __nv_saturatef(__a); }
__DEVICE__ int __signbitd(double __a) { return __nv_signbitd(__a); }
__DEVICE__ int __signbitf(float __a) { return __nv_signbitf(__a); }
__DEVICE__ void __sincosf(float __a, float *__s, float *__c) {
return __nv_fast_sincosf(__a, __s, __c);
}
__DEVICE__ float __sinf(float __a) { return __nv_fast_sinf(__a); }
__DEVICE__ int __syncthreads_and(int __a) { return __nvvm_bar0_and(__a); }
__DEVICE__ int __syncthreads_count(int __a) { return __nvvm_bar0_popc(__a); }
__DEVICE__ int __syncthreads_or(int __a) { return __nvvm_bar0_or(__a); }
__DEVICE__ float __tanf(float __a) { return __nv_fast_tanf(__a); }
__DEVICE__ void __threadfence(void) { __nvvm_membar_gl(); }
__DEVICE__ void __threadfence_block(void) { __nvvm_membar_cta(); };
__DEVICE__ void __threadfence_system(void) { __nvvm_membar_sys(); };
__DEVICE__ void __trap(void) { __asm__ __volatile__("trap;"); }
__DEVICE__ unsigned int __uAtomicAdd(unsigned int *__p, unsigned int __v) {
return __nvvm_atom_add_gen_i((int *)__p, __v);
}
__DEVICE__ unsigned int __uAtomicAdd_block(unsigned int *__p,
unsigned int __v) {
return __nvvm_atom_cta_add_gen_i((int *)__p, __v);
}
__DEVICE__ unsigned int __uAtomicAdd_system(unsigned int *__p,
unsigned int __v) {
return __nvvm_atom_sys_add_gen_i((int *)__p, __v);
}
__DEVICE__ unsigned int __uAtomicAnd(unsigned int *__p, unsigned int __v) {
return __nvvm_atom_and_gen_i((int *)__p, __v);
}
__DEVICE__ unsigned int __uAtomicAnd_block(unsigned int *__p,
unsigned int __v) {
return __nvvm_atom_cta_and_gen_i((int *)__p, __v);
}
__DEVICE__ unsigned int __uAtomicAnd_system(unsigned int *__p,
unsigned int __v) {
return __nvvm_atom_sys_and_gen_i((int *)__p, __v);
}
__DEVICE__ unsigned int __uAtomicCAS(unsigned int *__p, unsigned int __cmp,
unsigned int __v) {
return __nvvm_atom_cas_gen_i((int *)__p, __cmp, __v);
}
__DEVICE__ unsigned int
__uAtomicCAS_block(unsigned int *__p, unsigned int __cmp, unsigned int __v) {
return __nvvm_atom_cta_cas_gen_i((int *)__p, __cmp, __v);
}
__DEVICE__ unsigned int
__uAtomicCAS_system(unsigned int *__p, unsigned int __cmp, unsigned int __v) {
return __nvvm_atom_sys_cas_gen_i((int *)__p, __cmp, __v);
}
__DEVICE__ unsigned int __uAtomicDec(unsigned int *__p, unsigned int __v) {
return __nvvm_atom_dec_gen_ui(__p, __v);
}
__DEVICE__ unsigned int __uAtomicDec_block(unsigned int *__p,
unsigned int __v) {
return __nvvm_atom_cta_dec_gen_ui(__p, __v);
}
__DEVICE__ unsigned int __uAtomicDec_system(unsigned int *__p,
unsigned int __v) {
return __nvvm_atom_sys_dec_gen_ui(__p, __v);
}
__DEVICE__ unsigned int __uAtomicExch(unsigned int *__p, unsigned int __v) {
return __nvvm_atom_xchg_gen_i((int *)__p, __v);
}
__DEVICE__ unsigned int __uAtomicExch_block(unsigned int *__p,
unsigned int __v) {
return __nvvm_atom_cta_xchg_gen_i((int *)__p, __v);
}
__DEVICE__ unsigned int __uAtomicExch_system(unsigned int *__p,
unsigned int __v) {
return __nvvm_atom_sys_xchg_gen_i((int *)__p, __v);
}
__DEVICE__ unsigned int __uAtomicInc(unsigned int *__p, unsigned int __v) {
return __nvvm_atom_inc_gen_ui(__p, __v);
}
__DEVICE__ unsigned int __uAtomicInc_block(unsigned int *__p,
unsigned int __v) {
return __nvvm_atom_cta_inc_gen_ui(__p, __v);
}
__DEVICE__ unsigned int __uAtomicInc_system(unsigned int *__p,
unsigned int __v) {
return __nvvm_atom_sys_inc_gen_ui(__p, __v);
}
__DEVICE__ unsigned int __uAtomicMax(unsigned int *__p, unsigned int __v) {
return __nvvm_atom_max_gen_ui(__p, __v);
}
__DEVICE__ unsigned int __uAtomicMax_block(unsigned int *__p,
unsigned int __v) {
return __nvvm_atom_cta_max_gen_ui(__p, __v);
}
__DEVICE__ unsigned int __uAtomicMax_system(unsigned int *__p,
unsigned int __v) {
return __nvvm_atom_sys_max_gen_ui(__p, __v);
}
__DEVICE__ unsigned int __uAtomicMin(unsigned int *__p, unsigned int __v) {
return __nvvm_atom_min_gen_ui(__p, __v);
}
__DEVICE__ unsigned int __uAtomicMin_block(unsigned int *__p,
unsigned int __v) {
return __nvvm_atom_cta_min_gen_ui(__p, __v);
}
__DEVICE__ unsigned int __uAtomicMin_system(unsigned int *__p,
unsigned int __v) {
return __nvvm_atom_sys_min_gen_ui(__p, __v);
}
__DEVICE__ unsigned int __uAtomicOr(unsigned int *__p, unsigned int __v) {
return __nvvm_atom_or_gen_i((int *)__p, __v);
}
__DEVICE__ unsigned int __uAtomicOr_block(unsigned int *__p, unsigned int __v) {
return __nvvm_atom_cta_or_gen_i((int *)__p, __v);
}
__DEVICE__ unsigned int __uAtomicOr_system(unsigned int *__p,
unsigned int __v) {
return __nvvm_atom_sys_or_gen_i((int *)__p, __v);
}
__DEVICE__ unsigned int __uAtomicXor(unsigned int *__p, unsigned int __v) {
return __nvvm_atom_xor_gen_i((int *)__p, __v);
}
__DEVICE__ unsigned int __uAtomicXor_block(unsigned int *__p,
unsigned int __v) {
return __nvvm_atom_cta_xor_gen_i((int *)__p, __v);
}
__DEVICE__ unsigned int __uAtomicXor_system(unsigned int *__p,
unsigned int __v) {
return __nvvm_atom_sys_xor_gen_i((int *)__p, __v);
}
__DEVICE__ unsigned int __uhadd(unsigned int __a, unsigned int __b) {
return __nv_uhadd(__a, __b);
}
__DEVICE__ double __uint2double_rn(unsigned int __a) {
return __nv_uint2double_rn(__a);
}
__DEVICE__ float __uint2float_rd(unsigned int __a) {
return __nv_uint2float_rd(__a);
}
__DEVICE__ float __uint2float_rn(unsigned int __a) {
return __nv_uint2float_rn(__a);
}
__DEVICE__ float __uint2float_ru(unsigned int __a) {
return __nv_uint2float_ru(__a);
}
__DEVICE__ float __uint2float_rz(unsigned int __a) {
return __nv_uint2float_rz(__a);
}
__DEVICE__ float __uint_as_float(unsigned int __a) {
return __nv_uint_as_float(__a);
} //
__DEVICE__ double __ull2double_rd(unsigned long long __a) {
return __nv_ull2double_rd(__a);
}
__DEVICE__ double __ull2double_rn(unsigned long long __a) {
return __nv_ull2double_rn(__a);
}
__DEVICE__ double __ull2double_ru(unsigned long long __a) {
return __nv_ull2double_ru(__a);
}
__DEVICE__ double __ull2double_rz(unsigned long long __a) {
return __nv_ull2double_rz(__a);
}
__DEVICE__ float __ull2float_rd(unsigned long long __a) {
return __nv_ull2float_rd(__a);
}
__DEVICE__ float __ull2float_rn(unsigned long long __a) {
return __nv_ull2float_rn(__a);
}
__DEVICE__ float __ull2float_ru(unsigned long long __a) {
return __nv_ull2float_ru(__a);
}
__DEVICE__ float __ull2float_rz(unsigned long long __a) {
return __nv_ull2float_rz(__a);
}
__DEVICE__ unsigned long long __ullAtomicAdd(unsigned long long *__p,
unsigned long long __v) {
return __nvvm_atom_add_gen_ll((long long *)__p, __v);
}
__DEVICE__ unsigned long long __ullAtomicAdd_block(unsigned long long *__p,
unsigned long long __v) {
return __nvvm_atom_cta_add_gen_ll((long long *)__p, __v);
}
__DEVICE__ unsigned long long __ullAtomicAdd_system(unsigned long long *__p,
unsigned long long __v) {
return __nvvm_atom_sys_add_gen_ll((long long *)__p, __v);
}
__DEVICE__ unsigned long long __ullAtomicAnd(unsigned long long *__p,
unsigned long long __v) {
return __nvvm_atom_and_gen_ll((long long *)__p, __v);
}
__DEVICE__ unsigned long long __ullAtomicAnd_block(unsigned long long *__p,
unsigned long long __v) {
return __nvvm_atom_cta_and_gen_ll((long long *)__p, __v);
}
__DEVICE__ unsigned long long __ullAtomicAnd_system(unsigned long long *__p,
unsigned long long __v) {
return __nvvm_atom_sys_and_gen_ll((long long *)__p, __v);
}
__DEVICE__ unsigned long long __ullAtomicCAS(unsigned long long *__p,
unsigned long long __cmp,
unsigned long long __v) {
return __nvvm_atom_cas_gen_ll((long long *)__p, __cmp, __v);
}
__DEVICE__ unsigned long long __ullAtomicCAS_block(unsigned long long *__p,
unsigned long long __cmp,
unsigned long long __v) {
return __nvvm_atom_cta_cas_gen_ll((long long *)__p, __cmp, __v);
}
__DEVICE__ unsigned long long __ullAtomicCAS_system(unsigned long long *__p,
unsigned long long __cmp,
unsigned long long __v) {
return __nvvm_atom_sys_cas_gen_ll((long long *)__p, __cmp, __v);
}
__DEVICE__ unsigned long long __ullAtomicExch(unsigned long long *__p,
unsigned long long __v) {
return __nvvm_atom_xchg_gen_ll((long long *)__p, __v);
}
__DEVICE__ unsigned long long __ullAtomicExch_block(unsigned long long *__p,
unsigned long long __v) {
return __nvvm_atom_cta_xchg_gen_ll((long long *)__p, __v);
}
__DEVICE__ unsigned long long __ullAtomicExch_system(unsigned long long *__p,
unsigned long long __v) {
return __nvvm_atom_sys_xchg_gen_ll((long long *)__p, __v);
}
__DEVICE__ unsigned long long __ullAtomicMax(unsigned long long *__p,
unsigned long long __v) {
return __nvvm_atom_max_gen_ull(__p, __v);
}
__DEVICE__ unsigned long long __ullAtomicMax_block(unsigned long long *__p,
unsigned long long __v) {
return __nvvm_atom_cta_max_gen_ull(__p, __v);
}
__DEVICE__ unsigned long long __ullAtomicMax_system(unsigned long long *__p,
unsigned long long __v) {
return __nvvm_atom_sys_max_gen_ull(__p, __v);
}
__DEVICE__ unsigned long long __ullAtomicMin(unsigned long long *__p,
unsigned long long __v) {
return __nvvm_atom_min_gen_ull(__p, __v);
}
__DEVICE__ unsigned long long __ullAtomicMin_block(unsigned long long *__p,
unsigned long long __v) {
return __nvvm_atom_cta_min_gen_ull(__p, __v);
}
__DEVICE__ unsigned long long __ullAtomicMin_system(unsigned long long *__p,
unsigned long long __v) {
return __nvvm_atom_sys_min_gen_ull(__p, __v);
}
__DEVICE__ unsigned long long __ullAtomicOr(unsigned long long *__p,
unsigned long long __v) {
return __nvvm_atom_or_gen_ll((long long *)__p, __v);
}
__DEVICE__ unsigned long long __ullAtomicOr_block(unsigned long long *__p,
unsigned long long __v) {
return __nvvm_atom_cta_or_gen_ll((long long *)__p, __v);
}
__DEVICE__ unsigned long long __ullAtomicOr_system(unsigned long long *__p,
unsigned long long __v) {
return __nvvm_atom_sys_or_gen_ll((long long *)__p, __v);
}
__DEVICE__ unsigned long long __ullAtomicXor(unsigned long long *__p,
unsigned long long __v) {
return __nvvm_atom_xor_gen_ll((long long *)__p, __v);
}
__DEVICE__ unsigned long long __ullAtomicXor_block(unsigned long long *__p,
unsigned long long __v) {
return __nvvm_atom_cta_xor_gen_ll((long long *)__p, __v);
}
__DEVICE__ unsigned long long __ullAtomicXor_system(unsigned long long *__p,
unsigned long long __v) {
return __nvvm_atom_sys_xor_gen_ll((long long *)__p, __v);
}
__DEVICE__ unsigned int __umul24(unsigned int __a, unsigned int __b) {
return __nv_umul24(__a, __b);
}
__DEVICE__ unsigned long long __umul64hi(unsigned long long __a,
unsigned long long __b) {
return __nv_umul64hi(__a, __b);
}
__DEVICE__ unsigned int __umulhi(unsigned int __a, unsigned int __b) {
return __nv_umulhi(__a, __b);
}
__DEVICE__ unsigned int __urhadd(unsigned int __a, unsigned int __b) {
return __nv_urhadd(__a, __b);
}
__DEVICE__ unsigned int __usad(unsigned int __a, unsigned int __b,
unsigned int __c) {
return __nv_usad(__a, __b, __c);
}
#if CUDA_VERSION >= 9000 && CUDA_VERSION < 9020
__DEVICE__ unsigned int __vabs2(unsigned int __a) { return __nv_vabs2(__a); }
__DEVICE__ unsigned int __vabs4(unsigned int __a) { return __nv_vabs4(__a); }
__DEVICE__ unsigned int __vabsdiffs2(unsigned int __a, unsigned int __b) {
return __nv_vabsdiffs2(__a, __b);
}
__DEVICE__ unsigned int __vabsdiffs4(unsigned int __a, unsigned int __b) {
return __nv_vabsdiffs4(__a, __b);
}
__DEVICE__ unsigned int __vabsdiffu2(unsigned int __a, unsigned int __b) {
return __nv_vabsdiffu2(__a, __b);
}
__DEVICE__ unsigned int __vabsdiffu4(unsigned int __a, unsigned int __b) {
return __nv_vabsdiffu4(__a, __b);
}
__DEVICE__ unsigned int __vabsss2(unsigned int __a) {
return __nv_vabsss2(__a);
}
__DEVICE__ unsigned int __vabsss4(unsigned int __a) {
return __nv_vabsss4(__a);
}
__DEVICE__ unsigned int __vadd2(unsigned int __a, unsigned int __b) {
return __nv_vadd2(__a, __b);
}
__DEVICE__ unsigned int __vadd4(unsigned int __a, unsigned int __b) {
return __nv_vadd4(__a, __b);
}
__DEVICE__ unsigned int __vaddss2(unsigned int __a, unsigned int __b) {
return __nv_vaddss2(__a, __b);
}
__DEVICE__ unsigned int __vaddss4(unsigned int __a, unsigned int __b) {
return __nv_vaddss4(__a, __b);
}
__DEVICE__ unsigned int __vaddus2(unsigned int __a, unsigned int __b) {
return __nv_vaddus2(__a, __b);
}
__DEVICE__ unsigned int __vaddus4(unsigned int __a, unsigned int __b) {
return __nv_vaddus4(__a, __b);
}
__DEVICE__ unsigned int __vavgs2(unsigned int __a, unsigned int __b) {
return __nv_vavgs2(__a, __b);
}
__DEVICE__ unsigned int __vavgs4(unsigned int __a, unsigned int __b) {
return __nv_vavgs4(__a, __b);
}
__DEVICE__ unsigned int __vavgu2(unsigned int __a, unsigned int __b) {
return __nv_vavgu2(__a, __b);
}
__DEVICE__ unsigned int __vavgu4(unsigned int __a, unsigned int __b) {
return __nv_vavgu4(__a, __b);
}
__DEVICE__ unsigned int __vcmpeq2(unsigned int __a, unsigned int __b) {
return __nv_vcmpeq2(__a, __b);
}
__DEVICE__ unsigned int __vcmpeq4(unsigned int __a, unsigned int __b) {
return __nv_vcmpeq4(__a, __b);
}
__DEVICE__ unsigned int __vcmpges2(unsigned int __a, unsigned int __b) {
return __nv_vcmpges2(__a, __b);
}
__DEVICE__ unsigned int __vcmpges4(unsigned int __a, unsigned int __b) {
return __nv_vcmpges4(__a, __b);
}
__DEVICE__ unsigned int __vcmpgeu2(unsigned int __a, unsigned int __b) {
return __nv_vcmpgeu2(__a, __b);
}
__DEVICE__ unsigned int __vcmpgeu4(unsigned int __a, unsigned int __b) {
return __nv_vcmpgeu4(__a, __b);
}
__DEVICE__ unsigned int __vcmpgts2(unsigned int __a, unsigned int __b) {
return __nv_vcmpgts2(__a, __b);
}
__DEVICE__ unsigned int __vcmpgts4(unsigned int __a, unsigned int __b) {
return __nv_vcmpgts4(__a, __b);
}
__DEVICE__ unsigned int __vcmpgtu2(unsigned int __a, unsigned int __b) {
return __nv_vcmpgtu2(__a, __b);
}
__DEVICE__ unsigned int __vcmpgtu4(unsigned int __a, unsigned int __b) {
return __nv_vcmpgtu4(__a, __b);
}
__DEVICE__ unsigned int __vcmples2(unsigned int __a, unsigned int __b) {
return __nv_vcmples2(__a, __b);
}
__DEVICE__ unsigned int __vcmples4(unsigned int __a, unsigned int __b) {
return __nv_vcmples4(__a, __b);
}
__DEVICE__ unsigned int __vcmpleu2(unsigned int __a, unsigned int __b) {
return __nv_vcmpleu2(__a, __b);
}
__DEVICE__ unsigned int __vcmpleu4(unsigned int __a, unsigned int __b) {
return __nv_vcmpleu4(__a, __b);
}
__DEVICE__ unsigned int __vcmplts2(unsigned int __a, unsigned int __b) {
return __nv_vcmplts2(__a, __b);
}
__DEVICE__ unsigned int __vcmplts4(unsigned int __a, unsigned int __b) {
return __nv_vcmplts4(__a, __b);
}
__DEVICE__ unsigned int __vcmpltu2(unsigned int __a, unsigned int __b) {
return __nv_vcmpltu2(__a, __b);
}
__DEVICE__ unsigned int __vcmpltu4(unsigned int __a, unsigned int __b) {
return __nv_vcmpltu4(__a, __b);
}
__DEVICE__ unsigned int __vcmpne2(unsigned int __a, unsigned int __b) {
return __nv_vcmpne2(__a, __b);
}
__DEVICE__ unsigned int __vcmpne4(unsigned int __a, unsigned int __b) {
return __nv_vcmpne4(__a, __b);
}
__DEVICE__ unsigned int __vhaddu2(unsigned int __a, unsigned int __b) {
return __nv_vhaddu2(__a, __b);
}
__DEVICE__ unsigned int __vhaddu4(unsigned int __a, unsigned int __b) {
return __nv_vhaddu4(__a, __b);
}
__DEVICE__ unsigned int __vmaxs2(unsigned int __a, unsigned int __b) {
return __nv_vmaxs2(__a, __b);
}
__DEVICE__ unsigned int __vmaxs4(unsigned int __a, unsigned int __b) {
return __nv_vmaxs4(__a, __b);
}
__DEVICE__ unsigned int __vmaxu2(unsigned int __a, unsigned int __b) {
return __nv_vmaxu2(__a, __b);
}
__DEVICE__ unsigned int __vmaxu4(unsigned int __a, unsigned int __b) {
return __nv_vmaxu4(__a, __b);
}
__DEVICE__ unsigned int __vmins2(unsigned int __a, unsigned int __b) {
return __nv_vmins2(__a, __b);
}
__DEVICE__ unsigned int __vmins4(unsigned int __a, unsigned int __b) {
return __nv_vmins4(__a, __b);
}
__DEVICE__ unsigned int __vminu2(unsigned int __a, unsigned int __b) {
return __nv_vminu2(__a, __b);
}
__DEVICE__ unsigned int __vminu4(unsigned int __a, unsigned int __b) {
return __nv_vminu4(__a, __b);
}
__DEVICE__ unsigned int __vneg2(unsigned int __a) { return __nv_vneg2(__a); }
__DEVICE__ unsigned int __vneg4(unsigned int __a) { return __nv_vneg4(__a); }
__DEVICE__ unsigned int __vnegss2(unsigned int __a) {
return __nv_vnegss2(__a);
}
__DEVICE__ unsigned int __vnegss4(unsigned int __a) {
return __nv_vnegss4(__a);
}
__DEVICE__ unsigned int __vsads2(unsigned int __a, unsigned int __b) {
return __nv_vsads2(__a, __b);
}
__DEVICE__ unsigned int __vsads4(unsigned int __a, unsigned int __b) {
return __nv_vsads4(__a, __b);
}
__DEVICE__ unsigned int __vsadu2(unsigned int __a, unsigned int __b) {
return __nv_vsadu2(__a, __b);
}
__DEVICE__ unsigned int __vsadu4(unsigned int __a, unsigned int __b) {
return __nv_vsadu4(__a, __b);
}
__DEVICE__ unsigned int __vseteq2(unsigned int __a, unsigned int __b) {
return __nv_vseteq2(__a, __b);
}
__DEVICE__ unsigned int __vseteq4(unsigned int __a, unsigned int __b) {
return __nv_vseteq4(__a, __b);
}
__DEVICE__ unsigned int __vsetges2(unsigned int __a, unsigned int __b) {
return __nv_vsetges2(__a, __b);
}
__DEVICE__ unsigned int __vsetges4(unsigned int __a, unsigned int __b) {
return __nv_vsetges4(__a, __b);
}
__DEVICE__ unsigned int __vsetgeu2(unsigned int __a, unsigned int __b) {
return __nv_vsetgeu2(__a, __b);
}
__DEVICE__ unsigned int __vsetgeu4(unsigned int __a, unsigned int __b) {
return __nv_vsetgeu4(__a, __b);
}
__DEVICE__ unsigned int __vsetgts2(unsigned int __a, unsigned int __b) {
return __nv_vsetgts2(__a, __b);
}
__DEVICE__ unsigned int __vsetgts4(unsigned int __a, unsigned int __b) {
return __nv_vsetgts4(__a, __b);
}
__DEVICE__ unsigned int __vsetgtu2(unsigned int __a, unsigned int __b) {
return __nv_vsetgtu2(__a, __b);
}
__DEVICE__ unsigned int __vsetgtu4(unsigned int __a, unsigned int __b) {
return __nv_vsetgtu4(__a, __b);
}
__DEVICE__ unsigned int __vsetles2(unsigned int __a, unsigned int __b) {
return __nv_vsetles2(__a, __b);
}
__DEVICE__ unsigned int __vsetles4(unsigned int __a, unsigned int __b) {
return __nv_vsetles4(__a, __b);
}