ACloudViewer  3.9.4
A Modern Library for 3D Data Processing
PyramidGL.cpp
Go to the documentation of this file.
1 // File: PyramidGL.cpp
3 // Author: Changchang Wu
4 // Description : implementation of PyramidGL/PyramidNaive/PyramidPackdc .
5 //
6 //
7 // Copyright (c) 2007 University of North Carolina at Chapel Hill
8 // All Rights Reserved
9 //
10 // Permission to use, copy, modify and distribute this software and its
11 // documentation for educational, research and non-profit purposes, without
12 // fee, and without a written agreement is hereby granted, provided that the
13 // above copyright notice and the following paragraph appear in all copies.
14 //
15 // The University of North Carolina at Chapel Hill make no representations
16 // about the suitability of this software for any purpose. It is provided
17 // 'as is' without express or implied warranty.
18 //
19 // Please send BUG REPORTS to ccwu@cs.unc.edu
20 //
22 
23 #include "GL/glew.h"
24 #include <iostream>
25 #include <iomanip>
26 #include <vector>
27 #include <algorithm>
28 #include <fstream>
29 #include <math.h>
30 #include <string.h>
31 using namespace std;
32 
33 #include "GlobalUtil.h"
34 #include "GLTexImage.h"
35 #include "SiftGPU.h"
36 #include "ShaderMan.h"
37 #include "SiftPyramid.h"
38 #include "ProgramGLSL.h"
39 #include "PyramidGL.h"
40 #include "FrameBufferObject.h"
41 
42 #ifdef SIFTGPU_USE_SSE_FOR
43 #ifndef __SSE__
44 #error Compiling SSE functions but SSE is not supported by the compiler.
45 #endif
46 #include <xmmintrin.h>
47 #endif
48 
49 
50 #define USE_TIMING() double t, t0, tt;
51 #define OCTAVE_START() if(GlobalUtil::_timingO){ t = t0 = CLOCK(); cout<<"#"<<i+_down_sample_factor<<"\t"; }
52 #define LEVEL_FINISH() if(GlobalUtil::_timingL){ glFinish(); tt = CLOCK();cout<<(tt-t)<<"\t"; t = CLOCK();}
53 #define OCTAVE_FINISH() if(GlobalUtil::_timingO)cout<<"|\t"<<(CLOCK()-t0)<<endl;
54 
55 
57 // Construction/Destruction
60 {
61  _texPyramid = NULL;
62  _auxPyramid = NULL;
63 }
64 
66 {
68 }
69 
70 //align must be 2^i
71 void PyramidGL:: GetAlignedStorageSize(int num, int align, int &fw, int &fh)
72 {
73  if(num <=0)
74  {
75  fw = fh = 0;
76  }else if(num < align*align)
77  {
78  fw = align;
79  fh = (int)ceil(double(num) / fw);
81  {
82  double dn = double(num);
83  int nb = (int) ceil(dn/GlobalUtil::_texMaxDim/align);
84  fw = align * nb;
85  fh = (int)ceil(dn /fw);
86  }else
87  {
88  double dn = double(num);
89  int nb = (int) ceil(dn/GlobalUtil::_texMaxDim/align);
90  fh = align * nb;
91  if(nb <=1)
92  {
93  fw = (int)ceil(dn / fh);
94  //align this dimension to blocksize
95  fw = ((int) ceil(double(fw) /align))*align;
96  }else
97  {
99  }
100 
101  }
102 
103 
104 }
105 
106 void PyramidGL::GetTextureStorageSize(int num, int &fw, int& fh)
107 {
108  if(num <=0)
109  {
110  fw = fh = 0;
111  }else if(num <= GlobalUtil::_FeatureTexBlock)
112  {
113  fw = num;
114  fh = 1;
116  {
117  double dn = double(num);
119  fw = GlobalUtil::_FeatureTexBlock * nb;
120  fh = (int)ceil(dn /fw);
121  }else
122  {
123  double dn = double(num);
126  if(nb <=1)
127  {
128  fw = (int)ceil(dn / fh);
129 
130  //align this dimension to blocksize
131 
132  //
133  if( fw < fh)
134  {
135  int temp = fh;
136  fh = fw;
137  fw = temp;
138  }
139  }else
140  {
142  }
143  }
144 }
145 
147 {
148  if(_texPyramid)
149  {
150  delete [] _texPyramid;
151  _texPyramid = NULL;
152  }
153  if(_auxPyramid)
154  {
155  delete [] _auxPyramid;
156  _auxPyramid = NULL;
157  }
158 }
160 {
161  _featureTex = NULL;
167 }
168 
170 {
174 }
175 
177 {
179  if(!GlobalUtil::_GoodOpenGL) return;
180 
183 }
184 
186 {
187  //integers vector to store the feature numbers.
188  if(_levelFeatureNum)
189  {
190  delete [] _levelFeatureNum;
192  }
193  //texture used to store features
194  if( _featureTex)
195  {
196  delete [] _featureTex;
197  _featureTex = NULL;
198  }
199  //texture used for multi-orientation
200  if(_orientationTex)
201  {
202  delete [] _orientationTex;
204  }
205  int no = _octave_num* param._dog_level_num;
206 
207  //two sets of vbos used to display the features
209  {
210  glDeleteBuffers(no, _featureDisplayVBO);
211  delete [] _featureDisplayVBO;
213  }
214  if( _featurePointVBO)
215  {
216  glDeleteBuffers(no, _featurePointVBO);
217  delete [] _featurePointVBO;
219  }
220 
221 }
222 
224 {
225  //histogram reduction
226  if(_histoPyramidTex)
227  {
228  delete[] _histoPyramidTex;
229  _hpLevelNum = 0;
231  }
232 
233  //descriptor storage shared by all levels
234  if(_descriptorTex)
235  {
236  delete [] _descriptorTex;
238  }
239  //cpu reduction buffer.
240  if(_histo_buffer)
241  {
242  delete[] _histo_buffer;
243  _histo_buffer = 0;
244  }
245 }
246 
248 {
249  GLTexImage * tex, *htex;
250  int hist_level_num = _hpLevelNum - _pyramid_octave_first;
251 
253  htex = _histoPyramidTex + hist_level_num - 1;
254  int w = tex->GetImgWidth() >> 1;
255  int h = tex->GetImgHeight() >> 1;
256 
257  for(int k = 0; k <hist_level_num -1; k++, htex--)
258  {
259  if(htex->GetImgHeight()!= h || htex->GetImgWidth() != w)
260  {
261  htex->SetImageSize(w, h);
262  htex->ZeroHistoMargin();
263  }
264 
265  w = (w + 1)>>1; h = (h + 1) >> 1;
266  }
267 }
268 
269 void PyramidNaive::FitPyramid(int w, int h)
270 {
271  //(w, h) <= (_pyramid_width, _pyramid_height);
272 
274  //
276 
277  int _octave_num_max = GetRequiredOctaveNum(min(w, h));
278 
279  if(_octave_num < 1 || _octave_num > _octave_num_max)
280  {
281  _octave_num = _octave_num_max;
282  }
283 
284 
285  int pw = _pyramid_width>>1, ph = _pyramid_height>>1;
287  pw >= w && ph >= h)
288  {
290  pw >>= 1;
291  ph >>= 1;
292  }
293 
294  for(int i = 0; i < _octave_num; i++)
295  {
296  GLTexImage * tex = GetBaseLevel(i + _octave_min);
298  for(int j = param._level_min; j <= param._level_max; j++, tex++, aux++)
299  {
300  tex->SetImageSize(w, h);
301  aux->SetImageSize(w, h);
302  }
303  w>>=1;
304  h>>=1;
305  }
306 }
307 void PyramidNaive::InitPyramid(int w, int h, int ds)
308 {
309  int wp, hp, toobig = 0;
310  if(ds == 0)
311  {
314  {
317  }else
318  {
319  wp = w << (-GlobalUtil::_octave_min_default);
320  hp = h << (-GlobalUtil::_octave_min_default);
321  }
323  }else
324  {
325  //must use 0 as _octave_min;
326  _octave_min = 0;
327  _down_sample_factor = ds;
328  w >>= ds;
329  h >>= ds;
330  wp = w;
331  hp = h;
332 
333  }
334 
336  {
337  _octave_min ++;
338  wp >>= 1;
339  hp >>= 1;
340  toobig = 1;
341  }
342 
344  max(max(wp, hp), max(_pyramid_width, _pyramid_height)) > 1024 * sqrt(GlobalUtil::_MemCapGPU / 140.0))
345  {
346  _octave_min ++;
347  wp >>= 1;
348  hp >>= 1;
349  toobig = 2;
350  }
351 
352  if(toobig && GlobalUtil::_verbose)
353  {
354  std::cout<<(toobig == 2 ? "[**SKIP OCTAVES**]:\tExceeding Memory Cap (-nomc)\n" :
355  "[**SKIP OCTAVES**]:\tReaching the dimension limit (-maxd)!\n");
356  }
357 
358  if( wp == _pyramid_width && hp == _pyramid_height && _allocated )
359  {
360  FitPyramid(wp, hp);
362  {
363  ResizePyramid(wp, hp);
364  }
365  else if( wp > _pyramid_width || hp > _pyramid_height )
366  {
368  if(wp < _pyramid_width || hp < _pyramid_height) FitPyramid(wp, hp);
369  }
370  else
371  {
372  //try use the pyramid allocated for large image on small input images
373  FitPyramid(wp, hp);
374  }
375 
376  //select the initial smoothing filter according to the new _octave_min
378 }
379 
380 void PyramidNaive::ResizePyramid( int w, int h)
381 {
382  //
383  unsigned int totalkb = 0;
384  int _octave_num_new, input_sz;
385  int i, j;
386  GLTexImage * tex, *aux;
387  //
388 
389  if(_pyramid_width == w && _pyramid_height == h && _allocated) return;
390 
391  if(w > GlobalUtil::_texMaxDim || h > GlobalUtil::_texMaxDim) return ;
392 
393  if(GlobalUtil::_verbose && GlobalUtil::_timingS) std::cout<<"[Allocate Pyramid]:\t" <<w<<"x"<<h<<endl;
394  //first octave does not change
396 
397 
398  //compute # of octaves
399 
400  input_sz = min(w,h) ;
401 
402 
403  _pyramid_width = w;
404  _pyramid_height = h;
405 
406  //reset to preset parameters
407  _octave_num_new = GlobalUtil::_octave_num_default;
408 
409  if(_octave_num_new < 1) _octave_num_new = GetRequiredOctaveNum(input_sz) ;
410 
411  if(_pyramid_octave_num != _octave_num_new)
412  {
413  //destroy the original pyramid if the # of octave changes
414  if(_octave_num >0)
415  {
418  }
419  _pyramid_octave_num = _octave_num_new;
420  }
421 
423 
424  int noct = _octave_num;
425  int nlev = param._level_num;
426 
427  // //initialize the pyramid
428  if(_texPyramid==NULL) _texPyramid = new GLTexImage[ noct* nlev ];
429  if(_auxPyramid==NULL) _auxPyramid = new GLTexImage[ noct* nlev ];
430 
431 
434  for(i = 0; i< noct; i++)
435  {
436  totalkb += (nlev * w * h * 16 / 1024);
437  for( j = 0; j< nlev; j++, tex++)
438  {
439  tex->InitTexture(w, h);
440  //tex->AttachToFBO(0);
441  }
442  //several auxilary textures are not actually required
443  totalkb += ((nlev - 3) * w * h * 16 /1024);
444  for( j = 0; j< nlev ; j++, aux++)
445  {
446  if(j < 2) continue;
447  if(j >= nlev - 1) continue;
448  aux->InitTexture(w, h, 0);
449  //aux->AttachToFBO(0);
450  }
451 
452  w>>=1;
453  h>>=1;
454  }
455 
456  totalkb += ResizeFeatureStorage();
457 
458 
459  //
460  _allocated = 1;
461 
462  if(GlobalUtil::_verbose && GlobalUtil::_timingS) std::cout<<"[Allocate Pyramid]:\t" <<(totalkb/1024)<<"MB\n";
463 
464 }
465 
466 
468 {
469  int totalkb = 0;
472 
473  int wmax = GetBaseLevel(_octave_min)->GetDrawWidth();
474  int hmax = GetBaseLevel(_octave_min)->GetDrawHeight();
475  int w ,h, i;
476 
477  //use a fbo to initialize textures..
478  FrameBufferObject fbo;
479 
480  //
481  if(_histo_buffer == NULL) _histo_buffer = new float[((size_t)1) << (2 + 2 * GlobalUtil::_ListGenSkipGPU)];
482  //histogram for feature detection
483 
484  int num = (int)ceil(log(double(max(wmax, hmax)))/log(2.0));
485 
486  if( _hpLevelNum != num)
487  {
488  _hpLevelNum = num;
490  {
491  if(_histoPyramidTex ) delete [] _histoPyramidTex;
493  w = h = 1 ;
494  for(i = 0; i < _hpLevelNum; i++)
495  {
496  _histoPyramidTex[i].InitTexture(w, h, 0);
498  w<<=1;
499  h<<=1;
500  }
501  }
502  }
503 
504  // (4 ^ (_hpLevelNum) -1 / 3) pixels
505  if(GlobalUtil::_ListGenGPU) totalkb += (((1 << (2 * _hpLevelNum)) -1) / 3 * 16 / 1024);
506 
507 
508 
509  //initialize the feature texture
510 
511  int idx = 0, n = _octave_num * param._dog_level_num;
512  if(_featureTex==NULL) _featureTex = new GLTexImage[n];
514  {
516  }
517 
518 
519  for(i = 0; i < _octave_num; i++)
520  {
522  int fmax = int(tex->GetImgWidth()*tex->GetImgHeight()*GlobalUtil::_MaxFeaturePercent);
523  int fw, fh;
524  //
526  else if(fmax < 32) fmax = 32; //give it at least a space of 32 feature
527 
528  GetTextureStorageSize(fmax, fw, fh);
529 
530  for(int j = 0; j < param._dog_level_num; j++, idx++)
531  {
532 
533  _featureTex[idx].InitTexture(fw, fh, 0);
534  _featureTex[idx].AttachToFBO(0);
535  //
536  if(_orientationTex)
537  {
538  _orientationTex[idx].InitTexture(fw, fh, 0);
540  }
541  }
542  totalkb += fw * fh * 16 * param._dog_level_num * (_orientationTex? 2 : 1) /1024;
543  }
544 
545 
546  //this just need be initialized once
547  if(_descriptorTex==NULL)
548  {
549  //initialize feature texture pyramid
550  wmax = _featureTex->GetImgWidth();
551  hmax = _featureTex->GetImgHeight();
552 
553  int nf, ns;
555  {
556  //32*4 = 128.
557  nf = 32 / GlobalUtil::_DescriptorPPT; // how many textures we need
558  ns = max(4, GlobalUtil::_DescriptorPPT); // how many point in one texture for one descriptor
559  }else
560  {
561  //at least one, resue for visualization and other work
562  nf = 1; ns = 4;
563  }
564  //
565  _alignment = ns;
566  //
567  _descriptorTex = new GLTexImage[nf];
568 
569  int fw, fh;
570  GetAlignedStorageSize(hmax*wmax* max(ns, 10), _alignment, fw, fh);
571 
572  if(fh < hmax ) fh = hmax;
573  if(fw < wmax ) fw = wmax;
574 
575  totalkb += ( fw * fh * nf * 16 /1024);
576  for(i =0; i < nf; i++)
577  {
578  _descriptorTex[i].InitTexture(fw, fh);
579  }
580  }else
581  {
583  totalkb += nf * _descriptorTex[0].GetTexWidth() * _descriptorTex[0].GetTexHeight() * 16 /1024;
584  }
585  return totalkb;
586 }
587 
588 
590 {
591  USE_TIMING();
592  GLTexPacked * tex;
593  FilterProgram** filter;
594  FrameBufferObject fbo;
595 
596  glDrawBuffer(GL_COLOR_ATTACHMENT0_EXT);
597  input->FitTexViewPort();
598 
599  for (int i = _octave_min; i < _octave_min + _octave_num; i++)
600  {
601 
602  tex = (GLTexPacked*)GetBaseLevel(i);
604 
605  OCTAVE_START();
606 
607  if( i == _octave_min )
608  {
609  if(i < 0) TextureUpSample(tex, input, 1<<(-i) );
610  else TextureDownSample(tex, input, 1<<i);
612  }else
613  {
616  }
617  LEVEL_FINISH();
618 
619  for(int j = param._level_min + 1; j <= param._level_max ; j++, tex++, filter++)
620  {
621  // filtering
622  ShaderMan::FilterImage(*filter, tex+1, tex, NULL);
623  LEVEL_FINISH();
624  }
625  OCTAVE_FINISH();
626 
627  }
628  if(GlobalUtil::_timingS) glFinish();
629  UnloadProgram();
630 }
631 
632 
633 
634 
635 
636 
637 GLTexImage* PyramidNaive::GetLevelTexture(int octave, int level, int dataName)
638 {
639  if(octave <_octave_min || octave > _octave_min + _octave_num) return NULL;
640  switch(dataName)
641  {
642  case DATA_GAUSSIAN:
643  case DATA_DOG:
644  case DATA_GRAD:
645  case DATA_ROT:
646  return _texPyramid+ (_pyramid_octave_first + octave - _octave_min) * param._level_num + (level - param._level_min);
647  case DATA_KEYPOINT:
648  return _auxPyramid + (_pyramid_octave_first + octave - _octave_min) * param._level_num + (level - param._level_min);
649  default:
650  return NULL;
651  }
652 }
653 
655 {
657  + (level - param._level_min);
658 }
659 
660 //in the packed implementation
661 // DATA_GAUSSIAN, DATA_DOG, DATA_GAD will be stored in different textures.
662 GLTexImage* PyramidNaive::GetBaseLevel(int octave, int dataName)
663 {
664  if(octave <_octave_min || octave > _octave_min + _octave_num) return NULL;
665  switch(dataName)
666  {
667  case DATA_GAUSSIAN:
668  case DATA_DOG:
669  case DATA_GRAD:
670  case DATA_ROT:
672  case DATA_KEYPOINT:
674  default:
675  return NULL;
676  }
677 }
678 
679 
680 
681 
682 
683 
684 
685 
686 
688 {
689 
690  int i, j;
691  double ts, t1;
692  GLTexImage * tex;
693  FrameBufferObject fbo;
694 
695 
697 
698  glDrawBuffer(GL_COLOR_ATTACHMENT0_EXT);
699 
700  for ( i = _octave_min; i < _octave_min + _octave_num; i++)
701  {
702  for( j = param._level_min + 1 ; j < param._level_max ; j++)
703  {
704  tex = GetLevelTexture(i, j);
705  tex->FitTexViewPort();
706  tex->AttachToFBO(0);
707  tex->BindTex();
709  tex->DrawQuadMT4();
710  }
711  }
712 
714  {
715  glFinish();
716  t1 = CLOCK();
717  std::cout<<"<Compute Gradient>\t"<<(t1-ts)<<"\n";
718  }
719 
720  UnloadProgram();
722  fbo.UnattachTex(GL_COLOR_ATTACHMENT1_EXT);
723 }
724 
725 
726 //keypoint detection with subpixel localization
728 {
729  int i, j;
730  double t0, t, ts, t1, t2;
731  GLTexImage * tex, *aux;
732  FrameBufferObject fbo;
733 
735 
736  glDrawBuffer(GL_COLOR_ATTACHMENT0_EXT);
737  //extra gradient data required for visualization
738  int gradient_only_levels[2] = {param._level_min +1, param._level_max};
739  int n_gradient_only_level = GlobalUtil::_UseSiftGPUEX ? 2 : 1;
740  for ( i = _octave_min; i < _octave_min + _octave_num; i++)
741  {
742  for( j =0; j < n_gradient_only_level ; j++)
743  {
744  tex = GetLevelTexture(i, gradient_only_levels[j]);
745  tex->FitTexViewPort();
746  tex->AttachToFBO(0);
747  tex->BindTex();
749  tex->DrawQuadMT4();
750  }
751  }
752 
754  {
755  glFinish();
756  t1 = CLOCK();
757  }
758 
759  GLenum buffers[] = { GL_COLOR_ATTACHMENT0_EXT, GL_COLOR_ATTACHMENT1_EXT };
760  glDrawBuffers(2, buffers);
761  for ( i = _octave_min; i < _octave_min + _octave_num; i++)
762  {
764  {
765  t0 = CLOCK();
766  std::cout<<"#"<<(i + _down_sample_factor)<<"\t";
767  }
768  tex = GetBaseLevel(i) + 2;
769  aux = GetBaseLevel(i, DATA_KEYPOINT) +2;
770  aux->FitTexViewPort();
771 
772  for( j = param._level_min + 2; j < param._level_max ; j++, aux++, tex++)
773  {
774  if(GlobalUtil::_timingL)t = CLOCK();
775  tex->AttachToFBO(0);
776  aux->AttachToFBO(1);
777  glActiveTexture(GL_TEXTURE0);
778  tex->BindTex();
779  glActiveTexture(GL_TEXTURE1);
780  (tex+1)->BindTex();
781  glActiveTexture(GL_TEXTURE2);
782  (tex-1)->BindTex();
783  ShaderMan::UseShaderKeypoint((tex+1)->GetTexID(), (tex-1)->GetTexID());
784  aux->DrawQuadMT8();
785 
787  {
788  glFinish();
789  std::cout<<(CLOCK()-t)<<"\t";
790  }
791  tex->DetachFBO(0);
792  aux->DetachFBO(1);
793  }
795  {
796  std::cout<<"|\t"<<(CLOCK()-t0)<<"\n";
797  }
798  }
799 
801  {
802  glFinish();
803  t2 = CLOCK();
805  std::cout <<"<Get Keypoints .. >\t"<<(t2-t1)<<"\n"
806  <<"<Extra Gradient.. >\t"<<(t1-ts)<<"\n";
807  }
808  UnloadProgram();
810  fbo.UnattachTex(GL_COLOR_ATTACHMENT1_EXT);
811 
812 
813 }
814 
815 void PyramidNaive::GenerateFeatureList(int i, int j)
816 {
817  int hist_level_num = _hpLevelNum - _pyramid_octave_first;
818  int hist_skip_gpu = GlobalUtil::_ListGenSkipGPU;
819  int idx = i * param._dog_level_num + j;
820  GLTexImage* htex, *ftex, *tex;
821  tex = GetBaseLevel(_octave_min + i, DATA_KEYPOINT) + 2 + j;
822  ftex = _featureTex + idx;
823  htex = _histoPyramidTex + hist_level_num - 1 - i;
824 
826  glActiveTexture(GL_TEXTURE0);
827  tex->BindTex();
828  htex->AttachToFBO(0);
829  int tight = ((htex->GetImgWidth() * 2 == tex->GetImgWidth() -1 || tex->GetTexWidth() == tex->GetImgWidth()) &&
830  (htex->GetImgHeight() *2 == tex->GetImgHeight()-1 || tex->GetTexHeight() == tex->GetImgHeight()));
832  htex->FitTexViewPort();
833  //this uses the fact that no feature is on the edge.
834  htex->DrawQuadReduction();
835 
836  //reduction..
837  htex--;
838 
839  //this part might have problems on several GPUS
840  //because the output of one pass is the input of the next pass
841  //need to call glFinish to make it right
842  //but too much glFinish makes it slow
843  for(int k = 0; k <hist_level_num - i - 1 - hist_skip_gpu; k++, htex--)
844  {
845  htex->AttachToFBO(0);
846  htex->FitTexViewPort();
847  (htex+1)->BindTex();
849  htex->DrawQuadReduction();
850  }
851 
852  //
853  if(hist_skip_gpu == 0)
854  {
855  //read back one pixel
856  float fn[4], fcount;
857  glReadPixels(0, 0, 1, 1, GL_RGBA , GL_FLOAT, fn);
858  fcount = (fn[0] + fn[1] + fn[2] + fn[3]);
859  if(fcount < 1) fcount = 0;
860 
861 
862  _levelFeatureNum[ idx] = (int)(fcount);
863  SetLevelFeatureNum(idx, (int)fcount);
864  _featureNum += int(fcount);
865 
866  //
867  if(fcount < 1.0) return;
868 
869 
871 
872  htex= _histoPyramidTex;
873 
874  htex->BindTex();
875 
876  //first pass
877  ftex->AttachToFBO(0);
879  {
880  //this is very important...
881  ftex->FitRealTexViewPort();
882  glClear(GL_COLOR_BUFFER_BIT);
883  glFinish();
884  }else
885  {
886  ftex->FitTexViewPort();
887  //glFinish();
888  }
889 
890 
891  ShaderMan::UseShaderGenListStart((float)ftex->GetImgWidth(), htex->GetTexID());
892 
893  ftex->DrawQuad();
894  //make sure it finishes before the next step
895  ftex->DetachFBO(0);
896 
897  //pass on each pyramid level
898  htex++;
899  }else
900  {
901 
902  int tw = htex[1].GetDrawWidth(), th = htex[1].GetDrawHeight();
903  int fc = 0;
904  glReadPixels(0, 0, tw, th, GL_RGBA , GL_FLOAT, _histo_buffer);
905  _keypoint_buffer.resize(0);
906  for(int y = 0, pos = 0; y < th; y++)
907  {
908  for(int x= 0; x < tw; x++)
909  {
910  for(int c = 0; c < 4; c++, pos++)
911  {
912  int ss = (int) _histo_buffer[pos];
913  if(ss == 0) continue;
914  float ft[4] = {2 * x + (c%2? 1.5f: 0.5f), 2 * y + (c>=2? 1.5f: 0.5f), 0, 1 };
915  for(int t = 0; t < ss; t++)
916  {
917  ft[2] = (float) t;
918  _keypoint_buffer.insert(_keypoint_buffer.end(), ft, ft+4);
919  }
920  fc += (int)ss;
921  }
922  }
923  }
924  _levelFeatureNum[ idx] = fc;
925  SetLevelFeatureNum(idx, fc);
926  if(fc == 0) return;
927  _featureNum += fc;
929  ftex->AttachToFBO(0);
931  {
932  ftex->FitRealTexViewPort();
933  glClear(GL_COLOR_BUFFER_BIT);
934  glFlush();
935  }else
936  {
937  ftex->FitTexViewPort();
938  glFlush();
939  }
940  _keypoint_buffer.resize(ftex->GetDrawWidth() * ftex->GetDrawHeight()*4, 0);
942  glActiveTexture(GL_TEXTURE0);
943  ftex->BindTex();
944  glTexSubImage2D(GlobalUtil::_texTarget, 0, 0, 0, ftex->GetDrawWidth(),
945  ftex->GetDrawHeight(), GL_RGBA, GL_FLOAT, &_keypoint_buffer[0]);
946  htex += 2;
947  }
948 
949  for(int lev = 1 + hist_skip_gpu; lev < hist_level_num - i; lev++, htex++)
950  {
951 
952  glActiveTexture(GL_TEXTURE0);
953  ftex->BindTex();
954  ftex->AttachToFBO(0);
955  glActiveTexture(GL_TEXTURE1);
956  htex->BindTex();
958  ftex->DrawQuad();
959  ftex->DetachFBO(0);
960  }
962 
963 }
964 
965 //generate feature list on GPU
967 {
968  //generate the histogram0pyramid
969  FrameBufferObject fbo;
970  glReadBuffer(GL_COLOR_ATTACHMENT0_EXT);
971  glDrawBuffer(GL_COLOR_ATTACHMENT0_EXT);
972  double t1, t2;
973  int ocount, reverse = (GlobalUtil::_TruncateMethod == 1);
974  _featureNum = 0;
975 
977 
978  //for(int i = 0, idx = 0; i < _octave_num; i++)
979  FOR_EACH_OCTAVE(i, reverse)
980  {
981  //output
983  {
984  t1= CLOCK();
985  ocount = 0;
986  std::cout<<"#"<<i+_octave_min + _down_sample_factor<<":\t";
987  }
988  //for(int j = 0; j < param._dog_level_num; j++, idx++)
989  FOR_EACH_LEVEL(j, reverse)
990  {
991 
994  {
996  continue;
997  }else
998  {
999  GenerateFeatureList(i, j);
1001  {
1002  int idx = i * param._dog_level_num + j;
1003  std::cout<< _levelFeatureNum[idx] <<"\t";
1004  ocount += _levelFeatureNum[idx];
1005  }
1006  }
1007  }
1009  {
1010  t2 = CLOCK();
1011  std::cout << "| \t" << int(ocount) << " :\t(" << (t2 - t1) << ")\n";
1012  }
1013  }
1014  if(GlobalUtil::_timingS)glFinish();
1016  {
1017  std::cout<<"#Features:\t"<<_featureNum<<"\n";
1018  }
1019 }
1020 
1021 
1023 {
1024  //use a big VBO to save all the SIFT box vertices
1025  int w, h, esize; GLint bsize;
1026  int nvbo = _octave_num * param._dog_level_num;
1027  //initialize the vbos
1029  {
1030  _featureDisplayVBO = new GLuint[nvbo];
1031  glGenBuffers( nvbo, _featureDisplayVBO );
1032  }
1033  if(_featurePointVBO == NULL)
1034  {
1035  _featurePointVBO = new GLuint[nvbo];
1036  glGenBuffers(nvbo, _featurePointVBO);
1037  }
1038 
1039  FrameBufferObject fbo;
1040  glReadBuffer(GL_COLOR_ATTACHMENT0_EXT);
1041  glDrawBuffer(GL_COLOR_ATTACHMENT0_EXT);
1042  glActiveTexture(GL_TEXTURE0);
1043  //
1044  GLTexImage & tempTex = *_descriptorTex;
1045  //
1046  for(int i = 0, idx = 0; i < _octave_num; i++)
1047  {
1048  for(int j = 0; j < param._dog_level_num; j ++, idx++)
1049  {
1050  GLTexImage * ftex = _featureTex + idx;
1051  if(_levelFeatureNum[idx]<=0)continue;
1052 
1053  //copy the texture into vbo
1054  fbo.BindFBO();
1055  tempTex.AttachToFBO(0);
1056 
1057  ftex->BindTex();
1058  ftex->FitTexViewPort();
1060  ftex->DrawQuad();
1061 
1062  glBindBuffer(GL_PIXEL_PACK_BUFFER_ARB, _featurePointVBO[ idx]);
1063  glGetBufferParameteriv(GL_PIXEL_PACK_BUFFER_ARB, GL_BUFFER_SIZE, &bsize);
1064  esize = ftex->GetImgHeight() * ftex->GetImgWidth()*sizeof(float) *4;
1065 
1066  //increase size when necessary
1067  if(bsize < esize)
1068  {
1069  glBufferData(GL_PIXEL_PACK_BUFFER_ARB, esize*3/2 , NULL, GL_STATIC_DRAW_ARB);
1070  glGetBufferParameteriv(GL_PIXEL_PACK_BUFFER_ARB, GL_BUFFER_SIZE, &bsize);
1071  }
1072 
1073  //read back if we have enough buffer
1074  if(bsize >= esize) glReadPixels(0, 0, ftex->GetImgWidth(), ftex->GetImgHeight(), GL_RGBA, GL_FLOAT, 0);
1075  else glBufferData(GL_PIXEL_PACK_BUFFER_ARB, 0, NULL, GL_STATIC_DRAW_ARB);
1076  glBindBuffer(GL_PIXEL_PACK_BUFFER_ARB, 0);
1077 
1078 
1079  //box display vbo
1080  int count = _levelFeatureNum[idx]* 10;
1082  w = (int)ceil(double(count)/ h);
1083 
1084  //input
1085  fbo.BindFBO();
1086  ftex->BindTex();
1087 
1088  //output
1089  tempTex.AttachToFBO(0);
1091  //shader
1092  ShaderMan::UseShaderGenVBO( (float)ftex->GetImgWidth(), (float) w,
1093  param.GetLevelSigma(j + param._level_min + 1));
1094  GLTexImage::DrawQuad(0, (float)w, 0, (float)h);
1095 
1096  //
1097  glBindBuffer(GL_PIXEL_PACK_BUFFER_ARB, _featureDisplayVBO[ idx]);
1098  glGetBufferParameteriv(GL_PIXEL_PACK_BUFFER_ARB, GL_BUFFER_SIZE, &bsize);
1099  esize = w*h * sizeof(float)*4;
1100  //increase size when necessary
1101  if(bsize < esize)
1102  {
1103  glBufferData(GL_PIXEL_PACK_BUFFER_ARB, esize*3/2, NULL, GL_STATIC_DRAW_ARB);
1104  glGetBufferParameteriv(GL_PIXEL_PACK_BUFFER_ARB, GL_BUFFER_SIZE, &bsize);
1105  }
1106 
1107  //read back if we have enough buffer
1108  if(bsize >= esize) glReadPixels(0, 0, w, h, GL_RGBA, GL_FLOAT, 0);
1109  else glBufferData(GL_PIXEL_PACK_BUFFER_ARB, 0, NULL, GL_STATIC_DRAW_ARB);
1110  glBindBuffer(GL_PIXEL_PACK_BUFFER_ARB, 0);
1111 
1112 
1113 
1114 
1115  }
1116  }
1117  glReadBuffer(GL_NONE);
1118  glFinish();
1119 
1120 }
1121 
1122 
1123 
1124 
1125 
1127 {
1128  GLTexImage * gtex;
1129  GLTexImage * stex = NULL;
1130  GLTexImage * ftex = _featureTex;
1131  GLTexImage * otex = _orientationTex;
1132  int sid = 0;
1133  int * count = _levelFeatureNum;
1134  float sigma, sigma_step = powf(2.0f, 1.0f/param._dog_level_num);
1135  FrameBufferObject fbo;
1136  if(_orientationTex)
1137  {
1138  GLenum buffers[] = { GL_COLOR_ATTACHMENT0_EXT, GL_COLOR_ATTACHMENT1_EXT };
1139  glDrawBuffers(2, buffers);
1140  }else
1141  {
1142  glDrawBuffer(GL_COLOR_ATTACHMENT0_EXT);
1143  }
1144  for(int i = 0; i < _octave_num; i++)
1145  {
1148  stex = GetBaseLevel(i+_octave_min, DATA_KEYPOINT) + 2;
1149 
1150  for(int j = 0; j < param._dog_level_num; j++, ftex++, otex++, count++, gtex++, stex++)
1151  {
1152  if(*count<=0)continue;
1153 
1154  sigma = param.GetLevelSigma(j+param._level_min+1);
1155 
1156  //
1157  ftex->FitTexViewPort();
1158 
1159  glActiveTexture(GL_TEXTURE0);
1160  ftex->BindTex();
1161  glActiveTexture(GL_TEXTURE1);
1162  gtex->BindTex();
1163  //
1164  ftex->AttachToFBO(0);
1165  if(_orientationTex) otex->AttachToFBO(1);
1167  {
1168  glActiveTexture(GL_TEXTURE2);
1169  stex->BindTex();
1170  sid = * stex;
1171  }
1172  ShaderMan::UseShaderOrientation(gtex->GetTexID(),
1173  gtex->GetImgWidth(), gtex->GetImgHeight(),
1174  sigma, sid, sigma_step, _existing_keypoints);
1175  ftex->DrawQuad();
1176  // glFinish();
1177 
1178  }
1179  }
1180 
1182  if(GlobalUtil::_timingS)glFinish();
1183 
1184  if(_orientationTex) fbo.UnattachTex(GL_COLOR_ATTACHMENT1_EXT);
1185 
1186 }
1187 
1188 
1189 
1190 //to compare with GPU feature list generation
1192 {
1193 
1194  FrameBufferObject fbo;
1195  _featureNum = 0;
1197  float * mem = new float [tex->GetTexWidth()*tex->GetTexHeight()];
1198  vector<float> list;
1199  int idx = 0;
1200  for(int i = 0; i < _octave_num; i++)
1201  {
1202  for(int j = 0; j < param._dog_level_num; j++, idx++)
1203  {
1204  tex = GetBaseLevel(_octave_min + i, DATA_KEYPOINT) + j + 2;
1205  tex->BindTex();
1206  glGetTexImage(GlobalUtil::_texTarget, 0, GL_RED, GL_FLOAT, mem);
1207  //tex->AttachToFBO(0);
1208  //tex->FitTexViewPort();
1209  //glReadPixels(0, 0, tex->GetTexWidth(), tex->GetTexHeight(), GL_RED, GL_FLOAT, mem);
1210  //
1211  //make a list of
1212  list.resize(0);
1213  float * p = mem;
1214  int fcount = 0 ;
1215  for(int k = 0; k < tex->GetTexHeight(); k++)
1216  {
1217  for( int m = 0; m < tex->GetTexWidth(); m ++, p++)
1218  {
1219  if(*p==0)continue;
1220  if(m ==0 || k ==0 || k >= tex->GetImgHeight() -1 || m >= tex->GetImgWidth() -1 ) continue;
1221  list.push_back(m+0.5f);
1222  list.push_back(k+0.5f);
1223  list.push_back(0);
1224  list.push_back(1);
1225  fcount ++;
1226 
1227 
1228  }
1229  }
1230  if(fcount==0)continue;
1231 
1232 
1233 
1234  GLTexImage * ftex = _featureTex+idx;
1235  _levelFeatureNum[idx] = (fcount);
1236  SetLevelFeatureNum(idx, fcount);
1237 
1238  _featureNum += (fcount);
1239 
1240 
1241  int fw = ftex->GetImgWidth();
1242  int fh = ftex->GetImgHeight();
1243 
1244  list.resize(4*fh*fw);
1245 
1246  ftex->BindTex();
1247  ftex->AttachToFBO(0);
1248  // glTexImage2D(GlobalUtil::_texTarget, 0, GlobalUtil::_iTexFormat, fw, fh, 0, GL_BGRA, GL_FLOAT, &list[0]);
1249  glTexSubImage2D(GlobalUtil::_texTarget, 0, 0, 0, fw, fh, GL_RGBA, GL_FLOAT, &list[0]);
1250  //
1251  }
1252  }
1254  delete[] mem;
1256  {
1257  std::cout<<"#Features:\t"<<_featureNum<<"\n";
1258  }
1259 }
1260 
1261 #define FEATURELIST_USE_PBO
1262 
1264 {
1265  //make a compact feature list, each with only one orientation
1266  //download orientations and the featue list
1267  //reshape it and upload it
1268 
1269  FrameBufferObject fbo;
1270  int i, szmax =0, sz;
1272  for( i = 0; i < n; i++)
1273  {
1275  if(sz > szmax ) szmax = sz;
1276  }
1277  float * buffer = new float[szmax*24];
1278  float * buffer1 = buffer;
1279  float * buffer2 = buffer + szmax*4;
1280  float * buffer3 = buffer + szmax*8;
1281 
1282  glReadBuffer(GL_COLOR_ATTACHMENT0_EXT);
1283  glDrawBuffer(GL_COLOR_ATTACHMENT0_EXT);
1284 
1285 #ifdef FEATURELIST_USE_PBO
1286  GLuint ListUploadPBO;
1287  glGenBuffers(1, &ListUploadPBO);
1288  glBindBuffer(GL_PIXEL_UNPACK_BUFFER_ARB, ListUploadPBO);
1289  glBufferData(GL_PIXEL_UNPACK_BUFFER_ARB, szmax * 8 * sizeof(float), NULL, GL_STREAM_DRAW);
1290 #endif
1291 
1292  _featureNum = 0;
1293 
1294 #ifdef NO_DUPLICATE_DOWNLOAD
1295  const double twopi = 2.0*3.14159265358979323846;
1296  _keypoint_buffer.resize(0);
1297  float os = _octave_min>=0? float(1<<_octave_min): 1.0f/(1<<(-_octave_min));
1298  if(_down_sample_factor>0) os *= float(1<<_down_sample_factor);
1299  float offset = GlobalUtil::_LoweOrigin? 0 : 0.5f;
1300 #endif
1301 
1302  for(i = 0; i < n; i++)
1303  {
1304  if(_levelFeatureNum[i]==0)continue;
1305 
1306  _featureTex[i].AttachToFBO(0);
1308  glReadPixels(0, 0, _featureTex[i].GetImgWidth(), _featureTex[i].GetImgHeight(),GL_RGBA, GL_FLOAT, buffer1);
1309 
1310  int fcount =0, ocount;
1311  float * src = buffer1;
1312  float * orientation = buffer2;
1313  float * des = buffer3;
1315  {
1316  //read back orientations from another texture
1318  glReadPixels(0, 0, _orientationTex[i].GetImgWidth(), _orientationTex[i].GetImgHeight(),GL_RGBA, GL_FLOAT, buffer2);
1319  //make the feature list
1320  for(int j = 0; j < _levelFeatureNum[i]; j++, src+=4, orientation+=4)
1321  {
1322  if(_existing_keypoints)
1323  {
1324  des[0] = src[0];
1325  des[1] = src[1];
1326  des[2] = orientation[0];
1327  des[3] = src[3];
1328  fcount++;
1329  des += 4;
1330  }else
1331  {
1332  ocount = (int)src[2];
1333  for(int k = 0 ; k < ocount; k++, des+=4)
1334  {
1335  des[0] = src[0];
1336  des[1] = src[1];
1337  des[2] = orientation[k];
1338  des[3] = src[3];
1339  fcount++;
1340  }
1341  }
1342  }
1343  }else
1344  {
1345  _featureTex[i].DetachFBO(0);
1346  const static double factor = 2.0*3.14159265358979323846/65535.0;
1347  for(int j = 0; j < _levelFeatureNum[i]; j++, src+=4)
1348  {
1349  unsigned short * orientations = (unsigned short*) (&src[2]);
1350  if(_existing_keypoints)
1351  {
1352  des[0] = src[0];
1353  des[1] = src[1];
1354  des[2] = float( factor* orientations[0]);
1355  des[3] = src[3];
1356  fcount++;
1357  des += 4;
1358  }else
1359  {
1360  if(orientations[0] != 65535)
1361  {
1362  des[0] = src[0];
1363  des[1] = src[1];
1364  des[2] = float( factor* orientations[0]);
1365  des[3] = src[3];
1366  fcount++;
1367  des += 4;
1368 
1369  if(orientations[1] != 65535)
1370  {
1371  des[0] = src[0];
1372  des[1] = src[1];
1373  des[2] = float(factor* orientations[1]);
1374  des[3] = src[3];
1375  fcount++;
1376  des += 4;
1377  }
1378  }
1379  }
1380  }
1381  }
1382 
1383  if (fcount == 0){ _levelFeatureNum[i] = 0; continue; }
1384 
1385  //texture size --------------
1386  SetLevelFeatureNum(i, fcount);
1387  int nfw = _featureTex[i].GetImgWidth();
1388  int nfh = _featureTex[i].GetImgHeight();
1389  int sz = nfh * nfw;
1390  if(sz > fcount) memset(des, 0, sizeof(float) * (sz - fcount) * 4);
1391 
1392 #ifndef FEATURELIST_USE_PBO
1393  _featureTex[i].BindTex();
1394  glTexSubImage2D(GlobalUtil::_texTarget, 0, 0, 0, nfw, nfh, GL_RGBA, GL_FLOAT, buffer3);
1395  _featureTex[i].UnbindTex();
1396 #else
1397  float* mem = (float*) glMapBuffer(GL_PIXEL_UNPACK_BUFFER_ARB, GL_WRITE_ONLY);
1398  memcpy(mem, buffer3, sz * 4 * sizeof(float) );
1399  glUnmapBuffer(GL_PIXEL_UNPACK_BUFFER_ARB);
1400  _featureTex[i].BindTex();
1401  glTexSubImage2D(GlobalUtil::_texTarget, 0, 0, 0, nfw, nfh, GL_RGBA, GL_FLOAT, 0);
1402  _featureTex[i].UnbindTex();
1403 #endif
1404 
1405 #ifdef NO_DUPLICATE_DOWNLOAD
1406  if(fcount > 0)
1407  {
1408  float oss = os * (1 << (i / param._dog_level_num));
1409  _keypoint_buffer.resize((_featureNum + fcount) * 4);
1410  float* ds = &_keypoint_buffer[_featureNum * 4];
1411  float* fs = buffer3;
1412  for(int k = 0; k < fcount; k++, ds+=4, fs+=4)
1413  {
1414  ds[0] = oss*(fs[0]-0.5f) + offset; //x
1415  ds[1] = oss*(fs[1]-0.5f) + offset; //y
1416  ds[3] = (float)fmod(twopi-fs[2], twopi); //orientation, mirrored
1417  ds[2] = oss*fs[3]; //scale
1418  }
1419  }
1420 #endif
1421  _levelFeatureNum[i] = fcount;
1422  _featureNum += fcount;
1423  }
1424 
1425  delete[] buffer;
1427  {
1428  std::cout<<"#Features MO:\t"<<_featureNum<<endl;
1429  }
1431 #ifdef FEATURELIST_USE_PBO
1432  glBindBuffer(GL_PIXEL_UNPACK_BUFFER_ARB, 0);
1433  glDeleteBuffers(1, &ListUploadPBO);
1434 #endif
1435 }
1436 
1437 
1438 
1439 inline void PyramidGL::SetLevelFeatureNum(int idx, int fcount)
1440 {
1441  int fw, fh;
1442  GLTexImage * ftex = _featureTex + idx;
1443  //set feature texture size. normally fh will be one
1444  GetTextureStorageSize(fcount, fw, fh);
1445  if(fcount > ftex->GetTexWidth()*ftex->GetTexHeight())
1446  {
1447  ftex->InitTexture(fw, fh, 0);
1448  if(_orientationTex) _orientationTex[idx].InitTexture(fw, fh, 0);
1449 
1450  }
1452  fh = fcount ==0? 0:(int)ceil(double(fcount)/fw);
1453  else
1454  fw = fcount ==0? 0:(int)ceil(double(fcount)/fh);
1455  ftex->SetImageSize(fw, fh);
1456  if(_orientationTex) _orientationTex[idx].SetImageSize(fw, fh);
1457 }
1458 
1460 {
1465 }
1466 
1468 {
1469  //
1470  int idx = 0;
1471 // int n = _octave_num * param._dog_level_num;
1472  float sigma, sigma_step = powf(2.0f, 1.0f/param._dog_level_num);
1473  GLTexImage * ftex = _featureTex;
1474 
1475  FrameBufferObject fbo;
1476  glDrawBuffer(GL_COLOR_ATTACHMENT0_EXT);
1477  for(int i = 0; i < _octave_num; i++)
1478  {
1480  for(int j = 0; j < param._dog_level_num; j++, ftex++, gtex++, idx ++)
1481  {
1482  if(_levelFeatureNum[idx]<=0)continue;
1483  sigma = param.GetLevelSigma(j+param._level_min+1);
1484 
1485  //
1486  ftex->AttachToFBO(0);
1487  ftex->FitTexViewPort();
1488 
1489  glActiveTexture(GL_TEXTURE0);
1490  ftex->BindTex();
1491  glActiveTexture(GL_TEXTURE1);
1492  gtex->BindTex();
1493 
1494  ShaderMan::UseShaderSimpleOrientation(gtex->GetTexID(), sigma, sigma_step);
1495  ftex->DrawQuad();
1496  }
1497  }
1498 
1500 
1501 }
1502 
1503 
1504 #ifdef SIFTGPU_USE_SSE_FOR
1505  static inline float dotproduct_128d(float * p)
1506  {
1507  float z = 0.0f;
1508  __m128 sse =_mm_load_ss(&z);
1509  float* pf = (float*) (&sse);
1510  for( int i = 0; i < 32; i++, p+=4)
1511  {
1512  __m128 ps = _mm_loadu_ps(p);
1513  sse = _mm_add_ps(sse, _mm_mul_ps(ps, ps));
1514  }
1515  return pf[0] + pf[1] + pf[2] + pf[3];
1516 
1517  }
1518  static inline void multiply_and_truncate_128d(float* p, float m)
1519  {
1520  float z = 0.2f;
1521  __m128 t = _mm_load_ps1(&z);
1522  __m128 r = _mm_load_ps1(&m);
1523  for(int i = 0; i < 32; i++, p+=4)
1524  {
1525  __m128 ps = _mm_loadu_ps(p);
1526  _mm_storeu_ps(p, _mm_min_ps(_mm_mul_ps(ps, r), t));
1527  }
1528  }
1529  static inline void multiply_128d(float* p, float m)
1530  {
1531  __m128 r = _mm_load_ps1(&m);
1532  for(int i = 0; i < 32; i++, p+=4)
1533  {
1534  __m128 ps = _mm_loadu_ps(p);
1535  _mm_storeu_ps(p, _mm_mul_ps(ps, r));
1536  }
1537  }
1538 #endif
1539 
1540 
1541 inline void PyramidGL::NormalizeDescriptor(int num, float*pd)
1542 {
1543 
1544 #ifdef SIFTGPU_USE_SSE_FOR
1545  for(int k = 0; k < num; k++, pd +=128)
1546  {
1547  float sq;
1548  //normalize and truncate to .2
1549  sq = dotproduct_128d(pd); sq = 1.0f / sqrtf(sq);
1550  multiply_and_truncate_128d(pd, sq);
1551 
1552  //renormalize
1553  sq = dotproduct_128d(pd); sq = 1.0f / sqrtf(sq);
1554  multiply_128d(pd, sq);
1555  }
1556 #else
1557  //descriptor normalization runs on cpu for OpenGL implemenations
1558  for(int k = 0; k < num; k++, pd +=128)
1559  {
1560  int v;
1561  float* ppd, sq = 0;
1562  //int v;
1563  //normalize
1564  ppd = pd;
1565  for(v = 0 ; v < 128; v++, ppd++) sq += (*ppd)*(*ppd);
1566  sq = 1.0f / sqrtf(sq);
1567  //truncate to .2
1568  ppd = pd;
1569  for(v = 0; v < 128; v ++, ppd++) *ppd = min(*ppd*sq, 0.2f);
1570 
1571  //renormalize
1572  ppd = pd; sq = 0;
1573  for(v = 0; v < 128; v++, ppd++) sq += (*ppd)*(*ppd);
1574  sq = 1.0f / sqrtf(sq);
1575 
1576  ppd = pd;
1577  for(v = 0; v < 128; v ++, ppd++) *ppd = *ppd*sq;
1578  }
1579 
1580 #endif
1581 }
1582 
1583 inline void PyramidGL::InterlaceDescriptorF2(int w, int h, float* buf, float* pd, int step)
1584 {
1585  /*
1586  if(GlobalUtil::_DescriptorPPR == 8)
1587  {
1588  const int dstep = w * 128;
1589  float* pp1 = buf;
1590  float* pp2 = buf + step;
1591 
1592  for(int u = 0; u < h ; u++, pd+=dstep)
1593  {
1594  int v;
1595  float* ppd = pd;
1596  for(v= 0; v < w; v++)
1597  {
1598  for(int t = 0; t < 8; t++)
1599  {
1600  *ppd++ = *pp1++;*ppd++ = *pp1++;*ppd++ = *pp1++;*ppd++ = *pp1++;
1601  *ppd++ = *pp2++;*ppd++ = *pp2++;*ppd++ = *pp2++;*ppd++ = *pp2++;
1602  }
1603  ppd += 64;
1604  }
1605  ppd = pd + 64;
1606  for(v= 0; v < w; v++)
1607  {
1608  for(int t = 0; t < 8; t++)
1609  {
1610  *ppd++ = *pp1++;*ppd++ = *pp1++;*ppd++ = *pp1++;*ppd++ = *pp1++;
1611  *ppd++ = *pp2++;*ppd++ = *pp2++;*ppd++ = *pp2++;*ppd++ = *pp2++;
1612  }
1613  ppd += 64;
1614  }
1615  }
1616 
1617  }else */
1619  {
1620  //interlace
1621  for(int k = 0; k < 2; k++)
1622  {
1623  float* pp = buf + k * step;
1624  float* ppd = pd + k * 4;
1625  for(int u = 0; u < h ; u++)
1626  {
1627  int v;
1628  for(v= 0; v < w; v++)
1629  {
1630  for(int t = 0; t < 8; t++)
1631  {
1632  ppd[0] = pp[0];
1633  ppd[1] = pp[1];
1634  ppd[2] = pp[2];
1635  ppd[3] = pp[3];
1636  ppd += 8;
1637  pp+= 4;
1638  }
1639  ppd += 64;
1640  }
1641  ppd += ( 64 - 128 * w );
1642  for(v= 0; v < w; v++)
1643  {
1644  for(int t = 0; t < 8; t++)
1645  {
1646  ppd[0] = pp[0];
1647  ppd[1] = pp[1];
1648  ppd[2] = pp[2];
1649  ppd[3] = pp[3];
1650 
1651  ppd += 8;
1652  pp+= 4;
1653  }
1654  ppd += 64;
1655  }
1656  ppd -=64;
1657  }
1658  }
1659  }else if(GlobalUtil::_DescriptorPPR == 4)
1660  {
1661 
1662  }
1663 
1664 
1665 
1666 }
1668 {
1669  //descriptors...
1670  float sigma;
1671  int idx, i, j, k, w, h;
1672  int ndf = 32 / GlobalUtil::_DescriptorPPT; //number of textures
1673  int block_width = GlobalUtil::_DescriptorPPR;
1675  float* pd = &_descriptor_buffer[0], * pbuf = NULL;
1676  vector<float>read_buffer, descriptor_buffer2;
1677 
1678  //use another buffer, if we need to re-order the descriptors
1679  if(_keypoint_index.size() > 0)
1680  {
1681  descriptor_buffer2.resize(_descriptor_buffer.size());
1682  pd = &descriptor_buffer2[0];
1683  }
1684  FrameBufferObject fbo;
1685 
1686  GLTexImage * gtex, *otex, * ftex;
1687  GLenum buffers[8] = {
1688  GL_COLOR_ATTACHMENT0_EXT, GL_COLOR_ATTACHMENT1_EXT ,
1689  GL_COLOR_ATTACHMENT2_EXT, GL_COLOR_ATTACHMENT3_EXT ,
1690  GL_COLOR_ATTACHMENT4_EXT, GL_COLOR_ATTACHMENT5_EXT ,
1691  GL_COLOR_ATTACHMENT6_EXT, GL_COLOR_ATTACHMENT7_EXT ,
1692  };
1693 
1694  glDrawBuffers(ndf, buffers);
1695  glReadBuffer(GL_COLOR_ATTACHMENT0_EXT);
1696 
1697 
1698  for( i = 0, idx = 0, ftex = _featureTex; i < _octave_num; i++)
1699  {
1700  gtex = GetBaseLevel(i + _octave_min, DATA_GRAD) + 1;
1701  otex = GetBaseLevel(i + _octave_min, DATA_ROT) + 1;
1702  for( j = 0; j < param._dog_level_num; j++, ftex++, idx++, gtex++, otex++)
1703  {
1704  if(_levelFeatureNum[idx]==0)continue;
1705 
1707  int count = _levelFeatureNum[idx] * block_width;
1708  GetAlignedStorageSize(count, block_width, w, h);
1709  h = ((int)ceil(double(count) / w)) * block_height;
1710 
1711  //not enought space for holding the descriptor data
1712  if(w > _descriptorTex[0].GetTexWidth() || h > _descriptorTex[0].GetTexHeight())
1713  {
1714  for(k = 0; k < ndf; k++)_descriptorTex[k].InitTexture(w, h);
1715  }
1716  for(k = 0; k < ndf; k++) _descriptorTex[k].AttachToFBO(k);
1718  glActiveTexture(GL_TEXTURE0);
1719  ftex->BindTex();
1720  glActiveTexture(GL_TEXTURE1);
1721  gtex->BindTex();
1722  if(otex!=gtex)
1723  {
1724  glActiveTexture(GL_TEXTURE2);
1725  otex->BindTex();
1726  }
1727 
1728  ShaderMan::UseShaderDescriptor(gtex->GetTexID(), otex->GetTexID(),
1729  w, ftex->GetImgWidth(), gtex->GetImgWidth(), gtex->GetImgHeight(), sigma);
1730  GLTexImage::DrawQuad(0, (float)w, 0, (float)h);
1731 
1732  //read back float format descriptors and do normalization on CPU
1733  int step = w*h*4;
1734  if((unsigned int)step*ndf > read_buffer.size())
1735  {
1736  read_buffer.resize(ndf*step);
1737  }
1738  pbuf = &read_buffer[0];
1739 
1740  //read back
1741  for(k = 0; k < ndf; k++, pbuf+=step)
1742  {
1743  glReadBuffer(GL_COLOR_ATTACHMENT0_EXT + k);
1744  if(GlobalUtil::_IsNvidia || w * h <= 16384) //were
1745  {
1746  glReadPixels(0, 0, w, h, GL_RGBA, GL_FLOAT, pbuf);
1747  }else
1748  {
1749  int hstep = 16384 / w;
1750  for(int kk = 0; kk < h; kk += hstep)
1751  glReadPixels(0, kk, w, min(hstep, h - kk), GL_RGBA, GL_FLOAT, pbuf + w * kk * 4);
1752  }
1753  }
1754 
1755  //the following two steps run on cpu, so better cpu better speed
1756  //and release version can be a lot faster than debug version
1757  //interlace data on the two texture to get the descriptor
1758  InterlaceDescriptorF2(w / block_width, h / block_height, &read_buffer[0], pd, step);
1759 
1760  //need to do normalization
1761  //the new version uses SSE to speed up this part
1763 
1764  pd += 128*_levelFeatureNum[idx];
1765  glReadBuffer(GL_NONE);
1766  }
1767  }
1768 
1769 
1770  //finally, put the descriptor back to their original order for existing keypoint list.
1771  if(_keypoint_index.size() > 0)
1772  {
1773  for(i = 0; i < _featureNum; ++i)
1774  {
1775  int index = _keypoint_index[i];
1776  memcpy(&_descriptor_buffer[index*128], &descriptor_buffer2[i*128], 128 * sizeof(float));
1777  }
1778  }
1779 
1782  glDrawBuffer(GL_NONE);
1784  if(GlobalUtil::_timingS)glFinish();
1785  for(i = 0; i < ndf; i++) fbo.UnattachTex(GL_COLOR_ATTACHMENT0_EXT +i);
1786 
1787 }
1788 
1789 
1791 {
1792  const double twopi = 2.0*3.14159265358979323846;
1793  int idx = 0;
1794  float * buffer = &_keypoint_buffer[0];
1795  vector<float> keypoint_buffer2;
1796  //use a different keypoint buffer when processing with an exisint features list
1797  //without orientation information.
1798  if(_keypoint_index.size() > 0)
1799  {
1800  keypoint_buffer2.resize(_keypoint_buffer.size());
1801  buffer = &keypoint_buffer2[0];
1802  }
1803  float * p = buffer, *ps, sigma;
1804  GLTexImage * ftex = _featureTex;
1805  FrameBufferObject fbo;
1806  ftex->FitRealTexViewPort();
1808  float os = _octave_min>=0? float(1<<_octave_min): 1.0f/(1<<(-_octave_min));
1809  if(_down_sample_factor>0) os *= float(1<<_down_sample_factor);
1810  float offset = GlobalUtil::_LoweOrigin? 0 : 0.5f;
1812  for(int i = 0; i < _octave_num; i++, os *= 2.0f)
1813  {
1814 
1815  for(int j = 0; j < param._dog_level_num; j++, idx++, ftex++)
1816  {
1817 
1818  if(_levelFeatureNum[idx]>0)
1819  {
1820  ftex->AttachToFBO(0);
1821  glReadPixels(0, 0, ftex->GetImgWidth(), ftex->GetImgHeight(),GL_RGBA, GL_FLOAT, p);
1822  ps = p;
1823  for(int k = 0; k < _levelFeatureNum[idx]; k++, ps+=4)
1824  {
1825  ps[0] = os*(ps[0]-0.5f) + offset; //x
1826  ps[1] = os*(ps[1]-0.5f) + offset; //y
1827  sigma = os*ps[3];
1828  ps[3] = (float)fmod(twopi-ps[2], twopi); //orientation, mirrored
1829  ps[2] = sigma; //scale
1830  }
1831  p+= 4* _levelFeatureNum[idx];
1832  }
1833  }
1834  }
1835 
1836  //put the feature into their original order
1837 
1838  if(_keypoint_index.size() > 0)
1839  {
1840  for(int i = 0; i < _featureNum; ++i)
1841  {
1842  int index = _keypoint_index[i];
1843  memcpy(&_keypoint_buffer[index*4], &keypoint_buffer2[i*4], 4 * sizeof(float));
1844  }
1845  }
1846 }
1847 
1848 
1850 {
1851  //generate feature list texture from existing keypoints
1852  //do feature sorting in the same time?
1853 
1854  FrameBufferObject fbo;
1855  vector<float> list;
1856  int idx = 0;
1857  const double twopi = 2.0*3.14159265358979323846;
1858  float sigma_half_step = powf(2.0f, 0.5f / param._dog_level_num);
1859  float octave_sigma = _octave_min>=0? float(1<<_octave_min): 1.0f/(1<<(-_octave_min));
1860  float offset = GlobalUtil::_LoweOrigin? 0 : 0.5f;
1861  if(_down_sample_factor>0) octave_sigma *= float(1<<_down_sample_factor);
1862 
1863 
1865 
1866  _keypoint_index.resize(0); // should already be 0
1867  for(int i = 0; i < _octave_num; i++, octave_sigma*= 2.0f)
1868  {
1869  for(int j = 0; j < param._dog_level_num; j++, idx++)
1870  {
1871  list.resize(0);
1872  float level_sigma = param.GetLevelSigma(j + param._level_min + 1) * octave_sigma;
1873  float sigma_min = level_sigma / sigma_half_step;
1874  float sigma_max = level_sigma * sigma_half_step;
1875  int fcount = 0 ;
1876  for(int k = 0; k < _featureNum; k++)
1877  {
1878  float * key = &_keypoint_buffer[k*4];
1879  float sigmak = key[2];
1880 
1882  if(IsUsingRectDescription()) sigmak = min(key[2], key[3]) / 12.0f;
1883 
1884  if( (sigmak >= sigma_min && sigmak < sigma_max)
1885  ||(sigmak < sigma_min && i ==0 && j == 0)
1886  ||(sigmak > sigma_max && j == param._dog_level_num - 1&&
1888  {
1889  //add this keypoint to the list
1890  list.push_back((key[0] - offset) / octave_sigma + 0.5f);
1891  list.push_back((key[1] - offset) / octave_sigma + 0.5f);
1893  {
1894  list.push_back(key[2] / octave_sigma);
1895  list.push_back(key[3] / octave_sigma);
1896  }else
1897  {
1898  list.push_back((float)fmod(twopi-key[3], twopi));
1899  list.push_back(key[2] / octave_sigma);
1900  }
1901  fcount ++;
1902  //save the index of keypoints
1903  _keypoint_index.push_back(k);
1904  }
1905  }
1906 
1907  _levelFeatureNum[idx] = fcount;
1908  if(fcount==0)continue;
1909  GLTexImage * ftex = _featureTex+idx;
1910 
1911  SetLevelFeatureNum(idx, fcount);
1912 
1913  int fw = ftex->GetImgWidth();
1914  int fh = ftex->GetImgHeight();
1915 
1916  list.resize(4*fh*fw);
1917 
1918  ftex->BindTex();
1919  ftex->AttachToFBO(0);
1920  glTexSubImage2D(GlobalUtil::_texTarget, 0, 0, 0, fw, fh, GL_RGBA, GL_FLOAT, &list[0]);
1921 
1922  if( fcount == _featureNum) _keypoint_index.resize(0);
1923  }
1925  }
1928  {
1929  std::cout<<"#Features:\t"<<_featureNum<<"\n";
1930  }
1931 }
1932 
1933 
1934 
1936 {
1937  _allPyramid = NULL;
1938 }
1939 
1941 {
1943 }
1944 
1945 
1946 //build the gaussian pyrmaid
1947 
1949 {
1950  //
1951  USE_TIMING();
1952  GLTexImage * tex, *tmp;
1953  FilterProgram ** filter;
1954  FrameBufferObject fbo;
1955 
1956  glDrawBuffer(GL_COLOR_ATTACHMENT0_EXT);
1957  input->FitTexViewPort();
1958 
1959  for (int i = _octave_min; i < _octave_min + _octave_num; i++)
1960  {
1961  tex = GetBaseLevel(i);
1962  tmp = GetBaseLevel(i, DATA_DOG) + 2; //use this as a temperory texture
1963 
1964 
1966 
1967  OCTAVE_START();
1968 
1969  if( i == _octave_min )
1970  {
1971  if(i < 0) TextureUpSample(tex, input, 1<<(-i-1));
1972  else TextureDownSample(tex, input, 1<<(i+1));
1973  ShaderMan::FilterInitialImage(tex, tmp);
1974  }else
1975  {
1977  ShaderMan::FilterSampledImage(tex, tmp);
1978  }
1979  LEVEL_FINISH();
1980 
1981  for(int j = param._level_min + 1; j <= param._level_max ; j++, tex++, filter++)
1982  {
1983  // filtering
1984  ShaderMan::FilterImage(*filter, tex+1, tex, tmp);
1985  LEVEL_FINISH();
1986  }
1987 
1988  OCTAVE_FINISH();
1989 
1990  }
1991  if(GlobalUtil::_timingS) glFinish();
1992  UnloadProgram();
1993 }
1994 
1996 {
1997 
1998  //first pass, compute dog, gradient, orientation
1999  GLenum buffers[4] = {
2000  GL_COLOR_ATTACHMENT0_EXT, GL_COLOR_ATTACHMENT1_EXT ,
2001  GL_COLOR_ATTACHMENT2_EXT, GL_COLOR_ATTACHMENT3_EXT
2002  };
2003 
2004  int i, j;
2005  double ts, t1;
2006  FrameBufferObject fbo;
2007 
2009 
2010  for(i = _octave_min; i < _octave_min + _octave_num; i++)
2011  {
2012  GLTexImage * gus = GetBaseLevel(i) + 1;
2013  GLTexImage * dog = GetBaseLevel(i, DATA_DOG) + 1;
2014  GLTexImage * grd = GetBaseLevel(i, DATA_GRAD) + 1;
2015  GLTexImage * rot = GetBaseLevel(i, DATA_ROT) + 1;
2016  glDrawBuffers(3, buffers);
2017  gus->FitTexViewPort();
2018  //compute the gradient
2019  for(j = 0; j < param._dog_level_num ; j++, gus++, dog++, grd++, rot++)
2020  {
2021  //gradient, dog, orientation
2022  glActiveTexture(GL_TEXTURE0);
2023  gus->BindTex();
2024  glActiveTexture(GL_TEXTURE1);
2025  (gus-1)->BindTex();
2026  //output
2027  dog->AttachToFBO(0);
2028  grd->AttachToFBO(1);
2029  rot->AttachToFBO(2);
2030  ShaderMan::UseShaderGradientPass((gus-1)->GetTexID());
2031  //compute
2032  dog->DrawQuadMT4();
2033  }
2034  }
2036  {
2037  glFinish();
2039  {
2040  t1 = CLOCK();
2041  std::cout <<"<Gradient, DOG >\t"<<(t1-ts)<<"\n";
2042  }
2043  }
2046  UnloadProgram();
2048  fbo.UnattachTex(GL_COLOR_ATTACHMENT1_EXT);
2049 }
2050 
2052 {
2053 
2054  //first pass, compute dog, gradient, orientation
2055  GLenum buffers[4] = {
2056  GL_COLOR_ATTACHMENT0_EXT, GL_COLOR_ATTACHMENT1_EXT ,
2057  GL_COLOR_ATTACHMENT2_EXT, GL_COLOR_ATTACHMENT3_EXT
2058  };
2059 
2060  int i, j;
2061  double t0, t, ts, t1, t2;
2062  FrameBufferObject fbo;
2063 
2065 
2066  for(i = _octave_min; i < _octave_min + _octave_num; i++)
2067  {
2068  GLTexImage * gus = GetBaseLevel(i) + 1;
2069  GLTexImage * dog = GetBaseLevel(i, DATA_DOG) + 1;
2070  GLTexImage * grd = GetBaseLevel(i, DATA_GRAD) + 1;
2071  GLTexImage * rot = GetBaseLevel(i, DATA_ROT) + 1;
2072  glDrawBuffers(3, buffers);
2073  gus->FitTexViewPort();
2074  //compute the gradient
2075  for(j = param._level_min +1; j <= param._level_max ; j++, gus++, dog++, grd++, rot++)
2076  {
2077  //gradient, dog, orientation
2078  glActiveTexture(GL_TEXTURE0);
2079  gus->BindTex();
2080  glActiveTexture(GL_TEXTURE1);
2081  (gus-1)->BindTex();
2082  //output
2083  dog->AttachToFBO(0);
2084  grd->AttachToFBO(1);
2085  rot->AttachToFBO(2);
2086  ShaderMan::UseShaderGradientPass((gus-1)->GetTexID());
2087  //compute
2088  dog->DrawQuadMT4();
2089  }
2090  }
2092  {
2093  glFinish();
2094  t1 = CLOCK();
2095  }
2098  //glDrawBuffers(1, buffers);
2099  glDrawBuffer(GL_COLOR_ATTACHMENT0_EXT);
2100 
2101 
2103 
2104  for ( i = _octave_min; i < _octave_min + _octave_num; i++)
2105  {
2107  {
2108  t0 = CLOCK();
2109  std::cout<<"#"<<(i + _down_sample_factor)<<"\t";
2110  }
2111  GLTexImage * dog = GetBaseLevel(i, DATA_DOG) + 2;
2112  GLTexImage * key = GetBaseLevel(i, DATA_KEYPOINT) +2;
2113  GLTexImage * gus = GetBaseLevel(i) + 2;
2114  key->FitTexViewPort();
2115 
2116  for( j = param._level_min +2; j < param._level_max ; j++, dog++, key++, gus++)
2117  {
2118  if(GlobalUtil::_timingL)t = CLOCK();
2119  key->AttachToFBO(0);
2120  glActiveTexture(GL_TEXTURE0);
2121  dog->BindTex();
2122  glActiveTexture(GL_TEXTURE1);
2123  (dog+1)->BindTex();
2124  glActiveTexture(GL_TEXTURE2);
2125  (dog-1)->BindTex();
2127  {
2128  glActiveTexture(GL_TEXTURE3);
2129  gus->BindTex();
2130  }
2131  ShaderMan::UseShaderKeypoint((dog+1)->GetTexID(), (dog-1)->GetTexID());
2132  key->DrawQuadMT8();
2134  {
2135  glFinish();
2136  std::cout<<(CLOCK()-t)<<"\t";
2137  }
2138  }
2140  {
2141  glFinish();
2142  std::cout<<"|\t"<<(CLOCK()-t0)<<"\n";
2143  }
2144  }
2145 
2147  {
2148  glFinish();
2150  {
2151  t2 = CLOCK();
2152  std::cout <<"<Gradient, DOG >\t"<<(t1-ts)<<"\n"
2153  <<"<Get Keypoints >\t"<<(t2-t1)<<"\n";
2154  }
2155 
2156  }
2157  UnloadProgram();
2159  fbo.UnattachTex(GL_COLOR_ATTACHMENT1_EXT);
2160 }
2161 
2162 
2163 void PyramidPacked::GenerateFeatureList(int i, int j)
2164 {
2165  float fcount = 0.0f;
2166  int hist_skip_gpu = GlobalUtil::_ListGenSkipGPU;
2167  int idx = i * param._dog_level_num + j;
2168  int hist_level_num = _hpLevelNum - _pyramid_octave_first;
2169  GLTexImage * htex, * ftex, * tex;
2170  htex = _histoPyramidTex + hist_level_num - 1 - i;
2171  ftex = _featureTex + idx;
2172  tex = GetBaseLevel(_octave_min + i, DATA_KEYPOINT) + 2 + j;
2173 
2174 
2175  //fill zero to an extra row/col if the height/width is odd
2176  glActiveTexture(GL_TEXTURE0);
2177  tex->BindTex();
2178  htex->AttachToFBO(0);
2179  int tight = (htex->GetImgWidth() * 4 == tex->GetImgWidth() -1 && htex->GetImgHeight() *4 == tex->GetImgHeight()-1 );
2181  htex->FitTexViewPort();
2182  //this uses the fact that no feature is on the edge.
2183  htex->DrawQuadReduction();
2184  //reduction..
2185  htex--;
2186 
2187  //this part might have problems on several GPUS
2188  //because the output of one pass is the input of the next pass
2189  //may require glFinish to make it right, but too much glFinish makes it slow
2190  for(int k = 0; k <hist_level_num - i-1 - hist_skip_gpu; k++, htex--)
2191  {
2192  htex->AttachToFBO(0);
2193  htex->FitTexViewPort();
2194  (htex+1)->BindTex();
2196  htex->DrawQuadReduction();
2197  }
2198 
2199  if(hist_skip_gpu == 0)
2200  {
2201  //read back one pixel
2202  float fn[4];
2203  glReadPixels(0, 0, 1, 1, GL_RGBA , GL_FLOAT, fn);
2204  fcount = (fn[0] + fn[1] + fn[2] + fn[3]);
2205  if(fcount < 1) fcount = 0;
2206 
2207  _levelFeatureNum[ idx] = (int)(fcount);
2208  SetLevelFeatureNum(idx, (int)fcount);
2209 
2210  //save number of features
2211  _featureNum += int(fcount);
2212 
2213  //
2214  if(fcount < 1.0) return;;
2215 
2216 
2218  htex= _histoPyramidTex;
2219 
2220  htex->BindTex();
2221 
2222  //first pass
2223  ftex->AttachToFBO(0);
2225  {
2226  //this is very important...
2227  ftex->FitRealTexViewPort();
2228  glClear(GL_COLOR_BUFFER_BIT);
2229  glFinish();
2230  }else
2231  {
2232  ftex->FitTexViewPort();
2233  //glFinish();
2234  }
2235 
2236 
2237  ShaderMan::UseShaderGenListStart((float)ftex->GetImgWidth(), htex->GetTexID());
2238 
2239  ftex->DrawQuad();
2240  //make sure it finishes before the next step
2241  ftex->DetachFBO(0);
2242  //pass on each pyramid level
2243  htex++;
2244  }else
2245  {
2246 
2247  int tw = htex[1].GetDrawWidth(), th = htex[1].GetDrawHeight();
2248  int fc = 0;
2249  glReadPixels(0, 0, tw, th, GL_RGBA , GL_FLOAT, _histo_buffer);
2250  _keypoint_buffer.resize(0);
2251  for(int y = 0, pos = 0; y < th; y++)
2252  {
2253  for(int x= 0; x < tw; x++)
2254  {
2255  for(int c = 0; c < 4; c++, pos++)
2256  {
2257  int ss = (int) _histo_buffer[pos];
2258  if(ss == 0) continue;
2259  float ft[4] = {2 * x + (c%2? 1.5f: 0.5f), 2 * y + (c>=2? 1.5f: 0.5f), 0, 1 };
2260  for(int t = 0; t < ss; t++)
2261  {
2262  ft[2] = (float) t;
2263  _keypoint_buffer.insert(_keypoint_buffer.end(), ft, ft+4);
2264  }
2265  fc += (int)ss;
2266  }
2267  }
2268  }
2269  _levelFeatureNum[ idx] = fc;
2270  SetLevelFeatureNum(idx, fc);
2271  if(fc == 0) return;
2272 
2273  fcount = (float) fc;
2274  _featureNum += fc;
2276  ftex->AttachToFBO(0);
2278  {
2279  ftex->FitRealTexViewPort();
2280  glClear(GL_COLOR_BUFFER_BIT);
2281  glFlush();
2282  }else
2283  {
2284  ftex->FitTexViewPort();
2285  glFlush();
2286  }
2287  _keypoint_buffer.resize(ftex->GetDrawWidth() * ftex->GetDrawHeight()*4, 0);
2289  glActiveTexture(GL_TEXTURE0);
2290  ftex->BindTex();
2291  glTexSubImage2D(GlobalUtil::_texTarget, 0, 0, 0, ftex->GetDrawWidth(),
2292  ftex->GetDrawHeight(), GL_RGBA, GL_FLOAT, &_keypoint_buffer[0]);
2293  htex += 2;
2294  }
2295 
2296  for(int lev = 1 + hist_skip_gpu; lev < hist_level_num - i; lev++, htex++)
2297  {
2298  glActiveTexture(GL_TEXTURE0);
2299  ftex->BindTex();
2300  ftex->AttachToFBO(0);
2301  glActiveTexture(GL_TEXTURE1);
2302  htex->BindTex();
2304  ftex->DrawQuad();
2305  ftex->DetachFBO(0);
2306  }
2307 
2308  ftex->AttachToFBO(0);
2309  glActiveTexture(GL_TEXTURE1);
2310  tex->BindTex();
2312  ftex->DrawQuad();
2314 
2315 }
2316 
2318 {
2319  //generate the histogram pyramid
2320  FrameBufferObject fbo;
2321  glReadBuffer(GL_COLOR_ATTACHMENT0_EXT);
2322  glDrawBuffer(GL_COLOR_ATTACHMENT0_EXT);
2323  double t1, t2;
2324  int ocount= 0, reverse = (GlobalUtil::_TruncateMethod == 1);
2325  _featureNum = 0;
2326 
2328  //for(int i = 0, idx = 0; i < _octave_num; i++)
2329  FOR_EACH_OCTAVE(i, reverse)
2330  {
2332  {
2333  t1= CLOCK();
2334  ocount = 0;
2335  std::cout<<"#"<<i+_octave_min + _down_sample_factor<<":\t";
2336  }
2337  //for(int j = 0; j < param._dog_level_num; j++, idx++)
2338  FOR_EACH_LEVEL(j, reverse)
2339  {
2342  {
2343  _levelFeatureNum[i * param._dog_level_num + j] = 0;
2344  continue;
2345  }
2346 
2347  GenerateFeatureList(i, j);
2348 
2350  {
2351  int idx = i * param._dog_level_num + j;
2352  ocount += _levelFeatureNum[idx];
2353  std::cout<< _levelFeatureNum[idx] <<"\t";
2354  }
2355  }
2357  {
2358  t2 = CLOCK();
2359  std::cout << "| \t" << int(ocount) << " :\t(" << (t2 - t1) << ")\n";
2360  }
2361  }
2362  if(GlobalUtil::_timingS)glFinish();
2364  {
2365  std::cout<<"#Features:\t"<<_featureNum<<"\n";
2366  }
2367 
2368 }
2369 
2371 {
2372  FrameBufferObject fbo;
2373  _featureNum = 0;
2375  float * mem = new float [tex->GetTexWidth()*tex->GetTexHeight()*4];
2376  vector<float> list;
2377  int idx = 0;
2378  for(int i = 0; i < _octave_num; i++)
2379  {
2380  for(int j = 0; j < param._dog_level_num; j++, idx++)
2381  {
2382  tex = GetBaseLevel(_octave_min + i, DATA_KEYPOINT) + j + 2;
2383  tex->BindTex();
2384  glGetTexImage(GlobalUtil::_texTarget, 0, GL_RGBA, GL_FLOAT, mem);
2385  //tex->AttachToFBO(0);
2386  //tex->FitTexViewPort();
2387  //glReadPixels(0, 0, tex->GetTexWidth(), tex->GetTexHeight(), GL_RED, GL_FLOAT, mem);
2388  //
2389  //make a list of
2390  list.resize(0);
2391  float *pl = mem;
2392  int fcount = 0 ;
2393  for(int k = 0; k < tex->GetDrawHeight(); k++)
2394  {
2395  float * p = pl;
2396  pl += tex->GetTexWidth() * 4;
2397  for( int m = 0; m < tex->GetDrawWidth(); m ++, p+=4)
2398  {
2399  // if(m ==0 || k ==0 || k == tex->GetDrawHeight() -1 || m == tex->GetDrawWidth() -1) continue;
2400  // if(*p == 0) continue;
2401  int t = ((int) fabs(p[0])) - 1;
2402  if(t < 0) continue;
2403  int xx = m + m + ( (t %2)? 1 : 0);
2404  int yy = k + k + ( (t <2)? 0 : 1);
2405  if(xx ==0 || yy == 0) continue;
2406  if(xx >= tex->GetImgWidth() - 1 || yy >= tex->GetImgHeight() - 1)continue;
2407  list.push_back(xx + 0.5f + p[1]);
2408  list.push_back(yy + 0.5f + p[2]);
2409  list.push_back(GlobalUtil::_KeepExtremumSign && p[0] < 0 ? -1.0f : 1.0f);
2410  list.push_back(p[3]);
2411  fcount ++;
2412  }
2413  }
2414  if(fcount==0)continue;
2415 
2416  if(GlobalUtil::_timingL) std::cout<<fcount<<".";
2417 
2418  GLTexImage * ftex = _featureTex+idx;
2419  _levelFeatureNum[idx] = (fcount);
2420  SetLevelFeatureNum(idx, fcount);
2421 
2422  _featureNum += (fcount);
2423 
2424 
2425  int fw = ftex->GetImgWidth();
2426  int fh = ftex->GetImgHeight();
2427 
2428  list.resize(4*fh*fw);
2429 
2430  ftex->BindTex();
2431  ftex->AttachToFBO(0);
2432  // glTexImage2D(GlobalUtil::_texTarget, 0, GlobalUtil::_iTexFormat, fw, fh, 0, GL_BGRA, GL_FLOAT, &list[0]);
2433  glTexSubImage2D(GlobalUtil::_texTarget, 0, 0, 0, fw, fh, GL_RGBA, GL_FLOAT, &list[0]);
2434  //
2435  }
2436  }
2438  delete[] mem;
2440  {
2441  std::cout<<"#Features:\t"<<_featureNum<<"\n";
2442  }
2443 }
2444 
2445 
2446 
2448 {
2449  GLTexImage * gtex, * otex;
2450  GLTexImage * ftex = _featureTex;
2451  GLTexImage * fotex = _orientationTex;
2452  int * count = _levelFeatureNum;
2453  float sigma, sigma_step = powf(2.0f, 1.0f/param._dog_level_num);
2454 
2455 
2456  FrameBufferObject fbo;
2457  if(_orientationTex)
2458  {
2459  GLenum buffers[] = { GL_COLOR_ATTACHMENT0_EXT, GL_COLOR_ATTACHMENT1_EXT };
2460  glDrawBuffers(2, buffers);
2461  }else
2462  {
2463  glDrawBuffer(GL_COLOR_ATTACHMENT0_EXT);
2464  }
2465 
2466  for(int i = 0; i < _octave_num; i++)
2467  {
2468  gtex = GetBaseLevel(i+_octave_min, DATA_GRAD) + 1;
2469  otex = GetBaseLevel(i+_octave_min, DATA_ROT) + 1;
2470 
2471  for(int j = 0; j < param._dog_level_num; j++, ftex++, otex++, count++, gtex++, fotex++)
2472  {
2473  if(*count<=0)continue;
2474 
2475  sigma = param.GetLevelSigma(j+param._level_min+1);
2476 
2477 
2478  ftex->FitTexViewPort();
2479 
2480  glActiveTexture(GL_TEXTURE0);
2481  ftex->BindTex();
2482  glActiveTexture(GL_TEXTURE1);
2483  gtex->BindTex();
2484  glActiveTexture(GL_TEXTURE2);
2485  otex->BindTex();
2486  //
2487  ftex->AttachToFBO(0);
2488  if(_orientationTex) fotex->AttachToFBO(1);
2489 
2491 
2492  ShaderMan::UseShaderOrientation(gtex->GetTexID(),
2493  gtex->GetImgWidth(), gtex->GetImgHeight(),
2494  sigma, otex->GetTexID(), sigma_step, _existing_keypoints);
2495  ftex->DrawQuad();
2496  }
2497  }
2498 
2500  if(GlobalUtil::_timingS)glFinish();
2501  if(_orientationTex) fbo.UnattachTex(GL_COLOR_ATTACHMENT1_EXT);
2502 
2503 }
2504 
2505 
2507 {
2508  //
2509  int idx = 0;
2510 // int n = _octave_num * param._dog_level_num;
2511  float sigma, sigma_step = powf(2.0f, 1.0f/param._dog_level_num);
2512  GLTexImage * ftex = _featureTex;
2513 
2514  FrameBufferObject fbo;
2515  glDrawBuffer(GL_COLOR_ATTACHMENT0_EXT);
2516  for(int i = 0; i < _octave_num; i++)
2517  {
2518  GLTexImage *otex = GetBaseLevel(i + _octave_min, DATA_ROT) + 2;
2519  for(int j = 0; j < param._dog_level_num; j++, ftex++, otex++, idx ++)
2520  {
2521  if(_levelFeatureNum[idx]<=0)continue;
2522  sigma = param.GetLevelSigma(j+param._level_min+1);
2523  //
2524  ftex->AttachToFBO(0);
2525  ftex->FitTexViewPort();
2526 
2527  glActiveTexture(GL_TEXTURE0);
2528  ftex->BindTex();
2529  glActiveTexture(GL_TEXTURE1);
2530  otex->BindTex();
2531 
2532  ShaderMan::UseShaderSimpleOrientation(otex->GetTexID(), sigma, sigma_step);
2533  ftex->DrawQuad();
2534  }
2535  }
2537 }
2538 
2539 void PyramidPacked::InitPyramid(int w, int h, int ds)
2540 {
2541  int wp, hp, toobig = 0;
2542  if(ds == 0)
2543  {
2544  _down_sample_factor = 0;
2546  {
2549  }else
2550  {
2551  wp = w << (-GlobalUtil::_octave_min_default);
2552  hp = h << (-GlobalUtil::_octave_min_default);
2553  }
2555  }else
2556  {
2557  //must use 0 as _octave_min;
2558  _octave_min = 0;
2559  _down_sample_factor = ds;
2560  w >>= ds;
2561  h >>= ds;
2562  wp = w;
2563  hp = h;
2564  }
2565 
2566  while(wp > GlobalUtil::_texMaxDim || hp > GlobalUtil::_texMaxDim )
2567  {
2568  _octave_min ++;
2569  wp >>= 1;
2570  hp >>= 1;
2571  toobig = 1;
2572  }
2573 
2575  max(max(wp, hp), max(_pyramid_width, _pyramid_height)) > 1024 * sqrt(GlobalUtil::_MemCapGPU / 96.0) )
2576  {
2577  _octave_min ++;
2578  wp >>= 1;
2579  hp >>= 1;
2580  toobig = 2;
2581  }
2582 
2583  if(toobig && GlobalUtil::_verbose)
2584  {
2585  std::cout<<(toobig == 2 ? "[**SKIP OCTAVES**]:\tExceeding Memory Cap (-nomc)\n" :
2586  "[**SKIP OCTAVES**]:\tReaching the dimension limit (-maxd)!\n");
2587  }
2588 
2589  if( wp == _pyramid_width && hp == _pyramid_height && _allocated )
2590  {
2591  FitPyramid(wp, hp);
2593  {
2594  ResizePyramid(wp, hp);
2595  }
2596  else if( wp > _pyramid_width || hp > _pyramid_height )
2597  {
2599  if(wp < _pyramid_width || hp < _pyramid_height) FitPyramid(wp, hp);
2600  }
2601  else
2602  {
2603  //try use the pyramid allocated for large image on small input images
2604  FitPyramid(wp, hp);
2605  }
2606 
2607  //select the initial smoothing filter according to the new _octave_min
2609 }
2610 
2611 
2612 
2613 void PyramidPacked::FitPyramid(int w, int h)
2614 {
2615  //(w, h) <= (_pyramid_width, _pyramid_height);
2616 
2618  //
2620 
2621  int _octave_num_max = GetRequiredOctaveNum(min(w, h));
2622 
2623  if(_octave_num < 1 || _octave_num > _octave_num_max)
2624  {
2625  _octave_num = _octave_num_max;
2626  }
2627 
2628 
2629  int pw = _pyramid_width>>1, ph = _pyramid_height>>1;
2631  pw >= w && ph >= h)
2632  {
2634  pw >>= 1;
2635  ph >>= 1;
2636  }
2637 
2638  for(int i = 0; i < _octave_num; i++)
2639  {
2640  GLTexImage * tex = GetBaseLevel(i + _octave_min);
2645  for(int j = param._level_min; j <= param._level_max; j++, tex++, dog++, grd++, rot++, key++)
2646  {
2647  tex->SetImageSize(w, h);
2648  if(j == param._level_min) continue;
2649  dog->SetImageSize(w, h);
2650  grd->SetImageSize(w, h);
2651  rot->SetImageSize(w, h);
2652  if(j == param._level_min + 1 || j == param._level_max) continue;
2653  key->SetImageSize(w, h);
2654  }
2655  w>>=1;
2656  h>>=1;
2657  }
2658 }
2659 
2660 
2661 void PyramidPacked::ResizePyramid( int w, int h)
2662 {
2663  //
2664  unsigned int totalkb = 0;
2665  int _octave_num_new, input_sz, i, j;
2666  //
2667 
2668  if(_pyramid_width == w && _pyramid_height == h && _allocated) return;
2669 
2670  if(w > GlobalUtil::_texMaxDim || h > GlobalUtil::_texMaxDim) return ;
2671 
2672  if(GlobalUtil::_verbose && GlobalUtil::_timingS) std::cout<<"[Allocate Pyramid]:\t" <<w<<"x"<<h<<endl;
2673  //first octave does not change
2675 
2676 
2677  //compute # of octaves
2678 
2679  input_sz = min(w,h) ;
2680  _pyramid_width = w;
2681  _pyramid_height = h;
2682 
2683  //reset to preset parameters
2684  _octave_num_new = GlobalUtil::_octave_num_default;
2685 
2686  if(_octave_num_new < 1) _octave_num_new = GetRequiredOctaveNum(input_sz) ;
2687 
2688 
2689  if(_pyramid_octave_num != _octave_num_new)
2690  {
2691  //destroy the original pyramid if the # of octave changes
2692  if(_octave_num >0)
2693  {
2696  }
2697  _pyramid_octave_num = _octave_num_new;
2698  }
2699 
2701 
2702  int noct = _octave_num;
2703  int nlev = param._level_num;
2704 
2705  // //initialize the pyramid
2706  if(_allPyramid==NULL) _allPyramid = new GLTexPacked[ noct* nlev * DATA_NUM];
2707 
2708 
2714 
2715 
2717 
2718  for(i = 0; i< noct; i++)
2719  {
2720  for( j = 0; j< nlev; j++, gus++, dog++, grd++, rot++, key++)
2721  {
2722  gus->InitTexture(w, h);
2723  if(j==0)continue;
2724  dog->InitTexture(w, h);
2725  grd->InitTexture(w, h, 0);
2726  rot->InitTexture(w, h);
2727  if(j<=1 || j >=nlev -1) continue;
2728  key->InitTexture(w, h, 0);
2729  }
2730  int tsz = (gus -1)->GetTexPixelCount() * 16;
2731  totalkb += ((nlev *5 -6)* tsz / 1024);
2732  //several auxilary textures are not actually required
2733  w>>=1;
2734  h>>=1;
2735  }
2736 
2737  totalkb += ResizeFeatureStorage();
2738 
2739  _allocated = 1;
2740 
2741  if(GlobalUtil::_verbose && GlobalUtil::_timingS) std::cout<<"[Allocate Pyramid]:\t" <<(totalkb/1024)<<"MB\n";
2742 
2743 }
2744 
2746 {
2747  if(_allPyramid)
2748  {
2749  delete [] _allPyramid;
2750  _allPyramid = NULL;
2751  }
2752 }
2753 
2754 
2755 GLTexImage* PyramidPacked::GetLevelTexture(int octave, int level, int dataName)
2756 {
2757  return _allPyramid+ (_pyramid_octave_first + octave - _octave_min) * param._level_num
2758  + param._level_num * _pyramid_octave_num * dataName
2759  + (level - param._level_min);
2760 
2761 }
2762 
2764 {
2765  return _allPyramid+ (_pyramid_octave_first + octave - _octave_min) * param._level_num
2766  + (level - param._level_min);
2767 }
2768 
2769 //in the packed implementation( still in progress)
2770 // DATA_GAUSSIAN, DATA_DOG, DATA_GAD will be stored in different textures.
2771 
2772 GLTexImage* PyramidPacked::GetBaseLevel(int octave, int dataName)
2773 {
2774  if(octave <_octave_min || octave > _octave_min + _octave_num) return NULL;
2776  int num = param._level_num * _pyramid_octave_num;
2777  return _allPyramid + num *dataName + offset;
2778 }
2779 
2780 
2782 {
2783  GLTexImage * tex, *htex;
2784  int hist_level_num = _hpLevelNum - _pyramid_octave_first;
2785 
2786  tex = GetBaseLevel(_octave_min , DATA_KEYPOINT) + 2;
2787  htex = _histoPyramidTex + hist_level_num - 1;
2788  int w = (tex->GetImgWidth() + 2) >> 2;
2789  int h = (tex->GetImgHeight() + 2)>> 2;
2790 
2791 
2792  //4n+1 -> n; 4n+2,2, 3 -> n+1
2793  for(int k = 0; k <hist_level_num -1; k++, htex--)
2794  {
2795  if(htex->GetImgHeight()!= h || htex->GetImgWidth() != w)
2796  {
2797  htex->SetImageSize(w, h);
2798  htex->ZeroHistoMargin();
2799  }
2800 
2801  w = (w + 1)>>1; h = (h + 1) >> 1;
2802  }
2803 }
2804 
int count
int offset
#define USE_TIMING()
Definition: PyramidGL.cpp:50
#define LEVEL_FINISH()
Definition: PyramidGL.cpp:52
#define OCTAVE_START()
Definition: PyramidGL.cpp:51
#define OCTAVE_FINISH()
Definition: PyramidGL.cpp:53
#define FOR_EACH_OCTAVE(i, R)
Definition: SiftPyramid.h:183
#define FOR_EACH_LEVEL(j, R)
Definition: SiftPyramid.h:184
#define NULL
static void DeleteGlobalFBO()
static void UnattachTex(GLenum attachment)
void DrawQuadReduction()
Definition: GLTexImage.cpp:464
int GetImgHeight()
Definition: GLTexImage.h:80
virtual void InitTexture(int width, int height, int clamp_to_edge=1)
Definition: GLTexImage.cpp:94
int GetDrawHeight()
Definition: GLTexImage.h:84
void FitTexViewPort()
Definition: GLTexImage.cpp:353
int GetTexWidth()
Definition: GLTexImage.h:81
GLuint GetTexID()
Definition: GLTexImage.h:76
virtual void SetImageSize(int width, int height)
Definition: GLTexImage.cpp:88
void DrawQuadMT8()
Definition: GLTexImage.cpp:294
void BindTex()
Definition: GLTexImage.cpp:159
static void DetachFBO(int i)
Definition: GLTexImage.cpp:368
void DrawQuadMT4()
Definition: GLTexImage.cpp:257
static void DrawQuad(float x1, float x2, float y1, float y2)
Definition: GLTexImage.cpp:374
static void UnbindTex()
Definition: GLTexImage.cpp:164
int GetDrawWidth()
Definition: GLTexImage.h:83
virtual void ZeroHistoMargin()
Definition: GLTexImage.cpp:196
int GetTexHeight()
Definition: GLTexImage.h:82
int GetImgWidth()
Definition: GLTexImage.h:79
void FitRealTexViewPort()
Definition: GLTexImage.cpp:358
void DrawQuad()
Definition: GLTexImage.cpp:170
void AttachToFBO(int i)
Definition: GLTexImage.cpp:363
static void UnbindMultiTex(int n)
Definition: GLTexImage.cpp:724
virtual void InitTexture(int width, int height, int clamp_to_edge=1)
Definition: GLTexImage.cpp:534
static int _TruncateMethod
Definition: GlobalUtil.h:74
static int _octave_num_default
Definition: GlobalUtil.h:79
static int _UseSiftGPUEX
Definition: GlobalUtil.h:76
static int _SubpixelLocalization
Definition: GlobalUtil.h:72
static int _octave_min_default
Definition: GlobalUtil.h:78
static int _timingL
Definition: GlobalUtil.h:47
static int _OrientationPack2
Definition: GlobalUtil.h:60
static int _DescriptorPPR
Definition: GlobalUtil.h:68
static int _verbose
Definition: GlobalUtil.h:44
static int _FeatureCountThreshold
Definition: GlobalUtil.h:90
static int _LoweOrigin
Definition: GlobalUtil.h:85
static int _NarrowFeatureTex
Definition: GlobalUtil.h:71
static int _ForceTightPyramid
Definition: GlobalUtil.h:77
static int _texMaxDim
Definition: GlobalUtil.h:39
static int _KeepExtremumSign
Definition: GlobalUtil.h:89
static GLuint _texTarget
Definition: GlobalUtil.h:37
static int _timingS
Definition: GlobalUtil.h:45
static float _MaxFeaturePercent
Definition: GlobalUtil.h:66
static int _DarknessAdaption
Definition: GlobalUtil.h:92
static int _FitMemoryCap
Definition: GlobalUtil.h:43
static int _MaxLevelFeatureNum
Definition: GlobalUtil.h:67
static int _MemCapGPU
Definition: GlobalUtil.h:42
static int _NormalizedSIFT
Definition: GlobalUtil.h:87
static int _ListGenGPU
Definition: GlobalUtil.h:61
static int _FeatureTexBlock
Definition: GlobalUtil.h:70
static int _IsNvidia
Definition: GlobalUtil.h:49
static int _DescriptorPPT
Definition: GlobalUtil.h:69
static int _KeyPointListForceLevel0
Definition: GlobalUtil.h:91
static int _ListGenSkipGPU
Definition: GlobalUtil.h:62
static int _MaxOrientation
Definition: GlobalUtil.h:59
static int _timingO
Definition: GlobalUtil.h:46
static int _GoodOpenGL
Definition: GlobalUtil.h:83
static bool CheckFramebufferStatus()
Definition: GlobalUtil.cpp:250
static void CleanupOpenGL()
Definition: GlobalUtil.cpp:156
static void CheckErrorsGL(const char *location=NULL)
Definition: GlobalUtil.h:137
static void FitViewPort(int width, int height)
Definition: GlobalUtil.cpp:234
static double CLOCK()
Definition: GlobalUtil.h:126
static void InitGLParam(int NotTargetGL=0)
Definition: GlobalUtil.cpp:319
PyramidGL(SiftParam &sp)
Definition: PyramidGL.cpp:159
void SetLevelFeatureNum(int idx, int num)
Definition: PyramidGL.cpp:1439
GLTexImage * _orientationTex
Definition: PyramidGL.h:41
virtual GLTexImage * GetBaseLevel(int octave, int dataName=DATA_GAUSSIAN)=0
virtual void GetFeatureDescriptors()
Definition: PyramidGL.cpp:1667
virtual void DownloadKeypoints()
Definition: PyramidGL.cpp:1790
virtual void GenerateFeatureListTex()
Definition: PyramidGL.cpp:1849
virtual void ReshapeFeatureListCPU()
Definition: PyramidGL.cpp:1263
virtual int ResizeFeatureStorage()
Definition: PyramidGL.cpp:467
virtual void DestroyPerLevelData()
Definition: PyramidGL.cpp:185
GLTexImage * _histoPyramidTex
Definition: PyramidGL.h:38
void GetAlignedStorageSize(int num, int align, int &fw, int &fh)
Definition: PyramidGL.cpp:71
virtual void GenerateFeatureDisplayVBO()
Definition: PyramidGL.cpp:1022
virtual ~PyramidGL()
Definition: PyramidGL.cpp:169
virtual void CleanUpAfterSIFT()
Definition: PyramidGL.cpp:1459
static void InterlaceDescriptorF2(int w, int h, float *buf, float *pd, int step)
Definition: PyramidGL.cpp:1583
static void NormalizeDescriptor(int num, float *pd)
Definition: PyramidGL.cpp:1541
GLTexImage * _featureTex
Definition: PyramidGL.h:39
virtual void DestroySharedData()
Definition: PyramidGL.cpp:223
void GetTextureStorageSize(int num, int &fw, int &fh)
Definition: PyramidGL.cpp:106
GLTexImage * _descriptorTex
Definition: PyramidGL.h:40
void InitializeContext()
Definition: PyramidGL.cpp:176
GLTexImage * _auxPyramid
Definition: PyramidGL.h:69
GLTexImage * GetLevelTexture(int octave, int level)
Definition: PyramidGL.cpp:654
void ComputeGradient()
Definition: PyramidGL.cpp:687
void ResizePyramid(int w, int h)
Definition: PyramidGL.cpp:380
virtual void GenerateFeatureList()
Definition: PyramidGL.cpp:966
GLTexImage * GetBaseLevel(int octave, int dataName=DATA_GAUSSIAN)
Definition: PyramidGL.cpp:662
virtual void GetFeatureOrientations()
Definition: PyramidGL.cpp:1126
void GenerateFeatureListCPU()
Definition: PyramidGL.cpp:1191
GLTexImage * _texPyramid
Definition: PyramidGL.h:68
void InitPyramid(int w, int h, int ds)
Definition: PyramidGL.cpp:307
void DestroyPyramidData()
Definition: PyramidGL.cpp:146
void FitHistogramPyramid()
Definition: PyramidGL.cpp:247
void BuildPyramid(GLTexInput *input)
Definition: PyramidGL.cpp:589
void FitPyramid(int w, int h)
Definition: PyramidGL.cpp:269
PyramidNaive(SiftParam &sp)
Definition: PyramidGL.cpp:59
void GetSimplifiedOrientation()
Definition: PyramidGL.cpp:1467
void DetectKeypointsEX()
Definition: PyramidGL.cpp:727
GLTexImage * GetLevelTexture(int octave, int level)
Definition: PyramidGL.cpp:2763
PyramidPacked(SiftParam &sp)
Definition: PyramidGL.cpp:1935
void FitPyramid(int w, int h)
Definition: PyramidGL.cpp:2613
void GenerateFeatureListCPU()
Definition: PyramidGL.cpp:2370
void DetectKeypointsEX()
Definition: PyramidGL.cpp:2051
void DestroyPyramidData()
Definition: PyramidGL.cpp:2745
void GetFeatureOrientations()
Definition: PyramidGL.cpp:2447
void ComputeGradient()
Definition: PyramidGL.cpp:1995
void GetSimplifiedOrientation()
Definition: PyramidGL.cpp:2506
void FitHistogramPyramid()
Definition: PyramidGL.cpp:2781
void GenerateFeatureList()
Definition: PyramidGL.cpp:2317
void ResizePyramid(int w, int h)
Definition: PyramidGL.cpp:2661
GLTexImage * GetBaseLevel(int octave, int dataName=DATA_GAUSSIAN)
Definition: PyramidGL.cpp:2772
void InitPyramid(int w, int h, int ds)
Definition: PyramidGL.cpp:2539
void BuildPyramid(GLTexInput *input)
Definition: PyramidGL.cpp:1948
FilterProgram ** f_gaussian_step
Definition: ProgramGLSL.h:144
static void UseShaderGenListHisto()
Definition: ShaderMan.cpp:266
static void UseShaderGenListEnd(int ktex)
Definition: ShaderMan.cpp:287
static void FilterSampledImage(GLTexImage *tex, GLTexImage *buf)
Definition: ShaderMan.cpp:134
static void FilterInitialImage(GLTexImage *tex, GLTexImage *buf)
Definition: ShaderMan.cpp:129
static void UseShaderGenListStart(float fw, int tex0)
Definition: ShaderMan.cpp:275
static void UseShaderCopyKeypoint()
Definition: ShaderMan.cpp:314
static void InitShaderMan(SiftParam &param)
Definition: ShaderMan.cpp:48
static void FilterImage(FilterProgram *filter, GLTexImage *dst, GLTexImage *src, GLTexImage *tmp)
Definition: ShaderMan.cpp:76
static void DestroyShaders()
Definition: ShaderMan.cpp:65
static void UnloadProgram()
Definition: ShaderMan.cpp:71
static void UseShaderGenListInit(int w, int h, int tight=1)
Definition: ShaderMan.cpp:254
static void UseShaderKeypoint(int texU, int texD)
Definition: ShaderMan.cpp:246
static void UseShaderGenVBO(float width, float fwidth, float size)
Definition: ShaderMan.cpp:303
static void TextureUpSample(GLTexImage *dst, GLTexImage *src, int scale)
Definition: ShaderMan.cpp:177
static void SelectInitialSmoothingFilter(int octave_min, SiftParam &param)
Definition: ShaderMan.cpp:346
static void UseShaderDescriptor(int gtex, int otex, int dwidth, int fwidth, int width, int height, float sigma)
Definition: ShaderMan.cpp:340
static void UseShaderGradientPass(int texP=0)
Definition: ShaderMan.cpp:239
static void TextureDownSample(GLTexImage *dst, GLTexImage *src, int scale=2)
Definition: ShaderMan.cpp:155
static void UseShaderOrientation(int gtex, int width, int height, float sigma, int auxtex, float step, int keypoint_list)
Definition: ShaderMan.cpp:327
static void UseShaderSimpleOrientation(int oTex, float sigma, float sigma_step)
Definition: ShaderMan.cpp:319
static ShaderBag * s_bag
Definition: ShaderMan.h:43
static void UseShaderGenListStep(int tex, int tex0)
Definition: ShaderMan.cpp:281
int _level_num
Definition: SiftGPU.h:86
int _level_ds
Definition: SiftGPU.h:91
int _level_min
Definition: SiftGPU.h:89
int _dog_level_num
Definition: SiftGPU.h:85
int _level_max
Definition: SiftGPU.h:90
float GetLevelSigma(int lev)
Definition: SiftGPU.cpp:1190
int _octave_num
Definition: SiftPyramid.h:85
int _pyramid_width
Definition: SiftPyramid.h:89
int _octave_min
Definition: SiftPyramid.h:83
int _pyramid_octave_num
Definition: SiftPyramid.h:87
int _pyramid_octave_first
Definition: SiftPyramid.h:88
SiftParam & param
Definition: SiftPyramid.h:67
int _featureNum
Definition: SiftPyramid.h:70
int _existing_keypoints
Definition: SiftPyramid.h:73
GLuint * _featurePointVBO
Definition: SiftPyramid.h:77
int _pyramid_height
Definition: SiftPyramid.h:90
int * _levelFeatureNum
Definition: SiftPyramid.h:69
GLuint * _featureDisplayVBO
Definition: SiftPyramid.h:76
static int GetRequiredOctaveNum(int inputsz)
int _allocated
Definition: SiftPyramid.h:92
int _down_sample_factor
Definition: SiftPyramid.h:91
float * _histo_buffer
Definition: SiftPyramid.h:71
vector< int > _keypoint_index
Definition: SiftPyramid.h:74
virtual int IsUsingRectDescription()
Definition: SiftPyramid.h:127
int _hpLevelNum
Definition: SiftPyramid.h:68
vector< float > _descriptor_buffer
Definition: SiftPyramid.h:97
int _alignment
Definition: SiftPyramid.h:93
vector< float > _keypoint_buffer
Definition: SiftPyramid.h:96
__host__ __device__ float2 fabs(float2 v)
Definition: cutil_math.h:1254
int min(int a, int b)
Definition: cutil_math.h:53
int max(int a, int b)
Definition: cutil_math.h:48
QTextStream & endl(QTextStream &stream)
Definition: QtCompat.h:718
MiniVec< float, N > ceil(const MiniVec< float, N > &a)
Definition: MiniVec.h:89
Definition: Eigen.h:85