third_party/opencv
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

fix test failure of cudafilters Median_Accuracy

Browse Source 
* avoid race condition
This commit is contained in:
Tomoaki Teshima 2018-10-03 00:59:08 +09:00
parent a1ffc00ba8
commit fa8684b5c8
1 changed files with 22 additions and 20 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

42
modules/cudafilters/src/cuda/median_filter.cu
View File

@ -127,14 +127,14 @@ namespace cv { namespace cuda { namespace device
Hscan[tx]=H[tx];
}
__syncthreads();
if(tx<8){
if(tx>=1 )
Hscan[tx]+=Hscan[tx-1];
if(tx>=2)
Hscan[tx]+=Hscan[tx-2];
if(tx>=4)
Hscan[tx]+=Hscan[tx-4];
}
if (1 <= tx && tx < 8 )
Hscan[tx]+=Hscan[tx-1];
__syncthreads();
if (2 <= tx && tx < 8 )
Hscan[tx]+=Hscan[tx-2];
__syncthreads();
if (4 <= tx && tx < 8 )
Hscan[tx]+=Hscan[tx-4];
__syncthreads();
if(tx<7){
@ -158,18 +158,20 @@ namespace cv { namespace cuda { namespace device
Hscan[tx]=H[tx];
}
__syncthreads();
if(tx<32){
if(tx>=1)
Hscan[tx]+=Hscan[tx-1];
if(tx>=2)
Hscan[tx]+=Hscan[tx-2];
if(tx>=4)
Hscan[tx]+=Hscan[tx-4];
if(tx>=8)
Hscan[tx]+=Hscan[tx-8];
if(tx>=16)
Hscan[tx]+=Hscan[tx-16];
}
if ( 1 <= tx && tx < 32 )
Hscan[tx]+=Hscan[tx-1];
__syncthreads();
if ( 2 <= tx && tx < 32 )
Hscan[tx]+=Hscan[tx-2];
__syncthreads();
if ( 4 <= tx && tx < 32 )
Hscan[tx]+=Hscan[tx-4];
__syncthreads();
if ( 8 <= tx && tx < 32 )
Hscan[tx]+=Hscan[tx-8];
__syncthreads();
if ( 16 <= tx && tx < 32 )
Hscan[tx]+=Hscan[tx-16];
__syncthreads();
if(tx<31){
if(Hscan[tx+1] > medPos && Hscan[tx] < medPos){