fix GPU WARP border mode in CUDA 7.0 and Maxwell architecture

This commit is contained in:
Vladislav Vinogradov 2015-04-07 13:04:31 +03:00
parent d948b2a235
commit 27302c367c

View File

@ -626,12 +626,12 @@ namespace cv { namespace gpu { namespace device
__device__ __forceinline__ int idx_row_low(int y) const __device__ __forceinline__ int idx_row_low(int y) const
{ {
return (y >= 0) * y + (y < 0) * (y - ((y - height + 1) / height) * height); return (y >= 0) ? y : (y - ((y - height + 1) / height) * height);
} }
__device__ __forceinline__ int idx_row_high(int y) const __device__ __forceinline__ int idx_row_high(int y) const
{ {
return (y < height) * y + (y >= height) * (y % height); return (y < height) ? y : (y % height);
} }
__device__ __forceinline__ int idx_row(int y) const __device__ __forceinline__ int idx_row(int y) const
@ -641,12 +641,12 @@ namespace cv { namespace gpu { namespace device
__device__ __forceinline__ int idx_col_low(int x) const __device__ __forceinline__ int idx_col_low(int x) const
{ {
return (x >= 0) * x + (x < 0) * (x - ((x - width + 1) / width) * width); return (x >= 0) ? x : (x - ((x - width + 1) / width) * width);
} }
__device__ __forceinline__ int idx_col_high(int x) const __device__ __forceinline__ int idx_col_high(int x) const
{ {
return (x < width) * x + (x >= width) * (x % width); return (x < width) ? x : (x % width);
} }
__device__ __forceinline__ int idx_col(int x) const __device__ __forceinline__ int idx_col(int x) const