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
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
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