diff --git a/projects/clr/hipamd/docs/markdown/device_md_gen.py b/projects/clr/hipamd/docs/markdown/device_md_gen.py new file mode 100644 index 0000000000..4795ea98e0 --- /dev/null +++ b/projects/clr/hipamd/docs/markdown/device_md_gen.py @@ -0,0 +1,508 @@ +""" +Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +""" + +""" +1. This files uses Python3 to run + +List of device functions: +acosf +acoshf +asinf +asinhf +atan2f +atanf +atanhf +cbrtf +ceilf +copysignf +cosf +coshf +cospif +cyl_bessel_i0f +cyl_bessel_i1f +erfcf +erfcinvf +erfcxf +erff +erfinvf +exp10f +exp2f +expf +expm1f +fabsf +fdimf +fdividef +floorf +fmaf +fmaxf +fminf +fmodf +frexpf +hypotf +ilogbf +isfinite +isinf +isnan +j0f +j1f +jnf +ldexpf +lgammaf +llrintf +llroundf +log10f +log1pf +logbf +lrintf +lroundf +modff +nanf +nearbyintf +nextafterf +norm3df +norm4df +normcdff +normcdfinvf +normf +powf +rcbrtf +remainderf +remquof +rhypotf +rintf +rnorm3df +rnorm4df +rnormf +roundf +rsqrtf +scalblnf +scalbnf +signbit +sincosf +sincospif +sinf +sinhf +sinpif +sqrtf +tanf +tanhf +tgammaf +truncf +y0f +y1f +ynf +acos +acosh +asin +asinh +atan +atan2 +atanh +cbrt +ceil +copysign +cos +cosh +cospi +cyl_bessel_i0 +cyl_bessel_i1 +erf +erfc +erfcinv +erfcx +erfinv +exp +exp10 +exp2 +expm1 +fabs +fdim +floor +fma +fmax +fmin +fmod +frexp +hypot +ilogb +isfinite +isinf +isnan +j0 +j1 +jn +ldexp +lgamma +llrint +llround +log +log10 +log1p +log2 +logb +lrint +lround +modf +nan +nearbyint +nextafter +norm +norm3d +norm4d +normcdf +normcdfinv +pow +rcbrt +remainder +remquo +rhypot +rint +rnorm +rnorm3d +rnorm4d +round +rsqrt +scalbln +scalbn +signbit +sin +sincos +sincospi +sinh +sinpi +sqrt +tan +tanh +tgamma +trunc +y0 +y1 +yn +__cosf +__exp10f +__expf +__fadd_rd +__fadd_rn +__fadd_ru +__fadd_rz +__fdiv_rd +__fdiv_rn +__fdiv_ru +__fdiv_rz +__fdividef +__fmaf_rd +__fmaf_rn +__fmaf_ru +__fmaf_rz +__fmul_rd +__fmul_rn +__fmul_ru +__fmul_rz +__frcp_rd +__frcp_rn +__frcp_ru +__frcp_rz +__frsqrt_rn +__fsqrt_rd +__fsqrt_rn +__fsqrt_ru +__fsqrt_rz +__fsub_rd +__fsub_rn +__fsub_ru +__log10f +__log2f +__logf +__powf +__saturatef +__sincosf +__sinf +__tanf +__dadd_rd +__dadd_rn +__dadd_ru +__dadd_rz +__ddiv_rd +__ddiv_rn +__ddiv_ru +__ddiv_rz +__dmul_rd +__dmul_rn +__dmul_ru +__dmul_rz +__drcp_rd +__drcp_rn +__drcp_ru +__drcp_rz +__dsqrt_rd +__dsqrt_rn +__dsqrt_ru +__dsqrt_rz +__dsub_rd +__dsub_rn +__dsub_ru +__dsub_rz +__fma_rd +__fma_rn +__fma_ru +__fma_rz +__brev +__brevll +__byte_perm +__clz +__clzll +__ffs +__ffsll +__hadd +__mul24 +__mul64hi +__mulhi +__popc +__popcll +__rhadd +__sad +__uhadd +__umul24 +__umul64hi +__umulhi +__urhadd +__usad +__double2float_rd +__double2float_rn +__double2float_ru +__double2float_rz +__double2hiint +__double2int_rd +__double2int_rn +__double2int_ru +__double2int_rz +__double2ll_rd +__double2ll_rn +__double2ll_ru +__double2ll_rz +__double2loint +__double2uint_rd +__double2uint_rn +__double2uint_ru +__double2uint_rz +__double2ull_rd +__double2ull_rn +__double2ull_ru +__double2ull_rz +__double_as_longlong +__float2half_rn +__half2float +__float2half_rn +__half2float +__float2int_rd +__float2int_rn +__float2int_ru +__float2int_rz +__float2ll_rd +__float2ll_rn +__float2ll_ru +__float2ll_rz +__float2uint_rd +__float2uint_rn +__float2uint_ru +__float2uint_rz +__float2ull_rd +__float2ull_rn +__float2ull_ru +__float2ull_rz +__float_as_int +__float_as_uint +__hiloint2double +__int2double_rn +__int2float_rd +__int2float_rn +__int2float_ru +__int2float_rz +__int_as_float +__ll2double_rd +__ll2double_rn +__ll2double_ru +__ll2double_rz +__ll2float_rd +__ll2float_rn +__ll2float_ru +__ll2float_rz +__longlong_as_double +__uint2double_rn +__uint2float_rd +__uint2float_rn +__uint2float_ru +__uint2float_rz +__uint_as_float +__ull2double_rd +__ull2double_rn +__ull2double_ru +__ull2double_rz +__ull2float_rd +__ull2float_rn +__ull2float_ru +__ull2float_rz +__heq +__hge +__hgt +__hisinf +__hisnan +__hle +__hlt +__hne +__hbeq2 +__hbge2 +__hbgt2 +__hble2 +__hblt2 +__hbne2 +__heq2 +__hge2 +__hgt2 +__hisnan2 +__hle2 +__hlt2 +__hne2 +__float22half2_rn +__float2half +__float2half2_rn +__float2half_rd +__float2half_rn +__float2half_ru +__float2half_rz +__floats2half2_rn +__half22float2 +__half2float +half2half2 +__half2int_rd +__half2int_rn +__half2int_ru +__half2int_rz +__half2ll_rd +__half2ll_rn +__half2ll_ru +__half2ll_rz +__half2short_rd +__half2short_rn +__half2short_ru +__half2short_rz +__half2uint_rd +__half2uint_rn +__half2uint_ru +__half2uint_rz +__half2ull_rd +__half2ull_rn +__half2ull_ru +__half2ull_rz +__half2ushort_rd +__half2ushort_rn +__half2ushort_ru +__half2ushort_rz +__half_as_short +__half_as_ushort +__halves2half2 +__high2float +__high2half +__high2half2 +__highs2half2 +__int2half_rd +__int2half_rn +__int2half_ru +__int2half_rz +__ll2half_rd +__ll2half_rn +__ll2half_ru +__ll2half_rz +__low2float +__low2half +__low2half2 +__low2half2 +__lowhigh2highlow +__lows2half2 +__short2half_rd +__short2half_rn +__short2half_ru +__short2half_rz +__uint2half_rd +__uint2half_rn +__uint2half_ru +__uint2half_rz +__ull2half_rd +__ull2half_rn +__ull2half_ru +__ull2half_rz +__ushort2half_rd +__ushort2half_rn +__ushort2half_ru +__ushort2half_rz +__ushort_as_half +""" +# The dictionary is to place description of each device function. Expand it to all the device functions +deviceFuncDesc = {'acosf': "This function returns floating point of arc cosine from a floating point input"} + +fnames = ["../../include/hip/hcc_detail/math_functions.h","../../include/hip/hcc_detail/device_functions.h","../../include/hip/hcc_detail/hip_fp16.h"] +markdownFileName = "./hip-math-api.md" + +preamble = "# HIP MATH APIs Documentation \n"+\ +"HIP supports most of the device functions supported by CUDA. Way to find the unsupported one is to search for the function and check its description\n" + \ +"Note: This document is not human generated. Any changes to this file will be discarded. Please make changes to Python3 script docs/markdown/device_md_gen.py\n\n" + \ +"## For Developers \n" + \ +"If you add or fixed a device function, make sure to add a signature of the function and definition later.\n" + \ +"For example, if you want to add `__device__ float __dotf(float4, float4)`, which does a dot product on 4 float vector components \n" + \ +"The way to add to the header is, \n" + \ +"```cpp \n" + \ +"__device__ static float __dotf(float4, float4); \n" + \ +"/*Way down in the file....*/\n" + \ +"__device__ static inline float __dotf(float4 x, float4 y) { \n" + \ +" /*implementation*/\n}\n" + \ +"```\n\n" + \ +"This helps python script to add the device function newly declared into markdown documentation (as it looks at functions with `;` at the end and `__device__` at the beginning)\n\n" + \ +"The next step would be to add Description to `deviceFuncDesc`.\n" + \ +"From the above example, it can be writtern as,\n`deviceFuncDesc['__dotf'] = 'This functions takes 2 4 component float vector and outputs dot product across them'`\n\n" + +def generateSnippet(name, description, signature): + return "### " + name + "\n" + \ + "```cpp \n" + signature + "\n```\n" + \ + "**Description:** " + description + "\n\n\n" + +def getName(line): + l1 = line.split('(') + l2 = l1[0].split(' ') + return l2[-1] + +with open(markdownFileName, 'w') as mdfd: + mdfd.truncate() + mdfd.write(preamble) + for fname in fnames: + with open(fname) as fd: + lines = fd.readlines() + for line in lines: + if line.find('HIP_FAST_MATH') != -1: + break; + if line.find('__device__') != -1 and line.find(';') != -1 and line.find('hip') == -1: + name = getName(line) + if line.find('//') == -1: + if name in deviceFuncDesc: + mdfd.write(generateSnippet(name, deviceFuncDesc[name], line)) + else: + mdfd.write(generateSnippet(name, "Supported", line)) + else: + mdfd.write(generateSnippet(name, "**NOT Supported**", line)) + fd.close() + mdfd.close() diff --git a/projects/clr/hipamd/docs/markdown/hip-math-api.md b/projects/clr/hipamd/docs/markdown/hip-math-api.md new file mode 100644 index 0000000000..daf34fbfa3 --- /dev/null +++ b/projects/clr/hipamd/docs/markdown/hip-math-api.md @@ -0,0 +1,3470 @@ +# HIP MATH APIs Documentation +HIP supports most of the device functions supported by CUDA. Way to find the unsupported one is to search for the function and check its description +Note: This document is not human generated. Any changes to this file will be discarded. Please make changes to Python3 script docs/markdown/device_md_gen.py + +## For Developers +If you add or fixed a device function, make sure to add a signature of the function and definition later. +For example, if you want to add `__device__ float __dotf(float4, float4)`, which does a dot product on 4 float vector components +The way to add to the header is, +```cpp +__device__ static float __dotf(float4, float4); +/*Way down in the file....*/ +__device__ static inline float __dotf(float4 x, float4 y) { + /*implementation*/ +} +``` + +This helps python script to add the device function newly declared into markdown documentation (as it looks at functions with `;` at the end and `__device__` at the beginning) + +The next step would be to add Description to `deviceFuncDesc`. +From the above example, it can be writtern as, +`deviceFuncDesc['__dotf'] = 'This functions takes 2 4 component float vector and outputs dot product across them'` + +### acosf +```cpp +__device__ float acosf(float x); + +``` +**Description:** This function returns floating point of arc cosine from a floating point input + + +### acoshf +```cpp +__device__ float acoshf(float x); + +``` +**Description:** Supported + + +### asinf +```cpp +__device__ float asinf(float x); + +``` +**Description:** Supported + + +### asinhf +```cpp +__device__ float asinhf(float x); + +``` +**Description:** Supported + + +### atan2f +```cpp +__device__ float atan2f(float y, float x); + +``` +**Description:** Supported + + +### atanf +```cpp +__device__ float atanf(float x); + +``` +**Description:** Supported + + +### atanhf +```cpp +__device__ float atanhf(float x); + +``` +**Description:** Supported + + +### cbrtf +```cpp +__device__ float cbrtf(float x); + +``` +**Description:** Supported + + +### ceilf +```cpp +__device__ float ceilf(float x); + +``` +**Description:** Supported + + +### copysignf +```cpp +__device__ float copysignf(float x, float y); + +``` +**Description:** Supported + + +### cosf +```cpp +__device__ float cosf(float x); + +``` +**Description:** Supported + + +### coshf +```cpp +__device__ float coshf(float x); + +``` +**Description:** Supported + + +### cospif +```cpp +__device__ __host__ float cospif(float x); + +``` +**Description:** Supported + + +### cyl_bessel_i0f +```cpp +//__device__ float cyl_bessel_i0f(float x); + +``` +**Description:** **NOT Supported** + + +### cyl_bessel_i1f +```cpp +//__device__ float cyl_bessel_i1f(float x); + +``` +**Description:** **NOT Supported** + + +### erfcf +```cpp +__device__ float erfcf(float x); + +``` +**Description:** Supported + + +### erfcinvf +```cpp +__device__ float erfcinvf(float y); + +``` +**Description:** Supported + + +### erfcxf +```cpp +__device__ float erfcxf(float x); + +``` +**Description:** Supported + + +### erff +```cpp +__device__ float erff(float x); + +``` +**Description:** Supported + + +### erfinvf +```cpp +__device__ float erfinvf(float y); + +``` +**Description:** Supported + + +### exp10f +```cpp +__device__ float exp10f(float x); + +``` +**Description:** Supported + + +### exp2f +```cpp +__device__ float exp2f(float x); + +``` +**Description:** Supported + + +### expf +```cpp +__device__ float expf(float x); + +``` +**Description:** Supported + + +### expm1f +```cpp +__device__ float expm1f(float x); + +``` +**Description:** Supported + + +### fabsf +```cpp +__device__ float fabsf(float x); + +``` +**Description:** Supported + + +### fdimf +```cpp +__device__ float fdimf(float x, float y); + +``` +**Description:** Supported + + +### fdividef +```cpp +__device__ __host__ float fdividef(float x, float y); + +``` +**Description:** Supported + + +### floorf +```cpp +__device__ float floorf(float x); + +``` +**Description:** Supported + + +### fmaf +```cpp +__device__ float fmaf(float x, float y, float z); + +``` +**Description:** Supported + + +### fmaxf +```cpp +__device__ float fmaxf(float x, float y); + +``` +**Description:** Supported + + +### fminf +```cpp +__device__ float fminf(float x, float y); + +``` +**Description:** Supported + + +### fmodf +```cpp +__device__ float fmodf(float x, float y); + +``` +**Description:** Supported + + +### frexpf +```cpp +//__device__ float frexpf(float x, int* nptr); + +``` +**Description:** **NOT Supported** + + +### hypotf +```cpp +__device__ float hypotf(float x, float y); + +``` +**Description:** Supported + + +### ilogbf +```cpp +__device__ float ilogbf(float x); + +``` +**Description:** Supported + + +### isfinite +```cpp +__device__ __host__ int isfinite(float a); + +``` +**Description:** Supported + + +### isinf +```cpp +__device__ unsigned isinf(float a); + +``` +**Description:** Supported + + +### isnan +```cpp +__device__ unsigned isnan(float a); + +``` +**Description:** Supported + + +### j0f +```cpp +__device__ float j0f(float x); + +``` +**Description:** Supported + + +### j1f +```cpp +__device__ float j1f(float x); + +``` +**Description:** Supported + + +### jnf +```cpp +__device__ float jnf(int n, float x); + +``` +**Description:** Supported + + +### ldexpf +```cpp +__device__ float ldexpf(float x, int exp); + +``` +**Description:** Supported + + +### lgammaf +```cpp +//__device__ float lgammaf(float x); + +``` +**Description:** **NOT Supported** + + +### llrintf +```cpp +__device__ long long int llrintf(float x); + +``` +**Description:** Supported + + +### llroundf +```cpp +__device__ long long int llroundf(float x); + +``` +**Description:** Supported + + +### log10f +```cpp +__device__ float log10f(float x); + +``` +**Description:** Supported + + +### log1pf +```cpp +__device__ float log1pf(float x); + +``` +**Description:** Supported + + +### logbf +```cpp +__device__ float logbf(float x); + +``` +**Description:** Supported + + +### lrintf +```cpp +__device__ long int lrintf(float x); + +``` +**Description:** Supported + + +### lroundf +```cpp +__device__ long int lroundf(float x); + +``` +**Description:** Supported + + +### modff +```cpp +//__device__ float modff(float x, float *iptr); + +``` +**Description:** **NOT Supported** + + +### nanf +```cpp +__device__ float nanf(const char* tagp); + +``` +**Description:** Supported + + +### nearbyintf +```cpp +__device__ float nearbyintf(float x); + +``` +**Description:** Supported + + +### nextafterf +```cpp +//__device__ float nextafterf(float x, float y); + +``` +**Description:** **NOT Supported** + + +### norm3df +```cpp +__device__ float norm3df(float a, float b, float c); + +``` +**Description:** Supported + + +### norm4df +```cpp +__device__ float norm4df(float a, float b, float c, float d); + +``` +**Description:** Supported + + +### normcdff +```cpp +__device__ float normcdff(float y); + +``` +**Description:** Supported + + +### normcdfinvf +```cpp +__device__ float normcdfinvf(float y); + +``` +**Description:** Supported + + +### normf +```cpp +__device__ float normf(int dim, const float *a); + +``` +**Description:** Supported + + +### powf +```cpp +__device__ float powf(float x, float y); + +``` +**Description:** Supported + + +### rcbrtf +```cpp +__device__ float rcbrtf(float x); + +``` +**Description:** Supported + + +### remainderf +```cpp +__device__ float remainderf(float x, float y); + +``` +**Description:** Supported + + +### remquof +```cpp +__device__ float remquof(float x, float y, int *quo); + +``` +**Description:** Supported + + +### rhypotf +```cpp +__device__ float rhypotf(float x, float y); + +``` +**Description:** Supported + + +### rintf +```cpp +__device__ float rintf(float x); + +``` +**Description:** Supported + + +### rnorm3df +```cpp +__device__ float rnorm3df(float a, float b, float c); + +``` +**Description:** Supported + + +### rnorm4df +```cpp +__device__ float rnorm4df(float a, float b, float c, float d); + +``` +**Description:** Supported + + +### rnormf +```cpp +__device__ float rnormf(int dim, const float* a); + +``` +**Description:** Supported + + +### roundf +```cpp +__device__ float roundf(float x); + +``` +**Description:** Supported + + +### rsqrtf +```cpp +__device__ __host__ float rsqrtf(float x); + +``` +**Description:** Supported + + +### scalblnf +```cpp +__device__ float scalblnf(float x, long int n); + +``` +**Description:** Supported + + +### scalbnf +```cpp +__device__ float scalbnf(float x, int n); + +``` +**Description:** Supported + + +### signbit +```cpp +__device__ __host__ unsigned signbit(float a); + +``` +**Description:** Supported + + +### sincosf +```cpp +__device__ void sincosf(float x, float *sptr, float *cptr); + +``` +**Description:** Supported + + +### sincospif +```cpp +__device__ void sincospif(float x, float *sptr, float *cptr); + +``` +**Description:** Supported + + +### sinf +```cpp +__device__ float sinf(float x); + +``` +**Description:** Supported + + +### sinhf +```cpp +__device__ float sinhf(float x); + +``` +**Description:** Supported + + +### sinpif +```cpp +__device__ __host__ float sinpif(float x); + +``` +**Description:** Supported + + +### sqrtf +```cpp +__device__ float sqrtf(float x); + +``` +**Description:** Supported + + +### tanf +```cpp +__device__ float tanf(float x); + +``` +**Description:** Supported + + +### tanhf +```cpp +__device__ float tanhf(float x); + +``` +**Description:** Supported + + +### tgammaf +```cpp +__device__ float tgammaf(float x); + +``` +**Description:** Supported + + +### truncf +```cpp +__device__ float truncf(float x); + +``` +**Description:** Supported + + +### y0f +```cpp +__device__ float y0f(float x); + +``` +**Description:** Supported + + +### y1f +```cpp +__device__ float y1f(float x); + +``` +**Description:** Supported + + +### ynf +```cpp +__device__ float ynf(int n, float x); + +``` +**Description:** Supported + + +### acos +```cpp +__device__ double acos(double x); + +``` +**Description:** Supported + + +### acosh +```cpp +__device__ double acosh(double x); + +``` +**Description:** Supported + + +### asin +```cpp +__device__ double asin(double x); + +``` +**Description:** Supported + + +### asinh +```cpp +__device__ double asinh(double x); + +``` +**Description:** Supported + + +### atan +```cpp +__device__ double atan(double x); + +``` +**Description:** Supported + + +### atan2 +```cpp +__device__ double atan2(double y, double x); + +``` +**Description:** Supported + + +### atanh +```cpp +__device__ double atanh(double x); + +``` +**Description:** Supported + + +### cbrt +```cpp +__device__ double cbrt(double x); + +``` +**Description:** Supported + + +### ceil +```cpp +__device__ double ceil(double x); + +``` +**Description:** Supported + + +### copysign +```cpp +__device__ double copysign(double x, double y); + +``` +**Description:** Supported + + +### cos +```cpp +__device__ double cos(double x); + +``` +**Description:** Supported + + +### cosh +```cpp +__device__ double cosh(double x); + +``` +**Description:** Supported + + +### cospi +```cpp +__device__ __host__ double cospi(double x); + +``` +**Description:** Supported + + +### cyl_bessel_i0 +```cpp +//__device__ double cyl_bessel_i0(double x); + +``` +**Description:** **NOT Supported** + + +### cyl_bessel_i1 +```cpp +//__device__ double cyl_bessel_i1(double x); + +``` +**Description:** **NOT Supported** + + +### erf +```cpp +__device__ double erf(double x); + +``` +**Description:** Supported + + +### erfc +```cpp +__device__ double erfc(double x); + +``` +**Description:** Supported + + +### erfcinv +```cpp +__device__ double erfcinv(double y); + +``` +**Description:** Supported + + +### erfcx +```cpp +__device__ double erfcx(double x); + +``` +**Description:** Supported + + +### erfinv +```cpp +__device__ double erfinv(double x); + +``` +**Description:** Supported + + +### exp +```cpp +__device__ double exp(double x); + +``` +**Description:** Supported + + +### exp10 +```cpp +__device__ double exp10(double x); + +``` +**Description:** Supported + + +### exp2 +```cpp +__device__ double exp2(double x); + +``` +**Description:** Supported + + +### expm1 +```cpp +__device__ double expm1(double x); + +``` +**Description:** Supported + + +### fabs +```cpp +__device__ double fabs(double x); + +``` +**Description:** Supported + + +### fdim +```cpp +__device__ double fdim(double x, double y); + +``` +**Description:** Supported + + +### floor +```cpp +__device__ double floor(double x); + +``` +**Description:** Supported + + +### fma +```cpp +__device__ double fma(double x, double y, double z); + +``` +**Description:** Supported + + +### fmax +```cpp +__device__ double fmax(double x, double y); + +``` +**Description:** Supported + + +### fmin +```cpp +__device__ double fmin(double x, double y); + +``` +**Description:** Supported + + +### fmod +```cpp +__device__ double fmod(double x, double y); + +``` +**Description:** Supported + + +### frexp +```cpp +//__device__ double frexp(double x, int *nptr); + +``` +**Description:** **NOT Supported** + + +### hypot +```cpp +__device__ double hypot(double x, double y); + +``` +**Description:** Supported + + +### ilogb +```cpp +__device__ double ilogb(double x); + +``` +**Description:** Supported + + +### isfinite +```cpp +__device__ __host__ unsigned isfinite(double x); + +``` +**Description:** Supported + + +### isinf +```cpp +__device__ unsigned isinf(double x); + +``` +**Description:** Supported + + +### isnan +```cpp +__device__ unsigned isnan(double x); + +``` +**Description:** Supported + + +### j0 +```cpp +__device__ double j0(double x); + +``` +**Description:** Supported + + +### j1 +```cpp +__device__ double j1(double x); + +``` +**Description:** Supported + + +### jn +```cpp +__device__ double jn(int n, double x); + +``` +**Description:** Supported + + +### ldexp +```cpp +__device__ double ldexp(double x, int exp); + +``` +**Description:** Supported + + +### lgamma +```cpp +__device__ double lgamma(double x); + +``` +**Description:** Supported + + +### llrint +```cpp +__device__ long long llrint(double x); + +``` +**Description:** Supported + + +### llround +```cpp +__device__ long long llround(double x); + +``` +**Description:** Supported + + +### log +```cpp +__device__ double log(double x); + +``` +**Description:** Supported + + +### log10 +```cpp +__device__ double log10(double x); + +``` +**Description:** Supported + + +### log1p +```cpp +__device__ double log1p(double x); + +``` +**Description:** Supported + + +### log2 +```cpp +__device__ double log2(double x); + +``` +**Description:** Supported + + +### logb +```cpp +__device__ double logb(double x); + +``` +**Description:** Supported + + +### lrint +```cpp +__device__ long int lrint(double x); + +``` +**Description:** Supported + + +### lround +```cpp +__device__ long int lround(double x); + +``` +**Description:** Supported + + +### modf +```cpp +//__device__ double modf(double x, double *iptr); + +``` +**Description:** **NOT Supported** + + +### nan +```cpp +__device__ double nan(const char* tagp); + +``` +**Description:** Supported + + +### nearbyint +```cpp +__device__ double nearbyint(double x); + +``` +**Description:** Supported + + +### nextafter +```cpp +__device__ double nextafter(double x, double y); + +``` +**Description:** Supported + + +### norm +```cpp +__device__ double norm(int dim, const double* t); + +``` +**Description:** Supported + + +### norm3d +```cpp +__device__ double norm3d(double a, double b, double c); + +``` +**Description:** Supported + + +### norm4d +```cpp +__device__ double norm4d(double a, double b, double c, double d); + +``` +**Description:** Supported + + +### normcdf +```cpp +__device__ double normcdf(double y); + +``` +**Description:** Supported + + +### normcdfinv +```cpp +__device__ double normcdfinv(double y); + +``` +**Description:** Supported + + +### pow +```cpp +__device__ double pow(double x, double y); + +``` +**Description:** Supported + + +### rcbrt +```cpp +__device__ double rcbrt(double x); + +``` +**Description:** Supported + + +### remainder +```cpp +__device__ double remainder(double x, double y); + +``` +**Description:** Supported + + +### remquo +```cpp +//__device__ double remquo(double x, double y, int *quo); + +``` +**Description:** **NOT Supported** + + +### rhypot +```cpp +__device__ double rhypot(double x, double y); + +``` +**Description:** Supported + + +### rint +```cpp +__device__ double rint(double x); + +``` +**Description:** Supported + + +### rnorm +```cpp +__device__ double rnorm(int dim, const double* t); + +``` +**Description:** Supported + + +### rnorm3d +```cpp +__device__ double rnorm3d(double a, double b, double c); + +``` +**Description:** Supported + + +### rnorm4d +```cpp +__device__ double rnorm4d(double a, double b, double c, double d); + +``` +**Description:** Supported + + +### round +```cpp +__device__ double round(double x); + +``` +**Description:** Supported + + +### rsqrt +```cpp +__device__ __host__ double rsqrt(double x); + +``` +**Description:** Supported + + +### scalbln +```cpp +__device__ double scalbln(double x, long int n); + +``` +**Description:** Supported + + +### scalbn +```cpp +__device__ double scalbn(double x, int n); + +``` +**Description:** Supported + + +### signbit +```cpp +__device__ __host__ unsigned signbit(double a); + +``` +**Description:** Supported + + +### sin +```cpp +__device__ double sin(double a); + +``` +**Description:** Supported + + +### sincos +```cpp +__device__ void sincos(double x, double *sptr, double *cptr); + +``` +**Description:** Supported + + +### sincospi +```cpp +__device__ void sincospi(double x, double *sptr, double *cptr); + +``` +**Description:** Supported + + +### sinh +```cpp +__device__ double sinh(double x); + +``` +**Description:** Supported + + +### sinpi +```cpp +__device__ __host__ double sinpi(double x); + +``` +**Description:** Supported + + +### sqrt +```cpp +__device__ double sqrt(double x); + +``` +**Description:** Supported + + +### tan +```cpp +__device__ double tan(double x); + +``` +**Description:** Supported + + +### tanh +```cpp +__device__ double tanh(double x); + +``` +**Description:** Supported + + +### tgamma +```cpp +__device__ double tgamma(double x); + +``` +**Description:** Supported + + +### trunc +```cpp +__device__ double trunc(double x); + +``` +**Description:** Supported + + +### y0 +```cpp +__device__ double y0(double x); + +``` +**Description:** Supported + + +### y1 +```cpp +__device__ double y1(double y); + +``` +**Description:** Supported + + +### yn +```cpp +__device__ double yn(int n, double x); + +``` +**Description:** Supported + + +### __cosf +```cpp +__device__ float __cosf(float x); + +``` +**Description:** Supported + + +### __exp10f +```cpp +__device__ float __exp10f(float x); + +``` +**Description:** Supported + + +### __expf +```cpp +__device__ float __expf(float x); + +``` +**Description:** Supported + + +### __fadd_rd +```cpp +__device__ static float __fadd_rd(float x, float y); + +``` +**Description:** Supported + + +### __fadd_rn +```cpp +__device__ static float __fadd_rn(float x, float y); + +``` +**Description:** Supported + + +### __fadd_ru +```cpp +__device__ static float __fadd_ru(float x, float y); + +``` +**Description:** Supported + + +### __fadd_rz +```cpp +__device__ static float __fadd_rz(float x, float y); + +``` +**Description:** Supported + + +### __fdiv_rd +```cpp +__device__ static float __fdiv_rd(float x, float y); + +``` +**Description:** Supported + + +### __fdiv_rn +```cpp +__device__ static float __fdiv_rn(float x, float y); + +``` +**Description:** Supported + + +### __fdiv_ru +```cpp +__device__ static float __fdiv_ru(float x, float y); + +``` +**Description:** Supported + + +### __fdiv_rz +```cpp +__device__ static float __fdiv_rz(float x, float y); + +``` +**Description:** Supported + + +### __fdividef +```cpp +__device__ static float __fdividef(float x, float y); + +``` +**Description:** Supported + + +### __fmaf_rd +```cpp +__device__ float __fmaf_rd(float x, float y, float z); + +``` +**Description:** Supported + + +### __fmaf_rn +```cpp +__device__ float __fmaf_rn(float x, float y, float z); + +``` +**Description:** Supported + + +### __fmaf_ru +```cpp +__device__ float __fmaf_ru(float x, float y, float z); + +``` +**Description:** Supported + + +### __fmaf_rz +```cpp +__device__ float __fmaf_rz(float x, float y, float z); + +``` +**Description:** Supported + + +### __fmul_rd +```cpp +__device__ static float __fmul_rd(float x, float y); + +``` +**Description:** Supported + + +### __fmul_rn +```cpp +__device__ static float __fmul_rn(float x, float y); + +``` +**Description:** Supported + + +### __fmul_ru +```cpp +__device__ static float __fmul_ru(float x, float y); + +``` +**Description:** Supported + + +### __fmul_rz +```cpp +__device__ static float __fmul_rz(float x, float y); + +``` +**Description:** Supported + + +### __frcp_rd +```cpp +__device__ float __frcp_rd(float x); + +``` +**Description:** Supported + + +### __frcp_rn +```cpp +__device__ float __frcp_rn(float x); + +``` +**Description:** Supported + + +### __frcp_ru +```cpp +__device__ float __frcp_ru(float x); + +``` +**Description:** Supported + + +### __frcp_rz +```cpp +__device__ float __frcp_rz(float x); + +``` +**Description:** Supported + + +### __frsqrt_rn +```cpp +__device__ float __frsqrt_rn(float x); + +``` +**Description:** Supported + + +### __fsqrt_rd +```cpp +__device__ float __fsqrt_rd(float x); + +``` +**Description:** Supported + + +### __fsqrt_rn +```cpp +__device__ float __fsqrt_rn(float x); + +``` +**Description:** Supported + + +### __fsqrt_ru +```cpp +__device__ float __fsqrt_ru(float x); + +``` +**Description:** Supported + + +### __fsqrt_rz +```cpp +__device__ float __fsqrt_rz(float x); + +``` +**Description:** Supported + + +### __fsub_rd +```cpp +__device__ static float __fsub_rd(float x, float y); + +``` +**Description:** Supported + + +### __fsub_rn +```cpp +__device__ static float __fsub_rn(float x, float y); + +``` +**Description:** Supported + + +### __fsub_ru +```cpp +__device__ static float __fsub_ru(float x, float y); + +``` +**Description:** Supported + + +### __log10f +```cpp +__device__ float __log10f(float x); + +``` +**Description:** Supported + + +### __log2f +```cpp +__device__ float __log2f(float x); + +``` +**Description:** Supported + + +### __logf +```cpp +__device__ float __logf(float x); + +``` +**Description:** Supported + + +### __powf +```cpp +__device__ float __powf(float base, float exponent); + +``` +**Description:** Supported + + +### __saturatef +```cpp +__device__ static float __saturatef(float x); + +``` +**Description:** Supported + + +### __sincosf +```cpp +__device__ void __sincosf(float x, float *s, float *c); + +``` +**Description:** Supported + + +### __sinf +```cpp +__device__ float __sinf(float x); + +``` +**Description:** Supported + + +### __tanf +```cpp +__device__ float __tanf(float x); + +``` +**Description:** Supported + + +### __dadd_rd +```cpp +__device__ static double __dadd_rd(double x, double y); + +``` +**Description:** Supported + + +### __dadd_rn +```cpp +__device__ static double __dadd_rn(double x, double y); + +``` +**Description:** Supported + + +### __dadd_ru +```cpp +__device__ static double __dadd_ru(double x, double y); + +``` +**Description:** Supported + + +### __dadd_rz +```cpp +__device__ static double __dadd_rz(double x, double y); + +``` +**Description:** Supported + + +### __ddiv_rd +```cpp +__device__ static double __ddiv_rd(double x, double y); + +``` +**Description:** Supported + + +### __ddiv_rn +```cpp +__device__ static double __ddiv_rn(double x, double y); + +``` +**Description:** Supported + + +### __ddiv_ru +```cpp +__device__ static double __ddiv_ru(double x, double y); + +``` +**Description:** Supported + + +### __ddiv_rz +```cpp +__device__ static double __ddiv_rz(double x, double y); + +``` +**Description:** Supported + + +### __dmul_rd +```cpp +__device__ static double __dmul_rd(double x, double y); + +``` +**Description:** Supported + + +### __dmul_rn +```cpp +__device__ static double __dmul_rn(double x, double y); + +``` +**Description:** Supported + + +### __dmul_ru +```cpp +__device__ static double __dmul_ru(double x, double y); + +``` +**Description:** Supported + + +### __dmul_rz +```cpp +__device__ static double __dmul_rz(double x, double y); + +``` +**Description:** Supported + + +### __drcp_rd +```cpp +__device__ double __drcp_rd(double x); + +``` +**Description:** Supported + + +### __drcp_rn +```cpp +__device__ double __drcp_rn(double x); + +``` +**Description:** Supported + + +### __drcp_ru +```cpp +__device__ double __drcp_ru(double x); + +``` +**Description:** Supported + + +### __drcp_rz +```cpp +__device__ double __drcp_rz(double x); + +``` +**Description:** Supported + + +### __dsqrt_rd +```cpp +__device__ double __dsqrt_rd(double x); + +``` +**Description:** Supported + + +### __dsqrt_rn +```cpp +__device__ double __dsqrt_rn(double x); + +``` +**Description:** Supported + + +### __dsqrt_ru +```cpp +__device__ double __dsqrt_ru(double x); + +``` +**Description:** Supported + + +### __dsqrt_rz +```cpp +__device__ double __dsqrt_rz(double x); + +``` +**Description:** Supported + + +### __dsub_rd +```cpp +__device__ static double __dsub_rd(double x, double y); + +``` +**Description:** Supported + + +### __dsub_rn +```cpp +__device__ static double __dsub_rn(double x, double y); + +``` +**Description:** Supported + + +### __dsub_ru +```cpp +__device__ static double __dsub_ru(double x, double y); + +``` +**Description:** Supported + + +### __dsub_rz +```cpp +__device__ static double __dsub_rz(double x, double y); + +``` +**Description:** Supported + + +### __fma_rd +```cpp +__device__ double __fma_rd(double x, double y, double z); + +``` +**Description:** Supported + + +### __fma_rn +```cpp +__device__ double __fma_rn(double x, double y, double z); + +``` +**Description:** Supported + + +### __fma_ru +```cpp +__device__ double __fma_ru(double x, double y, double z); + +``` +**Description:** Supported + + +### __fma_rz +```cpp +__device__ double __fma_rz(double x, double y, double z); + +``` +**Description:** Supported + + +### __brev +```cpp +__device__ unsigned int __brev( unsigned int x); + +``` +**Description:** Supported + + +### __brevll +```cpp +__device__ unsigned long long int __brevll( unsigned long long int x); + +``` +**Description:** Supported + + +### __byte_perm +```cpp +__device__ unsigned int __byte_perm(unsigned int x, unsigned int y, unsigned int s); + +``` +**Description:** Supported + + +### __clz +```cpp +__device__ unsigned int __clz(int x); + +``` +**Description:** Supported + + +### __clzll +```cpp +__device__ unsigned int __clzll(long long int x); + +``` +**Description:** Supported + + +### __ffs +```cpp +__device__ unsigned int __ffs(int x); + +``` +**Description:** Supported + + +### __ffsll +```cpp +__device__ unsigned int __ffsll(long long int x); + +``` +**Description:** Supported + + +### __hadd +```cpp +__device__ static unsigned int __hadd(int x, int y); + +``` +**Description:** Supported + + +### __mul24 +```cpp +__device__ static int __mul24(int x, int y); + +``` +**Description:** Supported + + +### __mul64hi +```cpp +__device__ long long int __mul64hi(long long int x, long long int y); + +``` +**Description:** Supported + + +### __mulhi +```cpp +__device__ static int __mulhi(int x, int y); + +``` +**Description:** Supported + + +### __popc +```cpp +__device__ unsigned int __popc(unsigned int x); + +``` +**Description:** Supported + + +### __popcll +```cpp +__device__ unsigned int __popcll(unsigned long long int x); + +``` +**Description:** Supported + + +### __rhadd +```cpp +__device__ static int __rhadd(int x, int y); + +``` +**Description:** Supported + + +### __sad +```cpp +__device__ static unsigned int __sad(int x, int y, int z); + +``` +**Description:** Supported + + +### __uhadd +```cpp +__device__ static unsigned int __uhadd(unsigned int x, unsigned int y); + +``` +**Description:** Supported + + +### __umul24 +```cpp +__device__ static int __umul24(unsigned int x, unsigned int y); + +``` +**Description:** Supported + + +### __umul64hi +```cpp +__device__ unsigned long long int __umul64hi(unsigned long long int x, unsigned long long int y); + +``` +**Description:** Supported + + +### __umulhi +```cpp +__device__ static unsigned int __umulhi(unsigned int x, unsigned int y); + +``` +**Description:** Supported + + +### __urhadd +```cpp +__device__ static unsigned int __urhadd(unsigned int x, unsigned int y); + +``` +**Description:** Supported + + +### __usad +```cpp +__device__ static unsigned int __usad(unsigned int x, unsigned int y, unsigned int z); + +``` +**Description:** Supported + + +### __double2float_rd +```cpp +__device__ float __double2float_rd(double x); + +``` +**Description:** Supported + + +### __double2float_rn +```cpp +__device__ float __double2float_rn(double x); + +``` +**Description:** Supported + + +### __double2float_ru +```cpp +__device__ float __double2float_ru(double x); + +``` +**Description:** Supported + + +### __double2float_rz +```cpp +__device__ float __double2float_rz(double x); + +``` +**Description:** Supported + + +### __double2hiint +```cpp +__device__ int __double2hiint(double x); + +``` +**Description:** Supported + + +### __double2int_rd +```cpp +__device__ int __double2int_rd(double x); + +``` +**Description:** Supported + + +### __double2int_rn +```cpp +__device__ int __double2int_rn(double x); + +``` +**Description:** Supported + + +### __double2int_ru +```cpp +__device__ int __double2int_ru(double x); + +``` +**Description:** Supported + + +### __double2int_rz +```cpp +__device__ int __double2int_rz(double x); + +``` +**Description:** Supported + + +### __double2ll_rd +```cpp +__device__ long long int __double2ll_rd(double x); + +``` +**Description:** Supported + + +### __double2ll_rn +```cpp +__device__ long long int __double2ll_rn(double x); + +``` +**Description:** Supported + + +### __double2ll_ru +```cpp +__device__ long long int __double2ll_ru(double x); + +``` +**Description:** Supported + + +### __double2ll_rz +```cpp +__device__ long long int __double2ll_rz(double x); + +``` +**Description:** Supported + + +### __double2loint +```cpp +__device__ int __double2loint(double x); + +``` +**Description:** Supported + + +### __double2uint_rd +```cpp +__device__ unsigned int __double2uint_rd(double x); + +``` +**Description:** Supported + + +### __double2uint_rn +```cpp +__device__ unsigned int __double2uint_rn(double x); + +``` +**Description:** Supported + + +### __double2uint_ru +```cpp +__device__ unsigned int __double2uint_ru(double x); + +``` +**Description:** Supported + + +### __double2uint_rz +```cpp +__device__ unsigned int __double2uint_rz(double x); + +``` +**Description:** Supported + + +### __double2ull_rd +```cpp +__device__ unsigned long long int __double2ull_rd(double x); + +``` +**Description:** Supported + + +### __double2ull_rn +```cpp +__device__ unsigned long long int __double2ull_rn(double x); + +``` +**Description:** Supported + + +### __double2ull_ru +```cpp +__device__ unsigned long long int __double2ull_ru(double x); + +``` +**Description:** Supported + + +### __double2ull_rz +```cpp +__device__ unsigned long long int __double2ull_rz(double x); + +``` +**Description:** Supported + + +### __double_as_longlong +```cpp +__device__ long long int __double_as_longlong(double x); + +``` +**Description:** Supported + + +### __float2half_rn +```cpp +__device__ unsigned short __float2half_rn(float x); + +``` +**Description:** Supported + + +### __half2float +```cpp +__device__ float __half2float(unsigned short); + +``` +**Description:** Supported + + +### __float2half_rn +```cpp +__device__ __half __float2half_rn(float x); + +``` +**Description:** Supported + + +### __half2float +```cpp +__device__ float __half2float(__half); + +``` +**Description:** Supported + + +### __float2int_rd +```cpp +__device__ int __float2int_rd(float x); + +``` +**Description:** Supported + + +### __float2int_rn +```cpp +__device__ int __float2int_rn(float x); + +``` +**Description:** Supported + + +### __float2int_ru +```cpp +__device__ int __float2int_ru(float x); + +``` +**Description:** Supported + + +### __float2int_rz +```cpp +__device__ int __float2int_rz(float x); + +``` +**Description:** Supported + + +### __float2ll_rd +```cpp +__device__ long long int __float2ll_rd(float x); + +``` +**Description:** Supported + + +### __float2ll_rn +```cpp +__device__ long long int __float2ll_rn(float x); + +``` +**Description:** Supported + + +### __float2ll_ru +```cpp +__device__ long long int __float2ll_ru(float x); + +``` +**Description:** Supported + + +### __float2ll_rz +```cpp +__device__ long long int __float2ll_rz(float x); + +``` +**Description:** Supported + + +### __float2uint_rd +```cpp +__device__ unsigned int __float2uint_rd(float x); + +``` +**Description:** Supported + + +### __float2uint_rn +```cpp +__device__ unsigned int __float2uint_rn(float x); + +``` +**Description:** Supported + + +### __float2uint_ru +```cpp +__device__ unsigned int __float2uint_ru(float x); + +``` +**Description:** Supported + + +### __float2uint_rz +```cpp +__device__ unsigned int __float2uint_rz(float x); + +``` +**Description:** Supported + + +### __float2ull_rd +```cpp +__device__ unsigned long long int __float2ull_rd(float x); + +``` +**Description:** Supported + + +### __float2ull_rn +```cpp +__device__ unsigned long long int __float2ull_rn(float x); + +``` +**Description:** Supported + + +### __float2ull_ru +```cpp +__device__ unsigned long long int __float2ull_ru(float x); + +``` +**Description:** Supported + + +### __float2ull_rz +```cpp +__device__ unsigned long long int __float2ull_rz(float x); + +``` +**Description:** Supported + + +### __float_as_int +```cpp +__device__ int __float_as_int(float x); + +``` +**Description:** Supported + + +### __float_as_uint +```cpp +__device__ unsigned int __float_as_uint(float x); + +``` +**Description:** Supported + + +### __hiloint2double +```cpp +__device__ double __hiloint2double(int hi, int lo); + +``` +**Description:** Supported + + +### __int2double_rn +```cpp +__device__ double __int2double_rn(int x); + +``` +**Description:** Supported + + +### __int2float_rd +```cpp +__device__ float __int2float_rd(int x); + +``` +**Description:** Supported + + +### __int2float_rn +```cpp +__device__ float __int2float_rn(int x); + +``` +**Description:** Supported + + +### __int2float_ru +```cpp +__device__ float __int2float_ru(int x); + +``` +**Description:** Supported + + +### __int2float_rz +```cpp +__device__ float __int2float_rz(int x); + +``` +**Description:** Supported + + +### __int_as_float +```cpp +__device__ float __int_as_float(int x); + +``` +**Description:** Supported + + +### __ll2double_rd +```cpp +__device__ double __ll2double_rd(long long int x); + +``` +**Description:** Supported + + +### __ll2double_rn +```cpp +__device__ double __ll2double_rn(long long int x); + +``` +**Description:** Supported + + +### __ll2double_ru +```cpp +__device__ double __ll2double_ru(long long int x); + +``` +**Description:** Supported + + +### __ll2double_rz +```cpp +__device__ double __ll2double_rz(long long int x); + +``` +**Description:** Supported + + +### __ll2float_rd +```cpp +__device__ float __ll2float_rd(long long int x); + +``` +**Description:** Supported + + +### __ll2float_rn +```cpp +__device__ float __ll2float_rn(long long int x); + +``` +**Description:** Supported + + +### __ll2float_ru +```cpp +__device__ float __ll2float_ru(long long int x); + +``` +**Description:** Supported + + +### __ll2float_rz +```cpp +__device__ float __ll2float_rz(long long int x); + +``` +**Description:** Supported + + +### __longlong_as_double +```cpp +__device__ double __longlong_as_double(long long int x); + +``` +**Description:** Supported + + +### __uint2double_rn +```cpp +__device__ double __uint2double_rn(int x); + +``` +**Description:** Supported + + +### __uint2float_rd +```cpp +__device__ float __uint2float_rd(unsigned int x); + +``` +**Description:** Supported + + +### __uint2float_rn +```cpp +__device__ float __uint2float_rn(unsigned int x); + +``` +**Description:** Supported + + +### __uint2float_ru +```cpp +__device__ float __uint2float_ru(unsigned int x); + +``` +**Description:** Supported + + +### __uint2float_rz +```cpp +__device__ float __uint2float_rz(unsigned int x); + +``` +**Description:** Supported + + +### __uint_as_float +```cpp +__device__ float __uint_as_float(unsigned int x); + +``` +**Description:** Supported + + +### __ull2double_rd +```cpp +__device__ double __ull2double_rd(unsigned long long int x); + +``` +**Description:** Supported + + +### __ull2double_rn +```cpp +__device__ double __ull2double_rn(unsigned long long int x); + +``` +**Description:** Supported + + +### __ull2double_ru +```cpp +__device__ double __ull2double_ru(unsigned long long int x); + +``` +**Description:** Supported + + +### __ull2double_rz +```cpp +__device__ double __ull2double_rz(unsigned long long int x); + +``` +**Description:** Supported + + +### __ull2float_rd +```cpp +__device__ float __ull2float_rd(unsigned long long int x); + +``` +**Description:** Supported + + +### __ull2float_rn +```cpp +__device__ float __ull2float_rn(unsigned long long int x); + +``` +**Description:** Supported + + +### __ull2float_ru +```cpp +__device__ float __ull2float_ru(unsigned long long int x); + +``` +**Description:** Supported + + +### __ull2float_rz +```cpp +__device__ float __ull2float_rz(unsigned long long int x); + +``` +**Description:** Supported + + +### __heq +```cpp +__device__ bool __heq(__half a, __half b); + +``` +**Description:** Supported + + +### __hge +```cpp +__device__ bool __hge(__half a, __half b); + +``` +**Description:** Supported + + +### __hgt +```cpp +__device__ bool __hgt(__half a, __half b); + +``` +**Description:** Supported + + +### __hisinf +```cpp +__device__ bool __hisinf(__half a); + +``` +**Description:** Supported + + +### __hisnan +```cpp +__device__ bool __hisnan(__half a); + +``` +**Description:** Supported + + +### __hle +```cpp +__device__ bool __hle(__half a, __half b); + +``` +**Description:** Supported + + +### __hlt +```cpp +__device__ bool __hlt(__half a, __half b); + +``` +**Description:** Supported + + +### __hne +```cpp +__device__ bool __hne(__half a, __half b); + +``` +**Description:** Supported + + +### __hbeq2 +```cpp +__device__ bool __hbeq2(__half2 a, __half2 b); + +``` +**Description:** Supported + + +### __hbge2 +```cpp +__device__ bool __hbge2(__half2 a, __half2 b); + +``` +**Description:** Supported + + +### __hbgt2 +```cpp +__device__ bool __hbgt2(__half2 a, __half2 b); + +``` +**Description:** Supported + + +### __hble2 +```cpp +__device__ bool __hble2(__half2 a, __half2 b); + +``` +**Description:** Supported + + +### __hblt2 +```cpp +__device__ bool __hblt2(__half2 a, __half2 b); + +``` +**Description:** Supported + + +### __hbne2 +```cpp +__device__ bool __hbne2(__half2 a, __half2 b); + +``` +**Description:** Supported + + +### __heq2 +```cpp +__device__ __half2 __heq2(__half2 a, __half2 b); + +``` +**Description:** Supported + + +### __hge2 +```cpp +__device__ __half2 __hge2(__half2 a, __half2 b); + +``` +**Description:** Supported + + +### __hgt2 +```cpp +__device__ __half2 __hgt2(__half2 a, __half2 b); + +``` +**Description:** Supported + + +### __hisnan2 +```cpp +__device__ __half2 __hisnan2(__half2 a); + +``` +**Description:** Supported + + +### __hle2 +```cpp +__device__ __half2 __hle2(__half2 a, __half2 b); + +``` +**Description:** Supported + + +### __hlt2 +```cpp +__device__ __half2 __hlt2(__half2 a, __half2 b); + +``` +**Description:** Supported + + +### __hne2 +```cpp +__device__ __half2 __hne2(__half2 a, __half2 b); + +``` +**Description:** Supported + + +### __float22half2_rn +```cpp +__device__ __half2 __float22half2_rn(const float2 a); + +``` +**Description:** Supported + + +### __float2half +```cpp +__device__ __half __float2half(const float a); + +``` +**Description:** Supported + + +### __float2half2_rn +```cpp +__device__ __half2 __float2half2_rn(const float a); + +``` +**Description:** Supported + + +### __float2half_rd +```cpp +__device__ __half __float2half_rd(const float a); + +``` +**Description:** Supported + + +### __float2half_rn +```cpp +__device__ __half __float2half_rn(const float a); + +``` +**Description:** Supported + + +### __float2half_ru +```cpp +__device__ __half __float2half_ru(const float a); + +``` +**Description:** Supported + + +### __float2half_rz +```cpp +__device__ __half __float2half_rz(const float a); + +``` +**Description:** Supported + + +### __floats2half2_rn +```cpp +__device__ __half2 __floats2half2_rn(const float a, const float b); + +``` +**Description:** Supported + + +### __half22float2 +```cpp +__device__ float2 __half22float2(const __half2 a); + +``` +**Description:** Supported + + +### __half2float +```cpp +__device__ float __half2float(const __half a); + +``` +**Description:** Supported + + +### half2half2 +```cpp +__device__ __half2 half2half2(const __half a); + +``` +**Description:** Supported + + +### __half2int_rd +```cpp +__device__ int __half2int_rd(__half h); + +``` +**Description:** Supported + + +### __half2int_rn +```cpp +__device__ int __half2int_rn(__half h); + +``` +**Description:** Supported + + +### __half2int_ru +```cpp +__device__ int __half2int_ru(__half h); + +``` +**Description:** Supported + + +### __half2int_rz +```cpp +__device__ int __half2int_rz(__half h); + +``` +**Description:** Supported + + +### __half2ll_rd +```cpp +__device__ long long int __half2ll_rd(__half h); + +``` +**Description:** Supported + + +### __half2ll_rn +```cpp +__device__ long long int __half2ll_rn(__half h); + +``` +**Description:** Supported + + +### __half2ll_ru +```cpp +__device__ long long int __half2ll_ru(__half h); + +``` +**Description:** Supported + + +### __half2ll_rz +```cpp +__device__ long long int __half2ll_rz(__half h); + +``` +**Description:** Supported + + +### __half2short_rd +```cpp +__device__ short __half2short_rd(__half h); + +``` +**Description:** Supported + + +### __half2short_rn +```cpp +__device__ short __half2short_rn(__half h); + +``` +**Description:** Supported + + +### __half2short_ru +```cpp +__device__ short __half2short_ru(__half h); + +``` +**Description:** Supported + + +### __half2short_rz +```cpp +__device__ short __half2short_rz(__half h); + +``` +**Description:** Supported + + +### __half2uint_rd +```cpp +__device__ unsigned int __half2uint_rd(__half h); + +``` +**Description:** Supported + + +### __half2uint_rn +```cpp +__device__ unsigned int __half2uint_rn(__half h); + +``` +**Description:** Supported + + +### __half2uint_ru +```cpp +__device__ unsigned int __half2uint_ru(__half h); + +``` +**Description:** Supported + + +### __half2uint_rz +```cpp +__device__ unsigned int __half2uint_rz(__half h); + +``` +**Description:** Supported + + +### __half2ull_rd +```cpp +__device__ unsigned long long int __half2ull_rd(__half h); + +``` +**Description:** Supported + + +### __half2ull_rn +```cpp +__device__ unsigned long long int __half2ull_rn(__half h); + +``` +**Description:** Supported + + +### __half2ull_ru +```cpp +__device__ unsigned long long int __half2ull_ru(__half h); + +``` +**Description:** Supported + + +### __half2ull_rz +```cpp +__device__ unsigned long long int __half2ull_rz(__half h); + +``` +**Description:** Supported + + +### __half2ushort_rd +```cpp +__device__ unsigned short int __half2ushort_rd(__half h); + +``` +**Description:** Supported + + +### __half2ushort_rn +```cpp +__device__ unsigned short int __half2ushort_rn(__half h); + +``` +**Description:** Supported + + +### __half2ushort_ru +```cpp +__device__ unsigned short int __half2ushort_ru(__half h); + +``` +**Description:** Supported + + +### __half2ushort_rz +```cpp +__device__ unsigned short int __half2ushort_rz(__half h); + +``` +**Description:** Supported + + +### __half_as_short +```cpp +__device__ short int __half_as_short(const __half h); + +``` +**Description:** Supported + + +### __half_as_ushort +```cpp +__device__ unsigned short int __half_as_ushort(const __half h); + +``` +**Description:** Supported + + +### __halves2half2 +```cpp +__device__ __half2 __halves2half2(const __half a, const __half b); + +``` +**Description:** Supported + + +### __high2float +```cpp +__device__ float __high2float(const __half2 a); + +``` +**Description:** Supported + + +### __high2half +```cpp +__device__ __half __high2half(const __half2 a); + +``` +**Description:** Supported + + +### __high2half2 +```cpp +__device__ __half2 __high2half2(const __half2 a); + +``` +**Description:** Supported + + +### __highs2half2 +```cpp +__device__ __half2 __highs2half2(const __half2 a, const __half2 b); + +``` +**Description:** Supported + + +### __int2half_rd +```cpp +__device__ __half __int2half_rd(int i); + +``` +**Description:** Supported + + +### __int2half_rn +```cpp +__device__ __half __int2half_rn(int i); + +``` +**Description:** Supported + + +### __int2half_ru +```cpp +__device__ __half __int2half_ru(int i); + +``` +**Description:** Supported + + +### __int2half_rz +```cpp +__device__ __half __int2half_rz(int i); + +``` +**Description:** Supported + + +### __ll2half_rd +```cpp +__device__ __half __ll2half_rd(long long int i); + +``` +**Description:** Supported + + +### __ll2half_rn +```cpp +__device__ __half __ll2half_rn(long long int i); + +``` +**Description:** Supported + + +### __ll2half_ru +```cpp +__device__ __half __ll2half_ru(long long int i); + +``` +**Description:** Supported + + +### __ll2half_rz +```cpp +__device__ __half __ll2half_rz(long long int i); + +``` +**Description:** Supported + + +### __low2float +```cpp +__device__ float __low2float(const __half2 a); + +``` +**Description:** Supported + + +### __low2half +```cpp +__device__ __half __low2half(const __half2 a); + +``` +**Description:** Supported + + +### __low2half2 +```cpp +__device__ __half2 __low2half2(const __half2 a, const __half2 b); + +``` +**Description:** Supported + + +### __low2half2 +```cpp +__device__ __half2 __low2half2(const __half2 a); + +``` +**Description:** Supported + + +### __lowhigh2highlow +```cpp +__device__ __half2 __lowhigh2highlow(const __half2 a); + +``` +**Description:** Supported + + +### __lows2half2 +```cpp +__device__ __half2 __lows2half2(const __half2 a, const __half2 b); + +``` +**Description:** Supported + + +### __short2half_rd +```cpp +__device__ __half __short2half_rd(short int i); + +``` +**Description:** Supported + + +### __short2half_rn +```cpp +__device__ __half __short2half_rn(short int i); + +``` +**Description:** Supported + + +### __short2half_ru +```cpp +__device__ __half __short2half_ru(short int i); + +``` +**Description:** Supported + + +### __short2half_rz +```cpp +__device__ __half __short2half_rz(short int i); + +``` +**Description:** Supported + + +### __uint2half_rd +```cpp +__device__ __half __uint2half_rd(unsigned int i); + +``` +**Description:** Supported + + +### __uint2half_rn +```cpp +__device__ __half __uint2half_rn(unsigned int i); + +``` +**Description:** Supported + + +### __uint2half_ru +```cpp +__device__ __half __uint2half_ru(unsigned int i); + +``` +**Description:** Supported + + +### __uint2half_rz +```cpp +__device__ __half __uint2half_rz(unsigned int i); + +``` +**Description:** Supported + + +### __ull2half_rd +```cpp +__device__ __half __ull2half_rd(unsigned long long int i); + +``` +**Description:** Supported + + +### __ull2half_rn +```cpp +__device__ __half __ull2half_rn(unsigned long long int i); + +``` +**Description:** Supported + + +### __ull2half_ru +```cpp +__device__ __half __ull2half_ru(unsigned long long int i); + +``` +**Description:** Supported + + +### __ull2half_rz +```cpp +__device__ __half __ull2half_rz(unsigned long long int i); + +``` +**Description:** Supported + + +### __ushort2half_rd +```cpp +__device__ __half __ushort2half_rd(unsigned short int i); + +``` +**Description:** Supported + + +### __ushort2half_rn +```cpp +__device__ __half __ushort2half_rn(unsigned short int i); + +``` +**Description:** Supported + + +### __ushort2half_ru +```cpp +__device__ __half __ushort2half_ru(unsigned short int i); + +``` +**Description:** Supported + + +### __ushort2half_rz +```cpp +__device__ __half __ushort2half_rz(unsigned short int i); + +``` +**Description:** Supported + + +### __ushort_as_half +```cpp +__device__ __half __ushort_as_half(const unsigned short int i); + +``` +**Description:** Supported + + diff --git a/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h b/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h index a2894f3d9b..212f22d5dc 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h @@ -23,6 +23,86 @@ THE SOFTWARE. #include #include + + + + +// Single Precision Fast Math +__device__ float __cosf(float x); +__device__ float __exp10f(float x); +__device__ float __expf(float x); +__device__ static float __fadd_rd(float x, float y); +__device__ static float __fadd_rn(float x, float y); +__device__ static float __fadd_ru(float x, float y); +__device__ static float __fadd_rz(float x, float y); +__device__ static float __fdiv_rd(float x, float y); +__device__ static float __fdiv_rn(float x, float y); +__device__ static float __fdiv_ru(float x, float y); +__device__ static float __fdiv_rz(float x, float y); +__device__ static float __fdividef(float x, float y); +__device__ float __fmaf_rd(float x, float y, float z); +__device__ float __fmaf_rn(float x, float y, float z); +__device__ float __fmaf_ru(float x, float y, float z); +__device__ float __fmaf_rz(float x, float y, float z); +__device__ static float __fmul_rd(float x, float y); +__device__ static float __fmul_rn(float x, float y); +__device__ static float __fmul_ru(float x, float y); +__device__ static float __fmul_rz(float x, float y); +__device__ float __frcp_rd(float x); +__device__ float __frcp_rn(float x); +__device__ float __frcp_ru(float x); +__device__ float __frcp_rz(float x); +__device__ float __frsqrt_rn(float x); +__device__ float __fsqrt_rd(float x); +__device__ float __fsqrt_rn(float x); +__device__ float __fsqrt_ru(float x); +__device__ float __fsqrt_rz(float x); +__device__ static float __fsub_rd(float x, float y); +__device__ static float __fsub_rn(float x, float y); +__device__ static float __fsub_ru(float x, float y); +__device__ float __log10f(float x); +__device__ float __log2f(float x); +__device__ float __logf(float x); +__device__ float __powf(float base, float exponent); +__device__ static float __saturatef(float x); +__device__ void __sincosf(float x, float *s, float *c); +__device__ float __sinf(float x); +__device__ float __tanf(float x); + + +/* +Double Precision Intrinsics +*/ + +__device__ static double __dadd_rd(double x, double y); +__device__ static double __dadd_rn(double x, double y); +__device__ static double __dadd_ru(double x, double y); +__device__ static double __dadd_rz(double x, double y); +__device__ static double __ddiv_rd(double x, double y); +__device__ static double __ddiv_rn(double x, double y); +__device__ static double __ddiv_ru(double x, double y); +__device__ static double __ddiv_rz(double x, double y); +__device__ static double __dmul_rd(double x, double y); +__device__ static double __dmul_rn(double x, double y); +__device__ static double __dmul_ru(double x, double y); +__device__ static double __dmul_rz(double x, double y); +__device__ double __drcp_rd(double x); +__device__ double __drcp_rn(double x); +__device__ double __drcp_ru(double x); +__device__ double __drcp_rz(double x); +__device__ double __dsqrt_rd(double x); +__device__ double __dsqrt_rn(double x); +__device__ double __dsqrt_ru(double x); +__device__ double __dsqrt_rz(double x); +__device__ static double __dsub_rd(double x, double y); +__device__ static double __dsub_rn(double x, double y); +__device__ static double __dsub_ru(double x, double y); +__device__ static double __dsub_rz(double x, double y); +__device__ double __fma_rd(double x, double y, double z); +__device__ double __fma_rn(double x, double y, double z); +__device__ double __fma_ru(double x, double y, double z); +__device__ double __fma_rz(double x, double y, double z); + // Single Precision Fast Math extern __attribute__((const)) float __hip_fast_cosf(float) __asm("llvm.cos.f32"); extern __attribute__((const)) float __hip_fast_exp2f(float) __asm("llvm.exp2.f32"); @@ -349,6 +429,21 @@ __device__ unsigned int __clz(int x); __device__ unsigned int __clzll(long long int x); __device__ unsigned int __ffs(int x); __device__ unsigned int __ffsll(long long int x); +__device__ static unsigned int __hadd(int x, int y); +__device__ static int __mul24(int x, int y); +__device__ long long int __mul64hi(long long int x, long long int y); +__device__ static int __mulhi(int x, int y); +__device__ unsigned int __popc(unsigned int x); +__device__ unsigned int __popcll(unsigned long long int x); +__device__ static int __rhadd(int x, int y); +__device__ static unsigned int __sad(int x, int y, int z); +__device__ static unsigned int __uhadd(unsigned int x, unsigned int y); +__device__ static int __umul24(unsigned int x, unsigned int y); +__device__ unsigned long long int __umul64hi(unsigned long long int x, unsigned long long int y); +__device__ static unsigned int __umulhi(unsigned int x, unsigned int y); +__device__ static unsigned int __urhadd(unsigned int x, unsigned int y); +__device__ static unsigned int __usad(unsigned int x, unsigned int y, unsigned int z); + __device__ static inline unsigned int __hadd(int x, int y) { int z = x + y; int sign = z & 0x8000000; @@ -358,12 +453,9 @@ __device__ static inline unsigned int __hadd(int x, int y) { __device__ static inline int __mul24(int x, int y) { return __hip_hc_ir_mul24_int(x, y); } -__device__ long long int __mul64hi(long long int x, long long int y); __device__ static inline int __mulhi(int x, int y) { return __hip_hc_ir_mulhi_int(x, y); } -__device__ unsigned int __popc( unsigned int x); -__device__ unsigned int __popcll( unsigned long long int x); __device__ static inline int __rhadd(int x, int y) { int z = x + y + 1; int sign = z & 0x8000000; @@ -373,14 +465,12 @@ __device__ static inline int __rhadd(int x, int y) { __device__ static inline unsigned int __sad(int x, int y, int z) { return x > y ? x - y + z : y - x + z; } - __device__ static inline unsigned int __uhadd(unsigned int x, unsigned int y) { return (x + y) >> 1; } __device__ static inline int __umul24(unsigned int x, unsigned int y) { return __hip_hc_ir_umul24_int(x, y); } -__device__ unsigned long long int __umul64hi(unsigned long long int x, unsigned long long int y); __device__ static inline unsigned int __umulhi(unsigned int x, unsigned int y) { return __hip_hc_ir_umulhi_int(x, y); } @@ -440,10 +530,10 @@ CUDA implements half as unsigned short whereas, HIP doesn't. */ -__device__ int float2int_rd(float x); -__device__ int float2int_rn(float x); -__device__ int float2int_ru(float x); -__device__ int float2int_rz(float x); +__device__ int __float2int_rd(float x); +__device__ int __float2int_rn(float x); +__device__ int __float2int_ru(float x); +__device__ int __float2int_rz(float x); __device__ long long int __float2ll_rd(float x); __device__ long long int __float2ll_rn(float x); diff --git a/projects/clr/hipamd/include/hip/hcc_detail/math_functions.h b/projects/clr/hipamd/include/hip/hcc_detail/math_functions.h index 21ec4510c6..b07830958f 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/math_functions.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/math_functions.h @@ -141,8 +141,8 @@ __device__ double copysign(double x, double y); __device__ double cos(double x); __device__ double cosh(double x); __device__ __host__ double cospi(double x); -__device__ double cyl_bessel_i0(double x); -__device__ double cyl_bessel_i1(double x); +//__device__ double cyl_bessel_i0(double x); +//__device__ double cyl_bessel_i1(double x); __device__ double erf(double x); __device__ double erfc(double x); __device__ double erfcinv(double y); @@ -232,6 +232,8 @@ __host__ double rnorm3d(double a, double b, double c); __host__ double rnorm4d(double a, double b, double c, double d); __host__ void sincospi(double x, double *sptr, double *cptr); +// ENDPARSER + #ifdef HIP_FAST_MATH // Single Precision Precise Math when enabled