cudaarithm: cuda::abs does not saturate the most negative 8S/16S/32S value (abs(-128) = -128), and cuda::norm(NORM_INF) can return 0 or a negative value for such inputs
System information (version)
- OpenCV => 5.x, opencv f8ef5fd (2026-09-06) + opencv_contrib 17af220 (2026-08-14), built from source with CUDA 12.6, CUDA_ARCH_BIN=8.7
- Operating System / Platform => Ubuntu 22.04.5 aarch64 (NVIDIA Jetson Orin Nano, JetPack 6.2.3); only this device was available, but nothing in the code path depends on the device
- Compiler => GCC 11.4.0, nvcc 12.6
Detailed description
cv::abs saturates: for 8S, 16S and 32S input the most negative element maps to the most positive one (127, 32767, 2147483647; the cpu columns below), and cudaarithm.hpp refers to cv::abs and cv::norm for the semantics (@sa abs, line 241; @sa norm, line 822). cv::cuda::abs goes through absMat in modules/cudaarithm/src/cuda/math.cu (lines 77-82, dispatched at 84-107) with abs_func<T> from modules/cudev/include/opencv2/cudev/functional/functional.hpp:
template <> struct abs_func<schar> : unary_function<schar, schar>
{
__device__ __forceinline__ schar operator ()(schar x) const
{
return ::abs((int) x);
}The int result is converted back to the element type without saturation, so -128 comes back as -128; short is the same at lines 384-389, and abs_func<int> (lines 400-406) is return ::abs(x);, which returns INT_MIN for INT_MIN here. The consequences in the run below:
cuda::abs: 8SC1-128 -> -128, 16SC1-32768 -> -32768, 32SC1INT_MIN -> INT_MIN(hostcv::abs: 127, 32767, 2147483647).cuda::norm(NORM_INF)isgridFindMaxVal(abs_(src), ...)inmodules/cudaarithm/src/cuda/minmax.cu(lines 147-157): the maximum is taken over those unsaturated values in the element type, so the most negative element only wins when there is nothing else. In the run below: 127 instead of 128 (8S), 32767 instead of 32768 (16S), 2147483647 instead of 2147483648 (32S); for the input{-128, 0}the GPU returns 0, for{-128, 3, -3}it returns 3, and for an input consisting only of -128 it returns -128, a negative norm (16S the same; 32S: 0 for{INT_MIN, 0}and -2147483647 for all-INT_MIN, last block below).cuda::absdiff(a, b)for 32SC1:|INT_MIN - 0| -> INT_MIN(modules/cudaarithm/src/cuda/absdiff_mat.culines 72-78,saturate_cast<T>(_abs(a - b)), which is the identity forint); 8S and 16S are correct. The scalar overloadcuda::absdiff(a, Scalar::all(0))computes infloatfor CV_32S (absDiffScalarImpl<int, float>,absdiff_scalar.culine 113) and returns 2147483647 for this input, so the two overloads disagree on the same CV_32S input. (core.hpplines 1405-1406 say thatabsdiffdoes not saturate for CV_32S; the 5.x host implementation does saturate,modules/core/src/arithm.simd.hpplines 522-530, and so does the CUDA scalar overload.)cuda::norm(NORM_L1)andcuda::absSumare correct for all three types:modules/cudaarithm/src/cuda/sum.culine 86 widens before the absolute value,abs_(cvt_<res_type>(src)).
(The cpu value on the 32SC1 norm L1 row, -2147482644, is a host-side overflow in cv::norm, unrelated to this report; the correct sum is the absSum value 6442451948.)
The existing test CUDA_Arithm/Abs (modules/cudaarithm/test/test_element_operations.cpp lines 1504-1508) is instantiated for CV_16S and CV_32F only, and randomMat defaults to values in [0, 255] (modules/ts/include/opencv2/ts/cuda_test.hpp line 60), so the boundary value is never tested.
Suggested fix: saturate in abs_func<schar> / abs_func<short> (saturate_cast<T>(::abs((int) x))) and return INT_MAX for INT_MIN in abs_func<int>; in findMaxAbsImpl widen before taking the absolute value (abs_(cvt_<R>(src))), otherwise norm(NORM_INF) on 8S would still give 127 against the host's 128; for CV_32S that alone is not enough, because R is int there (minmax.cu line 169) and the result depth is CV_32S (line 181), which cannot hold 2147483648; in AbsDiffOp1<int> compute the difference and its absolute value in a wider type.
Steps to reproduce
// repro_E2_cuda_abs_norm_saturation.cpp — cv::cuda::abs, cuda::absdiff (matrix and scalar overloads), cuda::norm
// (NORM_INF, NORM_L1) and cuda::absSum on 8S / 16S / 32S inputs that contain the most negative value of the type,
// against the host functions. All values are printed with %.0f (no %g rounding).
#include <opencv2/core.hpp>
#include <opencv2/core/cuda.hpp>
#include <opencv2/cudaarithm.hpp>
#include <cstdio>
#include <climits>
#include <string>
template <typename T> static void row(const char* label, const cv::Mat& c, const cv::Mat& g) {
printf("%-30s cpu:", label); for (int i = 0; i < c.cols; i++) printf(" %.0f", (double)c.at<T>(0, i));
printf(" gpu:"); for (int i = 0; i < g.cols; i++) printf(" %.0f", (double)g.at<T>(0, i)); printf("\n");
}
template <typename T> static void one(const char* name, const cv::Mat& a) {
cv::cuda::GpuMat ga(a), gd; cv::Mat cpu, gpu; std::string n = name;
cpu = cv::abs(a); cv::cuda::abs(ga, gd); gd.download(gpu); row<T>((n + " abs").c_str(), cpu, gpu);
cv::Mat z = cv::Mat::zeros(a.size(), a.type()), cd, gdh; cv::cuda::GpuMat gz(z), gdd;
cv::absdiff(a, z, cd); cv::cuda::absdiff(ga, gz, gdd); gdd.download(gdh); row<T>((n + " absdiff(a, zeros)").c_str(), cd, gdh);
cv::absdiff(a, cv::Scalar::all(0), cd); cv::cuda::absdiff(ga, cv::Scalar::all(0), gdd); gdd.download(gdh); row<T>((n + " absdiff(a, Scalar(0))").c_str(), cd, gdh);
printf("%-30s cpu: %.0f gpu: %.0f\n", (n + " norm INF").c_str(), cv::norm(a, cv::NORM_INF), cv::cuda::norm(ga, cv::NORM_INF));
printf("%-30s cpu: %.0f gpu: %.0f\n", (n + " norm L1").c_str(), cv::norm(a, cv::NORM_L1), cv::cuda::norm(ga, cv::NORM_L1));
cv::Mat a64; a.convertTo(a64, CV_64F);
printf("%-30s cpu: %.0f gpu: %.0f\n", (n + " absSum").c_str(), cv::sum(cv::abs(a64))[0], cv::cuda::absSum(ga)[0]);
}
template <typename T> static void normInf(const char* name, const cv::Mat& a) {
cv::cuda::GpuMat ga(a); printf("%-30s input:", name); for (int i = 0; i < a.cols; i++) printf(" %.0f", (double)a.at<T>(0, i));
printf(" cpu norm INF: %.0f gpu norm INF: %.0f\n", cv::norm(a, cv::NORM_INF), cv::cuda::norm(ga, cv::NORM_INF));
}
int main() {
cv::cuda::printShortCudaDeviceInfo(cv::cuda::getDevice());
one<schar>("8SC1", (cv::Mat_<schar>(1, 6) << 0, -128, -127, 5, -128, 100));
one<short>("16SC1", (cv::Mat_<short>(1, 6) << 0, -32768, -32767, 5, -32768, 1000));
one<int>("32SC1", (cv::Mat_<int>(1, 6) << 0, INT_MIN, INT_MIN + 1, 5, INT_MIN, 1000));
printf("--- norm INF with fewer other values:\n");
normInf<schar>("8SC1 {-128, 0}", (cv::Mat_<schar>(1, 2) << -128, 0));
normInf<schar>("8SC1 {-128, -128, -128}", (cv::Mat_<schar>(1, 3) << -128, -128, -128));
normInf<schar>("8SC1 {-128, 3, -3}", (cv::Mat_<schar>(1, 3) << -128, 3, -3));
normInf<short>("16SC1 {-32768, 0}", (cv::Mat_<short>(1, 2) << -32768, 0));
normInf<short>("16SC1 {-32768, -32768}", (cv::Mat_<short>(1, 2) << -32768, -32768));
normInf<int>("32SC1 {INT_MIN, 0}", (cv::Mat_<int>(1, 2) << INT_MIN, 0));
normInf<int>("32SC1 {INT_MIN, INT_MIN}", (cv::Mat_<int>(1, 2) << INT_MIN, INT_MIN));
return 0;
}Built with g++ -std=c++17 -O2 repro_E2.cpp -o repro_E2 -I<build> -I<opencv>/include -I<opencv>/modules/core/include -I<contrib>/modules/cudaarithm/include -I<contrib>/modules/cudev/include -I/usr/local/cuda/include -L<build>/lib -lopencv_cudaarithm -lopencv_core. Output:
Device 0: "Orin" 7607Mb, sm_87, Driver/Runtime ver.12.60/12.60
8SC1 abs cpu: 0 127 127 5 127 100 gpu: 0 -128 127 5 -128 100
8SC1 absdiff(a, zeros) cpu: 0 127 127 5 127 100 gpu: 0 127 127 5 127 100
8SC1 absdiff(a, Scalar(0)) cpu: 0 127 127 5 127 100 gpu: 0 127 127 5 127 100
8SC1 norm INF cpu: 128 gpu: 127
8SC1 norm L1 cpu: 488 gpu: 488
8SC1 absSum cpu: 488 gpu: 488
16SC1 abs cpu: 0 32767 32767 5 32767 1000 gpu: 0 -32768 32767 5 -32768 1000
16SC1 absdiff(a, zeros) cpu: 0 32767 32767 5 32767 1000 gpu: 0 32767 32767 5 32767 1000
16SC1 absdiff(a, Scalar(0)) cpu: 0 32767 32767 5 32767 1000 gpu: 0 32767 32767 5 32767 1000
16SC1 norm INF cpu: 32768 gpu: 32767
16SC1 norm L1 cpu: 99308 gpu: 99308
16SC1 absSum cpu: 99308 gpu: 99308
32SC1 abs cpu: 0 2147483647 2147483647 5 2147483647 1000 gpu: 0 -2147483648 2147483647 5 -2147483648 1000
32SC1 absdiff(a, zeros) cpu: 0 2147483647 2147483647 5 2147483647 1000 gpu: 0 -2147483648 2147483647 5 -2147483648 1000
32SC1 absdiff(a, Scalar(0)) cpu: 0 2147483647 2147483647 5 2147483647 1000 gpu: 0 2147483647 2147483647 5 2147483647 1000
32SC1 norm INF cpu: 2147483648 gpu: 2147483647
32SC1 norm L1 cpu: -2147482644 gpu: 6442451948
32SC1 absSum cpu: 6442451948 gpu: 6442451948
--- norm INF with fewer other values:
8SC1 {-128, 0} input: -128 0 cpu norm INF: 128 gpu norm INF: 0
8SC1 {-128, -128, -128} input: -128 -128 -128 cpu norm INF: 128 gpu norm INF: -128
8SC1 {-128, 3, -3} input: -128 3 -3 cpu norm INF: 128 gpu norm INF: 3
16SC1 {-32768, 0} input: -32768 0 cpu norm INF: 32768 gpu norm INF: 0
16SC1 {-32768, -32768} input: -32768 -32768 cpu norm INF: 32768 gpu norm INF: -32768
32SC1 {INT_MIN, 0} input: -2147483648 0 cpu norm INF: 2147483648 gpu norm INF: 0
32SC1 {INT_MIN, INT_MIN} input: -2147483648 -2147483648 cpu norm INF: 2147483648 gpu norm INF: -2147483647Note: the investigation and this report were prepared with an AI assistant; all runs are from my own board and the output blocks are the programs' own output, pasted line for line.
Issue submission checklist
- I report the issue, it's not a question
- I checked the problem with documentation, FAQ, open issues, forum.opencv.org, Stack Overflow, etc and have not found any solution
- I updated to the latest OpenCV version and the issue is still there
- There is reproducer code and related data files: videos, images, onnx, etc
Source: opencv/opencv_contrib