1/*===---- __clang_hip_libdevice_declares.h - HIP device library decls -------=== 2 * 3 * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. 4 * See https://llvm.org/LICENSE.txt for license information. 5 * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception 6 * 7 *===-----------------------------------------------------------------------=== 8 */ 9 10#ifndef __CLANG_HIP_LIBDEVICE_DECLARES_H__ 11#define __CLANG_HIP_LIBDEVICE_DECLARES_H__ 12 13extern "C" { 14 15// BEGIN FLOAT 16__device__ __attribute__((const)) float __ocml_acos_f32(float); 17__device__ __attribute__((pure)) float __ocml_acosh_f32(float); 18__device__ __attribute__((const)) float __ocml_asin_f32(float); 19__device__ __attribute__((pure)) float __ocml_asinh_f32(float); 20__device__ __attribute__((const)) float __ocml_atan2_f32(float, float); 21__device__ __attribute__((const)) float __ocml_atan_f32(float); 22__device__ __attribute__((pure)) float __ocml_atanh_f32(float); 23__device__ __attribute__((pure)) float __ocml_cbrt_f32(float); 24__device__ __attribute__((const)) float __ocml_ceil_f32(float); 25__device__ __attribute__((const)) __device__ float __ocml_copysign_f32(float, 26 float); 27__device__ float __ocml_cos_f32(float); 28__device__ float __ocml_native_cos_f32(float); 29__device__ __attribute__((pure)) __device__ float __ocml_cosh_f32(float); 30__device__ float __ocml_cospi_f32(float); 31__device__ float __ocml_i0_f32(float); 32__device__ float __ocml_i1_f32(float); 33__device__ __attribute__((pure)) float __ocml_erfc_f32(float); 34__device__ __attribute__((pure)) float __ocml_erfcinv_f32(float); 35__device__ __attribute__((pure)) float __ocml_erfcx_f32(float); 36__device__ __attribute__((pure)) float __ocml_erf_f32(float); 37__device__ __attribute__((pure)) float __ocml_erfinv_f32(float); 38__device__ __attribute__((pure)) float __ocml_exp10_f32(float); 39__device__ __attribute__((pure)) float __ocml_native_exp10_f32(float); 40__device__ __attribute__((pure)) float __ocml_exp2_f32(float); 41__device__ __attribute__((pure)) float __ocml_exp_f32(float); 42__device__ __attribute__((pure)) float __ocml_native_exp_f32(float); 43__device__ __attribute__((pure)) float __ocml_expm1_f32(float); 44__device__ __attribute__((const)) float __ocml_fabs_f32(float); 45__device__ __attribute__((const)) float __ocml_fdim_f32(float, float); 46__device__ __attribute__((const)) float __ocml_floor_f32(float); 47__device__ __attribute__((const)) float __ocml_fma_f32(float, float, float); 48__device__ __attribute__((const)) float __ocml_fmax_f32(float, float); 49__device__ __attribute__((const)) float __ocml_fmin_f32(float, float); 50__device__ __attribute__((const)) __device__ float __ocml_fmod_f32(float, 51 float); 52__device__ float __ocml_frexp_f32(float, 53 __attribute__((address_space(5))) int *); 54__device__ __attribute__((const)) float __ocml_hypot_f32(float, float); 55__device__ __attribute__((const)) int __ocml_ilogb_f32(float); 56__device__ __attribute__((const)) int __ocml_isfinite_f32(float); 57__device__ __attribute__((const)) int __ocml_isinf_f32(float); 58__device__ __attribute__((const)) int __ocml_isnan_f32(float); 59__device__ float __ocml_j0_f32(float); 60__device__ float __ocml_j1_f32(float); 61__device__ __attribute__((const)) float __ocml_ldexp_f32(float, int); 62__device__ float __ocml_lgamma_f32(float); 63__device__ __attribute__((pure)) float __ocml_log10_f32(float); 64__device__ __attribute__((pure)) float __ocml_native_log10_f32(float); 65__device__ __attribute__((pure)) float __ocml_log1p_f32(float); 66__device__ __attribute__((pure)) float __ocml_log2_f32(float); 67__device__ __attribute__((pure)) float __ocml_native_log2_f32(float); 68__device__ __attribute__((const)) float __ocml_logb_f32(float); 69__device__ __attribute__((pure)) float __ocml_log_f32(float); 70__device__ __attribute__((pure)) float __ocml_native_log_f32(float); 71__device__ float __ocml_modf_f32(float, 72 __attribute__((address_space(5))) float *); 73__device__ __attribute__((const)) float __ocml_nearbyint_f32(float); 74__device__ __attribute__((const)) float __ocml_nextafter_f32(float, float); 75__device__ __attribute__((const)) float __ocml_len3_f32(float, float, float); 76__device__ __attribute__((const)) float __ocml_len4_f32(float, float, float, 77 float); 78__device__ __attribute__((pure)) float __ocml_ncdf_f32(float); 79__device__ __attribute__((pure)) float __ocml_ncdfinv_f32(float); 80__device__ __attribute__((pure)) float __ocml_pow_f32(float, float); 81__device__ __attribute__((pure)) float __ocml_rcbrt_f32(float); 82__device__ __attribute__((const)) float __ocml_remainder_f32(float, float); 83__device__ float __ocml_remquo_f32(float, float, 84 __attribute__((address_space(5))) int *); 85__device__ __attribute__((const)) float __ocml_rhypot_f32(float, float); 86__device__ __attribute__((const)) float __ocml_rint_f32(float); 87__device__ __attribute__((const)) float __ocml_rlen3_f32(float, float, float); 88__device__ __attribute__((const)) float __ocml_rlen4_f32(float, float, float, 89 float); 90__device__ __attribute__((const)) float __ocml_round_f32(float); 91__device__ __attribute__((pure)) float __ocml_rsqrt_f32(float); 92__device__ __attribute__((const)) float __ocml_scalb_f32(float, float); 93__device__ __attribute__((const)) float __ocml_scalbn_f32(float, int); 94__device__ __attribute__((const)) int __ocml_signbit_f32(float); 95__device__ float __ocml_sincos_f32(float, 96 __attribute__((address_space(5))) float *); 97__device__ float __ocml_sincospi_f32(float, 98 __attribute__((address_space(5))) float *); 99__device__ float __ocml_sin_f32(float); 100__device__ float __ocml_native_sin_f32(float); 101__device__ __attribute__((pure)) float __ocml_sinh_f32(float); 102__device__ float __ocml_sinpi_f32(float); 103__device__ __attribute__((const)) float __ocml_sqrt_f32(float); 104__device__ __attribute__((const)) float __ocml_native_sqrt_f32(float); 105__device__ float __ocml_tan_f32(float); 106__device__ __attribute__((pure)) float __ocml_tanh_f32(float); 107__device__ float __ocml_tgamma_f32(float); 108__device__ __attribute__((const)) float __ocml_trunc_f32(float); 109__device__ float __ocml_y0_f32(float); 110__device__ float __ocml_y1_f32(float); 111 112// BEGIN INTRINSICS 113__device__ __attribute__((const)) float __ocml_add_rte_f32(float, float); 114__device__ __attribute__((const)) float __ocml_add_rtn_f32(float, float); 115__device__ __attribute__((const)) float __ocml_add_rtp_f32(float, float); 116__device__ __attribute__((const)) float __ocml_add_rtz_f32(float, float); 117__device__ __attribute__((const)) float __ocml_sub_rte_f32(float, float); 118__device__ __attribute__((const)) float __ocml_sub_rtn_f32(float, float); 119__device__ __attribute__((const)) float __ocml_sub_rtp_f32(float, float); 120__device__ __attribute__((const)) float __ocml_sub_rtz_f32(float, float); 121__device__ __attribute__((const)) float __ocml_mul_rte_f32(float, float); 122__device__ __attribute__((const)) float __ocml_mul_rtn_f32(float, float); 123__device__ __attribute__((const)) float __ocml_mul_rtp_f32(float, float); 124__device__ __attribute__((const)) float __ocml_mul_rtz_f32(float, float); 125__device__ __attribute__((const)) float __ocml_div_rte_f32(float, float); 126__device__ __attribute__((const)) float __ocml_div_rtn_f32(float, float); 127__device__ __attribute__((const)) float __ocml_div_rtp_f32(float, float); 128__device__ __attribute__((const)) float __ocml_div_rtz_f32(float, float); 129__device__ __attribute__((const)) float __ocml_sqrt_rte_f32(float, float); 130__device__ __attribute__((const)) float __ocml_sqrt_rtn_f32(float, float); 131__device__ __attribute__((const)) float __ocml_sqrt_rtp_f32(float, float); 132__device__ __attribute__((const)) float __ocml_sqrt_rtz_f32(float, float); 133__device__ __attribute__((const)) float __ocml_fma_rte_f32(float, float, float); 134__device__ __attribute__((const)) float __ocml_fma_rtn_f32(float, float, float); 135__device__ __attribute__((const)) float __ocml_fma_rtp_f32(float, float, float); 136__device__ __attribute__((const)) float __ocml_fma_rtz_f32(float, float, float); 137 138__device__ __attribute__((const)) float 139__llvm_amdgcn_cos_f32(float) __asm("llvm.amdgcn.cos.f32"); 140__device__ __attribute__((const)) float 141__llvm_amdgcn_rcp_f32(float) __asm("llvm.amdgcn.rcp.f32"); 142__device__ __attribute__((const)) float 143__llvm_amdgcn_rsq_f32(float) __asm("llvm.amdgcn.rsq.f32"); 144__device__ __attribute__((const)) float 145__llvm_amdgcn_sin_f32(float) __asm("llvm.amdgcn.sin.f32"); 146// END INTRINSICS 147// END FLOAT 148 149// BEGIN DOUBLE 150__device__ __attribute__((const)) double __ocml_acos_f64(double); 151__device__ __attribute__((pure)) double __ocml_acosh_f64(double); 152__device__ __attribute__((const)) double __ocml_asin_f64(double); 153__device__ __attribute__((pure)) double __ocml_asinh_f64(double); 154__device__ __attribute__((const)) double __ocml_atan2_f64(double, double); 155__device__ __attribute__((const)) double __ocml_atan_f64(double); 156__device__ __attribute__((pure)) double __ocml_atanh_f64(double); 157__device__ __attribute__((pure)) double __ocml_cbrt_f64(double); 158__device__ __attribute__((const)) double __ocml_ceil_f64(double); 159__device__ __attribute__((const)) double __ocml_copysign_f64(double, double); 160__device__ double __ocml_cos_f64(double); 161__device__ __attribute__((pure)) double __ocml_cosh_f64(double); 162__device__ double __ocml_cospi_f64(double); 163__device__ double __ocml_i0_f64(double); 164__device__ double __ocml_i1_f64(double); 165__device__ __attribute__((pure)) double __ocml_erfc_f64(double); 166__device__ __attribute__((pure)) double __ocml_erfcinv_f64(double); 167__device__ __attribute__((pure)) double __ocml_erfcx_f64(double); 168__device__ __attribute__((pure)) double __ocml_erf_f64(double); 169__device__ __attribute__((pure)) double __ocml_erfinv_f64(double); 170__device__ __attribute__((pure)) double __ocml_exp10_f64(double); 171__device__ __attribute__((pure)) double __ocml_exp2_f64(double); 172__device__ __attribute__((pure)) double __ocml_exp_f64(double); 173__device__ __attribute__((pure)) double __ocml_expm1_f64(double); 174__device__ __attribute__((const)) double __ocml_fabs_f64(double); 175__device__ __attribute__((const)) double __ocml_fdim_f64(double, double); 176__device__ __attribute__((const)) double __ocml_floor_f64(double); 177__device__ __attribute__((const)) double __ocml_fma_f64(double, double, double); 178__device__ __attribute__((const)) double __ocml_fmax_f64(double, double); 179__device__ __attribute__((const)) double __ocml_fmin_f64(double, double); 180__device__ __attribute__((const)) double __ocml_fmod_f64(double, double); 181__device__ double __ocml_frexp_f64(double, 182 __attribute__((address_space(5))) int *); 183__device__ __attribute__((const)) double __ocml_hypot_f64(double, double); 184__device__ __attribute__((const)) int __ocml_ilogb_f64(double); 185__device__ __attribute__((const)) int __ocml_isfinite_f64(double); 186__device__ __attribute__((const)) int __ocml_isinf_f64(double); 187__device__ __attribute__((const)) int __ocml_isnan_f64(double); 188__device__ double __ocml_j0_f64(double); 189__device__ double __ocml_j1_f64(double); 190__device__ __attribute__((const)) double __ocml_ldexp_f64(double, int); 191__device__ double __ocml_lgamma_f64(double); 192__device__ __attribute__((pure)) double __ocml_log10_f64(double); 193__device__ __attribute__((pure)) double __ocml_log1p_f64(double); 194__device__ __attribute__((pure)) double __ocml_log2_f64(double); 195__device__ __attribute__((const)) double __ocml_logb_f64(double); 196__device__ __attribute__((pure)) double __ocml_log_f64(double); 197__device__ double __ocml_modf_f64(double, 198 __attribute__((address_space(5))) double *); 199__device__ __attribute__((const)) double __ocml_nearbyint_f64(double); 200__device__ __attribute__((const)) double __ocml_nextafter_f64(double, double); 201__device__ __attribute__((const)) double __ocml_len3_f64(double, double, 202 double); 203__device__ __attribute__((const)) double __ocml_len4_f64(double, double, double, 204 double); 205__device__ __attribute__((pure)) double __ocml_ncdf_f64(double); 206__device__ __attribute__((pure)) double __ocml_ncdfinv_f64(double); 207__device__ __attribute__((pure)) double __ocml_pow_f64(double, double); 208__device__ __attribute__((pure)) double __ocml_rcbrt_f64(double); 209__device__ __attribute__((const)) double __ocml_remainder_f64(double, double); 210__device__ double __ocml_remquo_f64(double, double, 211 __attribute__((address_space(5))) int *); 212__device__ __attribute__((const)) double __ocml_rhypot_f64(double, double); 213__device__ __attribute__((const)) double __ocml_rint_f64(double); 214__device__ __attribute__((const)) double __ocml_rlen3_f64(double, double, 215 double); 216__device__ __attribute__((const)) double __ocml_rlen4_f64(double, double, 217 double, double); 218__device__ __attribute__((const)) double __ocml_round_f64(double); 219__device__ __attribute__((pure)) double __ocml_rsqrt_f64(double); 220__device__ __attribute__((const)) double __ocml_scalb_f64(double, double); 221__device__ __attribute__((const)) double __ocml_scalbn_f64(double, int); 222__device__ __attribute__((const)) int __ocml_signbit_f64(double); 223__device__ double __ocml_sincos_f64(double, 224 __attribute__((address_space(5))) double *); 225__device__ double 226__ocml_sincospi_f64(double, __attribute__((address_space(5))) double *); 227__device__ double __ocml_sin_f64(double); 228__device__ __attribute__((pure)) double __ocml_sinh_f64(double); 229__device__ double __ocml_sinpi_f64(double); 230__device__ __attribute__((const)) double __ocml_sqrt_f64(double); 231__device__ double __ocml_tan_f64(double); 232__device__ __attribute__((pure)) double __ocml_tanh_f64(double); 233__device__ double __ocml_tgamma_f64(double); 234__device__ __attribute__((const)) double __ocml_trunc_f64(double); 235__device__ double __ocml_y0_f64(double); 236__device__ double __ocml_y1_f64(double); 237 238// BEGIN INTRINSICS 239__device__ __attribute__((const)) double __ocml_add_rte_f64(double, double); 240__device__ __attribute__((const)) double __ocml_add_rtn_f64(double, double); 241__device__ __attribute__((const)) double __ocml_add_rtp_f64(double, double); 242__device__ __attribute__((const)) double __ocml_add_rtz_f64(double, double); 243__device__ __attribute__((const)) double __ocml_sub_rte_f64(double, double); 244__device__ __attribute__((const)) double __ocml_sub_rtn_f64(double, double); 245__device__ __attribute__((const)) double __ocml_sub_rtp_f64(double, double); 246__device__ __attribute__((const)) double __ocml_sub_rtz_f64(double, double); 247__device__ __attribute__((const)) double __ocml_mul_rte_f64(double, double); 248__device__ __attribute__((const)) double __ocml_mul_rtn_f64(double, double); 249__device__ __attribute__((const)) double __ocml_mul_rtp_f64(double, double); 250__device__ __attribute__((const)) double __ocml_mul_rtz_f64(double, double); 251__device__ __attribute__((const)) double __ocml_div_rte_f64(double, double); 252__device__ __attribute__((const)) double __ocml_div_rtn_f64(double, double); 253__device__ __attribute__((const)) double __ocml_div_rtp_f64(double, double); 254__device__ __attribute__((const)) double __ocml_div_rtz_f64(double, double); 255__device__ __attribute__((const)) double __ocml_sqrt_rte_f64(double, double); 256__device__ __attribute__((const)) double __ocml_sqrt_rtn_f64(double, double); 257__device__ __attribute__((const)) double __ocml_sqrt_rtp_f64(double, double); 258__device__ __attribute__((const)) double __ocml_sqrt_rtz_f64(double, double); 259__device__ __attribute__((const)) double __ocml_fma_rte_f64(double, double, 260 double); 261__device__ __attribute__((const)) double __ocml_fma_rtn_f64(double, double, 262 double); 263__device__ __attribute__((const)) double __ocml_fma_rtp_f64(double, double, 264 double); 265__device__ __attribute__((const)) double __ocml_fma_rtz_f64(double, double, 266 double); 267 268__device__ __attribute__((const)) double 269__llvm_amdgcn_rcp_f64(double) __asm("llvm.amdgcn.rcp.f64"); 270__device__ __attribute__((const)) double 271__llvm_amdgcn_rsq_f64(double) __asm("llvm.amdgcn.rsq.f64"); 272 273__device__ __attribute__((const)) _Float16 __ocml_ceil_f16(_Float16); 274__device__ _Float16 __ocml_cos_f16(_Float16); 275__device__ __attribute__((pure)) _Float16 __ocml_exp_f16(_Float16); 276__device__ __attribute__((pure)) _Float16 __ocml_exp10_f16(_Float16); 277__device__ __attribute__((pure)) _Float16 __ocml_exp2_f16(_Float16); 278__device__ __attribute__((const)) _Float16 __ocml_floor_f16(_Float16); 279__device__ __attribute__((const)) _Float16 __ocml_fma_f16(_Float16, _Float16, 280 _Float16); 281__device__ __attribute__((const)) _Float16 __ocml_fabs_f16(_Float16); 282__device__ __attribute__((const)) int __ocml_isinf_f16(_Float16); 283__device__ __attribute__((const)) int __ocml_isnan_f16(_Float16); 284__device__ __attribute__((pure)) _Float16 __ocml_log_f16(_Float16); 285__device__ __attribute__((pure)) _Float16 __ocml_log10_f16(_Float16); 286__device__ __attribute__((pure)) _Float16 __ocml_log2_f16(_Float16); 287__device__ __attribute__((const)) _Float16 __llvm_amdgcn_rcp_f16(_Float16); 288__device__ __attribute__((const)) _Float16 __ocml_rint_f16(_Float16); 289__device__ __attribute__((const)) _Float16 __ocml_rsqrt_f16(_Float16); 290__device__ _Float16 __ocml_sin_f16(_Float16); 291__device__ __attribute__((const)) _Float16 __ocml_sqrt_f16(_Float16); 292__device__ __attribute__((const)) _Float16 __ocml_trunc_f16(_Float16); 293 294typedef _Float16 __2f16 __attribute__((ext_vector_type(2))); 295typedef short __2i16 __attribute__((ext_vector_type(2))); 296 297__device__ __attribute__((const)) float __ockl_fdot2(__2f16 a, __2f16 b, 298 float c, bool s); 299__device__ __attribute__((const)) __2f16 __ocml_ceil_2f16(__2f16); 300__device__ __attribute__((const)) __2f16 __ocml_fabs_2f16(__2f16); 301__device__ __2f16 __ocml_cos_2f16(__2f16); 302__device__ __attribute__((pure)) __2f16 __ocml_exp_2f16(__2f16); 303__device__ __attribute__((pure)) __2f16 __ocml_exp10_2f16(__2f16); 304__device__ __attribute__((pure)) __2f16 __ocml_exp2_2f16(__2f16); 305__device__ __attribute__((const)) __2f16 __ocml_floor_2f16(__2f16); 306__device__ __attribute__((const)) 307__2f16 __ocml_fma_2f16(__2f16, __2f16, __2f16); 308__device__ __attribute__((const)) __2i16 __ocml_isinf_2f16(__2f16); 309__device__ __attribute__((const)) __2i16 __ocml_isnan_2f16(__2f16); 310__device__ __attribute__((pure)) __2f16 __ocml_log_2f16(__2f16); 311__device__ __attribute__((pure)) __2f16 __ocml_log10_2f16(__2f16); 312__device__ __attribute__((pure)) __2f16 __ocml_log2_2f16(__2f16); 313__device__ inline __2f16 314__llvm_amdgcn_rcp_2f16(__2f16 __x) // Not currently exposed by ROCDL. 315{ 316 return __2f16{__llvm_amdgcn_rcp_f16(__x.x), __llvm_amdgcn_rcp_f16(__x.y)}; 317} 318__device__ __attribute__((const)) __2f16 __ocml_rint_2f16(__2f16); 319__device__ __attribute__((const)) __2f16 __ocml_rsqrt_2f16(__2f16); 320__device__ __2f16 __ocml_sin_2f16(__2f16); 321__device__ __attribute__((const)) __2f16 __ocml_sqrt_2f16(__2f16); 322__device__ __attribute__((const)) __2f16 __ocml_trunc_2f16(__2f16); 323 324} // extern "C" 325 326#endif // __CLANG_HIP_LIBDEVICE_DECLARES_H__ 327