|
|
▶ With other host compilers: __RETURN_TYPE is ‘int’. Returns a nonzero value if
and only if a is negative.
__device__ double sin(double x)
Calculate the sine of the input argument.
Calculate the sine of the input argument x (measured in radians).
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Double-Precision Floating-Point Functions section.
Returns
▶ sin( ±0 ) returns ±0.
▶ sin( ±∞ ) returns NaN.
▶ sin(NaN) returns NaN.
__device__ void sincos(double x, double *sptr, double *cptr)
Calculate the sine and cosine of the first input argument.
Calculate the 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.
See also:
sin() and cos().
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Double-Precision Floating-Point Functions section.
__device__ void sincospi(double x, double *sptr, double *cptr)
Calculate the sine and cosine of the first input argument ×π.
Calculate the 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.
See also:
sinpi() and cospi().
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Double-Precision Floating-Point Functions section.
__device__ double sinh(double x)
Calculate the hyperbolic sine of the input argument.
Calculate the hyperbolic sine of the input argument x.
8.1. Functions
315
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Double-Precision Floating-Point Functions section.
Returns
▶ sinh( ±0 ) returns ±0.
▶ sinh( ±∞ ) returns ±∞.
▶ sinh(NaN) returns NaN.
__device__ double sinpi(double x)
Calculate the sine of the input argument ×π.
Calculate the sine of x ×π (measured in radians), where x is the input argument.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Double-Precision Floating-Point Functions section.
Returns
▶ sinpi( ±0 ) returns ±0.
▶ sinpi( ±∞ ) returns NaN.
▶ sinpi(NaN) returns NaN.
__device__ double sqrt(double x)
Calculate the square root of the input argument.
Calculate the nonnegative square root of x,
√x.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Double-Precision Floating-Point Functions section.
Returns
Returns
√x.
▶ sqrt( ±0 ) returns ±0.
▶ sqrt( +∞ ) returns +∞.
▶ sqrt(x) returns NaN if x is less than 0.
▶ sqrt(NaN) returns NaN.
__device__ double tan(double x)
Calculate the tangent of the input argument.
Calculate the tangent of the input argument x (measured in radians).
316
Chapter 8. Double Precision Mathematical Functions
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Double-Precision Floating-Point Functions section.
Returns
▶ tan( ±0 ) returns ±0.
▶ tan( ±∞ ) returns NaN.
▶ tan(NaN) returns NaN.
__device__ double tanh(double x)
Calculate the hyperbolic tangent of the input argument.
Calculate the hyperbolic tangent of the input argument x.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Double-Precision Floating-Point Functions section.
Returns
▶ tanh( ±0 ) returns ±0.
▶ tanh( ±∞ ) returns ±1.
▶ tanh(NaN) returns NaN.
__device__ double tgamma(double x)
Calculate the gamma function of the input argument.
Calculate the gamma function of the input argument x, namely the value of Γ(x).
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Double-Precision Floating-Point Functions section.
Returns
▶ tgamma( ±0 ) returns ±∞.
▶ tgamma(x) returns NaN if x < 0 and x is an integer.
▶ tgamma( -∞ ) returns NaN.
▶ tgamma( +∞ ) returns +∞.
▶ tgamma(NaN) returns NaN.
__device__ double trunc(double x)
Truncate input argument to the integral part.
Round x to the nearest integer value that does not exceed x in magnitude.
8.1. Functions
317
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Double-Precision Floating-Point Functions section.
Returns
Returns truncated integer value.
▶ trunc( ±0 ) returns ±0.
▶ trunc( ±∞ ) returns ±∞.
▶ trunc(NaN) returns NaN.
__device__ double y0(double x)
Calculate the value of the Bessel function of the second kind of order 0 for the input argument.
Calculate the value of the Bessel function of the second kind of order 0 for the input argument
x, Y0(x).
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Double-Precision Floating-Point Functions section.
Returns
Returns the value of the Bessel function of the second kind of order 0.
▶ y0( ±0 ) returns -∞.
▶ y0(x) returns NaN for x < 0.
▶ y0( +∞ ) returns +0.
▶ y0(NaN) returns NaN.
__device__ double y1(double x)
Calculate the value of the Bessel function of the second kind of order 1 for the input argument.
Calculate the value of the Bessel function of the second kind of order 1 for the input argument
x, Y1(x).
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Double-Precision Floating-Point Functions section.
Returns
Returns the value of the Bessel function of the second kind of order 1.
▶ y1( ±0 ) returns -∞.
▶ y1(x) returns NaN for x < 0.
▶ y1( +∞ ) returns +0.
▶ y1(NaN) returns NaN.
318
Chapter 8. Double Precision Mathematical Functions
__device__ double yn(int n, double x)
Calculate the value of the Bessel function of the second kind of order n for the input argument.
Calculate the value of the Bessel function of the second kind of order n for the input argument
x, Yn(x).
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Double-Precision Floating-Point Functions section.
Returns
Returns the value of the Bessel function of the second kind of order n.
▶ yn(n, x) returns NaN for n < 0.
▶ yn(n, ±0 ) returns -∞.
▶ yn(n, x) returns NaN for x < 0.
▶ yn(n, +∞ ) returns +0.
▶ yn(n, NaN) returns NaN.
8.1. Functions
319
320
Chapter 8. Double Precision Mathematical Functions
Chapter 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)
Compute1x in round-down mode.
__device__ double __drcp_rn(double x)
Compute1x in round-to-nearest-even mode.
321
__device__ double __drcp_ru(double x)
Compute1x in round-up mode.
__device__ double __drcp_rz(double x)
Compute1x in round-towards-zero mode.
__device__ dou√e __dsqrt_rd(double x)
Compute
x in round-down mode.
__device__ dou√e __dsqrt_rn(double x)
Compute
x in round-to-nearest-even mode.
__device__ dou√e __dsqrt_ru(double x)
Compute
x in round-up mode.
__device__ dou√e __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.
9.1. Functions
__device__ double __dadd_rd(double x, double y)
Add two floating-point values in round-down mode.
Adds two floating-point values x and y 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.
322
Chapter 9. Double Precision Intrinsics
Returns
Returns x + y.
▶ __dadd_rd(x, y) is equivalent to __dadd_rd(y, x).
▶ __dadd_rd(x, ±∞) returns ±∞ for finite x.
▶ __dadd_rd( ±∞, ±∞) returns ±∞.
▶ __dadd_rd( ±∞, ∓∞) returns NaN.
▶ __dadd_rd( ±0, ±0) returns ±0.
▶ __dadd_rd(x, -x) returns -0 for finite x, including ±0.
▶ If either argument is NaN, NaN is returned.
__device__ double __dadd_rn(double x, double y)
Add two floating-point values in round-to-nearest-even mode.
Adds two floating-point values x and y 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
Returns x + y.
▶ __dadd_rn(x, y) is equivalent to __dadd_rn(y, x).
▶ __dadd_rn(x, ±∞) returns ±∞ for finite x.
▶ __dadd_rn( ±∞, ±∞) returns ±∞.
▶ __dadd_rn( ±∞, ∓∞) returns NaN.
▶ __dadd_rn( ±0, ±0) returns ±0.
▶ __dadd_rn(x, -x) returns +0 for finite x, including ±0.
▶ If either argument is NaN, NaN is returned.
__device__ double __dadd_ru(double x, double y)
Add two floating-point values in round-up mode.
Adds two floating-point values x and y 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.
9.1. Functions
323
Returns
Returns x + y.
▶ __dadd_ru(x, y) is equivalent to __dadd_ru(y, x).
▶ __dadd_ru(x, ±∞) returns ±∞ for finite x.
▶ __dadd_ru( ±∞, ±∞) returns ±∞.
▶ __dadd_ru( ±∞, ∓∞) returns NaN.
▶ __dadd_ru( ±0, ±0) returns ±0.
▶ __dadd_ru(x, -x) returns +0 for finite x, including ±0.
▶ If either argument is NaN, NaN is returned.
__device__ double __dadd_rz(double x, double y)
Add two floating-point values in round-towards-zero mode.
Adds two floating-point values x and y 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
Returns x + y.
▶ __dadd_rz(x, y) is equivalent to __dadd_rz(y, x).
▶ __dadd_rz(x, ±∞) returns ±∞ for finite x.
▶ __dadd_rz( ±∞, ±∞) returns ±∞.
▶ __dadd_rz( ±∞, ∓∞) returns NaN.
▶ __dadd_rz( ±0, ±0) returns ±0.
▶ __dadd_rz(x, -x) returns +0 for finite x, including ±0.
▶ If either argument is NaN, NaN is returned.
__device__ double __ddiv_rd(double x, double y)
Divide two floating-point values in round-down mode.
Divides two floating-point values x by y 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.
324
Chapter 9. Double Precision Intrinsics
Returns
Returns x / y.
▶ sign of the quotient x / y is XOR of the signs of x and y when neither inputs nor
result are NaN.
▶ __ddiv_rd( ±0, ±0) returns NaN.
▶ __ddiv_rd( ±∞, ±∞) returns NaN.
▶ __ddiv_rd(x, ±∞) returns 0 of appropriate sign for finite x.
▶ __ddiv_rd( ±∞, y) returns ∞ of appropriate sign for finite y.
▶ __ddiv_rd(x, ±0) returns ∞ of appropriate sign for x = 0.
▶ __ddiv_rd( ±0, y) returns 0 of appropriate sign for y = 0.
▶ If either argument is NaN, NaN is returned.
__device__ double __ddiv_rn(double x, double y)
Divide two floating-point values in round-to-nearest-even mode.
Divides two floating-point values x by y 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 x / y.
▶ sign of the quotient x / y is XOR of the signs of x and y when neither inputs nor
result are NaN.
▶ __ddiv_rn( ±0, ±0) returns NaN.
▶ __ddiv_rn( ±∞, ±∞) returns NaN.
▶ __ddiv_rn(x, ±∞) returns 0 of appropriate sign for finite x.
▶ __ddiv_rn( ±∞, y) returns ∞ of appropriate sign for finite y.
▶ __ddiv_rn(x, ±0) returns ∞ of appropriate sign for x = 0.
▶ __ddiv_rn( ±0, y) returns 0 of appropriate sign for y = 0.
▶ If either argument is NaN, NaN is returned.
__device__ double __ddiv_ru(double x, double y)
Divide two floating-point values in round-up mode.
Divides two floating-point values x by y in round-up (to positive infinity) mode.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Intrinsic Functions section.
9.1. Functions
325
Note: Requires compute capability >= 2.0.
Returns
Returns x / y.
▶ sign of the quotient x / y is XOR of the signs of x and y when neither inputs nor
result are NaN.
▶ __ddiv_ru( ±0, ±0) returns NaN.
▶ __ddiv_ru( ±∞, ±∞) returns NaN.
▶ __ddiv_ru(x, ±∞) returns 0 of appropriate sign for finite x.
▶ __ddiv_ru( ±∞, y) returns ∞ of appropriate sign for finite y.
▶ __ddiv_ru(x, ±0) returns ∞ of appropriate sign for x = 0.
▶ __ddiv_ru( ±0, y) returns 0 of appropriate sign for y = 0.
▶ If either argument is NaN, NaN is returned.
__device__ double __ddiv_rz(double x, double y)
Divide two floating-point values in round-towards-zero mode.
Divides two floating-point values x by y 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 x / y.
▶ sign of the quotient x / y is XOR of the signs of x and y when neither inputs nor
result are NaN.
▶ __ddiv_rz( ±0, ±0) returns NaN.
▶ __ddiv_rz( ±∞, ±∞) returns NaN.
▶ __ddiv_rz(x, ±∞) returns 0 of appropriate sign for finite x.
▶ __ddiv_rz( ±∞, y) returns ∞ of appropriate sign for finite y.
▶ __ddiv_rz(x, ±0) returns ∞ of appropriate sign for x = 0.
▶ __ddiv_rz( ±0, y) returns 0 of appropriate sign for y = 0.
▶ If either argument is NaN, NaN is returned.
__device__ double __dmul_rd(double x, double y)
Multiply two floating-point values in round-down mode.
Multiplies two floating-point values x and y in round-down (to negative infinity) mode.
326
Chapter 9. Double Precision Intrinsics
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
Returns x * y.
▶ sign of the product x * y is XOR of the signs of x and y when neither inputs nor
result are NaN.
▶ __dmul_rd(x, y) is equivalent to __dmul_rd(y, x).
▶ __dmul_rd(x, ±∞) returns ∞ of appropriate sign for x = 0.
▶ __dmul_rd( ±0, ±∞) returns NaN.
▶ __dmul_rd( ±0, y) returns 0 of appropriate sign for finite y.
▶ If either argument is NaN, NaN is returned.
__device__ double __dmul_rn(double x, double y)
Multiply two floating-point values in round-to-nearest-even mode.
Multiplies two floating-point values x and y 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
Returns x * y.
▶ sign of the product x * y is XOR of the signs of x and y when neither inputs nor
result are NaN.
▶ __dmul_rn(x, y) is equivalent to __dmul_rn(y, x).
▶ __dmul_rn(x, ±∞) returns ∞ of appropriate sign for x = 0.
▶ __dmul_rn( ±0, ±∞) returns NaN.
▶ __dmul_rn( ±0, y) returns 0 of appropriate sign for finite y.
▶ If either argument is NaN, NaN is returned.
__device__ double __dmul_ru(double x, double y)
Multiply two floating-point values in round-up mode.
Multiplies two floating-point values x and y in round-up (to positive infinity) mode.
9.1. Functions
327
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
Returns x * y.
▶ sign of the product x * y is XOR of the signs of x and y when neither inputs nor
result are NaN.
▶ __dmul_ru(x, y) is equivalent to __dmul_ru(y, x).
▶ __dmul_ru(x, ±∞) returns ∞ of appropriate sign for x = 0.
▶ __dmul_ru( ±0, ±∞) returns NaN.
▶ __dmul_ru( ±0, y) returns 0 of appropriate sign for finite y.
▶ If either argument is NaN, NaN is returned.
__device__ double __dmul_rz(double x, double y)
Multiply two floating-point values in round-towards-zero mode.
Multiplies two floating-point values x and y 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
Returns x * y.
▶ sign of the product x * y is XOR of the signs of x and y when neither inputs nor
result are NaN.
▶ __dmul_rz(x, y) is equivalent to __dmul_rz(y, x).
▶ __dmul_rz(x, ±∞) returns ∞ of appropriate sign for x = 0.
▶ __dmul_rz( ±0, ±∞) returns NaN.
▶ __dmul_rz( ±0, y) returns 0 of appropriate sign for finite y.
▶ If either argument is NaN, NaN is returned.
__device__ double __drcp_rd(double x)
Compute1x in round-down mode.
Compute the reciprocal of x in round-down (to negative infinity) mode.
328
Chapter 9. Double Precision Intrinsics
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Intrinsic Functions section.
Note: Requires compute capability >= 2.0.
Returns
Returns1x .
__device__ double __drcp_rn(double x)
Compute1x in round-to-nearest-even mode.
Compute the reciprocal of x 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
Returns1x .
__device__ double __drcp_ru(double x)
Compute1x in round-up mode.
Compute the reciprocal of x 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
Returns1x .
__device__ double __drcp_rz(double x)
Compute1x in round-towards-zero mode.
Compute the reciprocal of x in round-towards-zero mode.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Intrinsic Functions section.
9.1. Functions
329
Note: Requires compute capability >= 2.0.
Returns
Returns1x .
__device__ double __dsqrt_rd(double x)
Compute
√x in round-down mode.
Compute the square root of x 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
√x.
__device__ double __dsqrt_rn(double x)
Compute
√x in round-to-nearest-even mode.
Compute the square root of x 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
√x.
__device__ double __dsqrt_ru(double x)
Compute
√x in round-up mode.
Compute the square root of x 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
√x.
330
Chapter 9. Double Precision Intrinsics
__device__ double __dsqrt_rz(double x)
Compute
√x in round-towards-zero mode.
Compute the square root of x 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
√x.
__device__ double __dsub_rd(double x, double y)
Subtract two floating-point values in round-down mode.
Subtracts two floating-point values x and y 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
Returns x - y.
▶ __dsub_rd( ±∞, y) returns ±∞ for finite y.
▶ __dsub_rd(x, ±∞) returns ∓∞ for finite x.
▶ __dsub_rd( ±∞, ±∞) returns NaN.
▶ __dsub_rd( ±∞, ∓∞) returns ±∞.
▶ __dsub_rd( ±0, ∓0) returns ±0.
▶ __dsub_rd(x, x) returns -0 for finite x, including ±0.
▶ If either argument is NaN, NaN is returned.
__device__ double __dsub_rn(double x, double y)
Subtract two floating-point values in round-to-nearest-even mode.
Subtracts two floating-point values x and y 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.
9.1. Functions
331
Returns
Returns x - y.
▶ __dsub_rn( ±∞, y) returns ±∞ for finite y.
▶ __dsub_rn(x, ±∞) returns ∓∞ for finite x.
▶ __dsub_rn( ±∞, ±∞) returns NaN.
▶ __dsub_rn( ±∞, ∓∞) returns ±∞.
▶ __dsub_rn( ±0, ∓0) returns ±0.
▶ __dsub_rn(x, x) returns +0 for finite x, including ±0.
▶ If either argument is NaN, NaN is returned.
__device__ double __dsub_ru(double x, double y)
Subtract two floating-point values in round-up mode.
Subtracts two floating-point values x and y 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
Returns x - y.
▶ __dsub_ru( ±∞, y) returns ±∞ for finite y.
▶ __dsub_ru(x, ±∞) returns ∓∞ for finite x.
▶ __dsub_ru( ±∞, ±∞) returns NaN.
▶ __dsub_ru( ±∞, ∓∞) returns ±∞.
▶ __dsub_ru( ±0, ∓0) returns ±0.
▶ __dsub_ru(x, x) returns +0 for finite x, including ±0.
▶ If either argument is NaN, NaN is returned.
__device__ double __dsub_rz(double x, double y)
Subtract two floating-point values in round-towards-zero mode.
Subtracts two floating-point values x and y 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.
332
Chapter 9. Double Precision Intrinsics
Returns
Returns x - y.
▶ __dsub_rz( ±∞, y) returns ±∞ for finite y.
▶ __dsub_rz(x, ±∞) returns ∓∞ for finite x.
▶ __dsub_rz( ±∞, ±∞) returns NaN.
▶ __dsub_rz( ±∞, ∓∞) returns ±∞.
▶ __dsub_rz( ±0, ∓0) returns ±0.
▶ __dsub_rz(x, x) returns +0 for finite x, including ±0.
▶ If either argument is NaN, NaN is returned.
__device__ double __fma_rd(double x, double y, double z)
Compute x × y + z as a single operation in round-down mode.
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, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Intrinsic Functions section.
Returns
Returns the rounded value of x × y + z as a single operation.
▶ __fma_rd( ±∞ , ±0 , z) returns NaN.
▶ __fma_rd( ±0 , ±∞ , z) returns NaN.
▶ __fma_rd(x, y, -∞ ) returns NaN if x × y is an exact +∞.
▶ __fma_rd(x, y, +∞ ) returns NaN if x × y is an exact -∞.
▶ __fma_rd(x, y, ±0) returns ±0 if x × y is exact ±0.
▶ __fma_rd(x, y, ∓0) returns -0 if x × y is exact ±0.
▶ __fma_rd(x, y, z) returns -0 if x × y + z is exactly zero and z = 0.
▶ If either argument is NaN, NaN is returned.
__device__ double __fma_rn(double x, double y, double z)
Compute x × y + z as a single operation in round-to-nearest-even mode.
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, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Intrinsic Functions section.
Returns
Returns the rounded value of x × y + z as a single operation.
▶ __fma_rn( ±∞ , ±0 , z) returns NaN.
▶ __fma_rn( ±0 , ±∞ , z) returns NaN.
9.1. Functions
333
▶ __fma_rn(x, y, -∞ ) returns NaN if x × y is an exact +∞.
▶ __fma_rn(x, y, +∞ ) returns NaN if x × y is an exact -∞.
▶ __fma_rn(x, y, ±0) returns ±0 if x × y is exact ±0.
▶ __fma_rn(x, y, ∓0) returns +0 if x × y is exact ±0.
▶ __fma_rn(x, y, z) returns +0 if x × y + z is exactly zero and z = 0.
▶ If either argument is NaN, NaN is returned.
__device__ double __fma_ru(double x, double y, double z)
Compute x × y + z as a single operation in round-up mode.
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, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Intrinsic Functions section.
Returns
Returns the rounded value of x × y + z as a single operation.
▶ __fma_ru( ±∞ , ±0 , z) returns NaN.
▶ __fma_ru( ±0 , ±∞ , z) returns NaN.
▶ __fma_ru(x, y, -∞ ) returns NaN if x × y is an exact +∞.
▶ __fma_ru(x, y, +∞ ) returns NaN if x × y is an exact -∞.
▶ __fma_ru(x, y, ±0) returns ±0 if x × y is exact ±0.
▶ __fma_ru(x, y, ∓0) returns +0 if x × y is exact ±0.
▶ __fma_ru(x, y, z) returns +0 if x × y + z is exactly zero and z = 0.
▶ If either argument is NaN, NaN is returned.
__device__ double __fma_rz(double x, double y, double z)
Compute x × y + z as a single operation in round-towards-zero mode.
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, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Intrinsic Functions section.
Returns
Returns the rounded value of x × y + z as a single operation.
▶ __fma_rz( ±∞ , ±0 , z) returns NaN.
▶ __fma_rz( ±0 , ±∞ , z) returns NaN.
▶ __fma_rz(x, y, -∞ ) returns NaN if x × y is an exact +∞.
▶ __fma_rz(x, y, +∞ ) returns NaN if x × y is an exact -∞.
334
Chapter 9. Double Precision Intrinsics
▶ __fma_rz(x, y, ±0) returns ±0 if x × y is exact ±0.
▶ __fma_rz(x, y, ∓0) returns +0 if x × y is exact ±0.
▶ __fma_rz(x, y, z) returns +0 if x × y + z is exactly zero and z = 0.
▶ If either argument is NaN, NaN is returned.
9.1. Functions
335
336
Chapter 9. Double Precision Intrinsics
Chapter 10. FP128 Quad Precision
Mathematical Functions
This section describes quad precision mathematical functions.
To use these functions, include the header file device_fp128_functions.h in your program.
Functions declared here have __nv_fp128_ prefix to distinguish them from other global namespace
symbols.
Note that FP128 CUDA Math functions are only available to device programs on platforms where host
compiler supports the basic quad precision datatype __float128 or _Float128.
Every FP128 CUDA Math function name is overloaded to support either of these host-compiler-
specific types, whenever the types are available. See for example:
#ifdef __FLOAT128_CPP_SPELLING_ENABLED__
__float128 __nv_fp128_sqrt(__float128 x);
#endif
#ifdef __FLOAT128_C_SPELLING_ENABLED__
_Float128 __nv_fp128_sqrt(_Float128 x);
#endif
Note: FP128 device computations require compute capability >= 10.0.
Functions
__device__ __float128 __nv_fp128_acos(__float128 x)
Calculate cos-1 x , the arc cosine of input argument.
__device__ __float128 __nv_fp128_acosh(__float128 x)
Calculate cosh-1 x , the nonnegative inverse hyperbolic cosine of the input argument.
__device__ __float128 __nv_fp128_add(__float128 x, __float128 y)
Compute x + y , the sum of the two floating-point inputs using round-to-nearest-even rounding
mode.
__device__ __float128 __nv_fp128_asin(__float128 x)
Calculate sin-1 x , the arc sine of input argument.
__device__ __float128 __nv_fp128_asinh(__float128 x)
Calculate sinh-1 x , the inverse hyperbolic sine of the input argument.
337
__device__ __float128 __nv_fp128_atan(__float128 x)
Calculate tan-1 x , the arc tangent of input argument.
__device__ __float128 __nv_fp128_atanh(__float128 x)
Calculate tanh-1 x , the inverse hyperbolic tangent of the input argument.
__device__ __float128 __nv_fp128_ceil(__float128 x)
Calculate ⌈x⌉ , the smallest integer greater than or equal to x .
__device__ __float128 __nv_fp128_copysign(__float128 x, __float128 y)
Create value with the magnitude of the first agument x , and the sign of the second argument y
__device__ __float128 __nv_fp128_cos(__float128 x)
Calculate cos x , the cosine of input argument (measured in radians).
__device__ __float128 __nv_fp128_cosh(__float128 x)
Calculate cosh x , the hyperbolic cosine of the input argument.
__device__ __float128 __nv_fp128_div(__float128 x, __float128 y)
Computexy , the quotient of the two floating-point inputs using round-to-nearest-even rounding
mode.
__device__ __float128 __nv_fp128_exp(__float128 x)
Calculate ex , the base e exponential of the input argument.
__device__ __float128 __nv_fp128_exp10(__float128 x)
Calculate 10x , the base 10 exponential of the input argument.
__device__ __float128 __nv_fp128_exp2(__float128 x)
Calculate 2x , the base 2 exponential of the input argument.
__device__ __float128 __nv_fp128_expm1(__float128 x)
Calculate ex - 1 , the base e exponential of the input argument, minus 1.
__device__ __float128 __nv_fp128_fabs(__float128 x)
Calculate |x| , the absolute value of the input argument.
__device__ __float128 __nv_fp128_fdim(__float128 x, __float128 y)
Compute the positive difference between x and y .
__device__ __float128 __nv_fp128_floor(__float128 x)
Calculate ⌊x⌋ , the largest integer less than or equal to x .
__device__ __float128 __nv_fp128_fma(__float128 x, __float128 y, __float128 c)
Compute x × y + z as a single operation using round-to-nearest-even rounding mode.
__device__ __float128 __nv_fp128_fmax(__float128 x, __float128 y)
Determine the maximum numeric value of the arguments.
__device__ __float128 __nv_fp128_fmin(__float128 x, __float128 y)
Determine the minimum numeric value of the arguments.
__device__ __float128 __nv_fp128_fmod(__float128 x, __float128 y)
Calculate the floating-point remainder of x / y .
__device__ __float128 __nv_fp128_frexp(__float128 x, int *nptr)
Extract mantissa and exponent of the floating-point input argument.
__device__ __fl√t128 __nv_fp128_hypot(__float128 x, __float128 y)
Calculate
x2 + y2 , the square root of the sum of squares of two arguments.
__device__ int __nv_fp128_ilogb(__float128 x)
Compute the unbiased integer exponent of the input argument.
338
Chapter 10. FP128 Quad Precision Mathematical Functions
__device__ int __nv_fp128_isnan(__float128 x)
Determine whether the input argument is a NaN.
__device__ int __nv_fp128_isunordered(__float128 x, __float128 y)
Determine whether the pair of inputs is unordered.
__device__ __float128 __nv_fp128_ldexp(__float128 x, int exp)
Calculate the value of x · 2exp .
__device__ __float128 __nv_fp128_log(__float128 x)
Calculate loge x , the base e logarithm of the input argument.
__device__ __float128 __nv_fp128_log10(__float128 x)
Calculate log10 x , the base 10 logarithm of the input argument.
__device__ __float128 __nv_fp128_log1p(__float128 x)
Calculate the value of loge(1 + x) .
__device__ __float128 __nv_fp128_log2(__float128 x)
Calculate log2 x , the base 2 logarithm of the input argument.
__device__ __float128 __nv_fp128_modf(__float128 x, __float128 *iptr)
Break down the input argument into fractional and integral parts.
__device__ __float128 __nv_fp128_mul(__float128 x, __float128 y)
Compute x·y , the product of the two floating-point inputs using round-to-nearest-even rounding
mode.
__device__ __float128 __nv_fp128_pow(__float128 x, __float128 y)
Calculate the value of xy , first argument to the power of second argument.
__device__ __float128 __nv_fp128_remainder(__float128 x, __float128 y)
Compute the floating-point remainder function.
__device__ __float128 __nv_fp128_rint(__float128 x)
Round to nearest integer value in floating-point format, with halfway cases rounded to the near-
est even integer value.
__device__ __float128 __nv_fp128_round(__float128 x)
Round to nearest integer value in floating-point format, with halfway cases rounded away from
zero.
__device__ __float128 __nv_fp128_sin(__float128 x)
Calculate sin x , the sine of input argument (measured in radians).
__device__ __float128 __nv_fp128_sinh(__float128 x)
Calculate sinh x , the hyperbolic sine of the input argument.
__device__ __fl√t128 __nv_fp128_sqrt(__float128 x)
Calculate
x , the square root of the input argument.
__device__ __float128 __nv_fp128_sub(__float128 x, __float128 y)
Compute x - y , the difference of the two floating-point inputs using round-to-nearest-even
rounding mode.
__device__ __float128 __nv_fp128_tan(__float128 x)
Calculate tan x , the tangent of input argument (measured in radians).
__device__ __float128 __nv_fp128_tanh(__float128 x)
Calculate tanh x , the hyperbolic tangent of the input argument.
__device__ __float128 __nv_fp128_trunc(__float128 x)
Truncate input argument to the integral part.
339
10.1. Functions
__device__ __float128 __nv_fp128_acos(__float128 x)
Calculate cos-1 x, the arc cosine of input argument.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
Note: FP128 device computations require compute capability >= 10.0.
Returns
The principal value of the arc cosine of the input argument x. Result will be in radi-
ans, in the interval [0, π ] for x inside [-1, +1].
▶ __nv_fp128_acos(1) returns +0.
▶ __nv_fp128_acos(x) returns NaN for x outside [-1, +1].
▶ __nv_fp128_acos(NaN) returns NaN.
__device__ __float128 __nv_fp128_acosh(__float128 x)
Calculate cosh-1 x, the nonnegative inverse hyperbolic cosine of the input argument.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
Note: FP128 device computations require compute capability >= 10.0.
Returns
Result will be in the interval [0, +∞ ].
▶ __nv_fp128_acosh(1) returns 0.
▶ __nv_fp128_acosh(x) returns NaN for x in the interval [ -∞ , 1).
▶ __nv_fp128_acosh( +∞ ) returns +∞.
▶ __nv_fp128_acosh(NaN) returns NaN.
__device__ __float128 __nv_fp128_add(__float128 x, __float128 y)
Compute x + y, the sum of the two floating-point inputs using round-to-nearest-even rounding
mode.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
340
Chapter 10. FP128 Quad Precision Mathematical Functions
Note: FP128 device computations require compute capability >= 10.0.
Returns
Returns x + y.
▶ __nv_fp128_add(x, y) is equivalent to __nv_fp128_add(y, x).
▶ __nv_fp128_add(x, ±∞) returns ±∞ for finite x.
▶ __nv_fp128_add( ±∞, ±∞) returns ±∞.
▶ __nv_fp128_add( ±∞, ∓∞) returns NaN.
▶ __nv_fp128_add( ±0, ±0) returns ±0.
▶ __nv_fp128_add(x, -x) returns +0 for finite x, including ±0.
▶ If either argument is NaN, NaN is returned.
__device__ __float128 __nv_fp128_asin(__float128 x)
Calculate sin-1 x, the arc sine of input argument.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
Note: FP128 device computations require compute capability >= 10.0.
Returns
The principal value of the arc sine of the input argument x. Result will be in radians,
in the interval [- π/2 , + π/2 ] for x inside [-1, +1].
▶ __nv_fp128_asin( ±0 ) returns ±0.
▶ __nv_fp128_asin(x) returns NaN for x outside [-1, +1].
▶ __nv_fp128_asin(NaN) returns NaN.
__device__ __float128 __nv_fp128_asinh(__float128 x)
Calculate sinh-1 x, the inverse hyperbolic sine of the input argument.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
Note: FP128 device computations require compute capability >= 10.0.
Returns
▶ __nv_fp128_asinh( ±0 ) returns ±0.
▶ __nv_fp128_asinh( ±∞ ) returns ±∞.
10.1. Functions
341
▶ __nv_fp128_asinh(NaN) returns NaN.
__device__ __float128 __nv_fp128_atan(__float128 x)
Calculate tan-1 x, the arc tangent of input argument.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
Note: FP128 device computations require compute capability >= 10.0.
Returns
The principal value of the arc tangent of the input argument x. Result will be in
radians, in the interval [- π/2 , + π/2 ].
▶ __nv_fp128_atan( ±0 ) returns ±0.
▶ __nv_fp128_atan( ±∞ ) returns ±π /2.
▶ __nv_fp128_atan(NaN) returns NaN.
__device__ __float128 __nv_fp128_atanh(__float128 x)
Calculate tanh-1 x, the inverse hyperbolic tangent of the input argument.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
Note: FP128 device computations require compute capability >= 10.0.
Returns
▶ __nv_fp128_atanh( ±0 ) returns ±0.
▶ __nv_fp128_atanh( ±1 ) returns ±∞.
▶ __nv_fp128_atanh(x) returns NaN for x outside interval [-1, 1].
▶ __nv_fp128_atanh(NaN) returns NaN.
__device__ __float128 __nv_fp128_ceil(__float128 x)
Calculate ⌈x⌉, the smallest integer greater than or equal to x.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
Note: FP128 device computations require compute capability >= 10.0.
342
Chapter 10. FP128 Quad Precision Mathematical Functions
Returns
⌈x⌉ expressed as a floating-point number.
▶ __nv_fp128_ceil( ±∞ ) returns ±∞.
▶ __nv_fp128_ceil( ±0 ) returns ±0.
▶ __nv_fp128_ceil(NaN) returns NaN.
__device__ __float128 __nv_fp128_copysign(__float128 x, __float128 y)
Create value with the magnitude of the first agument x, and the sign of the second argument y.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
Note: FP128 device computations require compute capability >= 10.0.
Returns
▶ copysign(NaN, y) returns a NaN with the sign of y.
__device__ __float128 __nv_fp128_cos(__float128 x)
Calculate cos x, the cosine of input argument (measured in radians).
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
Note: FP128 device computations require compute capability >= 10.0.
Returns
cos x.
▶ __nv_fp128_cos( ±0 ) returns 1.
▶ __nv_fp128_cos( ±∞ ) returns NaN.
▶ __nv_fp128_cos(NaN) returns NaN.
__device__ __float128 __nv_fp128_cosh(__float128 x)
Calculate cosh x, the hyperbolic cosine of the input argument.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
Note: FP128 device computations require compute capability >= 10.0.
10.1. Functions
343
Returns
▶ __nv_fp128_cosh( ±0 ) returns 1.
▶ __nv_fp128_cosh( ±∞ ) returns +∞.
▶ __nv_fp128_cosh(NaN) returns NaN.
__device__ __float128 __nv_fp128_div(__float128 x, __float128 y)
Computexy , the quotient of the two floating-point inputs using round-to-nearest-even rounding
mode.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
Note: FP128 device computations require compute capability >= 10.0.
Returns
▶ sign of the quotient x / y is XOR of the signs of x and y when neither inputs nor
result are NaN.
▶ __nv_fp128_div( ±0, ±0) returns NaN.
▶ __nv_fp128_div( ±∞, ±∞) returns NaN.
▶ __nv_fp128_div(x, ±∞) returns 0 of appropriate sign for finite x.
▶ __nv_fp128_div( ±∞, y) returns ∞ of appropriate sign for finite y.
▶ __nv_fp128_div(x, ±0) returns ∞ of appropriate sign for x = 0.
▶ __nv_fp128_div( ±0, y) returns 0 of appropriate sign for y = 0.
▶ If either argument is NaN, NaN is returned.
__device__ __float128 __nv_fp128_exp(__float128 x)
Calculate ex, the base e exponential of the input argument.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
Note: FP128 device computations require compute capability >= 10.0.
Returns
▶ __nv_fp128_exp( ±0 ) returns 1.
▶ __nv_fp128_exp( -∞ ) returns +0.
▶ __nv_fp128_exp( +∞ ) returns +∞.
344
Chapter 10. FP128 Quad Precision Mathematical Functions
▶ __nv_fp128_exp(NaN) returns NaN.
__device__ __float128 __nv_fp128_exp10(__float128 x)
Calculate 10x, the base 10 exponential of the input argument.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
Note: FP128 device computations require compute capability >= 10.0.
Returns
▶ __nv_fp128_exp10( ±0 ) returns 1.
▶ __nv_fp128_exp10( -∞ ) returns +0.
▶ __nv_fp128_exp10( +∞ ) returns +∞.
▶ __nv_fp128_exp10(NaN) returns NaN.
__device__ __float128 __nv_fp128_exp2(__float128 x)
Calculate 2x, the base 2 exponential of the input argument.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
Note: FP128 device computations require compute capability >= 10.0.
Returns
▶ __nv_fp128_exp2( ±0 ) returns 1.
▶ ex__nv_fp128_exp2p2f( -∞ ) returns +0.
▶ __nv_fp128_exp2( +∞ ) returns +∞.
▶ __nv_fp128_exp2(NaN) returns NaN.
__device__ __float128 __nv_fp128_expm1(__float128 x)
Calculate ex - 1, the base e exponential of the input argument, minus 1.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
Note: FP128 device computations require compute capability >= 10.0.
10.1. Functions
345
Returns
▶ __nv_fp128_expm1( ±0 ) returns ±0.
▶ __nv_fp128_expm1( -∞ ) returns -1.
▶ __nv_fp128_expm1( +∞ ) returns +∞.
▶ __nv_fp128_expm1(NaN) returns NaN.
__device__ __float128 __nv_fp128_fabs(__float128 x)
Calculate |x|, the absolute value of the input argument.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
Note: FP128 device computations require compute capability >= 10.0.
Returns
▶ __nv_fp128_fabs( ±∞ ) returns +∞.
▶ __nv_fp128_fabs( ±0 ) returns +0.
▶ __nv_fp128_fabs(NaN) returns an unspecified NaN.
__device__ __float128 __nv_fp128_fdim(__float128 x, __float128 y)
Compute the positive difference between x and y.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
Note: FP128 device computations require compute capability >= 10.0.
Returns
▶ __nv_fp128_fdim(x, y) returns x - y if x > y.
▶ __nv_fp128_fdim(x, y) returns +0 if x ≤ y.
▶ If either argument is NaN, NaN is returned.
__device__ __float128 __nv_fp128_floor(__float128 x)
Calculate ⌊x⌋, the largest integer less than or equal to x.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
346
Chapter 10. FP128 Quad Precision Mathematical Functions
Note: FP128 device computations require compute capability >= 10.0.
Returns
⌊x⌋ expressed as a floating-point number.
▶ __nv_fp128_floor( ±∞ ) returns ±∞.
▶ __nv_fp128_floor( ±0 ) returns ±0.
▶ __nv_fp128_floor(NaN) returns NaN.
__device__ __float128 __nv_fp128_fma(__float128 x, __float128 y, __float128 c)
Compute x × y + z as a single operation using round-to-nearest-even rounding mode.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
Note: FP128 device computations require compute capability >= 10.0.
Returns
The value of x × y + z as a single ternary operation, rounded once using round-to-
nearest, ties-to-even rounding mode.
▶ __nv_fp128_fma( ±∞ , ±0 , z) returns NaN.
▶ __nv_fp128_fma( ±0 , ±∞ , z) returns NaN.
▶ __nv_fp128_fma(x, y, -∞ ) returns NaN if x × y is an exact +∞.
▶ __nv_fp128_fma(x, y, +∞ ) returns NaN if x × y is an exact -∞.
▶ __nv_fp128_fma(x, y, ±0) returns ±0 if x × y is exact ±0.
▶ __nv_fp128_fma(x, y, ∓0) returns +0 if x × y is exact ±0.
▶ __nv_fp128_fma(x, y, z) returns +0 if x × y + z is exactly zero and z = 0.
▶ If either argument is NaN, NaN is returned.
__device__ __float128 __nv_fp128_fmax(__float128 x, __float128 y)
Determine the maximum numeric value of the arguments.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
Note: FP128 device computations require compute capability >= 10.0.
Returns
The maximum numeric value of the arguments x and y. Treats NaN arguments as
missing data.
10.1. Functions
347
▶ If both arguments are NaN, returns NaN.
▶ If one argument is NaN, returns the numeric argument.
__device__ __float128 __nv_fp128_fmin(__float128 x, __float128 y)
Determine the minimum numeric value of the arguments.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
Note: FP128 device computations require compute capability >= 10.0.
Returns
The minimum numeric value of the arguments x and y. Treats NaN arguments as
missing data.
▶ If both arguments are NaN, returns NaN.
▶ If one argument is NaN, returns the numeric argument.
__device__ __float128 __nv_fp128_fmod(__float128 x, __float128 y)
Calculate the floating-point remainder of x / y.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
Note: FP128 device computations require compute capability >= 10.0.
Returns
The floating-point remainder of the division operation x / y calculated by this func-
tion is exactly the value x - n*y, where n is x / y with its fractional part truncated.
▶ The computed value will have the same sign as x, and its magnitude will be less
than the magnitude of y.
▶ __nv_fp128_fmod( ±0 , y) returns ±0 if y is not zero.
▶ __nv_fp128_fmod(x, ±∞ ) returns x if x is finite.
▶ __nv_fp128_fmod(x, y) returns NaN if x is ±∞ or y is zero.
▶ If either argument is NaN, NaN is returned.
__device__ __float128 __nv_fp128_frexp(__float128 x, int *nptr)
Extract mantissa and exponent of the floating-point input argument.
Decompose the floating-point value x into a component m for the normalized fraction element
and an integral term n for the exponent. The absolute value of m will be greater than or equal to
0.5 and less than 1.0 or it will be equal to 0; x = m · 2n. The integer exponent n will be stored in
the location to which nptr points.
348
Chapter 10. FP128 Quad Precision Mathematical Functions
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
Note: FP128 device computations require compute capability >= 10.0.
Returns
The fractional component m.
▶ __nv_fp128_frexp( ±0 , nptr) returns ±0 and stores zero in the location pointed
to by nptr.
▶ __nv_fp128_frexp( ±∞ , nptr) returns ±∞ and stores an unspecified value in
the location to which nptr points.
▶ __nv_fp128_frexp(NaN, y) returns a NaN and stores an unspecified value in the
location to which nptr points.
__device__ __float128 __nv_fp128_hypot(__float128 x, __float128 y)
√
Calculate
x2 + y2, the square root of the sum of squares of two arguments.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
Note: FP128 device computations require compute capability >= 10.0.
Returns
The length of the hypotenuse of a right triangle whose two sides have lengths |x|
and |y| without undue overflow or underflow.
▶ __nv_fp128_hypot(x,y), __nv_fp128_hypot(y,x), and __nv_fp128_hypot(x, -y)
are equivalent.
▶ __nv_fp128_hypot(x, ±0 ) is equivalent to __nv_fp128_fabs(x).
▶ __nv_fp128_hypot( ±∞ ,y) returns +∞, even if y is a NaN.
▶ __nv_fp128_hypot(NaN, y) returns NaN, when y is not ±∞.
__device__ int __nv_fp128_ilogb(__float128 x)
Compute the unbiased integer exponent of the input argument.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
Note: FP128 device computations require compute capability >= 10.0.
10.1. Functions
349
Returns
▶ If successful, returns the unbiased exponent of the argument.
▶ __nv_fp128_ilogb( ±0 ) returns INT_MIN.
▶ __nv_fp128_ilogb(NaN) returns INT_MIN.
▶ __nv_fp128_ilogb( ±∞ ) returns INT_MAX.
▶ Note: above behavior does not take into account FP_ILOGB0 nor FP_ILOGBNAN.
__device__ int __nv_fp128_isnan(__float128 x)
Determine whether the input argument is a NaN.
Note: FP128 device computations require compute capability >= 10.0.
Returns
A nonzero value if and only if x is a NaN value.
__device__ int __nv_fp128_isunordered(__float128 x, __float128 y)
Determine whether the pair of inputs is unordered.
Note: FP128 device computations require compute capability >= 10.0.
Returns
▶ nonzero value if at least one of input values is a NaN.
▶ zero otherwise
__device__ __float128 __nv_fp128_ldexp(__float128 x, int exp)
Calculate the value of x · 2exp.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
Note: FP128 device computations require compute capability >= 10.0.
Returns
▶ __nv_fp128_ldexp( ±0 , exp) returns ±0.
▶ __nv_fp128_ldexp(x, 0) returns x.
▶ __nv_fp128_ldexp( ±∞ , exp) returns ±∞.
▶ __nv_fp128_ldexp(NaN, exp) returns NaN.
350
Chapter 10. FP128 Quad Precision Mathematical Functions
__device__ __float128 __nv_fp128_log(__float128 x)
Calculate loge x, the base e logarithm of the input argument.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
Note: FP128 device computations require compute capability >= 10.0.
Returns
▶ __nv_fp128_log( ±0 ) returns -∞.
▶ __nv_fp128_log(1) returns +0.
▶ __nv_fp128_log(x) returns NaN for x < 0.
▶ __nv_fp128_log( +∞ ) returns +∞.
▶ __nv_fp128_log(NaN) returns NaN.
__device__ __float128 __nv_fp128_log10(__float128 x)
Calculate log10 x, the base 10 logarithm of the input argument.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
Note: FP128 device computations require compute capability >= 10.0.
Returns
▶ __nv_fp128_log10( ±0 ) returns -∞.
▶ __nv_fp128_log10(1) returns +0.
▶ __nv_fp128_log10(x) returns NaN for x < 0.
▶ __nv_fp128_log10( +∞ ) returns +∞.
▶ __nv_fp128_log10(NaN) returns NaN.
__device__ __float128 __nv_fp128_log1p(__float128 x)
Calculate the value of loge(1 + x).
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
Note: FP128 device computations require compute capability >= 10.0.
10.1. Functions
351
Returns
▶ __nv_fp128_log1p( ±0 ) returns ±0.
▶ __nv_fp128_log1p(-1) returns -∞.
▶ __nv_fp128_log1p(x) returns NaN for x < -1.
▶ __nv_fp128_log1p( +∞ ) returns +∞.
▶ __nv_fp128_log1p(NaN) returns NaN.
__device__ __float128 __nv_fp128_log2(__float128 x)
Calculate log2 x, the base 2 logarithm of the input argument.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
Note: FP128 device computations require compute capability >= 10.0.
Returns
▶ __nv_fp128_log2( ±0 ) returns -∞.
▶ __nv_fp128_log2(1) returns +0.
▶ __nv_fp128_log2(x) returns NaN for x < 0.
▶ __nv_fp128_log2( +∞ ) returns +∞.
▶ __nv_fp128_log2(NaN) returns NaN.
__device__ __float128 __nv_fp128_modf(__float128 x, __float128 *iptr)
Break down the input argument into fractional and integral parts.
Break down the argument x into fractional and integral parts. The integral part is stored in
floating-point format in the location to which iptr points. Fractional and integral parts are given
the same sign as the argument x.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
Note: FP128 device computations require compute capability >= 10.0.
Returns
▶ __nv_fp128_modf( ±x , iptr) returns a result with the same sign as x.
▶ __nv_fp128_modf( ±∞ , iptr) returns ±0 and stores ±∞ in the object pointed
to by iptr.
352
Chapter 10. FP128 Quad Precision Mathematical Functions
▶ __nv_fp128_modf(NaN, iptr) stores a NaN in the object pointed to by iptr
and returns a NaN.
__device__ __float128 __nv_fp128_mul(__float128 x, __float128 y)
Compute x·y, the product of the two floating-point inputs using round-to-nearest-even rounding
mode.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
Note: FP128 device computations require compute capability >= 10.0.
Returns
Returns x * y.
▶ sign of the product x * y is XOR of the signs of x and y when neither inputs nor
result are NaN.
▶ __nv_fp128_mul(x, y) is equivalent to __nv_fp128_mul(y, x).
▶ __nv_fp128_mul(x, ±∞) returns ∞ of appropriate sign for x = 0.
▶ __nv_fp128_mul( ±0, ±∞) returns NaN.
▶ __nv_fp128_mul( ±0, y) returns 0 of appropriate sign for finite y.
▶ If either argument is NaN, NaN is returned.
__device__ __float128 __nv_fp128_pow(__float128 x, __float128 y)
Calculate the value of xy, first argument to the power of second argument.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
Note: FP128 device computations require compute capability >= 10.0.
Returns
▶ __nv_fp128_pow( ±0 , y) returns ±∞ for y an odd integer less than 0.
▶ __nv_fp128_pow( ±0 , y) returns +∞ for y less than 0 and not an odd integer.
▶ __nv_fp128_pow( ±0 , y) returns ±0 for y an odd integer greater than 0.
▶ __nv_fp128_pow( ±0 , y) returns +0 for y > 0 and not an odd integer.
▶ __nv_fp128_pow(-1, ±∞ ) returns 1.
▶ __nv_fp128_pow(+1, y) returns 1 for any y, even a NaN.
▶ __nv_fp128_pow(x, ±0 ) returns 1 for any x, even a NaN.
▶ __nv_fp128_pow(x, y) returns a NaN for finite x < 0 and finite non-integer y.
10.1. Functions
353
▶ __nv_fp128_pow(x, -∞ ) returns +∞ for |x| < 1.
▶ __nv_fp128_pow(x, -∞ ) returns +0 for |x| > 1.
▶ __nv_fp128_pow(x, +∞ ) returns +0 for |x| < 1.
▶ __nv_fp128_pow(x, +∞ ) returns +∞ for |x| > 1.
▶ __nv_fp128_pow( -∞ , y) returns -0 for y an odd integer less than 0.
▶ __nv_fp128_pow( -∞ , y) returns +0 for y < 0 and not an odd integer.
▶ __nv_fp128_pow( -∞ , y) returns -∞ for y an odd integer greater than 0.
▶ __nv_fp128_pow( -∞ , y) returns +∞ for y > 0 and not an odd integer.
▶ __nv_fp128_pow( +∞ , y) returns +0 for y < 0.
▶ __nv_fp128_pow( +∞ , y) returns +∞ for y > 0.
▶ __nv_fp128_pow(x, y) returns NaN if either x or y or both are NaN and x = +1
and y = ±0.
__device__ __float128 __nv_fp128_remainder(__float128 x, __float128 y)
Compute the floating-point remainder function.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
Note: FP128 device computations require compute capability >= 10.0.
Returns
The floating-point remainder r of dividing x by y for nonzero y is defined as r = x -
ny. The value n is the integer value nearestxy . In the halfway cases when |n -xy | =12
, the even n value is chosen.
▶ __nv_fp128_remainder(x, ±0 ) returns NaN.
▶ __nv_fp128_remainder( ±∞ , y) returns NaN.
▶ __nv_fp128_remainder(x, ±∞ ) returns x for finite x.
▶ If either argument is NaN, NaN is returned.
__device__ __float128 __nv_fp128_rint(__float128 x)
Round to nearest integer value in floating-point format, with halfway cases rounded to the near-
est even integer value.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
Note: FP128 device computations require compute capability >= 10.0.
354
Chapter 10. FP128 Quad Precision Mathematical Functions
Returns
▶ __nv_fp128_rint( ±0 ) returns ±0.
▶ __nv_fp128_rint( ±∞ ) returns ±∞.
▶ __nv_fp128_rint(NaN) returns NaN.
__device__ __float128 __nv_fp128_round(__float128 x)
Round to nearest integer value in floating-point format, with halfway cases rounded away from
zero.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
Note: FP128 device computations require compute capability >= 10.0.
Returns
▶ __nv_fp128_round( ±0 ) returns ±0.
▶ __nv_fp128_round( ±∞ ) returns ±∞.
▶ __nv_fp128_round(NaN) returns NaN.
__device__ __float128 __nv_fp128_sin(__float128 x)
Calculate sin x, the sine of input argument (measured in radians).
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
Note: FP128 device computations require compute capability >= 10.0.
Returns
sin x.
▶ __nv_fp128_sin( ±0 ) returns ±0.
▶ __nv_fp128_sin( ±∞ ) returns NaN.
▶ __nv_fp128_sin(NaN) returns NaN.
__device__ __float128 __nv_fp128_sinh(__float128 x)
Calculate sinh x, the hyperbolic sine of the input argument.
Calculate sinh x, the hyperbolic sine of the input argument x.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
10.1. Functions
355
Note: FP128 device computations require compute capability >= 10.0.
Returns
▶ __nv_fp128_sinhinh( ±0 ) returns ±0.
▶ __nv_fp128_sinh( ±∞ ) returns ±∞.
▶ __nv_fp128_sinh(NaN) returns NaN.
__device__ __float128 __nv_fp128_sqrt(__float128 x)
Calculate
√x, the square root of the input argument.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
Note: FP128 device computations require compute capability >= 10.0.
Returns√
x.
▶ __nv_fp128_sqrt( ±0 ) returns ±0.
▶ __nv_fp128_sqrt( +∞ ) returns +∞.
▶ __nv_fp128_sqrt(x) returns NaN if x is less than 0.
▶ __nv_fp128_sqrt(NaN) returns NaN.
__device__ __float128 __nv_fp128_sub(__float128 x, __float128 y)
Compute x - y, the difference of the two floating-point inputs using round-to-nearest-even
rounding mode.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
Note: FP128 device computations require compute capability >= 10.0.
Returns
Returns x - y.
▶ __nv_fp128_sub( ±∞, y) returns ±∞ for finite y.
▶ __nv_fp128_sub(x, ±∞) returns ∓∞ for finite x.
▶ __nv_fp128_sub( ±∞, ±∞) returns NaN.
▶ __nv_fp128_sub( ±∞, ∓∞) returns ±∞.
▶ __nv_fp128_sub( ±0, ∓0) returns ±0.
356
Chapter 10. FP128 Quad Precision Mathematical Functions
▶ __nv_fp128_sub(x, x) returns +0 for finite x, including ±0.
▶ If either argument is NaN, NaN is returned.
__device__ __float128 __nv_fp128_tan(__float128 x)
Calculate tan x, the tangent of input argument (measured in radians).
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
Note: FP128 device computations require compute capability >= 10.0.
Returns
tan x.
▶ __nv_fp128_tan( ±0 ) returns ±0.
▶ __nv_fp128_tan( ±∞ ) returns NaN.
▶ __nv_fp128_tan(NaN) returns NaN.
__device__ __float128 __nv_fp128_tanh(__float128 x)
Calculate tanh x, the hyperbolic tangent of the input argument.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
Note: FP128 device computations require compute capability >= 10.0.
Returns
▶ __nv_fp128_tanh( ±0 ) returns ±0.
▶ __nv_fp128_tanh( ±∞ ) returns ±1.
▶ __nv_fp128_tanh(NaN) returns NaN.
__device__ __float128 __nv_fp128_trunc(__float128 x)
Truncate input argument to the integral part.
Note: For accuracy information, see the CUDA C++ Programming Guide, Mathematical Functions
Appendix, Quad-Precision Floating-Point Functions section.
Note: FP128 device computations require compute capability >= 10.0.
10.1. Functions
357
Returns
Rounded x to the nearest integer value in floating-point format, that does not ex-
ceed x in magnitude.
▶ __nv_fp128_trunc( ±0 ) returns ±0.
▶ __nv_fp128_trunc( ±∞ ) returns ±∞.
▶ __nv_fp128_trunc(NaN) returns NaN.
358
Chapter 10. FP128 Quad Precision Mathematical Functions
Chapter 11. Type Casting Intrinsics
This section describes type casting 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__ float __double2float_rd(double x)
Convert a double to a float in round-down mode.
__device__ float __double2float_rn(double x)
Convert a double to a float in round-to-nearest-even mode.
__device__ float __double2float_ru(double x)
Convert a double to a float in round-up mode.
__device__ float __double2float_rz(double x)
Convert a double to a float in round-towards-zero mode.
__device__ int __double2hiint(double x)
Reinterpret high 32 bits in a double as a signed integer.
__device__ int __double2int_rd(double x)
Convert a double to a signed int in round-down mode.
__device__ int __double2int_rn(double x)
Convert a double to a signed int in round-to-nearest-even mode.
__device__ int __double2int_ru(double x)
Convert a double to a signed int in round-up mode.
__device__ int __double2int_rz(double x)
Convert a double to a signed int in round-towards-zero mode.
__device__ long long int __double2ll_rd(double x)
Convert a double to a signed 64-bit int in round-down mode.
__device__ long long int __double2ll_rn(double x)
Convert a double to a signed 64-bit int in round-to-nearest-even mode.
__device__ long long int __double2ll_ru(double x)
Convert a double to a signed 64-bit int in round-up mode.
__device__ long long int __double2ll_rz(double x)
Convert a double to a signed 64-bit int in round-towards-zero mode.
__device__ int __double2loint(double x)
Reinterpret low 32 bits in a double as a signed integer.
359
__device__ unsigned int __double2uint_rd(double x)
Convert a double to an unsigned int in round-down mode.
__device__ unsigned int __double2uint_rn(double x)
Convert a double to an unsigned int in round-to-nearest-even mode.
__device__ unsigned int __double2uint_ru(double x)
Convert a double to an unsigned int in round-up mode.
__device__ unsigned int __double2uint_rz(double x)
Convert a double to an unsigned int in round-towards-zero mode.
__device__ unsigned long long int __double2ull_rd(double x)
Convert a double to an unsigned 64-bit int in round-down mode.
__device__ unsigned long long int __double2ull_rn(double x)
Convert a double to an unsigned 64-bit int in round-to-nearest-even mode.
__device__ unsigned long long int __double2ull_ru(double x)
Convert a double to an unsigned 64-bit int in round-up mode.
__device__ unsigned long long int __double2ull_rz(double x)
Convert a double to an unsigned 64-bit int in round-towards-zero mode.
__device__ long long int __double_as_longlong(double x)
Reinterpret bits in a double as a 64-bit signed integer.
__device__ int __float2int_rd(float x)
Convert a float to a signed integer in round-down mode.
__device__ int __float2int_rn(float x)
Convert a float to a signed integer in round-to-nearest-even mode.
__device__ int __float2int_ru(float)
Convert a float to a signed integer in round-up mode.
__device__ int __float2int_rz(float x)
Convert a float to a signed integer in round-towards-zero mode.
__device__ long long int __float2ll_rd(float x)
Convert a float to a signed 64-bit integer in round-down mode.
__device__ long long int __float2ll_rn(float x)
Convert a float to a signed 64-bit integer in round-to-nearest-even mode.
__device__ long long int __float2ll_ru(float x)
Convert a float to a signed 64-bit integer in round-up mode.
__device__ long long int __float2ll_rz(float x)
Convert a float to a signed 64-bit integer in round-towards-zero mode.
__device__ unsigned int __float2uint_rd(float x)
Convert a float to an unsigned integer in round-down mode.
__device__ unsigned int __float2uint_rn(float x)
Convert a float to an unsigned integer in round-to-nearest-even mode.
__device__ unsigned int __float2uint_ru(float x)
Convert a float to an unsigned integer in round-up mode.
__device__ unsigned int __float2uint_rz(float x)
Convert a float to an unsigned integer in round-towards-zero mode.
360
Chapter 11. Type Casting Intrinsics
__device__ unsigned long long int __float2ull_rd(float x)
Convert a float to an unsigned 64-bit integer in round-down mode.
__device__ unsigned long long int __float2ull_rn(float x)
Convert a float to an unsigned 64-bit integer in round-to-nearest-even mode.
__device__ unsigned long long int __float2ull_ru(float x)
Convert a float to an unsigned 64-bit integer in round-up mode.
__device__ unsigned long long int __float2ull_rz(float x)
Convert a float to an unsigned 64-bit integer in round-towards-zero mode.
__device__ int __float_as_int(float x)
Reinterpret bits in a float as a signed integer.
__device__ unsigned int __float_as_uint(float x)
Reinterpret bits in a float as a unsigned integer.
__device__ double __hiloint2double(int hi, int lo)
Reinterpret high and low 32-bit integer values as a double.
__device__ double __int2double_rn(int x)
Convert a signed int to a double.
__device__ float __int2float_rd(int x)
Convert a signed integer to a float in round-down mode.
__device__ float __int2float_rn(int x)
Convert a signed integer to a float in round-to-nearest-even mode.
__device__ float __int2float_ru(int x)
Convert a signed integer to a float in round-up mode.
__device__ float __int2float_rz(int x)
Convert a signed integer to a float in round-towards-zero mode.
__device__ float __int_as_float(int x)
Reinterpret bits in an integer as a float.
__device__ double __ll2double_rd(long long int x)
Convert a signed 64-bit int to a double in round-down mode.
__device__ double __ll2double_rn(long long int x)
Convert a signed 64-bit int to a double in round-to-nearest-even mode.
__device__ double __ll2double_ru(long long int x)
Convert a signed 64-bit int to a double in round-up mode.
__device__ double __ll2double_rz(long long int x)
Convert a signed 64-bit int to a double in round-towards-zero mode.
__device__ float __ll2float_rd(long long int x)
Convert a signed integer to a float in round-down mode.
__device__ float __ll2float_rn(long long int x)
Convert a signed 64-bit integer to a float in round-to-nearest-even mode.
__device__ float __ll2float_ru(long long int x)
Convert a signed integer to a float in round-up mode.
__device__ float __ll2float_rz(long long int x)
Convert a signed integer to a float in round-towards-zero mode.
361
__device__ double __longlong_as_double(long long int x)
Reinterpret bits in a 64-bit signed integer as a double.
__device__ double __uint2double_rn(unsigned int x)
Convert an unsigned int to a double.
__device__ float __uint2float_rd(unsigned int x)
Convert an unsigned integer to a float in round-down mode.
__device__ float __uint2float_rn(unsigned int x)
Convert an unsigned integer to a float in round-to-nearest-even mode.
__device__ float __uint2float_ru(unsigned int x)
Convert an unsigned integer to a float in round-up mode.
__device__ float __uint2float_rz(unsigned int x)
Convert an unsigned integer to a float in round-towards-zero mode.
__device__ float __uint_as_float(unsigned int x)
Reinterpret bits in an unsigned integer as a float.
__device__ double __ull2double_rd(unsigned long long int x)
Convert an unsigned 64-bit int to a double in round-down mode.
__device__ double __ull2double_rn(unsigned long long int x)
Convert an unsigned 64-bit int to a double in round-to-nearest-even mode.
__device__ double __ull2double_ru(unsigned long long int x)
Convert an unsigned 64-bit int to a double in round-up mode.
__device__ double __ull2double_rz(unsigned long long int x)
Convert an unsigned 64-bit int to a double in round-towards-zero mode.
__device__ float __ull2float_rd(unsigned long long int x)
Convert an unsigned integer to a float in round-down mode.
__device__ float __ull2float_rn(unsigned long long int x)
Convert an unsigned integer to a float in round-to-nearest-even mode.
__device__ float __ull2float_ru(unsigned long long int x)
Convert an unsigned integer to a float in round-up mode.
__device__ float __ull2float_rz(unsigned long long int x)
Convert an unsigned integer to a float in round-towards-zero mode.
11.1. Functions
__device__ float __double2float_rd(double x)
Convert a double to a float in round-down mode.
Convert the double-precision floating-point value x to a single-precision floating-point value in
round-down (to negative infinity) mode.
Returns
Returns converted value.
362
Chapter 11. Type Casting Intrinsics
__device__ float __double2float_rn(double x)
Convert a double to a float in round-to-nearest-even mode.
Convert the double-precision floating-point value x to a single-precision floating-point value in
round-to-nearest-even mode.
Returns
Returns converted value.
__device__ float __double2float_ru(double x)
Convert a double to a float in round-up mode.
Convert the double-precision floating-point value x to a single-precision floating-point value in
round-up (to positive infinity) mode.
Returns
Returns converted value.
__device__ float __double2float_rz(double x)
Convert a double to a float in round-towards-zero mode.
Convert the double-precision floating-point value x to a single-precision floating-point value in
round-towards-zero mode.
Returns
Returns converted value.
__device__ int __double2hiint(double x)
Reinterpret high 32 bits in a double as a signed integer.
Reinterpret the high 32 bits in the double-precision floating-point value x as a signed integer.
Returns
Returns reinterpreted value.
__device__ int __double2int_rd(double x)
Convert a double to a signed int in round-down mode.
Convert the double-precision floating-point value x to a signed integer value in round-down (to
negative infinity) mode.
Note: When the floating-point input rounded to integral is outside the range of the return type,
the behavior is undefined.
Returns
Returns converted value.
__device__ int __double2int_rn(double x)
Convert a double to a signed int in round-to-nearest-even mode.
Convert the double-precision floating-point value x to a signed integer value in round-to-nearest-
even mode.
Note: When the floating-point input rounded to integral is outside the range of the return type,
the behavior is undefined.
11.1. Functions
363
Returns
Returns converted value.
__device__ int __double2int_ru(double x)
Convert a double to a signed int in round-up mode.
Convert the double-precision floating-point value x to a signed integer value in round-up (to pos-
itive infinity) mode.
Note: When the floating-point input rounded to integral is outside the range of the return type,
the behavior is undefined.
Returns
Returns converted value.
__device__ int __double2int_rz(double x)
Convert a double to a signed int in round-towards-zero mode.
Convert the double-precision floating-point value x to a signed integer value in round-towards-
zero mode.
Note: When the floating-point input rounded to integral is outside the range of the return type,
the behavior is undefined.
Returns
Returns converted value.
__device__ long long int __double2ll_rd(double x)
Convert a double to a signed 64-bit int in round-down mode.
Convert the double-precision floating-point value x to a signed 64-bit integer value in round-
down (to negative infinity) mode.
Note: When the floating-point input rounded to integral is outside the range of the return type,
the behavior is undefined.
Returns
Returns converted value.
__device__ long long int __double2ll_rn(double x)
Convert a double to a signed 64-bit int in round-to-nearest-even mode.
Convert the double-precision floating-point value x to a signed 64-bit integer value in round-to-
nearest-even mode.
Note: When the floating-point input rounded to integral is outside the range of the return type,
the behavior is undefined.
364
Chapter 11. Type Casting Intrinsics
Returns
Returns converted value.
__device__ long long int __double2ll_ru(double x)
Convert a double to a signed 64-bit int in round-up mode.
Convert the double-precision floating-point value x to a signed 64-bit integer value in round-up
(to positive infinity) mode.
Note: When the floating-point input rounded to integral is outside the range of the return type,
the behavior is undefined.
Returns
Returns converted value.
__device__ long long int __double2ll_rz(double x)
Convert a double to a signed 64-bit int in round-towards-zero mode.
Convert the double-precision floating-point value x to a signed 64-bit integer value in round-
towards-zero mode.
Note: When the floating-point input rounded to integral is outside the range of the return type,
the behavior is undefined.
Returns
Returns converted value.
__device__ int __double2loint(double x)
Reinterpret low 32 bits in a double as a signed integer.
Reinterpret the low 32 bits in the double-precision floating-point value x as a signed integer.
Returns
Returns reinterpreted value.
__device__ unsigned int __double2uint_rd(double x)
Convert a double to an unsigned int in round-down mode.
Convert the double-precision floating-point value x to an unsigned integer value in round-down
(to negative infinity) mode.
Note: When the floating-point input rounded to integral is outside the range of the return type,
the behavior is undefined.
Returns
Returns converted value.
__device__ unsigned int __double2uint_rn(double x)
Convert a double to an unsigned int in round-to-nearest-even mode.
Convert the double-precision floating-point value x to an unsigned integer value in round-to-
nearest-even mode.
11.1. Functions
365
Note: When the floating-point input rounded to integral is outside the range of the return type,
the behavior is undefined.
Returns
Returns converted value.
__device__ unsigned int __double2uint_ru(double x)
Convert a double to an unsigned int in round-up mode.
Convert the double-precision floating-point value x to an unsigned integer value in round-up (to
positive infinity) mode.
Note: When the floating-point input rounded to integral is outside the range of the return type,
the behavior is undefined.
Returns
Returns converted value.
__device__ unsigned int __double2uint_rz(double x)
Convert a double to an unsigned int in round-towards-zero mode.
Convert the double-precision floating-point value x to an unsigned integer value in round-
towards-zero mode.
Note: When the floating-point input rounded to integral is outside the range of the return type,
the behavior is undefined.
Returns
Returns converted value.
__device__ unsigned long long int __double2ull_rd(double x)
Convert a double to an unsigned 64-bit int in round-down mode.
Convert the double-precision floating-point value x to an unsigned 64-bit integer value in round-
down (to negative infinity) mode.
Note: When the floating-point input rounded to integral is outside the range of the return type,
the behavior is undefined.
Returns
Returns converted value.
__device__ unsigned long long int __double2ull_rn(double x)
Convert a double to an unsigned 64-bit int in round-to-nearest-even mode.
Convert the double-precision floating-point value x to an unsigned 64-bit integer value in round-
to-nearest-even mode.
366
Chapter 11. Type Casting Intrinsics
Note: When the floating-point input rounded to integral is outside the range of the return type,
the behavior is undefined.
Returns
Returns converted value.
__device__ unsigned long long int __double2ull_ru(double x)
Convert a double to an unsigned 64-bit int in round-up mode.
Convert the double-precision floating-point value x to an unsigned 64-bit integer value in round-
up (to positive infinity) mode.
Note: When the floating-point input rounded to integral is outside the range of the return type,
the behavior is undefined.
Returns
Returns converted value.
__device__ unsigned long long int __double2ull_rz(double x)
Convert a double to an unsigned 64-bit int in round-towards-zero mode.
Convert the double-precision floating-point value x to an unsigned 64-bit integer value in round-
towards-zero mode.
Note: When the floating-point input rounded to integral is outside the range of the return type,
the behavior is undefined.
Returns
Returns converted value.
__device__ long long int __double_as_longlong(double x)
Reinterpret bits in a double as a 64-bit signed integer.
Reinterpret the bits in the double-precision floating-point value x as a signed 64-bit integer.
Returns
Returns reinterpreted value.
__device__ int __float2int_rd(float x)
Convert a float to a signed integer in round-down mode.
Convert the single-precision floating-point value x to a signed integer in round-down (to negative
infinity) mode.
Note: When the floating-point input rounded to integral is outside the range of the return type,
the behavior is undefined.
Returns
Returns converted value.
11.1. Functions
367
__device__ int __float2int_rn(float x)
Convert a float to a signed integer in round-to-nearest-even mode.
Convert the single-precision floating-point value x to a signed integer in round-to-nearest-even
mode.
Note: When the floating-point input rounded to integral is outside the range of the return type,
the behavior is undefined.
Returns
Returns converted value.
__device__ int __float2int_ru(float)
Convert a float to a signed integer in round-up mode.
Convert the single-precision floating-point value x to a signed integer in round-up (to positive
infinity) mode.
Note: When the floating-point input rounded to integral is outside the range of the return type,
the behavior is undefined.
Returns
Returns converted value.
__device__ int __float2int_rz(float x)
Convert a float to a signed integer in round-towards-zero mode.
Convert the single-precision floating-point value x to a signed integer in round-towards-zero
mode.
Note: When the floating-point input rounded to integral is outside the range of the return type,
the behavior is undefined.
Returns
Returns converted value.
__device__ long long int __float2ll_rd(float x)
Convert a float to a signed 64-bit integer in round-down mode.
Convert the single-precision floating-point value x to a signed 64-bit integer in round-down (to
negative infinity) mode.
Note: When the floating-point input rounded to integral is outside the range of the return type,
the behavior is undefined.
Returns
Returns converted value.
368
Chapter 11. Type Casting Intrinsics
__device__ long long int __float2ll_rn(float x)
Convert a float to a signed 64-bit integer in round-to-nearest-even mode.
Convert the single-precision floating-point value x to a signed 64-bit integer in round-to-nearest-
even mode.
Note: When the floating-point input rounded to integral is outside the range of the return type,
the behavior is undefined.
Returns
Returns converted value.
__device__ long long int __float2ll_ru(float x)
Convert a float to a signed 64-bit integer in round-up mode.
Convert the single-precision floating-point value x to a signed 64-bit integer in round-up (to pos-
itive infinity) mode.
Note: When the floating-point input rounded to integral is outside the range of the return type,
the behavior is undefined.
Returns
Returns converted value.
__device__ long long int __float2ll_rz(float x)
Convert a float to a signed 64-bit integer in round-towards-zero mode.
Convert the single-precision floating-point value x to a signed 64-bit integer in round-towards-
zero mode.
Note: When the floating-point input rounded to integral is outside the range of the return type,
the behavior is undefined.
Returns
Returns converted value.
__device__ unsigned int __float2uint_rd(float x)
Convert a float to an unsigned integer in round-down mode.
Convert the single-precision floating-point value x to an unsigned integer in round-down (to neg-
ative infinity) mode.
Note: When the floating-point input rounded to integral is outside the range of the return type,
the behavior is undefined.
Returns
Returns converted value.
11.1. Functions
369
__device__ unsigned int __float2uint_rn(float x)
Convert a float to an unsigned integer in round-to-nearest-even mode.
Convert the single-precision floating-point value x to an unsigned integer in round-to-nearest-
even mode.
Note: When the floating-point input rounded to integral is outside the range of the return type,
the behavior is undefined.
Returns
Returns converted value.
__device__ unsigned int __float2uint_ru(float x)
Convert a float to an unsigned integer in round-up mode.
Convert the single-precision floating-point value x to an unsigned integer in round-up (to positive
infinity) mode.
Note: When the floating-point input rounded to integral is outside the range of the return type,
the behavior is undefined.
Returns
Returns converted value.
__device__ unsigned int __float2uint_rz(float x)
Convert a float to an unsigned integer in round-towards-zero mode.
Convert the single-precision floating-point value x to an unsigned integer in round-towards-zero
mode.
Note: When the floating-point input rounded to integral is outside the range of the return type,
the behavior is undefined.
Returns
Returns converted value.
__device__ unsigned long long int __float2ull_rd(float x)
Convert a float to an unsigned 64-bit integer in round-down mode.
Convert the single-precision floating-point value x to an unsigned 64-bit integer in round-down
(to negative infinity) mode.
Note: When the floating-point input rounded to integral is outside the range of the return type,
the behavior is undefined.
Returns
Returns converted value.
370
Chapter 11. Type Casting Intrinsics
__device__ unsigned long long int __float2ull_rn(float x)
Convert a float to an unsigned 64-bit integer in round-to-nearest-even mode.
Convert the single-precision floating-point value x to an unsigned 64-bit integer in round-to-
nearest-even mode.
Note: When the floating-point input rounded to integral is outside the range of the return type,
the behavior is undefined.
Returns
Returns converted value.
__device__ unsigned long long int __float2ull_ru(float x)
Convert a float to an unsigned 64-bit integer in round-up mode.
Convert the single-precision floating-point value x to an unsigned 64-bit integer in round-up (to
positive infinity) mode.
Note: When the floating-point input rounded to integral is outside the range of the return type,
the behavior is undefined.
Returns
Returns converted value.
__device__ unsigned long long int __float2ull_rz(float x)
Convert a float to an unsigned 64-bit integer in round-towards-zero mode.
Convert the single-precision floating-point value x to an unsigned 64-bit integer in round-
towards-zero mode.
Note: When the floating-point input rounded to integral is outside the range of the return type,
the behavior is undefined.
Returns
Returns converted value.
__device__ int __float_as_int(float x)
Reinterpret bits in a float as a signed integer.
Reinterpret the bits in the single-precision floating-point value x as a signed integer.
Returns
Returns reinterpreted value.
__device__ unsigned int __float_as_uint(float x)
Reinterpret bits in a float as a unsigned integer.
Reinterpret the bits in the single-precision floating-point value x as a unsigned integer.
Returns
Returns reinterpreted value.
11.1. Functions
371
__device__ double __hiloint2double(int hi, int lo)
Reinterpret high and low 32-bit integer values as a double.
Reinterpret the integer value of hi as the high 32 bits of a double-precision floating-point value
and the integer value of lo as the low 32 bits of the same double-precision floating-point value.
Returns
Returns reinterpreted value.
__device__ double __int2double_rn(int x)
Convert a signed int to a double.
Convert the signed integer value x to a double-precision floating-point value.
Returns
Returns converted value.
__device__ float __int2float_rd(int x)
Convert a signed integer to a float in round-down mode.
Convert the signed integer value x to a single-precision floating-point value in round-down (to
negative infinity) mode.
Returns
Returns converted value.
__device__ float __int2float_rn(int x)
Convert a signed integer to a float in round-to-nearest-even mode.
Convert the signed integer value x to a single-precision floating-point value in round-to-nearest-
even mode.
Returns
Returns converted value.
__device__ float __int2float_ru(int x)
Convert a signed integer to a float in round-up mode.
Convert the signed integer value x to a single-precision floating-point value in round-up (to pos-
itive infinity) mode.
Returns
Returns converted value.
__device__ float __int2float_rz(int x)
Convert a signed integer to a float in round-towards-zero mode.
Convert the signed integer value x to a single-precision floating-point value in round-towards-
zero mode.
Returns
Returns converted value.
__device__ float __int_as_float(int x)
Reinterpret bits in an integer as a float.
Reinterpret the bits in the signed integer value x as a single-precision floating-point value.
Returns
Returns reinterpreted value.
372
Chapter 11. Type Casting Intrinsics
__device__ double __ll2double_rd(long long int x)
Convert a signed 64-bit int to a double in round-down mode.
Convert the signed 64-bit integer value x to a double-precision floating-point value in round-
down (to negative infinity) mode.
Returns
Returns converted value.
__device__ double __ll2double_rn(long long int x)
Convert a signed 64-bit int to a double in round-to-nearest-even mode.
Convert the signed 64-bit integer value x to a double-precision floating-point value in round-to-
nearest-even mode.
Returns
Returns converted value.
__device__ double __ll2double_ru(long long int x)
Convert a signed 64-bit int to a double in round-up mode.
Convert the signed 64-bit integer value x to a double-precision floating-point value in round-up
(to positive infinity) mode.
Returns
Returns converted value.
__device__ double __ll2double_rz(long long int x)
Convert a signed 64-bit int to a double in round-towards-zero mode.
Convert the signed 64-bit integer value x to a double-precision floating-point value in round-
towards-zero mode.
Returns
Returns converted value.
__device__ float __ll2float_rd(long long int x)
Convert a signed integer to a float in round-down mode.
Convert the signed integer value x to a single-precision floating-point value in round-down (to
negative infinity) mode.
Returns
Returns converted value.
__device__ float __ll2float_rn(long long int x)
Convert a signed 64-bit integer to a float in round-to-nearest-even mode.
Convert the signed 64-bit integer value x to a single-precision floating-point value in round-to-
nearest-even mode.
Returns
Returns converted value.
__device__ float __ll2float_ru(long long int x)
Convert a signed integer to a float in round-up mode.
Convert the signed integer value x to a single-precision floating-point value in round-up (to pos-
itive infinity) mode.
Returns
Returns converted value.
11.1. Functions
373
__device__ float __ll2float_rz(long long int x)
Convert a signed integer to a float in round-towards-zero mode.
Convert the signed integer value x to a single-precision floating-point value in round-towards-
zero mode.
Returns
Returns converted value.
__device__ double __longlong_as_double(long long int x)
Reinterpret bits in a 64-bit signed integer as a double.
Reinterpret the bits in the 64-bit signed integer value x as a double-precision floating-point value.
Returns
Returns reinterpreted value.
__device__ double __uint2double_rn(unsigned int x)
Convert an unsigned int to a double.
Convert the unsigned integer value x to a double-precision floating-point value.
Returns
Returns converted value.
__device__ float __uint2float_rd(unsigned int x)
Convert an unsigned integer to a float in round-down mode.
Convert the unsigned integer value x to a single-precision floating-point value in round-down (to
negative infinity) mode.
Returns
Returns converted value.
__device__ float __uint2float_rn(unsigned int x)
Convert an unsigned integer to a float in round-to-nearest-even mode.
Convert the unsigned integer value x to a single-precision floating-point value in round-to-
nearest-even mode.
Returns
Returns converted value.
__device__ float __uint2float_ru(unsigned int x)
Convert an unsigned integer to a float in round-up mode.
Convert the unsigned integer value x to a single-precision floating-point value in round-up (to
positive infinity) mode.
Returns
Returns converted value.
__device__ float __uint2float_rz(unsigned int x)
Convert an unsigned integer to a float in round-towards-zero mode.
Convert the unsigned integer value x to a single-precision floating-point value in round-towards-
zero mode.
Returns
Returns converted value.
374
Chapter 11. Type Casting Intrinsics
__device__ float __uint_as_float(unsigned int x)
Reinterpret bits in an unsigned integer as a float.
Reinterpret the bits in the unsigned integer value x as a single-precision floating-point value.
Returns
Returns reinterpreted value.
__device__ double __ull2double_rd(unsigned long long int x)
Convert an unsigned 64-bit int to a double in round-down mode.
Convert the unsigned 64-bit integer value x to a double-precision floating-point value in round-
down (to negative infinity) mode.
Returns
Returns converted value.
__device__ double __ull2double_rn(unsigned long long int x)
Convert an unsigned 64-bit int to a double in round-to-nearest-even mode.
Convert the unsigned 64-bit integer value x to a double-precision floating-point value in round-
to-nearest-even mode.
Returns
Returns converted value.
__device__ double __ull2double_ru(unsigned long long int x)
Convert an unsigned 64-bit int to a double in round-up mode.
Convert the unsigned 64-bit integer value x to a double-precision floating-point value in round-up
(to positive infinity) mode.
Returns
Returns converted value.
__device__ double __ull2double_rz(unsigned long long int x)
Convert an unsigned 64-bit int to a double in round-towards-zero mode.
Convert the unsigned 64-bit integer value x to a double-precision floating-point value in round-
towards-zero mode.
Returns
Returns converted value.
__device__ float __ull2float_rd(unsigned long long int x)
Convert an unsigned integer to a float in round-down mode.
Convert the unsigned integer value x to a single-precision floating-point value in round-down (to
negative infinity) mode.
Returns
Returns converted value.
__device__ float __ull2float_rn(unsigned long long int x)
Convert an unsigned integer to a float in round-to-nearest-even mode.
Convert the unsigned integer value x to a single-precision floating-point value in round-to-
nearest-even mode.
Returns
Returns converted value.
11.1. Functions
375
__device__ float __ull2float_ru(unsigned long long int x)
Convert an unsigned integer to a float in round-up mode.
Convert the unsigned integer value x to a single-precision floating-point value in round-up (to
positive infinity) mode.
Returns
Returns converted value.
__device__ float __ull2float_rz(unsigned long long int x)
Convert an unsigned integer to a float in round-towards-zero mode.
Convert the unsigned integer value x to a single-precision floating-point value in round-towards-
zero mode.
Returns
Returns converted value.
376
Chapter 11. Type Casting Intrinsics
Chapter 12. Integer Mathematical
Functions
This section describes integer mathematical functions.
To use these functions, you do not need to include any additional header file in your program.
Functions
__device__ long int abs(long int a)
Calculate the absolute value of the input long int argument.
__device__ int abs(int a)
Calculate the absolute value of the input int argument.
__device__ long long int abs(long long int a)
Calculate the absolute value of the input long long int argument.
__device__ long int labs(long int a)
Calculate the absolute value of the input long int argument.
__device__ long long int llabs(long long int a)
Calculate the absolute value of the input long long int argument.
__device__ long long int llmax(const long long int a, const long long int b)
Calculate the maximum value of the input long long int arguments.
__device__ long long int llmin(const long long int a, const long long int b)
Calculate the minimum value of the input long long int arguments.
__device__ unsigned long int max(const long int a, const unsigned long int b)
Calculate the maximum value of the input long int and unsigned long int arguments.
__device__ unsigned long long int max(const unsigned long long int a, const unsigned long long int
b)
Calculate the maximum value of the input unsigned long long int arguments.
__device__ unsigned int max(const unsigned int a, const int b)
Calculate the maximum value of the input unsigned int and int arguments.
__device__ unsigned long long int max(const long long int a, const unsigned long long int b)
Calculate the maximum value of the input long long int and unsigned long long int argu-
ments.
__device__ unsigned long int max(const unsigned long int a, const unsigned long int b)
Calculate the maximum value of the input unsigned long int arguments.
377
__device__ long long int max(const long long int a, const long long int b)
Calculate the maximum value of the input long long int arguments.
__device__ unsigned long long int max(const unsigned long long int a, const long long int b)
Calculate the maximum value of the input unsigned long long int and long long int argu-
ments.
__device__ unsigned long int max(const unsigned long int a, const long int b)
Calculate the maximum value of the input unsigned long int and long int arguments.
__device__ long int max(const long int a, const long int b)
Calculate the maximum value of the input long int arguments.
__device__ int max(const int a, const int b)
Calculate the maximum value of the input int arguments.
__device__ unsigned int max(const unsigned int a, const unsigned int b)
Calculate the maximum value of the input unsigned int arguments.
__device__ unsigned int max(const int a, const unsigned int b)
Calculate the maximum value of the input int and unsigned int arguments.
__device__ unsigned long int min(const long int a, const unsigned long int b)
Calculate the minimum value of the input long int and unsigned long int arguments.
__device__ unsigned long long int min(const unsigned long long int a, const unsigned long long int
b)
Calculate the minimum value of the input unsigned long long int arguments.
__device__ unsigned long long int min(const unsigned long long int a, const long long int b)
Calculate the minimum value of the input unsigned long long int and long long int argu-
ments.
__device__ int min(const int a, const int b)
Calculate the minimum value of the input int arguments.
__device__ unsigned int min(const unsigned int a, const int b)
Calculate the minimum value of the input unsigned int and int arguments.
__device__ unsigned long long int min(const long long int a, const unsigned long long int b)
Calculate the minimum value of the input long long int and unsigned long long int argu-
ments.
__device__ long long int min(const long long int a, const long long int b)
Calculate the minimum value of the input long long int arguments.
__device__ unsigned int min(const int a, const unsigned int b)
Calculate the minimum value of the input int and unsigned int arguments.
__device__ long int min(const long int a, const long int b)
Calculate the minimum value of the input long int arguments.
__device__ unsigned int min(const unsigned int a, const unsigned int b)
Calculate the minimum value of the input unsigned int arguments.
__device__ unsigned long int min(const unsigned long int a, const long int b)
Calculate the minimum value of the input unsigned long int and long int arguments.
__device__ unsigned long int min(const unsigned long int a, const unsigned long int b)
Calculate the minimum value of the input unsigned long int arguments.
378
Chapter 12. Integer Mathematical Functions
__device__ unsigned long long int ullmax(const unsigned long long int a, const unsigned long long
int b)
Calculate the maximum value of the input unsigned long long int arguments.
__device__ unsigned long long int ullmin(const unsigned long long int a, const unsigned long long
int b)
Calculate the minimum value of the input unsigned long long int arguments.
__device__ unsigned int umax(const unsigned int a, const unsigned int b)
Calculate the maximum value of the input unsigned int arguments.
__device__ unsigned int umin(const unsigned int a, const unsigned int b)
Calculate the minimum value of the input unsigned int arguments.
12.1. Functions
__device__ long int abs(long int a)
Calculate the absolute value of the input long int argument.
Calculate the absolute value of the input argument a.
Returns
Returns the absolute value of the input argument.
▶ abs(LONG_MIN) is Undefined
__device__ int abs(int a)
Calculate the absolute value of the input int argument.
Calculate the absolute value of the input argument a.
Returns
Returns the absolute value of the input argument.
▶ abs(INT_MIN) is Undefined
__device__ long long int abs(long long int a)
Calculate the absolute value of the input long long int argument.
Calculate the absolute value of the input argument a.
Returns
Returns the absolute value of the input argument.
▶ abs(LLONG_MIN) is Undefined
__device__ long int labs(long int a)
Calculate the absolute value of the input long int argument.
Calculate the absolute value of the input argument a.
Returns
Returns the absolute value of the input argument.
▶ labs(LONG_MIN) is Undefined
__device__ long long int llabs(long long int a)
Calculate the absolute value of the input long long int argument.
Calculate the absolute value of the input argument a.
12.1. Functions
379
Returns
Returns the absolute value of the input argument.
▶ llabs(LLONG_MIN) is Undefined
__device__ long long int llmax(const long long int a, const long long int b)
Calculate the maximum value of the input long long int arguments.
Calculate the maximum value of the arguments a and b.
__device__ long long int llmin(const long long int a, const long long int b)
Calculate the minimum value of the input long long int arguments.
Calculate the minimum value of the arguments a and b.
__device__ unsigned long int max(const long int a, const unsigned long int b)
Calculate the maximum value of the input long int and unsigned long int arguments.
Calculate the maximum value of the arguments a and b, perform integer promotion first.
__device__ unsigned long long int max(const unsigned long long int a, const unsigned long long int
b)
Calculate the maximum value of the input unsigned long long int arguments.
Calculate the maximum value of the arguments a and b.
__device__ unsigned int max(const unsigned int a, const int b)
Calculate the maximum value of the input unsigned int and int arguments.
Calculate the maximum value of the arguments a and b, perform integer promotion first.
__device__ unsigned long long int max(const long long int a, const unsigned long long int b)
Calculate the maximum value of the input long long int and unsigned long long int argu-
ments.
Calculate the maximum value of the arguments a and b, perform integer promotion first.
__device__ unsigned long int max(const unsigned long int a, const unsigned long int b)
Calculate the maximum value of the input unsigned long int arguments.
Calculate the maximum value of the arguments a and b.
__device__ long long int max(const long long int a, const long long int b)
Calculate the maximum value of the input long long int arguments.
Calculate the maximum value of the arguments a and b.
__device__ unsigned long long int max(const unsigned long long int a, const long long int b)
Calculate the maximum value of the input unsigned long long int and long long int argu-
ments.
Calculate the maximum value of the arguments a and b, perform integer promotion first.
__device__ unsigned long int max(const unsigned long int a, const long int b)
Calculate the maximum value of the input unsigned long int and long int arguments.
Calculate the maximum value of the arguments a and b, perform integer promotion first.
__device__ long int max(const long int a, const long int b)
Calculate the maximum value of the input long int arguments.
Calculate the maximum value of the arguments a and b.
380
Chapter 12. Integer Mathematical Functions
__device__ int max(const int a, const int b)
Calculate the maximum value of the input int arguments.
Calculate the maximum value of the arguments a and b.
__device__ unsigned int max(const unsigned int a, const unsigned int b)
Calculate the maximum value of the input unsigned int arguments.
Calculate the maximum value of the arguments a and b.
__device__ unsigned int max(const int a, const unsigned int b)
Calculate the maximum value of the input int and unsigned int arguments.
Calculate the maximum value of the arguments a and b, perform integer promotion first.
__device__ unsigned long int min(const long int a, const unsigned long int b)
Calculate the minimum value of the input long int and unsigned long int arguments.
Calculate the minimum value of the arguments a and b, perform integer promotion first.
__device__ unsigned long long int min(const unsigned long long int a, const unsigned long long int
b)
Calculate the minimum value of the input unsigned long long int arguments.
Calculate the minimum value of the arguments a and b.
__device__ unsigned long long int min(const unsigned long long int a, const long long int b)
Calculate the minimum value of the input unsigned long long int and long long int argu-
ments.
Calculate the minimum value of the arguments a and b, perform integer promotion first.
__device__ int min(const int a, const int b)
Calculate the minimum value of the input int arguments.
Calculate the minimum value of the arguments a and b.
__device__ unsigned int min(const unsigned int a, const int b)
Calculate the minimum value of the input unsigned int and int arguments.
Calculate the minimum value of the arguments a and b, perform integer promotion first.
__device__ unsigned long long int min(const long long int a, const unsigned long long int b)
Calculate the minimum value of the input long long int and unsigned long long int argu-
ments.
Calculate the minimum value of the arguments a and b, perform integer promotion first.
__device__ long long int min(const long long int a, const long long int b)
Calculate the minimum value of the input long long int arguments.
Calculate the minimum value of the arguments a and b.
__device__ unsigned int min(const int a, const unsigned int b)
Calculate the minimum value of the input int and unsigned int arguments.
Calculate the minimum value of the arguments a and b, perform integer promotion first.
__device__ long int min(const long int a, const long int b)
Calculate the minimum value of the input long int arguments.
Calculate the minimum value of the arguments a and b.
12.1. Functions
381
__device__ unsigned int min(const unsigned int a, const unsigned int b)
Calculate the minimum value of the input unsigned int arguments.
Calculate the minimum value of the arguments a and b.
__device__ unsigned long int min(const unsigned long int a, const long int b)
Calculate the minimum value of the input unsigned long int and long int arguments.
Calculate the minimum value of the arguments a and b, perform integer promotion first.
__device__ unsigned long int min(const unsigned long int a, const unsigned long int b)
Calculate the minimum value of the input unsigned long int arguments.
Calculate the minimum value of the arguments a and b.
__device__ unsigned long long int ullmax(const unsigned long long int a, const unsigned long long
int b)
Calculate the maximum value of the input unsigned long long int arguments.
Calculate the maximum value of the arguments a and b.
__device__ unsigned long long int ullmin(const unsigned long long int a, const unsigned long long
int b)
Calculate the minimum value of the input unsigned long long int arguments.
Calculate the minimum value of the arguments a and b.
__device__ unsigned int umax(const unsigned int a, const unsigned int b)
Calculate the maximum value of the input unsigned int arguments.
Calculate the maximum value of the arguments a and b.
__device__ unsigned int umin(const unsigned int a, const unsigned int b)
Calculate the minimum value of the input unsigned int arguments.
Calculate the minimum value of the arguments a and b.
382
Chapter 12. Integer Mathematical Functions
Chapter 13. Integer Intrinsics
This section describes integer intrinsic functions.
All of these functions are supported in device code. For some of the functions, host-specific imple-
mentations are also provided. For example, see __nv_bswap16(). To use these functions, you do not
need to include any additional header file in your program.
Functions
__device__ unsigned int __brev(unsigned int x)
Reverse the bit order of a 32-bit unsigned integer.
__device__ unsigned long long int __brevll(unsigned long long int x)
Reverse the bit order of a 64-bit unsigned integer.
__device__ unsigned int __byte_perm(unsigned int x, unsigned int y, unsigned int s)
Return selected bytes from two 32-bit unsigned integers.
__device__ int __clz(int x)
Return the number of consecutive high-order zero bits in a 32-bit integer.
__device__ int __clzll(long long int x)
Count the number of consecutive high-order zero bits in a 64-bit integer.
__device__ int __dp2a_hi(int srcA, int srcB, int c)
Two-way signed int16 by int8 dot product with int32 accumulate, taking the upper half of
the second input.
__device__ unsigned int __dp2a_hi(unsigned int srcA, unsigned int srcB, unsigned int c)
Two-way unsigned int16 by int8 dot product with unsigned int32 accumulate, taking the
upper half of the second input.
__device__ unsigned int __dp2a_hi(ushort2 srcA, uchar4 srcB, unsigned int c)
Two-way unsigned int16 by int8 dot product with unsigned int32 accumulate, taking the
upper half of the second input.
__device__ int __dp2a_hi(short2 srcA, char4 srcB, int c)
Two-way signed int16 by int8 dot product with int32 accumulate, taking the upper half of
the second input.
__device__ unsigned int __dp2a_lo(ushort2 srcA, uchar4 srcB, unsigned int c)
Two-way unsigned int16 by int8 dot product with unsigned int32 accumulate, taking the
lower half of the second input.
__device__ int __dp2a_lo(short2 srcA, char4 srcB, int c)
Two-way signed int16 by int8 dot product with int32 accumulate, taking the lower half of
the second input.
383
__device__ unsigned int __dp2a_lo(unsigned int srcA, unsigned int srcB, unsigned int c)
Two-way unsigned int16 by int8 dot product with unsigned int32 accumulate, taking the
lower half of the second input.
__device__ int __dp2a_lo(int srcA, int srcB, int c)
Two-way signed int16 by int8 dot product with int32 accumulate, taking the lower half of
the second input.
__device__ unsigned int __dp4a(uchar4 srcA, uchar4 srcB, unsigned int c)
Four-way unsigned int8 dot product with unsigned int32 accumulate.
__device__ unsigned int __dp4a(unsigned int srcA, unsigned int srcB, unsigned int c)
Four-way unsigned int8 dot product with unsigned int32 accumulate.
__device__ int __dp4a(int srcA, int srcB, int c)
Four-way signed int8 dot product with int32 accumulate.
__device__ int __dp4a(char4 srcA, char4 srcB, int c)
Four-way signed int8 dot product with int32 accumulate.
__device__ int __ffs(int x)
Find the position of the least significant bit set to 1 in a 32-bit integer.
__device__ int __ffsll(long long int x)
Find the position of the least significant bit set to 1 in a 64-bit integer.
__device__ unsigned __fns(unsigned mask, unsigned base, int offset)
Find the position of the n-th set to 1 bit in a 32-bit integer.
__device__ unsigned int __funnelshift_l(unsigned int lo, unsigned int hi, unsigned int shift)
Concatenate hi : lo , shift left by shift & 31 bits, return the most significant 32 bits.
__device__ unsigned int __funnelshift_lc(unsigned int lo, unsigned int hi, unsigned int shift)
Concatenate hi : lo , shift left by min( shift , 32) bits, return the most significant 32 bits.
__device__ unsigned int __funnelshift_r(unsigned int lo, unsigned int hi, unsigned int shift)
Concatenate hi : lo , shift right by shift & 31 bits, return the least significant 32 bits.
__device__ unsigned int __funnelshift_rc(unsigned int lo, unsigned int hi, unsigned int shift)
Concatenate hi : lo , shift right by min( shift , 32) bits, return the least significant 32 bits.
__device__ int __hadd(int x, int y)
Compute average of signed input arguments, avoiding overflow in the intermediate sum.
__device__ int __mul24(int x, int y)
Calculate the least significant 32 bits of the product of the least significant 24 bits of two inte-
gers.
__device__ long long int __mul64hi(long long int x, long long int y)
Calculate the most significant 64 bits of the product of the two 64-bit integers.
__device__ int __mulhi(int x, int y)
Calculate the most significant 32 bits of the product of the two 32-bit integers.
__host__ __device__ unsigned short __nv_bswap16(unsigned short x)
Reverse the order of bytes of the 16-bit unsigned integer.
__host__ __device__ unsigned int __nv_bswap32(unsigned int x)
Reverse the order of bytes of the 32-bit unsigned integer.
__host__ __device__ unsigned long long __nv_bswap64(unsigned long long x)
Reverse the order of bytes of the 64-bit unsigned integer.
384
Chapter 13. Integer Intrinsics
__device__ int __popc(unsigned int x)
Count the number of bits that are set to 1 in a 32-bit integer.
__device__ int __popcll(unsigned long long int x)
Count the number of bits that are set to 1 in a 64-bit integer.
__device__ int __rhadd(int x, int y)
Compute rounded average of signed input arguments, avoiding overflow in the intermediate sum.
__device__ unsigned int __sad(int x, int y, unsigned int z)
Calculate |x - y| + z , the sum of absolute difference.
__device__ unsigned int __uhadd(unsigned int x, unsigned int y)
Compute average of unsigned input arguments, avoiding overflow in the intermediate sum.
__device__ unsigned int __umul24(unsigned int x, unsigned int y)
Calculate the least significant 32 bits of the product of the least significant 24 bits of two un-
signed integers.
__device__ unsigned long long int __umul64hi(unsigned long long int x, unsigned long long int y)
Calculate the most significant 64 bits of the product of the two 64 unsigned bit integers.
__device__ unsigned int __umulhi(unsigned int x, unsigned int y)
Calculate the most significant 32 bits of the product of the two 32-bit unsigned integers.
__device__ unsigned int __urhadd(unsigned int x, unsigned int y)
Compute rounded average of unsigned input arguments, avoiding overflow in the intermediate
sum.
__device__ unsigned int __usad(unsigned int x, unsigned int y, unsigned int z)
Calculate |x - y| + z , the sum of absolute difference.
13.1. Functions
__device__ unsigned int __brev(unsigned int x)
Reverse the bit order of a 32-bit unsigned integer.
Reverses the bit order of the 32-bit unsigned integer x.
Returns
Returns the bit-reversed value of x. i.e. bit N of the return value corresponds to bit
31-N of x.
__device__ unsigned long long int __brevll(unsigned long long int x)
Reverse the bit order of a 64-bit unsigned integer.
Reverses the bit order of the 64-bit unsigned integer x.
Returns
Returns the bit-reversed value of x. i.e. bit N of the return value corresponds to bit
63-N of x.
__device__ unsigned int __byte_perm(unsigned int x, unsigned int y, unsigned int s)
Return selected bytes from two 32-bit unsigned integers.
Create 8-byte source
13.1. Functions
385
▶ uint64_t tmp64 = ((uint64_t)y << 32) | x;
Extract selector bits
▶ selector0 = (s >> 0) & 0x7;
▶ selector1 = (s >> 4) & 0x7;
▶ selector2 = (s >> 8) & 0x7;
▶ selector3 = (s >> 12) & 0x7;
Return 4 selected bytes from 8-byte source:
▶ res[07:00] = tmp64[selector0];
▶ res[15:08] = tmp64[selector1];
▶ res[23:16] = tmp64[selector2];
▶ res[31:24] = tmp64[selector3];
Returns
Returns a 32-bit integer consisting of four bytes from eight input bytes provided
in the two input integers x and y, as specified by a selector, s.
__device__ int __clz(int x)
Return the number of consecutive high-order zero bits in a 32-bit integer.
Count the number of consecutive leading zero bits, starting at the most significant bit (bit 31)
of x.
Returns
Returns a value between 0 and 32 inclusive representing the number of zero bits.
__device__ int __clzll(long long int x)
Count the number of consecutive high-order zero bits in a 64-bit integer.
Count the number of consecutive leading zero bits, starting at the most significant bit (bit 63)
of x.
Returns
Returns a value between 0 and 64 inclusive representing the number of zero bits.
__device__ int __dp2a_hi(int srcA, int srcB, int c)
Two-way signed int16 by int8 dot product with int32 accumulate, taking the upper half of
the second input.
Extracts two packed 16-bit integers from scrA and two packed 8-bit integers from the upper
16 bits of srcB, then creates two pairwise 8x16 products and adds them together to a signed
32-bit integer c.
__device__ unsigned int __dp2a_hi(unsigned int srcA, unsigned int srcB, unsigned int c)
Two-way unsigned int16 by int8 dot product with unsigned int32 accumulate, taking the
upper half of the second input.
Extracts two packed 16-bit integers from scrA and two packed 8-bit integers from the upper 16
bits of srcB, then creates two pairwise 8x16 products and adds them together to an unsigned
32-bit integer c.
386
Chapter 13. Integer Intrinsics
__device__ unsigned int __dp2a_hi(ushort2 srcA, uchar4 srcB, unsigned int c)
Two-way unsigned int16 by int8 dot product with unsigned int32 accumulate, taking the
upper half of the second input.
Takes two packed 16-bit integers from scrA vector and two packed 8-bit integers from the upper
16 bits of srcB vector, then creates two pairwise 8x16 products and adds them together to an
unsigned 32-bit integer c.
__device__ int __dp2a_hi(short2 srcA, char4 srcB, int c)
Two-way signed int16 by int8 dot product with int32 accumulate, taking the upper half of
the second input.
Takes two packed 16-bit integers from scrA vector and two packed 8-bit integers from the upper
16 bits of srcB vector, then creates two pairwise 8x16 products and adds them together to a
signed 32-bit integer c.
__device__ unsigned int __dp2a_lo(ushort2 srcA, uchar4 srcB, unsigned int c)
Two-way unsigned int16 by int8 dot product with unsigned int32 accumulate, taking the
lower half of the second input.
Takes two packed 16-bit integers from scrA vector and two packed 8-bit integers from the lower
16 bits of srcB vector, then creates two pairwise 8x16 products and adds them together to an
unsigned 32-bit integer c.
__device__ int __dp2a_lo(short2 srcA, char4 srcB, int c)
Two-way signed int16 by int8 dot product with int32 accumulate, taking the lower half of
the second input.
Takes two packed 16-bit integers from scrA vector and two packed 8-bit integers from the lower
16 bits of srcB vector, then creates two pairwise 8x16 products and adds them together to a
signed 32-bit integer c.
__device__ unsigned int __dp2a_lo(unsigned int srcA, unsigned int srcB, unsigned int c)
Two-way unsigned int16 by int8 dot product with unsigned int32 accumulate, taking the
lower half of the second input.
Extracts two packed 16-bit integers from scrA and two packed 8-bit integers from the lower 16
bits of srcB, then creates two pairwise 8x16 products and adds them together to an unsigned
32-bit integer c.
__device__ int __dp2a_lo(int srcA, int srcB, int c)
Two-way signed int16 by int8 dot product with int32 accumulate, taking the lower half of
the second input.
Extracts two packed 16-bit integers from scrA and two packed 8-bit integers from the lower
16 bits of srcB, then creates two pairwise 8x16 products and adds them together to a signed
32-bit integer c.
__device__ unsigned int __dp4a(uchar4 srcA, uchar4 srcB, unsigned int c)
Four-way unsigned int8 dot product with unsigned int32 accumulate.
Takes four pairs of packed byte-sized integers from scrA and srcB vectors, then creates four
pairwise products and adds them together to an unsigned 32-bit integer c.
__device__ unsigned int __dp4a(unsigned int srcA, unsigned int srcB, unsigned int c)
Four-way unsigned int8 dot product with unsigned int32 accumulate.
Extracts four pairs of packed byte-sized integers from scrA and srcB, then creates four pairwise
products and adds them together to an unsigned 32-bit integer c.
13.1. Functions
387
__device__ int __dp4a(int srcA, int srcB, int c)
Four-way signed int8 dot product with int32 accumulate.
Extracts four pairs of packed byte-sized integers from scrA and srcB, then creates four pairwise
products and adds them together to a signed 32-bit integer c.
__device__ int __dp4a(char4 srcA, char4 srcB, int c)
Four-way signed int8 dot product with int32 accumulate.
Takes four pairs of packed byte-sized integers from scrA and srcB vectors, then creates four
pairwise products and adds them together to a signed 32-bit integer c.
__device__ int __ffs(int x)
Find the position of the least significant bit set to 1 in a 32-bit integer.
Find the position of the first (least significant) bit set to 1 in x, where the least significant bit
position is 1.
Returns
Returns a value between 0 and 32 inclusive representing the position of the first
bit set.
▶ __ffs(0) returns 0.
__device__ int __ffsll(long long int x)
Find the position of the least significant bit set to 1 in a 64-bit integer.
Find the position of the first (least significant) bit set to 1 in x, where the least significant bit
position is 1.
Returns
Returns a value between 0 and 64 inclusive representing the position of the first
bit set.
▶ __ffsll(0) returns 0.
__device__ unsigned __fns(unsigned mask, unsigned base, int offset)
Find the position of the n-th set to 1 bit in a 32-bit integer.
Given a 32-bit value mask and an integer value base (between 0 and 31), find the n-th (given by
offset) set bit in mask from the base bit. If not found, return 0xFFFFFFFF.
See
also
integer-arithmetic-instructions-fns for more information.
Returns
Returns a value between 0 and 32 inclusive representing the position of the n-th
set bit.
▶ parameter base must be <=31, otherwise behavior is undefined.
__device__ unsigned int __funnelshift_l(unsigned int lo, unsigned int hi, unsigned int shift)
Concatenate hi : lo, shift left by shift & 31 bits, return the most significant 32 bits.
Shift the 64-bit value formed by concatenating argument lo and hi left by the amount specified
by the argument shift. Argument lo holds bits 31:0 and argument hi holds bits 63:32 of the
64-bit source value. The source is shifted left by the wrapped value of shift (shift & 31). The
most significant 32-bits of the result are returned.
Returns
Returns the most significant 32 bits of the shifted 64-bit value.
388
Chapter 13. Integer Intrinsics
__device__ unsigned int __funnelshift_lc(unsigned int lo, unsigned int hi, unsigned int shift)
Concatenate hi : lo, shift left by min(shift, 32) bits, return the most significant 32 bits.
Shift the 64-bit value formed by concatenating argument lo and hi left by the amount specified
by the argument shift. Argument lo holds bits 31:0 and argument hi holds bits 63:32 of the
64-bit source value. The source is shifted left by the clamped value of shift (min(shift, 32)).
The most significant 32-bits of the result are returned.
Returns
Returns the most significant 32 bits of the shifted 64-bit value.
__device__ unsigned int __funnelshift_r(unsigned int lo, unsigned int hi, unsigned int shift)
Concatenate hi : lo, shift right by shift & 31 bits, return the least significant 32 bits.
Shift the 64-bit value formed by concatenating argument lo and hi right by the amount specified
by the argument shift. Argument lo holds bits 31:0 and argument hi holds bits 63:32 of the
64-bit source value. The source is shifted right by the wrapped value of shift (shift & 31). The
least significant 32-bits of the result are returned.
Returns
Returns the least significant 32 bits of the shifted 64-bit value.
__device__ unsigned int __funnelshift_rc(unsigned int lo, unsigned int hi, unsigned int shift)
Concatenate hi : lo, shift right by min(shift, 32) bits, return the least significant 32 bits.
Shift the 64-bit value formed by concatenating argument lo and hi right by the amount specified
by the argument shift. Argument lo holds bits 31:0 and argument hi holds bits 63:32 of the
64-bit source value. The source is shifted right by the clamped value of shift (min(shift, 32)).
The least significant 32-bits of the result are returned.
Returns
Returns the least significant 32 bits of the shifted 64-bit value.
__device__ int __hadd(int x, int y)
Compute average of signed input arguments, avoiding overflow in the intermediate sum.
Compute average of signed input arguments x and y as ( x + y ) >> 1, avoiding overflow in the
intermediate sum.
Returns
Returns a signed integer value representing the signed average value of the two
inputs.
__device__ int __mul24(int x, int y)
Calculate the least significant 32 bits of the product of the least significant 24 bits of two inte-
gers.
Calculate the least significant 32 bits of the product of the least significant 24 bits of x and y.
The high order 8 bits of x and y are ignored.
Returns
Returns the least significant 32 bits of the product x * y.
__device__ long long int __mul64hi(long long int x, long long int y)
Calculate the most significant 64 bits of the product of the two 64-bit integers.
Calculate the most significant 64 bits of the 128-bit product x * y, where x and y are 64-bit
integers.
Returns
Returns the most significant 64 bits of the product x * y.
13.1. Functions
389
__device__ int __mulhi(int x, int y)
Calculate the most significant 32 bits of the product of the two 32-bit integers.
Calculate the most significant 32 bits of the 64-bit product x * y, where x and y are 32-bit inte-
gers.
Returns
Returns the most significant 32 bits of the product x * y.
__host__ __device__ unsigned short __nv_bswap16(unsigned short x)
Reverse the order of bytes of the 16-bit unsigned integer.
Reverse the order of bytes of x . Only supported in MSVC and other host compilers which define
the __GNUC__ macro, such as GCC and CLANG.
Returns
Returns x with the order of bytes reversed.
__host__ __device__ unsigned int __nv_bswap32(unsigned int x)
Reverse the order of bytes of the 32-bit unsigned integer.
Reverse the order of bytes of x . Only supported in MSVC and other host compilers which define
the __GNUC__ macro, such as GCC and CLANG.
Returns
Returns x with the order of bytes reversed.
__host__ __device__ unsigned long long __nv_bswap64(unsigned long long x)
Reverse the order of bytes of the 64-bit unsigned integer.
Reverse the order of bytes of x . Only supported in MSVC and other host compilers which define
the __GNUC__ macro, such as GCC and CLANG.
Returns
Returns x with the order of bytes reversed.
__device__ int __popc(unsigned int x)
Count the number of bits that are set to 1 in a 32-bit integer.
Count the number of bits that are set to 1 in x.
Returns
Returns a value between 0 and 32 inclusive representing the number of set bits.
__device__ int __popcll(unsigned long long int x)
Count the number of bits that are set to 1 in a 64-bit integer.
Count the number of bits that are set to 1 in x.
Returns
Returns a value between 0 and 64 inclusive representing the number of set bits.
__device__ int __rhadd(int x, int y)
Compute rounded average of signed input arguments, avoiding overflow in the intermediate sum.
Compute average of signed input arguments x and y as ( x + y + 1 ) >> 1, avoiding overflow in the
intermediate sum.
Returns
Returns a signed integer value representing the signed rounded average value of
the two inputs.
390
Chapter 13. Integer Intrinsics
__device__ unsigned int __sad(int x, int y, unsigned int z)
Calculate |x - y| + z , the sum of absolute difference.
Calculate |x - y| + z , the 32-bit sum of the third argument z plus and the absolute value of the
difference between the first argument, x, and second argument, y.
Inputs x and y are signed 32-bit integers, input z is a 32-bit unsigned integer.
Returns
Returns |x - y| + z.
__device__ unsigned int __uhadd(unsigned int x, unsigned int y)
Compute average of unsigned input arguments, avoiding overflow in the intermediate sum.
Compute average of unsigned input arguments x and y as ( x + y ) >> 1, avoiding overflow in the
intermediate sum.
Returns
Returns an unsigned integer value representing the unsigned average value of the
two inputs.
__device__ unsigned int __umul24(unsigned int x, unsigned int y)
Calculate the least significant 32 bits of the product of the least significant 24 bits of two un-
signed integers.
Calculate the least significant 32 bits of the product of the least significant 24 bits of x and y.
The high order 8 bits of x and y are ignored.
Returns
Returns the least significant 32 bits of the product x * y.
__device__ unsigned long long int __umul64hi(unsigned long long int x, unsigned long long int y)
Calculate the most significant 64 bits of the product of the two 64 unsigned bit integers.
Calculate the most significant 64 bits of the 128-bit product x * y, where x and y are 64-bit
unsigned integers.
Returns
Returns the most significant 64 bits of the product x * y.
__device__ unsigned int __umulhi(unsigned int x, unsigned int y)
Calculate the most significant 32 bits of the product of the two 32-bit unsigned integers.
Calculate the most significant 32 bits of the 64-bit product x * y, where x and y are 32-bit un-
signed integers.
Returns
Returns the most significant 32 bits of the product x * y.
__device__ unsigned int __urhadd(unsigned int x, unsigned int y)
Compute rounded average of unsigned input arguments, avoiding overflow in the intermediate
sum.
Compute average of unsigned input arguments x and y as ( x + y + 1 ) >> 1, avoiding overflow in
the intermediate sum.
Returns
Returns an unsigned integer value representing the unsigned rounded average
value of the two inputs.
13.1. Functions
391
__device__ unsigned int __usad(unsigned int x, unsigned int y, unsigned int z)
Calculate |x - y| + z , the sum of absolute difference.
Calculate |x - y| + z , the 32-bit sum of the third argument z plus and the absolute value of the
difference between the first argument, x, and second argument, y.
Inputs x, y, and z are unsigned 32-bit integers.
Returns
Returns |x - y| + z.
392
Chapter 13. Integer Intrinsics
Chapter 14. SIMD Intrinsics
This section describes SIMD 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__ unsigned int __vabs2(unsigned int a)
Computes per-halfword absolute value: |a|.
__device__ unsigned int __vabs4(unsigned int a)
Computes per-byte absolute value: |a|.
__device__ unsigned int __vabsdiffs2(unsigned int a, unsigned int b)
Computes per-halfword absolute difference of signed integer: |a - b|.
__device__ unsigned int __vabsdiffs4(unsigned int a, unsigned int b)
Computes per-byte absolute difference of signed integer: |a - b|.
__device__ unsigned int __vabsdiffu2(unsigned int a, unsigned int b)
Computes per-halfword absolute difference of unsigned integer: |a - b|.
__device__ unsigned int __vabsdiffu4(unsigned int a, unsigned int b)
Computes per-byte absolute difference of unsigned integer: |a - b|.
__device__ unsigned int __vabsss2(unsigned int a)
Computes per-halfword absolute value with signed saturation: |a|.
__device__ unsigned int __vabsss4(unsigned int a)
Computes per-byte absolute value with signed saturation: |a|.
__device__ unsigned int __vadd2(unsigned int a, unsigned int b)
Performs per-halfword (un)signed addition, with wrap-around: a + b.
__device__ unsigned int __vadd4(unsigned int a, unsigned int b)
Performs per-byte (un)signed addition: a + b.
__device__ unsigned int __vaddss2(unsigned int a, unsigned int b)
Performs per-halfword addition with signed saturation: a + b.
__device__ unsigned int __vaddss4(unsigned int a, unsigned int b)
Performs per-byte addition with signed saturation: a + b.
__device__ unsigned int __vaddus2(unsigned int a, unsigned int b)
Performs per-halfword addition with unsigned saturation: a + b.
__device__ unsigned int __vaddus4(unsigned int a, unsigned int b)
Performs per-byte addition with unsigned saturation: a + b.
393
__device__ unsigned int __vavgs2(unsigned int a, unsigned int b)
Performs per-halfword signed rounded average computation.
__device__ unsigned int __vavgs4(unsigned int a, unsigned int b)
Computes per-byte signed rounded average.
__device__ unsigned int __vavgu2(unsigned int a, unsigned int b)
Performs per-halfword unsigned rounded average computation.
__device__ unsigned int __vavgu4(unsigned int a, unsigned int b)
Performs per-byte unsigned rounded average.
__device__ unsigned int __vcmpeq2(unsigned int a, unsigned int b)
Performs per-halfword (un)signed comparison: a == b ? 0xffff : 0.
__device__ unsigned int __vcmpeq4(unsigned int a, unsigned int b)
Performs per-byte (un)signed comparison: a == b ? 0xff : 0.
__device__ unsigned int __vcmpges2(unsigned int a, unsigned int b)
Performs per-halfword signed comparison: a >= b ? 0xffff : 0.
__device__ unsigned int __vcmpges4(unsigned int a, unsigned int b)
Performs per-byte signed comparison: a >= b ? 0xff : 0.
__device__ unsigned int __vcmpgeu2(unsigned int a, unsigned int b)
Performs per-halfword unsigned comparison: a >= b ? 0xffff : 0.
__device__ unsigned int __vcmpgeu4(unsigned int a, unsigned int b)
Performs per-byte unsigned comparison: a >= b ? 0xff : 0.
__device__ unsigned int __vcmpgts2(unsigned int a, unsigned int b)
Performs per-halfword signed comparison: a > b ? 0xffff : 0.
__device__ unsigned int __vcmpgts4(unsigned int a, unsigned int b)
Performs per-byte signed comparison: a > b ? 0xff : 0.
__device__ unsigned int __vcmpgtu2(unsigned int a, unsigned int b)
Performs per-halfword unsigned comparison: a > b ? 0xffff : 0.
__device__ unsigned int __vcmpgtu4(unsigned int a, unsigned int b)
Performs per-byte unsigned comparison: a > b ? 0xff : 0.
__device__ unsigned int __vcmples2(unsigned int a, unsigned int b)
Performs per-halfword signed comparison: a <= b ? 0xffff : 0.
__device__ unsigned int __vcmples4(unsigned int a, unsigned int b)
Performs per-byte signed comparison: a <= b ? 0xff : 0.
__device__ unsigned int __vcmpleu2(unsigned int a, unsigned int b)
Performs per-halfword unsigned comparison: a <= b ? 0xffff : 0.
__device__ unsigned int __vcmpleu4(unsigned int a, unsigned int b)
Performs per-byte unsigned comparison: a <= b ? 0xff : 0.
__device__ unsigned int __vcmplts2(unsigned int a, unsigned int b)
Performs per-halfword signed comparison: a < b ? 0xffff : 0.
__device__ unsigned int __vcmplts4(unsigned int a, unsigned int b)
Performs per-byte signed comparison: a < b ? 0xff : 0.
__device__ unsigned int __vcmpltu2(unsigned int a, unsigned int b)
Performs per-halfword unsigned comparison: a < b ? 0xffff : 0.
394
Chapter 14. SIMD Intrinsics
|
||
|
|
|