 __global__ void stereoKernel(unsigned char *left, unsigned char *right, size_t img_step, unsigned char* disp, size_t disp_pitch, int maxdisp)
 {
     extern __shared__ unsigned int col_ssd_cache[];
-    unsigned int *col_ssd = col_ssd_cache + BLOCK_W + threadIdx.x;
-    unsigned int *col_ssd_extra = threadIdx.x < (2 * RADIUS) ? col_ssd + BLOCK_W : 0;  //#define N_DIRTY_PIXELS (2 * RADIUS)
+    volatile unsigned int *col_ssd = col_ssd_cache + BLOCK_W + threadIdx.x;
+    volatile unsigned int *col_ssd_extra = threadIdx.x < (2 * RADIUS) ? col_ssd + BLOCK_W : 0;  //#define N_DIRTY_PIXELS (2 * RADIUS)
 
     //#define X (blockIdx.x * BLOCK_W + threadIdx.x + STEREO_MAXD)
-    int X = (blockIdx.x * BLOCK_W + threadIdx.x + maxdisp);
+    int X = (blockIdx.x * BLOCK_W + threadIdx.x + maxdisp + RADIUS);
     //#define Y (__mul24(blockIdx.y, ROWSperTHREAD) + RADIUS)
     #define Y (blockIdx.y * ROWSperTHREAD + RADIUS)
     //int Y = blockIdx.y * ROWSperTHREAD + RADIUS;

}

template<int RADIUS>
__device__ unsigned int CalcSSD(unsigned int *col_ssd_cache, unsigned int *col_ssd)
{
    unsigned int cache = 0;
    unsigned int cache2 = 0;


}

template<int RADIUS>
__device__ void InitColSSD(int x_tex, int y_tex, int im_pitch, unsigned char* imageL, unsigned char* imageR, int d, volatile unsigned int *col_ssd)
{
    unsigned char leftPixel1;
    int idx;
