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