     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 );
+    __shared__ float normalized_y[32];
+
+    float norm;
+    if( threadIdx.y == 0 ) {
+        // rnormf() is powerful: reciprocal L2 norm of N 
+        // consecutive floats, here N=128
+        norm = rnormf( 128, ptr1 );
+        normalized_y[threadIdx.x] = norm;
+    }
+    __syncthreads();
+    norm = normalized_y[threadIdx.y];
 
     descr.x *= norm;
     descr.y *= norm;

}

__global__
void normalize_histogram( Descriptor* descs )
{
    Descriptor* desc = &descs[blockIdx.x];

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

#include "sift_pyramid.h"
#include "sift_constants.h"
#include "s_gradiant.h"
#include "assist.h"

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