Skip to content

O
opencv_contrib
Commit b8dc51ad authored by Tomoaki Teshima's avatar Tomoaki Teshima
Browse files
Options

[moved from opencv] fix test failure of cudafilters Median_Accuracy 

  * avoid race condition

original commit: https://github.com/opencv/opencv/commit/fa8684b5c8267f1c0c89f6279b9b0ff87500480c
parent cfab48b4
Hide whitespace changes
Inline Side-by-side
Showing with 22 additions and 20 deletions
modules/cudafilters/src/cuda/median_filter.cu
View file @ b8dc51ad
...@@ -127,14 +127,14 @@ namespace cv { namespace cuda { namespace device ...@@ -127,14 +127,14 @@ namespace cv { namespace cuda { namespace device
Hscan[tx]=H[tx]; Hscan[tx]=H[tx];
} }
__syncthreads(); __syncthreads();
if(tx<8){ if (1 <= tx && tx < 8 )
if(tx>=1 ) Hscan[tx]+=Hscan[tx-1];
Hscan[tx]+=Hscan[tx-1]; __syncthreads();
if(tx>=2) if (2 <= tx && tx < 8 )
Hscan[tx]+=Hscan[tx-2]; Hscan[tx]+=Hscan[tx-2];
if(tx>=4) __syncthreads();
Hscan[tx]+=Hscan[tx-4]; if (4 <= tx && tx < 8 )
} Hscan[tx]+=Hscan[tx-4];
__syncthreads(); __syncthreads();
if(tx<7){ if(tx<7){
...@@ -158,18 +158,20 @@ namespace cv { namespace cuda { namespace device ...@@ -158,18 +158,20 @@ namespace cv { namespace cuda { namespace device
Hscan[tx]=H[tx]; Hscan[tx]=H[tx];
} }
__syncthreads(); __syncthreads();
if(tx<32){ if ( 1 <= tx && tx < 32 )
if(tx>=1) Hscan[tx]+=Hscan[tx-1];
Hscan[tx]+=Hscan[tx-1]; __syncthreads();
if(tx>=2) if ( 2 <= tx && tx < 32 )
Hscan[tx]+=Hscan[tx-2]; Hscan[tx]+=Hscan[tx-2];
if(tx>=4) __syncthreads();
Hscan[tx]+=Hscan[tx-4]; if ( 4 <= tx && tx < 32 )
if(tx>=8) Hscan[tx]+=Hscan[tx-4];
Hscan[tx]+=Hscan[tx-8]; __syncthreads();
if(tx>=16) if ( 8 <= tx && tx < 32 )
Hscan[tx]+=Hscan[tx-16]; Hscan[tx]+=Hscan[tx-8];
} __syncthreads();
if ( 16 <= tx && tx < 32 )
Hscan[tx]+=Hscan[tx-16];
__syncthreads(); __syncthreads();
if(tx<31){ if(tx<31){
if(Hscan[tx+1] > medPos && Hscan[tx] < medPos){ if(Hscan[tx+1] > medPos && Hscan[tx] < medPos){
......
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment