datamov_utils.hpp
Go to the documentation of this file.
1 /*M///////////////////////////////////////////////////////////////////////////////////////
2 //
3 // IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
4 //
5 // By downloading, copying, installing or using the software you agree to this license.
6 // If you do not agree to this license, do not download, install,
7 // copy or use the software.
8 //
9 //
10 // License Agreement
11 // For Open Source Computer Vision Library
12 //
13 // Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
14 // Copyright (C) 2009, Willow Garage Inc., all rights reserved.
15 // Third party copyrights are property of their respective owners.
16 //
17 // Redistribution and use in source and binary forms, with or without modification,
18 // are permitted provided that the following conditions are met:
19 //
20 // * Redistribution's of source code must retain the above copyright notice,
21 // this list of conditions and the following disclaimer.
22 //
23 // * Redistribution's in binary form must reproduce the above copyright notice,
24 // this list of conditions and the following disclaimer in the documentation
25 // and/or other materials provided with the distribution.
26 //
27 // * The name of the copyright holders may not be used to endorse or promote products
28 // derived from this software without specific prior written permission.
29 //
30 // This software is provided by the copyright holders and contributors "as is" and
31 // any express or implied warranties, including, but not limited to, the implied
32 // warranties of merchantability and fitness for a particular purpose are disclaimed.
33 // In no event shall the Intel Corporation or contributors be liable for any direct,
34 // indirect, incidental, special, exemplary, or consequential damages
35 // (including, but not limited to, procurement of substitute goods or services;
36 // loss of use, data, or profits; or business interruption) however caused
37 // and on any theory of liability, whether in contract, strict liability,
38 // or tort (including negligence or otherwise) arising in any way out of
39 // the use of this software, even if advised of the possibility of such damage.
40 //
41 //M*/
42 
43 #ifndef __OPENCV_GPU_DATAMOV_UTILS_HPP__
44 #define __OPENCV_GPU_DATAMOV_UTILS_HPP__
45 
46 #include "common.hpp"
47 
48 namespace cv { namespace gpu { namespace device
49 {
50  #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 200
51 
52  // for Fermi memory space is detected automatically
53  template <typename T> struct ForceGlob
54  {
55  __device__ __forceinline__ static void Load(const T* ptr, int offset, T& val) { val = ptr[offset]; }
56  };
57 
58  #else // __CUDA_ARCH__ >= 200
59 
60  #if defined(_WIN64) || defined(__LP64__)
61  // 64-bit register modifier for inlined asm
62  #define OPENCV_GPU_ASM_PTR "l"
63  #else
64  // 32-bit register modifier for inlined asm
65  #define OPENCV_GPU_ASM_PTR "r"
66  #endif
67 
68  template<class T> struct ForceGlob;
69 
70  #define OPENCV_GPU_DEFINE_FORCE_GLOB(base_type, ptx_type, reg_mod) \
71  template <> struct ForceGlob<base_type> \
72  { \
73  __device__ __forceinline__ static void Load(const base_type* ptr, int offset, base_type& val) \
74  { \
75  asm("ld.global."#ptx_type" %0, [%1];" : "="#reg_mod(val) : OPENCV_GPU_ASM_PTR(ptr + offset)); \
76  } \
77  };
78 
79  #define OPENCV_GPU_DEFINE_FORCE_GLOB_B(base_type, ptx_type) \
80  template <> struct ForceGlob<base_type> \
81  { \
82  __device__ __forceinline__ static void Load(const base_type* ptr, int offset, base_type& val) \
83  { \
84  asm("ld.global."#ptx_type" %0, [%1];" : "=r"(*reinterpret_cast<uint*>(&val)) : OPENCV_GPU_ASM_PTR(ptr + offset)); \
85  } \
86  };
87 
92  OPENCV_GPU_DEFINE_FORCE_GLOB (short, s16, h)
96  OPENCV_GPU_DEFINE_FORCE_GLOB (double, f64, d)
97 
98  #undef OPENCV_GPU_DEFINE_FORCE_GLOB
99  #undef OPENCV_GPU_DEFINE_FORCE_GLOB_B
100  #undef OPENCV_GPU_ASM_PTR
101 
102  #endif // __CUDA_ARCH__ >= 200
103 }}} // namespace cv { namespace gpu { namespace device
104 
105 #endif // __OPENCV_GPU_DATAMOV_UTILS_HPP__
GLdouble GLdouble GLdouble r
signed char schar
Definition: types_c.h:174
CvArr CvPoint offset
Definition: imgproc_c.h:77
s8 h u32
Definition: datamov_utils.hpp:93
s8 u16
Definition: datamov_utils.hpp:91
s8 h OPENCV_GPU_DEFINE_FORCE_GLOB(short, s16, h) OPENCV_GPU_DEFINE_FORCE_GLOB(uint
int d
Definition: legacy.hpp:3064
const CvMat const CvMat const CvMat CvMat CvMat CvMat CvMat CvSize CvMat CvMat * T
Definition: calib3d.hpp:270
GLuint GLfloat * val
GLintptr offset
unsigned int uint
Definition: common.hpp:104
Definition: datamov_utils.hpp:53
unsigned short ushort
Definition: common.hpp:101
int int int int int int h
const char * ptr
Definition: core_c.h:942
OPENCV_GPU_DEFINE_FORCE_GLOB_B(uchar, u8) OPENCV_GPU_DEFINE_FORCE_GLOB_B(schar
unsigned char uchar
Definition: types_c.h:170
s8 h r f32
Definition: datamov_utils.hpp:95
GLclampf f
__device__ static __forceinline__ void Load(const T *ptr, int offset, T &val)
Definition: datamov_utils.hpp:55