Skip to content

Commit

Permalink
Merge moved code from opencv/3.4
Browse files Browse the repository at this point in the history
  • Loading branch information
@alalek
alalek committed Feb 21, 2021
2 parents 3ea4443 + fdf80d7 commit 8ab0020
Show file tree
Hide file tree
Showing 6 changed files with 51 additions and 51 deletions.
89 changes: 38 additions & 51 deletions modules/cudafilters/src/cuda/median_filter.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -55,9 +55,6 @@

namespace cv { namespace cuda { namespace device
{
// // namespace imgproc
// {

__device__ void histogramAddAndSub8(int* H, const int * hist_colAdd,const int * hist_colSub){
int tx = threadIdx.x;
if (tx<8){
Expand Down Expand Up @@ -125,35 +122,40 @@ namespace cv { namespace cuda { namespace device
luc[tx]=0;
}

#define scanNeighbor(array, range, index, threadIndex) \
{ \
int v = 0; \
if (index <= threadIndex && threadIndex < range) \
v = array[threadIndex] + array[threadIndex-index]; \
__syncthreads(); \
if (index <= threadIndex && threadIndex < range) \
array[threadIndex] = v; \
}
#define findMedian(array, range, threadIndex, result, count, position) \
if (threadIndex < range) \
{ \
if (array[threadIndex+1] > position && array[threadIndex] <= position) \
{ \
*result = threadIndex+1; \
*count = array[threadIndex]; \
} \
}

__device__ void histogramMedianPar8LookupOnly(int* H,int* Hscan, const int medPos,int* retval, int* countAtMed){
int tx=threadIdx.x;
*retval=*countAtMed=0;
if(tx<8){
Hscan[tx]=H[tx];
}
__syncthreads();
if (1 <= tx && tx < 8 )
Hscan[tx]+=Hscan[tx-1];
scanNeighbor(Hscan, 8, 1, tx);
__syncthreads();
if (2 <= tx && tx < 8 )
Hscan[tx]+=Hscan[tx-2];
scanNeighbor(Hscan, 8, 2, tx);
__syncthreads();
if (4 <= tx && tx < 8 )
Hscan[tx]+=Hscan[tx-4];
scanNeighbor(Hscan, 8, 4, tx);
__syncthreads();

if(tx<7){
if(Hscan[tx+1] > medPos && Hscan[tx] < medPos){
*retval=tx+1;
*countAtMed=Hscan[tx];
}
else if(Hscan[tx]==medPos){
if(Hscan[tx+1]>medPos){
*retval=tx+1;
*countAtMed=Hscan[tx];
}
}
}
findMedian(Hscan, 7, tx, retval, countAtMed, medPos);
}

__device__ void histogramMedianPar32LookupOnly(int* H,int* Hscan, const int medPos,int* retval, int* countAtMed){
Expand All @@ -163,33 +165,18 @@ namespace cv { namespace cuda { namespace device
Hscan[tx]=H[tx];
}
__syncthreads();
if ( 1 <= tx && tx < 32 )
Hscan[tx]+=Hscan[tx-1];
scanNeighbor(Hscan, 32, 1, tx);
__syncthreads();
if ( 2 <= tx && tx < 32 )
Hscan[tx]+=Hscan[tx-2];
scanNeighbor(Hscan, 32, 2, tx);
__syncthreads();
if ( 4 <= tx && tx < 32 )
Hscan[tx]+=Hscan[tx-4];
scanNeighbor(Hscan, 32, 4, tx);
__syncthreads();
if ( 8 <= tx && tx < 32 )
Hscan[tx]+=Hscan[tx-8];
scanNeighbor(Hscan, 32, 8, tx);
__syncthreads();
if ( 16 <= tx && tx < 32 )
Hscan[tx]+=Hscan[tx-16];
scanNeighbor(Hscan, 32, 16, tx);
__syncthreads();
if(tx<31){
if(Hscan[tx+1] > medPos && Hscan[tx] < medPos){
*retval=tx+1;
*countAtMed=Hscan[tx];
}
else if(Hscan[tx]==medPos){
if(Hscan[tx+1]>medPos){
*retval=tx+1;
*countAtMed=Hscan[tx];
}
}
}

findMedian(Hscan, 31, tx, retval, countAtMed, medPos);
}

__global__ void cuMedianFilterMultiBlock(PtrStepSzb src, PtrStepSzb dest, PtrStepSzi histPar, PtrStepSzi coarseHistGrid,int r, int medPos_)
Expand Down Expand Up @@ -288,7 +275,6 @@ namespace cv { namespace cuda { namespace device
__syncthreads();

histogramMultipleAdd8(HCoarse,histCoarse, 2*r+1);
// __syncthreads();
int cols_m_1=cols-1;

for(int j=r;j<cols-r;j++){
Expand All @@ -300,23 +286,24 @@ namespace cv { namespace cuda { namespace device
histogramMedianPar8LookupOnly(HCoarse,HCoarseScan,medPos, &firstBin,&countAtMed);
__syncthreads();

if ( luc[firstBin] <= (j-r))
int loopIndex = luc[firstBin];
if (loopIndex <= (j-r))
{
histogramClear32(HFine[firstBin]);
for ( luc[firstBin] = j-r; luc[firstBin] < ::min(j+r+1,cols); luc[firstBin]++ ){
histogramAdd32(HFine[firstBin], hist+(luc[firstBin]*256+(firstBin<<5) ) );
for ( loopIndex = j-r; loopIndex < ::min(j+r+1,cols); loopIndex++ ){
histogramAdd32(HFine[firstBin], hist+(loopIndex*256+(firstBin<<5) ) );
}
}
else{
for ( ; luc[firstBin] < (j+r+1);luc[firstBin]++ ) {
for ( ; loopIndex < (j+r+1);loopIndex++ ) {
histogramAddAndSub32(HFine[firstBin],
hist+(::min(luc[firstBin],cols_m_1)*256+(firstBin<<5) ),
hist+(::max(luc[firstBin]-2*r-1,0)*256+(firstBin<<5) ) );
hist+(::min(loopIndex,cols_m_1)*256+(firstBin<<5) ),
hist+(::max(loopIndex-2*r-1,0)*256+(firstBin<<5) ) );
__syncthreads();

}
}
__syncthreads();
luc[firstBin] = loopIndex;

int leftOver=medPos-countAtMed;
if(leftOver>=0){
Expand Down
2 changes: 2 additions & 0 deletions modules/cudaoptflow/src/brox.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -64,6 +64,8 @@ namespace {
{
}

virtual String getDefaultName() const { return "DenseOpticalFlow.BroxOpticalFlow"; }

virtual void calc(InputArray I0, InputArray I1, InputOutputArray flow, Stream& stream);

virtual double getFlowSmoothness() const { return alpha_; }
Expand Down
2 changes: 2 additions & 0 deletions modules/cudaoptflow/src/farneback.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -129,6 +129,8 @@ namespace

virtual void calc(InputArray I0, InputArray I1, InputOutputArray flow, Stream& stream);

virtual String getDefaultName() const { return "DenseOpticalFlow.FarnebackOpticalFlow"; }

private:
int numLevels_;
double pyrScale_;
Expand Down
4 changes: 4 additions & 0 deletions modules/cudaoptflow/src/pyrlk.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -347,6 +347,8 @@ namespace
sparse(prevImg, nextImg, prevPts, nextPts, status, err, stream);
}
}

virtual String getDefaultName() const { return "SparseOpticalFlow.SparsePyrLKOpticalFlow"; }
};

class DensePyrLKOpticalFlowImpl : public DensePyrLKOpticalFlow, private PyrLKOpticalFlowBase
Expand Down Expand Up @@ -388,6 +390,8 @@ namespace
GpuMat flows[] = {u, v};
cuda::merge(flows, 2, _flow, stream);
}

virtual String getDefaultName() const { return "DenseOpticalFlow.DensePyrLKOpticalFlow"; }
};
}

Expand Down
3 changes: 3 additions & 0 deletions modules/cudaoptflow/src/tvl1flow.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -119,6 +119,9 @@ namespace

virtual void calc(InputArray I0, InputArray I1, InputOutputArray flow, Stream& stream);

virtual String getDefaultName() const { return "DenseOpticalFlow.OpticalFlowDual_TVL1"; }


private:
double tau_;
double lambda_;
Expand Down
2 changes: 2 additions & 0 deletions modules/optflow/src/tvl1flow.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -99,6 +99,8 @@ class OpticalFlowDual_TVL1 : public DualTVL1OpticalFlow
}
OpticalFlowDual_TVL1();

virtual String getDefaultName() const CV_OVERRIDE { return "DenseOpticalFlow.DualTVL1OpticalFlow"; }

void calc(InputArray I0, InputArray I1, InputOutputArray flow) CV_OVERRIDE;
void collectGarbage() CV_OVERRIDE;

Expand Down

0 comments on commit 8ab0020

Please sign in to comment.