 
 }
 
-__device__ int reduceSum(int tid, int data[], int N, int pow2){
-
-  if (tid<N){
-
-    for(unsigned int s=pow2/2; s>0; s>>=1) {
-      if ((tid < s) && (tid + s < pow2)){
-          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 block_tid = blockDim.x * threadIdx.y + threadIdx.x;
-  int tid = block_tid + (blockIdx.x*blockDim.x) + (blockIdx.y*blockDim.x*blockDim.y*gridDim.x);
-  
-  ///pixel index of this thread
-  int2 pos;
-  pos.x = blockIdx.x * blockDim.x + threadIdx.x;
-  pos.y = blockIdx.y * blockDim.y + threadIdx.y;
-  int pixIdx = pos.y * size.x + pos.x;
-  int2 sliceIdx;
-  sliceIdx.x = threadIdx.x+1;
-  sliceIdx.y = threadIdx.y+1;
-  int edge;
-  int i = 0;
-
-  /// load center
-  s_slice[sliceIdx.y][sliceIdx.x] = hys_img[pixIdx];
-
-  if (!block_tid) m = 1;
-
-  do{
-
-    if(m){
-      /// load top
-      if(!threadIdx.y){
-        if(!threadIdx.x){
-          s_slice[0][0] = ((pos.x>0)||(pos.y>0)) * hys_img[pixIdx-size.x-1];///<TL
-        }
-        s_slice[0][sliceIdx.x] = (pos.y>0) * hys_img[pixIdx-size.x];
-        if(threadIdx.x == (blockDim.x-1)){
-           s_slice[0][blockDim.x+1] = ((pos.x<(size.x-1))&&(pos.y>0)) * hys_img[pixIdx-size.x+1];///<TR
-        }
-      }
-      /// load bottom
-      if(threadIdx.y == (blockDim.y-1)){
-        if(!threadIdx.x){
-          s_slice[blockDim.y+1][0] = ((pos.x>0)&&(pos.y<(size.y-1))) * hys_img[pixIdx+size.x-1];///<BL
-        }
-        s_slice[blockDim.y+1][sliceIdx.x] = (pos.y<(size.y-1)) * hys_img[pixIdx+size.x];
-        if(threadIdx.x == (blockDim.x-1)){
-          s_slice[blockDim.y+1][blockDim.x+1] = ((pos.x<(size.x-1))&&(pos.y<(size.y-1))) * hys_img[pixIdx+size.x+1];///<BR
-        }
-      }
-      /// load left
-      if(!threadIdx.x){
-        s_slice[sliceIdx.y][0] = (pos.x>0) * hys_img[pixIdx-1];
-      }
-      /// load right
-      if(threadIdx.x == blockDim.x-1){
-        s_slice[sliceIdx.y][blockDim.x+1] = (pos.x<(size.x-1)) * hys_img[pixIdx+1];
-      }
-    }
-
-    if (!block_tid) m = 0;
-
-    for(i=0;i<1;i++){
-//    while(reduceSum256(block_tid, s_modified)){
-
-      s_modified[block_tid] = 0;
-
-      if(s_slice[sliceIdx.y][sliceIdx.x] == 128){
-
-        __syncthreads();
-
-        /// edge == 1 if at last one pixel's neighbour is a definitive edge 
-        /// and edge == 0 if doesn't
-        edge = (!(s_slice[sliceIdx.y-1][sliceIdx.x-1] != 255) *\
-                 (s_slice[sliceIdx.y-1][sliceIdx.x] != 255) *\
-                 (s_slice[sliceIdx.y-1][sliceIdx.x+1] != 255) *\
-                 (s_slice[sliceIdx.y][sliceIdx.x-1] != 255) *\
-                 (s_slice[sliceIdx.y][sliceIdx.x+1] != 255) *\
-                 (s_slice[sliceIdx.y+1][sliceIdx.x-1] != 255) *\
-                 (s_slice[sliceIdx.y+1][sliceIdx.x] != 255) *\
-                 (s_slice[sliceIdx.y+1][sliceIdx.x+1] != 255));
-
-        s_modified[block_tid] = (edge);
-        s_slice[sliceIdx.y][sliceIdx.x] = (float) 128 + (edge)*127;
-
-      }
-
-      /// this means that the inner loop was executed for this block.
-      /// it works because all threads of a block execute the inner loop if at
-      /// last one pixel is modified on it's image slice.
-      if (!block_tid) m = 1;
-
-      __syncthreads();
-
-      if (!reduceSum256(block_tid,s_modified)) break;
-      
-    }
-
-    if (!block_tid) modified[blockIdx.y*gridDim.x + blockIdx.x] = m;
-
-    if (m){
-      if((pos.x < (size.x)) && (pos.y < (size.y))){
-        /// store center
-        hys_img[pixIdx] = s_slice[sliceIdx.x][sliceIdx.y];
-      }
-    }
-
-    __syncthreads();
-
-//  if (tid<pow2){
-
-//    for(unsigned int s=pow2/2; s>0; s>>=1) {
-//      if ((tid < s) && ((tid+s)<N)){
-//          modified[tid] += modified[tid+s];
-//      }
-//      __syncthreads();
-//    }
-
-//  }
-
-
-if (!block_tid) reduceSum(blockIdx.y*gridDim.x + blockIdx.x,modified,N,pow2);// this function is doing wrong!
-
-//  }while(reduceSum(tid,modified,N,pow2));
-  }while(0);
-
-  if((pos.x < (size.x)) && (pos.y < (size.y))){ 
-    hys_img[pixIdx] = s_slice[sliceIdx.y][sliceIdx.x];
-  }
-
-}
-
-__global__ void kernel_hysteresis_glm1D(int *hys_img, int3 size, int *modified){
+__global__ void kernel_hysteresis_glm1D(float *hys_img, int3 size, int *modified){
 
   __shared__ float s_slice[324];
   __shared__ int s_modified[256];
-  int edge;
+  float edge;
   int i;
 
   // thread index

  cutCreateTimer( &timer );
  cutStartTimer( timer );  /// Start timer

  if (maxKernelWidth < 1) maxKernelWidth = 1;
  if (maxKernelWidth%2 == 0) maxKernelWidth--;
  short halfkernelsize = maxKernelWidth >> 1;

///calculate gaussian mask
  float *cudaGaussKernel;
  cudaMalloc((void**)&cudaGaussKernel,kernelSize);

  calculateGaussianKernel<<<1,maxKernelWidth,kernelSize>>>(cudaGaussKernel, gaussianVariance, maxKernelWidth);

///use calculated mask to gaussian filter

  ///allocate temporary memory to image data
  float *d_tmp;
  cudaMalloc((void**) &d_tmp, size.z*sizeof(float));

  /// bind a texture to the CUDA array
  cudaBindTexture (NULL, gaussTexRef, cudaGaussKernel);
  CUT_CHECK_ERROR("Texture bind failed");

  /// host side settable texture attributes
  gaussTexRef.normalized = false;
  gaussTexRef.filterMode = cudaFilterModePoint;

  /// bind a texture to the CUDA array
  cudaBindTexture (NULL, texRef, d_img);
  CUT_CHECK_ERROR("Texture bind failed");

  /// host side settable texture attributes
  texRef.normalized = false;
  texRef.filterMode = cudaFilterModePoint;

  kernel_1DConvolutionH_texture<<<DimGrid,DimBlock,kernelSize>>>(d_tmp,size,halfkernelsize);

  ///bind temporary data texture
  cudaUnbindTexture(texRef);
  cudaBindTexture (NULL ,texRef, d_tmp);

  kernel_1DConvolutionV_texture<<<DimGrid,DimBlock,kernelSize>>>(d_img,size,halfkernelsize);

  ///free allocated memory
  cudaFree(d_tmp);
  cudaUnbindTexture(texRef);
  cudaFree(cudaGaussKernel);

               (s_slice[sliceIdx+18] != 255) *\
               (s_slice[sliceIdx+19] != 255));
      s_modified[threadIdx.x] = (edge);
      s_slice[sliceIdx] = 128 + (edge)*127;

    }

