fixed warnings on old compute capabilities

diff --git a/modules/gpu/include/opencv2/gpu/device/detail/reduce.hpp b/modules/gpu/include/opencv2/gpu/device/detail/reduce.hpp
index 628129e..2b2ba67 100644 (file)
--- a/modules/gpu/include/opencv2/gpu/device/detail/reduce.hpp
+++ b/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

diff --git a/modules/gpu/include/opencv2/gpu/device/detail/reduce_key_val.hpp b/modules/gpu/include/opencv2/gpu/device/detail/reduce_key_val.hpp
index f7531da..f1aa285 100644 (file)
--- a/modules/gpu/include/opencv2/gpu/device/detail/reduce_key_val.hpp
+++ b/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

diff --git a/modules/gpu/src/cuda/surf.cu b/modules/gpu/src/cuda/surf.cu
index c121925..73fcbc3 100644 (file)
--- a/modules/gpu/src/cuda/surf.cu
+++ b/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;