minor update of device layer
This commit is contained in:
Anatoly Baksheev
parent
0f5f57561e
commit
c19f88853a
69
modules/gpu/src/opencv2/gpu/device/funcattrib.hpp
Normal file
69
modules/gpu/src/opencv2/gpu/device/funcattrib.hpp
Normal file
@ -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<cstdio>
|
||||||
|
|
||||||
|
namespace pcl
|
||||||
|
{
|
||||||
|
namespace device
|
||||||
|
{
|
||||||
|
template<class Func>
|
||||||
|
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_ */
|
||||||
@ -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 <class Worker, typename T>
|
|
||||||
__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 <class Worker, typename InIter, typename OutIter>
|
|
||||||
__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 <class Worker, typename T>
|
|
||||||
__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 <class Worker, typename OutIter>
|
|
||||||
__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__ */
|
|
||||||
@ -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_ */
|
|
||||||
66
modules/gpu/src/opencv2/gpu/device/static_check.hpp
Normal file
66
modules/gpu/src/opencv2/gpu/device/static_check.hpp
Normal file
@ -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<bool expr> struct Static {};
|
||||||
|
|
||||||
|
template<> struct Static<true>
|
||||||
|
{
|
||||||
|
__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_ */
|
||||||
99
modules/gpu/src/opencv2/gpu/device/warp.hpp
Normal file
99
modules/gpu/src/opencv2/gpu/device/warp.hpp
Normal file
@ -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<typename It, typename T>
|
||||||
|
static __device__ __forceinline__ void fill(It beg, It end, const T& value)
|
||||||
|
{
|
||||||
|
for(It t = beg + laneId(); t < end; t += STRIDE)
|
||||||
|
*t = value;
|
||||||
|
}
|
||||||
|
|
||||||
|
template<typename InIt, typename OutIt>
|
||||||
|
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<typename InIt, typename OutIt, class UnOp>
|
||||||
|
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<typename InIt1, typename InIt2, typename OutIt, class BinOp>
|
||||||
|
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_ */
|
||||||
Loading…
x
Reference in New Issue
Block a user