1 /*M///////////////////////////////////////////////////////////////////////////////////////
3 // IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
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.
11 // For Open Source Computer Vision Library
13 // Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
14 // Copyright (C) 2009, Willow Garage Inc., all rights reserved.
15 // Copyright (C) 1993-2011, NVIDIA Corporation, all rights reserved.
16 // Third party copyrights are property of their respective owners.
18 // Redistribution and use in source and binary forms, with or without modification,
19 // are permitted provided that the following conditions are met:
21 // * Redistribution's of source code must retain the above copyright notice,
22 // this list of conditions and the following disclaimer.
24 // * Redistribution's in binary form must reproduce the above copyright notice,
25 // this list of conditions and the following disclaimer in the documentation
26 // and/or other materials provided with the distribution.
28 // * The name of the copyright holders may not be used to endorse or promote products
29 // derived from this software without specific prior written permission.
31 // This software is provided by the copyright holders and contributors "as is" and
32 // any express or implied warranties, including, but not limited to, the implied
33 // warranties of merchantability and fitness for a particular purpose are disclaimed.
34 // In no event shall the Intel Corporation or contributors be liable for any direct,
35 // indirect, incidental, special, exemplary, or consequential damages
36 // (including, but not limited to, procurement of substitute goods or services;
37 // loss of use, data, or profits; or business interruption) however caused
38 // and on any theory of liability, whether in contract, strict liability,
39 // or tort (including negligence or otherwise) arising in any way out of
40 // the use of this software, even if advised of the possibility of such damage.
44 #if !defined CUDA_DISABLER
46 #include "internal_shared.hpp"
47 #include "opencv2/gpu/device/saturate_cast.hpp"
48 #include "opencv2/gpu/device/vec_math.hpp"
49 #include "opencv2/gpu/device/limits.hpp"
50 #include "opencv2/gpu/device/border_interpolate.hpp"
51 #include "opencv2/gpu/device/static_check.hpp"
53 namespace cv { namespace gpu { namespace device
55 namespace column_filter
57 #define MAX_KERNEL_SIZE 32
59 __constant__ float c_kernel[MAX_KERNEL_SIZE];
61 void loadKernel(const float* kernel, int ksize, cudaStream_t stream)
64 cudaSafeCall( cudaMemcpyToSymbol(c_kernel, kernel, ksize * sizeof(float), 0, cudaMemcpyDeviceToDevice) );
66 cudaSafeCall( cudaMemcpyToSymbolAsync(c_kernel, kernel, ksize * sizeof(float), 0, cudaMemcpyDeviceToDevice, stream) );
69 template <int KSIZE, typename T, typename D, typename B>
70 __global__ void linearColumnFilter(const PtrStepSz<T> src, PtrStep<D> dst, const int anchor, const B brd)
72 #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 200)
73 const int BLOCK_DIM_X = 16;
74 const int BLOCK_DIM_Y = 16;
75 const int PATCH_PER_BLOCK = 4;
76 const int HALO_SIZE = KSIZE <= 16 ? 1 : 2;
78 const int BLOCK_DIM_X = 16;
79 const int BLOCK_DIM_Y = 8;
80 const int PATCH_PER_BLOCK = 2;
81 const int HALO_SIZE = 2;
84 typedef typename TypeVec<float, VecTraits<T>::cn>::vec_type sum_t;
86 __shared__ sum_t smem[(PATCH_PER_BLOCK + 2 * HALO_SIZE) * BLOCK_DIM_Y][BLOCK_DIM_X];
88 const int x = blockIdx.x * BLOCK_DIM_X + threadIdx.x;
93 const T* src_col = src.ptr() + x;
95 const int yStart = blockIdx.y * (BLOCK_DIM_Y * PATCH_PER_BLOCK) + threadIdx.y;
101 for (int j = 0; j < HALO_SIZE; ++j)
102 smem[threadIdx.y + j * BLOCK_DIM_Y][threadIdx.x] = saturate_cast<sum_t>(src(yStart - (HALO_SIZE - j) * BLOCK_DIM_Y, x));
108 for (int j = 0; j < HALO_SIZE; ++j)
109 smem[threadIdx.y + j * BLOCK_DIM_Y][threadIdx.x] = saturate_cast<sum_t>(brd.at_low(yStart - (HALO_SIZE - j) * BLOCK_DIM_Y, src_col, src.step));
112 if (blockIdx.y + 2 < gridDim.y)
116 for (int j = 0; j < PATCH_PER_BLOCK; ++j)
117 smem[threadIdx.y + HALO_SIZE * BLOCK_DIM_Y + j * BLOCK_DIM_Y][threadIdx.x] = saturate_cast<sum_t>(src(yStart + j * BLOCK_DIM_Y, x));
121 for (int j = 0; j < HALO_SIZE; ++j)
122 smem[threadIdx.y + (PATCH_PER_BLOCK + HALO_SIZE) * BLOCK_DIM_Y + j * BLOCK_DIM_Y][threadIdx.x] = saturate_cast<sum_t>(src(yStart + (PATCH_PER_BLOCK + j) * BLOCK_DIM_Y, x));
128 for (int j = 0; j < PATCH_PER_BLOCK; ++j)
129 smem[threadIdx.y + HALO_SIZE * BLOCK_DIM_Y + j * BLOCK_DIM_Y][threadIdx.x] = saturate_cast<sum_t>(brd.at_high(yStart + j * BLOCK_DIM_Y, src_col, src.step));
133 for (int j = 0; j < HALO_SIZE; ++j)
134 smem[threadIdx.y + (PATCH_PER_BLOCK + HALO_SIZE) * BLOCK_DIM_Y + j * BLOCK_DIM_Y][threadIdx.x] = saturate_cast<sum_t>(brd.at_high(yStart + (PATCH_PER_BLOCK + j) * BLOCK_DIM_Y, src_col, src.step));
140 for (int j = 0; j < PATCH_PER_BLOCK; ++j)
142 const int y = yStart + j * BLOCK_DIM_Y;
146 sum_t sum = VecTraits<sum_t>::all(0);
149 for (int k = 0; k < KSIZE; ++k)
150 sum = sum + smem[threadIdx.y + HALO_SIZE * BLOCK_DIM_Y + j * BLOCK_DIM_Y - anchor + k][threadIdx.x] * c_kernel[k];
152 dst(y, x) = saturate_cast<D>(sum);
157 template <int KSIZE, typename T, typename D, template<typename> class B>
158 void linearColumnFilter_caller(PtrStepSz<T> src, PtrStepSz<D> dst, int anchor, int cc, cudaStream_t stream)
177 const dim3 block(BLOCK_DIM_X, BLOCK_DIM_Y);
178 const dim3 grid(divUp(src.cols, BLOCK_DIM_X), divUp(src.rows, BLOCK_DIM_Y * PATCH_PER_BLOCK));
182 linearColumnFilter<KSIZE, T, D><<<grid, block, 0, stream>>>(src, dst, anchor, brd);
184 cudaSafeCall( cudaGetLastError() );
187 cudaSafeCall( cudaDeviceSynchronize() );
190 template <typename T, typename D>
191 void linearColumnFilter_gpu(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream)
193 typedef void (*caller_t)(PtrStepSz<T> src, PtrStepSz<D> dst, int anchor, int cc, cudaStream_t stream);
195 static const caller_t callers[5][33] =
199 linearColumnFilter_caller< 1, T, D, BrdColReflect101>,
200 linearColumnFilter_caller< 2, T, D, BrdColReflect101>,
201 linearColumnFilter_caller< 3, T, D, BrdColReflect101>,
202 linearColumnFilter_caller< 4, T, D, BrdColReflect101>,
203 linearColumnFilter_caller< 5, T, D, BrdColReflect101>,
204 linearColumnFilter_caller< 6, T, D, BrdColReflect101>,
205 linearColumnFilter_caller< 7, T, D, BrdColReflect101>,
206 linearColumnFilter_caller< 8, T, D, BrdColReflect101>,
207 linearColumnFilter_caller< 9, T, D, BrdColReflect101>,
208 linearColumnFilter_caller<10, T, D, BrdColReflect101>,
209 linearColumnFilter_caller<11, T, D, BrdColReflect101>,
210 linearColumnFilter_caller<12, T, D, BrdColReflect101>,
211 linearColumnFilter_caller<13, T, D, BrdColReflect101>,
212 linearColumnFilter_caller<14, T, D, BrdColReflect101>,
213 linearColumnFilter_caller<15, T, D, BrdColReflect101>,
214 linearColumnFilter_caller<16, T, D, BrdColReflect101>,
215 linearColumnFilter_caller<17, T, D, BrdColReflect101>,
216 linearColumnFilter_caller<18, T, D, BrdColReflect101>,
217 linearColumnFilter_caller<19, T, D, BrdColReflect101>,
218 linearColumnFilter_caller<20, T, D, BrdColReflect101>,
219 linearColumnFilter_caller<21, T, D, BrdColReflect101>,
220 linearColumnFilter_caller<22, T, D, BrdColReflect101>,
221 linearColumnFilter_caller<23, T, D, BrdColReflect101>,
222 linearColumnFilter_caller<24, T, D, BrdColReflect101>,
223 linearColumnFilter_caller<25, T, D, BrdColReflect101>,
224 linearColumnFilter_caller<26, T, D, BrdColReflect101>,
225 linearColumnFilter_caller<27, T, D, BrdColReflect101>,
226 linearColumnFilter_caller<28, T, D, BrdColReflect101>,
227 linearColumnFilter_caller<29, T, D, BrdColReflect101>,
228 linearColumnFilter_caller<30, T, D, BrdColReflect101>,
229 linearColumnFilter_caller<31, T, D, BrdColReflect101>,
230 linearColumnFilter_caller<32, T, D, BrdColReflect101>
234 linearColumnFilter_caller< 1, T, D, BrdColReplicate>,
235 linearColumnFilter_caller< 2, T, D, BrdColReplicate>,
236 linearColumnFilter_caller< 3, T, D, BrdColReplicate>,
237 linearColumnFilter_caller< 4, T, D, BrdColReplicate>,
238 linearColumnFilter_caller< 5, T, D, BrdColReplicate>,
239 linearColumnFilter_caller< 6, T, D, BrdColReplicate>,
240 linearColumnFilter_caller< 7, T, D, BrdColReplicate>,
241 linearColumnFilter_caller< 8, T, D, BrdColReplicate>,
242 linearColumnFilter_caller< 9, T, D, BrdColReplicate>,
243 linearColumnFilter_caller<10, T, D, BrdColReplicate>,
244 linearColumnFilter_caller<11, T, D, BrdColReplicate>,
245 linearColumnFilter_caller<12, T, D, BrdColReplicate>,
246 linearColumnFilter_caller<13, T, D, BrdColReplicate>,
247 linearColumnFilter_caller<14, T, D, BrdColReplicate>,
248 linearColumnFilter_caller<15, T, D, BrdColReplicate>,
249 linearColumnFilter_caller<16, T, D, BrdColReplicate>,
250 linearColumnFilter_caller<17, T, D, BrdColReplicate>,
251 linearColumnFilter_caller<18, T, D, BrdColReplicate>,
252 linearColumnFilter_caller<19, T, D, BrdColReplicate>,
253 linearColumnFilter_caller<20, T, D, BrdColReplicate>,
254 linearColumnFilter_caller<21, T, D, BrdColReplicate>,
255 linearColumnFilter_caller<22, T, D, BrdColReplicate>,
256 linearColumnFilter_caller<23, T, D, BrdColReplicate>,
257 linearColumnFilter_caller<24, T, D, BrdColReplicate>,
258 linearColumnFilter_caller<25, T, D, BrdColReplicate>,
259 linearColumnFilter_caller<26, T, D, BrdColReplicate>,
260 linearColumnFilter_caller<27, T, D, BrdColReplicate>,
261 linearColumnFilter_caller<28, T, D, BrdColReplicate>,
262 linearColumnFilter_caller<29, T, D, BrdColReplicate>,
263 linearColumnFilter_caller<30, T, D, BrdColReplicate>,
264 linearColumnFilter_caller<31, T, D, BrdColReplicate>,
265 linearColumnFilter_caller<32, T, D, BrdColReplicate>
269 linearColumnFilter_caller< 1, T, D, BrdColConstant>,
270 linearColumnFilter_caller< 2, T, D, BrdColConstant>,
271 linearColumnFilter_caller< 3, T, D, BrdColConstant>,
272 linearColumnFilter_caller< 4, T, D, BrdColConstant>,
273 linearColumnFilter_caller< 5, T, D, BrdColConstant>,
274 linearColumnFilter_caller< 6, T, D, BrdColConstant>,
275 linearColumnFilter_caller< 7, T, D, BrdColConstant>,
276 linearColumnFilter_caller< 8, T, D, BrdColConstant>,
277 linearColumnFilter_caller< 9, T, D, BrdColConstant>,
278 linearColumnFilter_caller<10, T, D, BrdColConstant>,
279 linearColumnFilter_caller<11, T, D, BrdColConstant>,
280 linearColumnFilter_caller<12, T, D, BrdColConstant>,
281 linearColumnFilter_caller<13, T, D, BrdColConstant>,
282 linearColumnFilter_caller<14, T, D, BrdColConstant>,
283 linearColumnFilter_caller<15, T, D, BrdColConstant>,
284 linearColumnFilter_caller<16, T, D, BrdColConstant>,
285 linearColumnFilter_caller<17, T, D, BrdColConstant>,
286 linearColumnFilter_caller<18, T, D, BrdColConstant>,
287 linearColumnFilter_caller<19, T, D, BrdColConstant>,
288 linearColumnFilter_caller<20, T, D, BrdColConstant>,
289 linearColumnFilter_caller<21, T, D, BrdColConstant>,
290 linearColumnFilter_caller<22, T, D, BrdColConstant>,
291 linearColumnFilter_caller<23, T, D, BrdColConstant>,
292 linearColumnFilter_caller<24, T, D, BrdColConstant>,
293 linearColumnFilter_caller<25, T, D, BrdColConstant>,
294 linearColumnFilter_caller<26, T, D, BrdColConstant>,
295 linearColumnFilter_caller<27, T, D, BrdColConstant>,
296 linearColumnFilter_caller<28, T, D, BrdColConstant>,
297 linearColumnFilter_caller<29, T, D, BrdColConstant>,
298 linearColumnFilter_caller<30, T, D, BrdColConstant>,
299 linearColumnFilter_caller<31, T, D, BrdColConstant>,
300 linearColumnFilter_caller<32, T, D, BrdColConstant>
304 linearColumnFilter_caller< 1, T, D, BrdColReflect>,
305 linearColumnFilter_caller< 2, T, D, BrdColReflect>,
306 linearColumnFilter_caller< 3, T, D, BrdColReflect>,
307 linearColumnFilter_caller< 4, T, D, BrdColReflect>,
308 linearColumnFilter_caller< 5, T, D, BrdColReflect>,
309 linearColumnFilter_caller< 6, T, D, BrdColReflect>,
310 linearColumnFilter_caller< 7, T, D, BrdColReflect>,
311 linearColumnFilter_caller< 8, T, D, BrdColReflect>,
312 linearColumnFilter_caller< 9, T, D, BrdColReflect>,
313 linearColumnFilter_caller<10, T, D, BrdColReflect>,
314 linearColumnFilter_caller<11, T, D, BrdColReflect>,
315 linearColumnFilter_caller<12, T, D, BrdColReflect>,
316 linearColumnFilter_caller<13, T, D, BrdColReflect>,
317 linearColumnFilter_caller<14, T, D, BrdColReflect>,
318 linearColumnFilter_caller<15, T, D, BrdColReflect>,
319 linearColumnFilter_caller<16, T, D, BrdColReflect>,
320 linearColumnFilter_caller<17, T, D, BrdColReflect>,
321 linearColumnFilter_caller<18, T, D, BrdColReflect>,
322 linearColumnFilter_caller<19, T, D, BrdColReflect>,
323 linearColumnFilter_caller<20, T, D, BrdColReflect>,
324 linearColumnFilter_caller<21, T, D, BrdColReflect>,
325 linearColumnFilter_caller<22, T, D, BrdColReflect>,
326 linearColumnFilter_caller<23, T, D, BrdColReflect>,
327 linearColumnFilter_caller<24, T, D, BrdColReflect>,
328 linearColumnFilter_caller<25, T, D, BrdColReflect>,
329 linearColumnFilter_caller<26, T, D, BrdColReflect>,
330 linearColumnFilter_caller<27, T, D, BrdColReflect>,
331 linearColumnFilter_caller<28, T, D, BrdColReflect>,
332 linearColumnFilter_caller<29, T, D, BrdColReflect>,
333 linearColumnFilter_caller<30, T, D, BrdColReflect>,
334 linearColumnFilter_caller<31, T, D, BrdColReflect>,
335 linearColumnFilter_caller<32, T, D, BrdColReflect>
339 linearColumnFilter_caller< 1, T, D, BrdColWrap>,
340 linearColumnFilter_caller< 2, T, D, BrdColWrap>,
341 linearColumnFilter_caller< 3, T, D, BrdColWrap>,
342 linearColumnFilter_caller< 4, T, D, BrdColWrap>,
343 linearColumnFilter_caller< 5, T, D, BrdColWrap>,
344 linearColumnFilter_caller< 6, T, D, BrdColWrap>,
345 linearColumnFilter_caller< 7, T, D, BrdColWrap>,
346 linearColumnFilter_caller< 8, T, D, BrdColWrap>,
347 linearColumnFilter_caller< 9, T, D, BrdColWrap>,
348 linearColumnFilter_caller<10, T, D, BrdColWrap>,
349 linearColumnFilter_caller<11, T, D, BrdColWrap>,
350 linearColumnFilter_caller<12, T, D, BrdColWrap>,
351 linearColumnFilter_caller<13, T, D, BrdColWrap>,
352 linearColumnFilter_caller<14, T, D, BrdColWrap>,
353 linearColumnFilter_caller<15, T, D, BrdColWrap>,
354 linearColumnFilter_caller<16, T, D, BrdColWrap>,
355 linearColumnFilter_caller<17, T, D, BrdColWrap>,
356 linearColumnFilter_caller<18, T, D, BrdColWrap>,
357 linearColumnFilter_caller<19, T, D, BrdColWrap>,
358 linearColumnFilter_caller<20, T, D, BrdColWrap>,
359 linearColumnFilter_caller<21, T, D, BrdColWrap>,
360 linearColumnFilter_caller<22, T, D, BrdColWrap>,
361 linearColumnFilter_caller<23, T, D, BrdColWrap>,
362 linearColumnFilter_caller<24, T, D, BrdColWrap>,
363 linearColumnFilter_caller<25, T, D, BrdColWrap>,
364 linearColumnFilter_caller<26, T, D, BrdColWrap>,
365 linearColumnFilter_caller<27, T, D, BrdColWrap>,
366 linearColumnFilter_caller<28, T, D, BrdColWrap>,
367 linearColumnFilter_caller<29, T, D, BrdColWrap>,
368 linearColumnFilter_caller<30, T, D, BrdColWrap>,
369 linearColumnFilter_caller<31, T, D, BrdColWrap>,
370 linearColumnFilter_caller<32, T, D, BrdColWrap>
374 loadKernel(kernel, ksize, stream);
376 callers[brd_type][ksize]((PtrStepSz<T>)src, (PtrStepSz<D>)dst, anchor, cc, stream);
379 template void linearColumnFilter_gpu<float , uchar >(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
380 template void linearColumnFilter_gpu<float3, uchar3>(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
381 template void linearColumnFilter_gpu<float4, uchar4>(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
382 template void linearColumnFilter_gpu<float3, short3>(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
383 template void linearColumnFilter_gpu<float , int >(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
384 template void linearColumnFilter_gpu<float , float >(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
385 template void linearColumnFilter_gpu<float3, float3>(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
386 template void linearColumnFilter_gpu<float4, float4>(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);
387 } // namespace column_filter
388 }}} // namespace cv { namespace gpu { namespace device
391 #endif /* CUDA_DISABLER */