Cinder

  • Main Page
  • Related Pages
  • Modules
  • Namespaces
  • Classes
  • Files
  • File List
  • File Members

include/opencv2/gpu/devmem2d.hpp

Go to the documentation of this file.
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__ */