Conflicts: doc/tutorials/definitions/noContent.rst doc/tutorials/gpu/gpu-basics-similarity/gpu-basics-similarity.rst doc/tutorials/introduction/android_binary_package/dev_with_OCV_on_Android.rst doc/tutorials/introduction/how_to_write_a_tutorial/how_to_write_a_tutorial.rst modules/core/include/opencv2/core/core.hpp modules/core/include/opencv2/core/internal.hpp modules/core/include/opencv2/core/version.hpp modules/gpu/CMakeLists.txt modules/highgui/perf/perf_output.cpp modules/highgui/test/test_video_io.cpp modules/ocl/include/opencv2/ocl/ocl.hpp modules/ocl/perf/main.cpp modules/ocl/src/hog.cpp modules/ocl/src/initialization.cpp modules/ocl/src/moments.cpp modules/ocl/src/opencl/moments.cl modules/ocl/test/main.cpp modules/ocl/test/test_moments.cpp modules/python/test/test.py modules/ts/include/opencv2/ts/ts_perf.hpp modules/ts/src/precomp.hpp modules/ts/src/ts_perf.cpppull/1081/head
Roman Donchenko
12 years ago
commit
4cf9990d4e
146 changed files with 11629 additions and 2732 deletions
File diff suppressed because it is too large
Load Diff
@ -1 +0,0 @@ |
|||||||
All Android specific sources are moved to platforms/android. |
|
||||||
@ -1,3 +1,3 @@ |
|||||||
|
|
||||||
.. note:: |
.. note:: |
||||||
Unfortunetly we have no tutorials into this section. Nevertheless, our tutorial writting team is working on it. If you have a tutorial suggestion or you have writen yourself a tutorial (or coded a sample code) that you would like to see here please contact us via our :opencv_group:`user group <>`. |
Unfortunetly we have no tutorials into this section. And you can help us with that, since OpenCV is a community effort. If you have a tutorial suggestion or you have written a tutorial yourself (or coded a sample code) that you would like to see here, please contact follow these instructions: :ref:`howToWriteTutorial` and :how_to_contribute:`How to contribute <>`. |
||||||
|
|||||||
File diff suppressed because one or more lines are too long
File diff suppressed because one or more lines are too long
@ -1 +0,0 @@ |
|||||||
cmake -GXcode -DCMAKE_TOOLCHAIN_FILE=../opencv/ios/cmake/Toolchains/Toolchain-iPhoneOS_Xcode.cmake -DCMAKE_INSTALL_PREFIX=../OpenCV_iPhoneOS ../opencv |
|
||||||
@ -1 +0,0 @@ |
|||||||
cmake -GXcode -DCMAKE_TOOLCHAIN_FILE=../opencv/ios/cmake/Toolchains/Toolchain-iPhoneSimulator_Xcode.cmake -DCMAKE_INSTALL_PREFIX=../OpenCV_iPhoneSimulator ../opencv |
|
||||||
@ -1,15 +0,0 @@ |
|||||||
Assuming that your build directory is on the same level that opencv source, |
|
||||||
From the build directory run |
|
||||||
../opencv/ios/configure-device_xcode.sh |
|
||||||
or |
|
||||||
../opencv/ios/configure-simulator_xcode.sh |
|
||||||
|
|
||||||
Then from the same folder invoke |
|
||||||
|
|
||||||
xcodebuild -sdk iphoneos -configuration Release -target ALL_BUILD |
|
||||||
xcodebuild -sdk iphoneos -configuration Release -target install install |
|
||||||
|
|
||||||
or |
|
||||||
|
|
||||||
xcodebuild -sdk iphonesimulator -configuration Release -target ALL_BUILD |
|
||||||
xcodebuild -sdk iphonesimulator -configuration Release -target install install |
|
||||||
@ -0,0 +1,49 @@ |
|||||||
|
|
||||||
|
/*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 "opencv2/core/core.hpp" |
||||||
|
#import "opencv2/highgui/cap_ios.h" |
||||||
|
|
||||||
|
UIImage* MatToUIImage(const cv::Mat& image); |
||||||
|
void UIImageToMat(const UIImage* image, |
||||||
|
cv::Mat& m, bool alphaExist = false); |
||||||
File diff suppressed because it is too large
Load Diff
@ -0,0 +1,117 @@ |
|||||||
|
/*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*/ |
||||||
|
|
||||||
|
#import "opencv2/highgui/cap_ios.h" |
||||||
|
#include "precomp.hpp" |
||||||
|
|
||||||
|
UIImage* MatToUIImage(const cv::Mat& image) { |
||||||
|
|
||||||
|
NSData *data = [NSData dataWithBytes:image.data |
||||||
|
length:image.elemSize()*image.total()]; |
||||||
|
|
||||||
|
CGColorSpaceRef colorSpace; |
||||||
|
|
||||||
|
if (image.elemSize() == 1) { |
||||||
|
colorSpace = CGColorSpaceCreateDeviceGray(); |
||||||
|
} else { |
||||||
|
colorSpace = CGColorSpaceCreateDeviceRGB(); |
||||||
|
} |
||||||
|
|
||||||
|
CGDataProviderRef provider = |
||||||
|
CGDataProviderCreateWithCFData((__bridge CFDataRef)data); |
||||||
|
|
||||||
|
// Creating CGImage from cv::Mat |
||||||
|
CGImageRef imageRef = CGImageCreate(image.cols, |
||||||
|
image.rows, |
||||||
|
8, |
||||||
|
8 * image.elemSize(), |
||||||
|
image.step.p[0], |
||||||
|
colorSpace, |
||||||
|
kCGImageAlphaNone| |
||||||
|
kCGBitmapByteOrderDefault, |
||||||
|
provider, |
||||||
|
NULL, |
||||||
|
false, |
||||||
|
kCGRenderingIntentDefault |
||||||
|
); |
||||||
|
|
||||||
|
|
||||||
|
// Getting UIImage from CGImage |
||||||
|
UIImage *finalImage = [UIImage imageWithCGImage:imageRef]; |
||||||
|
CGImageRelease(imageRef); |
||||||
|
CGDataProviderRelease(provider); |
||||||
|
CGColorSpaceRelease(colorSpace); |
||||||
|
|
||||||
|
return finalImage; |
||||||
|
} |
||||||
|
|
||||||
|
void UIImageToMat(const UIImage* image, |
||||||
|
cv::Mat& m, bool alphaExist) { |
||||||
|
CGColorSpaceRef colorSpace = CGImageGetColorSpace(image.CGImage); |
||||||
|
CGFloat cols = image.size.width, rows = image.size.height; |
||||||
|
CGContextRef contextRef; |
||||||
|
CGBitmapInfo bitmapInfo = kCGImageAlphaPremultipliedLast; |
||||||
|
if (CGColorSpaceGetModel(colorSpace) == 0) |
||||||
|
{ |
||||||
|
m.create(rows, cols, CV_8UC1); // 8 bits per component, 1 channel |
||||||
|
bitmapInfo = kCGImageAlphaNone; |
||||||
|
if (!alphaExist) |
||||||
|
bitmapInfo = kCGImageAlphaNone; |
||||||
|
contextRef = CGBitmapContextCreate(m.data, m.cols, m.rows, 8, |
||||||
|
m.step[0], colorSpace, |
||||||
|
bitmapInfo); |
||||||
|
} |
||||||
|
else |
||||||
|
{ |
||||||
|
m.create(rows, cols, CV_8UC4); // 8 bits per component, 4 channels |
||||||
|
if (!alphaExist) |
||||||
|
bitmapInfo = kCGImageAlphaNoneSkipLast | |
||||||
|
kCGBitmapByteOrderDefault; |
||||||
|
contextRef = CGBitmapContextCreate(m.data, m.cols, m.rows, 8, |
||||||
|
m.step[0], colorSpace, |
||||||
|
bitmapInfo); |
||||||
|
} |
||||||
|
CGContextDrawImage(contextRef, CGRectMake(0, 0, cols, rows), |
||||||
|
image.CGImage); |
||||||
|
CGContextRelease(contextRef); |
||||||
|
CGColorSpaceRelease(colorSpace); |
||||||
|
} |
||||||
@ -1,4 +1,11 @@ |
|||||||
#include "perf_precomp.hpp" |
#include "perf_precomp.hpp" |
||||||
#include "opencv2/ts/gpu_perf.hpp" |
#include "opencv2/ts/gpu_perf.hpp" |
||||||
|
|
||||||
CV_PERF_TEST_MAIN(nonfree, perf::printCudaInfo()) |
static const char * impls[] = { |
||||||
|
#ifdef HAVE_CUDA |
||||||
|
"cuda", |
||||||
|
#endif |
||||||
|
"plain" |
||||||
|
}; |
||||||
|
|
||||||
|
CV_PERF_TEST_MAIN_WITH_IMPLS(nonfree, impls, perf::printCudaInfo()) |
||||||
|
|||||||
@ -0,0 +1,450 @@ |
|||||||
|
/*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 |
||||||
|
// Sen Liu, swjtuls1987@126.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 oclMaterials 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*/ |
||||||
|
|
||||||
|
|
||||||
|
#define tx get_local_id(0) |
||||||
|
#define ty get_local_id(1) |
||||||
|
#define bx get_group_id(0) |
||||||
|
#define bdx get_local_size(0) |
||||||
|
|
||||||
|
#define BORDER_SIZE 5 |
||||||
|
#define MAX_KSIZE_HALF 100 |
||||||
|
|
||||||
|
#ifndef polyN |
||||||
|
#define polyN 5 |
||||||
|
#endif |
||||||
|
|
||||||
|
__kernel void polynomialExpansion(__global float * dst, |
||||||
|
__global __const float * src, |
||||||
|
__global __const float * c_g, |
||||||
|
__global __const float * c_xg, |
||||||
|
__global __const float * c_xxg, |
||||||
|
__local float * smem, |
||||||
|
const float4 ig, |
||||||
|
const int height, const int width, |
||||||
|
int dstStep, int srcStep) |
||||||
|
{ |
||||||
|
const int y = get_global_id(1); |
||||||
|
const int x = bx * (bdx - 2*polyN) + tx - polyN; |
||||||
|
|
||||||
|
dstStep /= sizeof(*dst); |
||||||
|
srcStep /= sizeof(*src); |
||||||
|
|
||||||
|
int xWarped; |
||||||
|
__local float *row = smem + tx; |
||||||
|
|
||||||
|
if (y < height && y >= 0) |
||||||
|
{ |
||||||
|
xWarped = min(max(x, 0), width - 1); |
||||||
|
|
||||||
|
row[0] = src[mad24(y, srcStep, xWarped)] * c_g[0]; |
||||||
|
row[bdx] = 0.f; |
||||||
|
row[2*bdx] = 0.f; |
||||||
|
|
||||||
|
#pragma unroll |
||||||
|
for (int k = 1; k <= polyN; ++k) |
||||||
|
{ |
||||||
|
float t0 = src[mad24(max(y - k, 0), srcStep, xWarped)]; |
||||||
|
float t1 = src[mad24(min(y + k, height - 1), srcStep, xWarped)]; |
||||||
|
|
||||||
|
row[0] += c_g[k] * (t0 + t1); |
||||||
|
row[bdx] += c_xg[k] * (t1 - t0); |
||||||
|
row[2*bdx] += c_xxg[k] * (t0 + t1); |
||||||
|
} |
||||||
|
} |
||||||
|
|
||||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||||
|
|
||||||
|
if (y < height && y >= 0 && tx >= polyN && tx + polyN < bdx && x < width) |
||||||
|
{ |
||||||
|
float b1 = c_g[0] * row[0]; |
||||||
|
float b3 = c_g[0] * row[bdx]; |
||||||
|
float b5 = c_g[0] * row[2*bdx]; |
||||||
|
float b2 = 0, b4 = 0, b6 = 0; |
||||||
|
|
||||||
|
#pragma unroll |
||||||
|
for (int k = 1; k <= polyN; ++k) |
||||||
|
{ |
||||||
|
b1 += (row[k] + row[-k]) * c_g[k]; |
||||||
|
b4 += (row[k] + row[-k]) * c_xxg[k]; |
||||||
|
b2 += (row[k] - row[-k]) * c_xg[k]; |
||||||
|
b3 += (row[k + bdx] + row[-k + bdx]) * c_g[k]; |
||||||
|
b6 += (row[k + bdx] - row[-k + bdx]) * c_xg[k]; |
||||||
|
b5 += (row[k + 2*bdx] + row[-k + 2*bdx]) * c_g[k]; |
||||||
|
} |
||||||
|
|
||||||
|
dst[mad24(y, dstStep, xWarped)] = b3*ig.s0; |
||||||
|
dst[mad24(height + y, dstStep, xWarped)] = b2*ig.s0; |
||||||
|
dst[mad24(2*height + y, dstStep, xWarped)] = b1*ig.s1 + b5*ig.s2; |
||||||
|
dst[mad24(3*height + y, dstStep, xWarped)] = b1*ig.s1 + b4*ig.s2; |
||||||
|
dst[mad24(4*height + y, dstStep, xWarped)] = b6*ig.s3; |
||||||
|
} |
||||||
|
} |
||||||
|
|
||||||
|
inline int idx_row_low(const int y, const int last_row) |
||||||
|
{ |
||||||
|
return abs(y) % (last_row + 1); |
||||||
|
} |
||||||
|
|
||||||
|
inline int idx_row_high(const int y, const int last_row) |
||||||
|
{ |
||||||
|
return abs(last_row - abs(last_row - y)) % (last_row + 1); |
||||||
|
} |
||||||
|
|
||||||
|
inline int idx_row(const int y, const int last_row) |
||||||
|
{ |
||||||
|
return idx_row_low(idx_row_high(y, last_row), last_row); |
||||||
|
} |
||||||
|
|
||||||
|
inline int idx_col_low(const int x, const int last_col) |
||||||
|
{ |
||||||
|
return abs(x) % (last_col + 1); |
||||||
|
} |
||||||
|
|
||||||
|
inline int idx_col_high(const int x, const int last_col) |
||||||
|
{ |
||||||
|
return abs(last_col - abs(last_col - x)) % (last_col + 1); |
||||||
|
} |
||||||
|
|
||||||
|
inline int idx_col(const int x, const int last_col) |
||||||
|
{ |
||||||
|
return idx_col_low(idx_col_high(x, last_col), last_col); |
||||||
|
} |
||||||
|
|
||||||
|
__kernel void gaussianBlur(__global float * dst, |
||||||
|
__global const float * src, |
||||||
|
__global const float * c_gKer, |
||||||
|
__local float * smem, |
||||||
|
const int height, const int width, |
||||||
|
int dstStep, int srcStep, |
||||||
|
const int ksizeHalf) |
||||||
|
{ |
||||||
|
const int y = get_global_id(1); |
||||||
|
const int x = get_global_id(0); |
||||||
|
|
||||||
|
dstStep /= sizeof(*dst); |
||||||
|
srcStep /= sizeof(*src); |
||||||
|
|
||||||
|
__local float *row = smem + ty * (bdx + 2*ksizeHalf); |
||||||
|
|
||||||
|
if (y < height) |
||||||
|
{ |
||||||
|
// Vertical pass |
||||||
|
for (int i = tx; i < bdx + 2*ksizeHalf; i += bdx) |
||||||
|
{ |
||||||
|
int xExt = (int)(bx * bdx) + i - ksizeHalf; |
||||||
|
xExt = idx_col(xExt, width - 1); |
||||||
|
row[i] = src[mad24(y, srcStep, xExt)] * c_gKer[0]; |
||||||
|
for (int j = 1; j <= ksizeHalf; ++j) |
||||||
|
row[i] += (src[mad24(idx_row_low(y - j, height - 1), srcStep, xExt)] |
||||||
|
+ src[mad24(idx_row_high(y + j, height - 1), srcStep, xExt)]) * c_gKer[j]; |
||||||
|
} |
||||||
|
} |
||||||
|
|
||||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||||
|
|
||||||
|
if (y < height && y >= 0 && x < width && x >= 0) |
||||||
|
{ |
||||||
|
// Horizontal pass |
||||||
|
row += tx + ksizeHalf; |
||||||
|
float res = row[0] * c_gKer[0]; |
||||||
|
for (int i = 1; i <= ksizeHalf; ++i) |
||||||
|
res += (row[-i] + row[i]) * c_gKer[i]; |
||||||
|
|
||||||
|
dst[mad24(y, dstStep, x)] = res; |
||||||
|
} |
||||||
|
} |
||||||
|
|
||||||
|
__constant float c_border[BORDER_SIZE + 1] = { 0.14f, 0.14f, 0.4472f, 0.4472f, 0.4472f, 1.f }; |
||||||
|
|
||||||
|
__kernel void updateMatrices(__global float * M, |
||||||
|
__global const float * flowx, __global const float * flowy, |
||||||
|
__global const float * R0, __global const float * R1, |
||||||
|
const int height, const int width, |
||||||
|
int mStep, int xStep, int yStep, int R0Step, int R1Step) |
||||||
|
{ |
||||||
|
const int y = get_global_id(1); |
||||||
|
const int x = get_global_id(0); |
||||||
|
|
||||||
|
mStep /= sizeof(*M); |
||||||
|
xStep /= sizeof(*flowx); |
||||||
|
yStep /= sizeof(*flowy); |
||||||
|
R0Step /= sizeof(*R0); |
||||||
|
R1Step /= sizeof(*R1); |
||||||
|
|
||||||
|
if (y < height && y >= 0 && x < width && x >= 0) |
||||||
|
{ |
||||||
|
float dx = flowx[mad24(y, xStep, x)]; |
||||||
|
float dy = flowy[mad24(y, yStep, x)]; |
||||||
|
float fx = x + dx; |
||||||
|
float fy = y + dy; |
||||||
|
|
||||||
|
int x1 = convert_int(floor(fx)); |
||||||
|
int y1 = convert_int(floor(fy)); |
||||||
|
fx -= x1; |
||||||
|
fy -= y1; |
||||||
|
|
||||||
|
float r2, r3, r4, r5, r6; |
||||||
|
|
||||||
|
if (x1 >= 0 && y1 >= 0 && x1 < width - 1 && y1 < height - 1) |
||||||
|
{ |
||||||
|
float a00 = (1.f - fx) * (1.f - fy); |
||||||
|
float a01 = fx * (1.f - fy); |
||||||
|
float a10 = (1.f - fx) * fy; |
||||||
|
float a11 = fx * fy; |
||||||
|
|
||||||
|
r2 = a00 * R1[mad24(y1, R1Step, x1)] + |
||||||
|
a01 * R1[mad24(y1, R1Step, x1 + 1)] + |
||||||
|
a10 * R1[mad24(y1 + 1, R1Step, x1)] + |
||||||
|
a11 * R1[mad24(y1 + 1, R1Step, x1 + 1)]; |
||||||
|
|
||||||
|
r3 = a00 * R1[mad24(height + y1, R1Step, x1)] + |
||||||
|
a01 * R1[mad24(height + y1, R1Step, x1 + 1)] + |
||||||
|
a10 * R1[mad24(height + y1 + 1, R1Step, x1)] + |
||||||
|
a11 * R1[mad24(height + y1 + 1, R1Step, x1 + 1)]; |
||||||
|
|
||||||
|
r4 = a00 * R1[mad24(2*height + y1, R1Step, x1)] + |
||||||
|
a01 * R1[mad24(2*height + y1, R1Step, x1 + 1)] + |
||||||
|
a10 * R1[mad24(2*height + y1 + 1, R1Step, x1)] + |
||||||
|
a11 * R1[mad24(2*height + y1 + 1, R1Step, x1 + 1)]; |
||||||
|
|
||||||
|
r5 = a00 * R1[mad24(3*height + y1, R1Step, x1)] + |
||||||
|
a01 * R1[mad24(3*height + y1, R1Step, x1 + 1)] + |
||||||
|
a10 * R1[mad24(3*height + y1 + 1, R1Step, x1)] + |
||||||
|
a11 * R1[mad24(3*height + y1 + 1, R1Step, x1 + 1)]; |
||||||
|
|
||||||
|
r6 = a00 * R1[mad24(4*height + y1, R1Step, x1)] + |
||||||
|
a01 * R1[mad24(4*height + y1, R1Step, x1 + 1)] + |
||||||
|
a10 * R1[mad24(4*height + y1 + 1, R1Step, x1)] + |
||||||
|
a11 * R1[mad24(4*height + y1 + 1, R1Step, x1 + 1)]; |
||||||
|
|
||||||
|
r4 = (R0[mad24(2*height + y, R0Step, x)] + r4) * 0.5f; |
||||||
|
r5 = (R0[mad24(3*height + y, R0Step, x)] + r5) * 0.5f; |
||||||
|
r6 = (R0[mad24(4*height + y, R0Step, x)] + r6) * 0.25f; |
||||||
|
} |
||||||
|
else |
||||||
|
{ |
||||||
|
r2 = r3 = 0.f; |
||||||
|
r4 = R0[mad24(2*height + y, R0Step, x)]; |
||||||
|
r5 = R0[mad24(3*height + y, R0Step, x)]; |
||||||
|
r6 = R0[mad24(4*height + y, R0Step, x)] * 0.5f; |
||||||
|
} |
||||||
|
|
||||||
|
r2 = (R0[mad24(y, R0Step, x)] - r2) * 0.5f; |
||||||
|
r3 = (R0[mad24(height + y, R0Step, x)] - r3) * 0.5f; |
||||||
|
|
||||||
|
r2 += r4*dy + r6*dx; |
||||||
|
r3 += r6*dy + r5*dx; |
||||||
|
|
||||||
|
float scale = |
||||||
|
c_border[min(x, BORDER_SIZE)] * |
||||||
|
c_border[min(y, BORDER_SIZE)] * |
||||||
|
c_border[min(width - x - 1, BORDER_SIZE)] * |
||||||
|
c_border[min(height - y - 1, BORDER_SIZE)]; |
||||||
|
|
||||||
|
r2 *= scale; |
||||||
|
r3 *= scale; |
||||||
|
r4 *= scale; |
||||||
|
r5 *= scale; |
||||||
|
r6 *= scale; |
||||||
|
|
||||||
|
M[mad24(y, mStep, x)] = r4*r4 + r6*r6; |
||||||
|
M[mad24(height + y, mStep, x)] = (r4 + r5)*r6; |
||||||
|
M[mad24(2*height + y, mStep, x)] = r5*r5 + r6*r6; |
||||||
|
M[mad24(3*height + y, mStep, x)] = r4*r2 + r6*r3; |
||||||
|
M[mad24(4*height + y, mStep, x)] = r6*r2 + r5*r3; |
||||||
|
} |
||||||
|
} |
||||||
|
|
||||||
|
__kernel void boxFilter5(__global float * dst, |
||||||
|
__global const float * src, |
||||||
|
__local float * smem, |
||||||
|
const int height, const int width, |
||||||
|
int dstStep, int srcStep, |
||||||
|
const int ksizeHalf) |
||||||
|
{ |
||||||
|
const int y = get_global_id(1); |
||||||
|
const int x = get_global_id(0); |
||||||
|
|
||||||
|
const float boxAreaInv = 1.f / ((1 + 2*ksizeHalf) * (1 + 2*ksizeHalf)); |
||||||
|
const int smw = bdx + 2*ksizeHalf; // shared memory "width" |
||||||
|
__local float *row = smem + 5 * ty * smw; |
||||||
|
|
||||||
|
dstStep /= sizeof(*dst); |
||||||
|
srcStep /= sizeof(*src); |
||||||
|
|
||||||
|
if (y < height) |
||||||
|
{ |
||||||
|
// Vertical pass |
||||||
|
for (int i = tx; i < bdx + 2*ksizeHalf; i += bdx) |
||||||
|
{ |
||||||
|
int xExt = (int)(bx * bdx) + i - ksizeHalf; |
||||||
|
xExt = min(max(xExt, 0), width - 1); |
||||||
|
|
||||||
|
#pragma unroll |
||||||
|
for (int k = 0; k < 5; ++k) |
||||||
|
row[k*smw + i] = src[mad24(k*height + y, srcStep, xExt)]; |
||||||
|
|
||||||
|
for (int j = 1; j <= ksizeHalf; ++j) |
||||||
|
#pragma unroll |
||||||
|
for (int k = 0; k < 5; ++k) |
||||||
|
row[k*smw + i] += |
||||||
|
src[mad24(k*height + max(y - j, 0), srcStep, xExt)] + |
||||||
|
src[mad24(k*height + min(y + j, height - 1), srcStep, xExt)]; |
||||||
|
} |
||||||
|
} |
||||||
|
|
||||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||||
|
|
||||||
|
if (y < height && y >= 0 && x < width && x >= 0) |
||||||
|
{ |
||||||
|
// Horizontal pass |
||||||
|
|
||||||
|
row += tx + ksizeHalf; |
||||||
|
float res[5]; |
||||||
|
|
||||||
|
#pragma unroll |
||||||
|
for (int k = 0; k < 5; ++k) |
||||||
|
res[k] = row[k*smw]; |
||||||
|
|
||||||
|
for (int i = 1; i <= ksizeHalf; ++i) |
||||||
|
#pragma unroll |
||||||
|
for (int k = 0; k < 5; ++k) |
||||||
|
res[k] += row[k*smw - i] + row[k*smw + i]; |
||||||
|
|
||||||
|
#pragma unroll |
||||||
|
for (int k = 0; k < 5; ++k) |
||||||
|
dst[mad24(k*height + y, dstStep, x)] = res[k] * boxAreaInv; |
||||||
|
} |
||||||
|
} |
||||||
|
|
||||||
|
__kernel void updateFlow(__global float4 * flowx, __global float4 * flowy, |
||||||
|
__global const float4 * M, |
||||||
|
const int height, const int width, |
||||||
|
int xStep, int yStep, int mStep) |
||||||
|
{ |
||||||
|
const int y = get_global_id(1); |
||||||
|
const int x = get_global_id(0); |
||||||
|
|
||||||
|
xStep /= sizeof(*flowx); |
||||||
|
yStep /= sizeof(*flowy); |
||||||
|
mStep /= sizeof(*M); |
||||||
|
|
||||||
|
if (y < height && y >= 0 && x < width && x >= 0) |
||||||
|
{ |
||||||
|
float4 g11 = M[mad24(y, mStep, x)]; |
||||||
|
float4 g12 = M[mad24(height + y, mStep, x)]; |
||||||
|
float4 g22 = M[mad24(2*height + y, mStep, x)]; |
||||||
|
float4 h1 = M[mad24(3*height + y, mStep, x)]; |
||||||
|
float4 h2 = M[mad24(4*height + y, mStep, x)]; |
||||||
|
|
||||||
|
float4 detInv = (float4)(1.f) / (g11*g22 - g12*g12 + (float4)(1e-3f)); |
||||||
|
|
||||||
|
flowx[mad24(y, xStep, x)] = (g11*h2 - g12*h1) * detInv; |
||||||
|
flowy[mad24(y, yStep, x)] = (g22*h1 - g12*h2) * detInv; |
||||||
|
} |
||||||
|
} |
||||||
|
|
||||||
|
__kernel void gaussianBlur5(__global float * dst, |
||||||
|
__global const float * src, |
||||||
|
__global const float * c_gKer, |
||||||
|
__local float * smem, |
||||||
|
const int height, const int width, |
||||||
|
int dstStep, int srcStep, |
||||||
|
const int ksizeHalf) |
||||||
|
{ |
||||||
|
const int y = get_global_id(1); |
||||||
|
const int x = get_global_id(0); |
||||||
|
|
||||||
|
const int smw = bdx + 2*ksizeHalf; // shared memory "width" |
||||||
|
__local volatile float *row = smem + 5 * ty * smw; |
||||||
|
|
||||||
|
dstStep /= sizeof(*dst); |
||||||
|
srcStep /= sizeof(*src); |
||||||
|
|
||||||
|
if (y < height) |
||||||
|
{ |
||||||
|
// Vertical pass |
||||||
|
for (int i = tx; i < bdx + 2*ksizeHalf; i += bdx) |
||||||
|
{ |
||||||
|
int xExt = (int)(bx * bdx) + i - ksizeHalf; |
||||||
|
xExt = idx_col(xExt, width - 1); |
||||||
|
|
||||||
|
#pragma unroll |
||||||
|
for (int k = 0; k < 5; ++k) |
||||||
|
row[k*smw + i] = src[mad24(k*height + y, srcStep, xExt)] * c_gKer[0]; |
||||||
|
|
||||||
|
for (int j = 1; j <= ksizeHalf; ++j) |
||||||
|
#pragma unroll |
||||||
|
for (int k = 0; k < 5; ++k) |
||||||
|
row[k*smw + i] += |
||||||
|
(src[mad24(k*height + idx_row_low(y - j, height - 1), srcStep, xExt)] + |
||||||
|
src[mad24(k*height + idx_row_high(y + j, height - 1), srcStep, xExt)]) * c_gKer[j]; |
||||||
|
} |
||||||
|
} |
||||||
|
|
||||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||||
|
|
||||||
|
if (y < height && y >= 0 && x < width && x >= 0) |
||||||
|
{ |
||||||
|
// Horizontal pass |
||||||
|
|
||||||
|
row += tx + ksizeHalf; |
||||||
|
float res[5]; |
||||||
|
|
||||||
|
#pragma unroll |
||||||
|
for (int k = 0; k < 5; ++k) |
||||||
|
res[k] = row[k*smw] * c_gKer[0]; |
||||||
|
|
||||||
|
for (int i = 1; i <= ksizeHalf; ++i) |
||||||
|
#pragma unroll |
||||||
|
for (int k = 0; k < 5; ++k) |
||||||
|
res[k] += (row[k*smw - i] + row[k*smw + i]) * c_gKer[i]; |
||||||
|
|
||||||
|
#pragma unroll |
||||||
|
for (int k = 0; k < 5; ++k) |
||||||
|
dst[mad24(k*height + y, dstStep, x)] = res[k]; |
||||||
|
} |
||||||
|
} |
||||||
@ -0,0 +1,540 @@ |
|||||||
|
/*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
|
||||||
|
// Sen Liu, swjtuls1987@126.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 oclMaterials 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 "precomp.hpp" |
||||||
|
#include "opencv2/video/tracking.hpp" |
||||||
|
|
||||||
|
using namespace std; |
||||||
|
using namespace cv; |
||||||
|
using namespace cv::ocl; |
||||||
|
|
||||||
|
#define MIN_SIZE 32 |
||||||
|
|
||||||
|
namespace cv |
||||||
|
{ |
||||||
|
namespace ocl |
||||||
|
{ |
||||||
|
///////////////////////////OpenCL kernel strings///////////////////////////
|
||||||
|
extern const char *optical_flow_farneback; |
||||||
|
} |
||||||
|
} |
||||||
|
|
||||||
|
namespace cv { |
||||||
|
namespace ocl { |
||||||
|
namespace optflow_farneback |
||||||
|
{ |
||||||
|
oclMat g; |
||||||
|
oclMat xg; |
||||||
|
oclMat xxg; |
||||||
|
oclMat gKer; |
||||||
|
|
||||||
|
float ig[4]; |
||||||
|
|
||||||
|
inline int divUp(int total, int grain) |
||||||
|
{ |
||||||
|
return (total + grain - 1) / grain; |
||||||
|
} |
||||||
|
|
||||||
|
inline void setGaussianBlurKernel(const float *c_gKer, int ksizeHalf) |
||||||
|
{ |
||||||
|
cv::Mat t_gKer(1, ksizeHalf + 1, CV_32FC1, const_cast<float *>(c_gKer)); |
||||||
|
gKer.upload(t_gKer); |
||||||
|
} |
||||||
|
|
||||||
|
static void gaussianBlurOcl(const oclMat &src, int ksizeHalf, oclMat &dst) |
||||||
|
{ |
||||||
|
string kernelName("gaussianBlur"); |
||||||
|
size_t localThreads[3] = { 256, 1, 1 }; |
||||||
|
size_t globalThreads[3] = { divUp(src.cols, localThreads[0]) * localThreads[0], src.rows, 1 }; |
||||||
|
int smem_size = (localThreads[0] + 2*ksizeHalf) * sizeof(float); |
||||||
|
|
||||||
|
CV_Assert(dst.size() == src.size()); |
||||||
|
std::vector< std::pair<size_t, const void *> > args; |
||||||
|
args.push_back(std::make_pair(sizeof(cl_mem), (void *)&dst.data)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_mem), (void *)&src.data)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_mem), (void *)&gKer.data)); |
||||||
|
args.push_back(std::make_pair(smem_size, (void *)NULL)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_int), (void *)&dst.rows)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_int), (void *)&dst.cols)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_int), (void *)&dst.step)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_int), (void *)&src.step)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_int), (void *)&ksizeHalf)); |
||||||
|
|
||||||
|
openCLExecuteKernel(Context::getContext(), &optical_flow_farneback, kernelName, |
||||||
|
globalThreads, localThreads, args, -1, -1); |
||||||
|
} |
||||||
|
|
||||||
|
static void polynomialExpansionOcl(const oclMat &src, int polyN, oclMat &dst) |
||||||
|
{ |
||||||
|
string kernelName("polynomialExpansion"); |
||||||
|
size_t localThreads[3] = { 256, 1, 1 }; |
||||||
|
size_t globalThreads[3] = { divUp(src.cols, localThreads[0] - 2*polyN) * localThreads[0], src.rows, 1 }; |
||||||
|
int smem_size = 3 * localThreads[0] * sizeof(float); |
||||||
|
|
||||||
|
std::vector< std::pair<size_t, const void *> > args; |
||||||
|
args.push_back(std::make_pair(sizeof(cl_mem), (void *)&dst.data)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_mem), (void *)&src.data)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_mem), (void *)&g.data)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_mem), (void *)&xg.data)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_mem), (void *)&xxg.data)); |
||||||
|
args.push_back(std::make_pair(smem_size, (void *)NULL)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_float4), (void *)&ig)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_int), (void *)&src.rows)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_int), (void *)&src.cols)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_int), (void *)&dst.step)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_int), (void *)&src.step)); |
||||||
|
|
||||||
|
char opt [128]; |
||||||
|
sprintf(opt, "-D polyN=%d", polyN); |
||||||
|
|
||||||
|
openCLExecuteKernel(Context::getContext(), &optical_flow_farneback, kernelName, |
||||||
|
globalThreads, localThreads, args, -1, -1, opt); |
||||||
|
} |
||||||
|
|
||||||
|
static void updateMatricesOcl(const oclMat &flowx, const oclMat &flowy, const oclMat &R0, const oclMat &R1, oclMat &M) |
||||||
|
{ |
||||||
|
string kernelName("updateMatrices"); |
||||||
|
size_t localThreads[3] = { 32, 8, 1 }; |
||||||
|
size_t globalThreads[3] = { divUp(flowx.cols, localThreads[0]) * localThreads[0], |
||||||
|
divUp(flowx.rows, localThreads[1]) * localThreads[1], |
||||||
|
1 |
||||||
|
}; |
||||||
|
|
||||||
|
std::vector< std::pair<size_t, const void *> > args; |
||||||
|
args.push_back(std::make_pair(sizeof(cl_mem), (void *)&M.data)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_mem), (void *)&flowx.data)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_mem), (void *)&flowy.data)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_mem), (void *)&R0.data)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_mem), (void *)&R1.data)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_int), (void *)&flowx.rows)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_int), (void *)&flowx.cols)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_int), (void *)&M.step)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_int), (void *)&flowx.step)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_int), (void *)&flowy.step)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_int), (void *)&R0.step)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_int), (void *)&R1.step)); |
||||||
|
|
||||||
|
openCLExecuteKernel(Context::getContext(), &optical_flow_farneback, kernelName, |
||||||
|
globalThreads, localThreads, args, -1, -1); |
||||||
|
} |
||||||
|
|
||||||
|
static void boxFilter5Ocl(const oclMat &src, int ksizeHalf, oclMat &dst) |
||||||
|
{ |
||||||
|
string kernelName("boxFilter5"); |
||||||
|
int height = src.rows / 5; |
||||||
|
size_t localThreads[3] = { 256, 1, 1 }; |
||||||
|
size_t globalThreads[3] = { divUp(src.cols, localThreads[0]) * localThreads[0], height, 1 }; |
||||||
|
int smem_size = (localThreads[0] + 2*ksizeHalf) * 5 * sizeof(float); |
||||||
|
|
||||||
|
std::vector< std::pair<size_t, const void *> > args; |
||||||
|
args.push_back(std::make_pair(sizeof(cl_mem), (void *)&dst.data)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_mem), (void *)&src.data)); |
||||||
|
args.push_back(std::make_pair(smem_size, (void *)NULL)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_int), (void *)&height)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_int), (void *)&src.cols)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_int), (void *)&dst.step)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_int), (void *)&src.step)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_int), (void *)&ksizeHalf)); |
||||||
|
|
||||||
|
openCLExecuteKernel(Context::getContext(), &optical_flow_farneback, kernelName, |
||||||
|
globalThreads, localThreads, args, -1, -1); |
||||||
|
} |
||||||
|
|
||||||
|
static void updateFlowOcl(const oclMat &M, oclMat &flowx, oclMat &flowy) |
||||||
|
{ |
||||||
|
string kernelName("updateFlow"); |
||||||
|
int cols = divUp(flowx.cols, 4); |
||||||
|
size_t localThreads[3] = { 32, 8, 1 }; |
||||||
|
size_t globalThreads[3] = { divUp(cols, localThreads[0]) * localThreads[0], |
||||||
|
divUp(flowx.rows, localThreads[1]) * localThreads[0], |
||||||
|
1 |
||||||
|
}; |
||||||
|
|
||||||
|
std::vector< std::pair<size_t, const void *> > args; |
||||||
|
args.push_back(std::make_pair(sizeof(cl_mem), (void *)&flowx.data)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_mem), (void *)&flowy.data)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_mem), (void *)&M.data)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_int), (void *)&flowx.rows)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_int), (void *)&cols)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_int), (void *)&flowx.step)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_int), (void *)&flowy.step)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_int), (void *)&M.step)); |
||||||
|
|
||||||
|
openCLExecuteKernel(Context::getContext(), &optical_flow_farneback, kernelName, |
||||||
|
globalThreads, localThreads, args, -1, -1); |
||||||
|
} |
||||||
|
|
||||||
|
static void gaussianBlur5Ocl(const oclMat &src, int ksizeHalf, oclMat &dst) |
||||||
|
{ |
||||||
|
string kernelName("gaussianBlur5"); |
||||||
|
int height = src.rows / 5; |
||||||
|
int width = src.cols; |
||||||
|
size_t localThreads[3] = { 256, 1, 1 }; |
||||||
|
size_t globalThreads[3] = { divUp(width, localThreads[0]) * localThreads[0], height, 1 }; |
||||||
|
int smem_size = (localThreads[0] + 2*ksizeHalf) * 5 * sizeof(float); |
||||||
|
|
||||||
|
std::vector< std::pair<size_t, const void *> > args; |
||||||
|
args.push_back(std::make_pair(sizeof(cl_mem), (void *)&dst.data)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_mem), (void *)&src.data)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_mem), (void *)&gKer.data)); |
||||||
|
args.push_back(std::make_pair(smem_size, (void *)NULL)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_int), (void *)&height)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_int), (void *)&width)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_int), (void *)&dst.step)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_int), (void *)&src.step)); |
||||||
|
args.push_back(std::make_pair(sizeof(cl_int), (void *)&ksizeHalf)); |
||||||
|
|
||||||
|
openCLExecuteKernel(Context::getContext(), &optical_flow_farneback, kernelName, |
||||||
|
globalThreads, localThreads, args, -1, -1); |
||||||
|
} |
||||||
|
} |
||||||
|
} |
||||||
|
} // namespace cv { namespace ocl { namespace optflow_farneback
|
||||||
|
|
||||||
|
static oclMat allocMatFromBuf(int rows, int cols, int type, oclMat &mat) |
||||||
|
{ |
||||||
|
if (!mat.empty() && mat.type() == type && mat.rows >= rows && mat.cols >= cols) |
||||||
|
return mat(Rect(0, 0, cols, rows)); |
||||||
|
return mat = oclMat(rows, cols, type); |
||||||
|
} |
||||||
|
|
||||||
|
cv::ocl::FarnebackOpticalFlow::FarnebackOpticalFlow() |
||||||
|
{ |
||||||
|
numLevels = 5; |
||||||
|
pyrScale = 0.5; |
||||||
|
fastPyramids = false; |
||||||
|
winSize = 13; |
||||||
|
numIters = 10; |
||||||
|
polyN = 5; |
||||||
|
polySigma = 1.1; |
||||||
|
flags = 0; |
||||||
|
} |
||||||
|
|
||||||
|
void cv::ocl::FarnebackOpticalFlow::releaseMemory() |
||||||
|
{ |
||||||
|
frames_[0].release(); |
||||||
|
frames_[1].release(); |
||||||
|
pyrLevel_[0].release(); |
||||||
|
pyrLevel_[1].release(); |
||||||
|
M_.release(); |
||||||
|
bufM_.release(); |
||||||
|
R_[0].release(); |
||||||
|
R_[1].release(); |
||||||
|
blurredFrame_[0].release(); |
||||||
|
blurredFrame_[1].release(); |
||||||
|
pyramid0_.clear(); |
||||||
|
pyramid1_.clear(); |
||||||
|
} |
||||||
|
|
||||||
|
void cv::ocl::FarnebackOpticalFlow::prepareGaussian( |
||||||
|
int n, double sigma, float *g, float *xg, float *xxg, |
||||||
|
double &ig11, double &ig03, double &ig33, double &ig55) |
||||||
|
{ |
||||||
|
double s = 0.; |
||||||
|
for (int x = -n; x <= n; x++) |
||||||
|
{ |
||||||
|
g[x] = (float)std::exp(-x*x/(2*sigma*sigma)); |
||||||
|
s += g[x]; |
||||||
|
} |
||||||
|
|
||||||
|
s = 1./s; |
||||||
|
for (int x = -n; x <= n; x++) |
||||||
|
{ |
||||||
|
g[x] = (float)(g[x]*s); |
||||||
|
xg[x] = (float)(x*g[x]); |
||||||
|
xxg[x] = (float)(x*x*g[x]); |
||||||
|
} |
||||||
|
|
||||||
|
Mat_<double> G(6, 6); |
||||||
|
G.setTo(0); |
||||||
|
|
||||||
|
for (int y = -n; y <= n; y++) |
||||||
|
{ |
||||||
|
for (int x = -n; x <= n; x++) |
||||||
|
{ |
||||||
|
G(0,0) += g[y]*g[x]; |
||||||
|
G(1,1) += g[y]*g[x]*x*x; |
||||||
|
G(3,3) += g[y]*g[x]*x*x*x*x; |
||||||
|
G(5,5) += g[y]*g[x]*x*x*y*y; |
||||||
|
} |
||||||
|
} |
||||||
|
|
||||||
|
//G[0][0] = 1.;
|
||||||
|
G(2,2) = G(0,3) = G(0,4) = G(3,0) = G(4,0) = G(1,1); |
||||||
|
G(4,4) = G(3,3); |
||||||
|
G(3,4) = G(4,3) = G(5,5); |
||||||
|
|
||||||
|
// invG:
|
||||||
|
// [ x e e ]
|
||||||
|
// [ y ]
|
||||||
|
// [ y ]
|
||||||
|
// [ e z ]
|
||||||
|
// [ e z ]
|
||||||
|
// [ u ]
|
||||||
|
Mat_<double> invG = G.inv(DECOMP_CHOLESKY); |
||||||
|
|
||||||
|
ig11 = invG(1,1); |
||||||
|
ig03 = invG(0,3); |
||||||
|
ig33 = invG(3,3); |
||||||
|
ig55 = invG(5,5); |
||||||
|
} |
||||||
|
|
||||||
|
void cv::ocl::FarnebackOpticalFlow::setPolynomialExpansionConsts(int n, double sigma) |
||||||
|
{ |
||||||
|
vector<float> buf(n*6 + 3); |
||||||
|
float* g = &buf[0] + n; |
||||||
|
float* xg = g + n*2 + 1; |
||||||
|
float* xxg = xg + n*2 + 1; |
||||||
|
|
||||||
|
if (sigma < FLT_EPSILON) |
||||||
|
sigma = n*0.3; |
||||||
|
|
||||||
|
double ig11, ig03, ig33, ig55; |
||||||
|
prepareGaussian(n, sigma, g, xg, xxg, ig11, ig03, ig33, ig55); |
||||||
|
|
||||||
|
cv::Mat t_g(1, n + 1, CV_32FC1, g); |
||||||
|
cv::Mat t_xg(1, n + 1, CV_32FC1, xg); |
||||||
|
cv::Mat t_xxg(1, n + 1, CV_32FC1, xxg); |
||||||
|
|
||||||
|
optflow_farneback::g.upload(t_g); |
||||||
|
optflow_farneback::xg.upload(t_xg); |
||||||
|
optflow_farneback::xxg.upload(t_xxg); |
||||||
|
|
||||||
|
optflow_farneback::ig[0] = static_cast<float>(ig11); |
||||||
|
optflow_farneback::ig[1] = static_cast<float>(ig03); |
||||||
|
optflow_farneback::ig[2] = static_cast<float>(ig33); |
||||||
|
optflow_farneback::ig[3] = static_cast<float>(ig55); |
||||||
|
} |
||||||
|
|
||||||
|
void cv::ocl::FarnebackOpticalFlow::updateFlow_boxFilter( |
||||||
|
const oclMat& R0, const oclMat& R1, oclMat& flowx, oclMat &flowy, |
||||||
|
oclMat& M, oclMat &bufM, int blockSize, bool updateMatrices) |
||||||
|
{ |
||||||
|
optflow_farneback::boxFilter5Ocl(M, blockSize/2, bufM); |
||||||
|
|
||||||
|
swap(M, bufM); |
||||||
|
|
||||||
|
finish(); |
||||||
|
|
||||||
|
optflow_farneback::updateFlowOcl(M, flowx, flowy); |
||||||
|
|
||||||
|
if (updateMatrices) |
||||||
|
optflow_farneback::updateMatricesOcl(flowx, flowy, R0, R1, M); |
||||||
|
} |
||||||
|
|
||||||
|
|
||||||
|
void cv::ocl::FarnebackOpticalFlow::updateFlow_gaussianBlur( |
||||||
|
const oclMat& R0, const oclMat& R1, oclMat& flowx, oclMat& flowy, |
||||||
|
oclMat& M, oclMat &bufM, int blockSize, bool updateMatrices) |
||||||
|
{ |
||||||
|
optflow_farneback::gaussianBlur5Ocl(M, blockSize/2, bufM); |
||||||
|
|
||||||
|
swap(M, bufM); |
||||||
|
|
||||||
|
optflow_farneback::updateFlowOcl(M, flowx, flowy); |
||||||
|
|
||||||
|
if (updateMatrices) |
||||||
|
optflow_farneback::updateMatricesOcl(flowx, flowy, R0, R1, M); |
||||||
|
} |
||||||
|
|
||||||
|
|
||||||
|
void cv::ocl::FarnebackOpticalFlow::operator ()( |
||||||
|
const oclMat &frame0, const oclMat &frame1, oclMat &flowx, oclMat &flowy) |
||||||
|
{ |
||||||
|
CV_Assert(frame0.channels() == 1 && frame1.channels() == 1); |
||||||
|
CV_Assert(frame0.size() == frame1.size()); |
||||||
|
CV_Assert(polyN == 5 || polyN == 7); |
||||||
|
CV_Assert(!fastPyramids || std::abs(pyrScale - 0.5) < 1e-6); |
||||||
|
|
||||||
|
Size size = frame0.size(); |
||||||
|
oclMat prevFlowX, prevFlowY, curFlowX, curFlowY; |
||||||
|
|
||||||
|
flowx.create(size, CV_32F); |
||||||
|
flowy.create(size, CV_32F); |
||||||
|
oclMat flowx0 = flowx; |
||||||
|
oclMat flowy0 = flowy; |
||||||
|
|
||||||
|
// Crop unnecessary levels
|
||||||
|
double scale = 1; |
||||||
|
int numLevelsCropped = 0; |
||||||
|
for (; numLevelsCropped < numLevels; numLevelsCropped++) |
||||||
|
{ |
||||||
|
scale *= pyrScale; |
||||||
|
if (size.width*scale < MIN_SIZE || size.height*scale < MIN_SIZE) |
||||||
|
break; |
||||||
|
} |
||||||
|
|
||||||
|
frame0.convertTo(frames_[0], CV_32F); |
||||||
|
frame1.convertTo(frames_[1], CV_32F); |
||||||
|
|
||||||
|
if (fastPyramids) |
||||||
|
{ |
||||||
|
// Build Gaussian pyramids using pyrDown()
|
||||||
|
pyramid0_.resize(numLevelsCropped + 1); |
||||||
|
pyramid1_.resize(numLevelsCropped + 1); |
||||||
|
pyramid0_[0] = frames_[0]; |
||||||
|
pyramid1_[0] = frames_[1]; |
||||||
|
for (int i = 1; i <= numLevelsCropped; ++i) |
||||||
|
{ |
||||||
|
pyrDown(pyramid0_[i - 1], pyramid0_[i]); |
||||||
|
pyrDown(pyramid1_[i - 1], pyramid1_[i]); |
||||||
|
} |
||||||
|
} |
||||||
|
|
||||||
|
setPolynomialExpansionConsts(polyN, polySigma); |
||||||
|
|
||||||
|
for (int k = numLevelsCropped; k >= 0; k--) |
||||||
|
{ |
||||||
|
scale = 1; |
||||||
|
for (int i = 0; i < k; i++) |
||||||
|
scale *= pyrScale; |
||||||
|
|
||||||
|
double sigma = (1./scale - 1) * 0.5; |
||||||
|
int smoothSize = cvRound(sigma*5) | 1; |
||||||
|
smoothSize = std::max(smoothSize, 3); |
||||||
|
|
||||||
|
int width = cvRound(size.width*scale); |
||||||
|
int height = cvRound(size.height*scale); |
||||||
|
|
||||||
|
if (fastPyramids) |
||||||
|
{ |
||||||
|
width = pyramid0_[k].cols; |
||||||
|
height = pyramid0_[k].rows; |
||||||
|
} |
||||||
|
|
||||||
|
if (k > 0) |
||||||
|
{ |
||||||
|
curFlowX.create(height, width, CV_32F); |
||||||
|
curFlowY.create(height, width, CV_32F); |
||||||
|
} |
||||||
|
else |
||||||
|
{ |
||||||
|
curFlowX = flowx0; |
||||||
|
curFlowY = flowy0; |
||||||
|
} |
||||||
|
|
||||||
|
if (!prevFlowX.data) |
||||||
|
{ |
||||||
|
if (flags & cv::OPTFLOW_USE_INITIAL_FLOW) |
||||||
|
{ |
||||||
|
resize(flowx0, curFlowX, Size(width, height), 0, 0, INTER_LINEAR); |
||||||
|
resize(flowy0, curFlowY, Size(width, height), 0, 0, INTER_LINEAR); |
||||||
|
multiply(scale, curFlowX, curFlowX); |
||||||
|
multiply(scale, curFlowY, curFlowY); |
||||||
|
} |
||||||
|
else |
||||||
|
{ |
||||||
|
curFlowX.setTo(0); |
||||||
|
curFlowY.setTo(0); |
||||||
|
} |
||||||
|
} |
||||||
|
else |
||||||
|
{ |
||||||
|
resize(prevFlowX, curFlowX, Size(width, height), 0, 0, INTER_LINEAR); |
||||||
|
resize(prevFlowY, curFlowY, Size(width, height), 0, 0, INTER_LINEAR); |
||||||
|
multiply(1./pyrScale, curFlowX, curFlowX); |
||||||
|
multiply(1./pyrScale, curFlowY, curFlowY); |
||||||
|
} |
||||||
|
|
||||||
|
oclMat M = allocMatFromBuf(5*height, width, CV_32F, M_); |
||||||
|
oclMat bufM = allocMatFromBuf(5*height, width, CV_32F, bufM_); |
||||||
|
oclMat R[2] = |
||||||
|
{ |
||||||
|
allocMatFromBuf(5*height, width, CV_32F, R_[0]), |
||||||
|
allocMatFromBuf(5*height, width, CV_32F, R_[1]) |
||||||
|
}; |
||||||
|
|
||||||
|
if (fastPyramids) |
||||||
|
{ |
||||||
|
optflow_farneback::polynomialExpansionOcl(pyramid0_[k], polyN, R[0]); |
||||||
|
optflow_farneback::polynomialExpansionOcl(pyramid1_[k], polyN, R[1]); |
||||||
|
} |
||||||
|
else |
||||||
|
{ |
||||||
|
oclMat blurredFrame[2] = |
||||||
|
{ |
||||||
|
allocMatFromBuf(size.height, size.width, CV_32F, blurredFrame_[0]), |
||||||
|
allocMatFromBuf(size.height, size.width, CV_32F, blurredFrame_[1]) |
||||||
|
}; |
||||||
|
oclMat pyrLevel[2] = |
||||||
|
{ |
||||||
|
allocMatFromBuf(height, width, CV_32F, pyrLevel_[0]), |
||||||
|
allocMatFromBuf(height, width, CV_32F, pyrLevel_[1]) |
||||||
|
}; |
||||||
|
|
||||||
|
Mat g = getGaussianKernel(smoothSize, sigma, CV_32F); |
||||||
|
optflow_farneback::setGaussianBlurKernel(g.ptr<float>(smoothSize/2), smoothSize/2); |
||||||
|
|
||||||
|
for (int i = 0; i < 2; i++) |
||||||
|
{ |
||||||
|
optflow_farneback::gaussianBlurOcl(frames_[i], smoothSize/2, blurredFrame[i]); |
||||||
|
resize(blurredFrame[i], pyrLevel[i], Size(width, height), INTER_LINEAR); |
||||||
|
optflow_farneback::polynomialExpansionOcl(pyrLevel[i], polyN, R[i]); |
||||||
|
} |
||||||
|
} |
||||||
|
|
||||||
|
optflow_farneback::updateMatricesOcl(curFlowX, curFlowY, R[0], R[1], M); |
||||||
|
|
||||||
|
if (flags & OPTFLOW_FARNEBACK_GAUSSIAN) |
||||||
|
{ |
||||||
|
Mat g = getGaussianKernel(winSize, winSize/2*0.3f, CV_32F); |
||||||
|
optflow_farneback::setGaussianBlurKernel(g.ptr<float>(winSize/2), winSize/2); |
||||||
|
} |
||||||
|
for (int i = 0; i < numIters; i++) |
||||||
|
{ |
||||||
|
if (flags & OPTFLOW_FARNEBACK_GAUSSIAN) |
||||||
|
updateFlow_gaussianBlur(R[0], R[1], curFlowX, curFlowY, M, bufM, winSize, i < numIters-1); |
||||||
|
else |
||||||
|
updateFlow_boxFilter(R[0], R[1], curFlowX, curFlowY, M, bufM, winSize, i < numIters-1); |
||||||
|
} |
||||||
|
|
||||||
|
prevFlowX = curFlowX; |
||||||
|
prevFlowY = curFlowY; |
||||||
|
} |
||||||
|
|
||||||
|
flowx = curFlowX; |
||||||
|
flowy = curFlowY; |
||||||
|
} |
||||||
@ -0,0 +1,7 @@ |
|||||||
|
Building OpenCV from Source, using CMake and Command Line |
||||||
|
========================================================= |
||||||
|
|
||||||
|
cd ~/<my_working_directory> |
||||||
|
python opencv/platforms/ios/build_framework.py ios |
||||||
|
|
||||||
|
If everything's fine, a few minutes later you will get ~/<my_working_directory>/ios/opencv2.framework. You can add this framework to your Xcode projects. |
||||||
@ -1 +1,3 @@ |
|||||||
This folder contains toolchains and additional files that are needed for cross compitation. |
This folder contains toolchains and additional files that are needed for cross compilation. |
||||||
|
For more information see introduction tutorials for target platform in documentation: |
||||||
|
http://docs.opencv.org/doc/tutorials/introduction/table_of_content_introduction/table_of_content_introduction.html#table-of-content-introduction |
||||||
Some files were not shown because too many files have changed in this diff Show More
Loading…
Reference in new issue