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_cmath.h - HIP cmath decls -----------------------------=== * * 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_CMATH_H__ #define __CLANG_HIP_CMATH_H__ #if !defined(__HIP__) && !defined(__OPENMP_AMDGCN__) #error "This file is for HIP and OpenMP AMDGCN device compilation only." #endif #if !defined(__HIPCC_RTC__) #if defined(__cplusplus) #include <limits> #include <type_traits> #include <utility> #endif #include <limits.h> #include <stdint.h> #endif // !defined(__HIPCC_RTC__) #pragma push_macro("__DEVICE__") #pragma push_macro("__CONSTEXPR__") #ifdef __OPENMP_AMDGCN__ #define __DEVICE__ static __attribute__((always_inline, nothrow)) #define __CONSTEXPR__ constexpr #else #define __DEVICE__ static __device__ inline __attribute__((always_inline)) #define __CONSTEXPR__ #endif // __OPENMP_AMDGCN__ // Start with functions that cannot be defined by DEF macros below. #if defined(__cplusplus) #if defined __OPENMP_AMDGCN__ __DEVICE__ __CONSTEXPR__ float fabs(float __x) { return ::fabsf(__x); } __DEVICE__ __CONSTEXPR__ float sin(float __x) { return ::sinf(__x); } __DEVICE__ __CONSTEXPR__ float cos(float __x) { return ::cosf(__x); } #endif __DEVICE__ __CONSTEXPR__ double abs(double __x) { return ::fabs(__x); } __DEVICE__ __CONSTEXPR__ float abs(float __x) { return ::fabsf(__x); } __DEVICE__ __CONSTEXPR__ long long abs(long long __n) { return ::llabs(__n); } __DEVICE__ __CONSTEXPR__ long abs(long __n) { return ::labs(__n); } __DEVICE__ __CONSTEXPR__ float fma(float __x, float __y, float __z) { return ::fmaf(__x, __y, __z); } #if !defined(__HIPCC_RTC__) // The value returned by fpclassify is platform dependent, therefore it is not // supported by hipRTC. __DEVICE__ __CONSTEXPR__ int fpclassify(float __x) { return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL, FP_ZERO, __x); } __DEVICE__ __CONSTEXPR__ int fpclassify(double __x) { return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL, FP_ZERO, __x); } #endif // !defined(__HIPCC_RTC__) __DEVICE__ __CONSTEXPR__ float frexp(float __arg, int *__exp) { return ::frexpf(__arg, __exp); } #if defined(__OPENMP_AMDGCN__) // For OpenMP we work around some old system headers that have non-conforming // `isinf(float)` and `isnan(float)` implementations that return an `int`. We do // this by providing two versions of these functions, differing only in the // return type. To avoid conflicting definitions we disable implicit base // function generation. That means we will end up with two specializations, one // per type, but only one has a base function defined by the system header. #pragma omp begin declare variant match( \ implementation = {extension(disable_implicit_base)}) // FIXME: We lack an extension to customize the mangling of the variants, e.g., // add a suffix. This means we would clash with the names of the variants // (note that we do not create implicit base functions here). To avoid // this clash we add a new trait to some of them that is always true // (this is LLVM after all ;)). It will only influence the mangled name // of the variants inside the inner region and avoid the clash. #pragma omp begin declare variant match(implementation = {vendor(llvm)}) __DEVICE__ __CONSTEXPR__ int isinf(float __x) { return ::__isinff(__x); } __DEVICE__ __CONSTEXPR__ int isinf(double __x) { return ::__isinf(__x); } __DEVICE__ __CONSTEXPR__ int isfinite(float __x) { return ::__finitef(__x); } __DEVICE__ __CONSTEXPR__ int isfinite(double __x) { return ::__finite(__x); } __DEVICE__ __CONSTEXPR__ int isnan(float __x) { return ::__isnanf(__x); } __DEVICE__ __CONSTEXPR__ int isnan(double __x) { return ::__isnan(__x); } #pragma omp end declare variant #endif // defined(__OPENMP_AMDGCN__) __DEVICE__ __CONSTEXPR__ bool isinf(float __x) { return ::__isinff(__x); } __DEVICE__ __CONSTEXPR__ bool isinf(double __x) { return ::__isinf(__x); } __DEVICE__ __CONSTEXPR__ bool isfinite(float __x) { return ::__finitef(__x); } __DEVICE__ __CONSTEXPR__ bool isfinite(double __x) { return ::__finite(__x); } __DEVICE__ __CONSTEXPR__ bool isnan(float __x) { return ::__isnanf(__x); } __DEVICE__ __CONSTEXPR__ bool isnan(double __x) { return ::__isnan(__x); } #if defined(__OPENMP_AMDGCN__) #pragma omp end declare variant #endif // defined(__OPENMP_AMDGCN__) __DEVICE__ __CONSTEXPR__ bool isgreater(float __x, float __y) { return __builtin_isgreater(__x, __y); } __DEVICE__ __CONSTEXPR__ bool isgreater(double __x, double __y) { return __builtin_isgreater(__x, __y); } __DEVICE__ __CONSTEXPR__ bool isgreaterequal(float __x, float __y) { return __builtin_isgreaterequal(__x, __y); } __DEVICE__ __CONSTEXPR__ bool isgreaterequal(double __x, double __y) { return __builtin_isgreaterequal(__x, __y); } __DEVICE__ __CONSTEXPR__ bool isless(float __x, float __y) { return __builtin_isless(__x, __y); } __DEVICE__ __CONSTEXPR__ bool isless(double __x, double __y) { return __builtin_isless(__x, __y); } __DEVICE__ __CONSTEXPR__ bool islessequal(float __x, float __y) { return __builtin_islessequal(__x, __y); } __DEVICE__ __CONSTEXPR__ bool islessequal(double __x, double __y) { return __builtin_islessequal(__x, __y); } __DEVICE__ __CONSTEXPR__ bool islessgreater(float __x, float __y) { return __builtin_islessgreater(__x, __y); } __DEVICE__ __CONSTEXPR__ bool islessgreater(double __x, double __y) { return __builtin_islessgreater(__x, __y); } __DEVICE__ __CONSTEXPR__ bool isnormal(float __x) { return __builtin_isnormal(__x); } __DEVICE__ __CONSTEXPR__ bool isnormal(double __x) { return __builtin_isnormal(__x); } __DEVICE__ __CONSTEXPR__ bool isunordered(float __x, float __y) { return __builtin_isunordered(__x, __y); } __DEVICE__ __CONSTEXPR__ bool isunordered(double __x, double __y) { return __builtin_isunordered(__x, __y); } __DEVICE__ __CONSTEXPR__ float modf(float __x, float *__iptr) { return ::modff(__x, __iptr); } __DEVICE__ __CONSTEXPR__ float pow(float __base, int __iexp) { return ::powif(__base, __iexp); } __DEVICE__ __CONSTEXPR__ double pow(double __base, int __iexp) { return ::powi(__base, __iexp); } __DEVICE__ __CONSTEXPR__ float remquo(float __x, float __y, int *__quo) { return ::remquof(__x, __y, __quo); } __DEVICE__ __CONSTEXPR__ float scalbln(float __x, long int __n) { return ::scalblnf(__x, __n); } __DEVICE__ __CONSTEXPR__ bool signbit(float __x) { return ::__signbitf(__x); } __DEVICE__ __CONSTEXPR__ bool signbit(double __x) { return ::__signbit(__x); } // Notably missing above is nexttoward. We omit it because // ocml doesn't provide an implementation, and we don't want to be in the // business of implementing tricky libm functions in this header. // Other functions. __DEVICE__ __CONSTEXPR__ _Float16 fma(_Float16 __x, _Float16 __y, _Float16 __z) { return __builtin_fmaf16(__x, __y, __z); } __DEVICE__ __CONSTEXPR__ _Float16 pow(_Float16 __base, int __iexp) { return __ocml_pown_f16(__base, __iexp); } #ifndef __OPENMP_AMDGCN__ // BEGIN DEF_FUN and HIP_OVERLOAD // BEGIN DEF_FUN #pragma push_macro("__DEF_FUN1") #pragma push_macro("__DEF_FUN2") #pragma push_macro("__DEF_FUN2_FI") // Define cmath functions with float argument and returns __retty. #define __DEF_FUN1(__retty, __func) \ __DEVICE__ __CONSTEXPR__ __retty __func(float __x) { return __func##f(__x); } // Define cmath functions with two float arguments and returns __retty. #define __DEF_FUN2(__retty, __func) \ __DEVICE__ __CONSTEXPR__ __retty __func(float __x, float __y) { \ return __func##f(__x, __y); \ } // Define cmath functions with a float and an int argument and returns __retty. #define __DEF_FUN2_FI(__retty, __func) \ __DEVICE__ __CONSTEXPR__ __retty __func(float __x, int __y) { \ return __func##f(__x, __y); \ } __DEF_FUN1(float, acos) __DEF_FUN1(float, acosh) __DEF_FUN1(float, asin) __DEF_FUN1(float, asinh) __DEF_FUN1(float, atan) __DEF_FUN2(float, atan2) __DEF_FUN1(float, atanh) __DEF_FUN1(float, cbrt) __DEF_FUN1(float, ceil) __DEF_FUN2(float, copysign) __DEF_FUN1(float, cos) __DEF_FUN1(float, cosh) __DEF_FUN1(float, erf) __DEF_FUN1(float, erfc) __DEF_FUN1(float, exp) __DEF_FUN1(float, exp2) __DEF_FUN1(float, expm1) __DEF_FUN1(float, fabs) __DEF_FUN2(float, fdim) __DEF_FUN1(float, floor) __DEF_FUN2(float, fmax) __DEF_FUN2(float, fmin) __DEF_FUN2(float, fmod) __DEF_FUN2(float, hypot) __DEF_FUN1(int, ilogb) __DEF_FUN2_FI(float, ldexp) __DEF_FUN1(float, lgamma) __DEF_FUN1(float, log) __DEF_FUN1(float, log10) __DEF_FUN1(float, log1p) __DEF_FUN1(float, log2) __DEF_FUN1(float, logb) __DEF_FUN1(long long, llrint) __DEF_FUN1(long long, llround) __DEF_FUN1(long, lrint) __DEF_FUN1(long, lround) __DEF_FUN1(float, nearbyint) __DEF_FUN2(float, nextafter) __DEF_FUN2(float, pow) __DEF_FUN2(float, remainder) __DEF_FUN1(float, rint) __DEF_FUN1(float, round) __DEF_FUN2_FI(float, scalbn) __DEF_FUN1(float, sin) __DEF_FUN1(float, sinh) __DEF_FUN1(float, sqrt) __DEF_FUN1(float, tan) __DEF_FUN1(float, tanh) __DEF_FUN1(float, tgamma) __DEF_FUN1(float, trunc) #pragma pop_macro("__DEF_FUN1") #pragma pop_macro("__DEF_FUN2") #pragma pop_macro("__DEF_FUN2_FI") // END DEF_FUN // BEGIN HIP_OVERLOAD #pragma push_macro("__HIP_OVERLOAD1") #pragma push_macro("__HIP_OVERLOAD2") // __hip_enable_if::type is a type function which returns __T if __B is true. template <bool __B, class __T = void> struct __hip_enable_if {}; template <class __T> struct __hip_enable_if<true, __T> { typedef __T type; }; namespace __hip { template <class _Tp> struct is_integral { enum { value = 0 }; }; template <> struct is_integral<bool> { enum { value = 1 }; }; template <> struct is_integral<char> { enum { value = 1 }; }; template <> struct is_integral<signed char> { enum { value = 1 }; }; template <> struct is_integral<unsigned char> { enum { value = 1 }; }; template <> struct is_integral<wchar_t> { enum { value = 1 }; }; template <> struct is_integral<short> { enum { value = 1 }; }; template <> struct is_integral<unsigned short> { enum { value = 1 }; }; template <> struct is_integral<int> { enum { value = 1 }; }; template <> struct is_integral<unsigned int> { enum { value = 1 }; }; template <> struct is_integral<long> { enum { value = 1 }; }; template <> struct is_integral<unsigned long> { enum { value = 1 }; }; template <> struct is_integral<long long> { enum { value = 1 }; }; template <> struct is_integral<unsigned long long> { enum { value = 1 }; }; // ToDo: specializes is_arithmetic<_Float16> template <class _Tp> struct is_arithmetic { enum { value = 0 }; }; template <> struct is_arithmetic<bool> { enum { value = 1 }; }; template <> struct is_arithmetic<char> { enum { value = 1 }; }; template <> struct is_arithmetic<signed char> { enum { value = 1 }; }; template <> struct is_arithmetic<unsigned char> { enum { value = 1 }; }; template <> struct is_arithmetic<wchar_t> { enum { value = 1 }; }; template <> struct is_arithmetic<short> { enum { value = 1 }; }; template <> struct is_arithmetic<unsigned short> { enum { value = 1 }; }; template <> struct is_arithmetic<int> { enum { value = 1 }; }; template <> struct is_arithmetic<unsigned int> { enum { value = 1 }; }; template <> struct is_arithmetic<long> { enum { value = 1 }; }; template <> struct is_arithmetic<unsigned long> { enum { value = 1 }; }; template <> struct is_arithmetic<long long> { enum { value = 1 }; }; template <> struct is_arithmetic<unsigned long long> { enum { value = 1 }; }; template <> struct is_arithmetic<float> { enum { value = 1 }; }; template <> struct is_arithmetic<double> { enum { value = 1 }; }; struct true_type { static const __constant__ bool value = true; }; struct false_type { static const __constant__ bool value = false; }; template <typename __T, typename __U> struct is_same : public false_type {}; template <typename __T> struct is_same<__T, __T> : public true_type {}; template <typename __T> struct add_rvalue_reference { typedef __T &&type; }; template <typename __T> typename add_rvalue_reference<__T>::type declval(); // decltype is only available in C++11 and above. #if __cplusplus >= 201103L // __hip_promote template <class _Tp> struct __numeric_type { static void __test(...); static _Float16 __test(_Float16); static float __test(float); static double __test(char); static double __test(int); static double __test(unsigned); static double __test(long); static double __test(unsigned long); static double __test(long long); static double __test(unsigned long long); static double __test(double); // No support for long double, use double instead. static double __test(long double); template <typename _U> static auto __test_impl(int) -> decltype(__test(declval<_U>())); template <typename _U> static void __test_impl(...); typedef decltype(__test_impl<_Tp>(0)) type; static const bool value = !is_same<type, void>::value; }; template <> struct __numeric_type<void> { static const bool value = true; }; template <class _A1, class _A2 = void, class _A3 = void, bool = __numeric_type<_A1>::value &&__numeric_type<_A2>::value &&__numeric_type<_A3>::value> class __promote_imp { public: static const bool value = false; }; template <class _A1, class _A2, class _A3> class __promote_imp<_A1, _A2, _A3, true> { private: typedef typename __promote_imp<_A1>::type __type1; typedef typename __promote_imp<_A2>::type __type2; typedef typename __promote_imp<_A3>::type __type3; public: typedef decltype(__type1() + __type2() + __type3()) type; static const bool value = true; }; template <class _A1, class _A2> class __promote_imp<_A1, _A2, void, true> { private: typedef typename __promote_imp<_A1>::type __type1; typedef typename __promote_imp<_A2>::type __type2; public: typedef decltype(__type1() + __type2()) type; static const bool value = true; }; template <class _A1> class __promote_imp<_A1, void, void, true> { public: typedef typename __numeric_type<_A1>::type type; static const bool value = true; }; template <class _A1, class _A2 = void, class _A3 = void> class __promote : public __promote_imp<_A1, _A2, _A3> {}; #endif //__cplusplus >= 201103L } // namespace __hip // __HIP_OVERLOAD1 is used to resolve function calls with integer argument to // avoid compilation error due to ambibuity. e.g. floor(5) is resolved with // floor(double). #define __HIP_OVERLOAD1(__retty, __fn) \ template <typename __T> \ __DEVICE__ __CONSTEXPR__ \ typename __hip_enable_if<__hip::is_integral<__T>::value, __retty>::type \ __fn(__T __x) { \ return ::__fn((double)__x); \ } // __HIP_OVERLOAD2 is used to resolve function calls with mixed float/double // or integer argument to avoid compilation error due to ambibuity. e.g. // max(5.0f, 6.0) is resolved with max(double, double). #if __cplusplus >= 201103L #define __HIP_OVERLOAD2(__retty, __fn) \ template <typename __T1, typename __T2> \ __DEVICE__ __CONSTEXPR__ typename __hip_enable_if< \ __hip::is_arithmetic<__T1>::value && __hip::is_arithmetic<__T2>::value, \ typename __hip::__promote<__T1, __T2>::type>::type \ __fn(__T1 __x, __T2 __y) { \ typedef typename __hip::__promote<__T1, __T2>::type __result_type; \ return __fn((__result_type)__x, (__result_type)__y); \ } #else #define __HIP_OVERLOAD2(__retty, __fn) \ template <typename __T1, typename __T2> \ __DEVICE__ __CONSTEXPR__ \ typename __hip_enable_if<__hip::is_arithmetic<__T1>::value && \ __hip::is_arithmetic<__T2>::value, \ __retty>::type \ __fn(__T1 __x, __T2 __y) { \ return __fn((double)__x, (double)__y); \ } #endif __HIP_OVERLOAD1(double, acos) __HIP_OVERLOAD1(double, acosh) __HIP_OVERLOAD1(double, asin) __HIP_OVERLOAD1(double, asinh) __HIP_OVERLOAD1(double, atan) __HIP_OVERLOAD2(double, atan2) __HIP_OVERLOAD1(double, atanh) __HIP_OVERLOAD1(double, cbrt) __HIP_OVERLOAD1(double, ceil) __HIP_OVERLOAD2(double, copysign) __HIP_OVERLOAD1(double, cos) __HIP_OVERLOAD1(double, cosh) __HIP_OVERLOAD1(double, erf) __HIP_OVERLOAD1(double, erfc) __HIP_OVERLOAD1(double, exp) __HIP_OVERLOAD1(double, exp2) __HIP_OVERLOAD1(double, expm1) __HIP_OVERLOAD1(double, fabs) __HIP_OVERLOAD2(double, fdim) __HIP_OVERLOAD1(double, floor) __HIP_OVERLOAD2(double, fmax) __HIP_OVERLOAD2(double, fmin) __HIP_OVERLOAD2(double, fmod) #if !defined(__HIPCC_RTC__) __HIP_OVERLOAD1(int, fpclassify) #endif // !defined(__HIPCC_RTC__) __HIP_OVERLOAD2(double, hypot) __HIP_OVERLOAD1(int, ilogb) __HIP_OVERLOAD1(bool, isfinite) __HIP_OVERLOAD2(bool, isgreater) __HIP_OVERLOAD2(bool, isgreaterequal) __HIP_OVERLOAD1(bool, isinf) __HIP_OVERLOAD2(bool, isless) __HIP_OVERLOAD2(bool, islessequal) __HIP_OVERLOAD2(bool, islessgreater) __HIP_OVERLOAD1(bool, isnan) __HIP_OVERLOAD1(bool, isnormal) __HIP_OVERLOAD2(bool, isunordered) __HIP_OVERLOAD1(double, lgamma) __HIP_OVERLOAD1(double, log) __HIP_OVERLOAD1(double, log10) __HIP_OVERLOAD1(double, log1p) __HIP_OVERLOAD1(double, log2) __HIP_OVERLOAD1(double, logb) __HIP_OVERLOAD1(long long, llrint) __HIP_OVERLOAD1(long long, llround) __HIP_OVERLOAD1(long, lrint) __HIP_OVERLOAD1(long, lround) __HIP_OVERLOAD1(double, nearbyint) __HIP_OVERLOAD2(double, nextafter) __HIP_OVERLOAD2(double, pow) __HIP_OVERLOAD2(double, remainder) __HIP_OVERLOAD1(double, rint) __HIP_OVERLOAD1(double, round) __HIP_OVERLOAD1(bool, signbit) __HIP_OVERLOAD1(double, sin) __HIP_OVERLOAD1(double, sinh) __HIP_OVERLOAD1(double, sqrt) __HIP_OVERLOAD1(double, tan) __HIP_OVERLOAD1(double, tanh) __HIP_OVERLOAD1(double, tgamma) __HIP_OVERLOAD1(double, trunc) // Overload these but don't add them to std, they are not part of cmath. __HIP_OVERLOAD2(double, max) __HIP_OVERLOAD2(double, min) // Additional Overloads that don't quite match HIP_OVERLOAD. #if __cplusplus >= 201103L template <typename __T1, typename __T2, typename __T3> __DEVICE__ __CONSTEXPR__ typename __hip_enable_if< __hip::is_arithmetic<__T1>::value && __hip::is_arithmetic<__T2>::value && __hip::is_arithmetic<__T3>::value, typename __hip::__promote<__T1, __T2, __T3>::type>::type fma(__T1 __x, __T2 __y, __T3 __z) { typedef typename __hip::__promote<__T1, __T2, __T3>::type __result_type; return ::fma((__result_type)__x, (__result_type)__y, (__result_type)__z); } #else template <typename __T1, typename __T2, typename __T3> __DEVICE__ __CONSTEXPR__ typename __hip_enable_if<__hip::is_arithmetic<__T1>::value && __hip::is_arithmetic<__T2>::value && __hip::is_arithmetic<__T3>::value, double>::type fma(__T1 __x, __T2 __y, __T3 __z) { return ::fma((double)__x, (double)__y, (double)__z); } #endif template <typename __T> __DEVICE__ __CONSTEXPR__ typename __hip_enable_if<__hip::is_integral<__T>::value, double>::type frexp(__T __x, int *__exp) { return ::frexp((double)__x, __exp); } template <typename __T> __DEVICE__ __CONSTEXPR__ typename __hip_enable_if<__hip::is_integral<__T>::value, double>::type ldexp(__T __x, int __exp) { return ::ldexp((double)__x, __exp); } template <typename __T> __DEVICE__ __CONSTEXPR__ typename __hip_enable_if<__hip::is_integral<__T>::value, double>::type modf(__T __x, double *__exp) { return ::modf((double)__x, __exp); } #if __cplusplus >= 201103L template <typename __T1, typename __T2> __DEVICE__ __CONSTEXPR__ typename __hip_enable_if<__hip::is_arithmetic<__T1>::value && __hip::is_arithmetic<__T2>::value, typename __hip::__promote<__T1, __T2>::type>::type remquo(__T1 __x, __T2 __y, int *__quo) { typedef typename __hip::__promote<__T1, __T2>::type __result_type; return ::remquo((__result_type)__x, (__result_type)__y, __quo); } #else template <typename __T1, typename __T2> __DEVICE__ __CONSTEXPR__ typename __hip_enable_if<__hip::is_arithmetic<__T1>::value && __hip::is_arithmetic<__T2>::value, double>::type remquo(__T1 __x, __T2 __y, int *__quo) { return ::remquo((double)__x, (double)__y, __quo); } #endif template <typename __T> __DEVICE__ __CONSTEXPR__ typename __hip_enable_if<__hip::is_integral<__T>::value, double>::type scalbln(__T __x, long int __exp) { return ::scalbln((double)__x, __exp); } template <typename __T> __DEVICE__ __CONSTEXPR__ typename __hip_enable_if<__hip::is_integral<__T>::value, double>::type scalbn(__T __x, int __exp) { return ::scalbn((double)__x, __exp); } #pragma pop_macro("__HIP_OVERLOAD1") #pragma pop_macro("__HIP_OVERLOAD2") // END HIP_OVERLOAD // END DEF_FUN and HIP_OVERLOAD #endif // ifndef __OPENMP_AMDGCN__ #endif // defined(__cplusplus) #ifndef __OPENMP_AMDGCN__ // Define these overloads inside the namespace our standard library uses. #if !defined(__HIPCC_RTC__) #ifdef _LIBCPP_BEGIN_NAMESPACE_STD _LIBCPP_BEGIN_NAMESPACE_STD #else namespace std { #ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION _GLIBCXX_BEGIN_NAMESPACE_VERSION #endif // _GLIBCXX_BEGIN_NAMESPACE_VERSION #endif // _LIBCPP_BEGIN_NAMESPACE_STD // Pull the new overloads we defined above into namespace std. // using ::abs; - This may be considered for C++. using ::acos; using ::acosh; using ::asin; using ::asinh; using ::atan; using ::atan2; using ::atanh; using ::cbrt; using ::ceil; using ::copysign; using ::cos; using ::cosh; using ::erf; using ::erfc; using ::exp; using ::exp2; using ::expm1; using ::fabs; using ::fdim; using ::floor; using ::fma; using ::fmax; using ::fmin; using ::fmod; using ::fpclassify; using ::frexp; using ::hypot; using ::ilogb; using ::isfinite; using ::isgreater; using ::isgreaterequal; using ::isless; using ::islessequal; using ::islessgreater; using ::isnormal; using ::isunordered; using ::ldexp; using ::lgamma; using ::llrint; using ::llround; using ::log; using ::log10; using ::log1p; using ::log2; using ::logb; using ::lrint; using ::lround; using ::modf; // using ::nan; - This may be considered for C++. // using ::nanf; - This may be considered for C++. // using ::nanl; - This is not yet defined. using ::nearbyint; using ::nextafter; // using ::nexttoward; - Omit this since we do not have a definition. using ::pow; using ::remainder; using ::remquo; using ::rint; using ::round; using ::scalbln; using ::scalbn; using ::signbit; using ::sin; using ::sinh; using ::sqrt; using ::tan; using ::tanh; using ::tgamma; using ::trunc; // Well this is fun: We need to pull these symbols in for libc++, but we can't // pull them in with libstdc++, because its ::isinf and ::isnan are different // than its std::isinf and std::isnan. #ifndef __GLIBCXX__ using ::isinf; using ::isnan; #endif // Finally, pull the "foobarf" functions that HIP defines into std. using ::acosf; using ::acoshf; using ::asinf; using ::asinhf; using ::atan2f; using ::atanf; using ::atanhf; using ::cbrtf; using ::ceilf; using ::copysignf; using ::cosf; using ::coshf; using ::erfcf; using ::erff; using ::exp2f; using ::expf; using ::expm1f; using ::fabsf; using ::fdimf; using ::floorf; using ::fmaf; using ::fmaxf; using ::fminf; using ::fmodf; using ::frexpf; using ::hypotf; using ::ilogbf; using ::ldexpf; using ::lgammaf; using ::llrintf; using ::llroundf; using ::log10f; using ::log1pf; using ::log2f; using ::logbf; using ::logf; using ::lrintf; using ::lroundf; using ::modff; using ::nearbyintf; using ::nextafterf; // using ::nexttowardf; - Omit this since we do not have a definition. using ::powf; using ::remainderf; using ::remquof; using ::rintf; using ::roundf; using ::scalblnf; using ::scalbnf; using ::sinf; using ::sinhf; using ::sqrtf; using ::tanf; using ::tanhf; using ::tgammaf; using ::truncf; #ifdef _LIBCPP_END_NAMESPACE_STD _LIBCPP_END_NAMESPACE_STD #else #ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION _GLIBCXX_END_NAMESPACE_VERSION #endif // _GLIBCXX_BEGIN_NAMESPACE_VERSION } // namespace std #endif // _LIBCPP_END_NAMESPACE_STD #endif // !defined(__HIPCC_RTC__) // Define device-side math functions from <ymath.h> on MSVC. #if !defined(__HIPCC_RTC__) #if defined(_MSC_VER) // Before VS2019, `<ymath.h>` is also included in `<limits>` and other headers. // But, from VS2019, it's only included in `<complex>`. Need to include // `<ymath.h>` here to ensure C functions declared there won't be markded as // `__host__` and `__device__` through `<complex>` wrapper. #include <ymath.h> #if defined(__cplusplus) extern "C" { #endif // defined(__cplusplus) __DEVICE__ __CONSTEXPR__ __attribute__((overloadable)) double _Cosh(double x, double y) { return cosh(x) * y; } __DEVICE__ __CONSTEXPR__ __attribute__((overloadable)) float _FCosh(float x, float y) { return coshf(x) * y; } __DEVICE__ __CONSTEXPR__ __attribute__((overloadable)) short _Dtest(double *p) { return fpclassify(*p); } __DEVICE__ __CONSTEXPR__ __attribute__((overloadable)) short _FDtest(float *p) { return fpclassify(*p); } __DEVICE__ __CONSTEXPR__ __attribute__((overloadable)) double _Sinh(double x, double y) { return sinh(x) * y; } __DEVICE__ __CONSTEXPR__ __attribute__((overloadable)) float _FSinh(float x, float y) { return sinhf(x) * y; } #if defined(__cplusplus) } #endif // defined(__cplusplus) #endif // defined(_MSC_VER) #endif // !defined(__HIPCC_RTC__) #endif // ifndef __OPENMP_AMDGCN__ #pragma pop_macro("__DEVICE__") #pragma pop_macro("__CONSTEXPR__") #endif // __CLANG_HIP_CMATH_H__