From 39373cd9f991cd112bcc1a7eef3291d45d36d476 Mon Sep 17 00:00:00 2001 From: Anatoly Baksheev Date: Tue, 30 Aug 2011 09:08:41 +0000 Subject: [PATCH] minor update of device layer --- modules/gpu/src/opencv2/gpu/device/funcattrib.hpp | 121 +++++++------ .../gpu/src/opencv2/gpu/device/static_check.hpp | 110 ++++++------ modules/gpu/src/opencv2/gpu/device/warp.hpp | 191 +++++++++++---------- 3 files changed, 228 insertions(+), 194 deletions(-) diff --git a/modules/gpu/src/opencv2/gpu/device/funcattrib.hpp b/modules/gpu/src/opencv2/gpu/device/funcattrib.hpp index d38990b..a3fe1c2 100644 --- a/modules/gpu/src/opencv2/gpu/device/funcattrib.hpp +++ b/modules/gpu/src/opencv2/gpu/device/funcattrib.hpp @@ -1,69 +1,78 @@ -/* -* Software License Agreement (BSD License) -* -* Copyright (c) 2011, Willow Garage, Inc. -* All rights reserved. -* -* Redistribution and use in source and binary forms, with or without -* modification, are permitted provided that the following conditions -* are met: -* -* * Redistributions of source code must retain the above copyright -* notice, this list of conditions and the following disclaimer. -* * Redistributions in binary form must reproduce the above -* copyright notice, this list of conditions and the following -* disclaimer in the documentation and/or other materials provided -* with the distribution. -* * Neither the name of Willow Garage, Inc. nor the names of its -* contributors may be used to endorse or promote products derived -* from this software without specific prior written permission. -* -* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS -* "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT -* LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS -* FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE -* COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, -* INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, -* BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; -* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER -* CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT -* LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN -* ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE -* POSSIBILITY OF SUCH DAMAGE. -* -* Author: Anatoly Baskeheev, Itseez Ltd, (myname.mysurname@mycompany.com) -*/ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors "as is" and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ -#ifndef PCL_DEVICE_FUNCATTRIB_HPP_ -#define PCL_DEVICE_FUNCATTRIB_HPP_ + +#ifndef __OPENCV_GPU_DEVICE_FUNCATTRIB_HPP_ +#define __OPENCV_GPU_DEVICE_FUNCATTRIB_HPP_ #include -namespace pcl +namespace cv { - namespace device + namespace gpu { - template - void printFuncAttrib(Func& func) + namespace device { + template + void printFuncAttrib(Func& func) + { - cudaFuncAttributes attrs; - cudaFuncGetAttributes(&attrs, func); + cudaFuncAttributes attrs; + cudaFuncGetAttributes(&attrs, func); - printf("=== Function stats ===\n"); - printf("Name: \n"); - printf("sharedSizeBytes = %d\n", attrs.sharedSizeBytes); - printf("constSizeBytes = %d\n", attrs.constSizeBytes); - printf("localSizeBytes = %d\n", attrs.localSizeBytes); - printf("maxThreadsPerBlock = %d\n", attrs.maxThreadsPerBlock); - printf("numRegs = %d\n", attrs.numRegs); - printf("ptxVersion = %d\n", attrs.ptxVersion); - printf("binaryVersion = %d\n", attrs.binaryVersion); - printf("\n"); - fflush(stdout); + printf("=== Function stats ===\n"); + printf("Name: \n"); + printf("sharedSizeBytes = %d\n", attrs.sharedSizeBytes); + printf("constSizeBytes = %d\n", attrs.constSizeBytes); + printf("localSizeBytes = %d\n", attrs.localSizeBytes); + printf("maxThreadsPerBlock = %d\n", attrs.maxThreadsPerBlock); + printf("numRegs = %d\n", attrs.numRegs); + printf("ptxVersion = %d\n", attrs.ptxVersion); + printf("binaryVersion = %d\n", attrs.binaryVersion); + printf("\n"); + fflush(stdout); + } } } - } -#endif /* PCL_DEVICE_FUNCATTRIB_HPP_ */ \ No newline at end of file +#endif /* __OPENCV_GPU_DEVICE_FUNCATTRIB_HPP_ */ \ No newline at end of file diff --git a/modules/gpu/src/opencv2/gpu/device/static_check.hpp b/modules/gpu/src/opencv2/gpu/device/static_check.hpp index 7e3a139..28da9a9 100644 --- a/modules/gpu/src/opencv2/gpu/device/static_check.hpp +++ b/modules/gpu/src/opencv2/gpu/device/static_check.hpp @@ -1,66 +1,72 @@ -/* -* Software License Agreement (BSD License) -* -* Copyright (c) 2011, Willow Garage, Inc. -* All rights reserved. -* -* Redistribution and use in source and binary forms, with or without -* modification, are permitted provided that the following conditions -* are met: -* -* * Redistributions of source code must retain the above copyright -* notice, this list of conditions and the following disclaimer. -* * Redistributions in binary form must reproduce the above -* copyright notice, this list of conditions and the following -* disclaimer in the documentation and/or other materials provided -* with the distribution. -* * Neither the name of Willow Garage, Inc. nor the names of its -* contributors may be used to endorse or promote products derived -* from this software without specific prior written permission. -* -* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS -* "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT -* LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS -* FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE -* COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, -* INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, -* BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; -* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER -* CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT -* LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN -* ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE -* POSSIBILITY OF SUCH DAMAGE. -* -* Author: Anatoly Baskeheev, Itseez Ltd, (myname.mysurname@mycompany.com) -*/ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors "as is" and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ -#ifndef PCL_GPU_DEVICE_STATIC_CHECK_HPP_ -#define PCL_GPU_DEVICE_STATIC_CHECK_HPP_ +#ifndef __OPENCV_GPU_GPU_DEVICE_STATIC_CHECK_HPP__ +#define __OPENCV_GPU_GPU_DEVICE_STATIC_CHECK_HPP__ #if defined(__CUDACC__) - #define __PCL_GPU_HOST_DEVICE__ __host__ __device__ __forceinline__ + #define __OPENCV_GPU_HOST_DEVICE__ __host__ __device__ __forceinline__ #else - #define __PCL_GPU_HOST_DEVICE__ + #define __OPENCV_GPU_HOST_DEVICE__ #endif -namespace pcl +namespace cv { - namespace device - { - template struct Static {}; - - template<> struct Static - { - __PCL_GPU_HOST_DEVICE__ static void check() {}; - }; - } - namespace gpu { - using pcl::device::Static; + namespace device + { + template struct Static {}; + + template<> struct Static + { + __OPENCV_GPU_HOST_DEVICE__ static void check() {}; + }; + } + + using cv::gpu::device::Static; } } #undef __PCL_GPU_HOST_DEVICE__ -#endif /* PCL_GPU_DEVICE_STATIC_CHECK_HPP_ */ \ No newline at end of file +#endif /* __OPENCV_GPU_GPU_DEVICE_STATIC_CHECK_HPP__ */ \ No newline at end of file diff --git a/modules/gpu/src/opencv2/gpu/device/warp.hpp b/modules/gpu/src/opencv2/gpu/device/warp.hpp index 987e5ba..cd9baf2 100644 --- a/modules/gpu/src/opencv2/gpu/device/warp.hpp +++ b/modules/gpu/src/opencv2/gpu/device/warp.hpp @@ -1,99 +1,118 @@ -/* -* Software License Agreement (BSD License) -* -* Copyright (c) 2011, Willow Garage, Inc. -* All rights reserved. -* -* Redistribution and use in source and binary forms, with or without -* modification, are permitted provided that the following conditions -* are met: -* -* * Redistributions of source code must retain the above copyright -* notice, this list of conditions and the following disclaimer. -* * Redistributions in binary form must reproduce the above -* copyright notice, this list of conditions and the following -* disclaimer in the documentation and/or other materials provided -* with the distribution. -* * Neither the name of Willow Garage, Inc. nor the names of its -* contributors may be used to endorse or promote products derived -* from this software without specific prior written permission. -* -* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS -* "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT -* LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS -* FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE -* COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, -* INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, -* BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; -* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER -* CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT -* LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN -* ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE -* POSSIBILITY OF SUCH DAMAGE. -* -* Author: Anatoly Baskeheev, Itseez Ltd, (myname.mysurname@mycompany.com) -*/ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors "as is" and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ -#ifndef PCL_DEVICE_UTILS_WARP_HPP_ -#define PCL_DEVICE_UTILS_WARP_HPP_ +#ifndef __OPENCV_GPU_DEVICE_WARP_HPP_ +#define __OPENCV_GPU_DEVICE_WARP_HPP_ -namespace pcl +namespace cv { - namespace device + namespace gpu { - struct Warp + namespace device { - enum + struct Warp { - LOG_WARP_SIZE = 5, - WARP_SIZE = 1 << LOG_WARP_SIZE, - STRIDE = WARP_SIZE - }; + enum + { + LOG_WARP_SIZE = 5, + WARP_SIZE = 1 << LOG_WARP_SIZE, + STRIDE = WARP_SIZE + }; - /** \brief Returns the warp lane ID of the calling thread. */ - static __device__ __forceinline__ unsigned int laneId() - { - unsigned int ret; - asm("mov.u32 %0, %laneid;" : "=r"(ret) ); - return ret; - } + /** \brief Returns the warp lane ID of the calling thread. */ + static __device__ __forceinline__ unsigned int laneId() + { + unsigned int ret; + asm("mov.u32 %0, %laneid;" : "=r"(ret) ); + return ret; + } - template - static __device__ __forceinline__ void fill(It beg, It end, const T& value) - { - for(It t = beg + laneId(); t < end; t += STRIDE) - *t = value; - } + template + static __device__ __forceinline__ void fill(It beg, It end, const T& value) + { + for(It t = beg + laneId(); t < end; t += STRIDE) + *t = value; + } - template - static __device__ __forceinline__ OutIt copy(InIt beg, InIt end, OutIt out) - { - for(InIt t = beg + laneId(); t < end; t += STRIDE, out += STRIDE) - *out = *t; - return out; - } - - template - static __device__ __forceinline__ OutIt transform(InIt beg, InIt end, OutIt out, UnOp op) - { - for(InIt t = beg + laneId(); t < end; t += STRIDE, out += STRIDE) - *out = op(*t); - return out; - } + template + static __device__ __forceinline__ OutIt copy(InIt beg, InIt end, OutIt out) + { + for(InIt t = beg + laneId(); t < end; t += STRIDE, out += STRIDE) + *out = *t; + return out; + } - template - static __device__ __forceinline__ OutIt transform(InIt1 beg1, InIt1 end1, InIt2 beg2, OutIt out, BinOp op) - { - unsigned int lane = laneId(); - - InIt1 t1 = beg1 + lane; - InIt2 t2 = beg2 + lane; - for(; t1 < end1; t1 += STRIDE, t2 += STRIDE, out += STRIDE) - *out = op(*t1, *t2); - return out; - } - }; + template + static __device__ __forceinline__ OutIt transform(InIt beg, InIt end, OutIt out, UnOp op) + { + for(InIt t = beg + laneId(); t < end; t += STRIDE, out += STRIDE) + *out = op(*t); + return out; + } + + template + static __device__ __forceinline__ OutIt transform(InIt1 beg1, InIt1 end1, InIt2 beg2, OutIt out, BinOp op) + { + unsigned int lane = laneId(); + + InIt1 t1 = beg1 + lane; + InIt2 t2 = beg2 + lane; + for(; t1 < end1; t1 += STRIDE, t2 += STRIDE, out += STRIDE) + *out = op(*t1, *t2); + return out; + } + + template + static __device__ __forceinline__ void yota(OutIt beg, OutIt end, T value) + { + unsigned int lane = laneId(); + value += lane; + + for(OutIt t = beg + lane; t < end; t += STRIDE, value += STRIDE) + *t = value; + } + }; + } } } -#endif /* PCL_DEVICE_UTILS_WARP_HPP_ */ \ No newline at end of file +#endif /* __OPENCV_GPU_DEVICE_WARP_HPP_ */ \ No newline at end of file -- 2.7.4