mirror/opencv
1
0
Fork 0
mirror of https://github.com/opencv/opencv.git synced 2025-06-07 17:44:04 +08:00
Code Issues Actions Packages Projects Releases Wiki Activity

Merge pull request #20172 from alalek:fixup_19334

Browse Source
This commit is contained in:
Alexander Alekhin 2021-05-28 14:09:52 +00:00
commit 450dc92452
2 changed files with 65 additions and 1 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
modules/core/src/ocl.cpp
View File

@ -2805,7 +2805,6 @@ struct Kernel::Impl
void registerImageArgument(int arg, const Image2D& image)
{
CV_CheckGE(arg, 0, "");
CV_CheckLT(arg, (int)MAX_ARRS, "");
if (arg < (int)shadow_images.size() && shadow_images[arg].ptr() != image.ptr()) // TODO future: replace ptr => impl (more strong check)
{
CV_Check(arg, !isInProgress, "ocl::Kernel: clearing of pending Image2D arguments is not allowed");

65
modules/core/test/ocl/test_opencl.cpp
View File

@ -132,4 +132,69 @@ 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);
}
}} // namespace