   cudaUnbindTexture(texRef);
   CUT_CHECK_ERROR("Memory unbind failed");
 
-  /// do not remove... ?
-  int *tt;
-  cudaMalloc((void**) &tt, (blocksPerGrid*sizeof(int)));
-
-  kernel_hysteresis_glm1D<<<DimGrid,DimBlock>>>(d_hys, size);
+  int *gridUpdate;
+  cudaMalloc((void**) &gridUpdate, (blocksPerGrid*sizeof(int)));
+
+  kernel_hysteresis_glm1D<<<DimGrid,DimBlock>>>(d_hys, size, gridUpdate);
   CUT_CHECK_ERROR("Hysteresis Kernel failed");
 
   /// bind a texture to the CUDA array


}

__global__ void kernel_hysteresis_glm1D(float *hys_img, int3 size){

  __shared__ float s_slice[SLICE_WIDTH*SLICE_WIDTH];
  __shared__ int modified_block_pixels; /// control inner loop
  __shared__ int modified_image_pixels; /// control outer loop
  float edge;

  // pixel position indexes on slice
  int2 slice_pos;


    }while(modified_block_pixels);// end inner loop

    if (!threadIdx.x) gridUpdate[blockIdx.x]++;
    gU = gridUpdate[blockIdx.x];

    if (modified_image_pixels) hys_img[pixIdx] = s_slice[sliceIdx];

    // barrier
    while (! ((blockIdx.x<gridWidth)+\
             ((((blockIdx.x%gridWidth) == 0)+(gU == gridUpdate[blockIdx.x-gridWidth-1])) *\
             (gU == gridUpdate[blockIdx.x-gridWidth]) *\
             (((blockIdx.x%gridWidth)==(gridWidth-1))+(gU == gridUpdate[blockIdx.x-gridWidth+1])))) *\
             (((blockIdx.x%gridWidth)==0)+(gU == gridUpdate[blockIdx.x-1])) *\
             (((blockIdx.x%gridWidth)==(gridWidth-1))+(gU == gridUpdate[blockIdx.x-1])) *\
             ((blockIdx.x>=(gridDim.x-gridWidth))+\
             ((((blockIdx.x%gridWidth) == 0)+(gU == gridUpdate[blockIdx.x+gridWidth-1])) *\
             (gU == gridUpdate[blockIdx.x+gridWidth]) *\
             (((blockIdx.x%gridWidth)==(gridWidth-1))+(gU == gridUpdate[blockIdx.x+gridWidth+1])))) ){

    }
    

  }while(modified_image_pixels);//end outer loop

