generic-library/opencv
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

minor update of device layer

Browse Source
This commit is contained in:
Anatoly Baksheev 2011-08-30 09:08:41 +00:00
parent c19f88853a
commit 39373cd9f9
3 changed files with 231 additions and 197 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

121
modules/gpu/src/opencv2/gpu/device/funcattrib.hpp
View File

@ -1,69 +1,78 @@
/* /*M///////////////////////////////////////////////////////////////////////////////////////
* Software License Agreement (BSD License) //
* // IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
* Copyright (c) 2011, Willow Garage, Inc. //
* All rights reserved. // 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,
* Redistribution and use in source and binary forms, with or without // copy or use the software.
* modification, are permitted provided that the following conditions //
* are met: //
* // License Agreement
* * Redistributions of source code must retain the above copyright // For Open Source Computer Vision Library
* notice, this list of conditions and the following disclaimer. //
* * Redistributions in binary form must reproduce the above // Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
* copyright notice, this list of conditions and the following // Copyright (C) 2009, Willow Garage Inc., all rights reserved.
* disclaimer in the documentation and/or other materials provided // Third party copyrights are property of their respective owners.
* with the distribution. //
* * Neither the name of Willow Garage, Inc. nor the names of its // Redistribution and use in source and binary forms, with or without modification,
* contributors may be used to endorse or promote products derived // are permitted provided that the following conditions are met:
* from this software without specific prior written permission. //
* // * Redistribution's of source code must retain the above copyright notice,
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS // this list of conditions and the following disclaimer.
* "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT //
* LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS // * Redistribution's in binary form must reproduce the above copyright notice,
* FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE // this list of conditions and the following disclaimer in the documentation
* COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, // and/or other materials provided with the distribution.
* INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, //
* BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; // * The name of the copyright holders may not be used to endorse or promote products
* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER // derived from this software without specific prior written permission.
* CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT //
* LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN // This software is provided by the copyright holders and contributors "as is" and
* ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE // any express or implied warranties, including, but not limited to, the implied
* POSSIBILITY OF SUCH DAMAGE. // 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,
* Author: Anatoly Baskeheev, Itseez Ltd, (myname.mysurname@mycompany.com) // 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<cstdio> #include<cstdio>
namespace pcl namespace cv
{ {
namespace device namespace gpu
{ {
template<class Func> namespace device
void printFuncAttrib(Func& func)
{ {
template<class Func>
void printFuncAttrib(Func& func)
{
cudaFuncAttributes attrs; cudaFuncAttributes attrs;
cudaFuncGetAttributes(&attrs, func); cudaFuncGetAttributes(&attrs, func);
printf("=== Function stats ===\n"); printf("=== Function stats ===\n");
printf("Name: \n"); printf("Name: \n");
printf("sharedSizeBytes = %d\n", attrs.sharedSizeBytes); printf("sharedSizeBytes = %d\n", attrs.sharedSizeBytes);
printf("constSizeBytes = %d\n", attrs.constSizeBytes); printf("constSizeBytes = %d\n", attrs.constSizeBytes);
printf("localSizeBytes = %d\n", attrs.localSizeBytes); printf("localSizeBytes = %d\n", attrs.localSizeBytes);
printf("maxThreadsPerBlock = %d\n", attrs.maxThreadsPerBlock); printf("maxThreadsPerBlock = %d\n", attrs.maxThreadsPerBlock);
printf("numRegs = %d\n", attrs.numRegs); printf("numRegs = %d\n", attrs.numRegs);
printf("ptxVersion = %d\n", attrs.ptxVersion); printf("ptxVersion = %d\n", attrs.ptxVersion);
printf("binaryVersion = %d\n", attrs.binaryVersion); printf("binaryVersion = %d\n", attrs.binaryVersion);
printf("\n"); printf("\n");
fflush(stdout); fflush(stdout);
}
} }
} }
} }
#endif /* PCL_DEVICE_FUNCATTRIB_HPP_ */ #endif /* __OPENCV_GPU_DEVICE_FUNCATTRIB_HPP_ */

110
modules/gpu/src/opencv2/gpu/device/static_check.hpp
View File

@ -1,66 +1,72 @@
/* /*M///////////////////////////////////////////////////////////////////////////////////////
* Software License Agreement (BSD License) //
* // IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
* Copyright (c) 2011, Willow Garage, Inc. //
* All rights reserved. // 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,
* Redistribution and use in source and binary forms, with or without // copy or use the software.
* modification, are permitted provided that the following conditions //
* are met: //
* // License Agreement
* * Redistributions of source code must retain the above copyright // For Open Source Computer Vision Library
* notice, this list of conditions and the following disclaimer. //
* * Redistributions in binary form must reproduce the above // Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
* copyright notice, this list of conditions and the following // Copyright (C) 2009, Willow Garage Inc., all rights reserved.
* disclaimer in the documentation and/or other materials provided // Third party copyrights are property of their respective owners.
* with the distribution. //
* * Neither the name of Willow Garage, Inc. nor the names of its // Redistribution and use in source and binary forms, with or without modification,
* contributors may be used to endorse or promote products derived // are permitted provided that the following conditions are met:
* from this software without specific prior written permission. //
* // * Redistribution's of source code must retain the above copyright notice,
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS // this list of conditions and the following disclaimer.
* "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT //
* LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS // * Redistribution's in binary form must reproduce the above copyright notice,
* FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE // this list of conditions and the following disclaimer in the documentation
* COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, // and/or other materials provided with the distribution.
* INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, //
* BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; // * The name of the copyright holders may not be used to endorse or promote products
* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER // derived from this software without specific prior written permission.
* CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT //
* LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN // This software is provided by the copyright holders and contributors "as is" and
* ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE // any express or implied warranties, including, but not limited to, the implied
* POSSIBILITY OF SUCH DAMAGE. // 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,
* Author: Anatoly Baskeheev, Itseez Ltd, (myname.mysurname@mycompany.com) // 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_ #ifndef __OPENCV_GPU_GPU_DEVICE_STATIC_CHECK_HPP__
#define PCL_GPU_DEVICE_STATIC_CHECK_HPP_ #define __OPENCV_GPU_GPU_DEVICE_STATIC_CHECK_HPP__
#if defined(__CUDACC__) #if defined(__CUDACC__)
#define __PCL_GPU_HOST_DEVICE__ __host__ __device__ __forceinline__ #define __OPENCV_GPU_HOST_DEVICE__ __host__ __device__ __forceinline__
#else #else
#define __PCL_GPU_HOST_DEVICE__ #define __OPENCV_GPU_HOST_DEVICE__
#endif #endif
namespace pcl namespace cv
{ {
namespace device
{
template<bool expr> struct Static {};
template<> struct Static<true>
{
__PCL_GPU_HOST_DEVICE__ static void check() {};
};
}
namespace gpu namespace gpu
{ {
using pcl::device::Static; namespace device
{
template<bool expr> struct Static {};
template<> struct Static<true>
{
__OPENCV_GPU_HOST_DEVICE__ static void check() {};
};
}
using cv::gpu::device::Static;
} }
} }
#undef __PCL_GPU_HOST_DEVICE__ #undef __PCL_GPU_HOST_DEVICE__
#endif /* PCL_GPU_DEVICE_STATIC_CHECK_HPP_ */ #endif /* __OPENCV_GPU_GPU_DEVICE_STATIC_CHECK_HPP__ */

197
modules/gpu/src/opencv2/gpu/device/warp.hpp
View File

@ -1,99 +1,118 @@
/* /*M///////////////////////////////////////////////////////////////////////////////////////
* Software License Agreement (BSD License) //
* // IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
* Copyright (c) 2011, Willow Garage, Inc. //
* All rights reserved. // 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,
* Redistribution and use in source and binary forms, with or without // copy or use the software.
* modification, are permitted provided that the following conditions //
* are met: //
* // License Agreement
* * Redistributions of source code must retain the above copyright // For Open Source Computer Vision Library
* notice, this list of conditions and the following disclaimer. //
* * Redistributions in binary form must reproduce the above // Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
* copyright notice, this list of conditions and the following // Copyright (C) 2009, Willow Garage Inc., all rights reserved.
* disclaimer in the documentation and/or other materials provided // Third party copyrights are property of their respective owners.
* with the distribution. //
* * Neither the name of Willow Garage, Inc. nor the names of its // Redistribution and use in source and binary forms, with or without modification,
* contributors may be used to endorse or promote products derived // are permitted provided that the following conditions are met:
* from this software without specific prior written permission. //
* // * Redistribution's of source code must retain the above copyright notice,
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS // this list of conditions and the following disclaimer.
* "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT //
* LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS // * Redistribution's in binary form must reproduce the above copyright notice,
* FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE // this list of conditions and the following disclaimer in the documentation
* COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, // and/or other materials provided with the distribution.
* INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, //
* BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; // * The name of the copyright holders may not be used to endorse or promote products
* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER // derived from this software without specific prior written permission.
* CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT //
* LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN // This software is provided by the copyright holders and contributors "as is" and
* ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE // any express or implied warranties, including, but not limited to, the implied
* POSSIBILITY OF SUCH DAMAGE. // 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,
* Author: Anatoly Baskeheev, Itseez Ltd, (myname.mysurname@mycompany.com) // 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_ #ifndef __OPENCV_GPU_DEVICE_WARP_HPP_
#define PCL_DEVICE_UTILS_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, enum
WARP_SIZE = 1 << LOG_WARP_SIZE, {
STRIDE = WARP_SIZE 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;
}
template<typename OutIt, typename T>
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;
}
}; };
}
/** \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_ */ #endif /* __OPENCV_GPU_DEVICE_WARP_HPP_ */