   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;


  // pixel position at the hysteresis image
  int pixIdx = pos.y * size.x + pos.x;

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

  // load center
  s_slice[sliceIdx] = hys_img[pixIdx];
