         int fh = 120;
 
         dim3 block(32, 8);
-        dim3 grid(fw / 32, fh / 8, 47);
+        dim3 grid(fw, fh / 8, 47);
 
         const Level* l = (const Level*)levels.ptr();
         const Octave* oct = ((const Octave*)octaves.ptr());

        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;

    __device__ __forceinline__ int get(const int x, int y, uchar4 area)
    {

        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("%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("%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,
//         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;

        Level level = levels[blockIdx.z];

        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 + 1024;

        float confidence = 0.f;

        for(; st < stEnd; st += 32)
        {

            const int nId = (st + threadIdx.x) * 3;
            dprintf("\n\n%d: stage: %d %d\n",threadIdx.x, st, nId);
            Node node = nodes[nId];

            float threshold = rescale(level, node.rect, node);
            int sum = get(x, y + (node.threshold >> 28) * 121, node.rect);

            int next = 1 + (int)(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 + 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(__any((confidence <= stages[(st + threadIdx.x)]))) break;
        }

        if(st == stEnd && !threadIdx.x)
        {
            int idx = atomicInc(ctr, ndetections);
            // store detection
