43 #ifndef __OPENCV_GPU_SATURATE_CAST_HPP__
44 #define __OPENCV_GPU_SATURATE_CAST_HPP__
48 namespace cv {
namespace gpu {
namespace device
53 template<
typename _Tp> __device__ __forceinline__ _Tp
saturate_cast(
short v) {
return _Tp(v); }
55 template<
typename _Tp> __device__ __forceinline__ _Tp
saturate_cast(
int v) {
return _Tp(v); }
56 template<
typename _Tp> __device__ __forceinline__ _Tp
saturate_cast(
float v) {
return _Tp(v); }
57 template<
typename _Tp> __device__ __forceinline__ _Tp
saturate_cast(
double v) {
return _Tp(v); }
63 asm(
"cvt.sat.u8.s8 %0, %1;" :
"=r"(res) :
"r"(vi));
69 asm(
"cvt.sat.u8.s16 %0, %1;" :
"=r"(res) :
"h"(
v));
75 asm(
"cvt.sat.u8.u16 %0, %1;" :
"=r"(res) :
"h"(
v));
81 asm(
"cvt.sat.u8.s32 %0, %1;" :
"=r"(res) :
"r"(
v));
87 asm(
"cvt.sat.u8.u32 %0, %1;" :
"=r"(res) :
"r"(
v));
93 asm(
"cvt.rni.sat.u8.f32 %0, %1;" :
"=r"(res) :
"f"(
v));
98 #if __CUDA_ARCH__ >= 130
100 asm(
"cvt.rni.sat.u8.f64 %0, %1;" :
"=r"(res) :
"d"(
v));
111 asm(
"cvt.sat.s8.u8 %0, %1;" :
"=r"(res) :
"r"(vi));
117 asm(
"cvt.sat.s8.s16 %0, %1;" :
"=r"(res) :
"h"(
v));
123 asm(
"cvt.sat.s8.u16 %0, %1;" :
"=r"(res) :
"h"(
v));
129 asm(
"cvt.sat.s8.s32 %0, %1;" :
"=r"(res) :
"r"(
v));
135 asm(
"cvt.sat.s8.u32 %0, %1;" :
"=r"(res) :
"r"(
v));
141 asm(
"cvt.rni.sat.s8.f32 %0, %1;" :
"=r"(res) :
"f"(
v));
146 #if __CUDA_ARCH__ >= 130
148 asm(
"cvt.rni.sat.s8.f64 %0, %1;" :
"=r"(res) :
"d"(
v));
159 asm(
"cvt.sat.u16.s8 %0, %1;" :
"=h"(res) :
"r"(vi));
165 asm(
"cvt.sat.u16.s16 %0, %1;" :
"=h"(res) :
"h"(
v));
171 asm(
"cvt.sat.u16.s32 %0, %1;" :
"=h"(res) :
"r"(
v));
177 asm(
"cvt.sat.u16.u32 %0, %1;" :
"=h"(res) :
"r"(
v));
183 asm(
"cvt.rni.sat.u16.f32 %0, %1;" :
"=h"(res) :
"f"(
v));
188 #if __CUDA_ARCH__ >= 130
190 asm(
"cvt.rni.sat.u16.f64 %0, %1;" :
"=h"(res) :
"d"(
v));
200 asm(
"cvt.sat.s16.u16 %0, %1;" :
"=h"(res) :
"h"(
v));
206 asm(
"cvt.sat.s16.s32 %0, %1;" :
"=h"(res) :
"r"(
v));
212 asm(
"cvt.sat.s16.u32 %0, %1;" :
"=h"(res) :
"r"(
v));
218 asm(
"cvt.rni.sat.s16.f32 %0, %1;" :
"=h"(res) :
"f"(
v));
223 #if __CUDA_ARCH__ >= 130
225 asm(
"cvt.rni.sat.s16.f64 %0, %1;" :
"=h"(res) :
"d"(
v));
235 asm(
"cvt.sat.s32.u32 %0, %1;" :
"=r"(res) :
"r"(
v));
240 return __float2int_rn(
v);
244 #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130
245 return __double2int_rn(
v);
255 asm(
"cvt.sat.u32.s8 %0, %1;" :
"=r"(res) :
"r"(vi));
261 asm(
"cvt.sat.u32.s16 %0, %1;" :
"=r"(res) :
"h"(
v));
267 asm(
"cvt.sat.u32.s32 %0, %1;" :
"=r"(res) :
"r"(
v));
272 return __float2uint_rn(
v);
276 #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130
277 return __double2uint_rn(
v);
signed char schar
Definition: common.hpp:102
unsigned char uchar
Definition: common.hpp:100
unsigned int uint
Definition: common.hpp:104
unsigned short ushort
Definition: common.hpp:101
__device__ __forceinline__ _Tp saturate_cast(uchar v)
Definition: saturate_cast.hpp:50
::max::max::max float
Definition: functional.hpp:326