 
             int tid = threadIdx.x;
 
-            const T* selected_disparity = selected_disp_pyr + y_out/2 * cmsg_step + x_out/2;
-            T* data_cost = data_cost_ + y_out * cmsg_step + x_out;
+            const T* selected_disparity = selected_disp_pyr + y_out/2 * msg_step + x_out/2;
+            T* data_cost = data_cost_ + y_out * msg_step + x_out;
 
             if (d < nr_plane)
             {
-                int sel_disp = selected_disparity[d * cdisp_step2];
+                int sel_disp = selected_disparity[d * disp_step2];
 
                 int x0 = x_out << level;
                 int y0 = y_out << level;

                int x0 = x << level;
                int xt = (x + 1) << level;

                const T* selected_disparity = selected_disp_pyr + y/2 * cmsg_step + x/2;
                T* data_cost = data_cost_ + y * cmsg_step + x;

                for(int d = 0; d < nr_plane; d++)
                {

                            }
                        }
                    }
                    data_cost[disp_step1 * d] = saturate_cast<T>(val);
                }
            }
        }

        template <typename T, int winsz, int channels>
        __global__ void compute_data_cost_reduce(const uchar *cleft, const uchar *cright, size_t cimg_step, const T* selected_disp_pyr, T* data_cost_, int level, int rows, int cols, int h, int nr_plane, float data_weight, float max_data_term, int min_disp, size_t msg_step, size_t disp_step1, size_t disp_step2)
        {
            int x_out = blockIdx.x;
            int y_out = blockIdx.y % h;
