cuda::sub_sat_overflow

Defined in <cuda/numeric> header.

namespace cuda {

template <class T>
struct overflow_result;

template <class T>
[[nodiscard]] __host__ __device__ constexpr
overflow_result<T> sub_sat_overflow(T lhs, T rhs) noexcept; // (1)

template <class T>
[[nodiscard]] __host__ __device__ constexpr
bool sub_sat_overflow(T& result, T lhs, T rhs) noexcept; // (2)

} // namespace cuda

The function cuda::sub_sat_overflow performs saturating subtraction of two values lhs and rhs with overflow detection.

Parameters

• result: The result of the saturating subtraction. (2)

• lhs: The left-hand side operand. (1, 2)

• rhs: The right-hand side operand. (1, 2)

Return value

1. Returns an overflow_result object containing the result of the saturating subtraction and a boolean flag indicating whether an overflow or underflow occurred.

2. Returns true if an overflow or underflow occurred, false otherwise.

Constraints

• T must be an integer type.

Performance considerations

• Functionality is implemented by correcting the cuda::sub_overflow result in case of overflow/underflow.

• Unsigned computations are generally faster than signed computations.

Example

#include <cuda/numeric>
#include <cuda/std/cassert>
#include <cuda/std/limits>

__global__ void kernel()
{
    constexpr auto uint_max = cuda::std::numeric_limits<unsigned>::max();

    const auto result = cuda::sub_sat_overflow(42, uint_max); // saturated
    assert(result.value == 0);
    assert(result.overflow);

    int value;
    if (cuda::sub_sat_overflow(value, 42, 1024))
    {
        assert(false); // shouldn't be reached
    }
    assert(value == 982);
}

int main()
{
    kernel<<<1, 1>>>();
    cudaDeviceSynchronize();
}

See it on Godbolt 🔗