 #include "sift_pyramid.h"
 #include "sift_constants.h"
 #include "s_gradiant.h"
+#include "assist.h"
 
 /*************************************************************
  * V1: device side

    descr.w *= norm;
#else // not DESC_USE_ROOT_SIFT
    // L2 norm
    descr.x *= descr.x;
    descr.y *= descr.y;
    descr.z *= descr.z;
    descr.w *= descr.w;

    float norm = descr.x + descr.y + descr.z + descr.w;

    norm += __shfl_down( norm, 16 );
    norm += __shfl_down( norm,  8 );
    norm += __shfl_down( norm,  4 );
    norm += __shfl_down( norm,  2 );
    norm += __shfl_down( norm,  1 );

    if( threadIdx.x == 0 ) {
        norm = __fsqrt_rn(norm) + DESC_MIN_FLOAT;
        norm = __frcp_rn(norm); // reciprocal
    }

    norm = __shfl( norm,  0 );

    descr.x *= norm;
    descr.y *= norm;

}

__global__
void normalize_histogram( Descriptor* descs, int num_orientations )
{
    int offset = ( blockIdx.x < num_orientations ) ? blockIdx.x
                                                   : num_orientations-1;
    Descriptor* desc = &descs[offset];

    float*  ptr1 = desc->features;
    float4* ptr4 = (float4*)ptr1;
