minor

diff --git a/modules/gpu/src/cuda/stereocsbp.cu b/modules/gpu/src/cuda/stereocsbp.cu
index 0a4435c..d9222c3 100644 (file)
--- a/modules/gpu/src/cuda/stereocsbp.cu
+++ b/modules/gpu/src/cuda/stereocsbp.cu
@@ -644,8 +644,8 @@ namespace cv { namespace gpu { namespace device
             {\r
                 const T* u_cur = u_cur_ + ::min(h2-1, y/2 + 1) * cmsg_step2 + x/2;\r
                 const T* d_cur = d_cur_ + ::max(0, y/2 - 1)    * cmsg_step2 + x/2;\r
-                const T* l_cur = l_cur_ + y/2                  * cmsg_step2 + ::min(w2-1, x/2 + 1);\r
-                const T* r_cur = r_cur_ + y/2                  * cmsg_step2 + ::max(0, x/2 - 1);\r
+                const T* l_cur = l_cur_ + (y/2)                * cmsg_step2 + ::min(w2-1, x/2 + 1);\r
+                const T* r_cur = r_cur_ + (y/2)                * cmsg_step2 + ::max(0, x/2 - 1);\r
 \r
                 T* data_cost_new = (T*)ctemp + y * cmsg_step1 + x;\r
 \r
@@ -731,7 +731,7 @@ namespace cv { namespace gpu { namespace device
 \r
         template <typename T>\r
         __device__ void message_per_pixel(const T* data, T* msg_dst, const T* msg1, const T* msg2, const T* msg3,\r
-                                          const T* dst_disp, const T* src_disp, int nr_plane, T* temp)\r
+                                          const T* dst_disp, const T* src_disp, int nr_plane, volatile T* temp)\r
         {\r
             T minimum = numeric_limits<T>::max();\r
 \r
@@ -808,11 +808,10 @@ namespace cv { namespace gpu { namespace device
             for(int t = 0; t < iters; ++t)\r
             {\r
                 compute_message<<<grid, threads, 0, stream>>>(u, d, l, r, data_cost_selected, selected_disp_pyr_cur, h, w, nr_plane, t & 1);\r
-                cudaSafeCall( cudaGetLastError() );\r
-\r
-                if (stream == 0)\r
-                    cudaSafeCall( cudaDeviceSynchronize() );\r
+                cudaSafeCall( cudaGetLastError() );                \r
             }\r
+                       if (stream == 0)\r
+                    cudaSafeCall( cudaDeviceSynchronize() );\r
         };\r
 \r
         template void calc_all_iterations(short* u, short* d, short* l, short* r, const short* data_cost_selected, const short* selected_disp_pyr_cur, size_t msg_step,\r
@@ -830,12 +829,12 @@ namespace cv { namespace gpu { namespace device
         template <typename T>\r
         __global__ void compute_disp(const T* u_, const T* d_, const T* l_, const T* r_,\r
                                      const T* data_cost_selected, const T* disp_selected_pyr,\r
-                                     short* disp, size_t res_step, int cols, int rows, int nr_plane)\r
+                                     PtrStepSz<short> disp, int nr_plane)\r
         {\r
             int x = blockIdx.x * blockDim.x + threadIdx.x;\r
             int y = blockIdx.y * blockDim.y + threadIdx.y;\r
 \r
-            if (y > 0 && y < rows - 1 && x > 0 && x < cols - 1)\r
+            if (y > 0 && y < disp.rows - 1 && x > 0 && x < disp.cols - 1)\r
             {\r
                 const T* data = data_cost_selected + y * cmsg_step1 + x;\r
                 const T* disp_selected = disp_selected_pyr + y * cmsg_step1 + x;\r
@@ -858,7 +857,7 @@ namespace cv { namespace gpu { namespace device
                         best = saturate_cast<short>(disp_selected[idx]);\r
                     }\r
                 }\r
-                disp[res_step * y + x] = best;\r
+                               disp(y, x) = best;\r
             }\r
         }\r
 \r
@@ -876,8 +875,7 @@ namespace cv { namespace gpu { namespace device
             grid.x = divUp(disp.cols, threads.x);\r
             grid.y = divUp(disp.rows, threads.y);\r
 \r
-            compute_disp<<<grid, threads, 0, stream>>>(u, d, l, r, data_cost_selected, disp_selected,\r
-                                                       disp.data, disp.step / disp.elemSize(), disp.cols, disp.rows, nr_plane);\r
+            compute_disp<<<grid, threads, 0, stream>>>(u, d, l, r, data_cost_selected, disp_selected, disp, nr_plane);\r
             cudaSafeCall( cudaGetLastError() );\r
 \r
             if (stream == 0)\r