1 /*M/////////////////////////////////////////////////////////////////////////////////////// 2 // 3 // IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. 4 // 5 // By downloading, copying, installing or using the software you agree to this license. 6 // If you do not agree to this license, do not download, install, 7 // copy or use the software. 8 // 9 // 10 // License Agreement 11 // For Open Source Computer Vision Library 12 // 13 // Copyright (C) 2000-2008, Intel Corporation, all rights reserved. 14 // Copyright (C) 2009, Willow Garage Inc., all rights reserved. 15 // Third party copyrights are property of their respective owners. 16 // 17 // Redistribution and use in source and binary forms, with or without modification, 18 // are permitted provided that the following conditions are met: 19 // 20 // * Redistribution's of source code must retain the above copyright notice, 21 // this list of conditions and the following disclaimer. 22 // 23 // * Redistribution's in binary form must reproduce the above copyright notice, 24 // this list of conditions and the following disclaimer in the documentation 25 // and/or other materials provided with the distribution. 26 // 27 // * The name of the copyright holders may not be used to endorse or promote products 28 // derived from this software without specific prior written permission. 29 // 30 // This software is provided by the copyright holders and contributors "as is" and 31 // any express or implied warranties, including, but not limited to, the implied 32 // warranties of merchantability and fitness for a particular purpose are disclaimed. 33 // In no event shall the Intel Corporation or contributors be liable for any direct, 34 // indirect, incidental, special, exemplary, or consequential damages 35 // (including, but not limited to, procurement of substitute goods or services; 36 // loss of use, data, or profits; or business interruption) however caused 37 // and on any theory of liability, whether in contract, strict liability, 38 // or tort (including negligence or otherwise) arising in any way out of 39 // the use of this software, even if advised of the possibility of such damage. 40 // 41 //M*/ 42 43 #if !defined CUDA_DISABLER 44 45 #include "opencv2/core/cuda/common.hpp" 46 47 namespace cv { namespace cuda { namespace device 48 { 49 namespace stereobm 50 { 51 ////////////////////////////////////////////////////////////////////////////////////////////////// 52 /////////////////////////////////////// Stereo BM //////////////////////////////////////////////// 53 ////////////////////////////////////////////////////////////////////////////////////////////////// 54 55 #define ROWSperTHREAD 21 // the number of rows a thread will process 56 57 #define BLOCK_W 128 // the thread block width (464) 58 #define N_DISPARITIES 8 59 60 #define STEREO_MIND 0 // The minimum d range to check 61 #define STEREO_DISP_STEP N_DISPARITIES // the d step, must be <= 1 to avoid aliasing 62 63 __constant__ unsigned int* cminSSDImage; 64 __constant__ size_t cminSSD_step; 65 __constant__ int cwidth; 66 __constant__ int cheight; 67 SQ(int a)68 __device__ __forceinline__ int SQ(int a) 69 { 70 return a * a; 71 } 72 73 template<int RADIUS> CalcSSD(volatile unsigned int * col_ssd_cache,volatile unsigned int * col_ssd)74 __device__ unsigned int CalcSSD(volatile unsigned int *col_ssd_cache, volatile unsigned int *col_ssd) 75 { 76 unsigned int cache = 0; 77 unsigned int cache2 = 0; 78 79 for(int i = 1; i <= RADIUS; i++) 80 cache += col_ssd[i]; 81 82 col_ssd_cache[0] = cache; 83 84 __syncthreads(); 85 86 if (threadIdx.x < BLOCK_W - RADIUS) 87 cache2 = col_ssd_cache[RADIUS]; 88 else 89 for(int i = RADIUS + 1; i < (2 * RADIUS + 1); i++) 90 cache2 += col_ssd[i]; 91 92 return col_ssd[0] + cache + cache2; 93 } 94 95 template<int RADIUS> MinSSD(volatile unsigned int * col_ssd_cache,volatile unsigned int * col_ssd)96 __device__ uint2 MinSSD(volatile unsigned int *col_ssd_cache, volatile unsigned int *col_ssd) 97 { 98 unsigned int ssd[N_DISPARITIES]; 99 100 //See above: #define COL_SSD_SIZE (BLOCK_W + 2 * RADIUS) 101 ssd[0] = CalcSSD<RADIUS>(col_ssd_cache, col_ssd + 0 * (BLOCK_W + 2 * RADIUS)); 102 __syncthreads(); 103 ssd[1] = CalcSSD<RADIUS>(col_ssd_cache, col_ssd + 1 * (BLOCK_W + 2 * RADIUS)); 104 __syncthreads(); 105 ssd[2] = CalcSSD<RADIUS>(col_ssd_cache, col_ssd + 2 * (BLOCK_W + 2 * RADIUS)); 106 __syncthreads(); 107 ssd[3] = CalcSSD<RADIUS>(col_ssd_cache, col_ssd + 3 * (BLOCK_W + 2 * RADIUS)); 108 __syncthreads(); 109 ssd[4] = CalcSSD<RADIUS>(col_ssd_cache, col_ssd + 4 * (BLOCK_W + 2 * RADIUS)); 110 __syncthreads(); 111 ssd[5] = CalcSSD<RADIUS>(col_ssd_cache, col_ssd + 5 * (BLOCK_W + 2 * RADIUS)); 112 __syncthreads(); 113 ssd[6] = CalcSSD<RADIUS>(col_ssd_cache, col_ssd + 6 * (BLOCK_W + 2 * RADIUS)); 114 __syncthreads(); 115 ssd[7] = CalcSSD<RADIUS>(col_ssd_cache, col_ssd + 7 * (BLOCK_W + 2 * RADIUS)); 116 117 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]))); 118 119 int bestIdx = 0; 120 for (int i = 0; i < N_DISPARITIES; i++) 121 { 122 if (mssd == ssd[i]) 123 bestIdx = i; 124 } 125 126 return make_uint2(mssd, bestIdx); 127 } 128 129 template<int RADIUS> StepDown(int idx1,int idx2,unsigned char * imageL,unsigned char * imageR,int d,volatile unsigned int * col_ssd)130 __device__ void StepDown(int idx1, int idx2, unsigned char* imageL, unsigned char* imageR, int d, volatile unsigned int *col_ssd) 131 { 132 unsigned char leftPixel1; 133 unsigned char leftPixel2; 134 unsigned char rightPixel1[8]; 135 unsigned char rightPixel2[8]; 136 unsigned int diff1, diff2; 137 138 leftPixel1 = imageL[idx1]; 139 leftPixel2 = imageL[idx2]; 140 141 idx1 = idx1 - d; 142 idx2 = idx2 - d; 143 144 rightPixel1[7] = imageR[idx1 - 7]; 145 rightPixel1[0] = imageR[idx1 - 0]; 146 rightPixel1[1] = imageR[idx1 - 1]; 147 rightPixel1[2] = imageR[idx1 - 2]; 148 rightPixel1[3] = imageR[idx1 - 3]; 149 rightPixel1[4] = imageR[idx1 - 4]; 150 rightPixel1[5] = imageR[idx1 - 5]; 151 rightPixel1[6] = imageR[idx1 - 6]; 152 153 rightPixel2[7] = imageR[idx2 - 7]; 154 rightPixel2[0] = imageR[idx2 - 0]; 155 rightPixel2[1] = imageR[idx2 - 1]; 156 rightPixel2[2] = imageR[idx2 - 2]; 157 rightPixel2[3] = imageR[idx2 - 3]; 158 rightPixel2[4] = imageR[idx2 - 4]; 159 rightPixel2[5] = imageR[idx2 - 5]; 160 rightPixel2[6] = imageR[idx2 - 6]; 161 162 //See above: #define COL_SSD_SIZE (BLOCK_W + 2 * RADIUS) 163 diff1 = leftPixel1 - rightPixel1[0]; 164 diff2 = leftPixel2 - rightPixel2[0]; 165 col_ssd[0 * (BLOCK_W + 2 * RADIUS)] += SQ(diff2) - SQ(diff1); 166 167 diff1 = leftPixel1 - rightPixel1[1]; 168 diff2 = leftPixel2 - rightPixel2[1]; 169 col_ssd[1 * (BLOCK_W + 2 * RADIUS)] += SQ(diff2) - SQ(diff1); 170 171 diff1 = leftPixel1 - rightPixel1[2]; 172 diff2 = leftPixel2 - rightPixel2[2]; 173 col_ssd[2 * (BLOCK_W + 2 * RADIUS)] += SQ(diff2) - SQ(diff1); 174 175 diff1 = leftPixel1 - rightPixel1[3]; 176 diff2 = leftPixel2 - rightPixel2[3]; 177 col_ssd[3 * (BLOCK_W + 2 * RADIUS)] += SQ(diff2) - SQ(diff1); 178 179 diff1 = leftPixel1 - rightPixel1[4]; 180 diff2 = leftPixel2 - rightPixel2[4]; 181 col_ssd[4 * (BLOCK_W + 2 * RADIUS)] += SQ(diff2) - SQ(diff1); 182 183 diff1 = leftPixel1 - rightPixel1[5]; 184 diff2 = leftPixel2 - rightPixel2[5]; 185 col_ssd[5 * (BLOCK_W + 2 * RADIUS)] += SQ(diff2) - SQ(diff1); 186 187 diff1 = leftPixel1 - rightPixel1[6]; 188 diff2 = leftPixel2 - rightPixel2[6]; 189 col_ssd[6 * (BLOCK_W + 2 * RADIUS)] += SQ(diff2) - SQ(diff1); 190 191 diff1 = leftPixel1 - rightPixel1[7]; 192 diff2 = leftPixel2 - rightPixel2[7]; 193 col_ssd[7 * (BLOCK_W + 2 * RADIUS)] += SQ(diff2) - SQ(diff1); 194 } 195 196 template<int RADIUS> InitColSSD(int x_tex,int y_tex,int im_pitch,unsigned char * imageL,unsigned char * imageR,int d,volatile unsigned int * col_ssd)197 __device__ void InitColSSD(int x_tex, int y_tex, int im_pitch, unsigned char* imageL, unsigned char* imageR, int d, volatile unsigned int *col_ssd) 198 { 199 unsigned char leftPixel1; 200 int idx; 201 unsigned int diffa[] = {0, 0, 0, 0, 0, 0, 0, 0}; 202 203 for(int i = 0; i < (2 * RADIUS + 1); i++) 204 { 205 idx = y_tex * im_pitch + x_tex; 206 leftPixel1 = imageL[idx]; 207 idx = idx - d; 208 209 diffa[0] += SQ(leftPixel1 - imageR[idx - 0]); 210 diffa[1] += SQ(leftPixel1 - imageR[idx - 1]); 211 diffa[2] += SQ(leftPixel1 - imageR[idx - 2]); 212 diffa[3] += SQ(leftPixel1 - imageR[idx - 3]); 213 diffa[4] += SQ(leftPixel1 - imageR[idx - 4]); 214 diffa[5] += SQ(leftPixel1 - imageR[idx - 5]); 215 diffa[6] += SQ(leftPixel1 - imageR[idx - 6]); 216 diffa[7] += SQ(leftPixel1 - imageR[idx - 7]); 217 218 y_tex += 1; 219 } 220 //See above: #define COL_SSD_SIZE (BLOCK_W + 2 * RADIUS) 221 col_ssd[0 * (BLOCK_W + 2 * RADIUS)] = diffa[0]; 222 col_ssd[1 * (BLOCK_W + 2 * RADIUS)] = diffa[1]; 223 col_ssd[2 * (BLOCK_W + 2 * RADIUS)] = diffa[2]; 224 col_ssd[3 * (BLOCK_W + 2 * RADIUS)] = diffa[3]; 225 col_ssd[4 * (BLOCK_W + 2 * RADIUS)] = diffa[4]; 226 col_ssd[5 * (BLOCK_W + 2 * RADIUS)] = diffa[5]; 227 col_ssd[6 * (BLOCK_W + 2 * RADIUS)] = diffa[6]; 228 col_ssd[7 * (BLOCK_W + 2 * RADIUS)] = diffa[7]; 229 } 230 231 template<int RADIUS> stereoKernel(unsigned char * left,unsigned char * right,size_t img_step,PtrStepb disp,int maxdisp)232 __global__ void stereoKernel(unsigned char *left, unsigned char *right, size_t img_step, PtrStepb disp, int maxdisp) 233 { 234 extern __shared__ unsigned int col_ssd_cache[]; 235 volatile unsigned int *col_ssd = col_ssd_cache + BLOCK_W + threadIdx.x; 236 volatile unsigned int *col_ssd_extra = threadIdx.x < (2 * RADIUS) ? col_ssd + BLOCK_W : 0; //#define N_DIRTY_PIXELS (2 * RADIUS) 237 238 //#define X (blockIdx.x * BLOCK_W + threadIdx.x + STEREO_MAXD) 239 int X = (blockIdx.x * BLOCK_W + threadIdx.x + maxdisp + RADIUS); 240 //#define Y (__mul24(blockIdx.y, ROWSperTHREAD) + RADIUS) 241 #define Y (blockIdx.y * ROWSperTHREAD + RADIUS) 242 //int Y = blockIdx.y * ROWSperTHREAD + RADIUS; 243 244 unsigned int* minSSDImage = cminSSDImage + X + Y * cminSSD_step; 245 unsigned char* disparImage = disp.data + X + Y * disp.step; 246 /* if (X < cwidth) 247 { 248 unsigned int *minSSDImage_end = minSSDImage + min(ROWSperTHREAD, cheight - Y) * minssd_step; 249 for(uint *ptr = minSSDImage; ptr != minSSDImage_end; ptr += minssd_step ) 250 *ptr = 0xFFFFFFFF; 251 }*/ 252 int end_row = ::min(ROWSperTHREAD, cheight - Y - RADIUS); 253 int y_tex; 254 int x_tex = X - RADIUS; 255 256 if (x_tex >= cwidth) 257 return; 258 259 for(int d = STEREO_MIND; d < maxdisp; d += STEREO_DISP_STEP) 260 { 261 y_tex = Y - RADIUS; 262 263 InitColSSD<RADIUS>(x_tex, y_tex, img_step, left, right, d, col_ssd); 264 265 if (col_ssd_extra > 0) 266 if (x_tex + BLOCK_W < cwidth) 267 InitColSSD<RADIUS>(x_tex + BLOCK_W, y_tex, img_step, left, right, d, col_ssd_extra); 268 269 __syncthreads(); //before MinSSD function 270 271 if (X < cwidth - RADIUS && Y < cheight - RADIUS) 272 { 273 uint2 minSSD = MinSSD<RADIUS>(col_ssd_cache + threadIdx.x, col_ssd); 274 if (minSSD.x < minSSDImage[0]) 275 { 276 disparImage[0] = (unsigned char)(d + minSSD.y); 277 minSSDImage[0] = minSSD.x; 278 } 279 } 280 281 for(int row = 1; row < end_row; row++) 282 { 283 int idx1 = y_tex * img_step + x_tex; 284 int idx2 = (y_tex + (2 * RADIUS + 1)) * img_step + x_tex; 285 286 __syncthreads(); 287 288 StepDown<RADIUS>(idx1, idx2, left, right, d, col_ssd); 289 290 if (col_ssd_extra) 291 if (x_tex + BLOCK_W < cwidth) 292 StepDown<RADIUS>(idx1, idx2, left + BLOCK_W, right + BLOCK_W, d, col_ssd_extra); 293 294 y_tex += 1; 295 296 __syncthreads(); //before MinSSD function 297 298 if (X < cwidth - RADIUS && row < cheight - RADIUS - Y) 299 { 300 int idx = row * cminSSD_step; 301 uint2 minSSD = MinSSD<RADIUS>(col_ssd_cache + threadIdx.x, col_ssd); 302 if (minSSD.x < minSSDImage[idx]) 303 { 304 disparImage[disp.step * row] = (unsigned char)(d + minSSD.y); 305 minSSDImage[idx] = minSSD.x; 306 } 307 } 308 } // for row loop 309 } // for d loop 310 } 311 312 kernel_caller(const PtrStepSzb & left,const PtrStepSzb & right,const PtrStepSzb & disp,int maxdisp,cudaStream_t & stream)313 template<int RADIUS> void kernel_caller(const PtrStepSzb& left, const PtrStepSzb& right, const PtrStepSzb& disp, int maxdisp, cudaStream_t & stream) 314 { 315 dim3 grid(1,1,1); 316 dim3 threads(BLOCK_W, 1, 1); 317 318 grid.x = divUp(left.cols - maxdisp - 2 * RADIUS, BLOCK_W); 319 grid.y = divUp(left.rows - 2 * RADIUS, ROWSperTHREAD); 320 321 //See above: #define COL_SSD_SIZE (BLOCK_W + 2 * RADIUS) 322 size_t smem_size = (BLOCK_W + N_DISPARITIES * (BLOCK_W + 2 * RADIUS)) * sizeof(unsigned int); 323 324 stereoKernel<RADIUS><<<grid, threads, smem_size, stream>>>(left.data, right.data, left.step, disp, maxdisp); 325 cudaSafeCall( cudaGetLastError() ); 326 327 if (stream == 0) 328 cudaSafeCall( cudaDeviceSynchronize() ); 329 }; 330 331 typedef void (*kernel_caller_t)(const PtrStepSzb& left, const PtrStepSzb& right, const PtrStepSzb& disp, int maxdisp, cudaStream_t & stream); 332 333 const static kernel_caller_t callers[] = 334 { 335 0, 336 kernel_caller< 1>, kernel_caller< 2>, kernel_caller< 3>, kernel_caller< 4>, kernel_caller< 5>, 337 kernel_caller< 6>, kernel_caller< 7>, kernel_caller< 8>, kernel_caller< 9>, kernel_caller<10>, 338 kernel_caller<11>, kernel_caller<12>, kernel_caller<13>, kernel_caller<14>, kernel_caller<15>, 339 kernel_caller<16>, kernel_caller<17>, kernel_caller<18>, kernel_caller<19>, kernel_caller<20>, 340 kernel_caller<21>, kernel_caller<22>, kernel_caller<23>, kernel_caller<24>, kernel_caller<25> 341 342 //0,0,0, 0,0,0, 0,0,kernel_caller<9> 343 }; 344 const int calles_num = sizeof(callers)/sizeof(callers[0]); 345 stereoBM_CUDA(const PtrStepSzb & left,const PtrStepSzb & right,const PtrStepSzb & disp,int maxdisp,int winsz,const PtrStepSz<unsigned int> & minSSD_buf,cudaStream_t & stream)346 void stereoBM_CUDA(const PtrStepSzb& left, const PtrStepSzb& right, const PtrStepSzb& disp, int maxdisp, int winsz, const PtrStepSz<unsigned int>& minSSD_buf, cudaStream_t& stream) 347 { 348 int winsz2 = winsz >> 1; 349 350 if (winsz2 == 0 || winsz2 >= calles_num) 351 CV_Error(cv::Error::StsBadArg, "Unsupported window size"); 352 353 //cudaSafeCall( cudaFuncSetCacheConfig(&stereoKernel, cudaFuncCachePreferL1) ); 354 //cudaSafeCall( cudaFuncSetCacheConfig(&stereoKernel, cudaFuncCachePreferShared) ); 355 356 cudaSafeCall( cudaMemset2D(disp.data, disp.step, 0, disp.cols, disp.rows) ); 357 cudaSafeCall( cudaMemset2D(minSSD_buf.data, minSSD_buf.step, 0xFF, minSSD_buf.cols * minSSD_buf.elemSize(), disp.rows) ); 358 359 cudaSafeCall( cudaMemcpyToSymbol( cwidth, &left.cols, sizeof(left.cols) ) ); 360 cudaSafeCall( cudaMemcpyToSymbol( cheight, &left.rows, sizeof(left.rows) ) ); 361 cudaSafeCall( cudaMemcpyToSymbol( cminSSDImage, &minSSD_buf.data, sizeof(minSSD_buf.data) ) ); 362 363 size_t minssd_step = minSSD_buf.step/minSSD_buf.elemSize(); 364 cudaSafeCall( cudaMemcpyToSymbol( cminSSD_step, &minssd_step, sizeof(minssd_step) ) ); 365 366 callers[winsz2](left, right, disp, maxdisp, stream); 367 } 368 369 ////////////////////////////////////////////////////////////////////////////////////////////////// 370 /////////////////////////////////////// Sobel Prefiler /////////////////////////////////////////// 371 ////////////////////////////////////////////////////////////////////////////////////////////////// 372 373 texture<unsigned char, 2, cudaReadModeElementType> texForSobel; 374 prefilter_kernel(PtrStepSzb output,int prefilterCap)375 __global__ void prefilter_kernel(PtrStepSzb output, int prefilterCap) 376 { 377 int x = blockDim.x * blockIdx.x + threadIdx.x; 378 int y = blockDim.y * blockIdx.y + threadIdx.y; 379 380 if (x < output.cols && y < output.rows) 381 { 382 int conv = (int)tex2D(texForSobel, x - 1, y - 1) * (-1) + (int)tex2D(texForSobel, x + 1, y - 1) * (1) + 383 (int)tex2D(texForSobel, x - 1, y ) * (-2) + (int)tex2D(texForSobel, x + 1, y ) * (2) + 384 (int)tex2D(texForSobel, x - 1, y + 1) * (-1) + (int)tex2D(texForSobel, x + 1, y + 1) * (1); 385 386 387 conv = ::min(::min(::max(-prefilterCap, conv), prefilterCap) + prefilterCap, 255); 388 output.ptr(y)[x] = conv & 0xFF; 389 } 390 } 391 prefilter_xsobel(const PtrStepSzb & input,const PtrStepSzb & output,int prefilterCap,cudaStream_t & stream)392 void prefilter_xsobel(const PtrStepSzb& input, const PtrStepSzb& output, int prefilterCap, cudaStream_t & stream) 393 { 394 cudaChannelFormatDesc desc = cudaCreateChannelDesc<unsigned char>(); 395 cudaSafeCall( cudaBindTexture2D( 0, texForSobel, input.data, desc, input.cols, input.rows, input.step ) ); 396 397 dim3 threads(16, 16, 1); 398 dim3 grid(1, 1, 1); 399 400 grid.x = divUp(input.cols, threads.x); 401 grid.y = divUp(input.rows, threads.y); 402 403 prefilter_kernel<<<grid, threads, 0, stream>>>(output, prefilterCap); 404 cudaSafeCall( cudaGetLastError() ); 405 406 if (stream == 0) 407 cudaSafeCall( cudaDeviceSynchronize() ); 408 409 cudaSafeCall( cudaUnbindTexture (texForSobel ) ); 410 } 411 412 413 ////////////////////////////////////////////////////////////////////////////////////////////////// 414 /////////////////////////////////// Textureness filtering //////////////////////////////////////// 415 ////////////////////////////////////////////////////////////////////////////////////////////////// 416 417 texture<unsigned char, 2, cudaReadModeNormalizedFloat> texForTF; 418 sobel(int x,int y)419 __device__ __forceinline__ float sobel(int x, int y) 420 { 421 float conv = tex2D(texForTF, x - 1, y - 1) * (-1) + tex2D(texForTF, x + 1, y - 1) * (1) + 422 tex2D(texForTF, x - 1, y ) * (-2) + tex2D(texForTF, x + 1, y ) * (2) + 423 tex2D(texForTF, x - 1, y + 1) * (-1) + tex2D(texForTF, x + 1, y + 1) * (1); 424 return fabs(conv); 425 } 426 CalcSums(float * cols,float * cols_cache,int winsz)427 __device__ float CalcSums(float *cols, float *cols_cache, int winsz) 428 { 429 float cache = 0; 430 float cache2 = 0; 431 int winsz2 = winsz/2; 432 433 for(int i = 1; i <= winsz2; i++) 434 cache += cols[i]; 435 436 cols_cache[0] = cache; 437 438 __syncthreads(); 439 440 if (threadIdx.x < blockDim.x - winsz2) 441 cache2 = cols_cache[winsz2]; 442 else 443 for(int i = winsz2 + 1; i < winsz; i++) 444 cache2 += cols[i]; 445 446 return cols[0] + cache + cache2; 447 } 448 449 #define RpT (2 * ROWSperTHREAD) // got experimentally 450 textureness_kernel(PtrStepSzb disp,int winsz,float threshold)451 __global__ void textureness_kernel(PtrStepSzb disp, int winsz, float threshold) 452 { 453 int winsz2 = winsz/2; 454 int n_dirty_pixels = (winsz2) * 2; 455 456 extern __shared__ float cols_cache[]; 457 float *cols = cols_cache + blockDim.x + threadIdx.x; 458 float *cols_extra = threadIdx.x < n_dirty_pixels ? cols + blockDim.x : 0; 459 460 int x = blockIdx.x * blockDim.x + threadIdx.x; 461 int beg_row = blockIdx.y * RpT; 462 int end_row = ::min(beg_row + RpT, disp.rows); 463 464 if (x < disp.cols) 465 { 466 int y = beg_row; 467 468 float sum = 0; 469 float sum_extra = 0; 470 471 for(int i = y - winsz2; i <= y + winsz2; ++i) 472 { 473 sum += sobel(x - winsz2, i); 474 if (cols_extra) 475 sum_extra += sobel(x + blockDim.x - winsz2, i); 476 } 477 *cols = sum; 478 if (cols_extra) 479 *cols_extra = sum_extra; 480 481 __syncthreads(); 482 483 float sum_win = CalcSums(cols, cols_cache + threadIdx.x, winsz) * 255; 484 if (sum_win < threshold) 485 disp.data[y * disp.step + x] = 0; 486 487 __syncthreads(); 488 489 for(int y = beg_row + 1; y < end_row; ++y) 490 { 491 sum = sum - sobel(x - winsz2, y - winsz2 - 1) + sobel(x - winsz2, y + winsz2); 492 *cols = sum; 493 494 if (cols_extra) 495 { 496 sum_extra = sum_extra - sobel(x + blockDim.x - winsz2, y - winsz2 - 1) + sobel(x + blockDim.x - winsz2, y + winsz2); 497 *cols_extra = sum_extra; 498 } 499 500 __syncthreads(); 501 float sum_win = CalcSums(cols, cols_cache + threadIdx.x, winsz) * 255; 502 if (sum_win < threshold) 503 disp.data[y * disp.step + x] = 0; 504 505 __syncthreads(); 506 } 507 } 508 } 509 postfilter_textureness(const PtrStepSzb & input,int winsz,float avgTexturenessThreshold,const PtrStepSzb & disp,cudaStream_t & stream)510 void postfilter_textureness(const PtrStepSzb& input, int winsz, float avgTexturenessThreshold, const PtrStepSzb& disp, cudaStream_t & stream) 511 { 512 avgTexturenessThreshold *= winsz * winsz; 513 514 texForTF.filterMode = cudaFilterModeLinear; 515 texForTF.addressMode[0] = cudaAddressModeWrap; 516 texForTF.addressMode[1] = cudaAddressModeWrap; 517 518 cudaChannelFormatDesc desc = cudaCreateChannelDesc<unsigned char>(); 519 cudaSafeCall( cudaBindTexture2D( 0, texForTF, input.data, desc, input.cols, input.rows, input.step ) ); 520 521 dim3 threads(128, 1, 1); 522 dim3 grid(1, 1, 1); 523 524 grid.x = divUp(input.cols, threads.x); 525 grid.y = divUp(input.rows, RpT); 526 527 size_t smem_size = (threads.x + threads.x + (winsz/2) * 2 ) * sizeof(float); 528 textureness_kernel<<<grid, threads, smem_size, stream>>>(disp, winsz, avgTexturenessThreshold); 529 cudaSafeCall( cudaGetLastError() ); 530 531 if (stream == 0) 532 cudaSafeCall( cudaDeviceSynchronize() ); 533 534 cudaSafeCall( cudaUnbindTexture (texForTF) ); 535 } 536 } // namespace stereobm 537 }}} // namespace cv { namespace cuda { namespace cudev 538 539 540 #endif /* CUDA_DISABLER */ 541