From aabc33c7724fa7ebc0546012efa4fafc8fc1278d Mon Sep 17 00:00:00 2001
From: Vladislav Vinogradov <vlad.vinogradov@itseez.com>
Date: Wed, 5 Dec 2012 10:36:54 +0400
Subject: [PATCH] temporarily disabled reduceKeyVal implementation for kepler

used fermi implementation
---
 .../include/opencv2/gpu/device/detail/reduce_key_val.hpp    | 6 +++---
 1 file changed, 3 insertions(+), 3 deletions(-)

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 ca2c43127..a84e0c2fd 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
@@ -399,7 +399,7 @@ namespace cv { namespace gpu { namespace device
             template <class KP, class KR, class VP, class VR, class Cmp>
             static __device__ void reduce(KP skeys, KR key, VP svals, VR val, unsigned int tid, Cmp cmp)
             {
-            #if __CUDA_ARCH__ >= 300
+            #if 0 // __CUDA_ARCH__ >= 300
                 (void) skeys;
                 (void) svals;
                 (void) tid;
@@ -424,7 +424,7 @@ namespace cv { namespace gpu { namespace device
             {
                 const unsigned int laneId = Warp::laneId();
 
-            #if __CUDA_ARCH__ >= 300
+            #if 0 // __CUDA_ARCH__ >= 300
                 Unroll<16, KP, KR, VP, VR, Cmp>::loopShfl(key, val, cmp, warpSize);
 
                 if (laneId == 0)
@@ -454,7 +454,7 @@ namespace cv { namespace gpu { namespace device
 
                 if (tid < 32)
                 {
-                #if __CUDA_ARCH__ >= 300
+                #if 0 // __CUDA_ARCH__ >= 300
                     loadFromSmem(svals, val, tid);
 
                     Unroll<M / 2, KP, KR, VP, VR, Cmp>::loopShfl(key, val, cmp, M);