mirror of https://github.com/opencv/opencv.git
Michael Pratt
10 years ago
committed by
arc
commit
c9876c3cd8
453 changed files with 11924 additions and 14701 deletions
@ -0,0 +1,61 @@ |
||||
# Main variables: |
||||
# OPENNI2_LIBRARY and OPENNI2_INCLUDES to link OpenCV modules with OpenNI2 |
||||
# HAVE_OPENNI2 for conditional compilation OpenCV with/without OpenNI2 |
||||
|
||||
if(NOT "${OPENNI2_LIB_DIR}" STREQUAL "${OPENNI2_LIB_DIR_INTERNAL}") |
||||
unset(OPENNI2_LIBRARY CACHE) |
||||
unset(OPENNI2_LIB_DIR CACHE) |
||||
endif() |
||||
|
||||
if(NOT "${OPENNI2_INCLUDE_DIR}" STREQUAL "${OPENNI2_INCLUDE_DIR_INTERNAL}") |
||||
unset(OPENNI2_INCLUDES CACHE) |
||||
unset(OPENNI2_INCLUDE_DIR CACHE) |
||||
endif() |
||||
|
||||
if(WIN32) |
||||
if(NOT (MSVC64 OR MINGW64)) |
||||
find_file(OPENNI2_INCLUDES "OpenNI.h" PATHS "$ENV{OPEN_NI_INSTALL_PATH}Include" DOC "OpenNI2 c++ interface header") |
||||
find_library(OPENNI2_LIBRARY "OpenNI2" PATHS $ENV{OPENNI2_LIB} DOC "OpenNI2 library") |
||||
else() |
||||
find_file(OPENNI2_INCLUDES "OpenNI.h" PATHS "$ENV{OPEN_NI_INSTALL_PATH64}Include" DOC "OpenNI2 c++ interface header") |
||||
find_library(OPENNI2_LIBRARY "OpenNI2" PATHS $ENV{OPENNI2_LIB64} DOC "OpenNI2 library") |
||||
endif() |
||||
elseif(UNIX OR APPLE) |
||||
find_file(OPENNI2_INCLUDES "OpenNI.h" PATHS "/usr/include/ni2" "/usr/include/openni2" DOC "OpenNI2 c++ interface header") |
||||
find_library(OPENNI2_LIBRARY "OpenNI2" PATHS "/usr/lib" DOC "OpenNI2 library") |
||||
endif() |
||||
|
||||
if(OPENNI2_LIBRARY AND OPENNI2_INCLUDES) |
||||
set(HAVE_OPENNI2 TRUE) |
||||
endif() #if(OPENNI2_LIBRARY AND OPENNI2_INCLUDES) |
||||
|
||||
get_filename_component(OPENNI2_LIB_DIR "${OPENNI2_LIBRARY}" PATH) |
||||
get_filename_component(OPENNI2_INCLUDE_DIR ${OPENNI2_INCLUDES} PATH) |
||||
|
||||
if(HAVE_OPENNI2) |
||||
set(OPENNI2_LIB_DIR "${OPENNI2_LIB_DIR}" CACHE PATH "Path to OpenNI2 libraries" FORCE) |
||||
set(OPENNI2_INCLUDE_DIR "${OPENNI2_INCLUDE_DIR}" CACHE PATH "Path to OpenNI2 headers" FORCE) |
||||
endif() |
||||
|
||||
if(OPENNI2_LIBRARY) |
||||
set(OPENNI2_LIB_DIR_INTERNAL "${OPENNI2_LIB_DIR}" CACHE INTERNAL "This is the value of the last time OPENNI_LIB_DIR was set successfully." FORCE) |
||||
else() |
||||
message( WARNING, " OpenNI2 library directory (set by OPENNI2_LIB_DIR variable) is not found or does not have OpenNI2 libraries." ) |
||||
endif() |
||||
|
||||
if(OPENNI2_INCLUDES) |
||||
set(OPENNI2_INCLUDE_DIR_INTERNAL "${OPENNI2_INCLUDE_DIR}" CACHE INTERNAL "This is the value of the last time OPENNI2_INCLUDE_DIR was set successfully." FORCE) |
||||
else() |
||||
message( WARNING, " OpenNI2 include directory (set by OPENNI2_INCLUDE_DIR variable) is not found or does not have OpenNI2 include files." ) |
||||
endif() |
||||
|
||||
mark_as_advanced(FORCE OPENNI2_LIBRARY) |
||||
mark_as_advanced(FORCE OPENNI2_INCLUDES) |
||||
|
||||
if(HAVE_OPENNI2) |
||||
ocv_parse_header("${OPENNI2_INCLUDE_DIR}/OniVersion.h" ONI_VERSION_LINE ONI_VERSION_MAJOR ONI_VERSION_MINOR ONI_VERSION_MAINTENANCE ONI_VERSION_BUILD) |
||||
if(ONI_VERSION_MAJOR) |
||||
set(OPENNI2_VERSION_STRING ${ONI_VERSION_MAJOR}.${ONI_VERSION_MINOR}.${ONI_VERSION_MAINTENANCE} CACHE INTERNAL "OpenNI2 version") |
||||
set(OPENNI2_VERSION_BUILD ${ONI_VERSION_BUILD} CACHE INTERNAL "OpenNI2 build version") |
||||
endif() |
||||
endif() |
||||
@ -1,4 +1,7 @@ |
||||
if(IOS) |
||||
configure_file("${OpenCV_SOURCE_DIR}/platforms/ios/Info.plist.in" |
||||
"${CMAKE_BINARY_DIR}/ios/Info.plist") |
||||
elseif(APPLE) |
||||
configure_file("${OpenCV_SOURCE_DIR}/platforms/osx/Info.plist.in" |
||||
"${CMAKE_BINARY_DIR}/osx/Info.plist") |
||||
endif() |
||||
|
||||
@ -0,0 +1,161 @@ |
||||
.. _akazeMatching: |
||||
|
||||
|
||||
AKAZE local features matching |
||||
****************************** |
||||
|
||||
Introduction |
||||
------------------ |
||||
|
||||
In this tutorial we will learn how to use [AKAZE]_ local features to detect and match keypoints on two images. |
||||
|
||||
We will find keypoints on a pair of images with given homography matrix, |
||||
match them and count the number of inliers (i. e. matches that fit in the given homography). |
||||
|
||||
You can find expanded version of this example here: https://github.com/pablofdezalc/test_kaze_akaze_opencv |
||||
|
||||
.. [AKAZE] Fast Explicit Diffusion for Accelerated Features in Nonlinear Scale Spaces. Pablo F. Alcantarilla, Jesús Nuevo and Adrien Bartoli. In British Machine Vision Conference (BMVC), Bristol, UK, September 2013. |
||||
|
||||
Data |
||||
------------------ |
||||
We are going to use images 1 and 3 from *Graffity* sequence of Oxford dataset. |
||||
|
||||
.. image:: images/graf.png |
||||
:height: 200pt |
||||
:width: 320pt |
||||
:alt: Graffity |
||||
:align: center |
||||
|
||||
Homography is given by a 3 by 3 matrix: |
||||
|
||||
.. code-block:: none |
||||
|
||||
7.6285898e-01 -2.9922929e-01 2.2567123e+02 |
||||
3.3443473e-01 1.0143901e+00 -7.6999973e+01 |
||||
3.4663091e-04 -1.4364524e-05 1.0000000e+00 |
||||
|
||||
You can find the images (*graf1.png*, *graf3.png*) and homography (*H1to3p.xml*) in *opencv/samples/cpp*. |
||||
|
||||
Source Code |
||||
=========== |
||||
.. literalinclude:: ../../../../samples/cpp/tutorial_code/features2D/AKAZE_match.cpp |
||||
:language: cpp |
||||
:linenos: |
||||
:tab-width: 4 |
||||
|
||||
Explanation |
||||
=========== |
||||
|
||||
1. **Load images and homography** |
||||
|
||||
.. code-block:: cpp |
||||
|
||||
Mat img1 = imread("graf1.png", IMREAD_GRAYSCALE); |
||||
Mat img2 = imread("graf3.png", IMREAD_GRAYSCALE); |
||||
|
||||
Mat homography; |
||||
FileStorage fs("H1to3p.xml", FileStorage::READ); |
||||
fs.getFirstTopLevelNode() >> homography; |
||||
|
||||
We are loading grayscale images here. Homography is stored in the xml created with FileStorage. |
||||
|
||||
2. **Detect keypoints and compute descriptors using AKAZE** |
||||
|
||||
.. code-block:: cpp |
||||
|
||||
vector<KeyPoint> kpts1, kpts2; |
||||
Mat desc1, desc2; |
||||
|
||||
AKAZE akaze; |
||||
akaze(img1, noArray(), kpts1, desc1); |
||||
akaze(img2, noArray(), kpts2, desc2); |
||||
|
||||
We create AKAZE object and use it's *operator()* functionality. Since we don't need the *mask* parameter, *noArray()* is used. |
||||
|
||||
3. **Use brute-force matcher to find 2-nn matches** |
||||
|
||||
.. code-block:: cpp |
||||
|
||||
BFMatcher matcher(NORM_HAMMING); |
||||
vector< vector<DMatch> > nn_matches; |
||||
matcher.knnMatch(desc1, desc2, nn_matches, 2); |
||||
|
||||
We use Hamming distance, because AKAZE uses binary descriptor by default. |
||||
|
||||
4. **Use 2-nn matches to find correct keypoint matches** |
||||
|
||||
.. code-block:: cpp |
||||
|
||||
for(size_t i = 0; i < nn_matches.size(); i++) { |
||||
DMatch first = nn_matches[i][0]; |
||||
float dist1 = nn_matches[i][0].distance; |
||||
float dist2 = nn_matches[i][1].distance; |
||||
|
||||
if(dist1 < nn_match_ratio * dist2) { |
||||
matched1.push_back(kpts1[first.queryIdx]); |
||||
matched2.push_back(kpts2[first.trainIdx]); |
||||
} |
||||
} |
||||
|
||||
If the closest match is *ratio* closer than the second closest one, then the match is correct. |
||||
|
||||
5. **Check if our matches fit in the homography model** |
||||
|
||||
.. code-block:: cpp |
||||
|
||||
for(int i = 0; i < matched1.size(); i++) { |
||||
Mat col = Mat::ones(3, 1, CV_64F); |
||||
col.at<double>(0) = matched1[i].pt.x; |
||||
col.at<double>(1) = matched1[i].pt.y; |
||||
|
||||
col = homography * col; |
||||
col /= col.at<double>(2); |
||||
float dist = sqrt( pow(col.at<double>(0) - matched2[i].pt.x, 2) + |
||||
pow(col.at<double>(1) - matched2[i].pt.y, 2)); |
||||
|
||||
if(dist < inlier_threshold) { |
||||
int new_i = inliers1.size(); |
||||
inliers1.push_back(matched1[i]); |
||||
inliers2.push_back(matched2[i]); |
||||
good_matches.push_back(DMatch(new_i, new_i, 0)); |
||||
} |
||||
} |
||||
|
||||
If the distance from first keypoint's projection to the second keypoint is less than threshold, then it it fits in the homography. |
||||
|
||||
We create a new set of matches for the inliers, because it is required by the drawing function. |
||||
|
||||
6. **Output results** |
||||
|
||||
.. code-block:: cpp |
||||
|
||||
Mat res; |
||||
drawMatches(img1, inliers1, img2, inliers2, good_matches, res); |
||||
imwrite("res.png", res); |
||||
... |
||||
|
||||
Here we save the resulting image and print some statistics. |
||||
|
||||
Results |
||||
======= |
||||
|
||||
Found matches |
||||
-------------- |
||||
|
||||
.. image:: images/res.png |
||||
:height: 200pt |
||||
:width: 320pt |
||||
:alt: Matches |
||||
:align: center |
||||
|
||||
A-KAZE Matching Results |
||||
-------------------------- |
||||
Keypoints 1: 2943 |
||||
|
||||
Keypoints 2: 3511 |
||||
|
||||
Matches: 447 |
||||
|
||||
Inliers: 308 |
||||
|
||||
Inliers Ratio: 0.689038 |
||||
|
After Width: | Height: | Size: 2.0 MiB |
|
After Width: | Height: | Size: 1.8 MiB |
|
After Width: | Height: | Size: 63 KiB |
@ -0,0 +1,864 @@ |
||||
// This file is part of OpenCV project. |
||||
// It is subject to the license terms in the LICENSE file found in the top-level directory |
||||
// of this distribution and at http://opencv.org/license.html. |
||||
|
||||
// Copyright (C) 2014, Itseez, Inc., all rights reserved. |
||||
// Third party copyrights are property of their respective owners. |
||||
|
||||
#define SQRT_2 0.707106781188f |
||||
#define sin_120 0.866025403784f |
||||
#define fft5_2 0.559016994374f |
||||
#define fft5_3 -0.951056516295f |
||||
#define fft5_4 -1.538841768587f |
||||
#define fft5_5 0.363271264002f |
||||
|
||||
__attribute__((always_inline)) |
||||
float2 mul_float2(float2 a, float2 b) { |
||||
return (float2)(fma(a.x, b.x, -a.y * b.y), fma(a.x, b.y, a.y * b.x)); |
||||
} |
||||
|
||||
__attribute__((always_inline)) |
||||
float2 twiddle(float2 a) { |
||||
return (float2)(a.y, -a.x); |
||||
} |
||||
|
||||
__attribute__((always_inline)) |
||||
void butterfly2(float2 a0, float2 a1, __local float2* smem, __global const float2* twiddles, |
||||
const int x, const int block_size) |
||||
{ |
||||
const int k = x & (block_size - 1); |
||||
a1 = mul_float2(twiddles[k], a1); |
||||
const int dst_ind = (x << 1) - k; |
||||
|
||||
smem[dst_ind] = a0 + a1; |
||||
smem[dst_ind+block_size] = a0 - a1; |
||||
} |
||||
|
||||
__attribute__((always_inline)) |
||||
void butterfly4(float2 a0, float2 a1, float2 a2, float2 a3, __local float2* smem, __global const float2* twiddles, |
||||
const int x, const int block_size) |
||||
{ |
||||
const int k = x & (block_size - 1); |
||||
a1 = mul_float2(twiddles[k], a1); |
||||
a2 = mul_float2(twiddles[k + block_size], a2); |
||||
a3 = mul_float2(twiddles[k + 2*block_size], a3); |
||||
|
||||
const int dst_ind = ((x - k) << 2) + k; |
||||
|
||||
float2 b0 = a0 + a2; |
||||
a2 = a0 - a2; |
||||
float2 b1 = a1 + a3; |
||||
a3 = twiddle(a1 - a3); |
||||
|
||||
smem[dst_ind] = b0 + b1; |
||||
smem[dst_ind + block_size] = a2 + a3; |
||||
smem[dst_ind + 2*block_size] = b0 - b1; |
||||
smem[dst_ind + 3*block_size] = a2 - a3; |
||||
} |
||||
|
||||
__attribute__((always_inline)) |
||||
void butterfly3(float2 a0, float2 a1, float2 a2, __local float2* smem, __global const float2* twiddles, |
||||
const int x, const int block_size) |
||||
{ |
||||
const int k = x % block_size; |
||||
a1 = mul_float2(twiddles[k], a1); |
||||
a2 = mul_float2(twiddles[k+block_size], a2); |
||||
const int dst_ind = ((x - k) * 3) + k; |
||||
|
||||
float2 b1 = a1 + a2; |
||||
a2 = twiddle(sin_120*(a1 - a2)); |
||||
float2 b0 = a0 - (float2)(0.5f)*b1; |
||||
|
||||
smem[dst_ind] = a0 + b1; |
||||
smem[dst_ind + block_size] = b0 + a2; |
||||
smem[dst_ind + 2*block_size] = b0 - a2; |
||||
} |
||||
|
||||
__attribute__((always_inline)) |
||||
void butterfly5(float2 a0, float2 a1, float2 a2, float2 a3, float2 a4, __local float2* smem, __global const float2* twiddles, |
||||
const int x, const int block_size) |
||||
{ |
||||
const int k = x % block_size; |
||||
a1 = mul_float2(twiddles[k], a1); |
||||
a2 = mul_float2(twiddles[k + block_size], a2); |
||||
a3 = mul_float2(twiddles[k+2*block_size], a3); |
||||
a4 = mul_float2(twiddles[k+3*block_size], a4); |
||||
|
||||
const int dst_ind = ((x - k) * 5) + k; |
||||
__local float2* dst = smem + dst_ind; |
||||
|
||||
float2 b0, b1, b5; |
||||
|
||||
b1 = a1 + a4; |
||||
a1 -= a4; |
||||
|
||||
a4 = a3 + a2; |
||||
a3 -= a2; |
||||
|
||||
a2 = b1 + a4; |
||||
b0 = a0 - (float2)0.25f * a2; |
||||
|
||||
b1 = fft5_2 * (b1 - a4); |
||||
a4 = fft5_3 * (float2)(-a1.y - a3.y, a1.x + a3.x); |
||||
b5 = (float2)(a4.x - fft5_5 * a1.y, a4.y + fft5_5 * a1.x); |
||||
|
||||
a4.x += fft5_4 * a3.y; |
||||
a4.y -= fft5_4 * a3.x; |
||||
|
||||
a1 = b0 + b1; |
||||
b0 -= b1; |
||||
|
||||
dst[0] = a0 + a2; |
||||
dst[block_size] = a1 + a4; |
||||
dst[2 * block_size] = b0 + b5; |
||||
dst[3 * block_size] = b0 - b5; |
||||
dst[4 * block_size] = a1 - a4; |
||||
} |
||||
|
||||
__attribute__((always_inline)) |
||||
void fft_radix2(__local float2* smem, __global const float2* twiddles, const int x, const int block_size, const int t) |
||||
{ |
||||
float2 a0, a1; |
||||
|
||||
if (x < t) |
||||
{ |
||||
a0 = smem[x]; |
||||
a1 = smem[x+t]; |
||||
} |
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
|
||||
if (x < t) |
||||
butterfly2(a0, a1, smem, twiddles, x, block_size); |
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
} |
||||
|
||||
__attribute__((always_inline)) |
||||
void fft_radix2_B2(__local float2* smem, __global const float2* twiddles, const int x1, const int block_size, const int t) |
||||
{ |
||||
const int x2 = x1 + t/2; |
||||
float2 a0, a1, a2, a3; |
||||
|
||||
if (x1 < t/2) |
||||
{ |
||||
a0 = smem[x1]; a1 = smem[x1+t]; |
||||
a2 = smem[x2]; a3 = smem[x2+t]; |
||||
} |
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
|
||||
if (x1 < t/2) |
||||
{ |
||||
butterfly2(a0, a1, smem, twiddles, x1, block_size); |
||||
butterfly2(a2, a3, smem, twiddles, x2, block_size); |
||||
} |
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
} |
||||
|
||||
__attribute__((always_inline)) |
||||
void fft_radix2_B3(__local float2* smem, __global const float2* twiddles, const int x1, const int block_size, const int t) |
||||
{ |
||||
const int x2 = x1 + t/3; |
||||
const int x3 = x1 + 2*t/3; |
||||
float2 a0, a1, a2, a3, a4, a5; |
||||
|
||||
if (x1 < t/3) |
||||
{ |
||||
a0 = smem[x1]; a1 = smem[x1+t]; |
||||
a2 = smem[x2]; a3 = smem[x2+t]; |
||||
a4 = smem[x3]; a5 = smem[x3+t]; |
||||
} |
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
|
||||
if (x1 < t/3) |
||||
{ |
||||
butterfly2(a0, a1, smem, twiddles, x1, block_size); |
||||
butterfly2(a2, a3, smem, twiddles, x2, block_size); |
||||
butterfly2(a4, a5, smem, twiddles, x3, block_size); |
||||
} |
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
} |
||||
|
||||
__attribute__((always_inline)) |
||||
void fft_radix2_B4(__local float2* smem, __global const float2* twiddles, const int x1, const int block_size, const int t) |
||||
{ |
||||
const int thread_block = t/4; |
||||
const int x2 = x1 + thread_block; |
||||
const int x3 = x1 + 2*thread_block; |
||||
const int x4 = x1 + 3*thread_block; |
||||
float2 a0, a1, a2, a3, a4, a5, a6, a7; |
||||
|
||||
if (x1 < t/4) |
||||
{ |
||||
a0 = smem[x1]; a1 = smem[x1+t]; |
||||
a2 = smem[x2]; a3 = smem[x2+t]; |
||||
a4 = smem[x3]; a5 = smem[x3+t]; |
||||
a6 = smem[x4]; a7 = smem[x4+t]; |
||||
} |
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
|
||||
if (x1 < t/4) |
||||
{ |
||||
butterfly2(a0, a1, smem, twiddles, x1, block_size); |
||||
butterfly2(a2, a3, smem, twiddles, x2, block_size); |
||||
butterfly2(a4, a5, smem, twiddles, x3, block_size); |
||||
butterfly2(a6, a7, smem, twiddles, x4, block_size); |
||||
} |
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
} |
||||
|
||||
__attribute__((always_inline)) |
||||
void fft_radix2_B5(__local float2* smem, __global const float2* twiddles, const int x1, const int block_size, const int t) |
||||
{ |
||||
const int thread_block = t/5; |
||||
const int x2 = x1 + thread_block; |
||||
const int x3 = x1 + 2*thread_block; |
||||
const int x4 = x1 + 3*thread_block; |
||||
const int x5 = x1 + 4*thread_block; |
||||
float2 a0, a1, a2, a3, a4, a5, a6, a7, a8, a9; |
||||
|
||||
if (x1 < t/5) |
||||
{ |
||||
a0 = smem[x1]; a1 = smem[x1+t]; |
||||
a2 = smem[x2]; a3 = smem[x2+t]; |
||||
a4 = smem[x3]; a5 = smem[x3+t]; |
||||
a6 = smem[x4]; a7 = smem[x4+t]; |
||||
a8 = smem[x5]; a9 = smem[x5+t]; |
||||
} |
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
|
||||
if (x1 < t/5) |
||||
{ |
||||
butterfly2(a0, a1, smem, twiddles, x1, block_size); |
||||
butterfly2(a2, a3, smem, twiddles, x2, block_size); |
||||
butterfly2(a4, a5, smem, twiddles, x3, block_size); |
||||
butterfly2(a6, a7, smem, twiddles, x4, block_size); |
||||
butterfly2(a8, a9, smem, twiddles, x5, block_size); |
||||
} |
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
} |
||||
|
||||
__attribute__((always_inline)) |
||||
void fft_radix4(__local float2* smem, __global const float2* twiddles, const int x, const int block_size, const int t) |
||||
{ |
||||
float2 a0, a1, a2, a3; |
||||
|
||||
if (x < t) |
||||
{ |
||||
a0 = smem[x]; a1 = smem[x+t]; a2 = smem[x+2*t]; a3 = smem[x+3*t]; |
||||
} |
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
|
||||
if (x < t) |
||||
butterfly4(a0, a1, a2, a3, smem, twiddles, x, block_size); |
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
} |
||||
|
||||
__attribute__((always_inline)) |
||||
void fft_radix4_B2(__local float2* smem, __global const float2* twiddles, const int x1, const int block_size, const int t) |
||||
{ |
||||
const int x2 = x1 + t/2; |
||||
float2 a0, a1, a2, a3, a4, a5, a6, a7; |
||||
|
||||
if (x1 < t/2) |
||||
{ |
||||
a0 = smem[x1]; a1 = smem[x1+t]; a2 = smem[x1+2*t]; a3 = smem[x1+3*t]; |
||||
a4 = smem[x2]; a5 = smem[x2+t]; a6 = smem[x2+2*t]; a7 = smem[x2+3*t]; |
||||
} |
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
|
||||
if (x1 < t/2) |
||||
{ |
||||
butterfly4(a0, a1, a2, a3, smem, twiddles, x1, block_size); |
||||
butterfly4(a4, a5, a6, a7, smem, twiddles, x2, block_size); |
||||
} |
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
} |
||||
|
||||
__attribute__((always_inline)) |
||||
void fft_radix4_B3(__local float2* smem, __global const float2* twiddles, const int x1, const int block_size, const int t) |
||||
{ |
||||
const int x2 = x1 + t/3; |
||||
const int x3 = x2 + t/3; |
||||
float2 a0, a1, a2, a3, a4, a5, a6, a7, a8, a9, a10, a11; |
||||
|
||||
if (x1 < t/3) |
||||
{ |
||||
a0 = smem[x1]; a1 = smem[x1+t]; a2 = smem[x1+2*t]; a3 = smem[x1+3*t]; |
||||
a4 = smem[x2]; a5 = smem[x2+t]; a6 = smem[x2+2*t]; a7 = smem[x2+3*t]; |
||||
a8 = smem[x3]; a9 = smem[x3+t]; a10 = smem[x3+2*t]; a11 = smem[x3+3*t]; |
||||
} |
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
|
||||
if (x1 < t/3) |
||||
{ |
||||
butterfly4(a0, a1, a2, a3, smem, twiddles, x1, block_size); |
||||
butterfly4(a4, a5, a6, a7, smem, twiddles, x2, block_size); |
||||
butterfly4(a8, a9, a10, a11, smem, twiddles, x3, block_size); |
||||
} |
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
} |
||||
|
||||
__attribute__((always_inline)) |
||||
void fft_radix8(__local float2* smem, __global const float2* twiddles, const int x, const int block_size, const int t) |
||||
{ |
||||
const int k = x % block_size; |
||||
float2 a0, a1, a2, a3, a4, a5, a6, a7; |
||||
|
||||
if (x < t) |
||||
{ |
||||
int tw_ind = block_size / 8; |
||||
|
||||
a0 = smem[x]; |
||||
a1 = mul_float2(twiddles[k], smem[x + t]); |
||||
a2 = mul_float2(twiddles[k + block_size],smem[x+2*t]); |
||||
a3 = mul_float2(twiddles[k+2*block_size],smem[x+3*t]); |
||||
a4 = mul_float2(twiddles[k+3*block_size],smem[x+4*t]); |
||||
a5 = mul_float2(twiddles[k+4*block_size],smem[x+5*t]); |
||||
a6 = mul_float2(twiddles[k+5*block_size],smem[x+6*t]); |
||||
a7 = mul_float2(twiddles[k+6*block_size],smem[x+7*t]); |
||||
|
||||
float2 b0, b1, b6, b7; |
||||
|
||||
b0 = a0 + a4; |
||||
a4 = a0 - a4; |
||||
b1 = a1 + a5; |
||||
a5 = a1 - a5; |
||||
a5 = (float2)(SQRT_2) * (float2)(a5.x + a5.y, -a5.x + a5.y); |
||||
b6 = twiddle(a2 - a6); |
||||
a2 = a2 + a6; |
||||
b7 = a3 - a7; |
||||
b7 = (float2)(SQRT_2) * (float2)(-b7.x + b7.y, -b7.x - b7.y); |
||||
a3 = a3 + a7; |
||||
|
||||
a0 = b0 + a2; |
||||
a2 = b0 - a2; |
||||
a1 = b1 + a3; |
||||
a3 = twiddle(b1 - a3); |
||||
a6 = a4 - b6; |
||||
a4 = a4 + b6; |
||||
a7 = twiddle(a5 - b7); |
||||
a5 = a5 + b7; |
||||
|
||||
} |
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
|
||||
if (x < t) |
||||
{ |
||||
const int dst_ind = ((x - k) << 3) + k; |
||||
__local float2* dst = smem + dst_ind; |
||||
|
||||
dst[0] = a0 + a1; |
||||
dst[block_size] = a4 + a5; |
||||
dst[2 * block_size] = a2 + a3; |
||||
dst[3 * block_size] = a6 + a7; |
||||
dst[4 * block_size] = a0 - a1; |
||||
dst[5 * block_size] = a4 - a5; |
||||
dst[6 * block_size] = a2 - a3; |
||||
dst[7 * block_size] = a6 - a7; |
||||
} |
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
} |
||||
|
||||
__attribute__((always_inline)) |
||||
void fft_radix3(__local float2* smem, __global const float2* twiddles, const int x, const int block_size, const int t) |
||||
{ |
||||
float2 a0, a1, a2; |
||||
|
||||
if (x < t) |
||||
{ |
||||
a0 = smem[x]; a1 = smem[x+t]; a2 = smem[x+2*t]; |
||||
} |
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
|
||||
if (x < t) |
||||
butterfly3(a0, a1, a2, smem, twiddles, x, block_size); |
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
} |
||||
|
||||
__attribute__((always_inline)) |
||||
void fft_radix3_B2(__local float2* smem, __global const float2* twiddles, const int x1, const int block_size, const int t) |
||||
{ |
||||
const int x2 = x1 + t/2; |
||||
float2 a0, a1, a2, a3, a4, a5; |
||||
|
||||
if (x1 < t/2) |
||||
{ |
||||
a0 = smem[x1]; a1 = smem[x1+t]; a2 = smem[x1+2*t]; |
||||
a3 = smem[x2]; a4 = smem[x2+t]; a5 = smem[x2+2*t]; |
||||
} |
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
|
||||
if (x1 < t/2) |
||||
{ |
||||
butterfly3(a0, a1, a2, smem, twiddles, x1, block_size); |
||||
butterfly3(a3, a4, a5, smem, twiddles, x2, block_size); |
||||
} |
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
} |
||||
|
||||
__attribute__((always_inline)) |
||||
void fft_radix3_B3(__local float2* smem, __global const float2* twiddles, const int x1, const int block_size, const int t) |
||||
{ |
||||
const int x2 = x1 + t/3; |
||||
const int x3 = x2 + t/3; |
||||
float2 a0, a1, a2, a3, a4, a5, a6, a7, a8; |
||||
|
||||
if (x1 < t/2) |
||||
{ |
||||
a0 = smem[x1]; a1 = smem[x1+t]; a2 = smem[x1+2*t]; |
||||
a3 = smem[x2]; a4 = smem[x2+t]; a5 = smem[x2+2*t]; |
||||
a6 = smem[x3]; a7 = smem[x3+t]; a8 = smem[x3+2*t]; |
||||
} |
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
|
||||
if (x1 < t/2) |
||||
{ |
||||
butterfly3(a0, a1, a2, smem, twiddles, x1, block_size); |
||||
butterfly3(a3, a4, a5, smem, twiddles, x2, block_size); |
||||
butterfly3(a6, a7, a8, smem, twiddles, x3, block_size); |
||||
} |
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
} |
||||
|
||||
__attribute__((always_inline)) |
||||
void fft_radix3_B4(__local float2* smem, __global const float2* twiddles, const int x1, const int block_size, const int t) |
||||
{ |
||||
const int thread_block = t/4; |
||||
const int x2 = x1 + thread_block; |
||||
const int x3 = x1 + 2*thread_block; |
||||
const int x4 = x1 + 3*thread_block; |
||||
float2 a0, a1, a2, a3, a4, a5, a6, a7, a8, a9, a10, a11; |
||||
|
||||
if (x1 < t/4) |
||||
{ |
||||
a0 = smem[x1]; a1 = smem[x1+t]; a2 = smem[x1+2*t]; |
||||
a3 = smem[x2]; a4 = smem[x2+t]; a5 = smem[x2+2*t]; |
||||
a6 = smem[x3]; a7 = smem[x3+t]; a8 = smem[x3+2*t]; |
||||
a9 = smem[x4]; a10 = smem[x4+t]; a11 = smem[x4+2*t]; |
||||
} |
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
|
||||
if (x1 < t/4) |
||||
{ |
||||
butterfly3(a0, a1, a2, smem, twiddles, x1, block_size); |
||||
butterfly3(a3, a4, a5, smem, twiddles, x2, block_size); |
||||
butterfly3(a6, a7, a8, smem, twiddles, x3, block_size); |
||||
butterfly3(a9, a10, a11, smem, twiddles, x4, block_size); |
||||
} |
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
} |
||||
|
||||
__attribute__((always_inline)) |
||||
void fft_radix5(__local float2* smem, __global const float2* twiddles, const int x, const int block_size, const int t) |
||||
{ |
||||
const int k = x % block_size; |
||||
float2 a0, a1, a2, a3, a4; |
||||
|
||||
if (x < t) |
||||
{ |
||||
a0 = smem[x]; a1 = smem[x + t]; a2 = smem[x+2*t]; a3 = smem[x+3*t]; a4 = smem[x+4*t]; |
||||
} |
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
|
||||
if (x < t) |
||||
butterfly5(a0, a1, a2, a3, a4, smem, twiddles, x, block_size); |
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
} |
||||
|
||||
__attribute__((always_inline)) |
||||
void fft_radix5_B2(__local float2* smem, __global const float2* twiddles, const int x1, const int block_size, const int t) |
||||
{ |
||||
const int x2 = x1+t/2; |
||||
float2 a0, a1, a2, a3, a4, a5, a6, a7, a8, a9; |
||||
|
||||
if (x1 < t/2) |
||||
{ |
||||
a0 = smem[x1]; a1 = smem[x1 + t]; a2 = smem[x1+2*t]; a3 = smem[x1+3*t]; a4 = smem[x1+4*t]; |
||||
a5 = smem[x2]; a6 = smem[x2 + t]; a7 = smem[x2+2*t]; a8 = smem[x2+3*t]; a9 = smem[x2+4*t]; |
||||
} |
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
|
||||
if (x1 < t/2) |
||||
{ |
||||
butterfly5(a0, a1, a2, a3, a4, smem, twiddles, x1, block_size); |
||||
butterfly5(a5, a6, a7, a8, a9, smem, twiddles, x2, block_size); |
||||
} |
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
} |
||||
|
||||
#ifdef DFT_SCALE |
||||
#define SCALE_VAL(x, scale) x*scale |
||||
#else |
||||
#define SCALE_VAL(x, scale) x |
||||
#endif |
||||
|
||||
__kernel void fft_multi_radix_rows(__global const uchar* src_ptr, int src_step, int src_offset, int src_rows, int src_cols, |
||||
__global uchar* dst_ptr, int dst_step, int dst_offset, int dst_rows, int dst_cols, |
||||
__global float2* twiddles_ptr, const int t, const int nz) |
||||
{ |
||||
const int x = get_global_id(0); |
||||
const int y = get_group_id(1); |
||||
const int block_size = LOCAL_SIZE/kercn; |
||||
if (y < nz) |
||||
{ |
||||
__local float2 smem[LOCAL_SIZE]; |
||||
__global const float2* twiddles = (__global float2*) twiddles_ptr; |
||||
const int ind = x; |
||||
#ifdef IS_1D |
||||
float scale = 1.f/dst_cols; |
||||
#else |
||||
float scale = 1.f/(dst_cols*dst_rows); |
||||
#endif |
||||
|
||||
#ifdef COMPLEX_INPUT |
||||
__global const float2* src = (__global const float2*)(src_ptr + mad24(y, src_step, mad24(x, (int)(sizeof(float)*2), src_offset))); |
||||
#pragma unroll |
||||
for (int i=0; i<kercn; i++) |
||||
smem[x+i*block_size] = src[i*block_size]; |
||||
#else |
||||
__global const float* src = (__global const float*)(src_ptr + mad24(y, src_step, mad24(x, (int)sizeof(float), src_offset))); |
||||
#pragma unroll |
||||
for (int i=0; i<kercn; i++) |
||||
smem[x+i*block_size] = (float2)(src[i*block_size], 0.f); |
||||
#endif |
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
|
||||
RADIX_PROCESS; |
||||
|
||||
#ifdef COMPLEX_OUTPUT |
||||
#ifdef NO_CONJUGATE |
||||
// copy result without complex conjugate |
||||
const int cols = dst_cols/2 + 1; |
||||
#else |
||||
const int cols = dst_cols; |
||||
#endif |
||||
|
||||
__global float2* dst = (__global float2*)(dst_ptr + mad24(y, dst_step, dst_offset)); |
||||
#pragma unroll |
||||
for (int i=x; i<cols; i+=block_size) |
||||
dst[i] = SCALE_VAL(smem[i], scale); |
||||
#else |
||||
// pack row to CCS |
||||
__local float* smem_1cn = (__local float*) smem; |
||||
__global float* dst = (__global float*)(dst_ptr + mad24(y, dst_step, dst_offset)); |
||||
for (int i=x; i<dst_cols-1; i+=block_size) |
||||
dst[i+1] = SCALE_VAL(smem_1cn[i+2], scale); |
||||
if (x == 0) |
||||
dst[0] = SCALE_VAL(smem_1cn[0], scale); |
||||
#endif |
||||
} |
||||
else |
||||
{ |
||||
// fill with zero other rows |
||||
#ifdef COMPLEX_OUTPUT |
||||
__global float2* dst = (__global float2*)(dst_ptr + mad24(y, dst_step, dst_offset)); |
||||
#else |
||||
__global float* dst = (__global float*)(dst_ptr + mad24(y, dst_step, dst_offset)); |
||||
#endif |
||||
#pragma unroll |
||||
for (int i=x; i<dst_cols; i+=block_size) |
||||
dst[i] = 0.f; |
||||
} |
||||
} |
||||
|
||||
__kernel void fft_multi_radix_cols(__global const uchar* src_ptr, int src_step, int src_offset, int src_rows, int src_cols, |
||||
__global uchar* dst_ptr, int dst_step, int dst_offset, int dst_rows, int dst_cols, |
||||
__global float2* twiddles_ptr, const int t, const int nz) |
||||
{ |
||||
const int x = get_group_id(0); |
||||
const int y = get_global_id(1); |
||||
|
||||
if (x < nz) |
||||
{ |
||||
__local float2 smem[LOCAL_SIZE]; |
||||
__global const uchar* src = src_ptr + mad24(y, src_step, mad24(x, (int)(sizeof(float)*2), src_offset)); |
||||
__global const float2* twiddles = (__global float2*) twiddles_ptr; |
||||
const int ind = y; |
||||
const int block_size = LOCAL_SIZE/kercn; |
||||
float scale = 1.f/(dst_rows*dst_cols); |
||||
|
||||
#pragma unroll |
||||
for (int i=0; i<kercn; i++) |
||||
smem[y+i*block_size] = *((__global const float2*)(src + i*block_size*src_step)); |
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
|
||||
RADIX_PROCESS; |
||||
|
||||
#ifdef COMPLEX_OUTPUT |
||||
__global uchar* dst = dst_ptr + mad24(y, dst_step, mad24(x, (int)(sizeof(float)*2), dst_offset)); |
||||
#pragma unroll |
||||
for (int i=0; i<kercn; i++) |
||||
*((__global float2*)(dst + i*block_size*dst_step)) = SCALE_VAL(smem[y + i*block_size], scale); |
||||
#else |
||||
if (x == 0) |
||||
{ |
||||
// pack first column to CCS |
||||
__local float* smem_1cn = (__local float*) smem; |
||||
__global uchar* dst = dst_ptr + mad24(y+1, dst_step, dst_offset); |
||||
for (int i=y; i<dst_rows-1; i+=block_size, dst+=dst_step*block_size) |
||||
*((__global float*) dst) = SCALE_VAL(smem_1cn[i+2], scale); |
||||
if (y == 0) |
||||
*((__global float*) (dst_ptr + dst_offset)) = SCALE_VAL(smem_1cn[0], scale); |
||||
} |
||||
else if (x == (dst_cols+1)/2) |
||||
{ |
||||
// pack last column to CCS (if needed) |
||||
__local float* smem_1cn = (__local float*) smem; |
||||
__global uchar* dst = dst_ptr + mad24(dst_cols-1, (int)sizeof(float), mad24(y+1, dst_step, dst_offset)); |
||||
for (int i=y; i<dst_rows-1; i+=block_size, dst+=dst_step*block_size) |
||||
*((__global float*) dst) = SCALE_VAL(smem_1cn[i+2], scale); |
||||
if (y == 0) |
||||
*((__global float*) (dst_ptr + mad24(dst_cols-1, (int)sizeof(float), dst_offset))) = SCALE_VAL(smem_1cn[0], scale); |
||||
} |
||||
else |
||||
{ |
||||
__global uchar* dst = dst_ptr + mad24(x, (int)sizeof(float)*2, mad24(y, dst_step, dst_offset - (int)sizeof(float))); |
||||
#pragma unroll |
||||
for (int i=y; i<dst_rows; i+=block_size, dst+=block_size*dst_step) |
||||
vstore2(SCALE_VAL(smem[i], scale), 0, (__global float*) dst); |
||||
} |
||||
#endif |
||||
} |
||||
} |
||||
|
||||
__kernel void ifft_multi_radix_rows(__global const uchar* src_ptr, int src_step, int src_offset, int src_rows, int src_cols, |
||||
__global uchar* dst_ptr, int dst_step, int dst_offset, int dst_rows, int dst_cols, |
||||
__global float2* twiddles_ptr, const int t, const int nz) |
||||
{ |
||||
const int x = get_global_id(0); |
||||
const int y = get_group_id(1); |
||||
const int block_size = LOCAL_SIZE/kercn; |
||||
#ifdef IS_1D |
||||
const float scale = 1.f/dst_cols; |
||||
#else |
||||
const float scale = 1.f/(dst_cols*dst_rows); |
||||
#endif |
||||
|
||||
if (y < nz) |
||||
{ |
||||
__local float2 smem[LOCAL_SIZE]; |
||||
__global const float2* twiddles = (__global float2*) twiddles_ptr; |
||||
const int ind = x; |
||||
|
||||
#if defined(COMPLEX_INPUT) && !defined(NO_CONJUGATE) |
||||
__global const float2* src = (__global const float2*)(src_ptr + mad24(y, src_step, mad24(x, (int)(sizeof(float)*2), src_offset))); |
||||
#pragma unroll |
||||
for (int i=0; i<kercn; i++) |
||||
{ |
||||
smem[x+i*block_size].x = src[i*block_size].x; |
||||
smem[x+i*block_size].y = -src[i*block_size].y; |
||||
} |
||||
#else |
||||
|
||||
#if !defined(REAL_INPUT) && defined(NO_CONJUGATE) |
||||
__global const float2* src = (__global const float2*)(src_ptr + mad24(y, src_step, mad24(2, (int)sizeof(float), src_offset))); |
||||
|
||||
#pragma unroll |
||||
for (int i=x; i<(LOCAL_SIZE-1)/2; i+=block_size) |
||||
{ |
||||
smem[i+1].x = src[i].x; |
||||
smem[i+1].y = -src[i].y; |
||||
smem[LOCAL_SIZE-i-1] = src[i]; |
||||
} |
||||
#else |
||||
|
||||
#pragma unroll |
||||
for (int i=x; i<(LOCAL_SIZE-1)/2; i+=block_size) |
||||
{ |
||||
float2 src = vload2(0, (__global const float*)(src_ptr + mad24(y, src_step, mad24(2*i+1, (int)sizeof(float), src_offset)))); |
||||
|
||||
smem[i+1].x = src.x; |
||||
smem[i+1].y = -src.y; |
||||
smem[LOCAL_SIZE-i-1] = src; |
||||
} |
||||
|
||||
#endif |
||||
|
||||
if (x==0) |
||||
{ |
||||
smem[0].x = *(__global const float*)(src_ptr + mad24(y, src_step, src_offset)); |
||||
smem[0].y = 0.f; |
||||
|
||||
if(LOCAL_SIZE % 2 ==0) |
||||
{ |
||||
#if !defined(REAL_INPUT) && defined(NO_CONJUGATE) |
||||
smem[LOCAL_SIZE/2].x = src[LOCAL_SIZE/2-1].x; |
||||
#else |
||||
smem[LOCAL_SIZE/2].x = *(__global const float*)(src_ptr + mad24(y, src_step, mad24(LOCAL_SIZE-1, (int)sizeof(float), src_offset))); |
||||
#endif |
||||
smem[LOCAL_SIZE/2].y = 0.f; |
||||
} |
||||
} |
||||
#endif |
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
|
||||
RADIX_PROCESS; |
||||
|
||||
// copy data to dst |
||||
#ifdef COMPLEX_OUTPUT |
||||
__global float2* dst = (__global float*)(dst_ptr + mad24(y, dst_step, mad24(x, (int)(sizeof(float)*2), dst_offset))); |
||||
#pragma unroll |
||||
for (int i=0; i<kercn; i++) |
||||
{ |
||||
dst[i*block_size].x = SCALE_VAL(smem[x + i*block_size].x, scale); |
||||
dst[i*block_size].y = SCALE_VAL(-smem[x + i*block_size].y, scale); |
||||
} |
||||
#else |
||||
__global float* dst = (__global float*)(dst_ptr + mad24(y, dst_step, mad24(x, (int)(sizeof(float)), dst_offset))); |
||||
#pragma unroll |
||||
for (int i=0; i<kercn; i++) |
||||
{ |
||||
dst[i*block_size] = SCALE_VAL(smem[x + i*block_size].x, scale); |
||||
} |
||||
#endif |
||||
} |
||||
else |
||||
{ |
||||
// fill with zero other rows |
||||
#ifdef COMPLEX_OUTPUT |
||||
__global float2* dst = (__global float2*)(dst_ptr + mad24(y, dst_step, dst_offset)); |
||||
#else |
||||
__global float* dst = (__global float*)(dst_ptr + mad24(y, dst_step, dst_offset)); |
||||
#endif |
||||
#pragma unroll |
||||
for (int i=x; i<dst_cols; i+=block_size) |
||||
dst[i] = 0.f; |
||||
} |
||||
} |
||||
|
||||
__kernel void ifft_multi_radix_cols(__global const uchar* src_ptr, int src_step, int src_offset, int src_rows, int src_cols, |
||||
__global uchar* dst_ptr, int dst_step, int dst_offset, int dst_rows, int dst_cols, |
||||
__global float2* twiddles_ptr, const int t, const int nz) |
||||
{ |
||||
const int x = get_group_id(0); |
||||
const int y = get_global_id(1); |
||||
|
||||
#ifdef COMPLEX_INPUT |
||||
if (x < nz) |
||||
{ |
||||
__local float2 smem[LOCAL_SIZE]; |
||||
__global const uchar* src = src_ptr + mad24(y, src_step, mad24(x, (int)(sizeof(float)*2), src_offset)); |
||||
__global uchar* dst = dst_ptr + mad24(y, dst_step, mad24(x, (int)(sizeof(float)*2), dst_offset)); |
||||
__global const float2* twiddles = (__global float2*) twiddles_ptr; |
||||
const int ind = y; |
||||
const int block_size = LOCAL_SIZE/kercn; |
||||
|
||||
#pragma unroll |
||||
for (int i=0; i<kercn; i++) |
||||
{ |
||||
float2 temp = *((__global const float2*)(src + i*block_size*src_step)); |
||||
smem[y+i*block_size].x = temp.x; |
||||
smem[y+i*block_size].y = -temp.y; |
||||
} |
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
|
||||
RADIX_PROCESS; |
||||
|
||||
// copy data to dst |
||||
#pragma unroll |
||||
for (int i=0; i<kercn; i++) |
||||
{ |
||||
__global float2* res = (__global float2*)(dst + i*block_size*dst_step); |
||||
res[0].x = smem[y + i*block_size].x; |
||||
res[0].y = -smem[y + i*block_size].y; |
||||
} |
||||
} |
||||
#else |
||||
if (x < nz) |
||||
{ |
||||
__global const float2* twiddles = (__global float2*) twiddles_ptr; |
||||
const int ind = y; |
||||
const int block_size = LOCAL_SIZE/kercn; |
||||
|
||||
__local float2 smem[LOCAL_SIZE]; |
||||
#ifdef EVEN |
||||
if (x!=0 && (x!=(nz-1))) |
||||
#else |
||||
if (x!=0) |
||||
#endif |
||||
{ |
||||
__global const uchar* src = src_ptr + mad24(y, src_step, mad24(2*x-1, (int)sizeof(float), src_offset)); |
||||
#pragma unroll |
||||
for (int i=0; i<kercn; i++) |
||||
{ |
||||
float2 temp = vload2(0, (__global const float*)(src + i*block_size*src_step)); |
||||
smem[y+i*block_size].x = temp.x; |
||||
smem[y+i*block_size].y = -temp.y; |
||||
} |
||||
} |
||||
else |
||||
{ |
||||
int ind = x==0 ? 0: 2*x-1; |
||||
__global const float* src = (__global const float*)(src_ptr + mad24(1, src_step, mad24(ind, (int)sizeof(float), src_offset))); |
||||
int step = src_step/(int)sizeof(float); |
||||
|
||||
#pragma unroll |
||||
for (int i=y; i<(LOCAL_SIZE-1)/2; i+=block_size) |
||||
{ |
||||
smem[i+1].x = src[2*i*step]; |
||||
smem[i+1].y = -src[(2*i+1)*step]; |
||||
|
||||
smem[LOCAL_SIZE-i-1].x = src[2*i*step];; |
||||
smem[LOCAL_SIZE-i-1].y = src[(2*i+1)*step]; |
||||
} |
||||
if (y==0) |
||||
{ |
||||
smem[0].x = *(__global const float*)(src_ptr + mad24(ind, (int)sizeof(float), src_offset)); |
||||
smem[0].y = 0.f; |
||||
|
||||
if(LOCAL_SIZE % 2 ==0) |
||||
{ |
||||
smem[LOCAL_SIZE/2].x = src[(LOCAL_SIZE-2)*step]; |
||||
smem[LOCAL_SIZE/2].y = 0.f; |
||||
} |
||||
} |
||||
} |
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
|
||||
RADIX_PROCESS; |
||||
|
||||
// copy data to dst |
||||
__global uchar* dst = dst_ptr + mad24(y, dst_step, mad24(x, (int)(sizeof(float2)), dst_offset)); |
||||
|
||||
#pragma unroll |
||||
for (int i=0; i<kercn; i++) |
||||
{ |
||||
__global float2* res = (__global float2*)(dst + i*block_size*dst_step); |
||||
res[0].x = smem[y + i*block_size].x; |
||||
res[0].y = -smem[y + i*block_size].y; |
||||
} |
||||
} |
||||
#endif |
||||
} |
||||
@ -0,0 +1,8 @@ |
||||
namespace cv { namespace cuda { namespace device |
||||
{ |
||||
namespace disp_bilateral_filter |
||||
{ |
||||
template<typename T> |
||||
void disp_bilateral_filter(PtrStepSz<T> disp, PtrStepSzb img, int channels, int iters, const float *, const float *, size_t, int radius, short edge_disc, short max_disc, cudaStream_t stream); |
||||
} |
||||
}}} |
||||
@ -0,0 +1,29 @@ |
||||
namespace cv { namespace cuda { namespace device |
||||
{ |
||||
namespace stereocsbp |
||||
{ |
||||
template<class T> |
||||
void init_data_cost(const uchar *left, const uchar *right, uchar *ctemp, size_t cimg_step, int rows, int cols, T* disp_selected_pyr, T* data_cost_selected, size_t msg_step, |
||||
int h, int w, int level, int nr_plane, int ndisp, int channels, float data_weight, float max_data_term, int min_disp, bool use_local_init_data_cost, cudaStream_t stream); |
||||
|
||||
template<class T> |
||||
void compute_data_cost(const uchar *left, const uchar *right, size_t cimg_step, const T* disp_selected_pyr, T* data_cost, size_t msg_step, |
||||
int rows, int cols, int h, int w, int h2, int level, int nr_plane, int channels, float data_weight, float max_data_term, |
||||
int min_disp, cudaStream_t stream); |
||||
|
||||
template<class T> |
||||
void init_message(uchar *ctemp, T* u_new, T* d_new, T* l_new, T* r_new, |
||||
const T* u_cur, const T* d_cur, const T* l_cur, const T* r_cur, |
||||
T* selected_disp_pyr_new, const T* selected_disp_pyr_cur, |
||||
T* data_cost_selected, const T* data_cost, size_t msg_step, |
||||
int h, int w, int nr_plane, int h2, int w2, int nr_plane2, cudaStream_t stream); |
||||
|
||||
template<class T> |
||||
void calc_all_iterations(uchar *ctemp, T* u, T* d, T* l, T* r, const T* data_cost_selected, |
||||
const T* selected_disp_pyr_cur, size_t msg_step, int h, int w, int nr_plane, int iters, int max_disc_term, float disc_single_jump, cudaStream_t stream); |
||||
|
||||
template<class T> |
||||
void compute_disp(const T* u, const T* d, const T* l, const T* r, const T* data_cost_selected, const T* disp_selected, size_t msg_step, |
||||
const PtrStepSz<short>& disp, int nr_plane, cudaStream_t stream); |
||||
} |
||||
}}} |
||||
File diff suppressed because it is too large
Load Diff
@ -1,65 +0,0 @@ |
||||
/**
|
||||
* @file AKAZE.h |
||||
* @brief Main class for detecting and computing binary descriptors in an |
||||
* accelerated nonlinear scale space |
||||
* @date Mar 27, 2013 |
||||
* @author Pablo F. Alcantarilla, Jesus Nuevo |
||||
*/ |
||||
|
||||
#pragma once |
||||
|
||||
/* ************************************************************************* */ |
||||
// Includes
|
||||
#include "precomp.hpp" |
||||
#include "AKAZEConfig.h" |
||||
|
||||
/* ************************************************************************* */ |
||||
// AKAZE Class Declaration
|
||||
class AKAZEFeatures { |
||||
|
||||
private: |
||||
|
||||
AKAZEOptions options_; ///< Configuration options for AKAZE
|
||||
std::vector<TEvolution> evolution_; ///< Vector of nonlinear diffusion evolution
|
||||
|
||||
/// FED parameters
|
||||
int ncycles_; ///< Number of cycles
|
||||
bool reordering_; ///< Flag for reordering time steps
|
||||
std::vector<std::vector<float > > tsteps_; ///< Vector of FED dynamic time steps
|
||||
std::vector<int> nsteps_; ///< Vector of number of steps per cycle
|
||||
|
||||
/// Matrices for the M-LDB descriptor computation
|
||||
cv::Mat descriptorSamples_; // List of positions in the grids to sample LDB bits from.
|
||||
cv::Mat descriptorBits_; |
||||
cv::Mat bitMask_; |
||||
|
||||
public: |
||||
|
||||
/// Constructor with input arguments
|
||||
AKAZEFeatures(const AKAZEOptions& options); |
||||
|
||||
/// Scale Space methods
|
||||
void Allocate_Memory_Evolution(); |
||||
int Create_Nonlinear_Scale_Space(const cv::Mat& img); |
||||
void Feature_Detection(std::vector<cv::KeyPoint>& kpts); |
||||
void Compute_Determinant_Hessian_Response(void); |
||||
void Compute_Multiscale_Derivatives(void); |
||||
void Find_Scale_Space_Extrema(std::vector<cv::KeyPoint>& kpts); |
||||
void Do_Subpixel_Refinement(std::vector<cv::KeyPoint>& kpts); |
||||
|
||||
// Feature description methods
|
||||
void Compute_Descriptors(std::vector<cv::KeyPoint>& kpts, cv::Mat& desc); |
||||
|
||||
static void Compute_Main_Orientation(cv::KeyPoint& kpt, const std::vector<TEvolution>& evolution_); |
||||
}; |
||||
|
||||
/* ************************************************************************* */ |
||||
// Inline functions
|
||||
|
||||
// Inline functions
|
||||
void generateDescriptorSubsample(cv::Mat& sampleList, cv::Mat& comparisons, |
||||
int nbits, int pattern_size, int nchannels); |
||||
float get_angle(float x, float y); |
||||
float gaussian(float x, float y, float sigma); |
||||
void check_descriptor_limits(int& x, int& y, int width, int height); |
||||
int fRound(float flt); |
||||
File diff suppressed because it is too large
Load Diff
@ -0,0 +1,62 @@ |
||||
/**
|
||||
* @file AKAZE.h |
||||
* @brief Main class for detecting and computing binary descriptors in an |
||||
* accelerated nonlinear scale space |
||||
* @date Mar 27, 2013 |
||||
* @author Pablo F. Alcantarilla, Jesus Nuevo |
||||
*/ |
||||
|
||||
#ifndef __OPENCV_FEATURES_2D_AKAZE_FEATURES_H__ |
||||
#define __OPENCV_FEATURES_2D_AKAZE_FEATURES_H__ |
||||
|
||||
/* ************************************************************************* */ |
||||
// Includes
|
||||
#include "precomp.hpp" |
||||
#include "AKAZEConfig.h" |
||||
#include "TEvolution.h" |
||||
|
||||
/* ************************************************************************* */ |
||||
// AKAZE Class Declaration
|
||||
class AKAZEFeatures { |
||||
|
||||
private: |
||||
|
||||
AKAZEOptions options_; ///< Configuration options for AKAZE
|
||||
std::vector<TEvolution> evolution_; ///< Vector of nonlinear diffusion evolution
|
||||
|
||||
/// FED parameters
|
||||
int ncycles_; ///< Number of cycles
|
||||
bool reordering_; ///< Flag for reordering time steps
|
||||
std::vector<std::vector<float > > tsteps_; ///< Vector of FED dynamic time steps
|
||||
std::vector<int> nsteps_; ///< Vector of number of steps per cycle
|
||||
|
||||
/// Matrices for the M-LDB descriptor computation
|
||||
cv::Mat descriptorSamples_; // List of positions in the grids to sample LDB bits from.
|
||||
cv::Mat descriptorBits_; |
||||
cv::Mat bitMask_; |
||||
|
||||
public: |
||||
|
||||
/// Constructor with input arguments
|
||||
AKAZEFeatures(const AKAZEOptions& options); |
||||
|
||||
/// Scale Space methods
|
||||
void Allocate_Memory_Evolution(); |
||||
int Create_Nonlinear_Scale_Space(const cv::Mat& img); |
||||
void Feature_Detection(std::vector<cv::KeyPoint>& kpts); |
||||
void Compute_Determinant_Hessian_Response(void); |
||||
void Compute_Multiscale_Derivatives(void); |
||||
void Find_Scale_Space_Extrema(std::vector<cv::KeyPoint>& kpts); |
||||
void Do_Subpixel_Refinement(std::vector<cv::KeyPoint>& kpts); |
||||
|
||||
/// Feature description methods
|
||||
void Compute_Descriptors(std::vector<cv::KeyPoint>& kpts, cv::Mat& desc); |
||||
static void Compute_Main_Orientation(cv::KeyPoint& kpt, const std::vector<TEvolution>& evolution_); |
||||
}; |
||||
|
||||
/* ************************************************************************* */ |
||||
/// Inline functions
|
||||
void generateDescriptorSubsample(cv::Mat& sampleList, cv::Mat& comparisons, |
||||
int nbits, int pattern_size, int nchannels); |
||||
|
||||
#endif |
||||
File diff suppressed because it is too large
Load Diff
@ -0,0 +1,35 @@ |
||||
/**
|
||||
* @file TEvolution.h |
||||
* @brief Header file with the declaration of the TEvolution struct |
||||
* @date Jun 02, 2014 |
||||
* @author Pablo F. Alcantarilla |
||||
*/ |
||||
|
||||
#ifndef __OPENCV_FEATURES_2D_TEVOLUTION_H__ |
||||
#define __OPENCV_FEATURES_2D_TEVOLUTION_H__ |
||||
|
||||
/* ************************************************************************* */ |
||||
/// KAZE/A-KAZE nonlinear diffusion filtering evolution
|
||||
struct TEvolution { |
||||
|
||||
TEvolution() { |
||||
etime = 0.0f; |
||||
esigma = 0.0f; |
||||
octave = 0; |
||||
sublevel = 0; |
||||
sigma_size = 0; |
||||
} |
||||
|
||||
cv::Mat Lx, Ly; ///< First order spatial derivatives
|
||||
cv::Mat Lxx, Lxy, Lyy; ///< Second order spatial derivatives
|
||||
cv::Mat Lt; ///< Evolution image
|
||||
cv::Mat Lsmooth; ///< Smoothed image
|
||||
cv::Mat Ldet; ///< Detector response
|
||||
float etime; ///< Evolution time
|
||||
float esigma; ///< Evolution sigma. For linear diffusion t = sigma^2 / 2
|
||||
int octave; ///< Image octave
|
||||
int sublevel; ///< Image sublevel in each octave
|
||||
int sigma_size; ///< Integer esigma. For computing the feature detector responses
|
||||
}; |
||||
|
||||
#endif |
||||
@ -0,0 +1,77 @@ |
||||
#ifndef __OPENCV_FEATURES_2D_KAZE_UTILS_H__ |
||||
#define __OPENCV_FEATURES_2D_KAZE_UTILS_H__ |
||||
|
||||
/* ************************************************************************* */ |
||||
/**
|
||||
* @brief This function computes the angle from the vector given by (X Y). From 0 to 2*Pi |
||||
*/ |
||||
inline float getAngle(float x, float y) { |
||||
|
||||
if (x >= 0 && y >= 0) { |
||||
return atanf(y / x); |
||||
} |
||||
|
||||
if (x < 0 && y >= 0) { |
||||
return static_cast<float>(CV_PI)-atanf(-y / x); |
||||
} |
||||
|
||||
if (x < 0 && y < 0) { |
||||
return static_cast<float>(CV_PI)+atanf(y / x); |
||||
} |
||||
|
||||
if (x >= 0 && y < 0) { |
||||
return static_cast<float>(2.0 * CV_PI) - atanf(-y / x); |
||||
} |
||||
|
||||
return 0; |
||||
} |
||||
|
||||
/* ************************************************************************* */ |
||||
/**
|
||||
* @brief This function computes the value of a 2D Gaussian function |
||||
* @param x X Position |
||||
* @param y Y Position |
||||
* @param sig Standard Deviation |
||||
*/ |
||||
inline float gaussian(float x, float y, float sigma) { |
||||
return expf(-(x*x + y*y) / (2.0f*sigma*sigma)); |
||||
} |
||||
|
||||
/* ************************************************************************* */ |
||||
/**
|
||||
* @brief This function checks descriptor limits |
||||
* @param x X Position |
||||
* @param y Y Position |
||||
* @param width Image width |
||||
* @param height Image height |
||||
*/ |
||||
inline void checkDescriptorLimits(int &x, int &y, int width, int height) { |
||||
|
||||
if (x < 0) { |
||||
x = 0; |
||||
} |
||||
|
||||
if (y < 0) { |
||||
y = 0; |
||||
} |
||||
|
||||
if (x > width - 1) { |
||||
x = width - 1; |
||||
} |
||||
|
||||
if (y > height - 1) { |
||||
y = height - 1; |
||||
} |
||||
} |
||||
|
||||
/* ************************************************************************* */ |
||||
/**
|
||||
* @brief This funtion rounds float to nearest integer |
||||
* @param flt Input float |
||||
* @return dst Nearest integer |
||||
*/ |
||||
inline int fRound(float flt) { |
||||
return (int)(flt + 0.5f); |
||||
} |
||||
|
||||
#endif |
||||
Some files were not shown because too many files have changed in this diff Show More
Loading…
Reference in new issue