OCV_OPTION(WITH_OPENCL "Include OpenCL Runtime support" ON IF (NOT IOS) )
OCV_OPTION(WITH_OPENCLAMDFFT "Include AMD OpenCL FFT library support" ON IF (NOT ANDROID AND NOT IOS) )
OCV_OPTION(WITH_OPENCLAMDBLAS "Include AMD OpenCL BLAS library support" ON IF (NOT ANDROID AND NOT IOS) )
+OCV_OPTION(WITH_DIRECTX "Include DirectX support" ON IF WIN32 )
# OpenCV build components
include(cmake/OpenCVDetectOpenCL.cmake)
endif()
+# --- DirectX ---
+if(WITH_DIRECTX)
+ include(cmake/OpenCVDetectDirectX.cmake)
+endif()
+
# --- Matlab/Octave ---
include(cmake/OpenCVFindMatlab.cmake)
--- /dev/null
+if(WIN32)
+ try_compile(__VALID_DIRECTX
+ "${OpenCV_BINARY_DIR}"
+ "${OpenCV_SOURCE_DIR}/cmake/checks/directx.cpp"
+ OUTPUT_VARIABLE TRY_OUT
+ )
+ if(NOT __VALID_DIRECTX)
+ return()
+ endif()
+ set(HAVE_DIRECTX ON)
+ set(HAVE_D3D11 ON)
+ set(HAVE_D3D10 ON)
+ set(HAVE_D3D9 ON)
+endif()
--- /dev/null
+#include <windows.h>
+
+#include <d3d11.h>
+#pragma comment (lib, "d3d11.lib")
+
+HINSTANCE g_hInst = NULL;
+D3D_DRIVER_TYPE g_driverType = D3D_DRIVER_TYPE_NULL;
+D3D_FEATURE_LEVEL g_featureLevel = D3D_FEATURE_LEVEL_11_0;
+ID3D11Device* g_pd3dDevice = NULL;
+ID3D11DeviceContext* g_pImmediateContext = NULL;
+IDXGISwapChain* g_pSwapChain = NULL;
+
+static HRESULT InitDevice()
+{
+ HRESULT hr = S_OK;
+
+ UINT width = 640;
+ UINT height = 480;
+
+ UINT createDeviceFlags = 0;
+
+ D3D_DRIVER_TYPE driverTypes[] =
+ {
+ D3D_DRIVER_TYPE_HARDWARE,
+ D3D_DRIVER_TYPE_WARP,
+ D3D_DRIVER_TYPE_REFERENCE,
+ };
+ UINT numDriverTypes = ARRAYSIZE(driverTypes);
+
+ D3D_FEATURE_LEVEL featureLevels[] =
+ {
+ D3D_FEATURE_LEVEL_11_0,
+ D3D_FEATURE_LEVEL_10_1,
+ D3D_FEATURE_LEVEL_10_0,
+ };
+ UINT numFeatureLevels = ARRAYSIZE(featureLevels);
+
+ DXGI_SWAP_CHAIN_DESC sd;
+ ZeroMemory( &sd, sizeof( sd ) );
+ sd.BufferCount = 1;
+ sd.BufferDesc.Width = width;
+ sd.BufferDesc.Height = height;
+ sd.BufferDesc.Format = DXGI_FORMAT_R8G8B8A8_UNORM;
+ sd.BufferDesc.RefreshRate.Numerator = 60;
+ sd.BufferDesc.RefreshRate.Denominator = 1;
+ sd.BufferUsage = DXGI_USAGE_RENDER_TARGET_OUTPUT;
+ sd.OutputWindow = NULL; //g_hWnd;
+ sd.SampleDesc.Count = 1;
+ sd.SampleDesc.Quality = 0;
+ sd.Windowed = TRUE;
+
+ for (UINT driverTypeIndex = 0; driverTypeIndex < numDriverTypes; driverTypeIndex++)
+ {
+ g_driverType = driverTypes[driverTypeIndex];
+ hr = D3D11CreateDeviceAndSwapChain(NULL, g_driverType, NULL, createDeviceFlags, featureLevels, numFeatureLevels,
+ D3D11_SDK_VERSION, &sd, &g_pSwapChain, &g_pd3dDevice, &g_featureLevel, &g_pImmediateContext);
+ if (SUCCEEDED(hr))
+ break;
+ }
+ if (FAILED(hr))
+ return hr;
+
+ return S_OK;
+}
+
+int main(int /*argc*/, char** /*argv*/)
+{
+ InitDevice();
+ return 0;
+}
/* IEEE1394 capturing support - libdc1394 v2.x */
#cmakedefine HAVE_DC1394_2
+/* DirectX */
+#cmakedefine HAVE_DIRECTX
+#cmakedefine HAVE_D3D11
+#cmakedefine HAVE_D3D10
+#cmakedefine HAVE_D3D9
+
/* DirectShow Video Capture library */
#cmakedefine HAVE_DSHOW
.. |Author_EricCh| unicode:: Eric U+0020 Christiansen
.. |Author_AndreyP| unicode:: Andrey U+0020 Pavlenko
.. |Author_AlexS| unicode:: Alexander U+0020 Smorkalov
+.. |Author_MimmoC| unicode:: Mimmo U+0020 Cosenza
.. |Author_BarisD| unicode:: Bar U+0131 U+015F U+0020 Evrim U+0020 Demir U+00F6 z
.. |Author_DomenicoB| unicode:: Domenico U+0020 Daniele U+0020 Bloisi
--- /dev/null
+.. _clojure_dev_intro:
+
+Introduction to OpenCV Development with Clojure
+***********************************************
+
+As of OpenCV 2.4.4, OpenCV supports desktop Java development using
+nearly the same interface as for Android development.
+
+`Clojure <http://clojure.org/>`_ is a contemporary LISP dialect hosted
+by the Java Virtual Machine and it offers a complete interoperability
+with the underlying JVM. This means that we should even be able to use
+the Clojure REPL (Read Eval Print Loop) as and interactive programmable
+interface to the underlying OpenCV engine.
+
+What we'll do in this tutorial
+==============================
+
+This tutorial will help you in setting up a basic Clojure environment
+for interactively learning OpenCV within the fully programmable
+CLojure REPL.
+
+Tutorial source code
+--------------------
+
+You can find a runnable source code of the sample in the
+:file:`samples/java/clojure/simple-sample` folder of the OpenCV
+repository. After having installed OpenCV and Clojure as explained in
+the tutorial, issue the following command to run the sample from the
+command line.
+
+.. code:: bash
+
+ cd path/to/samples/java/clojure/simple-sample
+ lein run
+
+Preamble
+========
+
+For detailed instruction on installing OpenCV with desktop Java support
+refer to the `corresponding tutorial <http://docs.opencv.org/2.4.4-beta/doc/tutorials/introduction/desktop_java/java_dev_intro.html>`_.
+
+If you are in hurry, here is a minimum quick start guide to install
+OpenCV on Mac OS X:
+
+ NOTE 1: I'm assuming you already installed
+ `xcode <https://developer.apple.com/xcode/>`_,
+ `jdk <http://www.oracle.com/technetwork/java/javase/downloads/index.html>`_
+ and `Cmake <http://www.cmake.org/cmake/resources/software.html>`_.
+
+.. code:: bash
+
+ cd ~/
+ mkdir opt
+ git clone https://github.com/Itseez/opencv.git
+ cd opencv
+ git checkout 2.4
+ mkdir build
+ cd build
+ cmake -DBUILD_SHARED_LIBS=OFF ..
+ ...
+ ...
+ make -j8
+ # optional
+ # make install
+
+Install Leiningen
+=================
+
+Once you installed OpenCV with desktop java support the only other
+requirement is to install
+`Leiningeng <https://github.com/technomancy/leiningen>`_ which allows
+you to manage the entire life cycle of your CLJ projects.
+
+The available `installation guide <https://github.com/technomancy/leiningen#installation>`_ is very easy to be followed:
+
+1. `Download the script <https://raw.github.com/technomancy/leiningen/stable/bin/lein>`_
+2. Place it on your ``$PATH`` (cf. ``~/bin`` is a good choice if it is
+ on your ``path``.)
+3. Set the script to be executable. (i.e. ``chmod 755 ~/bin/lein``).
+
+If you work on Windows, follow `this instruction <https://github.com/technomancy/leiningen#windows>`_
+
+You now have both the OpenCV library and a fully installed basic Clojure
+environment. What is now needed is to configure the Clojure environment
+to interact with the OpenCV library.
+
+Install the localrepo Leiningen plugin
+=======================================
+
+The set of commands (tasks in Leiningen parlance) natively supported by
+Leiningen can be very easily extended by various plugins. One of them is
+the `lein-localrepo <https://github.com/kumarshantanu/lein-localrepo>`_
+plugin which allows to install any jar lib as an artifact in the local
+maven repository of your machine (typically in the ``~/.m2/repository``
+directory of your username).
+
+We're going to use this ``lein`` plugin to add to the local maven
+repository the opencv components needed by Java and Clojure to use the
+opencv lib.
+
+Generally speaking, if you want to use a plugin on project base only, it
+can be added directly to a CLJ project created by ``lein``.
+
+Instead, when you want a plugin to be available to any CLJ project in
+your username space, you can add it to the ``profiles.clj`` in the
+``~/.lein/`` directory.
+
+The ``lein-localrepo`` plugin will be useful to me in other CLJ
+projects where I need to call native libs wrapped by a Java interface.
+So I decide to make it available to any CLJ project:
+
+.. code:: bash
+
+ mkdir ~/.lein
+
+Create a file named ``profiles.clj`` in the ``~/.lein`` directory and
+copy into it the following content:
+
+.. code:: clojure
+
+ {:user {:plugins [[lein-localrepo "0.5.2"]]}}
+
+Here we're saying that the version release ``"0.5.2"`` of the
+``lein-localrepo`` plugin will be available to the ``:user`` profile for
+any CLJ project created by ``lein``.
+
+You do not need to do anything else to install the plugin because it
+will be automatically downloaded from a remote repository the very first
+time you issue any ``lein`` task.
+
+Install the java specific libs as local repository
+==================================================
+
+If you followed the standard documentation for installing OpenCV on your
+computer, you should find the following two libs under the directory
+where you built OpenCV:
+
+- the ``build/bin/opencv-247.jar`` java lib
+- the ``build/lib/libopencv_java247.dylib`` native lib (or ``.so`` in
+ you built OpenCV a GNU/Linux OS)
+
+They are the only opencv libs needed by the JVM to interact with OpenCV.
+
+Take apart the needed opencv libs
+---------------------------------
+
+Create a new directory to store in the above two libs. Start by copying
+into it the ``opencv-247.jar`` lib.
+
+.. code:: bash
+
+ cd ~/opt
+ mkdir clj-opencv
+ cd clj-opencv
+ cp ~/opt/opencv/build/bin/opencv-247.jar .
+
+First lib done.
+
+Now, to be able to add the ``libopencv_java247.dylib`` shared native lib
+to the local maven repository, we first need to package it as a jar
+file.
+
+The native lib has to be copied into a directories layout which mimics
+the names of your operating system and architecture. I'm using a Mac OS
+X with a X86 64 bit architecture. So my layout will be the following:
+
+.. code:: bash
+
+ mkdir -p native/macosx/x86_64
+
+Copy into the ``x86_64`` directory the ``libopencv_java247.dylib`` lib.
+
+.. code:: bash
+
+ cp ~/opt/opencv/build/lib/libopencv_java247.dylib native/macosx/x86_64/
+
+If you're running OpenCV from a different OS/Architecture pair, here
+is a summary of the mapping you can choose from.
+
+.. code:: bash
+
+ OS
+
+ Mac OS X -> macosx
+ Windows -> windows
+ Linux -> linux
+ SunOS -> solaris
+
+ Architectures
+
+ amd64 -> x86_64
+ x86_64 -> x86_64
+ x86 -> x86
+ i386 -> x86
+ arm -> arm
+ sparc -> sparc
+
+Package the native lib as a jar
+-------------------------------
+
+Next you need to package the native lib in a jar file by using the
+``jar`` command to create a new jar file from a directory.
+
+.. code:: bash
+
+ jar -cMf opencv-native-247.jar native
+
+Note that ehe ``M`` option instructs the ``jar`` command to not create
+a MANIFEST file for the artifact.
+
+Your directories layout should look like the following:
+
+.. code:: bash
+
+ tree
+ .
+ |__ native
+ | |__ macosx
+ | |__ x86_64
+ | |__ libopencv_java247.dylib
+ |
+ |__ opencv-247.jar
+ |__ opencv-native-247.jar
+
+ 3 directories, 3 files
+
+Locally install the jars
+------------------------
+
+We are now ready to add the two jars as artifacts to the local maven
+repository with the help of the ``lein-localrepo`` plugin.
+
+.. code:: bash
+
+ lein localrepo install opencv-247.jar opencv/opencv 2.4.7
+
+Here the ``localrepo install`` task creates the ``2.4.7.`` release of
+the ``opencv/opencv`` maven artifact from the ``opencv-247.jar`` lib and
+then installs it into the local maven repository. The ``opencv/opencv``
+artifact will then be available to any maven compliant project
+(Leiningen is internally based on maven).
+
+Do the same thing with the native lib previously wrapped in a new jar
+file.
+
+.. code:: bash
+
+ lein localrepo install opencv-native-247.jar opencv/opencv-native 2.4.7
+
+Note that the groupId, ``opencv``, of the two artifacts is the same. We
+are now ready to create a new CLJ project to start interacting with
+OpenCV.
+
+Create a project
+----------------
+
+Create a new CLJ project by using the ``lein new`` task from the
+terminal.
+
+.. code:: bash
+
+ # cd in the directory where you work with your development projects (e.g. ~/devel)
+ lein new simple-sample
+ Generating a project called simple-sample based on the 'default' template.
+ To see other templates (app, lein plugin, etc), try `lein help new`.
+
+The above task creates the following ``simple-sample`` directories
+layout:
+
+.. code:: bash
+
+ tree simple-sample/
+ simple-sample/
+ |__ LICENSE
+ |__ README.md
+ |__ doc
+ | |__ intro.md
+ |
+ |__ project.clj
+ |__ resources
+ |__ src
+ | |__ simple_sample
+ | |__ core.clj
+ |__ test
+ |__ simple_sample
+ |__ core_test.clj
+
+ 6 directories, 6 files
+
+We need to add the two ``opencv`` artifacts as dependencies of the newly
+created project. Open the ``project.clj`` and modify its dependencies
+section as follows:
+
+.. code:: bash
+
+ (defproject simple-sample "0.1.0-SNAPSHOT"
+ :description "FIXME: write description"
+ :url "http://example.com/FIXME"
+ :license {:name "Eclipse Public License"
+ :url "http://www.eclipse.org/legal/epl-v10.html"}
+ :dependencies [[org.clojure/clojure "1.5.1"]
+ [opencv/opencv "2.4.7"] ; added line
+ [opencv/opencv-native "2.4.7"]]) ;added line
+
+
+Note that The Clojure Programming Language is a jar artifact too. This
+is why Clojure is called an hosted language.
+
+To verify that everything went right issue the ``lein deps`` task. The
+very first time you run a ``lein`` task it will take sometime to
+download all the required dependencies before executing the task
+itself.
+
+.. code:: bash
+
+ cd simple-sample
+ lein deps
+ ...
+
+The ``deps`` task reads and merges from the ``project.clj`` and the
+``~/.lein/profiles.clj`` files all the dependencies of the
+``simple-sample`` project and verifies if they have already been
+cached in the local maven repository. If the task returns without
+messages about not being able to retrieve the two new artifacts your
+installation is correct, otherwise go back and double check that you
+did everything right.
+
+REPLing with OpenCV
+-------------------
+
+Now ``cd`` in the ``simple-sample`` directory and issue the following
+``lein`` task:
+
+.. code:: bash
+
+ cd simple-sample
+ lein repl
+ ...
+ ...
+ nREPL server started on port 50907 on host 127.0.0.1
+ REPL-y 0.3.0
+ Clojure 1.5.1
+ Docs: (doc function-name-here)
+ (find-doc "part-of-name-here")
+ Source: (source function-name-here)
+ Javadoc: (javadoc java-object-or-class-here)
+ Exit: Control+D or (exit) or (quit)
+ Results: Stored in vars *1, *2, *3, an exception in *e
+
+ user=>
+
+You can immediately interact with the REPL by issuing any CLJ expression
+to be evaluated.
+
+.. code:: clojure
+
+ user=> (+ 41 1)
+ 42
+ user=> (println "Hello, OpenCV!")
+ Hello, OpenCV!
+ nil
+ user=> (defn foo [] (str "bar"))
+ #'user/foo
+ user=> (foo)
+ "bar"
+
+When ran from the home directory of a lein based project, even if the
+``lein repl`` task automatically loads all the project dependencies, you
+still need to load the opencv native library to be able to interact with
+the OpenCV.
+
+.. code:: clojure
+
+ user=> (clojure.lang.RT/loadLibrary org.opencv.core.Core/NATIVE_LIBRARY_NAME)
+ nil
+
+Then you can start interacting with OpenCV by just referencing the fully
+qualified names of its classes.
+
+ NOTE 2: `Here <http://docs.opencv.org/java/>`_ you can find the
+ full OpenCV Java API.
+
+.. code:: clojure
+
+ user=> (org.opencv.core.Point. 0 0)
+ #<Point {0.0, 0.0}>
+
+Here we created a two dimensions opencv ``Point`` instance. Even if all
+the java packages included within the java interface to OpenCV are
+immediately available from the CLJ REPL, it's very annoying to prefix
+the ``Point.`` instance constructors with the fully qualified package
+name.
+
+Fortunately CLJ offer a very easy way to overcome this annoyance by
+directly importing the ``Point`` class.
+
+.. code:: clojure
+
+ user=> (import 'org.opencv.core.Point)
+ org.opencv.core.Point
+ user=> (def p1 (Point. 0 0))
+ #'user/p1
+ user=> p1
+ #<Point {0.0, 0.0}>
+ user=> (def p2 (Point. 100 100))
+ #'user/p2
+
+We can even inspect the class of an instance and verify if the value of
+a symbol is an instance of a ``Point`` java class.
+
+.. code:: clojure
+
+ user=> (class p1)
+ org.opencv.core.Point
+ user=> (instance? org.opencv.core.Point p1)
+ true
+
+If we now want to use the opencv ``Rect`` class to create a rectangle,
+we again have to fully qualify its constructor even if it leaves in
+the same ``org.opencv.core`` package of the ``Point`` class.
+
+.. code:: clojure
+
+ user=> (org.opencv.core.Rect. p1 p2)
+ #<Rect {0, 0, 100x100}>
+
+Again, the CLJ importing facilities is very handy and let you to map
+more symbols in one shot.
+
+.. code:: clojure
+
+ user=> (import '[org.opencv.core Point Rect Size])
+ org.opencv.core.Size
+ user=> (def r1 (Rect. p1 p2))
+ #'user/r1
+ user=> r1
+ #<Rect {0, 0, 100x100}>
+ user=> (class r1)
+ org.opencv.core.Rect
+ user=> (instance? org.opencv.core.Rect r1)
+ true
+ user=> (Size. 100 100)
+ #<Size 100x100>
+ user=> (def sq-100 (Size. 100 100))
+ #'user/sq-100
+ user=> (class sq-100)
+ org.opencv.core.Size
+ user=> (instance? org.opencv.core.Size sq-100)
+ true
+
+Obviously you can call methods on instances as well.
+
+.. code:: clojure
+
+ user=> (.area r1)
+ 10000.0
+ user=> (.area sq-100)
+ 10000.0
+
+Or modify the value of a member field.
+
+.. code:: clojure
+
+ user=> (set! (.x p1) 10)
+ 10
+ user=> p1
+ #<Point {10.0, 0.0}>
+ user=> (set! (.width sq-100) 10)
+ 10
+ user=> (set! (.height sq-100) 10)
+ 10
+ user=> (.area sq-100)
+ 100.0
+
+If you find yourself not remembering a OpenCV class behavior, the
+REPL gives you the opportunity to easily search the corresponding
+javadoc documention:
+
+.. code:: clojure
+
+ user=> (javadoc Rect)
+ "http://www.google.com/search?btnI=I%27m%20Feeling%20Lucky&q=allinurl:org/opencv/core/Rect.html"
+
+Mimic the OpenCV Java Tutorial Sample in the REPL
+-------------------------------------------------
+
+Let's now try to port to Clojure the `opencv java tutorial sample <http://docs.opencv.org/2.4.4-beta/doc/tutorials/introduction/desktop_java/java_dev_intro.html>`_.
+Instead of writing it in a source file we're going to evaluate it at the
+REPL.
+
+Following is the original Java source code of the cited sample.
+
+.. code:: java
+
+ import org.opencv.core.Mat;
+ import org.opencv.core.CvType;
+ import org.opencv.core.Scalar;
+
+ class SimpleSample {
+
+ static{ System.loadLibrary("opencv_java244"); }
+
+ public static void main(String[] args) {
+ Mat m = new Mat(5, 10, CvType.CV_8UC1, new Scalar(0));
+ System.out.println("OpenCV Mat: " + m);
+ Mat mr1 = m.row(1);
+ mr1.setTo(new Scalar(1));
+ Mat mc5 = m.col(5);
+ mc5.setTo(new Scalar(5));
+ System.out.println("OpenCV Mat data:\n" + m.dump());
+ }
+
+ }
+
+Add injections to the project
+-----------------------------
+
+Before start coding, we'd like to eliminate the boring need of
+interactively loading the native opencv lib any time we start a new REPL
+to interact with it.
+
+First, stop the REPL by evaluating the ``(exit)`` expression at the REPL
+prompt.
+
+.. code:: clojure
+
+ user=> (exit)
+ Bye for now!
+
+Then open your ``project.clj`` file and edit it as follows:
+
+.. code:: clojure
+
+ (defproject simple-sample "0.1.0-SNAPSHOT"
+ ...
+ :injections [(clojure.lang.RT/loadLibrary org.opencv.core.Core/NATIVE_LIBRARY_NAME)])
+
+Here we're saying to load the opencv native lib anytime we run the REPL
+in such a way that we have not anymore to remember to manually do it.
+
+Rerun the ``lein repl`` task
+
+.. code:: bash
+
+ lein repl
+ nREPL server started on port 51645 on host 127.0.0.1
+ REPL-y 0.3.0
+ Clojure 1.5.1
+ Docs: (doc function-name-here)
+ (find-doc "part-of-name-here")
+ Source: (source function-name-here)
+ Javadoc: (javadoc java-object-or-class-here)
+ Exit: Control+D or (exit) or (quit)
+ Results: Stored in vars *1, *2, *3, an exception in *e
+
+ user=>
+
+Import the interested OpenCV java interfaces.
+
+.. code:: clojure
+
+ user=> (import '[org.opencv.core Mat CvType Scalar])
+ org.opencv.core.Scalar
+
+We're going to mimic almost verbatim the original OpenCV java tutorial
+to:
+
+- create a 5x10 matrix with all its elements intialized to 0
+- change the value of every element of the second row to 1
+- change the value of every element of the 6th column to 5
+- print the content of the obtained matrix
+
+.. code:: clojure
+
+ user=> (def m (Mat. 5 10 CvType/CV_8UC1 (Scalar. 0 0)))
+ #'user/m
+ user=> (def mr1 (.row m 1))
+ #'user/mr1
+ user=> (.setTo mr1 (Scalar. 1 0))
+ #<Mat Mat [ 1*10*CV_8UC1, isCont=true, isSubmat=true, nativeObj=0x7fc9dac49880, dataAddr=0x7fc9d9c98d5a ]>
+ user=> (def mc5 (.col m 5))
+ #'user/mc5
+ user=> (.setTo mc5 (Scalar. 5 0))
+ #<Mat Mat [ 5*1*CV_8UC1, isCont=false, isSubmat=true, nativeObj=0x7fc9d9c995a0, dataAddr=0x7fc9d9c98d55 ]>
+ user=> (println (.dump m))
+ [0, 0, 0, 0, 0, 5, 0, 0, 0, 0;
+ 1, 1, 1, 1, 1, 5, 1, 1, 1, 1;
+ 0, 0, 0, 0, 0, 5, 0, 0, 0, 0;
+ 0, 0, 0, 0, 0, 5, 0, 0, 0, 0;
+ 0, 0, 0, 0, 0, 5, 0, 0, 0, 0]
+ nil
+
+If you are accustomed to a functional language all those abused and
+mutating nouns are going to irritate your preference for verbs. Even
+if the CLJ interop syntax is very handy and complete, there is still
+an impedance mismatch between any OOP language and any FP language
+(bein Scala a mixed paradigms programming language).
+
+To exit the REPL type ``(exit)``, ``ctr-D`` or ``(quit)`` at the REPL
+prompt.
+
+.. code:: clojure
+
+ user=> (exit)
+ Bye for now!
+
+Interactively load and blur an image
+------------------------------------
+
+In the next sample you will learn how to interactively load and blur and
+image from the REPL by using the following OpenCV methods:
+
+- the ``imread`` static method from the ``Highgui`` class to read an
+ image from a file
+- the ``imwrite`` static method from the ``Highgui`` class to write an
+ image to a file
+- the ``GaussianBlur`` static method from the ``Imgproc`` class to
+ apply to blur the original image
+
+We're also going to use the ``Mat`` class which is returned from the
+``imread`` method and accpeted as the main argument to both the
+``GaussianBlur`` and the ``imwrite`` methods.
+
+Add an image to the project
+---------------------------
+
+First we want to add an image file to a newly create directory for
+storing static resources of the project.
+
+.. image:: images/lena.png
+ :alt: Original Image
+ :align: center
+
+.. code:: bash
+
+ mkdir -p resources/images
+ cp ~/opt/opencv/doc/tutorials/introduction/desktop_java/images/lena.png resource/images/
+
+Read the image
+--------------
+
+Now launch the REPL as usual and start by importing all the OpenCV
+classes we're going to use:
+
+.. code:: clojure
+
+ lein repl
+ nREPL server started on port 50624 on host 127.0.0.1
+ REPL-y 0.3.0
+ Clojure 1.5.1
+ Docs: (doc function-name-here)
+ (find-doc "part-of-name-here")
+ Source: (source function-name-here)
+ Javadoc: (javadoc java-object-or-class-here)
+ Exit: Control+D or (exit) or (quit)
+ Results: Stored in vars *1, *2, *3, an exception in *e
+
+ user=> (import '[org.opencv.core Mat Size CvType]
+ '[org.opencv.highgui Highgui]
+ '[org.opencv.imgproc Imgproc])
+ org.opencv.imgproc.Imgproc
+
+Now read the image from the ``resources/images/lena.png`` file.
+
+.. code:: clojure
+
+ user=> (def lena (Highgui/imread "resources/images/lena.png"))
+ #'user/lena
+ user=> lena
+ #<Mat Mat [ 512*512*CV_8UC3, isCont=true, isSubmat=false, nativeObj=0x7f9ab3054c40, dataAddr=0x19fea9010 ]>
+
+As you see, by simply evaluating the ``lena`` symbol we know that
+``lena.png`` is a ``512x512`` matrix of ``CV_8UC3`` elements type. Let's
+create a new ``Mat`` instance of the same dimensions and elements type.
+
+.. code:: clojure
+
+ user=> (def blurred (Mat. 512 512 CvType/CV_8UC3))
+ #'user/blurred
+ user=>
+
+Now apply a ``GaussianBlur`` filter using ``lena`` as the source matrix
+and ``blurred`` as the destination matrix.
+
+.. code:: clojure
+
+ user=> (Imgproc/GaussianBlur lena blurred (Size. 5 5) 3 3)
+ nil
+
+As a last step just save the ``blurred`` matrix in a new image file.
+
+.. code:: clojure
+
+ user=> (Highgui/imwrite "resources/images/blurred.png" blurred)
+ true
+ user=> (exit)
+ Bye for now!
+
+Following is the new blurred image of Lena.
+
+.. image:: images/blurred.png
+ :alt: Blurred Image
+ :align: center
+
+Next Steps
+==========
+
+This tutorial only introduces the very basic environment set up to be
+able to interact with OpenCV in a CLJ REPL.
+
+I recommend any Clojure newbie to read the `Clojure Java Interop chapter <http://clojure.org/java_interop>`_ to get all you need to know
+to interoperate with any plain java lib that has not been wrapped in
+Clojure to make it usable in a more idiomatic and functional way within
+Clojure.
+
+The OpenCV Java API does not wrap the ``highgui`` module
+functionalities depending on ``Qt`` (e.g. ``namedWindow`` and
+``imshow``. If you want to create windows and show images into them
+while interacting with OpenCV from the REPL, at the moment you're left
+at your own. You could use Java Swing to fill the gap.
+
+
+License
+-------
+
+Copyright © 2013 Giacomo (Mimmo) Cosenza aka Magomimmo
+
+Distributed under the BSD 3-clause License, the same of OpenCV.
:height: 90pt
:width: 90pt
+ ================ =================================================
+ |ClojureLogo| **Title:** :ref:`clojure_dev_intro`
+
+ *Compatibility:* > OpenCV 2.4.4
+
+ *Author:* |Author_MimmoC|
+
+ A tutorial on how to interactively use OpenCV from the Clojure REPL.
+
+ ================ =================================================
+
+ .. |ClojureLogo| image:: images/clojure-logo.png
+ :height: 90pt
+ :width: 90pt
+
* **Android**
.. tabularcolumns:: m{100pt} m{300pt}
../windows_visual_studio_image_watch/windows_visual_studio_image_watch
../desktop_java/java_dev_intro
../java_eclipse/java_eclipse
+ ../clojure_dev_intro/clojure_dev_intro
../android_binary_package/android_dev_intro
../android_binary_package/O4A_SDK
../android_binary_package/dev_with_OCV_on_Android
#include "opencv2/core/operations.hpp"
#include "opencv2/core/cvstd.inl.hpp"
-
#endif /*__OPENCV_CORE_HPP__*/
CV_EXPORTS void error(int _code, const String& _err, const char* _func, const char* _file, int _line);
#ifdef __GNUC__
-# if defined __clang__ || defined __APPLE__
-# pragma GCC diagnostic push
-# pragma GCC diagnostic ignored "-Winvalid-noreturn"
-# endif
+# if defined __clang__ || defined __APPLE__
+# pragma GCC diagnostic push
+# pragma GCC diagnostic ignored "-Winvalid-noreturn"
+# endif
#endif
CV_INLINE CV_NORETURN void errorNoReturn(int _code, const String& _err, const char* _func, const char* _file, int _line)
{
error(_code, _err, _func, _file, _line);
+#ifdef __GNUC__
+# if !defined __clang__ && !defined __APPLE__
+ // this suppresses this warning: "noreturn" function does return [enabled by default]
+ __builtin_trap();
+ // or use infinite loop: for (;;) {}
+# endif
+#endif
}
#ifdef __GNUC__
-# if defined __clang__ || defined __APPLE__
-# pragma GCC diagnostic pop
-# endif
+# if defined __clang__ || defined __APPLE__
+# pragma GCC diagnostic pop
+# endif
#endif
--- /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-2013, Advanced Micro Devices, Inc., 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*/
+
+#ifndef __OPENCV_CORE_DIRECTX_HPP__
+#define __OPENCV_CORE_DIRECTX_HPP__
+
+#include "mat.hpp"
+#include "ocl.hpp"
+
+#if !defined(__d3d11_h__)
+struct ID3D11Device;
+struct ID3D11Texture2D;
+#endif
+
+#if !defined(__d3d10_h__)
+struct ID3D10Device;
+struct ID3D10Texture2D;
+#endif
+
+#if !defined(_D3D9_H_)
+struct IDirect3DDevice9;
+struct IDirect3DDevice9Ex;
+struct IDirect3DSurface9;
+#endif
+
+namespace cv { namespace directx {
+
+namespace ocl {
+using namespace cv::ocl;
+
+// TODO static functions in the Context class
+CV_EXPORTS Context2& initializeContextFromD3D11Device(ID3D11Device* pD3D11Device);
+CV_EXPORTS Context2& initializeContextFromD3D10Device(ID3D10Device* pD3D10Device);
+CV_EXPORTS Context2& initializeContextFromDirect3DDevice9Ex(IDirect3DDevice9Ex* pDirect3DDevice9Ex);
+CV_EXPORTS Context2& initializeContextFromDirect3DDevice9(IDirect3DDevice9* pDirect3DDevice9);
+
+} // namespace cv::directx::ocl
+
+CV_EXPORTS void convertToD3D11Texture2D(InputArray src, ID3D11Texture2D* pD3D11Texture2D);
+CV_EXPORTS void convertFromD3D11Texture2D(ID3D11Texture2D* pD3D11Texture2D, OutputArray dst);
+
+CV_EXPORTS void convertToD3D10Texture2D(InputArray src, ID3D10Texture2D* pD3D10Texture2D);
+CV_EXPORTS void convertFromD3D10Texture2D(ID3D10Texture2D* pD3D10Texture2D, OutputArray dst);
+
+CV_EXPORTS void convertToDirect3DSurface9(InputArray src, IDirect3DSurface9* pDirect3DSurface9, void* surfaceSharedHandle = NULL);
+CV_EXPORTS void convertFromDirect3DSurface9(IDirect3DSurface9* pDirect3DSurface9, OutputArray dst, void* surfaceSharedHandle = NULL);
+
+// Get OpenCV type from DirectX type, return -1 if there is no equivalent
+CV_EXPORTS int getTypeFromDXGI_FORMAT(const int iDXGI_FORMAT); // enum DXGI_FORMAT for D3D10/D3D11
+
+// Get OpenCV type from DirectX type, return -1 if there is no equivalent
+CV_EXPORTS int getTypeFromD3DFORMAT(const int iD3DFORMAT); // enum D3DTYPE for D3D9
+
+
+} } // namespace cv::directx
+
+#endif // __OPENCV_CORE_DIRECTX_HPP__
{
if( u && CV_XADD(&u->refcount, -1) == 1 )
deallocate();
+ u = NULL;
data = datastart = dataend = datalimit = 0;
size.p[0] = 0;
- u = 0;
}
inline
CV_EXPORTS bool haveOpenCL();
CV_EXPORTS bool useOpenCL();
+CV_EXPORTS bool haveAmdBlas();
CV_EXPORTS void setUseOpenCL(bool flag);
CV_EXPORTS void finish2();
Program getProg(const ProgramSource2& prog,
const String& buildopt, String& errmsg);
- static Context2& getDefault();
+ static Context2& getDefault(bool initialize = true);
void* ptr() const;
+
+ struct Impl;
+ inline struct Impl* _getImpl() const { return p; };
protected:
+ Impl* p;
+};
+
+
+// TODO Move to internal header
+void initializeContextFromHandle(Context2& ctx, void* platform, void* context, void* device);
+
+class CV_EXPORTS Platform
+{
+public:
+ Platform();
+ ~Platform();
+ Platform(const Platform& p);
+ Platform& operator = (const Platform& p);
+
+ void* ptr() const;
+ static Platform& getDefault();
+
struct Impl;
+ inline struct Impl* _getImpl() const { return p; };
+protected:
Impl* p;
};
AutoLock& operator = (const AutoLock&);
};
+class CV_EXPORTS TLSDataContainer
+{
+private:
+ int key_;
+protected:
+ TLSDataContainer();
+ virtual ~TLSDataContainer();
+public:
+ virtual void* createDataInstance() const = 0;
+ virtual void deleteDataInstance(void* data) const = 0;
+
+ void* getData() const;
+};
+
+template <typename T>
+class TLSData : protected TLSDataContainer
+{
+public:
+ inline TLSData() {}
+ inline ~TLSData() {}
+ inline T* get() const { return (T*)getData(); }
+private:
+ virtual void* createDataInstance() const { return new T; }
+ virtual void deleteDataInstance(void* data) const { delete (T*)data; }
+};
+
// The CommandLineParser class is designed for command line arguments parsing
class CV_EXPORTS CommandLineParser
--- /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-2013, Advanced Micro Devices, Inc., 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*/
+
+#include "precomp.hpp"
+
+#include "opencv2/core.hpp"
+#include "opencv2/core/ocl.hpp"
+#include "opencv2/core/directx.hpp"
+
+#ifdef HAVE_DIRECTX
+#include <vector>
+# include "directx.inc.hpp"
+#else // HAVE_DIRECTX
+#define NO_DIRECTX_SUPPORT_ERROR CV_ErrorNoReturn(cv::Error::StsBadFunc, "OpenCV was build without DirectX support")
+#endif
+
+#ifndef HAVE_OPENCL
+# define NO_OPENCL_SUPPORT_ERROR CV_ErrorNoReturn(cv::Error::StsBadFunc, "OpenCV was build without OpenCL support")
+#endif // HAVE_OPENCL
+
+namespace cv { namespace directx {
+
+int getTypeFromDXGI_FORMAT(const int iDXGI_FORMAT)
+{
+ (void)iDXGI_FORMAT;
+#if !defined(HAVE_DIRECTX)
+ NO_DIRECTX_SUPPORT_ERROR;
+#else
+ const int errorType = -1;
+ switch ((enum DXGI_FORMAT)iDXGI_FORMAT)
+ {
+ //case DXGI_FORMAT_UNKNOWN:
+ //case DXGI_FORMAT_R32G32B32A32_TYPELESS:
+ case DXGI_FORMAT_R32G32B32A32_FLOAT: return CV_32FC4;
+ case DXGI_FORMAT_R32G32B32A32_UINT:
+ case DXGI_FORMAT_R32G32B32A32_SINT: return CV_32SC4;
+ //case DXGI_FORMAT_R32G32B32_TYPELESS:
+ case DXGI_FORMAT_R32G32B32_FLOAT: return CV_32FC3;
+ case DXGI_FORMAT_R32G32B32_UINT:
+ case DXGI_FORMAT_R32G32B32_SINT: return CV_32SC3;
+ //case DXGI_FORMAT_R16G16B16A16_TYPELESS:
+ //case DXGI_FORMAT_R16G16B16A16_FLOAT:
+ case DXGI_FORMAT_R16G16B16A16_UNORM:
+ case DXGI_FORMAT_R16G16B16A16_UINT: return CV_16UC4;
+ case DXGI_FORMAT_R16G16B16A16_SNORM:
+ case DXGI_FORMAT_R16G16B16A16_SINT: return CV_16SC4;
+ //case DXGI_FORMAT_R32G32_TYPELESS:
+ //case DXGI_FORMAT_R32G32_FLOAT:
+ //case DXGI_FORMAT_R32G32_UINT:
+ //case DXGI_FORMAT_R32G32_SINT:
+ //case DXGI_FORMAT_R32G8X24_TYPELESS:
+ //case DXGI_FORMAT_D32_FLOAT_S8X24_UINT:
+ //case DXGI_FORMAT_R32_FLOAT_X8X24_TYPELESS:
+ //case DXGI_FORMAT_X32_TYPELESS_G8X24_UINT:
+ //case DXGI_FORMAT_R10G10B10A2_TYPELESS:
+ //case DXGI_FORMAT_R10G10B10A2_UNORM:
+ //case DXGI_FORMAT_R10G10B10A2_UINT:
+ //case DXGI_FORMAT_R11G11B10_FLOAT:
+ //case DXGI_FORMAT_R8G8B8A8_TYPELESS:
+ case DXGI_FORMAT_R8G8B8A8_UNORM:
+ case DXGI_FORMAT_R8G8B8A8_UNORM_SRGB:
+ case DXGI_FORMAT_R8G8B8A8_UINT: return CV_8UC4;
+ case DXGI_FORMAT_R8G8B8A8_SNORM:
+ case DXGI_FORMAT_R8G8B8A8_SINT: return CV_8SC4;
+ //case DXGI_FORMAT_R16G16_TYPELESS:
+ //case DXGI_FORMAT_R16G16_FLOAT:
+ case DXGI_FORMAT_R16G16_UNORM:
+ case DXGI_FORMAT_R16G16_UINT: return CV_16UC2;
+ case DXGI_FORMAT_R16G16_SNORM:
+ case DXGI_FORMAT_R16G16_SINT: return CV_16SC2;
+ //case DXGI_FORMAT_R32_TYPELESS:
+ //case DXGI_FORMAT_D32_FLOAT:
+ case DXGI_FORMAT_R32_FLOAT: return CV_32FC1;
+ case DXGI_FORMAT_R32_UINT:
+ case DXGI_FORMAT_R32_SINT: return CV_32SC1;
+ //case DXGI_FORMAT_R24G8_TYPELESS:
+ //case DXGI_FORMAT_D24_UNORM_S8_UINT:
+ //case DXGI_FORMAT_R24_UNORM_X8_TYPELESS:
+ //case DXGI_FORMAT_X24_TYPELESS_G8_UINT:
+ //case DXGI_FORMAT_R8G8_TYPELESS:
+ case DXGI_FORMAT_R8G8_UNORM:
+ case DXGI_FORMAT_R8G8_UINT: return CV_8UC2;
+ case DXGI_FORMAT_R8G8_SNORM:
+ case DXGI_FORMAT_R8G8_SINT: return CV_8SC2;
+ //case DXGI_FORMAT_R16_TYPELESS:
+ //case DXGI_FORMAT_R16_FLOAT:
+ case DXGI_FORMAT_D16_UNORM:
+ case DXGI_FORMAT_R16_UNORM:
+ case DXGI_FORMAT_R16_UINT: return CV_16UC1;
+ case DXGI_FORMAT_R16_SNORM:
+ case DXGI_FORMAT_R16_SINT: return CV_16SC1;
+ //case DXGI_FORMAT_R8_TYPELESS:
+ case DXGI_FORMAT_R8_UNORM:
+ case DXGI_FORMAT_R8_UINT: return CV_8UC1;
+ case DXGI_FORMAT_R8_SNORM:
+ case DXGI_FORMAT_R8_SINT: return CV_8SC1;
+ case DXGI_FORMAT_A8_UNORM: return CV_8UC1;
+ //case DXGI_FORMAT_R1_UNORM:
+ //case DXGI_FORMAT_R9G9B9E5_SHAREDEXP:
+ //case DXGI_FORMAT_R8G8_B8G8_UNORM:
+ //case DXGI_FORMAT_G8R8_G8B8_UNORM:
+ //case DXGI_FORMAT_BC1_TYPELESS:
+ //case DXGI_FORMAT_BC1_UNORM:
+ //case DXGI_FORMAT_BC1_UNORM_SRGB:
+ //case DXGI_FORMAT_BC2_TYPELESS:
+ //case DXGI_FORMAT_BC2_UNORM:
+ //case DXGI_FORMAT_BC2_UNORM_SRGB:
+ //case DXGI_FORMAT_BC3_TYPELESS:
+ //case DXGI_FORMAT_BC3_UNORM:
+ //case DXGI_FORMAT_BC3_UNORM_SRGB:
+ //case DXGI_FORMAT_BC4_TYPELESS:
+ //case DXGI_FORMAT_BC4_UNORM:
+ //case DXGI_FORMAT_BC4_SNORM:
+ //case DXGI_FORMAT_BC5_TYPELESS:
+ //case DXGI_FORMAT_BC5_UNORM:
+ //case DXGI_FORMAT_BC5_SNORM:
+ //case DXGI_FORMAT_B5G6R5_UNORM:
+ //case DXGI_FORMAT_B5G5R5A1_UNORM:
+ case DXGI_FORMAT_B8G8R8A8_UNORM:
+ case DXGI_FORMAT_B8G8R8X8_UNORM: return CV_8UC4;
+ //case DXGI_FORMAT_R10G10B10_XR_BIAS_A2_UNORM:
+ //case DXGI_FORMAT_B8G8R8A8_TYPELESS:
+ case DXGI_FORMAT_B8G8R8A8_UNORM_SRGB: return CV_8UC4;
+ //case DXGI_FORMAT_B8G8R8X8_TYPELESS:
+ case DXGI_FORMAT_B8G8R8X8_UNORM_SRGB: return CV_8UC4;
+ //case DXGI_FORMAT_BC6H_TYPELESS:
+ //case DXGI_FORMAT_BC6H_UF16:
+ //case DXGI_FORMAT_BC6H_SF16:
+ //case DXGI_FORMAT_BC7_TYPELESS:
+ //case DXGI_FORMAT_BC7_UNORM:
+ //case DXGI_FORMAT_BC7_UNORM_SRGB:
+ default: break;
+ }
+ return errorType;
+#endif
+}
+
+int getTypeFromD3DFORMAT(const int iD3DFORMAT)
+{
+ (void)iD3DFORMAT;
+#if !defined(HAVE_DIRECTX)
+ NO_DIRECTX_SUPPORT_ERROR;
+#else
+ const int errorType = -1;
+ switch ((enum _D3DFORMAT)iD3DFORMAT)
+ {
+ //case D3DFMT_UNKNOWN:
+ case D3DFMT_R8G8B8: return CV_8UC3;
+ case D3DFMT_A8R8G8B8:
+ case D3DFMT_X8R8G8B8: return CV_8UC4;
+ //case D3DFMT_R5G6B5:
+ //case D3DFMT_X1R5G5B5:
+ //case D3DFMT_A1R5G5B5:
+ //case D3DFMT_A4R4G4B4:
+ //case D3DFMT_R3G3B2:
+ case D3DFMT_A8: return CV_8UC1;
+ //case D3DFMT_A8R3G3B2:
+ //case D3DFMT_X4R4G4B4:
+ //case D3DFMT_A2B10G10R10:
+ case D3DFMT_A8B8G8R8:
+ case D3DFMT_X8B8G8R8: return CV_8UC4;
+ //case D3DFMT_G16R16:
+ //case D3DFMT_A2R10G10B10:
+ //case D3DFMT_A16B16G16R16:
+
+ case D3DFMT_A8P8: return CV_8UC2;
+ case D3DFMT_P8: return CV_8UC1;
+
+ case D3DFMT_L8: return CV_8UC1;
+ case D3DFMT_A8L8: return CV_8UC2;
+ //case D3DFMT_A4L4:
+
+ case D3DFMT_V8U8: return CV_8UC2;
+ //case D3DFMT_L6V5U5:
+ case D3DFMT_X8L8V8U8:
+ case D3DFMT_Q8W8V8U8: return CV_8UC4;
+ case D3DFMT_V16U16: return CV_16UC4; // TODO 16SC4 ?
+ //case D3DFMT_A2W10V10U10:
+
+ case D3DFMT_D16_LOCKABLE: return CV_16UC1;
+ case D3DFMT_D32: return CV_32SC1;
+ //case D3DFMT_D15S1:
+ //case D3DFMT_D24S8:
+ //case D3DFMT_D24X8:
+ //case D3DFMT_D24X4S4:
+ case D3DFMT_D16: return CV_16UC1;
+
+ case D3DFMT_D32F_LOCKABLE: return CV_32FC1;
+ default: break;
+ }
+ return errorType;
+#endif
+}
+
+namespace ocl {
+
+#if defined(HAVE_DIRECTX) && defined(HAVE_OPENCL)
+static bool g_isDirect3DDevice9Ex = false; // Direct3DDevice9Ex or Direct3DDevice9 was used
+#endif
+
+Context2& initializeContextFromD3D11Device(ID3D11Device* pD3D11Device)
+{
+ (void)pD3D11Device;
+#if !defined(HAVE_DIRECTX)
+ NO_DIRECTX_SUPPORT_ERROR;
+#elif !defined(HAVE_OPENCL)
+ NO_OPENCL_SUPPORT_ERROR;
+#else
+ cl_uint numPlatforms;
+ cl_int status = clGetPlatformIDs(0, NULL, &numPlatforms);
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLInitError, "OpenCL: Can't get number of platforms");
+ if (numPlatforms == 0)
+ CV_Error(cv::Error::OpenCLInitError, "OpenCL: No available platforms");
+
+ std::vector<cl_platform_id> platforms(numPlatforms);
+ status = clGetPlatformIDs(numPlatforms, &platforms[0], NULL);
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLInitError, "OpenCL: Can't get number of platforms");
+
+ // TODO Filter platforms by name from OPENCV_OPENCL_DEVICE
+
+ int found = -1;
+ cl_device_id device = NULL;
+ cl_uint numDevices = 0;
+ cl_context context = NULL;
+
+ // try with CL_PREFERRED_DEVICES_FOR_D3D11_KHR
+ for (int i = 0; i < (int)numPlatforms; i++)
+ {
+ clGetDeviceIDsFromD3D11KHR_fn clGetDeviceIDsFromD3D11KHR = (clGetDeviceIDsFromD3D11KHR_fn)
+ clGetExtensionFunctionAddressForPlatform(platforms[i], "clGetDeviceIDsFromD3D11KHR");
+ if (!clGetDeviceIDsFromD3D11KHR)
+ continue;
+
+ device = NULL;
+ numDevices = 0;
+ status = clGetDeviceIDsFromD3D11KHR(platforms[i], CL_D3D11_DEVICE_KHR, pD3D11Device,
+ CL_PREFERRED_DEVICES_FOR_D3D11_KHR, 1, &device, &numDevices);
+ if (status != CL_SUCCESS)
+ continue;
+ if (numDevices > 0)
+ {
+ cl_context_properties properties[] = {
+ CL_CONTEXT_PLATFORM, (cl_context_properties)platforms[i],
+ CL_CONTEXT_D3D11_DEVICE_KHR, (cl_context_properties)(pD3D11Device),
+ CL_CONTEXT_INTEROP_USER_SYNC, CL_FALSE,
+ NULL, NULL
+ };
+ context = clCreateContext(properties, 1, &device, NULL, NULL, &status);
+ if (status != CL_SUCCESS)
+ {
+ clReleaseDevice(device);
+ }
+ else
+ {
+ found = i;
+ break;
+ }
+ }
+ }
+ if (found < 0)
+ {
+ // try with CL_ALL_DEVICES_FOR_D3D11_KHR
+ for (int i = 0; i < (int)numPlatforms; i++)
+ {
+ clGetDeviceIDsFromD3D11KHR_fn clGetDeviceIDsFromD3D11KHR = (clGetDeviceIDsFromD3D11KHR_fn)
+ clGetExtensionFunctionAddressForPlatform(platforms[i], "clGetDeviceIDsFromD3D11KHR");
+ if (!clGetDeviceIDsFromD3D11KHR)
+ continue;
+
+ device = NULL;
+ numDevices = 0;
+ status = clGetDeviceIDsFromD3D11KHR(platforms[i], CL_D3D11_DEVICE_KHR, pD3D11Device,
+ CL_ALL_DEVICES_FOR_D3D11_KHR, 1, &device, &numDevices);
+ if (status != CL_SUCCESS)
+ continue;
+ if (numDevices > 0)
+ {
+ cl_context_properties properties[] = {
+ CL_CONTEXT_PLATFORM, (cl_context_properties)platforms[i],
+ CL_CONTEXT_D3D11_DEVICE_KHR, (cl_context_properties)(pD3D11Device),
+ CL_CONTEXT_INTEROP_USER_SYNC, CL_FALSE,
+ NULL, NULL
+ };
+ context = clCreateContext(properties, 1, &device, NULL, NULL, &status);
+ if (status != CL_SUCCESS)
+ {
+ clReleaseDevice(device);
+ }
+ else
+ {
+ found = i;
+ break;
+ }
+ }
+ }
+ if (found < 0)
+ CV_Error(cv::Error::OpenCLInitError, "OpenCL: Can't create context for DirectX interop");
+ }
+
+
+ Context2& ctx = Context2::getDefault(false);
+ initializeContextFromHandle(ctx, platforms[found], context, device);
+ return ctx;
+#endif
+}
+
+Context2& initializeContextFromD3D10Device(ID3D10Device* pD3D10Device)
+{
+ (void)pD3D10Device;
+#if !defined(HAVE_DIRECTX)
+ NO_DIRECTX_SUPPORT_ERROR;
+#elif !defined(HAVE_OPENCL)
+ NO_OPENCL_SUPPORT_ERROR;
+#else
+ cl_uint numPlatforms;
+ cl_int status = clGetPlatformIDs(0, NULL, &numPlatforms);
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLInitError, "OpenCL: Can't get number of platforms");
+ if (numPlatforms == 0)
+ CV_Error(cv::Error::OpenCLInitError, "OpenCL: No available platforms");
+
+ std::vector<cl_platform_id> platforms(numPlatforms);
+ status = clGetPlatformIDs(numPlatforms, &platforms[0], NULL);
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLInitError, "OpenCL: Can't get number of platforms");
+
+ // TODO Filter platforms by name from OPENCV_OPENCL_DEVICE
+
+ int found = -1;
+ cl_device_id device = NULL;
+ cl_uint numDevices = 0;
+ cl_context context = NULL;
+
+ // try with CL_PREFERRED_DEVICES_FOR_D3D10_KHR
+ for (int i = 0; i < (int)numPlatforms; i++)
+ {
+ clGetDeviceIDsFromD3D10KHR_fn clGetDeviceIDsFromD3D10KHR = (clGetDeviceIDsFromD3D10KHR_fn)
+ clGetExtensionFunctionAddressForPlatform(platforms[i], "clGetDeviceIDsFromD3D10KHR");
+ if (!clGetDeviceIDsFromD3D10KHR)
+ continue;
+
+ device = NULL;
+ numDevices = 0;
+ status = clGetDeviceIDsFromD3D10KHR(platforms[i], CL_D3D10_DEVICE_KHR, pD3D10Device,
+ CL_PREFERRED_DEVICES_FOR_D3D10_KHR, 1, &device, &numDevices);
+ if (status != CL_SUCCESS)
+ continue;
+ if (numDevices > 0)
+ {
+ cl_context_properties properties[] = {
+ CL_CONTEXT_PLATFORM, (cl_context_properties)platforms[i],
+ CL_CONTEXT_D3D10_DEVICE_KHR, (cl_context_properties)(pD3D10Device),
+ CL_CONTEXT_INTEROP_USER_SYNC, CL_FALSE,
+ NULL, NULL
+ };
+ context = clCreateContext(properties, 1, &device, NULL, NULL, &status);
+ if (status != CL_SUCCESS)
+ {
+ clReleaseDevice(device);
+ }
+ else
+ {
+ found = i;
+ break;
+ }
+ }
+ }
+ if (found < 0)
+ {
+ // try with CL_ALL_DEVICES_FOR_D3D10_KHR
+ for (int i = 0; i < (int)numPlatforms; i++)
+ {
+ clGetDeviceIDsFromD3D10KHR_fn clGetDeviceIDsFromD3D10KHR = (clGetDeviceIDsFromD3D10KHR_fn)
+ clGetExtensionFunctionAddressForPlatform(platforms[i], "clGetDeviceIDsFromD3D10KHR");
+ if (!clGetDeviceIDsFromD3D10KHR)
+ continue;
+
+ device = NULL;
+ numDevices = 0;
+ status = clGetDeviceIDsFromD3D10KHR(platforms[i], CL_D3D10_DEVICE_KHR, pD3D10Device,
+ CL_ALL_DEVICES_FOR_D3D10_KHR, 1, &device, &numDevices);
+ if (status != CL_SUCCESS)
+ continue;
+ if (numDevices > 0)
+ {
+ cl_context_properties properties[] = {
+ CL_CONTEXT_PLATFORM, (cl_context_properties)platforms[i],
+ CL_CONTEXT_D3D10_DEVICE_KHR, (cl_context_properties)(pD3D10Device),
+ CL_CONTEXT_INTEROP_USER_SYNC, CL_FALSE,
+ NULL, NULL
+ };
+ context = clCreateContext(properties, 1, &device, NULL, NULL, &status);
+ if (status != CL_SUCCESS)
+ {
+ clReleaseDevice(device);
+ }
+ else
+ {
+ found = i;
+ break;
+ }
+ }
+ }
+ if (found < 0)
+ CV_Error(cv::Error::OpenCLInitError, "OpenCL: Can't create context for DirectX interop");
+ }
+
+
+ Context2& ctx = Context2::getDefault(false);
+ initializeContextFromHandle(ctx, platforms[found], context, device);
+ return ctx;
+#endif
+}
+
+Context2& initializeContextFromDirect3DDevice9Ex(IDirect3DDevice9Ex* pDirect3DDevice9Ex)
+{
+ (void)pDirect3DDevice9Ex;
+#if !defined(HAVE_DIRECTX)
+ NO_DIRECTX_SUPPORT_ERROR;
+#elif !defined(HAVE_OPENCL)
+ NO_OPENCL_SUPPORT_ERROR;
+#else
+ cl_uint numPlatforms;
+ cl_int status = clGetPlatformIDs(0, NULL, &numPlatforms);
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLInitError, "OpenCL: Can't get number of platforms");
+ if (numPlatforms == 0)
+ CV_Error(cv::Error::OpenCLInitError, "OpenCL: No available platforms");
+
+ std::vector<cl_platform_id> platforms(numPlatforms);
+ status = clGetPlatformIDs(numPlatforms, &platforms[0], NULL);
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLInitError, "OpenCL: Can't get number of platforms");
+
+ // TODO Filter platforms by name from OPENCV_OPENCL_DEVICE
+
+ int found = -1;
+ cl_device_id device = NULL;
+ cl_uint numDevices = 0;
+ cl_context context = NULL;
+
+ // try with CL_PREFERRED_DEVICES_FOR_DX9_MEDIA_ADAPTER_KHR
+ for (int i = 0; i < (int)numPlatforms; i++)
+ {
+ clGetDeviceIDsFromDX9MediaAdapterKHR_fn clGetDeviceIDsFromDX9MediaAdapterKHR = (clGetDeviceIDsFromDX9MediaAdapterKHR_fn)
+ clGetExtensionFunctionAddressForPlatform(platforms[i], "clGetDeviceIDsFromDX9MediaAdapterKHR");
+ if (!clGetDeviceIDsFromDX9MediaAdapterKHR)
+ continue;
+
+ device = NULL;
+ numDevices = 0;
+ cl_dx9_media_adapter_type_khr type = CL_ADAPTER_D3D9EX_KHR;
+ status = clGetDeviceIDsFromDX9MediaAdapterKHR(platforms[i], 1, &type, &pDirect3DDevice9Ex,
+ CL_PREFERRED_DEVICES_FOR_DX9_MEDIA_ADAPTER_KHR, 1, &device, &numDevices);
+ if (status != CL_SUCCESS)
+ continue;
+ if (numDevices > 0)
+ {
+ cl_context_properties properties[] = {
+ CL_CONTEXT_PLATFORM, (cl_context_properties)platforms[i],
+ CL_CONTEXT_ADAPTER_D3D9EX_KHR, (cl_context_properties)(pDirect3DDevice9Ex),
+ CL_CONTEXT_INTEROP_USER_SYNC, CL_FALSE,
+ NULL, NULL
+ };
+ context = clCreateContext(properties, 1, &device, NULL, NULL, &status);
+ if (status != CL_SUCCESS)
+ {
+ clReleaseDevice(device);
+ }
+ else
+ {
+ found = i;
+ break;
+ }
+ }
+ }
+ if (found < 0)
+ {
+ // try with CL_ALL_DEVICES_FOR_DX9_MEDIA_ADAPTER_KHR
+ for (int i = 0; i < (int)numPlatforms; i++)
+ {
+ clGetDeviceIDsFromDX9MediaAdapterKHR_fn clGetDeviceIDsFromDX9MediaAdapterKHR = (clGetDeviceIDsFromDX9MediaAdapterKHR_fn)
+ clGetExtensionFunctionAddressForPlatform(platforms[i], "clGetDeviceIDsFromDX9MediaAdapterKHR");
+ if (!clGetDeviceIDsFromDX9MediaAdapterKHR)
+ continue;
+
+ device = NULL;
+ numDevices = 0;
+ cl_dx9_media_adapter_type_khr type = CL_ADAPTER_D3D9EX_KHR;
+ status = clGetDeviceIDsFromDX9MediaAdapterKHR(platforms[i], 1, &type, &pDirect3DDevice9Ex,
+ CL_ALL_DEVICES_FOR_DX9_MEDIA_ADAPTER_KHR, 1, &device, &numDevices);
+ if (status != CL_SUCCESS)
+ continue;
+ if (numDevices > 0)
+ {
+ cl_context_properties properties[] = {
+ CL_CONTEXT_PLATFORM, (cl_context_properties)platforms[i],
+ CL_CONTEXT_ADAPTER_D3D9EX_KHR, (cl_context_properties)(pDirect3DDevice9Ex),
+ CL_CONTEXT_INTEROP_USER_SYNC, CL_FALSE,
+ NULL, NULL
+ };
+ context = clCreateContext(properties, 1, &device, NULL, NULL, &status);
+ if (status != CL_SUCCESS)
+ {
+ clReleaseDevice(device);
+ }
+ else
+ {
+ found = i;
+ break;
+ }
+ }
+ }
+ if (found < 0)
+ CV_Error(cv::Error::OpenCLInitError, "OpenCL: Can't create context for DirectX interop");
+ }
+
+ Context2& ctx = Context2::getDefault(false);
+ initializeContextFromHandle(ctx, platforms[found], context, device);
+ g_isDirect3DDevice9Ex = true;
+ return ctx;
+#endif
+}
+
+Context2& initializeContextFromDirect3DDevice9(IDirect3DDevice9* pDirect3DDevice9)
+{
+ (void)pDirect3DDevice9;
+#if !defined(HAVE_DIRECTX)
+ NO_DIRECTX_SUPPORT_ERROR;
+#elif !defined(HAVE_OPENCL)
+ NO_OPENCL_SUPPORT_ERROR;
+#else
+ cl_uint numPlatforms;
+ cl_int status = clGetPlatformIDs(0, NULL, &numPlatforms);
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLInitError, "OpenCL: Can't get number of platforms");
+ if (numPlatforms == 0)
+ CV_Error(cv::Error::OpenCLInitError, "OpenCL: No available platforms");
+
+ std::vector<cl_platform_id> platforms(numPlatforms);
+ status = clGetPlatformIDs(numPlatforms, &platforms[0], NULL);
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLInitError, "OpenCL: Can't get number of platforms");
+
+ // TODO Filter platforms by name from OPENCV_OPENCL_DEVICE
+
+ int found = -1;
+ cl_device_id device = NULL;
+ cl_uint numDevices = 0;
+ cl_context context = NULL;
+
+ // try with CL_PREFERRED_DEVICES_FOR_DX9_MEDIA_ADAPTER_KHR
+ for (int i = 0; i < (int)numPlatforms; i++)
+ {
+ clGetDeviceIDsFromDX9MediaAdapterKHR_fn clGetDeviceIDsFromDX9MediaAdapterKHR = (clGetDeviceIDsFromDX9MediaAdapterKHR_fn)
+ clGetExtensionFunctionAddressForPlatform(platforms[i], "clGetDeviceIDsFromDX9MediaAdapterKHR");
+ if (!clGetDeviceIDsFromDX9MediaAdapterKHR)
+ continue;
+
+ device = NULL;
+ numDevices = 0;
+ cl_dx9_media_adapter_type_khr type = CL_ADAPTER_D3D9_KHR;
+ status = clGetDeviceIDsFromDX9MediaAdapterKHR(platforms[i], 1, &type, &pDirect3DDevice9,
+ CL_PREFERRED_DEVICES_FOR_DX9_MEDIA_ADAPTER_KHR, 1, &device, &numDevices);
+ if (status != CL_SUCCESS)
+ continue;
+ if (numDevices > 0)
+ {
+ cl_context_properties properties[] = {
+ CL_CONTEXT_PLATFORM, (cl_context_properties)platforms[i],
+ CL_CONTEXT_ADAPTER_D3D9_KHR, (cl_context_properties)(pDirect3DDevice9),
+ CL_CONTEXT_INTEROP_USER_SYNC, CL_FALSE,
+ NULL, NULL
+ };
+ context = clCreateContext(properties, 1, &device, NULL, NULL, &status);
+ if (status != CL_SUCCESS)
+ {
+ clReleaseDevice(device);
+ }
+ else
+ {
+ found = i;
+ break;
+ }
+ }
+ }
+ if (found < 0)
+ {
+ // try with CL_ALL_DEVICES_FOR_DX9_MEDIA_ADAPTER_KHR
+ for (int i = 0; i < (int)numPlatforms; i++)
+ {
+ clGetDeviceIDsFromDX9MediaAdapterKHR_fn clGetDeviceIDsFromDX9MediaAdapterKHR = (clGetDeviceIDsFromDX9MediaAdapterKHR_fn)
+ clGetExtensionFunctionAddressForPlatform(platforms[i], "clGetDeviceIDsFromDX9MediaAdapterKHR");
+ if (!clGetDeviceIDsFromDX9MediaAdapterKHR)
+ continue;
+
+ device = NULL;
+ numDevices = 0;
+ cl_dx9_media_adapter_type_khr type = CL_ADAPTER_D3D9_KHR;
+ status = clGetDeviceIDsFromDX9MediaAdapterKHR(platforms[i], 1, &type, &pDirect3DDevice9,
+ CL_ALL_DEVICES_FOR_DX9_MEDIA_ADAPTER_KHR, 1, &device, &numDevices);
+ if (status != CL_SUCCESS)
+ continue;
+ if (numDevices > 0)
+ {
+ cl_context_properties properties[] = {
+ CL_CONTEXT_PLATFORM, (cl_context_properties)platforms[i],
+ CL_CONTEXT_ADAPTER_D3D9_KHR, (cl_context_properties)(pDirect3DDevice9),
+ CL_CONTEXT_INTEROP_USER_SYNC, CL_FALSE,
+ NULL, NULL
+ };
+ context = clCreateContext(properties, 1, &device, NULL, NULL, &status);
+ if (status != CL_SUCCESS)
+ {
+ clReleaseDevice(device);
+ }
+ else
+ {
+ found = i;
+ break;
+ }
+ }
+ }
+ if (found < 0)
+ CV_Error(cv::Error::OpenCLInitError, "OpenCL: Can't create context for DirectX interop");
+ }
+
+ Context2& ctx = Context2::getDefault(false);
+ initializeContextFromHandle(ctx, platforms[found], context, device);
+ g_isDirect3DDevice9Ex = false;
+ return ctx;
+#endif
+}
+
+} // namespace cv::ocl
+
+#if defined(HAVE_DIRECTX) && defined(HAVE_OPENCL)
+clCreateFromD3D11Texture2DKHR_fn clCreateFromD3D11Texture2DKHR = NULL;
+clEnqueueAcquireD3D11ObjectsKHR_fn clEnqueueAcquireD3D11ObjectsKHR = NULL;
+clEnqueueReleaseD3D11ObjectsKHR_fn clEnqueueReleaseD3D11ObjectsKHR = NULL;
+
+static void __OpenCLinitializeD3D11()
+{
+ using namespace cv::ocl;
+ static cl_platform_id initializedPlatform = NULL;
+ cl_platform_id platform = (cl_platform_id)Platform::getDefault().ptr();
+ if (initializedPlatform != platform)
+ {
+ clCreateFromD3D11Texture2DKHR = (clCreateFromD3D11Texture2DKHR_fn)
+ clGetExtensionFunctionAddressForPlatform(platform, "clCreateFromD3D11Texture2DKHR");
+ clEnqueueAcquireD3D11ObjectsKHR = (clEnqueueAcquireD3D11ObjectsKHR_fn)
+ clGetExtensionFunctionAddressForPlatform(platform, "clEnqueueAcquireD3D11ObjectsKHR");
+ clEnqueueReleaseD3D11ObjectsKHR = (clEnqueueReleaseD3D11ObjectsKHR_fn)
+ clGetExtensionFunctionAddressForPlatform(platform, "clEnqueueReleaseD3D11ObjectsKHR");
+ initializedPlatform = platform;
+ }
+ if (!clCreateFromD3D11Texture2DKHR || !clEnqueueAcquireD3D11ObjectsKHR || !clEnqueueReleaseD3D11ObjectsKHR)
+ {
+ CV_Error(cv::Error::OpenCLInitError, "OpenCL: Can't find functions for D3D11");
+ }
+}
+#endif // defined(HAVE_DIRECTX) && defined(HAVE_OPENCL)
+
+void convertToD3D11Texture2D(InputArray src, ID3D11Texture2D* pD3D11Texture2D)
+{
+ (void)src; (void)pD3D11Texture2D;
+#if !defined(HAVE_DIRECTX)
+ NO_DIRECTX_SUPPORT_ERROR;
+#elif defined(HAVE_OPENCL)
+ __OpenCLinitializeD3D11();
+
+ D3D11_TEXTURE2D_DESC desc = { 0 };
+ pD3D11Texture2D->GetDesc(&desc);
+
+ int srcType = src.type();
+ int textureType = getTypeFromDXGI_FORMAT(desc.Format);
+ CV_Assert(textureType == srcType);
+
+ Size srcSize = src.size();
+ CV_Assert(srcSize.width == (int)desc.Width && srcSize.height == (int)desc.Height);
+
+ using namespace cv::ocl;
+ Context2& ctx = Context2::getDefault();
+ cl_context context = (cl_context)ctx.ptr();
+
+ UMat u = src.getUMat();
+
+ // TODO Add support for roi
+ CV_Assert(u.offset == 0);
+ CV_Assert(u.isContinuous());
+
+ cl_int status = 0;
+ cl_mem clImage = clCreateFromD3D11Texture2DKHR(context, CL_MEM_WRITE_ONLY, pD3D11Texture2D, 0, &status);
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clCreateFromD3D11Texture2DKHR failed");
+
+ cl_mem clBuffer = (cl_mem)u.handle(ACCESS_READ);
+
+ cl_command_queue q = (cl_command_queue)Queue::getDefault().ptr();
+ status = clEnqueueAcquireD3D11ObjectsKHR(q, 1, &clImage, 0, NULL, NULL);
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clEnqueueAcquireD3D11ObjectsKHR failed");
+ size_t offset = 0; // TODO
+ size_t dst_origin[3] = {0, 0, 0};
+ size_t region[3] = {u.cols, u.rows, 1};
+ status = clEnqueueCopyBufferToImage(q, clBuffer, clImage, offset, dst_origin, region, 0, NULL, NULL);
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clEnqueueCopyBufferToImage failed");
+ status = clEnqueueReleaseD3D11ObjectsKHR(q, 1, &clImage, 0, NULL, NULL);
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clEnqueueReleaseD3D11ObjectsKHR failed");
+
+ status = clFinish(q); // TODO Use events
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clFinish failed");
+
+ status = clReleaseMemObject(clImage); // TODO RAII
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clReleaseMem failed");
+#else
+ // TODO memcpy
+ NO_OPENCL_SUPPORT_ERROR;
+#endif
+}
+void convertFromD3D11Texture2D(ID3D11Texture2D* pD3D11Texture2D, OutputArray dst)
+{
+ (void)pD3D11Texture2D; (void)dst;
+#if !defined(HAVE_DIRECTX)
+ NO_DIRECTX_SUPPORT_ERROR;
+#elif defined(HAVE_OPENCL)
+ __OpenCLinitializeD3D11();
+
+ D3D11_TEXTURE2D_DESC desc = { 0 };
+ pD3D11Texture2D->GetDesc(&desc);
+
+ int textureType = getTypeFromDXGI_FORMAT(desc.Format);
+ CV_Assert(textureType >= 0);
+
+ using namespace cv::ocl;
+ Context2& ctx = Context2::getDefault();
+ cl_context context = (cl_context)ctx.ptr();
+
+ // TODO Need to specify ACCESS_WRITE here somehow to prevent useless data copying!
+ dst.create(Size(desc.Width, desc.Height), textureType);
+ UMat u = dst.getUMat();
+
+ // TODO Add support for roi
+ CV_Assert(u.offset == 0);
+ CV_Assert(u.isContinuous());
+
+ cl_int status = 0;
+ cl_mem clImage = clCreateFromD3D11Texture2DKHR(context, CL_MEM_READ_ONLY, pD3D11Texture2D, 0, &status);
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clCreateFromD3D11Texture2DKHR failed");
+
+ cl_mem clBuffer = (cl_mem)u.handle(ACCESS_READ);
+
+ cl_command_queue q = (cl_command_queue)Queue::getDefault().ptr();
+ status = clEnqueueAcquireD3D11ObjectsKHR(q, 1, &clImage, 0, NULL, NULL);
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clEnqueueAcquireD3D11ObjectsKHR failed");
+ size_t offset = 0; // TODO
+ size_t src_origin[3] = {0, 0, 0};
+ size_t region[3] = {u.cols, u.rows, 1};
+ status = clEnqueueCopyImageToBuffer(q, clImage, clBuffer, src_origin, region, offset, 0, NULL, NULL);
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clEnqueueCopyImageToBuffer failed");
+ status = clEnqueueReleaseD3D11ObjectsKHR(q, 1, &clImage, 0, NULL, NULL);
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clEnqueueReleaseD3D11ObjectsKHR failed");
+
+ status = clFinish(q); // TODO Use events
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clFinish failed");
+
+ status = clReleaseMemObject(clImage); // TODO RAII
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clReleaseMem failed");
+#else
+ // TODO memcpy
+ NO_OPENCL_SUPPORT_ERROR;
+#endif
+}
+
+#if defined(HAVE_DIRECTX) && defined(HAVE_OPENCL)
+clCreateFromD3D10Texture2DKHR_fn clCreateFromD3D10Texture2DKHR = NULL;
+clEnqueueAcquireD3D10ObjectsKHR_fn clEnqueueAcquireD3D10ObjectsKHR = NULL;
+clEnqueueReleaseD3D10ObjectsKHR_fn clEnqueueReleaseD3D10ObjectsKHR = NULL;
+
+static void __OpenCLinitializeD3D10()
+{
+ using namespace cv::ocl;
+ static cl_platform_id initializedPlatform = NULL;
+ cl_platform_id platform = (cl_platform_id)Platform::getDefault().ptr();
+ if (initializedPlatform != platform)
+ {
+ clCreateFromD3D10Texture2DKHR = (clCreateFromD3D10Texture2DKHR_fn)
+ clGetExtensionFunctionAddressForPlatform(platform, "clCreateFromD3D10Texture2DKHR");
+ clEnqueueAcquireD3D10ObjectsKHR = (clEnqueueAcquireD3D10ObjectsKHR_fn)
+ clGetExtensionFunctionAddressForPlatform(platform, "clEnqueueAcquireD3D10ObjectsKHR");
+ clEnqueueReleaseD3D10ObjectsKHR = (clEnqueueReleaseD3D10ObjectsKHR_fn)
+ clGetExtensionFunctionAddressForPlatform(platform, "clEnqueueReleaseD3D10ObjectsKHR");
+ initializedPlatform = platform;
+ }
+ if (!clCreateFromD3D10Texture2DKHR || !clEnqueueAcquireD3D10ObjectsKHR || !clEnqueueReleaseD3D10ObjectsKHR)
+ {
+ CV_Error(cv::Error::OpenCLInitError, "OpenCL: Can't find functions for D3D10");
+ }
+}
+#endif // defined(HAVE_DIRECTX) && defined(HAVE_OPENCL)
+
+void convertToD3D10Texture2D(InputArray src, ID3D10Texture2D* pD3D10Texture2D)
+{
+ (void)src; (void)pD3D10Texture2D;
+#if !defined(HAVE_DIRECTX)
+ NO_DIRECTX_SUPPORT_ERROR;
+#elif defined(HAVE_OPENCL)
+ __OpenCLinitializeD3D10();
+
+ D3D10_TEXTURE2D_DESC desc = { 0 };
+ pD3D10Texture2D->GetDesc(&desc);
+
+ int srcType = src.type();
+ int textureType = getTypeFromDXGI_FORMAT(desc.Format);
+ CV_Assert(textureType == srcType);
+
+ Size srcSize = src.size();
+ CV_Assert(srcSize.width == (int)desc.Width && srcSize.height == (int)desc.Height);
+
+ using namespace cv::ocl;
+ Context2& ctx = Context2::getDefault();
+ cl_context context = (cl_context)ctx.ptr();
+
+ UMat u = src.getUMat();
+
+ // TODO Add support for roi
+ CV_Assert(u.offset == 0);
+ CV_Assert(u.isContinuous());
+
+ cl_int status = 0;
+ cl_mem clImage = clCreateFromD3D10Texture2DKHR(context, CL_MEM_WRITE_ONLY, pD3D10Texture2D, 0, &status);
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clCreateFromD3D10Texture2DKHR failed");
+
+ cl_mem clBuffer = (cl_mem)u.handle(ACCESS_READ);
+
+ cl_command_queue q = (cl_command_queue)Queue::getDefault().ptr();
+ status = clEnqueueAcquireD3D10ObjectsKHR(q, 1, &clImage, 0, NULL, NULL);
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clEnqueueAcquireD3D10ObjectsKHR failed");
+ size_t offset = 0; // TODO
+ size_t dst_origin[3] = {0, 0, 0};
+ size_t region[3] = {u.cols, u.rows, 1};
+ status = clEnqueueCopyBufferToImage(q, clBuffer, clImage, offset, dst_origin, region, 0, NULL, NULL);
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clEnqueueCopyBufferToImage failed");
+ status = clEnqueueReleaseD3D10ObjectsKHR(q, 1, &clImage, 0, NULL, NULL);
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clEnqueueReleaseD3D10ObjectsKHR failed");
+
+ status = clFinish(q); // TODO Use events
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clFinish failed");
+
+ status = clReleaseMemObject(clImage); // TODO RAII
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clReleaseMem failed");
+#else
+ // TODO memcpy
+ NO_OPENCL_SUPPORT_ERROR;
+#endif
+}
+void convertFromD3D10Texture2D(ID3D10Texture2D* pD3D10Texture2D, OutputArray dst)
+{
+ (void)pD3D10Texture2D; (void)dst;
+#if !defined(HAVE_DIRECTX)
+ NO_DIRECTX_SUPPORT_ERROR;
+#elif defined(HAVE_OPENCL)
+ __OpenCLinitializeD3D10();
+
+ D3D10_TEXTURE2D_DESC desc = { 0 };
+ pD3D10Texture2D->GetDesc(&desc);
+
+ int textureType = getTypeFromDXGI_FORMAT(desc.Format);
+ CV_Assert(textureType >= 0);
+
+ using namespace cv::ocl;
+ Context2& ctx = Context2::getDefault();
+ cl_context context = (cl_context)ctx.ptr();
+
+ // TODO Need to specify ACCESS_WRITE here somehow to prevent useless data copying!
+ dst.create(Size(desc.Width, desc.Height), textureType);
+ UMat u = dst.getUMat();
+
+ // TODO Add support for roi
+ CV_Assert(u.offset == 0);
+ CV_Assert(u.isContinuous());
+
+ cl_int status = 0;
+ cl_mem clImage = clCreateFromD3D10Texture2DKHR(context, CL_MEM_READ_ONLY, pD3D10Texture2D, 0, &status);
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clCreateFromD3D10Texture2DKHR failed");
+
+ cl_mem clBuffer = (cl_mem)u.handle(ACCESS_READ);
+
+ cl_command_queue q = (cl_command_queue)Queue::getDefault().ptr();
+ status = clEnqueueAcquireD3D10ObjectsKHR(q, 1, &clImage, 0, NULL, NULL);
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clEnqueueAcquireD3D10ObjectsKHR failed");
+ size_t offset = 0; // TODO
+ size_t src_origin[3] = {0, 0, 0};
+ size_t region[3] = {u.cols, u.rows, 1};
+ status = clEnqueueCopyImageToBuffer(q, clImage, clBuffer, src_origin, region, offset, 0, NULL, NULL);
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clEnqueueCopyImageToBuffer failed");
+ status = clEnqueueReleaseD3D10ObjectsKHR(q, 1, &clImage, 0, NULL, NULL);
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clEnqueueReleaseD3D10ObjectsKHR failed");
+
+ status = clFinish(q); // TODO Use events
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clFinish failed");
+
+ status = clReleaseMemObject(clImage); // TODO RAII
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clReleaseMem failed");
+#else
+ // TODO memcpy
+ NO_OPENCL_SUPPORT_ERROR;
+#endif
+}
+
+#if defined(HAVE_DIRECTX) && defined(HAVE_OPENCL)
+clCreateFromDX9MediaSurfaceKHR_fn clCreateFromDX9MediaSurfaceKHR = NULL;
+clEnqueueAcquireDX9MediaSurfacesKHR_fn clEnqueueAcquireDX9MediaSurfacesKHR = NULL;
+clEnqueueReleaseDX9MediaSurfacesKHR_fn clEnqueueReleaseDX9MediaSurfacesKHR = NULL;
+
+static void __OpenCLinitializeD3D9()
+{
+ using namespace cv::ocl;
+ static cl_platform_id initializedPlatform = NULL;
+ cl_platform_id platform = (cl_platform_id)Platform::getDefault().ptr();
+ if (initializedPlatform != platform)
+ {
+ clCreateFromDX9MediaSurfaceKHR = (clCreateFromDX9MediaSurfaceKHR_fn)
+ clGetExtensionFunctionAddressForPlatform(platform, "clCreateFromDX9MediaSurfaceKHR");
+ clEnqueueAcquireDX9MediaSurfacesKHR = (clEnqueueAcquireDX9MediaSurfacesKHR_fn)
+ clGetExtensionFunctionAddressForPlatform(platform, "clEnqueueAcquireDX9MediaSurfacesKHR");
+ clEnqueueReleaseDX9MediaSurfacesKHR = (clEnqueueReleaseDX9MediaSurfacesKHR_fn)
+ clGetExtensionFunctionAddressForPlatform(platform, "clEnqueueReleaseDX9MediaSurfacesKHR");
+ initializedPlatform = platform;
+ }
+ if (!clCreateFromDX9MediaSurfaceKHR || !clEnqueueAcquireDX9MediaSurfacesKHR || !clEnqueueReleaseDX9MediaSurfacesKHR)
+ {
+ CV_Error(cv::Error::OpenCLInitError, "OpenCL: Can't find functions for D3D9");
+ }
+}
+#endif // defined(HAVE_DIRECTX) && defined(HAVE_OPENCL)
+
+void convertToDirect3DSurface9(InputArray src, IDirect3DSurface9* pDirect3DSurface9, void* surfaceSharedHandle)
+{
+ (void)src; (void)pDirect3DSurface9; (void)surfaceSharedHandle;
+#if !defined(HAVE_DIRECTX)
+ NO_DIRECTX_SUPPORT_ERROR;
+#elif defined(HAVE_OPENCL)
+ __OpenCLinitializeD3D9();
+
+ D3DSURFACE_DESC desc;
+ if (FAILED(pDirect3DSurface9->GetDesc(&desc)))
+ {
+ CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: Can't get D3D surface description");
+ }
+
+ int srcType = src.type();
+ int surfaceType = getTypeFromD3DFORMAT(desc.Format);
+ CV_Assert(surfaceType == srcType);
+
+ Size srcSize = src.size();
+ CV_Assert(srcSize.width == (int)desc.Width && srcSize.height == (int)desc.Height);
+
+ using namespace cv::ocl;
+ Context2& ctx = Context2::getDefault();
+ cl_context context = (cl_context)ctx.ptr();
+
+ UMat u = src.getUMat();
+
+ // TODO Add support for roi
+ CV_Assert(u.offset == 0);
+ CV_Assert(u.isContinuous());
+
+ cl_int status = 0;
+ cl_dx9_surface_info_khr surfaceInfo = {pDirect3DSurface9, (HANDLE)surfaceSharedHandle};
+ cl_mem clImage = clCreateFromDX9MediaSurfaceKHR(context, CL_MEM_WRITE_ONLY,
+ ocl::g_isDirect3DDevice9Ex ? CL_ADAPTER_D3D9EX_KHR : CL_ADAPTER_D3D9_KHR,
+ &surfaceInfo, 0, &status);
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clCreateFromDX9MediaSurfaceKHR failed");
+
+ cl_mem clBuffer = (cl_mem)u.handle(ACCESS_READ);
+
+ cl_command_queue q = (cl_command_queue)Queue::getDefault().ptr();
+ status = clEnqueueAcquireDX9MediaSurfacesKHR(q, 1, &clImage, 0, NULL, NULL);
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clEnqueueAcquireDX9MediaSurfacesKHR failed");
+ size_t offset = 0; // TODO
+ size_t dst_origin[3] = {0, 0, 0};
+ size_t region[3] = {u.cols, u.rows, 1};
+ status = clEnqueueCopyBufferToImage(q, clBuffer, clImage, offset, dst_origin, region, 0, NULL, NULL);
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clEnqueueCopyBufferToImage failed");
+ status = clEnqueueReleaseDX9MediaSurfacesKHR(q, 1, &clImage, 0, NULL, NULL);
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clEnqueueReleaseDX9MediaSurfacesKHR failed");
+
+ status = clFinish(q); // TODO Use events
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clFinish failed");
+
+ status = clReleaseMemObject(clImage); // TODO RAII
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clReleaseMem failed");
+#else
+ // TODO pDirect3DSurface9->LockRect() + memcpy + Unlock()
+ NO_OPENCL_SUPPORT_ERROR;
+#endif
+}
+
+void convertFromDirect3DSurface9(IDirect3DSurface9* pDirect3DSurface9, OutputArray dst, void* surfaceSharedHandle)
+{
+ (void)pDirect3DSurface9; (void)dst; (void)surfaceSharedHandle;
+#if !defined(HAVE_DIRECTX)
+ NO_DIRECTX_SUPPORT_ERROR;
+#elif defined(HAVE_OPENCL)
+ __OpenCLinitializeD3D9();
+
+ D3DSURFACE_DESC desc;
+ if (FAILED(pDirect3DSurface9->GetDesc(&desc)))
+ {
+ CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: Can't get D3D surface description");
+ }
+
+ int surfaceType = getTypeFromD3DFORMAT(desc.Format);
+ CV_Assert(surfaceType >= 0);
+
+ using namespace cv::ocl;
+ Context2& ctx = Context2::getDefault();
+ cl_context context = (cl_context)ctx.ptr();
+
+ // TODO Need to specify ACCESS_WRITE here somehow to prevent useless data copying!
+ dst.create(Size(desc.Width, desc.Height), surfaceType);
+ UMat u = dst.getUMat();
+
+ // TODO Add support for roi
+ CV_Assert(u.offset == 0);
+ CV_Assert(u.isContinuous());
+
+ cl_int status = 0;
+ cl_dx9_surface_info_khr surfaceInfo = {pDirect3DSurface9, (HANDLE)surfaceSharedHandle};
+ cl_mem clImage = clCreateFromDX9MediaSurfaceKHR(context, CL_MEM_READ_ONLY,
+ ocl::g_isDirect3DDevice9Ex ? CL_ADAPTER_D3D9EX_KHR : CL_ADAPTER_D3D9_KHR,
+ &surfaceInfo, 0, &status);
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clCreateFromDX9MediaSurfaceKHR failed");
+
+ cl_mem clBuffer = (cl_mem)u.handle(ACCESS_WRITE);
+
+ cl_command_queue q = (cl_command_queue)Queue::getDefault().ptr();
+ status = clEnqueueAcquireDX9MediaSurfacesKHR(q, 1, &clImage, 0, NULL, NULL);
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clEnqueueAcquireDX9MediaSurfacesKHR failed");
+ size_t offset = 0; // TODO
+ size_t src_origin[3] = {0, 0, 0};
+ size_t region[3] = {u.cols, u.rows, 1};
+ status = clEnqueueCopyImageToBuffer(q, clImage, clBuffer, src_origin, region, offset, 0, NULL, NULL);
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clEnqueueCopyImageToBuffer failed");
+ status = clEnqueueReleaseDX9MediaSurfacesKHR(q, 1, &clImage, 0, NULL, NULL);
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clEnqueueReleaseDX9MediaSurfacesKHR failed");
+
+ status = clFinish(q); // TODO Use events
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clFinish failed");
+
+ status = clReleaseMemObject(clImage); // TODO RAII
+ if (status != CL_SUCCESS)
+ CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clReleaseMem failed");
+#else
+ // TODO pDirect3DSurface9->LockRect() + memcpy + Unlock()
+ NO_OPENCL_SUPPORT_ERROR;
+#endif
+}
+
+} } // namespace cv::directx
--- /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-2013, Advanced Micro Devices, Inc., 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*/
+
+#if defined(HAVE_DIRECTX)
+#include <d3d11.h>
+#include <d3d10.h>
+#include <d3d9.h>
+
+#ifdef HAVE_OPENCL
+#include "opencv2/core/opencl/runtime/opencl_core.hpp"
+
+#include <CL/cl_d3d11.h>
+#include <CL/cl_d3d10.h>
+#include <CL/cl_dx9_media_sharing.h>
+#endif // HAVE_OPENCL
+
+#endif // HAVE_DIRECTX
//M*/
#include "precomp.hpp"
+#include "opencv2/core/opencl/runtime/opencl_clamdblas.hpp"
#ifdef HAVE_IPP
#include "ippversion.h"
GEMMStore(c_data, c_step, d_buf, d_buf_step, d_data, d_step, d_size, alpha, beta, flags);
}
+#ifdef HAVE_CLAMDBLAS
+
+static bool ocl_gemm( InputArray matA, InputArray matB, double alpha,
+ InputArray matC, double beta, OutputArray matD, int flags )
+{
+ int type = matA.type(), esz = CV_ELEM_SIZE(type);
+ bool haveC = matC.kind() != cv::_InputArray::NONE;
+ Size sizeA = matA.size(), sizeB = matB.size(), sizeC = haveC ? matC.size() : Size(0, 0);
+ bool atrans = (flags & GEMM_1_T) != 0, btrans = (flags & GEMM_2_T) != 0, ctrans = (flags & GEMM_3_T) != 0;
+
+ if (atrans)
+ sizeA = Size(sizeA.height, sizeA.width);
+ if (btrans)
+ sizeB = Size(sizeB.height, sizeB.width);
+ if (haveC && ctrans)
+ sizeC = Size(sizeC.height, sizeC.width);
+
+ Size sizeD(sizeB.width, sizeA.height);
+
+ CV_Assert( matB.type() == type && (!haveC || matC.type() == type) );
+ CV_Assert( sizeA.width == sizeB.height && (!haveC || sizeC == sizeD) );
+
+ matD.create(sizeD, type);
+ if ( matA.offset() % esz != 0 || matA.step() % esz != 0 ||
+ matB.offset() % esz != 0 || matB.step() % esz != 0 ||
+ (haveC && (matC.offset() % esz != 0 || matC.step() % esz != 0)) )
+ return false;
+
+ UMat A = matA.getUMat(), B = matB.getUMat(), D = matD.getUMat();
+ if (haveC)
+ ctrans ? transpose(matC, D) : matC.getMat().copyTo(D); // TODO fix it as soon as .copyTo works as expected
+ else
+ D.setTo(Scalar::all(0));
+
+ int M = sizeD.height, N = sizeD.width, K = sizeA.width;
+ int lda = (int)A.step / esz, ldb = (int)B.step / esz, ldc = (int)D.step / esz;
+ int offa = (int)A.offset / esz, offb = (int)B.offset / esz, offc = (int)D.offset / esz;
+
+ cl_command_queue clq = (cl_command_queue)ocl::Queue::getDefault().ptr();
+ clAmdBlasTranspose transA = atrans ? clAmdBlasTrans : clAmdBlasNoTrans;
+ clAmdBlasTranspose transB = btrans ? clAmdBlasTrans : clAmdBlasNoTrans;
+ clAmdBlasOrder order = clAmdBlasRowMajor;
+ clAmdBlasStatus status = clAmdBlasSuccess;
+
+ if (type == CV_32FC1)
+ status = clAmdBlasSgemmEx(order, transA, transB, M, N, K,
+ (cl_float)alpha, (const cl_mem)A.handle(ACCESS_READ), offa, lda,
+ (const cl_mem)B.handle(ACCESS_READ), offb, ldb,
+ (cl_float)beta, (cl_mem)D.handle(ACCESS_RW), offc, ldc,
+ 1, &clq, 0, NULL, NULL);
+ else if (type == CV_64FC1)
+ status = clAmdBlasDgemmEx(order, transA, transB, M, N, K,
+ alpha, (const cl_mem)A.handle(ACCESS_READ), offa, lda,
+ (const cl_mem)B.handle(ACCESS_READ), offb, ldb,
+ beta, (cl_mem)D.handle(ACCESS_RW), offc, ldc,
+ 1, &clq, 0, NULL, NULL);
+ else if (type == CV_32FC2)
+ {
+ cl_float2 alpha_2 = { { (cl_float)alpha, 0 } };
+ cl_float2 beta_2 = { { (cl_float)beta, 0 } };
+ status = clAmdBlasCgemmEx(order, transA, transB, M, N, K,
+ alpha_2, (const cl_mem)A.handle(ACCESS_READ), offa, lda,
+ (const cl_mem)B.handle(ACCESS_READ), offb, ldb,
+ beta_2, (cl_mem)D.handle(ACCESS_RW), offc, ldc,
+ 1, &clq, 0, NULL, NULL);
+ }
+ else if (type == CV_64FC2)
+ {
+ cl_double2 alpha_2 = { { alpha, 0 } };
+ cl_double2 beta_2 = { { beta, 0 } };
+ status = clAmdBlasZgemmEx(order, transA, transB, M, N, K,
+ alpha_2, (const cl_mem)A.handle(ACCESS_READ), offa, lda,
+ (const cl_mem)B.handle(ACCESS_READ), offb, ldb,
+ beta_2, (cl_mem)D.handle(ACCESS_RW), offc, ldc,
+ 1, &clq, 0, NULL, NULL);
+ }
+ else
+ CV_Error(Error::StsUnsupportedFormat, "");
+
+ return status == clAmdBlasSuccess;
+}
+
+#endif
+
}
void cv::gemm( InputArray matA, InputArray matB, double alpha,
InputArray matC, double beta, OutputArray _matD, int flags )
{
+#ifdef HAVE_CLAMDBLAS
+ if (ocl::haveAmdBlas() && matA.dims() <= 2 && matB.dims() <= 2 && matC.dims() <= 2 &&
+ ocl::useOpenCL() && _matD.isUMat() &&
+ ocl_gemm(matA, matB, alpha, matC, beta, _matD, flags))
+ return;
+#endif
+
const int block_lin_size = 128;
const int block_size = block_lin_size * block_lin_size;
void MatAllocator::unmap(UMatData* u) const
{
if(u->urefcount == 0 && u->refcount == 0)
+ {
deallocate(u);
+ u = NULL;
+ }
}
void MatAllocator::download(UMatData* u, void* dstptr,
{
CV_Assert( sz[i] <= (size_t)INT_MAX );
if( sz[i] == 0 )
- return;
+ return;
if( srcofs )
- srcptr += srcofs[i]*(i <= dims-2 ? srcstep[i] : 1);
+ srcptr += srcofs[i]*(i <= dims-2 ? srcstep[i] : 1);
if( dstofs )
- dstptr += dstofs[i]*(i <= dims-2 ? dststep[i] : 1);
+ dstptr += dstofs[i]*(i <= dims-2 ? dststep[i] : 1);
isz[i] = (int)sz[i];
}
UMatData* u = new UMatData(this);
u->data = u->origdata = data;
u->size = total;
- u->refcount = data0 == 0;
if(data0)
u->flags |= UMatData::USER_ALLOCATED;
void deallocate(UMatData* u) const
{
+ CV_Assert(u->urefcount >= 0);
+ CV_Assert(u->refcount >= 0);
if(u && u->refcount == 0)
{
if( !(u->flags & UMatData::USER_ALLOCATED) )
CV_Assert( step[dims-1] == (size_t)CV_ELEM_SIZE(flags) );
}
+ addref();
finalizeHdr(*this);
}
{
if(u)
(u->currAllocator ? u->currAllocator : allocator ? allocator : getStdAllocator())->unmap(u);
+ u = NULL;
}
Mat::Mat(const Mat& m, const Range& _rowRange, const Range& _colRange)
#include "precomp.hpp"
#include <map>
+#include "opencv2/core/opencl/runtime/opencl_clamdblas.hpp"
+
#ifdef HAVE_OPENCL
#include "opencv2/core/opencl/runtime/opencl_core.hpp"
#else
return h1.a < h2.a || (h1.a == h2.a && h1.b < h2.b);
}
-static bool g_isInitialized = false;
+static bool g_isOpenCLInitialized = false;
static bool g_isOpenCLAvailable = false;
+
bool haveOpenCL()
{
- if (!g_isInitialized)
+ if (!g_isOpenCLInitialized)
{
- if (!g_isInitialized)
+ try
{
- try
- {
- cl_uint n = 0;
- cl_int err = ::clGetPlatformIDs(0, NULL, &n);
- if (err != CL_SUCCESS)
- g_isOpenCLAvailable = false;
- else
- g_isOpenCLAvailable = true;
- }
- catch (...)
- {
- g_isOpenCLAvailable = false;
- }
- g_isInitialized = true;
+ cl_uint n = 0;
+ g_isOpenCLAvailable = ::clGetPlatformIDs(0, NULL, &n) == CL_SUCCESS;
}
+ catch (...)
+ {
+ g_isOpenCLAvailable = false;
+ }
+ g_isOpenCLInitialized = true;
}
return g_isOpenCLAvailable;
}
bool useOpenCL()
{
- TLSData* data = TLSData::get();
+ CoreTLSData* data = coreTlsData.get();
if( data->useOpenCL < 0 )
data->useOpenCL = (int)haveOpenCL();
return data->useOpenCL > 0;
{
if( haveOpenCL() )
{
- TLSData* data = TLSData::get();
+ CoreTLSData* data = coreTlsData.get();
data->useOpenCL = flag ? 1 : 0;
}
}
+#ifdef HAVE_CLAMDBLAS
+
+class AmdBlasHelper
+{
+public:
+ static AmdBlasHelper & getInstance()
+ {
+ static AmdBlasHelper amdBlas;
+ return amdBlas;
+ }
+
+ bool isAvailable() const
+ {
+ return g_isAmdBlasAvailable;
+ }
+
+ ~AmdBlasHelper()
+ {
+ try
+ {
+ clAmdBlasTeardown();
+ }
+ catch (...) { }
+ }
+
+protected:
+ AmdBlasHelper()
+ {
+ if (!g_isAmdBlasInitialized)
+ {
+ AutoLock lock(m);
+
+ if (!g_isAmdBlasInitialized && haveOpenCL())
+ {
+ try
+ {
+ g_isAmdBlasAvailable = clAmdBlasSetup() == clAmdBlasSuccess;
+ }
+ catch (...)
+ {
+ g_isAmdBlasAvailable = false;
+ }
+ }
+ else
+ g_isAmdBlasAvailable = false;
+
+ g_isAmdBlasInitialized = true;
+ }
+ }
+
+private:
+ static Mutex m;
+ static bool g_isAmdBlasInitialized;
+ static bool g_isAmdBlasAvailable;
+};
+
+bool AmdBlasHelper::g_isAmdBlasAvailable = false;
+bool AmdBlasHelper::g_isAmdBlasInitialized = false;
+Mutex AmdBlasHelper::m;
+
+bool haveAmdBlas()
+{
+ return AmdBlasHelper::getInstance().isAvailable();
+}
+
+#else
+
+bool haveAmdBlas()
+{
+ return false;
+}
+
+#endif
+
void finish2()
{
Queue::getDefault().finish();
void release() { if( CV_XADD(&refcount, -1) == 1 ) delete this; } \
int refcount
-class Platform
-{
-public:
- Platform();
- ~Platform();
- Platform(const Platform& p);
- Platform& operator = (const Platform& p);
-
- void* ptr() const;
- static Platform& getDefault();
-protected:
- struct Impl;
- Impl* p;
-};
-
struct Platform::Impl
{
Impl()
Impl(void* d)
{
handle = (cl_device_id)d;
+ refcount = 1;
}
template<typename _TpCL, typename _TpOut>
const Device& Device::getDefault()
{
const Context2& ctx = Context2::getDefault();
- int idx = TLSData::get()->device;
+ int idx = coreTlsData.get()->device;
return ctx.device(idx);
}
struct Context2::Impl
{
+ Impl()
+ {
+ refcount = 1;
+ handle = 0;
+ }
+
Impl(int dtype0)
{
refcount = 1;
cl_context handle;
std::vector<Device> devices;
- bool initialized;
typedef ProgramSource2::hash_t hash_t;
return !p || idx >= p->devices.size() ? dummy : p->devices[idx];
}
-Context2& Context2::getDefault()
+Context2& Context2::getDefault(bool initialize)
{
static Context2 ctx;
- if( !ctx.p && haveOpenCL() )
+ if(!ctx.p && haveOpenCL())
{
- // do not create new Context2 right away.
- // 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);
- if(!ctx.p)
- ctx.create(Device::TYPE_DGPU);
- if(!ctx.p)
- ctx.create(Device::TYPE_IGPU);
- if(!ctx.p)
- ctx.create(Device::TYPE_CPU);
+ if (initialize)
+ {
+ // do not create new Context2 right away.
+ // 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);
+ if(!ctx.p)
+ ctx.create(Device::TYPE_DGPU);
+ if(!ctx.p)
+ ctx.create(Device::TYPE_IGPU);
+ if(!ctx.p)
+ ctx.create(Device::TYPE_CPU);
+ }
+ else
+ {
+ ctx.p = new Impl();
+ }
}
return ctx;
return p ? p->getProg(prog, buildopts, errmsg) : Program();
}
+void initializeContextFromHandle(Context2& ctx, void* platform, void* _context, void* _device)
+{
+ cl_context context = (cl_context)_context;
+ cl_device_id device = (cl_device_id)_device;
+
+ // cleanup old context
+ Context2::Impl* impl = ctx._getImpl();
+ if (impl->handle)
+ {
+ cl_int status = clReleaseContext(impl->handle);
+ (void)status;
+ }
+ impl->devices.clear();
+
+ impl->handle = context;
+ impl->devices.resize(1);
+ impl->devices[0].set(device);
+
+ Platform& p = Platform::getDefault();
+ Platform::Impl* pImpl = p._getImpl();
+ pImpl->handle = (cl_platform_id)platform;
+}
+
+
struct Queue::Impl
{
Impl(const Context2& c, const Device& d)
Queue& Queue::getDefault()
{
- Queue& q = TLSData::get()->oclQueue;
+ Queue& q = coreTlsData.get()->oclQueue;
if( !q.p && haveOpenCL() )
q.create(Context2::getDefault());
return q;
UMatData* defaultAllocate(int dims, const int* sizes, int type, void* data, size_t* step, int flags) const
{
UMatData* u = matStdAllocator->allocate(dims, sizes, type, data, step, flags);
- u->urefcount = 1;
- u->refcount = 0;
return u;
}
u->data = 0;
u->size = total;
u->handle = handle;
- u->urefcount = 1;
u->flags = flags0;
return u;
}
if(accessFlags & ACCESS_WRITE)
u->markHostCopyObsolete(true);
- CV_XADD(&u->urefcount, 1);
return true;
}
if(!u)
return;
+ CV_Assert(u->urefcount >= 0);
+ CV_Assert(u->refcount >= 0);
+
// TODO: !!! when we add Shared Virtual Memory Support,
// this function (as well as the others) should be corrected
CV_Assert(u->handle != 0 && u->urefcount == 0);
--- /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 DOUBLE_SUPPORT
+#ifdef cl_amd_fp64
+#pragma OPENCL EXTENSION cl_amd_fp64:enable
+#elif defined (cl_khr_fp64)
+#pragma OPENCL EXTENSION cl_khr_fp64:enable
+#endif
+#endif
+
+#define noconvert
+
+__kernel void convertTo(__global const uchar * srcptr, int src_step, int src_offset,
+ __global uchar * dstptr, int dst_step, int dst_offset, int dst_rows, int dst_cols,
+ float alpha, float beta )
+{
+ int x = get_global_id(0);
+ int y = get_global_id(1);
+
+ if (x < dst_cols && y < dst_rows)
+ {
+ int src_index = mad24(y, src_step, src_offset + x * (int)sizeof(srcT) );
+ int dst_index = mad24(y, dst_step, dst_offset + x * (int)sizeof(dstT) );
+
+ __global const srcT * src = (__global const srcT *)(srcptr + src_index);
+ __global dstT * dst = (__global dstT *)(dstptr + dst_index);
+
+ dst[0] = convertToDT( src[0] * alpha + beta );
+ }
+}
#endif
#define noconvert
+#define EXTRA_PARAMS
#if defined OP_SUM || defined OP_SUM_ABS || defined OP_SUM_SQR
#if OP_SUM
#elif OP_SUM_SQR
#define FUNC(a, b) a += b * b
#endif
+#define DECLARE_LOCAL_MEM \
+ __local dstT localmem[WGS2_ALIGNED]
#define DEFINE_ACCUMULATOR \
dstT accumulator = (dstT)(0)
#define REDUCE_GLOBAL \
dstT temp = convertToDT(src[0]); \
FUNC(accumulator, temp)
+#define SET_LOCAL_1 \
+ localmem[lid] = accumulator
#define REDUCE_LOCAL_1 \
localmem[lid - WGS2_ALIGNED] += accumulator
#define REDUCE_LOCAL_2 \
localmem[lid] += localmem[lid2]
+#define CALC_RESULT \
+ __global dstT * dst = (__global dstT *)(dstptr + (int)sizeof(dstT) * gid); \
+ dst[0] = localmem[0]
#elif defined OP_COUNT_NON_ZERO
#define dstT int
+#define DECLARE_LOCAL_MEM \
+ __local dstT localmem[WGS2_ALIGNED]
#define DEFINE_ACCUMULATOR \
dstT accumulator = (dstT)(0); \
srcT zero = (srcT)(0), one = (srcT)(1)
#define REDUCE_GLOBAL \
accumulator += src[0] == zero ? zero : one
+#define SET_LOCAL_1 \
+ localmem[lid] = accumulator
#define REDUCE_LOCAL_1 \
localmem[lid - WGS2_ALIGNED] += accumulator
#define REDUCE_LOCAL_2 \
localmem[lid] += localmem[lid2]
+#define CALC_RESULT \
+ __global dstT * dst = (__global dstT *)(dstptr + (int)sizeof(dstT) * gid); \
+ dst[0] = localmem[0]
+
+#elif defined OP_MIN_MAX_LOC || defined OP_MIN_MAX_LOC_MASK
+
+#if defined (DEPTH_0)
+#define srcT uchar
+#define MIN_VAL 0
+#define MAX_VAL 255
+#endif
+#if defined (DEPTH_1)
+#define srcT char
+#define MIN_VAL -128
+#define MAX_VAL 127
+#endif
+#if defined (DEPTH_2)
+#define srcT ushort
+#define MIN_VAL 0
+#define MAX_VAL 65535
+#endif
+#if defined (DEPTH_3)
+#define srcT short
+#define MIN_VAL -32768
+#define MAX_VAL 32767
+#endif
+#if defined (DEPTH_4)
+#define srcT int
+#define MIN_VAL INT_MIN
+#define MAX_VAL INT_MAX
+#endif
+#if defined (DEPTH_5)
+#define srcT float
+#define MIN_VAL (-FLT_MAX)
+#define MAX_VAL FLT_MAX
+#endif
+#if defined (DEPTH_6)
+#define srcT double
+#define MIN_VAL (-DBL_MAX)
+#define MAX_VAL DBL_MAX
+#endif
+
+#define DECLARE_LOCAL_MEM \
+ __local srcT localmem_min[WGS2_ALIGNED]; \
+ __local srcT localmem_max[WGS2_ALIGNED]; \
+ __local int localmem_minloc[WGS2_ALIGNED]; \
+ __local int localmem_maxloc[WGS2_ALIGNED]
+#define DEFINE_ACCUMULATOR \
+ srcT minval = MAX_VAL; \
+ srcT maxval = MIN_VAL; \
+ int negative = -1; \
+ int minloc = negative; \
+ int maxloc = negative; \
+ srcT temp; \
+ int temploc
+#define REDUCE_GLOBAL \
+ temp = src[0]; \
+ temploc = id; \
+ srcT temp_minval = minval, temp_maxval = maxval; \
+ minval = min(minval, temp); \
+ maxval = max(maxval, temp); \
+ minloc = (minval == temp_minval) ? (temp_minval == MAX_VAL) ? temploc : minloc : temploc; \
+ maxloc = (maxval == temp_maxval) ? (temp_maxval == MIN_VAL) ? temploc : maxloc : temploc
+#define SET_LOCAL_1 \
+ localmem_min[lid] = minval; \
+ localmem_max[lid] = maxval; \
+ localmem_minloc[lid] = minloc; \
+ localmem_maxloc[lid] = maxloc
+#define REDUCE_LOCAL_1 \
+ srcT oldmin = localmem_min[lid-WGS2_ALIGNED]; \
+ srcT oldmax = localmem_max[lid-WGS2_ALIGNED]; \
+ localmem_min[lid - WGS2_ALIGNED] = min(minval,localmem_min[lid-WGS2_ALIGNED]); \
+ localmem_max[lid - WGS2_ALIGNED] = max(maxval,localmem_max[lid-WGS2_ALIGNED]); \
+ srcT minv = localmem_min[lid - WGS2_ALIGNED], maxv = localmem_max[lid - WGS2_ALIGNED]; \
+ localmem_minloc[lid - WGS2_ALIGNED] = (minv == minval) ? (minv == oldmin) ? \
+ min(minloc, localmem_minloc[lid-WGS2_ALIGNED]) : minloc : localmem_minloc[lid-WGS2_ALIGNED]; \
+ localmem_maxloc[lid - WGS2_ALIGNED] = (maxv == maxval) ? (maxv == oldmax) ? \
+ min(maxloc, localmem_maxloc[lid-WGS2_ALIGNED]) : maxloc : localmem_maxloc[lid-WGS2_ALIGNED]
+#define REDUCE_LOCAL_2 \
+ srcT oldmin = localmem_min[lid]; \
+ srcT oldmax = localmem_max[lid]; \
+ localmem_min[lid] = min(localmem_min[lid], localmem_min[lid2]); \
+ localmem_max[lid] = max(localmem_max[lid], localmem_max[lid2]); \
+ srcT min1 = localmem_min[lid], min2 = localmem_min[lid2]; \
+ localmem_minloc[lid] = (localmem_minloc[lid] == negative) ? localmem_minloc[lid2] : (localmem_minloc[lid2] == negative) ? \
+ localmem_minloc[lid] : (min1 == min2) ? (min1 == oldmin) ? min(localmem_minloc[lid2],localmem_minloc[lid]) : \
+ localmem_minloc[lid2] : localmem_minloc[lid]; \
+ srcT max1 = localmem_max[lid], max2 = localmem_max[lid2]; \
+ localmem_maxloc[lid] = (localmem_maxloc[lid] == negative) ? localmem_maxloc[lid2] : (localmem_maxloc[lid2] == negative) ? \
+ localmem_maxloc[lid] : (max1 == max2) ? (max1 == oldmax) ? min(localmem_maxloc[lid2],localmem_maxloc[lid]) : \
+ localmem_maxloc[lid2] : localmem_maxloc[lid]
+#define CALC_RESULT \
+ __global srcT * dstminval = (__global srcT *)(dstptr + (int)sizeof(srcT) * gid); \
+ __global srcT * dstmaxval = (__global srcT *)(dstptr2 + (int)sizeof(srcT) * gid); \
+ dstminval[0] = localmem_min[0]; \
+ dstmaxval[0] = localmem_max[0]; \
+ dstlocptr[gid] = localmem_minloc[0]; \
+ dstlocptr2[gid] = localmem_maxloc[0]
+
+#if defined OP_MIN_MAX_LOC_MASK
+#undef DEFINE_ACCUMULATOR
+#define DEFINE_ACCUMULATOR \
+ srcT minval = MAX_VAL; \
+ srcT maxval = MIN_VAL; \
+ int negative = -1; \
+ int minloc = negative; \
+ int maxloc = negative; \
+ srcT temp, temp_mask, zeroVal = (srcT)(0); \
+ int temploc
+#undef REDUCE_GLOBAL
+#define REDUCE_GLOBAL \
+ temp = src[0]; \
+ temploc = id; \
+ int mask_index = mad24(id / cols, mask_step, mask_offset + (id % cols) * (int)sizeof(uchar)); \
+ __global const uchar * mask = (__global const uchar *)(maskptr + mask_index); \
+ temp_mask = mask[0]; \
+ srcT temp_minval = minval, temp_maxval = maxval; \
+ minval = (temp_mask == zeroVal) ? minval : min(minval, temp); \
+ maxval = (temp_mask == zeroVal) ? maxval : max(maxval, temp); \
+ minloc = (temp_mask == zeroVal) ? minloc : (minval == temp_minval) ? (temp_minval == MAX_VAL) ? temploc : minloc : temploc; \
+ maxloc = (temp_mask == zeroVal) ? maxloc : (maxval == temp_maxval) ? (temp_maxval == MIN_VAL) ? temploc : maxloc : temploc
+#endif
#else
#error "No operation"
+#endif
+#if defined OP_MIN_MAX_LOC
+#undef EXTRA_PARAMS
+#define EXTRA_PARAMS , __global uchar * dstptr2, __global int * dstlocptr, __global int * dstlocptr2
+#endif
+#if defined OP_MIN_MAX_LOC_MASK
+#undef EXTRA_PARAMS
+#define EXTRA_PARAMS , __global uchar * dstptr2, __global int * dstlocptr, __global int * dstlocptr2, \
+ __global const uchar * maskptr, int mask_step, int mask_offset, __global int * test
#endif
__kernel void reduce(__global const uchar * srcptr, int step, int offset, int cols,
- int total, int groupnum, __global uchar * dstptr)
+ int total, int groupnum, __global uchar * dstptr EXTRA_PARAMS)
{
int lid = get_local_id(0);
int gid = get_group_id(0);
int id = get_global_id(0);
- __local dstT localmem[WGS2_ALIGNED];
+ DECLARE_LOCAL_MEM;
DEFINE_ACCUMULATOR;
for (int grain = groupnum * WGS; id < total; id += grain)
}
if (lid < WGS2_ALIGNED)
- localmem[lid] = accumulator;
+ {
+ SET_LOCAL_1;
+ }
barrier(CLK_LOCAL_MEM_FENCE);
- if (lid >= WGS2_ALIGNED)
+ if (lid >= WGS2_ALIGNED && total >= WGS2_ALIGNED)
+ {
REDUCE_LOCAL_1;
+ }
barrier(CLK_LOCAL_MEM_FENCE);
for (int lsize = WGS2_ALIGNED >> 1; lsize > 0; lsize >>= 1)
if (lid == 0)
{
- __global dstT * dst = (__global dstT *)(dstptr + (int)sizeof(dstT) * gid);
- dst[0] = localmem[0];
+ CALC_RESULT;
}
}
#ifndef __OPENCV_PRECOMP_H__
#define __OPENCV_PRECOMP_H__
+#include "opencv2/opencv_modules.hpp"
+#include "cvconfig.h"
+
#include "opencv2/core/utility.hpp"
#include "opencv2/core/core_c.h"
#include "opencv2/core/cuda.hpp"
#if defined WIN32 || defined _WIN32
void deleteThreadAllocData();
-void deleteThreadData();
#endif
template<typename T1, typename T2=T1, typename T3=T1> struct OpAdd
void convertAndUnrollScalar( const Mat& sc, int buftype, uchar* scbuf, size_t blocksize );
-struct TLSData
+struct CoreTLSData
{
- TLSData();
+ CoreTLSData() : device(0), useOpenCL(-1)
+ {}
+
RNG rng;
int device;
ocl::Queue oclQueue;
int useOpenCL; // 1 - use, 0 - do not use, -1 - auto/not initialized
-
- static TLSData* get();
};
+extern TLSData<CoreTLSData> coreTlsData;
+
#if defined(BUILD_SHARED_LIBS)
#if defined WIN32 || defined _WIN32 || defined WINCE
#define CL_RUNTIME_EXPORT __declspec(dllexport)
cv::RNG& cv::theRNG()
{
- return TLSData::get()->rng;
+ return coreTlsData.get()->rng;
}
void cv::randu(InputOutputArray dst, InputArray low, InputArray high)
}
+namespace cv
+{
+
+template <typename T>
+void getMinMaxRes(const Mat &minv, const Mat &maxv, const Mat &minl, const Mat &maxl, double* minVal,
+ double* maxVal, int* minLoc, int* maxLoc, const int groupnum, const int cn, const int cols)
+{
+ T min = std::numeric_limits<T>::max();
+ T max = std::numeric_limits<T>::min() > 0 ? -std::numeric_limits<T>::max() : std::numeric_limits<T>::min();
+ int minloc = INT_MAX, maxloc = INT_MAX;
+ for( int i = 0; i < groupnum; i++)
+ {
+ T current_min = minv.at<T>(0,i);
+ T current_max = maxv.at<T>(0,i);
+ T oldmin = min, oldmax = max;
+ min = std::min(min, current_min);
+ max = std::max(max, current_max);
+ if (cn == 1)
+ {
+ int current_minloc = minl.at<int>(0,i);
+ int current_maxloc = maxl.at<int>(0,i);
+ if(current_minloc < 0 || current_maxloc < 0) continue;
+ minloc = (oldmin == current_min) ? std::min(minloc, current_minloc) : (oldmin < current_min) ? minloc : current_minloc;
+ maxloc = (oldmax == current_max) ? std::min(maxloc, current_maxloc) : (oldmax > current_max) ? maxloc : current_maxloc;
+ }
+ }
+ bool zero_mask = (maxloc == INT_MAX) || (minloc == INT_MAX);
+ if(minVal)
+ *minVal = zero_mask ? 0 : (double)min;
+ if(maxVal)
+ *maxVal = zero_mask ? 0 : (double)max;
+ if(minLoc)
+ {
+ minLoc[0] = zero_mask ? -1 : minloc/cols;
+ minLoc[1] = zero_mask ? -1 : minloc%cols;
+ }
+ if(maxLoc)
+ {
+ maxLoc[0] = zero_mask ? -1 : maxloc/cols;
+ maxLoc[1] = zero_mask ? -1 : maxloc%cols;
+ }
+}
+
+typedef void (*getMinMaxResFunc)(const Mat &minv, const Mat &maxv, const Mat &minl, const Mat &maxl, double *minVal,
+ double *maxVal, int *minLoc, int *maxLoc, const int gropunum, const int cn, const int cols);
+
+static bool ocl_minMaxIdx( InputArray _src, double* minVal, double* maxVal, int* minLoc, int* maxLoc, InputArray _mask)
+{
+ CV_Assert( (_src.channels() == 1 && (_mask.empty() || _mask.type() == CV_8U)) ||
+ (_src.channels() >= 1 && _mask.empty() && !minLoc && !maxLoc) );
+
+ int type = _src.type(), depth = CV_MAT_DEPTH(type);
+ bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0;
+
+ if (depth == CV_64F && !doubleSupport)
+ return false;
+
+ int groupnum = ocl::Device::getDefault().maxComputeUnits();
+ size_t wgs = ocl::Device::getDefault().maxWorkGroupSize();
+
+ int wgs2_aligned = 1;
+ while (wgs2_aligned < (int)wgs)
+ wgs2_aligned <<= 1;
+ wgs2_aligned >>= 1;
+
+ String opts = format("-D DEPTH_%d -D OP_MIN_MAX_LOC%s -D WGS=%d -D WGS2_ALIGNED=%d %s",
+ depth, _mask.empty() ? "" : "_MASK", (int)wgs, wgs2_aligned, doubleSupport ? "-D DOUBLE_SUPPORT" : "");
+
+ ocl::Kernel k("reduce", ocl::core::reduce_oclsrc, opts);
+ if (k.empty())
+ return false;
+
+ UMat src = _src.getUMat(), minval(1, groupnum, src.type()),
+ maxval(1, groupnum, src.type()), minloc( 1, groupnum, CV_32SC1),
+ maxloc( 1, groupnum, CV_32SC1), mask;
+ if(!_mask.empty())
+ mask = _mask.getUMat();
+
+ if(src.channels()>1)
+ src = src.reshape(1);
+
+ if(mask.empty())
+ k.args(ocl::KernelArg::ReadOnlyNoSize(src), src.cols, (int)src.total(),
+ groupnum, ocl::KernelArg::PtrWriteOnly(minval), ocl::KernelArg::PtrWriteOnly(maxval),
+ ocl::KernelArg::PtrWriteOnly(minloc), ocl::KernelArg::PtrWriteOnly(maxloc));
+ else
+ k.args(ocl::KernelArg::ReadOnlyNoSize(src), src.cols, (int)src.total(), groupnum,
+ ocl::KernelArg::PtrWriteOnly(minval), ocl::KernelArg::PtrWriteOnly(maxval),
+ ocl::KernelArg::PtrWriteOnly(minloc), ocl::KernelArg::PtrWriteOnly(maxloc), ocl::KernelArg::ReadOnlyNoSize(mask));
+
+ size_t globalsize = groupnum * wgs;
+ if (!k.run(1, &globalsize, &wgs, true))
+ return false;
+
+ Mat minv = minval.getMat(ACCESS_READ), maxv = maxval.getMat(ACCESS_READ),
+ minl = minloc.getMat(ACCESS_READ), maxl = maxloc.getMat(ACCESS_READ);
+
+ static getMinMaxResFunc functab[7] =
+ {
+ getMinMaxRes<uchar>,
+ getMinMaxRes<char>,
+ getMinMaxRes<ushort>,
+ getMinMaxRes<short>,
+ getMinMaxRes<int>,
+ getMinMaxRes<float>,
+ getMinMaxRes<double>
+ };
+
+ getMinMaxResFunc func;
+
+ func = functab[depth];
+ func(minv, maxv, minl, maxl, minVal, maxVal, minLoc, maxLoc, groupnum, src.channels(), src.cols);
+
+ return true;
+}
+}
+
void cv::minMaxIdx(InputArray _src, double* minVal,
double* maxVal, int* minIdx, int* maxIdx,
InputArray _mask)
{
+ CV_Assert( (_src.channels() == 1 && (_mask.empty() || _mask.type() == CV_8U)) ||
+ (_src.channels() >= 1 && _mask.empty() && !minIdx && !maxIdx) );
+
+ if( ocl::useOpenCL() && _src.isUMat() && _src.dims() <= 2 && ( _mask.empty() || _src.size() == _mask.size() )
+ && ocl_minMaxIdx(_src, minVal, maxVal, minIdx, maxIdx, _mask) )
+ return;
+
Mat src = _src.getMat(), mask = _mask.getMat();
int depth = src.depth(), cn = src.channels();
- CV_Assert( (cn == 1 && (mask.empty() || mask.type() == CV_8U)) ||
- (cn >= 1 && mask.empty() && !minIdx && !maxIdx) );
-
#if defined (HAVE_IPP) && (IPP_VERSION_MAJOR >= 7)
size_t total_size = src.total();
int rows = src.size[0], cols = (int)(total_size/rows);
void cv::minMaxLoc( InputArray _img, double* minVal, double* maxVal,
Point* minLoc, Point* maxLoc, InputArray mask )
{
- Mat img = _img.getMat();
- CV_Assert(img.dims <= 2);
+ CV_Assert(_img.dims() <= 2);
minMaxIdx(_img, minVal, maxVal, (int*)minLoc, (int*)maxLoc, mask);
if( minLoc )
bool __termination = false;
}
-#if defined CVAPI_EXPORTS && defined WIN32 && !defined WINCE
-#ifdef HAVE_WINRT
- #pragma warning(disable:4447) // Disable warning 'main' signature found without threading model
-#endif
-
-BOOL WINAPI DllMain( HINSTANCE, DWORD, LPVOID );
-
-BOOL WINAPI DllMain( HINSTANCE, DWORD fdwReason, LPVOID lpReserved )
-{
- if( fdwReason == DLL_THREAD_DETACH || fdwReason == DLL_PROCESS_DETACH )
- {
- if (lpReserved != NULL) // called after ExitProcess() call
- cv::__termination = true;
- cv::deleteThreadAllocData();
- cv::deleteThreadData();
- }
- return TRUE;
-}
-#endif
-
namespace cv
{
void Mutex::unlock() { impl->unlock(); }
bool Mutex::trylock() { return impl->trylock(); }
-}
//////////////////////////////// thread-local storage ////////////////////////////////
-namespace cv
+class TLSStorage
{
+ std::vector<void*> tlsData_;
+public:
+ TLSStorage() { tlsData_.reserve(16); }
+ ~TLSStorage();
+ inline void* getData(int key) const
+ {
+ CV_DbgAssert(key >= 0);
+ return (key < (int)tlsData_.size()) ? tlsData_[key] : NULL;
+ }
+ inline void setData(int key, void* data)
+ {
+ CV_DbgAssert(key >= 0);
+ if (key >= (int)tlsData_.size())
+ {
+ tlsData_.resize(key + 1, NULL);
+ }
+ tlsData_[key] = data;
+ }
-TLSData::TLSData()
-{
- device = 0;
- useOpenCL = -1;
-}
+ inline static TLSStorage* get();
+};
#ifdef WIN32
+#pragma warning(disable:4505) // unreferenced local function has been removed
#ifdef HAVE_WINRT
// using C++11 thread attribute for local thread data
- static __declspec( thread ) TLSData* g_tlsdata = NULL;
+ static __declspec( thread ) TLSStorage* g_tlsdata = NULL;
- static void deleteThreadRNGData()
+ static void deleteThreadData()
{
if (g_tlsdata)
+ {
delete g_tlsdata;
+ g_tlsdata = NULL;
+ }
}
- TLSData* TLSData::get()
+ inline TLSStorage* TLSStorage::get()
{
if (!g_tlsdata)
{
- g_tlsdata = new TLSData;
+ g_tlsdata = new TLSStorage;
}
return g_tlsdata;
}
#endif
static DWORD tlsKey = TLS_OUT_OF_INDEXES;
- void deleteThreadData()
+ static void deleteThreadData()
{
- if( tlsKey != TLS_OUT_OF_INDEXES )
- delete (TLSData*)TlsGetValue( tlsKey );
+ if(tlsKey != TLS_OUT_OF_INDEXES)
+ {
+ delete (TLSStorage*)TlsGetValue(tlsKey);
+ TlsSetValue(tlsKey, NULL);
+ }
}
- TLSData* TLSData::get()
+ inline TLSStorage* TLSStorage::get()
{
- if( tlsKey == TLS_OUT_OF_INDEXES )
+ if (tlsKey == TLS_OUT_OF_INDEXES)
{
tlsKey = TlsAlloc();
CV_Assert(tlsKey != TLS_OUT_OF_INDEXES);
}
- TLSData* d = (TLSData*)TlsGetValue( tlsKey );
- if( !d )
+ TLSStorage* d = (TLSStorage*)TlsGetValue(tlsKey);
+ if (!d)
{
- d = new TLSData;
- TlsSetValue( tlsKey, d );
+ d = new TLSStorage;
+ TlsSetValue(tlsKey, d);
}
return d;
}
#endif //HAVE_WINRT
+
+#if defined CVAPI_EXPORTS && defined WIN32 && !defined WINCE
+#ifdef HAVE_WINRT
+ #pragma warning(disable:4447) // Disable warning 'main' signature found without threading model
+#endif
+
+BOOL WINAPI DllMain(HINSTANCE, DWORD fdwReason, LPVOID);
+
+BOOL WINAPI DllMain(HINSTANCE, DWORD fdwReason, LPVOID lpReserved)
+{
+ if (fdwReason == DLL_THREAD_DETACH || fdwReason == DLL_PROCESS_DETACH)
+ {
+ if (lpReserved != NULL) // called after ExitProcess() call
+ cv::__termination = true;
+ cv::deleteThreadAllocData();
+ cv::deleteThreadData();
+ }
+ return TRUE;
+}
+#endif
+
#else
static pthread_key_t tlsKey = 0;
static pthread_once_t tlsKeyOnce = PTHREAD_ONCE_INIT;
- static void deleteTLSData(void* data)
+ static void deleteTLSStorage(void* data)
{
- delete (TLSData*)data;
+ delete (TLSStorage*)data;
}
static void makeKey()
{
- int errcode = pthread_key_create(&tlsKey, deleteTLSData);
+ int errcode = pthread_key_create(&tlsKey, deleteTLSStorage);
CV_Assert(errcode == 0);
}
- TLSData* TLSData::get()
+ inline TLSStorage* TLSStorage::get()
{
pthread_once(&tlsKeyOnce, makeKey);
- TLSData* d = (TLSData*)pthread_getspecific(tlsKey);
+ TLSStorage* d = (TLSStorage*)pthread_getspecific(tlsKey);
if( !d )
{
- d = new TLSData;
+ d = new TLSStorage;
pthread_setspecific(tlsKey, d);
}
return d;
}
#endif
+
+class TLSContainerStorage
+{
+ cv::Mutex mutex_;
+ std::vector<TLSDataContainer*> tlsContainers_;
+public:
+ TLSContainerStorage() { }
+ ~TLSContainerStorage()
+ {
+ for (size_t i = 0; i < tlsContainers_.size(); i++)
+ {
+ CV_DbgAssert(tlsContainers_[i] == NULL); // not all keys released
+ tlsContainers_[i] = NULL;
+ }
+ }
+
+ int allocateKey(TLSDataContainer* pContainer)
+ {
+ cv::AutoLock lock(mutex_);
+ tlsContainers_.push_back(pContainer);
+ return (int)tlsContainers_.size() - 1;
+ }
+ void releaseKey(int id, TLSDataContainer* pContainer)
+ {
+ cv::AutoLock lock(mutex_);
+ CV_Assert(tlsContainers_[id] == pContainer);
+ tlsContainers_[id] = NULL;
+ // currently, we don't go into thread's TLSData and release data for this key
+ }
+
+ void destroyData(int key, void* data)
+ {
+ cv::AutoLock lock(mutex_);
+ TLSDataContainer* k = tlsContainers_[key];
+ if (!k)
+ return;
+ try
+ {
+ k->deleteDataInstance(data);
+ }
+ catch (...)
+ {
+ CV_DbgAssert(k == NULL); // Debug this!
+ }
+ }
+};
+static TLSContainerStorage tlsContainerStorage;
+
+TLSDataContainer::TLSDataContainer()
+ : key_(-1)
+{
+ key_ = tlsContainerStorage.allocateKey(this);
+}
+
+TLSDataContainer::~TLSDataContainer()
+{
+ tlsContainerStorage.releaseKey(key_, this);
+ key_ = -1;
+}
+
+void* TLSDataContainer::getData() const
+{
+ CV_Assert(key_ >= 0);
+ TLSStorage* tlsData = TLSStorage::get();
+ void* data = tlsData->getData(key_);
+ if (!data)
+ {
+ data = this->createDataInstance();
+ CV_DbgAssert(data != NULL);
+ tlsData->setData(key_, data);
+ }
+ return data;
}
+TLSStorage::~TLSStorage()
+{
+ for (int i = 0; i < (int)tlsData_.size(); i++)
+ {
+ void*& data = tlsData_[i];
+ if (data)
+ {
+ tlsContainerStorage.destroyData(i, data);
+ data = NULL;
+ }
+ }
+ tlsData_.clear();
+}
+
+TLSData<CoreTLSData> coreTlsData;
+
+} // namespace cv
+
/* End of file. */
if(!a)
a = a0;
temp_u = a->allocate(dims, size.p, type(), data, step.p, accessFlags);
+ temp_u->refcount = 1;
}
UMat::getStdAllocator()->allocate(temp_u, accessFlags);
hdr.flags = flags;
finalizeHdr(hdr);
hdr.u = temp_u;
hdr.offset = data - datastart;
+ hdr.addref();
return hdr;
}
}
finalizeHdr(*this);
+ addref();
}
void UMat::copySize(const UMat& m)
void UMat::deallocate()
{
u->currAllocator->deallocate(u);
+ u = NULL;
}
if( _dst.kind() == _InputArray::UMAT )
{
UMat dst = _dst.getUMat();
- void* srchandle = handle(ACCESS_READ);
- void* dsthandle = dst.handle(ACCESS_WRITE);
- if( srchandle == dsthandle && dst.offset == offset )
+ if( u == dst.u && dst.offset == offset )
return;
dst.ndoffset(dstofs);
+ dstofs[dims-1] *= esz;
CV_Assert(u->currAllocator == dst.u->currAllocator);
u->currAllocator->copy(u, dst.u, dims, sz, srcofs, step.p, dstofs, dst.step.p, false);
}
}
}
-void UMat::convertTo(OutputArray, int, double, double) const
+void UMat::convertTo(OutputArray _dst, int _type, double alpha, double beta) const
{
- CV_Error(Error::StsNotImplemented, "");
+ bool noScale = std::fabs(alpha - 1) < DBL_EPSILON && std::fabs(beta) < DBL_EPSILON;
+ int stype = type(), cn = CV_MAT_CN(stype);
+
+ if( _type < 0 )
+ _type = _dst.fixedType() ? _dst.type() : stype;
+ else
+ _type = CV_MAKETYPE(CV_MAT_DEPTH(_type), cn);
+
+ int sdepth = CV_MAT_DEPTH(stype), ddepth = CV_MAT_DEPTH(_type);
+ if( sdepth == ddepth && noScale )
+ {
+ copyTo(_dst);
+ return;
+ }
+
+ bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0;
+ bool needDouble = sdepth == CV_64F || ddepth == CV_64F;
+ if( dims <= 2 && cn && _dst.isUMat() && ocl::useOpenCL() &&
+ ((needDouble && doubleSupport) || !needDouble) )
+ {
+ char cvt[40];
+ ocl::Kernel k("convertTo", ocl::core::convert_oclsrc,
+ format("-D srcT=%s -D dstT=%s -D convertToDT=%s%s", ocl::typeToStr(sdepth),
+ ocl::typeToStr(ddepth), ocl::convertTypeStr(CV_32F, ddepth, 1, cvt),
+ doubleSupport ? " -D DOUBLE_SUPPORT" : ""));
+ if (!k.empty())
+ {
+ _dst.create( size(), _type );
+ UMat dst = _dst.getUMat();
+
+ float alphaf = (float)alpha, betaf = (float)beta;
+ k.args(ocl::KernelArg::ReadOnlyNoSize(*this), ocl::KernelArg::WriteOnly(dst, cn), alphaf, betaf);
+
+ size_t globalsize[2] = { dst.cols * cn, dst.rows };
+ if (k.run(2, globalsize, NULL, false))
+ return;
+ }
+ }
+
+ Mat m = getMat(ACCESS_READ);
+ m.convertTo(_dst, _type, alpha, beta);
}
UMat& UMat::setTo(InputArray _value, InputArray _mask)
{
bool haveMask = !_mask.empty();
int tp = type(), cn = CV_MAT_CN(tp);
- if( dims <= 2 && cn <= 4 && ocl::useOpenCL() )
+ if( dims <= 2 && cn <= 4 && cn != 3 && ocl::useOpenCL() )
{
Mat value = _value.getMat();
CV_Assert( checkScalar(value, type(), _value.kind(), _InputArray::UMAT) );
return *this;
}
-UMat& UMat::operator = (const Scalar&)
+UMat& UMat::operator = (const Scalar& s)
{
- CV_Error(Error::StsNotImplemented, "");
+ setTo(s);
return *this;
}
Near(0);
}
}
+//////////////////////////////////////// minMaxIdx /////////////////////////////////////////
+
+typedef ArithmTestBase MinMaxIdx;
+
+OCL_TEST_P(MinMaxIdx, Mat)
+{
+ for (int j = 0; j < test_loop_times; j++)
+ {
+ generateTestData();
+
+ int p1[2], p2[2], up1[2], up2[2];
+ double minv, maxv, uminv, umaxv;
+
+ if(src1_roi.channels() > 1)
+ {
+ OCL_OFF(cv::minMaxIdx(src2_roi, &minv, &maxv) );
+ OCL_ON(cv::minMaxIdx(usrc2_roi, &uminv, &umaxv));
+
+ EXPECT_DOUBLE_EQ(minv, uminv);
+ EXPECT_DOUBLE_EQ(maxv, umaxv);
+ }
+ else
+ {
+ OCL_OFF(cv::minMaxIdx(src2_roi, &minv, &maxv, p1, p2, noArray()));
+ OCL_ON(cv::minMaxIdx(usrc2_roi, &uminv, &umaxv, up1, up2, noArray()));
+
+ EXPECT_DOUBLE_EQ(minv, uminv);
+ EXPECT_DOUBLE_EQ(maxv, umaxv);
+ for( int i = 0; i < 2; i++)
+ {
+ EXPECT_EQ(p1[i], up1[i]);
+ EXPECT_EQ(p2[i], up2[i]);
+ }
+ }
+ }
+}
+
+typedef ArithmTestBase MinMaxIdx_Mask;
+
+OCL_TEST_P(MinMaxIdx_Mask, Mat)
+{
+ for (int j = 0; j < test_loop_times; j++)
+ {
+ generateTestData();
+
+ int p1[2], p2[2], up1[2], up2[2];
+ double minv, maxv, uminv, umaxv;
+
+ OCL_OFF(cv::minMaxIdx(src2_roi, &minv, &maxv, p1, p2, mask_roi));
+ OCL_ON(cv::minMaxIdx(usrc2_roi, &uminv, &umaxv, up1, up2, umask_roi));
+
+ EXPECT_DOUBLE_EQ(minv, uminv);
+ EXPECT_DOUBLE_EQ(maxv, umaxv);
+ for( int i = 0; i < 2; i++)
+ {
+ EXPECT_EQ(p1[i], up1[i]);
+ EXPECT_EQ(p2[i], up2[i]);
+ }
+
+ }
+}
//////////////////////////////// Norm /////////////////////////////////////////////////
OCL_INSTANTIATE_TEST_CASE_P(Arithm, Phase, Combine(::testing::Values(CV_32F, CV_64F), OCL_ALL_CHANNELS, Bool()));
OCL_INSTANTIATE_TEST_CASE_P(Arithm, Magnitude, Combine(::testing::Values(CV_32F, CV_64F), OCL_ALL_CHANNELS, Bool()));
OCL_INSTANTIATE_TEST_CASE_P(Arithm, Flip, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(Arithm, MinMaxIdx, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
+OCL_INSTANTIATE_TEST_CASE_P(Arithm, MinMaxIdx_Mask, Combine(OCL_ALL_DEPTHS, ::testing::Values(Channels(1)), Bool()));
OCL_INSTANTIATE_TEST_CASE_P(Arithm, Norm, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
OCL_INSTANTIATE_TEST_CASE_P(Arithm, Sqrt, Combine(::testing::Values(CV_32F, CV_64F), OCL_ALL_CHANNELS, Bool()));
--- /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, Multicoreware, Inc., all rights reserved.
+// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// @Authors
+// Peng Xiao, pengxiao@multicorewareinc.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"
+
+#ifdef HAVE_OPENCL
+
+namespace cvtest {
+namespace ocl {
+
+////////////////////////////////////////////////////////////////////////////
+// GEMM
+
+PARAM_TEST_CASE(Gemm,
+ MatType,
+ bool, // GEMM_1_T
+ bool, // GEMM_2_T
+ bool, // GEMM_3_T
+ bool // ROI
+ )
+{
+ bool use_roi;
+ int type, flags;
+ bool atrans, btrans, ctrans;
+
+ double alpha, beta;
+
+ TEST_DECLARE_INPUT_PARAMETER(A)
+ TEST_DECLARE_INPUT_PARAMETER(B)
+ TEST_DECLARE_INPUT_PARAMETER(C)
+ TEST_DECLARE_OUTPUT_PARAMETER(D)
+
+ virtual void SetUp()
+ {
+ atrans = btrans = ctrans = false;
+
+ type = GET_PARAM(0);
+ use_roi = GET_PARAM(4);
+
+ flags = 0;
+ if (GET_PARAM(1))
+ flags |= GEMM_1_T, atrans = true;
+ if (GET_PARAM(2))
+ flags |= GEMM_2_T, btrans = true;
+ if (GET_PARAM(3))
+ flags |= GEMM_3_T, ctrans = true;
+ }
+
+ void generateTestData()
+ {
+ Size ARoiSize = randomSize(1, MAX_VALUE);
+ Border ABorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
+ randomSubMat(A, A_roi, ARoiSize, ABorder, type, -11, 11);
+
+ if (atrans)
+ ARoiSize = Size(ARoiSize.height, ARoiSize.width);
+
+ Size BRoiSize = randomSize(1, MAX_VALUE);
+ if (btrans)
+ BRoiSize.width = ARoiSize.width;
+ else
+ BRoiSize.height = ARoiSize.width;
+
+ Border BBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
+ randomSubMat(B, B_roi, BRoiSize, BBorder, type, -11, 11);
+
+ if (btrans)
+ BRoiSize = Size(BRoiSize.height, BRoiSize.width);
+
+ Size DRoiSize = Size(BRoiSize.width, ARoiSize.height), CRoiSizeT(DRoiSize.height, DRoiSize.width);
+ Border CBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
+ randomSubMat(C, C_roi, ctrans ? CRoiSizeT : DRoiSize, CBorder, type, -11, 11);
+
+ Border DBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
+ randomSubMat(D, D_roi, DRoiSize, DBorder, type, -11, 11);
+
+ alpha = randomDouble(-4, 4);
+ beta = randomDouble(-4, 4);
+
+ UMAT_UPLOAD_INPUT_PARAMETER(A)
+ UMAT_UPLOAD_INPUT_PARAMETER(B)
+ UMAT_UPLOAD_INPUT_PARAMETER(C)
+ UMAT_UPLOAD_OUTPUT_PARAMETER(D)
+ }
+};
+
+OCL_TEST_P(Gemm, Accuracy)
+{
+ for (int i = 0; i < test_loop_times; ++i)
+ {
+ generateTestData();
+
+ OCL_OFF(cv::gemm(A_roi, B_roi, alpha, C_roi, beta, D_roi, flags));
+ OCL_ON(cv::gemm(uA_roi, uB_roi, alpha, uC_roi, beta, uD_roi, flags));
+
+ double eps = D_roi.size().area() * 1e-4;
+ OCL_EXPECT_MATS_NEAR(D, eps);
+ }
+}
+
+OCL_INSTANTIATE_TEST_CASE_P(Core, Gemm, ::testing::Combine(
+ testing::Values(CV_32FC1, CV_32FC2, CV_64FC1, CV_64FC2),
+ Bool(), Bool(), Bool(), 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"
+
+#ifdef HAVE_OPENCL
+
+namespace cvtest {
+namespace ocl {
+
+////////////////////////////////converto/////////////////////////////////////////////////
+
+PARAM_TEST_CASE(MatrixTestBase, MatDepth, MatDepth, Channels, bool)
+{
+ int src_depth, cn, dstType;
+ bool use_roi;
+
+ TEST_DECLARE_INPUT_PARAMETER(src)
+ TEST_DECLARE_OUTPUT_PARAMETER(dst)
+
+ virtual void SetUp()
+ {
+ src_depth = GET_PARAM(0);
+ cn = GET_PARAM(2);
+ dstType = CV_MAKE_TYPE(GET_PARAM(1), cn);
+
+ use_roi = GET_PARAM(3);
+ }
+
+ virtual void generateTestData()
+ {
+ Size roiSize = randomSize(1, MAX_VALUE);
+ Border srcBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
+ randomSubMat(src, src_roi, roiSize, srcBorder, CV_MAKE_TYPE(src_depth, cn), -MAX_VALUE, MAX_VALUE);
+
+ Border dstBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
+ randomSubMat(dst, dst_roi, roiSize, dstBorder, dstType, 5, 16);
+
+ UMAT_UPLOAD_INPUT_PARAMETER(src)
+ UMAT_UPLOAD_OUTPUT_PARAMETER(dst)
+ }
+};
+
+typedef MatrixTestBase ConvertTo;
+
+OCL_TEST_P(ConvertTo, Accuracy)
+{
+ for (int j = 0; j < test_loop_times; j++)
+ {
+ generateTestData();
+
+ double alpha = randomDouble(-4, 4), beta = randomDouble(-4, 4);
+
+ OCL_OFF(src_roi.convertTo(dst_roi, dstType, alpha, beta));
+ OCL_ON(usrc_roi.convertTo(udst_roi, dstType, alpha, beta));
+
+ double eps = src_depth >= CV_32F || CV_MAT_DEPTH(dstType) >= CV_32F ? 1e-4 : 1;
+ OCL_EXPECT_MATS_NEAR(dst, eps);
+ }
+}
+
+typedef MatrixTestBase CopyTo;
+
+OCL_TEST_P(CopyTo, Accuracy)
+{
+ for (int j = 0; j < test_loop_times; j++)
+ {
+ generateTestData();
+
+ OCL_OFF(src_roi.copyTo(dst_roi));
+ OCL_ON(usrc_roi.copyTo(udst_roi));
+
+ OCL_EXPECT_MATS_NEAR(dst, 0);
+ }
+}
+
+OCL_INSTANTIATE_TEST_CASE_P(MatrixOperation, ConvertTo, Combine(
+ OCL_ALL_DEPTHS, OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
+
+OCL_INSTANTIATE_TEST_CASE_P(MatrixOperation, CopyTo, Combine(
+ OCL_ALL_DEPTHS, Values((MatDepth)0), OCL_ALL_CHANNELS, Bool()));
+
+} } // namespace cvtest::ocl
+
+#endif
#include <string>
#include <iostream>
-#include <fstream>
-#include <iterator>
-#include <limits>
-#include <numeric>
#include "opencv2/core/ocl.hpp"
using namespace cv;
using namespace std;
-class CV_UMatTest : public cvtest::BaseTest
+class CV_UMatTest :
+ public cvtest::BaseTest
{
public:
CV_UMatTest() {}
TEST(Core_UMat, getUMat)
{
{
- int a[3] = { 1, 2, 3 };
- Mat m = Mat(1, 1, CV_32SC3, a);
- UMat u = m.getUMat(ACCESS_READ);
- EXPECT_NE((void*)NULL, u.u);
+ int a[3] = { 1, 2, 3 };
+ Mat m = Mat(1, 1, CV_32SC3, a);
+ UMat u = m.getUMat(ACCESS_READ);
+ EXPECT_NE((void*)NULL, u.u);
}
{
- Mat m(10, 10, CV_8UC1), ref;
- for (int y = 0; y < m.rows; ++y)
- {
- uchar * const ptr = m.ptr<uchar>(y);
- for (int x = 0; x < m.cols; ++x)
- ptr[x] = (uchar)(x + y * 2);
- }
+ Mat m(10, 10, CV_8UC1), ref;
+ for (int y = 0; y < m.rows; ++y)
+ {
+ uchar * const ptr = m.ptr<uchar>(y);
+ for (int x = 0; x < m.cols; ++x)
+ ptr[x] = (uchar)(x + y * 2);
+ }
- ref = m.clone();
- Rect r(1, 1, 8, 8);
- ref(r).setTo(17);
+ ref = m.clone();
+ Rect r(1, 1, 8, 8);
+ ref(r).setTo(17);
- {
- UMat um = m(r).getUMat(ACCESS_WRITE);
- um.setTo(17);
- }
+ {
+ UMat um = m(r).getUMat(ACCESS_WRITE);
+ um.setTo(17);
+ }
- double err = norm(m, ref, NORM_INF);
- if(err > 0)
- {
- std::cout << "m: " << m << std::endl;
- std::cout << "ref: " << ref << std::endl;
- }
- EXPECT_EQ(err, 0.);
+ double err = norm(m, ref, NORM_INF);
+ if (err > 0)
+ {
+ std::cout << "m: " << std::endl << m << std::endl;
+ std::cout << "ref: " << std::endl << ref << std::endl;
+ }
+ EXPECT_EQ(0., err);
}
}
virtual int descriptorSize() const = 0;
virtual int descriptorType() const = 0;
+ virtual int defaultNorm() const = 0;
static Ptr<DescriptorExtractor> create( const String& descriptorExtractorType );
virtual void write( FileStorage& ) const;
virtual int descriptorSize() const;
virtual int descriptorType() const;
+ virtual int defaultNorm() const;
protected:
...
};
virtual void write( FileStorage& ) const;
virtual int descriptorSize() const;
virtual int descriptorType() const;
+ virtual int defaultNorm() const;
protected:
...
};
void add( const Mat& descriptors );
const vector<Mat>& getDescriptors() const;
- int descripotorsCount() const;
+ int descriptorsCount() const;
virtual void clear();
-BOWTrainer::descripotorsCount
+BOWTrainer::descriptorsCount
---------------------------------
Returns the count of all descriptors stored in the training set.
-.. ocv:function:: int BOWTrainer::descripotorsCount() const
+.. ocv:function:: int BOWTrainer::descriptorsCount() const
CV_WRAP virtual int descriptorSize() const = 0;
CV_WRAP virtual int descriptorType() const = 0;
+ CV_WRAP virtual int defaultNorm() const = 0;
CV_WRAP virtual bool empty() const;
int descriptorSize() const;
// returns the descriptor type
int descriptorType() const;
+ // returns the default norm type
+ int defaultNorm() const;
// Compute the BRISK features on an image
void operator()(InputArray image, InputArray mask, std::vector<KeyPoint>& keypoints) const;
int descriptorSize() const;
// returns the descriptor type
int descriptorType() const;
+ // returns the default norm type
+ int defaultNorm() const;
// Compute the ORB features and descriptors on an image
void operator()(InputArray image, InputArray mask, std::vector<KeyPoint>& keypoints) const;
/** returns the descriptor type */
virtual int descriptorType() const;
+ /** returns the default norm type */
+ virtual int defaultNorm() const;
+
/** select the 512 "best description pairs"
* @param images grayscale images set
* @param keypoints set of detected keypoints
virtual int descriptorSize() const;
virtual int descriptorType() const;
+ virtual int defaultNorm() const;
virtual bool empty() const;
virtual int descriptorSize() const;
virtual int descriptorType() const;
+ virtual int defaultNorm() const;
/// @todo read and write for brief
void add( const Mat& descriptors );
const std::vector<Mat>& getDescriptors() const;
- int descripotorsCount() const;
+ int descriptorsCount() const;
virtual void clear();
return descriptors;
}
-int BOWTrainer::descripotorsCount() const
+int BOWTrainer::descriptorsCount() const
{
return descriptors.empty() ? 0 : size;
}
return CV_8UC1;
}
+int BriefDescriptorExtractor::defaultNorm() const
+{
+ return NORM_HAMMING;
+}
+
void BriefDescriptorExtractor::read( const FileNode& fn)
{
int dSize = fn["descriptorSize"];
return CV_8U;
}
+int
+BRISK::defaultNorm() const
+{
+ return NORM_HAMMING;
+}
+
BRISK::~BRISK()
{
delete[] patternPoints_;
return descriptorExtractor->descriptorType();
}
+int OpponentColorDescriptorExtractor::defaultNorm() const
+{
+ return descriptorExtractor->defaultNorm();
+}
+
bool OpponentColorDescriptorExtractor::empty() const
{
return !descriptorExtractor || descriptorExtractor->empty();
static const int FREAK_NB_PAIRS = FREAK::NB_PAIRS;
static const int FREAK_NB_ORIENPAIRS = FREAK::NB_ORIENPAIRS;
+// default pairs
static const int FREAK_DEF_PAIRS[FREAK::NB_PAIRS] =
-{ // default pairs
+{
404,431,818,511,181,52,311,874,774,543,719,230,417,205,11,
560,149,265,39,306,165,857,250,8,61,15,55,717,44,412,
592,134,761,695,660,782,625,487,549,516,271,665,762,392,178,
670,249,36,581,389,605,331,518,442,822
};
+// used to sort pairs during pairs selection
struct PairStat
-{ // used to sort pairs during pairs selection
+{
double mean;
int idx;
};
struct sortMean
{
- bool operator()( const PairStat& a, const PairStat& b ) const {
+ bool operator()( const PairStat& a, const PairStat& b ) const
+ {
return a.mean < b.mean;
}
};
radius[6]/2.0, radius[6]/2.0
};
// fill the lookup table
- for( int scaleIdx=0; scaleIdx < FREAK_NB_SCALES; ++scaleIdx ) {
+ for( int scaleIdx=0; scaleIdx < FREAK_NB_SCALES; ++scaleIdx )
+ {
patternSizes[scaleIdx] = 0; // proper initialization
scalingFactor = std::pow(scaleStep,scaleIdx); //scale of the pattern, scaleStep ^ scaleIdx
- for( int orientationIdx = 0; orientationIdx < FREAK_NB_ORIENTATION; ++orientationIdx ) {
+ for( int orientationIdx = 0; orientationIdx < FREAK_NB_ORIENTATION; ++orientationIdx )
+ {
theta = double(orientationIdx)* 2*CV_PI/double(FREAK_NB_ORIENTATION); // orientation of the pattern
int pointIdx = 0;
PatternPoint* patternLookupPtr = &patternLookup[0];
- for( size_t i = 0; i < 8; ++i ) {
- for( int k = 0 ; k < n[i]; ++k ) {
+ for( size_t i = 0; i < 8; ++i )
+ {
+ for( int k = 0 ; k < n[i]; ++k )
+ {
beta = CV_PI/n[i] * (i%2); // orientation offset so that groups of points on each circles are staggered
alpha = double(k)* 2*CV_PI/double(n[i])+beta+theta;
orientationPairs[39].i=30; orientationPairs[39].j=33; orientationPairs[40].i=31; orientationPairs[40].j=34; orientationPairs[41].i=32; orientationPairs[41].j=35;
orientationPairs[42].i=36; orientationPairs[42].j=39; orientationPairs[43].i=37; orientationPairs[43].j=40; orientationPairs[44].i=38; orientationPairs[44].j=41;
- for( unsigned m = FREAK_NB_ORIENPAIRS; m--; ) {
+ for( unsigned m = FREAK_NB_ORIENPAIRS; m--; )
+ {
const float dx = patternLookup[orientationPairs[m].i].x-patternLookup[orientationPairs[m].j].x;
const float dy = patternLookup[orientationPairs[m].i].y-patternLookup[orientationPairs[m].j].y;
const float norm_sq = (dx*dx+dy*dy);
// build the list of description pairs
std::vector<DescriptionPair> allPairs;
- for( unsigned int i = 1; i < (unsigned int)FREAK_NB_POINTS; ++i ) {
+ for( unsigned int i = 1; i < (unsigned int)FREAK_NB_POINTS; ++i )
+ {
// (generate all the pairs)
- for( unsigned int j = 0; (unsigned int)j < i; ++j ) {
+ for( unsigned int j = 0; (unsigned int)j < i; ++j )
+ {
DescriptionPair pair = {(uchar)i,(uchar)j};
allPairs.push_back(pair);
}
}
// Input vector provided
- if( !selectedPairs0.empty() ) {
- if( (int)selectedPairs0.size() == FREAK_NB_PAIRS ) {
+ if( !selectedPairs0.empty() )
+ {
+ if( (int)selectedPairs0.size() == FREAK_NB_PAIRS )
+ {
for( int i = 0; i < FREAK_NB_PAIRS; ++i )
descriptionPairs[i] = allPairs[selectedPairs0.at(i)];
}
- else {
+ else
+ {
CV_Error(Error::StsVecLengthErr, "Input vector does not match the required size");
}
}
- else { // default selected pairs
+ else // default selected pairs
+ {
for( int i = 0; i < FREAK_NB_PAIRS; ++i )
descriptionPairs[i] = allPairs[FREAK_DEF_PAIRS[i]];
}
}
-void FREAK::computeImpl( const Mat& image, std::vector<KeyPoint>& keypoints, Mat& descriptors ) const {
+void FREAK::computeImpl( const Mat& image, std::vector<KeyPoint>& keypoints, Mat& descriptors ) const
+{
if( image.empty() )
return;
int direction1;
// compute the scale index corresponding to the keypoint size and remove keypoints close to the border
- if( scaleNormalized ) {
- for( size_t k = keypoints.size(); k--; ) {
+ if( scaleNormalized )
+ {
+ for( size_t k = keypoints.size(); k--; )
+ {
//Is k non-zero? If so, decrement it and continue"
kpScaleIdx[k] = std::max( (int)(std::log(keypoints[k].size/FREAK_SMALLEST_KP_SIZE)*sizeCst+0.5) ,0);
if( kpScaleIdx[k] >= FREAK_NB_SCALES )
keypoints[k].pt.y <= patternSizes[kpScaleIdx[k]] ||
keypoints[k].pt.x >= image.cols-patternSizes[kpScaleIdx[k]] ||
keypoints[k].pt.y >= image.rows-patternSizes[kpScaleIdx[k]]
- ) {
+ )
+ {
keypoints.erase(kpBegin+k);
kpScaleIdx.erase(ScaleIdxBegin+k);
}
}
}
- else {
+ else
+ {
const int scIdx = std::max( (int)(1.0986122886681*sizeCst+0.5) ,0);
- for( size_t k = keypoints.size(); k--; ) {
+ for( size_t k = keypoints.size(); k--; )
+ {
kpScaleIdx[k] = scIdx; // equivalent to the formule when the scale is normalized with a constant size of keypoints[k].size=3*SMALLEST_KP_SIZE
- if( kpScaleIdx[k] >= FREAK_NB_SCALES ) {
+ if( kpScaleIdx[k] >= FREAK_NB_SCALES )
+ {
kpScaleIdx[k] = FREAK_NB_SCALES-1;
}
if( keypoints[k].pt.x <= patternSizes[kpScaleIdx[k]] ||
keypoints[k].pt.y <= patternSizes[kpScaleIdx[k]] ||
keypoints[k].pt.x >= image.cols-patternSizes[kpScaleIdx[k]] ||
keypoints[k].pt.y >= image.rows-patternSizes[kpScaleIdx[k]]
- ) {
+ )
+ {
keypoints.erase(kpBegin+k);
kpScaleIdx.erase(ScaleIdxBegin+k);
}
}
// allocate descriptor memory, estimate orientations, extract descriptors
- if( !extAll ) {
+ if( !extAll )
+ {
// extract the best comparisons only
descriptors = cv::Mat::zeros((int)keypoints.size(), FREAK_NB_PAIRS/8, CV_8U);
#if CV_SSE2
#else
std::bitset<FREAK_NB_PAIRS>* ptr = (std::bitset<FREAK_NB_PAIRS>*) (descriptors.data+(keypoints.size()-1)*descriptors.step[0]);
#endif
- for( size_t k = keypoints.size(); k--; ) {
+ for( size_t k = keypoints.size(); k--; )
+ {
// estimate orientation (gradient)
- if( !orientationNormalized ) {
+ if( !orientationNormalized )
+ {
thetaIdx = 0; // assign 0° to all keypoints
keypoints[k].angle = 0.0;
}
- else {
+ else
+ {
// get the points intensity value in the un-rotated pattern
- for( int i = FREAK_NB_POINTS; i--; ) {
+ for( int i = FREAK_NB_POINTS; i--; )
+ {
pointsValue[i] = meanIntensity(image, imgIntegral, keypoints[k].pt.x,keypoints[k].pt.y, kpScaleIdx[k], 0, i);
}
direction0 = 0;
direction1 = 0;
- for( int m = 45; m--; ) {
+ for( int m = 45; m--; )
+ {
//iterate through the orientation pairs
const int delta = (pointsValue[ orientationPairs[m].i ]-pointsValue[ orientationPairs[m].j ]);
direction0 += delta*(orientationPairs[m].weight_dx)/2048;
thetaIdx -= FREAK_NB_ORIENTATION;
}
// extract descriptor at the computed orientation
- for( int i = FREAK_NB_POINTS; i--; ) {
+ for( int i = FREAK_NB_POINTS; i--; )
+ {
pointsValue[i] = meanIntensity(image, imgIntegral, keypoints[k].pt.x,keypoints[k].pt.y, kpScaleIdx[k], thetaIdx, i);
}
#if CV_SSE2
#endif
}
}
- else { // extract all possible comparisons for selection
+ else // extract all possible comparisons for selection
+ {
descriptors = cv::Mat::zeros((int)keypoints.size(), 128, CV_8U);
std::bitset<1024>* ptr = (std::bitset<1024>*) (descriptors.data+(keypoints.size()-1)*descriptors.step[0]);
- for( size_t k = keypoints.size(); k--; ) {
+ for( size_t k = keypoints.size(); k--; )
+ {
//estimate orientation (gradient)
- if( !orientationNormalized ) {
+ if( !orientationNormalized )
+ {
thetaIdx = 0;//assign 0° to all keypoints
keypoints[k].angle = 0.0;
}
- else {
+ else
+ {
//get the points intensity value in the un-rotated pattern
for( int i = FREAK_NB_POINTS;i--; )
pointsValue[i] = meanIntensity(image, imgIntegral, keypoints[k].pt.x,keypoints[k].pt.y, kpScaleIdx[k], 0, i);
direction0 = 0;
direction1 = 0;
- for( int m = 45; m--; ) {
+ for( int m = 45; m--; )
+ {
//iterate through the orientation pairs
const int delta = (pointsValue[ orientationPairs[m].i ]-pointsValue[ orientationPairs[m].j ]);
direction0 += delta*(orientationPairs[m].weight_dx)/2048;
thetaIdx -= FREAK_NB_ORIENTATION;
}
// get the points intensity value in the rotated pattern
- for( int i = FREAK_NB_POINTS; i--; ) {
+ for( int i = FREAK_NB_POINTS; i--; )
+ {
pointsValue[i] = meanIntensity(image, imgIntegral, keypoints[k].pt.x,
keypoints[k].pt.y, kpScaleIdx[k], thetaIdx, i);
}
int cnt(0);
- for( int i = 1; i < FREAK_NB_POINTS; ++i ) {
+ for( int i = 1; i < FREAK_NB_POINTS; ++i )
+ {
//(generate all the pairs)
- for( int j = 0; j < i; ++j ) {
+ for( int j = 0; j < i; ++j )
+ {
ptr->set(cnt, pointsValue[i] >= pointsValue[j] );
++cnt;
}
const float kp_y,
const unsigned int scale,
const unsigned int rot,
- const unsigned int point) const {
+ const unsigned int point) const
+{
// get point position in image
const PatternPoint& FreakPoint = patternLookup[scale*FREAK_NB_ORIENTATION*FREAK_NB_POINTS + rot*FREAK_NB_POINTS + point];
const float xf = FreakPoint.x+kp_x;
const float radius = FreakPoint.sigma;
// calculate output:
- if( radius < 0.5 ) {
+ if( radius < 0.5 )
+ {
// interpolation multipliers:
const int r_x = static_cast<int>((xf-x)*1024);
const int r_y = static_cast<int>((yf-y)*1024);
if( verbose )
std::cout << "Number of images: " << images.size() << std::endl;
- for( size_t i = 0;i < images.size(); ++i ) {
+ for( size_t i = 0;i < images.size(); ++i )
+ {
Mat descriptorsTmp;
computeImpl(images[i],keypoints[i],descriptorsTmp);
descriptors.push_back(descriptorsTmp);
Mat descriptorsFloat = Mat::zeros(descriptors.rows, 903, CV_32F);
std::bitset<1024>* ptr = (std::bitset<1024>*) (descriptors.data+(descriptors.rows-1)*descriptors.step[0]);
- for( int m = descriptors.rows; m--; ) {
- for( int n = 903; n--; ) {
+ for( int m = descriptors.rows; m--; )
+ {
+ for( int n = 903; n--; )
+ {
if( ptr->test(n) == true )
descriptorsFloat.at<float>(m,n)=1.0f;
}
}
std::vector<PairStat> pairStat;
- for( int n = 903; n--; ) {
+ for( int n = 903; n--; )
+ {
// the higher the variance, the better --> mean = 0.5
PairStat tmp = { fabs( mean(descriptorsFloat.col(n))[0]-0.5 ) ,n};
pairStat.push_back(tmp);
std::sort( pairStat.begin(),pairStat.end(), sortMean() );
std::vector<PairStat> bestPairs;
- for( int m = 0; m < 903; ++m ) {
+ for( int m = 0; m < 903; ++m )
+ {
if( verbose )
std::cout << m << ":" << bestPairs.size() << " " << std::flush;
double corrMax(0);
- for( size_t n = 0; n < bestPairs.size(); ++n ) {
+ for( size_t n = 0; n < bestPairs.size(); ++n )
+ {
int idxA = bestPairs[n].idx;
int idxB = pairStat[m].idx;
double corr(0);
// compute correlation between 2 pairs
corr = fabs(compareHist(descriptorsFloat.col(idxA), descriptorsFloat.col(idxB), HISTCMP_CORREL));
- if( corr > corrMax ) {
+ if( corr > corrMax )
+ {
corrMax = corr;
if( corrMax >= corrTresh )
break;
if( corrMax < corrTresh/*0.7*/ )
bestPairs.push_back(pairStat[m]);
- if( bestPairs.size() >= 512 ) {
+ if( bestPairs.size() >= 512 )
+ {
if( verbose )
std::cout << m << std::endl;
break;
}
std::vector<int> idxBestPairs;
- if( (int)bestPairs.size() >= FREAK_NB_PAIRS ) {
+ if( (int)bestPairs.size() >= FREAK_NB_PAIRS )
+ {
for( int i = 0; i < FREAK_NB_PAIRS; ++i )
idxBestPairs.push_back(bestPairs[i].idx);
}
- else {
+ else
+ {
if( verbose )
std::cout << "correlation threshold too small (restrictive)" << std::endl;
CV_Error(Error::StsError, "correlation threshold too small (restrictive)");
/*
+// create an image showing the brisk pattern
void FREAKImpl::drawPattern()
-{ // create an image showing the brisk pattern
+{
Mat pattern = Mat::zeros(1000, 1000, CV_8UC3) + Scalar(255,255,255);
int sFac = 500 / patternScale;
- for( int n = 0; n < kNB_POINTS; ++n ) {
+ for( int n = 0; n < kNB_POINTS; ++n )
+ {
PatternPoint& pt = patternLookup[n];
circle(pattern, Point( pt.x*sFac,pt.y*sFac)+Point(500,500), pt.sigma*sFac, Scalar(0,0,255),2);
// rectangle(pattern, Point( (pt.x-pt.sigma)*sFac,(pt.y-pt.sigma)*sFac)+Point(500,500), Point( (pt.x+pt.sigma)*sFac,(pt.y+pt.sigma)*sFac)+Point(500,500), Scalar(0,0,255),2);
{
}
-int FREAK::descriptorSize() const {
+int FREAK::descriptorSize() const
+{
return FREAK_NB_PAIRS / 8; // descriptor length in bytes
}
-int FREAK::descriptorType() const {
+int FREAK::descriptorType() const
+{
return CV_8U;
}
+int FREAK::defaultNorm() const
+{
+ return NORM_HAMMING;
+}
+
} // END NAMESPACE CV
return CV_8U;
}
+int ORB::defaultNorm() const
+{
+ return NORM_HAMMING;
+}
+
/** Compute the ORB features and descriptors on an image
* @param img the image to compute the features and descriptors on
* @param mask the mask to apply
TEST(Features2d_RotationInvariance_Descriptor_BRISK, regression)
{
DescriptorRotationInvarianceTest test(Algorithm::create<FeatureDetector>("Feature2D.BRISK"),
- Algorithm::create<DescriptorExtractor>("Feature2D.BRISK"),
- NORM_HAMMING,
+ Algorithm::create<DescriptorExtractor>("Feature2D.BRISK"),
+ Algorithm::create<DescriptorExtractor>("Feature2D.BRISK")->defaultNorm(),
0.99f);
test.safe_run();
}
{
DescriptorRotationInvarianceTest test(Algorithm::create<FeatureDetector>("Feature2D.ORB"),
Algorithm::create<DescriptorExtractor>("Feature2D.ORB"),
- NORM_HAMMING,
+ Algorithm::create<DescriptorExtractor>("Feature2D.ORB")->defaultNorm(),
0.99f);
test.safe_run();
}
//{
// DescriptorRotationInvarianceTest test(Algorithm::create<FeatureDetector>("Feature2D.ORB"),
// Algorithm::create<DescriptorExtractor>("Feature2D.FREAK"),
-// NORM_HAMMING,
+// Algorithm::create<DescriptorExtractor>("Feature2D.FREAK")->defaultNorm(),
// 0.f);
// test.safe_run();
//}
//TEST(Features2d_ScaleInvariance_Descriptor_BRISK, regression)
//{
// DescriptorScaleInvarianceTest test(Algorithm::create<FeatureDetector>("Feature2D.BRISK"),
-// Algorithm::create<DescriptorExtractor>("Feature2D.BRISK"),
-// NORM_HAMMING,
-// 0.99f);
+// Algorithm::create<DescriptorExtractor>("Feature2D.BRISK"),
+// Algorithm::create<DescriptorExtractor>("Feature2D.BRISK")->defaultNorm(),
+// 0.99f);
// test.safe_run();
//}
//TEST(Features2d_ScaleInvariance_Descriptor_ORB, regression)
//{
// DescriptorScaleInvarianceTest test(Algorithm::create<FeatureDetector>("Feature2D.ORB"),
-// Algorithm::create<DescriptorExtractor>("Feature2D.ORB"),
-// NORM_HAMMING,
-// 0.01f);
+// Algorithm::create<DescriptorExtractor>("Feature2D.ORB"),
+// Algorithm::create<DescriptorExtractor>("Feature2D.ORB")->defaultNorm(),
+// 0.01f);
// test.safe_run();
//}
//TEST(Features2d_ScaleInvariance_Descriptor_FREAK, regression)
//{
// DescriptorScaleInvarianceTest test(Algorithm::create<FeatureDetector>("Feature2D.ORB"),
-// Algorithm::create<DescriptorExtractor>("Feature2D.FREAK"),
-// NORM_HAMMING,
-// 0.01f);
+// Algorithm::create<DescriptorExtractor>("Feature2D.FREAK"),
+// Algorithm::create<DescriptorExtractor>("Feature2D.FREAK")->defaultNorm(),
+// 0.01f);
// test.safe_run();
//}
--- /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-2013, Advanced Micro Devices, Inc., 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 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*/
+
+///////////////////////////////////////////////////////////////////////////////////////////////////
+/////////////////////////////////Macro for border type////////////////////////////////////////////
+/////////////////////////////////////////////////////////////////////////////////////////////////
+#ifdef BORDER_REPLICATE
+//BORDER_REPLICATE: aaaaaa|abcdefgh|hhhhhhh
+#define ADDR_L(i, l_edge, r_edge) ((i) < (l_edge) ? (l_edge) : (i))
+#define ADDR_R(i, r_edge, addr) ((i) >= (r_edge) ? (r_edge)-1 : (addr))
+#define ADDR_H(i, t_edge, b_edge) ((i) < (t_edge) ? (t_edge) :(i))
+#define ADDR_B(i, b_edge, addr) ((i) >= (b_edge) ? (b_edge)-1 :(addr))
+#endif
+
+#ifdef BORDER_REFLECT
+//BORDER_REFLECT: fedcba|abcdefgh|hgfedcb
+#define ADDR_L(i, l_edge, r_edge) ((i) < (l_edge) ? -(i)-1 : (i))
+#define ADDR_R(i, r_edge, addr) ((i) >= (r_edge) ? -(i)-1+((r_edge)<<1) : (addr))
+#define ADDR_H(i, t_edge, b_edge) ((i) < (t_edge) ? -(i)-1 : (i))
+#define ADDR_B(i, b_edge, addr) ((i) >= (b_edge) ? -(i)-1+((b_edge)<<1) : (addr))
+#endif
+
+#ifdef BORDER_REFLECT_101
+//BORDER_REFLECT_101: gfedcb|abcdefgh|gfedcba
+#define ADDR_L(i, l_edge, r_edge) ((i) < (l_edge) ? -(i) : (i))
+#define ADDR_R(i, r_edge, addr) ((i) >= (r_edge) ? -(i)-2+((r_edge)<<1) : (addr))
+#define ADDR_H(i, t_edge, b_edge) ((i) < (t_edge) ? -(i) : (i))
+#define ADDR_B(i, b_edge, addr) ((i) >= (b_edge) ? -(i)-2+((b_edge)<<1) : (addr))
+#endif
+
+//blur function does not support BORDER_WRAP
+#ifdef BORDER_WRAP
+//BORDER_WRAP: cdefgh|abcdefgh|abcdefg
+#define ADDR_L(i, l_edge, r_edge) ((i) < (l_edge) ? (i)+(r_edge) : (i))
+#define ADDR_R(i, r_edge, addr) ((i) >= (r_edge) ? (i)-(r_edge) : (addr))
+#define ADDR_H(i, t_edge, b_edge) ((i) < (t_edge) ? (i)+(b_edge) : (i))
+#define ADDR_B(i, b_edge, addr) ((i) >= (b_edge) ? (i)-(b_edge) : (addr))
+#endif
+
+#ifdef EXTRA_EXTRAPOLATION // border > src image size
+#ifdef BORDER_CONSTANT
+// None
+#elif defined BORDER_REPLICATE
+#define EXTRAPOLATE(x, y, minX, minY, maxX, maxY) \
+ { \
+ x = max(min(x, maxX - 1), minX); \
+ y = max(min(y, maxY - 1), minY); \
+ }
+#elif defined BORDER_WRAP
+#define EXTRAPOLATE(x, y, minX, minY, maxX, maxY) \
+ { \
+ if (x < minX) \
+ x -= ((x - maxX + 1) / maxX) * maxX; \
+ if (x >= maxX) \
+ x %= maxX; \
+ if (y < minY) \
+ y -= ((y - maxY + 1) / maxY) * maxY; \
+ if (y >= maxY) \
+ y %= maxY; \
+ }
+#elif defined(BORDER_REFLECT) || defined(BORDER_REFLECT_101)
+#define EXTRAPOLATE_(x, y, minX, minY, maxX, maxY, delta) \
+ { \
+ if (maxX - minX == 1) \
+ x = minX; \
+ else \
+ do \
+ { \
+ if (x < minX) \
+ x = minX - (x - minX) - 1 + delta; \
+ else \
+ x = maxX - 1 - (x - maxX) - delta; \
+ } \
+ while (x >= maxX || x < minX); \
+ \
+ if (maxY - minY == 1) \
+ y = minY; \
+ else \
+ do \
+ { \
+ if (y < minY) \
+ y = minY - (y - minY) - 1 + delta; \
+ else \
+ y = maxY - 1 - (y - maxY) - delta; \
+ } \
+ while (y >= maxY || y < minY); \
+ }
+#ifdef BORDER_REFLECT
+#define EXTRAPOLATE(x, y, minX, minY, maxX, maxY) EXTRAPOLATE_(x, y, minX, minY, maxX, maxY, 0)
+#elif defined(BORDER_REFLECT_101)
+#define EXTRAPOLATE(x, y, minX, minY, maxX, maxY) EXTRAPOLATE_(x, y, minX, minY, maxX, maxY, 1)
+#endif
+#else
+#error No extrapolation method
+#endif
+#else
+#define EXTRAPOLATE(x, y, minX, minY, maxX, maxY) \
+ { \
+ int _row = y - minY, _col = x - minX; \
+ _row = ADDR_H(_row, 0, maxY - minY); \
+ _row = ADDR_B(_row, maxY - minY, _row); \
+ y = _row + minY; \
+ \
+ _col = ADDR_L(_col, 0, maxX - minX); \
+ _col = ADDR_R(_col, maxX - minX, _col); \
+ x = _col + minX; \
+ }
+#endif
+
+#if USE_DOUBLE
+#ifdef cl_amd_fp64
+#pragma OPENCL EXTENSION cl_amd_fp64:enable
+#elif defined (cl_khr_fp64)
+#pragma OPENCL EXTENSION cl_khr_fp64:enable
+#endif
+#define FPTYPE double
+#define CONVERT_TO_FPTYPE CAT(convert_double, VEC_SIZE)
+#else
+#define FPTYPE float
+#define CONVERT_TO_FPTYPE CAT(convert_float, VEC_SIZE)
+#endif
+
+#if DATA_DEPTH == 0
+#define BASE_TYPE uchar
+#elif DATA_DEPTH == 1
+#define BASE_TYPE char
+#elif DATA_DEPTH == 2
+#define BASE_TYPE ushort
+#elif DATA_DEPTH == 3
+#define BASE_TYPE short
+#elif DATA_DEPTH == 4
+#define BASE_TYPE int
+#elif DATA_DEPTH == 5
+#define BASE_TYPE float
+#elif DATA_DEPTH == 6
+#define BASE_TYPE double
+#else
+#error data_depth
+#endif
+
+#define __CAT(x, y) x##y
+#define CAT(x, y) __CAT(x, y)
+
+#define uchar1 uchar
+#define char1 char
+#define ushort1 ushort
+#define short1 short
+#define int1 int
+#define float1 float
+#define double1 double
+
+#define convert_uchar1_sat_rte convert_uchar_sat_rte
+#define convert_char1_sat_rte convert_char_sat_rte
+#define convert_ushort1_sat_rte convert_ushort_sat_rte
+#define convert_short1_sat_rte convert_short_sat_rte
+#define convert_int1_sat_rte convert_int_sat_rte
+#define convert_float1
+#define convert_double1
+
+#if DATA_DEPTH == 5 || DATA_DEPTH == 6
+#define CONVERT_TO_TYPE CAT(CAT(convert_, BASE_TYPE), VEC_SIZE)
+#else
+#define CONVERT_TO_TYPE CAT(CAT(CAT(convert_, BASE_TYPE), VEC_SIZE), _sat_rte)
+#endif
+
+#define VEC_SIZE DATA_CHAN
+
+#define VEC_TYPE CAT(BASE_TYPE, VEC_SIZE)
+#define TYPE VEC_TYPE
+
+#define SCALAR_TYPE CAT(FPTYPE, VEC_SIZE)
+
+#define INTERMEDIATE_TYPE CAT(FPTYPE, VEC_SIZE)
+
+#define TYPE_SIZE (VEC_SIZE*sizeof(BASE_TYPE))
+
+struct RectCoords
+{
+ int x1, y1, x2, y2;
+};
+
+//#define DEBUG
+#ifdef DEBUG
+#define DEBUG_ONLY(x) x
+#define ASSERT(condition) do { if (!(condition)) { printf("BUG in boxFilter kernel (global=%d,%d): " #condition "\n", get_global_id(0), get_global_id(1)); } } while (0)
+#else
+#define DEBUG_ONLY(x)
+#define ASSERT(condition)
+#endif
+
+
+inline INTERMEDIATE_TYPE readSrcPixel(int2 pos, __global const uchar* srcptr, int srcstep, const struct RectCoords srcCoords
+#ifdef BORDER_CONSTANT
+ , SCALAR_TYPE borderValue
+#endif
+ )
+{
+#ifdef BORDER_ISOLATED
+ if(pos.x >= srcCoords.x1 && pos.y >= srcCoords.y1 && pos.x < srcCoords.x2 && pos.y < srcCoords.y2)
+#else
+ if(pos.x >= 0 && pos.y >= 0 && pos.x < srcCoords.x2 && pos.y < srcCoords.y2)
+#endif
+ {
+ __global TYPE* ptr = (__global TYPE*)(srcptr + pos.y * srcstep + pos.x * sizeof(TYPE));
+ return CONVERT_TO_FPTYPE(*ptr);
+ }
+ else
+ {
+#ifdef BORDER_CONSTANT
+ return borderValue;
+#else
+ int selected_col = pos.x;
+ int selected_row = pos.y;
+
+ EXTRAPOLATE(selected_col, selected_row,
+#ifdef BORDER_ISOLATED
+ srcCoords.x1, srcCoords.y1,
+#else
+ 0, 0,
+#endif
+ srcCoords.x2, srcCoords.y2
+ );
+
+ // debug border mapping
+ //printf("pos=%d,%d --> %d, %d\n", pos.x, pos.y, selected_col, selected_row);
+
+ pos = (int2)(selected_col, selected_row);
+ if(pos.x >= 0 && pos.y >= 0 && pos.x < srcCoords.x2 && pos.y < srcCoords.y2)
+ {
+ __global TYPE* ptr = (__global TYPE*)(srcptr + pos.y * srcstep + pos.x * sizeof(TYPE));
+ return CONVERT_TO_FPTYPE(*ptr);
+ }
+ else
+ {
+ // for debug only
+ DEBUG_ONLY(printf("BUG in boxFilter kernel\n"));
+ return (FPTYPE)(0.0f);
+ }
+#endif
+ }
+}
+
+// INPUT PARAMETER: BLOCK_SIZE_Y (via defines)
+
+__kernel
+__attribute__((reqd_work_group_size(LOCAL_SIZE, 1, 1)))
+void boxFilter(__global const uchar* srcptr, int srcstep, int srcOffsetX, int srcOffsetY, int srcEndX, int srcEndY,
+ __global uchar* dstptr, int dststep, int dstoffset,
+ int rows, int cols,
+#ifdef BORDER_CONSTANT
+ SCALAR_TYPE borderValue,
+#endif
+ FPTYPE alpha
+ )
+{
+ const struct RectCoords srcCoords = {srcOffsetX, srcOffsetY, srcEndX, srcEndY}; // for non-isolated border: offsetX, offsetY, wholeX, wholeY
+
+ const int x = get_local_id(0) + (LOCAL_SIZE - (KERNEL_SIZE_X - 1)) * get_group_id(0) - ANCHOR_X;
+ const int y = get_global_id(1) * BLOCK_SIZE_Y;
+
+ const int local_id = get_local_id(0);
+
+ INTERMEDIATE_TYPE data[KERNEL_SIZE_Y];
+ __local INTERMEDIATE_TYPE sumOfCols[LOCAL_SIZE];
+
+ int2 srcPos = (int2)(srcCoords.x1 + x, srcCoords.y1 + y - ANCHOR_Y);
+ for(int sy = 0; sy < KERNEL_SIZE_Y; sy++, srcPos.y++)
+ {
+ data[sy] = readSrcPixel(srcPos, srcptr, srcstep, srcCoords
+#ifdef BORDER_CONSTANT
+ , borderValue
+#endif
+ );
+ }
+
+ INTERMEDIATE_TYPE tmp_sum = 0;
+ for(int sy = 0; sy < KERNEL_SIZE_Y; sy++)
+ {
+ tmp_sum += (data[sy]);
+ }
+
+ sumOfCols[local_id] = tmp_sum;
+ barrier(CLK_LOCAL_MEM_FENCE);
+
+ int2 pos = (int2)(x, y);
+ __global TYPE* dstPtr = (__global TYPE*)(dstptr + pos.y * dststep + dstoffset + pos.x * TYPE_SIZE/*sizeof(TYPE)*/); // Pointer can be out of bounds!
+
+ int sy_index = 0; // current index in data[] array
+ int stepsY = min(rows - pos.y, BLOCK_SIZE_Y);
+ ASSERT(stepsY > 0);
+ for (; ;)
+ {
+ ASSERT(pos.y < rows);
+
+ if(local_id >= ANCHOR_X && local_id < LOCAL_SIZE - (KERNEL_SIZE_X - 1 - ANCHOR_X) &&
+ pos.x >= 0 && pos.x < cols)
+ {
+ ASSERT(pos.y >= 0 && pos.y < rows);
+
+ INTERMEDIATE_TYPE total_sum = 0;
+#pragma unroll
+ for (int sx = 0; sx < KERNEL_SIZE_X; sx++)
+ {
+ total_sum += sumOfCols[local_id + sx - ANCHOR_X];
+ }
+ *dstPtr = CONVERT_TO_TYPE(((INTERMEDIATE_TYPE)alpha) * total_sum);
+ }
+
+#if BLOCK_SIZE_Y == 1
+ break;
+#else
+ if (--stepsY == 0)
+ break;
+
+ barrier(CLK_LOCAL_MEM_FENCE);
+
+ tmp_sum = sumOfCols[local_id]; // TODO FIX IT: workaround for BUG in OpenCL compiler
+ // only works with scalars: ASSERT(fabs(tmp_sum - sumOfCols[local_id]) < (INTERMEDIATE_TYPE)1e-6);
+ tmp_sum -= data[sy_index];
+
+ data[sy_index] = readSrcPixel(srcPos, srcptr, srcstep, srcCoords
+#ifdef BORDER_CONSTANT
+ , borderValue
+#endif
+ );
+ srcPos.y++;
+
+ tmp_sum += data[sy_index];
+ sumOfCols[local_id] = tmp_sum;
+
+ sy_index = (sy_index + 1 < KERNEL_SIZE_Y) ? sy_index + 1 : 0;
+
+ barrier(CLK_LOCAL_MEM_FENCE);
+
+ // next line
+ DEBUG_ONLY(pos.y++);
+ dstPtr = (__global TYPE*)((__global char*)dstPtr + dststep); // Pointer can be out of bounds!
+#endif // BLOCK_SIZE_Y == 1
+ }
+}
--- /dev/null
+// 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.
+// 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 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.
+
+#define DATA_TYPE type
+
+#define scnbytes ((int)sizeof(type))
+
+#define op(a,b) { mid=a; a=min(a,b); b=max(mid,b);}
+
+__kernel void medianFilter3(__global const uchar* srcptr, int srcStep, int srcOffset,
+ __global uchar* dstptr, int dstStep, int dstOffset,
+ int rows, int cols)
+{
+ __local DATA_TYPE data[18][18];
+
+ int x = get_local_id(0);
+ int y = get_local_id(1);
+
+ int gx= get_global_id(0);
+ int gy= get_global_id(1);
+
+ int dx = gx - x - 1;
+ int dy = gy - y - 1;
+
+ const int id = min((int)(x*16+y), 9*18-1);
+
+ int dr = id / 18;
+ int dc = id % 18;
+
+ int c = clamp(dx+dc, 0, cols-1);
+
+ int r = clamp(dy+dr, 0, rows-1);
+ int index1 = mad24(r, srcStep, srcOffset + c*scnbytes);
+
+ r = clamp(dy+dr+9, 0, rows-1);
+ int index9 = mad24(r, srcStep, srcOffset + c*scnbytes);
+
+ __global DATA_TYPE * src = (__global DATA_TYPE *)(srcptr + index1);
+ data[dr][dc] = src[0];
+
+ src = (__global DATA_TYPE *)(srcptr + index9);
+ data[dr+9][dc] = src[0];
+
+ barrier(CLK_LOCAL_MEM_FENCE);
+
+ DATA_TYPE p0=data[y][x], p1=data[y][(x+1)], p2=data[y][(x+2)];
+ DATA_TYPE p3=data[y+1][x], p4=data[y+1][(x+1)], p5=data[y+1][(x+2)];
+ DATA_TYPE p6=data[y+2][x], p7=data[y+2][(x+1)], p8=data[y+2][(x+2)];
+ DATA_TYPE mid;
+
+ op(p1, p2); op(p4, p5); op(p7, p8); op(p0, p1);
+ op(p3, p4); op(p6, p7); op(p1, p2); op(p4, p5);
+ op(p7, p8); op(p0, p3); op(p5, p8); op(p4, p7);
+ op(p3, p6); op(p1, p4); op(p2, p5); op(p4, p7);
+ op(p4, p2); op(p6, p4); op(p4, p2);
+
+ int dst_index = mad24( gy, dstStep, dstOffset + gx * scnbytes);
+
+ if( gy < rows && gx < cols)
+ {
+ __global DATA_TYPE* dst = (__global DATA_TYPE *)(dstptr + dst_index);
+ dst[0] = p4;
+ }
+}
+
+__kernel void medianFilter5(__global const uchar* srcptr, int srcStep, int srcOffset,
+ __global uchar* dstptr, int dstStep, int dstOffset,
+ int rows, int cols)
+{
+ __local DATA_TYPE data[20][20];
+
+ int x =get_local_id(0);
+ int y =get_local_id(1);
+
+ int gx=get_global_id(0);
+ int gy=get_global_id(1);
+
+ int dx = gx - x - 2;
+ int dy = gy - y - 2;
+
+ const int id = min((int)(x*16+y), 10*20-1);
+
+ int dr=id/20;
+ int dc=id%20;
+
+ int c=clamp(dx+dc, 0, cols-1);
+
+ int r = clamp(dy+dr, 0, rows-1);
+ int index1 = mad24(r, srcStep, srcOffset + c*scnbytes);
+
+ r = clamp(dy+dr+10, 0, rows-1);
+ int index10 = mad24(r, srcStep, srcOffset + c*scnbytes);
+
+ __global DATA_TYPE * src = (__global DATA_TYPE *)(srcptr + index1);
+ data[dr][dc] = src[0];
+ src = (__global DATA_TYPE *)(srcptr + index10);
+ data[dr+10][dc] = src[0];
+
+ barrier(CLK_LOCAL_MEM_FENCE);
+
+ DATA_TYPE p0=data[y][x], p1=data[y][x+1], p2=data[y][x+2], p3=data[y][x+3], p4=data[y][x+4];
+ DATA_TYPE p5=data[y+1][x], p6=data[y+1][x+1], p7=data[y+1][x+2], p8=data[y+1][x+3], p9=data[y+1][x+4];
+ DATA_TYPE p10=data[y+2][x], p11=data[y+2][x+1], p12=data[y+2][x+2], p13=data[y+2][x+3], p14=data[y+2][x+4];
+ DATA_TYPE p15=data[y+3][x], p16=data[y+3][x+1], p17=data[y+3][x+2], p18=data[y+3][x+3], p19=data[y+3][x+4];
+ DATA_TYPE p20=data[y+4][x], p21=data[y+4][x+1], p22=data[y+4][x+2], p23=data[y+4][x+3], p24=data[y+4][x+4];
+ DATA_TYPE mid;
+
+ op(p1, p2); op(p0, p1); op(p1, p2); op(p4, p5); op(p3, p4);
+ op(p4, p5); op(p0, p3); op(p2, p5); op(p2, p3); op(p1, p4);
+ op(p1, p2); op(p3, p4); op(p7, p8); op(p6, p7); op(p7, p8);
+ op(p10, p11); op(p9, p10); op(p10, p11); op(p6, p9); op(p8, p11);
+ op(p8, p9); op(p7, p10); op(p7, p8); op(p9, p10); op(p0, p6);
+ op(p4, p10); op(p4, p6); op(p2, p8); op(p2, p4); op(p6, p8);
+ op(p1, p7); op(p5, p11); op(p5, p7); op(p3, p9); op(p3, p5);
+ op(p7, p9); op(p1, p2); op(p3, p4); op(p5, p6); op(p7, p8);
+ op(p9, p10); op(p13, p14); op(p12, p13); op(p13, p14); op(p16, p17);
+ op(p15, p16); op(p16, p17); op(p12, p15); op(p14, p17); op(p14, p15);
+ op(p13, p16); op(p13, p14); op(p15, p16); op(p19, p20); op(p18, p19);
+ op(p19, p20); op(p21, p22); op(p23, p24); op(p21, p23); op(p22, p24);
+ op(p22, p23); op(p18, p21); op(p20, p23); op(p20, p21); op(p19, p22);
+ op(p22, p24); op(p19, p20); op(p21, p22); op(p23, p24); op(p12, p18);
+ op(p16, p22); op(p16, p18); op(p14, p20); op(p20, p24); op(p14, p16);
+ op(p18, p20); op(p22, p24); op(p13, p19); op(p17, p23); op(p17, p19);
+ op(p15, p21); op(p15, p17); op(p19, p21); op(p13, p14); op(p15, p16);
+ op(p17, p18); op(p19, p20); op(p21, p22); op(p23, p24); op(p0, p12);
+ op(p8, p20); op(p8, p12); op(p4, p16); op(p16, p24); op(p12, p16);
+ op(p2, p14); op(p10, p22); op(p10, p14); op(p6, p18); op(p6, p10);
+ op(p10, p12); op(p1, p13); op(p9, p21); op(p9, p13); op(p5, p17);
+ op(p13, p17); op(p3, p15); op(p11, p23); op(p11, p15); op(p7, p19);
+ op(p7, p11); op(p11, p13); op(p11, p12);
+
+ int dst_index = mad24( gy, dstStep, dstOffset + gx * scnbytes);
+
+ if( gy < rows && gx < cols)
+ {
+ __global DATA_TYPE* dst = (__global DATA_TYPE *)(dstptr + dst_index);
+ dst[0] = p12;
+ }
+}
\ No newline at end of file
std::vector<int> sum;
};
+#define DIVUP(total, grain) ((total + grain - 1) / (grain))
+
+static bool ocl_boxFilter( InputArray _src, OutputArray _dst, int ddepth,
+ Size ksize, Point anchor, int borderType )
+{
+ int type = _src.type();
+ int cn = CV_MAT_CN(type);
+ if ((1 != cn) && (2 != cn) && (4 != cn))
+ return false;//TODO
+
+ int sdepth = CV_MAT_DEPTH(type);
+ if( ddepth < 0 )
+ ddepth = sdepth;
+ else if (ddepth != sdepth)
+ return false;
+ if( anchor.x < 0 )
+ anchor.x = ksize.width / 2;
+ if( anchor.y < 0 )
+ anchor.y = ksize.height / 2;
+
+ ocl::Kernel kernel;
+
+ //Normalize the result by default
+ float alpha = 1.0f / (ksize.height * ksize.width);
+ bool isIsolatedBorder = (borderType & BORDER_ISOLATED) != 0;
+ bool useDouble = (CV_64F == sdepth);
+ const cv::ocl::Device &device = cv::ocl::Device::getDefault();
+ int doubleFPConfig = device.doubleFPConfig();
+ if (useDouble && (0 == doubleFPConfig))
+ return false;// may be we have to check is (0 != (CL_FP_SOFT_FLOAT & doubleFPConfig)) ?
+
+ const char* btype = NULL;
+ switch (borderType & ~BORDER_ISOLATED)
+ {
+ case BORDER_CONSTANT:
+ btype = "BORDER_CONSTANT";
+ break;
+ case BORDER_REPLICATE:
+ btype = "BORDER_REPLICATE";
+ break;
+ case BORDER_REFLECT:
+ btype = "BORDER_REFLECT";
+ break;
+ case BORDER_WRAP:
+ //CV_Error(CV_StsUnsupportedFormat, "BORDER_WRAP is not supported!");
+ return false;
+ case BORDER_REFLECT101:
+ btype = "BORDER_REFLECT_101";
+ break;
+ }
+
+ cv::Size sz = _src.size();
+
+ size_t globalsize[2] = {sz.width, sz.height};
+ size_t localsize[2] = {0, 1};
+
+ UMat src; Size wholeSize;
+ if (!isIsolatedBorder)
+ {
+ src = _src.getUMat();
+ Point ofs;
+ src.locateROI(wholeSize, ofs);
+ }
+
+ size_t maxWorkItemSizes[32]; device.maxWorkItemSizes(maxWorkItemSizes);
+ size_t tryWorkItems = maxWorkItemSizes[0];
+ for (;;)
+ {
+ size_t BLOCK_SIZE = tryWorkItems;
+ while (BLOCK_SIZE > 32 && BLOCK_SIZE >= (size_t)ksize.width * 2 && BLOCK_SIZE > (size_t)sz.width * 2)
+ BLOCK_SIZE /= 2;
+ size_t BLOCK_SIZE_Y = 8; // TODO Check heuristic value on devices
+ while (BLOCK_SIZE_Y < BLOCK_SIZE / 8 && BLOCK_SIZE_Y * device.maxComputeUnits() * 32 < (size_t)sz.height)
+ BLOCK_SIZE_Y *= 2;
+
+ if ((size_t)ksize.width > BLOCK_SIZE)
+ return false;
+
+ int requiredTop = anchor.y;
+ int requiredLeft = (int)BLOCK_SIZE; // not this: anchor.x;
+ int requiredBottom = ksize.height - 1 - anchor.y;
+ int requiredRight = (int)BLOCK_SIZE; // not this: ksize.width - 1 - anchor.x;
+ int h = isIsolatedBorder ? sz.height : wholeSize.height;
+ int w = isIsolatedBorder ? sz.width : wholeSize.width;
+
+ bool extra_extrapolation = h < requiredTop || h < requiredBottom || w < requiredLeft || w < requiredRight;
+
+ if ((w < ksize.width) || (h < ksize.height))
+ return false;
+
+ char build_options[1024];
+ sprintf(build_options, "-D LOCAL_SIZE=%d -D BLOCK_SIZE_Y=%d -D DATA_DEPTH=%d -D DATA_CHAN=%d -D USE_DOUBLE=%d -D ANCHOR_X=%d -D ANCHOR_Y=%d -D KERNEL_SIZE_X=%d -D KERNEL_SIZE_Y=%d -D %s -D %s -D %s",
+ (int)BLOCK_SIZE, (int)BLOCK_SIZE_Y,
+ sdepth, cn, useDouble ? 1 : 0,
+ anchor.x, anchor.y, ksize.width, ksize.height,
+ btype,
+ extra_extrapolation ? "EXTRA_EXTRAPOLATION" : "NO_EXTRA_EXTRAPOLATION",
+ isIsolatedBorder ? "BORDER_ISOLATED" : "NO_BORDER_ISOLATED");
+
+ localsize[0] = BLOCK_SIZE;
+ globalsize[0] = DIVUP(sz.width, BLOCK_SIZE - (ksize.width - 1)) * BLOCK_SIZE;
+ globalsize[1] = DIVUP(sz.height, BLOCK_SIZE_Y);
+
+ cv::String errmsg;
+ kernel.create("boxFilter", cv::ocl::imgproc::boxFilter_oclsrc, build_options);
+
+ size_t kernelWorkGroupSize = kernel.workGroupSize();
+ if (localsize[0] <= kernelWorkGroupSize)
+ break;
+
+ if (BLOCK_SIZE < kernelWorkGroupSize)
+ return false;
+ tryWorkItems = kernelWorkGroupSize;
+ }
+
+ _dst.create(sz, CV_MAKETYPE(ddepth, cn));
+ UMat dst = _dst.getUMat();
+ if (src.empty())
+ src = _src.getUMat();
+ int idxArg = 0;
+ idxArg = kernel.set(idxArg, ocl::KernelArg::PtrReadOnly(src));
+ idxArg = kernel.set(idxArg, (int)src.step);
+ int srcOffsetX = (int)((src.offset % src.step) / src.elemSize());
+ int srcOffsetY = (int)(src.offset / src.step);
+ int srcEndX = (isIsolatedBorder ? (srcOffsetX + sz.width) : wholeSize.width);
+ int srcEndY = (isIsolatedBorder ? (srcOffsetY + sz.height) : wholeSize.height);
+ idxArg = kernel.set(idxArg, srcOffsetX);
+ idxArg = kernel.set(idxArg, srcOffsetY);
+ idxArg = kernel.set(idxArg, srcEndX);
+ idxArg = kernel.set(idxArg, srcEndY);
+ idxArg = kernel.set(idxArg, ocl::KernelArg::WriteOnly(dst));
+ float borderValue[4] = {0, 0, 0, 0};
+ double borderValueDouble[4] = {0, 0, 0, 0};
+ if ((borderType & ~BORDER_ISOLATED) == BORDER_CONSTANT)
+ {
+ int cnocl = (3 == cn) ? 4 : cn;
+ if (useDouble)
+ idxArg = kernel.set(idxArg, (void *)&borderValueDouble[0], sizeof(double) * cnocl);
+ else
+ idxArg = kernel.set(idxArg, (void *)&borderValue[0], sizeof(float) * cnocl);
+ }
+ if (useDouble)
+ idxArg = kernel.set(idxArg, (double)alpha);
+ else
+ idxArg = kernel.set(idxArg, (float)alpha);
+
+ return kernel.run(2, globalsize, localsize, true);
+}
}
+
cv::Ptr<cv::BaseRowFilter> cv::getRowSumFilter(int srcType, int sumType, int ksize, int anchor)
{
int sdepth = CV_MAT_DEPTH(srcType), ddepth = CV_MAT_DEPTH(sumType);
Size ksize, Point anchor,
bool normalize, int borderType )
{
+ bool use_opencl = ocl::useOpenCL() && _dst.isUMat() && normalize;
+ if( use_opencl && ocl_boxFilter(_src, _dst, ddepth, ksize, anchor, borderType) )
+ return;
+
Mat src = _src.getMat();
int sdepth = src.depth(), cn = src.channels();
if( ddepth < 0 )
}
+namespace cv
+{
+ static bool ocl_medianFilter ( InputArray _src, OutputArray _dst, int m)
+ {
+ int type = _src.type();
+ int depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type);
+
+ if (!((depth == CV_8U || depth == CV_16U || depth == CV_16S || depth == CV_32F) && (cn != 3 && cn <= 4)))
+ return false;
+
+ const char * kernelName;
+
+ if (m==3)
+ kernelName = "medianFilter3";
+ else if (m==5)
+ kernelName = "medianFilter5";
+ else
+ return false;
+
+ ocl::Kernel k(kernelName,ocl::imgproc::medianFilter_oclsrc,format("-D type=%s",ocl::typeToStr(type)));
+ if (k.empty())
+ return false;
+
+ _dst.create(_src.size(),type);
+ UMat src = _src.getUMat(), dst = _dst.getUMat();
+
+ size_t globalsize[2] = {(src.cols + 18) / 16 * 16, (src.rows + 15) / 16 * 16};
+ size_t localsize[2] = {16, 16};
+
+ return k.args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::WriteOnly(dst)).run(2,globalsize,localsize,false);
+ }
+}
+
void cv::medianBlur( InputArray _src0, OutputArray _dst, int ksize )
{
- Mat src0 = _src0.getMat();
- _dst.create( src0.size(), src0.type() );
- Mat dst = _dst.getMat();
+ CV_Assert( (ksize % 2 == 1) && (_src0.dims() <= 2 ));
if( ksize <= 1 )
{
+ Mat src0 = _src0.getMat();
+ _dst.create( src0.size(), src0.type() );
+ Mat dst = _dst.getMat();
src0.copyTo(dst);
return;
}
- CV_Assert( ksize % 2 == 1 );
+ bool use_opencl = ocl::useOpenCL() && _dst.isUMat();
+ if ( use_opencl && ocl_medianFilter(_src0,_dst, ksize))
+ return;
+
+ Mat src0 = _src0.getMat();
+ _dst.create( src0.size(), src0.type() );
+ Mat dst = _dst.getMat();
#ifdef HAVE_TEGRA_OPTIMIZATION
if (tegra::medianBlur(src0, dst, ksize))
--- /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.
+//
+// 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"
+
+#ifdef HAVE_OPENCL
+
+namespace cvtest {
+namespace ocl {
+
+enum
+{
+ noType = -1
+};
+
+/////////////////////////////////////////////////////////////////////////////////////////////////
+// boxFilter
+PARAM_TEST_CASE(BoxFilter, MatDepth, Channels, BorderType, bool)
+{
+ static const int kernelMinSize = 2;
+ static const int kernelMaxSize = 10;
+
+ int type;
+ Size ksize;
+ Size dsize;
+ Point anchor;
+ int borderType;
+ bool useRoi;
+
+ TEST_DECLARE_INPUT_PARAMETER(src)
+ TEST_DECLARE_OUTPUT_PARAMETER(dst)
+
+ virtual void SetUp()
+ {
+ type = CV_MAKE_TYPE(GET_PARAM(0), GET_PARAM(1));
+ borderType = GET_PARAM(2); // only not isolated border tested, because CPU module doesn't support isolated border case.
+ useRoi = GET_PARAM(3);
+ }
+
+ void random_roi()
+ {
+ dsize = randomSize(1, MAX_VALUE);
+
+ ksize = randomSize(kernelMinSize, kernelMaxSize);
+
+ Size roiSize = randomSize(ksize.width, MAX_VALUE, ksize.height, MAX_VALUE);
+ Border srcBorder = randomBorder(0, useRoi ? MAX_VALUE : 0);
+ randomSubMat(src, src_roi, roiSize, srcBorder, type, -MAX_VALUE, MAX_VALUE);
+
+ Border dstBorder = randomBorder(0, useRoi ? MAX_VALUE : 0);
+ randomSubMat(dst, dst_roi, dsize, dstBorder, type, -MAX_VALUE, MAX_VALUE);
+
+ anchor.x = randomInt(-1, ksize.width);
+ anchor.y = randomInt(-1, ksize.height);
+
+ UMAT_UPLOAD_INPUT_PARAMETER(src)
+ UMAT_UPLOAD_OUTPUT_PARAMETER(dst)
+ }
+
+ void Near(double threshold = 0.0)
+ {
+ EXPECT_MAT_NEAR(dst, udst, threshold);
+ EXPECT_MAT_NEAR(dst_roi, udst_roi, threshold);
+ }
+};
+
+OCL_TEST_P(BoxFilter, Mat)
+{
+ for (int j = 0; j < test_loop_times; j++)
+ {
+ random_roi();
+
+ OCL_OFF(cv::boxFilter(src_roi, dst_roi, -1, ksize, anchor, true, borderType));
+ OCL_ON(cv::boxFilter(usrc_roi, udst_roi, -1, ksize, anchor, true, borderType));
+
+ Near(1.0);
+ }
+}
+
+
+OCL_INSTANTIATE_TEST_CASE_P(ImageProc, BoxFilter,
+ Combine(
+ Values(CV_8U, CV_16U, CV_16S, CV_32S, CV_32F),
+ Values(1, 2, 4),
+ Values((BorderType)BORDER_CONSTANT,
+ (BorderType)BORDER_REPLICATE,
+ (BorderType)BORDER_REFLECT,
+ (BorderType)BORDER_REFLECT_101),
+ Bool() // ROI
+ )
+ );
+
+
+} } // 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, Multicoreware, Inc., all rights reserved.
+// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., 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 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"
+
+#ifdef HAVE_OPENCL
+
+namespace cvtest {
+namespace ocl {
+
+/////////////////////////////////////////////medianFilter//////////////////////////////////////////////////////////
+
+PARAM_TEST_CASE(MedianFilter, MatDepth, Channels, int, bool)
+{
+ int type;
+ int ksize;
+ bool use_roi;
+
+ TEST_DECLARE_INPUT_PARAMETER(src)
+ TEST_DECLARE_OUTPUT_PARAMETER(dst)
+
+ virtual void SetUp()
+ {
+ type = CV_MAKE_TYPE(GET_PARAM(0), GET_PARAM(1));
+ ksize = GET_PARAM(2);
+ use_roi = GET_PARAM(3);
+ }
+
+ virtual void generateTestData()
+ {
+ Size roiSize = randomSize(1, MAX_VALUE);
+ Border srcBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
+ randomSubMat(src, src_roi, roiSize, srcBorder, type, -MAX_VALUE, MAX_VALUE);
+
+ Border dstBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
+ randomSubMat(dst, dst_roi, roiSize, dstBorder, type, -MAX_VALUE, MAX_VALUE);
+
+ UMAT_UPLOAD_INPUT_PARAMETER(src)
+ UMAT_UPLOAD_OUTPUT_PARAMETER(dst)
+ }
+
+ void Near(double threshold = 0.0)
+ {
+ EXPECT_MAT_NEAR(dst, udst, threshold);
+ EXPECT_MAT_NEAR(dst_roi, udst_roi, threshold);
+ }
+};
+
+OCL_TEST_P(MedianFilter, Mat)
+{
+ for (int j = 0; j < test_loop_times; j++)
+ {
+ generateTestData();
+
+ OCL_OFF(cv::medianBlur(src_roi, dst_roi, ksize));
+ OCL_ON(cv::medianBlur(usrc_roi, udst_roi, ksize));
+
+ Near(0);
+ }
+}
+
+OCL_INSTANTIATE_TEST_CASE_P(ImageProc, MedianFilter, Combine(
+ Values(CV_8U, CV_16U, CV_16S, CV_32F),
+ Values(1, 2, 4),
+ Values(3, 5),
+ Bool())
+ );
+} } // namespace cvtest::ocl
+
+#endif
virtual void write( FileStorage &fs ) const;
virtual int descriptorSize() const;
virtual int descriptorType() const;
+ virtual int defaultNorm() const;
protected:
...
}
virtual int descriptorSize() const { return classifier_.classes(); }
virtual int descriptorType() const { return DataType<T>::type; }
+ virtual int defaultNorm() const { return NORM_L1; }
virtual bool empty() const;
//! returns the descriptor size in float's (64 or 128)
int descriptorSize() const;
+ //! returns the default norm type
+ int defaultNorm() const;
//! upload host keypoints to device memory
void uploadKeypoints(const std::vector<KeyPoint>& keypoints, GpuMat& keypointsGPU);
//! returns the descriptor type
CV_WRAP int descriptorType() const;
+ //! returns the default norm type
+ CV_WRAP int defaultNorm() const;
+
//! finds the keypoints using SIFT algorithm
void operator()(InputArray img, InputArray mask,
std::vector<KeyPoint>& keypoints) const;
//! returns the descriptor type
CV_WRAP int descriptorType() const;
+ //! returns the descriptor type
+ CV_WRAP int defaultNorm() const;
+
//! finds the keypoints using fast hessian detector used in SURF
void operator()(InputArray img, InputArray mask,
CV_OUT std::vector<KeyPoint>& keypoints) const;
//! returns the descriptor size in float's (64 or 128)
int descriptorSize() const;
+ //! returns the default norm type
+ int defaultNorm() const;
//! upload host keypoints to device memory
void uploadKeypoints(const std::vector<cv::KeyPoint> &keypoints, oclMat &keypointsocl);
//! download keypoints from device to host memory
return CV_32F;
}
+int SIFT::defaultNorm() const
+{
+ return NORM_L2;
+}
+
void SIFT::operator()(InputArray _image, InputArray _mask,
std::vector<KeyPoint>& keypoints) const
int SURF::descriptorSize() const { return extended ? 128 : 64; }
int SURF::descriptorType() const { return CV_32F; }
+int SURF::defaultNorm() const { return NORM_L2; }
void SURF::operator()(InputArray imgarg, InputArray maskarg,
CV_OUT std::vector<KeyPoint>& keypoints) const
return extended ? 128 : 64;
}
+int cv::cuda::SURF_CUDA::defaultNorm() const
+{
+ return NORM_L2;
+}
+
void cv::cuda::SURF_CUDA::uploadKeypoints(const std::vector<KeyPoint>& keypoints, GpuMat& keypointsGPU)
{
if (keypoints.empty())
return extended ? 128 : 64;
}
+int cv::ocl::SURF_OCL::defaultNorm() const
+{
+ return NORM_L2;
+}
+
void cv::ocl::SURF_OCL::uploadKeypoints(const std::vector<KeyPoint> &keypoints, oclMat &keypointsGPU)
{
if (keypoints.empty())
CV_Assert(kpt2[i].response > 0 );
vector<DMatch> matches;
- BFMatcher(NORM_L2, true).match(d1, d2, matches);
+ BFMatcher(f->defaultNorm(), true).match(d1, d2, matches);
vector<Point2f> pt1, pt2;
for( size_t i = 0; i < matches.size(); i++ ) {
cv::Mat descriptors_gold;
surf_gold(image, cv::noArray(), keypoints, descriptors_gold, true);
- cv::BFMatcher matcher(cv::NORM_L2);
+ cv::BFMatcher matcher(surf.defaultNorm());
std::vector<cv::DMatch> matches;
matcher.match(descriptors_gold, cv::Mat(descriptors), matches);
cv::Mat descriptors_gold;
surf_gold(image, cv::noArray(), keypoints, descriptors_gold, true);
- cv::BFMatcher matcher(cv::NORM_L2);
+ cv::BFMatcher matcher(surf.defaultNorm());
std::vector<cv::DMatch> matches;
matcher.match(descriptors_gold, cv::Mat(descriptors), matches);
set(the_description "OpenCL-accelerated Computer Vision")
ocv_define_module(ocl opencv_core opencv_imgproc opencv_features2d opencv_objdetect opencv_video opencv_calib3d opencv_ml "${OPENCL_LIBRARIES}")
+if(TARGET opencv_test_ocl)
+ target_link_libraries(opencv_test_ocl "${OPENCL_LIBRARIES}")
+endif()
ocv_warnings_disable(CMAKE_CXX_FLAGS -Wshadow)
ocl::setDevice
--------------
-Returns void
+Initialize OpenCL computation context
.. ocv:function:: void ocl::setDevice( const DeviceInfo* info )
:param info: device info
+ocl::initializeContext
+--------------------------------
+Alternative way to initialize OpenCL computation context.
+
+.. ocv:function:: void ocl::initializeContext(void* pClPlatform, void* pClContext, void* pClDevice)
+
+ :param pClPlatform: selected ``platform_id`` (via pointer, parameter type is ``cl_platform_id*``)
+
+ :param pClContext: selected ``cl_context`` (via pointer, parameter type is ``cl_context*``)
+
+ :param pClDevice: selected ``cl_device_id`` (via pointer, parameter type is ``cl_device_id*``)
+
+This function can be used for context initialization with D3D/OpenGL interoperability.
+
ocl::setBinaryPath
------------------
Returns void
const PlatformInfo* platform;
DeviceInfo();
+ ~DeviceInfo();
};
struct PlatformInfo
std::vector<const DeviceInfo*> devices;
PlatformInfo();
+ ~PlatformInfo();
};
//////////////////////////////// Initialization & Info ////////////////////////
// set device you want to use
CV_EXPORTS void setDevice(const DeviceInfo* info);
+ // Initialize from OpenCL handles directly.
+ // Argument types is (pointers): cl_platform_id*, cl_context*, cl_device_id*
+ CV_EXPORTS void initializeContext(void* pClPlatform, void* pClContext, void* pClDevice);
+
enum FEATURE_TYPE
{
FEATURE_CL_DOUBLE = 1,
data = m.data;
datastart = m.datastart;
dataend = m.dataend;
+ clCxt = m.clCxt;
wholerows = m.wholerows;
wholecols = m.wholecols;
offset = m.offset;
namespace cv {
namespace ocl {
+using namespace cl_utils;
+
+#if defined(WIN32)
+static bool __termination = false;
+#endif
+
struct __Module
{
__Module();
return __module.initializationMutex;
}
-
-struct PlatformInfoImpl
-{
- cl_platform_id platform_id;
-
- std::vector<int> deviceIDs;
-
- PlatformInfo info;
-
- PlatformInfoImpl()
- : platform_id(NULL)
- {
- }
-};
-
-struct DeviceInfoImpl
+static cv::Mutex& getCurrentContextMutex()
{
- cl_platform_id platform_id;
- cl_device_id device_id;
-
- DeviceInfo info;
-
- DeviceInfoImpl()
- : platform_id(NULL), device_id(NULL)
- {
- }
-};
-
-static std::vector<PlatformInfoImpl> global_platforms;
-static std::vector<DeviceInfoImpl> global_devices;
+ return __module.currentContextMutex;
+}
static bool parseOpenCLVersion(const std::string& versionStr, int& major, int& minor)
{
return true;
}
+struct PlatformInfoImpl : public PlatformInfo
+{
+ cl_platform_id platform_id;
+
+ std::vector<int> deviceIDs;
+
+ PlatformInfoImpl()
+ : platform_id(NULL)
+ {
+ }
+
+ void init(int id, cl_platform_id platform)
+ {
+ CV_Assert(platform_id == NULL);
+
+ this->_id = id;
+ platform_id = platform;
+
+ openCLSafeCall(getStringInfo(clGetPlatformInfo, platform, CL_PLATFORM_PROFILE, this->platformProfile));
+ openCLSafeCall(getStringInfo(clGetPlatformInfo, platform, CL_PLATFORM_VERSION, this->platformVersion));
+ openCLSafeCall(getStringInfo(clGetPlatformInfo, platform, CL_PLATFORM_NAME, this->platformName));
+ openCLSafeCall(getStringInfo(clGetPlatformInfo, platform, CL_PLATFORM_VENDOR, this->platformVendor));
+ openCLSafeCall(getStringInfo(clGetPlatformInfo, platform, CL_PLATFORM_EXTENSIONS, this->platformExtensons));
+
+ parseOpenCLVersion(this->platformVersion,
+ this->platformVersionMajor, this->platformVersionMinor);
+ }
+
+};
+
+struct DeviceInfoImpl: public DeviceInfo
+{
+ cl_platform_id platform_id;
+ cl_device_id device_id;
+
+ DeviceInfoImpl()
+ : platform_id(NULL), device_id(NULL)
+ {
+ }
+
+ void init(int id, PlatformInfoImpl& platformInfoImpl, cl_device_id device)
+ {
+ CV_Assert(device_id == NULL);
+
+ this->_id = id;
+ platform_id = platformInfoImpl.platform_id;
+ device_id = device;
+
+ this->platform = &platformInfoImpl;
+
+ cl_device_type type = cl_device_type(-1);
+ openCLSafeCall(getScalarInfo(clGetDeviceInfo, device, CL_DEVICE_TYPE, type));
+ this->deviceType = DeviceType(type);
+
+ openCLSafeCall(getStringInfo(clGetDeviceInfo, device, CL_DEVICE_PROFILE, this->deviceProfile));
+ openCLSafeCall(getStringInfo(clGetDeviceInfo, device, CL_DEVICE_VERSION, this->deviceVersion));
+ openCLSafeCall(getStringInfo(clGetDeviceInfo, device, CL_DEVICE_NAME, this->deviceName));
+ openCLSafeCall(getStringInfo(clGetDeviceInfo, device, CL_DEVICE_VENDOR, this->deviceVendor));
+ cl_uint vendorID = 0;
+ openCLSafeCall(getScalarInfo(clGetDeviceInfo, device, CL_DEVICE_VENDOR_ID, vendorID));
+ this->deviceVendorId = vendorID;
+ openCLSafeCall(getStringInfo(clGetDeviceInfo, device, CL_DRIVER_VERSION, this->deviceDriverVersion));
+ openCLSafeCall(getStringInfo(clGetDeviceInfo, device, CL_DEVICE_EXTENSIONS, this->deviceExtensions));
+
+ parseOpenCLVersion(this->deviceVersion,
+ this->deviceVersionMajor, this->deviceVersionMinor);
+
+ size_t maxWorkGroupSize = 0;
+ openCLSafeCall(getScalarInfo(clGetDeviceInfo, device, CL_DEVICE_MAX_WORK_GROUP_SIZE, maxWorkGroupSize));
+ this->maxWorkGroupSize = maxWorkGroupSize;
+
+ cl_uint maxDimensions = 0;
+ openCLSafeCall(getScalarInfo(clGetDeviceInfo, device, CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS, maxDimensions));
+ std::vector<size_t> maxWorkItemSizes(maxDimensions);
+ openCLSafeCall(clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_ITEM_SIZES, sizeof(size_t) * maxDimensions,
+ (void *)&maxWorkItemSizes[0], 0));
+ this->maxWorkItemSizes = maxWorkItemSizes;
+
+ cl_uint maxComputeUnits = 0;
+ openCLSafeCall(getScalarInfo(clGetDeviceInfo, device, CL_DEVICE_MAX_COMPUTE_UNITS, maxComputeUnits));
+ this->maxComputeUnits = maxComputeUnits;
+
+ cl_ulong localMemorySize = 0;
+ openCLSafeCall(getScalarInfo(clGetDeviceInfo, device, CL_DEVICE_LOCAL_MEM_SIZE, localMemorySize));
+ this->localMemorySize = (size_t)localMemorySize;
+
+ cl_ulong maxMemAllocSize = 0;
+ openCLSafeCall(getScalarInfo(clGetDeviceInfo, device, CL_DEVICE_MAX_MEM_ALLOC_SIZE, maxMemAllocSize));
+ this->maxMemAllocSize = (size_t)maxMemAllocSize;
+
+ cl_bool unifiedMemory = false;
+ openCLSafeCall(getScalarInfo(clGetDeviceInfo, device, CL_DEVICE_HOST_UNIFIED_MEMORY, unifiedMemory));
+ this->isUnifiedMemory = unifiedMemory != 0;
+
+ //initialize extra options for compilation. Currently only fp64 is included.
+ //Assume 4KB is enough to store all possible extensions.
+ openCLSafeCall(getStringInfo(clGetDeviceInfo, device, CL_DEVICE_EXTENSIONS, this->deviceExtensions));
+
+ size_t fp64_khr = this->deviceExtensions.find("cl_khr_fp64");
+ if(fp64_khr != std::string::npos)
+ {
+ this->compilationExtraOptions += "-D DOUBLE_SUPPORT";
+ this->haveDoubleSupport = true;
+ }
+ else
+ {
+ this->haveDoubleSupport = false;
+ }
+
+ size_t intel_platform = platformInfoImpl.platformVendor.find("Intel");
+ if(intel_platform != std::string::npos)
+ {
+ this->compilationExtraOptions += " -D INTEL_DEVICE";
+ this->isIntelDevice = true;
+ }
+ else
+ {
+ this->isIntelDevice = false;
+ }
+
+ if (id < 0)
+ {
+#ifdef CL_VERSION_1_2
+ if (this->deviceVersionMajor > 1 || (this->deviceVersionMajor == 1 && this->deviceVersionMinor >= 2))
+ {
+ ::clRetainDevice(device);
+ }
+#endif
+ }
+ }
+};
+
+static std::vector<PlatformInfoImpl> global_platforms;
+static std::vector<DeviceInfoImpl> global_devices;
+
static void split(const std::string &s, char delim, std::vector<std::string> &elems) {
std::stringstream ss(s);
std::string item;
static bool __initialized = false;
static int initializeOpenCLDevices()
{
- using namespace cl_utils;
-
assert(!__initialized);
__initialized = true;
for (size_t i = 0; i < platforms.size(); ++i)
{
PlatformInfoImpl& platformInfo = global_platforms[i];
- platformInfo.info._id = i;
cl_platform_id platform = platforms[i];
-
- platformInfo.platform_id = platform;
- openCLSafeCall(getStringInfo(clGetPlatformInfo, platform, CL_PLATFORM_PROFILE, platformInfo.info.platformProfile));
- openCLSafeCall(getStringInfo(clGetPlatformInfo, platform, CL_PLATFORM_VERSION, platformInfo.info.platformVersion));
- openCLSafeCall(getStringInfo(clGetPlatformInfo, platform, CL_PLATFORM_NAME, platformInfo.info.platformName));
- openCLSafeCall(getStringInfo(clGetPlatformInfo, platform, CL_PLATFORM_VENDOR, platformInfo.info.platformVendor));
- openCLSafeCall(getStringInfo(clGetPlatformInfo, platform, CL_PLATFORM_EXTENSIONS, platformInfo.info.platformExtensons));
-
- parseOpenCLVersion(platformInfo.info.platformVersion,
- platformInfo.info.platformVersionMajor, platformInfo.info.platformVersionMinor);
+ platformInfo.init(i, platform);
std::vector<cl_device_id> devices;
cl_int status = getDevices(platform, CL_DEVICE_TYPE_ALL, devices);
int baseIndx = global_devices.size();
global_devices.resize(baseIndx + devices.size());
platformInfo.deviceIDs.resize(devices.size());
- platformInfo.info.devices.resize(devices.size());
+ platformInfo.devices.resize(devices.size());
for(size_t j = 0; j < devices.size(); ++j)
{
cl_device_id device = devices[j];
DeviceInfoImpl& deviceInfo = global_devices[baseIndx + j];
- deviceInfo.info._id = baseIndx + j;
- deviceInfo.platform_id = platform;
- deviceInfo.device_id = device;
-
- deviceInfo.info.platform = &platformInfo.info;
- platformInfo.deviceIDs[j] = deviceInfo.info._id;
-
- cl_device_type type = cl_device_type(-1);
- openCLSafeCall(getScalarInfo(clGetDeviceInfo, device, CL_DEVICE_TYPE, type));
- deviceInfo.info.deviceType = DeviceType(type);
-
- openCLSafeCall(getStringInfo(clGetDeviceInfo, device, CL_DEVICE_PROFILE, deviceInfo.info.deviceProfile));
- openCLSafeCall(getStringInfo(clGetDeviceInfo, device, CL_DEVICE_VERSION, deviceInfo.info.deviceVersion));
- openCLSafeCall(getStringInfo(clGetDeviceInfo, device, CL_DEVICE_NAME, deviceInfo.info.deviceName));
- openCLSafeCall(getStringInfo(clGetDeviceInfo, device, CL_DEVICE_VENDOR, deviceInfo.info.deviceVendor));
- cl_uint vendorID = 0;
- openCLSafeCall(getScalarInfo(clGetDeviceInfo, device, CL_DEVICE_VENDOR_ID, vendorID));
- deviceInfo.info.deviceVendorId = vendorID;
- openCLSafeCall(getStringInfo(clGetDeviceInfo, device, CL_DRIVER_VERSION, deviceInfo.info.deviceDriverVersion));
- openCLSafeCall(getStringInfo(clGetDeviceInfo, device, CL_DEVICE_EXTENSIONS, deviceInfo.info.deviceExtensions));
-
- parseOpenCLVersion(deviceInfo.info.deviceVersion,
- deviceInfo.info.deviceVersionMajor, deviceInfo.info.deviceVersionMinor);
-
- size_t maxWorkGroupSize = 0;
- openCLSafeCall(getScalarInfo(clGetDeviceInfo, device, CL_DEVICE_MAX_WORK_GROUP_SIZE, maxWorkGroupSize));
- deviceInfo.info.maxWorkGroupSize = maxWorkGroupSize;
-
- cl_uint maxDimensions = 0;
- openCLSafeCall(getScalarInfo(clGetDeviceInfo, device, CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS, maxDimensions));
- std::vector<size_t> maxWorkItemSizes(maxDimensions);
- openCLSafeCall(clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_ITEM_SIZES, sizeof(size_t) * maxDimensions,
- (void *)&maxWorkItemSizes[0], 0));
- deviceInfo.info.maxWorkItemSizes = maxWorkItemSizes;
-
- cl_uint maxComputeUnits = 0;
- openCLSafeCall(getScalarInfo(clGetDeviceInfo, device, CL_DEVICE_MAX_COMPUTE_UNITS, maxComputeUnits));
- deviceInfo.info.maxComputeUnits = maxComputeUnits;
-
- cl_ulong localMemorySize = 0;
- openCLSafeCall(getScalarInfo(clGetDeviceInfo, device, CL_DEVICE_LOCAL_MEM_SIZE, localMemorySize));
- deviceInfo.info.localMemorySize = (size_t)localMemorySize;
-
- cl_ulong maxMemAllocSize = 0;
- openCLSafeCall(getScalarInfo(clGetDeviceInfo, device, CL_DEVICE_MAX_MEM_ALLOC_SIZE, maxMemAllocSize));
- deviceInfo.info.maxMemAllocSize = (size_t)maxMemAllocSize;
-
- cl_bool unifiedMemory = false;
- openCLSafeCall(getScalarInfo(clGetDeviceInfo, device, CL_DEVICE_HOST_UNIFIED_MEMORY, unifiedMemory));
- deviceInfo.info.isUnifiedMemory = unifiedMemory != 0;
-
- //initialize extra options for compilation. Currently only fp64 is included.
- //Assume 4KB is enough to store all possible extensions.
- openCLSafeCall(getStringInfo(clGetDeviceInfo, device, CL_DEVICE_EXTENSIONS, deviceInfo.info.deviceExtensions));
-
- size_t fp64_khr = deviceInfo.info.deviceExtensions.find("cl_khr_fp64");
- if(fp64_khr != std::string::npos)
- {
- deviceInfo.info.compilationExtraOptions += "-D DOUBLE_SUPPORT";
- deviceInfo.info.haveDoubleSupport = true;
- }
- else
- {
- deviceInfo.info.haveDoubleSupport = false;
- }
-
- size_t intel_platform = platformInfo.info.platformVendor.find("Intel");
- if(intel_platform != std::string::npos)
- {
- deviceInfo.info.compilationExtraOptions += " -D INTEL_DEVICE";
- deviceInfo.info.isIntelDevice = true;
- }
- else
- {
- deviceInfo.info.isIntelDevice = false;
- }
+ platformInfo.deviceIDs[j] = baseIndx + j;
+ deviceInfo.init(baseIndx + j, platformInfo, device);
}
}
}
for(size_t j = 0; j < platformInfo.deviceIDs.size(); ++j)
{
DeviceInfoImpl& deviceInfo = global_devices[platformInfo.deviceIDs[j]];
- platformInfo.info.devices[j] = &deviceInfo.info;
+ platformInfo.devices[j] = &deviceInfo;
}
}
// nothing
}
+DeviceInfo::~DeviceInfo() { }
+
PlatformInfo::PlatformInfo()
: _id(-1),
platformVersionMajor(0), platformVersionMinor(0)
// nothing
}
+PlatformInfo::~PlatformInfo() { }
+
+class ContextImpl;
+
+struct CommandQueue
+{
+ ContextImpl* context_;
+ cl_command_queue clQueue_;
+
+ CommandQueue() : context_(NULL), clQueue_(NULL) { }
+ ~CommandQueue() { release(); }
+
+ void create(ContextImpl* context_);
+ void release()
+ {
+#ifdef WIN32
+ // if process is on termination stage (ExitProcess was called and other threads were terminated)
+ // then disable command queue release because it may cause program hang
+ if (!__termination)
+#endif
+ {
+ if(clQueue_)
+ {
+ openCLSafeCall(clReleaseCommandQueue(clQueue_)); // some cleanup problems are here
+ }
+
+ }
+ clQueue_ = NULL;
+ context_ = NULL;
+ }
+};
+
+cv::TLSData<CommandQueue> commandQueueTLSData;
+
//////////////////////////////// OpenCL context ////////////////////////
//This is a global singleton class used to represent a OpenCL context.
class ContextImpl : public Context
{
public:
- const cl_device_id clDeviceID;
+ cl_device_id clDeviceID;
cl_context clContext;
- cl_command_queue clCmdQueue;
- const DeviceInfo& deviceInfo;
+ const DeviceInfoImpl& deviceInfoImpl;
protected:
- ContextImpl(const DeviceInfo& deviceInfo, cl_device_id clDeviceID)
- : clDeviceID(clDeviceID), clContext(NULL), clCmdQueue(NULL), deviceInfo(deviceInfo)
+ ContextImpl(const DeviceInfoImpl& _deviceInfoImpl, cl_context context)
+ : clDeviceID(_deviceInfoImpl.device_id), clContext(context), deviceInfoImpl(_deviceInfoImpl)
{
- // nothing
+#ifdef CL_VERSION_1_2
+ if (supportsFeature(FEATURE_CL_VER_1_2))
+ {
+ openCLSafeCall(clRetainDevice(clDeviceID));
+ }
+#endif
+ openCLSafeCall(clRetainContext(clContext));
+
+ ContextImpl* old = NULL;
+ {
+ cv::AutoLock lock(getCurrentContextMutex());
+ old = currentContext;
+ currentContext = this;
+ }
+ if (old != NULL)
+ {
+ delete old;
+ }
+ }
+ ~ContextImpl()
+ {
+ CV_Assert(this != currentContext);
+
+#ifdef CL_VERSION_1_2
+ if (supportsFeature(FEATURE_CL_VER_1_2))
+ {
+ openCLSafeCall(clReleaseDevice(clDeviceID));
+ }
+#endif
+ if (deviceInfoImpl._id < 0) // not in the global registry, so we should cleanup it
+ {
+#ifdef CL_VERSION_1_2
+ if (supportsFeature(FEATURE_CL_VER_1_2))
+ {
+ openCLSafeCall(clReleaseDevice(deviceInfoImpl.device_id));
+ }
+#endif
+ PlatformInfoImpl* platformImpl = (PlatformInfoImpl*)(deviceInfoImpl.platform);
+ delete platformImpl;
+ delete const_cast<DeviceInfoImpl*>(&deviceInfoImpl);
+ }
+ clDeviceID = NULL;
+
+#ifdef WIN32
+ // if process is on termination stage (ExitProcess was called and other threads were terminated)
+ // then disable command queue release because it may cause program hang
+ if (!__termination)
+#endif
+ {
+ if(clContext)
+ {
+ openCLSafeCall(clReleaseContext(clContext));
+ }
+ }
+ clContext = NULL;
}
- ~ContextImpl();
public:
static void setContext(const DeviceInfo* deviceInfo);
+ static void initializeContext(void* pClPlatform, void* pClContext, void* pClDevice);
bool supportsFeature(FEATURE_TYPE featureType) const;
static void cleanupContext(void);
+ static ContextImpl* getContext();
private:
ContextImpl(const ContextImpl&); // disabled
ContextImpl& operator=(const ContextImpl&); // disabled
+
+ static ContextImpl* currentContext;
};
-static ContextImpl* currentContext = NULL;
+ContextImpl* ContextImpl::currentContext = NULL;
static bool __deviceSelected = false;
Context* Context::getContext()
{
+ return ContextImpl::getContext();
+}
+
+ContextImpl* ContextImpl::getContext()
+{
if (currentContext == NULL)
{
static bool defaultInitiaization = false;
const DeviceInfo& Context::getDeviceInfo() const
{
- return ((ContextImpl*)this)->deviceInfo;
+ return ((ContextImpl*)this)->deviceInfoImpl;
}
const void* Context::getOpenCLContextPtr() const
const void* Context::getOpenCLCommandQueuePtr() const
{
- return &(((ContextImpl*)this)->clCmdQueue);
+ ContextImpl* pThis = (ContextImpl*)this;
+ CommandQueue* commandQueue = commandQueueTLSData.get();
+ if (commandQueue->context_ != pThis)
+ {
+ commandQueue->create(pThis);
+ }
+ return &commandQueue->clQueue_;
}
const void* Context::getOpenCLDeviceIDPtr() const
switch (featureType)
{
case FEATURE_CL_INTEL_DEVICE:
- return deviceInfo.isIntelDevice;
+ return deviceInfoImpl.isIntelDevice;
case FEATURE_CL_DOUBLE:
- return deviceInfo.haveDoubleSupport;
+ return deviceInfoImpl.haveDoubleSupport;
case FEATURE_CL_UNIFIED_MEM:
- return deviceInfo.isUnifiedMemory;
+ return deviceInfoImpl.isUnifiedMemory;
case FEATURE_CL_VER_1_2:
- return deviceInfo.deviceVersionMajor > 1 || (deviceInfo.deviceVersionMajor == 1 && deviceInfo.deviceVersionMinor >= 2);
+ return deviceInfoImpl.deviceVersionMajor > 1 || (deviceInfoImpl.deviceVersionMajor == 1 && deviceInfoImpl.deviceVersionMinor >= 2);
}
CV_Error(CV_StsBadArg, "Invalid feature type");
return false;
}
-#if defined(WIN32)
-static bool __termination = false;
-#endif
-
-ContextImpl::~ContextImpl()
-{
-#ifdef WIN32
- // if process is on termination stage (ExitProcess was called and other threads were terminated)
- // then disable command queue release because it may cause program hang
- if (!__termination)
-#endif
- {
- if(clCmdQueue)
- {
- openCLSafeCall(clReleaseCommandQueue(clCmdQueue)); // some cleanup problems are here
- }
-
- if(clContext)
- {
- openCLSafeCall(clReleaseContext(clContext));
- }
- }
- clCmdQueue = NULL;
- clContext = NULL;
-}
-
void fft_teardown();
void clBlasTeardown();
fft_teardown();
clBlasTeardown();
- cv::AutoLock lock(__module.currentContextMutex);
+ cv::AutoLock lock(getCurrentContextMutex());
if (currentContext)
- delete currentContext;
- currentContext = NULL;
+ {
+ ContextImpl* ctx = currentContext;
+ currentContext = NULL;
+ delete ctx;
+ }
}
void ContextImpl::setContext(const DeviceInfo* deviceInfo)
{
- CV_Assert(deviceInfo->_id >= 0 && deviceInfo->_id < (int)global_devices.size());
+ CV_Assert(deviceInfo->_id >= 0); // we can't specify custom devices
+ CV_Assert(deviceInfo->_id < (int)global_devices.size());
{
- cv::AutoLock lock(__module.currentContextMutex);
+ cv::AutoLock lock(getCurrentContextMutex());
if (currentContext)
{
- if (currentContext->deviceInfo._id == deviceInfo->_id)
+ if (currentContext->deviceInfoImpl._id == deviceInfo->_id)
return;
}
}
DeviceInfoImpl& infoImpl = global_devices[deviceInfo->_id];
- CV_Assert(deviceInfo == &infoImpl.info);
+ CV_Assert(deviceInfo == &infoImpl);
cl_int status = 0;
cl_context_properties cps[3] = { CL_CONTEXT_PLATFORM, (cl_context_properties)(infoImpl.platform_id), 0 };
cl_context clContext = clCreateContext(cps, 1, &infoImpl.device_id, NULL, NULL, &status);
openCLVerifyCall(status);
-#ifdef PRINT_KERNEL_RUN_TIME
- cl_command_queue clCmdQueue = clCreateCommandQueue(clContext, infoImpl.device_id, CL_QUEUE_PROFILING_ENABLE, &status);
-#else /*PRINT_KERNEL_RUN_TIME*/
- cl_command_queue clCmdQueue = clCreateCommandQueue(clContext, infoImpl.device_id, 0, &status);
-#endif /*PRINT_KERNEL_RUN_TIME*/
- openCLVerifyCall(status);
- ContextImpl* ctx = new ContextImpl(infoImpl.info, infoImpl.device_id);
- ctx->clCmdQueue = clCmdQueue;
- ctx->clContext = clContext;
+ ContextImpl* ctx = new ContextImpl(infoImpl, clContext);
+ clReleaseContext(clContext);
+ (void)ctx;
+}
- ContextImpl* old = NULL;
- {
- cv::AutoLock lock(__module.currentContextMutex);
- old = currentContext;
- currentContext = ctx;
- }
- if (old != NULL)
- {
- delete old;
- }
+void ContextImpl::initializeContext(void* pClPlatform, void* pClContext, void* pClDevice)
+{
+ CV_Assert(pClPlatform != NULL);
+ CV_Assert(pClContext != NULL);
+ CV_Assert(pClDevice != NULL);
+ cl_platform_id platform = *(cl_platform_id*)pClPlatform;
+ cl_context context = *(cl_context*)pClContext;
+ cl_device_id device = *(cl_device_id*)pClDevice;
+
+ PlatformInfoImpl* platformInfoImpl = new PlatformInfoImpl();
+ platformInfoImpl->init(-1, platform);
+ DeviceInfoImpl* deviceInfoImpl = new DeviceInfoImpl();
+ deviceInfoImpl->init(-1, *platformInfoImpl, device);
+
+ ContextImpl* ctx = new ContextImpl(*deviceInfoImpl, context);
+ (void)ctx;
+}
+
+void CommandQueue::create(ContextImpl* context)
+{
+ release();
+ cl_int status = 0;
+ // TODO add CL_QUEUE_PROFILING_ENABLE
+ cl_command_queue clCmdQueue = clCreateCommandQueue(context->clContext, context->clDeviceID, 0, &status);
+ openCLVerifyCall(status);
+ context_ = context;
+ clQueue_ = clCmdQueue;
}
int getOpenCLPlatforms(PlatformsInfo& platforms)
for (size_t id = 0; id < global_platforms.size(); ++id)
{
PlatformInfoImpl& impl = global_platforms[id];
- platforms.push_back(&impl.info);
+ platforms.push_back(&impl);
}
return platforms.size();
for (size_t id = 0; id < global_devices.size(); ++id)
{
DeviceInfoImpl& deviceInfo = global_devices[id];
- if (((int)deviceInfo.info.deviceType & deviceType) != 0)
+ if (((int)deviceInfo.deviceType & deviceType) != 0)
{
- devices.push_back(&deviceInfo.info);
+ devices.push_back(&deviceInfo);
}
}
}
}
}
+void initializeContext(void* pClPlatform, void* pClContext, void* pClDevice)
+{
+ try
+ {
+ ContextImpl::initializeContext(pClPlatform, pClContext, pClDevice);
+ __deviceSelected = true;
+ }
+ catch (...)
+ {
+ __deviceSelected = true;
+ throw;
+ }
+}
+
bool supportsFeature(FEATURE_TYPE featureType)
{
return Context::getContext()->supportsFeature(featureType);
//M*/
#include "test_precomp.hpp"
-#include "opencv2/core/opencl/runtime/opencl_core.hpp" // for OpenCL types: cl_mem
+#include "opencv2/core/opencl/runtime/opencl_core.hpp" // for OpenCL types & functions
#include "opencv2/core/ocl.hpp"
TEST(TestAPI, openCLExecuteKernelInterop)
t = (double)cv::getTickCount() - t;
printf("cpu exec time = %gms per iter\n", t*1000./niters/cv::getTickFrequency());
}
+
+// This test must be DISABLED by default!
+// (We can't restore original context for other tests)
+TEST(TestAPI, DISABLED_InitializationFromHandles)
+{
+#define MAX_PLATFORMS 16
+ cl_platform_id platforms[MAX_PLATFORMS] = { NULL };
+ cl_uint numPlatforms = 0;
+ cl_int status = ::clGetPlatformIDs(MAX_PLATFORMS, &platforms[0], &numPlatforms);
+ ASSERT_EQ(CL_SUCCESS, status) << "clGetPlatformIDs";
+ ASSERT_NE(0, (int)numPlatforms);
+
+ int selectedPlatform = 0;
+ cl_platform_id platform = platforms[selectedPlatform];
+
+ ASSERT_NE((void*)NULL, platform);
+
+ cl_device_id device = NULL;
+ status = ::clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 1, &device, NULL);
+ ASSERT_EQ(CL_SUCCESS, status) << "clGetDeviceIDs";
+ ASSERT_NE((void*)NULL, device);
+
+ cl_context_properties cps[3] = { CL_CONTEXT_PLATFORM, (cl_context_properties)(platform), 0 };
+ cl_context context = ::clCreateContext(cps, 1, &device, NULL, NULL, &status);
+ ASSERT_EQ(CL_SUCCESS, status) << "clCreateContext";
+ ASSERT_NE((void*)NULL, context);
+
+ ASSERT_NO_THROW(cv::ocl::initializeContext(&platform, &context, &device));
+
+ status = ::clReleaseContext(context);
+ ASSERT_EQ(CL_SUCCESS, status) << "clReleaseContext";
+
+#ifdef CL_VERSION_1_2
+#if 1
+ {
+ cv::ocl::Context* ctx = cv::ocl::Context::getContext();
+ ASSERT_NE((void*)NULL, ctx);
+ if (ctx->supportsFeature(cv::ocl::FEATURE_CL_VER_1_2)) // device supports OpenCL 1.2+
+ {
+ status = ::clReleaseDevice(device);
+ ASSERT_EQ(CL_SUCCESS, status) << "clReleaseDevice";
+ }
+ }
+#else // code below doesn't work on Linux (SEGFAULTs on 1.1- devices are not handled via exceptions)
+ try
+ {
+ status = ::clReleaseDevice(device); // NOTE This works only with !DEVICES! that supports OpenCL 1.2
+ (void)status; // no check
+ }
+ catch (...)
+ {
+ // nothing, there is no problem
+ }
+#endif
+#endif
+
+ // print the name of current device
+ cv::ocl::Context* ctx = cv::ocl::Context::getContext();
+ ASSERT_NE((void*)NULL, ctx);
+ const cv::ocl::DeviceInfo& deviceInfo = ctx->getDeviceInfo();
+ std::cout << "Device name: " << deviceInfo.deviceName << std::endl;
+ std::cout << "Platform name: " << deviceInfo.platform->platformName << std::endl;
+
+ ASSERT_EQ(context, *(cl_context*)ctx->getOpenCLContextPtr());
+ ASSERT_EQ(device, *(cl_device_id*)ctx->getOpenCLDeviceIDPtr());
+
+ // do some calculations and check results
+ cv::RNG rng;
+ Size sz(100, 100);
+ cv::Mat srcMat = cvtest::randomMat(rng, sz, CV_32FC4, -10, 10, false);
+ cv::Mat dstMat;
+
+ cv::ocl::oclMat srcGpuMat(srcMat);
+ cv::ocl::oclMat dstGpuMat;
+
+ cv::Scalar v = cv::Scalar::all(1);
+ cv::add(srcMat, v, dstMat);
+ cv::ocl::add(srcGpuMat, v, dstGpuMat);
+
+ cv::Mat dstGpuMatMap;
+ dstGpuMat.download(dstGpuMatMap);
+
+ EXPECT_LE(checkNorm(dstMat, dstGpuMatMap), 1e-3);
+}
add_subdirectory(gpu)
add_subdirectory(ocl)
+if(WIN32 AND HAVE_DIRECTX)
+ add_subdirectory(directx)
+endif()
+
if(ANDROID AND BUILD_ANDROID_EXAMPLES)
add_subdirectory(android)
endif()
add_subdirectory(ocl)
# FIXIT: can't use cvconfig.h in samples: add_subdirectory(gpu)
+if(WIN32)
+ add_subdirectory(directx)
+endif()
+
#
# END OF BUILD CASE 2: Build samples with library binaries
#
Mat descriptors2; de.compute( img2, keypoints2, descriptors2 );
// Match descriptors
- BFMatcher matcher(NORM_L1);
+ BFMatcher matcher(de.defaultNorm());
vector<DMatch> matches;
matcher.match( descriptors1, descriptors2, matches );
while( images.size() > 0 )
{
- if( bowTrainer.descripotorsCount() > maxDescCount )
+ if( bowTrainer.descriptorsCount() > maxDescCount )
{
#ifdef DEBUG_DESC_PROGRESS
- cout << "Breaking due to full memory ( descriptors count = " << bowTrainer.descripotorsCount()
+ cout << "Breaking due to full memory ( descriptors count = " << bowTrainer.descriptorsCount()
<< "; descriptor size in bytes = " << descByteSize << "; all used memory = "
- << bowTrainer.descripotorsCount()*descByteSize << endl;
+ << bowTrainer.descriptorsCount()*descByteSize << endl;
#endif
break;
}
for( int i = 0; i < descCount; i++ )
{
- if( usedMask[i] && bowTrainer.descripotorsCount() < maxDescCount )
+ if( usedMask[i] && bowTrainer.descriptorsCount() < maxDescCount )
bowTrainer.add( imageDescriptors.row(i) );
}
}
#ifdef DEBUG_DESC_PROGRESS
cout << images.size() << " images left, " << images[randImgIdx].id << " processed - "
<</* descs_extracted << "/" << image_descriptors.rows << " extracted - " << */
- cvRound((static_cast<double>(bowTrainer.descripotorsCount())/static_cast<double>(maxDescCount))*100.0)
+ cvRound((static_cast<double>(bowTrainer.descriptorsCount())/static_cast<double>(maxDescCount))*100.0)
<< " % memory used" << ( imageDescriptors.empty() ? " -> no descriptors extracted, skipping" : "") << endl;
#endif
images.erase( images.begin() + randImgIdx );
}
- cout << "Maximum allowed descriptor count: " << maxDescCount << ", Actual descriptor count: " << bowTrainer.descripotorsCount() << endl;
+ cout << "Maximum allowed descriptor count: " << maxDescCount << ", Actual descriptor count: " << bowTrainer.descriptorsCount() << endl;
cout << "Training vocabulary..." << endl;
vocabulary = bowTrainer.cluster();
//Do matching using features2d
cout << "matching with BruteForceMatcher<Hamming>" << endl;
- BFMatcher matcher_popcount(NORM_HAMMING);
+ BFMatcher matcher_popcount(extractor.defaultNorm());
vector<DMatch> matches_popcount;
double pop_time = match(kpts_1, kpts_2, matcher_popcount, desc_1, desc_2, matches_popcount);
cout << "done BruteForceMatcher<Hamming> matching. took " << pop_time << " seconds" << endl;
virtual void readAlgorithm( )
{
string classifierFile = data_path + "/features2d/calonder_classifier.rtc";
+ Ptr<DescriptorExtractor> extractor = makePtr<CalonderDescriptorExtractor<float> >( classifierFile );
defaultDescMatcher = makePtr<VectorDescriptorMatch>(
- makePtr<CalonderDescriptorExtractor<float> >( classifierFile ),
- makePtr<BFMatcher>(int(NORM_L2)));
+ extractor,
+ makePtr<BFMatcher>(extractor->defaultNorm()));
specificDescMatcher = defaultDescMatcher;
}
};
// The standard Hamming distance can be used such as
// BFMatcher matcher(NORM_HAMMING);
// or the proposed cascade of hamming distance using SSSE3
- BFMatcher matcher(NORM_HAMMING);
+ BFMatcher matcher(extractor.defaultNorm());
// detect
double t = (double)getTickCount();
extractor.compute(img2, keypoints2, descriptors2);
// matching descriptors
- BFMatcher matcher(NORM_L2);
+ BFMatcher matcher(extractor.defaultNorm());
vector<DMatch> matches;
matcher.match(descriptors1, descriptors2, matches);
extractor.compute(img2, keypoints2, descriptors2);
// matching descriptors
- BFMatcher matcher(NORM_L2);
+ BFMatcher matcher(extractor.defaultNorm());
vector<DMatch> matches;
matcher.match(descriptors1, descriptors2, matches);
const char* algorithm_opt = "--algorithm=";
const char* maxdisp_opt = "--max-disparity=";
const char* blocksize_opt = "--blocksize=";
- const char* nodisplay_opt = "--no-display=";
+ const char* nodisplay_opt = "--no-display";
const char* scale_opt = "--scale=";
if(argc < 3)
extractor.compute( img_2, keypoints_2, descriptors_2 );
//-- Step 3: Matching descriptor vectors with a brute force matcher
- BFMatcher matcher(NORM_L2);
+ BFMatcher matcher(extractor.defaultNorm());
std::vector< DMatch > matches;
matcher.match( descriptors_1, descriptors_2, matches );
vector<DMatch> matches;
- BFMatcher desc_matcher(NORM_HAMMING);
+ BFMatcher desc_matcher(brief.defaultNorm());
vector<Point2f> train_pts, query_pts;
vector<KeyPoint> train_kpts, query_kpts;
--- /dev/null
+SET(OPENCV_DIRECTX_SAMPLES_REQUIRED_DEPS opencv_core opencv_imgproc opencv_highgui)
+
+ocv_check_dependencies(${OPENCV_DIRECTX_SAMPLES_REQUIRED_DEPS})
+
+if(BUILD_EXAMPLES AND OCV_DEPENDENCIES_FOUND)
+ set(project "directx")
+ string(TOUPPER "${project}" project_upper)
+
+ project("${project}_samples")
+
+ ocv_include_modules(${OPENCV_DIRECTX_SAMPLES_REQUIRED_DEPS})
+
+ # ---------------------------------------------
+ # Define executable targets
+ # ---------------------------------------------
+ MACRO(OPENCV_DEFINE_DIRECTX_EXAMPLE name srcs)
+ set(the_target "example_${project}_${name}")
+ add_executable(${the_target} ${srcs})
+
+ target_link_libraries(${the_target} ${OPENCV_LINKER_LIBS} ${OPENCV_DIRECTX_SAMPLES_REQUIRED_DEPS})
+
+ set_target_properties(${the_target} PROPERTIES
+ OUTPUT_NAME "${project}-example-${name}"
+ PROJECT_LABEL "(EXAMPLE_${project_upper}) ${name}")
+
+ if(ENABLE_SOLUTION_FOLDERS)
+ set_target_properties(${the_target} PROPERTIES FOLDER "samples//${project}")
+ endif()
+
+ if(WIN32)
+ if(MSVC AND NOT BUILD_SHARED_LIBS)
+ set_target_properties(${the_target} PROPERTIES LINK_FLAGS "/NODEFAULTLIB:atlthunk.lib /NODEFAULTLIB:atlsd.lib /DEBUG")
+ endif()
+ install(TARGETS ${the_target} RUNTIME DESTINATION "${OPENCV_SAMPLES_BIN_INSTALL_PATH}/${project}" COMPONENT main)
+ endif()
+ ENDMACRO()
+
+ file(GLOB all_samples RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} *.cpp)
+
+ foreach(sample_filename ${all_samples})
+ get_filename_component(sample ${sample_filename} NAME_WE)
+ file(GLOB sample_srcs RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${sample}.*)
+ OPENCV_DEFINE_DIRECTX_EXAMPLE(${sample} ${sample_srcs})
+ endforeach()
+endif()
--- /dev/null
+#include <windows.h>
+#include <d3d10.h>
+#pragma comment (lib, "d3d10.lib")
+
+#define USE_D3D10
+#define WINDOW_NAME "OpenCV Direct3D 10 Sample"
+
+IDXGISwapChain *swapchain = NULL;
+ID3D10Device *dev = NULL;
+ID3D10Texture2D *pBackBufferTexture = NULL;
+ID3D10Texture2D *pCPUWriteTexture = NULL;
+ID3D10Texture2D *pInputTexture = NULL;
+ID3D10RenderTargetView *backbuffer = NULL;
+
+#include "d3d_base.inl.hpp"
+
+bool initDirect3D()
+{
+ DXGI_SWAP_CHAIN_DESC scd;
+
+ ZeroMemory(&scd, sizeof(DXGI_SWAP_CHAIN_DESC));
+
+ scd.BufferCount = 1; // one back buffer
+ scd.BufferDesc.Format = DXGI_FORMAT_R8G8B8A8_UNORM; // use 32-bit color
+ scd.BufferDesc.Width = WIDTH; // set the back buffer width
+ scd.BufferDesc.Height = HEIGHT; // set the back buffer height
+ scd.BufferUsage = DXGI_USAGE_RENDER_TARGET_OUTPUT; // how swap chain is to be used
+ scd.OutputWindow = hWnd; // the window to be used
+ scd.SampleDesc.Count = 1; // how many multisamples
+ scd.Windowed = TRUE; // windowed/full-screen mode
+ scd.SwapEffect = DXGI_SWAP_EFFECT_DISCARD;
+ scd.Flags = DXGI_SWAP_CHAIN_FLAG_ALLOW_MODE_SWITCH; // allow full-screen switching
+
+ if (FAILED(D3D10CreateDeviceAndSwapChain(
+ NULL,
+ D3D10_DRIVER_TYPE_HARDWARE,
+ NULL,
+ 0,
+ D3D10_SDK_VERSION,
+ &scd,
+ &swapchain,
+ &dev)))
+ {
+ return false;
+ }
+
+ if (FAILED(swapchain->GetBuffer(0, __uuidof(ID3D10Texture2D), (LPVOID*)&pBackBufferTexture)))
+ {
+ return false;
+ }
+
+ if (FAILED(dev->CreateRenderTargetView(pBackBufferTexture, NULL, &backbuffer)))
+ {
+ return false;
+ }
+
+ dev->OMSetRenderTargets(1, &backbuffer, NULL);
+
+ D3D10_VIEWPORT viewport;
+ ZeroMemory(&viewport, sizeof(D3D10_VIEWPORT));
+ viewport.Width = WIDTH;
+ viewport.Height = HEIGHT;
+ viewport.MinDepth = 0.0f;
+ viewport.MaxDepth = 0.0f;
+ dev->RSSetViewports(1, &viewport);
+
+ return true;
+}
+
+bool initDirect3DTextures()
+{
+ { // Create texture for demo 0
+ D3D10_TEXTURE2D_DESC desc = { 0 };
+ desc.Width = WIDTH;
+ desc.Height = HEIGHT;
+ desc.MipLevels = desc.ArraySize = 1;
+ desc.Format = DXGI_FORMAT_R8G8B8A8_UNORM;
+ desc.SampleDesc.Count = 1;
+ desc.BindFlags = D3D10_BIND_SHADER_RESOURCE;
+ desc.Usage = D3D10_USAGE_DYNAMIC;
+ desc.CPUAccessFlags = D3D10_CPU_ACCESS_WRITE;
+ if (FAILED(dev->CreateTexture2D(&desc, NULL, &pCPUWriteTexture)))
+ {
+ std::cerr << "Can't create texture for CPU write sample" << std::endl;
+ return false;
+ }
+ }
+
+ { // Create Read-only texture
+ cv::Mat inputMat = getInputTexture();
+
+ D3D10_TEXTURE2D_DESC desc = { 0 };
+ desc.Width = inputMat.size().width;
+ desc.Height = inputMat.size().height;
+ desc.MipLevels = desc.ArraySize = 1;
+ desc.Format = DXGI_FORMAT_R8G8B8A8_UNORM;
+ desc.SampleDesc.Count = 1;
+ desc.BindFlags = D3D10_BIND_SHADER_RESOURCE;
+ desc.Usage = D3D10_USAGE_IMMUTABLE;
+ desc.CPUAccessFlags = cv::ocl::useOpenCL() ? 0 : D3D10_CPU_ACCESS_READ;
+
+ D3D10_SUBRESOURCE_DATA srInitData;
+ srInitData.pSysMem = inputMat.data;
+ srInitData.SysMemPitch = (UINT)inputMat.step[0];
+
+ if (FAILED(dev->CreateTexture2D(&desc, &srInitData, &pInputTexture)))
+ {
+ std::cerr << "Can't create texture with input image" << std::endl;
+ return false;
+ }
+ }
+
+ return true;
+}
+
+void cleanUp(void)
+{
+ if (swapchain) swapchain->SetFullscreenState(FALSE, NULL); // switch to windowed mode
+
+ SAFE_RELEASE(swapchain);
+ SAFE_RELEASE(pCPUWriteTexture);
+ SAFE_RELEASE(pInputTexture);
+ SAFE_RELEASE(pBackBufferTexture);
+ SAFE_RELEASE(backbuffer);
+ SAFE_RELEASE(dev);
+}
+
+
+void render(void)
+{
+ // check to make sure you have a valid Direct3D device
+ CV_Assert(dev);
+
+ renderToD3DObject();
+
+ // switch the back buffer and the front buffer
+ swapchain->Present(0, 0);
+}
--- /dev/null
+#include <windows.h>
+#include <d3d11.h>
+#pragma comment (lib, "d3d11.lib")
+
+#define USE_D3D11
+#define WINDOW_NAME "OpenCV Direct3D 11 Sample"
+
+IDXGISwapChain *swapchain = NULL;
+ID3D11Device *dev = NULL;
+ID3D11DeviceContext *devcon = NULL;
+ID3D11Texture2D *pBackBufferTexture = NULL;
+ID3D11Texture2D *pCPUWriteTexture = NULL;
+ID3D11Texture2D *pInputTexture = NULL;
+ID3D11RenderTargetView *backbuffer = NULL;
+
+#include "d3d_base.inl.hpp"
+
+bool initDirect3D()
+{
+ DXGI_SWAP_CHAIN_DESC scd;
+
+ ZeroMemory(&scd, sizeof(DXGI_SWAP_CHAIN_DESC));
+
+ scd.BufferCount = 1; // one back buffer
+ scd.BufferDesc.Format = DXGI_FORMAT_R8G8B8A8_UNORM; // use 32-bit color
+ scd.BufferDesc.Width = WIDTH; // set the back buffer width
+ scd.BufferDesc.Height = HEIGHT; // set the back buffer height
+ scd.BufferUsage = DXGI_USAGE_RENDER_TARGET_OUTPUT; // how swap chain is to be used
+ scd.OutputWindow = hWnd; // the window to be used
+ scd.SampleDesc.Count = 1; // how many multisamples
+ scd.Windowed = TRUE; // windowed/full-screen mode
+ scd.SwapEffect = DXGI_SWAP_EFFECT_DISCARD;
+ scd.Flags = DXGI_SWAP_CHAIN_FLAG_ALLOW_MODE_SWITCH; // allow full-screen switching
+
+ if (FAILED(D3D11CreateDeviceAndSwapChain(
+ NULL,
+ D3D_DRIVER_TYPE_HARDWARE,
+ NULL,
+ 0,
+ NULL,
+ 0,
+ D3D11_SDK_VERSION,
+ &scd,
+ &swapchain,
+ &dev,
+ NULL,
+ &devcon)))
+ {
+ return false;
+ }
+
+ if (FAILED(swapchain->GetBuffer(0, __uuidof(ID3D11Texture2D), (LPVOID*)&pBackBufferTexture)))
+ {
+ return false;
+ }
+
+ if (FAILED(dev->CreateRenderTargetView(pBackBufferTexture, NULL, &backbuffer)))
+ {
+ return false;
+ }
+
+ devcon->OMSetRenderTargets(1, &backbuffer, NULL);
+
+ D3D11_VIEWPORT viewport = { 0 };
+ viewport.Width = WIDTH;
+ viewport.Height = HEIGHT;
+ viewport.MinDepth = 0.0f;
+ viewport.MaxDepth = 0.0f;
+ devcon->RSSetViewports(1, &viewport);
+
+ return true;
+}
+
+bool initDirect3DTextures()
+{
+ { // Create texture for demo 0
+ D3D11_TEXTURE2D_DESC desc = { 0 };
+ desc.Width = WIDTH;
+ desc.Height = HEIGHT;
+ desc.MipLevels = desc.ArraySize = 1;
+ desc.Format = DXGI_FORMAT_R8G8B8A8_UNORM;
+ desc.SampleDesc.Count = 1;
+ desc.BindFlags = D3D11_BIND_SHADER_RESOURCE;
+ desc.Usage = D3D11_USAGE_DYNAMIC;
+ desc.CPUAccessFlags = D3D11_CPU_ACCESS_WRITE;
+ if (FAILED(dev->CreateTexture2D(&desc, NULL, &pCPUWriteTexture)))
+ {
+ std::cerr << "Can't create texture for CPU write sample" << std::endl;
+ return false;
+ }
+ }
+
+ { // Create Read-only texture
+ cv::Mat inputMat = getInputTexture();
+
+ D3D11_TEXTURE2D_DESC desc = { 0 };
+ desc.Width = inputMat.size().width;
+ desc.Height = inputMat.size().height;
+ desc.MipLevels = desc.ArraySize = 1;
+ desc.Format = DXGI_FORMAT_R8G8B8A8_UNORM;
+ desc.SampleDesc.Count = 1;
+ desc.BindFlags = D3D11_BIND_SHADER_RESOURCE;
+ desc.Usage = D3D11_USAGE_IMMUTABLE;
+ desc.CPUAccessFlags = cv::ocl::useOpenCL() ? 0 : D3D11_CPU_ACCESS_READ;
+
+ D3D11_SUBRESOURCE_DATA srInitData;
+ srInitData.pSysMem = inputMat.data;
+ srInitData.SysMemPitch = (UINT)inputMat.step[0];
+
+ if (FAILED(dev->CreateTexture2D(&desc, &srInitData, &pInputTexture)))
+ {
+ std::cerr << "Can't create texture with input image" << std::endl;
+ return false;
+ }
+ }
+
+ return true;
+}
+
+void cleanUp(void)
+{
+ if (swapchain) swapchain->SetFullscreenState(FALSE, NULL); // switch to windowed mode
+
+ SAFE_RELEASE(swapchain);
+ SAFE_RELEASE(pCPUWriteTexture);
+ SAFE_RELEASE(pInputTexture);
+ SAFE_RELEASE(pBackBufferTexture);
+ SAFE_RELEASE(backbuffer);
+ SAFE_RELEASE(dev);
+ SAFE_RELEASE(devcon);
+}
+
+
+void render(void)
+{
+ // check to make sure you have a valid Direct3D device
+ CV_Assert(dev);
+
+ renderToD3DObject();
+
+ // switch the back buffer and the front buffer
+ swapchain->Present(0, 0);
+}
--- /dev/null
+#include <windows.h>
+#include <d3d9.h>
+#pragma comment (lib, "d3d9.lib")
+
+#define USE_D3D9
+#define WINDOW_NAME "OpenCV Direct3D 9 Sample"
+
+IDirect3D9 *pD3D = NULL;
+IDirect3DDevice9 *dev = NULL;
+IDirect3DSurface9 *pBackBuffer = NULL;
+IDirect3DSurface9 *pCPUWriteSurface = NULL; // required name
+IDirect3DSurface9 *pReadOnlySurface = NULL; // required name
+HANDLE readOnlySurfaceShared = 0; // required name
+IDirect3DSurface9 *pSurface = NULL; // required name
+HANDLE surfaceShared = 0; // required name
+
+#include "d3d_base.inl.hpp"
+
+bool initDirect3D(void)
+{
+ if (NULL == (pD3D = Direct3DCreate9(D3D_SDK_VERSION)))
+ {
+ return false;
+ }
+
+ D3DPRESENT_PARAMETERS d3dpp;
+ ZeroMemory(&d3dpp,sizeof(D3DPRESENT_PARAMETERS));
+
+ DWORD flags = D3DCREATE_HARDWARE_VERTEXPROCESSING | D3DCREATE_PUREDEVICE | D3DCREATE_NOWINDOWCHANGES
+ | D3DCREATE_MULTITHREADED;
+
+ d3dpp.Windowed = true;
+ d3dpp.Flags = 0;
+ d3dpp.BackBufferCount = 0;
+ d3dpp.BackBufferFormat = D3DFMT_A8R8G8B8;
+ d3dpp.BackBufferHeight = HEIGHT;
+ d3dpp.BackBufferWidth = WIDTH;
+ d3dpp.MultiSampleType = D3DMULTISAMPLE_NONE;
+ d3dpp.SwapEffect = D3DSWAPEFFECT_DISCARD;
+ d3dpp.hDeviceWindow = hWnd;
+ d3dpp.PresentationInterval = D3DPRESENT_INTERVAL_IMMEDIATE;
+ d3dpp.FullScreen_RefreshRateInHz = D3DPRESENT_RATE_DEFAULT;
+
+ if (FAILED(pD3D->CreateDevice(D3DADAPTER_DEFAULT, D3DDEVTYPE_HAL, hWnd, flags, &d3dpp, &dev)))
+ {
+ return false;
+ }
+
+ if (FAILED(dev->GetBackBuffer(0, 0, D3DBACKBUFFER_TYPE_MONO, &pBackBuffer)))
+ {
+ return false;
+ }
+
+ return true;
+}
+
+bool initDirect3DTextures()
+{
+ // Note: sharing is not supported on some platforms
+ if (FAILED(dev->CreateOffscreenPlainSurface(WIDTH, HEIGHT, D3DFMT_A8R8G8B8, D3DPOOL_DEFAULT, &pSurface, NULL/*&surfaceShared*/)))
+ {
+ std::cerr << "Can't create surface for result" << std::endl;
+ return false;
+ }
+
+ // Note: sharing is not supported on some platforms
+ if (FAILED(dev->CreateOffscreenPlainSurface(WIDTH, HEIGHT, D3DFMT_A8R8G8B8, D3DPOOL_DEFAULT, &pReadOnlySurface, NULL/*&readOnlySurfaceShared*/)))
+ {
+ std::cerr << "Can't create read only surface" << std::endl;
+ return false;
+ }
+ else
+ {
+ IDirect3DSurface9* pTmpSurface;
+ if (FAILED(dev->CreateOffscreenPlainSurface(WIDTH, HEIGHT, D3DFMT_A8R8G8B8, D3DPOOL_DEFAULT, &pTmpSurface, NULL)))
+ {
+ std::cerr << "Can't create temp surface for CPU write" << std::endl;
+ return false;
+ }
+
+ D3DLOCKED_RECT memDesc = {0, NULL};
+ RECT rc = {0, 0, WIDTH, HEIGHT};
+ if (SUCCEEDED(pTmpSurface->LockRect(&memDesc, &rc, 0)))
+ {
+ cv::Mat m(cv::Size(WIDTH, HEIGHT), CV_8UC4, memDesc.pBits, (int)memDesc.Pitch);
+ getInputTexture().copyTo(m);
+ pTmpSurface->UnlockRect();
+ dev->StretchRect(pTmpSurface, NULL, pReadOnlySurface, NULL, D3DTEXF_NONE);
+ }
+ else
+ {
+ std::cerr << "Can't LockRect() on surface" << std::endl;
+ }
+ pTmpSurface->Release();
+ }
+
+ if (FAILED(dev->CreateOffscreenPlainSurface(WIDTH, HEIGHT, D3DFMT_A8R8G8B8, D3DPOOL_DEFAULT, &pCPUWriteSurface, NULL)))
+ {
+ std::cerr << "Can't create surface for CPU write" << std::endl;
+ return false;
+ }
+
+ return true;
+}
+
+void render(void)
+{
+ // check to make sure you have a valid Direct3D device
+ CV_Assert(dev);
+
+ renderToD3DObject();
+
+ if (g_sampleType == 0)
+ {
+ // nothing
+ }
+ else if (g_sampleType == 1)
+ {
+ if (FAILED(dev->StretchRect(pCPUWriteSurface, NULL, pBackBuffer, NULL, D3DTEXF_NONE)))
+ {
+ std::cerr << "Can't StretchRect()" << std::endl;
+ }
+ }
+ else
+ {
+ if (FAILED(dev->StretchRect(pSurface, NULL, pBackBuffer, NULL, D3DTEXF_NONE)))
+ {
+ std::cerr << "Can't StretchRect()" << std::endl;
+ }
+ }
+
+ if (SUCCEEDED(dev -> BeginScene()))
+ {
+ // end the scene
+ dev -> EndScene();
+ }
+
+ // present the back buffer contents to the display
+ dev->Present(NULL, NULL, NULL, NULL);
+}
+
+void cleanUp (void)
+{
+ SAFE_RELEASE(pCPUWriteSurface);
+ SAFE_RELEASE(pReadOnlySurface);
+ SAFE_RELEASE(pSurface);
+ SAFE_RELEASE(pBackBuffer);
+ SAFE_RELEASE(dev);
+ SAFE_RELEASE(pD3D);}
--- /dev/null
+#include <windows.h>
+#include <d3d9.h>
+#pragma comment (lib, "d3d9.lib")
+
+#define USE_D3DEX
+#define WINDOW_NAME "OpenCV Direct3D 9 Ex Sample"
+
+IDirect3D9Ex *pD3D = NULL;
+IDirect3DDevice9Ex *dev = NULL;
+IDirect3DSurface9 *pBackBuffer = NULL;
+IDirect3DSurface9 *pCPUWriteSurface = NULL; // required name
+IDirect3DSurface9 *pReadOnlySurface = NULL; // required name
+HANDLE readOnlySurfaceShared = 0; // required name
+IDirect3DSurface9 *pSurface = NULL; // required name
+HANDLE surfaceShared = 0; // required name
+
+#include "d3d_base.inl.hpp"
+
+bool initDirect3D(void)
+{
+ if (FAILED(Direct3DCreate9Ex(D3D_SDK_VERSION, &pD3D)))
+ {
+ return false;
+ }
+
+ D3DDISPLAYMODEEX ddm;
+ ZeroMemory(&ddm, sizeof(ddm));
+ ddm.Size = sizeof(D3DDISPLAYMODEEX);
+ D3DDISPLAYROTATION rotation;
+ if (FAILED(pD3D->GetAdapterDisplayModeEx(D3DADAPTER_DEFAULT, &ddm, &rotation)))
+ {
+ return false;
+ }
+
+ D3DPRESENT_PARAMETERS d3dpp;
+ ZeroMemory(&d3dpp,sizeof(D3DPRESENT_PARAMETERS));
+
+ DWORD flags = D3DCREATE_HARDWARE_VERTEXPROCESSING | D3DCREATE_PUREDEVICE | D3DCREATE_NOWINDOWCHANGES
+ | D3DCREATE_MULTITHREADED;
+
+ d3dpp.Windowed = true;
+ d3dpp.Flags = 0;
+ d3dpp.BackBufferCount = 0;
+ d3dpp.BackBufferFormat = ddm.Format;
+ d3dpp.BackBufferHeight = HEIGHT;
+ d3dpp.BackBufferWidth = WIDTH;
+ d3dpp.MultiSampleType = D3DMULTISAMPLE_NONE;
+ d3dpp.SwapEffect = D3DSWAPEFFECT_DISCARD;
+ d3dpp.hDeviceWindow = hWnd;
+ d3dpp.PresentationInterval = D3DPRESENT_INTERVAL_IMMEDIATE;
+ d3dpp.FullScreen_RefreshRateInHz = D3DPRESENT_RATE_DEFAULT;
+
+ if (FAILED(pD3D->CreateDeviceEx(D3DADAPTER_DEFAULT, D3DDEVTYPE_HAL, hWnd, flags, &d3dpp, NULL, &dev)))
+ {
+ return false;
+ }
+
+ if (FAILED(dev->GetBackBuffer(0, 0, D3DBACKBUFFER_TYPE_MONO, &pBackBuffer)))
+ {
+ return false;
+ }
+
+ return true;
+}
+
+bool initDirect3DTextures()
+{
+ if (FAILED(dev->CreateOffscreenPlainSurface(WIDTH, HEIGHT, D3DFMT_A8R8G8B8, D3DPOOL_DEFAULT, &pSurface, &surfaceShared)))
+ {
+ std::cerr << "Can't create surface for result" << std::endl;
+ return false;
+ }
+
+ if (FAILED(dev->CreateOffscreenPlainSurface(WIDTH, HEIGHT, D3DFMT_A8R8G8B8, D3DPOOL_DEFAULT, &pReadOnlySurface, &readOnlySurfaceShared)))
+ {
+ std::cerr << "Can't create read only surface" << std::endl;
+ return false;
+ }
+ else
+ {
+ IDirect3DSurface9* pTmpSurface;
+ if (FAILED(dev->CreateOffscreenPlainSurface(WIDTH, HEIGHT, D3DFMT_A8R8G8B8, D3DPOOL_DEFAULT, &pTmpSurface, NULL)))
+ {
+ std::cerr << "Can't create temp surface for CPU write" << std::endl;
+ return false;
+ }
+
+ D3DLOCKED_RECT memDesc = {0, NULL};
+ RECT rc = {0, 0, WIDTH, HEIGHT};
+ if (SUCCEEDED(pTmpSurface->LockRect(&memDesc, &rc, 0)))
+ {
+ cv::Mat m(cv::Size(WIDTH, HEIGHT), CV_8UC4, memDesc.pBits, (int)memDesc.Pitch);
+ getInputTexture().copyTo(m);
+ pTmpSurface->UnlockRect();
+ dev->StretchRect(pTmpSurface, NULL, pReadOnlySurface, NULL, D3DTEXF_NONE);
+ }
+ else
+ {
+ std::cerr << "Can't LockRect() on surface" << std::endl;
+ }
+ pTmpSurface->Release();
+ }
+
+ if (FAILED(dev->CreateOffscreenPlainSurface(WIDTH, HEIGHT, D3DFMT_A8R8G8B8, D3DPOOL_DEFAULT, &pCPUWriteSurface, NULL)))
+ {
+ std::cerr << "Can't create surface for CPU write" << std::endl;
+ return false;
+ }
+
+ return true;
+}
+
+
+void render(void)
+{
+ // check to make sure you have a valid Direct3D device
+ CV_Assert(dev);
+
+ renderToD3DObject();
+
+ if (g_sampleType == 0)
+ {
+ // nothing
+ }
+ else if (g_sampleType == 1)
+ {
+ if (FAILED(dev->StretchRect(pCPUWriteSurface, NULL, pBackBuffer, NULL, D3DTEXF_NONE)))
+ {
+ std::cerr << "Can't StretchRect()" << std::endl;
+ }
+ }
+ else
+ {
+ if (FAILED(dev->StretchRect(pSurface, NULL, pBackBuffer, NULL, D3DTEXF_NONE)))
+ {
+ std::cerr << "Can't StretchRect()" << std::endl;
+ }
+ }
+
+ if (SUCCEEDED(dev -> BeginScene()))
+ {
+ // end the scene
+ dev -> EndScene();
+ }
+
+ // present the back buffer contents to the display
+ dev->Present(NULL, NULL, NULL, NULL);
+}
+
+void cleanUp (void)
+{
+ SAFE_RELEASE(pCPUWriteSurface);
+ SAFE_RELEASE(pReadOnlySurface);
+ SAFE_RELEASE(pSurface);
+ SAFE_RELEASE(pBackBuffer);
+ SAFE_RELEASE(dev);
+ SAFE_RELEASE(pD3D);
+}
--- /dev/null
+//
+// Don't use as a standalone file
+//
+
+#include "opencv2/core.hpp"
+#include "opencv2/core/utility.hpp" // cv::format
+#include "opencv2/imgproc.hpp" // cvtColor
+#include "opencv2/imgproc/types_c.h" // cvtColor
+#include "opencv2/highgui.hpp" // imread
+#include "opencv2/core/directx.hpp"
+
+#include <iostream>
+#include <queue>
+
+using namespace cv;
+using namespace cv::directx;
+static const int fontFace = cv::FONT_HERSHEY_DUPLEX;
+#if !defined(USE_D3D9) && !defined(USE_D3DEX)
+const cv::Scalar frameColor(255,128,0,255);
+#else
+const cv::Scalar frameColor(0,128,255,255); // BGRA for D3D9
+#endif
+
+#define SAFE_RELEASE(p) if (p) { p->Release(); p = NULL; }
+
+const int WIDTH = 1024;
+const int HEIGHT = 768;
+
+HINSTANCE hInstance;
+HWND hWnd;
+
+// external declaration
+bool initDirect3D(void);
+bool initDirect3DTextures(void);
+void render(void);
+void cleanUp (void);
+
+#define USAGE_DESCRIPTION_0 "1 - CPU write via LockRect/Map"
+#define USAGE_DESCRIPTION_1 "2* - Mat->D3D"
+#define USAGE_DESCRIPTION_2 "3* - D3D->UMat / change UMat / UMat->D3D"
+#define USAGE_DESCRIPTION_3 "0 - show input texture without any processing"
+#define USAGE_DESCRIPTION_SPACE "SPACE - toggle frame processing (only data transfers)"
+
+static int g_sampleType = 0;
+static int g_disableProcessing = false;
+
+// forward declaration
+static LRESULT CALLBACK WndProc(HWND, UINT, WPARAM, LPARAM);
+
+static bool initWindow()
+{
+ WNDCLASSEX wcex;
+
+ wcex.cbSize = sizeof(WNDCLASSEX);
+ wcex.style = CS_HREDRAW | CS_VREDRAW;
+ wcex.lpfnWndProc = WndProc;
+ wcex.cbClsExtra = 0;
+ wcex.cbWndExtra = 0;
+ wcex.hInstance = hInstance;
+ wcex.hIcon = LoadIcon(0, IDI_APPLICATION);
+ wcex.hCursor = LoadCursor(0, IDC_ARROW);
+ wcex.hbrBackground = 0;
+ wcex.lpszMenuName = 0L;
+ wcex.lpszClassName = "OpenCVDirectX";
+ wcex.hIconSm = 0;
+
+ RegisterClassEx(&wcex);
+
+ RECT rc = {0, 0, WIDTH, HEIGHT};
+ AdjustWindowRect(&rc, WS_OVERLAPPEDWINDOW, false);
+ hWnd = CreateWindow("OpenCVDirectX", WINDOW_NAME,
+ WS_OVERLAPPEDWINDOW, CW_USEDEFAULT, CW_USEDEFAULT, rc.right - rc.left, rc.bottom - rc.top, NULL, NULL, hInstance, NULL);
+
+ if (!hWnd)
+ return false;
+
+ ShowWindow(hWnd, SW_SHOW);
+ UpdateWindow(hWnd);
+
+ return true;
+}
+
+static LRESULT CALLBACK WndProc(HWND hWnd, UINT message, WPARAM wParam, LPARAM lParam)
+{
+ switch (message)
+ {
+ case WM_DESTROY:
+ PostQuitMessage(0);
+ return 0;
+
+ case WM_CHAR:
+ if (wParam >= '0' && wParam <= '3')
+ {
+ g_sampleType = (char)wParam - '0';
+ return 0;
+ }
+ else if (wParam == ' ')
+ {
+ g_disableProcessing = !g_disableProcessing;
+ return 0;
+ }
+ else if (wParam == VK_ESCAPE)
+ {
+ DestroyWindow(hWnd);
+ return 0;
+ }
+ break;
+ }
+ return DefWindowProc(hWnd, message, wParam, lParam);
+}
+
+static float getFps()
+{
+ static std::queue<int64> time_queue;
+
+ int64 now = cv::getTickCount(), then = 0;
+ time_queue.push(now);
+
+ if (time_queue.size() >= 2)
+ then = time_queue.front();
+
+ if (time_queue.size() >= 25)
+ time_queue.pop();
+
+ return time_queue.size() * (float)cv::getTickFrequency() / (now - then);
+}
+
+static int bgColor[4] = {0, 0, 0, 0};
+static cv::Mat* inputMat = NULL;
+
+static void renderToD3DObject(void)
+{
+ static int frame = 0;
+
+ const float fps = getFps();
+
+ String deviceName = cv::ocl::useOpenCL() ? cv::ocl::Context2::getDefault().device(0).name() : "No OpenCL device";
+
+ if ((frame % std::max(1, (int)(fps / 25))) == 0)
+ {
+ String msg = format("%s%s: %s, Sample %d, Frame %d, fps %g (%g ms)",
+ g_disableProcessing ? "(FRAME PROCESSING DISABLED) " : "",
+ WINDOW_NAME, deviceName.c_str(), g_sampleType,
+ frame, fps, (int(10 * 1000.0 / fps)) * 0.1);
+ SetWindowText(hWnd, msg.c_str());
+ }
+
+ // 0..255
+ int c[4] =
+ {
+ std::abs((frame & 0x1ff) - 0x100),
+ std::abs(((frame * 2) & 0x1ff) - 0x100),
+ std::abs(((frame / 2) & 0x1ff) - 0x100),
+ 0
+ };
+
+ int c1 = c[0] / 2 - 0x40 - bgColor[0];
+ int c2 = c[1] / 2 - 0x40 - bgColor[1];
+ int c3 = c[2] / 2 - 0x40 - bgColor[2];
+
+ switch (g_sampleType)
+ {
+ case 0:
+#if defined(USE_D3D9) || defined (USE_D3DEX)
+ if (FAILED(dev->StretchRect(pReadOnlySurface, NULL, pBackBuffer, NULL, D3DTEXF_NONE)))
+ {
+ std::cerr << "Can't StretchRect()" << std::endl;
+ }
+#elif defined(USE_D3D10)
+ dev->CopyResource(pBackBufferTexture, pInputTexture);
+#elif defined(USE_D3D11)
+ devcon->CopyResource(pBackBufferTexture, pInputTexture);
+#else
+#error "Invalid USE_D3D value"
+#endif
+ break;
+
+ case 1:
+ {
+ int BOXSIZE = 50;
+ int x = std::abs(((frame * 1) % (2 * (WIDTH - BOXSIZE))) - (WIDTH - BOXSIZE));
+ int y = std::abs(((frame / 2) % (2 * (HEIGHT - BOXSIZE))) - (HEIGHT - BOXSIZE));
+ cv::Rect boxRect(x, y, BOXSIZE, BOXSIZE);
+#if defined(USE_D3D9) || defined (USE_D3DEX)
+ D3DLOCKED_RECT memDesc = {0, NULL};
+ RECT rc = {0, 0, WIDTH, HEIGHT};
+ if (SUCCEEDED(pCPUWriteSurface->LockRect(&memDesc, &rc, 0)))
+ {
+ if (!g_disableProcessing)
+ {
+ Mat m(Size(WIDTH, HEIGHT), CV_8UC4, memDesc.pBits, (int)memDesc.Pitch);
+ inputMat->copyTo(m);
+ m(boxRect).setTo(Scalar(c[0], c[1], c[2], 255));
+ }
+ pCPUWriteSurface->UnlockRect();
+ }
+ else
+ {
+ std::cerr << "Can't LockRect() on surface" << std::endl;
+ }
+#elif defined(USE_D3D10)
+ D3D10_MAPPED_TEXTURE2D mappedTex;
+ if (SUCCEEDED(pCPUWriteTexture->Map( D3D10CalcSubresource(0, 0, 1), D3D10_MAP_WRITE_DISCARD, 0, &mappedTex)))
+ {
+ if (!g_disableProcessing)
+ {
+ Mat m(Size(WIDTH, HEIGHT), CV_8UC4, mappedTex.pData, (int)mappedTex.RowPitch);
+ inputMat->copyTo(m);
+ m(boxRect).setTo(Scalar(c[0], c[1], c[2], 255));
+ }
+ pCPUWriteTexture->Unmap(D3D10CalcSubresource(0, 0, 1));
+ dev->CopyResource(pBackBufferTexture, pCPUWriteTexture);
+ }
+ else
+ {
+ std::cerr << "Can't Map() texture" << std::endl;
+ }
+#elif defined(USE_D3D11)
+ D3D11_MAPPED_SUBRESOURCE mappedTex;
+ if (SUCCEEDED(devcon->Map(pCPUWriteTexture, D3D11CalcSubresource(0, 0, 1), D3D11_MAP_WRITE_DISCARD, 0, &mappedTex)))
+ {
+ if (!g_disableProcessing)
+ {
+ Mat m(Size(WIDTH, HEIGHT), CV_8UC4, mappedTex.pData, (int)mappedTex.RowPitch);
+ inputMat->copyTo(m);
+ m(boxRect).setTo(Scalar(c[0], c[1], c[2], 255));
+ }
+ devcon->Unmap(pCPUWriteTexture, D3D11CalcSubresource(0, 0, 1));
+ devcon->CopyResource(pBackBufferTexture, pCPUWriteTexture);
+ }
+ else
+ {
+ std::cerr << "Can't Map() texture" << std::endl;
+ }
+#else
+#error "Invalid USE_D3D value"
+#endif
+ break;
+ }
+ case 2:
+ {
+ static Mat m;
+ if (!g_disableProcessing)
+ {
+#if 1
+ cv::add(*inputMat, Scalar(c1, c2, c3, 255), m);
+#else
+ inputMat->copyTo(m);
+#endif
+ cv::putText(m,
+ cv::format("Frame %d, fps %g (%g ms)",
+ frame, fps, (int(10 * 1000.0 / fps)) * 0.1),
+ cv::Point(8, 80), fontFace, 1, frameColor, 2);
+ }
+ else
+ {
+ m.create(Size(WIDTH, HEIGHT), CV_8UC4);
+ }
+ try
+ {
+ #if defined(USE_D3D9) || defined (USE_D3DEX)
+ convertToDirect3DSurface9(m, pSurface, (void*)surfaceShared);
+ #elif defined(USE_D3D10)
+ convertToD3D10Texture2D(m, pBackBufferTexture);
+ #elif defined(USE_D3D11)
+ convertToD3D11Texture2D(m, pBackBufferTexture);
+ #else
+ #error "Invalid USE_D3D value"
+ #endif
+ }
+ catch (cv::Exception& e)
+ {
+ std::cerr << "Can't convert to D3D object: exception: " << e.what() << std::endl;
+ }
+ catch (...)
+ {
+ std::cerr << "Can't convert to D3D object" << std::endl;
+ }
+ break;
+ }
+ case 3:
+ {
+ static UMat tmp;
+ try
+ {
+#if defined(USE_D3D9) || defined (USE_D3DEX)
+ convertFromDirect3DSurface9(pReadOnlySurface, tmp, (void*)readOnlySurfaceShared);
+#elif defined(USE_D3D10)
+ convertFromD3D10Texture2D(pInputTexture, tmp);
+#elif defined(USE_D3D11)
+ convertFromD3D11Texture2D(pInputTexture, tmp);
+#else
+#error "Invalid USE_D3D value"
+#endif
+ }
+ catch (cv::Exception& e)
+ {
+ std::cerr << "Can't convert from D3D object: exception: " << e.what() << std::endl;
+ }
+ catch (...)
+ {
+ std::cerr << "Can't convert from D3D object" << std::endl;
+ }
+ static UMat res;
+ if (!g_disableProcessing)
+ {
+ cv::add(tmp, Scalar(c1, c2, c3, 255), res);
+ }
+ else
+ {
+ res = tmp;
+ }
+ try
+ {
+#if defined(USE_D3D9) || defined (USE_D3DEX)
+ convertToDirect3DSurface9(res, pSurface, (void*)surfaceShared);
+#elif defined(USE_D3D10)
+ convertToD3D10Texture2D(res, pBackBufferTexture);
+#elif defined(USE_D3D11)
+ convertToD3D11Texture2D(res, pBackBufferTexture);
+#else
+#error "Invalid USE_D3D value"
+#endif
+ }
+ catch (cv::Exception& e)
+ {
+ std::cerr << "Can't convert to D3D object: exception: " << e.what() << std::endl;
+ }
+ catch (...)
+ {
+ std::cerr << "Can't convert to D3D object" << std::endl;
+ }
+ break;
+ }
+ }
+ frame++;
+}
+
+
+static cv::Mat getInputTexture()
+{
+ cv::Mat inputMat = cv::imread("input.bmp", cv::IMREAD_COLOR);
+
+ if (inputMat.depth() != CV_8U)
+ {
+ inputMat.convertTo(inputMat, CV_8U);
+ }
+ if (inputMat.type() == CV_8UC3)
+ {
+ cv::cvtColor(inputMat, inputMat, CV_RGB2BGRA);
+ }
+ if (inputMat.type() != CV_8UC4 || inputMat.size().area() == 0)
+ {
+ std::cerr << "Invalid input image format. Generate other" << std::endl;
+ inputMat.create(cv::Size(WIDTH, HEIGHT), CV_8UC4);
+ inputMat.setTo(cv::Scalar(0, 0, 255, 255));
+ bgColor[0] = -128; bgColor[1] = -128; bgColor[2] = 127; bgColor[3] = -128;
+ }
+ if (inputMat.size().width != WIDTH || inputMat.size().height != HEIGHT)
+ {
+ cv::resize(inputMat, inputMat, cv::Size(WIDTH, HEIGHT));
+ }
+ String deviceName = cv::ocl::useOpenCL() ? cv::ocl::Context2::getDefault().device(0).name() : "No OpenCL device";
+ cv::Scalar color(64, 255, 64, 255);
+ cv::putText(inputMat,
+ cv::format("OpenCL Device name: %s", deviceName.c_str()),
+ cv::Point(8,32), fontFace, 1, color);
+ cv::putText(inputMat, WINDOW_NAME, cv::Point(50, HEIGHT - 32), fontFace, 1, color);
+ cv::putText(inputMat, USAGE_DESCRIPTION_0, cv::Point(30, 128), fontFace, 1, color);
+ cv::putText(inputMat, USAGE_DESCRIPTION_1, cv::Point(30, 192), fontFace, 1, color);
+ cv::putText(inputMat, USAGE_DESCRIPTION_2, cv::Point(30, 256), fontFace, 1, color);
+ cv::putText(inputMat, USAGE_DESCRIPTION_3, cv::Point(30, 320), fontFace, 1, color);
+ cv::putText(inputMat, USAGE_DESCRIPTION_SPACE, cv::Point(30, 448), fontFace, 1, color);
+
+#if defined(USE_D3D9) || defined (USE_D3DEX)
+ cv::cvtColor(inputMat, inputMat, CV_RGBA2BGRA);
+ std::swap(bgColor[0], bgColor[2]);
+#endif
+
+ // Make a global copy
+ ::inputMat = new cv::Mat(inputMat);
+
+ return inputMat;
+}
+
+static int mainLoop()
+{
+ hInstance = GetModuleHandle(NULL);
+
+ if (!initWindow())
+ CV_Error(cv::Error::StsError, "Can't create window");
+
+ if (!initDirect3D())
+ CV_Error(cv::Error::StsError, "Can't create D3D object");
+
+ if (cv::ocl::haveOpenCL())
+ {
+#if defined(USE_D3D9)
+ cv::ocl::Context2& ctx = cv::directx::ocl::initializeContextFromDirect3DDevice9(dev);
+#elif defined (USE_D3DEX)
+ cv::ocl::Context2& ctx = cv::directx::ocl::initializeContextFromDirect3DDevice9Ex(dev);
+#elif defined(USE_D3D10)
+ cv::ocl::Context2& ctx = cv::directx::ocl::initializeContextFromD3D10Device(dev);
+#elif defined(USE_D3D11)
+ cv::ocl::Context2& ctx = cv::directx::ocl::initializeContextFromD3D11Device(dev);
+#else
+#error "Invalid USE_D3D value"
+#endif
+ std::cout << "Selected device: " << ctx.device(0).name().c_str() << std::endl;
+ g_sampleType = 2;
+ }
+ else
+ {
+ std::cerr << "OpenCL is not available. DirectX - OpenCL interop will not work" << std::endl;
+ }
+
+ if (!initDirect3DTextures())
+ CV_Error(cv::Error::StsError, "Can't create D3D texture object");
+
+ MSG msg;
+ ZeroMemory(&msg, sizeof(msg));
+
+ while (msg.message != WM_QUIT)
+ {
+ if (PeekMessage(&msg, NULL, 0U, 0U, PM_REMOVE))
+ {
+ TranslateMessage(&msg);
+ DispatchMessage(&msg);
+ }
+ else
+ {
+ render();
+ }
+ }
+
+ cleanUp();
+
+ return static_cast<int>(msg.wParam);
+}
+
+int main(int /*argc*/, char ** /*argv*/)
+{
+ try
+ {
+ return mainLoop();
+ }
+ catch (cv::Exception& e)
+ {
+ std::cerr << "Exception: " << e.what() << std::endl;
+ return 10;
+ }
+ catch (...)
+ {
+ std::cerr << "FATAL ERROR: Unknown exception" << std::endl;
+ return 11;
+ }
+}
cout << "FOUND " << keypoints2GPU.cols << " keypoints on second image" << endl;
// matching descriptors
- BFMatcher_CUDA matcher(NORM_L2);
+ BFMatcher_CUDA matcher(surf.defaultNorm());
GpuMat trainIdx, distance;
matcher.matchSingle(descriptors1GPU, descriptors2GPU, trainIdx, distance);
--- /dev/null
+/target
+/classes
+/checkouts
+pom.xml
+pom.xml.asc
+*.jar
+*.class
+/.lein-*
+/.nrepl-port
--- /dev/null
+(defproject simple-sample "0.1.0-SNAPSHOT"
+ :pom-addition [:developers [:developer {:id "magomimmo"}
+ [:name "Mimmo Cosenza"]
+ [:url "https://github.com/magomimmoo"]]]
+
+ :description "A simple project to start REPLing with OpenCV"
+ :url "http://example.com/FIXME"
+ :license {:name "BSD 3-Clause License"
+ :url "http://opensource.org/licenses/BSD-3-Clause"}
+ :dependencies [[org.clojure/clojure "1.5.1"]
+ [opencv/opencv "2.4.7"]
+ [opencv/opencv-native "2.4.7"]]
+ :main simple-sample.core
+ :injections [(clojure.lang.RT/loadLibrary org.opencv.core.Core/NATIVE_LIBRARY_NAME)])
--- /dev/null
+;;; to run this code from the terminal: "$ lein run". It will save a
+;;; blurred image version of resources/images/lena.png as
+;;; resources/images/blurred.png
+
+(ns simple-sample.core
+ (:import [org.opencv.core Point Rect Mat CvType Size Scalar]
+ org.opencv.highgui.Highgui
+ org.opencv.imgproc.Imgproc))
+
+(defn -main [& args]
+ (let [lena (Highgui/imread "resources/images/lena.png")
+ blurred (Mat. 512 512 CvType/CV_8UC3)]
+ (print "Blurring...")
+ (Imgproc/GaussianBlur lena blurred (Size. 5 5) 3 3)
+ (Highgui/imwrite "resources/images/blurred.png" blurred)
+ (println "done!")))
--- /dev/null
+(ns simple-sample.core-test
+ (:require [clojure.test :refer :all]
+ [simple-sample.core :refer :all]))
+
+(deftest a-test
+ (testing "FIXME, I fail."
+ (is (= 0 1))))
projects / platform / upstream / opencv.git / commitdiff
