fix for brute-force matcher with CUDA 4.2.

Marina Kolpakova · Marina Kolpakova · commit ea3fe441df5e · 2012-04-30T19:06:37.000Z
diff --git 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)
         {

Original file line number	Diff line number	Diff line change
`@@ -97,7 +97,7 @@ namespace cv { namespace gpu { namespace device`
`97`	`97`	`}`
`98`	`98`
`99`	`99`	`template <int BLOCK_SIZE, int MAX_DESC_LEN, typename Dist, typename T, typename Mask>`
`100`		`- __device__ void loopUnrolledCached(int queryIdx, const DevMem2D_<T>& query, int imgIdx, const DevMem2D_<T>& train, const Mask& mask,`
	`100`	`+ __device__ void loopUnrolledCached(int queryIdx, const DevMem2D_<T>& query,volatile int imgIdx, const DevMem2D_<T>& train, const Mask& mask,`
`101`	`101`	`typename Dist::value_type* s_query, typename Dist::value_type* s_train,`
`102`	`102`	`float& bestDistance, int& bestTrainIdx, int& bestImgIdx)`
`103`	`103`	`{`
`@@ -253,7 +253,7 @@ namespace cv { namespace gpu { namespace device`
`253`	`253`	`// Match Unrolled`
`254`	`254`
`255`	`255`	`template <int BLOCK_SIZE, int MAX_DESC_LEN, typename Dist, typename T, typename Mask>`
`256`		`- __device__ void loopUnrolled(int queryIdx, const DevMem2D_<T>& query, int imgIdx, const DevMem2D_<T>& train, const Mask& mask,`
	`256`	`+ __device__ void loopUnrolled(int queryIdx, const DevMem2D_<T>& query,volatile int imgIdx, const DevMem2D_<T>& train, const Mask& mask,`
`257`	`257`	`typename Dist::value_type* s_query, typename Dist::value_type* s_train,`
`258`	`258`	`float& bestDistance, int& bestTrainIdx, int& bestImgIdx)`
`259`	`259`	`{`
`@@ -409,7 +409,7 @@ namespace cv { namespace gpu { namespace device`
`409`	`409`	`// Match`
`410`	`410`
`411`	`411`	`template <int BLOCK_SIZE, typename Dist, typename T, typename Mask>`
`412`		`- __device__ void loop(int queryIdx, const DevMem2D_<T>& query, int imgIdx, const DevMem2D_<T>& train, const Mask& mask,`
	`412`	`+ __device__ void loop(int queryIdx, const DevMem2D_<T>& query, volatile int imgIdx, const DevMem2D_<T>& train, const Mask& mask,`
`413`	`413`	`typename Dist::value_type* s_query, typename Dist::value_type* s_train,`
`414`	`414`	`float& bestDistance, int& bestTrainIdx, int& bestImgIdx)`
`415`	`415`	`{`