core(OpenCL): thread-local OpenCL execution context

modules/core/include/opencv2/core/mat.hpp

@@ -565,6 +565,7 @@ struct CV_EXPORTS UMatData
     int allocatorFlags_;
     int mapcount;
     UMatData* originalUMatData;
+    std::shared_ptr<void> allocatorContext;
 };
 CV_ENUM_FLAGS(UMatData::MemoryFlag)
 

modules/core/include/opencv2/core/ocl.hpp

@@ -229,8 +229,15 @@ public:
 
     CV_WRAP static const Device& getDefault();
 
-protected:
+    /**
+     * @param d OpenCL handle (cl_device_id). clRetainDevice() is called on success.
+     */
+    static Device fromHandle(void* d);
+
     struct Impl;
+    inline Impl* getImpl() const { return (Impl*)p; }
+    inline bool empty() const { return !p; }
+protected:
     Impl* p;
 };
 
@@ -239,33 +246,55 @@ class CV_EXPORTS Context
 {
 public:
     Context();
-    explicit Context(int dtype);
+    explicit Context(int dtype);  //!< @deprecated
     ~Context();
     Context(const Context& c);
-    Context& operator = (const Context& c);
+    Context& operator= (const Context& c);
 
+    /** @deprecated */
     bool create();
+    /** @deprecated */
     bool create(int dtype);
+
     size_t ndevices() const;
-    const Device& device(size_t idx) const;
+    Device& device(size_t idx) const;
     Program getProg(const ProgramSource& prog,
                     const String& buildopt, String& errmsg);
     void unloadProg(Program& prog);
 
+
+    /** Get thread-local OpenCL context (initialize if necessary) */
+#if 0  // OpenCV 5.0
+    static Context& getDefault();
+#else
     static Context& getDefault(bool initialize = true);
+#endif
+
+    /** @returns cl_context value */
     void* ptr() const;
 
-    friend void initializeContextFromHandle(Context& ctx, void* platform, void* context, void* device);
 
     bool useSVM() const;
     void setUseSVM(bool enabled);
 
+    /**
+     * @param context OpenCL handle (cl_context). clRetainContext() is called on success
+     */
+    static Context fromHandle(void* context);
+    static Context fromDevice(const ocl::Device& device);
+    static Context create(const std::string& configuration);
+
+    void release();
+
     struct Impl;
     inline Impl* getImpl() const { return (Impl*)p; }
+    inline bool empty() const { return !p; }
+// TODO OpenCV 5.0
 //protected:
     Impl* p;
 };
 
+/** @deprecated */
 class CV_EXPORTS Platform
 {
 public:
@@ -275,11 +304,14 @@ public:
     Platform& operator = (const Platform& p);
 
     void* ptr() const;
+
+    /** @deprecated */
     static Platform& getDefault();
 
-    friend void initializeContextFromHandle(Context& ctx, void* platform, void* context, void* device);
-protected:
     struct Impl;
+    inline Impl* getImpl() const { return (Impl*)p; }
+    inline bool empty() const { return !p; }
+protected:
     Impl* p;
 };
 
@@ -319,6 +351,7 @@ CV_EXPORTS void convertFromBuffer(void* cl_mem_buffer, size_t step, int rows, in
 CV_EXPORTS void convertFromImage(void* cl_mem_image, UMat& dst);
 
 // TODO Move to internal header
+/// @deprecated
 void initializeContextFromHandle(Context& ctx, void* platform, void* context, void* device);
 
 class CV_EXPORTS Queue
@@ -340,6 +373,7 @@ public:
 
     struct Impl; friend struct Impl;
     inline Impl* getImpl() const { return p; }
+    inline bool empty() const { return !p; }
 protected:
     Impl* p;
 };
@@ -490,6 +524,7 @@ public:
 
     struct Impl; friend struct Impl;
     inline Impl* getImpl() const { return (Impl*)p; }
+    inline bool empty() const { return !p; }
 protected:
     Impl* p;
 public:
@@ -571,6 +606,7 @@ public:
 
     struct Impl; friend struct Impl;
     inline Impl* getImpl() const { return (Impl*)p; }
+    inline bool empty() const { return !p; }
 protected:
     Impl* p;
 };
@@ -579,6 +615,9 @@ class CV_EXPORTS PlatformInfo
 {
 public:
     PlatformInfo();
+    /**
+     * @param id pointer cl_platform_id (cl_platform_id*)
+     */
     explicit PlatformInfo(void* id);
     ~PlatformInfo();
 
@@ -591,8 +630,9 @@ public:
     int deviceNumber() const;
     void getDevice(Device& device, int d) const;
 
-protected:
     struct Impl;
+    bool empty() const { return !p; }
+protected:
     Impl* p;
 };
 
@@ -689,6 +729,106 @@ private:
 CV_EXPORTS MatAllocator* getOpenCLAllocator();
 
 
+class CV_EXPORTS_W OpenCLExecutionContext
+{
+public:
+    OpenCLExecutionContext() = default;
+    ~OpenCLExecutionContext() = default;
+
+    OpenCLExecutionContext(const OpenCLExecutionContext& other) = default;
+    OpenCLExecutionContext(OpenCLExecutionContext&& other) = default;
+
+    OpenCLExecutionContext& operator=(const OpenCLExecutionContext& other) = default;
+    OpenCLExecutionContext& operator=(OpenCLExecutionContext&& other) = default;
+
+    /** Get associated ocl::Context */
+    Context& getContext() const;
+    /** Get associated ocl::Device */
+    Device& getDevice() const;
+    /** Get associated ocl::Queue */
+    Queue& getQueue() const;
+
+    bool useOpenCL() const;
+    void setUseOpenCL(bool flag);
+
+    /** Get OpenCL execution context of current thread.
+     *
+     * Initialize OpenCL execution context if it is empty
+     * - create new
+     * - reuse context of the main thread (threadID = 0)
+     */
+    static OpenCLExecutionContext& getCurrent();
+
+    /** Get OpenCL execution context of current thread (can be empty) */
+    static OpenCLExecutionContext& getCurrentRef();
+
+    /** Bind this OpenCL execution context to current thread.
+     *
+     * Context can't be empty.
+     *
+     * @note clFinish is not called for queue of previous execution context
+     */
+    void bind() const;
+
+    /** Creates new execution context with same OpenCV context and device
+     *
+     * @param q OpenCL queue
+     */
+    OpenCLExecutionContext cloneWithNewQueue(const ocl::Queue& q) const;
+    /** @overload */
+    OpenCLExecutionContext cloneWithNewQueue() const;
+
+    /** @brief Creates OpenCL execution context
+     * OpenCV will check if available OpenCL platform has platformName name, then assign context to
+     * OpenCV and call `clRetainContext` function. The deviceID device will be used as target device and
+     * new command queue will be created.
+     *
+     * @note Lifetime of passed handles is transferred to OpenCV wrappers on success
+     *
+     * @param platformName name of OpenCL platform to attach, this string is used to check if platform is available to OpenCV at runtime
+     * @param platformID ID of platform attached context was created for (cl_platform_id)
+     * @param context OpenCL context to be attached to OpenCV (cl_context)
+     * @param deviceID OpenCL device (cl_device_id)
+     */
+    static OpenCLExecutionContext create(const std::string& platformName, void* platformID, void* context, void* deviceID);
+
+    /** @brief Creates OpenCL execution context
+     *
+     * @param context non-empty OpenCL context
+     * @param device non-empty OpenCL device (must be a part of context)
+     * @param queue non-empty OpenCL queue for provided context and device
+     */
+    static OpenCLExecutionContext create(const Context& context, const Device& device, const ocl::Queue& queue);
+    /** @overload */
+    static OpenCLExecutionContext create(const Context& context, const Device& device);
+
+    struct Impl;
+    inline bool empty() const { return !p; }
+    void release();
+protected:
+    std::shared_ptr<Impl> p;
+};
+
+class OpenCLExecutionContextScope
+{
+    OpenCLExecutionContext ctx_;
+public:
+    inline OpenCLExecutionContextScope(const OpenCLExecutionContext& ctx)
+    {
+        CV_Assert(!ctx.empty());
+        ctx_ = OpenCLExecutionContext::getCurrentRef();
+        ctx.bind();
+    }
+
+    inline ~OpenCLExecutionContextScope()
+    {
+        if (!ctx_.empty())
+        {
+            ctx_.bind();
+        }
+    }
+};
+
 #ifdef __OPENCV_BUILD
 namespace internal {
 

modules/core/src/directx.cpp

@@ -458,9 +458,22 @@ Context& initializeContextFromD3D11Device(ID3D11Device* pD3D11Device)
         CV_Error(cv::Error::OpenCLInitError, "OpenCL: Can't create context for DirectX interop");
     }
 
-    Context& ctx = Context::getDefault(false);
+    cl_platform_id platform = platforms[found];
-    initializeContextFromHandle(ctx, platforms[found], context, device);
+    std::string platformName = PlatformInfo(platform).name();
-    return ctx;
+
+    OpenCLExecutionContext clExecCtx;
+    try
+    {
+        clExecCtx = OpenCLExecutionContext::create(platformName, platform, context, device);
+    }
+    catch (...)
+    {
+        clReleaseDevice(device);
+        clReleaseContext(context);
+        throw;
+    }
+    clExecCtx.bind();
+    return const_cast<Context&>(clExecCtx.getContext());
 #endif
 }
 
@@ -565,10 +578,22 @@ Context& initializeContextFromD3D10Device(ID3D10Device* pD3D10Device)
             CV_Error(cv::Error::OpenCLInitError, "OpenCL: Can't create context for DirectX interop");
     }
 
+    cl_platform_id platform = platforms[found];
+    std::string platformName = PlatformInfo(platform).name();
 
-    Context& ctx = Context::getDefault(false);
+    OpenCLExecutionContext clExecCtx;
-    initializeContextFromHandle(ctx, platforms[found], context, device);
+    try
-    return ctx;
+    {
+        clExecCtx = OpenCLExecutionContext::create(platformName, platform, context, device);
+    }
+    catch (...)
+    {
+        clReleaseDevice(device);
+        clReleaseContext(context);
+        throw;
+    }
+    clExecCtx.bind();
+    return const_cast<Context&>(clExecCtx.getContext());
 #endif
 }
 
@@ -675,10 +700,23 @@ Context& initializeContextFromDirect3DDevice9Ex(IDirect3DDevice9Ex* pDirect3DDev
             CV_Error(cv::Error::OpenCLInitError, "OpenCL: Can't create context for DirectX interop");
     }
 
-    Context& ctx = Context::getDefault(false);
+    cl_platform_id platform = platforms[found];
-    initializeContextFromHandle(ctx, platforms[found], context, device);
+    std::string platformName = PlatformInfo(platform).name();
+
+    OpenCLExecutionContext clExecCtx;
+    try
+    {
+        clExecCtx = OpenCLExecutionContext::create(platformName, platform, context, device);
+    }
+    catch (...)
+    {
+        clReleaseDevice(device);
+        clReleaseContext(context);
+        throw;
+    }
+    clExecCtx.bind();
     g_isDirect3DDevice9Ex = true;
-    return ctx;
+    return const_cast<Context&>(clExecCtx.getContext());
 #endif
 }
 
@@ -785,10 +823,23 @@ Context& initializeContextFromDirect3DDevice9(IDirect3DDevice9* pDirect3DDevice9
             CV_Error(cv::Error::OpenCLInitError, "OpenCL: Can't create context for DirectX interop");
     }
 
-    Context& ctx = Context::getDefault(false);
+    cl_platform_id platform = platforms[found];
-    initializeContextFromHandle(ctx, platforms[found], context, device);
+    std::string platformName = PlatformInfo(platform).name();
+
+    OpenCLExecutionContext clExecCtx;
+    try
+    {
+        clExecCtx = OpenCLExecutionContext::create(platformName, platform, context, device);
+    }
+    catch (...)
+    {
+        clReleaseDevice(device);
+        clReleaseContext(context);
+        throw;
+    }
+    clExecCtx.bind();
     g_isDirect3DDevice9Ex = false;
-    return ctx;
+    return const_cast<Context&>(clExecCtx.getContext());
 #endif
 }
 

modules/core/src/ocl_disabled.impl.hpp

@@ -144,6 +144,8 @@ const Device& Device::getDefault()
     return dummy;
 }
 
+/* static */ Device Device::fromHandle(void* d) { OCL_NOT_AVAILABLE(); }
+
 
 Context::Context() : p(NULL) { }
 Context::Context(int dtype) : p(NULL) { }
@@ -154,7 +156,7 @@ Context& Context::operator=(const Context& c) { return *this; }
 bool Context::create() { return false; }
 bool Context::create(int dtype) { return false; }
 size_t Context::ndevices() const { return 0; }
-const Device& Context::device(size_t idx) const { OCL_NOT_AVAILABLE(); }
+Device& Context::device(size_t idx) const { OCL_NOT_AVAILABLE(); }
 Program Context::getProg(const ProgramSource& prog, const String& buildopt, String& errmsg) { OCL_NOT_AVAILABLE(); }
 void Context::unloadProg(Program& prog) { }
 
@@ -169,6 +171,13 @@ void* Context::ptr() const { return NULL; }
 bool Context::useSVM() const { return false; }
 void Context::setUseSVM(bool enabled) { }
 
+/* static */ Context Context::fromHandle(void* context) { OCL_NOT_AVAILABLE(); }
+/* static */ Context Context::fromDevice(const ocl::Device& device) { OCL_NOT_AVAILABLE(); }
+/* static */ Context Context::create(const std::string& configuration) { OCL_NOT_AVAILABLE(); }
+
+void Context::release() { }
+
+
 Platform::Platform() : p(NULL) { }
 Platform::~Platform() { }
 Platform::Platform(const Platform&) : p(NULL) { }
@@ -355,6 +364,43 @@ MatAllocator* getOpenCLAllocator() { return NULL; }
 
 internal::ProgramEntry::operator ProgramSource&() const { OCL_NOT_AVAILABLE(); }
 
+
+struct OpenCLExecutionContext::Impl
+{
+    Impl() = default;
+};
+
+Context& OpenCLExecutionContext::getContext() const { OCL_NOT_AVAILABLE(); }
+Device& OpenCLExecutionContext::getDevice() const { OCL_NOT_AVAILABLE(); }
+Queue& OpenCLExecutionContext::getQueue() const { OCL_NOT_AVAILABLE(); }
+
+bool OpenCLExecutionContext::useOpenCL() const { return false; }
+void OpenCLExecutionContext::setUseOpenCL(bool flag) { }
+
+static
+OpenCLExecutionContext& getDummyOpenCLExecutionContext()
+{
+    static OpenCLExecutionContext dummy;
+    return dummy;
+}
+
+/* static */
+OpenCLExecutionContext& OpenCLExecutionContext::getCurrent() { return getDummyOpenCLExecutionContext(); }
+
+/* static */
+OpenCLExecutionContext& OpenCLExecutionContext::getCurrentRef() { return getDummyOpenCLExecutionContext(); }
+
+void OpenCLExecutionContext::bind() const { OCL_NOT_AVAILABLE(); }
+
+OpenCLExecutionContext OpenCLExecutionContext::cloneWithNewQueue(const ocl::Queue& q) const { OCL_NOT_AVAILABLE(); }
+OpenCLExecutionContext OpenCLExecutionContext::cloneWithNewQueue() const { OCL_NOT_AVAILABLE(); }
+
+/* static */ OpenCLExecutionContext OpenCLExecutionContext::create(const std::string& platformName, void* platformID, void* context, void* deviceID) { OCL_NOT_AVAILABLE(); }
+/* static */ OpenCLExecutionContext OpenCLExecutionContext::create(const Context& context, const Device& device, const ocl::Queue& queue) { OCL_NOT_AVAILABLE(); }
+/* static */ OpenCLExecutionContext OpenCLExecutionContext::create(const Context& context, const Device& device) { OCL_NOT_AVAILABLE(); }
+
+void OpenCLExecutionContext::release() { }
+
 }}
 
 #if defined(_MSC_VER)

modules/core/src/opengl.cpp

@@ -1689,9 +1689,14 @@ Context& initializeContextFromGL()
     if (found < 0)
         CV_Error(cv::Error::OpenCLInitError, "OpenCL: Can't create context for OpenGL interop");
 
-    Context& ctx = Context::getDefault(false);
+    cl_platform_id platform = platforms[found];
-    initializeContextFromHandle(ctx, platforms[found], context, device);
+    std::string platformName = PlatformInfo(platform).name();
-    return ctx;
+
+    OpenCLExecutionContext clExecCtx = OpenCLExecutionContext::create(platformName, platform, context, deviceID);
+    clReleaseDevice(device);
+    clReleaseContext(context);
+    clExecCtx.bind();
+    return const_cast<Context&>(clExecCtx.getContext());
 #endif
 }
 

modules/core/src/precomp.hpp

@@ -322,7 +322,7 @@ struct CoreTLSData
 {
     CoreTLSData() :
 //#ifdef HAVE_OPENCL
-        device(0), useOpenCL(-1),
+        oclExecutionContextInitialized(false), useOpenCL(-1),
 //#endif
         useIPP(-1),
         useIPP_NE(-1)
@@ -333,8 +333,8 @@ struct CoreTLSData
 
     RNG rng;
 //#ifdef HAVE_OPENCL
-    int device; // device index of an array of devices in a context, see also Device::getDefault
+    ocl::OpenCLExecutionContext oclExecutionContext;
-    ocl::Queue oclQueue; // the queue used for running a kernel, see also getQueue, Kernel::run
+    bool oclExecutionContextInitialized;
     int useOpenCL; // 1 - use, 0 - do not use, -1 - auto/not initialized
 //#endif
     int useIPP;    // 1 - use, 0 - do not use, -1 - auto/not initialized

modules/core/src/va_intel.cpp

@@ -106,7 +106,7 @@ Context& initializeContextFromVA(VADisplay display, bool tryInterop)
                                                                CL_PREFERRED_DEVICES_FOR_VA_API_INTEL, 0, NULL, &numDevices);
             if ((status != CL_SUCCESS) || !(numDevices > 0))
                 continue;
-            numDevices = 1; // initializeContextFromHandle() expects only 1 device
+            numDevices = 1; // OpenCV expects only 1 device
             status = clGetDeviceIDsFromVA_APIMediaAdapterINTEL(platforms[i], CL_VA_API_DISPLAY_INTEL, display,
                                                                CL_PREFERRED_DEVICES_FOR_VA_API_INTEL, numDevices, &device, NULL);
             if (status != CL_SUCCESS)
@@ -135,9 +135,23 @@ Context& initializeContextFromVA(VADisplay display, bool tryInterop)
         if (found >= 0)
         {
             contextInitialized = true;
-            Context& ctx = Context::getDefault(false);
+
-            initializeContextFromHandle(ctx, platforms[found], context, device);
+            cl_platform_id platform = platforms[found];
-            return ctx;
+            std::string platformName = PlatformInfo(platform).name();
+
+            OpenCLExecutionContext clExecCtx;
+            try
+            {
+                clExecCtx = OpenCLExecutionContext::create(platformName, platform, context, device);
+            }
+            catch (...)
+            {
+                clReleaseDevice(device);
+                clReleaseContext(context);
+                throw;
+            }
+            clExecCtx.bind();
+            return const_cast<Context&>(clExecCtx.getContext());
         }
     }
 # endif // HAVE_VA_INTEL && HAVE_OPENCL

modules/core/test/test_opencl.cpp

@@ -0,0 +1,191 @@
+// This file is part of OpenCV project.
+// It is subject to the license terms in the LICENSE file found in the top-level directory
+// of this distribution and at http://opencv.org/license.html.
+
+#include "test_precomp.hpp"
+#include "opencv2/ts/ocl_test.hpp"
+
+namespace opencv_test {
+namespace ocl {
+
+static void executeUMatCall(bool requireOpenCL = true)
+{
+    UMat a(100, 100, CV_8UC1, Scalar::all(0));
+    UMat b;
+    cv::add(a, Scalar::all(1), b);
+    Mat b_cpu = b.getMat(ACCESS_READ);
+    EXPECT_EQ(0, cv::norm(b_cpu - 1, NORM_INF));
+
+    if (requireOpenCL)
+    {
+        EXPECT_TRUE(cv::ocl::useOpenCL());
+    }
+}
+
+TEST(OCL_Context, createFromDevice)
+{
+    bool useOCL = cv::ocl::useOpenCL();
+
+    OpenCLExecutionContext ctx = OpenCLExecutionContext::getCurrent();
+
+    if (!useOCL)
+    {
+        ASSERT_TRUE(ctx.empty());  // Other tests should not broke global state
+        throw SkipTestException("OpenCL is not available / disabled");
+    }
+
+    ASSERT_FALSE(ctx.empty());
+
+    ocl::Device device = ctx.getDevice();
+    ASSERT_FALSE(device.empty());
+
+    ocl::Context context = ocl::Context::fromDevice(device);
+    ocl::Context context2 = ocl::Context::fromDevice(device);
+
+    EXPECT_TRUE(context.getImpl() == context2.getImpl()) << "Broken cache for OpenCL context (device)";
+}
+
+TEST(OCL_OpenCLExecutionContext, basic)
+{
+    bool useOCL = cv::ocl::useOpenCL();
+
+    OpenCLExecutionContext ctx = OpenCLExecutionContext::getCurrent();
+
+    if (!useOCL)
+    {
+        ASSERT_TRUE(ctx.empty());  // Other tests should not broke global state
+        throw SkipTestException("OpenCL is not available / disabled");
+    }
+
+    ASSERT_FALSE(ctx.empty());
+
+    ocl::Context context = ctx.getContext();
+    ocl::Context context2 = ocl::Context::getDefault();
+    EXPECT_TRUE(context.getImpl() == context2.getImpl());
+
+    ocl::Device device = ctx.getDevice();
+    ocl::Device device2 = ocl::Device::getDefault();
+    EXPECT_TRUE(device.getImpl() == device2.getImpl());
+
+    ocl::Queue queue = ctx.getQueue();
+    ocl::Queue queue2 = ocl::Queue::getDefault();
+    EXPECT_TRUE(queue.getImpl() == queue2.getImpl());
+}
+
+TEST(OCL_OpenCLExecutionContext, createAndBind)
+{
+    bool useOCL = cv::ocl::useOpenCL();
+
+    OpenCLExecutionContext ctx = OpenCLExecutionContext::getCurrent();
+
+    if (!useOCL)
+    {
+        ASSERT_TRUE(ctx.empty());  // Other tests should not broke global state
+        throw SkipTestException("OpenCL is not available / disabled");
+    }
+
+    ASSERT_FALSE(ctx.empty());
+
+    ocl::Context context = ctx.getContext();
+    ocl::Device device = ctx.getDevice();
+
+    OpenCLExecutionContext ctx2 = OpenCLExecutionContext::create(context, device);
+    ASSERT_FALSE(ctx2.empty());
+
+    try
+    {
+        ctx2.bind();
+        executeUMatCall();
+        ctx.bind();
+        executeUMatCall();
+    }
+    catch (...)
+    {
+        ctx.bind();  // restore
+        throw;
+    }
+}
+
+TEST(OCL_OpenCLExecutionContext, createGPU)
+{
+    bool useOCL = cv::ocl::useOpenCL();
+
+    OpenCLExecutionContext ctx = OpenCLExecutionContext::getCurrent();
+
+    if (!useOCL)
+    {
+        ASSERT_TRUE(ctx.empty());  // Other tests should not broke global state
+        throw SkipTestException("OpenCL is not available / disabled");
+    }
+
+    ASSERT_FALSE(ctx.empty());
+
+    ocl::Context context = ocl::Context::create(":GPU:1");
+    if (context.empty())
+    {
+        context = ocl::Context::create(":CPU:");
+        if (context.empty())
+            throw SkipTestException("OpenCL GPU1/CPU devices are not available");
+    }
+
+    ocl::Device device = context.device(0);
+
+    OpenCLExecutionContext ctx2 = OpenCLExecutionContext::create(context, device);
+    ASSERT_FALSE(ctx2.empty());
+
+    try
+    {
+        ctx2.bind();
+        executeUMatCall();
+        ctx.bind();
+        executeUMatCall();
+    }
+    catch (...)
+    {
+        ctx.bind();  // restore
+        throw;
+    }
+}
+
+TEST(OCL_OpenCLExecutionContext, ScopeTest)
+{
+    bool useOCL = cv::ocl::useOpenCL();
+
+    OpenCLExecutionContext ctx = OpenCLExecutionContext::getCurrent();
+
+    if (!useOCL)
+    {
+        ASSERT_TRUE(ctx.empty());  // Other tests should not broke global state
+        throw SkipTestException("OpenCL is not available / disabled");
+    }
+
+    ASSERT_FALSE(ctx.empty());
+
+    ocl::Context context = ocl::Context::create(":GPU:1");
+    if (context.empty())
+    {
+        context = ocl::Context::create(":CPU:");
+        if (context.empty())
+            context = ctx.getContext();
+    }
+
+    ocl::Device device = context.device(0);
+
+    OpenCLExecutionContext ctx2 = OpenCLExecutionContext::create(context, device);
+    ASSERT_FALSE(ctx2.empty());
+
+    try
+    {
+        OpenCLExecutionContextScope ctx_scope(ctx2);
+        executeUMatCall();
+    }
+    catch (...)
+    {
+        ctx.bind();  // restore
+        throw;
+    }
+
+    executeUMatCall();
+}
+
+} } // namespace opencv_test::ocl

samples/CMakeLists.txt

@@ -26,6 +26,7 @@ add_subdirectory(dnn)
 add_subdirectory(gpu)
 add_subdirectory(tapi)
 add_subdirectory(opencl)
+add_subdirectory(sycl)
 if(WIN32 AND HAVE_DIRECTX)
   add_subdirectory(directx)
 endif()
@@ -122,6 +123,7 @@ endif()
 add_subdirectory(dnn)
 # add_subdirectory(gpu)
 add_subdirectory(opencl)
+add_subdirectory(sycl)
 # add_subdirectory(opengl)
 # add_subdirectory(openvx)
 add_subdirectory(tapi)

samples/sycl/CMakeLists.txt

@@ -0,0 +1,80 @@
+if(OPENCV_SKIP_SAMPLES_SYCL)
+  return()
+endif()
+
+ocv_install_example_src(opencl *.cpp *.hpp CMakeLists.txt)
+
+set(OPENCV_SYCL_SAMPLES_REQUIRED_DEPS
+  opencv_core
+  opencv_imgproc
+  opencv_imgcodecs
+  opencv_videoio
+  opencv_highgui)
+ocv_check_dependencies(${OPENCV_SYCL_SAMPLES_REQUIRED_DEPS})
+
+if(NOT BUILD_EXAMPLES OR NOT OCV_DEPENDENCIES_FOUND OR OPENCV_SKIP_SAMPLES_BUILD_SYCL)
+  return()
+endif()
+
+if(CMAKE_VERSION VERSION_LESS "3.5")
+  message(STATUS "SYCL samples require CMake 3.5+")
+  return()
+endif()
+
+cmake_policy(VERSION 3.5)
+
+find_package(SYCL QUIET)  # will oneAPI support this straightforward way?
+
+if(NOT SYCL_FOUND AND NOT OPENCV_SKIP_SAMPLES_SYCL_ONEDNN)
+  # lets try scripts from oneAPI:oneDNN component
+  if(NOT DEFINED DNNLROOT AND DEFINED ENV{DNNLROOT})
+    set(DNNLROOT "$ENV{DNNLROOT}")
+  endif()
+  # Some verions of called script violate CMake policy and may emit unrecoverable CMake errors
+  # Use OPENCV_SKIP_SAMPLES_SYCL=1 / OPENCV_SKIP_SAMPLES_SYCL_ONEDNN to bypass this
+  find_package(dnnl CONFIG QUIET HINTS "${DNNLROOT}")
+endif()
+
+if(NOT SYCL_FOUND AND NOT OPENCV_SKIP_SAMPLES_SYCL_COMPUTECPP)
+  # lets try this SYCL SDK too: https://github.com/codeplaysoftware/computecpp-sdk
+  find_package(ComputeCpp QUIET)
+  if(ComputeCpp_FOUND)
+    set(SYCL_TARGET ComputeCpp::ComputeCpp)
+    set(SYCL_FLAGS ${ComputeCpp_FLAGS})
+    set(SYCL_INCLUDE_DIRS ${ComputeCpp_INCLUDE_DIRS})
+    set(SYCL_LIBRARIES ${ComputeCpp_LIBRARIES})
+  endif()
+endif()
+
+if(OPENCV_CMAKE_DEBUG_SYCL)
+  ocv_cmake_dump_vars("SYCL")  # OpenCV source tree is required
+endif()
+
+if(NOT SYCL_TARGET)
+  message(STATUS "SYCL/OpenCL samples are skipped: SYCL SDK is required")
+  message(STATUS "   - check configuration of SYCL_DIR/SYCL_ROOT/CMAKE_MODULE_PATH")
+  message(STATUS "   - ensure that right compiler is selected from SYCL SDK (e.g, clang++): CMAKE_CXX_COMPILER=${CMAKE_CXX_COMPILER}")
+  return()
+endif()
+
+project(sycl_samples)
+
+if(SYCL_FLAGS)  # "target_link_libraries(... ${SYCL_TARGET})" is not enough. Hacking...
+  set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${SYCL_FLAGS}")
+  set(CMAKE_EXE_LINKER_FLAGS "${CMAKE_EXE_LINKER_FLAGS} ${SYCL_FLAGS}")
+endif()
+
+ocv_include_modules_recurse(${OPENCV_SYCL_SAMPLES_REQUIRED_DEPS})
+ocv_include_directories(${OpenCL_INCLUDE_DIR})
+file(GLOB all_samples RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} *.cpp)
+foreach(sample_filename ${all_samples})
+  ocv_define_sample(tgt ${sample_filename} sycl)
+  ocv_target_link_libraries(${tgt} PRIVATE
+    ${OPENCV_LINKER_LIBS}
+    ${OPENCV_SYCL_SAMPLES_REQUIRED_DEPS}
+    ${SYCL_TARGET})
+
+  if(COMMAND add_sycl_to_target)  # ComputeCpp
+    add_sycl_to_target(TARGET ${tgt} SOURCES ${sample_filename})
+  endif()
+endforeach()

samples/sycl/sycl-opencv-interop.cpp

@@ -0,0 +1,351 @@
+/*
+ * The example of interoperability between SYCL/OpenCL and OpenCV.
+ * - SYCL: https://www.khronos.org/sycl/
+ * - SYCL runtime parameters: https://github.com/intel/llvm/blob/sycl/sycl/doc/EnvironmentVariables.md
+ */
+#include <CL/sycl.hpp>
+
+#include <opencv2/core.hpp>
+#include <opencv2/highgui.hpp>
+#include <opencv2/videoio.hpp>
+#include <opencv2/imgproc.hpp>
+
+#include <opencv2/core/ocl.hpp>
+
+
+class sycl_inverse_kernel;  // can be omitted - modern SYCL versions doesn't require this
+
+using namespace cv;
+
+
+class App
+{
+public:
+    App(const CommandLineParser& cmd);
+    ~App();
+
+    void initVideoSource();
+
+    void initSYCL();
+
+    void process_frame(cv::Mat& frame);
+
+    /// to check result with CPU-only reference code
+    Mat process_frame_reference(const cv::Mat& frame);
+
+    int run();
+
+    bool isRunning() { return m_running; }
+    bool doProcess() { return m_process; }
+
+    void setRunning(bool running)      { m_running = running; }
+    void setDoProcess(bool process)    { m_process = process; }
+
+protected:
+    void handleKey(char key);
+
+private:
+    bool                        m_running;
+    bool                        m_process;
+    bool                        m_show_ui;
+
+    int64                       m_t0;
+    int64                       m_t1;
+    float                       m_time;
+    float                       m_frequency;
+
+    std::string                 m_file_name;
+    int                         m_camera_id;
+    cv::VideoCapture            m_cap;
+    cv::Mat                     m_frame;
+
+    cl::sycl::queue sycl_queue;
+};
+
+
+App::App(const CommandLineParser& cmd)
+{
+    m_camera_id  = cmd.get<int>("camera");
+    m_file_name  = cmd.get<std::string>("video");
+
+    m_running    = false;
+    m_process    = false;
+} // ctor
+
+
+App::~App()
+{
+    // nothing
+}
+
+
+void App::initSYCL()
+{
+    using namespace cl::sycl;
+
+    // Configuration details: https://github.com/intel/llvm/blob/sycl/sycl/doc/EnvironmentVariables.md
+    cl::sycl::default_selector selector;
+
+    sycl_queue = cl::sycl::queue(selector, [](cl::sycl::exception_list l)
+    {
+        // exception_handler
+        for (auto ep : l)
+        {
+            try
+            {
+                std::rethrow_exception(ep);
+            }
+            catch (const cl::sycl::exception& e)
+            {
+                std::cerr << "SYCL exception: " << e.what() << std::endl;
+            }
+        }
+    });
+
+    auto device = sycl_queue.get_device();
+    auto platform = device.get_platform();
+    std::cout << "SYCL device: " << device.get_info<info::device::name>()
+        << " @ " << device.get_info<info::device::driver_version>()
+        << " (platform: " << platform.get_info<info::platform::name>() << ")" << std::endl;
+
+    if (device.is_host())
+    {
+        std::cerr << "SYCL can't select OpenCL device. Host is used for computations, interoperability is not available" << std::endl;
+    }
+    else
+    {
+        // bind OpenCL context/device/queue from SYCL to OpenCV
+        try
+        {
+            auto ctx = cv::ocl::OpenCLExecutionContext::create(
+                    platform.get_info<info::platform::name>(),
+                    platform.get(),
+                    sycl_queue.get_context().get(),
+                    device.get()
+                );
+            ctx.bind();
+        }
+        catch (const cv::Exception& e)
+        {
+            std::cerr << "OpenCV: Can't bind SYCL OpenCL context/device/queue: " << e.what() << std::endl;
+        }
+        std::cout << "OpenCV uses OpenCL: " << (cv::ocl::useOpenCL() ? "True" : "False") << std::endl;
+    }
+} // initSYCL()
+
+
+void App::initVideoSource()
+{
+    if (!m_file_name.empty() && m_camera_id == -1)
+    {
+        m_cap.open(samples::findFileOrKeep(m_file_name));
+        if (!m_cap.isOpened())
+            throw std::runtime_error(std::string("can't open video stream: ") + m_file_name);
+    }
+    else if (m_camera_id != -1)
+    {
+        m_cap.open(m_camera_id);
+        if (!m_cap.isOpened())
+            throw std::runtime_error(std::string("can't open camera: ") + std::to_string(m_camera_id));
+    }
+    else
+        throw std::runtime_error(std::string("specify video source"));
+} // initVideoSource()
+
+
+void App::process_frame(cv::Mat& frame)
+{
+    using namespace cl::sycl;
+
+    // cv::Mat => cl::sycl::buffer
+    {
+        CV_Assert(frame.isContinuous());
+        CV_CheckTypeEQ(frame.type(), CV_8UC1, "");
+
+        buffer<uint8_t, 2> frame_buffer(frame.data, range<2>(frame.rows, frame.cols));
+
+        // done automatically: frame_buffer.set_write_back(true);
+
+        sycl_queue.submit([&](handler& cgh) {
+          auto pixels = frame_buffer.get_access<access::mode::read_write>(cgh);
+
+          cgh.parallel_for<class sycl_inverse_kernel>(range<2>(frame.rows, frame.cols), [=](item<2> item) {
+              uint8_t v = pixels[item];
+              pixels[item] = ~v;
+          });
+        });
+
+        sycl_queue.wait_and_throw();
+    }
+
+    // No way to extract cl_mem from cl::sycl::buffer (ref: 3.6.11 "Interfacing with OpenCL" of SYCL 1.2.1)
+    // We just reusing OpenCL context/device/queue from SYCL here (see initSYCL() bind part) and call UMat processing
+    {
+        UMat blurResult;
+        {
+            UMat umat_buffer = frame.getUMat(ACCESS_RW);
+            cv::blur(umat_buffer, blurResult, Size(3, 3));  // UMat doesn't support inplace
+        }
+        Mat result;
+        blurResult.copyTo(result);
+        swap(result, frame);
+    }
+}
+
+Mat App::process_frame_reference(const cv::Mat& frame)
+{
+    Mat result;
+    cv::bitwise_not(frame, result);
+    Mat blurResult;
+    cv::blur(result, blurResult, Size(3, 3));  // avoid inplace
+    blurResult.copyTo(result);
+    return result;
+}
+
+int App::run()
+{
+    std::cout << "Initializing..." << std::endl;
+
+    initSYCL();
+    initVideoSource();
+
+    std::cout << "Press ESC to exit" << std::endl;
+    std::cout << "      'p' to toggle ON/OFF processing" << std::endl;
+
+    m_running = true;
+    m_process = true;
+    m_show_ui = true;
+
+    int processedFrames = 0;
+
+    cv::TickMeter timer;
+
+    // Iterate over all frames
+    while (isRunning() && m_cap.read(m_frame))
+    {
+        Mat m_frameGray;
+        cvtColor(m_frame, m_frameGray, COLOR_BGR2GRAY);
+
+        bool checkWithReference = (processedFrames == 0);
+        Mat reference_result;
+        if (checkWithReference)
+        {
+            reference_result = process_frame_reference(m_frameGray);
+        }
+
+        timer.reset();
+        timer.start();
+
+        if (m_process)
+        {
+            process_frame(m_frameGray);
+        }
+
+        timer.stop();
+
+        if (checkWithReference)
+        {
+            double diffInf = cv::norm(reference_result, m_frameGray, NORM_INF);
+            if (diffInf > 0)
+            {
+                std::cerr << "Result is not accurate. diffInf=" << diffInf << std::endl;
+                imwrite("reference.png", reference_result);
+                imwrite("actual.png", m_frameGray);
+            }
+        }
+
+        Mat img_to_show = m_frameGray;
+
+        std::ostringstream msg;
+        msg << "Frame " << processedFrames << " (" << m_frame.size
+            << ")   Time: " << cv::format("%.2f", timer.getTimeMilli()) << " msec"
+            << " (process: " << (m_process ? "True" : "False") << ")";
+        std::cout << msg.str() << std::endl;
+        putText(img_to_show, msg.str(), Point(5, 150), FONT_HERSHEY_SIMPLEX, 1., Scalar(255, 100, 0), 2);
+
+        if (m_show_ui)
+        {
+            try
+            {
+                imshow("sycl_interop", img_to_show);
+                int key = waitKey(1);
+                switch (key)
+                {
+                case 27:  // ESC
+                    m_running = false;
+                    break;
+
+                case 'p':  // fallthru
+                case 'P':
+                    m_process = !m_process;
+                    break;
+
+                default:
+                    break;
+                }
+            }
+            catch (const std::exception& e)
+            {
+                std::cerr << "ERROR(OpenCV UI): " << e.what() << std::endl;
+                if (processedFrames > 0)
+                    throw;
+                m_show_ui = false;  // UI is not available
+            }
+        }
+
+        processedFrames++;
+
+        if (!m_show_ui)
+        {
+            if (processedFrames > 100)
+                m_running = false;
+        }
+    }
+
+    return 0;
+}
+
+
+int main(int argc, char** argv)
+{
+    const char* keys =
+        "{ help h ?    |          | print help message }"
+        "{ camera c    | -1       | use camera as input }"
+        "{ video  v    |          | use video as input }";
+
+    CommandLineParser cmd(argc, argv, keys);
+    if (cmd.has("help"))
+    {
+        cmd.printMessage();
+        return EXIT_SUCCESS;
+    }
+
+    try
+    {
+        App app(cmd);
+        if (!cmd.check())
+        {
+            cmd.printErrors();
+            return 1;
+        }
+        app.run();
+    }
+    catch (const cv::Exception& e)
+    {
+        std::cout << "FATAL: OpenCV error: " << e.what() << std::endl;
+        return 1;
+    }
+    catch (const std::exception& e)
+    {
+        std::cout << "FATAL: C++ error: " << e.what() << std::endl;
+        return 1;
+    }
+
+    catch (...)
+    {
+        std::cout << "FATAL: unknown C++ exception" << std::endl;
+        return 1;
+    }
+
+    return EXIT_SUCCESS;
+} // main()