mirror of
https://github.com/opencv/opencv.git
synced 2024-11-30 06:10:02 +08:00
add move construct/assigns to cv::ocl main classes
- enables inline construct and assigns with r-values - enables compiler-created default move construct/assigns - includes test cases
This commit is contained in:
Dale Phurrough
parent
54d80d9168
commit
96a15434a2
@ -74,6 +74,8 @@ public:
|
||||
explicit Device(void* d);
|
||||
Device(const Device& d);
|
||||
Device& operator = (const Device& d);
|
||||
Device(Device&& d) CV_NOEXCEPT;
|
||||
Device& operator = (Device&& d) CV_NOEXCEPT;
|
||||
CV_WRAP ~Device();
|
||||
|
||||
void set(void* d);
|
||||
@ -250,6 +252,8 @@ public:
|
||||
~Context();
|
||||
Context(const Context& c);
|
||||
Context& operator= (const Context& c);
|
||||
Context(Context&& c) CV_NOEXCEPT;
|
||||
Context& operator = (Context&& c) CV_NOEXCEPT;
|
||||
|
||||
/** @deprecated */
|
||||
bool create();
|
||||
@ -302,6 +306,8 @@ public:
|
||||
~Platform();
|
||||
Platform(const Platform& p);
|
||||
Platform& operator = (const Platform& p);
|
||||
Platform(Platform&& p) CV_NOEXCEPT;
|
||||
Platform& operator = (Platform&& p) CV_NOEXCEPT;
|
||||
|
||||
void* ptr() const;
|
||||
|
||||
@ -362,6 +368,8 @@ public:
|
||||
~Queue();
|
||||
Queue(const Queue& q);
|
||||
Queue& operator = (const Queue& q);
|
||||
Queue(Queue&& q) CV_NOEXCEPT;
|
||||
Queue& operator = (Queue&& q) CV_NOEXCEPT;
|
||||
|
||||
bool create(const Context& c=Context(), const Device& d=Device());
|
||||
void finish();
|
||||
@ -384,7 +392,7 @@ class CV_EXPORTS KernelArg
|
||||
public:
|
||||
enum { LOCAL=1, READ_ONLY=2, WRITE_ONLY=4, READ_WRITE=6, CONSTANT=8, PTR_ONLY = 16, NO_SIZE=256 };
|
||||
KernelArg(int _flags, UMat* _m, int wscale=1, int iwscale=1, const void* _obj=0, size_t _sz=0);
|
||||
KernelArg();
|
||||
KernelArg() CV_NOEXCEPT;
|
||||
|
||||
static KernelArg Local(size_t localMemSize)
|
||||
{ return KernelArg(LOCAL, 0, 1, 1, 0, localMemSize); }
|
||||
@ -428,6 +436,8 @@ public:
|
||||
~Kernel();
|
||||
Kernel(const Kernel& k);
|
||||
Kernel& operator = (const Kernel& k);
|
||||
Kernel(Kernel&& k) CV_NOEXCEPT;
|
||||
Kernel& operator = (Kernel&& k) CV_NOEXCEPT;
|
||||
|
||||
bool empty() const;
|
||||
bool create(const char* kname, const Program& prog);
|
||||
@ -502,8 +512,9 @@ public:
|
||||
Program(const ProgramSource& src,
|
||||
const String& buildflags, String& errmsg);
|
||||
Program(const Program& prog);
|
||||
|
||||
Program& operator = (const Program& prog);
|
||||
Program(Program&& prog) CV_NOEXCEPT;
|
||||
Program& operator = (Program&& prog) CV_NOEXCEPT;
|
||||
~Program();
|
||||
|
||||
bool create(const ProgramSource& src,
|
||||
@ -551,6 +562,8 @@ public:
|
||||
~ProgramSource();
|
||||
ProgramSource(const ProgramSource& prog);
|
||||
ProgramSource& operator = (const ProgramSource& prog);
|
||||
ProgramSource(ProgramSource&& prog) CV_NOEXCEPT;
|
||||
ProgramSource& operator = (ProgramSource&& prog) CV_NOEXCEPT;
|
||||
|
||||
const String& source() const; // deprecated
|
||||
hash_t hash() const; // deprecated
|
||||
@ -623,6 +636,8 @@ public:
|
||||
|
||||
PlatformInfo(const PlatformInfo& i);
|
||||
PlatformInfo& operator =(const PlatformInfo& i);
|
||||
PlatformInfo(PlatformInfo&& i) CV_NOEXCEPT;
|
||||
PlatformInfo& operator = (PlatformInfo&& i) CV_NOEXCEPT;
|
||||
|
||||
String name() const;
|
||||
String vendor() const;
|
||||
@ -683,7 +698,7 @@ CV_EXPORTS void buildOptionsAddMatrixDescription(String& buildOptions, const Str
|
||||
class CV_EXPORTS Image2D
|
||||
{
|
||||
public:
|
||||
Image2D();
|
||||
Image2D() CV_NOEXCEPT;
|
||||
|
||||
/**
|
||||
@param src UMat object from which to get image properties and data
|
||||
@ -696,6 +711,8 @@ public:
|
||||
~Image2D();
|
||||
|
||||
Image2D & operator = (const Image2D & i);
|
||||
Image2D(Image2D &&) CV_NOEXCEPT;
|
||||
Image2D &operator=(Image2D &&) CV_NOEXCEPT;
|
||||
|
||||
/** Indicates if creating an aliased image should succeed.
|
||||
Depends on the underlying platform and the dimensions of the UMat.
|
||||
|
||||
@ -1480,6 +1480,23 @@ Platform& Platform::operator = (const Platform& pl)
|
||||
return *this;
|
||||
}
|
||||
|
||||
Platform::Platform(Platform&& pl) CV_NOEXCEPT
|
||||
{
|
||||
p = pl.p;
|
||||
pl.p = nullptr;
|
||||
}
|
||||
|
||||
Platform& Platform::operator = (Platform&& pl) CV_NOEXCEPT
|
||||
{
|
||||
if (this != &pl) {
|
||||
if(p)
|
||||
p->release();
|
||||
p = pl.p;
|
||||
pl.p = nullptr;
|
||||
}
|
||||
return *this;
|
||||
}
|
||||
|
||||
void* Platform::ptr() const
|
||||
{
|
||||
return p ? p->handle : 0;
|
||||
@ -1706,6 +1723,23 @@ Device& Device::operator = (const Device& d)
|
||||
return *this;
|
||||
}
|
||||
|
||||
Device::Device(Device&& d) CV_NOEXCEPT
|
||||
{
|
||||
p = d.p;
|
||||
d.p = nullptr;
|
||||
}
|
||||
|
||||
Device& Device::operator = (Device&& d) CV_NOEXCEPT
|
||||
{
|
||||
if (this != &d) {
|
||||
if(p)
|
||||
p->release();
|
||||
p = d.p;
|
||||
d.p = nullptr;
|
||||
}
|
||||
return *this;
|
||||
}
|
||||
|
||||
Device::~Device()
|
||||
{
|
||||
if(p)
|
||||
@ -2919,6 +2953,23 @@ Context& Context::operator = (const Context& c)
|
||||
return *this;
|
||||
}
|
||||
|
||||
Context::Context(Context&& c) CV_NOEXCEPT
|
||||
{
|
||||
p = c.p;
|
||||
c.p = nullptr;
|
||||
}
|
||||
|
||||
Context& Context::operator = (Context&& c) CV_NOEXCEPT
|
||||
{
|
||||
if (this != &c) {
|
||||
if(p)
|
||||
p->release();
|
||||
p = c.p;
|
||||
c.p = nullptr;
|
||||
}
|
||||
return *this;
|
||||
}
|
||||
|
||||
void* Context::ptr() const
|
||||
{
|
||||
return p == NULL ? NULL : p->handle;
|
||||
@ -3260,6 +3311,23 @@ Queue& Queue::operator = (const Queue& q)
|
||||
return *this;
|
||||
}
|
||||
|
||||
Queue::Queue(Queue&& q) CV_NOEXCEPT
|
||||
{
|
||||
p = q.p;
|
||||
q.p = nullptr;
|
||||
}
|
||||
|
||||
Queue& Queue::operator = (Queue&& q) CV_NOEXCEPT
|
||||
{
|
||||
if (this != &q) {
|
||||
if(p)
|
||||
p->release();
|
||||
p = q.p;
|
||||
q.p = nullptr;
|
||||
}
|
||||
return *this;
|
||||
}
|
||||
|
||||
Queue::~Queue()
|
||||
{
|
||||
if(p)
|
||||
@ -3315,7 +3383,7 @@ static cl_command_queue getQueue(const Queue& q)
|
||||
|
||||
/////////////////////////////////////////// KernelArg /////////////////////////////////////////////
|
||||
|
||||
KernelArg::KernelArg()
|
||||
KernelArg::KernelArg() CV_NOEXCEPT
|
||||
: flags(0), m(0), obj(0), sz(0), wscale(1), iwscale(1)
|
||||
{
|
||||
}
|
||||
@ -3493,6 +3561,23 @@ Kernel& Kernel::operator = (const Kernel& k)
|
||||
return *this;
|
||||
}
|
||||
|
||||
Kernel::Kernel(Kernel&& k) CV_NOEXCEPT
|
||||
{
|
||||
p = k.p;
|
||||
k.p = nullptr;
|
||||
}
|
||||
|
||||
Kernel& Kernel::operator = (Kernel&& k) CV_NOEXCEPT
|
||||
{
|
||||
if (this != &k) {
|
||||
if(p)
|
||||
p->release();
|
||||
p = k.p;
|
||||
k.p = nullptr;
|
||||
}
|
||||
return *this;
|
||||
}
|
||||
|
||||
Kernel::~Kernel()
|
||||
{
|
||||
if(p)
|
||||
@ -4091,6 +4176,23 @@ ProgramSource& ProgramSource::operator = (const ProgramSource& prog)
|
||||
return *this;
|
||||
}
|
||||
|
||||
ProgramSource::ProgramSource(ProgramSource&& prog) CV_NOEXCEPT
|
||||
{
|
||||
p = prog.p;
|
||||
prog.p = nullptr;
|
||||
}
|
||||
|
||||
ProgramSource& ProgramSource::operator = (ProgramSource&& prog) CV_NOEXCEPT
|
||||
{
|
||||
if (this != &prog) {
|
||||
if(p)
|
||||
p->release();
|
||||
p = prog.p;
|
||||
prog.p = nullptr;
|
||||
}
|
||||
return *this;
|
||||
}
|
||||
|
||||
const String& ProgramSource::source() const
|
||||
{
|
||||
CV_Assert(p);
|
||||
@ -4583,6 +4685,23 @@ Program& Program::operator = (const Program& prog)
|
||||
return *this;
|
||||
}
|
||||
|
||||
Program::Program(Program&& prog) CV_NOEXCEPT
|
||||
{
|
||||
p = prog.p;
|
||||
prog.p = nullptr;
|
||||
}
|
||||
|
||||
Program& Program::operator = (Program&& prog) CV_NOEXCEPT
|
||||
{
|
||||
if (this != &prog) {
|
||||
if(p)
|
||||
p->release();
|
||||
p = prog.p;
|
||||
prog.p = nullptr;
|
||||
}
|
||||
return *this;
|
||||
}
|
||||
|
||||
Program::~Program()
|
||||
{
|
||||
if(p)
|
||||
@ -6644,6 +6763,23 @@ PlatformInfo& PlatformInfo::operator =(const PlatformInfo& i)
|
||||
return *this;
|
||||
}
|
||||
|
||||
PlatformInfo::PlatformInfo(PlatformInfo&& i) CV_NOEXCEPT
|
||||
{
|
||||
p = i.p;
|
||||
i.p = nullptr;
|
||||
}
|
||||
|
||||
PlatformInfo& PlatformInfo::operator = (PlatformInfo&& i) CV_NOEXCEPT
|
||||
{
|
||||
if (this != &i) {
|
||||
if(p)
|
||||
p->release();
|
||||
p = i.p;
|
||||
i.p = nullptr;
|
||||
}
|
||||
return *this;
|
||||
}
|
||||
|
||||
int PlatformInfo::deviceNumber() const
|
||||
{
|
||||
return p ? (int)p->devices.size() : 0;
|
||||
@ -7184,7 +7320,7 @@ struct Image2D::Impl
|
||||
cl_mem handle;
|
||||
};
|
||||
|
||||
Image2D::Image2D()
|
||||
Image2D::Image2D() CV_NOEXCEPT
|
||||
{
|
||||
p = NULL;
|
||||
}
|
||||
@ -7242,6 +7378,23 @@ Image2D & Image2D::operator = (const Image2D & i)
|
||||
return *this;
|
||||
}
|
||||
|
||||
Image2D::Image2D(Image2D&& i) CV_NOEXCEPT
|
||||
{
|
||||
p = i.p;
|
||||
i.p = nullptr;
|
||||
}
|
||||
|
||||
Image2D& Image2D::operator = (Image2D&& i) CV_NOEXCEPT
|
||||
{
|
||||
if (this != &i) {
|
||||
if (p)
|
||||
p->release();
|
||||
p = i.p;
|
||||
i.p = nullptr;
|
||||
}
|
||||
return *this;
|
||||
}
|
||||
|
||||
Image2D::~Image2D()
|
||||
{
|
||||
if (p)
|
||||
|
||||
@ -38,6 +38,8 @@ Device::Device() : p(NULL) { }
|
||||
Device::Device(void* d) : p(NULL) { OCL_NOT_AVAILABLE(); }
|
||||
Device::Device(const Device& d) : p(NULL) { }
|
||||
Device& Device::operator=(const Device& d) { return *this; }
|
||||
Device::Device(Device&&) CV_NOEXCEPT : p(NULL) { }
|
||||
Device& Device::operator=(Device&&) CV_NOEXCEPT { return *this; }
|
||||
Device::~Device() { }
|
||||
|
||||
void Device::set(void* d) { OCL_NOT_AVAILABLE(); }
|
||||
@ -152,6 +154,8 @@ Context::Context(int dtype) : p(NULL) { }
|
||||
Context::~Context() { }
|
||||
Context::Context(const Context& c) : p(NULL) { }
|
||||
Context& Context::operator=(const Context& c) { return *this; }
|
||||
Context::Context(Context&&) CV_NOEXCEPT : p(NULL) { }
|
||||
Context& Context::operator=(Context&&) CV_NOEXCEPT { return *this; }
|
||||
|
||||
bool Context::create() { return false; }
|
||||
bool Context::create(int dtype) { return false; }
|
||||
@ -182,6 +186,8 @@ Platform::Platform() : p(NULL) { }
|
||||
Platform::~Platform() { }
|
||||
Platform::Platform(const Platform&) : p(NULL) { }
|
||||
Platform& Platform::operator=(const Platform&) { return *this; }
|
||||
Platform::Platform(Platform&&) CV_NOEXCEPT : p(NULL) { }
|
||||
Platform& Platform::operator=(Platform&&) CV_NOEXCEPT { return *this; }
|
||||
|
||||
void* Platform::ptr() const { return NULL; }
|
||||
|
||||
@ -203,6 +209,8 @@ Queue::Queue(const Context& c, const Device& d) : p(NULL) { OCL_NOT_AVAILABLE();
|
||||
Queue::~Queue() { }
|
||||
Queue::Queue(const Queue& q) {}
|
||||
Queue& Queue::operator=(const Queue& q) { return *this; }
|
||||
Queue::Queue(Queue&&) CV_NOEXCEPT : p(NULL) { }
|
||||
Queue& Queue::operator=(Queue&&) CV_NOEXCEPT { return *this; }
|
||||
|
||||
bool Queue::create(const Context& c, const Device& d) { OCL_NOT_AVAILABLE(); }
|
||||
void Queue::finish() {}
|
||||
@ -218,7 +226,7 @@ Queue& Queue::getDefault()
|
||||
const Queue& Queue::getProfilingQueue() const { OCL_NOT_AVAILABLE(); }
|
||||
|
||||
|
||||
KernelArg::KernelArg()
|
||||
KernelArg::KernelArg() CV_NOEXCEPT
|
||||
: flags(0), m(0), obj(0), sz(0), wscale(1), iwscale(1)
|
||||
{
|
||||
}
|
||||
@ -241,6 +249,8 @@ Kernel::Kernel(const char* kname, const ProgramSource& prog, const String& build
|
||||
Kernel::~Kernel() { }
|
||||
Kernel::Kernel(const Kernel& k) : p(NULL) { }
|
||||
Kernel& Kernel::operator=(const Kernel& k) { return *this; }
|
||||
Kernel::Kernel(Kernel&&) CV_NOEXCEPT : p(NULL) { }
|
||||
Kernel& Kernel::operator=(Kernel&&) CV_NOEXCEPT { return *this; }
|
||||
|
||||
bool Kernel::empty() const { return true; }
|
||||
bool Kernel::create(const char* kname, const Program& prog) { OCL_NOT_AVAILABLE(); }
|
||||
@ -268,6 +278,8 @@ Program::Program() : p(NULL) { }
|
||||
Program::Program(const ProgramSource& src, const String& buildflags, String& errmsg) : p(NULL) { OCL_NOT_AVAILABLE(); }
|
||||
Program::Program(const Program& prog) : p(NULL) { }
|
||||
Program& Program::operator=(const Program& prog) { return *this; }
|
||||
Program::Program(Program&&) CV_NOEXCEPT : p(NULL) { }
|
||||
Program& Program::operator=(Program&&) CV_NOEXCEPT { return *this; }
|
||||
Program::~Program() { }
|
||||
|
||||
bool Program::create(const ProgramSource& src, const String& buildflags, String& errmsg) { OCL_NOT_AVAILABLE(); }
|
||||
@ -290,6 +302,8 @@ ProgramSource::ProgramSource(const char* prog) : p(NULL) { }
|
||||
ProgramSource::~ProgramSource() { }
|
||||
ProgramSource::ProgramSource(const ProgramSource& prog) : p(NULL) { }
|
||||
ProgramSource& ProgramSource::operator=(const ProgramSource& prog) { return *this; }
|
||||
ProgramSource::ProgramSource(ProgramSource&&) CV_NOEXCEPT : p(NULL) { }
|
||||
ProgramSource& ProgramSource::operator=(ProgramSource&&) CV_NOEXCEPT { return *this; }
|
||||
|
||||
const String& ProgramSource::source() const { OCL_NOT_AVAILABLE(); }
|
||||
ProgramSource::hash_t ProgramSource::hash() const { OCL_NOT_AVAILABLE(); }
|
||||
@ -304,6 +318,8 @@ PlatformInfo::~PlatformInfo() { }
|
||||
|
||||
PlatformInfo::PlatformInfo(const PlatformInfo& i) : p(NULL) { }
|
||||
PlatformInfo& PlatformInfo::operator=(const PlatformInfo& i) { return *this; }
|
||||
PlatformInfo::PlatformInfo(PlatformInfo&&) CV_NOEXCEPT : p(NULL) { }
|
||||
PlatformInfo& PlatformInfo::operator=(PlatformInfo&&) CV_NOEXCEPT { return *this; }
|
||||
|
||||
String PlatformInfo::name() const { OCL_NOT_AVAILABLE(); }
|
||||
String PlatformInfo::vendor() const { OCL_NOT_AVAILABLE(); }
|
||||
@ -341,11 +357,13 @@ int predictOptimalVectorWidthMax(InputArray src1, InputArray src2, InputArray sr
|
||||
void buildOptionsAddMatrixDescription(String& buildOptions, const String& name, InputArray _m) { OCL_NOT_AVAILABLE(); }
|
||||
|
||||
|
||||
Image2D::Image2D() : p(NULL) { }
|
||||
Image2D::Image2D() CV_NOEXCEPT : p(NULL) { }
|
||||
Image2D::Image2D(const UMat &src, bool norm, bool alias) { OCL_NOT_AVAILABLE(); }
|
||||
Image2D::Image2D(const Image2D & i) : p(NULL) { OCL_NOT_AVAILABLE(); }
|
||||
Image2D::~Image2D() { }
|
||||
Image2D& Image2D::operator=(const Image2D & i) { return *this; }
|
||||
Image2D::Image2D(Image2D&&) CV_NOEXCEPT : p(NULL) { }
|
||||
Image2D& Image2D::operator=(Image2D&&) CV_NOEXCEPT { return *this; }
|
||||
|
||||
/* static */ bool Image2D::canCreateAlias(const UMat &u) { OCL_NOT_AVAILABLE(); }
|
||||
/* static */ bool Image2D::isFormatSupported(int depth, int cn, bool norm) { OCL_NOT_AVAILABLE(); }
|
||||
|
||||
@ -127,4 +127,120 @@ TEST(OpenCL, support_SPIR_programs)
|
||||
testOpenCLKernel(k);
|
||||
}
|
||||
|
||||
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
|
||||
|
||||
Loading…
Reference in New Issue
Block a user