From 742517c23679815863b35bde611b11ce82690d0e Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Mitja=20Puziga=C4=87a?= Date: Tue, 4 Jul 2023 12:14:09 +0200 Subject: [PATCH] fix illegal access when there are more descriptors than allocated --- src/popsift/s_desc_grid.cu | 8 +++++++- src/popsift/s_desc_igrid.cu | 2 ++ src/popsift/s_desc_iloop.cu | 9 ++++++++- src/popsift/s_desc_loop.cu | 9 ++++++++- src/popsift/s_desc_notile.cu | 1 + src/popsift/s_orientation.cu | 6 +++--- src/popsift/sift_pyramid.cu | 6 +++++- 7 files changed, 34 insertions(+), 7 deletions(-) diff --git a/src/popsift/s_desc_grid.cu b/src/popsift/s_desc_grid.cu index 099c4709..0c52d9de 100644 --- a/src/popsift/s_desc_grid.cu +++ b/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; diff --git a/src/popsift/s_desc_igrid.cu b/src/popsift/s_desc_igrid.cu index 9f77f12f..6e942d2a 100644 --- a/src/popsift/s_desc_igrid.cu +++ b/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]; diff --git a/src/popsift/s_desc_iloop.cu b/src/popsift/s_desc_iloop.cu index 84673a20..f990f845 100644 --- a/src/popsift/s_desc_iloop.cu +++ b/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; diff --git a/src/popsift/s_desc_loop.cu b/src/popsift/s_desc_loop.cu index 4c5f46c2..5f3e8b11 100644 --- a/src/popsift/s_desc_loop.cu +++ b/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; diff --git a/src/popsift/s_desc_notile.cu b/src/popsift/s_desc_notile.cu index 9ba8a927..175c1e63 100644 --- a/src/popsift/s_desc_notile.cu +++ b/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]; diff --git a/src/popsift/s_orientation.cu b/src/popsift/s_orientation.cu index f6b36fcd..b657bfda 100644 --- a/src/popsift/s_orientation.cu +++ b/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( 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); } } diff --git a/src/popsift/sift_pyramid.cu b/src/popsift/sift_pyramid.cu index 06060052..7105a718 100644 --- a/src/popsift/sift_pyramid.cu +++ b/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;