## 1.6. Double Precision Intrinsics

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

### Functions

__device__ ​ double __dadd_rd ( double  x, double  y )
Add two floating point values in round-down mode.
__device__ ​ double __dadd_rn ( double  x, double  y )
Add two floating point values in round-to-nearest-even mode.
__device__ ​ double __dadd_ru ( double  x, double  y )
Add two floating point values in round-up mode.
__device__ ​ double __dadd_rz ( double  x, double  y )
Add two floating point values in round-towards-zero mode.
__device__ ​ double __ddiv_rd ( double  x, double  y )
Divide two floating point values in round-down mode.
__device__ ​ double __ddiv_rn ( double  x, double  y )
Divide two floating point values in round-to-nearest-even mode.
__device__ ​ double __ddiv_ru ( double  x, double  y )
Divide two floating point values in round-up mode.
__device__ ​ double __ddiv_rz ( double  x, double  y )
Divide two floating point values in round-towards-zero mode.
__device__ ​ double __dmul_rd ( double  x, double  y )
Multiply two floating point values in round-down mode.
__device__ ​ double __dmul_rn ( double  x, double  y )
Multiply two floating point values in round-to-nearest-even mode.
__device__ ​ double __dmul_ru ( double  x, double  y )
Multiply two floating point values in round-up mode.
__device__ ​ double __dmul_rz ( double  x, double  y )
Multiply two floating point values in round-towards-zero mode.
__device__ ​ double __drcp_rd ( double  x )
Compute $\frac{1}{x}$ in round-down mode.
__device__ ​ double __drcp_rn ( double  x )
Compute $\frac{1}{x}$ in round-to-nearest-even mode.
__device__ ​ double __drcp_ru ( double  x )
Compute $\frac{1}{x}$ in round-up mode.
__device__ ​ double __drcp_rz ( double  x )
Compute $\frac{1}{x}$ in round-towards-zero mode.
__device__ ​ double __dsqrt_rd ( double  x )
Compute $\sqrt{x}$ in round-down mode.
__device__ ​ double __dsqrt_rn ( double  x )
Compute $\sqrt{x}$ in round-to-nearest-even mode.
__device__ ​ double __dsqrt_ru ( double  x )
Compute $\sqrt{x}$ in round-up mode.
__device__ ​ double __dsqrt_rz ( double  x )
Compute $\sqrt{x}$ in round-towards-zero mode.
__device__ ​ double __dsub_rd ( double  x, double  y )
Subtract two floating point values in round-down mode.
__device__ ​ double __dsub_rn ( double  x, double  y )
Subtract two floating point values in round-to-nearest-even mode.
__device__ ​ double __dsub_ru ( double  x, double  y )
Subtract two floating point values in round-up mode.
__device__ ​ double __dsub_rz ( double  x, double  y )
Subtract two floating point values in round-towards-zero mode.
__device__ ​ double __fma_rd ( double  x, double  y, double  z )
Compute $x×y+z$ as a single operation in round-down mode.
__device__ ​ double __fma_rn ( double  x, double  y, double  z )
Compute $x×y+z$ as a single operation in round-to-nearest-even mode.
__device__ ​ double __fma_ru ( double  x, double  y, double  z )
Compute $x×y+z$ as a single operation in round-up mode.
__device__ ​ double __fma_rz ( double  x, double  y, double  z )
Compute $x×y+z$ as a single operation in round-towards-zero mode.

### Functions

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

Returns x + y.

###### Description

Adds two floating point values 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 7.

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

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

Returns x + y.

###### Description

Adds two floating point values 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 7.

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

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

Returns x + y.

###### Description

Adds two floating point values 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 7.

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

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

Returns x + y.

###### Description

Adds two floating point values 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 7.

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

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

Returns x / y.

###### Description

Divides 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 7.

• Requires compute capability >= 2.0.

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

Returns x / y.

###### Description

Divides 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 7.

• Requires compute capability >= 2.0.

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

Returns x / y.

###### Description

Divides 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 7.

• Requires compute capability >= 2.0.

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

Returns x / y.

###### Description

Divides 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 7.

• Requires compute capability >= 2.0.

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

Returns x * y.

###### Description

Multiplies two floating point values 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 7.

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

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

Returns x * y.

###### Description

Multiplies two floating point values 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 7.

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

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

Returns x * y.

###### Description

Multiplies two floating point values 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 7.

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

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

Returns x * y.

###### Description

Multiplies two floating point values 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 7.

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

__device__ ​ double __drcp_rd ( double  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 7.

• Requires compute capability >= 2.0.

__device__ ​ double __drcp_rn ( double  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 7.

• Requires compute capability >= 2.0.

__device__ ​ double __drcp_ru ( double  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 7.

• Requires compute capability >= 2.0.

__device__ ​ double __drcp_rz ( double  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 7.

• Requires compute capability >= 2.0.

__device__ ​ double __dsqrt_rd ( double  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 7.

• Requires compute capability >= 2.0.

__device__ ​ double __dsqrt_rn ( double  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 7.

• Requires compute capability >= 2.0.

__device__ ​ double __dsqrt_ru ( double  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 7.

• Requires compute capability >= 2.0.

__device__ ​ double __dsqrt_rz ( double  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 7.

• Requires compute capability >= 2.0.

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

Returns x - y.

###### Description

Subtracts two floating point values 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 7.

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

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

Returns x - y.

###### Description

Subtracts two floating point values 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 7.

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

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

Returns x - y.

###### Description

Subtracts two floating point values 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 7.

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

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

Returns x - y.

###### Description

Subtracts two floating point values 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 7.

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

__device__ ​ double __fma_rd ( double  x, double  y, double  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 7.

__device__ ​ double __fma_rn ( double  x, double  y, double  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 7.

__device__ ​ double __fma_ru ( double  x, double  y, double  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 7.

__device__ ​ double __fma_rz ( double  x, double  y, double  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 7.