generic-library/opencv
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

fixed warnings on old compute capabilities

Browse Source
This commit is contained in:
Vladislav Vinogradov 2012-11-15 12:00:38 +04:00
parent e6b1ccdcdf
commit 76e8794e81
3 changed files with 14 additions and 0 deletions
modules/gpu
include/opencv2/gpu/device/detail
reduce.hppreduce_key_val.hpp
src/cuda
surf.cu

6
modules/gpu/include/opencv2/gpu/device/detail/reduce.hpp

@ -260,7 +260,9 @@ namespace cv { namespace gpu { namespace device
if (tid < N / 2)
{
#if __CUDA_ARCH__ >= 200
#pragma unroll
#endif
for (unsigned int i = N / 2; i >= 1; i /= 2)
merge(smem, val, tid, i, op);
}
@ -289,7 +291,9 @@ namespace cv { namespace gpu { namespace device
if (laneId < 16)
{
#if __CUDA_ARCH__ >= 200
#pragma unroll
#endif
for (int i = 16; i >= 1; i /= 2)
merge(smem, val, tid, i, op);
}
@ -311,7 +315,9 @@ namespace cv { namespace gpu { namespace device
for (int i = M / 2; i >= 1; i /= 2)
mergeShfl(val, i, M, op);
#else
#if __CUDA_ARCH__ >= 200
#pragma unroll
#endif
for (int i = M / 2; i >= 1; i /= 2)
merge(smem, val, tid, i, op);
#endif

6
modules/gpu/include/opencv2/gpu/device/detail/reduce_key_val.hpp

@ -388,7 +388,9 @@ namespace cv { namespace gpu { namespace device
if (tid < N / 2)
{
#if __CUDA_ARCH__ >= 200
#pragma unroll
#endif
for (unsigned int i = N / 2; i >= 1; i /= 2)
merge(skeys, key, svals, val, cmp, tid, i);
}
@ -421,7 +423,9 @@ namespace cv { namespace gpu { namespace device
if (laneId < 16)
{
#if __CUDA_ARCH__ >= 200
#pragma unroll
#endif
for (int i = 16; i >= 1; i /= 2)
merge(skeys, key, svals, val, cmp, tid, i);
}
@ -448,7 +452,9 @@ namespace cv { namespace gpu { namespace device
for (unsigned int i = M / 2; i >= 1; i /= 2)
mergeShfl(key, val, cml, i, M);
#else
#if __CUDA_ARCH__ >= 200
#pragma unroll
#endif
for (unsigned int i = M / 2; i >= 1; i /= 2)
merge(skeys, key, svals, val, cmp, tid, i);
#endif

2
modules/gpu/src/cuda/surf.cu

@ -568,7 +568,9 @@ namespace cv { namespace gpu { namespace device
float bestx = 0, besty = 0, best_mod = 0;
#if __CUDA_ARCH__ >= 200
#pragma unroll
#endif
for (int i = 0; i < 18; ++i)
{
const int dir = (i * 4 + threadIdx.y) * ORI_SEARCH_INC;