added PtrStep structure to pass in __global__ functions

diff --git a/modules/gpu/include/opencv2/gpu/devmem2d.hpp b/modules/gpu/include/opencv2/gpu/devmem2d.hpp
index 80290ed..f3293c0 100644 (file)
--- a/modules/gpu/include/opencv2/gpu/devmem2d.hpp
+++ b/modules/gpu/include/opencv2/gpu/devmem2d.hpp
@@ -47,30 +47,52 @@ namespace cv
 {    \r
     namespace gpu\r
     {\r
-        // Simple lightweight structure that encapsulates image ptr on device, its pitch and its sizes.\r
+        // Simple lightweight structures that encapsulates information about an image on device.\r
         // It is intended to pass to nvcc-compiled code. GpuMat depends on headers that nvcc can't compile\r
 \r
-        template <typename T>\r
-        struct DevMem2D_\r
+        template<typename T> struct PtrStep_\r
         {\r
-            typedef T elem_t;\r
-            enum { elem_size = sizeof(elem_t) };\r
+            T* ptr;\r
+            size_t step;\r
+\r
+            typedef T elem_type;\r
+            enum { elem_size = sizeof(elem_type) };\r
 \r
+#if defined(__CUDACC__) \r
+            __host__ __device__ \r
+#endif\r
+            size_t elemSize() const { return elem_size; }            \r
+        };\r
+        \r
+        template <typename T> struct DevMem2D_\r
+        {            \r
             int cols;\r
             int rows;\r
             T* ptr;\r
             size_t step;\r
             size_t elem_step;\r
 \r
+            /*__host__*/\r
             DevMem2D_() : cols(0), rows(0), ptr(0), step(0), elem_step(0) {}\r
 \r
+            /*__host__*/\r
             DevMem2D_(int rows_, int cols_, T *ptr_, size_t step_)\r
                 : cols(cols_), rows(rows_), ptr(ptr_), step(step_), elem_step(step_ / sizeof(T)) {}\r
             \r
             template <typename U>\r
+            /*__host__*/\r
             explicit DevMem2D_(const DevMem2D_<U>& d)\r
                 : cols(d.cols), rows(d.rows), ptr((T*)d.ptr), step(d.step), elem_step(d.step / sizeof(T)) {}\r
 \r
+            template <typename U>\r
+            /*__host__*/\r
+            operator PtrStep_<U>() const { PtrStep_<U> dt; dt.ptr = ptr; dt.step = step; return dt; }\r
+\r
+            typedef typename PtrStep_<T>::elem_type elem_type;\r
+            enum { elem_size = PtrStep_<T>::elem_size };\r
+#if defined(__CUDACC__) \r
+            __host__ __device__ \r
+#endif\r
             size_t elemSize() const { return elem_size; }\r
         };\r
 \r