00001 /*M/////////////////////////////////////////////////////////////////////////////////////// 00002 // 00003 // IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. 00004 // 00005 // By downloading, copying, installing or using the software you agree to this license. 00006 // If you do not agree to this license, do not download, install, 00007 // copy or use the software. 00008 // 00009 // 00010 // License Agreement 00011 // For Open Source Computer Vision Library 00012 // 00013 // Copyright (C) 2000-2008, Intel Corporation, all rights reserved. 00014 // Copyright (C) 2009, Willow Garage Inc., all rights reserved. 00015 // Third party copyrights are property of their respective owners. 00016 // 00017 // Redistribution and use in source and binary forms, with or without modification, 00018 // are permitted provided that the following conditions are met: 00019 // 00020 // * Redistribution's of source code must retain the above copyright notice, 00021 // this list of conditions and the following disclaimer. 00022 // 00023 // * Redistribution's in binary form must reproduce the above copyright notice, 00024 // this list of conditions and the following disclaimer in the documentation 00025 // and/or other GpuMaterials provided with the distribution. 00026 // 00027 // * The name of the copyright holders may not be used to endorse or promote products 00028 // derived from this software without specific prior written permission. 00029 // 00030 // This software is provided by the copyright holders and contributors "as is" and 00031 // any express or implied warranties, including, but not limited to, the implied 00032 // warranties of merchantability and fitness for a particular purpose are disclaimed. 00033 // In no event shall the Intel Corporation or contributors be liable for any direct, 00034 // indirect, incidental, special, exemplary, or consequential damages 00035 // (including, but not limited to, procurement of substitute goods or services; 00036 // loss of use, data, or profits; or business interruption) however caused 00037 // and on any theory of liability, whether in contract, strict liability, 00038 // or tort (including negligence or otherwise) arising in any way out of 00039 // the use of this software, even if advised of the possibility of such damage. 00040 // 00041 //M*/ 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 // Simple lightweight structures that encapsulates information about an image on device. 00056 // It is intended to pass to nvcc-compiled code. GpuMat depends on headers that nvcc can't compile 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 /* __OPENCV_GPU_DEVMEM2D_HPP__ */