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) - v8.0 (older) - Last updated September 4, 2016 - Send Feedback -

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