mirror/opencv
1
0
Fork 0
mirror of https://github.com/opencv/opencv.git synced 2024-11-29 05:29:54 +08:00
Code Issues Actions Packages Projects Releases Wiki Activity

added PtrStep structure to pass in __global__ functions

Browse Source
This commit is contained in:
Anatoly Baksheev 2010-10-22 16:57:22 +00:00
parent 17412aa54b
commit 5ac4b23b7c
1 changed files with 27 additions and 5 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

32
modules/gpu/include/opencv2/gpu/devmem2d.hpp
View File

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