Renesas GR-PEACH OpenCV Development
/
gr-peach-opencv-project-sd-card_update
Important changes to repositories hosted on mbed.com
Mbed hosted mercurial repositories are deprecated and are due to be permanently deleted in July 2026.
To keep a copy of this software download the repository Zip archive or clone locally using Mercurial.
It is also possible to export all your personal repositories from the account settings page.
Fork of
gr-peach-opencv-project-sd-card
by 
the do
Embed:
(wiki syntax)
Show/hide line numbers
reduce.hpp
Go to the documentation of this file.
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_REDUCE_HPP__ 00044 #define __OPENCV_CUDA_REDUCE_HPP__ 00045 00046 #include <thrust/tuple.h> 00047 #include "detail/reduce.hpp" 00048 #include "detail/reduce_key_val.hpp" 00049 00050 /** @file 00051 * @deprecated Use @ref cudev instead. 00052 */ 00053 00054 //! @cond IGNORED 00055 00056 namespace cv { namespace cuda { namespace device 00057 { 00058 template <int N, typename T, class Op> 00059 __device__ __forceinline__ void reduce(volatile T* smem, T& val, unsigned int tid, const Op& op) 00060 { 00061 reduce_detail::Dispatcher<N>::reductor::template reduce<volatile T*, T&, const Op&>(smem, val, tid, op); 00062 } 00063 template <int N, 00064 typename P0, typename P1, typename P2, typename P3, typename P4, typename P5, typename P6, typename P7, typename P8, typename P9, 00065 typename R0, typename R1, typename R2, typename R3, typename R4, typename R5, typename R6, typename R7, typename R8, typename R9, 00066 class Op0, class Op1, class Op2, class Op3, class Op4, class Op5, class Op6, class Op7, class Op8, class Op9> 00067 __device__ __forceinline__ void reduce(const thrust::tuple<P0, P1, P2, P3, P4, P5, P6, P7, P8, P9>& smem, 00068 const thrust::tuple<R0, R1, R2, R3, R4, R5, R6, R7, R8, R9>& val, 00069 unsigned int tid, 00070 const thrust::tuple<Op0, Op1, Op2, Op3, Op4, Op5, Op6, Op7, Op8, Op9>& op) 00071 { 00072 reduce_detail::Dispatcher<N>::reductor::template reduce< 00073 const thrust::tuple<P0, P1, P2, P3, P4, P5, P6, P7, P8, P9>&, 00074 const thrust::tuple<R0, R1, R2, R3, R4, R5, R6, R7, R8, R9>&, 00075 const thrust::tuple<Op0, Op1, Op2, Op3, Op4, Op5, Op6, Op7, Op8, Op9>&>(smem, val, tid, op); 00076 } 00077 00078 template <unsigned int N, typename K, typename V, class Cmp> 00079 __device__ __forceinline__ void reduceKeyVal(volatile K* skeys, K& key, volatile V* svals, V& val, unsigned int tid, const Cmp& cmp) 00080 { 00081 reduce_key_val_detail::Dispatcher<N>::reductor::template reduce<volatile K*, K&, volatile V*, V&, const Cmp&>(skeys, key, svals, val, tid, cmp); 00082 } 00083 template <unsigned int N, 00084 typename K, 00085 typename VP0, typename VP1, typename VP2, typename VP3, typename VP4, typename VP5, typename VP6, typename VP7, typename VP8, typename VP9, 00086 typename VR0, typename VR1, typename VR2, typename VR3, typename VR4, typename VR5, typename VR6, typename VR7, typename VR8, typename VR9, 00087 class Cmp> 00088 __device__ __forceinline__ void reduceKeyVal(volatile K* skeys, K& key, 00089 const thrust::tuple<VP0, VP1, VP2, VP3, VP4, VP5, VP6, VP7, VP8, VP9>& svals, 00090 const thrust::tuple<VR0, VR1, VR2, VR3, VR4, VR5, VR6, VR7, VR8, VR9>& val, 00091 unsigned int tid, const Cmp& cmp) 00092 { 00093 reduce_key_val_detail::Dispatcher<N>::reductor::template reduce<volatile K*, K&, 00094 const thrust::tuple<VP0, VP1, VP2, VP3, VP4, VP5, VP6, VP7, VP8, VP9>&, 00095 const thrust::tuple<VR0, VR1, VR2, VR3, VR4, VR5, VR6, VR7, VR8, VR9>&, 00096 const Cmp&>(skeys, key, svals, val, tid, cmp); 00097 } 00098 template <unsigned int N, 00099 typename KP0, typename KP1, typename KP2, typename KP3, typename KP4, typename KP5, typename KP6, typename KP7, typename KP8, typename KP9, 00100 typename KR0, typename KR1, typename KR2, typename KR3, typename KR4, typename KR5, typename KR6, typename KR7, typename KR8, typename KR9, 00101 typename VP0, typename VP1, typename VP2, typename VP3, typename VP4, typename VP5, typename VP6, typename VP7, typename VP8, typename VP9, 00102 typename VR0, typename VR1, typename VR2, typename VR3, typename VR4, typename VR5, typename VR6, typename VR7, typename VR8, typename VR9, 00103 class Cmp0, class Cmp1, class Cmp2, class Cmp3, class Cmp4, class Cmp5, class Cmp6, class Cmp7, class Cmp8, class Cmp9> 00104 __device__ __forceinline__ void reduceKeyVal(const thrust::tuple<KP0, KP1, KP2, KP3, KP4, KP5, KP6, KP7, KP8, KP9>& skeys, 00105 const thrust::tuple<KR0, KR1, KR2, KR3, KR4, KR5, KR6, KR7, KR8, KR9>& key, 00106 const thrust::tuple<VP0, VP1, VP2, VP3, VP4, VP5, VP6, VP7, VP8, VP9>& svals, 00107 const thrust::tuple<VR0, VR1, VR2, VR3, VR4, VR5, VR6, VR7, VR8, VR9>& val, 00108 unsigned int tid, 00109 const thrust::tuple<Cmp0, Cmp1, Cmp2, Cmp3, Cmp4, Cmp5, Cmp6, Cmp7, Cmp8, Cmp9>& cmp) 00110 { 00111 reduce_key_val_detail::Dispatcher<N>::reductor::template reduce< 00112 const thrust::tuple<KP0, KP1, KP2, KP3, KP4, KP5, KP6, KP7, KP8, KP9>&, 00113 const thrust::tuple<KR0, KR1, KR2, KR3, KR4, KR5, KR6, KR7, KR8, KR9>&, 00114 const thrust::tuple<VP0, VP1, VP2, VP3, VP4, VP5, VP6, VP7, VP8, VP9>&, 00115 const thrust::tuple<VR0, VR1, VR2, VR3, VR4, VR5, VR6, VR7, VR8, VR9>&, 00116 const thrust::tuple<Cmp0, Cmp1, Cmp2, Cmp3, Cmp4, Cmp5, Cmp6, Cmp7, Cmp8, Cmp9>& 00117 >(skeys, key, svals, val, tid, cmp); 00118 } 00119 00120 // smem_tuple 00121 00122 template <typename T0> 00123 __device__ __forceinline__ 00124 thrust::tuple<volatile T0*> 00125 smem_tuple(T0* t0) 00126 { 00127 return thrust::make_tuple((volatile T0*) t0); 00128 } 00129 00130 template <typename T0, typename T1> 00131 __device__ __forceinline__ 00132 thrust::tuple<volatile T0*, volatile T1*> 00133 smem_tuple(T0* t0, T1* t1) 00134 { 00135 return thrust::make_tuple((volatile T0*) t0, (volatile T1*) t1); 00136 } 00137 00138 template <typename T0, typename T1, typename T2> 00139 __device__ __forceinline__ 00140 thrust::tuple<volatile T0*, volatile T1*, volatile T2*> 00141 smem_tuple(T0* t0, T1* t1, T2* t2) 00142 { 00143 return thrust::make_tuple((volatile T0*) t0, (volatile T1*) t1, (volatile T2*) t2); 00144 } 00145 00146 template <typename T0, typename T1, typename T2, typename T3> 00147 __device__ __forceinline__ 00148 thrust::tuple<volatile T0*, volatile T1*, volatile T2*, volatile T3*> 00149 smem_tuple(T0* t0, T1* t1, T2* t2, T3* t3) 00150 { 00151 return thrust::make_tuple((volatile T0*) t0, (volatile T1*) t1, (volatile T2*) t2, (volatile T3*) t3); 00152 } 00153 00154 template <typename T0, typename T1, typename T2, typename T3, typename T4> 00155 __device__ __forceinline__ 00156 thrust::tuple<volatile T0*, volatile T1*, volatile T2*, volatile T3*, volatile T4*> 00157 smem_tuple(T0* t0, T1* t1, T2* t2, T3* t3, T4* t4) 00158 { 00159 return thrust::make_tuple((volatile T0*) t0, (volatile T1*) t1, (volatile T2*) t2, (volatile T3*) t3, (volatile T4*) t4); 00160 } 00161 00162 template <typename T0, typename T1, typename T2, typename T3, typename T4, typename T5> 00163 __device__ __forceinline__ 00164 thrust::tuple<volatile T0*, volatile T1*, volatile T2*, volatile T3*, volatile T4*, volatile T5*> 00165 smem_tuple(T0* t0, T1* t1, T2* t2, T3* t3, T4* t4, T5* t5) 00166 { 00167 return thrust::make_tuple((volatile T0*) t0, (volatile T1*) t1, (volatile T2*) t2, (volatile T3*) t3, (volatile T4*) t4, (volatile T5*) t5); 00168 } 00169 00170 template <typename T0, typename T1, typename T2, typename T3, typename T4, typename T5, typename T6> 00171 __device__ __forceinline__ 00172 thrust::tuple<volatile T0*, volatile T1*, volatile T2*, volatile T3*, volatile T4*, volatile T5*, volatile T6*> 00173 smem_tuple(T0* t0, T1* t1, T2* t2, T3* t3, T4* t4, T5* t5, T6* t6) 00174 { 00175 return thrust::make_tuple((volatile T0*) t0, (volatile T1*) t1, (volatile T2*) t2, (volatile T3*) t3, (volatile T4*) t4, (volatile T5*) t5, (volatile T6*) t6); 00176 } 00177 00178 template <typename T0, typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7> 00179 __device__ __forceinline__ 00180 thrust::tuple<volatile T0*, volatile T1*, volatile T2*, volatile T3*, volatile T4*, volatile T5*, volatile T6*, volatile T7*> 00181 smem_tuple(T0* t0, T1* t1, T2* t2, T3* t3, T4* t4, T5* t5, T6* t6, T7* t7) 00182 { 00183 return thrust::make_tuple((volatile T0*) t0, (volatile T1*) t1, (volatile T2*) t2, (volatile T3*) t3, (volatile T4*) t4, (volatile T5*) t5, (volatile T6*) t6, (volatile T7*) t7); 00184 } 00185 00186 template <typename T0, typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, typename T8> 00187 __device__ __forceinline__ 00188 thrust::tuple<volatile T0*, volatile T1*, volatile T2*, volatile T3*, volatile T4*, volatile T5*, volatile T6*, volatile T7*, volatile T8*> 00189 smem_tuple(T0* t0, T1* t1, T2* t2, T3* t3, T4* t4, T5* t5, T6* t6, T7* t7, T8* t8) 00190 { 00191 return thrust::make_tuple((volatile T0*) t0, (volatile T1*) t1, (volatile T2*) t2, (volatile T3*) t3, (volatile T4*) t4, (volatile T5*) t5, (volatile T6*) t6, (volatile T7*) t7, (volatile T8*) t8); 00192 } 00193 00194 template <typename T0, typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, typename T8, typename T9> 00195 __device__ __forceinline__ 00196 thrust::tuple<volatile T0*, volatile T1*, volatile T2*, volatile T3*, volatile T4*, volatile T5*, volatile T6*, volatile T7*, volatile T8*, volatile T9*> 00197 smem_tuple(T0* t0, T1* t1, T2* t2, T3* t3, T4* t4, T5* t5, T6* t6, T7* t7, T8* t8, T9* t9) 00198 { 00199 return thrust::make_tuple((volatile T0*) t0, (volatile T1*) t1, (volatile T2*) t2, (volatile T3*) t3, (volatile T4*) t4, (volatile T5*) t5, (volatile T6*) t6, (volatile T7*) t7, (volatile T8*) t8, (volatile T9*) t9); 00200 } 00201 }}} 00202 00203 //! @endcond 00204 00205 #endif // __OPENCV_CUDA_UTILITY_HPP__ 00206
Generated on Tue Jul 12 2022 14:47:34 by
1.7.2