mirror of https://github.com/opencv/opencv.git
Merge pull request #631 from bitwangyaoyao:2.4_stereo
Andrey Kamaev
12 years ago
committed by
OpenCV Buildbot
commit
18ca645fd2
7 changed files with 818 additions and 0 deletions
@ -0,0 +1,427 @@ |
|||||||
|
/*M/////////////////////////////////////////////////////////////////////////////////////// |
||||||
|
// |
||||||
|
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. |
||||||
|
// |
||||||
|
// By downloading, copying, installing or using the software you agree to this license. |
||||||
|
// If you do not agree to this license, do not download, install, |
||||||
|
// copy or use the software. |
||||||
|
// |
||||||
|
// |
||||||
|
// License Agreement |
||||||
|
// For Open Source Computer Vision Library |
||||||
|
// |
||||||
|
// Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved. |
||||||
|
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. |
||||||
|
// Third party copyrights are property of their respective owners. |
||||||
|
// |
||||||
|
// @Authors |
||||||
|
// Jia Haipeng, jiahaipeng95@gmail.com |
||||||
|
// |
||||||
|
// Redistribution and use in source and binary forms, with or without modification, |
||||||
|
// are permitted provided that the following conditions are met: |
||||||
|
// |
||||||
|
// * Redistribution's of source code must retain the above copyright notice, |
||||||
|
// this list of conditions and the following disclaimer. |
||||||
|
// |
||||||
|
// * Redistribution's in binary form must reproduce the above copyright notice, |
||||||
|
// this list of conditions and the following disclaimer in the documentation |
||||||
|
// and/or other 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 ROWSperTHREAD 21 // the number of rows a thread will process |
||||||
|
#define BLOCK_W 128 // the thread block width (464) |
||||||
|
#define N_DISPARITIES 8 |
||||||
|
|
||||||
|
#define STEREO_MIND 0 // The minimum d range to check |
||||||
|
#define STEREO_DISP_STEP N_DISPARITIES // the d step, must be <= 1 to avoid aliasing |
||||||
|
|
||||||
|
int SQ(int a) |
||||||
|
{ |
||||||
|
return a * a; |
||||||
|
} |
||||||
|
|
||||||
|
unsigned int CalcSSD(volatile __local unsigned int *col_ssd_cache, |
||||||
|
volatile __local unsigned int *col_ssd, int radius) |
||||||
|
{ |
||||||
|
unsigned int cache = 0; |
||||||
|
unsigned int cache2 = 0; |
||||||
|
|
||||||
|
for(int i = 1; i <= radius; i++) |
||||||
|
cache += col_ssd[i]; |
||||||
|
|
||||||
|
col_ssd_cache[0] = cache; |
||||||
|
|
||||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||||
|
|
||||||
|
if (get_local_id(0) < BLOCK_W - radius) |
||||||
|
cache2 = col_ssd_cache[radius]; |
||||||
|
else |
||||||
|
for(int i = radius + 1; i < (2 * radius + 1); i++) |
||||||
|
cache2 += col_ssd[i]; |
||||||
|
|
||||||
|
return col_ssd[0] + cache + cache2; |
||||||
|
} |
||||||
|
|
||||||
|
uint2 MinSSD(volatile __local unsigned int *col_ssd_cache, |
||||||
|
volatile __local unsigned int *col_ssd, int radius) |
||||||
|
{ |
||||||
|
unsigned int ssd[N_DISPARITIES]; |
||||||
|
|
||||||
|
//See above: #define COL_SSD_SIZE (BLOCK_W + 2 * radius) |
||||||
|
ssd[0] = CalcSSD(col_ssd_cache, col_ssd + 0 * (BLOCK_W + 2 * radius), radius); |
||||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||||
|
ssd[1] = CalcSSD(col_ssd_cache, col_ssd + 1 * (BLOCK_W + 2 * radius), radius); |
||||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||||
|
ssd[2] = CalcSSD(col_ssd_cache, col_ssd + 2 * (BLOCK_W + 2 * radius), radius); |
||||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||||
|
ssd[3] = CalcSSD(col_ssd_cache, col_ssd + 3 * (BLOCK_W + 2 * radius), radius); |
||||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||||
|
ssd[4] = CalcSSD(col_ssd_cache, col_ssd + 4 * (BLOCK_W + 2 * radius), radius); |
||||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||||
|
ssd[5] = CalcSSD(col_ssd_cache, col_ssd + 5 * (BLOCK_W + 2 * radius), radius); |
||||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||||
|
ssd[6] = CalcSSD(col_ssd_cache, col_ssd + 6 * (BLOCK_W + 2 * radius), radius); |
||||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||||
|
ssd[7] = CalcSSD(col_ssd_cache, col_ssd + 7 * (BLOCK_W + 2 * radius), radius); |
||||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||||
|
|
||||||
|
unsigned int mssd = min(min(min(ssd[0], ssd[1]), min(ssd[4], ssd[5])), min(min(ssd[2], ssd[3]), min(ssd[6], ssd[7]))); |
||||||
|
|
||||||
|
int bestIdx = 0; |
||||||
|
for (int i = 0; i < N_DISPARITIES; i++) |
||||||
|
{ |
||||||
|
if (mssd == ssd[i]) |
||||||
|
bestIdx = i; |
||||||
|
} |
||||||
|
|
||||||
|
return (uint2)(mssd, bestIdx); |
||||||
|
} |
||||||
|
|
||||||
|
void StepDown(int idx1, int idx2, __global unsigned char* imageL, |
||||||
|
__global unsigned char* imageR, int d, volatile __local unsigned int *col_ssd, int radius) |
||||||
|
{ |
||||||
|
unsigned char leftPixel1; |
||||||
|
unsigned char leftPixel2; |
||||||
|
unsigned char rightPixel1[8]; |
||||||
|
unsigned char rightPixel2[8]; |
||||||
|
unsigned int diff1, diff2; |
||||||
|
|
||||||
|
leftPixel1 = imageL[idx1]; |
||||||
|
leftPixel2 = imageL[idx2]; |
||||||
|
|
||||||
|
idx1 = idx1 - d; |
||||||
|
idx2 = idx2 - d; |
||||||
|
|
||||||
|
rightPixel1[7] = imageR[idx1 - 7]; |
||||||
|
rightPixel1[0] = imageR[idx1 - 0]; |
||||||
|
rightPixel1[1] = imageR[idx1 - 1]; |
||||||
|
rightPixel1[2] = imageR[idx1 - 2]; |
||||||
|
rightPixel1[3] = imageR[idx1 - 3]; |
||||||
|
rightPixel1[4] = imageR[idx1 - 4]; |
||||||
|
rightPixel1[5] = imageR[idx1 - 5]; |
||||||
|
rightPixel1[6] = imageR[idx1 - 6]; |
||||||
|
|
||||||
|
rightPixel2[7] = imageR[idx2 - 7]; |
||||||
|
rightPixel2[0] = imageR[idx2 - 0]; |
||||||
|
rightPixel2[1] = imageR[idx2 - 1]; |
||||||
|
rightPixel2[2] = imageR[idx2 - 2]; |
||||||
|
rightPixel2[3] = imageR[idx2 - 3]; |
||||||
|
rightPixel2[4] = imageR[idx2 - 4]; |
||||||
|
rightPixel2[5] = imageR[idx2 - 5]; |
||||||
|
rightPixel2[6] = imageR[idx2 - 6]; |
||||||
|
|
||||||
|
//See above: #define COL_SSD_SIZE (BLOCK_W + 2 * radius) |
||||||
|
diff1 = leftPixel1 - rightPixel1[0]; |
||||||
|
diff2 = leftPixel2 - rightPixel2[0]; |
||||||
|
col_ssd[0 * (BLOCK_W + 2 * radius)] += SQ(diff2) - SQ(diff1); |
||||||
|
|
||||||
|
diff1 = leftPixel1 - rightPixel1[1]; |
||||||
|
diff2 = leftPixel2 - rightPixel2[1]; |
||||||
|
col_ssd[1 * (BLOCK_W + 2 * radius)] += SQ(diff2) - SQ(diff1); |
||||||
|
|
||||||
|
diff1 = leftPixel1 - rightPixel1[2]; |
||||||
|
diff2 = leftPixel2 - rightPixel2[2]; |
||||||
|
col_ssd[2 * (BLOCK_W + 2 * radius)] += SQ(diff2) - SQ(diff1); |
||||||
|
|
||||||
|
diff1 = leftPixel1 - rightPixel1[3]; |
||||||
|
diff2 = leftPixel2 - rightPixel2[3]; |
||||||
|
col_ssd[3 * (BLOCK_W + 2 * radius)] += SQ(diff2) - SQ(diff1); |
||||||
|
|
||||||
|
diff1 = leftPixel1 - rightPixel1[4]; |
||||||
|
diff2 = leftPixel2 - rightPixel2[4]; |
||||||
|
col_ssd[4 * (BLOCK_W + 2 * radius)] += SQ(diff2) - SQ(diff1); |
||||||
|
|
||||||
|
diff1 = leftPixel1 - rightPixel1[5]; |
||||||
|
diff2 = leftPixel2 - rightPixel2[5]; |
||||||
|
col_ssd[5 * (BLOCK_W + 2 * radius)] += SQ(diff2) - SQ(diff1); |
||||||
|
|
||||||
|
diff1 = leftPixel1 - rightPixel1[6]; |
||||||
|
diff2 = leftPixel2 - rightPixel2[6]; |
||||||
|
col_ssd[6 * (BLOCK_W + 2 * radius)] += SQ(diff2) - SQ(diff1); |
||||||
|
|
||||||
|
diff1 = leftPixel1 - rightPixel1[7]; |
||||||
|
diff2 = leftPixel2 - rightPixel2[7]; |
||||||
|
col_ssd[7 * (BLOCK_W + 2 * radius)] += SQ(diff2) - SQ(diff1); |
||||||
|
} |
||||||
|
|
||||||
|
void InitColSSD(int x_tex, int y_tex, int im_pitch, __global unsigned char* imageL, |
||||||
|
__global unsigned char* imageR, int d, |
||||||
|
volatile __local unsigned int *col_ssd, int radius) |
||||||
|
{ |
||||||
|
unsigned char leftPixel1; |
||||||
|
int idx; |
||||||
|
unsigned int diffa[] = {0, 0, 0, 0, 0, 0, 0, 0}; |
||||||
|
|
||||||
|
for(int i = 0; i < (2 * radius + 1); i++) |
||||||
|
{ |
||||||
|
idx = y_tex * im_pitch + x_tex; |
||||||
|
leftPixel1 = imageL[idx]; |
||||||
|
idx = idx - d; |
||||||
|
|
||||||
|
diffa[0] += SQ(leftPixel1 - imageR[idx - 0]); |
||||||
|
diffa[1] += SQ(leftPixel1 - imageR[idx - 1]); |
||||||
|
diffa[2] += SQ(leftPixel1 - imageR[idx - 2]); |
||||||
|
diffa[3] += SQ(leftPixel1 - imageR[idx - 3]); |
||||||
|
diffa[4] += SQ(leftPixel1 - imageR[idx - 4]); |
||||||
|
diffa[5] += SQ(leftPixel1 - imageR[idx - 5]); |
||||||
|
diffa[6] += SQ(leftPixel1 - imageR[idx - 6]); |
||||||
|
diffa[7] += SQ(leftPixel1 - imageR[idx - 7]); |
||||||
|
|
||||||
|
y_tex += 1; |
||||||
|
} |
||||||
|
//See above: #define COL_SSD_SIZE (BLOCK_W + 2 * radius) |
||||||
|
col_ssd[0 * (BLOCK_W + 2 * radius)] = diffa[0]; |
||||||
|
col_ssd[1 * (BLOCK_W + 2 * radius)] = diffa[1]; |
||||||
|
col_ssd[2 * (BLOCK_W + 2 * radius)] = diffa[2]; |
||||||
|
col_ssd[3 * (BLOCK_W + 2 * radius)] = diffa[3]; |
||||||
|
col_ssd[4 * (BLOCK_W + 2 * radius)] = diffa[4]; |
||||||
|
col_ssd[5 * (BLOCK_W + 2 * radius)] = diffa[5]; |
||||||
|
col_ssd[6 * (BLOCK_W + 2 * radius)] = diffa[6]; |
||||||
|
col_ssd[7 * (BLOCK_W + 2 * radius)] = diffa[7]; |
||||||
|
} |
||||||
|
|
||||||
|
__kernel void stereoKernel(__global unsigned char *left, __global unsigned char *right, |
||||||
|
__global unsigned int *cminSSDImage, int cminSSD_step, |
||||||
|
__global unsigned char *disp, int disp_step,int cwidth, int cheight, |
||||||
|
int img_step, int maxdisp, int radius, |
||||||
|
__local unsigned int *col_ssd_cache) |
||||||
|
{ |
||||||
|
|
||||||
|
volatile __local unsigned int *col_ssd = col_ssd_cache + BLOCK_W + get_local_id(0); |
||||||
|
volatile __local unsigned int *col_ssd_extra = get_local_id(0) < (2 * radius) ? col_ssd + BLOCK_W : 0; |
||||||
|
|
||||||
|
int X = get_group_id(0) * BLOCK_W + get_local_id(0) + maxdisp + radius; |
||||||
|
// int Y = get_group_id(1) * ROWSperTHREAD + radius; |
||||||
|
|
||||||
|
#define Y (get_group_id(1) * ROWSperTHREAD + radius) |
||||||
|
|
||||||
|
volatile __global unsigned int* minSSDImage = cminSSDImage + X + Y * cminSSD_step; |
||||||
|
__global unsigned char* disparImage = disp + X + Y * disp_step; |
||||||
|
|
||||||
|
int end_row = ROWSperTHREAD < (cheight - Y) ? ROWSperTHREAD:(cheight - Y); |
||||||
|
int y_tex; |
||||||
|
int x_tex = X - radius; |
||||||
|
|
||||||
|
if (x_tex >= cwidth) |
||||||
|
return; |
||||||
|
|
||||||
|
for(int d = STEREO_MIND; d < maxdisp; d += STEREO_DISP_STEP) |
||||||
|
{ |
||||||
|
y_tex = Y - radius; |
||||||
|
|
||||||
|
InitColSSD(x_tex, y_tex, img_step, left, right, d, col_ssd, radius); |
||||||
|
if (col_ssd_extra > 0) |
||||||
|
if (x_tex + BLOCK_W < cwidth) |
||||||
|
InitColSSD(x_tex + BLOCK_W, y_tex, img_step, left, right, d, col_ssd_extra, radius); |
||||||
|
|
||||||
|
barrier(CLK_LOCAL_MEM_FENCE); //before MinSSD function |
||||||
|
|
||||||
|
if (X < cwidth - radius && Y < cheight - radius) |
||||||
|
{ |
||||||
|
uint2 minSSD = MinSSD(col_ssd_cache + get_local_id(0), col_ssd, radius); |
||||||
|
if (minSSD.x < minSSDImage[0]) |
||||||
|
{ |
||||||
|
disparImage[0] = (unsigned char)(d + minSSD.y); |
||||||
|
minSSDImage[0] = minSSD.x; |
||||||
|
} |
||||||
|
} |
||||||
|
|
||||||
|
for(int row = 1; row < end_row; row++) |
||||||
|
{ |
||||||
|
int idx1 = y_tex * img_step + x_tex; |
||||||
|
int idx2 = (y_tex + (2 * radius + 1)) * img_step + x_tex; |
||||||
|
|
||||||
|
barrier(CLK_GLOBAL_MEM_FENCE); |
||||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||||
|
|
||||||
|
StepDown(idx1, idx2, left, right, d, col_ssd, radius); |
||||||
|
if (col_ssd_extra > 0) |
||||||
|
if (x_tex + BLOCK_W < cwidth) |
||||||
|
StepDown(idx1, idx2, left + BLOCK_W, right + BLOCK_W, d, col_ssd_extra, radius); |
||||||
|
|
||||||
|
y_tex += 1; |
||||||
|
|
||||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||||
|
|
||||||
|
if (X < cwidth - radius && row < cheight - radius - Y) |
||||||
|
{ |
||||||
|
int idx = row * cminSSD_step; |
||||||
|
uint2 minSSD = MinSSD(col_ssd_cache + get_local_id(0), col_ssd, radius); |
||||||
|
if (minSSD.x < minSSDImage[idx]) |
||||||
|
{ |
||||||
|
disparImage[disp_step * row] = (unsigned char)(d + minSSD.y); |
||||||
|
minSSDImage[idx] = minSSD.x; |
||||||
|
} |
||||||
|
} |
||||||
|
} // for row loop |
||||||
|
} // for d loop |
||||||
|
} |
||||||
|
|
||||||
|
////////////////////////////////////////////////////////////////////////////////////////////////// |
||||||
|
//////////////////////////// Sobel Prefiler (signal channel)////////////////////////////////////// |
||||||
|
////////////////////////////////////////////////////////////////////////////////////////////////// |
||||||
|
|
||||||
|
__kernel void prefilter_xsobel(__global unsigned char *input, __global unsigned char *output, |
||||||
|
int rows, int cols, int prefilterCap) |
||||||
|
{ |
||||||
|
int x = get_global_id(0); |
||||||
|
int y = get_global_id(1); |
||||||
|
|
||||||
|
if(x < cols && y < rows) |
||||||
|
{ |
||||||
|
int cov = input[(y-1) * cols + (x-1)] * (-1) + input[(y-1) * cols + (x+1)] * (1) + |
||||||
|
input[(y) * cols + (x-1)] * (-2) + input[(y) * cols + (x+1)] * (2) + |
||||||
|
input[(y+1) * cols + (x-1)] * (-1) + input[(y+1) * cols + (x+1)] * (1); |
||||||
|
|
||||||
|
cov = min(min(max(-prefilterCap, cov), prefilterCap) + prefilterCap, 255); |
||||||
|
output[y * cols + x] = cov & 0xFF; |
||||||
|
} |
||||||
|
} |
||||||
|
|
||||||
|
|
||||||
|
////////////////////////////////////////////////////////////////////////////////////////////////// |
||||||
|
/////////////////////////////////// Textureness filtering //////////////////////////////////////// |
||||||
|
////////////////////////////////////////////////////////////////////////////////////////////////// |
||||||
|
|
||||||
|
float sobel(__global unsigned char *input, int x, int y, int rows, int cols) |
||||||
|
{ |
||||||
|
float conv = 0; |
||||||
|
int y1 = y==0? 0 : y-1; |
||||||
|
int x1 = x==0? 0 : x-1; |
||||||
|
if(x < cols && y < rows) |
||||||
|
{ |
||||||
|
conv = (float)input[(y1) * cols + (x1)] * (-1) + (float)input[(y1) * cols + (x+1)] * (1) + |
||||||
|
(float)input[(y) * cols + (x1)] * (-2) + (float)input[(y) * cols + (x+1)] * (2) + |
||||||
|
(float)input[(y+1) * cols + (x1)] * (-1) + (float)input[(y+1) * cols + (x+1)] * (1); |
||||||
|
|
||||||
|
} |
||||||
|
return fabs(conv); |
||||||
|
} |
||||||
|
|
||||||
|
float CalcSums(__local float *cols, __local float *cols_cache, int winsz) |
||||||
|
{ |
||||||
|
float cache = 0; |
||||||
|
float cache2 = 0; |
||||||
|
int winsz2 = winsz/2; |
||||||
|
|
||||||
|
int x = get_local_id(0); |
||||||
|
int group_size_x = get_local_size(0); |
||||||
|
|
||||||
|
for(int i = 1; i <= winsz2; i++) |
||||||
|
cache += cols[i]; |
||||||
|
|
||||||
|
cols_cache[0] = cache; |
||||||
|
|
||||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||||
|
|
||||||
|
if (x < group_size_x - winsz2) |
||||||
|
cache2 = cols_cache[winsz2]; |
||||||
|
else |
||||||
|
for(int i = winsz2 + 1; i < winsz; i++) |
||||||
|
cache2 += cols[i]; |
||||||
|
|
||||||
|
return cols[0] + cache + cache2; |
||||||
|
} |
||||||
|
|
||||||
|
#define RpT (2 * ROWSperTHREAD) // got experimentally |
||||||
|
__kernel void textureness_kernel(__global unsigned char *disp, int disp_rows, int disp_cols, |
||||||
|
int disp_step, __global unsigned char *input, int input_rows, |
||||||
|
int input_cols,int winsz, float threshold, |
||||||
|
__local float *cols_cache) |
||||||
|
{ |
||||||
|
int winsz2 = winsz/2; |
||||||
|
int n_dirty_pixels = (winsz2) * 2; |
||||||
|
|
||||||
|
int local_id_x = get_local_id(0); |
||||||
|
int group_size_x = get_local_size(0); |
||||||
|
int group_id_y = get_group_id(1); |
||||||
|
|
||||||
|
__local float *cols = cols_cache + group_size_x + local_id_x; |
||||||
|
__local float *cols_extra = local_id_x < n_dirty_pixels ? cols + group_size_x : 0; |
||||||
|
|
||||||
|
int x = get_global_id(0); |
||||||
|
int beg_row = group_id_y * RpT; |
||||||
|
int end_row = min(beg_row + RpT, disp_rows); |
||||||
|
|
||||||
|
// if (x < disp_cols) |
||||||
|
// { |
||||||
|
int y = beg_row; |
||||||
|
|
||||||
|
float sum = 0; |
||||||
|
float sum_extra = 0; |
||||||
|
|
||||||
|
for(int i = y - winsz2; i <= y + winsz2; ++i) |
||||||
|
{ |
||||||
|
sum += sobel(input, x - winsz2, i, input_rows, input_cols); |
||||||
|
if (cols_extra) |
||||||
|
sum_extra += sobel(input, x + group_size_x - winsz2, i, input_rows, input_cols); |
||||||
|
} |
||||||
|
*cols = sum; |
||||||
|
if (cols_extra) |
||||||
|
*cols_extra = sum_extra; |
||||||
|
|
||||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||||
|
|
||||||
|
float sum_win = CalcSums(cols, cols_cache + local_id_x, winsz) * 255; |
||||||
|
if (sum_win < threshold) |
||||||
|
disp[y * disp_step + x] = 0; |
||||||
|
|
||||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||||
|
|
||||||
|
for(int y = beg_row + 1; y < end_row; ++y) |
||||||
|
{ |
||||||
|
sum = sum - sobel(input, x - winsz2, y - winsz2 - 1, input_rows, input_cols) + |
||||||
|
sobel(input, x - winsz2, y + winsz2, input_rows, input_cols); |
||||||
|
*cols = sum; |
||||||
|
|
||||||
|
if (cols_extra) |
||||||
|
{ |
||||||
|
sum_extra = sum_extra - sobel(input, x + group_size_x - winsz2, y - winsz2 - 1,input_rows, input_cols) |
||||||
|
+ sobel(input, x + group_size_x - winsz2, y + winsz2, input_rows, input_cols); |
||||||
|
*cols_extra = sum_extra; |
||||||
|
} |
||||||
|
|
||||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||||
|
float sum_win = CalcSums(cols, cols_cache + local_id_x, winsz) * 255; |
||||||
|
if (sum_win < threshold) |
||||||
|
disp[y * disp_step + x] = 0; |
||||||
|
|
||||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||||
|
} |
||||||
|
// } |
||||||
|
} |
||||||
@ -0,0 +1,263 @@ |
|||||||
|
/*M///////////////////////////////////////////////////////////////////////////////////////
|
||||||
|
//
|
||||||
|
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
|
||||||
|
//
|
||||||
|
// By downloading, copying, installing or using the software you agree to this license.
|
||||||
|
// If you do not agree to this license, do not download, install,
|
||||||
|
// copy or use the software.
|
||||||
|
//
|
||||||
|
//
|
||||||
|
// License Agreement
|
||||||
|
// For Open Source Computer Vision Library
|
||||||
|
//
|
||||||
|
// Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved.
|
||||||
|
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
|
||||||
|
// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved.
|
||||||
|
// Third party copyrights are property of their respective owners.
|
||||||
|
//
|
||||||
|
// @Authors
|
||||||
|
// Jia Haipeng, jiahaipeng95@gmail.com
|
||||||
|
// Xiaopeng Fu, xiaopeng@multicorewareinc.com
|
||||||
|
//
|
||||||
|
// Redistribution and use in source and binary forms, with or without modification,
|
||||||
|
// are permitted provided that the following conditions are met:
|
||||||
|
//
|
||||||
|
// * Redistribution's of source code must retain the above copyright notice,
|
||||||
|
// this list of conditions and the following disclaimer.
|
||||||
|
//
|
||||||
|
// * Redistribution's in binary form must reproduce the above copyright notice,
|
||||||
|
// this list of conditions and the following disclaimer in the documentation
|
||||||
|
// and/or other 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 <vector> |
||||||
|
|
||||||
|
using namespace cv; |
||||||
|
using namespace cv::ocl; |
||||||
|
using namespace std; |
||||||
|
|
||||||
|
|
||||||
|
namespace cv |
||||||
|
{ |
||||||
|
namespace ocl |
||||||
|
{ |
||||||
|
|
||||||
|
///////////////////////////OpenCL kernel strings///////////////////////////
|
||||||
|
extern const char *stereobm; |
||||||
|
|
||||||
|
} |
||||||
|
} |
||||||
|
namespace cv |
||||||
|
{ |
||||||
|
namespace ocl |
||||||
|
{ |
||||||
|
namespace stereoBM |
||||||
|
{ |
||||||
|
/////////////////////////////////////////////////////////////////////////
|
||||||
|
//////////////////////////prefilter_xsbel////////////////////////////////
|
||||||
|
////////////////////////////////////////////////////////////////////////
|
||||||
|
static void prefilter_xsobel(const oclMat &input, oclMat &output, int prefilterCap) |
||||||
|
{ |
||||||
|
Context *clCxt = input.clCxt; |
||||||
|
|
||||||
|
string kernelName = "prefilter_xsobel"; |
||||||
|
cl_kernel kernel = openCLGetKernelFromSource(clCxt, &stereobm, kernelName); |
||||||
|
|
||||||
|
size_t blockSize = 1; |
||||||
|
size_t globalThreads[3] = { input.cols, input.rows, 1 }; |
||||||
|
size_t localThreads[3] = { blockSize, blockSize, 1 }; |
||||||
|
|
||||||
|
openCLVerifyKernel(clCxt, kernel, localThreads); |
||||||
|
openCLSafeCall(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&input.data)); |
||||||
|
openCLSafeCall(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&output.data)); |
||||||
|
openCLSafeCall(clSetKernelArg(kernel, 2, sizeof(cl_int), (void *)&input.rows)); |
||||||
|
openCLSafeCall(clSetKernelArg(kernel, 3, sizeof(cl_int), (void *)&input.cols)); |
||||||
|
openCLSafeCall(clSetKernelArg(kernel, 4, sizeof(cl_int), (void *)&prefilterCap)); |
||||||
|
|
||||||
|
openCLSafeCall(clEnqueueNDRangeKernel(clCxt->impl->clCmdQueue, kernel, 3, NULL, |
||||||
|
globalThreads, localThreads, 0, NULL, NULL)); |
||||||
|
|
||||||
|
clFinish(clCxt->impl->clCmdQueue); |
||||||
|
openCLSafeCall(clReleaseKernel(kernel)); |
||||||
|
|
||||||
|
} |
||||||
|
//////////////////////////////////////////////////////////////////////////
|
||||||
|
//////////////////////////////common////////////////////////////////////
|
||||||
|
////////////////////////////////////////////////////////////////////////
|
||||||
|
#define N_DISPARITIES 8 |
||||||
|
#define ROWSperTHREAD 21 |
||||||
|
#define BLOCK_W 128 |
||||||
|
static inline int divUp(int total, int grain) |
||||||
|
{ |
||||||
|
return (total + grain - 1) / grain; |
||||||
|
} |
||||||
|
////////////////////////////////////////////////////////////////////////////
|
||||||
|
///////////////////////////////stereoBM_GPU////////////////////////////////
|
||||||
|
////////////////////////////////////////////////////////////////////////////
|
||||||
|
static void stereo_bm(const oclMat &left, const oclMat &right, oclMat &disp, |
||||||
|
int maxdisp, int winSize, oclMat &minSSD_buf) |
||||||
|
{ |
||||||
|
int winsz2 = winSize >> 1; |
||||||
|
|
||||||
|
//if(winsz2 == 0 || winsz2 >= calles_num)
|
||||||
|
//cv::ocl:error("Unsupported window size", __FILE__, __LINE__, __FUNCTION__);
|
||||||
|
|
||||||
|
Context *clCxt = left.clCxt; |
||||||
|
|
||||||
|
string kernelName = "stereoKernel"; |
||||||
|
cl_kernel kernel = openCLGetKernelFromSource(clCxt, &stereobm, kernelName); |
||||||
|
|
||||||
|
disp.setTo(Scalar_<unsigned char>::all(0)); |
||||||
|
minSSD_buf.setTo(Scalar_<unsigned int>::all(0xFFFFFFFF)); |
||||||
|
|
||||||
|
size_t minssd_step = minSSD_buf.step / minSSD_buf.elemSize(); |
||||||
|
size_t local_mem_size = (BLOCK_W + N_DISPARITIES * (BLOCK_W + 2 * winsz2)) * |
||||||
|
sizeof(cl_uint); |
||||||
|
//size_t blockSize = 1;
|
||||||
|
size_t localThreads[] = { BLOCK_W, 1,1}; |
||||||
|
size_t globalThreads[] = { divUp(left.cols - maxdisp - 2 * winsz2, BLOCK_W) *BLOCK_W, |
||||||
|
divUp(left.rows - 2 * winsz2, ROWSperTHREAD), |
||||||
|
1 |
||||||
|
}; |
||||||
|
|
||||||
|
openCLVerifyKernel(clCxt, kernel, localThreads); |
||||||
|
openCLSafeCall(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&left.data)); |
||||||
|
openCLSafeCall(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&right.data)); |
||||||
|
openCLSafeCall(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *)&minSSD_buf.data)); |
||||||
|
openCLSafeCall(clSetKernelArg(kernel, 3, sizeof(cl_int), (void *)&minssd_step)); |
||||||
|
openCLSafeCall(clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *)&disp.data)); |
||||||
|
openCLSafeCall(clSetKernelArg(kernel, 5, sizeof(cl_int), (void *)&disp.step)); |
||||||
|
openCLSafeCall(clSetKernelArg(kernel, 6, sizeof(cl_int), (void *)&left.cols)); |
||||||
|
openCLSafeCall(clSetKernelArg(kernel, 7, sizeof(cl_int), (void *)&left.rows)); |
||||||
|
openCLSafeCall(clSetKernelArg(kernel, 8, sizeof(cl_int), (void *)&left.step)); |
||||||
|
openCLSafeCall(clSetKernelArg(kernel, 9, sizeof(cl_int), (void *)&maxdisp)); |
||||||
|
openCLSafeCall(clSetKernelArg(kernel, 10, sizeof(cl_int), (void *)&winsz2)); |
||||||
|
openCLSafeCall(clSetKernelArg(kernel, 11, local_mem_size, (void *)NULL)); |
||||||
|
|
||||||
|
openCLSafeCall(clEnqueueNDRangeKernel(clCxt->impl->clCmdQueue, kernel, 2, NULL, |
||||||
|
globalThreads, localThreads, 0, NULL, NULL)); |
||||||
|
|
||||||
|
|
||||||
|
clFinish(clCxt->impl->clCmdQueue); |
||||||
|
openCLSafeCall(clReleaseKernel(kernel)); |
||||||
|
} |
||||||
|
////////////////////////////////////////////////////////////////////////////
|
||||||
|
///////////////////////////////postfilter_textureness///////////////////////
|
||||||
|
////////////////////////////////////////////////////////////////////////////
|
||||||
|
static void postfilter_textureness(oclMat &left, int winSize, |
||||||
|
float avergeTexThreshold, oclMat &disparity) |
||||||
|
{ |
||||||
|
Context *clCxt = left.clCxt; |
||||||
|
|
||||||
|
string kernelName = "textureness_kernel"; |
||||||
|
cl_kernel kernel = openCLGetKernelFromSource(clCxt, &stereobm, kernelName); |
||||||
|
|
||||||
|
size_t blockSize = 1; |
||||||
|
size_t localThreads[] = { BLOCK_W, blockSize ,1}; |
||||||
|
size_t globalThreads[] = { divUp(left.cols, BLOCK_W) *BLOCK_W, |
||||||
|
divUp(left.rows, 2 * ROWSperTHREAD), |
||||||
|
1 |
||||||
|
}; |
||||||
|
|
||||||
|
size_t local_mem_size = (localThreads[0] + localThreads[0] + (winSize / 2) * 2) * sizeof(float); |
||||||
|
|
||||||
|
openCLVerifyKernel(clCxt, kernel, localThreads); |
||||||
|
openCLSafeCall(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&disparity.data)); |
||||||
|
openCLSafeCall(clSetKernelArg(kernel, 1, sizeof(cl_int), (void *)&disparity.rows)); |
||||||
|
openCLSafeCall(clSetKernelArg(kernel, 2, sizeof(cl_int), (void *)&disparity.cols)); |
||||||
|
openCLSafeCall(clSetKernelArg(kernel, 3, sizeof(cl_int), (void *)&disparity.step)); |
||||||
|
openCLSafeCall(clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *)&left.data)); |
||||||
|
openCLSafeCall(clSetKernelArg(kernel, 5, sizeof(cl_int), (void *)&left.rows)); |
||||||
|
openCLSafeCall(clSetKernelArg(kernel, 6, sizeof(cl_int), (void *)&left.cols)); |
||||||
|
openCLSafeCall(clSetKernelArg(kernel, 7, sizeof(cl_int), (void *)&winSize)); |
||||||
|
openCLSafeCall(clSetKernelArg(kernel, 8, sizeof(cl_float), (void *)&avergeTexThreshold)); |
||||||
|
openCLSafeCall(clSetKernelArg(kernel, 9, local_mem_size, NULL)); |
||||||
|
openCLSafeCall(clEnqueueNDRangeKernel(clCxt->impl->clCmdQueue, kernel, 2, NULL, |
||||||
|
globalThreads, localThreads, 0, NULL, NULL)); |
||||||
|
|
||||||
|
clFinish(clCxt->impl->clCmdQueue); |
||||||
|
openCLSafeCall(clReleaseKernel(kernel)); |
||||||
|
} |
||||||
|
//////////////////////////////////////////////////////////////////////////////
|
||||||
|
/////////////////////////////////////operator/////////////////////////////////
|
||||||
|
/////////////////////////////////////////////////////////////////////////////
|
||||||
|
static void operator_(oclMat &minSSD, oclMat &leBuf, oclMat &riBuf, int preset, int ndisp, |
||||||
|
int winSize, float avergeTexThreshold, const oclMat &left, |
||||||
|
const oclMat &right, oclMat &disparity) |
||||||
|
|
||||||
|
{ |
||||||
|
CV_DbgAssert(left.rows == right.rows && left.cols == right.cols); |
||||||
|
CV_DbgAssert(left.type() == CV_8UC1); |
||||||
|
CV_DbgAssert(right.type() == CV_8UC1); |
||||||
|
|
||||||
|
disparity.create(left.size(), CV_8UC1); |
||||||
|
minSSD.create(left.size(), CV_32SC1); |
||||||
|
|
||||||
|
oclMat le_for_bm = left; |
||||||
|
oclMat ri_for_bm = right; |
||||||
|
|
||||||
|
if (preset == cv::ocl::StereoBM_OCL::PREFILTER_XSOBEL) |
||||||
|
{ |
||||||
|
leBuf.create( left.size(), left.type()); |
||||||
|
riBuf.create(right.size(), right.type()); |
||||||
|
|
||||||
|
prefilter_xsobel( left, leBuf, 31); |
||||||
|
prefilter_xsobel(right, riBuf, 31); |
||||||
|
|
||||||
|
le_for_bm = leBuf; |
||||||
|
ri_for_bm = riBuf; |
||||||
|
} |
||||||
|
|
||||||
|
stereo_bm(le_for_bm, ri_for_bm, disparity, ndisp, winSize, minSSD); |
||||||
|
|
||||||
|
if (avergeTexThreshold) |
||||||
|
{ |
||||||
|
postfilter_textureness(le_for_bm, winSize, avergeTexThreshold, disparity); |
||||||
|
} |
||||||
|
} |
||||||
|
} |
||||||
|
} |
||||||
|
} |
||||||
|
const float defaultAvgTexThreshold = 3; |
||||||
|
|
||||||
|
cv::ocl::StereoBM_OCL::StereoBM_OCL() |
||||||
|
: preset(BASIC_PRESET), ndisp(DEFAULT_NDISP), winSize(DEFAULT_WINSZ), |
||||||
|
avergeTexThreshold(defaultAvgTexThreshold) {} |
||||||
|
|
||||||
|
cv::ocl::StereoBM_OCL::StereoBM_OCL(int preset_, int ndisparities_, int winSize_) |
||||||
|
: preset(preset_), ndisp(ndisparities_), winSize(winSize_), |
||||||
|
avergeTexThreshold(defaultAvgTexThreshold) |
||||||
|
{ |
||||||
|
const int max_supported_ndisp = 1 << (sizeof(unsigned char) * 8); |
||||||
|
CV_Assert(0 < ndisp && ndisp <= max_supported_ndisp); |
||||||
|
CV_Assert(ndisp % 8 == 0); |
||||||
|
CV_Assert(winSize % 2 == 1); |
||||||
|
} |
||||||
|
|
||||||
|
bool cv::ocl::StereoBM_OCL::checkIfGpuCallReasonable() |
||||||
|
{ |
||||||
|
return true; |
||||||
|
} |
||||||
|
|
||||||
|
void cv::ocl::StereoBM_OCL::operator() ( const oclMat &left, const oclMat &right, |
||||||
|
oclMat &disparity) |
||||||
|
{ |
||||||
|
cv::ocl::stereoBM::operator_(minSSD, leBuf, riBuf, preset, ndisp, winSize, avergeTexThreshold, left, right, disparity); |
||||||
|
} |
||||||
|
|
||||||
@ -0,0 +1,94 @@ |
|||||||
|
///////////////////////////////////////////////////////////////////////////////////////
|
||||||
|
//
|
||||||
|
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
|
||||||
|
//
|
||||||
|
// By downloading, copying, installing or using the software you agree to this license.
|
||||||
|
// If you do not agree to this license, do not download, install,
|
||||||
|
// copy or use the software.
|
||||||
|
//
|
||||||
|
//
|
||||||
|
// License Agreement
|
||||||
|
// For Open Source Computer Vision Library
|
||||||
|
//
|
||||||
|
// Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved.
|
||||||
|
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
|
||||||
|
// Third party copyrights are property of their respective owners.
|
||||||
|
//
|
||||||
|
// @Authors
|
||||||
|
|
||||||
|
//
|
||||||
|
// 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 <iomanip> |
||||||
|
|
||||||
|
#ifdef HAVE_OPENCL |
||||||
|
|
||||||
|
using namespace cv; |
||||||
|
|
||||||
|
extern std::string workdir; |
||||||
|
PARAM_TEST_CASE(StereoMatchBM, int, int) |
||||||
|
{ |
||||||
|
int n_disp; |
||||||
|
int winSize; |
||||||
|
|
||||||
|
virtual void SetUp() |
||||||
|
{ |
||||||
|
n_disp = GET_PARAM(0); |
||||||
|
winSize = GET_PARAM(1); |
||||||
|
} |
||||||
|
}; |
||||||
|
|
||||||
|
TEST_P(StereoMatchBM, Accuracy) |
||||||
|
{ |
||||||
|
|
||||||
|
Mat left_image = readImage(workdir + "../ocl/aloe-L.png", IMREAD_GRAYSCALE); |
||||||
|
Mat right_image = readImage(workdir + "../ocl/aloe-R.png", IMREAD_GRAYSCALE); |
||||||
|
Mat disp_gold = readImage(workdir + "../ocl/aloe-disp.png", IMREAD_GRAYSCALE); |
||||||
|
ocl::oclMat d_left, d_right; |
||||||
|
ocl::oclMat d_disp(left_image.size(), CV_8U); |
||||||
|
Mat disp; |
||||||
|
|
||||||
|
ASSERT_FALSE(left_image.empty()); |
||||||
|
ASSERT_FALSE(right_image.empty()); |
||||||
|
ASSERT_FALSE(disp_gold.empty()); |
||||||
|
d_left.upload(left_image); |
||||||
|
d_right.upload(right_image); |
||||||
|
|
||||||
|
ocl::StereoBM_OCL bm(0, n_disp, winSize); |
||||||
|
|
||||||
|
|
||||||
|
bm(d_left, d_right, d_disp); |
||||||
|
d_disp.download(disp); |
||||||
|
|
||||||
|
EXPECT_MAT_SIMILAR(disp_gold, disp, 1e-3); |
||||||
|
} |
||||||
|
|
||||||
|
INSTANTIATE_TEST_CASE_P(GPU_Calib3D, StereoMatchBM, testing::Combine(testing::Values(128), |
||||||
|
testing::Values(19))); |
||||||
|
|
||||||
|
#endif // HAVE_OPENCL
|
||||||
|
After Width: | Height: | Size: 720 KiB |
|
After Width: | Height: | Size: 722 KiB |
|
After Width: | Height: | Size: 59 KiB |
Loading…
Reference in new issue