mirror of https://github.com/opencv/opencv.git
Andrey Kamaev
12 years ago
commit
f77a375354
291 changed files with 5007 additions and 2222 deletions
|
After Width: | Height: | Size: 18 KiB |
|
After Width: | Height: | Size: 4.0 KiB |
|
After Width: | Height: | Size: 20 KiB |
|
After Width: | Height: | Size: 59 KiB |
|
After Width: | Height: | Size: 22 KiB |
|
After Width: | Height: | Size: 120 KiB |
|
After Width: | Height: | Size: 78 KiB |
|
After Width: | Height: | Size: 29 KiB |
|
After Width: | Height: | Size: 18 KiB |
|
After Width: | Height: | Size: 4.8 KiB |
@ -0,0 +1,144 @@ |
||||
.. _Windows_Visual_Studio_Image_Watch: |
||||
|
||||
Image Watch: viewing in-memory images in the Visual Studio debugger |
||||
******************************************************************* |
||||
|
||||
Image Watch is a plug-in for Microsoft Visual Studio that lets you to visualize in-memory images (*cv::Mat* or *IplImage\_* objects, for example) while debugging an application. This can be helpful for tracking down bugs, or for simply understanding what a given piece of code is doing. |
||||
|
||||
Prerequisites |
||||
============= |
||||
|
||||
This tutorial assumes that you have the following available: |
||||
|
||||
#. Visual Studio 2012 Professional (or better) with Update 1 installed. Update 1 can be downloaded `here <http://www.microsoft.com/en-us/download/details.aspx?id=35774>`_. |
||||
|
||||
#. An OpenCV installation on your Windows machine (Tutorial: :ref:`Windows_Installation`). |
||||
|
||||
#. Ability to create and build OpenCV projects in Visual Studio (Tutorial: :ref:`Windows_Visual_Studio_How_To`). |
||||
|
||||
Installation |
||||
============ |
||||
|
||||
`Download <http://go.microsoft.com/fwlink/?LinkId=285460>`_ the Image Watch installer. The installer comes in a single file with extension .vsix (*Visual Studio Extension*). To launch it, simply double-click on the .vsix file in Windows Explorer. When the installer has finished, make sure to restart Visual Studio to complete the installation. |
||||
|
||||
Example |
||||
======== |
||||
|
||||
Image Watch works with any existing project that uses OpenCV image objects (for example, *cv::Mat*). In this example, we use a minimal test program that loads an image from a file and runs an edge detector. To build the program, create a console application project in Visual Studio, name it "image-watch-demo", and insert the source code below. |
||||
|
||||
.. code-block:: c++ |
||||
|
||||
// Test application for the Visual Studio Image Watch Debugger extension |
||||
|
||||
#include <iostream> // std::cout |
||||
#include <opencv2/core/core.hpp> // cv::Mat |
||||
#include <opencv2/highgui/highgui.hpp> // cv::imread() |
||||
#include <opencv2/imgproc/imgproc.hpp> // cv::Canny() |
||||
|
||||
using namespace std; |
||||
using namespace cv; |
||||
|
||||
void help() |
||||
{ |
||||
cout |
||||
<< "----------------------------------------------------" << endl |
||||
<< "This is a test program for the Image Watch Debugger " << endl |
||||
<< "plug-in for Visual Studio. The program loads an " << endl |
||||
<< "image from a file and runs the Canny edge detector. " << endl |
||||
<< "No output is displayed or written to disk." |
||||
<< endl |
||||
<< "Usage:" << endl |
||||
<< "image-watch-demo inputimage" << endl |
||||
<< "----------------------------------------------------" << endl |
||||
<< endl; |
||||
} |
||||
|
||||
int main(int argc, char *argv[]) |
||||
{ |
||||
help(); |
||||
|
||||
if (argc != 2) |
||||
{ |
||||
cout << "Wrong number of parameters" << endl; |
||||
return -1; |
||||
} |
||||
|
||||
cout << "Loading input image: " << argv[1] << endl; |
||||
Mat input; |
||||
input = imread(argv[1], CV_LOAD_IMAGE_COLOR); |
||||
|
||||
cout << "Detecting edges in input image" << endl; |
||||
Mat edges; |
||||
Canny(input, edges, 10, 100); |
||||
|
||||
return 0; |
||||
} |
||||
|
||||
Make sure your active solution configuration (:menuselection:`Build --> Configuration Manager`) is set to a debug build (usually called "Debug"). This should disable compiler optimizations so that viewing variables in the debugger can work reliably. |
||||
|
||||
Build your solution (:menuselection:`Build --> Build Solution`, or press *F7*). |
||||
|
||||
Now set a breakpoint on the source line that says |
||||
|
||||
.. code-block:: c++ |
||||
|
||||
Mat edges; |
||||
|
||||
To set the breakpoint, right-click on the source line and select :menuselection:`Breakpoints --> Insert Breakpoint` from the context menu. |
||||
|
||||
Launch the program in the debugger (:menuselection:`Debug --> Start Debugging`, or hit *F5*). When the breakpoint is hit, the program is paused and Visual Studio displays a yellow instruction pointer at the breakpoint: |
||||
|
||||
.. image:: images/breakpoint.png |
||||
|
||||
Now you can inspect the state of you program. For example, you can bring up the *Locals* window (:menuselection:`Debug --> Windows --> Locals`), which will show the names and values of the variables in the current scope: |
||||
|
||||
.. image:: images/vs_locals.png |
||||
|
||||
Note that the built-in *Locals* window will display text only. This is where the Image Watch plug-in comes in. Image Watch is like another *Locals* window, but with an image viewer built into it. To bring up Image Watch, select :menuselection:`View --> Other Windows --> Image Watch`. Like Visual Studio's *Locals* window, Image Watch can dock to the Visual Studio IDE. Also, Visual Studio will remember whether you had Image Watch open, and where it was located between debugging sessions. This means you only have to do this once--the next time you start debugging, Image Watch will be back where you left it. Here's what the docked Image Watch window looks like at our breakpoint: |
||||
|
||||
.. image:: images/toolwindow.jpg |
||||
:height: 320pt |
||||
|
||||
The radio button at the top left (*Locals/Watch*) selects what is shown in the *Image List* below: *Locals* lists all OpenCV image objects in the current scope (this list is automatically populated). *Watch* shows image expressions that have been pinned for continuous inspection (not described here, see `Image Watch documentation <http://go.microsoft.com/fwlink/?LinkId=285461>`_ for details). The image list shows basic information such as width, height, number of channels, and, if available, a thumbnail. In our example, the image list contains our two local image variables, *input* and *edges*. |
||||
|
||||
If an image has a thumbnail, left-clicking on that image will select it for detailed viewing in the *Image Viewer* on the right. The viewer lets you pan (drag mouse) and zoom (mouse wheel). It also displays the pixel coordinate and value at the current mouse position. |
||||
|
||||
.. image:: images/viewer.jpg |
||||
:height: 160pt |
||||
|
||||
Note that the second image in the list, *edges*, is shown as "invalid". This indicates that some data members of this image object have corrupt or invalid values (for example, a negative image width). This is expected at this point in the program, since the C++ constructor for *edges* has not run yet, and so its members have undefined values (in debug mode they are usually filled with "0xCD" bytes). |
||||
|
||||
From here you can single-step through your code (:menuselection:`Debug->Step Over`, or press *F10*) and watch the pixels change: if you step once, over the *Mat edges;* statement, the *edges* image will change from "invalid" to "empty", which means that it is now in a valid state (default constructed), even though it has not been initialized yet (using *cv::Mat::create()*, for example). If you make one more step over the *cv::Canny()* call, you will see a thumbnail of the edge image appear in the image list. |
||||
|
||||
Now assume you want to do a visual sanity check of the *cv::Canny()* implementation. Bring the *edges* image into the viewer by selecting it in the *Image List* and zoom into a region with a clearly defined edge: |
||||
|
||||
.. image:: images/edges_zoom.png |
||||
:height: 160pt |
||||
|
||||
Right-click on the *Image Viewer* to bring up the view context menu and enable :menuselection:`Link Views` (a check box next to the menu item indicates whether the option is enabled). |
||||
|
||||
.. image:: images/viewer_context_menu.png |
||||
:height: 120pt |
||||
|
||||
The :menuselection:`Link Views` feature keeps the view region fixed when flipping between images of the same size. To see how this works, select the input image from the image list--you should now see the corresponding zoomed-in region in the input image: |
||||
|
||||
.. image:: images/input_zoom.png |
||||
:height: 160pt |
||||
|
||||
You may also switch back and forth between viewing input and edges with your up/down cursor keys. That way you can easily verify that the detected edges line up nicely with the data in the input image. |
||||
|
||||
More ... |
||||
==================== |
||||
|
||||
Image watch has a number of more advanced features, such as |
||||
|
||||
#. pinning images to a *Watch* list for inspection across scopes or between debugging sessions |
||||
|
||||
#. clamping, thresholding, or diff'ing images directly inside the Watch window |
||||
|
||||
#. comparing an in-memory image against a reference image from a file |
||||
|
||||
Please refer to the online `Image Watch Documentation <http://go.microsoft.com/fwlink/?LinkId=285461>`_ for details--you also can get to the documentation page by clicking on the *Help* link in the Image Watch window: |
||||
|
||||
.. image:: images/help_button.jpg |
||||
:height: 80pt |
||||
@ -1 +1,43 @@ |
||||
/*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) 2000-2008, Intel Corporation, all rights reserved.
|
||||
// Copyright (C) 2009, Willow Garage 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 "perf_precomp.hpp" |
||||
|
||||
@ -0,0 +1,186 @@ |
||||
/*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) 2000-2008, Intel Corporation, all rights reserved. |
||||
// Copyright (C) 2009, Willow Garage 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*/ |
||||
|
||||
#if !defined CUDA_DISABLER |
||||
|
||||
#include "opencv2/gpu/device/common.hpp" |
||||
#include "opencv2/gpu/device/functional.hpp" |
||||
#include "opencv2/gpu/device/emulation.hpp" |
||||
#include "opencv2/gpu/device/scan.hpp" |
||||
#include "opencv2/gpu/device/reduce.hpp" |
||||
#include "opencv2/gpu/device/saturate_cast.hpp" |
||||
|
||||
using namespace cv::gpu; |
||||
using namespace cv::gpu::device; |
||||
|
||||
namespace clahe |
||||
{ |
||||
__global__ void calcLutKernel(const PtrStepb src, PtrStepb lut, |
||||
const int2 tileSize, const int tilesX, |
||||
const int clipLimit, const float lutScale) |
||||
{ |
||||
__shared__ int smem[512]; |
||||
|
||||
const int tx = blockIdx.x; |
||||
const int ty = blockIdx.y; |
||||
const unsigned int tid = threadIdx.y * blockDim.x + threadIdx.x; |
||||
|
||||
smem[tid] = 0; |
||||
__syncthreads(); |
||||
|
||||
for (int i = threadIdx.y; i < tileSize.y; i += blockDim.y) |
||||
{ |
||||
const uchar* srcPtr = src.ptr(ty * tileSize.y + i) + tx * tileSize.x; |
||||
for (int j = threadIdx.x; j < tileSize.x; j += blockDim.x) |
||||
{ |
||||
const int data = srcPtr[j]; |
||||
Emulation::smem::atomicAdd(&smem[data], 1); |
||||
} |
||||
} |
||||
|
||||
__syncthreads(); |
||||
|
||||
int tHistVal = smem[tid]; |
||||
|
||||
__syncthreads(); |
||||
|
||||
if (clipLimit > 0) |
||||
{ |
||||
// clip histogram bar |
||||
|
||||
int clipped = 0; |
||||
if (tHistVal > clipLimit) |
||||
{ |
||||
clipped = tHistVal - clipLimit; |
||||
tHistVal = clipLimit; |
||||
} |
||||
|
||||
// find number of overall clipped samples |
||||
|
||||
reduce<256>(smem, clipped, tid, plus<int>()); |
||||
|
||||
// broadcast evaluated value |
||||
|
||||
__shared__ int totalClipped; |
||||
|
||||
if (tid == 0) |
||||
totalClipped = clipped; |
||||
__syncthreads(); |
||||
|
||||
// redistribute clipped samples evenly |
||||
|
||||
int redistBatch = totalClipped / 256; |
||||
tHistVal += redistBatch; |
||||
|
||||
int residual = totalClipped - redistBatch * 256; |
||||
if (tid < residual) |
||||
++tHistVal; |
||||
} |
||||
|
||||
const int lutVal = blockScanInclusive<256>(tHistVal, smem, tid); |
||||
|
||||
lut(ty * tilesX + tx, tid) = saturate_cast<uchar>(__float2int_rn(lutScale * lutVal)); |
||||
} |
||||
|
||||
void calcLut(PtrStepSzb src, PtrStepb lut, int tilesX, int tilesY, int2 tileSize, int clipLimit, float lutScale, cudaStream_t stream) |
||||
{ |
||||
const dim3 block(32, 8); |
||||
const dim3 grid(tilesX, tilesY); |
||||
|
||||
calcLutKernel<<<grid, block, 0, stream>>>(src, lut, tileSize, tilesX, clipLimit, lutScale); |
||||
|
||||
cudaSafeCall( cudaGetLastError() ); |
||||
|
||||
if (stream == 0) |
||||
cudaSafeCall( cudaDeviceSynchronize() ); |
||||
} |
||||
|
||||
__global__ void tranformKernel(const PtrStepSzb src, PtrStepb dst, const PtrStepb lut, const int2 tileSize, const int tilesX, const int tilesY) |
||||
{ |
||||
const int x = blockIdx.x * blockDim.x + threadIdx.x; |
||||
const int y = blockIdx.y * blockDim.y + threadIdx.y; |
||||
|
||||
if (x >= src.cols || y >= src.rows) |
||||
return; |
||||
|
||||
const float tyf = (static_cast<float>(y) / tileSize.y) - 0.5f; |
||||
int ty1 = __float2int_rd(tyf); |
||||
int ty2 = ty1 + 1; |
||||
const float ya = tyf - ty1; |
||||
ty1 = ::max(ty1, 0); |
||||
ty2 = ::min(ty2, tilesY - 1); |
||||
|
||||
const float txf = (static_cast<float>(x) / tileSize.x) - 0.5f; |
||||
int tx1 = __float2int_rd(txf); |
||||
int tx2 = tx1 + 1; |
||||
const float xa = txf - tx1; |
||||
tx1 = ::max(tx1, 0); |
||||
tx2 = ::min(tx2, tilesX - 1); |
||||
|
||||
const int srcVal = src(y, x); |
||||
|
||||
float res = 0; |
||||
|
||||
res += lut(ty1 * tilesX + tx1, srcVal) * ((1.0f - xa) * (1.0f - ya)); |
||||
res += lut(ty1 * tilesX + tx2, srcVal) * ((xa) * (1.0f - ya)); |
||||
res += lut(ty2 * tilesX + tx1, srcVal) * ((1.0f - xa) * (ya)); |
||||
res += lut(ty2 * tilesX + tx2, srcVal) * ((xa) * (ya)); |
||||
|
||||
dst(y, x) = saturate_cast<uchar>(res); |
||||
} |
||||
|
||||
void transform(PtrStepSzb src, PtrStepSzb dst, PtrStepb lut, int tilesX, int tilesY, int2 tileSize, cudaStream_t stream) |
||||
{ |
||||
const dim3 block(32, 8); |
||||
const dim3 grid(divUp(src.cols, block.x), divUp(src.rows, block.y)); |
||||
|
||||
cudaSafeCall( cudaFuncSetCacheConfig(tranformKernel, cudaFuncCachePreferL1) ); |
||||
|
||||
tranformKernel<<<grid, block, 0, stream>>>(src, dst, lut, tileSize, tilesX, tilesY); |
||||
cudaSafeCall( cudaGetLastError() ); |
||||
|
||||
if (stream == 0) |
||||
cudaSafeCall( cudaDeviceSynchronize() ); |
||||
} |
||||
} |
||||
|
||||
#endif // CUDA_DISABLER |
||||
Some files were not shown because too many files have changed in this diff Show More
Loading…
Reference in new issue