     __constant__ float c_NX[2][5] = {{0, 0, 2, 4, -1}, {2, 0, 4, 4, 1}};
     __constant__ float c_NY[2][5] = {{0, 0, 4, 2, 1}, {0, 2, 4, 4, -1}};
 
+    __device__ void reduceSum32(volatile float* v_sum, float& sum)
+    {
+        v_sum[threadIdx.x] = sum;
+
+        if (threadIdx.x < 16)
+        {
+            v_sum[threadIdx.x] = sum += v_sum[threadIdx.x + 16];
+            v_sum[threadIdx.x] = sum += v_sum[threadIdx.x + 8];
+            v_sum[threadIdx.x] = sum += v_sum[threadIdx.x + 4];
+            v_sum[threadIdx.x] = sum += v_sum[threadIdx.x + 2];
+            v_sum[threadIdx.x] = sum += v_sum[threadIdx.x + 1];
+        }
+    }
+
     __global__ void icvCalcOrientation(const float* featureX, const float* featureY, const float* featureSize, float* featureDir)
     {        
         #if defined (__CUDA_ARCH__) && __CUDA_ARCH__ >= 110

        if ((c_img_rows + 1) >= grad_wav_size && (c_img_cols + 1) >= grad_wav_size)
        {
            // Calc X, Y, angle and store it to shared memory
            {
                const int tid = threadIdx.y * blockDim.x + threadIdx.x;

                float X = 0.0f, Y = 0.0f, angle = 0.0f;

                if (tid < ORI_SAMPLES)
                {
                    const float margin = (float)(grad_wav_size - 1) / 2.0f;
                    const int x = __float2int_rn(featureX[blockIdx.x] + c_aptX[tid] * s - margin);
                    const int y = __float2int_rn(featureY[blockIdx.x] + c_aptY[tid] * s - margin);

                    if ((unsigned)y < (unsigned)((c_img_rows + 1) - grad_wav_size) && (unsigned)x < (unsigned)((c_img_cols + 1) - grad_wav_size))
                    {
                        X = c_aptW[tid] * icvCalcHaarPatternSum<2>(c_NX, 4, grad_wav_size, y, x);
                        Y = c_aptW[tid] * icvCalcHaarPatternSum<2>(c_NY, 4, grad_wav_size, y, x);
                    
                        angle = atan2f(Y, X);
                        if (angle < 0)
                            angle += 2.0f * CV_PI;
                        angle *= 180.0f / CV_PI;
                    }
                }
                if (tid < 128)
                {
                    s_X[tid] = X;
                    s_Y[tid] = Y;
                    s_angle[tid] = angle;
                }
            }
            __syncthreads();

            float bestx = 0, besty = 0, best_mod = 0;


        if (threadIdx.x < 5 && threadIdx.y < 5)
        {
            const int tid = threadIdx.y * 5 + threadIdx.x;

            const float dw = c_DW[yIndex * PATCH_SZ + xIndex];

