     norm += __shfl_down( norm,  4 );
     norm += __shfl_down( norm,  2 );
     norm += __shfl_down( norm,  1 );
+    norm += __shfl     ( norm,  0 );
+
+    norm = 512.0f / norm; /* multiplying with 512 is some scaling by convention */
+    descr.x *= norm;
+    descr.y *= norm;
+    descr.z *= norm;
+    descr.w *= norm;
+#else // not 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 );
 
     norm = sqrt(norm) + EPS_V1;
 
     norm += __shfl     ( norm,  0 );
-
 
     descr.x /= norm;
     descr.y /= norm;
     descr.z /= norm;
     descr.w /= norm;
+#endif // not USE_ROOT_SIFT
 
     ptr4[threadIdx.x] = descr;
 }

    float4 descr;
    descr = ptr4[threadIdx.x];

    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 );

#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
