From ea3fe441df5e874b229f7fafa2bc8042f97b97dc Mon Sep 17 00:00:00 2001
From: Marina Kolpakova <no@email>
Date: Mon, 30 Apr 2012 19:06:37 +0000
Subject: [PATCH] fix for brute-force matcher with CUDA 4.2.

---
 modules/gpu/src/cuda/bf_match.cu | 6 +++---
 1 file changed, 3 insertions(+), 3 deletions(-)

diff --git a/modules/gpu/src/cuda/bf_match.cu b/modules/gpu/src/cuda/bf_match.cu
index 7d6d62b530..ad8cc88179 100644
--- a/modules/gpu/src/cuda/bf_match.cu
+++ b/modules/gpu/src/cuda/bf_match.cu
@@ -97,7 +97,7 @@ namespace cv { namespace gpu { namespace device
         }
 
         template <int BLOCK_SIZE, int MAX_DESC_LEN, typename Dist, typename T, typename Mask> 
-        __device__ void loopUnrolledCached(int queryIdx, const DevMem2D_<T>& query, int imgIdx, const DevMem2D_<T>& train, const Mask& mask, 
+        __device__ void loopUnrolledCached(int queryIdx, const DevMem2D_<T>& query,volatile int imgIdx, const DevMem2D_<T>& train, const Mask& mask,
                                            typename Dist::value_type* s_query, typename Dist::value_type* s_train, 
                                            float& bestDistance, int& bestTrainIdx, int& bestImgIdx)
         {
@@ -253,7 +253,7 @@ namespace cv { namespace gpu { namespace device
         // Match Unrolled
 
         template <int BLOCK_SIZE, int MAX_DESC_LEN, typename Dist, typename T, typename Mask> 
-        __device__ void loopUnrolled(int queryIdx, const DevMem2D_<T>& query, int imgIdx, const DevMem2D_<T>& train, const Mask& mask, 
+        __device__ void loopUnrolled(int queryIdx, const DevMem2D_<T>& query,volatile int imgIdx, const DevMem2D_<T>& train, const Mask& mask,
                                      typename Dist::value_type* s_query, typename Dist::value_type* s_train, 
                                      float& bestDistance, int& bestTrainIdx, int& bestImgIdx)
         {
@@ -409,7 +409,7 @@ namespace cv { namespace gpu { namespace device
         // Match
 
         template <int BLOCK_SIZE, typename Dist, typename T, typename Mask> 
-        __device__ void loop(int queryIdx, const DevMem2D_<T>& query, int imgIdx, const DevMem2D_<T>& train, const Mask& mask, 
+        __device__ void loop(int queryIdx, const DevMem2D_<T>& query, volatile int imgIdx, const DevMem2D_<T>& train, const Mask& mask,
                              typename Dist::value_type* s_query, typename Dist::value_type* s_train, 
                              float& bestDistance, int& bestTrainIdx, int& bestImgIdx)
         {