converted old haar cascades to the new format; added the conversion function; added OpenCL optimization into CascadeClassfier; optimized the data structures and CPU code for the stump case.
Vadim Pisarevsky
11 years ago
commit
a1784b7320
31 changed files with 501415 additions and 584320 deletions
File diff suppressed because it is too large
Load Diff
File diff suppressed because it is too large
Load Diff
File diff suppressed because it is too large
Load Diff
File diff suppressed because it is too large
Load Diff
File diff suppressed because it is too large
Load Diff
File diff suppressed because it is too large
Load Diff
File diff suppressed because it is too large
Load Diff
File diff suppressed because it is too large
Load Diff
File diff suppressed because it is too large
Load Diff
File diff suppressed because it is too large
Load Diff
File diff suppressed because it is too large
Load Diff
File diff suppressed because it is too large
Load Diff
File diff suppressed because it is too large
Load Diff
File diff suppressed because it is too large
Load Diff
File diff suppressed because it is too large
Load Diff
File diff suppressed because it is too large
Load Diff
File diff suppressed because it is too large
Load Diff
File diff suppressed because it is too large
Load Diff
File diff suppressed because it is too large
Load Diff
File diff suppressed because it is too large
Load Diff
File diff suppressed because it is too large
Load Diff
File diff suppressed because it is too large
Load Diff
@ -0,0 +1,276 @@ |
||||
/*M///////////////////////////////////////////////////////////////////////////////////////
|
||||
//
|
||||
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
|
||||
//
|
||||
// By downloading, copying, installing or using the software you agree to this license.
|
||||
// If you do not agree to this license, do not download, install,
|
||||
// copy or use the software.
|
||||
//
|
||||
//
|
||||
// License Agreement
|
||||
// For Open Source Computer Vision Library
|
||||
//
|
||||
// Copyright (C) 2013, Itseez Inc, all rights reserved.
|
||||
// Third party copyrights are property of their respective owners.
|
||||
//
|
||||
// Redistribution and use in source and binary forms, with or without modification,
|
||||
// are permitted provided that the following conditions are met:
|
||||
//
|
||||
// * Redistribution's of source code must retain the above copyright notice,
|
||||
// this list of conditions and the following disclaimer.
|
||||
//
|
||||
// * Redistribution's in binary form must reproduce the above copyright notice,
|
||||
// this list of conditions and the following disclaimer in the documentation
|
||||
// and/or other materials provided with the distribution.
|
||||
//
|
||||
// * The name of Intel Corporation may not be used to endorse or promote products
|
||||
// derived from this software without specific prior written permission.
|
||||
//
|
||||
// This software is provided by the copyright holders and contributors "as is" and
|
||||
// any express or implied warranties, including, but not limited to, the implied
|
||||
// warranties of merchantability and fitness for a particular purpose are disclaimed.
|
||||
// In no event shall the Intel Corporation or contributors be liable for any direct,
|
||||
// indirect, incidental, special, exemplary, or consequential damages
|
||||
// (including, but not limited to, procurement of substitute goods or services;
|
||||
// loss of use, data, or profits; or business interruption) however caused
|
||||
// and on any theory of liability, whether in contract, strict liability,
|
||||
// or tort (including negligence or otherwise) arising in any way out of
|
||||
// the use of this software, even if advised of the possibility of such damage.
|
||||
//
|
||||
//M*/
|
||||
|
||||
/* Haar features calculation */ |
||||
|
||||
#include "precomp.hpp" |
||||
#include <stdio.h> |
||||
|
||||
namespace cv |
||||
{ |
||||
|
||||
/* field names */ |
||||
|
||||
#define ICV_HAAR_SIZE_NAME "size" |
||||
#define ICV_HAAR_STAGES_NAME "stages" |
||||
#define ICV_HAAR_TREES_NAME "trees" |
||||
#define ICV_HAAR_FEATURE_NAME "feature" |
||||
#define ICV_HAAR_RECTS_NAME "rects" |
||||
#define ICV_HAAR_TILTED_NAME "tilted" |
||||
#define ICV_HAAR_THRESHOLD_NAME "threshold" |
||||
#define ICV_HAAR_LEFT_NODE_NAME "left_node" |
||||
#define ICV_HAAR_LEFT_VAL_NAME "left_val" |
||||
#define ICV_HAAR_RIGHT_NODE_NAME "right_node" |
||||
#define ICV_HAAR_RIGHT_VAL_NAME "right_val" |
||||
#define ICV_HAAR_STAGE_THRESHOLD_NAME "stage_threshold" |
||||
#define ICV_HAAR_PARENT_NAME "parent" |
||||
#define ICV_HAAR_NEXT_NAME "next" |
||||
|
||||
namespace haar_cvt |
||||
{ |
||||
|
||||
struct HaarFeature |
||||
{ |
||||
enum { RECT_NUM = 3 }; |
||||
|
||||
HaarFeature() |
||||
{ |
||||
tilted = false; |
||||
for( int i = 0; i < RECT_NUM; i++ ) |
||||
{ |
||||
rect[i].r = Rect(0,0,0,0); |
||||
rect[i].weight = 0.f; |
||||
} |
||||
} |
||||
bool tilted; |
||||
struct
|
||||
{ |
||||
Rect r; |
||||
float weight; |
||||
} rect[RECT_NUM]; |
||||
}; |
||||
|
||||
struct HaarClassifierNode |
||||
{ |
||||
HaarClassifierNode() |
||||
{ |
||||
f = left = right = 0; |
||||
threshold = 0.f; |
||||
} |
||||
int f, left, right; |
||||
float threshold; |
||||
}; |
||||
|
||||
struct HaarClassifier |
||||
{ |
||||
std::vector<HaarClassifierNode> nodes; |
||||
std::vector<float> leaves; |
||||
}; |
||||
|
||||
struct HaarStageClassifier |
||||
{ |
||||
double threshold; |
||||
std::vector<HaarClassifier> weaks; |
||||
}; |
||||
|
||||
static bool convert(const String& oldcascade, const String& newcascade) |
||||
{ |
||||
FileStorage oldfs(oldcascade, FileStorage::READ); |
||||
if( !oldfs.isOpened() ) |
||||
return false; |
||||
FileNode oldroot = oldfs.getFirstTopLevelNode(); |
||||
|
||||
FileNode sznode = oldroot[ICV_HAAR_SIZE_NAME]; |
||||
if( sznode.empty() ) |
||||
return false; |
||||
int maxdepth = 0; |
||||
Size cascadesize; |
||||
cascadesize.width = (int)sznode[0]; |
||||
cascadesize.height = (int)sznode[1]; |
||||
std::vector<HaarFeature> features; |
||||
|
||||
size_t i, j, k, n; |
||||
|
||||
FileNode stages_seq = oldroot[ICV_HAAR_STAGES_NAME]; |
||||
size_t nstages = stages_seq.size(); |
||||
std::vector<HaarStageClassifier> stages(nstages); |
||||
|
||||
for( i = 0; i < nstages; i++ ) |
||||
{ |
||||
FileNode stagenode = stages_seq[i]; |
||||
HaarStageClassifier& stage = stages[i]; |
||||
stage.threshold = (double)stagenode[ICV_HAAR_STAGE_THRESHOLD_NAME]; |
||||
FileNode weaks_seq = stagenode[ICV_HAAR_TREES_NAME]; |
||||
size_t nweaks = weaks_seq.size(); |
||||
stage.weaks.resize(nweaks); |
||||
|
||||
for( j = 0; j < nweaks; j++ ) |
||||
{ |
||||
HaarClassifier& weak = stage.weaks[j]; |
||||
FileNode weaknode = weaks_seq[j]; |
||||
size_t nnodes = weaknode.size(); |
||||
|
||||
for( n = 0; n < nnodes; n++ ) |
||||
{ |
||||
FileNode nnode = weaknode[n]; |
||||
FileNode fnode = nnode[ICV_HAAR_FEATURE_NAME]; |
||||
HaarFeature f; |
||||
HaarClassifierNode node; |
||||
node.f = (int)features.size(); |
||||
f.tilted = (int)fnode[ICV_HAAR_TILTED_NAME] != 0; |
||||
FileNode rects_seq = fnode[ICV_HAAR_RECTS_NAME]; |
||||
size_t nrects = rects_seq.size(); |
||||
|
||||
for( k = 0; k < nrects; k++ ) |
||||
{ |
||||
FileNode rnode = rects_seq[k]; |
||||
f.rect[k].r.x = (int)rnode[0]; |
||||
f.rect[k].r.y = (int)rnode[1]; |
||||
f.rect[k].r.width = (int)rnode[2]; |
||||
f.rect[k].r.height = (int)rnode[3]; |
||||
f.rect[k].weight = (float)rnode[4]; |
||||
} |
||||
features.push_back(f); |
||||
node.threshold = nnode[ICV_HAAR_THRESHOLD_NAME]; |
||||
FileNode leftValNode = nnode[ICV_HAAR_LEFT_VAL_NAME]; |
||||
if( !leftValNode.empty() ) |
||||
{ |
||||
node.left = -(int)weak.leaves.size(); |
||||
weak.leaves.push_back((float)leftValNode); |
||||
} |
||||
else |
||||
{ |
||||
node.left = (int)nnode[ICV_HAAR_LEFT_NODE_NAME]; |
||||
} |
||||
FileNode rightValNode = nnode[ICV_HAAR_RIGHT_VAL_NAME]; |
||||
if( !rightValNode.empty() ) |
||||
{ |
||||
node.right = -(int)weak.leaves.size(); |
||||
weak.leaves.push_back((float)rightValNode); |
||||
} |
||||
else |
||||
{ |
||||
node.right = (int)nnode[ICV_HAAR_RIGHT_NODE_NAME]; |
||||
} |
||||
weak.nodes.push_back(node); |
||||
} |
||||
} |
||||
} |
||||
|
||||
FileStorage newfs(newcascade, FileStorage::WRITE); |
||||
if( !newfs.isOpened() ) |
||||
return false; |
||||
|
||||
size_t maxWeakCount = 0, nfeatures = features.size(); |
||||
for( i = 0; i < nstages; i++ ) |
||||
maxWeakCount = std::max(maxWeakCount, stages[i].weaks.size()); |
||||
|
||||
newfs << "cascade" << "{:opencv-cascade-classifier" |
||||
<< "stageType" << "BOOST" |
||||
<< "featureType" << "HAAR" |
||||
<< "height" << cascadesize.width |
||||
<< "width" << cascadesize.height |
||||
<< "stageParams" << "{" |
||||
<< "maxWeakCount" << (int)maxWeakCount |
||||
<< "}" |
||||
<< "featureParams" << "{" |
||||
<< "maxCatCount" << 0 |
||||
<< "}" |
||||
<< "stageNum" << (int)nstages |
||||
<< "stages" << "["; |
||||
|
||||
for( i = 0; i < nstages; i++ ) |
||||
{ |
||||
size_t nweaks = stages[i].weaks.size(); |
||||
newfs << "{" << "maxWeakCount" << (int)nweaks |
||||
<< "stageThreshold" << stages[i].threshold |
||||
<< "weakClassifiers" << "["; |
||||
for( j = 0; j < nweaks; j++ ) |
||||
{ |
||||
const HaarClassifier& c = stages[i].weaks[j]; |
||||
newfs << "{" << "internalNodes" << "["; |
||||
size_t nnodes = c.nodes.size(), nleaves = c.leaves.size(); |
||||
for( k = 0; k < nnodes; k++ ) |
||||
newfs << c.nodes[k].left << c.nodes[k].right |
||||
<< c.nodes[k].f << c.nodes[k].threshold; |
||||
newfs << "]" << "leafValues" << "["; |
||||
for( k = 0; k < nleaves; k++ ) |
||||
newfs << c.leaves[k]; |
||||
newfs << "]" << "}"; |
||||
} |
||||
newfs << "]" << "}"; |
||||
} |
||||
|
||||
newfs << "]" |
||||
<< "features" << "["; |
||||
|
||||
for( i = 0; i < nfeatures; i++ ) |
||||
{ |
||||
const HaarFeature& f = features[i]; |
||||
newfs << "{" << "rects" << "["; |
||||
for( j = 0; j < (size_t)HaarFeature::RECT_NUM; j++ ) |
||||
{ |
||||
if( j >= 2 && fabs(f.rect[j].weight) < FLT_EPSILON ) |
||||
break; |
||||
newfs << "[" << f.rect[j].r.x << f.rect[j].r.y << |
||||
f.rect[j].r.width << f.rect[j].r.height << f.rect[j].weight << "]"; |
||||
} |
||||
newfs << "]"; |
||||
if( f.tilted ) |
||||
newfs << "tilted" << 1; |
||||
newfs << "}"; |
||||
} |
||||
|
||||
newfs << "]" << "}"; |
||||
return true; |
||||
} |
||||
|
||||
} |
||||
|
||||
bool CascadeClassifier::convert(const String& oldcascade, const String& newcascade) |
||||
{ |
||||
bool ok = haar_cvt::convert(oldcascade, newcascade); |
||||
if( !ok && newcascade.size() > 0 ) |
||||
remove(newcascade.c_str()); |
||||
return ok; |
||||
} |
||||
|
||||
} |
||||
@ -0,0 +1,186 @@ |
||||
///////////////////////////// OpenCL kernels for face detection ////////////////////////////// |
||||
////////////////////////////// see the opencv/doc/license.txt /////////////////////////////// |
||||
|
||||
typedef struct __attribute__((aligned(4))) OptFeature |
||||
{ |
||||
int4 ofs[3] __attribute__((aligned (4))); |
||||
float4 weight __attribute__((aligned (4))); |
||||
} |
||||
OptFeature; |
||||
|
||||
typedef struct __attribute__((aligned(4))) Stump |
||||
{ |
||||
int featureIdx __attribute__((aligned (4))); |
||||
float threshold __attribute__((aligned (4))); // for ordered features only |
||||
float left __attribute__((aligned (4))); |
||||
float right __attribute__((aligned (4))); |
||||
} |
||||
Stump; |
||||
|
||||
typedef struct __attribute__((aligned (4))) Stage |
||||
{ |
||||
int first __attribute__((aligned (4))); |
||||
int ntrees __attribute__((aligned (4))); |
||||
float threshold __attribute__((aligned (4))); |
||||
} |
||||
Stage; |
||||
|
||||
__kernel void runHaarClassifierStump( |
||||
__global const int* sum, |
||||
int sumstep, int sumoffset, |
||||
__global const int* sqsum, |
||||
int sqsumstep, int sqsumoffset, |
||||
__global const OptFeature* optfeatures, |
||||
|
||||
int nstages, |
||||
__global const Stage* stages, |
||||
__global const Stump* stumps, |
||||
|
||||
volatile __global int* facepos, |
||||
int2 imgsize, int xyscale, float factor, |
||||
int4 normrect, int2 windowsize, int maxFaces) |
||||
{ |
||||
int ix = get_global_id(0)*xyscale; |
||||
int iy = get_global_id(1)*xyscale; |
||||
sumstep /= sizeof(int); |
||||
sqsumstep /= sizeof(int); |
||||
|
||||
if( ix < imgsize.x && iy < imgsize.y ) |
||||
{ |
||||
int ntrees; |
||||
int stageIdx, i; |
||||
float s = 0.f; |
||||
__global const Stump* stump = stumps; |
||||
__global const OptFeature* f; |
||||
|
||||
__global const int* psum = sum + mad24(iy, sumstep, ix); |
||||
__global const int* pnsum = psum + mad24(normrect.y, sumstep, normrect.x); |
||||
int normarea = normrect.z * normrect.w; |
||||
float invarea = 1.f/normarea; |
||||
float sval = (pnsum[0] - pnsum[normrect.z] - pnsum[mul24(normrect.w, sumstep)] + |
||||
pnsum[mad24(normrect.w, sumstep, normrect.z)])*invarea; |
||||
float sqval = (sqsum[mad24(iy + normrect.y, sqsumstep, ix + normrect.x)])*invarea; |
||||
float nf = (float)normarea * sqrt(max(sqval - sval * sval, 0.f)); |
||||
float4 weight, vsval; |
||||
int4 ofs, ofs0, ofs1, ofs2; |
||||
nf = nf > 0 ? nf : 1.f; |
||||
|
||||
for( stageIdx = 0; stageIdx < nstages; stageIdx++ ) |
||||
{ |
||||
ntrees = stages[stageIdx].ntrees; |
||||
s = 0.f; |
||||
for( i = 0; i < ntrees; i++, stump++ ) |
||||
{ |
||||
f = optfeatures + stump->featureIdx; |
||||
weight = f->weight; |
||||
|
||||
ofs = f->ofs[0]; |
||||
sval = (psum[ofs.x] - psum[ofs.y] - psum[ofs.z] + psum[ofs.w])*weight.x; |
||||
ofs = f->ofs[1]; |
||||
sval += (psum[ofs.x] - psum[ofs.y] - psum[ofs.z] + psum[ofs.w])*weight.y; |
||||
if( weight.z > 0 ) |
||||
{ |
||||
ofs = f->ofs[2]; |
||||
sval += (psum[ofs.x] - psum[ofs.y] - psum[ofs.z] + psum[ofs.w])*weight.z; |
||||
} |
||||
|
||||
s += (sval < stump->threshold*nf) ? stump->left : stump->right; |
||||
} |
||||
|
||||
if( s < stages[stageIdx].threshold ) |
||||
break; |
||||
} |
||||
|
||||
if( stageIdx == nstages ) |
||||
{ |
||||
int nfaces = atomic_inc(facepos); |
||||
if( nfaces < maxFaces ) |
||||
{ |
||||
volatile __global int* face = facepos + 1 + nfaces*4; |
||||
face[0] = convert_int_rte(ix*factor); |
||||
face[1] = convert_int_rte(iy*factor); |
||||
face[2] = convert_int_rte(windowsize.x*factor); |
||||
face[3] = convert_int_rte(windowsize.y*factor); |
||||
} |
||||
} |
||||
} |
||||
} |
||||
|
||||
#if 0 |
||||
__kernel void runLBPClassifierStump( |
||||
__global const int* sum, |
||||
int sumstep, int sumoffset, |
||||
__global const int* sqsum, |
||||
int sqsumstep, int sqsumoffset, |
||||
__global const OptFeature* optfeatures, |
||||
|
||||
int nstages, |
||||
__global const Stage* stages, |
||||
__global const Stump* stumps, |
||||
__global const int* bitsets, |
||||
int bitsetSize, |
||||
|
||||
volatile __global int* facepos, |
||||
int2 imgsize, int xyscale, float factor, |
||||
int4 normrect, int2 windowsize, int maxFaces) |
||||
{ |
||||
int ix = get_global_id(0)*xyscale*VECTOR_SIZE; |
||||
int iy = get_global_id(1)*xyscale; |
||||
sumstep /= sizeof(int); |
||||
sqsumstep /= sizeof(int); |
||||
|
||||
if( ix < imgsize.x && iy < imgsize.y ) |
||||
{ |
||||
int ntrees; |
||||
int stageIdx, i; |
||||
float s = 0.f; |
||||
__global const Stump* stump = stumps; |
||||
__global const int* bitset = bitsets; |
||||
__global const OptFeature* f; |
||||
|
||||
__global const int* psum = sum + mad24(iy, sumstep, ix); |
||||
__global const int* pnsum = psum + mad24(normrect.y, sumstep, normrect.x); |
||||
int normarea = normrect.z * normrect.w; |
||||
float invarea = 1.f/normarea; |
||||
float sval = (pnsum[0] - pnsum[normrect.z] - pnsum[mul24(normrect.w, sumstep)] + |
||||
pnsum[mad24(normrect.w, sumstep, normrect.z)])*invarea; |
||||
float sqval = (sqsum[mad24(iy + normrect.y, sqsumstep, ix + normrect.x)])*invarea; |
||||
float nf = (float)normarea * sqrt(max(sqval - sval * sval, 0.f)); |
||||
float4 weight; |
||||
int4 ofs; |
||||
nf = nf > 0 ? nf : 1.f; |
||||
|
||||
for( stageIdx = 0; stageIdx < nstages; stageIdx++ ) |
||||
{ |
||||
ntrees = stages[stageIdx].ntrees; |
||||
s = 0.f; |
||||
for( i = 0; i < ntrees; i++, stump++, bitset += bitsetSize ) |
||||
{ |
||||
f = optfeatures + stump->featureIdx; |
||||
|
||||
weight = f->weight; |
||||
|
||||
// compute LBP feature to val |
||||
s += (bitset[val >> 5] & (1 << (val & 31))) ? stump->left : stump->right; |
||||
} |
||||
|
||||
if( s < stages[stageIdx].threshold ) |
||||
break; |
||||
} |
||||
|
||||
if( stageIdx == nstages ) |
||||
{ |
||||
int nfaces = atomic_inc(facepos); |
||||
if( nfaces < maxFaces ) |
||||
{ |
||||
volatile __global int* face = facepos + 1 + nfaces*4; |
||||
face[0] = convert_int_rte(ix*factor); |
||||
face[1] = convert_int_rte(iy*factor); |
||||
face[2] = convert_int_rte(windowsize.x*factor); |
||||
face[3] = convert_int_rte(windowsize.y*factor); |
||||
} |
||||
} |
||||
} |
||||
} |
||||
#endif |
||||
|
||||
Loading…
Reference in new issue