## 1.5. Single Precision Intrinsics

This section describes single precision intrinsic functions that are only supported in device code.

### Functions

__device__ ​ float __cosf ( float  x )
Calculate the fast approximate cosine of the input argument.
__device__ ​ float __exp10f ( float  x )
Calculate the fast approximate base 10 exponential of the input argument.
__device__ ​ float __expf ( float  x )
Calculate the fast approximate base $e$ exponential of the input argument.
__device__ ​ float __fadd_rd ( float  x, float  y )
Add two floating point values in round-down mode.
__device__ ​ float __fadd_rn ( float  x, float  y )
Add two floating point values in round-to-nearest-even mode.
__device__ ​ float __fadd_ru ( float  x, float  y )
Add two floating point values in round-up mode.
__device__ ​ float __fadd_rz ( float  x, float  y )
Add two floating point values in round-towards-zero mode.
__device__ ​ float __fdiv_rd ( float  x, float  y )
Divide two floating point values in round-down mode.
__device__ ​ float __fdiv_rn ( float  x, float  y )
Divide two floating point values in round-to-nearest-even mode.
__device__ ​ float __fdiv_ru ( float  x, float  y )
Divide two floating point values in round-up mode.
__device__ ​ float __fdiv_rz ( float  x, float  y )
Divide two floating point values in round-towards-zero mode.
__device__ ​ float __fdividef ( float  x, float  y )
Calculate the fast approximate division of the input arguments.
__device__ ​ float __fmaf_rd ( float  x, float  y, float  z )
Compute $x×y+z$ as a single operation, in round-down mode.
__device__ ​ float __fmaf_rn ( float  x, float  y, float  z )
Compute $x×y+z$ as a single operation, in round-to-nearest-even mode.
__device__ ​ float __fmaf_ru ( float  x, float  y, float  z )
Compute $x×y+z$ as a single operation, in round-up mode.
__device__ ​ float __fmaf_rz ( float  x, float  y, float  z )
Compute $x×y+z$ as a single operation, in round-towards-zero mode.
__device__ ​ float __fmul_rd ( float  x, float  y )
Multiply two floating point values in round-down mode.
__device__ ​ float __fmul_rn ( float  x, float  y )
Multiply two floating point values in round-to-nearest-even mode.
__device__ ​ float __fmul_ru ( float  x, float  y )
Multiply two floating point values in round-up mode.
__device__ ​ float __fmul_rz ( float  x, float  y )
Multiply two floating point values in round-towards-zero mode.
__device__ ​ float __frcp_rd ( float  x )
Compute $\frac{1}{x}$ in round-down mode.
__device__ ​ float __frcp_rn ( float  x )
Compute $\frac{1}{x}$ in round-to-nearest-even mode.
__device__ ​ float __frcp_ru ( float  x )
Compute $\frac{1}{x}$ in round-up mode.
__device__ ​ float __frcp_rz ( float  x )
Compute $\frac{1}{x}$ in round-towards-zero mode.
__device__ ​ float __frsqrt_rn ( float  x )
Compute $1/\sqrt{x}$ in round-to-nearest-even mode.
__device__ ​ float __fsqrt_rd ( float  x )
Compute $\sqrt{x}$ in round-down mode.
__device__ ​ float __fsqrt_rn ( float  x )
Compute $\sqrt{x}$ in round-to-nearest-even mode.
__device__ ​ float __fsqrt_ru ( float  x )
Compute $\sqrt{x}$ in round-up mode.
__device__ ​ float __fsqrt_rz ( float  x )
Compute $\sqrt{x}$ in round-towards-zero mode.
__device__ ​ float __fsub_rd ( float  x, float  y )
Subtract two floating point values in round-down mode.
__device__ ​ float __fsub_rn ( float  x, float  y )
Subtract two floating point values in round-to-nearest-even mode.
__device__ ​ float __fsub_ru ( float  x, float  y )
Subtract two floating point values in round-up mode.
__device__ ​ float __fsub_rz ( float  x, float  y )
Subtract two floating point values in round-towards-zero mode.
__device__ ​ float __log10f ( float  x )
Calculate the fast approximate base 10 logarithm of the input argument.
__device__ ​ float __log2f ( float  x )
Calculate the fast approximate base 2 logarithm of the input argument.
__device__ ​ float __logf ( float  x )
Calculate the fast approximate base $e$ logarithm of the input argument.
__device__ ​ float __powf ( float  x, float  y )
Calculate the fast approximate of ${x}^{y}$ .
__device__ ​ float __saturatef ( float  x )
Clamp the input argument to [+0.0, 1.0].
__device__ ​ void __sincosf ( float  x, float* sptr, float* cptr )
Calculate the fast approximate of sine and cosine of the first input argument.
__device__ ​ float __sinf ( float  x )
Calculate the fast approximate sine of the input argument.
__device__ ​ float __tanf ( float  x )
Calculate the fast approximate tangent of the input argument.

### Functions

__device__ ​ float __cosf ( float  x )
Calculate the fast approximate cosine of the input argument.
###### Returns

Returns the approximate cosine of x.

###### Description

Calculate the fast approximate cosine of the input argument x, measured in radians.

Note:
• For accuracy information for this function see the CUDA C Programming Guide, Appendix D.2, Table 9.

• Input and output in the denormal range is flushed to sign preserving 0.0.

__device__ ​ float __exp10f ( float  x )
Calculate the fast approximate base 10 exponential of the input argument.
###### Returns

Returns an approximation to ${10}^{x}$ .

###### Description

Calculate the fast approximate base 10 exponential of the input argument x, ${10}^{x}$ .

Note:
• For accuracy information for this function see the CUDA C Programming Guide, Appendix D.2, Table 9.

• Most input and output values around denormal range are flushed to sign preserving 0.0.

__device__ ​ float __expf ( float  x )
Calculate the fast approximate base $e$ exponential of the input argument.
###### Returns

Returns an approximation to ${e}^{x}$ .

###### Description

Calculate the fast approximate base $e$ exponential of the input argument x, ${e}^{x}$ .

Note:
• For accuracy information for this function see the CUDA C Programming Guide, Appendix D.2, Table 9.

• Most input and output values around denormal range are flushed to sign preserving 0.0.

__device__ ​ float __fadd_rd ( float  x, float  y )
Add two floating point values in round-down mode.

Returns x + y.

###### Description

Compute the sum of x and y in round-down (to negative infinity) mode.

Note:
• For accuracy information for this function see the CUDA C Programming Guide, Appendix D.1, Table 6.

• This operation will never be merged into a single multiply-add instruction.

__device__ ​ float __fadd_rn ( float  x, float  y )
Add two floating point values in round-to-nearest-even mode.

Returns x + y.

###### Description

Compute the sum of x and y in round-to-nearest-even rounding mode.

Note:
• For accuracy information for this function see the CUDA C Programming Guide, Appendix D.1, Table 6.

• This operation will never be merged into a single multiply-add instruction.

__device__ ​ float __fadd_ru ( float  x, float  y )
Add two floating point values in round-up mode.

Returns x + y.

###### Description

Compute the sum of x and y in round-up (to positive infinity) mode.

Note:
• For accuracy information for this function see the CUDA C Programming Guide, Appendix D.1, Table 6.

• This operation will never be merged into a single multiply-add instruction.

__device__ ​ float __fadd_rz ( float  x, float  y )
Add two floating point values in round-towards-zero mode.

Returns x + y.

###### Description

Compute the sum of x and y in round-towards-zero mode.

Note:
• For accuracy information for this function see the CUDA C Programming Guide, Appendix D.1, Table 6.

• This operation will never be merged into a single multiply-add instruction.

__device__ ​ float __fdiv_rd ( float  x, float  y )
Divide two floating point values in round-down mode.

Returns x / y.

###### Description

Divide two floating point values x by y in round-down (to negative infinity) mode.

Note:

For accuracy information for this function see the CUDA C Programming Guide, Appendix D.1, Table 6.

__device__ ​ float __fdiv_rn ( float  x, float  y )
Divide two floating point values in round-to-nearest-even mode.

Returns x / y.

###### Description

Divide two floating point values x by y in round-to-nearest-even mode.

Note:

For accuracy information for this function see the CUDA C Programming Guide, Appendix D.1, Table 6.

__device__ ​ float __fdiv_ru ( float  x, float  y )
Divide two floating point values in round-up mode.

Returns x / y.

###### Description

Divide two floating point values x by y in round-up (to positive infinity) mode.

Note:

For accuracy information for this function see the CUDA C Programming Guide, Appendix D.1, Table 6.

__device__ ​ float __fdiv_rz ( float  x, float  y )
Divide two floating point values in round-towards-zero mode.

Returns x / y.

###### Description

Divide two floating point values x by y in round-towards-zero mode.

Note:

For accuracy information for this function see the CUDA C Programming Guide, Appendix D.1, Table 6.

__device__ ​ float __fdividef ( float  x, float  y )
Calculate the fast approximate division of the input arguments.
###### Returns

Returns x / y.

• __fdividef( $\mathrm{\infty }$ , y) returns NaN for ${2}^{126} .
• __fdividef(x, y) returns 0 for ${2}^{126} and $x\ne \mathrm{\infty }$ .

###### Description

Calculate the fast approximate division of x by y.

Note:

For accuracy information for this function see the CUDA C Programming Guide, Appendix D.2, Table 9.

__device__ ​ float __fmaf_rd ( float  x, float  y, float  z )
Compute $x×y+z$ as a single operation, in round-down mode.
###### Returns

Returns the rounded value of $x×y+z$ as a single operation.

• fmaf( $±\mathrm{\infty }$ , $±0$ , z) returns NaN.
• fmaf( $±0$ , $±\mathrm{\infty }$ , z) returns NaN.
• fmaf(x, y, $-\mathrm{\infty }$ ) returns NaN if $x×y$ is an exact $+\mathrm{\infty }$ .
• fmaf(x, y, $+\mathrm{\infty }$ ) returns NaN if $x×y$ is an exact $-\mathrm{\infty }$ .

###### Description

Computes the value of $x×y+z$ as a single ternary operation, rounding the result once in round-down (to negative infinity) mode.

Note:

For accuracy information for this function see the CUDA C Programming Guide, Appendix D.1, Table 6.

__device__ ​ float __fmaf_rn ( float  x, float  y, float  z )
Compute $x×y+z$ as a single operation, in round-to-nearest-even mode.
###### Returns

Returns the rounded value of $x×y+z$ as a single operation.

• fmaf( $±\mathrm{\infty }$ , $±0$ , z) returns NaN.
• fmaf( $±0$ , $±\mathrm{\infty }$ , z) returns NaN.
• fmaf(x, y, $-\mathrm{\infty }$ ) returns NaN if $x×y$ is an exact $+\mathrm{\infty }$ .
• fmaf(x, y, $+\mathrm{\infty }$ ) returns NaN if $x×y$ is an exact $-\mathrm{\infty }$ .

###### Description

Computes the value of $x×y+z$ as a single ternary operation, rounding the result once in round-to-nearest-even mode.

Note:

For accuracy information for this function see the CUDA C Programming Guide, Appendix D.1, Table 6.

__device__ ​ float __fmaf_ru ( float  x, float  y, float  z )
Compute $x×y+z$ as a single operation, in round-up mode.
###### Returns

Returns the rounded value of $x×y+z$ as a single operation.

• fmaf( $±\mathrm{\infty }$ , $±0$ , z) returns NaN.
• fmaf( $±0$ , $±\mathrm{\infty }$ , z) returns NaN.
• fmaf(x, y, $-\mathrm{\infty }$ ) returns NaN if $x×y$ is an exact $+\mathrm{\infty }$ .
• fmaf(x, y, $+\mathrm{\infty }$ ) returns NaN if $x×y$ is an exact $-\mathrm{\infty }$ .

###### Description

Computes the value of $x×y+z$ as a single ternary operation, rounding the result once in round-up (to positive infinity) mode.

Note:

For accuracy information for this function see the CUDA C Programming Guide, Appendix D.1, Table 6.

__device__ ​ float __fmaf_rz ( float  x, float  y, float  z )
Compute $x×y+z$ as a single operation, in round-towards-zero mode.
###### Returns

Returns the rounded value of $x×y+z$ as a single operation.

• fmaf( $±\mathrm{\infty }$ , $±0$ , z) returns NaN.
• fmaf( $±0$ , $±\mathrm{\infty }$ , z) returns NaN.
• fmaf(x, y, $-\mathrm{\infty }$ ) returns NaN if $x×y$ is an exact $+\mathrm{\infty }$ .
• fmaf(x, y, $+\mathrm{\infty }$ ) returns NaN if $x×y$ is an exact $-\mathrm{\infty }$ .

###### Description

Computes the value of $x×y+z$ as a single ternary operation, rounding the result once in round-towards-zero mode.

Note:

For accuracy information for this function see the CUDA C Programming Guide, Appendix D.1, Table 6.

__device__ ​ float __fmul_rd ( float  x, float  y )
Multiply two floating point values in round-down mode.

Returns x * y.

###### Description

Compute the product of x and y in round-down (to negative infinity) mode.

Note:
• For accuracy information for this function see the CUDA C Programming Guide, Appendix D.1, Table 6.

• This operation will never be merged into a single multiply-add instruction.

__device__ ​ float __fmul_rn ( float  x, float  y )
Multiply two floating point values in round-to-nearest-even mode.

Returns x * y.

###### Description

Compute the product of x and y in round-to-nearest-even mode.

Note:
• For accuracy information for this function see the CUDA C Programming Guide, Appendix D.1, Table 6.

• This operation will never be merged into a single multiply-add instruction.

__device__ ​ float __fmul_ru ( float  x, float  y )
Multiply two floating point values in round-up mode.

Returns x * y.

###### Description

Compute the product of x and y in round-up (to positive infinity) mode.

Note:
• For accuracy information for this function see the CUDA C Programming Guide, Appendix D.1, Table 6.

• This operation will never be merged into a single multiply-add instruction.

__device__ ​ float __fmul_rz ( float  x, float  y )
Multiply two floating point values in round-towards-zero mode.

Returns x * y.

###### Description

Compute the product of x and y in round-towards-zero mode.

Note:
• For accuracy information for this function see the CUDA C Programming Guide, Appendix D.1, Table 6.

• This operation will never be merged into a single multiply-add instruction.

__device__ ​ float __frcp_rd ( float  x )
Compute $\frac{1}{x}$ in round-down mode.
###### Returns

Returns $\frac{1}{x}$ .

###### Description

Compute the reciprocal of x in round-down (to negative infinity) mode.

Note:

For accuracy information for this function see the CUDA C Programming Guide, Appendix D.1, Table 6.

__device__ ​ float __frcp_rn ( float  x )
Compute $\frac{1}{x}$ in round-to-nearest-even mode.
###### Returns

Returns $\frac{1}{x}$ .

###### Description

Compute the reciprocal of x in round-to-nearest-even mode.

Note:

For accuracy information for this function see the CUDA C Programming Guide, Appendix D.1, Table 6.

__device__ ​ float __frcp_ru ( float  x )
Compute $\frac{1}{x}$ in round-up mode.
###### Returns

Returns $\frac{1}{x}$ .

###### Description

Compute the reciprocal of x in round-up (to positive infinity) mode.

Note:

For accuracy information for this function see the CUDA C Programming Guide, Appendix D.1, Table 6.

__device__ ​ float __frcp_rz ( float  x )
Compute $\frac{1}{x}$ in round-towards-zero mode.
###### Returns

Returns $\frac{1}{x}$ .

###### Description

Compute the reciprocal of x in round-towards-zero mode.

Note:

For accuracy information for this function see the CUDA C Programming Guide, Appendix D.1, Table 6.

__device__ ​ float __frsqrt_rn ( float  x )
Compute $1/\sqrt{x}$ in round-to-nearest-even mode.
###### Returns

Returns $1/\sqrt{x}$ .

###### Description

Compute the reciprocal square root of x in round-to-nearest-even mode.

Note:

For accuracy information for this function see the CUDA C Programming Guide, Appendix D.1, Table 6.

__device__ ​ float __fsqrt_rd ( float  x )
Compute $\sqrt{x}$ in round-down mode.
###### Returns

Returns $\sqrt{x}$ .

###### Description

Compute the square root of x in round-down (to negative infinity) mode.

Note:

For accuracy information for this function see the CUDA C Programming Guide, Appendix D.1, Table 6.

__device__ ​ float __fsqrt_rn ( float  x )
Compute $\sqrt{x}$ in round-to-nearest-even mode.
###### Returns

Returns $\sqrt{x}$ .

###### Description

Compute the square root of x in round-to-nearest-even mode.

Note:

For accuracy information for this function see the CUDA C Programming Guide, Appendix D.1, Table 6.

__device__ ​ float __fsqrt_ru ( float  x )
Compute $\sqrt{x}$ in round-up mode.
###### Returns

Returns $\sqrt{x}$ .

###### Description

Compute the square root of x in round-up (to positive infinity) mode.

Note:

For accuracy information for this function see the CUDA C Programming Guide, Appendix D.1, Table 6.

__device__ ​ float __fsqrt_rz ( float  x )
Compute $\sqrt{x}$ in round-towards-zero mode.
###### Returns

Returns $\sqrt{x}$ .

###### Description

Compute the square root of x in round-towards-zero mode.

Note:

For accuracy information for this function see the CUDA C Programming Guide, Appendix D.1, Table 6.

__device__ ​ float __fsub_rd ( float  x, float  y )
Subtract two floating point values in round-down mode.

Returns x - y.

###### Description

Compute the difference of x and y in round-down (to negative infinity) mode.

Note:
• For accuracy information for this function see the CUDA C Programming Guide, Appendix D.1, Table 6.

• This operation will never be merged into a single multiply-add instruction.

__device__ ​ float __fsub_rn ( float  x, float  y )
Subtract two floating point values in round-to-nearest-even mode.

Returns x - y.

###### Description

Compute the difference of x and y in round-to-nearest-even rounding mode.

Note:
• For accuracy information for this function see the CUDA C Programming Guide, Appendix D.1, Table 6.

• This operation will never be merged into a single multiply-add instruction.

__device__ ​ float __fsub_ru ( float  x, float  y )
Subtract two floating point values in round-up mode.

Returns x - y.

###### Description

Compute the difference of x and y in round-up (to positive infinity) mode.

Note:
• For accuracy information for this function see the CUDA C Programming Guide, Appendix D.1, Table 6.

• This operation will never be merged into a single multiply-add instruction.

__device__ ​ float __fsub_rz ( float  x, float  y )
Subtract two floating point values in round-towards-zero mode.

Returns x - y.

###### Description

Compute the difference of x and y in round-towards-zero mode.

Note:
• For accuracy information for this function see the CUDA C Programming Guide, Appendix D.1, Table 6.

• This operation will never be merged into a single multiply-add instruction.

__device__ ​ float __log10f ( float  x )
Calculate the fast approximate base 10 logarithm of the input argument.
###### Returns

Returns an approximation to ${\mathrm{log}}_{10}\left(x\right)$ .

###### Description

Calculate the fast approximate base 10 logarithm of the input argument x.

Note:
• For accuracy information for this function see the CUDA C Programming Guide, Appendix D.2, Table 9.

• Most input and output values around denormal range are flushed to sign preserving 0.0.

__device__ ​ float __log2f ( float  x )
Calculate the fast approximate base 2 logarithm of the input argument.
###### Returns

Returns an approximation to ${\mathrm{log}}_{2}\left(x\right)$ .

###### Description

Calculate the fast approximate base 2 logarithm of the input argument x.

Note:
• For accuracy information for this function see the CUDA C Programming Guide, Appendix D.2, Table 9.

• Input and output in the denormal range is flushed to sign preserving 0.0.

__device__ ​ float __logf ( float  x )
Calculate the fast approximate base $e$ logarithm of the input argument.
###### Returns

Returns an approximation to ${\mathrm{log}}_{e}\left(x\right)$ .

###### Description

Calculate the fast approximate base $e$ logarithm of the input argument x.

Note:
• For accuracy information for this function see the CUDA C Programming Guide, Appendix D.2, Table 9.

• Most input and output values around denormal range are flushed to sign preserving 0.0.

__device__ ​ float __powf ( float  x, float  y )
Calculate the fast approximate of ${x}^{y}$ .
###### Returns

Returns an approximation to ${x}^{y}$ .

###### Description

Calculate the fast approximate of x, the first input argument, raised to the power of y, the second input argument, ${x}^{y}$ .

Note:
• For accuracy information for this function see the CUDA C Programming Guide, Appendix D.2, Table 9.

• Most input and output values around denormal range are flushed to sign preserving 0.0.

__device__ ​ float __saturatef ( float  x )
Clamp the input argument to [+0.0, 1.0].
###### Returns

• __saturatef(x) returns 0 if x < 0.
• __saturatef(x) returns 1 if x > 1.
• __saturatef(x) returns x if $0\le x\le 1$ .
• __saturatef(NaN) returns 0.

###### Description

Clamp the input argument x to be within the interval [+0.0, 1.0].

__device__ ​ void __sincosf ( float  x, float* sptr, float* cptr )
Calculate the fast approximate of sine and cosine of the first input argument.

• none

###### Description

Calculate the fast approximate of sine and cosine of the first input argument x (measured in radians). The results for sine and cosine are written into the second argument, sptr, and, respectively, third argument, cptr.

Note:
• For accuracy information for this function see the CUDA C Programming Guide, Appendix D.2, Table 9.

• Denorm input/output is flushed to sign preserving 0.0.

__device__ ​ float __sinf ( float  x )
Calculate the fast approximate sine of the input argument.
###### Returns

Returns the approximate sine of x.

###### Description

Calculate the fast approximate sine of the input argument x, measured in radians.

Note:
• For accuracy information for this function see the CUDA C Programming Guide, Appendix D.2, Table 9.

• Input and output in the denormal range is flushed to sign preserving 0.0.

__device__ ​ float __tanf ( float  x )
Calculate the fast approximate tangent of the input argument.
###### Returns

Returns the approximate tangent of x.

###### Description

Calculate the fast approximate tangent of the input argument x, measured in radians.

Note:
• For accuracy information for this function see the CUDA C Programming Guide, Appendix D.2, Table 9.

• The result is computed as the fast divide of __sinf() by __cosf(). Denormal input and output are flushed to sign-preserving 0.0 at each step of the computation.