}
}
+namespace cv {
+
+static bool ocl_merge( InputArrayOfArrays _mv, OutputArray _dst )
+{
+ const std::vector<UMat> & src = *(const std::vector<UMat> *)(_mv.getObj());
+ CV_Assert(!src.empty());
+
+ int type = src[0].type(), depth = CV_MAT_DEPTH(type);
+ Size size = src[0].size();
+ bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0;
+
+ if (doubleSupport && depth == CV_64F)
+ return false;
+
+ size_t srcsize = src.size();
+ for (size_t i = 0; i < srcsize; ++i)
+ {
+ int itype = src[i].type(), icn = CV_MAT_CN(itype), idepth = CV_MAT_DEPTH(itype);
+ if (src[i].dims > 2 || icn != 1)
+ return false;
+ CV_Assert(size == src[i].size() && depth == idepth);
+ }
+
+ String srcargs, srcdecl, processelem;
+ for (size_t i = 0; i < srcsize; ++i)
+ {
+ srcargs += format("DECLARE_SRC_PARAM(%d)", i);
+ srcdecl += format("DECLARE_DATA(%d)", i);
+ processelem += format("PROCESS_ELEM(%d)", i);
+ }
+
+ ocl::Kernel k("merge", ocl::core::split_merge_oclsrc,
+ format("-D OP_MERGE -D cn=%d -D T=%s -D DECLARE_SRC_PARAMS_N=%s -D DECLARE_DATA_N=%s -D PROCESS_ELEMS_N=%s",
+ (int)srcsize, ocl::memopTypeToStr(depth), srcargs.c_str(), srcdecl.c_str(), processelem.c_str()));
+ if (k.empty())
+ return false;
+
+ _dst.create(size, CV_MAKE_TYPE(depth, srcsize));
+ UMat dst = _dst.getUMat();
+
+ int argidx = 0;
+ for (size_t i = 0; i < srcsize; ++i)
+ argidx = k.set(argidx, ocl::KernelArg::ReadOnlyNoSize(src[i]));
+ k.set(argidx, ocl::KernelArg::WriteOnly(dst));
+
+ size_t globalsize[2] = { dst.cols, dst.rows };
+ return k.run(2, globalsize, NULL, false);
+}
+
+}
+
void cv::merge(InputArrayOfArrays _mv, OutputArray _dst)
{
+ if (ocl::useOpenCL() && _mv.isUMatVector() && _dst.isUMat() && ocl_merge(_mv, _dst))
+ return;
+
std::vector<Mat> mv;
_mv.getMatVector(mv);
merge(!mv.empty() ? &mv[0] : 0, mv.size(), _dst);
return (size_t)(vv[i].data - vv[i].datastart);
}
+ if( k == STD_VECTOR_UMAT )
+ {
+ const std::vector<UMat>& vv = *(const std::vector<UMat>*)obj;
+ CV_Assert((size_t)i < vv.size());
+ return vv[i].offset;
+ }
+
if( k == GPU_MAT )
{
CV_Assert( i < 0 );
return vv[i].step;
}
+ if( k == STD_VECTOR_UMAT )
+ {
+ const std::vector<UMat>& vv = *(const std::vector<UMat>*)obj;
+ CV_Assert((size_t)i < vv.size());
+ return vv[i].step;
+ }
+
if( k == GPU_MAT )
{
CV_Assert( i < 0 );
// First, try to retrieve existing context of the same type.
// In its turn, Platform::getContext() may call Context2::create()
// if there is no such context.
- ctx.create(Device::TYPE_ACCELERATOR);
+ ctx.create(Device::TYPE_CPU);
if(!ctx.p)
ctx.create(Device::TYPE_DGPU);
if(!ctx.p)
CV_Assert(i >= 0);
if( i == 0 )
p->cleanupUMats();
- if( !p || !p->handle || clSetKernelArg(p->handle, (cl_uint)i, sz, value) < 0 )
+ cl_int retval;
+ if( !p || !p->handle || (retval = clSetKernelArg(p->handle, (cl_uint)i, sz, value)) < 0 )
+ {
+ printf("%d\n", retval);
return -1;
+ }
+ printf("%d\n", retval);
return i+1;
}
int Kernel::set(int i, const KernelArg& arg)
{
+ printf("Setting to index %d\n", i);
CV_Assert( i >= 0 );
if( !p || !p->handle )
return -1;
cl_mem h = (cl_mem)arg.m->handle(accessFlags);
if (ptronly)
- clSetKernelArg(p->handle, (cl_uint)i++, sizeof(h), &h);
+ printf("%d\n", clSetKernelArg(p->handle, (cl_uint)i++, sizeof(h), &h));
else if( arg.m->dims <= 2 )
{
UMat2D u2d(*arg.m);
- clSetKernelArg(p->handle, (cl_uint)i, sizeof(h), &h);
- clSetKernelArg(p->handle, (cl_uint)(i+1), sizeof(u2d.step), &u2d.step);
- clSetKernelArg(p->handle, (cl_uint)(i+2), sizeof(u2d.offset), &u2d.offset);
+ printf("setting ... \n");
+ printf("%d\n", clSetKernelArg(p->handle, (cl_uint)i, sizeof(h), &h));
+ printf("%d\n", clSetKernelArg(p->handle, (cl_uint)(i+1), sizeof(u2d.step), &u2d.step));
+ printf("%d\n", clSetKernelArg(p->handle, (cl_uint)(i+2), sizeof(u2d.offset), &u2d.offset));
i += 3;
if( !(arg.flags & KernelArg::NO_SIZE) )
{
int cols = u2d.cols*arg.wscale;
- clSetKernelArg(p->handle, (cl_uint)i, sizeof(u2d.rows), &u2d.rows);
- clSetKernelArg(p->handle, (cl_uint)(i+1), sizeof(cols), &cols);
+ printf("%d\n", clSetKernelArg(p->handle, (cl_uint)i, sizeof(u2d.rows), &u2d.rows));
+ printf("%d\n", clSetKernelArg(p->handle, (cl_uint)(i+1), sizeof(cols), &cols));
i += 2;
}
}
--- /dev/null
+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+// By downloading, copying, installing or using the software you agree to this license.
+// If you do not agree to this license, do not download, install,
+// copy or use the software.
+//
+//
+// License Agreement
+// For Open Source Computer Vision Library
+//
+// Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved.
+// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
+// Copyright (C) 2013, OpenCV Foundation, all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+// * Redistribution's of source code must retain the above copyright notice,
+// this list of conditions and the following disclaimer.
+//
+// * Redistribution's in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// * The name of the copyright holders may not be used to endorse or promote products
+// derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors as is and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the copyright holders or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#ifdef OP_MERGE
+
+#define DECLARE_SRC_PARAM(index) __global const uchar * src##index##ptr, int src##index##_step, int src##index##_offset,
+#define DECLARE_DATA(index) __global const T * src##index = \
+ (__global T *)(src##index##ptr + mad24(src##index##_step, y, x * (int)sizeof(T) + src##index##_offset));
+#define PROCESS_ELEM(index) dst[index] = src##index[0];
+
+__kernel void merge(DECLARE_SRC_PARAMS_N
+ __global uchar * dstptr, int dst_step, int dst_offset,
+ int rows, int cols)
+{
+ int x = get_global_id(0);
+ int y = get_global_id(1);
+
+ if (x < cols && y < rows)
+ {
+ DECLARE_DATA_N
+ __global T * dst = (__global T *)(dstptr + mad24(dst_step, y, x * (int)sizeof(T) * cn + dst_offset));
+ PROCESS_ELEMS_N
+ }
+}
+
+#elif defined OP_SPLIT
+
+__kernel void set(__global uchar* dstptr, int dststep, int dstoffset,
+ int rows, int cols, dstT value )
+{
+ int x = get_global_id(0);
+ int y = get_global_id(1);
+
+ if (x < cols && y < rows)
+ {
+ int dst_index = mad24(y, dststep, x*(int)sizeof(dstT) + dstoffset);
+ *(__global dstT*)(dstptr + dst_index) = value;
+ }
+}
+
+#else
+#error "No operation"
+#endif
#include "test_precomp.hpp"
#include "opencv2/ts/ocl_test.hpp"
-#ifdef HAVE_OPENCL
-
namespace cvtest {
namespace ocl {
OCL_INSTANTIATE_TEST_CASE_P(Arithm, Flip, Combine(Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), Values(1, 2, 3, 4), Bool()));
} } // namespace cvtest::ocl
-
-#endif // HAVE_OPENCL
--- /dev/null
+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+// By downloading, copying, installing or using the software you agree to this license.
+// If you do not agree to this license, do not download, install,
+// copy or use the software.
+//
+//
+// License Agreement
+// For Open Source Computer Vision Library
+//
+// Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved.
+// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
+// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// @Authors
+// Jia Haipeng, jiahaipeng95@gmail.com
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+// * Redistribution's of source code must retain the above copyright notice,
+// this list of conditions and the following disclaimer.
+//
+// * Redistribution's in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// * The name of the copyright holders may not be used to endorse or promote products
+// derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#include "test_precomp.hpp"
+#include "opencv2/ts/ocl_test.hpp"
+
+namespace cvtest {
+namespace ocl {
+
+PARAM_TEST_CASE(MergeTestBase, MatDepth, Channels, bool)
+{
+ int depth, cn;
+ bool use_roi;
+
+ TEST_DECLARE_INPUT_PARAMETER(src1)
+ TEST_DECLARE_INPUT_PARAMETER(src2)
+ TEST_DECLARE_INPUT_PARAMETER(src3)
+ TEST_DECLARE_INPUT_PARAMETER(src4)
+ TEST_DECLARE_OUTPUT_PARAMETER(dst)
+
+ std::vector<Mat> src_roi;
+ std::vector<UMat> usrc_roi;
+
+ virtual void SetUp()
+ {
+ depth = GET_PARAM(0);
+ cn = GET_PARAM(1);
+ use_roi = GET_PARAM(2);
+ }
+
+ virtual void random_roi()
+ {
+ CV_Assert(cn >= 1 && cn <= 4);
+ Size roiSize = randomSize(1, MAX_VALUE);
+
+ {
+ Border src1Border = randomBorder(0, use_roi ? MAX_VALUE : 0);
+ randomSubMat(src1, src1_roi, roiSize, src1Border, depth, 2, 11);
+
+ Border src2Border = randomBorder(0, use_roi ? MAX_VALUE : 0);
+ randomSubMat(src2, src2_roi, roiSize, src2Border, depth, -1540, 1740);
+
+ Border src3Border = randomBorder(0, use_roi ? MAX_VALUE : 0);
+ randomSubMat(src3, src3_roi, roiSize, src3Border, depth, -1540, 1740);
+
+ Border src4Border = randomBorder(0, use_roi ? MAX_VALUE : 0);
+ randomSubMat(src4, src4_roi, roiSize, src4Border, depth, -1540, 1740);
+ }
+
+ Border dstBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
+ randomSubMat(dst, dst_roi, roiSize, dstBorder, CV_MAKE_TYPE(depth, cn), 5, 16);
+
+ UMAT_UPLOAD_INPUT_PARAMETER(src1)
+ UMAT_UPLOAD_INPUT_PARAMETER(src2)
+ UMAT_UPLOAD_INPUT_PARAMETER(src3)
+ UMAT_UPLOAD_INPUT_PARAMETER(src4)
+ UMAT_UPLOAD_OUTPUT_PARAMETER(dst)
+
+ src_roi.push_back(src1_roi), usrc_roi.push_back(usrc1_roi);
+ if (cn >= 2)
+ src_roi.push_back(src2_roi), usrc_roi.push_back(usrc2_roi);
+ if (cn >= 3)
+ src_roi.push_back(src3_roi), usrc_roi.push_back(usrc3_roi);
+ if (cn >= 4)
+ src_roi.push_back(src4_roi), usrc_roi.push_back(usrc4_roi);
+ }
+
+ void Near(double threshold = 0.)
+ {
+ EXPECT_MAT_NEAR(dst, udst, threshold);
+ EXPECT_MAT_NEAR(dst_roi, udst_roi, threshold);
+ }
+};
+
+typedef MergeTestBase Merge;
+
+OCL_TEST_P(Merge, Accuracy)
+{
+ for(int j = 0; j < test_loop_times; j++)
+ {
+ random_roi();
+
+ OCL_OFF(cv::merge(src_roi, dst_roi));
+ OCL_ON(cv::merge(usrc_roi, udst_roi));
+
+ Near();
+ }
+}
+
+//PARAM_TEST_CASE(SplitTestBase, MatType, int, bool)
+//{
+// int type;
+// int channels;
+// bool use_roi;
+
+// cv::Mat src, src_roi;
+// cv::Mat dst[MAX_CHANNELS], dst_roi[MAX_CHANNELS];
+
+// cv::ocl::oclMat gsrc_whole, gsrc_roi;
+// cv::ocl::oclMat gdst_whole[MAX_CHANNELS], gdst_roi[MAX_CHANNELS];
+
+// virtual void SetUp()
+// {
+// type = GET_PARAM(0);
+// channels = GET_PARAM(1);
+// use_roi = GET_PARAM(2);
+// }
+
+// void random_roi()
+// {
+// Size roiSize = randomSize(1, MAX_VALUE);
+// Border srcBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
+// randomSubMat(src, src_roi, roiSize, srcBorder, CV_MAKETYPE(type, channels), 0, 256);
+// generateOclMat(gsrc_whole, gsrc_roi, src, roiSize, srcBorder);
+
+// for (int i = 0; i < channels; ++i)
+// {
+// Border dstBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
+// randomSubMat(dst[i], dst_roi[i], roiSize, dstBorder, CV_MAKETYPE(type, 1), 5, 16);
+// generateOclMat(gdst_whole[i], gdst_roi[i], dst[i], roiSize, dstBorder);
+// }
+// }
+//};
+
+//struct Split : SplitTestBase {};
+
+//#ifdef ANDROID
+//// NOTE: The test fail on Android is the top of the iceberg only
+//// The real fail reason is memory access vialation somewhere else
+//OCL_TEST_P(Split, DISABLED_Accuracy)
+//#else
+//OCL_TEST_P(Split, Accuracy)
+//#endif
+//{
+// for(int j = 0; j < LOOP_TIMES; j++)
+// {
+// random_roi();
+
+// cv::split(src_roi, dst_roi);
+// cv::ocl::split(gsrc_roi, gdst_roi);
+
+// for (int i = 0; i < channels; ++i)
+// {
+// EXPECT_MAT_NEAR(dst[i], gdst_whole[i], 0.0);
+// EXPECT_MAT_NEAR(dst_roi[i], gdst_roi[i], 0.0);
+// }
+// }
+//}
+
+
+OCL_INSTANTIATE_TEST_CASE_P(SplitMerge, Merge, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
+
+
+//INSTANTIATE_TEST_CASE_P(SplitMerge, Split , Combine(
+// Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F), Values(1, 2, 3, 4), Bool()));
+
+
+} } // namespace cvtest::ocl
projects / platform / upstream / opencv.git / commitdiff