From 0314e0e5d74e0e6c979505448c5eb103c2181989 Mon Sep 17 00:00:00 2001
From: "marina.kolpakova" <marina.kolpakova@itseez.com>
Date: Thu, 27 Sep 2012 16:54:37 +0400
Subject: [PATCH] add kind in detection representation

---
 modules/gpu/src/cuda/isf-sc.cu | 17 +++----
 modules/gpu/src/icf.hpp        | 81 ++++++++--------------------------
 2 files changed, 27 insertions(+), 71 deletions(-)

diff --git a/modules/gpu/src/cuda/isf-sc.cu b/modules/gpu/src/cuda/isf-sc.cu
index 4bf410fc59..adfc9edcb9 100644
--- a/modules/gpu/src/cuda/isf-sc.cu
+++ b/modules/gpu/src/cuda/isf-sc.cu
@@ -123,7 +123,6 @@ namespace icf {
     {
 
         dprintf("feature box %d %d %d %d ", area.x, area.y, area.z, area.w);
-        dprintf("get for channel %d\n", channel);
         dprintf("extract feature for: [%d %d] [%d %d] [%d %d] [%d %d]\n",
             x + area.x, y + area.y,  x + area.z, y + area.y,  x + area.z,y + area.w,
             x + area.x, y + area.w);
@@ -140,13 +139,13 @@ namespace icf {
     }
 
     __global__ void test_kernel(const Level* levels, const Octave* octaves, const float* stages,
-        const Node* nodes, const float* leaves, PtrStepSz<uchar4> objects, uint* ctr)
+        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 (x > 0 || y > 0 || blockIdx.z > 0) return;
+        // if (blockIdx.z != 31) return;
         if(x >= level.workRect.x || y >= level.workRect.y) return;
 
         Octave octave = octaves[level.octave];
@@ -191,10 +190,10 @@ namespace icf {
 
         if(st == stEnd)
         {
-            int idx = atomicInc(ctr, objects.cols);
-            uchar4 val;
-            val.x = x * 4;
-            objects(0, idx) = val;
+            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);
         }
     }
 
@@ -214,11 +213,13 @@ namespace icf {
         const Node* nd = (const Node*)nodes.ptr();
         const float* lf = (const float*)leaves.ptr();
         uint* ctr = (uint*)counter.ptr();
+        Detection* det = (Detection*)objects.ptr();
+        uint max_det = objects.cols / sizeof(Detection);
 
         cudaChannelFormatDesc desc = cudaCreateChannelDesc<int>();
         cudaSafeCall( cudaBindTexture2D(0, thogluv, hogluv.data, desc, hogluv.cols, hogluv.rows, hogluv.step));
 
-        test_kernel<<<grid, block>>>(l, oct, st, nd, lf, objects, ctr);
+        test_kernel<<<grid, block>>>(l, oct, st, nd, lf, det, max_det, ctr);
 
         cudaSafeCall( cudaGetLastError());
         cudaSafeCall( cudaDeviceSynchronize());
diff --git a/modules/gpu/src/icf.hpp b/modules/gpu/src/icf.hpp
index ecd1886d38..35658892f6 100644
--- a/modules/gpu/src/icf.hpp
+++ b/modules/gpu/src/icf.hpp
@@ -47,11 +47,11 @@
 #include <opencv2/gpu/device/common.hpp>
 #include <stdio.h>
 
-// #if defined __CUDACC__
-// # define __device __device__ __forceinline__
-// #else
-// # define __device
-// #endif
+#if defined __CUDACC__
+# define __device __device__ __forceinline__
+#else
+# define __device
+#endif
 
 
 namespace cv { namespace gpu { namespace device {
@@ -108,66 +108,21 @@ struct __align__(8) Node
     }
 };
 
-// struct __align__(8) Feature
-// {
-//     int channel;
-//     uchar4 rect;
+struct __align__(16) Detection
+{
+    ushort x;
+    ushort y;
+    ushort w;
+    ushort h;
+    float confidence;
+    int kind;
+
+    Detection(){}
+    __device Detection(int _x, int _y, uchar _w, uchar _h, float c)
+    : x(_x), y(_y), w(_w), h(_h), confidence(c), kind(0) {};
+};
 
-//     Feature(const int c, const uchar4 r) : channel(c), rect(r) {}
-// };
 }
 }}}
-// struct Cascade
-// {
-//     Cascade() {}
-//     Cascade(const cv::gpu::PtrStepSzb& octs, const cv::gpu::PtrStepSzf& sts, const cv::gpu::PtrStepSzb& nds,
-//         const cv::gpu::PtrStepSzf& lvs, const cv::gpu::PtrStepSzb& fts, const cv::gpu::PtrStepSzb& lls)
-//     : octaves(octs), stages(sts), nodes(nds), leaves(lvs), features(fts), levels(lls) {}
-
-//     void detect(const cv::gpu::PtrStepSzi& hogluv, cv::gpu::PtrStepSz<uchar4> objects, cudaStream_t stream) const;
-//     void __device detectAt(const int* __restrict__ hogluv, const int pitch, PtrStepSz<uchar4>& objects) const;
-//     float __device rescale(const icf::Level& level, uchar4& scaledRect,
-//                            const int channel, const float threshold) const;
-
-//     PtrStepSzb octaves;
-//     PtrStepSzf stages;
-//     PtrStepSzb nodes;
-//     PtrStepSzf leaves;
-//     PtrStepSzb features;
-
-//     PtrStepSzb levels;
-
-// };
-
-// struct ChannelStorage
-// {
-//     ChannelStorage(){}
-//     ChannelStorage(const cv::gpu::PtrStepSzb& buff, const cv::gpu::PtrStepSzb& shr,
-//         const cv::gpu::PtrStepSzb& itg, const int s)
-//     : dmem (buff), shrunk(shr), hogluv(itg), shrinkage(s) {}
-
-//     void frame(const cv::gpu::PtrStepSz<uchar3>& rgb, cudaStream_t stream){}
-
-//     PtrStepSzb dmem;
-//     PtrStepSzb shrunk;
-//     PtrStepSzb hogluv;
-
-//     enum
-//     {
-//         FRAME_WIDTH        = 640,
-//         FRAME_HEIGHT       = 480,
-//         TOTAL_SCALES       = 55,
-//         CLASSIFIERS        = 5,
-//         ORIG_OBJECT_WIDTH  = 64,
-//         ORIG_OBJECT_HEIGHT = 128,
-//         HOG_BINS           = 6,
-//         HOG_LUV_BINS       = 10
-//     };
-
-//     int shrinkage;
-//     static const float magnitudeScaling = 1.f ;// / sqrt(2);
-// };
-
-// }}}
 
 #endif
\ No newline at end of file