mirror/opencv
1
0
Fork 0
mirror of https://github.com/opencv/opencv.git synced 2025-01-19 15:04:01 +08:00
Code Issues Actions Packages Projects Releases Wiki Activity

add kind in detection representation

Browse Source
This commit is contained in:
marina.kolpakova 2012-09-27 16:54:37 +04:00
parent c0359ed5c5
commit 0314e0e5d7
2 changed files with 27 additions and 71 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

17
modules/gpu/src/cuda/isf-sc.cu
View File

@ -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());

81
modules/gpu/src/icf.hpp
View File

@ -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