fix illegal access when there are more orientations than allocated

src/popsift/s_desc_grid.cu

@@ -123,10 +123,16 @@ void ext_desc_grid_sub( const int           ix,
 
 __global__ void ext_desc_grid(int octave, cudaTextureObject_t layer_tex)
 {
-    const int   o_offset =  dct.ori_ps[octave] + blockIdx.x;
+    const int   num      = dct.ori_ct[octave];
+    const int   offset   = blockIdx.x;
+
+    const int   o_offset =  dct.ori_ps[octave] + offset;
     const int   ix       = threadIdx.y;
     const int   iy       = threadIdx.z;
 
+    if( offset >= num ) return;
+    if( o_offset >= dct.ori_total ) return;
+
     Descriptor* desc     = &dbuf.desc           [o_offset];
     const int   ext_idx  =  dobuf.feat_to_ext_map[o_offset];
     Extremum*   ext      =  dobuf.extrema + ext_idx;

src/popsift/s_desc_igrid.cu

@@ -79,8 +79,10 @@ __global__ void ext_desc_igrid(int octave, cudaTextureObject_t texLinear)
     const int   num      = dct.ori_ct[octave];
 
     const int   offset   = blockIdx.x * blockDim.z + threadIdx.z;
+
     const int   o_offset =  dct.ori_ps[octave] + offset;
     if( offset >= num ) return;
+    if( o_offset >= dct.ori_total ) return;
 
     Descriptor* desc     = &dbuf.desc           [o_offset];
     const int   ext_idx  =  dobuf.feat_to_ext_map[o_offset];

src/popsift/s_desc_iloop.cu

@@ -130,7 +130,14 @@ void ext_desc_iloop_sub( const float         ang,
 
 __global__ void ext_desc_iloop(int octave, cudaTextureObject_t layer_tex, int w, int h)
 {
-    const int   o_offset =  dct.ori_ps[octave] + blockIdx.x;
+    const int   num      = dct.ori_ct[octave];
+
+    const int   offset   = blockIdx.x;
+
+    const int   o_offset =  dct.ori_ps[octave] + offset;
+    if( offset >= num ) return;
+    if( o_offset >= dct.ori_total ) return;
+
     Descriptor* desc     = &dbuf.desc           [o_offset];
     const int   ext_idx  =  dobuf.feat_to_ext_map[o_offset];
     Extremum*   ext      =  dobuf.extrema + ext_idx;

src/popsift/s_desc_loop.cu

@@ -140,7 +140,14 @@ void ext_desc_loop_sub( const float         ang,
 
 __global__ void ext_desc_loop(int octave, cudaTextureObject_t layer_tex, int w, int h)
 {
-    const int   o_offset =  dct.ori_ps[octave] + blockIdx.x;
+    const int   num      = dct.ori_ct[octave];
+
+    const int   offset   = blockIdx.x;
+
+    const int   o_offset =  dct.ori_ps[octave] + offset;
+    if( offset >= num ) return;
+    if( o_offset >= dct.ori_total ) return;
+
     Descriptor* desc     = &dbuf.desc           [o_offset];
     const int   ext_idx  =  dobuf.feat_to_ext_map[o_offset];
     Extremum*   ext      =  dobuf.extrema + ext_idx;

src/popsift/s_desc_notile.cu

@@ -105,6 +105,7 @@ void ext_desc_notile( const int           octave,
 
     const int   o_offset =  dct.ori_ps[octave] + offset;
     if( offset >= num ) return;
+    if( o_offset >= dct.ori_total ) return;
 
     Descriptor* desc     = &dbuf.desc            [o_offset];
     const int   ext_idx  =  dobuf.feat_to_ext_map[o_offset];

src/popsift/s_orientation.cu

@@ -327,7 +327,7 @@ void ori_prefix_sum( const int total_ext_ct, const int num_octaves )
     ExtremaRead r( extremum );
     ExtremaWrt  w( extremum );
     ExtremaTot  t( total_ori );
-    ExtremaWrtMap wrtm( feat_to_ext_map, max( d_consts.max_orientations, dbuf.ori_allocated ) );
+    ExtremaWrtMap wrtm( feat_to_ext_map, dbuf.ori_allocated);
     ExclusivePrefixSum::Block<ExtremaRead,ExtremaWrt,ExtremaTot,ExtremaWrtMap>( total_ext_ct, r, w, t, wrtm );
 
     __syncthreads();
@@ -356,8 +356,8 @@ void ori_prefix_sum( const int total_ext_ct, const int num_octaves )
             dct.ori_ps[o] = dct.ori_ps[o-1] + dct.ori_ct[o-1];
         }
 
-        dct.ori_total = dct.ori_ps[MAX_OCTAVES-1] + dct.ori_ct[MAX_OCTAVES-1];
-        dct.ext_total = dct.ext_ps[MAX_OCTAVES-1] + dct.ext_ct[MAX_OCTAVES-1];
+        dct.ori_total = min(dct.ori_ps[MAX_OCTAVES-1] + dct.ori_ct[MAX_OCTAVES-1], dbuf.ori_allocated);
+        dct.ext_total = min(dct.ext_ps[MAX_OCTAVES-1] + dct.ext_ct[MAX_OCTAVES-1], dbuf.ext_allocated);
     }
 }
 

src/popsift/sift_pyramid.cu

@@ -259,7 +259,11 @@ void prep_features( Descriptor* descriptor_base, int up_fac )
     const float xpos    = ext.xpos  * powf(2.0f, float(octave - up_fac));
     const float ypos    = ext.ypos  * powf(2.0f, float(octave - up_fac));
     const float sigma   = ext.sigma * powf(2.0f, float(octave - up_fac));
-    const int   num_ori = ext.num_ori;
+    int   num_ori = ext.num_ori;
+
+    if( ext.idx_ori + num_ori > dct.ori_total ) {
+        num_ori = max(dct.ori_total - ext.idx_ori, 0);
+    }
 
     fet.xpos    = xpos;
     fet.ypos    = ypos;