Skip to content

Commit 18439b4

Browse files
ngleenglee
committed
[moved from opencv] __shfl_up_sync with mask for CUDA >= 9
* __shfl_up_sync with proper mask value for CUDA >= 9 * BlockScanInclusive for CUDA >= 9 * compatible_shfl_up for use in integral.hpp * Use CLAHE in cudev * Add tests for BlockScan original commit: opencv/opencv@970293a
1 parent 1c6b74d commit 18439b4

File tree

6 files changed

+385
-34
lines changed

6 files changed

+385
-34
lines changed

‎modules/cudaimgproc/src/cuda/clahe.cu

Lines changed: 9 additions & 15 deletions
Original file line numberDiff line numberDiff line change
@@ -42,15 +42,9 @@
4242

4343
#if !defined CUDA_DISABLER
4444

45-
#include "opencv2/core/cuda/common.hpp"
46-
#include "opencv2/core/cuda/functional.hpp"
47-
#include "opencv2/core/cuda/emulation.hpp"
48-
#include "opencv2/core/cuda/scan.hpp"
49-
#include "opencv2/core/cuda/reduce.hpp"
50-
#include "opencv2/core/cuda/saturate_cast.hpp"
45+
#include "opencv2/cudev.hpp"
5146

52-
using namespace cv::cuda;
53-
using namespace cv::cuda::device;
47+
using namespace cv::cudev;
5448

5549
namespace clahe
5650
{
@@ -73,7 +67,7 @@ namespace clahe
7367
for (int j = threadIdx.x; j < tileSize.x; j += blockDim.x)
7468
{
7569
const int data = srcPtr[j];
76-
Emulation::smem::atomicAdd(&smem[data], 1);
70+
::atomicAdd(&smem[data], 1);
7771
}
7872
}
7973

@@ -96,7 +90,7 @@ namespace clahe
9690

9791
// find number of overall clipped samples
9892

99-
reduce<256>(smem, clipped, tid, plus<int>());
93+
blockReduce<256>(smem, clipped, tid, plus<int>());
10094

10195
// broadcast evaluated value
10296

@@ -128,10 +122,10 @@ namespace clahe
128122

129123
calcLutKernel<<<grid, block, 0, stream>>>(src, lut, tileSize, tilesX, clipLimit, lutScale);
130124

131-
cudaSafeCall( cudaGetLastError() );
125+
CV_CUDEV_SAFE_CALL( cudaGetLastError() );
132126

133127
if (stream == 0)
134-
cudaSafeCall( cudaDeviceSynchronize() );
128+
CV_CUDEV_SAFE_CALL( cudaDeviceSynchronize() );
135129
}
136130

137131
__global__ void transformKernel(const PtrStepSzb src, PtrStepb dst, const PtrStepb lut, const int2 tileSize, const int tilesX, const int tilesY)
@@ -173,13 +167,13 @@ namespace clahe
173167
const dim3 block(32, 8);
174168
const dim3 grid(divUp(src.cols, block.x), divUp(src.rows, block.y));
175169

176-
cudaSafeCall( cudaFuncSetCacheConfig(transformKernel, cudaFuncCachePreferL1) );
170+
CV_CUDEV_SAFE_CALL( cudaFuncSetCacheConfig(transformKernel, cudaFuncCachePreferL1) );
177171

178172
transformKernel<<<grid, block, 0, stream>>>(src, dst, lut, tileSize, tilesX, tilesY);
179-
cudaSafeCall( cudaGetLastError() );
173+
CV_CUDEV_SAFE_CALL( cudaGetLastError() );
180174

181175
if (stream == 0)
182-
cudaSafeCall( cudaDeviceSynchronize() );
176+
CV_CUDEV_SAFE_CALL( cudaDeviceSynchronize() );
183177
}
184178
}
185179

‎modules/cudev/include/opencv2/cudev/block/scan.hpp

Lines changed: 141 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -48,12 +48,134 @@
4848

4949
#include "../common.hpp"
5050
#include "../warp/scan.hpp"
51+
#include "../warp/warp.hpp"
5152

5253
namespace cv { namespace cudev {
5354

5455
//! @addtogroup cudev
5556
//! @{
5657

58+
#if __CUDACC_VER_MAJOR__ >= 9
59+
60+
// Usage Note
61+
// - THREADS_NUM should be equal to the number of threads in this block.
62+
// - smem must be able to contain at least n elements of type T, where n is equal to the number
63+
// of warps in this block. The number can be calculated by divUp(THREADS_NUM, WARP_SIZE).
64+
//
65+
// Dev Note
66+
// - Starting from CUDA 9.0, support for Fermi is dropped. So CV_CUDEV_ARCH >= 300 is implied.
67+
// - "For Pascal and earlier architectures (CV_CUDEV_ARCH < 700), all threads in mask must execute
68+
// the same warp intrinsic instruction in convergence, and the union of all values in mask must
69+
// be equal to the warp's active mask."
70+
// (https://docs.nvidia.com/cuda/archive/10.0/cuda-c-programming-guide#independent-thread-scheduling-7-x)
71+
// - Above restriction does not apply starting from Volta (CV_CUDEV_ARCH >= 700). We just need to
72+
// take care so that "all non-exited threads named in mask must execute the same intrinsic with
73+
// the same mask."
74+
// (https://docs.nvidia.com/cuda/archive/10.0/cuda-c-programming-guide#warp-description)
75+
76+
template <int THREADS_NUM, typename T>
77+
__device__ T blockScanInclusive(T data, volatile T* smem, uint tid)
78+
{
79+
const int residual = THREADS_NUM & (WARP_SIZE - 1);
80+
81+
#if CV_CUDEV_ARCH < 700
82+
const uint residual_mask = (1U << residual) - 1;
83+
#endif
84+
85+
if (THREADS_NUM > WARP_SIZE)
86+
{
87+
// bottom-level inclusive warp scan
88+
#if CV_CUDEV_ARCH >= 700
89+
T warpResult = warpScanInclusive(0xFFFFFFFFU, data);
90+
#else
91+
T warpResult;
92+
93+
if (0 == residual)
94+
warpResult = warpScanInclusive(0xFFFFFFFFU, data);
95+
else
96+
{
97+
const int n_warps = divUp(THREADS_NUM, WARP_SIZE);
98+
const int warp_num = Warp::warpId();
99+
100+
if (warp_num < n_warps - 1)
101+
warpResult = warpScanInclusive(0xFFFFFFFFU, data);
102+
else
103+
{
104+
// We are at the last threads of a block whose number of threads
105+
// is not a multiple of the warp size
106+
warpResult = warpScanInclusive(residual_mask, data);
107+
}
108+
}
109+
#endif
110+
111+
__syncthreads();
112+
113+
// save top elements of each warp for exclusive warp scan
114+
// sync to wait for warp scans to complete (because smem is being overwritten)
115+
if ((tid & (WARP_SIZE - 1)) == (WARP_SIZE - 1))
116+
{
117+
smem[tid >> LOG_WARP_SIZE] = warpResult;
118+
}
119+
120+
__syncthreads();
121+
122+
int quot = THREADS_NUM / WARP_SIZE;
123+
124+
if (tid < quot)
125+
{
126+
// grab top warp elements
127+
T val = smem[tid];
128+
129+
uint mask = (1LLU << quot) - 1;
130+
131+
if (0 == residual)
132+
{
133+
// calculate exclusive scan and write back to shared memory
134+
smem[tid] = warpScanExclusive(mask, val);
135+
}
136+
else
137+
{
138+
// calculate inclusive scan and write back to shared memory with offset 1
139+
smem[tid + 1] = warpScanInclusive(mask, val);
140+
141+
if (tid == 0)
142+
smem[0] = 0;
143+
}
144+
}
145+
146+
__syncthreads();
147+
148+
// return updated warp scans
149+
return warpResult + smem[tid >> LOG_WARP_SIZE];
150+
}
151+
else
152+
{
153+
#if CV_CUDEV_ARCH >= 700
154+
return warpScanInclusive(0xFFFFFFFFU, data);
155+
#else
156+
if (THREADS_NUM == WARP_SIZE)
157+
return warpScanInclusive(0xFFFFFFFFU, data);
158+
else
159+
return warpScanInclusive(residual_mask, data);
160+
#endif
161+
}
162+
}
163+
164+
template <int THREADS_NUM, typename T>
165+
__device__ __forceinline__ T blockScanExclusive(T data, volatile T* smem, uint tid)
166+
{
167+
return blockScanInclusive<THREADS_NUM>(data, smem, tid) - data;
168+
}
169+
170+
#else // __CUDACC_VER_MAJOR__ >= 9
171+
172+
// Usage Note
173+
// - THREADS_NUM should be equal to the number of threads in this block.
174+
// - (>= Kepler) smem must be able to contain at least n elements of type T, where n is equal to the number
175+
// of warps in this block. The number can be calculated by divUp(THREADS_NUM, WARP_SIZE).
176+
// - (Fermi) smem must be able to contain at least n elements of type T, where n is equal to the number
177+
// of threads in this block (= THREADS_NUM).
178+
57179
template <int THREADS_NUM, typename T>
58180
__device__ T blockScanInclusive(T data, volatile T* smem, uint tid)
59181
{
@@ -73,18 +195,31 @@ __device__ T blockScanInclusive(T data, volatile T* smem, uint tid)
73195

74196
__syncthreads();
75197

76-
if (tid < (THREADS_NUM / WARP_SIZE))
198+
int quot = THREADS_NUM / WARP_SIZE;
199+
200+
if (tid < quot)
77201
{
78202
// grab top warp elements
79203
T val = smem[tid];
80204

81-
// calculate exclusive scan and write back to shared memory
82-
smem[tid] = warpScanExclusive(val, smem, tid);
205+
if (0 == (THREADS_NUM & (WARP_SIZE - 1)))
206+
{
207+
// calculate exclusive scan and write back to shared memory
208+
smem[tid] = warpScanExclusive(val, smem, tid);
209+
}
210+
else
211+
{
212+
// calculate inclusive scan and write back to shared memory with offset 1
213+
smem[tid + 1] = warpScanInclusive(val, smem, tid);
214+
215+
if (tid == 0)
216+
smem[0] = 0;
217+
}
83218
}
84219

85220
__syncthreads();
86221

87-
// return updated warp scans with exclusive scan results
222+
// return updated warp scans
88223
return warpResult + smem[tid >> LOG_WARP_SIZE];
89224
}
90225
else
@@ -99,6 +234,8 @@ __device__ __forceinline__ T blockScanExclusive(T data, volatile T* smem, uint t
99234
return blockScanInclusive<THREADS_NUM>(data, smem, tid) - data;
100235
}
101236

237+
#endif // __CUDACC_VER_MAJOR__ >= 9
238+
102239
//! @}
103240

104241
}}

‎modules/cudev/include/opencv2/cudev/grid/detail/integral.hpp

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -215,7 +215,7 @@ namespace integral_detail
215215
#pragma unroll
216216
for (int i = 1; i < 32; i *= 2)
217217
{
218-
const int n = shfl_up(sum, i, 32);
218+
const int n = compatible_shfl_up(sum, i, 32);
219219

220220
if (lane_id >= i)
221221
{
@@ -245,9 +245,9 @@ namespace integral_detail
245245
int warp_sum = sums[lane_id];
246246

247247
#pragma unroll
248-
for (int i = 1; i <= 32; i *= 2)
248+
for (int i = 1; i < 32; i *= 2)
249249
{
250-
const int n = shfl_up(warp_sum, i, 32);
250+
const int n = compatible_shfl_up(warp_sum, i, 32);
251251

252252
if (lane_id >= i)
253253
warp_sum += n;
@@ -453,7 +453,7 @@ namespace integral_detail
453453

454454
for (int i = 1; i <= 8; i *= 2)
455455
{
456-
T n = shfl_up(partial_sum, i, 32);
456+
T n = compatible_shfl_up(partial_sum, i, 32);
457457

458458
if (lane_id >= i)
459459
partial_sum += n;

‎modules/cudev/include/opencv2/cudev/warp/scan.hpp

Lines changed: 39 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -55,6 +55,36 @@ namespace cv { namespace cudev {
5555
//! @addtogroup cudev
5656
//! @{
5757

58+
#if __CUDACC_VER_MAJOR__ >= 9
59+
60+
// Starting from CUDA 9.0, support for Fermi is dropped.
61+
// So CV_CUDEV_ARCH >= 300 is implied.
62+
63+
template <typename T>
64+
__device__ T warpScanInclusive(uint mask, T data)
65+
{
66+
const uint laneId = Warp::laneId();
67+
68+
// scan on shufl functions
69+
#pragma unroll
70+
for (int i = 1; i <= (WARP_SIZE / 2); i *= 2)
71+
{
72+
const T val = shfl_up_sync(mask, data, i);
73+
if (laneId >= i)
74+
data += val;
75+
}
76+
77+
return data;
78+
}
79+
80+
template <typename T>
81+
__device__ __forceinline__ T warpScanExclusive(uint mask, T data)
82+
{
83+
return warpScanInclusive(mask, data) - data;
84+
}
85+
86+
#else // __CUDACC_VER_MAJOR__ >= 9
87+
5888
template <typename T>
5989
__device__ T warpScanInclusive(T data, volatile T* smem, uint tid)
6090
{
@@ -75,19 +105,16 @@ __device__ T warpScanInclusive(T data, volatile T* smem, uint tid)
75105

76106
return data;
77107
#else
78-
uint pos = 2 * tid - (tid & (WARP_SIZE - 1));
79-
smem[pos] = 0;
108+
const uint laneId = Warp::laneId();
80109

81-
pos += WARP_SIZE;
82-
smem[pos] = data;
110+
smem[tid] = data;
83111

84-
smem[pos] += smem[pos - 1];
85-
smem[pos] += smem[pos - 2];
86-
smem[pos] += smem[pos - 4];
87-
smem[pos] += smem[pos - 8];
88-
smem[pos] += smem[pos - 16];
112+
#pragma unroll
113+
for (int i = 1; i <= (WARP_SIZE / 2); i *= 2)
114+
if (laneId >= i)
115+
smem[tid] += smem[tid - i];
89116

90-
return smem[pos];
117+
return smem[tid];
91118
#endif
92119
}
93120

@@ -97,6 +124,8 @@ __device__ __forceinline__ T warpScanExclusive(T data, volatile T* smem, uint ti
97124
return warpScanInclusive(data, smem, tid) - data;
98125
}
99126

127+
#endif // __CUDACC_VER_MAJOR__ >= 9
128+
100129
//! @}
101130

102131
}}

0 commit comments

Comments
 (0)