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