optimized gpu::remap (use texture memory)

diff --git a/modules/gpu/src/cuda/imgproc.cu b/modules/gpu/src/cuda/imgproc.cu
index 032cffb..16fbb2b 100644 (file)
--- a/modules/gpu/src/cuda/imgproc.cu
+++ b/modules/gpu/src/cuda/imgproc.cu
@@ -77,8 +77,8 @@ namespace cv { namespace gpu { namespace imgproc
             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
+            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>>>(filter_src, mapx, mapy, dst);\r
             cudaSafeCall( cudaGetLastError() );\r
@@ -98,6 +98,23 @@ namespace cv { namespace gpu { namespace imgproc
             return tex2D(tex_remap_ ## type , x, y); \\r
         } \\r
     }; \\r
+    template <template <typename> class Filter, template <typename> class B> struct RemapDispatcherNonStream<Filter, B, type> \\r
+    { \\r
+        static void call(const DevMem2D_< type >& src, const DevMem2Df& mapx, const DevMem2Df& mapy, const DevMem2D_< type >& dst, const float* borderValue) \\r
+        { \\r
+            typedef typename TypeVec<float, VecTraits< type >::cn>::vec_type work_type; \\r
+            dim3 block(32, 8); \\r
+            dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y)); \\r
+            TextureBinder texHandler(&tex_remap_ ## type , src); \\r
+            tex_remap_ ## type ##_reader texSrc; \\r
+            B<work_type> brd(src.rows, src.cols, VecTraits<work_type>::make(borderValue)); \\r
+            BorderReader< tex_remap_ ## type ##_reader, B<work_type> > brdSrc(texSrc, brd); \\r
+            Filter< BorderReader< tex_remap_ ## type ##_reader, B<work_type> > > filter_src(brdSrc); \\r
+            remap<<<grid, block>>>(filter_src, mapx, mapy, dst); \\r
+            cudaSafeCall( cudaGetLastError() ); \\r
+            cudaSafeCall( cudaDeviceSynchronize() ); \\r
+        } \\r
+    }; \\r
     template <template <typename> class Filter> struct RemapDispatcherNonStream<Filter, BrdReplicate, type> \\r
     { \\r
         static void call(const DevMem2D_< type >& src, const DevMem2Df& mapx, const DevMem2Df& mapy, const DevMem2D_< type >& dst, const float*) \\r
@@ -106,7 +123,7 @@ namespace cv { namespace gpu { namespace imgproc
             dim3 grid(divUp(dst.cols, block.x), divUp(dst.rows, block.y)); \\r
             TextureBinder texHandler(&tex_remap_ ## type , src); \\r
             tex_remap_ ## type ##_reader texSrc; \\r
-            Filter<tex_remap_ ## type ##_reader> filter_src(texSrc); \\r
+            Filter< tex_remap_ ## type ##_reader > filter_src(texSrc); \\r
             remap<<<grid, block>>>(filter_src, mapx, mapy, dst); \\r
             cudaSafeCall( cudaGetLastError() ); \\r
             cudaSafeCall( cudaDeviceSynchronize() ); \\r