datamov_utils.hpp File Reference

#include "common.hpp"

Classes
struct	cv::gpu::device::ForceGlob< T >
struct	cv::gpu::device::ForceGlob< uchar >
struct	cv::gpu::device::ForceGlob< schar >
struct	cv::gpu::device::ForceGlob< char >
struct	cv::gpu::device::ForceGlob< ushort >
struct	cv::gpu::device::ForceGlob< short >
struct	cv::gpu::device::ForceGlob< uint >
struct	cv::gpu::device::ForceGlob< int >
struct	cv::gpu::device::ForceGlob< float >
struct	cv::gpu::device::ForceGlob< double >

Namespaces
	cv
	cv::gpu
	cv::gpu::device

Macros
#define	OPENCV_GPU_ASM_PTR   "r"
#define	OPENCV_GPU_DEFINE_FORCE_GLOB(base_type, ptx_type, reg_mod)
#define	OPENCV_GPU_DEFINE_FORCE_GLOB_B(base_type, ptx_type)

Macro Definition Documentation

OPENCV_GPU_ASM_PTR

#define OPENCV_GPU_ASM_PTR   "r"

OPENCV_GPU_DEFINE_FORCE_GLOB

#define OPENCV_GPU_DEFINE_FORCE_GLOB	(		base_type,
			ptx_type,
			reg_mod
	)

Value:

template <> struct ForceGlob<base_type> \
            { \
                __device__ __forceinline__ static void Load(const base_type* ptr, int offset, base_type& val) \
                { \
                    asm("ld.global."#ptx_type" %0, [%1];" : "="#reg_mod(val) : OPENCV_GPU_ASM_PTR(ptr + offset)); \
                } \
            };

OPENCV_GPU_DEFINE_FORCE_GLOB_B

#define OPENCV_GPU_DEFINE_FORCE_GLOB_B	(		base_type,
			ptx_type
	)

Value:

template <> struct ForceGlob<base_type> \
            { \
                __device__ __forceinline__ static void Load(const base_type* ptr, int offset, base_type& val) \
                { \
                    asm("ld.global."#ptx_type" %0, [%1];" : "=r"(*reinterpret_cast<uint*>(&val)) : OPENCV_GPU_ASM_PTR(ptr + offset)); \
                } \
            };