diff --git a/projects/hip/docs/old/reference/kernel_language.rst b/projects/hip/docs/old/reference/kernel_language.rst index 7f2f123ac5..469b7d6c97 100644 --- a/projects/hip/docs/old/reference/kernel_language.rst +++ b/projects/hip/docs/old/reference/kernel_language.rst @@ -183,6 +183,7 @@ To allow the host to dynamically allocate shared memory, you can specify ``exter launch parameter. .. note:: + Prior to the HIP-Clang compiler, dynamic shared memory had to be declared using the ``HIP_DYNAMIC_SHARED`` macro in order to ensure accuracy. This is because using static shared memory in the same kernel could've resulted in overlapping memory ranges and data-races. The @@ -226,8 +227,8 @@ The ``warpSize`` variable type is ``int``. It contains the warp size (in threads ``warpSize`` should only be used in device functions that develop portable wave-aware code. .. note:: - NVIDIA devices return 32 for this variable; AMD devices return 64 for gfx9 and 32 for gfx10 - and above. + + NVIDIA devices return 32 for this variable; AMD devices return 64 for gfx9 and 32 for gfx10 and above. Vector types ==================================================== @@ -302,20 +303,1003 @@ the HIP-Clang path, you can use the following workaround: ``hipHostMalloc()``. * Remove ``memcpy`` for all allocated fine-grained system memory regions. -.. _synchronization-functions: +.. _synchronization functions: Synchronization functions ==================================================== The ``__syncthreads()`` built-in function is supported in HIP. The ``__syncthreads_count(int)``, ``__syncthreads_and(int)``, and ``__syncthreads_or(int)`` functions are under development. + Math functions ==================================================== + HIP-Clang supports a set of math operations that are callable from the device. These are described in the following sections. -.. doxygengroup:: Math - :inner: +Single precision mathematical functions +-------------------------------------------------------------------------------------------- +Following is the list of supported single precision mathematical functions. + +.. list-table:: Single precision mathematical functions + + * - **Function** + - **Supported on Host** + - **Supported on Device** + + * - | float acosf ( float x ) + | Calculate the arc cosine of the input argument. + - ✓ + - ✓ + + * - | float acoshf ( float x ) + | Calculate the nonnegative arc hyperbolic cosine of the input argument. + - ✓ + - ✓ + + * - | float asinf ( float x ) + | Calculate the arc sine of the input argument + - ✓ + - ✓ + + * - | float asinhf ( float x ) + | Calculate the arc hyperbolic sine of the input argument. + - ✓ + - ✓ + + * - | float atan2f ( float y, float x ) + | Calculate the arc tangent of the ratio of first and second input arguments. + - ✓ + - ✓ + + * - | float atanf ( float x ) + | Calculate the arc tangent of the input argument. + - ✓ + - ✓ + + * - | float atanhf ( float x ) + | Calculate the arc hyperbolic tangent of the input argument. + - ✓ + - ✓ + + * - | float cbrtf ( float x ) + | Calculate the cube root of the input argument. + - ✓ + - ✓ + + * - | float ceilf ( float x ) + | Calculate ceiling of the input argument. + - ✓ + - ✓ + + * - | float copysignf ( float x, float y ) + | Create value with given magnitude, copying sign of second value. + - ✓ + - ✓ + + * - | float copysignf ( float x, float y ) + | Create value with given magnitude, copying sign of second value. + - ✓ + - ✓ + + * - | float cosf ( float x ) + | Calculate the cosine of the input argument. + - ✓ + - ✓ + + * - | float coshf ( float x ) + | Calculate the hyperbolic cosine of the input argument. + - ✓ + - ✓ + * - | float erfcf ( float x ) + | Calculate the complementary error function of the input argument. + - ✓ + - ✓ + + * - | float erff ( float x ) + | Calculate the error function of the input argument. + - ✓ + - ✓ + + * - | float exp10f ( float x ) + | Calculate the base 10 exponential of the input argument. + - ✓ + - ✓ + + * - | float exp2f ( float x ) + | Calculate the base 2 exponential of the input argument. + - ✓ + - ✓ + + * - | float expf ( float x ) + | Calculate the base e exponential of the input argument. + - ✓ + - ✓ + + * - | float expm1f ( float x ) + | Calculate the base e exponential of the input argument, minus 1. + - ✓ + - ✓ + + * - | float fabsf ( float x ) + | Calculate the absolute value of its argument. + - ✓ + - ✓ + + * - | float fdimf ( float x, float y ) + | Compute the positive difference between `x` and `y`. + - ✓ + - ✓ + + * - | float floorf ( float x ) + | Calculate the largest integer less than or equal to `x`. + - ✓ + - ✓ + + * - | float fmaf ( float x, float y, float z ) + | Compute `x × y + z` as a single operation. + - ✓ + - ✓ + + * - | float fmaxf ( float x, float y ) + | Determine the maximum numeric value of the arguments. + - ✓ + - ✓ + + * - | float fminf ( float x, float y ) + | Determine the minimum numeric value of the arguments. + - ✓ + - ✓ + + * - | float fmodf ( float x, float y ) + | Calculate the floating-point remainder of `x / y`. + - ✓ + - ✓ + + * - | float frexpf ( float x, int* nptr ) + | Extract mantissa and exponent of a floating-point value. + - ✓ + - ✗ + + * - | float hypotf ( float x, float y ) + | Calculate the square root of the sum of squares of two arguments. + - ✓ + - ✓ + + * - | int ilogbf ( float x ) + | Compute the unbiased integer exponent of the argument. + - ✓ + - ✓ + + * - | __RETURN_TYPE isfinite ( float a ) + | Determine whether argument is finite. + - ✓ + - ✓ + + * - | __RETURN_TYPE isinf ( float a ) + | Determine whether argument is infinite. + - ✓ + - ✓ + + * - | __RETURN_TYPE isnan ( float a ) + | Determine whether argument is a NaN. + - ✓ + - ✓ + + * - | float ldexpf ( float x, int exp ) + | Calculate the value of x ⋅ 2 of the exponent of the input argument. + - ✓ + - ✓ + + * - | loat log10f ( float x ) + | Calculate the base 10 logarithm of the input argument. + - ✓ + - ✓ + + * - | float log1pf ( float x ) + | Calculate the value of the exponent of the input argument + - ✓ + - ✓ + + * - | float logbf ( float x ) + | Calculate the floating point representation of the exponent of the input argument. + - ✓ + - ✓ + + * - | float log2f ( float x ) + | Calculate the base 2 logarithm of the input argument. + - ✓ + - ✓ + + * - | float logf ( float x ) + | Calculate the natural logarithm of the input argument. + - ✓ + - ✓ + + * - | float modff ( float x, float* iptr ) + | Break down the input argument into fractional and integral parts. + - ✓ + - ✗ + + * - | float nanf ( const char* tagp ) + | Returns "Not a Number" value. + - ✗ + - ✓ + + * - | float nearbyintf ( float x ) + | Round the input argument to the nearest integer. + - ✓ + - ✓ + + * - | float powf ( float x, float y ) + | Calculate the value of first argument to the power of second argument. + - ✓ + - ✓ + + * - | float remainderf ( float x, float y ) + | Compute single-precision floating-point remainder. + - ✓ + - ✓ + + * - | float remquof ( float x, float y, int* quo ) + | Compute single-precision floating-point remainder and part of quotient. + - ✓ + - ✗ + + * - | float roundf ( float x ) + | Round to nearest integer value in floating-point. + - ✓ + - ✓ + + * - | float scalbnf ( float x, int n ) + | Scale floating-point input by integer power of two. + - ✓ + - ✓ + + * - | __RETURN_TYPE signbit ( float a ) + | Return the sign bit of the input. + - ✓ + - ✓ + + * - | void sincosf ( float x, float* sptr, float* cptr ) + | Calculate the sine and cosine of the first input argument. + - ✓ + - ✗ + + * - | float sinf ( float x ) + | Calculate the sine of the input argument. + - ✓ + - ✓ + + * - | float sinhf ( float x ) + | Calculate the hyperbolic sine of the input argument. + - ✓ + - ✓ + + * - | float sqrtf ( float x ) + | Calculate the square root of the input argument. + - ✓ + - ✓ + + * - | float tanf ( float x ) + | Calculate the tangent of the input argument. + - ✓ + - ✓ + + * - | float tanhf ( float x ) + | Calculate the hyperbolic tangent of the input argument. + - ✓ + - ✓ + + * - | float truncf ( float x ) + | Truncate input argument to the integral part. + - ✓ + - ✓ + + * - | float tgammaf ( float x ) + | Calculate the gamma function of the input argument. + - ✓ + - ✓ + + * - | float erfcinvf ( float y ) + | Calculate the inverse complementary function of the input argument. + - ✓ + - ✓ + + * - | float erfcxf ( float x ) + | Calculate the scaled complementary error function of the input argument. + - ✓ + - ✓ + + * - | float erfinvf ( float y ) + | Calculate the inverse error function of the input argument. + - ✓ + - ✓ + + * - | float fdividef ( float x, float y ) + | Divide two floating point values. + - ✓ + - ✓ + + * - | float frexpf ( float x, `int *nptr` ) + | Extract mantissa and exponent of a floating-point value. + - ✓ + - ✓ + + * - | float j0f ( float x ) + | Calculate the value of the Bessel function of the first kind of order 0 for the input argument. + - ✓ + - ✓ + + * - | float j1f ( float x ) + | Calculate the value of the Bessel function of the first kind of order 1 for the input argument. + - ✓ + - ✓ + + * - | float jnf ( int n, float x ) + | Calculate the value of the Bessel function of the first kind of order n for the input argument. + - ✓ + - ✓ + + * - | float lgammaf ( float x ) + | Calculate the natural logarithm of the absolute value of the gamma function of the input argument. + - ✓ + - ✓ + + * - | long long int llrintf ( float x ) + | Round input to nearest integer value. + - ✓ + - ✓ + + * - | long long int llroundf ( float x ) + | Round to nearest integer value. + - ✓ + - ✓ + + * - | long int lrintf ( float x ) + | Round input to nearest integer value. + - ✓ + - ✓ + + * - | long int lroundf ( float x ) + | Round to nearest integer value. + - ✓ + - ✓ + + * - | float modff ( float x, `float *iptr` ) + | Break down the input argument into fractional and integral parts. + - ✓ + - ✓ + + * - | float nextafterf ( float x, float y ) + | Returns next representable single-precision floating-point value after argument. + - ✓ + - ✓ + + * - | float norm3df ( float a, float b, float c ) + | Calculate the square root of the sum of squares of three coordinates of the argument. + - ✓ + - ✓ + + * - | float norm4df ( float a, float b, float c, float d ) + | Calculate the square root of the sum of squares of four coordinates of the argument. + - ✓ + - ✓ + + * - | loat normcdff ( float y ) + | Calculate the standard normal cumulative distribution function. + - ✓ + - ✓ + + * - | float normcdfinvf ( float y ) + | Calculate the inverse of the standard normal cumulative distribution function. + - ✓ + - ✓ + + * - | float normf ( int dim, `const float *a` ) + | Calculate the square root of the sum of squares of any number of coordinates. + - ✓ + - ✓ + + * - | float rcbrtf ( float x ) + | Calculate the reciprocal cube root function. + - ✓ + - ✓ + + * - | float remquof ( float x, float y, `int *quo` ) + | Compute single-precision floating-point remainder and part of quotient. + - ✓ + - ✓ + + * - | float rhypotf ( float x, float y ) + | Calculate one over the square root of the sum of squares of two arguments. + - ✓ + - ✓ + + * - | float rintf ( float x ) + | Round input to nearest integer value in floating-point. + - ✓ + - ✓ + + * - | float rnorm3df ( float a, float b, float c ) + | Calculate one over the square root of the sum of squares of three coordinates of the argument. + - ✓ + - ✓ + + * - | float rnorm4df ( float a, float b, float c, float d ) + | Calculate one over the square root of the sum of squares of four coordinates of the argument. + - ✓ + - ✓ + + * - | float rnormf ( int dim, `const float *a` ) + | Calculate the reciprocal of square root of the sum of squares of any number of coordinates. + - ✓ + - ✓ + + * - | float scalblnf ( float x, long int n ) + | Scale floating-point input by integer power of two. + - ✓ + - ✓ + + * - | void sincosf ( float x, `float *sptr`, `float *cptr`) + | Calculate the sine and cosine of the first input argument. + - ✓ + - ✓ + + * - | void sincospif ( float x, `float *sptr`, `float *cptr`) + | Calculate the sine and cosine of the first input argument multiplied by PI. + - ✓ + - ✓ + + * - | float y0f ( float x ) + | Calculate the value of the Bessel function of the second kind of order 0 for the input argument. + - ✓ + - ✓ + + * - | float y1f ( float x ) + | Calculate the value of the Bessel function of the second kind of order 1 for the input argument. + - ✓ + - ✓ + + * - | float ynf ( int n, float x ) + | Calculate the value of the Bessel function of the second kind of order n for the input argument. + - ✓ + - ✓ + +Double precision mathematical functions +-------------------------------------------------------------------------------------------- +Following is the list of supported double precision mathematical functions. + +.. list-table:: Single precision mathematical functions + + * - **Function** + - **Supported on Host** + - **Supported on Device** + + * - | double acos ( double x ) + | Calculate the arc cosine of the input argument. + - ✓ + - ✓ + + * - | double acosh ( double x ) + | Calculate the nonnegative arc hyperbolic cosine of the input argument. + - ✓ + - ✓ + + * - | double asin ( double x ) + | Calculate the arc sine of the input argument. + - ✓ + - ✓ + + * - | double asinh ( double x ) + | Calculate the arc hyperbolic sine of the input argument. + - ✓ + - ✓ + + * - | double atan ( double x ) + | Calculate the arc tangent of the input argument. + - ✓ + - ✓ + + * - | double atan2 ( double y, double x ) + | Calculate the arc tangent of the ratio of first and second input arguments. + - ✓ + - ✓ + + * - | double atanh ( double x ) + | Calculate the arc hyperbolic tangent of the input argument. + - ✓ + - ✓ + + * - | double cbrt ( double x ) + | Calculate the cube root of the input argument. + - ✓ + - ✓ + + * - | double ceil ( double x ) + | Calculate ceiling of the input argument. + - ✓ + - ✓ + + * - | double copysign ( double x, double y ) + | Create value with given magnitude, copying sign of second value. + - ✓ + - ✓ + + * - | double cos ( double x ) + | Calculate the cosine of the input argument. + - ✓ + - ✓ + + * - | double cosh ( double x ) + | Calculate the hyperbolic cosine of the input argument. + - ✓ + - ✓ + + * - | double erf ( double x ) + | Calculate the error function of the input argument. + - ✓ + - ✓ + + * - | double erfc ( double x ) + | Calculate the complementary error function of the input argument. + - ✓ + - ✓ + + * - | double exp ( double x ) + | Calculate the base e exponential of the input argument. + - ✓ + - ✓ + + * - | double exp10 ( double x ) + | Calculate the base 10 exponential of the input argument. + - ✓ + - ✓ + + * - | double exp2 ( double x ) + | Calculate the base 2 exponential of the input argument. + - ✓ + - ✓ + + * - | double expm1 ( double x ) + | Calculate the base e exponential of the input argument, minus 1. + - ✓ + - ✓ + + * - | double fabs ( double x ) + | Calculate the absolute value of the input argument. + - ✓ + - ✓ + + * - | double fdim ( double x, double y ) + | Compute the positive difference between `x` and `y`. + - ✓ + - ✓ + + * - | double floor ( double x ) + | Calculate the largest integer less than or equal to `x`. + - ✓ + - ✓ + + * - | double fma ( double x, double y, double z ) + | Compute `x × y + z` as a single operation. + - ✓ + - ✓ + + * - | double fmax ( double , double ) + | Determine the maximum numeric value of the arguments. + - ✓ + - ✓ + + * - | double fmin ( double x, double y ) + | Determine the minimum numeric value of the arguments. + - ✓ + - ✓ + + * - | double fmod ( double x, double y ) + | Calculate the floating-point remainder of `x / y`. + - ✓ + - ✓ + + * - | double frexp ( double x, int* nptr ) + | Extract mantissa and exponent of a floating-point value. + - ✓ + - ✗ + + * - | double hypot ( double x, double y ) + | Calculate the square root of the sum of squares of two arguments. + - ✓ + - ✓ + + * - | int ilogb ( double x ) + | Compute the unbiased integer exponent of the argument. + - ✓ + - ✓ + + * - | __RETURN_TYPE isfinite ( double a ) + | Determine whether argument is finite. + - ✓ + - ✓ + + * - | __RETURN_TYPE isinf ( double a ) + | Determine whether argument is infinite. + - ✓ + - ✓ + + * - | __RETURN_TYPE isnan ( double a ) + | Determine whether argument is a NaN. + - ✓ + - ✓ + + * - | double ldexp ( double x, int exp ) + | Calculate the value of x ⋅ 2 exp. + - ✓ + - ✓ + + * - | double log ( double x ) + | Calculate the base e logarithm of the input argument. + - ✓ + - ✓ + + * - | double log10 ( double x ) + | Calculate the base 10 logarithm of the input argument. + - ✓ + - ✓ + + * - | double log1p ( double x ) + | Calculate the value of logarithm of exp ( 1 + x ). + - ✓ + - ✓ + + * - | double log2 ( double x ) + | Calculate the base 2 logarithm of the input argument. + - ✓ + - ✓ + + * - | double logb ( double x ) + | Calculate the floating point representation of the exponent of the input argument. + - ✓ + - ✓ + + * - | double modf ( double x, `double* iptr` ) + | Break down the input argument into fractional and integral parts. + - ✓ + - ✗ + + * - | double nan ( const `char* tagp`) + | Returns ``Not a Number`` value. + - ✗ + - ✓ + + * - | double nearbyint ( double x ) + | Round the input argument to the nearest integer. + - ✓ + - ✓ + + * - | double pow ( double x, double y ) + | Calculate the value of first argument to the power of second argument. + - ✓ + - ✓ + + * - | double remainder ( double x, double y ) + | Compute double-precision floating-point remainder. + - ✓ + - ✓ + + * - | double remquo ( double x, double y, `int* quo` ) + | Compute double-precision floating-point remainder and part of quotient. + - ✓ + - ✗ + + * - | double round ( double x ) + | Round to nearest integer value in floating-point. + - ✓ + - ✓ + + * - | double scalbn ( double x, int n ) + | Scale floating-point input by integer power of two. + - ✓ + - ✓ + + * - | __RETURN_TYPE signbit ( double a ) + | Return the sign bit of the input. + - ✓ + - ✓ + + * - | double sin ( double x ) + | Calculate the sine of the input argument. + - ✓ + - ✓ + + * - | void sincos ( double x, `double* sptr`, `double* cptr` ) + | Calculate the sine and cosine of the first input argument. + - ✓ + - ✗ + + * - | double sinh ( double x ) + | Calculate the hyperbolic sine of the input argument. + - ✓ + - ✓ + + * - | double sqrt ( double x ) + | Calculate the square root of the input argument. + - ✓ + - ✓ + + * - | double tan ( double x ) + | Calculate the tangent of the input argument. + - ✓ + - ✓ + + * - | double tanh ( double x ) + | Calculate the hyperbolic tangent of the input argument. + - ✓ + - ✓ + + * - | double tgamma ( double x ) + | Calculate the gamma function of the input argument. + - ✓ + - ✓ + + * - | double trunc ( double x ) + | Truncate input argument to the integral part. + - ✓ + - ✓ + + * - | double erfcinv ( double y ) + | Calculate the inverse complementary function of the input argument. + - ✓ + - ✓ + + * - | double erfcx ( double x ) + | Calculate the scaled complementary error function of the input argument. + - ✓ + - ✓ + + * - | double erfinv ( double y ) + | Calculate the inverse error function of the input argument. + - ✓ + - ✓ + + * - | double frexp ( float x, `int *nptr` ) + | Extract mantissa and exponent of a floating-point value. + - ✓ + - ✓ + + * - | double j0 ( double x ) + | Calculate the value of the Bessel function of the first kind of order 0 for the input argument. + - ✓ + - ✓ + + * - | double j1 ( double x ) + | Calculate the value of the Bessel function of the first kind of order 1 for the input argument. + - ✓ + - ✓ + + * - | double jn ( int n, double x ) + | Calculate the value of the Bessel function of the first kind of order n for the input argument. + - ✓ + - ✓ + + * - | double lgamma ( double x ) + | Calculate the natural logarithm of the absolute value of the gamma function of the input argument. + - ✓ + - ✓ + + * - | long long int llrint ( double x ) + | Round input to nearest integer value. + - ✓ + - ✓ + + + * - | long long int llround ( double x ) + | Round to nearest integer value. + - ✓ + - ✓ + + * - | long int lrint ( double x ) + | Round input to nearest integer value. + - ✓ + - ✓ + + * - | long int lround ( double x ) + | Round to nearest integer value. + - ✓ + - ✓ + + * - | double modf ( double x, `double *iptr` ) + | Break down the input argument into fractional and integral parts. + - ✓ + - ✓ + + * - | double nextafter ( double x, double y ) + | Returns next representable single-precision floating-point value after argument. + - ✓ + - ✓ + + * - | double norm3d ( double a, double b, double c ) + | Calculate the square root of the sum of squares of three coordinates of the argument. + - ✓ + - ✓ + + * - | float norm4d ( double a, double b, double c, double d ) + | Calculate the square root of the sum of squares of four coordinates of the argument. + - ✓ + - ✓ + + * - | double normcdf ( double y ) + | Calculate the standard normal cumulative distribution function. + - ✓ + - ✓ + + * - | double normcdfinv ( double y ) + | Calculate the inverse of the standard normal cumulative distribution function. + - ✓ + - ✓ + + * - | double rcbrt ( double x ) + | Calculate the reciprocal cube root function. + - ✓ + - ✓ + + * - | double remquo ( double x, `double y`, `int *quo` ) + | Compute single-precision floating-point remainder and part of quotient. + - ✓ + - ✓ + + * - | double rhypot ( double x, double y ) + | Calculate one over the square root of the sum of squares of two arguments. + - ✓ + - ✓ + + * - | double rint ( double x ) + | Round input to nearest integer value in floating-point. + - ✓ + - ✓ + + * - | double rnorm3d ( double a, double b, double c ) + | Calculate one over the square root of the sum of squares of three coordinates of the argument. + - ✓ + - ✓ + + * - | double rnorm4d ( double a, double b, double c, double d ) + | Calculate one over the square root of the sum of squares of four coordinates of the argument. + - ✓ + - ✓ + + * - | double rnorm ( int dim, `const double *a` ) + | Calculate the reciprocal of square root of the sum of squares of any number of coordinates. + - ✓ + - ✓ + + * - | double scalbln ( double x, long int n ) + | Scale floating-point input by integer power of two. + - ✓ + - ✓ + + * - | void sincos ( double x, `double *sptr`, `double *cptr` ) + | Calculate the sine and cosine of the first input argument. + - ✓ + - ✓ + + * - | void sincospi ( double x, `double *sptr`, `double *cptr` ) + | Calculate the sine and cosine of the first input argument multiplied by PI. + - ✓ + - ✓ + + * - | double y0f ( double x ) + | Calculate the value of the Bessel function of the second kind of order 0 for the input argument. + - ✓ + - ✓ + + * - | double y1 ( double x ) + | Calculate the value of the Bessel function of the second kind of order 1 for the input argument. + - ✓ + - ✓ + + * - | double yn ( int n, double x ) + | Calculate the value of the Bessel function of the second kind of order n for the input argument. + - ✓ + - ✓ + +Integer intrinsics +-------------------------------------------------------------------------------------------- +Following is the list of supported integer intrinsics. Note that intrinsics are supported on device only. + +.. list-table:: Single precision mathematical functions + + * - **Function** + + * - | double acos ( double x ) + | Calculate the arc cosine of the input argument. + + * - | unsigned int __brev ( unsigned int x ) + | Reverse the bit order of a 32 bit unsigned integer. + + * - | unsigned long long int __brevll ( unsigned long long int x ) + | Reverse the bit order of a 64 bit unsigned integer. + + * - | int __clz ( int x ) + | Return the number of consecutive high-order zero bits in a 32 bit integer. + + * - | unsigned int __clz(unsigned int x) + | Return the number of consecutive high-order zero bits in 32 bit unsigned integer. + + * - | int __clzll ( long long int x ) + | Count the number of consecutive high-order zero bits in a 64 bit integer. + + * - | unsigned int __clzll(long long int x) + | Return the number of consecutive high-order zero bits in 64 bit signed integer. + + * - | unsigned int __ffs(unsigned int x) + | Find the position of least signigicant bit set to 1 in a 32 bit unsigned integer. + + * - | unsigned int __ffs(int x) + | Find the position of least signigicant bit set to 1 in a 32 bit signed integer. + + * - | unsigned int __ffsll(unsigned long long int x) + | Find the position of least signigicant bit set to 1 in a 64 bit unsigned integer. + + * - | unsigned int __ffsll(long long int x) + | Find the position of least signigicant bit set to 1 in a 64 bit signed integer. + + * - | unsigned int __popc ( unsigned int x ) + | Count the number of bits that are set to 1 in a 32 bit integer. + + * - | unsigned int __popcll ( unsigned long long int x ) + | Count the number of bits that are set to 1 in a 64 bit integer. + + * - | int __mul24 ( int x, int y ) + | Multiply two 24bit integers. + + * - | unsigned int __umul24 ( unsigned int x, unsigned int y ) + | Multiply two 24bit unsigned integers. + +The HIP-Clang implementation of ``__ffs()`` and ``__ffsll()`` contains code to add a constant +1 to produce the ffs result format. +For the cases where this overhead is not acceptable and programmer is willing to specialize for the platform, +HIP-Clang provides `__lastbit_u32_u32(unsigned int input)` and `__lastbit_u32_u64(unsigned long long int input)`. +The index returned by ``__lastbit_`` instructions starts at -1, while for ffs the index starts at 0. + +Floating-point Intrinsics +-------------------------------------------------------------------------------------------- +Following is the list of supported floating-point intrinsics. Note that intrinsics are supported on device only. + +.. list-table:: Single precision mathematical functions + + * - **Function** + + * - | float __cosf ( float x ) + | Calculate the fast approximate cosine of the input argument. + + * - | float __expf ( float x ) + | Calculate the fast approximate base e exponential of the input argument. + + * - | float __frsqrt_rn ( float x ) + | Compute `1 / √x` in round-to-nearest-even mode. + + * - | float __fsqrt_rn ( float x ) + | Compute `√x` in round-to-nearest-even mode. + + * - | float __log10f ( float x ) + | Calculate the fast approximate base 10 logarithm of the input argument. + + * - | float __log2f ( float x ) + | Calculate the fast approximate base 2 logarithm of the input argument. + + * - | float __logf ( float x ) + | Calculate the fast approximate base e logarithm of the input argument. + + * - | float __powf ( float x, float y ) + | Calculate the fast approximate of xy. + + * - | float __sinf ( float x ) + | Calculate the fast approximate sine of the input argument. + + * - | float __tanf ( float x ) + | Calculate the fast approximate tangent of the input argument. + + * - | double __dsqrt_rn ( double x ) + | Compute `√x` in round-to-nearest-even mode. Texture functions =============================================== @@ -378,15 +1362,15 @@ For example: * ``atomicAnd``: This function is atomic and coherent within the GPU device running the function -* ``atomicAnd_system``: This function extends the atomic operation from the GPU device to other CPUs -and GPU devices in the system +* ``atomicAnd_system``: This function extends the atomic operation from the GPU device to other CPUs and GPU devices in the system. HIP supports the following atomic operations. -.. list-table:: +.. list-table:: Atomic operations + * - **Function** - - **Supported in HIP** - - **Supported in CUDA** + - **Supported in HIP** + - **Supported in CUDA** * - int atomicAdd(int* address, int val) - ✓ @@ -836,10 +1820,11 @@ way of communicating between groups of threads at a granularity that is differen HIP supports the following kernel language cooperative groups types and functions: -.. list-table:: +.. list-table:: Cooperative groups functions + * - **Function** - - **Supported in HIP** - - **Supported in CUDA** + - **Supported in HIP** + - **Supported in CUDA** * - void thread_group.sync(); - ✓ @@ -949,10 +1934,11 @@ spread over lanes in an unspecified manner. HIP does not support kernel language warp matrix types or functions. -.. list-table:: +.. list-table:: Warp matrix functions + * - **Function** - - **Supported in HIP** - - **Supported in CUDA** + - **Supported in HIP** + - **Supported in CUDA** * - void load_matrix_sync(fragment<...> &a, const T* mptr, unsigned lda) - ✗ @@ -1005,7 +1991,10 @@ HIP provides the function ``abort()`` which can be used to terminate the applica This function produces a similar effect of using ``asm("trap")`` in the CUDA code. -.. note:: In HIP, the function terminates the entire application, while in CUDA, `asm("trap")`only terminates the dispatch and the application continues to run. +.. note:: + + In HIP, the function terminates the entire application, while in CUDA, ``asm("trap")`` only terminates the dispatch and the application continues to run. + Printf ============================================================ @@ -1190,6 +2179,7 @@ The file format for binary is `.co` which means Code Object. The following comma [OUTPUT FILE] = Name of the generated code object file .. note:: + When using binary code objects is that the number of arguments to the kernel is different on HIP-Clang and NVCC path. Refer to the sample in samples/0_Intro/module_api for differences in the arguments to be passed to the kernel. gfx-arch-specific-kernel diff --git a/projects/hip/docs/old/user_guide/faq.md b/projects/hip/docs/old/user_guide/faq.md index fbc56a5c6d..0c5dd879e2 100644 --- a/projects/hip/docs/old/user_guide/faq.md +++ b/projects/hip/docs/old/user_guide/faq.md @@ -162,7 +162,7 @@ Yes. HIP generates the object code which conforms to the GCC ABI, and also link with GPU code compiled with HIP. Larger projects often contain a mixture of accelerator code (initially written in CUDA with nvcc) and host code (compiled with gcc, icc, or clang). These projects can convert the accelerator code to HIP, compile that code with hipcc, and link with object code from their preferred compiler. -## Can HIP API support C style application? What is the differentce between C and C++ ? +## Can HIP API support C style application? What is the difference between C and C++ ? HIP is C++ runtime API that supports C style applications as well. Some C style applications (and interfaces to other languages (Fortran, Python)) would call certain HIP APIs but not use kernel programming. diff --git a/projects/hip/include/hip/hip_runtime_api.h b/projects/hip/include/hip/hip_runtime_api.h index c31e1feae1..4541a39a88 100644 --- a/projects/hip/include/hip/hip_runtime_api.h +++ b/projects/hip/include/hip/hip_runtime_api.h @@ -102,7 +102,7 @@ typedef struct hipDeviceProp_t { char luid[8]; ///< 8-byte unique identifier. Only valid on windows unsigned int luidDeviceNodeMask; ///< LUID node mask size_t totalGlobalMem; ///< Size of global memory region (in bytes). - size_t sharedMemPerBlock; ///< Size of shared memory region (in bytes). + size_t sharedMemPerBlock; ///< Size of shared memory per block (in bytes). int regsPerBlock; ///< Registers per block. int warpSize; ///< Warp size. size_t memPitch; ///< Maximum pitch in bytes allowed by memory copies @@ -111,7 +111,8 @@ typedef struct hipDeviceProp_t { int maxThreadsDim[3]; ///< Max number of threads in each dimension (XYZ) of a block. int maxGridSize[3]; ///< Max grid dimensions (XYZ). int clockRate; ///< Max clock frequency of the multiProcessors in khz. - size_t totalConstMem; ///< Size of shared memory region (in bytes). + size_t totalConstMem; ///< Size of shared constant memory region on the device + ///< (in bytes). int major; ///< Major compute capability. On HCC, this is an approximation and features may ///< differ from CUDA CC. See the arch feature flags for portable ways to query ///< feature caps. @@ -6692,9 +6693,9 @@ hipError_t hipStreamBeginCapture(hipStream_t stream, hipStreamCaptureMode mode); * @brief Begins graph capture on a stream to an existing graph. * * @param [in] stream - Stream to initiate capture. -* @param [in] stream - Graph to capture into. +* @param [in] graph - Graph to capture into. * @param [in] dependencies - Dependencies of the first node captured in the stream. Can be NULL if -numDependencies is 0. +* numDependencies is 0. * @param [in] dependencyData - Optional array of data associated with each dependency. * @param [in] numDependencies - Number of dependencies. * @param [in] mode - Controls the interaction of this capture sequence with other API calls that