File Manager

X7ROOT File Manager

Current Path: /usr/lib64/llvm17/lib/clang/17/include

usr / lib64 / llvm17 / lib / clang / 17 / include /

ðŸ“ ..
ðŸ“„ __clang_cuda_builtin_vars.h(4.78 KB)
ðŸ“„ __clang_cuda_cmath.h(18.06 KB)
ðŸ“„ __clang_cuda_complex_builtins.h(9.36 KB)
ðŸ“„ __clang_cuda_device_functions.h(56.68 KB)
ðŸ“„ __clang_cuda_intrinsics.h(29.93 KB)
ðŸ“„ __clang_cuda_libdevice_declares.h(21.87 KB)
ðŸ“„ __clang_cuda_math.h(15.99 KB)
ðŸ“„ __clang_cuda_math_forward_declares.h(8.27 KB)
ðŸ“„ __clang_cuda_runtime_wrapper.h(17.61 KB)
ðŸ“„ __clang_cuda_texture_intrinsics.h(31.86 KB)
ðŸ“„ __clang_hip_cmath.h(26.34 KB)
ðŸ“„ __clang_hip_libdevice_declares.h(19.87 KB)
ðŸ“„ __clang_hip_math.h(31.96 KB)
ðŸ“„ __clang_hip_runtime_wrapper.h(4.65 KB)
ðŸ“„ __clang_hip_stdlib.h(1.19 KB)
ðŸ“„ __stddef_max_align_t.h(857 B)
ðŸ“„ __wmmintrin_aes.h(5.15 KB)
ðŸ“„ __wmmintrin_pclmul.h(1.99 KB)
ðŸ“„ adxintrin.h(7.37 KB)
ðŸ“„ altivec.h(697.32 KB)
ðŸ“„ ammintrin.h(7.54 KB)
ðŸ“„ amxcomplexintrin.h(6.81 KB)
ðŸ“„ amxfp16intrin.h(1.82 KB)
ðŸ“„ amxintrin.h(21.12 KB)
ðŸ“„ arm64intr.h(993 B)
ðŸ“„ arm_acle.h(25.66 KB)
ðŸ“„ arm_bf16.h(548 B)
ðŸ“„ arm_cde.h(32.67 KB)
ðŸ“„ arm_cmse.h(6.21 KB)
ðŸ“„ arm_fp16.h(16.92 KB)
ðŸ“„ arm_mve.h(1.48 MB)
ðŸ“„ arm_neon.h(2.45 MB)
ðŸ“„ arm_neon_sve_bridge.h(9.48 KB)
ðŸ“„ arm_sme_draft_spec_subject_to_change.h(60.2 KB)
ðŸ“„ arm_sve.h(1.51 MB)
ðŸ“„ armintr.h(843 B)
ðŸ“„ avx2intrin.h(186.96 KB)
ðŸ“„ avx512bf16intrin.h(10.51 KB)
ðŸ“„ avx512bitalgintrin.h(2.41 KB)
ðŸ“„ avx512bwintrin.h(75.33 KB)
ðŸ“„ avx512cdintrin.h(4.12 KB)
ðŸ“„ avx512dqintrin.h(58.75 KB)
ðŸ“„ avx512erintrin.h(11.83 KB)
ðŸ“„ avx512fintrin.h(382.64 KB)
ðŸ“„ avx512fp16intrin.h(156.63 KB)
ðŸ“„ avx512ifmaintrin.h(2.49 KB)
ðŸ“„ avx512ifmavlintrin.h(4.31 KB)
ðŸ“„ avx512pfintrin.h(4.53 KB)
ðŸ“„ avx512vbmi2intrin.h(13.17 KB)
ðŸ“„ avx512vbmiintrin.h(3.72 KB)
ðŸ“„ avx512vbmivlintrin.h(6.94 KB)
ðŸ“„ avx512vlbf16intrin.h(19.21 KB)
ðŸ“„ avx512vlbitalgintrin.h(4.23 KB)
ðŸ“„ avx512vlbwintrin.h(121.26 KB)
ðŸ“„ avx512vlcdintrin.h(7.66 KB)
ðŸ“„ avx512vldqintrin.h(46.41 KB)
ðŸ“„ avx512vlfp16intrin.h(85.51 KB)
ðŸ“„ avx512vlintrin.h(322.29 KB)
ðŸ“„ avx512vlvbmi2intrin.h(25.72 KB)
ðŸ“„ avx512vlvnniintrin.h(13.13 KB)
ðŸ“„ avx512vlvp2intersectintrin.h(4.44 KB)
ðŸ“„ avx512vnniintrin.h(4.21 KB)
ðŸ“„ avx512vp2intersectintrin.h(2.9 KB)
ðŸ“„ avx512vpopcntdqintrin.h(2 KB)
ðŸ“„ avx512vpopcntdqvlintrin.h(3.31 KB)
ðŸ“„ avxifmaintrin.h(5.75 KB)
ðŸ“„ avxintrin.h(195.41 KB)
ðŸ“„ avxneconvertintrin.h(14.09 KB)
ðŸ“„ avxvnniint16intrin.h(17.41 KB)
ðŸ“„ avxvnniint8intrin.h(18.67 KB)
ðŸ“„ avxvnniintrin.h(10.44 KB)
ðŸ“„ bmi2intrin.h(7.09 KB)
ðŸ“„ bmiintrin.h(14.12 KB)
ðŸ“„ builtins.h(741 B)
ðŸ“„ cet.h(1.49 KB)
ðŸ“„ cetintrin.h(3.27 KB)
ðŸ“„ cldemoteintrin.h(1.18 KB)
ðŸ“„ clflushoptintrin.h(1.17 KB)
ðŸ“„ clwbintrin.h(1.2 KB)
ðŸ“„ clzerointrin.h(1.19 KB)
ðŸ“„ cmpccxaddintrin.h(2.33 KB)
ðŸ“„ cpuid.h(11.01 KB)
ðŸ“„ crc32intrin.h(3.27 KB)
ðŸ“ cuda_wrappers
ðŸ“„ emmintrin.h(192.64 KB)
ðŸ“„ enqcmdintrin.h(2.12 KB)
ðŸ“„ f16cintrin.h(5.39 KB)
ðŸ“„ float.h(5.63 KB)
ðŸ“„ fma4intrin.h(6.82 KB)
ðŸ“„ fmaintrin.h(28.4 KB)
ðŸ“„ fxsrintrin.h(2.82 KB)
ðŸ“„ gfniintrin.h(7.57 KB)
ðŸ“„ hexagon_circ_brev_intrinsics.h(15.59 KB)
ðŸ“„ hexagon_protos.h(374.42 KB)
ðŸ“„ hexagon_types.h(130.33 KB)
ðŸ“„ hresetintrin.h(1.36 KB)
ðŸ“„ htmintrin.h(6.14 KB)
ðŸ“„ htmxlintrin.h(9.01 KB)
ðŸ“„ hvx_hexagon_protos.h(254.26 KB)
ðŸ“„ ia32intrin.h(12.72 KB)
ðŸ“„ immintrin.h(23.57 KB)
ðŸ“„ intrin.h(28.22 KB)
ðŸ“„ inttypes.h(2.26 KB)
ðŸ“„ invpcidintrin.h(764 B)
ðŸ“„ iso646.h(656 B)
ðŸ“„ keylockerintrin.h(17.98 KB)
ðŸ“„ larchintrin.h(7.8 KB)
ðŸ“„ limits.h(3.61 KB)
ðŸ“ llvm_libc_wrappers
ðŸ“„ lwpintrin.h(5 KB)
ðŸ“„ lzcntintrin.h(3.18 KB)
ðŸ“„ mm3dnow.h(4.5 KB)
ðŸ“„ mm_malloc.h(1.88 KB)
ðŸ“„ mmintrin.h(55.98 KB)
ðŸ“„ module.modulemap(3.33 KB)
ðŸ“„ movdirintrin.h(1.57 KB)
ðŸ“„ msa.h(25.01 KB)
ðŸ“„ mwaitxintrin.h(2.19 KB)
ðŸ“„ nmmintrin.h(709 B)
ðŸ“„ opencl-c-base.h(30.38 KB)
ðŸ“„ opencl-c.h(874.39 KB)
ðŸ“ openmp_wrappers
ðŸ“„ pconfigintrin.h(1.19 KB)
ðŸ“„ pkuintrin.h(934 B)
ðŸ“„ pmmintrin.h(10.5 KB)
ðŸ“„ popcntintrin.h(1.82 KB)
ðŸ“ ppc_wrappers
ðŸ“„ prfchiintrin.h(2.02 KB)
ðŸ“„ prfchwintrin.h(2.06 KB)
ðŸ“„ ptwriteintrin.h(1.05 KB)
ðŸ“„ raointintrin.h(6.59 KB)
ðŸ“„ rdpruintrin.h(1.59 KB)
ðŸ“„ rdseedintrin.h(2.85 KB)
ðŸ“„ riscv_ntlh.h(855 B)
ðŸ“„ rtmintrin.h(1.25 KB)
ðŸ“„ s390intrin.h(604 B)
ðŸ“„ serializeintrin.h(881 B)
ðŸ“„ sgxintrin.h(1.77 KB)
ðŸ“„ sha512intrin.h(5.95 KB)
ðŸ“„ shaintrin.h(7.37 KB)
ðŸ“„ sifive_vector.h(522 B)
ðŸ“„ sm3intrin.h(7.29 KB)
ðŸ“„ sm4intrin.h(8.2 KB)
ðŸ“„ smmintrin.h(99.32 KB)
ðŸ“„ stdalign.h(911 B)
ðŸ“„ stdarg.h(1.66 KB)
ðŸ“„ stdatomic.h(8.3 KB)
ðŸ“„ stdbool.h(1.04 KB)
ðŸ“„ stddef.h(4.16 KB)
ðŸ“„ stdint.h(32.49 KB)
ðŸ“„ stdnoreturn.h(1.17 KB)
ðŸ“„ tbmintrin.h(3.15 KB)
ðŸ“„ tgmath.h(29.68 KB)
ðŸ“„ tmmintrin.h(29.51 KB)
ðŸ“„ tsxldtrkintrin.h(1.97 KB)
ðŸ“„ uintrintrin.h(4.96 KB)
ðŸ“„ unwind.h(11.21 KB)
ðŸ“„ vadefs.h(1.39 KB)
ðŸ“„ vaesintrin.h(2.46 KB)
ðŸ“„ varargs.h(477 B)
ðŸ“„ vecintrin.h(360.82 KB)
ðŸ“„ velintrin.h(2.1 KB)
ðŸ“„ velintrin_approx.h(3.54 KB)
ðŸ“„ velintrin_gen.h(69.06 KB)
ðŸ“„ vpclmulqdqintrin.h(1.06 KB)
ðŸ“„ waitpkgintrin.h(1.33 KB)
ðŸ“„ wasm_simd128.h(76.25 KB)
ðŸ“„ wbnoinvdintrin.h(749 B)
ðŸ“„ wmmintrin.h(659 B)
ðŸ“„ x86gprintrin.h(2.32 KB)
ðŸ“„ x86intrin.h(1.81 KB)
ðŸ“„ xmmintrin.h(106.73 KB)
ðŸ“„ xopintrin.h(19.96 KB)
ðŸ“„ xsavecintrin.h(2.51 KB)
ðŸ“„ xsaveintrin.h(1.64 KB)
ðŸ“„ xsaveoptintrin.h(1 KB)
ðŸ“„ xsavesintrin.h(1.24 KB)
ðŸ“„ xtestintrin.h(873 B)

Editing: __clang_hip_math.h

/*===---- __clang_hip_math.h - Device-side HIP math support ----------------===
 *
 * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 * See https://llvm.org/LICENSE.txt for license information.
 * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 *
 *===-----------------------------------------------------------------------===
 */
#ifndef __CLANG_HIP_MATH_H__
#define __CLANG_HIP_MATH_H__

#if !defined(__HIP__) && !defined(__OPENMP_AMDGCN__)
#error "This file is for HIP and OpenMP AMDGCN device compilation only."
#endif

#if !defined(__HIPCC_RTC__)
#if defined(__cplusplus)
#include <algorithm>
#endif
#include <limits.h>
#include <stdint.h>
#ifdef __OPENMP_AMDGCN__
#include <omp.h>
#endif
#endif // !defined(__HIPCC_RTC__)

#pragma push_macro("__DEVICE__")

#ifdef __OPENMP_AMDGCN__
#define __DEVICE__ static inline __attribute__((always_inline, nothrow))
#else
#define __DEVICE__ static __device__ inline __attribute__((always_inline))
#endif

// A few functions return bool type starting only in C++11.
#pragma push_macro("__RETURN_TYPE")
#ifdef __OPENMP_AMDGCN__
#define __RETURN_TYPE int
#else
#if defined(__cplusplus)
#define __RETURN_TYPE bool
#else
#define __RETURN_TYPE int
#endif
#endif // __OPENMP_AMDGCN__

#if defined (__cplusplus) && __cplusplus < 201103L
// emulate static_assert on type sizes
template<bool>
struct __compare_result{};
template<>
struct __compare_result<true> {
  static const __device__ bool valid;
};

__DEVICE__
void __suppress_unused_warning(bool b){};
template <unsigned int S, unsigned int T>
__DEVICE__ void __static_assert_equal_size() {
  __suppress_unused_warning(__compare_result<S == T>::valid);
}

#define __static_assert_type_size_equal(A, B) \
  __static_assert_equal_size<A,B>()

#else
#define __static_assert_type_size_equal(A,B) \
  static_assert((A) == (B), "")

#endif

__DEVICE__
uint64_t __make_mantissa_base8(const char *__tagp __attribute__((nonnull))) {
  uint64_t __r = 0;
  while (*__tagp != '\0') {
    char __tmp = *__tagp;

if (__tmp >= '0' && __tmp <= '7')
      __r = (__r * 8u) + __tmp - '0';
    else
      return 0;

++__tagp;
  }

return __r;
}

__DEVICE__
uint64_t __make_mantissa_base10(const char *__tagp __attribute__((nonnull))) {
  uint64_t __r = 0;
  while (*__tagp != '\0') {
    char __tmp = *__tagp;

if (__tmp >= '0' && __tmp <= '9')
      __r = (__r * 10u) + __tmp - '0';
    else
      return 0;

++__tagp;
  }

return __r;
}

__DEVICE__
uint64_t __make_mantissa_base16(const char *__tagp __attribute__((nonnull))) {
  uint64_t __r = 0;
  while (*__tagp != '\0') {
    char __tmp = *__tagp;

if (__tmp >= '0' && __tmp <= '9')
      __r = (__r * 16u) + __tmp - '0';
    else if (__tmp >= 'a' && __tmp <= 'f')
      __r = (__r * 16u) + __tmp - 'a' + 10;
    else if (__tmp >= 'A' && __tmp <= 'F')
      __r = (__r * 16u) + __tmp - 'A' + 10;
    else
      return 0;

++__tagp;
  }

return __r;
}

__DEVICE__
uint64_t __make_mantissa(const char *__tagp __attribute__((nonnull))) {
  if (*__tagp == '0') {
    ++__tagp;

if (*__tagp == 'x' || *__tagp == 'X')
      return __make_mantissa_base16(__tagp);
    else
      return __make_mantissa_base8(__tagp);
  }

return __make_mantissa_base10(__tagp);
}

// BEGIN FLOAT
#if defined(__cplusplus)
__DEVICE__
int abs(int __x) {
  int __sgn = __x >> (sizeof(int) * CHAR_BIT - 1);
  return (__x ^ __sgn) - __sgn;
}
__DEVICE__
long labs(long __x) {
  long __sgn = __x >> (sizeof(long) * CHAR_BIT - 1);
  return (__x ^ __sgn) - __sgn;
}
__DEVICE__
long long llabs(long long __x) {
  long long __sgn = __x >> (sizeof(long long) * CHAR_BIT - 1);
  return (__x ^ __sgn) - __sgn;
}
#endif

__DEVICE__
float acosf(float __x) { return __ocml_acos_f32(__x); }

__DEVICE__
float acoshf(float __x) { return __ocml_acosh_f32(__x); }

__DEVICE__
float asinf(float __x) { return __ocml_asin_f32(__x); }

__DEVICE__
float asinhf(float __x) { return __ocml_asinh_f32(__x); }

__DEVICE__
float atan2f(float __x, float __y) { return __ocml_atan2_f32(__x, __y); }

__DEVICE__
float atanf(float __x) { return __ocml_atan_f32(__x); }

__DEVICE__
float atanhf(float __x) { return __ocml_atanh_f32(__x); }

__DEVICE__
float cbrtf(float __x) { return __ocml_cbrt_f32(__x); }

__DEVICE__
float ceilf(float __x) { return __builtin_ceilf(__x); }

__DEVICE__
float copysignf(float __x, float __y) { return __builtin_copysignf(__x, __y); }

__DEVICE__
float cosf(float __x) { return __ocml_cos_f32(__x); }

__DEVICE__
float coshf(float __x) { return __ocml_cosh_f32(__x); }

__DEVICE__
float cospif(float __x) { return __ocml_cospi_f32(__x); }

__DEVICE__
float cyl_bessel_i0f(float __x) { return __ocml_i0_f32(__x); }

__DEVICE__
float cyl_bessel_i1f(float __x) { return __ocml_i1_f32(__x); }

__DEVICE__
float erfcf(float __x) { return __ocml_erfc_f32(__x); }

__DEVICE__
float erfcinvf(float __x) { return __ocml_erfcinv_f32(__x); }

__DEVICE__
float erfcxf(float __x) { return __ocml_erfcx_f32(__x); }

__DEVICE__
float erff(float __x) { return __ocml_erf_f32(__x); }

__DEVICE__
float erfinvf(float __x) { return __ocml_erfinv_f32(__x); }

__DEVICE__
float exp10f(float __x) { return __ocml_exp10_f32(__x); }

__DEVICE__
float exp2f(float __x) { return __builtin_exp2f(__x); }

__DEVICE__
float expf(float __x) { return __builtin_expf(__x); }

__DEVICE__
float expm1f(float __x) { return __ocml_expm1_f32(__x); }

__DEVICE__
float fabsf(float __x) { return __builtin_fabsf(__x); }

__DEVICE__
float fdimf(float __x, float __y) { return __ocml_fdim_f32(__x, __y); }

__DEVICE__
float fdividef(float __x, float __y) { return __x / __y; }

__DEVICE__
float floorf(float __x) { return __builtin_floorf(__x); }

__DEVICE__
float fmaf(float __x, float __y, float __z) {
  return __builtin_fmaf(__x, __y, __z);
}

__DEVICE__
float fmaxf(float __x, float __y) { return __builtin_fmaxf(__x, __y); }

__DEVICE__
float fminf(float __x, float __y) { return __builtin_fminf(__x, __y); }

__DEVICE__
float fmodf(float __x, float __y) { return __ocml_fmod_f32(__x, __y); }

__DEVICE__
float frexpf(float __x, int *__nptr) {
  return __builtin_frexpf(__x, __nptr);
}

__DEVICE__
float hypotf(float __x, float __y) { return __ocml_hypot_f32(__x, __y); }

__DEVICE__
int ilogbf(float __x) { return __ocml_ilogb_f32(__x); }

__DEVICE__
__RETURN_TYPE __finitef(float __x) { return __builtin_isfinite(__x); }

__DEVICE__
__RETURN_TYPE __isinff(float __x) { return __builtin_isinf(__x); }

__DEVICE__
__RETURN_TYPE __isnanf(float __x) { return __builtin_isnan(__x); }

__DEVICE__
float j0f(float __x) { return __ocml_j0_f32(__x); }

__DEVICE__
float j1f(float __x) { return __ocml_j1_f32(__x); }

__DEVICE__
float jnf(int __n, float __x) { // TODO: we could use Ahmes multiplication
                                // and the Miller & Brown algorithm
  //       for linear recurrences to get O(log n) steps, but it's unclear if
  //       it'd be beneficial in this case.
  if (__n == 0)
    return j0f(__x);
  if (__n == 1)
    return j1f(__x);

float __x0 = j0f(__x);
  float __x1 = j1f(__x);
  for (int __i = 1; __i < __n; ++__i) {
    float __x2 = (2 * __i) / __x * __x1 - __x0;
    __x0 = __x1;
    __x1 = __x2;
  }

return __x1;
}

__DEVICE__
float ldexpf(float __x, int __e) { return __builtin_amdgcn_ldexpf(__x, __e); }

__DEVICE__
float lgammaf(float __x) { return __ocml_lgamma_f32(__x); }

__DEVICE__
long long int llrintf(float __x) { return __builtin_rintf(__x); }

__DEVICE__
long long int llroundf(float __x) { return __builtin_roundf(__x); }

__DEVICE__
float log10f(float __x) { return __builtin_log10f(__x); }

__DEVICE__
float log1pf(float __x) { return __ocml_log1p_f32(__x); }

__DEVICE__
float log2f(float __x) { return __builtin_log2f(__x); }

__DEVICE__
float logbf(float __x) { return __ocml_logb_f32(__x); }

__DEVICE__
float logf(float __x) { return __builtin_logf(__x); }

__DEVICE__
long int lrintf(float __x) { return __builtin_rintf(__x); }

__DEVICE__
long int lroundf(float __x) { return __builtin_roundf(__x); }

__DEVICE__
float modff(float __x, float *__iptr) {
  float __tmp;
#ifdef __OPENMP_AMDGCN__
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
#endif
  float __r =
      __ocml_modf_f32(__x, (__attribute__((address_space(5))) float *)&__tmp);
  *__iptr = __tmp;
  return __r;
}

__DEVICE__
float nanf(const char *__tagp __attribute__((nonnull))) {
  union {
    float val;
    struct ieee_float {
      unsigned int mantissa : 22;
      unsigned int quiet : 1;
      unsigned int exponent : 8;
      unsigned int sign : 1;
    } bits;
  } __tmp;
  __static_assert_type_size_equal(sizeof(__tmp.val), sizeof(__tmp.bits));

__tmp.bits.sign = 0u;
  __tmp.bits.exponent = ~0u;
  __tmp.bits.quiet = 1u;
  __tmp.bits.mantissa = __make_mantissa(__tagp);

return __tmp.val;
}

__DEVICE__
float nearbyintf(float __x) { return __builtin_nearbyintf(__x); }

__DEVICE__
float nextafterf(float __x, float __y) {
  return __ocml_nextafter_f32(__x, __y);
}

__DEVICE__
float norm3df(float __x, float __y, float __z) {
  return __ocml_len3_f32(__x, __y, __z);
}

__DEVICE__
float norm4df(float __x, float __y, float __z, float __w) {
  return __ocml_len4_f32(__x, __y, __z, __w);
}

__DEVICE__
float normcdff(float __x) { return __ocml_ncdf_f32(__x); }

__DEVICE__
float normcdfinvf(float __x) { return __ocml_ncdfinv_f32(__x); }

__DEVICE__
float normf(int __dim,
            const float *__a) { // TODO: placeholder until OCML adds support.
  float __r = 0;
  while (__dim--) {
    __r += __a[0] * __a[0];
    ++__a;
  }

return __ocml_sqrt_f32(__r);
}

__DEVICE__
float powf(float __x, float __y) { return __ocml_pow_f32(__x, __y); }

__DEVICE__
float powif(float __x, int __y) { return __ocml_pown_f32(__x, __y); }

__DEVICE__
float rcbrtf(float __x) { return __ocml_rcbrt_f32(__x); }

__DEVICE__
float remainderf(float __x, float __y) {
  return __ocml_remainder_f32(__x, __y);
}

__DEVICE__
float remquof(float __x, float __y, int *__quo) {
  int __tmp;
#ifdef __OPENMP_AMDGCN__
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
#endif
  float __r = __ocml_remquo_f32(
      __x, __y, (__attribute__((address_space(5))) int *)&__tmp);
  *__quo = __tmp;

return __r;
}

__DEVICE__
float rhypotf(float __x, float __y) { return __ocml_rhypot_f32(__x, __y); }

__DEVICE__
float rintf(float __x) { return __builtin_rintf(__x); }

__DEVICE__
float rnorm3df(float __x, float __y, float __z) {
  return __ocml_rlen3_f32(__x, __y, __z);
}

__DEVICE__
float rnorm4df(float __x, float __y, float __z, float __w) {
  return __ocml_rlen4_f32(__x, __y, __z, __w);
}

__DEVICE__
float rnormf(int __dim,
             const float *__a) { // TODO: placeholder until OCML adds support.
  float __r = 0;
  while (__dim--) {
    __r += __a[0] * __a[0];
    ++__a;
  }

return __ocml_rsqrt_f32(__r);
}

__DEVICE__
float roundf(float __x) { return __builtin_roundf(__x); }

__DEVICE__
float rsqrtf(float __x) { return __ocml_rsqrt_f32(__x); }

__DEVICE__
float scalblnf(float __x, long int __n) {
  return (__n < INT_MAX) ? __builtin_amdgcn_ldexpf(__x, __n)
                         : __ocml_scalb_f32(__x, __n);
}

__DEVICE__
float scalbnf(float __x, int __n) { return __builtin_amdgcn_ldexpf(__x, __n); }

__DEVICE__
__RETURN_TYPE __signbitf(float __x) { return __builtin_signbitf(__x); }

__DEVICE__
void sincosf(float __x, float *__sinptr, float *__cosptr) {
  float __tmp;
#ifdef __OPENMP_AMDGCN__
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
#endif
  *__sinptr =
      __ocml_sincos_f32(__x, (__attribute__((address_space(5))) float *)&__tmp);
  *__cosptr = __tmp;
}

__DEVICE__
void sincospif(float __x, float *__sinptr, float *__cosptr) {
  float __tmp;
#ifdef __OPENMP_AMDGCN__
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
#endif
  *__sinptr = __ocml_sincospi_f32(
      __x, (__attribute__((address_space(5))) float *)&__tmp);
  *__cosptr = __tmp;
}

__DEVICE__
float sinf(float __x) { return __ocml_sin_f32(__x); }

__DEVICE__
float sinhf(float __x) { return __ocml_sinh_f32(__x); }

__DEVICE__
float sinpif(float __x) { return __ocml_sinpi_f32(__x); }

__DEVICE__
float sqrtf(float __x) { return __ocml_sqrt_f32(__x); }

__DEVICE__
float tanf(float __x) { return __ocml_tan_f32(__x); }

__DEVICE__
float tanhf(float __x) { return __ocml_tanh_f32(__x); }

__DEVICE__
float tgammaf(float __x) { return __ocml_tgamma_f32(__x); }

__DEVICE__
float truncf(float __x) { return __builtin_truncf(__x); }

__DEVICE__
float y0f(float __x) { return __ocml_y0_f32(__x); }

__DEVICE__
float y1f(float __x) { return __ocml_y1_f32(__x); }

__DEVICE__
float ynf(int __n, float __x) { // TODO: we could use Ahmes multiplication
                                // and the Miller & Brown algorithm
  //       for linear recurrences to get O(log n) steps, but it's unclear if
  //       it'd be beneficial in this case. Placeholder until OCML adds
  //       support.
  if (__n == 0)
    return y0f(__x);
  if (__n == 1)
    return y1f(__x);

float __x0 = y0f(__x);
  float __x1 = y1f(__x);
  for (int __i = 1; __i < __n; ++__i) {
    float __x2 = (2 * __i) / __x * __x1 - __x0;
    __x0 = __x1;
    __x1 = __x2;
  }

return __x1;
}

// BEGIN INTRINSICS

__DEVICE__
float __cosf(float __x) { return __ocml_native_cos_f32(__x); }

__DEVICE__
float __exp10f(float __x) { return __ocml_native_exp10_f32(__x); }

__DEVICE__
float __expf(float __x) { return __ocml_native_exp_f32(__x); }

#if defined OCML_BASIC_ROUNDED_OPERATIONS
__DEVICE__
float __fadd_rd(float __x, float __y) { return __ocml_add_rtn_f32(__x, __y); }
__DEVICE__
float __fadd_rn(float __x, float __y) { return __ocml_add_rte_f32(__x, __y); }
__DEVICE__
float __fadd_ru(float __x, float __y) { return __ocml_add_rtp_f32(__x, __y); }
__DEVICE__
float __fadd_rz(float __x, float __y) { return __ocml_add_rtz_f32(__x, __y); }
#else
__DEVICE__
float __fadd_rn(float __x, float __y) { return __x + __y; }
#endif

#if defined OCML_BASIC_ROUNDED_OPERATIONS
__DEVICE__
float __fdiv_rd(float __x, float __y) { return __ocml_div_rtn_f32(__x, __y); }
__DEVICE__
float __fdiv_rn(float __x, float __y) { return __ocml_div_rte_f32(__x, __y); }
__DEVICE__
float __fdiv_ru(float __x, float __y) { return __ocml_div_rtp_f32(__x, __y); }
__DEVICE__
float __fdiv_rz(float __x, float __y) { return __ocml_div_rtz_f32(__x, __y); }
#else
__DEVICE__
float __fdiv_rn(float __x, float __y) { return __x / __y; }
#endif

__DEVICE__
float __fdividef(float __x, float __y) { return __x / __y; }

#if defined OCML_BASIC_ROUNDED_OPERATIONS
__DEVICE__
float __fmaf_rd(float __x, float __y, float __z) {
  return __ocml_fma_rtn_f32(__x, __y, __z);
}
__DEVICE__
float __fmaf_rn(float __x, float __y, float __z) {
  return __ocml_fma_rte_f32(__x, __y, __z);
}
__DEVICE__
float __fmaf_ru(float __x, float __y, float __z) {
  return __ocml_fma_rtp_f32(__x, __y, __z);
}
__DEVICE__
float __fmaf_rz(float __x, float __y, float __z) {
  return __ocml_fma_rtz_f32(__x, __y, __z);
}
#else
__DEVICE__
float __fmaf_rn(float __x, float __y, float __z) {
  return __builtin_fmaf(__x, __y, __z);
}
#endif

#if defined OCML_BASIC_ROUNDED_OPERATIONS
__DEVICE__
float __fmul_rd(float __x, float __y) { return __ocml_mul_rtn_f32(__x, __y); }
__DEVICE__
float __fmul_rn(float __x, float __y) { return __ocml_mul_rte_f32(__x, __y); }
__DEVICE__
float __fmul_ru(float __x, float __y) { return __ocml_mul_rtp_f32(__x, __y); }
__DEVICE__
float __fmul_rz(float __x, float __y) { return __ocml_mul_rtz_f32(__x, __y); }
#else
__DEVICE__
float __fmul_rn(float __x, float __y) { return __x * __y; }
#endif

#if defined OCML_BASIC_ROUNDED_OPERATIONS
__DEVICE__
float __frcp_rd(float __x) { return __ocml_div_rtn_f32(1.0f, __x); }
__DEVICE__
float __frcp_rn(float __x) { return __ocml_div_rte_f32(1.0f, __x); }
__DEVICE__
float __frcp_ru(float __x) { return __ocml_div_rtp_f32(1.0f, __x); }
__DEVICE__
float __frcp_rz(float __x) { return __ocml_div_rtz_f32(1.0f, __x); }
#else
__DEVICE__
float __frcp_rn(float __x) { return 1.0f / __x; }
#endif

__DEVICE__
float __frsqrt_rn(float __x) { return __builtin_amdgcn_rsqf(__x); }

#if defined OCML_BASIC_ROUNDED_OPERATIONS
__DEVICE__
float __fsqrt_rd(float __x) { return __ocml_sqrt_rtn_f32(__x); }
__DEVICE__
float __fsqrt_rn(float __x) { return __ocml_sqrt_rte_f32(__x); }
__DEVICE__
float __fsqrt_ru(float __x) { return __ocml_sqrt_rtp_f32(__x); }
__DEVICE__
float __fsqrt_rz(float __x) { return __ocml_sqrt_rtz_f32(__x); }
#else
__DEVICE__
float __fsqrt_rn(float __x) { return __ocml_native_sqrt_f32(__x); }
#endif

#if defined OCML_BASIC_ROUNDED_OPERATIONS
__DEVICE__
float __fsub_rd(float __x, float __y) { return __ocml_sub_rtn_f32(__x, __y); }
__DEVICE__
float __fsub_rn(float __x, float __y) { return __ocml_sub_rte_f32(__x, __y); }
__DEVICE__
float __fsub_ru(float __x, float __y) { return __ocml_sub_rtp_f32(__x, __y); }
__DEVICE__
float __fsub_rz(float __x, float __y) { return __ocml_sub_rtz_f32(__x, __y); }
#else
__DEVICE__
float __fsub_rn(float __x, float __y) { return __x - __y; }
#endif

__DEVICE__
float __log10f(float __x) { return __ocml_native_log10_f32(__x); }

__DEVICE__
float __log2f(float __x) { return __ocml_native_log2_f32(__x); }

__DEVICE__
float __logf(float __x) { return __ocml_native_log_f32(__x); }

__DEVICE__
float __powf(float __x, float __y) { return __ocml_pow_f32(__x, __y); }

__DEVICE__
float __saturatef(float __x) { return (__x < 0) ? 0 : ((__x > 1) ? 1 : __x); }

__DEVICE__
void __sincosf(float __x, float *__sinptr, float *__cosptr) {
  *__sinptr = __ocml_native_sin_f32(__x);
  *__cosptr = __ocml_native_cos_f32(__x);
}

__DEVICE__
float __sinf(float __x) { return __ocml_native_sin_f32(__x); }

__DEVICE__
float __tanf(float __x) { return __ocml_tan_f32(__x); }
// END INTRINSICS
// END FLOAT

// BEGIN DOUBLE
__DEVICE__
double acos(double __x) { return __ocml_acos_f64(__x); }

__DEVICE__
double acosh(double __x) { return __ocml_acosh_f64(__x); }

__DEVICE__
double asin(double __x) { return __ocml_asin_f64(__x); }

__DEVICE__
double asinh(double __x) { return __ocml_asinh_f64(__x); }

__DEVICE__
double atan(double __x) { return __ocml_atan_f64(__x); }

__DEVICE__
double atan2(double __x, double __y) { return __ocml_atan2_f64(__x, __y); }

__DEVICE__
double atanh(double __x) { return __ocml_atanh_f64(__x); }

__DEVICE__
double cbrt(double __x) { return __ocml_cbrt_f64(__x); }

__DEVICE__
double ceil(double __x) { return __builtin_ceil(__x); }

__DEVICE__
double copysign(double __x, double __y) {
  return __builtin_copysign(__x, __y);
}

__DEVICE__
double cos(double __x) { return __ocml_cos_f64(__x); }

__DEVICE__
double cosh(double __x) { return __ocml_cosh_f64(__x); }

__DEVICE__
double cospi(double __x) { return __ocml_cospi_f64(__x); }

__DEVICE__
double cyl_bessel_i0(double __x) { return __ocml_i0_f64(__x); }

__DEVICE__
double cyl_bessel_i1(double __x) { return __ocml_i1_f64(__x); }

__DEVICE__
double erf(double __x) { return __ocml_erf_f64(__x); }

__DEVICE__
double erfc(double __x) { return __ocml_erfc_f64(__x); }

__DEVICE__
double erfcinv(double __x) { return __ocml_erfcinv_f64(__x); }

__DEVICE__
double erfcx(double __x) { return __ocml_erfcx_f64(__x); }

__DEVICE__
double erfinv(double __x) { return __ocml_erfinv_f64(__x); }

__DEVICE__
double exp(double __x) { return __ocml_exp_f64(__x); }

__DEVICE__
double exp10(double __x) { return __ocml_exp10_f64(__x); }

__DEVICE__
double exp2(double __x) { return __ocml_exp2_f64(__x); }

__DEVICE__
double expm1(double __x) { return __ocml_expm1_f64(__x); }

__DEVICE__
double fabs(double __x) { return __builtin_fabs(__x); }

__DEVICE__
double fdim(double __x, double __y) { return __ocml_fdim_f64(__x, __y); }

__DEVICE__
double floor(double __x) { return __builtin_floor(__x); }

__DEVICE__
double fma(double __x, double __y, double __z) {
  return __builtin_fma(__x, __y, __z);
}

__DEVICE__
double fmax(double __x, double __y) { return __builtin_fmax(__x, __y); }

__DEVICE__
double fmin(double __x, double __y) { return __builtin_fmin(__x, __y); }

__DEVICE__
double fmod(double __x, double __y) { return __ocml_fmod_f64(__x, __y); }

__DEVICE__
double frexp(double __x, int *__nptr) {
  return __builtin_frexp(__x, __nptr);
}

__DEVICE__
double hypot(double __x, double __y) { return __ocml_hypot_f64(__x, __y); }

__DEVICE__
int ilogb(double __x) { return __ocml_ilogb_f64(__x); }

__DEVICE__
__RETURN_TYPE __finite(double __x) { return __builtin_isfinite(__x); }

__DEVICE__
__RETURN_TYPE __isinf(double __x) { return __builtin_isinf(__x); }

__DEVICE__
__RETURN_TYPE __isnan(double __x) { return __builtin_isnan(__x); }

__DEVICE__
double j0(double __x) { return __ocml_j0_f64(__x); }

__DEVICE__
double j1(double __x) { return __ocml_j1_f64(__x); }

__DEVICE__
double jn(int __n, double __x) { // TODO: we could use Ahmes multiplication
                                 // and the Miller & Brown algorithm
  //       for linear recurrences to get O(log n) steps, but it's unclear if
  //       it'd be beneficial in this case. Placeholder until OCML adds
  //       support.
  if (__n == 0)
    return j0(__x);
  if (__n == 1)
    return j1(__x);

double __x0 = j0(__x);
  double __x1 = j1(__x);
  for (int __i = 1; __i < __n; ++__i) {
    double __x2 = (2 * __i) / __x * __x1 - __x0;
    __x0 = __x1;
    __x1 = __x2;
  }
  return __x1;
}

__DEVICE__
double ldexp(double __x, int __e) { return __builtin_amdgcn_ldexp(__x, __e); }

__DEVICE__
double lgamma(double __x) { return __ocml_lgamma_f64(__x); }

__DEVICE__
long long int llrint(double __x) { return __builtin_rint(__x); }

__DEVICE__
long long int llround(double __x) { return __builtin_round(__x); }

__DEVICE__
double log(double __x) { return __ocml_log_f64(__x); }

__DEVICE__
double log10(double __x) { return __ocml_log10_f64(__x); }

__DEVICE__
double log1p(double __x) { return __ocml_log1p_f64(__x); }

__DEVICE__
double log2(double __x) { return __ocml_log2_f64(__x); }

__DEVICE__
double logb(double __x) { return __ocml_logb_f64(__x); }

__DEVICE__
long int lrint(double __x) { return __builtin_rint(__x); }

__DEVICE__
long int lround(double __x) { return __builtin_round(__x); }

__DEVICE__
double modf(double __x, double *__iptr) {
  double __tmp;
#ifdef __OPENMP_AMDGCN__
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
#endif
  double __r =
      __ocml_modf_f64(__x, (__attribute__((address_space(5))) double *)&__tmp);
  *__iptr = __tmp;

return __r;
}

__DEVICE__
double nan(const char *__tagp) {
#if !_WIN32
  union {
    double val;
    struct ieee_double {
      uint64_t mantissa : 51;
      uint32_t quiet : 1;
      uint32_t exponent : 11;
      uint32_t sign : 1;
    } bits;
  } __tmp;
  __static_assert_type_size_equal(sizeof(__tmp.val), sizeof(__tmp.bits));

__tmp.bits.sign = 0u;
  __tmp.bits.exponent = ~0u;
  __tmp.bits.quiet = 1u;
  __tmp.bits.mantissa = __make_mantissa(__tagp);

return __tmp.val;
#else
  __static_assert_type_size_equal(sizeof(uint64_t), sizeof(double));
  uint64_t __val = __make_mantissa(__tagp);
  __val |= 0xFFF << 51;
  return *reinterpret_cast<double *>(&__val);
#endif
}

__DEVICE__
double nearbyint(double __x) { return __builtin_nearbyint(__x); }

__DEVICE__
double nextafter(double __x, double __y) {
  return __ocml_nextafter_f64(__x, __y);
}

__DEVICE__
double norm(int __dim,
            const double *__a) { // TODO: placeholder until OCML adds support.
  double __r = 0;
  while (__dim--) {
    __r += __a[0] * __a[0];
    ++__a;
  }

return __ocml_sqrt_f64(__r);
}

__DEVICE__
double norm3d(double __x, double __y, double __z) {
  return __ocml_len3_f64(__x, __y, __z);
}

__DEVICE__
double norm4d(double __x, double __y, double __z, double __w) {
  return __ocml_len4_f64(__x, __y, __z, __w);
}

__DEVICE__
double normcdf(double __x) { return __ocml_ncdf_f64(__x); }

__DEVICE__
double normcdfinv(double __x) { return __ocml_ncdfinv_f64(__x); }

__DEVICE__
double pow(double __x, double __y) { return __ocml_pow_f64(__x, __y); }

__DEVICE__
double powi(double __x, int __y) { return __ocml_pown_f64(__x, __y); }

__DEVICE__
double rcbrt(double __x) { return __ocml_rcbrt_f64(__x); }

__DEVICE__
double remainder(double __x, double __y) {
  return __ocml_remainder_f64(__x, __y);
}

__DEVICE__
double remquo(double __x, double __y, int *__quo) {
  int __tmp;
#ifdef __OPENMP_AMDGCN__
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
#endif
  double __r = __ocml_remquo_f64(
      __x, __y, (__attribute__((address_space(5))) int *)&__tmp);
  *__quo = __tmp;

return __r;
}

__DEVICE__
double rhypot(double __x, double __y) { return __ocml_rhypot_f64(__x, __y); }

__DEVICE__
double rint(double __x) { return __builtin_rint(__x); }

__DEVICE__
double rnorm(int __dim,
             const double *__a) { // TODO: placeholder until OCML adds support.
  double __r = 0;
  while (__dim--) {
    __r += __a[0] * __a[0];
    ++__a;
  }

return __ocml_rsqrt_f64(__r);
}

__DEVICE__
double rnorm3d(double __x, double __y, double __z) {
  return __ocml_rlen3_f64(__x, __y, __z);
}

__DEVICE__
double rnorm4d(double __x, double __y, double __z, double __w) {
  return __ocml_rlen4_f64(__x, __y, __z, __w);
}

__DEVICE__
double round(double __x) { return __builtin_round(__x); }

__DEVICE__
double rsqrt(double __x) { return __ocml_rsqrt_f64(__x); }

__DEVICE__
double scalbln(double __x, long int __n) {
  return (__n < INT_MAX) ? __builtin_amdgcn_ldexp(__x, __n)
                         : __ocml_scalb_f64(__x, __n);
}
__DEVICE__
double scalbn(double __x, int __n) { return __builtin_amdgcn_ldexp(__x, __n); }

__DEVICE__
__RETURN_TYPE __signbit(double __x) { return __builtin_signbit(__x); }

__DEVICE__
double sin(double __x) { return __ocml_sin_f64(__x); }

__DEVICE__
void sincos(double __x, double *__sinptr, double *__cosptr) {
  double __tmp;
#ifdef __OPENMP_AMDGCN__
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
#endif
  *__sinptr = __ocml_sincos_f64(
      __x, (__attribute__((address_space(5))) double *)&__tmp);
  *__cosptr = __tmp;
}

__DEVICE__
void sincospi(double __x, double *__sinptr, double *__cosptr) {
  double __tmp;
#ifdef __OPENMP_AMDGCN__
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
#endif
  *__sinptr = __ocml_sincospi_f64(
      __x, (__attribute__((address_space(5))) double *)&__tmp);
  *__cosptr = __tmp;
}

__DEVICE__
double sinh(double __x) { return __ocml_sinh_f64(__x); }

__DEVICE__
double sinpi(double __x) { return __ocml_sinpi_f64(__x); }

__DEVICE__
double sqrt(double __x) { return __ocml_sqrt_f64(__x); }

__DEVICE__
double tan(double __x) { return __ocml_tan_f64(__x); }

__DEVICE__
double tanh(double __x) { return __ocml_tanh_f64(__x); }

__DEVICE__
double tgamma(double __x) { return __ocml_tgamma_f64(__x); }

__DEVICE__
double trunc(double __x) { return __builtin_trunc(__x); }

__DEVICE__
double y0(double __x) { return __ocml_y0_f64(__x); }

__DEVICE__
double y1(double __x) { return __ocml_y1_f64(__x); }

__DEVICE__
double yn(int __n, double __x) { // TODO: we could use Ahmes multiplication
                                 // and the Miller & Brown algorithm
  //       for linear recurrences to get O(log n) steps, but it's unclear if
  //       it'd be beneficial in this case. Placeholder until OCML adds
  //       support.
  if (__n == 0)
    return y0(__x);
  if (__n == 1)
    return y1(__x);

double __x0 = y0(__x);
  double __x1 = y1(__x);
  for (int __i = 1; __i < __n; ++__i) {
    double __x2 = (2 * __i) / __x * __x1 - __x0;
    __x0 = __x1;
    __x1 = __x2;
  }

return __x1;
}

// BEGIN INTRINSICS
#if defined OCML_BASIC_ROUNDED_OPERATIONS
__DEVICE__
double __dadd_rd(double __x, double __y) {
  return __ocml_add_rtn_f64(__x, __y);
}
__DEVICE__
double __dadd_rn(double __x, double __y) {
  return __ocml_add_rte_f64(__x, __y);
}
__DEVICE__
double __dadd_ru(double __x, double __y) {
  return __ocml_add_rtp_f64(__x, __y);
}
__DEVICE__
double __dadd_rz(double __x, double __y) {
  return __ocml_add_rtz_f64(__x, __y);
}
#else
__DEVICE__
double __dadd_rn(double __x, double __y) { return __x + __y; }
#endif

#if defined OCML_BASIC_ROUNDED_OPERATIONS
__DEVICE__
double __ddiv_rd(double __x, double __y) {
  return __ocml_div_rtn_f64(__x, __y);
}
__DEVICE__
double __ddiv_rn(double __x, double __y) {
  return __ocml_div_rte_f64(__x, __y);
}
__DEVICE__
double __ddiv_ru(double __x, double __y) {
  return __ocml_div_rtp_f64(__x, __y);
}
__DEVICE__
double __ddiv_rz(double __x, double __y) {
  return __ocml_div_rtz_f64(__x, __y);
}
#else
__DEVICE__
double __ddiv_rn(double __x, double __y) { return __x / __y; }
#endif

#if defined OCML_BASIC_ROUNDED_OPERATIONS
__DEVICE__
double __dmul_rd(double __x, double __y) {
  return __ocml_mul_rtn_f64(__x, __y);
}
__DEVICE__
double __dmul_rn(double __x, double __y) {
  return __ocml_mul_rte_f64(__x, __y);
}
__DEVICE__
double __dmul_ru(double __x, double __y) {
  return __ocml_mul_rtp_f64(__x, __y);
}
__DEVICE__
double __dmul_rz(double __x, double __y) {
  return __ocml_mul_rtz_f64(__x, __y);
}
#else
__DEVICE__
double __dmul_rn(double __x, double __y) { return __x * __y; }
#endif

#if defined OCML_BASIC_ROUNDED_OPERATIONS
__DEVICE__
double __drcp_rd(double __x) { return __ocml_div_rtn_f64(1.0, __x); }
__DEVICE__
double __drcp_rn(double __x) { return __ocml_div_rte_f64(1.0, __x); }
__DEVICE__
double __drcp_ru(double __x) { return __ocml_div_rtp_f64(1.0, __x); }
__DEVICE__
double __drcp_rz(double __x) { return __ocml_div_rtz_f64(1.0, __x); }
#else
__DEVICE__
double __drcp_rn(double __x) { return 1.0 / __x; }
#endif

#if defined OCML_BASIC_ROUNDED_OPERATIONS
__DEVICE__
double __dsqrt_rd(double __x) { return __ocml_sqrt_rtn_f64(__x); }
__DEVICE__
double __dsqrt_rn(double __x) { return __ocml_sqrt_rte_f64(__x); }
__DEVICE__
double __dsqrt_ru(double __x) { return __ocml_sqrt_rtp_f64(__x); }
__DEVICE__
double __dsqrt_rz(double __x) { return __ocml_sqrt_rtz_f64(__x); }
#else
__DEVICE__
double __dsqrt_rn(double __x) { return __ocml_sqrt_f64(__x); }
#endif

#if defined OCML_BASIC_ROUNDED_OPERATIONS
__DEVICE__
double __dsub_rd(double __x, double __y) {
  return __ocml_sub_rtn_f64(__x, __y);
}
__DEVICE__
double __dsub_rn(double __x, double __y) {
  return __ocml_sub_rte_f64(__x, __y);
}
__DEVICE__
double __dsub_ru(double __x, double __y) {
  return __ocml_sub_rtp_f64(__x, __y);
}
__DEVICE__
double __dsub_rz(double __x, double __y) {
  return __ocml_sub_rtz_f64(__x, __y);
}
#else
__DEVICE__
double __dsub_rn(double __x, double __y) { return __x - __y; }
#endif

#if defined OCML_BASIC_ROUNDED_OPERATIONS
__DEVICE__
double __fma_rd(double __x, double __y, double __z) {
  return __ocml_fma_rtn_f64(__x, __y, __z);
}
__DEVICE__
double __fma_rn(double __x, double __y, double __z) {
  return __ocml_fma_rte_f64(__x, __y, __z);
}
__DEVICE__
double __fma_ru(double __x, double __y, double __z) {
  return __ocml_fma_rtp_f64(__x, __y, __z);
}
__DEVICE__
double __fma_rz(double __x, double __y, double __z) {
  return __ocml_fma_rtz_f64(__x, __y, __z);
}
#else
__DEVICE__
double __fma_rn(double __x, double __y, double __z) {
  return __builtin_fma(__x, __y, __z);
}
#endif
// END INTRINSICS
// END DOUBLE

// C only macros
#if !defined(__cplusplus) && __STDC_VERSION__ >= 201112L
#define isfinite(__x) _Generic((__x), float : __finitef, double : __finite)(__x)
#define isinf(__x) _Generic((__x), float : __isinff, double : __isinf)(__x)
#define isnan(__x) _Generic((__x), float : __isnanf, double : __isnan)(__x)
#define signbit(__x)                                                           \
  _Generic((__x), float : __signbitf, double : __signbit)(__x)
#endif // !defined(__cplusplus) && __STDC_VERSION__ >= 201112L

#if defined(__cplusplus)
template <class T> __DEVICE__ T min(T __arg1, T __arg2) {
  return (__arg1 < __arg2) ? __arg1 : __arg2;
}

template <class T> __DEVICE__ T max(T __arg1, T __arg2) {
  return (__arg1 > __arg2) ? __arg1 : __arg2;
}

__DEVICE__ int min(int __arg1, int __arg2) {
  return (__arg1 < __arg2) ? __arg1 : __arg2;
}
__DEVICE__ int max(int __arg1, int __arg2) {
  return (__arg1 > __arg2) ? __arg1 : __arg2;
}

__DEVICE__
float max(float __x, float __y) { return __builtin_fmaxf(__x, __y); }

__DEVICE__
double max(double __x, double __y) { return __builtin_fmax(__x, __y); }

__DEVICE__
float min(float __x, float __y) { return __builtin_fminf(__x, __y); }

__DEVICE__
double min(double __x, double __y) { return __builtin_fmin(__x, __y); }

#if !defined(__HIPCC_RTC__) && !defined(__OPENMP_AMDGCN__)
__host__ inline static int min(int __arg1, int __arg2) {
  return std::min(__arg1, __arg2);
}

__host__ inline static int max(int __arg1, int __arg2) {
  return std::max(__arg1, __arg2);
}
#endif // !defined(__HIPCC_RTC__) && !defined(__OPENMP_AMDGCN__)
#endif

#pragma pop_macro("__DEVICE__")
#pragma pop_macro("__RETURN_TYPE")

#endif // __CLANG_HIP_MATH_H__

X7ROOT File Manager

Editing: __clang_hip_math.h

Upload File

Create Folder