clang/lib/Headers/__clang_hip_libdevice_declares.h

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