warp.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_DEVICE_WARP_HPP__
00044 #define __OPENCV_CUDA_DEVICE_WARP_HPP__
00045 
00046 /** @file
00047  * @deprecated Use @ref cudev instead.
00048  */
00049 
00050 //! @cond IGNORED
00051 
00052 namespace cv { namespace cuda { namespace device
00053 {
00054     struct Warp
00055     {
00056         enum
00057         {
00058             LOG_WARP_SIZE = 5,
00059             WARP_SIZE     = 1 << LOG_WARP_SIZE,
00060             STRIDE        = WARP_SIZE
00061         };
00062 
00063         /** \brief Returns the warp lane ID of the calling thread. */
00064         static __device__ __forceinline__ unsigned int laneId()
00065         {
00066             unsigned int ret;
00067             asm("mov.u32 %0, %laneid;" : "=r"(ret) );
00068             return ret;
00069         }
00070 
00071         template<typename It, typename T>
00072         static __device__ __forceinline__ void fill(It beg, It end, const T& value)
00073         {
00074             for(It t = beg + laneId(); t < end; t += STRIDE)
00075                 *t = value;
00076         }
00077 
00078         template<typename InIt, typename OutIt>
00079         static __device__ __forceinline__ OutIt copy(InIt beg, InIt end, OutIt out)
00080         {
00081             for(InIt t = beg + laneId(); t < end; t += STRIDE, out += STRIDE)
00082                 *out = *t;
00083             return out;
00084         }
00085 
00086         template<typename InIt, typename OutIt, class UnOp>
00087         static __device__ __forceinline__ OutIt transform(InIt beg, InIt end, OutIt out, UnOp op)
00088         {
00089             for(InIt t = beg + laneId(); t < end; t += STRIDE, out += STRIDE)
00090                 *out = op(*t);
00091             return out;
00092         }
00093 
00094         template<typename InIt1, typename InIt2, typename OutIt, class BinOp>
00095         static __device__ __forceinline__ OutIt transform(InIt1 beg1, InIt1 end1, InIt2 beg2, OutIt out, BinOp op)
00096         {
00097             unsigned int lane = laneId();
00098 
00099             InIt1 t1 = beg1 + lane;
00100             InIt2 t2 = beg2 + lane;
00101             for(; t1 < end1; t1 += STRIDE, t2 += STRIDE, out += STRIDE)
00102                 *out = op(*t1, *t2);
00103             return out;
00104         }
00105 
00106         template <class T, class BinOp>
00107         static __device__ __forceinline__ T reduce(volatile T *ptr, BinOp op)
00108         {
00109             const unsigned int lane = laneId();
00110 
00111             if (lane < 16)
00112             {
00113                 T partial = ptr[lane];
00114 
00115                 ptr[lane] = partial = op(partial, ptr[lane + 16]);
00116                 ptr[lane] = partial = op(partial, ptr[lane + 8]);
00117                 ptr[lane] = partial = op(partial, ptr[lane + 4]);
00118                 ptr[lane] = partial = op(partial, ptr[lane + 2]);
00119                 ptr[lane] = partial = op(partial, ptr[lane + 1]);
00120             }
00121 
00122             return *ptr;
00123         }
00124 
00125         template<typename OutIt, typename T>
00126         static __device__ __forceinline__ void yota(OutIt beg, OutIt end, T value)
00127         {
00128             unsigned int lane = laneId();
00129             value += lane;
00130 
00131             for(OutIt t = beg + lane; t < end; t += STRIDE, value += STRIDE)
00132                 *t = value;
00133         }
00134     };
00135 }}} // namespace cv { namespace cuda { namespace cudev
00136 
00137 //! @endcond
00138 
00139 #endif /* __OPENCV_CUDA_DEVICE_WARP_HPP__ */
00140