diff --git a/modules/gpu/src/nvidia/NCVHaarObjectDetection.cu b/modules/gpu/src/nvidia/NCVHaarObjectDetection.cu index 2c6f65727e..f17e169b83 100644 --- a/modules/gpu/src/nvidia/NCVHaarObjectDetection.cu +++ b/modules/gpu/src/nvidia/NCVHaarObjectDetection.cu @@ -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 -__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 @@ -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(curStageSum); + Ncv32f finalStageSum = subReduce, NUM_THREADS_CLASSIFIERPARALLEL>(curStageSum); if (finalStageSum < stageThreshold) { @@ -1125,7 +1118,7 @@ NCVStatus ncvApplyHaarClassifierCascade_device(NCVMatrix &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); diff --git a/modules/gpu/src/nvidia/core/NCV.cu b/modules/gpu/src/nvidia/core/NCV.cu index a5af00a8ea..f8be5c8bad 100644 --- a/modules/gpu/src/nvidia/core/NCV.cu +++ b/modules/gpu/src/nvidia/core/NCV.cu @@ -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 typedef struct _NcvTimeMoment { - struct timeval tv; + struct timeval tv; struct timezone tz; } NcvTimeMoment; diff --git a/modules/gpu/src/nvidia/core/NCVAlg.hpp b/modules/gpu/src/nvidia/core/NCVAlg.hpp index 0dda7c1dcc..7c6cebd08d 100644 --- a/modules/gpu/src/nvidia/core/NCVAlg.hpp +++ b/modules/gpu/src/nvidia/core/NCVAlg.hpp @@ -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 +struct functorAddValues +{ + static __device__ __inline__ void reduce(T &in1out, T &in2) + { + in1out += in2; + } +}; + + +template +struct functorMinValues +{ + static __device__ __inline__ void reduce(T &in1out, T &in2) + { + in1out = in1out > in2 ? in2 : in1out; + } +}; + + +template +struct functorMaxValues +{ + static __device__ __inline__ void reduce(T &in1out, T &in2) + { + in1out = in1out > in2 ? in1out : in2; + } +}; + + template static __device__ Tdata subReduce(Tdata threadElem) { diff --git a/modules/gpu/src/nvidia/core/NCVPyramid.cu b/modules/gpu/src/nvidia/core/NCVPyramid.cu index 20e52e580b..463178362c 100644 --- a/modules/gpu/src/nvidia/core/NCVPyramid.cu +++ b/modules/gpu/src/nvidia/core/NCVPyramid.cu @@ -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 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; template class NCVImagePyramid; template class NCVImagePyramid; template class NCVImagePyramid; -template class NCVImagePyramid; +template class NCVImagePyramid; + +#endif //_WIN32 diff --git a/modules/gpu/src/nvidia/core/NCVPyramid.hpp b/modules/gpu/src/nvidia/core/NCVPyramid.hpp index 92cf90fab9..1885b17235 100644 --- a/modules/gpu/src/nvidia/core/NCVPyramid.hpp +++ b/modules/gpu/src/nvidia/core/NCVPyramid.hpp @@ -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 #include #include "NCV.hpp" - + +#ifdef _WIN32 template class NCV_EXPORTS NCVMatrixStack @@ -91,7 +92,8 @@ private: const NCVMatrix *layer0; NCVMatrixStack pyramid; Ncv32u nLayers; -}; - +}; + +#endif //_WIN32 #endif //_ncvpyramid_hpp_