From 96a15434a28d2e66ee1a4e9f96300d4b9d1f4acc Mon Sep 17 00:00:00 2001 From: Dale Phurrough Date: Sat, 20 Feb 2021 18:56:04 +0100 Subject: [PATCH] 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 --- modules/core/include/opencv2/core/ocl.hpp | 23 +++- modules/core/src/ocl.cpp | 157 +++++++++++++++++++++- modules/core/src/ocl_disabled.impl.hpp | 22 ++- modules/core/test/ocl/test_opencl.cpp | 116 ++++++++++++++++ 4 files changed, 311 insertions(+), 7 deletions(-) diff --git a/modules/core/include/opencv2/core/ocl.hpp b/modules/core/include/opencv2/core/ocl.hpp index 66d7a20ad1..9b10bd608d 100644 --- a/modules/core/include/opencv2/core/ocl.hpp +++ b/modules/core/include/opencv2/core/ocl.hpp @@ -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. diff --git a/modules/core/src/ocl.cpp b/modules/core/src/ocl.cpp index 53b21e0f77..7ea1f129f7 100644 --- a/modules/core/src/ocl.cpp +++ b/modules/core/src/ocl.cpp @@ -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) diff --git a/modules/core/src/ocl_disabled.impl.hpp b/modules/core/src/ocl_disabled.impl.hpp index 97c3856b37..dcb46e7941 100644 --- a/modules/core/src/ocl_disabled.impl.hpp +++ b/modules/core/src/ocl_disabled.impl.hpp @@ -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(); } diff --git a/modules/core/test/ocl/test_opencl.cpp b/modules/core/test/ocl/test_opencl.cpp index 27cd82d424..9798273914 100644 --- a/modules/core/test/ocl/test_opencl.cpp +++ b/modules/core/test/ocl/test_opencl.cpp @@ -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 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