Merge pull request #29546 from uwezkhan/dtree-split-index-bound

bound categorical split indices in dtrees readSplit

README.md

@@ -7,7 +7,7 @@
   * Courses: <https://opencv.org/courses>
 * Docs: <https://docs.opencv.org/4.x/>
 * Q&A forum: <https://forum.opencv.org>
-  * previous forum (read only): <http://answers.opencv.org>
+  * previous forum (read only): <https://answers.opencv.org>
 * Issue tracking: <https://github.com/opencv/opencv/issues>
 * Additional OpenCV functionality: <https://github.com/opencv/opencv_contrib>
 * Donate to OpenCV: <https://opencv.org/support/>
@@ -28,9 +28,9 @@ Please read the [contribution guidelines](https://github.com/opencv/opencv/wiki/
 ### Additional Resources
 
 * [Submit your OpenCV-based project](https://form.jotform.com/233105358823151) for inclusion in Community Friday on opencv.org
-* [Subscribe to the OpenCV YouTube Channel](http://youtube.com/@opencvofficial) featuring OpenCV Live, an hour-long streaming show
-* [Follow OpenCV on LinkedIn](http://linkedin.com/company/opencv/) for daily posts showing the state-of-the-art in computer vision & AI
+* [Subscribe to the OpenCV YouTube Channel](https://youtube.com/@opencvofficial) featuring OpenCV Live, an hour-long streaming show
+* [Follow OpenCV on LinkedIn](https://linkedin.com/company/opencv/) for daily posts showing the state-of-the-art in computer vision & AI
 * [Apply to be an OpenCV Volunteer](https://form.jotform.com/232745316792159) to help organize events and online campaigns as well as amplify them
-* [Follow OpenCV on Mastodon](http://mastodon.social/@opencv) in the Fediverse
+* [Follow OpenCV on Mastodon](https://mastodon.social/@opencv) in the Fediverse
 * [Follow OpenCV on Twitter](https://twitter.com/opencvlive)
 * [OpenCV.ai](https://opencv.ai): Computer Vision and AI development services from the OpenCV team.

hal/ipp/CMakeLists.txt

@@ -46,6 +46,13 @@ if(WITH_IPP_CALLS_ENFORCED)
   message("WITH_IPP_CALLS_ENFORCED=${WITH_IPP_CALLS_ENFORCED}: enforced IPP calls are enabled in IPP HAL")
 endif()
 
+# Enable cv::instr tracing of IPP calls inside the HAL so they show up in the
+# instrumentation trace / IPP weight. Reuses the engine exported from
+# opencv_core (symbols resolved at link time as ipphal is archived into core).
+if(ENABLE_INSTRUMENTATION)
+  target_compile_definitions(ipphal PRIVATE ENABLE_INSTRUMENTATION)
+endif()
+
 target_include_directories(ipphal PRIVATE "${CMAKE_CURRENT_SOURCE_DIR}/include")
 ocv_warnings_disable(CMAKE_CXX_FLAGS -Wno-suggest-override)
 

hal/ipp/include/ipp_utils.hpp

@@ -1,6 +1,7 @@
 // This file is part of OpenCV project.
 // It is subject to the license terms in the LICENSE file found in the top-level directory
 // of this distribution and at http://opencv.org/license.html
+// Copyright (C) 2026, Intel Corporation, all rights reserved.
 
 #ifndef __IPP_HAL_UTILS_HPP__
 #define __IPP_HAL_UTILS_HPP__
@@ -23,7 +24,22 @@
 # include "ipp.h"
 #endif
 
+// IPP call instrumentation for the standalone IPP HAL.
+//
+// The HAL is built as a separate static library that does not include the
+// internal core/private.hpp, so by default its IPP calls never appear in the
+// cv::instr trace (reported IPP weight = 0%). When the parent build enables
+// ENABLE_INSTRUMENTATION, the shared instrumentation.private.hpp provides the
+// exact same CV_INSTRUMENT_FUN_IPP macro used by core (backed by symbols exported
+// from libopencv_core, resolved at link time), so HAL IPP calls are traced
+// identically to in-module ones. Otherwise the no-op passthrough is used.
+#if !defined(CV_INSTRUMENT_FUN_IPP)
+#if defined(ENABLE_INSTRUMENTATION)
+#include "opencv2/core/utils/instrumentation.private.hpp"
+#else
 #define CV_INSTRUMENT_FUN_IPP(FUN, ...) ((FUN)(__VA_ARGS__))
+#endif // defined(ENABLE_INSTRUMENTATION)
+#endif // !defined(CV_INSTRUMENT_FUN_IPP)
 
 #define CV_HAL_CHECK_USE_IPP() if(!cv::ipp::useIPP()) return CV_HAL_ERROR_NOT_IMPLEMENTED;
 

hal/riscv-rvv/include/core.hpp

@@ -99,9 +99,18 @@ int fast_atan_64(const double* y, const double* x, double* dst, size_t n, bool a
 /* ############ split ############ */
 
 int split8u(const uchar* src, uchar** dst, int len, int cn);
+int split16u(const ushort* src, ushort** dst, int len, int cn);
+int split32s(const int* src, int** dst, int len, int cn);
+int split64s(const int64* src, int64** dst, int len, int cn);
 
 #undef cv_hal_split8u
 #define cv_hal_split8u cv::rvv_hal::core::split8u
+#undef cv_hal_split16u
+#define cv_hal_split16u cv::rvv_hal::core::split16u
+#undef cv_hal_split32s
+#define cv_hal_split32s cv::rvv_hal::core::split32s
+#undef cv_hal_split64s
+#define cv_hal_split64s cv::rvv_hal::core::split64s
 
 /* ############ sqrt ############ */
 

hal/riscv-rvv/src/core/copy_mask.cpp

@@ -30,8 +30,27 @@ static int copyToMasked_e##X##c1(const uchar *src_data, size_t src_step, const u
     return CV_HAL_ERROR_OK; \
 }
 
-CV_HAL_RVV_COPY_MASK_eXc1(8,  8)
-CV_HAL_RVV_COPY_MASK_eXc1(16, 4)
+#define CV_HAL_RVV_COPY_MASK_MERGE_eXc1(X, data_lmul, mask_lmul) \
+static int copyToMasked_e##X##c1(const uchar *src_data, size_t src_step, const uchar *mask_data, size_t mask_step, \
+                                 uchar *dst_data, size_t dst_step, int width, int height) { \
+    for (; height--; mask_data += mask_step, src_data += src_step, dst_data += dst_step) { \
+        const uint##X##_t *src = (const uint##X##_t*)src_data; \
+        uint##X##_t *dst = (uint##X##_t*)dst_data; \
+        int vl; \
+        for (int i = 0; i < width; i += vl) { \
+            vl = __riscv_vsetvl_e##X##m##data_lmul(width - i); \
+            auto m = __riscv_vmsne(__riscv_vle8_v_u8m##mask_lmul(mask_data + i, vl), 0, vl); \
+            auto src_v = __riscv_vle##X##_v_u##X##m##data_lmul(src + i, vl); \
+            auto dst_v = __riscv_vle##X##_v_u##X##m##data_lmul(dst + i, vl); \
+            __riscv_vse##X##_v_u##X##m##data_lmul(dst + i, __riscv_vmerge(dst_v, src_v, m, vl), vl); \
+        } \
+    } \
+    return CV_HAL_ERROR_OK; \
+}
+
+CV_HAL_RVV_COPY_MASK_MERGE_eXc1(8,  8, 8)
+CV_HAL_RVV_COPY_MASK_MERGE_eXc1(16, 8, 4)
+
 CV_HAL_RVV_COPY_MASK_eXc1(32, 2)
 CV_HAL_RVV_COPY_MASK_eXc1(64, 1)
 

hal/riscv-rvv/src/core/merge.cpp

@@ -28,6 +28,12 @@
 #define __riscv_vcreate_v_u16m4x2(v0, v1) OPENCV_HAL_IMPL_RVV_VCREATE_x2(u16, 4, v0, v1)
 #define __riscv_vcreate_v_u16m2x3(v0, v1, v2) OPENCV_HAL_IMPL_RVV_VCREATE_x3(u16, 2, v0, v1, v2)
 #define __riscv_vcreate_v_u16m2x4(v0, v1, v2, v3) OPENCV_HAL_IMPL_RVV_VCREATE_x4(u16, 2, v0, v1, v2, v3)
+#define __riscv_vcreate_v_i32m4x2(v0, v1) OPENCV_HAL_IMPL_RVV_VCREATE_x2(i32, 4, v0, v1)
+#define __riscv_vcreate_v_i32m2x3(v0, v1, v2) OPENCV_HAL_IMPL_RVV_VCREATE_x3(i32, 2, v0, v1, v2)
+#define __riscv_vcreate_v_i32m2x4(v0, v1, v2, v3) OPENCV_HAL_IMPL_RVV_VCREATE_x4(i32, 2, v0, v1, v2, v3)
+#define __riscv_vcreate_v_i64m4x2(v0, v1) OPENCV_HAL_IMPL_RVV_VCREATE_x2(i64, 4, v0, v1)
+#define __riscv_vcreate_v_i64m2x3(v0, v1, v2) OPENCV_HAL_IMPL_RVV_VCREATE_x3(i64, 2, v0, v1, v2)
+#define __riscv_vcreate_v_i64m2x4(v0, v1, v2, v3) OPENCV_HAL_IMPL_RVV_VCREATE_x4(i64, 2, v0, v1, v2, v3)
 #endif  // clang < 18
 
 namespace cv { namespace rvv_hal { namespace core {
@@ -204,146 +210,177 @@ int merge16u(const ushort** src, ushort* dst, int len, int cn ) {
     return CV_HAL_ERROR_OK;
 }
 
-#if defined __GNUC__ && !defined(__clang__)
-__attribute__((optimize("no-tree-vectorize")))
-#endif
 int merge32s(const int** src, int* dst, int len, int cn ) {
-    int k = cn % 4 ? cn % 4 : 4;
-    int i, j;
-    if( k == 1 )
+    int vl = 0;
+    if (cn == 1)
     {
         const int* src0 = src[0];
-        #if defined(__clang__)
-        #pragma clang loop vectorize(disable)
-        #endif
-        for( i = j = 0; i < len; i++, j += cn )
-            dst[j] = src0[i];
-    }
-    else if( k == 2 )
-    {
-        const int *src0 = src[0], *src1 = src[1];
-        i = j = 0;
-        #if defined(__clang__)
-        #pragma clang loop vectorize(disable)
-        #endif
-        for( ; i < len; i++, j += cn )
+        for (int i = 0; i < len; i += vl)
         {
-            dst[j] = src0[i];
-            dst[j+1] = src1[i];
+            vl = __riscv_vsetvl_e32m8(len - i);
+            __riscv_vse32_v_i32m8(dst + i, __riscv_vle32_v_i32m8(src0 + i, vl), vl);
         }
     }
-    else if( k == 3 )
+    else if (cn == 2)
+    {
+        const int *src0 = src[0], *src1 = src[1];
+        for (int i = 0; i < len; i += vl)
+        {
+            vl = __riscv_vsetvl_e32m4(len - i);
+            vint32m4x2_t seg = __riscv_vcreate_v_i32m4x2(
+                __riscv_vle32_v_i32m4(src0 + i, vl),
+                __riscv_vle32_v_i32m4(src1 + i, vl)
+            );
+            __riscv_vsseg2e32_v_i32m4x2(dst + i * cn, seg, vl);
+        }
+    }
+    else if (cn == 3)
     {
         const int *src0 = src[0], *src1 = src[1], *src2 = src[2];
-        i = j = 0;
-        #if defined(__clang__)
-        #pragma clang loop vectorize(disable)
-        #endif
-        for( ; i < len; i++, j += cn )
+        for (int i = 0; i < len; i += vl)
         {
-            dst[j] = src0[i];
-            dst[j+1] = src1[i];
-            dst[j+2] = src2[i];
+            vl = __riscv_vsetvl_e32m2(len - i);
+            vint32m2x3_t seg = __riscv_vcreate_v_i32m2x3(
+                __riscv_vle32_v_i32m2(src0 + i, vl),
+                __riscv_vle32_v_i32m2(src1 + i, vl),
+                __riscv_vle32_v_i32m2(src2 + i, vl)
+            );
+            __riscv_vsseg3e32_v_i32m2x3(dst + i * cn, seg, vl);
+        }
+    }
+    else if (cn == 4)
+    {
+        const int *src0 = src[0], *src1 = src[1], *src2 = src[2], *src3 = src[3];
+        for (int i = 0; i < len; i += vl)
+        {
+            vl = __riscv_vsetvl_e32m2(len - i);
+            vint32m2x4_t seg = __riscv_vcreate_v_i32m2x4(
+                __riscv_vle32_v_i32m2(src0 + i, vl),
+                __riscv_vle32_v_i32m2(src1 + i, vl),
+                __riscv_vle32_v_i32m2(src2 + i, vl),
+                __riscv_vle32_v_i32m2(src3 + i, vl)
+            );
+            __riscv_vsseg4e32_v_i32m2x4(dst + i * cn, seg, vl);
         }
     }
     else
     {
-        const int *src0 = src[0], *src1 = src[1], *src2 = src[2], *src3 = src[3];
-        i = j = 0;
-        #if defined(__clang__)
-        #pragma clang loop vectorize(disable)
-        #endif
-        for( ; i < len; i++, j += cn )
+        int k = 0;
+        for (; k <= cn - 4; k += 4)
         {
-            dst[j] = src0[i]; dst[j+1] = src1[i];
-            dst[j+2] = src2[i]; dst[j+3] = src3[i];
+            const int *src0 = src[k], *src1 = src[k + 1], *src2 = src[k + 2], *src3 = src[k + 3];
+            for (int i = 0; i < len; i += vl)
+            {
+                vl = __riscv_vsetvl_e32m2(len - i);
+                vint32m2x4_t seg = __riscv_vcreate_v_i32m2x4(
+                    __riscv_vle32_v_i32m2(src0 + i, vl),
+                    __riscv_vle32_v_i32m2(src1 + i, vl),
+                    __riscv_vle32_v_i32m2(src2 + i, vl),
+                    __riscv_vle32_v_i32m2(src3 + i, vl)
+                );
+                __riscv_vssseg4e32_v_i32m2x4(dst + k + i * cn, cn * sizeof(int), seg, vl);
+            }
         }
-    }
-    #if defined(__clang__)
-    #pragma clang loop vectorize(disable)
-    #endif
-    for( ; k < cn; k += 4 )
-    {
-        const int *src0 = src[k], *src1 = src[k+1], *src2 = src[k+2], *src3 = src[k+3];
-        for( i = 0, j = k; i < len; i++, j += cn )
+        for (; k < cn; ++k)
         {
-            dst[j] = src0[i]; dst[j+1] = src1[i];
-            dst[j+2] = src2[i]; dst[j+3] = src3[i];
+            const int* srcK = src[k];
+            for (int i = 0; i < len; i += vl)
+            {
+                vl = __riscv_vsetvl_e32m2(len - i);
+                vint32m2_t seg = __riscv_vle32_v_i32m2(srcK + i, vl);
+                __riscv_vsse32_v_i32m2(dst + k + i * cn, cn * sizeof(int), seg, vl);
+            }
         }
     }
     return CV_HAL_ERROR_OK;
 }
 
-#if defined __GNUC__ && !defined(__clang__)
-__attribute__((optimize("no-tree-vectorize")))
-#endif
 int merge64s(const int64** src, int64* dst, int len, int cn ) {
-    int k = cn % 4 ? cn % 4 : 4;
-    int i, j;
-    if( k == 1 )
+    int vl = 0;
+    if (cn == 1)
     {
         const int64* src0 = src[0];
-        #if defined(__clang__)
-        #pragma clang loop vectorize(disable)
-        #endif
-        for( i = j = 0; i < len; i++, j += cn )
-            dst[j] = src0[i];
-    }
-    else if( k == 2 )
-    {
-        const int64 *src0 = src[0], *src1 = src[1];
-        i = j = 0;
-        #if defined(__clang__)
-        #pragma clang loop vectorize(disable)
-        #endif
-        for( ; i < len; i++, j += cn )
+        for (int i = 0; i < len; i += vl)
         {
-            dst[j] = src0[i];
-            dst[j+1] = src1[i];
+            vl = __riscv_vsetvl_e64m8(len - i);
+            __riscv_vse64_v_i64m8(dst + i, __riscv_vle64_v_i64m8(src0 + i, vl), vl);
         }
     }
-    else if( k == 3 )
+    else if (cn == 2)
+    {
+        const int64 *src0 = src[0], *src1 = src[1];
+        for (int i = 0; i < len; i += vl)
+        {
+            vl = __riscv_vsetvl_e64m4(len - i);
+            vint64m4x2_t seg = __riscv_vcreate_v_i64m4x2(
+                __riscv_vle64_v_i64m4(src0 + i, vl),
+                __riscv_vle64_v_i64m4(src1 + i, vl)
+            );
+            __riscv_vsseg2e64_v_i64m4x2(dst + i * cn, seg, vl);
+        }
+    }
+    else if (cn == 3)
     {
         const int64 *src0 = src[0], *src1 = src[1], *src2 = src[2];
-        i = j = 0;
-        #if defined(__clang__)
-        #pragma clang loop vectorize(disable)
-        #endif
-        for( ; i < len; i++, j += cn )
+        for (int i = 0; i < len; i += vl)
         {
-            dst[j] = src0[i];
-            dst[j+1] = src1[i];
-            dst[j+2] = src2[i];
+            vl = __riscv_vsetvl_e64m2(len - i);
+            vint64m2x3_t seg = __riscv_vcreate_v_i64m2x3(
+                __riscv_vle64_v_i64m2(src0 + i, vl),
+                __riscv_vle64_v_i64m2(src1 + i, vl),
+                __riscv_vle64_v_i64m2(src2 + i, vl)
+            );
+            __riscv_vsseg3e64_v_i64m2x3(dst + i * cn, seg, vl);
+        }
+    }
+    else if (cn == 4)
+    {
+        const int64 *src0 = src[0], *src1 = src[1], *src2 = src[2], *src3 = src[3];
+        for (int i = 0; i < len; i += vl)
+        {
+            vl = __riscv_vsetvl_e64m2(len - i);
+            vint64m2x4_t seg = __riscv_vcreate_v_i64m2x4(
+                __riscv_vle64_v_i64m2(src0 + i, vl),
+                __riscv_vle64_v_i64m2(src1 + i, vl),
+                __riscv_vle64_v_i64m2(src2 + i, vl),
+                __riscv_vle64_v_i64m2(src3 + i, vl)
+            );
+            __riscv_vsseg4e64_v_i64m2x4(dst + i * cn, seg, vl);
         }
     }
     else
     {
-        const int64 *src0 = src[0], *src1 = src[1], *src2 = src[2], *src3 = src[3];
-        i = j = 0;
-        #if defined(__clang__)
-        #pragma clang loop vectorize(disable)
-        #endif
-        for( ; i < len; i++, j += cn )
+        int k = 0;
+        for (; k <= cn - 4; k += 4)
         {
-            dst[j] = src0[i]; dst[j+1] = src1[i];
-            dst[j+2] = src2[i]; dst[j+3] = src3[i];
+            const int64 *src0 = src[k], *src1 = src[k + 1], *src2 = src[k + 2], *src3 = src[k + 3];
+            for (int i = 0; i < len; i += vl)
+            {
+                vl = __riscv_vsetvl_e64m2(len - i);
+                vint64m2x4_t seg = __riscv_vcreate_v_i64m2x4(
+                    __riscv_vle64_v_i64m2(src0 + i, vl),
+                    __riscv_vle64_v_i64m2(src1 + i, vl),
+                    __riscv_vle64_v_i64m2(src2 + i, vl),
+                    __riscv_vle64_v_i64m2(src3 + i, vl)
+                );
+                __riscv_vssseg4e64_v_i64m2x4(dst + k + i * cn, cn * sizeof(int64), seg, vl);
+            }
         }
-    }
-    #if defined(__clang__)
-    #pragma clang loop vectorize(disable)
-    #endif
-    for( ; k < cn; k += 4 )
-    {
-        const int64 *src0 = src[k], *src1 = src[k+1], *src2 = src[k+2], *src3 = src[k+3];
-        for( i = 0, j = k; i < len; i++, j += cn )
+        for (; k < cn; ++k)
         {
-            dst[j] = src0[i]; dst[j+1] = src1[i];
-            dst[j+2] = src2[i]; dst[j+3] = src3[i];
+            const int64* srcK = src[k];
+            for (int i = 0; i < len; i += vl)
+            {
+                vl = __riscv_vsetvl_e64m2(len - i);
+                vint64m2_t seg = __riscv_vle64_v_i64m2(srcK + i, vl);
+                __riscv_vsse64_v_i64m2(dst + k + i * cn, cn * sizeof(int64), seg, vl);
+            }
         }
     }
     return CV_HAL_ERROR_OK;
 }
 
+
 #endif // CV_HAL_RVV_1P0_ENABLED
 
 }}} // cv::rvv_hal::core

hal/riscv-rvv/src/core/split.cpp

@@ -8,84 +8,92 @@ namespace cv { namespace rvv_hal { namespace core {
 
 #if CV_HAL_RVV_1P0_ENABLED
 
-int split8u(const uchar* src, uchar** dst, int len, int cn)
-{
-    int vl = 0;
-    if (cn == 1)
-    {
-        uchar* dst0 = dst[0];
-        for (int i = 0; i < len; i += vl)
-        {
-            vl = __riscv_vsetvl_e8m8(len - i);
-            __riscv_vse8_v_u8m8(dst0 + i, __riscv_vle8_v_u8m8(src + i, vl), vl);
-        }
-    }
-    else if (cn == 2)
-    {
-        uchar *dst0 = dst[0], *dst1 = dst[1];
-        for (int i = 0; i < len; i += vl)
-        {
-            vl = __riscv_vsetvl_e8m4(len - i);
-            vuint8m4x2_t seg = __riscv_vlseg2e8_v_u8m4x2(src + i * cn, vl);
-            __riscv_vse8_v_u8m4(dst0 + i, __riscv_vget_v_u8m4x2_u8m4(seg, 0), vl);
-            __riscv_vse8_v_u8m4(dst1 + i, __riscv_vget_v_u8m4x2_u8m4(seg, 1), vl);
-        }
-    }
-    else if (cn == 3)
-    {
-        uchar *dst0 = dst[0], *dst1 = dst[1], *dst2 = dst[2];
-        for (int i = 0; i < len; i += vl)
-        {
-            vl = __riscv_vsetvl_e8m2(len - i);
-            vuint8m2x3_t seg = __riscv_vlseg3e8_v_u8m2x3(src + i * cn, vl);
-            __riscv_vse8_v_u8m2(dst0 + i, __riscv_vget_v_u8m2x3_u8m2(seg, 0), vl);
-            __riscv_vse8_v_u8m2(dst1 + i, __riscv_vget_v_u8m2x3_u8m2(seg, 1), vl);
-            __riscv_vse8_v_u8m2(dst2 + i, __riscv_vget_v_u8m2x3_u8m2(seg, 2), vl);
-        }
-    }
-    else if (cn == 4)
-    {
-        uchar *dst0 = dst[0], *dst1 = dst[1], *dst2 = dst[2], *dst3 = dst[3];
-        for (int i = 0; i < len; i += vl)
-        {
-            vl = __riscv_vsetvl_e8m2(len - i);
-            vuint8m2x4_t seg = __riscv_vlseg4e8_v_u8m2x4(src + i * cn, vl);
-            __riscv_vse8_v_u8m2(dst0 + i, __riscv_vget_v_u8m2x4_u8m2(seg, 0), vl);
-            __riscv_vse8_v_u8m2(dst1 + i, __riscv_vget_v_u8m2x4_u8m2(seg, 1), vl);
-            __riscv_vse8_v_u8m2(dst2 + i, __riscv_vget_v_u8m2x4_u8m2(seg, 2), vl);
-            __riscv_vse8_v_u8m2(dst3 + i, __riscv_vget_v_u8m2x4_u8m2(seg, 3), vl);
-        }
-    }
-    else
-    {
-        int k = 0;
-        for (; k <= cn - 4; k += 4)
-        {
-            uchar *dst0 = dst[k], *dst1 = dst[k + 1], *dst2 = dst[k + 2], *dst3 = dst[k + 3];
-            for (int i = 0; i < len; i += vl)
-            {
-                vl = __riscv_vsetvl_e8m2(len - i);
-                vuint8m2x4_t seg = __riscv_vlsseg4e8_v_u8m2x4(src + k + i * cn, cn, vl);
-                __riscv_vse8_v_u8m2(dst0 + i, __riscv_vget_v_u8m2x4_u8m2(seg, 0), vl);
-                __riscv_vse8_v_u8m2(dst1 + i, __riscv_vget_v_u8m2x4_u8m2(seg, 1), vl);
-                __riscv_vse8_v_u8m2(dst2 + i, __riscv_vget_v_u8m2x4_u8m2(seg, 2), vl);
-                __riscv_vse8_v_u8m2(dst3 + i, __riscv_vget_v_u8m2x4_u8m2(seg, 3), vl);
-            }
-        }
-        for (; k < cn; ++k)
-        {
-            uchar* dstK = dst[k];
-            for (int i = 0; i < len; i += vl)
-            {
-                vl = __riscv_vsetvl_e8m2(len - i);
-                vuint8m2_t seg = __riscv_vlse8_v_u8m2(src + k + i * cn, cn, vl);
-                __riscv_vse8_v_u8m2(dstK + i, seg, vl);
-            }
-        }
-    }
-    return CV_HAL_ERROR_OK;
+#define OPENCV_HAL_IMPL_RVV_SPLIT(func, T, vtype, suffix, width) \
+int func(const T* src, T** dst, int len, int cn) \
+{ \
+    int vl = 0; \
+    if (cn == 1) \
+    { \
+        T* dst0 = dst[0]; \
+        for (int i = 0; i < len; i += vl) \
+        { \
+            vl = __riscv_vsetvl_e##width##m8(len - i); \
+            __riscv_vse##width##_v_##suffix##m8(dst0 + i, __riscv_vle##width##_v_##suffix##m8(src + i, vl), vl); \
+        } \
+    } \
+    else if (cn == 2) \
+    { \
+        T *dst0 = dst[0], *dst1 = dst[1]; \
+        for (int i = 0; i < len; i += vl) \
+        { \
+            vl = __riscv_vsetvl_e##width##m4(len - i); \
+            v##vtype##m4x2_t seg = __riscv_vlseg2e##width##_v_##suffix##m4x2(src + i * cn, vl); \
+            __riscv_vse##width##_v_##suffix##m4(dst0 + i, __riscv_vget_v_##suffix##m4x2_##suffix##m4(seg, 0), vl); \
+            __riscv_vse##width##_v_##suffix##m4(dst1 + i, __riscv_vget_v_##suffix##m4x2_##suffix##m4(seg, 1), vl); \
+        } \
+    } \
+    else if (cn == 3) \
+    { \
+        T *dst0 = dst[0], *dst1 = dst[1], *dst2 = dst[2]; \
+        for (int i = 0; i < len; i += vl) \
+        { \
+            vl = __riscv_vsetvl_e##width##m2(len - i); \
+            v##vtype##m2x3_t seg = __riscv_vlseg3e##width##_v_##suffix##m2x3(src + i * cn, vl); \
+            __riscv_vse##width##_v_##suffix##m2(dst0 + i, __riscv_vget_v_##suffix##m2x3_##suffix##m2(seg, 0), vl); \
+            __riscv_vse##width##_v_##suffix##m2(dst1 + i, __riscv_vget_v_##suffix##m2x3_##suffix##m2(seg, 1), vl); \
+            __riscv_vse##width##_v_##suffix##m2(dst2 + i, __riscv_vget_v_##suffix##m2x3_##suffix##m2(seg, 2), vl); \
+        } \
+    } \
+    else if (cn == 4) \
+    { \
+        T *dst0 = dst[0], *dst1 = dst[1], *dst2 = dst[2], *dst3 = dst[3]; \
+        for (int i = 0; i < len; i += vl) \
+        { \
+            vl = __riscv_vsetvl_e##width##m2(len - i); \
+            v##vtype##m2x4_t seg = __riscv_vlseg4e##width##_v_##suffix##m2x4(src + i * cn, vl); \
+            __riscv_vse##width##_v_##suffix##m2(dst0 + i, __riscv_vget_v_##suffix##m2x4_##suffix##m2(seg, 0), vl); \
+            __riscv_vse##width##_v_##suffix##m2(dst1 + i, __riscv_vget_v_##suffix##m2x4_##suffix##m2(seg, 1), vl); \
+            __riscv_vse##width##_v_##suffix##m2(dst2 + i, __riscv_vget_v_##suffix##m2x4_##suffix##m2(seg, 2), vl); \
+            __riscv_vse##width##_v_##suffix##m2(dst3 + i, __riscv_vget_v_##suffix##m2x4_##suffix##m2(seg, 3), vl); \
+        } \
+    } \
+    else \
+    { \
+        int k = 0; \
+        for (; k <= cn - 4; k += 4) \
+        { \
+            T *dst0 = dst[k], *dst1 = dst[k + 1], *dst2 = dst[k + 2], *dst3 = dst[k + 3]; \
+            for (int i = 0; i < len; i += vl) \
+            { \
+                vl = __riscv_vsetvl_e##width##m2(len - i); \
+                v##vtype##m2x4_t seg = __riscv_vlsseg4e##width##_v_##suffix##m2x4(src + k + i * cn, cn * sizeof(T), vl); \
+                __riscv_vse##width##_v_##suffix##m2(dst0 + i, __riscv_vget_v_##suffix##m2x4_##suffix##m2(seg, 0), vl); \
+                __riscv_vse##width##_v_##suffix##m2(dst1 + i, __riscv_vget_v_##suffix##m2x4_##suffix##m2(seg, 1), vl); \
+                __riscv_vse##width##_v_##suffix##m2(dst2 + i, __riscv_vget_v_##suffix##m2x4_##suffix##m2(seg, 2), vl); \
+                __riscv_vse##width##_v_##suffix##m2(dst3 + i, __riscv_vget_v_##suffix##m2x4_##suffix##m2(seg, 3), vl); \
+            } \
+        } \
+        for (; k < cn; ++k) \
+        { \
+            T* dstK = dst[k]; \
+            for (int i = 0; i < len; i += vl) \
+            { \
+                vl = __riscv_vsetvl_e##width##m2(len - i); \
+                v##vtype##m2_t seg = __riscv_vlse##width##_v_##suffix##m2(src + k + i * cn, cn * sizeof(T), vl); \
+                __riscv_vse##width##_v_##suffix##m2(dstK + i, seg, vl); \
+            } \
+        } \
+    } \
+    return CV_HAL_ERROR_OK; \
 }
 
+OPENCV_HAL_IMPL_RVV_SPLIT(split8u, uchar, uint8, u8, 8)
+OPENCV_HAL_IMPL_RVV_SPLIT(split16u, ushort, uint16, u16, 16)
+OPENCV_HAL_IMPL_RVV_SPLIT(split32s, int, int32, i32, 32)
+OPENCV_HAL_IMPL_RVV_SPLIT(split64s, int64, int64, i64, 64)
+
+#undef OPENCV_HAL_IMPL_RVV_SPLIT
+
 #endif // CV_HAL_RVV_1P0_ENABLED
 
 }}} // cv::rvv_hal::core

modules/core/include/opencv2/core/private.hpp

@@ -690,23 +690,13 @@ public:
     TLSData<InstrTLSStruct> tlsStruct;
 };
 
-class CV_EXPORTS IntrumentationRegion
-{
-public:
-    IntrumentationRegion(const char* funName, const char* fileName, int lineNum, void *retAddress, bool alwaysExpand, TYPE instrType = TYPE_GENERAL, IMPL implType = IMPL_PLAIN);
-    ~IntrumentationRegion();
-
-private:
-    bool    m_disabled; // region status
-    uint64  m_regionTicks;
-};
-
 CV_EXPORTS InstrStruct& getInstrumentStruct();
 InstrTLSStruct&         getInstrumentTLSStruct();
-CV_EXPORTS InstrNode*   getCurrentNode();
 }
 }
 
+#include "opencv2/core/utils/instrumentation.private.hpp"
+
 #ifdef _WIN32
 #define CV_INSTRUMENT_GET_RETURN_ADDRESS _ReturnAddress()
 #else
@@ -718,45 +708,6 @@ CV_EXPORTS InstrNode*   getCurrentNode();
 #define CV_INSTRUMENT_REGION_CUSTOM_META(NAME, ALWAYS_EXPAND, TYPE, IMPL)\
     void *CVAUX_CONCAT(__curr_address__, __LINE__) = [&]() {return CV_INSTRUMENT_GET_RETURN_ADDRESS;}();\
     ::cv::instr::IntrumentationRegion CVAUX_CONCAT(__instr_region__, __LINE__) (NAME, __FILE__, __LINE__, CVAUX_CONCAT(__curr_address__, __LINE__), false, ::cv::instr::TYPE_GENERAL, ::cv::instr::IMPL_PLAIN);
-// Instrument functions with non-void return type
-#define CV_INSTRUMENT_FUN_RT_META(TYPE, IMPL, ERROR_COND, FUN, ...) ([&]()\
-{\
-    if(::cv::instr::useInstrumentation()){\
-        ::cv::instr::IntrumentationRegion __instr__(#FUN, __FILE__, __LINE__, NULL, false, TYPE, IMPL);\
-        try{\
-            auto instrStatus = ((FUN)(__VA_ARGS__));\
-            if(ERROR_COND){\
-                ::cv::instr::getCurrentNode()->m_payload.m_funError = true;\
-                CV_INSTRUMENT_MARK_META(IMPL, #FUN " - BadExit");\
-            }\
-            return instrStatus;\
-        }catch(...){\
-            ::cv::instr::getCurrentNode()->m_payload.m_funError = true;\
-            CV_INSTRUMENT_MARK_META(IMPL, #FUN " - BadExit");\
-            throw;\
-        }\
-    }else{\
-        return ((FUN)(__VA_ARGS__));\
-    }\
-}())
-// Instrument functions with void return type
-#define CV_INSTRUMENT_FUN_RV_META(TYPE, IMPL, FUN, ...) ([&]()\
-{\
-    if(::cv::instr::useInstrumentation()){\
-        ::cv::instr::IntrumentationRegion __instr__(#FUN, __FILE__, __LINE__, NULL, false, TYPE, IMPL);\
-        try{\
-            (FUN)(__VA_ARGS__);\
-        }catch(...){\
-            ::cv::instr::getCurrentNode()->m_payload.m_funError = true;\
-            CV_INSTRUMENT_MARK_META(IMPL, #FUN "- BadExit");\
-            throw;\
-        }\
-    }else{\
-        (FUN)(__VA_ARGS__);\
-    }\
-}())
-// Instrumentation information marker
-#define CV_INSTRUMENT_MARK_META(IMPL, NAME, ...) {::cv::instr::IntrumentationRegion __instr_mark__(NAME, __FILE__, __LINE__, NULL, false, ::cv::instr::TYPE_MARKER, IMPL);}
 
 ///// General instrumentation
 // General OpenCV region instrumentation macro
@@ -769,10 +720,6 @@ CV_EXPORTS InstrNode*   getCurrentNode();
 ///// IPP instrumentation
 // Wrapper region instrumentation macro
 #define CV_INSTRUMENT_REGION_IPP();          CV_INSTRUMENT_REGION_META(__FUNCTION__, false, ::cv::instr::TYPE_WRAPPER, ::cv::instr::IMPL_IPP)
-// Function instrumentation macro
-#define CV_INSTRUMENT_FUN_IPP(FUN, ...)     CV_INSTRUMENT_FUN_RT_META(::cv::instr::TYPE_FUN, ::cv::instr::IMPL_IPP, instrStatus < 0, FUN, __VA_ARGS__)
-// Diagnostic markers
-#define CV_INSTRUMENT_MARK_IPP(NAME)        CV_INSTRUMENT_MARK_META(::cv::instr::IMPL_IPP, NAME)
 
 ///// OpenCL instrumentation
 // Wrapper region instrumentation macro

modules/core/include/opencv2/core/utils/instrumentation.private.hpp

@@ -0,0 +1,88 @@
+// This file is part of OpenCV project.
+// It is subject to the license terms in the LICENSE file found in the top-level directory
+// of this distribution and at http://opencv.org/license.html.
+// Copyright (C) 2026, Intel Corporation, all rights reserved.
+
+#ifndef OPENCV_CORE_UTILS_INSTRUMENTATION_PRIVATE_HPP
+#define OPENCV_CORE_UTILS_INSTRUMENTATION_PRIVATE_HPP
+
+#include "opencv2/core/utils/instrumentation.hpp"
+
+// Region-based function-instrumentation primitives.
+//
+// Only meaningful when the build was configured with ENABLE_INSTRUMENTATION;
+// otherwise this header contributes nothing and each includer provides its own
+// no-op fallbacks. The IntrumentationRegion / getCurrentNode() definitions are
+// compiled into and exported (CV_EXPORTS) from the core library, so the HAL only
+// needs these declarations and resolves the symbols at link time.
+
+#ifdef ENABLE_INSTRUMENTATION
+
+namespace cv { namespace instr {
+
+// Scoped region: records one instrumentation node on construction and closes it
+// on destruction. Defined in modules/core/src/system.cpp.
+class CV_EXPORTS IntrumentationRegion
+{
+public:
+    IntrumentationRegion(const char* funName, const char* fileName, int lineNum, void *retAddress,
+                         bool alwaysExpand, TYPE instrType = TYPE_GENERAL, IMPL implType = IMPL_PLAIN);
+    ~IntrumentationRegion();
+
+private:
+    bool    m_disabled; // region status
+    uint64  m_regionTicks;
+};
+
+CV_EXPORTS InstrNode*   getCurrentNode();
+
+}} // namespace cv::instr
+
+// Instrumentation information marker
+#define CV_INSTRUMENT_MARK_META(IMPL, NAME, ...) {::cv::instr::IntrumentationRegion __instr_mark__(NAME, __FILE__, __LINE__, NULL, false, ::cv::instr::TYPE_MARKER, IMPL);}
+
+// Instrument functions with non-void return type
+#define CV_INSTRUMENT_FUN_RT_META(TYPE, IMPL, ERROR_COND, FUN, ...) ([&]()                              \
+{                                                                                                       \
+    if(::cv::instr::useInstrumentation()){                                                              \
+        ::cv::instr::IntrumentationRegion __instr__(#FUN, __FILE__, __LINE__, NULL, false, TYPE, IMPL); \
+        try{                                                                                            \
+            auto instrStatus = ((FUN)(__VA_ARGS__));                                                    \
+            if(ERROR_COND){                                                                             \
+                ::cv::instr::getCurrentNode()->m_payload.m_funError = true;                             \
+                CV_INSTRUMENT_MARK_META(IMPL, #FUN " - BadExit");                                       \
+            }                                                                                           \
+            return instrStatus;                                                                         \
+        }catch(...){                                                                                    \
+            ::cv::instr::getCurrentNode()->m_payload.m_funError = true;                                 \
+            CV_INSTRUMENT_MARK_META(IMPL, #FUN " - BadExit");                                           \
+            throw;                                                                                      \
+        }                                                                                               \
+    }else{                                                                                              \
+        return ((FUN)(__VA_ARGS__));                                                                    \
+    }                                                                                                   \
+}())
+// Instrument functions with void return type
+#define CV_INSTRUMENT_FUN_RV_META(TYPE, IMPL, FUN, ...) ([&]()                                          \
+{                                                                                                       \
+    if(::cv::instr::useInstrumentation()){                                                              \
+        ::cv::instr::IntrumentationRegion __instr__(#FUN, __FILE__, __LINE__, NULL, false, TYPE, IMPL); \
+        try{                                                                                            \
+            (FUN)(__VA_ARGS__);                                                                         \
+        }catch(...){                                                                                    \
+            ::cv::instr::getCurrentNode()->m_payload.m_funError = true;                                 \
+            CV_INSTRUMENT_MARK_META(IMPL, #FUN " - BadExit");                                           \
+            throw;                                                                                      \
+        }                                                                                               \
+    }else{                                                                                              \
+        (FUN)(__VA_ARGS__);                                                                             \
+    }                                                                                                   \
+}())
+
+// IPP function instrumentation macros
+#define CV_INSTRUMENT_FUN_IPP(FUN, ...)     CV_INSTRUMENT_FUN_RT_META(::cv::instr::TYPE_FUN, ::cv::instr::IMPL_IPP, instrStatus < 0, FUN, __VA_ARGS__)
+#define CV_INSTRUMENT_MARK_IPP(NAME)        CV_INSTRUMENT_MARK_META(::cv::instr::IMPL_IPP, NAME)
+
+#endif // ENABLE_INSTRUMENTATION
+
+#endif // OPENCV_CORE_UTILS_INSTRUMENTATION_PRIVATE_HPP

modules/imgcodecs/src/utils.cpp

@@ -169,13 +169,13 @@ void icvCvt_Gray2BGR_16u_C1C3R( const ushort* gray, int gray_step,
                               ushort* bgr, int bgr_step, Size size )
 {
     int i;
-    for( ; size.height--; gray += gray_step/sizeof(gray[0]) )
+    for( ; size.height--; gray += gray_step/(int)sizeof(gray[0]) )
     {
         for( i = 0; i < size.width; i++, bgr += 3 )
         {
             bgr[0] = bgr[1] = bgr[2] = gray[i];
         }
-        bgr += bgr_step/sizeof(bgr[0]) - size.width*3;
+        bgr += bgr_step/(int)sizeof(bgr[0]) - size.width*3;
     }
 }
 
@@ -214,8 +214,8 @@ void icvCvt_BGRA2BGR_16u_C4C3R( const ushort* bgra, int bgra_step,
             bgr[0] = t0; bgr[1] = t1;
             t0 = bgra[swap_rb^2]; bgr[2] = t0;
         }
-        bgr += bgr_step/sizeof(bgr[0]) - size.width*3;
-        bgra += bgra_step/sizeof(bgra[0]) - size.width*4;
+        bgr += bgr_step/(int)sizeof(bgr[0]) - size.width*3;
+        bgra += bgra_step/(int)sizeof(bgra[0]) - size.width*4;
     }
 }
 
@@ -252,8 +252,8 @@ void icvCvt_BGRA2RGBA_16u_C4R( const ushort* bgra, int bgra_step,
          rgba[0] = t2; rgba[1] = t1;
          rgba[2] = t0; rgba[3] = t3;
      }
-     bgra += bgra_step/sizeof(bgra[0]) - size.width*4;
-     rgba += rgba_step/sizeof(rgba[0]) - size.width*4;
+     bgra += bgra_step/(int)sizeof(bgra[0]) - size.width*4;
+     rgba += rgba_step/(int)sizeof(rgba[0]) - size.width*4;
  }
 }
 

modules/imgcodecs/test/test_tiff.cpp

@@ -689,6 +689,28 @@ TEST(Imgcodecs_Tiff, readWrite_unsigned)
     EXPECT_EQ(0, remove(filenameOutput.c_str()));
 }
 
+// See https://github.com/opencv/opencv/issues/29615
+// Decoding a 16-bit 4-channel TIFF used to form an out of range pointer in
+// icvCvt_BGRA2RGBA_16u_C4R because the byte step was divided by an unsigned
+// sizeof and then had size.width*4 subtracted, which wraps around for a zero
+// step. This just checks that a 16UC4 TIFF round-trips correctly; the value is
+// mainly that the sanitizer builds no longer report the pointer overflow.
+TEST(Imgcodecs_Tiff, regression_29615_16UC4)
+{
+    Mat img(4, 3, CV_16UC4);
+    randu(img, Scalar::all(0), Scalar::all(65535));
+
+    vector<uchar> buf;
+    ASSERT_NO_THROW(ASSERT_TRUE(imencode(".tiff", img, buf)));
+
+    Mat decoded;
+    ASSERT_NO_THROW(decoded = imdecode(buf, IMREAD_UNCHANGED));
+    ASSERT_FALSE(decoded.empty());
+    ASSERT_EQ(CV_16UC4, decoded.type());
+    ASSERT_EQ(img.size(), decoded.size());
+    EXPECT_EQ(0, cvtest::norm(img, decoded, NORM_INF));
+}
+
 TEST(Imgcodecs_Tiff, readWrite_32FC1)
 {
     const string root = cvtest::TS::ptr()->get_data_path();

modules/ml/src/tree.cpp

@@ -1839,12 +1839,14 @@ int DTreesImpl::readSplit( const FileNode& fn )
     Split split;
 
     int vi = (int)fn["var"];
-    CV_Assert( 0 <= vi && vi <= (int)varType.size() );
+    CV_Assert( 0 <= vi && vi < (int)varMapping.size() );
     vi = varMapping[vi]; // convert to varIdx if needed
+    CV_Assert( 0 <= vi && vi < (int)varType.size() );
     split.varIdx = vi;
 
     if( varType[vi] == VAR_CATEGORICAL ) // split on categorical var
     {
+        CV_Assert( vi < (int)catOfs.size() );
         int i, val, ssize = getSubsetSize(vi);
         split.subsetOfs = (int)subsets.size();
         for( i = 0; i < ssize; i++ )
@@ -1860,6 +1862,7 @@ int DTreesImpl::readSplit( const FileNode& fn )
         if( fns.isInt() )
         {
             val = (int)fns;
+            CV_Assert( 0 <= val && (val >> 5) < ssize );
             subset[val >> 5] |= 1 << (val & 31);
         }
         else
@@ -1869,6 +1872,7 @@ int DTreesImpl::readSplit( const FileNode& fn )
             for( i = 0; i < n; i++, ++it )
             {
                 val = (int)*it;
+                CV_Assert( 0 <= val && (val >> 5) < ssize );
                 subset[val >> 5] |= 1 << (val & 31);
             }
         }

modules/ml/test/test_save_load.cpp

@@ -96,6 +96,61 @@ ML_Legacy_Param param_list[] = {
 
 INSTANTIATE_TEST_CASE_P(/**/, ML_Legacy_Params, testing::ValuesIn(param_list));
 
+TEST(ML_DTrees, load_bad_categorical_split)
+{
+    // Train a tree with a categorical input so the model carries a
+    // categorical split serialized as an "in"/"not_in" value list.
+    const int n = 40;
+    Mat samples(n, 1, CV_32F);
+    Mat responses(n, 1, CV_32S);
+    for (int i = 0; i < n; i++)
+    {
+        int cat = i % 4;
+        samples.at<float>(i, 0) = (float)cat;
+        responses.at<int>(i, 0) = (cat == 1 || cat == 2) ? 1 : 0;
+    }
+    Mat varType(2, 1, CV_8U);
+    varType.at<uchar>(0) = ml::VAR_CATEGORICAL;
+    varType.at<uchar>(1) = ml::VAR_CATEGORICAL;
+
+    Ptr<ml::TrainData> td = ml::TrainData::create(samples, ml::ROW_SAMPLE, responses,
+                                                  noArray(), noArray(), noArray(), varType);
+    Ptr<ml::DTrees> dt = ml::DTrees::create();
+    dt->setMaxDepth(4);
+    dt->setCVFolds(0);
+    dt->setMaxCategories(4);
+    dt->setMinSampleCount(1);
+    dt->train(td);
+
+    const string filename = cv::tempfile(".yml");
+    dt->save(filename);
+
+    string model;
+    {
+        std::ifstream in(filename.c_str());
+        std::stringstream ss;
+        ss << in.rdbuf();
+        model = ss.str();
+    }
+
+    // Category values in the split list index a per-split subset bitmask.
+    // Injecting a value larger than the number of categories used to write
+    // past that bitmask; loading such a model must be rejected, not crash.
+    size_t pos = model.find("in:");
+    ASSERT_NE(pos, string::npos);
+    size_t br = model.find('[', pos);
+    ASSERT_NE(br, string::npos);
+    model.insert(br + 1, "1000000,");
+    {
+        std::ofstream out(filename.c_str());
+        out << model;
+    }
+
+    Ptr<ml::DTrees> bad;
+    EXPECT_THROW(bad = ml::DTrees::load(filename), Exception);
+    remove(filename.c_str());
+}
+
 /*TEST(ML_SVM, throw_exception_when_save_untrained_model)
 {
     Ptr<cv::ml::SVM> svm;

modules/objdetect/src/hog.cpp

@@ -723,7 +723,7 @@ void HOGCache::init(const HOGDescriptor* _descriptor,
     #if CV_SIMD128
         idx = v_float32x4(0.0f, 1.0f, 2.0f, 3.0f);
 
-        for (; j <= blockSize.height - 4; j += 4)
+        for (; j <= blockSize.width - 4; j += 4)
         {
             v_float32x4 t = v_sub(idx, _bw);
             t = v_mul(t, t);

modules/objdetect/test/test_cascadeandhog.cpp

@@ -1364,4 +1364,22 @@ TEST(Objdetect_CascadeDetector, small_img)
     }
 }
 
+// See https://github.com/opencv/opencv/issues/23580
+// HOGDescriptor::compute() overflowed the gaussian weights buffer when the block
+// was taller than it was wide: the SIMD store loop for the width buffer (_dj) was
+// bounded by blockSize.height instead of blockSize.width. The block is chosen wide
+// enough that the buffer is heap allocated, so the overflow is a real out-of-bounds
+// write. "Done" is simply that compute() runs without crashing.
+TEST(Objdetect_HOGDescriptor, issue_23580_tall_block_no_overflow)
+{
+    Size winSize(272, 2048); // height > width, width > AutoBuffer stack size
+    HOGDescriptor hog(winSize, /*blockSize*/ winSize, /*blockStride*/ Size(8, 8),
+                      /*cellSize*/ Size(8, 8), /*nbins*/ 9);
+
+    Mat src(winSize, CV_8UC1, Scalar::all(0));
+    std::vector<float> descriptors;
+    ASSERT_NO_THROW(hog.compute(src, descriptors));
+    EXPECT_FALSE(descriptors.empty());
+}
+
 }} // namespace