4 files changed, 215 insertions, 4 deletions

diff --git a/dynamic-layers/openembedded-layer/recipes-support/opencv/opencv/0110-core-implement-OpenCL-kernel-for-UMat-diag-to-avoid-.patch b/dynamic-layers/openembedded-layer/recipes-support/opencv/opencv/0110-core-implement-OpenCL-kernel-for-UMat-diag-to-avoid-.patch
new file mode 100644
index 000000000..0b8b90582
--- /dev/null
+++ b/dynamic-layers/openembedded-layer/recipes-support/opencv/opencv/0110-core-implement-OpenCL-kernel-for-UMat-diag-to-avoid-.patch
@@ -0,0 +1,155 @@
+From 60858fce9e0c1a4dc3421e479b5633c140aa2191 Mon Sep 17 00:00:00 2001
+From: jiajia Qian <jiajia.qian@nxp.com>
+Date: Wed, 29 Apr 2026 16:19:00 +0800
+Subject: [PATCH] core: implement OpenCL kernel for UMat::diag to avoid
+ aliasing races
+MIME-Version: 1.0
+Content-Type: text/plain; charset=UTF-8
+Content-Transfer-Encoding: 8bit
+
+The previous OpenCL implementation of UMat::diag() relied on a sequence of
+operations involving buffer initialization followed by copy/transpose on
+an aliased diag() view. Although these operations were enqueued on an
+in-order command queue, this implicitly assumed memory visibility between
+kernels operating on aliased regions of the same buffer.
+
+According to the OpenCL specification, in-order command queues only guarantee
+command scheduling order, while memory visibility between commands is only
+established through explicit command-level synchronization points (e.g.
+events, barriers, or clFinish).Under OpenCL’s relaxed memory model, such
+assumptions are not guaranteed without an explicit command-level synchronization
+point, and can lead to data races and incorrect results, especially for very
+small matrices (e.g. 1x1). The issue is more likely to be exposed on Mesa-based drivers
+when the GPU is running at lower frequencies (e.g. 500 MHz).
+
+This change introduces a dedicated OpenCL kernel to construct the diagonal
+matrix in a single kernel invocation, avoiding intermediate aliasing and
+eliminating the need for implicit ordering assumptions. If the OpenCL path
+is unavailable or unsupported, the implementation transparently falls back
+to the CPU path.
+
+Upstream-Status: Backport [https://github.com/opencv/opencv/commit/60858fce9e0c1a4dc3421e479b5633c140aa2191]
+Signed-off-by: jiajia Qian <jiajia.qian@nxp.com>
+---
+ modules/core/src/opencl/copyset.cl | 44 ++++++++++++++++++++++++++++
+ modules/core/src/umatrix.cpp       | 46 +++++++++++++++++++++++++++++-
+ 2 files changed, 89 insertions(+), 1 deletion(-)
+
+diff --git a/modules/core/src/opencl/copyset.cl b/modules/core/src/opencl/copyset.cl
+index beeccffed8..17aec71b5b 100644
+--- a/modules/core/src/opencl/copyset.cl
++++ b/modules/core/src/opencl/copyset.cl
+@@ -98,6 +98,50 @@ __kernel void copyToMask(__global const uchar * srcptr, int src_step, int src_of
+     }
+ }
+ 
++#elif defined(SET_DIAG)
++
++#ifdef cl_khr_fp64
++#pragma OPENCL EXTENSION cl_khr_fp64 : enable
++#endif
++
++__kernel void setDiag(__global uchar* dst, int dst_step, int dst_offset,
++                      int dst_rows, int dst_cols,
++                      __global const uchar* src, int src_step, int src_offset,
++                      int len)
++{
++    int x = get_global_id(0);
++    int y = get_global_id(1);
++
++    if (x < dst_cols && y < dst_rows)
++    {
++        int elem_size = (int)sizeof(T1) * cn;
++        int dst_idx = mad24(y, dst_step, mad24(x, elem_size, dst_offset));
++
++        if (x == y && x < len)
++        {
++#if IS_ROW_VECTOR
++            int src_idx = src_offset + x * elem_size;
++#else
++            int src_idx = mad24(x, src_step, src_offset);
++#endif
++            __global T1* dst_ptr = (__global T1*)(dst + dst_idx);
++            __global const T1* src_ptr = (__global const T1*)(src + src_idx);
++
++            #pragma unroll
++            for (int c = 0; c < cn; ++c)
++                dst_ptr[c] = src_ptr[c];
++        }
++        else
++        {
++            __global T1* dst_ptr = (__global T1*)(dst + dst_idx);
++
++            #pragma unroll
++            for (int c = 0; c < cn; ++c)
++                dst_ptr[c] = (T1)0;
++        }
++    }
++}
++
+ #else
+ 
+ #ifndef dstST
+diff --git a/modules/core/src/umatrix.cpp b/modules/core/src/umatrix.cpp
+index ceb4c064ec..f8845442d3 100644
+--- a/modules/core/src/umatrix.cpp
++++ b/modules/core/src/umatrix.cpp
+@@ -993,10 +993,54 @@ UMat UMat::reshape(int new_cn, int new_rows) const
+     return hdr;
+ }
+ 
++#ifdef HAVE_OPENCL
++namespace {
++static bool ocl_setDiag(const UMat& d, UMat& m, int len)
++{
++    int cn = d.channels();
++    int depth = d.depth();
++
++    if (depth == CV_64F && !ocl::Device::getDefault().doubleFPConfig())
++        return false;
++
++    String opts = format("-D SET_DIAG -D T1=%s -D cn=%d -D IS_ROW_VECTOR=%d",
++                         ocl::memopTypeToStr(depth),
++                         cn,
++                         (d.rows == 1) ? 1 : 0);
++
++    ocl::Kernel k("setDiag", ocl::core::copyset_oclsrc, opts);
++    if (k.empty())
++        return false;
++
++    k.args(ocl::KernelArg::WriteOnly(m),
++           ocl::KernelArg::ReadOnlyNoSize(d),
++           len);
++
++    size_t globalsize[2] = { (size_t)len, (size_t)len };
++    return k.run(2, globalsize, NULL, false);
++}
++}
++#endif
++
+ UMat UMat::diag(const UMat& d, UMatUsageFlags usageFlags)
+ {
+-    CV_Assert( d.cols == 1 || d.rows == 1 );
++    CV_INSTRUMENT_REGION();
++
++    CV_Assert(d.cols == 1 || d.rows == 1);
+     int len = d.rows + d.cols - 1;
++
++#ifdef HAVE_OPENCL
++    if (ocl::useOpenCL())
++    {
++        UMat m(len, len, d.type(), usageFlags);
++        if (ocl_setDiag(d, m, len))
++        {
++            CV_IMPL_ADD(CV_IMPL_OCL);
++            return m;
++        }
++    }
++#endif
++
+     UMat m(len, len, d.type(), Scalar(0), usageFlags);
+     UMat md = m.diag();
+     if( d.cols == 1 )
+-- 
+2.34.1
+
diff --git a/dynamic-layers/openembedded-layer/recipes-support/opencv/opencv/0111-lF-18405-opencv-example_tapi_clahe-sample-failed-wit.patch b/dynamic-layers/openembedded-layer/recipes-support/opencv/opencv/0111-lF-18405-opencv-example_tapi_clahe-sample-failed-wit.patch
new file mode 100644
index 000000000..d6e1c1e3d
--- /dev/null
+++ b/dynamic-layers/openembedded-layer/recipes-support/opencv/opencv/0111-lF-18405-opencv-example_tapi_clahe-sample-failed-wit.patch
@@ -0,0 +1,54 @@
+From 19be131ad53f835a2b569b05556bc4a12aa415c8 Mon Sep 17 00:00:00 2001
+From: Chingling Wang <chingling.wang@nxp.com>
+Date: Tue, 4 Aug 2026 15:55:15 -0700
+Subject: [PATCH] lF-18405 opencv example_tapi_clahe sample failed with
+ segmentation fault for imx8 platform
+
+gpu viv is asynchonous SVM mapping. imread() allocated memory is temp, may be deleted before mapping to
+GPU memory. need a temp Mat src to store the image data. Otheweise, may cause memory corruption leading to
+segmenation fault.
+
+Upstream-Status: Pending
+Signed-off-by: Chingling Wang <chingling.wang@nxp.com>
+---
+ samples/tapi/clahe.cpp | 10 +++++++---
+ 1 file changed, 7 insertions(+), 3 deletions(-)
+
+diff --git a/samples/tapi/clahe.cpp b/samples/tapi/clahe.cpp
+index ebef65f4e4..406102fd8d 100644
+--- a/samples/tapi/clahe.cpp
++++ b/samples/tapi/clahe.cpp
+@@ -52,6 +52,7 @@ int main(int argc, char** argv)
+     createTrackbar("Tile Size", "CLAHE", &tilesize, 32, (TrackbarCallback)TSize_Callback);
+     createTrackbar("Clip Limit", "CLAHE", &cliplimit, 20, (TrackbarCallback)Clip_Callback);
+ 
++    Mat src;
+     UMat frame, outframe;
+ 
+     int cur_clip;
+@@ -66,7 +67,8 @@ int main(int argc, char** argv)
+     if(!infile.empty())
+     {
+         infile = samples::findFile(infile);
+-        imread(infile).copyTo(frame);
++        src = imread(infile);
++        src.copyTo(frame);
+         if(frame.empty())
+         {
+             cout << "error read image: " << infile << endl;
+@@ -85,8 +87,10 @@ int main(int argc, char** argv)
+     {
+         if(capture.isOpened())
+             capture.read(frame);
+-        else
+-            imread(infile).copyTo(frame);
++        else{
++            src = imread(infile);
++            src.copyTo(frame);
++        }
+         if(frame.empty())
+         {
+             waitKey();
+-- 
+2.43.0
+
diff --git a/dynamic-layers/openembedded-layer/recipes-support/opencv/opencv_4.13.0.imx.bb b/dynamic-layers/openembedded-layer/recipes-support/opencv/opencv_4.13.0.imx.bb
index 5ed168bbe..2fad969a8 100644
--- a/dynamic-layers/openembedded-layer/recipes-support/opencv/opencv_4.13.0.imx.bb
+++ b/dynamic-layers/openembedded-layer/recipes-support/opencv/opencv_4.13.0.imx.bb
@@ -260,10 +260,10 @@ do_install:append() {
         rm -rf ${D}${bindir}/setup_vars_opencv4.sh
     fi
 
-    for fn in arithm.vsx3.cpp convert.vsx3.cpp; do
-        if [ -f ${B}/modules/core/$fn ]; then
-            sed -i -e 's,${S},/usr/src/debug/${PN}/${PV},g' ${B}/modules/core/$fn
-        fi
+    # Backported from meta-openembedded fdfca71954 (newer than the header's
+    # Upstream hash).
+    find ${B}/modules -name "*.cpp" | xargs -r grep -l "${S}" | while read -r fn; do
+        sed -i -e 's,${S},/usr/src/debug/${PN}/${PV},g' "$fn"
     done
 
     if [ -d ${D}/${bindir}/shape ]; then
diff --git a/dynamic-layers/openembedded-layer/recipes-support/opencv/opencv_4.13.0.imx.inc b/dynamic-layers/openembedded-layer/recipes-support/opencv/opencv_4.13.0.imx.inc
index 97803c8e5..95390c5aa 100644
--- a/dynamic-layers/openembedded-layer/recipes-support/opencv/opencv_4.13.0.imx.inc
+++ b/dynamic-layers/openembedded-layer/recipes-support/opencv/opencv_4.13.0.imx.inc
@@ -32,6 +32,8 @@ SRC_URI += "\
     file://0107-imgproc-perf-HoughLines-Fix-test-tolerance-and-compa.patch \
     file://0108-imgproc-perf-HoughLines-Fix-lower-bound-for-line-cou.patch \
     file://0109-core-ocl-fix-incorrect-results-for-in-place-flip-on-.patch \
+    file://0110-core-implement-OpenCL-kernel-for-UMat-diag-to-avoid-.patch \
+    file://0111-lF-18405-opencv-example_tapi_clahe-sample-failed-wit.patch \
 "
 
 # Add opencv_extra