Go to the documentation of this file.00001
00002
00003
00004
00005
00006
00007
00008
00009
00010
00011
00012
00013
00014
00015
00016
00017
00018
00019
00020
00021
00022
00023
00024
00025
00026
00027
00028
00029
00030
00031
00032
00033
00034
00035
00036
00037
00038
00039
00040
00041
00042
00043 #ifndef __OPENCV_GPU_DEVMEM2D_HPP__
00044 #define __OPENCV_GPU_DEVMEM2D_HPP__
00045
00046 #if defined(__DEVCLASES_ADD_THRUST_BEGIN_END__)
00047 #include "thrust/device_ptr.h"
00048 #endif
00049
00050
00051 namespace cv
00052 {
00053 namespace gpu
00054 {
00055
00056
00057
00058 #if defined(__CUDACC__)
00059 #define __CV_GPU_HOST_DEVICE__ __host__ __device__ __forceinline__
00060 #else
00061 #define __CV_GPU_HOST_DEVICE__
00062 #endif
00063
00064 template <bool expr> struct StaticAssert;
00065 template <> struct StaticAssert<true> {static __CV_GPU_HOST_DEVICE__ void check(){}};
00066
00067 template <typename T> struct DevMem2D_
00068 {
00069 int cols;
00070 int rows;
00071 T* data;
00072 size_t step;
00073
00074 DevMem2D_() : cols(0), rows(0), data(0), step(0) {}
00075
00076 DevMem2D_(int rows_, int cols_, T *data_, size_t step_)
00077 : cols(cols_), rows(rows_), data(data_), step(step_) {}
00078
00079 template <typename U>
00080 explicit DevMem2D_(const DevMem2D_<U>& d)
00081 : cols(d.cols), rows(d.rows), data((T*)d.data), step(d.step) {}
00082
00083 typedef T elem_type;
00084 enum { elem_size = sizeof(elem_type) };
00085
00086 __CV_GPU_HOST_DEVICE__ size_t elemSize() const { return elem_size; }
00087 __CV_GPU_HOST_DEVICE__ T* ptr(int y = 0) { return (T*)( (char*)data + y * step ); }
00088 __CV_GPU_HOST_DEVICE__ const T* ptr(int y = 0) const { return (const T*)( (const char*)data + y * step ); }
00089
00090 __CV_GPU_HOST_DEVICE__ operator T*() const { return data; }
00091
00092 #if defined(__DEVCLASES_ADD_THRUST_BEGIN_END__)
00093 thrust::device_ptr<T> begin() const { return thrust::device_ptr<T>(data); }
00094 thrust::device_ptr<T> end() const { return thrust::device_ptr<T>(data) + cols * rows; }
00095 #endif
00096 };
00097
00098 template<typename T> struct PtrStep_
00099 {
00100 T* data;
00101 size_t step;
00102
00103 PtrStep_() : data(0), step(0) {}
00104 PtrStep_(const DevMem2D_<T>& mem) : data(mem.data), step(mem.step) {}
00105
00106 typedef T elem_type;
00107 enum { elem_size = sizeof(elem_type) };
00108
00109 __CV_GPU_HOST_DEVICE__ size_t elemSize() const { return elem_size; }
00110 __CV_GPU_HOST_DEVICE__ T* ptr(int y = 0) { return (T*)( (char*)data + y * step); }
00111 __CV_GPU_HOST_DEVICE__ const T* ptr(int y = 0) const { return (const T*)( (const char*)data + y * step); }
00112
00113 #if defined(__DEVCLASES_ADD_THRUST_BEGIN_END__)
00114 thrust::device_ptr<T> begin() const { return thrust::device_ptr<T>(data); }
00115 #endif
00116 };
00117
00118 template<typename T> struct PtrElemStep_ : public PtrStep_<T>
00119 {
00120 PtrElemStep_(const DevMem2D_<T>& mem) : PtrStep_<T>(mem)
00121 {
00122 StaticAssert<256 % sizeof(T) == 0>::check();
00123
00124 PtrStep_<T>::step /= PtrStep_<T>::elem_size;
00125 }
00126 __CV_GPU_HOST_DEVICE__ T* ptr(int y = 0) { return PtrStep_<T>::data + y * PtrStep_<T>::step; }
00127 __CV_GPU_HOST_DEVICE__ const T* ptr(int y = 0) const { return PtrStep_<T>::data + y * PtrStep_<T>::step; }
00128 };
00129
00130 typedef DevMem2D_<unsigned char> DevMem2D;
00131 typedef DevMem2D_<float> DevMem2Df;
00132 typedef DevMem2D_<int> DevMem2Di;
00133
00134 typedef PtrStep_<unsigned char> PtrStep;
00135 typedef PtrStep_<float> PtrStepf;
00136 typedef PtrStep_<int> PtrStepi;
00137
00138 typedef PtrElemStep_<unsigned char> PtrElemStep;
00139 typedef PtrElemStep_<float> PtrElemStepf;
00140 typedef PtrElemStep_<int> PtrElemStepi;
00141
00142 #undef __CV_GPU_HOST_DEVICE__
00143 }
00144 }
00145
00146 #endif