cudaimgproc: cuda::demosaicing (and cuda::cvtColor with Bayer codes) on a GpuMat ROI whose x offset is not a multiple of 4 (8U) / 2 (16U) pixels fails with a sticky "misaligned address" CUDA error
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
The bilinear Bayer kernels in modules/cudaimgproc/src/cuda/debayer.cu read the source rows as 4-byte words: Bayer2BGR<uchar>::apply loads uchar4 (lines 67-77, e.g.
patch[1][1] = ((const uchar4*) src.ptr(s_y))[s_x];) and Bayer2BGR<ushort>::apply loads ushort2 the same way (lines 233-243). When src is a ROI of a larger GpuMat whose x offset in bytes is not a multiple of 4, i.e. an 8U x offset that is not a multiple of 4 or a 16U x offset that is odd, these loads are misaligned and the kernel fails with misaligned address. On the default stream the error is reported by the cudaSafeCall( cudaDeviceSynchronize() ) in Bayer2BGR_8u_gpu (line 357) and Bayer2BGR_16u_gpu (line 374); on a user stream the call returns and the error surfaces at the next synchronisation. The error is sticky: the context is left unusable and every later CUDA call that touches the device (allocations, copies, kernels, synchronisation) fails with the same error, hence one process per trial below. cv::cuda::cvtColor with the same Bayer codes goes through the same kernels (modules/cudaimgproc/src/color.cpp, line 1994 bayerBG_to_BGR) and fails identically (third run below).
Unaligned ROIs look like intended input for this module: the test helper createMat(..., useRoi) in modules/ts/src/cuda_test.cpp (lines 98-110) builds ROIs at an x offset of randomInt(5, 15) / 2, i.e. 2 to 7, the _MHT demosaicing path handles ROIs explicitly with locateROI (modules/cudaimgproc/src/color.cpp lines 2150 and 2189), and the same class of bug in cuda::calcHist was fixed in the kernel (#3473, PR #3475). I could not find a documented alignment requirement for cudaimgproc inputs.
A second consequence of the whole-vector loads: each row is processed as whole words (grid divUp(src.cols, 4 * block.x), line 349; thread guard (s_x << 2) >= src.cols, line 205; 16U: grid line 366, guard line 327), so when the width is not a multiple of 4 (8U) or is odd (16U) the last word, loaded on lines 67, 71 and 75 (16U: 233, 237, 241), extends past the end of the row; the right-hand neighbour is clamped to that same word (::min(s_x + 1, ((src.cols + 3) >> 2) - 1), lines 69, 73, 77; 16U ((src.cols + 1) >> 1) - 1, lines 235, 239, 243) and the out-of-row output columns are dropped (lines 217-222; 16U 339-340), but the last surviving output column is computed from the element at index cols, one past the end of the row. For a ROI that element is the parent pixel immediately to the right; for a GpuMat allocated by OpenCV that is not a ROI it is the row's pitch padding. In the second run below only the parent column immediately to the right of the ROI is changed between the two calls, and the last output column changes with it (8U widths 17, 18, 19, 113; 16U widths 17, 113; not for widths 16, 20 or 16U 18).
From the run below (one process per trial): for COLOR_BayerBG2BGR (46), COLOR_BayerGR2BGR (49) and COLOR_BayerBG2GRAY (86), widths 16, 17 and 113, 8U x offsets 1, 2, 3, 5 fail and 0, 4 work; 16U x offsets 1, 3, 5 fail and 0, 2, 4 work. Copying the ROI into a GpuMat of its own (roi.clone(), whose rows start at an aligned address) avoids the error (third run below); that is the workaround I use, with the over-read above then reading the clone's padding. (The max|cpu-gpu| value on the successful lines compares the GPU output with cv::cvtColor for the same code; the two use different Bayer code conventions, which is #4024 and not the subject of this report.)
The Demosaicing tests in modules/cudaimgproc/test/test_color.cpp are instantiated with WHOLE_SUBMAT (line 2469), but the bilinear cases call loadMat(src) without useRoi (e.g. line 2409); only the _MHT cases pass it, so the ROI half of the instantiation never reaches these kernels.
Suggested fix: handle the unaligned head and the partial last word of each row separately, the way histogram256Kernel was changed for cuda::calcHist in #3475, or shift the word loads to the previous aligned address and select the bytes and clamp the last word's elements to the row; failing that, an explicit assertion would still be better than a sticky CUDA error.
Steps to reproduce
// repro_H_cuda_demosaicing_misaligned_roi.cpp — cv::cuda::demosaicing on a GpuMat ROI whose x offset is not a
// multiple of 4 (8U) / 2 (16U) pixels: one trial per process (the CUDA "misaligned address" error is sticky).
// usage: repro_H <code> <8U|16U> <width> <xoff>
#include <opencv2/core.hpp>
#include <opencv2/core/cuda.hpp>
#include <opencv2/imgproc.hpp>
#include <opencv2/cudaimgproc.hpp>
#include <cstdio>
#include <cstdlib>
#include <cstring>
int main(int argc, char** argv) {
if (argc < 5) { fprintf(stderr, "usage: repro_H <code> <8U|16U> <width> <xoff>\n"); return 2; }
int code = atoi(argv[1]); int t = !strcmp(argv[2], "16U") ? CV_16UC1 : CV_8UC1; int w = atoi(argv[3]), xo = atoi(argv[4]), h = 8;
cv::Mat parent(h, w + xo + 3, t); cv::randu(parent, 0, 255); cv::Mat roi = parent(cv::Rect(xo, 0, w, h));
cv::cuda::GpuMat gparent(parent), d; cv::cuda::GpuMat groi = gparent(cv::Rect(xo, 0, w, h));
printf("demosaicing code=%d %s w=%d xoff=%d data%%4=%zu : ", code, argv[2], w, xo, (size_t)groi.data % 4); fflush(stdout);
try { cv::cuda::demosaicing(groi, d, code); cv::Mat g; d.download(g); cv::Mat c; cv::cvtColor(roi, c, code);
printf("ok, %s %dx%d, max|cpu-gpu|=%g\n", g.channels() == 3 ? "3ch" : "1ch", g.cols, g.rows, cv::norm(c, g, cv::NORM_INF)); }
catch (const cv::Exception& e) { std::string s = e.err; for (auto& ch : s) if (ch == '\n') ch = ' '; printf("EXC %s\n", s.c_str()); return 1; }
return 0;
}Built with g++ -std=c++17 -O2 repro_H.cpp -o repro_H -I<build> -I<opencv>/include -I<opencv>/modules/core/include -I<opencv>/modules/imgproc/include -I<contrib>/modules/cudaimgproc/include -I<contrib>/modules/cudev/include -I/usr/local/cuda/include -L<build>/lib -lopencv_cudaimgproc -lopencv_imgproc -lopencv_core, run as for code in 46 49 86; do for ty in 8U 16U; do for w in 16 17 113; do for xo in 0 1 2 3 4 5; do ./repro_H $code $ty $w $xo; done; done; done; done. Verbatim lines from that run (code 46 at width 16/17 for 8U, width 16 for 16U, and code 86 at width 16; the remaining widths and code 49 behave the same):
demosaicing code=46 8U w=16 xoff=0 data%4=0 : ok, 3ch 16x8, max|cpu-gpu|=168
demosaicing code=46 8U w=16 xoff=1 data%4=1 : EXC misaligned address
demosaicing code=46 8U w=16 xoff=2 data%4=2 : EXC misaligned address
demosaicing code=46 8U w=16 xoff=3 data%4=3 : EXC misaligned address
demosaicing code=46 8U w=16 xoff=4 data%4=0 : ok, 3ch 16x8, max|cpu-gpu|=179
demosaicing code=46 8U w=16 xoff=5 data%4=1 : EXC misaligned address
demosaicing code=46 8U w=17 xoff=0 data%4=0 : ok, 3ch 17x8, max|cpu-gpu|=179
demosaicing code=46 8U w=17 xoff=1 data%4=1 : EXC misaligned address
demosaicing code=46 8U w=17 xoff=2 data%4=2 : EXC misaligned address
demosaicing code=46 8U w=17 xoff=3 data%4=3 : EXC misaligned address
demosaicing code=46 8U w=17 xoff=4 data%4=0 : ok, 3ch 17x8, max|cpu-gpu|=179
demosaicing code=46 8U w=17 xoff=5 data%4=1 : EXC misaligned address
demosaicing code=46 16U w=16 xoff=0 data%4=0 : ok, 3ch 16x8, max|cpu-gpu|=168
demosaicing code=46 16U w=16 xoff=1 data%4=2 : EXC misaligned address
demosaicing code=46 16U w=16 xoff=2 data%4=0 : ok, 3ch 16x8, max|cpu-gpu|=168
demosaicing code=46 16U w=16 xoff=3 data%4=2 : EXC misaligned address
demosaicing code=46 16U w=16 xoff=4 data%4=0 : ok, 3ch 16x8, max|cpu-gpu|=179
demosaicing code=46 16U w=16 xoff=5 data%4=2 : EXC misaligned address
demosaicing code=86 8U w=16 xoff=0 data%4=0 : ok, 1ch 16x8, max|cpu-gpu|=109
demosaicing code=86 8U w=16 xoff=1 data%4=1 : EXC misaligned address
demosaicing code=86 8U w=16 xoff=2 data%4=2 : EXC misaligned address
demosaicing code=86 8U w=16 xoff=3 data%4=3 : EXC misaligned address
demosaicing code=86 8U w=16 xoff=4 data%4=0 : ok, 1ch 16x8, max|cpu-gpu|=69
demosaicing code=86 8U w=16 xoff=5 data%4=1 : EXC misaligned addressSecond run (repro_H2, below: ./repro_H2 <code> <8U|16U> <w>):
// repro_H2_cuda_demosaicing_row_overread.cpp — does cv::cuda::demosaicing read past the end of the row?
// A 4-byte aligned ROI of width w (w = 17: one more than a multiple of 4) is demosaiced twice; between the two
// runs only the parent pixel immediately to the right of the ROI (column w) is changed. usage: repro_H2 <code> <8U|16U> <w>
#include <opencv2/core.hpp>
#include <opencv2/core/cuda.hpp>
#include <opencv2/cudaimgproc.hpp>
#include <cstdio>
#include <cstdlib>
#include <cstring>
int main(int argc, char** argv) {
int code = argc > 1 ? atoi(argv[1]) : 46; int t = (argc > 2 && !strcmp(argv[2], "16U")) ? CV_16UC1 : CV_8UC1; int w = argc > 3 ? atoi(argv[3]) : 17, h = 8;
cv::Mat parent(h, w + 4, t); cv::randu(parent, 0, 255);
cv::Mat out[2];
for (int k = 0; k < 2; k++) {
parent.col(w).setTo(cv::Scalar(k ? 255 : 0)); // only the column just right of the ROI changes
cv::cuda::GpuMat gp(parent), d; cv::cuda::GpuMat roi = gp(cv::Rect(0, 0, w, h));
cv::cuda::demosaicing(roi, d, code); d.download(out[k]);
}
cv::Mat diff; cv::absdiff(out[0], out[1], diff); cv::Mat d1 = diff.reshape(1);
double mx; cv::Point p; cv::minMaxLoc(d1, nullptr, &mx, nullptr, &p);
int cn = out[0].channels(), lastcol = cv::countNonZero(d1.colRange((w - 1) * cn, w * cn)); // changed elements in the last output column
printf("code=%d %s w=%d: output pixels that change when parent column %d (outside the ROI) changes: %d, all in the last column: %s, max change %g at (x=%d,y=%d)\n",
code, t == CV_16UC1 ? "16U" : "8U", w, w, cv::countNonZero(d1), cv::countNonZero(d1) == lastcol ? "yes" : "no", mx, p.x / cn, p.y);
return 0;
}code=46 8U w=16: output pixels that change when parent column 16 (outside the ROI) changes: 0, all in the last column: yes, max change 0 at (x=0,y=0)
code=46 8U w=17: output pixels that change when parent column 17 (outside the ROI) changes: 12, all in the last column: yes, max change 128 at (x=16,y=2)
code=46 8U w=18: output pixels that change when parent column 18 (outside the ROI) changes: 12, all in the last column: yes, max change 128 at (x=17,y=0)
code=46 8U w=19: output pixels that change when parent column 19 (outside the ROI) changes: 12, all in the last column: yes, max change 128 at (x=18,y=0)
code=46 8U w=20: output pixels that change when parent column 20 (outside the ROI) changes: 0, all in the last column: yes, max change 0 at (x=0,y=0)
code=46 8U w=113: output pixels that change when parent column 113 (outside the ROI) changes: 12, all in the last column: yes, max change 128 at (x=112,y=0)
code=46 16U w=16: output pixels that change when parent column 16 (outside the ROI) changes: 0, all in the last column: yes, max change 0 at (x=0,y=0)
code=46 16U w=17: output pixels that change when parent column 17 (outside the ROI) changes: 12, all in the last column: yes, max change 128 at (x=16,y=2)
code=46 16U w=18: output pixels that change when parent column 18 (outside the ROI) changes: 0, all in the last column: yes, max change 0 at (x=0,y=0)
code=46 16U w=113: output pixels that change when parent column 113 (outside the ROI) changes: 12, all in the last column: yes, max change 128 at (x=112,y=0)
code=86 8U w=17: output pixels that change when parent column 17 (outside the ROI) changes: 8, all in the last column: yes, max change 76 at (x=16,y=2)Third run (repro_H3, below: for xo in 0 1 2 3 4 5; do ./repro_H3 $xo cvtColor; done; for xo in 0 1 2 3 4 5; do ./repro_H3 $xo clone; done, one process per trial):
// repro_H3_cuda_cvtColor_bayer_roi_and_clone.cpp — the same Bayer code through cv::cuda::cvtColor on an 8UC1 ROI
// at x offset <xoff> (one process per trial), and, in a separate process, cv::cuda::demosaicing on roi.clone().
// usage: repro_H3 <xoff> <cvtColor|clone>
#include <opencv2/core.hpp>
#include <opencv2/core/cuda.hpp>
#include <opencv2/imgproc.hpp>
#include <opencv2/cudaimgproc.hpp>
#include <cstdio>
#include <cstdlib>
#include <cstring>
int main(int argc, char** argv) {
int xo = argc > 1 ? atoi(argv[1]) : 1; const char* what = argc > 2 ? argv[2] : "cvtColor"; int w = 16, h = 8;
cv::Mat parent(h, w + xo + 3, CV_8UC1); cv::randu(parent, 0, 255);
cv::cuda::GpuMat gparent(parent), d; cv::cuda::GpuMat roi = gparent(cv::Rect(xo, 0, w, h));
printf("%-8s xoff=%d roi.data%%4=%zu : ", what, xo, (size_t)roi.data % 4); fflush(stdout);
try {
if (!strcmp(what, "clone")) { cv::cuda::GpuMat c = roi.clone(); printf("clone.data%%4=%zu ", (size_t)c.data % 4); cv::cuda::demosaicing(c, d, cv::COLOR_BayerBG2BGR); }
else cv::cuda::cvtColor(roi, d, cv::COLOR_BayerBG2BGR);
cv::Mat g; d.download(g); printf("ok %dx%d %dch\n", g.cols, g.rows, g.channels());
} catch (const cv::Exception& e) { std::string s = e.err; for (auto& ch : s) if (ch == '\n') ch = ' '; printf("EXC %s\n", s.c_str()); return 1; }
return 0;
}cvtColor xoff=0 roi.data%4=0 : ok 16x8 3ch
cvtColor xoff=1 roi.data%4=1 : EXC misaligned address
cvtColor xoff=2 roi.data%4=2 : EXC misaligned address
cvtColor xoff=3 roi.data%4=3 : EXC misaligned address
cvtColor xoff=4 roi.data%4=0 : ok 16x8 3ch
cvtColor xoff=5 roi.data%4=1 : EXC misaligned address
clone xoff=0 roi.data%4=0 : clone.data%4=0 ok 16x8 3ch
clone xoff=1 roi.data%4=1 : clone.data%4=0 ok 16x8 3ch
clone xoff=2 roi.data%4=2 : clone.data%4=0 ok 16x8 3ch
clone xoff=3 roi.data%4=3 : clone.data%4=0 ok 16x8 3ch
clone xoff=4 roi.data%4=0 : clone.data%4=0 ok 16x8 3ch
clone xoff=5 roi.data%4=1 : clone.data%4=0 ok 16x8 3chNote: 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