From: Alexander Alekhin <alexander.alekhin@intel.com>
Date: Sat, 29 May 2021 19:00:14 +0000 (+0000)
Subject: Merge remote-tracking branch 'upstream/3.4' into merge-3.4
X-Git-Tag: submit/tizen/20220120.021815~1^2~63
X-Git-Url: http://review.tizen.org/git/?a=commitdiff_plain;h=cb51a155b2e5887a8cf533a912c24fa943d0a52a;p=platform%2Fupstream%2Fopencv.git

Merge remote-tracking branch 'upstream/3.4' into merge-3.4
---

cb51a155b2e5887a8cf533a912c24fa943d0a52a
diff --cc modules/core/test/ocl/test_opencl.cpp
index e639f72948,0c79150d49..daa023534d
--- a/modules/core/test/ocl/test_opencl.cpp
+++ b/modules/core/test/ocl/test_opencl.cpp
@@@ -132,120 -132,69 +132,187 @@@ TEST(OpenCL, support_SPIR_programs
      testOpenCLKernel(k);
  }
  
++
+ TEST(OpenCL, image2Dcount_regression_19334)
+ {
+     cv::ocl::Context ctx = cv::ocl::Context::getDefault();
+     if (!ctx.ptr())
+     {
+         throw cvtest::SkipTestException("OpenCL is not available");
+     }
+     cv::ocl::Device device = cv::ocl::Device::getDefault();
+     if (!device.compilerAvailable())
+     {
+         throw cvtest::SkipTestException("OpenCL compiler is not available");
+     }
+ 
+     std::string module_name; // empty to disable OpenCL cache
+ 
+     static const char* opencl_kernel_src =
+ "__kernel void test_kernel(int a,\n"
+ "                          __global const uchar* src0, int src0_step, int src0_offset, int src0_rows, int src0_cols,\n"
+ "                          __global const uchar* src1, int src1_step, int src1_offset, int src1_rows, int src1_cols,\n"
+ "                          __global const uchar* src2, int src2_step, int src2_offset, int src2_rows, int src2_cols,\n"
+ "                          __read_only image2d_t image)\n"
+ "{\n"
+ "}";
+     cv::ocl::ProgramSource src(module_name, "test_opencl_image_arg", opencl_kernel_src, "");
+     cv::String errmsg;
+     cv::ocl::Program program(src, "", errmsg);
+     ASSERT_TRUE(program.ptr() != NULL);
+     cv::ocl::Kernel k("test_kernel", program);
+     ASSERT_FALSE(k.empty());
+ 
+     std::vector<UMat> images(4);
+     for (size_t i = 0; i < images.size(); ++i)
+         images[i] = UMat(10, 10, CV_8UC1);
+     cv::ocl::Image2D image;
+     try
+     {
+         cv::ocl::Image2D image_(images.back());
+         image = image_;
+     }
+     catch (const cv::Exception&)
+     {
+         throw cvtest::SkipTestException("OpenCL images are not supported");
+     }
+ 
+     int nargs = 0;
+     int a = 0;
+     nargs = k.set(nargs, a);
+     ASSERT_EQ(1, nargs);
+     nargs = k.set(nargs, images[0]);
+     ASSERT_EQ(6, nargs);
+     nargs = k.set(nargs, images[1]);
+     ASSERT_EQ(11, nargs);
+     nargs = k.set(nargs, images[2]);
+     ASSERT_EQ(16, nargs);
+ 
+     // do not throw (issue of #19334)
+     ASSERT_NO_THROW(nargs = k.set(nargs, image));
+     ASSERT_EQ(17, nargs);
+ 
+     // allow to replace image argument if kernel is not running
+     UMat image2(10, 10, CV_8UC1);
+     ASSERT_NO_THROW(nargs = k.set(16, cv::ocl::Image2D(image2)));
+     ASSERT_EQ(17, nargs);
+ }
+ 
++
 +TEST(OpenCL, move_construct_assign)
 +{
 +    cv::ocl::Context ctx1 = cv::ocl::Context::getDefault();
 +    if (!ctx1.ptr())
 +    {
 +        throw cvtest::SkipTestException("OpenCL is not available");
 +    }
 +    void* const ctx_ptr = ctx1.ptr();
 +    cv::ocl::Context ctx2(std::move(ctx1));
 +    ASSERT_EQ(ctx1.ptr(), nullptr);
 +    ASSERT_EQ(ctx2.ptr(), ctx_ptr);
 +    cv::ocl::Context ctx3 = std::move(ctx2);
 +    ASSERT_EQ(ctx2.ptr(), nullptr);
 +    ASSERT_EQ(ctx3.ptr(), ctx_ptr);
 +
 +    cv::ocl::Platform pl1 = cv::ocl::Platform::getDefault();
 +    void* const pl_ptr = pl1.ptr();
 +    cv::ocl::Platform pl2(std::move(pl1));
 +    ASSERT_EQ(pl1.ptr(), nullptr);
 +    ASSERT_EQ(pl2.ptr(), pl_ptr);
 +    cv::ocl::Platform pl3 = std::move(pl2);
 +    ASSERT_EQ(pl2.ptr(), nullptr);
 +    ASSERT_EQ(pl3.ptr(), pl_ptr);
 +
 +    std::vector<cv::ocl::PlatformInfo> platformInfos;
 +    cv::ocl::getPlatfomsInfo(platformInfos);
 +    const cv::String pi_name = platformInfos[0].name();
 +    cv::ocl::PlatformInfo pinfo2(std::move(platformInfos[0]));
 +    ASSERT_EQ(platformInfos[0].name(), cv::String());
 +    ASSERT_EQ(pinfo2.name(), pi_name);
 +    cv::ocl::PlatformInfo pinfo3 = std::move(pinfo2);
 +    ASSERT_EQ(pinfo2.name(), cv::String());
 +    ASSERT_EQ(pinfo3.name(), pi_name);
 +
 +    cv::ocl::Queue q1 = cv::ocl::Queue::getDefault();
 +    void* const q_ptr = q1.ptr();
 +    cv::ocl::Queue q2(std::move(q1));
 +    ASSERT_EQ(q1.ptr(), nullptr);
 +    ASSERT_EQ(q2.ptr(), q_ptr);
 +    cv::ocl::Queue q3 = std::move(q2);
 +    ASSERT_EQ(q2.ptr(), nullptr);
 +    ASSERT_EQ(q3.ptr(), q_ptr);
 +
 +    cv::ocl::Device d1 = cv::ocl::Device::getDefault();
 +    if (!d1.compilerAvailable())
 +    {
 +        throw cvtest::SkipTestException("OpenCL compiler is not available");
 +    }
 +    void* const d_ptr = d1.ptr();
 +    cv::ocl::Device d2(std::move(d1));
 +    ASSERT_EQ(d1.ptr(), nullptr);
 +    ASSERT_EQ(d2.ptr(), d_ptr);
 +    cv::ocl::Device d3 = std::move(d2);
 +    ASSERT_EQ(d2.ptr(), nullptr);
 +    ASSERT_EQ(d3.ptr(), d_ptr);
 +
 +    if (d3.imageSupport()) {
 +        cv::UMat umat1 = cv::UMat::ones(640, 480, CV_32FC1);
 +        cv::ocl::Image2D img1(umat1);
 +        void *const img_ptr = img1.ptr();
 +        cv::ocl::Image2D img2(std::move(img1));
 +        ASSERT_EQ(img1.ptr(), nullptr);
 +        ASSERT_EQ(img2.ptr(), img_ptr);
 +        cv::ocl::Image2D img3 = std::move(img2);
 +        ASSERT_EQ(img2.ptr(), nullptr);
 +        ASSERT_EQ(img3.ptr(), img_ptr);
 +    }
 +
 +    static const char* opencl_kernel_src =
 +"__kernel void test_kernel(__global const uchar* src, int src_step, int src_offset,\n"
 +"                          __global uchar* dst, int dst_step, int dst_offset, int dst_rows, int dst_cols,\n"
 +"                          int c)\n"
 +"{\n"
 +"   int x = get_global_id(0);\n"
 +"   int y = get_global_id(1);\n"
 +"   if (x < dst_cols && y < dst_rows)\n"
 +"   {\n"
 +"       int src_idx = y * src_step + x + src_offset;\n"
 +"       int dst_idx = y * dst_step + x + dst_offset;\n"
 +"       dst[dst_idx] = src[src_idx] + c;\n"
 +"   }\n"
 +"}\n";
 +    cv::String module_name; // empty to disable OpenCL cache
 +    cv::ocl::ProgramSource ps1(module_name, "move_construct_assign", opencl_kernel_src, "");
 +    cv::ocl::ProgramSource::Impl* const ps_ptr = ps1.getImpl();
 +    cv::ocl::ProgramSource ps2(std::move(ps1));
 +    ASSERT_EQ(ps1.getImpl(), nullptr);
 +    ASSERT_EQ(ps2.getImpl(), ps_ptr);
 +    cv::ocl::ProgramSource ps3 = std::move(ps2);
 +    ASSERT_EQ(ps2.getImpl(), nullptr);
 +    ASSERT_EQ(ps3.getImpl(), ps_ptr);
 +
 +    cv::String errmsg;
 +    cv::ocl::Program prog1(ps3, "", errmsg);
 +    void* const prog_ptr = prog1.ptr();
 +    ASSERT_NE(prog_ptr, nullptr);
 +    cv::ocl::Program prog2(std::move(prog1));
 +    ASSERT_EQ(prog1.ptr(), nullptr);
 +    ASSERT_EQ(prog2.ptr(), prog_ptr);
 +    cv::ocl::Program prog3 = std::move(prog2);
 +    ASSERT_EQ(prog2.ptr(), nullptr);
 +    ASSERT_EQ(prog3.ptr(), prog_ptr);
 +
 +    cv::ocl::Kernel k1("test_kernel", prog3);
 +    void* const k_ptr = k1.ptr();
 +    ASSERT_NE(k_ptr, nullptr);
 +    cv::ocl::Kernel k2(std::move(k1));
 +    ASSERT_EQ(k1.ptr(), nullptr);
 +    ASSERT_EQ(k2.ptr(), k_ptr);
 +    cv::ocl::Kernel k3 = std::move(k2);
 +    ASSERT_EQ(k2.ptr(), nullptr);
 +    ASSERT_EQ(k3.ptr(), k_ptr);
 +
 +    testOpenCLKernel(k3);
 +}
 +
  }} // namespace