   cudaMalloc((void**)&buffer3, nump);
   FilterExtrema_kernel << <blocks, threads>>> (newpts, pts, kptindices, width,
 					       buffer1, buffer2, buffer3);
-  cudaDeviceSynchronize();
-  safeCall(cudaGetLastError());
+  //cudaDeviceSynchronize();
+  //safeCall(cudaGetLastError());
   cudaFree(buffer1);
   cudaFree(buffer2);
   cudaFree(buffer3);

  }
}

#define FindNeighborsThreads 512
__global__ void FindNeighbors(cv::KeyPoint *pts, short *kptindices, int width) {
  __shared__ int gidx[1];

  // which scale?
  int scale = pts[blockIdx.x].class_id;

  int cmpIdx = scale < 2 ? 0 : d_ExtremaIdx[scale - 2];

  float size = pts[blockIdx.x].size;


  dim3 blocks(16, 1, 1);
  dim3 threads(BitonicSortThreads, 1, 1);
  bitonicSort << <blocks, threads>>> (pts, newpts);

  /*unsigned int extremaidx_h[16];
  cudaMemcpyFromSymbol(extremaidx_h,d_ExtremaIdx,16*sizeof(unsigned int));
  for (int i=0; i<16; ++i) {
      std::cout << i << ": " << extremaidx_h[i] << std::endl;
  }
  */
  
/*  cv::KeyPoint* newpts_h = new cv::KeyPoint[nump];
  cudaMemcpy(newpts_h,newpts,nump*sizeof(cv::KeyPoint),cudaMemcpyDeviceToHost);

  int scale = 0;
  for (int i=1; i<nump; ++i) {
      cv::KeyPoint &k0 = newpts_h[i-1];
      cv::KeyPoint &k1 = newpts_h[i];

      std::cout << i << ": " << newpts_h[i].class_id << ": " << newpts_h[i].pt.y << " " << newpts_h[i].pt.x << ", " << newpts_h[i].size;

      if (!(k0.pt.y<k1.pt.y || (k0.pt.y==k1.pt.y && k0.pt.x<k1.pt.x))) std::cout << "  <<<<";
      if (k1.size < 0 ) std::cout << "  ##############";
       

      std::cout << "\n";

  }
*/

    // Find all neighbors
  cudaStreamSynchronize(copyStream);
  blocks.x = nump;
  threads.x = FindNeighborsThreads;
  FindNeighbors << <blocks, threads>>> (newpts, kptindices, width);
  //cudaDeviceSynchronize();
  //safeCall(cudaGetLastError());
  
  // Filter extrema
  blocks.x = 1;
