datamov_utils.hpp

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 materials 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_CUDA_DATAMOV_UTILS_HPP__
00044 #define __OPENCV_CUDA_DATAMOV_UTILS_HPP__
00045 
00046 #include "common.hpp "
00047 
00048 /** @file
00049  * @deprecated Use @ref cudev instead.
00050  */
00051 
00052 //! @cond IGNORED
00053 
00054 namespace cv { namespace cuda { namespace device
00055 {
00056     #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 200
00057 
00058         // for Fermi memory space is detected automatically
00059         template <typename T> struct ForceGlob
00060         {
00061             __device__ __forceinline__ static void Load(const T* ptr, int offset, T& val)  { val = ptr[offset];  }
00062         };
00063 
00064     #else // __CUDA_ARCH__ >= 200
00065 
00066         #if defined(_WIN64) || defined(__LP64__)
00067             // 64-bit register modifier for inlined asm
00068             #define OPENCV_CUDA_ASM_PTR "l"
00069         #else
00070             // 32-bit register modifier for inlined asm
00071             #define OPENCV_CUDA_ASM_PTR "r"
00072         #endif
00073 
00074         template<class T> struct ForceGlob;
00075 
00076         #define OPENCV_CUDA_DEFINE_FORCE_GLOB(base_type, ptx_type, reg_mod) \
00077             template <> struct ForceGlob<base_type> \
00078             { \
00079                 __device__ __forceinline__ static void Load(const base_type* ptr, int offset, base_type& val) \
00080                 { \
00081                     asm("ld.global."#ptx_type" %0, [%1];" : "="#reg_mod(val) : OPENCV_CUDA_ASM_PTR(ptr + offset)); \
00082                 } \
00083             };
00084 
00085         #define OPENCV_CUDA_DEFINE_FORCE_GLOB_B(base_type, ptx_type) \
00086             template <> struct ForceGlob<base_type> \
00087             { \
00088                 __device__ __forceinline__ static void Load(const base_type* ptr, int offset, base_type& val) \
00089                 { \
00090                     asm("ld.global."#ptx_type" %0, [%1];" : "=r"(*reinterpret_cast<uint*>(&val)) : OPENCV_CUDA_ASM_PTR(ptr + offset)); \
00091                 } \
00092             };
00093 
00094             OPENCV_CUDA_DEFINE_FORCE_GLOB_B(uchar,  u8)
00095             OPENCV_CUDA_DEFINE_FORCE_GLOB_B(schar,  s8)
00096             OPENCV_CUDA_DEFINE_FORCE_GLOB_B(char,   b8)
00097             OPENCV_CUDA_DEFINE_FORCE_GLOB  (ushort, u16, h)
00098             OPENCV_CUDA_DEFINE_FORCE_GLOB  (short,  s16, h)
00099             OPENCV_CUDA_DEFINE_FORCE_GLOB  (uint,   u32, r)
00100             OPENCV_CUDA_DEFINE_FORCE_GLOB  (int,    s32, r)
00101             OPENCV_CUDA_DEFINE_FORCE_GLOB  (float,  f32, f)
00102             OPENCV_CUDA_DEFINE_FORCE_GLOB  (double, f64, d)
00103 
00104         #undef OPENCV_CUDA_DEFINE_FORCE_GLOB
00105         #undef OPENCV_CUDA_DEFINE_FORCE_GLOB_B
00106         #undef OPENCV_CUDA_ASM_PTR
00107 
00108     #endif // __CUDA_ARCH__ >= 200
00109 }}} // namespace cv { namespace cuda { namespace cudev
00110 
00111 //! @endcond
00112 
00113 #endif // __OPENCV_CUDA_DATAMOV_UTILS_HPP__
00114