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 // Copyright (C) 2013, OpenCV Foundation, all rights reserved.
16 // Third party copyrights are property of their respective owners.
17 //
18 // Redistribution and use in source and binary forms, with or without modification,
19 // are permitted provided that the following conditions are met:
20 //
21 // * Redistribution's of source code must retain the above copyright notice,
22 // this list of conditions and the following disclaimer.
23 //
24 // * Redistribution's in binary form must reproduce the above copyright notice,
25 // this list of conditions and the following disclaimer in the documentation
26 // and/or other materials provided with the distribution.
27 //
28 // * The name of the copyright holders may not be used to endorse or promote products
29 // derived from this software without specific prior written permission.
30 //
31 // This software is provided by the copyright holders and contributors "as is" and
32 // any express or implied warranties, including, but not limited to, the implied
33 // warranties of merchantability and fitness for a particular purpose are disclaimed.
34 // In no event shall the Intel Corporation or contributors be liable for any direct,
35 // indirect, incidental, special, exemplary, or consequential damages
36 // (including, but not limited to, procurement of substitute goods or services;
37 // loss of use, data, or profits; or business interruption) however caused
38 // and on any theory of liability, whether in contract, strict liability,
39 // or tort (including negligence or otherwise) arising in any way out of
40 // the use of this software, even if advised of the possibility of such damage.
41 //
42 //M*/
43
44 #pragma once
45
46 #ifndef __OPENCV_CUDEV_UTIL_SATURATE_CAST_HPP__
47 #define __OPENCV_CUDEV_UTIL_SATURATE_CAST_HPP__
48
49 #include "../common.hpp"
50
51 namespace cv { namespace cudev {
52
53 //! @addtogroup cudev
54 //! @{
55
saturate_cast(uchar v)56 template <typename T> __device__ __forceinline__ T saturate_cast(uchar v) { return T(v); }
saturate_cast(schar v)57 template <typename T> __device__ __forceinline__ T saturate_cast(schar v) { return T(v); }
saturate_cast(ushort v)58 template <typename T> __device__ __forceinline__ T saturate_cast(ushort v) { return T(v); }
saturate_cast(short v)59 template <typename T> __device__ __forceinline__ T saturate_cast(short v) { return T(v); }
saturate_cast(uint v)60 template <typename T> __device__ __forceinline__ T saturate_cast(uint v) { return T(v); }
saturate_cast(int v)61 template <typename T> __device__ __forceinline__ T saturate_cast(int v) { return T(v); }
saturate_cast(float v)62 template <typename T> __device__ __forceinline__ T saturate_cast(float v) { return T(v); }
saturate_cast(double v)63 template <typename T> __device__ __forceinline__ T saturate_cast(double v) { return T(v); }
64
saturate_cast(schar v)65 template <> __device__ __forceinline__ uchar saturate_cast<uchar>(schar v)
66 {
67 uint res = 0;
68 int vi = v;
69 asm("cvt.sat.u8.s8 %0, %1;" : "=r"(res) : "r"(vi));
70 return res;
71 }
saturate_cast(short v)72 template <> __device__ __forceinline__ uchar saturate_cast<uchar>(short v)
73 {
74 uint res = 0;
75 asm("cvt.sat.u8.s16 %0, %1;" : "=r"(res) : "h"(v));
76 return res;
77 }
saturate_cast(ushort v)78 template <> __device__ __forceinline__ uchar saturate_cast<uchar>(ushort v)
79 {
80 uint res = 0;
81 asm("cvt.sat.u8.u16 %0, %1;" : "=r"(res) : "h"(v));
82 return res;
83 }
saturate_cast(int v)84 template <> __device__ __forceinline__ uchar saturate_cast<uchar>(int v)
85 {
86 uint res = 0;
87 asm("cvt.sat.u8.s32 %0, %1;" : "=r"(res) : "r"(v));
88 return res;
89 }
saturate_cast(uint v)90 template <> __device__ __forceinline__ uchar saturate_cast<uchar>(uint v)
91 {
92 uint res = 0;
93 asm("cvt.sat.u8.u32 %0, %1;" : "=r"(res) : "r"(v));
94 return res;
95 }
saturate_cast(float v)96 template <> __device__ __forceinline__ uchar saturate_cast<uchar>(float v)
97 {
98 uint res = 0;
99 asm("cvt.rni.sat.u8.f32 %0, %1;" : "=r"(res) : "f"(v));
100 return res;
101 }
saturate_cast(double v)102 template <> __device__ __forceinline__ uchar saturate_cast<uchar>(double v)
103 {
104 uint res = 0;
105 asm("cvt.rni.sat.u8.f64 %0, %1;" : "=r"(res) : "d"(v));
106 return res;
107 }
108
saturate_cast(uchar v)109 template <> __device__ __forceinline__ schar saturate_cast<schar>(uchar v)
110 {
111 uint res = 0;
112 uint vi = v;
113 asm("cvt.sat.s8.u8 %0, %1;" : "=r"(res) : "r"(vi));
114 return res;
115 }
saturate_cast(short v)116 template <> __device__ __forceinline__ schar saturate_cast<schar>(short v)
117 {
118 uint res = 0;
119 asm("cvt.sat.s8.s16 %0, %1;" : "=r"(res) : "h"(v));
120 return res;
121 }
saturate_cast(ushort v)122 template <> __device__ __forceinline__ schar saturate_cast<schar>(ushort v)
123 {
124 uint res = 0;
125 asm("cvt.sat.s8.u16 %0, %1;" : "=r"(res) : "h"(v));
126 return res;
127 }
saturate_cast(int v)128 template <> __device__ __forceinline__ schar saturate_cast<schar>(int v)
129 {
130 uint res = 0;
131 asm("cvt.sat.s8.s32 %0, %1;" : "=r"(res) : "r"(v));
132 return res;
133 }
saturate_cast(uint v)134 template <> __device__ __forceinline__ schar saturate_cast<schar>(uint v)
135 {
136 uint res = 0;
137 asm("cvt.sat.s8.u32 %0, %1;" : "=r"(res) : "r"(v));
138 return res;
139 }
saturate_cast(float v)140 template <> __device__ __forceinline__ schar saturate_cast<schar>(float v)
141 {
142 uint res = 0;
143 asm("cvt.rni.sat.s8.f32 %0, %1;" : "=r"(res) : "f"(v));
144 return res;
145 }
saturate_cast(double v)146 template <> __device__ __forceinline__ schar saturate_cast<schar>(double v)
147 {
148 uint res = 0;
149 asm("cvt.rni.sat.s8.f64 %0, %1;" : "=r"(res) : "d"(v));
150 return res;
151 }
152
saturate_cast(schar v)153 template <> __device__ __forceinline__ ushort saturate_cast<ushort>(schar v)
154 {
155 ushort res = 0;
156 int vi = v;
157 asm("cvt.sat.u16.s8 %0, %1;" : "=h"(res) : "r"(vi));
158 return res;
159 }
saturate_cast(short v)160 template <> __device__ __forceinline__ ushort saturate_cast<ushort>(short v)
161 {
162 ushort res = 0;
163 asm("cvt.sat.u16.s16 %0, %1;" : "=h"(res) : "h"(v));
164 return res;
165 }
saturate_cast(int v)166 template <> __device__ __forceinline__ ushort saturate_cast<ushort>(int v)
167 {
168 ushort res = 0;
169 asm("cvt.sat.u16.s32 %0, %1;" : "=h"(res) : "r"(v));
170 return res;
171 }
saturate_cast(uint v)172 template <> __device__ __forceinline__ ushort saturate_cast<ushort>(uint v)
173 {
174 ushort res = 0;
175 asm("cvt.sat.u16.u32 %0, %1;" : "=h"(res) : "r"(v));
176 return res;
177 }
saturate_cast(float v)178 template <> __device__ __forceinline__ ushort saturate_cast<ushort>(float v)
179 {
180 ushort res = 0;
181 asm("cvt.rni.sat.u16.f32 %0, %1;" : "=h"(res) : "f"(v));
182 return res;
183 }
saturate_cast(double v)184 template <> __device__ __forceinline__ ushort saturate_cast<ushort>(double v)
185 {
186 ushort res = 0;
187 asm("cvt.rni.sat.u16.f64 %0, %1;" : "=h"(res) : "d"(v));
188 return res;
189 }
190
saturate_cast(ushort v)191 template <> __device__ __forceinline__ short saturate_cast<short>(ushort v)
192 {
193 short res = 0;
194 asm("cvt.sat.s16.u16 %0, %1;" : "=h"(res) : "h"(v));
195 return res;
196 }
saturate_cast(int v)197 template <> __device__ __forceinline__ short saturate_cast<short>(int v)
198 {
199 short res = 0;
200 asm("cvt.sat.s16.s32 %0, %1;" : "=h"(res) : "r"(v));
201 return res;
202 }
saturate_cast(uint v)203 template <> __device__ __forceinline__ short saturate_cast<short>(uint v)
204 {
205 short res = 0;
206 asm("cvt.sat.s16.u32 %0, %1;" : "=h"(res) : "r"(v));
207 return res;
208 }
saturate_cast(float v)209 template <> __device__ __forceinline__ short saturate_cast<short>(float v)
210 {
211 short res = 0;
212 asm("cvt.rni.sat.s16.f32 %0, %1;" : "=h"(res) : "f"(v));
213 return res;
214 }
saturate_cast(double v)215 template <> __device__ __forceinline__ short saturate_cast<short>(double v)
216 {
217 short res = 0;
218 asm("cvt.rni.sat.s16.f64 %0, %1;" : "=h"(res) : "d"(v));
219 return res;
220 }
221
saturate_cast(uint v)222 template <> __device__ __forceinline__ int saturate_cast<int>(uint v)
223 {
224 int res = 0;
225 asm("cvt.sat.s32.u32 %0, %1;" : "=r"(res) : "r"(v));
226 return res;
227 }
saturate_cast(float v)228 template <> __device__ __forceinline__ int saturate_cast<int>(float v)
229 {
230 return __float2int_rn(v);
231 }
saturate_cast(double v)232 template <> __device__ __forceinline__ int saturate_cast<int>(double v)
233 {
234 #if CV_CUDEV_ARCH >= 130
235 return __double2int_rn(v);
236 #else
237 return saturate_cast<int>((float) v);
238 #endif
239 }
240
saturate_cast(schar v)241 template <> __device__ __forceinline__ uint saturate_cast<uint>(schar v)
242 {
243 uint res = 0;
244 int vi = v;
245 asm("cvt.sat.u32.s8 %0, %1;" : "=r"(res) : "r"(vi));
246 return res;
247 }
saturate_cast(short v)248 template <> __device__ __forceinline__ uint saturate_cast<uint>(short v)
249 {
250 uint res = 0;
251 asm("cvt.sat.u32.s16 %0, %1;" : "=r"(res) : "h"(v));
252 return res;
253 }
saturate_cast(int v)254 template <> __device__ __forceinline__ uint saturate_cast<uint>(int v)
255 {
256 uint res = 0;
257 asm("cvt.sat.u32.s32 %0, %1;" : "=r"(res) : "r"(v));
258 return res;
259 }
saturate_cast(float v)260 template <> __device__ __forceinline__ uint saturate_cast<uint>(float v)
261 {
262 return __float2uint_rn(v);
263 }
saturate_cast(double v)264 template <> __device__ __forceinline__ uint saturate_cast<uint>(double v)
265 {
266 #if CV_CUDEV_ARCH >= 130
267 return __double2uint_rn(v);
268 #else
269 return saturate_cast<uint>((float) v);
270 #endif
271 }
272
273 //! @}
274
275 }}
276
277 #endif
278