     float4 descr;
     descr = ptr4[threadIdx.x];
 
-    descr.x *= descr.x;
-    descr.y *= descr.y;
-    descr.z *= descr.z;
-    descr.w *= descr.w;
-
+#ifdef USE_ROOT_SIFT
+    // L1 norm
     float norm = descr.x + descr.y + descr.z + descr.w;
 
     norm += __shfl_down( norm, 16 );

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

    norm = sqrt(norm) + EPS_V1;

    norm += __shfl     ( norm,  0 );


    descr.x /= norm;
    descr.y /= norm;
    descr.z /= norm;
    descr.w /= norm;

    ptr4[threadIdx.x] = descr;
}

#define DESCR_V1_NUM_THREADS 32
#define RPI_V1               (4.0f / M_PI)
#define EPS_V1               1E-15

#define USE_ROOT_SIFT

/*************************************************************
 * V1: device side
