saturate_cast.hpp

00001 /*M///////////////////////////////////////////////////////////////////////////////////////
00002 //
00003 //  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
00004 //
00005 //  By downloading, copying, installing or using the software you agree to this license.
00006 //  If you do not agree to this license, do not download, install,
00007 //  copy or use the software.
00008 //
00009 //
00010 //                           License Agreement
00011 //                For Open Source Computer Vision Library
00012 //
00013 // Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
00014 // Copyright (C) 2009, Willow Garage Inc., all rights reserved.
00015 // Third party copyrights are property of their respective owners.
00016 //
00017 // Redistribution and use in source and binary forms, with or without modification,
00018 // are permitted provided that the following conditions are met:
00019 //
00020 //   * Redistribution's of source code must retain the above copyright notice,
00021 //     this list of conditions and the following disclaimer.
00022 //
00023 //   * Redistribution's in binary form must reproduce the above copyright notice,
00024 //     this list of conditions and the following disclaimer in the documentation
00025 //     and/or other materials provided with the distribution.
00026 //
00027 //   * The name of the copyright holders may not be used to endorse or promote products
00028 //     derived from this software without specific prior written permission.
00029 //
00030 // This software is provided by the copyright holders and contributors "as is" and
00031 // any express or implied warranties, including, but not limited to, the implied
00032 // warranties of merchantability and fitness for a particular purpose are disclaimed.
00033 // In no event shall the Intel Corporation or contributors be liable for any direct,
00034 // indirect, incidental, special, exemplary, or consequential damages
00035 // (including, but not limited to, procurement of substitute goods or services;
00036 // loss of use, data, or profits; or business interruption) however caused
00037 // and on any theory of liability, whether in contract, strict liability,
00038 // or tort (including negligence or otherwise) arising in any way out of
00039 // the use of this software, even if advised of the possibility of such damage.
00040 //
00041 //M*/
00042 
00043 #ifndef __OPENCV_CUDA_SATURATE_CAST_HPP__
00044 #define __OPENCV_CUDA_SATURATE_CAST_HPP__
00045 
00046 #include "common.hpp "
00047 
00048 /** @file
00049  * @deprecated Use @ref cudev instead.
00050  */
00051 
00052 //! @cond IGNORED
00053 
00054 namespace cv { namespace cuda { namespace device
00055 {
00056     template<typename _Tp> __device__ __forceinline__ _Tp saturate_cast(uchar v) { return _Tp(v); }
00057     template<typename _Tp> __device__ __forceinline__ _Tp saturate_cast(schar v) { return _Tp(v); }
00058     template<typename _Tp> __device__ __forceinline__ _Tp saturate_cast(ushort v) { return _Tp(v); }
00059     template<typename _Tp> __device__ __forceinline__ _Tp saturate_cast(short v) { return _Tp(v); }
00060     template<typename _Tp> __device__ __forceinline__ _Tp saturate_cast(uint v) { return _Tp(v); }
00061     template<typename _Tp> __device__ __forceinline__ _Tp saturate_cast(int v) { return _Tp(v); }
00062     template<typename _Tp> __device__ __forceinline__ _Tp saturate_cast(float v) { return _Tp(v); }
00063     template<typename _Tp> __device__ __forceinline__ _Tp saturate_cast(double v) { return _Tp(v); }
00064 
00065     template<> __device__ __forceinline__ uchar saturate_cast<uchar>(schar v)
00066     {
00067         uint res = 0;
00068         int vi = v;
00069         asm("cvt.sat.u8.s8 %0, %1;" : "=r"(res) : "r"(vi));
00070         return res;
00071     }
00072     template<> __device__ __forceinline__ uchar saturate_cast<uchar>(short v)
00073     {
00074         uint res = 0;
00075         asm("cvt.sat.u8.s16 %0, %1;" : "=r"(res) : "h"(v));
00076         return res;
00077     }
00078     template<> __device__ __forceinline__ uchar saturate_cast<uchar>(ushort v)
00079     {
00080         uint res = 0;
00081         asm("cvt.sat.u8.u16 %0, %1;" : "=r"(res) : "h"(v));
00082         return res;
00083     }
00084     template<> __device__ __forceinline__ uchar saturate_cast<uchar>(int v)
00085     {
00086         uint res = 0;
00087         asm("cvt.sat.u8.s32 %0, %1;" : "=r"(res) : "r"(v));
00088         return res;
00089     }
00090     template<> __device__ __forceinline__ uchar saturate_cast<uchar>(uint v)
00091     {
00092         uint res = 0;
00093         asm("cvt.sat.u8.u32 %0, %1;" : "=r"(res) : "r"(v));
00094         return res;
00095     }
00096     template<> __device__ __forceinline__ uchar saturate_cast<uchar>(float v)
00097     {
00098         uint res = 0;
00099         asm("cvt.rni.sat.u8.f32 %0, %1;" : "=r"(res) : "f"(v));
00100         return res;
00101     }
00102     template<> __device__ __forceinline__ uchar saturate_cast<uchar>(double v)
00103     {
00104     #if __CUDA_ARCH__ >= 130
00105         uint res = 0;
00106         asm("cvt.rni.sat.u8.f64 %0, %1;" : "=r"(res) : "d"(v));
00107         return res;
00108     #else
00109         return saturate_cast<uchar>((float)v);
00110     #endif
00111     }
00112 
00113     template<> __device__ __forceinline__ schar saturate_cast<schar>(uchar v)
00114     {
00115         uint res = 0;
00116         uint vi = v;
00117         asm("cvt.sat.s8.u8 %0, %1;" : "=r"(res) : "r"(vi));
00118         return res;
00119     }
00120     template<> __device__ __forceinline__ schar saturate_cast<schar>(short v)
00121     {
00122         uint res = 0;
00123         asm("cvt.sat.s8.s16 %0, %1;" : "=r"(res) : "h"(v));
00124         return res;
00125     }
00126     template<> __device__ __forceinline__ schar saturate_cast<schar>(ushort v)
00127     {
00128         uint res = 0;
00129         asm("cvt.sat.s8.u16 %0, %1;" : "=r"(res) : "h"(v));
00130         return res;
00131     }
00132     template<> __device__ __forceinline__ schar saturate_cast<schar>(int v)
00133     {
00134         uint res = 0;
00135         asm("cvt.sat.s8.s32 %0, %1;" : "=r"(res) : "r"(v));
00136         return res;
00137     }
00138     template<> __device__ __forceinline__ schar saturate_cast<schar>(uint v)
00139     {
00140         uint res = 0;
00141         asm("cvt.sat.s8.u32 %0, %1;" : "=r"(res) : "r"(v));
00142         return res;
00143     }
00144     template<> __device__ __forceinline__ schar saturate_cast<schar>(float v)
00145     {
00146         uint res = 0;
00147         asm("cvt.rni.sat.s8.f32 %0, %1;" : "=r"(res) : "f"(v));
00148         return res;
00149     }
00150     template<> __device__ __forceinline__ schar saturate_cast<schar>(double v)
00151     {
00152     #if __CUDA_ARCH__ >= 130
00153         uint res = 0;
00154         asm("cvt.rni.sat.s8.f64 %0, %1;" : "=r"(res) : "d"(v));
00155         return res;
00156     #else
00157         return saturate_cast<schar>((float)v);
00158     #endif
00159     }
00160 
00161     template<> __device__ __forceinline__ ushort saturate_cast<ushort>(schar v)
00162     {
00163         ushort res = 0;
00164         int vi = v;
00165         asm("cvt.sat.u16.s8 %0, %1;" : "=h"(res) : "r"(vi));
00166         return res;
00167     }
00168     template<> __device__ __forceinline__ ushort saturate_cast<ushort>(short v)
00169     {
00170         ushort res = 0;
00171         asm("cvt.sat.u16.s16 %0, %1;" : "=h"(res) : "h"(v));
00172         return res;
00173     }
00174     template<> __device__ __forceinline__ ushort saturate_cast<ushort>(int v)
00175     {
00176         ushort res = 0;
00177         asm("cvt.sat.u16.s32 %0, %1;" : "=h"(res) : "r"(v));
00178         return res;
00179     }
00180     template<> __device__ __forceinline__ ushort saturate_cast<ushort>(uint v)
00181     {
00182         ushort res = 0;
00183         asm("cvt.sat.u16.u32 %0, %1;" : "=h"(res) : "r"(v));
00184         return res;
00185     }
00186     template<> __device__ __forceinline__ ushort saturate_cast<ushort>(float v)
00187     {
00188         ushort res = 0;
00189         asm("cvt.rni.sat.u16.f32 %0, %1;" : "=h"(res) : "f"(v));
00190         return res;
00191     }
00192     template<> __device__ __forceinline__ ushort saturate_cast<ushort>(double v)
00193     {
00194     #if __CUDA_ARCH__ >= 130
00195         ushort res = 0;
00196         asm("cvt.rni.sat.u16.f64 %0, %1;" : "=h"(res) : "d"(v));
00197         return res;
00198     #else
00199         return saturate_cast<ushort>((float)v);
00200     #endif
00201     }
00202 
00203     template<> __device__ __forceinline__ short saturate_cast<short>(ushort v)
00204     {
00205         short res = 0;
00206         asm("cvt.sat.s16.u16 %0, %1;" : "=h"(res) : "h"(v));
00207         return res;
00208     }
00209     template<> __device__ __forceinline__ short saturate_cast<short>(int v)
00210     {
00211         short res = 0;
00212         asm("cvt.sat.s16.s32 %0, %1;" : "=h"(res) : "r"(v));
00213         return res;
00214     }
00215     template<> __device__ __forceinline__ short saturate_cast<short>(uint v)
00216     {
00217         short res = 0;
00218         asm("cvt.sat.s16.u32 %0, %1;" : "=h"(res) : "r"(v));
00219         return res;
00220     }
00221     template<> __device__ __forceinline__ short saturate_cast<short>(float v)
00222     {
00223         short res = 0;
00224         asm("cvt.rni.sat.s16.f32 %0, %1;" : "=h"(res) : "f"(v));
00225         return res;
00226     }
00227     template<> __device__ __forceinline__ short saturate_cast<short>(double v)
00228     {
00229     #if __CUDA_ARCH__ >= 130
00230         short res = 0;
00231         asm("cvt.rni.sat.s16.f64 %0, %1;" : "=h"(res) : "d"(v));
00232         return res;
00233     #else
00234         return saturate_cast<short>((float)v);
00235     #endif
00236     }
00237 
00238     template<> __device__ __forceinline__ int saturate_cast<int>(uint v)
00239     {
00240         int res = 0;
00241         asm("cvt.sat.s32.u32 %0, %1;" : "=r"(res) : "r"(v));
00242         return res;
00243     }
00244     template<> __device__ __forceinline__ int saturate_cast<int>(float v)
00245     {
00246         return __float2int_rn(v);
00247     }
00248     template<> __device__ __forceinline__ int saturate_cast<int>(double v)
00249     {
00250     #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130
00251         return __double2int_rn(v);
00252     #else
00253         return saturate_cast<int>((float)v);
00254     #endif
00255     }
00256 
00257     template<> __device__ __forceinline__ uint saturate_cast<uint>(schar v)
00258     {
00259         uint res = 0;
00260         int vi = v;
00261         asm("cvt.sat.u32.s8 %0, %1;" : "=r"(res) : "r"(vi));
00262         return res;
00263     }
00264     template<> __device__ __forceinline__ uint saturate_cast<uint>(short v)
00265     {
00266         uint res = 0;
00267         asm("cvt.sat.u32.s16 %0, %1;" : "=r"(res) : "h"(v));
00268         return res;
00269     }
00270     template<> __device__ __forceinline__ uint saturate_cast<uint>(int v)
00271     {
00272         uint res = 0;
00273         asm("cvt.sat.u32.s32 %0, %1;" : "=r"(res) : "r"(v));
00274         return res;
00275     }
00276     template<> __device__ __forceinline__ uint saturate_cast<uint>(float v)
00277     {
00278         return __float2uint_rn(v);
00279     }
00280     template<> __device__ __forceinline__ uint saturate_cast<uint>(double v)
00281     {
00282     #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130
00283         return __double2uint_rn(v);
00284     #else
00285         return saturate_cast<uint>((float)v);
00286     #endif
00287     }
00288 }}}
00289 
00290 //! @endcond
00291 
00292 #endif /* __OPENCV_CUDA_SATURATE_CAST_HPP__ */
00293