From c19f88853a8e346ad129f4b58070540b6e7663fc Mon Sep 17 00:00:00 2001 From: Anatoly Baksheev Date: Tue, 30 Aug 2011 09:04:31 +0000 Subject: [PATCH] minor update of device layer --- .../gpu/src/opencv2/gpu/device/funcattrib.hpp | 69 ++++++++++ .../gpu/src/opencv2/gpu/device/kernels.hpp | 123 ------------------ modules/gpu/src/opencv2/gpu/device/laneid.hpp | 61 --------- .../src/opencv2/gpu/device/static_check.hpp | 66 ++++++++++ modules/gpu/src/opencv2/gpu/device/warp.hpp | 99 ++++++++++++++ 5 files changed, 234 insertions(+), 184 deletions(-) create mode 100644 modules/gpu/src/opencv2/gpu/device/funcattrib.hpp delete mode 100644 modules/gpu/src/opencv2/gpu/device/kernels.hpp delete mode 100644 modules/gpu/src/opencv2/gpu/device/laneid.hpp create mode 100644 modules/gpu/src/opencv2/gpu/device/static_check.hpp create mode 100644 modules/gpu/src/opencv2/gpu/device/warp.hpp diff --git a/modules/gpu/src/opencv2/gpu/device/funcattrib.hpp b/modules/gpu/src/opencv2/gpu/device/funcattrib.hpp new file mode 100644 index 0000000000..d38990bf50 --- /dev/null +++ b/modules/gpu/src/opencv2/gpu/device/funcattrib.hpp @@ -0,0 +1,69 @@ +/* +* 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) +*/ + +#ifndef PCL_DEVICE_FUNCATTRIB_HPP_ +#define PCL_DEVICE_FUNCATTRIB_HPP_ + +#include + +namespace pcl +{ + namespace device + { + template + void printFuncAttrib(Func& 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); + } + } + +} + +#endif /* PCL_DEVICE_FUNCATTRIB_HPP_ */ \ No newline at end of file diff --git a/modules/gpu/src/opencv2/gpu/device/kernels.hpp b/modules/gpu/src/opencv2/gpu/device/kernels.hpp deleted file mode 100644 index 5f3f04161f..0000000000 --- a/modules/gpu/src/opencv2/gpu/device/kernels.hpp +++ /dev/null @@ -1,123 +0,0 @@ -/*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 __OPENCV_GPU_KENELRS_HPP__ -#define __OPENCV_GPU_KENELRS_HPP__ - -namespace cv -{ - namespace gpu - { - namespace device - { - struct Grid1D - { - static __forceinline__ __device__ int STRIDE() { return gridDim.x * blockDim.x; } - static __forceinline__ __device__ int SHIFT() { return blockIdx.x * blockDim.x + threadIdx.x; } - }; - - struct Block1D - { - static __forceinline__ __device__ int STRIDE() { return blockDim.x; } - static __forceinline__ __device__ int SHIFT() { return threadIdx.x; } - }; - - struct Warp - { - static __forceinline__ __device__ int STRIDE() { return warpSize }; - static __forceinline__ __device__ int SHIFT() { return threadIdx.x & (warpSize - 1); } - }; - - template - __forceinline__ __device__ void Copy(const T* in, T *out, int length) - { - int STRIDE = Worker::STRIDE(); - int idx = Worker::SHIFT(); - - for (; idx < length; idx += STRIDE) - out[idx] = in[idx]; - } - - template - __forceinline__ __device__ void Copy(InIter beg, InIter end, OutIter out) - { - int STRIDE = Worker::STRIDE(); - int SHIFT = Worker::SHIFT(); - - beg += SHIFT; - out += SHIFT; - - for (; beg < end; beg += STRIDE, out += STRIDE) - *out = *beg; - } - - template - __forceinline__ __device__ void Yota(T* out, int beg, int end) - { - int STRIDE = Worker::STRIDE(); - int SHIFT = Worker::SHIFT(); - - int idx = SHIFT; - int cur = beg + SHIFT; - int length = end - beg; - - for (; idx < length; idx += STRIDE, cur += STRIDE) - out[idx] = cur; - } - - template - __forceinline__ __device__ void Yota(OutIter beg, OutIter end, int val) - { - int STRIDE = Worker::STRIDE(); - int SHIFT = Worker::SHIFT(); - - beg += SHIFT; - val += SHIFT; - - for (; beg < end; beg += STRIDE, val += STRIDE) - *beg = val; - } - } - } -} - -#endif /* __OPENCV_GPU_KENELRS_HPP__ */ \ No newline at end of file diff --git a/modules/gpu/src/opencv2/gpu/device/laneid.hpp b/modules/gpu/src/opencv2/gpu/device/laneid.hpp deleted file mode 100644 index 272773d2f1..0000000000 --- a/modules/gpu/src/opencv2/gpu/device/laneid.hpp +++ /dev/null @@ -1,61 +0,0 @@ -/*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 OPENCV_GPU_LANEID_HPP_ -#define OPENCV_GPU_LANEID_HPP_ - -namespace cv -{ - namespace device - { - // Returns the warp lane ID of the calling thread - __device__ __forceinline__ unsigned int LaneId() - { - unsigned int ret; - asm("mov.u32 %0, %laneid;" : "=r"(ret) ); - return ret; - } - } -} - -#endif /* OPENCV_GPU_LANEID_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 new file mode 100644 index 0000000000..7e3a139b64 --- /dev/null +++ b/modules/gpu/src/opencv2/gpu/device/static_check.hpp @@ -0,0 +1,66 @@ +/* +* 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) +*/ + +#ifndef PCL_GPU_DEVICE_STATIC_CHECK_HPP_ +#define PCL_GPU_DEVICE_STATIC_CHECK_HPP_ + +#if defined(__CUDACC__) + #define __PCL_GPU_HOST_DEVICE__ __host__ __device__ __forceinline__ +#else + #define __PCL_GPU_HOST_DEVICE__ +#endif + +namespace pcl +{ + namespace device + { + template struct Static {}; + + template<> struct Static + { + __PCL_GPU_HOST_DEVICE__ static void check() {}; + }; + } + + namespace gpu + { + using pcl::device::Static; + } +} + +#undef __PCL_GPU_HOST_DEVICE__ + +#endif /* PCL_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 new file mode 100644 index 0000000000..987e5babc1 --- /dev/null +++ b/modules/gpu/src/opencv2/gpu/device/warp.hpp @@ -0,0 +1,99 @@ +/* +* 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) +*/ + +#ifndef PCL_DEVICE_UTILS_WARP_HPP_ +#define PCL_DEVICE_UTILS_WARP_HPP_ + +namespace pcl +{ + namespace device + { + struct Warp + { + 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; + } + + 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 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; + } + }; + } +} + +#endif /* PCL_DEVICE_UTILS_WARP_HPP_ */ \ No newline at end of file -- GitLab