Linux sh52.ich-4.com 5.14.0-611.26.1.el9_7.x86_64 #1 SMP PREEMPT_DYNAMIC Thu Jan 29 05:24:47 EST 2026 x86_64
LiteSpeed
Server IP : 198.143.147.58 & Your IP : 216.73.217.21
Domains :
Cant Read [ /etc/named.conf ]
User : actualbuzz
Terminal
Auto Root
Create File
Create Folder
Localroot Suggester
Backdoor Destroyer
Readme
/
lib /
clang /
20 /
include /
Delete
Unzip
Name
Size
Permission
Date
Action
cuda_wrappers
[ DIR ]
drwxr-xr-x
2026-02-03 15:09
fuzzer
[ DIR ]
drwxr-xr-x
2026-02-03 15:08
llvm_libc_wrappers
[ DIR ]
drwxr-xr-x
2026-02-03 15:09
llvm_offload_wrappers
[ DIR ]
drwxr-xr-x
2026-02-03 15:09
openmp_wrappers
[ DIR ]
drwxr-xr-x
2026-02-03 15:09
orc
[ DIR ]
drwxr-xr-x
2026-02-03 15:08
ppc_wrappers
[ DIR ]
drwxr-xr-x
2026-02-03 15:09
profile
[ DIR ]
drwxr-xr-x
2026-02-03 15:08
sanitizer
[ DIR ]
drwxr-xr-x
2026-02-03 15:08
xray
[ DIR ]
drwxr-xr-x
2026-02-03 15:08
zos_wrappers
[ DIR ]
drwxr-xr-x
2026-02-03 15:09
__clang_cuda_builtin_vars.h
4.78
KB
-rw-r--r--
2025-07-08 23:06
__clang_cuda_cmath.h
18.06
KB
-rw-r--r--
2025-07-08 23:06
__clang_cuda_complex_builtins.h
9.36
KB
-rw-r--r--
2025-07-08 23:06
__clang_cuda_device_functions.h
57.36
KB
-rw-r--r--
2025-07-08 23:06
__clang_cuda_intrinsics.h
29.87
KB
-rw-r--r--
2025-07-08 23:06
__clang_cuda_libdevice_declares.h
21.89
KB
-rw-r--r--
2025-07-08 23:06
__clang_cuda_math.h
16.21
KB
-rw-r--r--
2025-07-08 23:06
__clang_cuda_math_forward_declares.h
8.27
KB
-rw-r--r--
2025-07-08 23:06
__clang_cuda_runtime_wrapper.h
17.6
KB
-rw-r--r--
2025-07-08 23:06
__clang_cuda_texture_intrinsics.h
31.86
KB
-rw-r--r--
2025-07-08 23:06
__clang_hip_cmath.h
26.48
KB
-rw-r--r--
2025-07-08 23:06
__clang_hip_libdevice_declares.h
19.87
KB
-rw-r--r--
2025-07-08 23:06
__clang_hip_math.h
32.78
KB
-rw-r--r--
2025-07-08 23:06
__clang_hip_runtime_wrapper.h
4.94
KB
-rw-r--r--
2025-07-08 23:06
__clang_hip_stdlib.h
1.19
KB
-rw-r--r--
2025-07-08 23:06
__stdarg___gnuc_va_list.h
467
B
-rw-r--r--
2025-07-08 23:06
__stdarg___va_copy.h
445
B
-rw-r--r--
2025-07-08 23:06
__stdarg_header_macro.h
417
B
-rw-r--r--
2025-07-08 23:06
__stdarg_va_arg.h
793
B
-rw-r--r--
2025-07-08 23:06
__stdarg_va_copy.h
451
B
-rw-r--r--
2025-07-08 23:06
__stdarg_va_list.h
448
B
-rw-r--r--
2025-07-08 23:06
__stddef_header_macro.h
417
B
-rw-r--r--
2025-07-08 23:06
__stddef_max_align_t.h
857
B
-rw-r--r--
2025-07-08 23:06
__stddef_null.h
875
B
-rw-r--r--
2025-07-08 23:06
__stddef_nullptr_t.h
958
B
-rw-r--r--
2025-07-08 23:06
__stddef_offsetof.h
708
B
-rw-r--r--
2025-07-08 23:06
__stddef_ptrdiff_t.h
717
B
-rw-r--r--
2025-07-08 23:06
__stddef_rsize_t.h
710
B
-rw-r--r--
2025-07-08 23:06
__stddef_size_t.h
708
B
-rw-r--r--
2025-07-08 23:06
__stddef_unreachable.h
735
B
-rw-r--r--
2025-07-08 23:06
__stddef_wchar_t.h
853
B
-rw-r--r--
2025-07-08 23:06
__stddef_wint_t.h
443
B
-rw-r--r--
2025-07-08 23:06
__wmmintrin_aes.h
5.15
KB
-rw-r--r--
2025-07-08 23:06
__wmmintrin_pclmul.h
1.99
KB
-rw-r--r--
2025-07-08 23:06
adcintrin.h
5.43
KB
-rw-r--r--
2025-07-08 23:06
adxintrin.h
3.37
KB
-rw-r--r--
2025-07-08 23:06
altivec.h
698.45
KB
-rw-r--r--
2025-07-08 23:06
amdgpuintrin.h
7.4
KB
-rw-r--r--
2025-07-08 23:06
ammintrin.h
7.56
KB
-rw-r--r--
2025-07-08 23:06
amxavx512intrin.h
12.67
KB
-rw-r--r--
2025-07-08 23:06
amxbf16transposeintrin.h
3.44
KB
-rw-r--r--
2025-07-08 23:06
amxcomplexintrin.h
6.81
KB
-rw-r--r--
2025-07-08 23:06
amxcomplextransposeintrin.h
11.78
KB
-rw-r--r--
2025-07-08 23:06
amxfp16intrin.h
3.25
KB
-rw-r--r--
2025-07-08 23:06
amxfp16transposeintrin.h
3.44
KB
-rw-r--r--
2025-07-08 23:06
amxfp8intrin.h
8.63
KB
-rw-r--r--
2025-07-08 23:06
amxintrin.h
19.83
KB
-rw-r--r--
2025-07-08 23:06
amxmovrsintrin.h
2.24
KB
-rw-r--r--
2025-07-08 23:06
amxmovrstransposeintrin.h
9.18
KB
-rw-r--r--
2025-07-08 23:06
amxtf32intrin.h
3.48
KB
-rw-r--r--
2025-07-08 23:06
amxtf32transposeintrin.h
3.56
KB
-rw-r--r--
2025-07-08 23:06
amxtransposeintrin.h
10.6
KB
-rw-r--r--
2025-07-08 23:06
arm64intr.h
993
B
-rw-r--r--
2025-07-08 23:06
arm_acle.h
29.97
KB
-rw-r--r--
2025-07-08 23:06
arm_bf16.h
548
B
-rw-r--r--
2025-09-17 20:10
arm_cde.h
32.67
KB
-rw-r--r--
2025-09-17 20:10
arm_cmse.h
6.21
KB
-rw-r--r--
2025-07-08 23:06
arm_fp16.h
16.94
KB
-rw-r--r--
2025-09-17 20:10
arm_mve.h
1.48
MB
-rw-r--r--
2025-09-17 20:10
arm_neon.h
2.66
MB
-rw-r--r--
2025-09-17 20:10
arm_neon_sve_bridge.h
9.48
KB
-rw-r--r--
2025-07-08 23:06
arm_sme.h
223.36
KB
-rw-r--r--
2025-09-17 20:10
arm_sve.h
2.03
MB
-rw-r--r--
2025-09-17 20:10
arm_vector_types.h
9.86
KB
-rw-r--r--
2025-09-17 20:10
armintr.h
843
B
-rw-r--r--
2025-07-08 23:06
avx10_2_512bf16intrin.h
22.87
KB
-rw-r--r--
2025-07-08 23:06
avx10_2_512convertintrin.h
12.88
KB
-rw-r--r--
2025-07-08 23:06
avx10_2_512minmaxintrin.h
7.35
KB
-rw-r--r--
2025-07-08 23:06
avx10_2_512niintrin.h
13.71
KB
-rw-r--r--
2025-07-08 23:06
avx10_2_512satcvtdsintrin.h
13.48
KB
-rw-r--r--
2025-07-08 23:06
avx10_2_512satcvtintrin.h
16.6
KB
-rw-r--r--
2025-07-08 23:06
avx10_2bf16intrin.h
43.63
KB
-rw-r--r--
2025-07-08 23:06
avx10_2convertintrin.h
22.75
KB
-rw-r--r--
2025-07-08 23:06
avx10_2copyintrin.h
2.35
KB
-rw-r--r--
2025-07-08 23:06
avx10_2minmaxintrin.h
13.53
KB
-rw-r--r--
2025-07-08 23:06
avx10_2niintrin.h
16.6
KB
-rw-r--r--
2025-07-08 23:06
avx10_2satcvtdsintrin.h
15.36
KB
-rw-r--r--
2025-07-08 23:06
avx10_2satcvtintrin.h
18.87
KB
-rw-r--r--
2025-07-08 23:06
avx2intrin.h
188.12
KB
-rw-r--r--
2025-07-08 23:06
avx512bf16intrin.h
10.58
KB
-rw-r--r--
2025-07-08 23:06
avx512bitalgintrin.h
2.56
KB
-rw-r--r--
2025-07-08 23:06
avx512bwintrin.h
76.02
KB
-rw-r--r--
2025-07-08 23:06
avx512cdintrin.h
4.23
KB
-rw-r--r--
2025-07-08 23:06
avx512dqintrin.h
58.87
KB
-rw-r--r--
2025-07-08 23:06
avx512fintrin.h
383.45
KB
-rw-r--r--
2025-07-08 23:06
avx512fp16intrin.h
156.92
KB
-rw-r--r--
2025-07-08 23:06
avx512ifmaintrin.h
2.6
KB
-rw-r--r--
2025-07-08 23:06
avx512ifmavlintrin.h
4.6
KB
-rw-r--r--
2025-07-08 23:06
avx512vbmi2intrin.h
13.18
KB
-rw-r--r--
2025-07-08 23:06
avx512vbmiintrin.h
3.83
KB
-rw-r--r--
2025-07-08 23:06
avx512vbmivlintrin.h
7.23
KB
-rw-r--r--
2025-07-08 23:06
avx512vlbf16intrin.h
19.46
KB
-rw-r--r--
2025-07-08 23:06
avx512vlbitalgintrin.h
4.52
KB
-rw-r--r--
2025-07-08 23:06
avx512vlbwintrin.h
121.56
KB
-rw-r--r--
2025-07-08 23:06
avx512vlcdintrin.h
7.95
KB
-rw-r--r--
2025-07-08 23:06
avx512vldqintrin.h
46.71
KB
-rw-r--r--
2025-07-08 23:06
avx512vlfp16intrin.h
85.51
KB
-rw-r--r--
2025-07-08 23:06
avx512vlintrin.h
322.6
KB
-rw-r--r--
2025-07-08 23:06
avx512vlvbmi2intrin.h
26.01
KB
-rw-r--r--
2025-07-08 23:06
avx512vlvnniintrin.h
13.41
KB
-rw-r--r--
2025-07-08 23:06
avx512vlvp2intersectintrin.h
4.67
KB
-rw-r--r--
2025-07-08 23:06
avx512vnniintrin.h
4.32
KB
-rw-r--r--
2025-07-08 23:06
avx512vp2intersectintrin.h
3.03
KB
-rw-r--r--
2025-07-08 23:06
avx512vpopcntdqintrin.h
2.3
KB
-rw-r--r--
2025-07-08 23:06
avx512vpopcntdqvlintrin.h
3.86
KB
-rw-r--r--
2025-07-08 23:06
avxifmaintrin.h
5.75
KB
-rw-r--r--
2025-07-08 23:06
avxintrin.h
197.12
KB
-rw-r--r--
2025-07-08 23:06
avxneconvertintrin.h
14.09
KB
-rw-r--r--
2025-07-08 23:06
avxvnniint16intrin.h
15.84
KB
-rw-r--r--
2025-07-08 23:06
avxvnniint8intrin.h
17.02
KB
-rw-r--r--
2025-07-08 23:06
avxvnniintrin.h
10.44
KB
-rw-r--r--
2025-07-08 23:06
bmi2intrin.h
7.37
KB
-rw-r--r--
2025-07-08 23:06
bmiintrin.h
20.02
KB
-rw-r--r--
2025-07-08 23:06
builtins.h
846
B
-rw-r--r--
2025-07-08 23:06
cet.h
1.49
KB
-rw-r--r--
2025-07-08 23:06
cetintrin.h
3.27
KB
-rw-r--r--
2025-07-08 23:06
cldemoteintrin.h
1.18
KB
-rw-r--r--
2025-07-08 23:06
clflushoptintrin.h
1.17
KB
-rw-r--r--
2025-07-08 23:06
clwbintrin.h
1.2
KB
-rw-r--r--
2025-07-08 23:06
clzerointrin.h
1.19
KB
-rw-r--r--
2025-07-08 23:06
cmpccxaddintrin.h
2.33
KB
-rw-r--r--
2025-07-08 23:06
cpuid.h
11.9
KB
-rw-r--r--
2025-07-08 23:06
crc32intrin.h
3.27
KB
-rw-r--r--
2025-07-08 23:06
emmintrin.h
196.89
KB
-rw-r--r--
2025-07-08 23:06
enqcmdintrin.h
2.12
KB
-rw-r--r--
2025-07-08 23:06
f16cintrin.h
5.39
KB
-rw-r--r--
2025-07-08 23:06
float.h
6.43
KB
-rw-r--r--
2025-07-08 23:06
fma4intrin.h
6.82
KB
-rw-r--r--
2025-07-08 23:06
fmaintrin.h
28.65
KB
-rw-r--r--
2025-07-08 23:06
fxsrintrin.h
2.82
KB
-rw-r--r--
2025-07-08 23:06
gfniintrin.h
9.59
KB
-rw-r--r--
2025-07-08 23:06
gpuintrin.h
9.07
KB
-rw-r--r--
2025-07-08 23:06
hexagon_circ_brev_intrinsics.h
15.59
KB
-rw-r--r--
2025-07-08 23:06
hexagon_protos.h
374.42
KB
-rw-r--r--
2025-07-08 23:06
hexagon_types.h
130.38
KB
-rw-r--r--
2025-07-08 23:06
hresetintrin.h
1.36
KB
-rw-r--r--
2025-07-08 23:06
htmintrin.h
6.14
KB
-rw-r--r--
2025-07-08 23:06
htmxlintrin.h
9.01
KB
-rw-r--r--
2025-07-08 23:06
hvx_hexagon_protos.h
275.18
KB
-rw-r--r--
2025-07-08 23:06
ia32intrin.h
25.38
KB
-rw-r--r--
2025-07-08 23:06
immintrin.h
25.46
KB
-rw-r--r--
2025-07-08 23:06
intrin.h
18.16
KB
-rw-r--r--
2025-07-08 23:06
intrin0.h
13.76
KB
-rw-r--r--
2025-07-08 23:06
inttypes.h
2.37
KB
-rw-r--r--
2025-07-08 23:06
invpcidintrin.h
764
B
-rw-r--r--
2025-07-08 23:06
iso646.h
763
B
-rw-r--r--
2025-07-08 23:06
keylockerintrin.h
17.85
KB
-rw-r--r--
2025-07-08 23:06
larchintrin.h
8.5
KB
-rw-r--r--
2025-07-08 23:06
lasxintrin.h
142.14
KB
-rw-r--r--
2025-07-08 23:06
limits.h
3.71
KB
-rw-r--r--
2025-07-08 23:06
lsxintrin.h
134.7
KB
-rw-r--r--
2025-07-08 23:06
lwpintrin.h
5
KB
-rw-r--r--
2025-07-08 23:06
lzcntintrin.h
3.45
KB
-rw-r--r--
2025-07-08 23:06
mm3dnow.h
729
B
-rw-r--r--
2025-07-08 23:06
mm_malloc.h
1.88
KB
-rw-r--r--
2025-07-08 23:06
mmintrin.h
58.6
KB
-rw-r--r--
2025-07-08 23:06
module.modulemap
6.47
KB
-rw-r--r--
2025-07-08 23:06
movdirintrin.h
1.57
KB
-rw-r--r--
2025-07-08 23:06
movrs_avx10_2_512intrin.h
3.88
KB
-rw-r--r--
2025-07-08 23:06
movrs_avx10_2intrin.h
6.87
KB
-rw-r--r--
2025-07-08 23:06
movrsintrin.h
1.98
KB
-rw-r--r--
2025-07-08 23:06
msa.h
25.01
KB
-rw-r--r--
2025-07-08 23:06
mwaitxintrin.h
2.19
KB
-rw-r--r--
2025-07-08 23:06
nmmintrin.h
709
B
-rw-r--r--
2025-07-08 23:06
nvptxintrin.h
7.42
KB
-rw-r--r--
2025-07-08 23:06
omp-tools.h
50.86
KB
-rw-r--r--
2025-09-17 20:37
omp.h
23.36
KB
-rw-r--r--
2025-09-17 20:37
ompt-multiplex.h
51.11
KB
-rw-r--r--
2025-07-08 23:06
ompt.h
50.86
KB
-rw-r--r--
2025-09-17 20:37
ompx.h
7.38
KB
-rw-r--r--
2025-09-17 20:37
opencl-c-base.h
30.94
KB
-rw-r--r--
2025-07-08 23:06
opencl-c.h
874.98
KB
-rw-r--r--
2025-07-08 23:06
pconfigintrin.h
1.19
KB
-rw-r--r--
2025-07-08 23:06
pkuintrin.h
934
B
-rw-r--r--
2025-07-08 23:06
pmmintrin.h
11.12
KB
-rw-r--r--
2025-07-08 23:06
popcntintrin.h
1.9
KB
-rw-r--r--
2025-07-08 23:06
prfchiintrin.h
2.02
KB
-rw-r--r--
2025-07-08 23:06
prfchwintrin.h
2.06
KB
-rw-r--r--
2025-07-08 23:06
ptrauth.h
15.61
KB
-rw-r--r--
2025-07-08 23:06
ptwriteintrin.h
1.05
KB
-rw-r--r--
2025-07-08 23:06
raointintrin.h
6.59
KB
-rw-r--r--
2025-07-08 23:06
rdpruintrin.h
1.59
KB
-rw-r--r--
2025-07-08 23:06
rdseedintrin.h
2.85
KB
-rw-r--r--
2025-07-08 23:06
riscv_bitmanip.h
5.59
KB
-rw-r--r--
2025-07-08 23:06
riscv_corev_alu.h
3.94
KB
-rw-r--r--
2025-07-08 23:06
riscv_crypto.h
5.09
KB
-rw-r--r--
2025-07-08 23:06
riscv_ntlh.h
744
B
-rw-r--r--
2025-07-08 23:06
riscv_vector.h
17.11
KB
-rw-r--r--
2025-09-17 20:11
rtmintrin.h
1.25
KB
-rw-r--r--
2025-07-08 23:06
s390intrin.h
604
B
-rw-r--r--
2025-07-08 23:06
serializeintrin.h
881
B
-rw-r--r--
2025-07-08 23:06
sgxintrin.h
1.77
KB
-rw-r--r--
2025-07-08 23:06
sha512intrin.h
5.95
KB
-rw-r--r--
2025-07-08 23:06
shaintrin.h
7.37
KB
-rw-r--r--
2025-07-08 23:06
sifive_vector.h
6.88
KB
-rw-r--r--
2025-07-08 23:06
sm3intrin.h
7.29
KB
-rw-r--r--
2025-07-08 23:06
sm4evexintrin.h
1.17
KB
-rw-r--r--
2025-07-08 23:06
sm4intrin.h
8.2
KB
-rw-r--r--
2025-07-08 23:06
smmintrin.h
99.34
KB
-rw-r--r--
2025-07-08 23:06
stdalign.h
756
B
-rw-r--r--
2025-07-08 23:06
stdarg.h
2.44
KB
-rw-r--r--
2025-07-08 23:06
stdatomic.h
8.46
KB
-rw-r--r--
2025-07-08 23:06
stdbool.h
1.14
KB
-rw-r--r--
2025-07-08 23:06
stdckdint.h
1.63
KB
-rw-r--r--
2025-07-08 23:06
stddef.h
4.93
KB
-rw-r--r--
2025-07-08 23:06
stdint.h
30.33
KB
-rw-r--r--
2025-07-08 23:06
stdnoreturn.h
1.29
KB
-rw-r--r--
2025-07-08 23:06
tbmintrin.h
3.42
KB
-rw-r--r--
2025-07-08 23:06
tgmath.h
29.68
KB
-rw-r--r--
2025-07-08 23:06
tmmintrin.h
31.18
KB
-rw-r--r--
2025-07-08 23:06
tsxldtrkintrin.h
1.97
KB
-rw-r--r--
2025-07-08 23:06
uintrintrin.h
4.96
KB
-rw-r--r--
2025-07-08 23:06
unwind.h
11.21
KB
-rw-r--r--
2025-07-08 23:06
usermsrintrin.h
1.54
KB
-rw-r--r--
2025-07-08 23:06
vadefs.h
1.39
KB
-rw-r--r--
2025-07-08 23:06
vaesintrin.h
2.61
KB
-rw-r--r--
2025-07-08 23:06
varargs.h
584
B
-rw-r--r--
2025-07-08 23:06
vecintrin.h
419.23
KB
-rw-r--r--
2025-07-08 23:06
velintrin.h
2.1
KB
-rw-r--r--
2025-07-08 23:06
velintrin_approx.h
3.54
KB
-rw-r--r--
2025-07-08 23:06
velintrin_gen.h
69.06
KB
-rw-r--r--
2025-07-08 23:06
vpclmulqdqintrin.h
1.06
KB
-rw-r--r--
2025-07-08 23:06
waitpkgintrin.h
1.33
KB
-rw-r--r--
2025-07-08 23:06
wasm_simd128.h
81.71
KB
-rw-r--r--
2025-07-08 23:06
wbnoinvdintrin.h
749
B
-rw-r--r--
2025-07-08 23:06
wmmintrin.h
659
B
-rw-r--r--
2025-07-08 23:06
x86gprintrin.h
2.22
KB
-rw-r--r--
2025-07-08 23:06
x86intrin.h
1.38
KB
-rw-r--r--
2025-07-08 23:06
xmmintrin.h
117.43
KB
-rw-r--r--
2025-07-08 23:06
xopintrin.h
19.96
KB
-rw-r--r--
2025-07-08 23:06
xsavecintrin.h
2.51
KB
-rw-r--r--
2025-07-08 23:06
xsaveintrin.h
1.64
KB
-rw-r--r--
2025-07-08 23:06
xsaveoptintrin.h
1
KB
-rw-r--r--
2025-07-08 23:06
xsavesintrin.h
1.24
KB
-rw-r--r--
2025-07-08 23:06
xtestintrin.h
873
B
-rw-r--r--
2025-07-08 23:06
yvals_core.h
687
B
-rw-r--r--
2025-07-08 23:06
Save
Rename
/*===---- __clang_hip_math.h - Device-side HIP math 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_HIP_MATH_H__ #define __CLANG_HIP_MATH_H__ #if !defined(__HIP__) && !defined(__OPENMP_AMDGCN__) #error "This file is for HIP and OpenMP AMDGCN device compilation only." #endif // The __CLANG_GPU_DISABLE_MATH_WRAPPERS macro provides a way to let standard // libcalls reach the link step instead of being eagerly replaced. #ifndef __CLANG_GPU_DISABLE_MATH_WRAPPERS #if !defined(__HIPCC_RTC__) #include <limits.h> #include <stdint.h> #ifdef __OPENMP_AMDGCN__ #include <omp.h> #endif #endif // !defined(__HIPCC_RTC__) #pragma push_macro("__DEVICE__") #ifdef __OPENMP_AMDGCN__ #define __DEVICE__ static inline __attribute__((always_inline, nothrow)) #else #define __DEVICE__ static __device__ inline __attribute__((always_inline)) #endif // Device library provides fast low precision and slow full-recision // implementations for some functions. Which one gets selected depends on // __CLANG_GPU_APPROX_TRANSCENDENTALS__ which gets defined by clang if // -ffast-math or -fgpu-approx-transcendentals are in effect. #pragma push_macro("__FAST_OR_SLOW") #if defined(__CLANG_GPU_APPROX_TRANSCENDENTALS__) #define __FAST_OR_SLOW(fast, slow) fast #else #define __FAST_OR_SLOW(fast, slow) slow #endif // A few functions return bool type starting only in C++11. #pragma push_macro("__RETURN_TYPE") #ifdef __OPENMP_AMDGCN__ #define __RETURN_TYPE int #else #if defined(__cplusplus) #define __RETURN_TYPE bool #else #define __RETURN_TYPE int #endif #endif // __OPENMP_AMDGCN__ #if defined (__cplusplus) && __cplusplus < 201103L // emulate static_assert on type sizes template<bool> struct __compare_result{}; template<> struct __compare_result<true> { static const __device__ bool valid; }; __DEVICE__ void __suppress_unused_warning(bool b){}; template <unsigned int S, unsigned int T> __DEVICE__ void __static_assert_equal_size() { __suppress_unused_warning(__compare_result<S == T>::valid); } #define __static_assert_type_size_equal(A, B) \ __static_assert_equal_size<A,B>() #else #define __static_assert_type_size_equal(A,B) \ static_assert((A) == (B), "") #endif __DEVICE__ uint64_t __make_mantissa_base8(const char *__tagp __attribute__((nonnull))) { uint64_t __r = 0; while (*__tagp != '\0') { char __tmp = *__tagp; if (__tmp >= '0' && __tmp <= '7') __r = (__r * 8u) + __tmp - '0'; else return 0; ++__tagp; } return __r; } __DEVICE__ uint64_t __make_mantissa_base10(const char *__tagp __attribute__((nonnull))) { uint64_t __r = 0; while (*__tagp != '\0') { char __tmp = *__tagp; if (__tmp >= '0' && __tmp <= '9') __r = (__r * 10u) + __tmp - '0'; else return 0; ++__tagp; } return __r; } __DEVICE__ uint64_t __make_mantissa_base16(const char *__tagp __attribute__((nonnull))) { uint64_t __r = 0; while (*__tagp != '\0') { char __tmp = *__tagp; if (__tmp >= '0' && __tmp <= '9') __r = (__r * 16u) + __tmp - '0'; else if (__tmp >= 'a' && __tmp <= 'f') __r = (__r * 16u) + __tmp - 'a' + 10; else if (__tmp >= 'A' && __tmp <= 'F') __r = (__r * 16u) + __tmp - 'A' + 10; else return 0; ++__tagp; } return __r; } __DEVICE__ uint64_t __make_mantissa(const char *__tagp __attribute__((nonnull))) { if (*__tagp == '0') { ++__tagp; if (*__tagp == 'x' || *__tagp == 'X') return __make_mantissa_base16(__tagp); else return __make_mantissa_base8(__tagp); } return __make_mantissa_base10(__tagp); } // BEGIN FLOAT // BEGIN INTRINSICS __DEVICE__ float __cosf(float __x) { return __ocml_native_cos_f32(__x); } __DEVICE__ float __exp10f(float __x) { const float __log2_10 = 0x1.a934f0p+1f; return __builtin_amdgcn_exp2f(__log2_10 * __x); } __DEVICE__ float __expf(float __x) { const float __log2_e = 0x1.715476p+0; return __builtin_amdgcn_exp2f(__log2_e * __x); } #if defined OCML_BASIC_ROUNDED_OPERATIONS __DEVICE__ float __fadd_rd(float __x, float __y) { return __ocml_add_rtn_f32(__x, __y); } __DEVICE__ float __fadd_rn(float __x, float __y) { return __ocml_add_rte_f32(__x, __y); } __DEVICE__ float __fadd_ru(float __x, float __y) { return __ocml_add_rtp_f32(__x, __y); } __DEVICE__ float __fadd_rz(float __x, float __y) { return __ocml_add_rtz_f32(__x, __y); } #else __DEVICE__ float __fadd_rn(float __x, float __y) { return __x + __y; } #endif #if defined OCML_BASIC_ROUNDED_OPERATIONS __DEVICE__ float __fdiv_rd(float __x, float __y) { return __ocml_div_rtn_f32(__x, __y); } __DEVICE__ float __fdiv_rn(float __x, float __y) { return __ocml_div_rte_f32(__x, __y); } __DEVICE__ float __fdiv_ru(float __x, float __y) { return __ocml_div_rtp_f32(__x, __y); } __DEVICE__ float __fdiv_rz(float __x, float __y) { return __ocml_div_rtz_f32(__x, __y); } #else __DEVICE__ float __fdiv_rn(float __x, float __y) { return __x / __y; } #endif __DEVICE__ float __fdividef(float __x, float __y) { return __x / __y; } #if defined OCML_BASIC_ROUNDED_OPERATIONS __DEVICE__ float __fmaf_rd(float __x, float __y, float __z) { return __ocml_fma_rtn_f32(__x, __y, __z); } __DEVICE__ float __fmaf_rn(float __x, float __y, float __z) { return __ocml_fma_rte_f32(__x, __y, __z); } __DEVICE__ float __fmaf_ru(float __x, float __y, float __z) { return __ocml_fma_rtp_f32(__x, __y, __z); } __DEVICE__ float __fmaf_rz(float __x, float __y, float __z) { return __ocml_fma_rtz_f32(__x, __y, __z); } #else __DEVICE__ float __fmaf_rn(float __x, float __y, float __z) { return __builtin_fmaf(__x, __y, __z); } #endif #if defined OCML_BASIC_ROUNDED_OPERATIONS __DEVICE__ float __fmul_rd(float __x, float __y) { return __ocml_mul_rtn_f32(__x, __y); } __DEVICE__ float __fmul_rn(float __x, float __y) { return __ocml_mul_rte_f32(__x, __y); } __DEVICE__ float __fmul_ru(float __x, float __y) { return __ocml_mul_rtp_f32(__x, __y); } __DEVICE__ float __fmul_rz(float __x, float __y) { return __ocml_mul_rtz_f32(__x, __y); } #else __DEVICE__ float __fmul_rn(float __x, float __y) { return __x * __y; } #endif #if defined OCML_BASIC_ROUNDED_OPERATIONS __DEVICE__ float __frcp_rd(float __x) { return __ocml_div_rtn_f32(1.0f, __x); } __DEVICE__ float __frcp_rn(float __x) { return __ocml_div_rte_f32(1.0f, __x); } __DEVICE__ float __frcp_ru(float __x) { return __ocml_div_rtp_f32(1.0f, __x); } __DEVICE__ float __frcp_rz(float __x) { return __ocml_div_rtz_f32(1.0f, __x); } #else __DEVICE__ float __frcp_rn(float __x) { return 1.0f / __x; } #endif __DEVICE__ float __frsqrt_rn(float __x) { return __builtin_amdgcn_rsqf(__x); } #if defined OCML_BASIC_ROUNDED_OPERATIONS __DEVICE__ float __fsqrt_rd(float __x) { return __ocml_sqrt_rtn_f32(__x); } __DEVICE__ float __fsqrt_rn(float __x) { return __ocml_sqrt_rte_f32(__x); } __DEVICE__ float __fsqrt_ru(float __x) { return __ocml_sqrt_rtp_f32(__x); } __DEVICE__ float __fsqrt_rz(float __x) { return __ocml_sqrt_rtz_f32(__x); } #else __DEVICE__ float __fsqrt_rn(float __x) { return __ocml_native_sqrt_f32(__x); } #endif #if defined OCML_BASIC_ROUNDED_OPERATIONS __DEVICE__ float __fsub_rd(float __x, float __y) { return __ocml_sub_rtn_f32(__x, __y); } __DEVICE__ float __fsub_rn(float __x, float __y) { return __ocml_sub_rte_f32(__x, __y); } __DEVICE__ float __fsub_ru(float __x, float __y) { return __ocml_sub_rtp_f32(__x, __y); } __DEVICE__ float __fsub_rz(float __x, float __y) { return __ocml_sub_rtz_f32(__x, __y); } #else __DEVICE__ float __fsub_rn(float __x, float __y) { return __x - __y; } #endif __DEVICE__ float __log10f(float __x) { return __builtin_log10f(__x); } __DEVICE__ float __log2f(float __x) { return __builtin_amdgcn_logf(__x); } __DEVICE__ float __logf(float __x) { return __builtin_logf(__x); } __DEVICE__ float __powf(float __x, float __y) { return __ocml_pow_f32(__x, __y); } __DEVICE__ float __saturatef(float __x) { return (__x < 0) ? 0 : ((__x > 1) ? 1 : __x); } __DEVICE__ void __sincosf(float __x, float *__sinptr, float *__cosptr) { *__sinptr = __ocml_native_sin_f32(__x); *__cosptr = __ocml_native_cos_f32(__x); } __DEVICE__ float __sinf(float __x) { return __ocml_native_sin_f32(__x); } __DEVICE__ float __tanf(float __x) { return __sinf(__x) * __builtin_amdgcn_rcpf(__cosf(__x)); } // END INTRINSICS #if defined(__cplusplus) __DEVICE__ int abs(int __x) { return __builtin_abs(__x); } __DEVICE__ long labs(long __x) { return __builtin_labs(__x); } __DEVICE__ long long llabs(long long __x) { return __builtin_llabs(__x); } #endif __DEVICE__ float acosf(float __x) { return __ocml_acos_f32(__x); } __DEVICE__ float acoshf(float __x) { return __ocml_acosh_f32(__x); } __DEVICE__ float asinf(float __x) { return __ocml_asin_f32(__x); } __DEVICE__ float asinhf(float __x) { return __ocml_asinh_f32(__x); } __DEVICE__ float atan2f(float __x, float __y) { return __ocml_atan2_f32(__x, __y); } __DEVICE__ float atanf(float __x) { return __ocml_atan_f32(__x); } __DEVICE__ float atanhf(float __x) { return __ocml_atanh_f32(__x); } __DEVICE__ float cbrtf(float __x) { return __ocml_cbrt_f32(__x); } __DEVICE__ float ceilf(float __x) { return __builtin_ceilf(__x); } __DEVICE__ float copysignf(float __x, float __y) { return __builtin_copysignf(__x, __y); } __DEVICE__ float cosf(float __x) { return __FAST_OR_SLOW(__cosf, __ocml_cos_f32)(__x); } __DEVICE__ float coshf(float __x) { return __ocml_cosh_f32(__x); } __DEVICE__ float cospif(float __x) { return __ocml_cospi_f32(__x); } __DEVICE__ float cyl_bessel_i0f(float __x) { return __ocml_i0_f32(__x); } __DEVICE__ float cyl_bessel_i1f(float __x) { return __ocml_i1_f32(__x); } __DEVICE__ float erfcf(float __x) { return __ocml_erfc_f32(__x); } __DEVICE__ float erfcinvf(float __x) { return __ocml_erfcinv_f32(__x); } __DEVICE__ float erfcxf(float __x) { return __ocml_erfcx_f32(__x); } __DEVICE__ float erff(float __x) { return __ocml_erf_f32(__x); } __DEVICE__ float erfinvf(float __x) { return __ocml_erfinv_f32(__x); } __DEVICE__ float exp10f(float __x) { return __ocml_exp10_f32(__x); } __DEVICE__ float exp2f(float __x) { return __builtin_exp2f(__x); } __DEVICE__ float expf(float __x) { return __builtin_expf(__x); } __DEVICE__ float expm1f(float __x) { return __ocml_expm1_f32(__x); } __DEVICE__ float fabsf(float __x) { return __builtin_fabsf(__x); } __DEVICE__ float fdimf(float __x, float __y) { return __ocml_fdim_f32(__x, __y); } __DEVICE__ float fdividef(float __x, float __y) { return __x / __y; } __DEVICE__ float floorf(float __x) { return __builtin_floorf(__x); } __DEVICE__ float fmaf(float __x, float __y, float __z) { return __builtin_fmaf(__x, __y, __z); } __DEVICE__ float fmaxf(float __x, float __y) { return __builtin_fmaxf(__x, __y); } __DEVICE__ float fminf(float __x, float __y) { return __builtin_fminf(__x, __y); } __DEVICE__ float fmodf(float __x, float __y) { return __ocml_fmod_f32(__x, __y); } __DEVICE__ float frexpf(float __x, int *__nptr) { return __builtin_frexpf(__x, __nptr); } __DEVICE__ float hypotf(float __x, float __y) { return __ocml_hypot_f32(__x, __y); } __DEVICE__ int ilogbf(float __x) { return __ocml_ilogb_f32(__x); } __DEVICE__ __RETURN_TYPE __finitef(float __x) { return __builtin_isfinite(__x); } __DEVICE__ __RETURN_TYPE __isinff(float __x) { return __builtin_isinf(__x); } __DEVICE__ __RETURN_TYPE __isnanf(float __x) { return __builtin_isnan(__x); } __DEVICE__ float j0f(float __x) { return __ocml_j0_f32(__x); } __DEVICE__ float j1f(float __x) { return __ocml_j1_f32(__x); } __DEVICE__ float jnf(int __n, float __x) { // TODO: we could use Ahmes multiplication // and the Miller & Brown algorithm // for linear recurrences to get O(log n) steps, but it's unclear if // it'd be beneficial in this case. if (__n == 0) return j0f(__x); if (__n == 1) return j1f(__x); float __x0 = j0f(__x); float __x1 = j1f(__x); for (int __i = 1; __i < __n; ++__i) { float __x2 = (2 * __i) / __x * __x1 - __x0; __x0 = __x1; __x1 = __x2; } return __x1; } __DEVICE__ float ldexpf(float __x, int __e) { return __builtin_amdgcn_ldexpf(__x, __e); } __DEVICE__ float lgammaf(float __x) { return __ocml_lgamma_f32(__x); } __DEVICE__ long long int llrintf(float __x) { return __builtin_rintf(__x); } __DEVICE__ long long int llroundf(float __x) { return __builtin_roundf(__x); } __DEVICE__ float log10f(float __x) { return __builtin_log10f(__x); } __DEVICE__ float log1pf(float __x) { return __ocml_log1p_f32(__x); } __DEVICE__ float log2f(float __x) { return __FAST_OR_SLOW(__log2f, __ocml_log2_f32)(__x); } __DEVICE__ float logbf(float __x) { return __ocml_logb_f32(__x); } __DEVICE__ float logf(float __x) { return __FAST_OR_SLOW(__logf, __ocml_log_f32)(__x); } __DEVICE__ long int lrintf(float __x) { return __builtin_rintf(__x); } __DEVICE__ long int lroundf(float __x) { return __builtin_roundf(__x); } __DEVICE__ float modff(float __x, float *__iptr) { float __tmp; #ifdef __OPENMP_AMDGCN__ #pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc) #endif float __r = __ocml_modf_f32(__x, (__attribute__((address_space(5))) float *)&__tmp); *__iptr = __tmp; return __r; } __DEVICE__ float nanf(const char *__tagp __attribute__((nonnull))) { union { float val; struct ieee_float { unsigned int mantissa : 22; unsigned int quiet : 1; unsigned int exponent : 8; unsigned int sign : 1; } bits; } __tmp; __static_assert_type_size_equal(sizeof(__tmp.val), sizeof(__tmp.bits)); __tmp.bits.sign = 0u; __tmp.bits.exponent = ~0u; __tmp.bits.quiet = 1u; __tmp.bits.mantissa = __make_mantissa(__tagp); return __tmp.val; } __DEVICE__ float nearbyintf(float __x) { return __builtin_nearbyintf(__x); } __DEVICE__ float nextafterf(float __x, float __y) { return __ocml_nextafter_f32(__x, __y); } __DEVICE__ float norm3df(float __x, float __y, float __z) { return __ocml_len3_f32(__x, __y, __z); } __DEVICE__ float norm4df(float __x, float __y, float __z, float __w) { return __ocml_len4_f32(__x, __y, __z, __w); } __DEVICE__ float normcdff(float __x) { return __ocml_ncdf_f32(__x); } __DEVICE__ float normcdfinvf(float __x) { return __ocml_ncdfinv_f32(__x); } __DEVICE__ float normf(int __dim, const float *__a) { // TODO: placeholder until OCML adds support. float __r = 0; while (__dim--) { __r += __a[0] * __a[0]; ++__a; } return __builtin_sqrtf(__r); } __DEVICE__ float powf(float __x, float __y) { return __ocml_pow_f32(__x, __y); } __DEVICE__ float powif(float __x, int __y) { return __ocml_pown_f32(__x, __y); } __DEVICE__ float rcbrtf(float __x) { return __ocml_rcbrt_f32(__x); } __DEVICE__ float remainderf(float __x, float __y) { return __ocml_remainder_f32(__x, __y); } __DEVICE__ float remquof(float __x, float __y, int *__quo) { int __tmp; #ifdef __OPENMP_AMDGCN__ #pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc) #endif float __r = __ocml_remquo_f32( __x, __y, (__attribute__((address_space(5))) int *)&__tmp); *__quo = __tmp; return __r; } __DEVICE__ float rhypotf(float __x, float __y) { return __ocml_rhypot_f32(__x, __y); } __DEVICE__ float rintf(float __x) { return __builtin_rintf(__x); } __DEVICE__ float rnorm3df(float __x, float __y, float __z) { return __ocml_rlen3_f32(__x, __y, __z); } __DEVICE__ float rnorm4df(float __x, float __y, float __z, float __w) { return __ocml_rlen4_f32(__x, __y, __z, __w); } __DEVICE__ float rnormf(int __dim, const float *__a) { // TODO: placeholder until OCML adds support. float __r = 0; while (__dim--) { __r += __a[0] * __a[0]; ++__a; } return __ocml_rsqrt_f32(__r); } __DEVICE__ float roundf(float __x) { return __builtin_roundf(__x); } __DEVICE__ float rsqrtf(float __x) { return __ocml_rsqrt_f32(__x); } __DEVICE__ float scalblnf(float __x, long int __n) { return (__n < INT_MAX) ? __builtin_amdgcn_ldexpf(__x, __n) : __ocml_scalb_f32(__x, __n); } __DEVICE__ float scalbnf(float __x, int __n) { return __builtin_amdgcn_ldexpf(__x, __n); } __DEVICE__ __RETURN_TYPE __signbitf(float __x) { return __builtin_signbitf(__x); } __DEVICE__ void sincosf(float __x, float *__sinptr, float *__cosptr) { float __tmp; #ifdef __OPENMP_AMDGCN__ #pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc) #endif #ifdef __CLANG_CUDA_APPROX_TRANSCENDENTALS__ __sincosf(__x, __sinptr, __cosptr); #else *__sinptr = __ocml_sincos_f32(__x, (__attribute__((address_space(5))) float *)&__tmp); *__cosptr = __tmp; #endif } __DEVICE__ void sincospif(float __x, float *__sinptr, float *__cosptr) { float __tmp; #ifdef __OPENMP_AMDGCN__ #pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc) #endif *__sinptr = __ocml_sincospi_f32( __x, (__attribute__((address_space(5))) float *)&__tmp); *__cosptr = __tmp; } __DEVICE__ float sinf(float __x) { return __FAST_OR_SLOW(__sinf, __ocml_sin_f32)(__x); } __DEVICE__ float sinhf(float __x) { return __ocml_sinh_f32(__x); } __DEVICE__ float sinpif(float __x) { return __ocml_sinpi_f32(__x); } __DEVICE__ float sqrtf(float __x) { return __builtin_sqrtf(__x); } __DEVICE__ float tanf(float __x) { return __ocml_tan_f32(__x); } __DEVICE__ float tanhf(float __x) { return __ocml_tanh_f32(__x); } __DEVICE__ float tgammaf(float __x) { return __ocml_tgamma_f32(__x); } __DEVICE__ float truncf(float __x) { return __builtin_truncf(__x); } __DEVICE__ float y0f(float __x) { return __ocml_y0_f32(__x); } __DEVICE__ float y1f(float __x) { return __ocml_y1_f32(__x); } __DEVICE__ float ynf(int __n, float __x) { // TODO: we could use Ahmes multiplication // and the Miller & Brown algorithm // for linear recurrences to get O(log n) steps, but it's unclear if // it'd be beneficial in this case. Placeholder until OCML adds // support. if (__n == 0) return y0f(__x); if (__n == 1) return y1f(__x); float __x0 = y0f(__x); float __x1 = y1f(__x); for (int __i = 1; __i < __n; ++__i) { float __x2 = (2 * __i) / __x * __x1 - __x0; __x0 = __x1; __x1 = __x2; } return __x1; } // END FLOAT // BEGIN DOUBLE __DEVICE__ double acos(double __x) { return __ocml_acos_f64(__x); } __DEVICE__ double acosh(double __x) { return __ocml_acosh_f64(__x); } __DEVICE__ double asin(double __x) { return __ocml_asin_f64(__x); } __DEVICE__ double asinh(double __x) { return __ocml_asinh_f64(__x); } __DEVICE__ double atan(double __x) { return __ocml_atan_f64(__x); } __DEVICE__ double atan2(double __x, double __y) { return __ocml_atan2_f64(__x, __y); } __DEVICE__ double atanh(double __x) { return __ocml_atanh_f64(__x); } __DEVICE__ double cbrt(double __x) { return __ocml_cbrt_f64(__x); } __DEVICE__ double ceil(double __x) { return __builtin_ceil(__x); } __DEVICE__ double copysign(double __x, double __y) { return __builtin_copysign(__x, __y); } __DEVICE__ double cos(double __x) { return __ocml_cos_f64(__x); } __DEVICE__ double cosh(double __x) { return __ocml_cosh_f64(__x); } __DEVICE__ double cospi(double __x) { return __ocml_cospi_f64(__x); } __DEVICE__ double cyl_bessel_i0(double __x) { return __ocml_i0_f64(__x); } __DEVICE__ double cyl_bessel_i1(double __x) { return __ocml_i1_f64(__x); } __DEVICE__ double erf(double __x) { return __ocml_erf_f64(__x); } __DEVICE__ double erfc(double __x) { return __ocml_erfc_f64(__x); } __DEVICE__ double erfcinv(double __x) { return __ocml_erfcinv_f64(__x); } __DEVICE__ double erfcx(double __x) { return __ocml_erfcx_f64(__x); } __DEVICE__ double erfinv(double __x) { return __ocml_erfinv_f64(__x); } __DEVICE__ double exp(double __x) { return __ocml_exp_f64(__x); } __DEVICE__ double exp10(double __x) { return __ocml_exp10_f64(__x); } __DEVICE__ double exp2(double __x) { return __ocml_exp2_f64(__x); } __DEVICE__ double expm1(double __x) { return __ocml_expm1_f64(__x); } __DEVICE__ double fabs(double __x) { return __builtin_fabs(__x); } __DEVICE__ double fdim(double __x, double __y) { return __ocml_fdim_f64(__x, __y); } __DEVICE__ double floor(double __x) { return __builtin_floor(__x); } __DEVICE__ double fma(double __x, double __y, double __z) { return __builtin_fma(__x, __y, __z); } __DEVICE__ double fmax(double __x, double __y) { return __builtin_fmax(__x, __y); } __DEVICE__ double fmin(double __x, double __y) { return __builtin_fmin(__x, __y); } __DEVICE__ double fmod(double __x, double __y) { return __ocml_fmod_f64(__x, __y); } __DEVICE__ double frexp(double __x, int *__nptr) { return __builtin_frexp(__x, __nptr); } __DEVICE__ double hypot(double __x, double __y) { return __ocml_hypot_f64(__x, __y); } __DEVICE__ int ilogb(double __x) { return __ocml_ilogb_f64(__x); } __DEVICE__ __RETURN_TYPE __finite(double __x) { return __builtin_isfinite(__x); } __DEVICE__ __RETURN_TYPE __isinf(double __x) { return __builtin_isinf(__x); } __DEVICE__ __RETURN_TYPE __isnan(double __x) { return __builtin_isnan(__x); } __DEVICE__ double j0(double __x) { return __ocml_j0_f64(__x); } __DEVICE__ double j1(double __x) { return __ocml_j1_f64(__x); } __DEVICE__ double jn(int __n, double __x) { // TODO: we could use Ahmes multiplication // and the Miller & Brown algorithm // for linear recurrences to get O(log n) steps, but it's unclear if // it'd be beneficial in this case. Placeholder until OCML adds // support. if (__n == 0) return j0(__x); if (__n == 1) return j1(__x); double __x0 = j0(__x); double __x1 = j1(__x); for (int __i = 1; __i < __n; ++__i) { double __x2 = (2 * __i) / __x * __x1 - __x0; __x0 = __x1; __x1 = __x2; } return __x1; } __DEVICE__ double ldexp(double __x, int __e) { return __builtin_amdgcn_ldexp(__x, __e); } __DEVICE__ double lgamma(double __x) { return __ocml_lgamma_f64(__x); } __DEVICE__ long long int llrint(double __x) { return __builtin_rint(__x); } __DEVICE__ long long int llround(double __x) { return __builtin_round(__x); } __DEVICE__ double log(double __x) { return __ocml_log_f64(__x); } __DEVICE__ double log10(double __x) { return __ocml_log10_f64(__x); } __DEVICE__ double log1p(double __x) { return __ocml_log1p_f64(__x); } __DEVICE__ double log2(double __x) { return __ocml_log2_f64(__x); } __DEVICE__ double logb(double __x) { return __ocml_logb_f64(__x); } __DEVICE__ long int lrint(double __x) { return __builtin_rint(__x); } __DEVICE__ long int lround(double __x) { return __builtin_round(__x); } __DEVICE__ double modf(double __x, double *__iptr) { double __tmp; #ifdef __OPENMP_AMDGCN__ #pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc) #endif double __r = __ocml_modf_f64(__x, (__attribute__((address_space(5))) double *)&__tmp); *__iptr = __tmp; return __r; } __DEVICE__ double nan(const char *__tagp) { #if !_WIN32 union { double val; struct ieee_double { uint64_t mantissa : 51; uint32_t quiet : 1; uint32_t exponent : 11; uint32_t sign : 1; } bits; } __tmp; __static_assert_type_size_equal(sizeof(__tmp.val), sizeof(__tmp.bits)); __tmp.bits.sign = 0u; __tmp.bits.exponent = ~0u; __tmp.bits.quiet = 1u; __tmp.bits.mantissa = __make_mantissa(__tagp); return __tmp.val; #else __static_assert_type_size_equal(sizeof(uint64_t), sizeof(double)); uint64_t __val = __make_mantissa(__tagp); __val |= 0xFFF << 51; return *reinterpret_cast<double *>(&__val); #endif } __DEVICE__ double nearbyint(double __x) { return __builtin_nearbyint(__x); } __DEVICE__ double nextafter(double __x, double __y) { return __ocml_nextafter_f64(__x, __y); } __DEVICE__ double norm(int __dim, const double *__a) { // TODO: placeholder until OCML adds support. double __r = 0; while (__dim--) { __r += __a[0] * __a[0]; ++__a; } return __builtin_sqrt(__r); } __DEVICE__ double norm3d(double __x, double __y, double __z) { return __ocml_len3_f64(__x, __y, __z); } __DEVICE__ double norm4d(double __x, double __y, double __z, double __w) { return __ocml_len4_f64(__x, __y, __z, __w); } __DEVICE__ double normcdf(double __x) { return __ocml_ncdf_f64(__x); } __DEVICE__ double normcdfinv(double __x) { return __ocml_ncdfinv_f64(__x); } __DEVICE__ double pow(double __x, double __y) { return __ocml_pow_f64(__x, __y); } __DEVICE__ double powi(double __x, int __y) { return __ocml_pown_f64(__x, __y); } __DEVICE__ double rcbrt(double __x) { return __ocml_rcbrt_f64(__x); } __DEVICE__ double remainder(double __x, double __y) { return __ocml_remainder_f64(__x, __y); } __DEVICE__ double remquo(double __x, double __y, int *__quo) { int __tmp; #ifdef __OPENMP_AMDGCN__ #pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc) #endif double __r = __ocml_remquo_f64( __x, __y, (__attribute__((address_space(5))) int *)&__tmp); *__quo = __tmp; return __r; } __DEVICE__ double rhypot(double __x, double __y) { return __ocml_rhypot_f64(__x, __y); } __DEVICE__ double rint(double __x) { return __builtin_rint(__x); } __DEVICE__ double rnorm(int __dim, const double *__a) { // TODO: placeholder until OCML adds support. double __r = 0; while (__dim--) { __r += __a[0] * __a[0]; ++__a; } return __ocml_rsqrt_f64(__r); } __DEVICE__ double rnorm3d(double __x, double __y, double __z) { return __ocml_rlen3_f64(__x, __y, __z); } __DEVICE__ double rnorm4d(double __x, double __y, double __z, double __w) { return __ocml_rlen4_f64(__x, __y, __z, __w); } __DEVICE__ double round(double __x) { return __builtin_round(__x); } __DEVICE__ double rsqrt(double __x) { return __ocml_rsqrt_f64(__x); } __DEVICE__ double scalbln(double __x, long int __n) { return (__n < INT_MAX) ? __builtin_amdgcn_ldexp(__x, __n) : __ocml_scalb_f64(__x, __n); } __DEVICE__ double scalbn(double __x, int __n) { return __builtin_amdgcn_ldexp(__x, __n); } __DEVICE__ __RETURN_TYPE __signbit(double __x) { return __builtin_signbit(__x); } __DEVICE__ double sin(double __x) { return __ocml_sin_f64(__x); } __DEVICE__ void sincos(double __x, double *__sinptr, double *__cosptr) { double __tmp; #ifdef __OPENMP_AMDGCN__ #pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc) #endif *__sinptr = __ocml_sincos_f64( __x, (__attribute__((address_space(5))) double *)&__tmp); *__cosptr = __tmp; } __DEVICE__ void sincospi(double __x, double *__sinptr, double *__cosptr) { double __tmp; #ifdef __OPENMP_AMDGCN__ #pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc) #endif *__sinptr = __ocml_sincospi_f64( __x, (__attribute__((address_space(5))) double *)&__tmp); *__cosptr = __tmp; } __DEVICE__ double sinh(double __x) { return __ocml_sinh_f64(__x); } __DEVICE__ double sinpi(double __x) { return __ocml_sinpi_f64(__x); } __DEVICE__ double sqrt(double __x) { return __builtin_sqrt(__x); } __DEVICE__ double tan(double __x) { return __ocml_tan_f64(__x); } __DEVICE__ double tanh(double __x) { return __ocml_tanh_f64(__x); } __DEVICE__ double tgamma(double __x) { return __ocml_tgamma_f64(__x); } __DEVICE__ double trunc(double __x) { return __builtin_trunc(__x); } __DEVICE__ double y0(double __x) { return __ocml_y0_f64(__x); } __DEVICE__ double y1(double __x) { return __ocml_y1_f64(__x); } __DEVICE__ double yn(int __n, double __x) { // TODO: we could use Ahmes multiplication // and the Miller & Brown algorithm // for linear recurrences to get O(log n) steps, but it's unclear if // it'd be beneficial in this case. Placeholder until OCML adds // support. if (__n == 0) return y0(__x); if (__n == 1) return y1(__x); double __x0 = y0(__x); double __x1 = y1(__x); for (int __i = 1; __i < __n; ++__i) { double __x2 = (2 * __i) / __x * __x1 - __x0; __x0 = __x1; __x1 = __x2; } return __x1; } // BEGIN INTRINSICS #if defined OCML_BASIC_ROUNDED_OPERATIONS __DEVICE__ double __dadd_rd(double __x, double __y) { return __ocml_add_rtn_f64(__x, __y); } __DEVICE__ double __dadd_rn(double __x, double __y) { return __ocml_add_rte_f64(__x, __y); } __DEVICE__ double __dadd_ru(double __x, double __y) { return __ocml_add_rtp_f64(__x, __y); } __DEVICE__ double __dadd_rz(double __x, double __y) { return __ocml_add_rtz_f64(__x, __y); } #else __DEVICE__ double __dadd_rn(double __x, double __y) { return __x + __y; } #endif #if defined OCML_BASIC_ROUNDED_OPERATIONS __DEVICE__ double __ddiv_rd(double __x, double __y) { return __ocml_div_rtn_f64(__x, __y); } __DEVICE__ double __ddiv_rn(double __x, double __y) { return __ocml_div_rte_f64(__x, __y); } __DEVICE__ double __ddiv_ru(double __x, double __y) { return __ocml_div_rtp_f64(__x, __y); } __DEVICE__ double __ddiv_rz(double __x, double __y) { return __ocml_div_rtz_f64(__x, __y); } #else __DEVICE__ double __ddiv_rn(double __x, double __y) { return __x / __y; } #endif #if defined OCML_BASIC_ROUNDED_OPERATIONS __DEVICE__ double __dmul_rd(double __x, double __y) { return __ocml_mul_rtn_f64(__x, __y); } __DEVICE__ double __dmul_rn(double __x, double __y) { return __ocml_mul_rte_f64(__x, __y); } __DEVICE__ double __dmul_ru(double __x, double __y) { return __ocml_mul_rtp_f64(__x, __y); } __DEVICE__ double __dmul_rz(double __x, double __y) { return __ocml_mul_rtz_f64(__x, __y); } #else __DEVICE__ double __dmul_rn(double __x, double __y) { return __x * __y; } #endif #if defined OCML_BASIC_ROUNDED_OPERATIONS __DEVICE__ double __drcp_rd(double __x) { return __ocml_div_rtn_f64(1.0, __x); } __DEVICE__ double __drcp_rn(double __x) { return __ocml_div_rte_f64(1.0, __x); } __DEVICE__ double __drcp_ru(double __x) { return __ocml_div_rtp_f64(1.0, __x); } __DEVICE__ double __drcp_rz(double __x) { return __ocml_div_rtz_f64(1.0, __x); } #else __DEVICE__ double __drcp_rn(double __x) { return 1.0 / __x; } #endif #if defined OCML_BASIC_ROUNDED_OPERATIONS __DEVICE__ double __dsqrt_rd(double __x) { return __ocml_sqrt_rtn_f64(__x); } __DEVICE__ double __dsqrt_rn(double __x) { return __ocml_sqrt_rte_f64(__x); } __DEVICE__ double __dsqrt_ru(double __x) { return __ocml_sqrt_rtp_f64(__x); } __DEVICE__ double __dsqrt_rz(double __x) { return __ocml_sqrt_rtz_f64(__x); } #else __DEVICE__ double __dsqrt_rn(double __x) { return __builtin_sqrt(__x); } #endif #if defined OCML_BASIC_ROUNDED_OPERATIONS __DEVICE__ double __dsub_rd(double __x, double __y) { return __ocml_sub_rtn_f64(__x, __y); } __DEVICE__ double __dsub_rn(double __x, double __y) { return __ocml_sub_rte_f64(__x, __y); } __DEVICE__ double __dsub_ru(double __x, double __y) { return __ocml_sub_rtp_f64(__x, __y); } __DEVICE__ double __dsub_rz(double __x, double __y) { return __ocml_sub_rtz_f64(__x, __y); } #else __DEVICE__ double __dsub_rn(double __x, double __y) { return __x - __y; } #endif #if defined OCML_BASIC_ROUNDED_OPERATIONS __DEVICE__ double __fma_rd(double __x, double __y, double __z) { return __ocml_fma_rtn_f64(__x, __y, __z); } __DEVICE__ double __fma_rn(double __x, double __y, double __z) { return __ocml_fma_rte_f64(__x, __y, __z); } __DEVICE__ double __fma_ru(double __x, double __y, double __z) { return __ocml_fma_rtp_f64(__x, __y, __z); } __DEVICE__ double __fma_rz(double __x, double __y, double __z) { return __ocml_fma_rtz_f64(__x, __y, __z); } #else __DEVICE__ double __fma_rn(double __x, double __y, double __z) { return __builtin_fma(__x, __y, __z); } #endif // END INTRINSICS // END DOUBLE // C only macros #if !defined(__cplusplus) && __STDC_VERSION__ >= 201112L #define isfinite(__x) _Generic((__x), float : __finitef, double : __finite)(__x) #define isinf(__x) _Generic((__x), float : __isinff, double : __isinf)(__x) #define isnan(__x) _Generic((__x), float : __isnanf, double : __isnan)(__x) #define signbit(__x) \ _Generic((__x), float : __signbitf, double : __signbit)(__x) #endif // !defined(__cplusplus) && __STDC_VERSION__ >= 201112L #if defined(__cplusplus) template <class T> __DEVICE__ T min(T __arg1, T __arg2) { return (__arg1 < __arg2) ? __arg1 : __arg2; } template <class T> __DEVICE__ T max(T __arg1, T __arg2) { return (__arg1 > __arg2) ? __arg1 : __arg2; } __DEVICE__ int min(int __arg1, int __arg2) { return (__arg1 < __arg2) ? __arg1 : __arg2; } __DEVICE__ int max(int __arg1, int __arg2) { return (__arg1 > __arg2) ? __arg1 : __arg2; } __DEVICE__ float max(float __x, float __y) { return __builtin_fmaxf(__x, __y); } __DEVICE__ double max(double __x, double __y) { return __builtin_fmax(__x, __y); } __DEVICE__ float min(float __x, float __y) { return __builtin_fminf(__x, __y); } __DEVICE__ double min(double __x, double __y) { return __builtin_fmin(__x, __y); } #if !defined(__HIPCC_RTC__) && !defined(__OPENMP_AMDGCN__) __host__ inline static int min(int __arg1, int __arg2) { return __arg1 < __arg2 ? __arg1 : __arg2; } __host__ inline static int max(int __arg1, int __arg2) { return __arg1 > __arg2 ? __arg1 : __arg2; } #endif // !defined(__HIPCC_RTC__) && !defined(__OPENMP_AMDGCN__) #endif #pragma pop_macro("__DEVICE__") #pragma pop_macro("__RETURN_TYPE") #pragma pop_macro("__FAST_OR_SLOW") #endif // __CLANG_GPU_DISABLE_MATH_WRAPPERS #endif // __CLANG_HIP_MATH_H__