fixed bug in gpu::remap under win32

diff --git a/modules/gpu/src/cuda/imgproc.cu b/modules/gpu/src/cuda/imgproc.cu
index 5373641..44a54ca 100644 (file)
--- a/modules/gpu/src/cuda/imgproc.cu
+++ b/modules/gpu/src/cuda/imgproc.cu
@@ -66,6 +66,24 @@ namespace cv { namespace gpu { namespace imgproc
             dst.ptr(y)[x] = saturate_cast<T>(src(ycoo, xcoo));\r
         }\r
     }\r
+\r
+    template <template <typename> class Filter, template <typename> class B, typename T> struct RemapDispatcherStream\r
+    {\r
+        static void call(const DevMem2D_<T>& src, const DevMem2Df& mapx, const DevMem2Df& mapy, const DevMem2D_<T>& dst, const float* borderValue, cudaStream_t stream)\r
+        {\r
+            typedef typename TypeVec<float, VecTraits<T>::cn>::vec_type work_type; \r
+            \r
+            dim3 block(32, 8);\r
+            dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y));\r
+\r
+            B<work_type> brd(src.rows, src.cols, VecTraits<work_type>::make(borderValue));\r
+            BorderReader< PtrStep_<T>, B<work_type> > brdSrc(src, brd);\r
+            Filter< BorderReader< PtrStep_<T>, B<work_type> > > filter_src(brdSrc);\r
+\r
+            remap<<<grid, block, 0, stream>>>(filter_src, mapx, mapy, dst);\r
+            cudaSafeCall( cudaGetLastError() );\r
+        }\r
+    };\r
     \r
     template <template <typename> class Filter, template <typename> class B, typename T> struct RemapDispatcherNonStream\r
     {\r
@@ -163,22 +181,7 @@ namespace cv { namespace gpu { namespace imgproc
             if (stream == 0)\r
                 RemapDispatcherNonStream<Filter, B, T>::call(src, mapx, mapy, dst, borderValue);\r
             else\r
-                callStream(src, mapx, mapy, dst, borderValue, stream);\r
-        }\r
-        \r
-        static void callStream(const DevMem2D_<T>& src, const DevMem2Df& mapx, const DevMem2Df& mapy, const DevMem2D_<T>& dst, const float* borderValue, cudaStream_t stream)\r
-        {\r
-            typedef typename TypeVec<float, VecTraits<T>::cn>::vec_type work_type; \r
-            \r
-            dim3 block(32, 8);\r
-            dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y));\r
-\r
-            B<work_type> brd(src.rows, src.cols, VecTraits<work_type>::make(borderValue));\r
-            BorderReader< PtrStep_<T>, B<work_type> > brd_src(src, brd);\r
-            Filter< BorderReader< PtrStep_<T>, B<work_type> > > filter_src(brd_src);\r
-\r
-            remap<<<grid, block, 0, stream>>>(filter_src, mapx, mapy, dst);\r
-            cudaSafeCall( cudaGetLastError() );\r
+                RemapDispatcherStream<Filter, B, T>::call(src, mapx, mapy, dst, borderValue, stream);\r
         }\r
     };\r
 \r

diff --git a/modules/gpu/src/opencv2/gpu/device/border_interpolate.hpp b/modules/gpu/src/opencv2/gpu/device/border_interpolate.hpp
index 50c929a..1f226b9 100644 (file)
--- a/modules/gpu/src/opencv2/gpu/device/border_interpolate.hpp
+++ b/modules/gpu/src/opencv2/gpu/device/border_interpolate.hpp
@@ -758,6 +758,29 @@ namespace cv { namespace gpu { namespace device
         const Ptr2D ptr;\r
         const B b;\r
     };\r
+\r
+    // under win32 there is some bug with templated types that passed as kernel parameters\r
+    // with this specialization all works fine\r
+    template <typename Ptr2D, typename D> struct BorderReader< Ptr2D, BrdConstant<D> >\r
+    {\r
+        typedef typename BrdConstant<D>::result_type elem_type;\r
+        typedef typename Ptr2D::index_type index_type;\r
+\r
+        __host__ __device__ __forceinline__ BorderReader(const Ptr2D& src_, const BrdConstant<D>& b) : \r
+            src(src_), height(b.height), width(b.width), val(b.val) \r
+        {\r
+        }\r
+\r
+        __device__ __forceinline__ D operator ()(index_type y, index_type x) const\r
+        {\r
+            return (x >= 0 && x < width && y >= 0 && y < height) ? saturate_cast<D>(src(y, x)) : val;\r
+        }\r
+\r
+        const Ptr2D src;\r
+        const int height;\r
+        const int width;\r
+        const D val;\r
+    };\r
 }}}\r
 \r
 #endif // __OPENCV_GPU_BORDER_INTERPOLATE_HPP__\r

diff --git a/samples/gpu/performance/tests.cpp b/samples/gpu/performance/tests.cpp
index e03988b..5f70236 100644 (file)
--- a/samples/gpu/performance/tests.cpp
+++ b/samples/gpu/performance/tests.cpp
@@ -80,7 +80,7 @@ TEST(remap)
     gpu::GpuMat d_src, d_dst, d_xmap, d_ymap;\r
     \r
     int interpolation = INTER_LINEAR;\r
-    int borderMode = BORDER_CONSTANT;\r
+    int borderMode = BORDER_REPLICATE;\r
 \r
     for (int size = 1000; size <= 4000; size *= 2)\r
     {\r