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 // Copyright (C) 2013, OpenCV Foundation, all rights reserved.
16 // Third party copyrights are property of their respective owners.
17 //
18 // Redistribution and use in source and binary forms, with or without modification,
19 // are permitted provided that the following conditions are met:
20 //
21 // * Redistribution's of source code must retain the above copyright notice,
22 // this list of conditions and the following disclaimer.
23 //
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.
27 //
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.
30 //
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.
41 //
42 //M*/
43
44 #pragma once
45
46 #ifndef __OPENCV_CUDEV_BLOCK_SCAN_HPP__
47 #define __OPENCV_CUDEV_BLOCK_SCAN_HPP__
48
49 #include "../common.hpp"
50 #include "../warp/scan.hpp"
51
52 namespace cv { namespace cudev {
53
54 //! @addtogroup cudev
55 //! @{
56
57 template <int THREADS_NUM, typename T>
blockScanInclusive(T data,volatile T * smem,uint tid)58 __device__ T blockScanInclusive(T data, volatile T* smem, uint tid)
59 {
60 if (THREADS_NUM > WARP_SIZE)
61 {
62 // bottom-level inclusive warp scan
63 T warpResult = warpScanInclusive(data, smem, tid);
64
65 __syncthreads();
66
67 // save top elements of each warp for exclusive warp scan
68 // sync to wait for warp scans to complete (because s_Data is being overwritten)
69 if ((tid & (WARP_SIZE - 1)) == (WARP_SIZE - 1))
70 {
71 smem[tid >> LOG_WARP_SIZE] = warpResult;
72 }
73
74 __syncthreads();
75
76 if (tid < (THREADS_NUM / WARP_SIZE))
77 {
78 // grab top warp elements
79 T val = smem[tid];
80
81 // calculate exclusive scan and write back to shared memory
82 smem[tid] = warpScanExclusive(val, smem, tid);
83 }
84
85 __syncthreads();
86
87 // return updated warp scans with exclusive scan results
88 return warpResult + smem[tid >> LOG_WARP_SIZE];
89 }
90 else
91 {
92 return warpScanInclusive(data, smem, tid);
93 }
94 }
95
96 template <int THREADS_NUM, typename T>
blockScanExclusive(T data,volatile T * smem,uint tid)97 __device__ __forceinline__ T blockScanExclusive(T data, volatile T* smem, uint tid)
98 {
99 return blockScanInclusive<THREADS_NUM>(data, smem, tid) - data;
100 }
101
102 //! @}
103
104 }}
105
106 #endif
107