Lines Matching refs:Ncv32u
92 __device__ Ncv32u warpScanInclusive(Ncv32u idata, volatile Ncv32u *s_Data) in warpScanInclusive()
101 const Ncv32u n = cv::cuda::device::shfl_up(idata, i); in warpScanInclusive()
108 Ncv32u pos = 2 * threadIdx.x - (threadIdx.x & (K_WARP_SIZE - 1)); in warpScanInclusive()
123 __device__ __forceinline__ Ncv32u warpScanExclusive(Ncv32u idata, volatile Ncv32u *s_Data) in warpScanExclusive()
128 template <Ncv32u tiNumScanThreads>
129 __device__ Ncv32u scan1Inclusive(Ncv32u idata, volatile Ncv32u *s_Data) in scan1Inclusive()
134 Ncv32u warpResult = warpScanInclusive(idata, s_Data); in scan1Inclusive()
150 Ncv32u val = s_Data[threadIdx.x]; in scan1Inclusive()
173 const Ncv32u MAX_GRID_DIM = 65535;
176 const Ncv32u NUM_THREADS_ANCHORSPARALLEL = 64;
198 texture<Ncv32u, 1, cudaReadModeElementType> texIImage;
201 __device__ HaarStage64 getStage(Ncv32u iStage, HaarStage64 *d_Stages) in getStage()
208 __device__ HaarClassifierNode128 getClassifierNode(Ncv32u iNode, HaarClassifierNode128 *d_Classifie… in getClassifierNode()
224 __device__ void getFeature(Ncv32u iFeature, HaarFeature64 *d_Features, in getFeature()
226 Ncv32u *rectX, Ncv32u *rectY, Ncv32u *rectWidth, Ncv32u *rectHeight) in getFeature()
243 __device__ Ncv32u getElemIImg(Ncv32u x, Ncv32u *d_IImg) in getElemIImg()
256 __device__ Ncv32u d_outMaskPosition;
259 __device__ void compactBlockWriteOutAnchorParallel(Ncv32u threadPassFlag, Ncv32u threadElem, Ncv32u… in compactBlockWriteOutAnchorParallel()
263 __shared__ Ncv32u shmem[NUM_THREADS_ANCHORSPARALLEL * 2]; in compactBlockWriteOutAnchorParallel()
264 __shared__ Ncv32u numPassed; in compactBlockWriteOutAnchorParallel()
265 __shared__ Ncv32u outMaskOffset; in compactBlockWriteOutAnchorParallel()
267 Ncv32u incScan = scan1Inclusive<NUM_THREADS_ANCHORSPARALLEL>(threadPassFlag, shmem); in compactBlockWriteOutAnchorParallel()
278 Ncv32u excScan = incScan - threadPassFlag; in compactBlockWriteOutAnchorParallel()
297 __global__ void applyHaarClassifierAnchorParallel(Ncv32u *d_IImg, Ncv32u IImgStride, in applyHaarClassifierAnchorParallel()
298 Ncv32f *d_weights, Ncv32u weightsStride, in applyHaarClassifierAnchorParallel()
300 Ncv32u *d_inMask, Ncv32u *d_outMask, in applyHaarClassifierAnchorParallel()
301 Ncv32u mask1Dlen, Ncv32u mask2Dstride, in applyHaarClassifierAnchorParallel()
302 … NcvSize32u anchorsRoi, Ncv32u startStageInc, Ncv32u endStageExc, Ncv32f scaleArea) in applyHaarClassifierAnchorParallel()
304 Ncv32u y_offs; in applyHaarClassifierAnchorParallel()
305 Ncv32u x_offs; in applyHaarClassifierAnchorParallel()
306 Ncv32u maskOffset; in applyHaarClassifierAnchorParallel()
307 Ncv32u outMaskVal; in applyHaarClassifierAnchorParallel()
370 for (Ncv32u iStage = startStageInc; iStage < endStageExc; iStage++) in applyHaarClassifierAnchorParallel()
375 Ncv32u numRootNodesInStage = curStage.getNumClassifierRootNodes(); in applyHaarClassifierAnchorParallel()
376 Ncv32u curRootNodeOffset = curStage.getStartClassifierRootNodeOffset(); in applyHaarClassifierAnchorParallel()
382 Ncv32u iNode = curRootNodeOffset; in applyHaarClassifierAnchorParallel()
390 Ncv32u curNodeFeaturesNum = featuresDesc.getNumFeatures(); in applyHaarClassifierAnchorParallel()
391 Ncv32u iFeature = featuresDesc.getFeaturesOffset(); in applyHaarClassifierAnchorParallel()
395 for (Ncv32u iRect=0; iRect<curNodeFeaturesNum; iRect++) in applyHaarClassifierAnchorParallel()
398 Ncv32u rectX, rectY, rectWidth, rectHeight; in applyHaarClassifierAnchorParallel()
403 Ncv32u iioffsTL = (y_offs + rectY) * IImgStride + (x_offs + rectX); in applyHaarClassifierAnchorParallel()
404 Ncv32u iioffsTR = iioffsTL + rectWidth; in applyHaarClassifierAnchorParallel()
405 Ncv32u iioffsBL = iioffsTL + rectHeight * IImgStride; in applyHaarClassifierAnchorParallel()
406 Ncv32u iioffsBR = iioffsBL + rectWidth; in applyHaarClassifierAnchorParallel()
408 Ncv32u rectSum = getElemIImg<tbCacheTextureIImg>(iioffsBR, d_IImg) - in applyHaarClassifierAnchorParallel()
485 __global__ void applyHaarClassifierClassifierParallel(Ncv32u *d_IImg, Ncv32u IImgStride, in applyHaarClassifierClassifierParallel()
486 Ncv32f *d_weights, Ncv32u weightsStride, in applyHaarClassifierClassifierParallel()
488 Ncv32u *d_inMask, Ncv32u *d_outMask, in applyHaarClassifierClassifierParallel()
489 Ncv32u mask1Dlen, Ncv32u mask2Dstride, in applyHaarClassifierClassifierParallel()
490 … NcvSize32u anchorsRoi, Ncv32u startStageInc, Ncv32u endStageExc, Ncv32f scaleArea) in applyHaarClassifierClassifierParallel()
492 Ncv32u maskOffset = MAX_GRID_DIM * blockIdx.y + blockIdx.x; in applyHaarClassifierClassifierParallel()
499 Ncv32u outMaskVal = d_inMask[maskOffset]; in applyHaarClassifierClassifierParallel()
500 Ncv32u y_offs = outMaskVal >> 16; in applyHaarClassifierClassifierParallel()
501 Ncv32u x_offs = outMaskVal & 0xFFFF; in applyHaarClassifierClassifierParallel()
506 for (Ncv32u iStage = startStageInc; iStage<endStageExc; iStage++) in applyHaarClassifierClassifierParallel()
513 Ncv32u curRootNodeOffset = curStage.getStartClassifierRootNodeOffset() + threadIdx.x; in applyHaarClassifierClassifierParallel()
516 …Ncv32u numRootChunks = (numRootNodesInStage + NUM_THREADS_CLASSIFIERPARALLEL - 1) >> NUM_THREADS_C… in applyHaarClassifierClassifierParallel()
518 for (Ncv32u chunkId=0; chunkId<numRootChunks; chunkId++) in applyHaarClassifierClassifierParallel()
524 Ncv32u iNode = curRootNodeOffset; in applyHaarClassifierClassifierParallel()
530 Ncv32u curNodeFeaturesNum = featuresDesc.getNumFeatures(); in applyHaarClassifierClassifierParallel()
531 Ncv32u iFeature = featuresDesc.getFeaturesOffset(); in applyHaarClassifierClassifierParallel()
535 for (Ncv32u iRect=0; iRect<curNodeFeaturesNum; iRect++) in applyHaarClassifierClassifierParallel()
538 Ncv32u rectX, rectY, rectWidth, rectHeight; in applyHaarClassifierClassifierParallel()
543 Ncv32u iioffsTL = (y_offs + rectY) * IImgStride + (x_offs + rectX); in applyHaarClassifierClassifierParallel()
544 Ncv32u iioffsTR = iioffsTL + rectWidth; in applyHaarClassifierClassifierParallel()
545 Ncv32u iioffsBL = iioffsTL + rectHeight * IImgStride; in applyHaarClassifierClassifierParallel()
546 Ncv32u iioffsBR = iioffsBL + rectWidth; in applyHaarClassifierClassifierParallel()
548 Ncv32u rectSum = getElemIImg<tbCacheTextureIImg>(iioffsBR, d_IImg) - in applyHaarClassifierClassifierParallel()
620 Ncv32u outMaskOffset = atomicAdd(&d_outMaskPosition, 1); in applyHaarClassifierClassifierParallel()
630 __global__ void initializeMaskVector(Ncv32u *d_inMask, Ncv32u *d_outMask, in initializeMaskVector()
631 Ncv32u mask1Dlen, Ncv32u mask2Dstride, in initializeMaskVector()
632 NcvSize32u anchorsRoi, Ncv32u step) in initializeMaskVector()
634 Ncv32u y_offs = blockIdx.y; in initializeMaskVector()
635 Ncv32u x_offs = blockIdx.x * NUM_THREADS_ANCHORSPARALLEL + threadIdx.x; in initializeMaskVector()
636 Ncv32u outMaskOffset = y_offs * gridDim.x * blockDim.x + x_offs; in initializeMaskVector()
638 Ncv32u y_offs_upsc = step * y_offs; in initializeMaskVector()
639 Ncv32u x_offs_upsc = step * x_offs; in initializeMaskVector()
640 Ncv32u inMaskOffset = y_offs_upsc * mask2Dstride + x_offs_upsc; in initializeMaskVector()
642 Ncv32u outElem = OBJDET_MASK_ELEMENT_INVALID_32U; in initializeMaskVector()
669 Ncv32u *d_IImg;
670 Ncv32u IImgStride;
672 Ncv32u weightsStride;
676 Ncv32u *d_inMask;
677 Ncv32u *d_outMask;
678 Ncv32u mask1Dlen;
679 Ncv32u mask2Dstride;
681 Ncv32u startStageInc;
682 Ncv32u endStageExc;
687 Ncv32u *_d_IImg, Ncv32u _IImgStride, in applyHaarClassifierAnchorParallelFunctor()
688 Ncv32f *_d_weights, Ncv32u _weightsStride, in applyHaarClassifierAnchorParallelFunctor()
690 Ncv32u *_d_inMask, Ncv32u *_d_outMask, in applyHaarClassifierAnchorParallelFunctor()
691 Ncv32u _mask1Dlen, Ncv32u _mask2Dstride, in applyHaarClassifierAnchorParallelFunctor()
692 NcvSize32u _anchorsRoi, Ncv32u _startStageInc, in applyHaarClassifierAnchorParallelFunctor()
693 Ncv32u _endStageExc, Ncv32f _scaleArea) : in applyHaarClassifierAnchorParallelFunctor()
744 Ncv32u *d_IImg, Ncv32u IImgStride, in applyHaarClassifierAnchorParallelDynTemplate()
745 Ncv32f *d_weights, Ncv32u weightsStride, in applyHaarClassifierAnchorParallelDynTemplate()
747 Ncv32u *d_inMask, Ncv32u *d_outMask, in applyHaarClassifierAnchorParallelDynTemplate()
748 Ncv32u mask1Dlen, Ncv32u mask2Dstride, in applyHaarClassifierAnchorParallelDynTemplate()
749 NcvSize32u anchorsRoi, Ncv32u startStageInc, in applyHaarClassifierAnchorParallelDynTemplate()
750 Ncv32u endStageExc, Ncv32f scaleArea) in applyHaarClassifierAnchorParallelDynTemplate()
779 Ncv32u *d_IImg;
780 Ncv32u IImgStride;
782 Ncv32u weightsStride;
786 Ncv32u *d_inMask;
787 Ncv32u *d_outMask;
788 Ncv32u mask1Dlen;
789 Ncv32u mask2Dstride;
791 Ncv32u startStageInc;
792 Ncv32u endStageExc;
797 Ncv32u *_d_IImg, Ncv32u _IImgStride, in applyHaarClassifierClassifierParallelFunctor()
798 Ncv32f *_d_weights, Ncv32u _weightsStride, in applyHaarClassifierClassifierParallelFunctor()
800 Ncv32u *_d_inMask, Ncv32u *_d_outMask, in applyHaarClassifierClassifierParallelFunctor()
801 Ncv32u _mask1Dlen, Ncv32u _mask2Dstride, in applyHaarClassifierClassifierParallelFunctor()
802 NcvSize32u _anchorsRoi, Ncv32u _startStageInc, in applyHaarClassifierClassifierParallelFunctor()
803 Ncv32u _endStageExc, Ncv32f _scaleArea) : in applyHaarClassifierClassifierParallelFunctor()
850 Ncv32u *d_IImg, Ncv32u IImgStride, in applyHaarClassifierClassifierParallelDynTemplate()
851 Ncv32f *d_weights, Ncv32u weightsStride, in applyHaarClassifierClassifierParallelDynTemplate()
853 Ncv32u *d_inMask, Ncv32u *d_outMask, in applyHaarClassifierClassifierParallelDynTemplate()
854 Ncv32u mask1Dlen, Ncv32u mask2Dstride, in applyHaarClassifierClassifierParallelDynTemplate()
855 NcvSize32u anchorsRoi, Ncv32u startStageInc, in applyHaarClassifierClassifierParallelDynTemplate()
856 Ncv32u endStageExc, Ncv32f scaleArea) in applyHaarClassifierClassifierParallelDynTemplate()
882 Ncv32u *d_inMask;
883 Ncv32u *d_outMask;
884 Ncv32u mask1Dlen;
885 Ncv32u mask2Dstride;
887 Ncv32u step;
891 Ncv32u *_d_inMask, Ncv32u *_d_outMask, in initializeMaskVectorFunctor()
892 Ncv32u _mask1Dlen, Ncv32u _mask2Dstride, in initializeMaskVectorFunctor()
893 NcvSize32u _anchorsRoi, Ncv32u _step) : in initializeMaskVectorFunctor()
925 Ncv32u *d_inMask, Ncv32u *d_outMask, in initializeMaskVectorDynTemplate()
926 Ncv32u mask1Dlen, Ncv32u mask2Dstride, in initializeMaskVectorDynTemplate()
927 NcvSize32u anchorsRoi, Ncv32u step) in initializeMaskVectorDynTemplate()
942 Ncv32u getStageNumWithNotLessThanNclassifiers(Ncv32u N, HaarClassifierCascadeDescriptor &haar, in getStageNumWithNotLessThanNclassifiers()
945 Ncv32u i = 0; in getStageNumWithNotLessThanNclassifiers()
957 NCVStatus ncvApplyHaarClassifierCascade_device(NCVMatrix<Ncv32u> &integral, in ncvApplyHaarClassifierCascade_device()
959 NCVMatrixAlloc<Ncv32u> &d_pixelMask, in ncvApplyHaarClassifierCascade_device()
960 Ncv32u &numDetections, in ncvApplyHaarClassifierCascade_device()
968 Ncv32u pixelStep, in ncvApplyHaarClassifierCascade_device()
1018 …NCVMatrixAlloc<Ncv32u> h_integralImage(cpuAllocator, integral.width, integral.height, integral.pit… in ncvApplyHaarClassifierCascade_device()
1022 …NCVMatrixAlloc<Ncv32u> h_pixelMask(cpuAllocator, d_pixelMask.width, d_pixelMask.height, d_pixelMas… in ncvApplyHaarClassifierCascade_device()
1029 …NCVMatrixAlloc<Ncv32u> h_pixelMask_d(cpuAllocator, d_pixelMask.width, d_pixelMask.height, d_pixelM… in ncvApplyHaarClassifierCascade_device()
1046 for (Ncv32u i=0; i<(Ncv32u)anchorsRoi.height; i++) in ncvApplyHaarClassifierCascade_device()
1048 for (Ncv32u j=0; j<d_pixelMask.stride(); j++) in ncvApplyHaarClassifierCascade_device()
1050 if ((i%pixelStep==0) && (j%pixelStep==0) && (j<(Ncv32u)anchorsRoi.width)) in ncvApplyHaarClassifierCascade_device()
1068 …NCVVectorReuse<Ncv32u> d_vecPixelMask(d_pixelMask.getSegment(), anchorsRoi.height * d_pixelMask.st… in ncvApplyHaarClassifierCascade_device()
1071 …NCVVectorAlloc<Ncv32u> d_vecPixelMaskTmp(gpuAllocator, static_cast<Ncv32u>(d_vecPixelMask.length()… in ncvApplyHaarClassifierCascade_device()
1074 NCVVectorAlloc<Ncv32u> hp_pool32u(cpuAllocator, 2); in ncvApplyHaarClassifierCascade_device()
1076 Ncv32u *hp_zero = &hp_pool32u.ptr()[0]; in ncvApplyHaarClassifierCascade_device()
1077 Ncv32u *hp_numDet = &hp_pool32u.ptr()[1]; in ncvApplyHaarClassifierCascade_device()
1091 NCVVector<Ncv32u> *d_ptrNowData = &d_vecPixelMask; in ncvApplyHaarClassifierCascade_device()
1092 NCVVector<Ncv32u> *d_ptrNowTmp = &d_vecPixelMaskTmp; in ncvApplyHaarClassifierCascade_device()
1094 Ncv32u szNppCompactTmpBuf; in ncvApplyHaarClassifierCascade_device()
1095 …nppsStCompactGetSize_32u(static_cast<Ncv32u>(d_vecPixelMask.length()), &szNppCompactTmpBuf, devPro… in ncvApplyHaarClassifierCascade_device()
1107 cfdTexIImage = cudaCreateChannelDesc<Ncv32u>(); in ncvApplyHaarClassifierCascade_device()
1131 Ncv32u stageStartAnchorParallel = 0; in ncvApplyHaarClassifierCascade_device()
1132 … Ncv32u stageMiddleSwitch = getStageNumWithNotLessThanNclassifiers(NUM_THREADS_CLASSIFIERPARALLEL, in ncvApplyHaarClassifierCascade_device()
1134 Ncv32u stageEndClassifierParallel = haar.NumStages; in ncvApplyHaarClassifierCascade_device()
1141 const Ncv32u compactEveryNstage = bDoAtomicCompaction ? 7 : 1; in ncvApplyHaarClassifierCascade_device()
1142 Ncv32u curStop = stageStartAnchorParallel; in ncvApplyHaarClassifierCascade_device()
1143 std::vector<Ncv32u> pixParallelStageStops; in ncvApplyHaarClassifierCascade_device()
1155 Ncv32u pixParallelStageStopsIndex = 0; in ncvApplyHaarClassifierCascade_device()
1161 ncvAssertCUDAReturn(cudaMemcpyToSymbolAsync(d_outMaskPosition, hp_zero, sizeof(Ncv32u), in ncvApplyHaarClassifierCascade_device()
1181 … static_cast<Ncv32u>(d_vecPixelMask.length()), d_pixelMask.stride(), in ncvApplyHaarClassifierCascade_device()
1188 … ncvAssertCUDAReturn(cudaMemcpyFromSymbolAsync(hp_numDet, d_outMaskPosition, sizeof(Ncv32u), in ncvApplyHaarClassifierCascade_device()
1196 … nppSt = nppsStCompact_32u(d_ptrNowTmp->ptr(), static_cast<Ncv32u>(d_vecPixelMask.length()), in ncvApplyHaarClassifierCascade_device()
1211 ncvAssertCUDAReturn(cudaMemcpyToSymbolAsync(d_outMaskPosition, hp_zero, sizeof(Ncv32u), in ncvApplyHaarClassifierCascade_device()
1244 … ncvAssertCUDAReturn(cudaMemcpyFromSymbolAsync(hp_numDet, d_outMaskPosition, sizeof(Ncv32u), in ncvApplyHaarClassifierCascade_device()
1251 … nppSt = nppsStCompact_32u(d_ptrNowData->ptr(), static_cast<Ncv32u>(d_vecPixelMask.length()), in ncvApplyHaarClassifierCascade_device()
1276 ncvAssertCUDAReturn(cudaMemcpyToSymbolAsync(d_outMaskPosition, hp_zero, sizeof(Ncv32u), in ncvApplyHaarClassifierCascade_device()
1314 … ncvAssertCUDAReturn(cudaMemcpyFromSymbolAsync(hp_numDet, d_outMaskPosition, sizeof(Ncv32u), in ncvApplyHaarClassifierCascade_device()
1339 ncvAssertCUDAReturn(cudaMemcpyToSymbolAsync(d_outMaskPosition, hp_zero, sizeof(Ncv32u), in ncvApplyHaarClassifierCascade_device()
1375 … ncvAssertCUDAReturn(cudaMemcpyFromSymbolAsync(hp_numDet, d_outMaskPosition, sizeof(Ncv32u), in ncvApplyHaarClassifierCascade_device()
1409 Ncv32u fpu_oldcw, fpu_cw; in ncvApplyHaarClassifierCascade_device()
1413 Ncv32u numDetGold; in ncvApplyHaarClassifierCascade_device()
1429 for (Ncv32u i=0; i<std::max(numDetGold, numDetections) && bPass; i++) in ncvApplyHaarClassifierCascade_device()
1455 const Ncv32u NUM_GROW_THREADS = 128;
1458 __device__ __host__ NcvRect32u pixelToRect(Ncv32u pixel, Ncv32u width, Ncv32u height, Ncv32f scale) in pixelToRect()
1461 res.x = (Ncv32u)(scale * (pixel & 0xFFFF)); in pixelToRect()
1462 res.y = (Ncv32u)(scale * (pixel >> 16)); in pixelToRect()
1463 res.width = (Ncv32u)(scale * width); in pixelToRect()
1464 res.height = (Ncv32u)(scale * height); in pixelToRect()
1469 __global__ void growDetectionsKernel(Ncv32u *pixelMask, Ncv32u numElements, in growDetectionsKernel()
1471 Ncv32u rectWidth, Ncv32u rectHeight, Ncv32f curScale) in growDetectionsKernel()
1473 Ncv32u blockId = blockIdx.y * 65535 + blockIdx.x; in growDetectionsKernel()
1474 Ncv32u elemAddr = blockId * NUM_GROW_THREADS + threadIdx.x; in growDetectionsKernel()
1483 NCVStatus ncvGrowDetectionsVector_device(NCVVector<Ncv32u> &pixelMask, in ncvGrowDetectionsVector_device()
1484 Ncv32u numPixelMaskDetections, in ncvGrowDetectionsVector_device()
1486 Ncv32u &totalDetections, in ncvGrowDetectionsVector_device()
1487 Ncv32u totalMaxDetections, in ncvGrowDetectionsVector_device()
1488 Ncv32u rectWidth, in ncvGrowDetectionsVector_device()
1489 Ncv32u rectHeight, in ncvGrowDetectionsVector_device()
1507 Ncv32u numDetsToCopy = numPixelMaskDetections; in ncvGrowDetectionsVector_device()
1547 Ncv32u &dstNumRects, in ncvDetectObjectsMultiScale_device()
1556 Ncv32u minNeighbors, //default 4 in ncvDetectObjectsMultiScale_device()
1558 Ncv32u pixelStep, //default 1 in ncvDetectObjectsMultiScale_device()
1559 … Ncv32u flags, //default NCVPipeObjDet_Default in ncvDetectObjectsMultiScale_device()
1605 Ncv32u integralWidth = d_srcImg.width() + 1; in ncvDetectObjectsMultiScale_device()
1606 Ncv32u integralHeight = d_srcImg.height() + 1; in ncvDetectObjectsMultiScale_device()
1608 NCVMatrixAlloc<Ncv32u> integral(gpuAllocator, integralWidth, integralHeight); in ncvDetectObjectsMultiScale_device()
1615 NCVMatrixAlloc<Ncv32u> d_pixelMask(gpuAllocator, d_srcImg.width(), d_srcImg.height()); in ncvDetectObjectsMultiScale_device()
1618 NCVMatrixAlloc<Ncv32u> d_scaledIntegralImage(gpuAllocator, integralWidth, integralHeight); in ncvDetectObjectsMultiScale_device()
1629 Ncv32u szTmpBufIntegral, szTmpBufSqIntegral; in ncvDetectObjectsMultiScale_device()
1655 Ncv32u lastCheckedScale = 0; in ncvDetectObjectsMultiScale_device()
1657 std::vector<Ncv32u> scalesVector; in ncvDetectObjectsMultiScale_device()
1663 Ncv32u scale = (Ncv32u)scaleIter; in ncvDetectObjectsMultiScale_device()
1709 for (Ncv32u i=0; i<scalesVector.size(); i++) in ncvDetectObjectsMultiScale_device()
1711 Ncv32u scale = scalesVector[i]; in ncvDetectObjectsMultiScale_device()
1753 Ncv32u detectionsOnThisScale; in ncvDetectObjectsMultiScale_device()
1764 NCVVectorReuse<Ncv32u> d_vecPixelMask(d_pixelMask.getSegment()); in ncvDetectObjectsMultiScale_device()
1770 static_cast<Ncv32u>(d_hypothesesIntermediate.length()), in ncvDetectObjectsMultiScale_device()
1793 Ncv32u numStrongHypothesesNow = dstNumRects; in ncvDetectObjectsMultiScale_device()
1805 for (Ncv32u j=1; j<numStrongHypothesesNow; j++) in ncvDetectObjectsMultiScale_device()
1867 dstNumRects = static_cast<Ncv32u>(d_dstRects.length()); in ncvDetectObjectsMultiScale_device()
1904 NCVStatus ncvApplyHaarClassifierCascade_host(NCVMatrix<Ncv32u> &h_integralImage, in ncvApplyHaarClassifierCascade_host()
1906 NCVMatrixAlloc<Ncv32u> &h_pixelMask, in ncvApplyHaarClassifierCascade_host()
1907 Ncv32u &numDetections, in ncvApplyHaarClassifierCascade_host()
1914 Ncv32u pixelStep, in ncvApplyHaarClassifierCascade_host()
1944 for (Ncv32u i=0; i<anchorsRoi.height; i++) in ncvApplyHaarClassifierCascade_host()
1946 for (Ncv32u j=0; j<h_pixelMask.stride(); j++) in ncvApplyHaarClassifierCascade_host()
1954 for (Ncv32u iStage = 0; iStage < haar.NumStages; iStage++) in ncvApplyHaarClassifierCascade_host()
1957 … Ncv32u numRootNodesInStage = h_HaarStages.ptr()[iStage].getNumClassifierRootNodes(); in ncvApplyHaarClassifierCascade_host()
1958 … Ncv32u curRootNodeOffset = h_HaarStages.ptr()[iStage].getStartClassifierRootNodeOffset(); in ncvApplyHaarClassifierCascade_host()
1979 Ncv32u curNodeOffset = curRootNodeOffset; in ncvApplyHaarClassifierCascade_host()
1985 Ncv32u curNodeFeaturesNum = curFeatDesc.getNumFeatures(); in ncvApplyHaarClassifierCascade_host()
1986 Ncv32u curNodeFeaturesOffs = curFeatDesc.getFeaturesOffset(); in ncvApplyHaarClassifierCascade_host()
1989 for (Ncv32u iRect=0; iRect<curNodeFeaturesNum; iRect++) in ncvApplyHaarClassifierCascade_host()
1992 Ncv32u rectX, rectY, rectWidth, rectHeight; in ncvApplyHaarClassifierCascade_host()
1995 … Ncv32u iioffsTL = (i + rectY) * h_integralImage.stride() + (j + rectX); in ncvApplyHaarClassifierCascade_host()
1996 Ncv32u iioffsTR = iioffsTL + rectWidth; in ncvApplyHaarClassifierCascade_host()
1997 Ncv32u iioffsBL = iioffsTL + rectHeight * h_integralImage.stride(); in ncvApplyHaarClassifierCascade_host()
1998 Ncv32u iioffsBR = iioffsBL + rectWidth; in ncvApplyHaarClassifierCascade_host()
2000 Ncv32u iivalTL = h_integralImage.ptr()[iioffsTL]; in ncvApplyHaarClassifierCascade_host()
2001 Ncv32u iivalTR = h_integralImage.ptr()[iioffsTR]; in ncvApplyHaarClassifierCascade_host()
2002 Ncv32u iivalBL = h_integralImage.ptr()[iioffsBL]; in ncvApplyHaarClassifierCascade_host()
2003 Ncv32u iivalBR = h_integralImage.ptr()[iioffsBR]; in ncvApplyHaarClassifierCascade_host()
2004 Ncv32u rectSum = iivalBR - iivalBL + iivalTL - iivalTR; in ncvApplyHaarClassifierCascade_host()
2054 Ncv32u i = 0; in ncvApplyHaarClassifierCascade_host()
2068 NCVStatus ncvGrowDetectionsVector_host(NCVVector<Ncv32u> &pixelMask, in ncvGrowDetectionsVector_host()
2069 Ncv32u numPixelMaskDetections, in ncvGrowDetectionsVector_host()
2071 Ncv32u &totalDetections, in ncvGrowDetectionsVector_host()
2072 Ncv32u totalMaxDetections, in ncvGrowDetectionsVector_host()
2073 Ncv32u rectWidth, in ncvGrowDetectionsVector_host()
2074 Ncv32u rectHeight, in ncvGrowDetectionsVector_host()
2087 Ncv32u numDetsToCopy = numPixelMaskDetections; in ncvGrowDetectionsVector_host()
2100 for (Ncv32u i=0; i<numDetsToCopy; i++) in ncvGrowDetectionsVector_host()
2134 Ncv32u curMaxTreeDepth = 0; in loadFromXML()
2154 curStage.setStartClassifierRootNodeOffset(static_cast<Ncv32u>(haarClassifierNodes.size())); in loadFromXML()
2161 Ncv32u nodeId = 0; in loadFromXML()
2185 Ncv32u leftNodeOffset = tree->left[n]; in loadFromXML()
2186 … nodeLeft.create((Ncv32u)(h_TmpClassifierNotRootNodes.size() + leftNodeOffset - 1)); in loadFromXML()
2201 Ncv32u rightNodeOffset = tree->right[n]; in loadFromXML()
2202 … nodeRight.create((Ncv32u)(h_TmpClassifierNotRootNodes.size() + rightNodeOffset - 1)); in loadFromXML()
2207 Ncv32u tiltedVal = feature->tilted; in loadFromXML()
2210 Ncv32u featureId = 0; in loadFromXML()
2213 Ncv32u rectX = feature->rect[l].r.x; in loadFromXML()
2214 Ncv32u rectY = feature->rect[l].r.y; in loadFromXML()
2215 Ncv32u rectWidth = feature->rect[l].r.width; in loadFromXML()
2216 Ncv32u rectHeight = feature->rect[l].r.height; in loadFromXML()
2234 featureId, static_cast<Ncv32u>(haarFeatures.size()) - featureId); in loadFromXML()
2260 haar.NumStages = static_cast<Ncv32u>(haarStages.size()); in loadFromXML()
2261 haar.NumClassifierRootNodes = static_cast<Ncv32u>(haarClassifierNodes.size()); in loadFromXML()
2262 …haar.NumClassifierTotalNodes = static_cast<Ncv32u>(haar.NumClassifierRootNodes + h_TmpClassifierNo… in loadFromXML()
2263 haar.NumFeatures = static_cast<Ncv32u>(haarFeatures.size()); in loadFromXML()
2266 Ncv32u offsetRoot = static_cast<Ncv32u>(haarClassifierNodes.size()); in loadFromXML()
2267 for (Ncv32u i=0; i<haarClassifierNodes.size(); i++) in loadFromXML()
2274 Ncv32u newOffset = nodeLeft.getNextNodeOffset() + offsetRoot; in loadFromXML()
2282 Ncv32u newOffset = nodeRight.getNextNodeOffset() + offsetRoot; in loadFromXML()
2288 for (Ncv32u i=0; i<h_TmpClassifierNotRootNodes.size(); i++) in loadFromXML()
2295 Ncv32u newOffset = nodeLeft.getNextNodeOffset() + offsetRoot; in loadFromXML()
2303 Ncv32u newOffset = nodeRight.getNextNodeOffset() + offsetRoot; in loadFromXML()
2329 Ncv32u fileVersion; in loadFromNVBIN()
2330 readCount = fread(&fileVersion, sizeof(Ncv32u), 1, fp); in loadFromNVBIN()
2333 Ncv32u fsize; in loadFromNVBIN()
2334 readCount = fread(&fsize, sizeof(Ncv32u), 1, fp); in loadFromNVBIN()
2337 Ncv32u fsizeActual = ftell(fp); in loadFromNVBIN()
2342 Ncv32u dataOffset = 0; in loadFromNVBIN()
2350 haar.NumStages = *(Ncv32u *)(&fdata[0]+dataOffset); in loadFromNVBIN()
2351 dataOffset += sizeof(Ncv32u); in loadFromNVBIN()
2352 haar.NumClassifierRootNodes = *(Ncv32u *)(&fdata[0]+dataOffset); in loadFromNVBIN()
2353 dataOffset += sizeof(Ncv32u); in loadFromNVBIN()
2354 haar.NumClassifierTotalNodes = *(Ncv32u *)(&fdata[0]+dataOffset); in loadFromNVBIN()
2355 dataOffset += sizeof(Ncv32u); in loadFromNVBIN()
2356 haar.NumFeatures = *(Ncv32u *)(&fdata[0]+dataOffset); in loadFromNVBIN()
2357 dataOffset += sizeof(Ncv32u); in loadFromNVBIN()
2369 Ncv32u szStages = haar.NumStages * sizeof(HaarStage64); in loadFromNVBIN()
2370 Ncv32u szClassifiers = haar.NumClassifierTotalNodes * sizeof(HaarClassifierNode128); in loadFromNVBIN()
2371 Ncv32u szFeatures = haar.NumFeatures * sizeof(HaarFeature64); in loadFromNVBIN()
2384 NCVStatus ncvHaarGetClassifierSize(const cv::String &filename, Ncv32u &numStages, in ncvHaarGetClassifierSize()
2385 Ncv32u &numNodes, Ncv32u &numFeatures) in ncvHaarGetClassifierSize()
2397 Ncv32u fileVersion; in ncvHaarGetClassifierSize()
2398 readCount = fread(&fileVersion, sizeof(Ncv32u), 1, fp); in ncvHaarGetClassifierSize()
2402 Ncv32u tmp; in ncvHaarGetClassifierSize()
2403 readCount = fread(&numStages, sizeof(Ncv32u), 1, fp); in ncvHaarGetClassifierSize()
2405 readCount = fread(&tmp, sizeof(Ncv32u), 1, fp); in ncvHaarGetClassifierSize()
2407 readCount = fread(&numNodes, sizeof(Ncv32u), 1, fp); in ncvHaarGetClassifierSize()
2409 readCount = fread(&numFeatures, sizeof(Ncv32u), 1, fp); in ncvHaarGetClassifierSize()
2495 Ncv32u szStages = haar.NumStages * sizeof(HaarStage64); in ncvHaarStoreNVBIN_host()
2496 Ncv32u szClassifiers = haar.NumClassifierTotalNodes * sizeof(HaarClassifierNode128); in ncvHaarStoreNVBIN_host()
2497 Ncv32u szFeatures = haar.NumFeatures * sizeof(HaarFeature64); in ncvHaarStoreNVBIN_host()
2499 Ncv32u dataOffset = 0; in ncvHaarStoreNVBIN_host()
2504 *(Ncv32u *)(&fdata[0]+dataOffset) = NVBIN_HAAR_VERSION; in ncvHaarStoreNVBIN_host()
2508 *(Ncv32u *)(&fdata[0]+dataOffset) = haar.NumStages; in ncvHaarStoreNVBIN_host()
2509 dataOffset += sizeof(Ncv32u); in ncvHaarStoreNVBIN_host()
2510 *(Ncv32u *)(&fdata[0]+dataOffset) = haar.NumClassifierRootNodes; in ncvHaarStoreNVBIN_host()
2511 dataOffset += sizeof(Ncv32u); in ncvHaarStoreNVBIN_host()
2512 *(Ncv32u *)(&fdata[0]+dataOffset) = haar.NumClassifierTotalNodes; in ncvHaarStoreNVBIN_host()
2513 dataOffset += sizeof(Ncv32u); in ncvHaarStoreNVBIN_host()
2514 *(Ncv32u *)(&fdata[0]+dataOffset) = haar.NumFeatures; in ncvHaarStoreNVBIN_host()
2515 dataOffset += sizeof(Ncv32u); in ncvHaarStoreNVBIN_host()
2529 Ncv32u fsize = dataOffset; in ncvHaarStoreNVBIN_host()
2534 dataOffset = sizeof(Ncv32u); in ncvHaarStoreNVBIN_host()
2535 *(Ncv32u *)(&fdata[0]+dataOffset) = fsize; in ncvHaarStoreNVBIN_host()