 
 }
 
-
-double ExtractDescriptors(cv::KeyPoint *d_pts, std::vector<CudaImage> &h_imgs, CudaImage *d_imgs,
-                          unsigned char *desc_d, float* vals_d, int patsize, int numPts) {
-  int size2 = patsize;
-  int size3 = ceil(2.0f * patsize / 3.0f);
-  int size4 = ceil(0.5f * patsize);
-  //int numPts;
-  //cudaMemcpyFromSymbol(&numPts, d_PointCounter, sizeof(int));
-
-  std::cout << "building descripor for " << numPts << " points\n";
-
-  // TimerGPU timer0(0);
-  dim3 blocks(numPts);
-  dim3 threads(EXTRACT_S);
-
-  ExtractDescriptors << <blocks, threads>>>(d_pts, d_imgs, vals_d, size2, size3, size4);
-
-
-  cudaMemsetAsync(desc_d, 0, numPts * 61);
-  BuildDescriptor << <blocks, 64>>> (vals_d, desc_d);
-
-
-  ////checkMsg("ExtractDescriptors() execution failed\n");
-  // safeCall(cudaThreadSynchronize());
-  double gpuTime = 0;  // timer0.read();
-#ifdef VERBOSE
-  printf("ExtractDescriptors time =     %.2f ms\n", gpuTime);
-#endif
-
-  return gpuTime;
-}
 
 __global__ void FindOrientation(cv::KeyPoint *d_pts, CudaImage *d_imgs) {
   __shared__ float resx[42], resy[42];

    offset += 2048;

    // How many did we add?
    if (minneighbor[2047] < 10000) {
      curridx[0] += 2048;
    } else {
      if (minneighbor[1024] < 10000) {
        if (threadIdx.x < 1023 && minneighbor[1024 + threadIdx.x] < 10000 &&
            minneighbor[1024 + threadIdx.x + 1] == 10001) {
          curridx[0] += 1024 + threadIdx.x + 1;
        }
      } else {
        if (minneighbor[threadIdx.x] < 10000 &&
            minneighbor[threadIdx.x + 1] == 10001) {
          curridx[0] += threadIdx.x + 1;
        }
      }

  }
}


double ExtractDescriptors(cv::KeyPoint *d_pts, std::vector<CudaImage> &h_imgs, CudaImage *d_imgs,
                          unsigned char *desc_d, float* vals_d, int patsize, int numPts) {
  int size2 = patsize;
  int size3 = ceil(2.0f * patsize / 3.0f);
  int size4 = ceil(0.5f * patsize);
  //int numPts;
  //cudaMemcpyFromSymbol(&numPts, d_PointCounter, sizeof(int));

  std::cout << "building descripor for " << numPts << " points\n";

  // TimerGPU timer0(0);
  dim3 blocks(numPts);
  dim3 threads(EXTRACT_S);

  ExtractDescriptors << <blocks, threads>>>(d_pts, d_imgs, vals_d, size2, size3, size4);


  cudaMemsetAsync(desc_d, 0, numPts * 61);
  BuildDescriptor << <blocks, 64>>> (vals_d, desc_d);


  ////checkMsg("ExtractDescriptors() execution failed\n");
  // safeCall(cudaThreadSynchronize());
  double gpuTime = 0;  // timer0.read();
#ifdef VERBOSE
  printf("ExtractDescriptors time =     %.2f ms\n", gpuTime);
#endif

  return gpuTime;
}




#define NTHREADS_MATCH 32
__global__ void MatchDescriptors(unsigned char *d1, unsigned char *d2,
                                 int pitch, int nkpts_2, cv::DMatch *matches) {
