saturate_cast.hpp
Go to the documentation of this file.
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_GPU_SATURATE_CAST_HPP__
44 #define __OPENCV_GPU_SATURATE_CAST_HPP__
45 
46 #include "common.hpp"
47 
48 namespace cv { namespace gpu { namespace device
49 {
50  template<typename _Tp> __device__ __forceinline__ _Tp saturate_cast(uchar v) { return _Tp(v); }
51  template<typename _Tp> __device__ __forceinline__ _Tp saturate_cast(schar v) { return _Tp(v); }
52  template<typename _Tp> __device__ __forceinline__ _Tp saturate_cast(ushort v) { return _Tp(v); }
53  template<typename _Tp> __device__ __forceinline__ _Tp saturate_cast(short v) { return _Tp(v); }
54  template<typename _Tp> __device__ __forceinline__ _Tp saturate_cast(uint 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); }
58 
59  template<> __device__ __forceinline__ uchar saturate_cast<uchar>(schar v)
60  {
61  uint res = 0;
62  int vi = v;
63  asm("cvt.sat.u8.s8 %0, %1;" : "=r"(res) : "r"(vi));
64  return res;
65  }
66  template<> __device__ __forceinline__ uchar saturate_cast<uchar>(short v)
67  {
68  uint res = 0;
69  asm("cvt.sat.u8.s16 %0, %1;" : "=r"(res) : "h"(v));
70  return res;
71  }
72  template<> __device__ __forceinline__ uchar saturate_cast<uchar>(ushort v)
73  {
74  uint res = 0;
75  asm("cvt.sat.u8.u16 %0, %1;" : "=r"(res) : "h"(v));
76  return res;
77  }
78  template<> __device__ __forceinline__ uchar saturate_cast<uchar>(int v)
79  {
80  uint res = 0;
81  asm("cvt.sat.u8.s32 %0, %1;" : "=r"(res) : "r"(v));
82  return res;
83  }
84  template<> __device__ __forceinline__ uchar saturate_cast<uchar>(uint v)
85  {
86  uint res = 0;
87  asm("cvt.sat.u8.u32 %0, %1;" : "=r"(res) : "r"(v));
88  return res;
89  }
90  template<> __device__ __forceinline__ uchar saturate_cast<uchar>(float v)
91  {
92  uint res = 0;
93  asm("cvt.rni.sat.u8.f32 %0, %1;" : "=r"(res) : "f"(v));
94  return res;
95  }
96  template<> __device__ __forceinline__ uchar saturate_cast<uchar>(double v)
97  {
98  #if __CUDA_ARCH__ >= 130
99  uint res = 0;
100  asm("cvt.rni.sat.u8.f64 %0, %1;" : "=r"(res) : "d"(v));
101  return res;
102  #else
103  return saturate_cast<uchar>((float)v);
104  #endif
105  }
106 
107  template<> __device__ __forceinline__ schar saturate_cast<schar>(uchar v)
108  {
109  uint res = 0;
110  uint vi = v;
111  asm("cvt.sat.s8.u8 %0, %1;" : "=r"(res) : "r"(vi));
112  return res;
113  }
114  template<> __device__ __forceinline__ schar saturate_cast<schar>(short v)
115  {
116  uint res = 0;
117  asm("cvt.sat.s8.s16 %0, %1;" : "=r"(res) : "h"(v));
118  return res;
119  }
120  template<> __device__ __forceinline__ schar saturate_cast<schar>(ushort v)
121  {
122  uint res = 0;
123  asm("cvt.sat.s8.u16 %0, %1;" : "=r"(res) : "h"(v));
124  return res;
125  }
126  template<> __device__ __forceinline__ schar saturate_cast<schar>(int v)
127  {
128  uint res = 0;
129  asm("cvt.sat.s8.s32 %0, %1;" : "=r"(res) : "r"(v));
130  return res;
131  }
132  template<> __device__ __forceinline__ schar saturate_cast<schar>(uint v)
133  {
134  uint res = 0;
135  asm("cvt.sat.s8.u32 %0, %1;" : "=r"(res) : "r"(v));
136  return res;
137  }
138  template<> __device__ __forceinline__ schar saturate_cast<schar>(float v)
139  {
140  uint res = 0;
141  asm("cvt.rni.sat.s8.f32 %0, %1;" : "=r"(res) : "f"(v));
142  return res;
143  }
144  template<> __device__ __forceinline__ schar saturate_cast<schar>(double v)
145  {
146  #if __CUDA_ARCH__ >= 130
147  uint res = 0;
148  asm("cvt.rni.sat.s8.f64 %0, %1;" : "=r"(res) : "d"(v));
149  return res;
150  #else
151  return saturate_cast<schar>((float)v);
152  #endif
153  }
154 
155  template<> __device__ __forceinline__ ushort saturate_cast<ushort>(schar v)
156  {
157  ushort res = 0;
158  int vi = v;
159  asm("cvt.sat.u16.s8 %0, %1;" : "=h"(res) : "r"(vi));
160  return res;
161  }
162  template<> __device__ __forceinline__ ushort saturate_cast<ushort>(short v)
163  {
164  ushort res = 0;
165  asm("cvt.sat.u16.s16 %0, %1;" : "=h"(res) : "h"(v));
166  return res;
167  }
168  template<> __device__ __forceinline__ ushort saturate_cast<ushort>(int v)
169  {
170  ushort res = 0;
171  asm("cvt.sat.u16.s32 %0, %1;" : "=h"(res) : "r"(v));
172  return res;
173  }
174  template<> __device__ __forceinline__ ushort saturate_cast<ushort>(uint v)
175  {
176  ushort res = 0;
177  asm("cvt.sat.u16.u32 %0, %1;" : "=h"(res) : "r"(v));
178  return res;
179  }
180  template<> __device__ __forceinline__ ushort saturate_cast<ushort>(float v)
181  {
182  ushort res = 0;
183  asm("cvt.rni.sat.u16.f32 %0, %1;" : "=h"(res) : "f"(v));
184  return res;
185  }
186  template<> __device__ __forceinline__ ushort saturate_cast<ushort>(double v)
187  {
188  #if __CUDA_ARCH__ >= 130
189  ushort res = 0;
190  asm("cvt.rni.sat.u16.f64 %0, %1;" : "=h"(res) : "d"(v));
191  return res;
192  #else
193  return saturate_cast<ushort>((float)v);
194  #endif
195  }
196 
197  template<> __device__ __forceinline__ short saturate_cast<short>(ushort v)
198  {
199  short res = 0;
200  asm("cvt.sat.s16.u16 %0, %1;" : "=h"(res) : "h"(v));
201  return res;
202  }
203  template<> __device__ __forceinline__ short saturate_cast<short>(int v)
204  {
205  short res = 0;
206  asm("cvt.sat.s16.s32 %0, %1;" : "=h"(res) : "r"(v));
207  return res;
208  }
209  template<> __device__ __forceinline__ short saturate_cast<short>(uint v)
210  {
211  short res = 0;
212  asm("cvt.sat.s16.u32 %0, %1;" : "=h"(res) : "r"(v));
213  return res;
214  }
215  template<> __device__ __forceinline__ short saturate_cast<short>(float v)
216  {
217  short res = 0;
218  asm("cvt.rni.sat.s16.f32 %0, %1;" : "=h"(res) : "f"(v));
219  return res;
220  }
221  template<> __device__ __forceinline__ short saturate_cast<short>(double v)
222  {
223  #if __CUDA_ARCH__ >= 130
224  short res = 0;
225  asm("cvt.rni.sat.s16.f64 %0, %1;" : "=h"(res) : "d"(v));
226  return res;
227  #else
228  return saturate_cast<short>((float)v);
229  #endif
230  }
231 
232  template<> __device__ __forceinline__ int saturate_cast<int>(uint v)
233  {
234  int res = 0;
235  asm("cvt.sat.s32.u32 %0, %1;" : "=r"(res) : "r"(v));
236  return res;
237  }
238  template<> __device__ __forceinline__ int saturate_cast<int>(float v)
239  {
240  return __float2int_rn(v);
241  }
242  template<> __device__ __forceinline__ int saturate_cast<int>(double v)
243  {
244  #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130
245  return __double2int_rn(v);
246  #else
247  return saturate_cast<int>((float)v);
248  #endif
249  }
250 
251  template<> __device__ __forceinline__ uint saturate_cast<uint>(schar v)
252  {
253  uint res = 0;
254  int vi = v;
255  asm("cvt.sat.u32.s8 %0, %1;" : "=r"(res) : "r"(vi));
256  return res;
257  }
258  template<> __device__ __forceinline__ uint saturate_cast<uint>(short v)
259  {
260  uint res = 0;
261  asm("cvt.sat.u32.s16 %0, %1;" : "=r"(res) : "h"(v));
262  return res;
263  }
264  template<> __device__ __forceinline__ uint saturate_cast<uint>(int v)
265  {
266  uint res = 0;
267  asm("cvt.sat.u32.s32 %0, %1;" : "=r"(res) : "r"(v));
268  return res;
269  }
270  template<> __device__ __forceinline__ uint saturate_cast<uint>(float v)
271  {
272  return __float2uint_rn(v);
273  }
274  template<> __device__ __forceinline__ uint saturate_cast<uint>(double v)
275  {
276  #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130
277  return __double2uint_rn(v);
278  #else
279  return saturate_cast<uint>((float)v);
280  #endif
281  }
282 }}}
283 
284 #endif /* __OPENCV_GPU_SATURATE_CAST_HPP__ */
GLuint res
signed char schar
Definition: common.hpp:102
unsigned char uchar
Definition: common.hpp:100
unsigned int uint
Definition: common.hpp:104
const GLdouble * v
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