9.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 file 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\(\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 \times y + z\) as a single operation in round-down mode.

__device__ double__fma_rn(double x, double y, double z)

Compute\(x \times y + z\) as a single operation in round-to-nearest-even mode.

__device__ double__fma_ru(double x, double y, double z)

Compute\(x \times y + z\) as a single operation in round-up mode.

__device__ double__fma_rz(double x, double y, double z)

Compute\(x \times y + z\) as a single operation in round-towards-zero mode.

9.1.Functions

__device__double__dadd_rd(doublex,doubley)

Add two floating-point values in round-down mode.

Adds two floating-point valuesx andy in round-down (to negative infinity) mode.

Note

For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions Appendix, Intrinsic Functions section.

Note

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

Returns

Returnsx +y.

  • __dadd_rd(x,y) is equivalent to __dadd_rd(y,x).

  • __dadd_rd(x,\( \pm\infty \)) returns\( \pm\infty \) for finitex.

  • __dadd_rd(\( \pm\infty \),\( \pm\infty \)) returns\( \pm\infty \).

  • __dadd_rd(\( \pm\infty \),\( \mp\infty \)) returns NaN.

  • __dadd_rd(\( \pm 0 \),\( \pm 0 \)) returns\( \pm 0 \).

  • __dadd_rd(x,-x) returns\( -0 \) for finitex, including\( \pm 0 \).

  • If either argument is NaN, NaN is returned.

__device__double__dadd_rn(doublex,doubley)

Add two floating-point values in round-to-nearest-even mode.

Adds two floating-point valuesx andy in round-to-nearest-even mode.

Note

For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions Appendix, Intrinsic Functions section.

Note

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

Returns

Returnsx +y.

  • __dadd_rn(x,y) is equivalent to __dadd_rn(y,x).

  • __dadd_rn(x,\( \pm\infty \)) returns\( \pm\infty \) for finitex.

  • __dadd_rn(\( \pm\infty \),\( \pm\infty \)) returns\( \pm\infty \).

  • __dadd_rn(\( \pm\infty \),\( \mp\infty \)) returns NaN.

  • __dadd_rn(\( \pm 0 \),\( \pm 0 \)) returns\( \pm 0 \).

  • __dadd_rn(x,-x) returns\( +0 \) for finitex, including\( \pm 0 \).

  • If either argument is NaN, NaN is returned.

__device__double__dadd_ru(doublex,doubley)

Add two floating-point values in round-up mode.

Adds two floating-point valuesx andy in round-up (to positive infinity) mode.

Note

For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions Appendix, Intrinsic Functions section.

Note

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

Returns

Returnsx +y.

  • __dadd_ru(x,y) is equivalent to __dadd_ru(y,x).

  • __dadd_ru(x,\( \pm\infty \)) returns\( \pm\infty \) for finitex.

  • __dadd_ru(\( \pm\infty \),\( \pm\infty \)) returns\( \pm\infty \).

  • __dadd_ru(\( \pm\infty \),\( \mp\infty \)) returns NaN.

  • __dadd_ru(\( \pm 0 \),\( \pm 0 \)) returns\( \pm 0 \).

  • __dadd_ru(x,-x) returns\( +0 \) for finitex, including\( \pm 0 \).

  • If either argument is NaN, NaN is returned.

__device__double__dadd_rz(doublex,doubley)

Add two floating-point values in round-towards-zero mode.

Adds two floating-point valuesx andy in round-towards-zero mode.

Note

For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions Appendix, Intrinsic Functions section.

Note

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

Returns

Returnsx +y.

  • __dadd_rz(x,y) is equivalent to __dadd_rz(y,x).

  • __dadd_rz(x,\( \pm\infty \)) returns\( \pm\infty \) for finitex.

  • __dadd_rz(\( \pm\infty \),\( \pm\infty \)) returns\( \pm\infty \).

  • __dadd_rz(\( \pm\infty \),\( \mp\infty \)) returns NaN.

  • __dadd_rz(\( \pm 0 \),\( \pm 0 \)) returns\( \pm 0 \).

  • __dadd_rz(x,-x) returns\( +0 \) for finitex, including\( \pm 0 \).

  • If either argument is NaN, NaN is returned.

__device__double__ddiv_rd(doublex,doubley)

Divide two floating-point values in round-down mode.

Divides two floating-point valuesx byy in round-down (to negative infinity) mode.

Note

For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions Appendix, Intrinsic Functions section.

Note

Requires compute capability >= 2.0.

Returns

Returnsx /y.

  • sign of the quotientx /y is XOR of the signs ofx andy when neither inputs nor result are NaN.

  • __ddiv_rd(\( \pm 0 \),\( \pm 0 \)) returns NaN.

  • __ddiv_rd(\( \pm\infty \),\( \pm\infty \)) returns NaN.

  • __ddiv_rd(x,\( \pm\infty \)) returns\( 0 \) of appropriate sign for finitex.

  • __ddiv_rd(\( \pm\infty \),y) returns\( \infty \) of appropriate sign for finitey.

  • __ddiv_rd(x,\( \pm 0 \)) returns\( \infty \) of appropriate sign forx\( \neq 0 \).

  • __ddiv_rd(\( \pm 0 \),y) returns\( 0 \) of appropriate sign fory\( \neq 0 \).

  • If either argument is NaN, NaN is returned.

__device__double__ddiv_rn(doublex,doubley)

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

Divides two floating-point valuesx byy in round-to-nearest-even mode.

Note

For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions Appendix, Intrinsic Functions section.

Note

Requires compute capability >= 2.0.

Returns

Returnsx /y.

  • sign of the quotientx /y is XOR of the signs ofx andy when neither inputs nor result are NaN.

  • __ddiv_rn(\( \pm 0 \),\( \pm 0 \)) returns NaN.

  • __ddiv_rn(\( \pm\infty \),\( \pm\infty \)) returns NaN.

  • __ddiv_rn(x,\( \pm\infty \)) returns\( 0 \) of appropriate sign for finitex.

  • __ddiv_rn(\( \pm\infty \),y) returns\( \infty \) of appropriate sign for finitey.

  • __ddiv_rn(x,\( \pm 0 \)) returns\( \infty \) of appropriate sign forx\( \neq 0 \).

  • __ddiv_rn(\( \pm 0 \),y) returns\( 0 \) of appropriate sign fory\( \neq 0 \).

  • If either argument is NaN, NaN is returned.

__device__double__ddiv_ru(doublex,doubley)

Divide two floating-point values in round-up mode.

Divides two floating-point valuesx byy in round-up (to positive infinity) mode.

Note

For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions Appendix, Intrinsic Functions section.

Note

Requires compute capability >= 2.0.

Returns

Returnsx /y.

  • sign of the quotientx /y is XOR of the signs ofx andy when neither inputs nor result are NaN.

  • __ddiv_ru(\( \pm 0 \),\( \pm 0 \)) returns NaN.

  • __ddiv_ru(\( \pm\infty \),\( \pm\infty \)) returns NaN.

  • __ddiv_ru(x,\( \pm\infty \)) returns\( 0 \) of appropriate sign for finitex.

  • __ddiv_ru(\( \pm\infty \),y) returns\( \infty \) of appropriate sign for finitey.

  • __ddiv_ru(x,\( \pm 0 \)) returns\( \infty \) of appropriate sign forx\( \neq 0 \).

  • __ddiv_ru(\( \pm 0 \),y) returns\( 0 \) of appropriate sign fory\( \neq 0 \).

  • If either argument is NaN, NaN is returned.

__device__double__ddiv_rz(doublex,doubley)

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

Divides two floating-point valuesx byy in round-towards-zero mode.

Note

For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions Appendix, Intrinsic Functions section.

Note

Requires compute capability >= 2.0.

Returns

Returnsx /y.

  • sign of the quotientx /y is XOR of the signs ofx andy when neither inputs nor result are NaN.

  • __ddiv_rz(\( \pm 0 \),\( \pm 0 \)) returns NaN.

  • __ddiv_rz(\( \pm\infty \),\( \pm\infty \)) returns NaN.

  • __ddiv_rz(x,\( \pm\infty \)) returns\( 0 \) of appropriate sign for finitex.

  • __ddiv_rz(\( \pm\infty \),y) returns\( \infty \) of appropriate sign for finitey.

  • __ddiv_rz(x,\( \pm 0 \)) returns\( \infty \) of appropriate sign forx\( \neq 0 \).

  • __ddiv_rz(\( \pm 0 \),y) returns\( 0 \) of appropriate sign fory\( \neq 0 \).

  • If either argument is NaN, NaN is returned.

__device__double__dmul_rd(doublex,doubley)

Multiply two floating-point values in round-down mode.

Multiplies two floating-point valuesx andy in round-down (to negative infinity) mode.

Note

For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions Appendix, Intrinsic Functions section.

Note

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

Returns

Returnsx *y.

  • sign of the productx *y is XOR of the signs ofx andy when neither inputs nor result are NaN.

  • __dmul_rd(x,y) is equivalent to __dmul_rd(y,x).

  • __dmul_rd(x,\( \pm\infty \)) returns\( \infty \) of appropriate sign forx\( \neq 0 \).

  • __dmul_rd(\( \pm 0 \),\( \pm\infty \)) returns NaN.

  • __dmul_rd(\( \pm 0 \),y) returns\( 0 \) of appropriate sign for finitey.

  • If either argument is NaN, NaN is returned.

__device__double__dmul_rn(doublex,doubley)

Multiply two floating-point values in round-to-nearest-even mode.

Multiplies two floating-point valuesx andy in round-to-nearest-even mode.

Note

For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions Appendix, Intrinsic Functions section.

Note

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

Returns

Returnsx *y.

  • sign of the productx *y is XOR of the signs ofx andy when neither inputs nor result are NaN.

  • __dmul_rn(x,y) is equivalent to __dmul_rn(y,x).

  • __dmul_rn(x,\( \pm\infty \)) returns\( \infty \) of appropriate sign forx\( \neq 0 \).

  • __dmul_rn(\( \pm 0 \),\( \pm\infty \)) returns NaN.

  • __dmul_rn(\( \pm 0 \),y) returns\( 0 \) of appropriate sign for finitey.

  • If either argument is NaN, NaN is returned.

__device__double__dmul_ru(doublex,doubley)

Multiply two floating-point values in round-up mode.

Multiplies two floating-point valuesx andy in round-up (to positive infinity) mode.

Note

For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions Appendix, Intrinsic Functions section.

Note

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

Returns

Returnsx *y.

  • sign of the productx *y is XOR of the signs ofx andy when neither inputs nor result are NaN.

  • __dmul_ru(x,y) is equivalent to __dmul_ru(y,x).

  • __dmul_ru(x,\( \pm\infty \)) returns\( \infty \) of appropriate sign forx\( \neq 0 \).

  • __dmul_ru(\( \pm 0 \),\( \pm\infty \)) returns NaN.

  • __dmul_ru(\( \pm 0 \),y) returns\( 0 \) of appropriate sign for finitey.

  • If either argument is NaN, NaN is returned.

__device__double__dmul_rz(doublex,doubley)

Multiply two floating-point values in round-towards-zero mode.

Multiplies two floating-point valuesx andy in round-towards-zero mode.

Note

For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions Appendix, Intrinsic Functions section.

Note

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

Returns

Returnsx *y.

  • sign of the productx *y is XOR of the signs ofx andy when neither inputs nor result are NaN.

  • __dmul_rz(x,y) is equivalent to __dmul_rz(y,x).

  • __dmul_rz(x,\( \pm\infty \)) returns\( \infty \) of appropriate sign forx\( \neq 0 \).

  • __dmul_rz(\( \pm 0 \),\( \pm\infty \)) returns NaN.

  • __dmul_rz(\( \pm 0 \),y) returns\( 0 \) of appropriate sign for finitey.

  • If either argument is NaN, NaN is returned.

__device__double__drcp_rd(doublex)

Compute\( \frac{1}{x} \) in round-down mode.

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

Note

For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions Appendix, Intrinsic Functions section.

Note

Requires compute capability >= 2.0.

Returns

Returns\( \frac{1}{x} \).

__device__double__drcp_rn(doublex)

Compute\( \frac{1}{x} \) in round-to-nearest-even mode.

Compute the reciprocal ofx in round-to-nearest-even mode.

Note

For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions Appendix, Intrinsic Functions section.

Note

Requires compute capability >= 2.0.

Returns

Returns\( \frac{1}{x} \).

__device__double__drcp_ru(doublex)

Compute\( \frac{1}{x} \) in round-up mode.

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

Note

For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions Appendix, Intrinsic Functions section.

Note

Requires compute capability >= 2.0.

Returns

Returns\( \frac{1}{x} \).

__device__double__drcp_rz(doublex)

Compute\( \frac{1}{x} \) in round-towards-zero mode.

Compute the reciprocal ofx in round-towards-zero mode.

Note

For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions Appendix, Intrinsic Functions section.

Note

Requires compute capability >= 2.0.

Returns

Returns\( \frac{1}{x} \).

__device__double__dsqrt_rd(doublex)

Compute\( \sqrt{x} \) in round-down mode.

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

Note

For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions Appendix, Intrinsic Functions section.

Note

Requires compute capability >= 2.0.

Returns

Returns\( \sqrt{x} \).

__device__double__dsqrt_rn(doublex)

Compute\( \sqrt{x} \) in round-to-nearest-even mode.

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

Note

For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions Appendix, Intrinsic Functions section.

Note

Requires compute capability >= 2.0.

Returns

Returns\( \sqrt{x} \).

__device__double__dsqrt_ru(doublex)

Compute\( \sqrt{x} \) in round-up mode.

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

Note

For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions Appendix, Intrinsic Functions section.

Note

Requires compute capability >= 2.0.

Returns

Returns\( \sqrt{x} \).

__device__double__dsqrt_rz(doublex)

Compute\( \sqrt{x} \) in round-towards-zero mode.

Compute the square root ofx in round-towards-zero mode.

Note

For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions Appendix, Intrinsic Functions section.

Note

Requires compute capability >= 2.0.

Returns

Returns\( \sqrt{x} \).

__device__double__dsub_rd(doublex,doubley)

Subtract two floating-point values in round-down mode.

Subtracts two floating-point valuesx andy in round-down (to negative infinity) mode.

Note

For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions Appendix, Intrinsic Functions section.

Note

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

Returns

Returnsx -y.

  • __dsub_rd(\( \pm\infty \),y) returns\( \pm\infty \) for finitey.

  • __dsub_rd(x,\( \pm\infty \)) returns\( \mp\infty \) for finitex.

  • __dsub_rd(\( \pm\infty \),\( \pm\infty \)) returns NaN.

  • __dsub_rd(\( \pm\infty \),\( \mp\infty \)) returns\( \pm\infty \).

  • __dsub_rd(\( \pm 0 \),\( \mp 0 \)) returns\( \pm 0 \).

  • __dsub_rd(x,x) returns\( -0 \) for finitex, including\( \pm 0 \).

  • If either argument is NaN, NaN is returned.

__device__double__dsub_rn(doublex,doubley)

Subtract two floating-point values in round-to-nearest-even mode.

Subtracts two floating-point valuesx andy in round-to-nearest-even mode.

Note

For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions Appendix, Intrinsic Functions section.

Note

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

Returns

Returnsx -y.

  • __dsub_rn(\( \pm\infty \),y) returns\( \pm\infty \) for finitey.

  • __dsub_rn(x,\( \pm\infty \)) returns\( \mp\infty \) for finitex.

  • __dsub_rn(\( \pm\infty \),\( \pm\infty \)) returns NaN.

  • __dsub_rn(\( \pm\infty \),\( \mp\infty \)) returns\( \pm\infty \).

  • __dsub_rn(\( \pm 0 \),\( \mp 0 \)) returns\( \pm 0 \).

  • __dsub_rn(x,x) returns\( +0 \) for finitex, including\( \pm 0 \).

  • If either argument is NaN, NaN is returned.

__device__double__dsub_ru(doublex,doubley)

Subtract two floating-point values in round-up mode.

Subtracts two floating-point valuesx andy in round-up (to positive infinity) mode.

Note

For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions Appendix, Intrinsic Functions section.

Note

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

Returns

Returnsx -y.

  • __dsub_ru(\( \pm\infty \),y) returns\( \pm\infty \) for finitey.

  • __dsub_ru(x,\( \pm\infty \)) returns\( \mp\infty \) for finitex.

  • __dsub_ru(\( \pm\infty \),\( \pm\infty \)) returns NaN.

  • __dsub_ru(\( \pm\infty \),\( \mp\infty \)) returns\( \pm\infty \).

  • __dsub_ru(\( \pm 0 \),\( \mp 0 \)) returns\( \pm 0 \).

  • __dsub_ru(x,x) returns\( +0 \) for finitex, including\( \pm 0 \).

  • If either argument is NaN, NaN is returned.

__device__double__dsub_rz(doublex,doubley)

Subtract two floating-point values in round-towards-zero mode.

Subtracts two floating-point valuesx andy in round-towards-zero mode.

Note

For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions Appendix, Intrinsic Functions section.

Note

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

Returns

Returnsx -y.

  • __dsub_rz(\( \pm\infty \),y) returns\( \pm\infty \) for finitey.

  • __dsub_rz(x,\( \pm\infty \)) returns\( \mp\infty \) for finitex.

  • __dsub_rz(\( \pm\infty \),\( \pm\infty \)) returns NaN.

  • __dsub_rz(\( \pm\infty \),\( \mp\infty \)) returns\( \pm\infty \).

  • __dsub_rz(\( \pm 0 \),\( \mp 0 \)) returns\( \pm 0 \).

  • __dsub_rz(x,x) returns\( +0 \) for finitex, including\( \pm 0 \).

  • If either argument is NaN, NaN is returned.

__device__double__fma_rd(doublex,doubley,doublez)

Compute\( x \times y + z \) as a single operation in round-down mode.

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

Note

For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions Appendix, Intrinsic Functions section.

Returns

Returns the rounded value of\( x \times y + z \) as a single operation.

  • __fma_rd(\( \pm \infty \) ,\( \pm 0 \) ,z) returns NaN.

  • __fma_rd(\( \pm 0 \) ,\( \pm \infty \) ,z) returns NaN.

  • __fma_rd(x,y,\( -\infty \) ) returns NaN if\( x \times y \) is an exact\( +\infty \).

  • __fma_rd(x,y,\( +\infty \) ) returns NaN if\( x \times y \) is an exact\( -\infty \).

  • __fma_rd(x,y,\( \pm 0 \)) returns\( \pm 0 \) if\( x \times y \) is exact\( \pm 0 \).

  • __fma_rd(x,y,\( \mp 0 \)) returns\( -0 \) if\( x \times y \) is exact\( \pm 0 \).

  • __fma_rd(x,y,z) returns\( -0 \) if\( x \times y + z \) is exactly zero and\( z \neq 0 \).

  • If either argument is NaN, NaN is returned.

__device__double__fma_rn(doublex,doubley,doublez)

Compute\( x \times y + z \) as a single operation in round-to-nearest-even mode.

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

Note

For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions Appendix, Intrinsic Functions section.

Returns

Returns the rounded value of\( x \times y + z \) as a single operation.

  • __fma_rn(\( \pm \infty \) ,\( \pm 0 \) ,z) returns NaN.

  • __fma_rn(\( \pm 0 \) ,\( \pm \infty \) ,z) returns NaN.

  • __fma_rn(x,y,\( -\infty \) ) returns NaN if\( x \times y \) is an exact\( +\infty \).

  • __fma_rn(x,y,\( +\infty \) ) returns NaN if\( x \times y \) is an exact\( -\infty \).

  • __fma_rn(x,y,\( \pm 0 \)) returns\( \pm 0 \) if\( x \times y \) is exact\( \pm 0 \).

  • __fma_rn(x,y,\( \mp 0 \)) returns\( +0 \) if\( x \times y \) is exact\( \pm 0 \).

  • __fma_rn(x,y,z) returns\( +0 \) if\( x \times y + z \) is exactly zero and\( z \neq 0 \).

  • If either argument is NaN, NaN is returned.

__device__double__fma_ru(doublex,doubley,doublez)

Compute\( x \times y + z \) as a single operation in round-up mode.

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

Note

For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions Appendix, Intrinsic Functions section.

Returns

Returns the rounded value of\( x \times y + z \) as a single operation.

  • __fma_ru(\( \pm \infty \) ,\( \pm 0 \) ,z) returns NaN.

  • __fma_ru(\( \pm 0 \) ,\( \pm \infty \) ,z) returns NaN.

  • __fma_ru(x,y,\( -\infty \) ) returns NaN if\( x \times y \) is an exact\( +\infty \).

  • __fma_ru(x,y,\( +\infty \) ) returns NaN if\( x \times y \) is an exact\( -\infty \).

  • __fma_ru(x,y,\( \pm 0 \)) returns\( \pm 0 \) if\( x \times y \) is exact\( \pm 0 \).

  • __fma_ru(x,y,\( \mp 0 \)) returns\( +0 \) if\( x \times y \) is exact\( \pm 0 \).

  • __fma_ru(x,y,z) returns\( +0 \) if\( x \times y + z \) is exactly zero and\( z \neq 0 \).

  • If either argument is NaN, NaN is returned.

__device__double__fma_rz(doublex,doubley,doublez)

Compute\( x \times y + z \) as a single operation in round-towards-zero mode.

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

Note

For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions Appendix, Intrinsic Functions section.

Returns

Returns the rounded value of\( x \times y + z \) as a single operation.

  • __fma_rz(\( \pm \infty \) ,\( \pm 0 \) ,z) returns NaN.

  • __fma_rz(\( \pm 0 \) ,\( \pm \infty \) ,z) returns NaN.

  • __fma_rz(x,y,\( -\infty \) ) returns NaN if\( x \times y \) is an exact\( +\infty \).

  • __fma_rz(x,y,\( +\infty \) ) returns NaN if\( x \times y \) is an exact\( -\infty \).

  • __fma_rz(x,y,\( \pm 0 \)) returns\( \pm 0 \) if\( x \times y \) is exact\( \pm 0 \).

  • __fma_rz(x,y,\( \mp 0 \)) returns\( +0 \) if\( x \times y \) is exact\( \pm 0 \).

  • __fma_rz(x,y,z) returns\( +0 \) if\( x \times y + z \) is exactly zero and\( z \neq 0 \).

  • If either argument is NaN, NaN is returned.