This document includes math equations (highlighted in red) which are best viewed with Firefox version 4.0 or higher, or another MathML-aware browser. There is also a PDF version of this document.
CUDA Math API (PDF) - v10.1.168 (older) - Last updated April 24, 2019 - Send Feedback

## 1.6. Double Precision Intrinsics

This section describes double precision intrinsic functions that are only supported in device code. To use these functions you do not need to include any additional header files in your program.

### 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 1 x in round-down mode.
__device__ ​ double __drcp_rn ( double x )
  Compute 1 x in round-to-nearest-even mode.
__device__ ​ double __drcp_ru ( double x )
  Compute 1 x in round-up mode.
__device__ ​ double __drcp_rz ( double x )
  Compute 1 x in round-towards-zero mode.
__device__ ​ double __dsqrt_rd ( double x )
  Compute x in round-down mode.
__device__ ​ double __dsqrt_rn ( double x )
  Compute x in round-to-nearest-even mode.
__device__ ​ double __dsqrt_ru ( double x )
  Compute x in round-up mode.
__device__ ​ double __dsqrt_rz ( double x )
  Compute 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

  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

  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

  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

  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

  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

  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

  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

  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

  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

  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

  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

  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 1 x in round-down mode.

  ###### Returns

  Returns 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 1 x in round-to-nearest-even mode.

  ###### Returns

  Returns 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 1 x in round-up mode.

  ###### Returns

  Returns 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 1 x in round-towards-zero mode.

  ###### Returns

  Returns 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 x in round-down mode.

  ###### Returns

  Returns 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 x in round-to-nearest-even mode.

  ###### Returns

  Returns 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 x in round-up mode.

  ###### Returns

  Returns 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 x in round-towards-zero mode.

  ###### Returns

  Returns 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

  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

  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

  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

  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( ± ∞ , ± 0 , `z`) returns NaN.
  - fmaf( ± 0 , ± ∞ , `z`) returns NaN.
  - fmaf(`x`, `y`, − ∞ ) returns NaN if x × y is an exact + ∞
  - fmaf(`x`, `y`, + ∞ ) returns NaN if x × y is an exact − ∞

  ###### 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( ± ∞ , ± 0 , `z`) returns NaN.
  - fmaf( ± 0 , ± ∞ , `z`) returns NaN.
  - fmaf(`x`, `y`, − ∞ ) returns NaN if x × y is an exact + ∞
  - fmaf(`x`, `y`, + ∞ ) returns NaN if x × y is an exact − ∞

  ###### 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( ± ∞ , ± 0 , `z`) returns NaN.
  - fmaf( ± 0 , ± ∞ , `z`) returns NaN.
  - fmaf(`x`, `y`, − ∞ ) returns NaN if x × y is an exact + ∞
  - fmaf(`x`, `y`, + ∞ ) returns NaN if x × y is an exact − ∞

  ###### 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( ± ∞ , ± 0 , `z`) returns NaN.
  - fmaf( ± 0 , ± ∞ , `z`) returns NaN.
  - fmaf(`x`, `y`, − ∞ ) returns NaN if x × y is an exact + ∞
  - fmaf(`x`, `y`, + ∞ ) returns NaN if x × y is an exact − ∞

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