   }
 }
 
+
+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) {

  upper *= 2048;

  int offset = 0;
  for (short i = threadIdx.x; i < upper; i += 2 * FilterExtremaThreads) {
    minneighbor[threadIdx.x] =
        i >= nump ? 10001 : (memberarray[i] < 0 ? 10001 : (kpts[memberarray[i]].size < 0 ? 10001 : memberarray[i]));
    minneighbor[threadIdx.x + 1024] =
        i + 1024 >= nump ? 10001
                         : (memberarray[i + 1024] < 0 ? 10001 : (kpts[memberarray[i+1024]].size < 0 ? 10001 : memberarray[i+1024]));

    __syncthreads();


    __syncthreads();

    for (int k = threadIdx.x; k < 2048; k += 1024) {
      if (minneighbor[k] < nump) {
          // Restore subpixel component
          float octsub = fabs(*(float*)(&kpts[minneighbor[k]].octave));
          int octave = (int)octsub;
