From 76e8794e81188a640bcc366527d6a2da1cffe072 Mon Sep 17 00:00:00 2001
From: Vladislav Vinogradov <vlad.vinogradov@itseez.com>
Date: Thu, 15 Nov 2012 12:00:38 +0400
Subject: [PATCH] fixed warnings on old compute capabilities

---
 modules/gpu/include/opencv2/gpu/device/detail/reduce.hpp    | 6 ++++++
 .../include/opencv2/gpu/device/detail/reduce_key_val.hpp    | 6 ++++++
 modules/gpu/src/cuda/surf.cu                                | 2 ++
 3 files changed, 14 insertions(+)

diff --git a/modules/gpu/include/opencv2/gpu/device/detail/reduce.hpp b/modules/gpu/include/opencv2/gpu/device/detail/reduce.hpp
index 628129ea3..2b2ba6773 100644
--- 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 f7531da24..f1aa285d3 100644
--- 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 c12192555..73fcbc3d1 100644
--- 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;