Skip to content

O
opencv
Commit 50d012b5 authored by Vladislav Vinogradov's avatar Vladislav Vinogradov
Browse files
Options

fixed gpu::pyrUp implementation

parent 5f9e47a9
Hide whitespace changes
Inline Side-by-side
Showing with 21 additions and 21 deletions
modules/gpu/src/cuda/imgproc.cu
View file @ 50d012b5
...@@ -1127,11 +1127,11 @@ namespace cv { namespace gpu { namespace imgproc ...@@ -1127,11 +1127,11 @@ namespace cv { namespace gpu { namespace imgproc
sum = VecTraits<value_type>::all(0); sum = VecTraits<value_type>::all(0);
sum = sum + 0.0625f * smem1[1 + threadIdx.y / 2][1 + ((tidx - 2) >> 1)]; sum = sum + (tidx % 2 == 0) * 0.0625f * smem1[1 + threadIdx.y / 2][1 + ((tidx - 2) >> 1)];
sum = sum + 0.25f * smem1[1 + threadIdx.y / 2][1 + ((tidx - 1) >> 1)]; sum = sum + (tidx % 2 != 0) * 0.25f * smem1[1 + threadIdx.y / 2][1 + ((tidx - 1) >> 1)];
sum = sum + 0.375f * smem1[1 + threadIdx.y / 2][1 + ((tidx ) >> 1)]; sum = sum + (tidx % 2 == 0) * 0.375f * smem1[1 + threadIdx.y / 2][1 + ((tidx ) >> 1)];
sum = sum + 0.25f * smem1[1 + threadIdx.y / 2][1 + ((tidx + 1) >> 1)]; sum = sum + (tidx % 2 != 0) * 0.25f * smem1[1 + threadIdx.y / 2][1 + ((tidx + 1) >> 1)];
sum = sum + 0.0625f * smem1[1 + threadIdx.y / 2][1 + ((tidx + 2) >> 1)]; sum = sum + (tidx % 2 == 0) * 0.0625f * smem1[1 + threadIdx.y / 2][1 + ((tidx + 2) >> 1)];
smem2[2 + threadIdx.y][tidx] = sum; smem2[2 + threadIdx.y][tidx] = sum;
...@@ -1139,11 +1139,11 @@ namespace cv { namespace gpu { namespace imgproc ...@@ -1139,11 +1139,11 @@ namespace cv { namespace gpu { namespace imgproc
{ {
sum = VecTraits<value_type>::all(0); sum = VecTraits<value_type>::all(0);
sum = sum + 0.0625f * smem1[0][1 + ((tidx - 2) >> 1)]; sum = sum + (tidx % 2 == 0) * 0.0625f * smem1[0][1 + ((tidx - 2) >> 1)];
sum = sum + 0.25f * smem1[0][1 + ((tidx - 1) >> 1)]; sum = sum + (tidx % 2 != 0) * 0.25f * smem1[0][1 + ((tidx - 1) >> 1)];
sum = sum + 0.375f * smem1[0][1 + ((tidx ) >> 1)]; sum = sum + (tidx % 2 == 0) * 0.375f * smem1[0][1 + ((tidx ) >> 1)];
sum = sum + 0.25f * smem1[0][1 + ((tidx + 1) >> 1)]; sum = sum + (tidx % 2 != 0) * 0.25f * smem1[0][1 + ((tidx + 1) >> 1)];
sum = sum + 0.0625f * smem1[0][1 + ((tidx + 2) >> 1)]; sum = sum + (tidx % 2 == 0) * 0.0625f * smem1[0][1 + ((tidx + 2) >> 1)];
smem2[threadIdx.y][tidx] = sum; smem2[threadIdx.y][tidx] = sum;
} }
...@@ -1152,11 +1152,11 @@ namespace cv { namespace gpu { namespace imgproc ...@@ -1152,11 +1152,11 @@ namespace cv { namespace gpu { namespace imgproc
{ {
sum = VecTraits<value_type>::all(0); sum = VecTraits<value_type>::all(0);
sum = sum + 0.0625f * smem1[9][1 + ((tidx - 2) >> 1)]; sum = sum + (tidx % 2 == 0) * 0.0625f * smem1[9][1 + ((tidx - 2) >> 1)];
sum = sum + 0.25f * smem1[9][1 + ((tidx - 1) >> 1)]; sum = sum + (tidx % 2 != 0) * 0.25f * smem1[9][1 + ((tidx - 1) >> 1)];
sum = sum + 0.375f * smem1[9][1 + ((tidx ) >> 1)]; sum = sum + (tidx % 2 == 0) * 0.375f * smem1[9][1 + ((tidx ) >> 1)];
sum = sum + 0.25f * smem1[9][1 + ((tidx + 1) >> 1)]; sum = sum + (tidx % 2 != 0) * 0.25f * smem1[9][1 + ((tidx + 1) >> 1)];
sum = sum + 0.0625f * smem1[9][1 + ((tidx + 2) >> 1)]; sum = sum + (tidx % 2 == 0) * 0.0625f * smem1[9][1 + ((tidx + 2) >> 1)];
smem2[4 + threadIdx.y][tidx] = sum; smem2[4 + threadIdx.y][tidx] = sum;
} }
...@@ -1165,14 +1165,14 @@ namespace cv { namespace gpu { namespace imgproc ...@@ -1165,14 +1165,14 @@ namespace cv { namespace gpu { namespace imgproc
sum = VecTraits<value_type>::all(0); sum = VecTraits<value_type>::all(0);
sum = sum + 0.0625f * smem2[2 + threadIdx.y - 2][tidx]; sum = sum + (tidx % 2 == 0) * 0.0625f * smem2[2 + threadIdx.y - 2][tidx];
sum = sum + 0.25f * smem2[2 + threadIdx.y - 1][tidx]; sum = sum + (tidx % 2 != 0) * 0.25f * smem2[2 + threadIdx.y - 1][tidx];
sum = sum + 0.375f * smem2[2 + threadIdx.y ][tidx]; sum = sum + (tidx % 2 == 0) * 0.375f * smem2[2 + threadIdx.y ][tidx];
sum = sum + 0.25f * smem2[2 + threadIdx.y + 1][tidx]; sum = sum + (tidx % 2 != 0) * 0.25f * smem2[2 + threadIdx.y + 1][tidx];
sum = sum + 0.0625f * smem2[2 + threadIdx.y + 2][tidx]; sum = sum + (tidx % 2 == 0) * 0.0625f * smem2[2 + threadIdx.y + 2][tidx];
if (x < dst.cols && y < dst.rows) if (x < dst.cols && y < dst.rows)
dst.ptr(y)[x] = saturate_cast<T>(sum); dst.ptr(y)[x] = saturate_cast<T>(4.0f * sum);
} }
template <typename T, int cn> void pyrUp_gpu(const DevMem2D& src, const DevMem2D& dst, cudaStream_t stream) template <typename T, int cn> void pyrUp_gpu(const DevMem2D& src, const DevMem2D& dst, cudaStream_t stream)
......
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