Lines Matching full:radius
73 template<int RADIUS>
79 for(int i = 1; i <= RADIUS; i++) in CalcSSD()
86 if (threadIdx.x < BLOCK_W - RADIUS) in CalcSSD()
87 cache2 = col_ssd_cache[RADIUS]; in CalcSSD()
89 for(int i = RADIUS + 1; i < (2 * RADIUS + 1); i++) in CalcSSD()
95 template<int RADIUS>
100 //See above: #define COL_SSD_SIZE (BLOCK_W + 2 * RADIUS) in MinSSD()
101 ssd[0] = CalcSSD<RADIUS>(col_ssd_cache, col_ssd + 0 * (BLOCK_W + 2 * RADIUS)); in MinSSD()
103 ssd[1] = CalcSSD<RADIUS>(col_ssd_cache, col_ssd + 1 * (BLOCK_W + 2 * RADIUS)); in MinSSD()
105 ssd[2] = CalcSSD<RADIUS>(col_ssd_cache, col_ssd + 2 * (BLOCK_W + 2 * RADIUS)); in MinSSD()
107 ssd[3] = CalcSSD<RADIUS>(col_ssd_cache, col_ssd + 3 * (BLOCK_W + 2 * RADIUS)); in MinSSD()
109 ssd[4] = CalcSSD<RADIUS>(col_ssd_cache, col_ssd + 4 * (BLOCK_W + 2 * RADIUS)); in MinSSD()
111 ssd[5] = CalcSSD<RADIUS>(col_ssd_cache, col_ssd + 5 * (BLOCK_W + 2 * RADIUS)); in MinSSD()
113 ssd[6] = CalcSSD<RADIUS>(col_ssd_cache, col_ssd + 6 * (BLOCK_W + 2 * RADIUS)); in MinSSD()
115 ssd[7] = CalcSSD<RADIUS>(col_ssd_cache, col_ssd + 7 * (BLOCK_W + 2 * RADIUS)); in MinSSD()
129 template<int RADIUS>
162 //See above: #define COL_SSD_SIZE (BLOCK_W + 2 * RADIUS) in StepDown()
165 col_ssd[0 * (BLOCK_W + 2 * RADIUS)] += SQ(diff2) - SQ(diff1); in StepDown()
169 col_ssd[1 * (BLOCK_W + 2 * RADIUS)] += SQ(diff2) - SQ(diff1); in StepDown()
173 col_ssd[2 * (BLOCK_W + 2 * RADIUS)] += SQ(diff2) - SQ(diff1); in StepDown()
177 col_ssd[3 * (BLOCK_W + 2 * RADIUS)] += SQ(diff2) - SQ(diff1); in StepDown()
181 col_ssd[4 * (BLOCK_W + 2 * RADIUS)] += SQ(diff2) - SQ(diff1); in StepDown()
185 col_ssd[5 * (BLOCK_W + 2 * RADIUS)] += SQ(diff2) - SQ(diff1); in StepDown()
189 col_ssd[6 * (BLOCK_W + 2 * RADIUS)] += SQ(diff2) - SQ(diff1); in StepDown()
193 col_ssd[7 * (BLOCK_W + 2 * RADIUS)] += SQ(diff2) - SQ(diff1); in StepDown()
196 template<int RADIUS>
203 for(int i = 0; i < (2 * RADIUS + 1); i++) in InitColSSD()
220 //See above: #define COL_SSD_SIZE (BLOCK_W + 2 * RADIUS) in InitColSSD()
221 col_ssd[0 * (BLOCK_W + 2 * RADIUS)] = diffa[0]; in InitColSSD()
222 col_ssd[1 * (BLOCK_W + 2 * RADIUS)] = diffa[1]; in InitColSSD()
223 col_ssd[2 * (BLOCK_W + 2 * RADIUS)] = diffa[2]; in InitColSSD()
224 col_ssd[3 * (BLOCK_W + 2 * RADIUS)] = diffa[3]; in InitColSSD()
225 col_ssd[4 * (BLOCK_W + 2 * RADIUS)] = diffa[4]; in InitColSSD()
226 col_ssd[5 * (BLOCK_W + 2 * RADIUS)] = diffa[5]; in InitColSSD()
227 col_ssd[6 * (BLOCK_W + 2 * RADIUS)] = diffa[6]; in InitColSSD()
228 col_ssd[7 * (BLOCK_W + 2 * RADIUS)] = diffa[7]; in InitColSSD()
231 template<int RADIUS>
236 … int *col_ssd_extra = threadIdx.x < (2 * RADIUS) ? col_ssd + BLOCK_W : 0; //#define N_DIRTY_PIXEL… in stereoKernel()
239 int X = (blockIdx.x * BLOCK_W + threadIdx.x + maxdisp + RADIUS); in stereoKernel()
240 //#define Y (__mul24(blockIdx.y, ROWSperTHREAD) + RADIUS) in stereoKernel()
241 #define Y (blockIdx.y * ROWSperTHREAD + RADIUS) in stereoKernel()
242 //int Y = blockIdx.y * ROWSperTHREAD + RADIUS; in stereoKernel()
252 int end_row = ::min(ROWSperTHREAD, cheight - Y - RADIUS); in stereoKernel()
254 int x_tex = X - RADIUS; in stereoKernel()
261 y_tex = Y - RADIUS; in stereoKernel()
263 InitColSSD<RADIUS>(x_tex, y_tex, img_step, left, right, d, col_ssd); in stereoKernel()
267 … InitColSSD<RADIUS>(x_tex + BLOCK_W, y_tex, img_step, left, right, d, col_ssd_extra); in stereoKernel()
271 if (X < cwidth - RADIUS && Y < cheight - RADIUS) in stereoKernel()
273 uint2 minSSD = MinSSD<RADIUS>(col_ssd_cache + threadIdx.x, col_ssd); in stereoKernel()
284 int idx2 = (y_tex + (2 * RADIUS + 1)) * img_step + x_tex; in stereoKernel()
288 StepDown<RADIUS>(idx1, idx2, left, right, d, col_ssd); in stereoKernel()
292 … StepDown<RADIUS>(idx1, idx2, left + BLOCK_W, right + BLOCK_W, d, col_ssd_extra); in stereoKernel()
298 if (X < cwidth - RADIUS && row < cheight - RADIUS - Y) in stereoKernel()
301 uint2 minSSD = MinSSD<RADIUS>(col_ssd_cache + threadIdx.x, col_ssd); in stereoKernel()
313 …template<int RADIUS> void kernel_caller(const PtrStepSzb& left, const PtrStepSzb& right, const Ptr… in kernel_caller()
318 grid.x = divUp(left.cols - maxdisp - 2 * RADIUS, BLOCK_W); in kernel_caller()
319 grid.y = divUp(left.rows - 2 * RADIUS, ROWSperTHREAD); in kernel_caller()
321 //See above: #define COL_SSD_SIZE (BLOCK_W + 2 * RADIUS) in kernel_caller()
322 … size_t smem_size = (BLOCK_W + N_DISPARITIES * (BLOCK_W + 2 * RADIUS)) * sizeof(unsigned int); in kernel_caller()
324 …stereoKernel<RADIUS><<<grid, threads, smem_size, stream>>>(left.data, right.data, left.step, disp,… in kernel_caller()