vec_distance_detail.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_VEC_DISTANCE_DETAIL_HPP__
00044 #define __OPENCV_CUDA_VEC_DISTANCE_DETAIL_HPP__
00045 
00046 #include "../datamov_utils.hpp"
00047 
00048 //! @cond IGNORED
00049 
00050 namespace cv { namespace cuda { namespace device
00051 {
00052     namespace vec_distance_detail
00053     {
00054         template <int THREAD_DIM, int N> struct UnrollVecDiffCached
00055         {
00056             template <typename Dist, typename T1, typename T2>
00057             static __device__ void calcCheck(const T1* vecCached, const T2* vecGlob, int len, Dist& dist, int ind)
00058             {
00059                 if (ind < len)
00060                 {
00061                     T1 val1 = *vecCached++;
00062 
00063                     T2 val2;
00064                     ForceGlob<T2>::Load(vecGlob, ind, val2);
00065 
00066                     dist.reduceIter(val1, val2);
00067 
00068                     UnrollVecDiffCached<THREAD_DIM, N - 1>::calcCheck(vecCached, vecGlob, len, dist, ind + THREAD_DIM);
00069                 }
00070             }
00071 
00072             template <typename Dist, typename T1, typename T2>
00073             static __device__ void calcWithoutCheck(const T1* vecCached, const T2* vecGlob, Dist& dist)
00074             {
00075                 T1 val1 = *vecCached++;
00076 
00077                 T2 val2;
00078                 ForceGlob<T2>::Load(vecGlob, 0, val2);
00079                 vecGlob += THREAD_DIM;
00080 
00081                 dist.reduceIter(val1, val2);
00082 
00083                 UnrollVecDiffCached<THREAD_DIM, N - 1>::calcWithoutCheck(vecCached, vecGlob, dist);
00084             }
00085         };
00086         template <int THREAD_DIM> struct UnrollVecDiffCached<THREAD_DIM, 0>
00087         {
00088             template <typename Dist, typename T1, typename T2>
00089             static __device__ __forceinline__ void calcCheck(const T1*, const T2*, int, Dist&, int)
00090             {
00091             }
00092 
00093             template <typename Dist, typename T1, typename T2>
00094             static __device__ __forceinline__ void calcWithoutCheck(const T1*, const T2*, Dist&)
00095             {
00096             }
00097         };
00098 
00099         template <int THREAD_DIM, int MAX_LEN, bool LEN_EQ_MAX_LEN> struct VecDiffCachedCalculator;
00100         template <int THREAD_DIM, int MAX_LEN> struct VecDiffCachedCalculator<THREAD_DIM, MAX_LEN, false>
00101         {
00102             template <typename Dist, typename T1, typename T2>
00103             static __device__ __forceinline__ void calc(const T1* vecCached, const T2* vecGlob, int len, Dist& dist, int tid)
00104             {
00105                 UnrollVecDiffCached<THREAD_DIM, MAX_LEN / THREAD_DIM>::calcCheck(vecCached, vecGlob, len, dist, tid);
00106             }
00107         };
00108         template <int THREAD_DIM, int MAX_LEN> struct VecDiffCachedCalculator<THREAD_DIM, MAX_LEN, true>
00109         {
00110             template <typename Dist, typename T1, typename T2>
00111             static __device__ __forceinline__ void calc(const T1* vecCached, const T2* vecGlob, int len, Dist& dist, int tid)
00112             {
00113                 UnrollVecDiffCached<THREAD_DIM, MAX_LEN / THREAD_DIM>::calcWithoutCheck(vecCached, vecGlob + tid, dist);
00114             }
00115         };
00116     } // namespace vec_distance_detail
00117 }}} // namespace cv { namespace cuda { namespace cudev
00118 
00119 //! @endcond
00120 
00121 #endif // __OPENCV_CUDA_VEC_DISTANCE_DETAIL_HPP__
00122