@@ -29,24 +29,6 @@ namespace popsift {
2929
3030namespace gauss {
3131
32- __global__
33- void get_by_2_interpolate ( cudaTextureObject_t src_data,
34- const int src_level,
35- cudaSurfaceObject_t dst_data,
36- const int dst_w,
37- const int dst_h )
38- {
39- const int idx = blockIdx .x * blockDim .x + threadIdx .x ;
40- const int idy = blockIdx .y * blockDim .y + threadIdx .y ;
41-
42- if ( idx >= dst_w ) return ;
43- if ( idy >= dst_h ) return ;
44-
45- const float val = readTex ( src_data, 2 .0f * idx + 1 .0f , 2 .0f * idy + 1 .0f , src_level );
46-
47- surf2DLayeredwrite ( val, dst_data, idx*4 , idy, 0 , cudaBoundaryModeZero ); // dst_data.ptr(idy)[idx] = val;
48- }
49-
5032__global__
5133void get_by_2_pick_every_second ( cudaTextureObject_t src_data,
5234 const int src_w,
@@ -214,35 +196,17 @@ inline void Pyramid::downscale_from_prev_octave( int octave, cudaStream_t stream
214196 h_grid.x = (unsigned int )grid_divide ( width, h_block.x );
215197 h_grid.y = (unsigned int )grid_divide ( height, h_block.y );
216198
217- switch ( mode )
218- {
219- case Config::PopSift :
220- case Config::VLFeat :
221- case Config::OpenCV :
222- gauss::get_by_2_pick_every_second
223- <<<h_grid,h_block,0 ,stream>>>
224- ( prev_oct_obj.getDataTexPoint ( ),
225- prev_oct_obj.getWidth (),
226- prev_oct_obj.getHeight (),
227- _levels-PREV_LEVEL,
228- oct_obj.getDataSurface ( ),
229- oct_obj.getWidth (),
230- oct_obj.getHeight () );
231-
232- POP_SYNC_CHK;
233- break ;
234- default :
235- gauss::get_by_2_interpolate
236- <<<h_grid,h_block,0 ,stream>>>
237- ( prev_oct_obj.getDataTexLinear ( ).tex ,
238- _levels-PREV_LEVEL,
239- oct_obj.getDataSurface ( ),
240- oct_obj.getWidth (),
241- oct_obj.getHeight () );
242-
243- POP_SYNC_CHK;
244- break ;
245- }
199+ gauss::get_by_2_pick_every_second
200+ <<<h_grid,h_block,0 ,stream>>>
201+ ( prev_oct_obj.getDataTexPoint ( ),
202+ prev_oct_obj.getWidth (),
203+ prev_oct_obj.getHeight (),
204+ _levels-PREV_LEVEL,
205+ oct_obj.getDataSurface ( ),
206+ oct_obj.getWidth (),
207+ oct_obj.getHeight () );
208+
209+ POP_SYNC_CHK;
246210}
247211
248212__host__
0 commit comments