{
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);
}
__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];
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);
}
}
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());
#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 {
}
};
-// 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