OpenCV 4.5.3(日本語機械翻訳)
saturate_cast.hpp
[詳解]
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_CUDA_SATURATE_CAST_HPP
44 #define OPENCV_CUDA_SATURATE_CAST_HPP
45
46 #include "common.hpp"
47
53
54 namespace cv { namespace cuda { namespace device
55{
56 template<typename _Tp> __device__ __forceinline__ _Tp saturate_cast(uchar v) { return _Tp(v); }
57 template<typename _Tp> __device__ __forceinline__ _Tp saturate_cast(schar v) { return _Tp(v); }
58 template<typename _Tp> __device__ __forceinline__ _Tp saturate_cast(ushort v) { return _Tp(v); }
59 template<typename _Tp> __device__ __forceinline__ _Tp saturate_cast(short v) { return _Tp(v); }
60 template<typename _Tp> __device__ __forceinline__ _Tp saturate_cast(uint v) { return _Tp(v); }
61 template<typename _Tp> __device__ __forceinline__ _Tp saturate_cast(int v) { return _Tp(v); }
62 template<typename _Tp> __device__ __forceinline__ _Tp saturate_cast(float v) { return _Tp(v); }
63 template<typename _Tp> __device__ __forceinline__ _Tp saturate_cast(double v) { return _Tp(v); }
64
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 }
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 }
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 }
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 }
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 }
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 }
102 template<> __device__ __forceinline__ uchar saturate_cast<uchar>(double v)
103 {
104 #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130
105 uint res = 0;
106 asm("cvt.rni.sat.u8.f64 %0, %1;" : "=r"(res) : "d"(v));
107 return res;
108 #else
109 return saturate_cast<uchar>((float)v);
110 #endif
111 }
112
113 template<> __device__ __forceinline__ schar saturate_cast<schar>(uchar v)
114 {
115 uint res = 0;
116 uint vi = v;
117 asm("cvt.sat.s8.u8 %0, %1;" : "=r"(res) : "r"(vi));
118 return res;
119 }
120 template<> __device__ __forceinline__ schar saturate_cast<schar>(short v)
121 {
122 uint res = 0;
123 asm("cvt.sat.s8.s16 %0, %1;" : "=r"(res) : "h"(v));
124 return res;
125 }
126 template<> __device__ __forceinline__ schar saturate_cast<schar>(ushort v)
127 {
128 uint res = 0;
129 asm("cvt.sat.s8.u16 %0, %1;" : "=r"(res) : "h"(v));
130 return res;
131 }
132 template<> __device__ __forceinline__ schar saturate_cast<schar>(int v)
133 {
134 uint res = 0;
135 asm("cvt.sat.s8.s32 %0, %1;" : "=r"(res) : "r"(v));
136 return res;
137 }
138 template<> __device__ __forceinline__ schar saturate_cast<schar>(uint v)
139 {
140 uint res = 0;
141 asm("cvt.sat.s8.u32 %0, %1;" : "=r"(res) : "r"(v));
142 return res;
143 }
144 template<> __device__ __forceinline__ schar saturate_cast<schar>(float v)
145 {
146 uint res = 0;
147 asm("cvt.rni.sat.s8.f32 %0, %1;" : "=r"(res) : "f"(v));
148 return res;
149 }
150 template<> __device__ __forceinline__ schar saturate_cast<schar>(double v)
151 {
152 #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130
153 uint res = 0;
154 asm("cvt.rni.sat.s8.f64 %0, %1;" : "=r"(res) : "d"(v));
155 return res;
156 #else
157 return saturate_cast<schar>((float)v);
158 #endif
159 }
160
161 template<> __device__ __forceinline__ ushort saturate_cast<ushort>(schar v)
162 {
163 ushort res = 0;
164 int vi = v;
165 asm("cvt.sat.u16.s8 %0, %1;" : "=h"(res) : "r"(vi));
166 return res;
167 }
168 template<> __device__ __forceinline__ ushort saturate_cast<ushort>(short v)
169 {
170 ushort res = 0;
171 asm("cvt.sat.u16.s16 %0, %1;" : "=h"(res) : "h"(v));
172 return res;
173 }
174 template<> __device__ __forceinline__ ushort saturate_cast<ushort>(int v)
175 {
176 ushort res = 0;
177 asm("cvt.sat.u16.s32 %0, %1;" : "=h"(res) : "r"(v));
178 return res;
179 }
180 template<> __device__ __forceinline__ ushort saturate_cast<ushort>(uint v)
181 {
182 ushort res = 0;
183 asm("cvt.sat.u16.u32 %0, %1;" : "=h"(res) : "r"(v));
184 return res;
185 }
186 template<> __device__ __forceinline__ ushort saturate_cast<ushort>(float v)
187 {
188 ushort res = 0;
189 asm("cvt.rni.sat.u16.f32 %0, %1;" : "=h"(res) : "f"(v));
190 return res;
191 }
192 template<> __device__ __forceinline__ ushort saturate_cast<ushort>(double v)
193 {
194 #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130
195 ushort res = 0;
196 asm("cvt.rni.sat.u16.f64 %0, %1;" : "=h"(res) : "d"(v));
197 return res;
198 #else
199 return saturate_cast<ushort>((float)v);
200 #endif
201 }
202
203 template<> __device__ __forceinline__ short saturate_cast<short>(ushort v)
204 {
205 short res = 0;
206 asm("cvt.sat.s16.u16 %0, %1;" : "=h"(res) : "h"(v));
207 return res;
208 }
209 template<> __device__ __forceinline__ short saturate_cast<short>(int v)
210 {
211 short res = 0;
212 asm("cvt.sat.s16.s32 %0, %1;" : "=h"(res) : "r"(v));
213 return res;
214 }
215 template<> __device__ __forceinline__ short saturate_cast<short>(uint v)
216 {
217 short res = 0;
218 asm("cvt.sat.s16.u32 %0, %1;" : "=h"(res) : "r"(v));
219 return res;
220 }
221 template<> __device__ __forceinline__ short saturate_cast<short>(float v)
222 {
223 short res = 0;
224 asm("cvt.rni.sat.s16.f32 %0, %1;" : "=h"(res) : "f"(v));
225 return res;
226 }
227 template<> __device__ __forceinline__ short saturate_cast<short>(double v)
228 {
229 #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130
230 short res = 0;
231 asm("cvt.rni.sat.s16.f64 %0, %1;" : "=h"(res) : "d"(v));
232 return res;
233 #else
234 return saturate_cast<short>((float)v);
235 #endif
236 }
237
238 template<> __device__ __forceinline__ int saturate_cast<int>(uint v)
239 {
240 int res = 0;
241 asm("cvt.sat.s32.u32 %0, %1;" : "=r"(res) : "r"(v));
242 return res;
243 }
244 template<> __device__ __forceinline__ int saturate_cast<int>(float v)
245 {
246 return __float2int_rn(v);
247 }
248 template<> __device__ __forceinline__ int saturate_cast<int>(double v)
249 {
250 #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130
251 return __double2int_rn(v);
252 #else
253 return saturate_cast<int>((float)v);
254 #endif
255 }
256
257 template<> __device__ __forceinline__ uint saturate_cast<uint>(schar v)
258 {
259 uint res = 0;
260 int vi = v;
261 asm("cvt.sat.u32.s8 %0, %1;" : "=r"(res) : "r"(vi));
262 return res;
263 }
264 template<> __device__ __forceinline__ uint saturate_cast<uint>(short v)
265 {
266 uint res = 0;
267 asm("cvt.sat.u32.s16 %0, %1;" : "=r"(res) : "h"(v));
268 return res;
269 }
270 template<> __device__ __forceinline__ uint saturate_cast<uint>(int v)
271 {
272 uint res = 0;
273 asm("cvt.sat.u32.s32 %0, %1;" : "=r"(res) : "r"(v));
274 return res;
275 }
276 template<> __device__ __forceinline__ uint saturate_cast<uint>(float v)
277 {
278 return __float2uint_rn(v);
279 }
280 template<> __device__ __forceinline__ uint saturate_cast<uint>(double v)
281 {
282 #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130
283 return __double2uint_rn(v);
284 #else
285 return saturate_cast<uint>((float)v);
286 #endif
287 }
288}}}
289
291
292 #endif /* OPENCV_CUDA_SATURATE_CAST_HPP */
cv::saturate_cast
static _Tp saturate_cast(uchar v)
Template function for accurate conversion from one primitive type to another.
Definition: saturate.hpp:80
cv
"black box" representation of the file storage associated with a file on disk.
Definition: aruco.hpp:75
構築: doxygen 1.9.2