cudaarithm: cuda::multiply / sqr compute the 16U product in int and do not saturate (65535*65535 = 0); a 32S product with dtype=CV_32F is wrapped too
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::multiply (which the cv::cuda::multiply documentation refers to) is documented as dst(I) = saturate(scale * src1(I) * src2(I)), with the note that saturation is not applied for CV_32S output. cv::cuda::multiply with scale == 1 uses MulOp from modules/cudaarithm/src/cuda/mul_mat.cu:
template <typename T, typename D> struct MulOp : binary_function<T, T, D>
{
__device__ __forceinline__ D operator ()(T a, T b) const
{
return saturate_cast<D>(a * b);
}
};For T = ushort or int the product a * b is an int and overflows before saturate_cast sees it. Measured (the expected rows are the exact product clamped to the type range, computed in double):
- 16U x 16U:
65535 * 65535 -> 0,46341 * 46341 -> 0(product just above INT_MAX),46340 * 46340 -> 65535(just below);cuda::sqron 16U the same.cuda::pow(src, 2)on the same values is correct. - 32S x 32S with
dtype = CV_32F(a float output, so the CV_32S note does not apply): the wrappedintproducts cast to float (1410065408,-2,0) instead of the float products1e10,4294967296,4294967296. - 32S x 32S with 32S output:
100000 * 100000 -> 1410065408(the product modulo 2^32),2147483647 * 2 -> -2,65536 * 65536 -> 0; this is covered by the CV_32S note and listed for completeness (thescale = 0.5path below saturates to 2147483647 for the same inputs). - With
scale = 0.5theMulScaleOppath (saturate_cast<D>(scale * a * b), lines 69-77) computes in the scale type and is correct (up to float rounding:46341 * 46341 * 0.5 -> 1073744128, exact 1073744140.5). - 16S x 16S is correct (the product fits in
int); 8U x 8U is correct. (cuda::pow(src, 3)on 16S gives -32767 for(-32768)^3, a different mechanism that is not part of this report.)
The same overflow on the CPU side was opencv/opencv#28557, fixed in 4.x by opencv/opencv#28579 with a uint16_t specialisation of the multiply op (saturate_cast<uint16_t>(static_cast<uint32_t>(a) * static_cast<uint32_t>(b))). The expected values above are computed in double rather than taken from the host, because the host cv::multiply of this 5.x build has the same 16U overflow for short arrays (its regression test is DISABLED_mul_overflow_28557 in 5.x; reported separately); for 16U arrays of 64 elements or more the host returns the saturated values on both aarch64 and x86.
Suggested fix: as in #28579, a T = ushort specialisation of MulOp (and of the sqr op) that widens to unsigned before the saturating cast, saturate_cast<D>((unsigned)a * (unsigned)b) (cudev's saturate_cast<ushort>(unsigned) clamps; the generic body must stay as it is for 16S), and for a float dtype a product computed in double (a float product would round differently from the host for non-overflowing inputs).
Steps to reproduce
// repro_F2_cuda_multiply_saturation.cpp — cv::cuda::multiply / sqr / pow on 16U, 16S and 32S inputs whose product
// does not fit in int. "expected" is the exact product (computed in double) clamped to the type range, i.e. what
// cv::multiply documents (saturation); the host OpenCV is not used as the reference.
#include <opencv2/core.hpp>
#include <opencv2/core/cuda.hpp>
#include <opencv2/cudaarithm.hpp>
#include <cstdio>
#include <cmath>
#include <algorithm>
template <typename T> static double clampT(double v) { double lo = (double)std::numeric_limits<T>::min(), hi = (double)std::numeric_limits<T>::max(); return std::min(std::max(std::round(v), lo), hi); }
template <typename T> static void show(const char* label, const cv::Mat& a, const cv::Mat& b, const cv::Mat& g, double scale, int power) {
printf("%-36s a:", label); for (int i = 0; i < a.cols; i++) printf(" %.0f", (double)a.at<T>(0, i));
printf(" b:"); for (int i = 0; i < b.cols; i++) printf(" %.0f", (double)b.at<T>(0, i));
printf("\n%-36s expected:", ""); for (int i = 0; i < a.cols; i++) { double v = power ? std::pow((double)a.at<T>(0, i), power) : (double)a.at<T>(0, i) * (double)b.at<T>(0, i) * scale; printf(" %.0f", clampT<T>(v)); }
printf("\n%-36s gpu: ", ""); for (int i = 0; i < g.cols; i++) printf(" %.0f", (double)g.at<T>(0, i)); printf("\n");
}
int main() {
cv::cuda::printShortCudaDeviceInfo(cv::cuda::getDevice());
{ cv::Mat a = (cv::Mat_<ushort>(1, 6) << 65535, 65535, 46341, 46340, 300, 2), b = (cv::Mat_<ushort>(1, 6) << 65535, 2, 46341, 46340, 300, 3), g;
cv::cuda::GpuMat ga(a), gb(b), gd;
cv::cuda::multiply(ga, gb, gd); gd.download(g); show<ushort>("multiply 16U x 16U", a, b, g, 1, 0);
cv::cuda::multiply(ga, gb, gd, 0.5); gd.download(g); show<ushort>("multiply 16U x 16U scale 0.5", a, b, g, 0.5, 0);
cv::cuda::sqr(ga, gd); gd.download(g); show<ushort>("sqr 16U", a, a, g, 1, 0);
cv::cuda::pow(ga, 2, gd); gd.download(g); show<ushort>("pow(16U, 2)", a, a, g, 1, 2); }
{ cv::Mat a = (cv::Mat_<short>(1, 6) << 32767, -32768, 32767, -30000, 300, 2), b = (cv::Mat_<short>(1, 6) << 32767, 32767, -32768, 30000, 300, 3), g;
cv::cuda::GpuMat ga(a), gb(b), gd;
cv::cuda::multiply(ga, gb, gd); gd.download(g); show<short>("multiply 16S x 16S", a, b, g, 1, 0);
cv::cuda::pow(ga, 3, gd); gd.download(g); show<short>("pow(16S, 3)", a, a, g, 1, 3); }
{ cv::Mat a = (cv::Mat_<int>(1, 6) << 100000, 2147483647, -100000, 65536, 46341, 2), b = (cv::Mat_<int>(1, 6) << 100000, 2, 100000, 65536, 46341, 3), g;
cv::cuda::GpuMat ga(a), gb(b), gd;
cv::cuda::multiply(ga, gb, gd); gd.download(g); show<int>("multiply 32S x 32S", a, b, g, 1, 0);
cv::cuda::multiply(ga, gb, gd, 0.5); gd.download(g); show<int>("multiply 32S x 32S scale 0.5", a, b, g, 0.5, 0);
cv::cuda::multiply(ga, gb, gd, 1.0, CV_32F); gd.download(g);
printf("%-36s expected (exact product as float):", "multiply 32S x 32S, dtype=CV_32F"); for (int i = 0; i < 6; i++) printf(" %.0f", (double)(float)((double)a.at<int>(0, i) * b.at<int>(0, i)));
printf("\n%-36s gpu: ", ""); for (int i = 0; i < 6; i++) printf(" %.0f", (double)g.at<float>(0, i)); printf("\n"); }
{ cv::Mat a = (cv::Mat_<uchar>(1, 4) << 255, 255, 16, 200), b = (cv::Mat_<uchar>(1, 4) << 255, 2, 16, 200), g;
cv::cuda::GpuMat ga(a), gb(b), gd; cv::cuda::multiply(ga, gb, gd); gd.download(g); show<uchar>("multiply 8U x 8U (control)", a, b, g, 1, 0); }
return 0;
}Built with g++ -std=c++17 -O2 repro_F2.cpp -o repro_F2 -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
multiply 16U x 16U a: 65535 65535 46341 46340 300 2 b: 65535 2 46341 46340 300 3
expected: 65535 65535 65535 65535 65535 6
gpu: 0 65535 0 65535 65535 6
multiply 16U x 16U scale 0.5 a: 65535 65535 46341 46340 300 2 b: 65535 2 46341 46340 300 3
expected: 65535 65535 65535 65535 45000 3
gpu: 65535 65535 65535 65535 45000 3
sqr 16U a: 65535 65535 46341 46340 300 2 b: 65535 65535 46341 46340 300 2
expected: 65535 65535 65535 65535 65535 4
gpu: 0 0 0 65535 65535 4
pow(16U, 2) a: 65535 65535 46341 46340 300 2 b: 65535 65535 46341 46340 300 2
expected: 65535 65535 65535 65535 65535 4
gpu: 65535 65535 65535 65535 65535 4
multiply 16S x 16S a: 32767 -32768 32767 -30000 300 2 b: 32767 32767 -32768 30000 300 3
expected: 32767 -32768 -32768 -32768 32767 6
gpu: 32767 -32768 -32768 -32768 32767 6
pow(16S, 3) a: 32767 -32768 32767 -30000 300 2 b: 32767 -32768 32767 -30000 300 2
expected: 32767 -32768 32767 -32768 32767 8
gpu: 32767 -32767 32767 -32767 32767 8
multiply 32S x 32S a: 100000 2147483647 -100000 65536 46341 2 b: 100000 2 100000 65536 46341 3
expected: 2147483647 2147483647 -2147483648 2147483647 2147483647 6
gpu: 1410065408 -2 -1410065408 0 -2147479015 6
multiply 32S x 32S scale 0.5 a: 100000 2147483647 -100000 65536 46341 2 b: 100000 2 100000 65536 46341 3
expected: 2147483647 2147483647 -2147483648 2147483647 1073744141 3
gpu: 2147483647 2147483647 -2147483648 2147483647 1073744128 3
multiply 32S x 32S, dtype=CV_32F expected (exact product as float): 10000000000 4294967296 -10000000000 4294967296 2147488256 6
gpu: 1410065408 -2 -1410065408 0 -2147479040 6
multiply 8U x 8U (control) a: 255 255 16 200 b: 255 2 16 200
expected: 255 255 255 255
gpu: 255 255 255 255Note: 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