     __device__ __forceinline__ int get(const int x, int y, uchar4 area)
     {
 
-        dprintf("feature box %d %d %d %d ", area.x, area.y, area.z, area.w);
-        dprintf("extract feature for: [%d %d] [%d %d] [%d %d] [%d %d]\n",
+        dprintf("%d: feature box %d %d %d %d\n",threadIdx.x, area.x, area.y, area.z, area.w);
+        dprintf("%d: extract feature for: [%d %d] [%d %d] [%d %d] [%d %d]\n",threadIdx.x,
             x + area.x, y + area.y,  x + area.z, y + area.y,  x + area.z,y + area.w,
             x + area.x, y + area.w);
-        dprintf("at point %d %d with offset %d\n", x, y, 0);
+        dprintf("%d: at point %d %d with offset %d\n", x, y, 0);
 
         int a = tex2D(thogluv, x + area.x, y + area.y);
         int b = tex2D(thogluv, x + area.z, y + area.y);
         int c = tex2D(thogluv, x + area.z, y + area.w);
         int d = tex2D(thogluv, x + area.x, y + area.w);
 
-        dprintf("    retruved integral values: %d %d %d %d\n", a, b, c, d);
+        dprintf("%d   retruved integral values: %d %d %d %d\n",threadIdx.x, a, b, c, d);
 
         return (a - b + c - d);
     }
 
-    __global__ void test_kernel(const Level* levels, const Octave* octaves, const float* stages,
+//     __global__ void test_kernel(const Level* levels, const Octave* octaves, const float* stages,
+//         const Node* nodes, const float* leaves, Detection* objects, const uint ndetections, uint* ctr)
+//     {
+//         const int y = blockIdx.y * blockDim.y + threadIdx.y;
+//         const int x = blockIdx.x * blockDim.x + threadIdx.x;
+//         Level level = levels[blockIdx.z];
+
+//         // if (blockIdx.z != 31) return;
+//         if(x >= level.workRect.x || y >= level.workRect.y) return;
+
+//         Octave octave = octaves[level.octave];
+
+//         int st = octave.index * octave.stages;
+//         const int stEnd = st + 1000;//octave.stages;
+
+//         float confidence = 0.f;
+
+// // #pragma unroll 2
+//         for(; st < stEnd; ++st)
+//         {
+//             dprintf("\n\nstage: %d\n", st);
+//             const int nId = st * 3;
+//             Node node = nodes[nId];
+
+//             dprintf("Node: [%d %d %d %d] %d %d\n", node.rect.x, node.rect.y, node.rect.z, node.rect.w,
+//                 node.threshold >> 28, node.threshold & 0x0FFFFFFFU);
+
+//             float threshold = rescale(level, node.rect, node);
+//             int sum = get(x, y + (node.threshold >> 28) * 121, node.rect);
+
+//             dprintf("Node: [%d %d %d %d] %f\n", node.rect.x, node.rect.y, node.rect.z,
+//                 node.rect.w, threshold);
+
+//             int next = 1 + (int)(sum >= threshold);
+//             dprintf("go: %d (%d >= %f)\n\n" ,next, sum, threshold);
+
+//             node = nodes[nId + next];
+//             threshold = rescale(level, node.rect, node);
+//             sum = get(x, y + (node.threshold >> 28) * 121, node.rect);
+
+//             const int lShift = (next - 1) * 2 + (int)(sum >= threshold);
+//             float impact = leaves[st * 4 + lShift];
+//             confidence += impact;
+
+//             if (confidence <= stages[st]) st = stEnd + 10;
+//             dprintf("decided: %d (%d >= %f) %d %f\n\n" ,next, sum, threshold, lShift, impact);
+//             dprintf("extracted stage: %f\n", stages[st]);
+//             dprintf("computed  score: %f\n\n", confidence);
+//         }
+
+//         if(st == stEnd)
+//         {
+//             int idx = atomicInc(ctr, ndetections);
+//             // store detection
+//             objects[idx] = Detection(__float2int_rn(x * octave.shrinkage),
+//                 __float2int_rn(y * octave.shrinkage), level.objSize.x, level.objSize.y, confidence);
+//         }
+//     }
+
+    __global__ void test_kernel_warp(const Level* levels, const Octave* octaves, const float* stages,
         const Node* nodes, const float* leaves, Detection* objects, const uint ndetections, uint* ctr)
     {
         const int y = blockIdx.y * blockDim.y + threadIdx.y;
-        const int x = blockIdx.x * blockDim.x + threadIdx.x;
+        const int x = blockIdx.x;
+
         Level level = levels[blockIdx.z];
 
-        // if (blockIdx.z != 31) return;
         if(x >= level.workRect.x || y >= level.workRect.y) return;
 
         Octave octave = octaves[level.octave];
-
         int st = octave.index * octave.stages;
-        const int stEnd = st + 1000;//octave.stages;
+        const int stEnd = st + 1024;
 
         float confidence = 0.f;
 
-// #pragma unroll 2
-        for(; st < stEnd; ++st)
+        for(; st < stEnd; st += 32)
         {
-            dprintf("\n\nstage: %d\n", st);
-            const int nId = st * 3;
+
+            const int nId = (st + threadIdx.x) * 3;
+            dprintf("\n\n%d: stage: %d %d\n",threadIdx.x, st, nId);
             Node node = nodes[nId];
-
-            dprintf("Node: [%d %d %d %d] %d %d\n", node.rect.x, node.rect.y, node.rect.z, node.rect.w,
-                node.threshold >> 28, node.threshold & 0x0FFFFFFFU);
 
             float threshold = rescale(level, node.rect, node);
             int sum = get(x, y + (node.threshold >> 28) * 121, node.rect);
 
-            dprintf("Node: [%d %d %d %d] %f\n", node.rect.x, node.rect.y, node.rect.z,
-                node.rect.w, threshold);
-
             int next = 1 + (int)(sum >= threshold);
-            dprintf("go: %d (%d >= %f)\n\n" ,next, sum, threshold);
+            dprintf("%d: go: %d (%d >= %f)\n\n" ,threadIdx.x, next, sum, threshold);
 
             node = nodes[nId + next];
             threshold = rescale(level, node.rect, node);
             sum = get(x, y + (node.threshold >> 28) * 121, node.rect);
 
             const int lShift = (next - 1) * 2 + (int)(sum >= threshold);
-            float impact = leaves[st * 4 + lShift];
+            float impact = leaves[(st + threadIdx.x) * 4 + lShift];
+
+            dprintf("%d: decided: %d (%d >= %f) %d %f\n\n" ,threadIdx.x, next, sum, threshold, lShift, impact);
+            dprintf("%d: extracted stage: %f\n",threadIdx.x, stages[(st + threadIdx.x)]);
+            dprintf("%d: computed  score: %f\n",threadIdx.x, impact);
+
+            // scan on shuffl functions
+            for (int i = 1; i < 32; i *= 2)
+            {
+                const float n = __shfl_up(impact, i, 32);
+
+                if (threadIdx.x >= i)
+                    impact += n;
+            }
+
+            dprintf("%d: impact scaned %f\n" ,threadIdx.x, impact);
+
             confidence += impact;
-
-            if (confidence <= stages[st]) st = stEnd + 10;
-            dprintf("decided: %d (%d >= %f) %d %f\n\n" ,next, sum, threshold, lShift, impact);
-            dprintf("extracted stage: %f\n", stages[st]);
-            dprintf("computed  score: %f\n\n", confidence);
+            if(__any((confidence <= stages[(st + threadIdx.x)]))) break;
         }
 
-        if(st == stEnd)
+        if(st == stEnd && !threadIdx.x)
         {
             int idx = atomicInc(ctr, ndetections);
             // store detection

        const float expected_new_area = farea * relScale * relScale;
        float approx =  sarea / expected_new_area;

        dprintf("new rect: %d box %d %d %d %d  rel areas %f %f\n", (node.threshold >> 28),
        scaledRect.x, scaledRect.y, scaledRect.z, scaledRect.w, farea * relScale * relScale, sarea);

        float rootThreshold = (node.threshold & 0x0FFFFFFFU) * approx;
        rootThreshold *= level.scaling[(node.threshold >> 28) > 6];

        dprintf("approximation %f %d -> %f %f\n", approx, (node.threshold & 0x0FFFFFFFU), rootThreshold,
            level.scaling[(node.threshold >> 28) > 6]);

        return rootThreshold;

        cudaChannelFormatDesc desc = cudaCreateChannelDesc<int>();
        cudaSafeCall( cudaBindTexture2D(0, thogluv, hogluv.data, desc, hogluv.cols, hogluv.rows, hogluv.step));

        test_kernel_warp<<<grid, block>>>(l, oct, st, nd, lf, det, max_det, ctr);

        cudaSafeCall( cudaGetLastError());
        cudaSafeCall( cudaDeviceSynchronize());
