   cudaMalloc((void**) &d_modified, (pow2*sizeof(int)));
   CUT_CHECK_ERROR("Memory hysteresis image creation failed");
 
-  kernel_hysteresis_glm3<<<TwoDimGrid,TwoDimBlock>>>(d_hys, size, d_modified, nBlocks, pow2);
+  int *d_modifsum;
+  cudaMalloc((void**) &d_modifsum, (nBlocks*sizeof(int)));
+  CUT_CHECK_ERROR("Memory hysteresis image creation failed");
+
+  CUDPPConfiguration config;
+  config.op = CUDPP_ADD;
+  config.datatype = CUDPP_INT;
+  config.algorithm = CUDPP_SCAN;
+  config.options = CUDPP_OPTION_BACKWARD | CUDPP_OPTION_INCLUSIVE;
+
+  CUDPPHandle scanplan = 0;
+  CUDPPResult result = cudppPlan(&scanplan, config, nBlocks, 1, 0);  
+
+  if (CUDPP_SUCCESS != result){
+    printf("Error creating CUDPPPlan\n");
+    exit(-1);
+  }
+
+  int a;
+  int cont[1];
+  for(a = 0; a<100; a++){
+
+    kernel_hysteresis_glm1D<<<DimGrid,DimBlock>>>(d_hys, size, d_modified);
+
+    cudppScan(scanplan, d_modifsum, d_modified, nBlocks);
+    cudaMemcpy(cont,d_modifsum,sizeof(int),cudaMemcpyDeviceToHost);
+
+    if (!cont[0]) break;
+
+  }
   CUT_CHECK_ERROR("Hysteresis Kernel failed");
+  //printf("Histeresis Interations: %d\n",a+1);
+
+  result = cudppDestroyPlan(scanplan);
+  if (CUDPP_SUCCESS != result){
+    printf("Error destroying CUDPPPlan\n");
+    exit(-1);
+  }
+
+/*********************** Teste ************************
+  int i;
+  int img[pow2];
+  cudaMemcpy(img,d_modified,pow2*sizeof(int),cudaMemcpyDeviceToHost);
+  for(i=0; i<pow2;i++){
+      printf("%d ",img[i]);
+  }
+  printf("\n");
+****************** Fim do Teste *********************/
+
+//printf("nBlocks: %d\n", nBlocks);
+//printf("pow2: %d\n", pow2);
+//printf("sum: %d\n", cont[0]);
+
 
   /// bind a texture to the CUDA array
   cudaBindTexture (NULL ,hysTexRef, d_hys);

  sliceIdx.x = threadIdx.x+1;
  sliceIdx.y = threadIdx.y+1;
  int edge;
  int i;

  ///load center
  s_slice[sliceIdx.y][sliceIdx.x] = hys_img[pixIdx];

  s_modified[block_tid] = 1;

  do{

    if ((threadIdx.x + threadIdx.y) == 0) m = 0;

    if (m){
      //store center
      hys_img[pixIdx] = s_slice[sliceIdx.x][sliceIdx.y];
    }

    if((!i) || 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[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;

