 
 }
 
-__global__ void kernel_hysteresis_glm2(int *hys_img, int3 size){
+__device__ int reduceSum(int tid, int data[], int N, int pow2){
+
+  if (tid<pow2){
+
+    if (tid>=N) data[tid] = 0;
+
+    for(unsigned int s=pow2/2; s>0; s>>=1) {
+      if (tid < s){
+      data[tid] += data[tid + s];
+    }
+      __syncthreads();
+    }
+
+  }
+  return(data[0]);
+
+}
+
+
+__global__ void kernel_hysteresis_glm3(int *hys_img, int3 size, int *modified, int N, int pow2){
 
   __shared__ float s_slice[18][18];
   __shared__ int s_modified[256];
   __shared__ int m;
 
-  int tid = blockDim.x * threadIdx.y + threadIdx.x;
+  int block_tid = blockDim.x * threadIdx.y + threadIdx.x;
+  int tid = block_tid + (blockIdx.x*blockDim.x*blockDim.y)+(blockIdx.y*blockDim.x*blockDim.y*gridDim.x);
   
   ///pixel index of this thread
   int2 pos;

  CUT_CHECK_ERROR("Memory unbind failed");

  int *d_modified;
  cudaMalloc((void**) &d_modified, (nBlocks*sizeof(int)));
  CUT_CHECK_ERROR("Memory hysteresis image creation failed");

  kernel_hysteresis_glm2<<<TwoDimGrid,TwoDimBlock>>>(d_hys, size);
  CUT_CHECK_ERROR("Hysteresis Kernel failed");

  /// bind a texture to the CUDA array


}

__device__ int reduceSum256(int tid, int data[]){
// data[] should be a 256 positios array

  if (tid < 128) { data[tid] += data[tid + 128]; }
  __syncthreads();
