Chiebot-Mirror
/
opencv
mirror of https://github.com/opencv/opencv.git
8
0
Fork 0
Code Issues Projects Releases Wiki Activity

[*] Fixed #1464

Browse Source 
[~] NCVPyramid uses tr1 and thus can be compiled with CL, commented out on linux
[+] Moved reduction functors to NCVAlg
[*] Warnings in NCV
pull/13383/head
Anton Obukhov 13 years ago
parent 325e0b1ab8
commit 2cb9192604
5 changed files with 66 additions and 33 deletions
Whitespace
Show all changes Ignore whitespace when comparing lines Ignore changes in amount of whitespace Ignore changes in whitespace at EOL
Split View
Diff Options
Show Stats Download Patch File Download Diff File
  1. 23
      modules/gpu/src/nvidia/NCVHaarObjectDetection.cu
  2. 21
      modules/gpu/src/nvidia/core/NCV.cu
  3. 34
      modules/gpu/src/nvidia/core/NCVAlg.hpp
  4. 9
      modules/gpu/src/nvidia/core/NCVPyramid.cu
  5. 12
      modules/gpu/src/nvidia/core/NCVPyramid.hpp

23
modules/gpu/src/nvidia/NCVHaarObjectDetection.cu
Unescape View File

@ -1,7 +1,7 @@
/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// 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.
@ -198,7 +198,7 @@ __device__ HaarClassifierNode128 getClassifierNode(Ncv32u iNode, HaarClassifierN
template <NcvBool tbCacheTextureCascade>
__device__ void getFeature(Ncv32u iFeature, HaarFeature64 *d_Features,
__device__ void getFeature(Ncv32u iFeature, HaarFeature64 *d_Features,
Ncv32f *weight,
Ncv32u *rectX, Ncv32u *rectY, Ncv32u *rectWidth, Ncv32u *rectHeight)
{
@ -321,9 +321,9 @@ __global__ void applyHaarClassifierAnchorParallel(Ncv32u *d_IImg, Ncv32u IImgStr
d_inMask != d_outMask &&
d_inMask[maskOffset] == OBJDET_MASK_ELEMENT_INVALID_32U))
{
if (tbDoAtomicCompaction)
if (tbDoAtomicCompaction)
{
bInactiveThread = true;
bInactiveThread = true;
}
else
{
@ -451,7 +451,7 @@ __global__ void applyHaarClassifierAnchorParallel(Ncv32u *d_IImg, Ncv32u IImgStr
}
}
template <NcvBool tbCacheTextureIImg,
NcvBool tbCacheTextureCascade,
NcvBool tbDoAtomicCompaction>
@ -565,14 +565,7 @@ __global__ void applyHaarClassifierClassifierParallel(Ncv32u *d_IImg, Ncv32u IIm
curRootNodeOffset += NUM_THREADS_CLASSIFIERPARALLEL;
}
struct functorAddValues
{
__device__ void reduce(Ncv32f &in1out, Ncv32f &in2)
{
in1out += in2;
}
};
Ncv32f finalStageSum = subReduce<Ncv32f, functorAddValues, NUM_THREADS_CLASSIFIERPARALLEL>(curStageSum);
Ncv32f finalStageSum = subReduce<Ncv32f, functorAddValues<Ncv32f>, NUM_THREADS_CLASSIFIERPARALLEL>(curStageSum);
if (finalStageSum < stageThreshold)
{
@ -1125,7 +1118,7 @@ NCVStatus ncvApplyHaarClassifierCascade_device(NCVMatrix<Ncv32u> &d_integralImag
}
if (curStop > compactEveryNstage && curStop - stageMiddleSwitch > compactEveryNstage / 2)
{
pixParallelStageStops[pixParallelStageStops.size()-1] =
pixParallelStageStops[pixParallelStageStops.size()-1] =
(stageMiddleSwitch - (curStop - 2 * compactEveryNstage)) / 2;
}
pixParallelStageStops.push_back(stageMiddleSwitch);

21
modules/gpu/src/nvidia/core/NCV.cu
Unescape View File

@ -1,7 +1,7 @@
/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// 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.
@ -278,6 +278,7 @@ NCVMemStackAllocator::NCVMemStackAllocator(NCVMemoryType memT, size_t capacity,
{
NcvBool bProperAlignment = (alignment & (alignment-1)) == 0;
ncvAssertPrintCheck(bProperAlignment, "NCVMemStackAllocator ctor:: _alignment not power of 2");
ncvAssertPrintCheck(memT != NCVMemoryTypeNone, "NCVMemStackAllocator ctor:: Incorrect allocator type");
allocBegin = NULL;
@ -294,7 +295,8 @@ NCVMemStackAllocator::NCVMemStackAllocator(NCVMemoryType memT, size_t capacity,
break;
case NCVMemoryTypeHostPageable:
allocBegin = (Ncv8u *)malloc(capacity);
break;
break;
default:;
}
}
else
@ -334,7 +336,8 @@ NCVMemStackAllocator::~NCVMemStackAllocator()
break;
case NCVMemoryTypeHostPageable:
free(allocBegin);
break;
break;
default:;
}
}
@ -454,7 +457,8 @@ NCVStatus NCVMemNativeAllocator::alloc(NCVMemSegment &seg, size_t size)
break;
case NCVMemoryTypeHostPageable:
seg.begin.ptr = (Ncv8u *)malloc(size);
break;
break;
default:;
}
this->currentSize += alignUp(size, this->_alignment);
@ -486,7 +490,8 @@ NCVStatus NCVMemNativeAllocator::dealloc(NCVMemSegment &seg)
break;
case NCVMemoryTypeHostPageable:
free(seg.begin.ptr);
break;
break;
default:;
}
seg.clear();
@ -568,13 +573,13 @@ typedef struct _NcvTimeMoment NcvTimeMoment;
return 1000.0 * 2 * ((t2->moment) - (t1->moment)) / (t1->freq + t2->freq);
}
#elif defined(__GNUC__)
#elif defined(__GNUC__)
#include <sys/time.h>
typedef struct _NcvTimeMoment
{
struct timeval tv;
struct timeval tv;
struct timezone tz;
} NcvTimeMoment;

34
modules/gpu/src/nvidia/core/NCVAlg.hpp
Unescape View File

@ -1,7 +1,7 @@
/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// 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.
@ -52,6 +52,36 @@ static T divUp(T a, T b)
}
template<typename T>
struct functorAddValues
{
static __device__ __inline__ void reduce(T &in1out, T &in2)
{
in1out += in2;
}
};
template<typename T>
struct functorMinValues
{
static __device__ __inline__ void reduce(T &in1out, T &in2)
{
in1out = in1out > in2 ? in2 : in1out;
}
};
template<typename T>
struct functorMaxValues
{
static __device__ __inline__ void reduce(T &in1out, T &in2)
{
in1out = in1out > in2 ? in1out : in2;
}
};
template<typename Tdata, class Tfunc, Ncv32u nThreads>
static __device__ Tdata subReduce(Tdata threadElem)
{

9
modules/gpu/src/nvidia/core/NCVPyramid.cu
Unescape View File

@ -1,7 +1,7 @@
/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// 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.
@ -46,6 +46,7 @@
#include "NCVPyramid.hpp"
#include "NCVPixelOperations.hpp"
#ifdef _WIN32
template<typename T, Ncv32u CN> struct __average4_CN {static T _average4_CN(const T &p00, const T &p01, const T &p10, const T &p11);};
@ -394,4 +395,6 @@ template class NCVImagePyramid<uint3>;
template class NCVImagePyramid<uint4>;
template class NCVImagePyramid<float1>;
template class NCVImagePyramid<float3>;
template class NCVImagePyramid<float4>;
template class NCVImagePyramid<float4>;
#endif //_WIN32

12
modules/gpu/src/nvidia/core/NCVPyramid.hpp
Unescape View File

@ -1,7 +1,7 @@
/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// 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.
@ -46,7 +46,8 @@
#include <memory>
#include <vector>
#include "NCV.hpp"
#ifdef _WIN32
template <class T>
class NCV_EXPORTS NCVMatrixStack
@ -91,7 +92,8 @@ private:
const NCVMatrix<T> *layer0;
NCVMatrixStack<T> pyramid;
Ncv32u nLayers;
};
};
#endif //_WIN32
#endif //_ncvpyramid_hpp_

Write Preview
Loading…
Cancel
Save