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 #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__