fixed compilation under linux

pull/13383/head
Vladislav Vinogradov 13 years ago
parent 8e061ac801
commit b873fa818f
  1. 4
      modules/gpu/src/cuda/copy_make_border.cu
  2. 4
      modules/gpu/src/cuda/pyr_down.cu
  3. 4
      modules/gpu/src/cuda/pyr_up.cu
  4. 4
      modules/gpu/src/cuda/remap.cu
  5. 4
      modules/gpu/src/cuda/resize.cu
  6. 20
      modules/gpu/src/imgproc.cpp
  7. 68
      modules/gpu/src/opencv2/gpu/device/functional.hpp

@ -45,7 +45,7 @@
BEGIN_OPENCV_DEVICE_NAMESPACE
namespace copy_make_border {
namespace imgproc {
template <typename Ptr2D, typename T> __global__ void copyMakeBorder(const Ptr2D src, DevMem2D_<T> dst, int top, int left)
{
@ -124,6 +124,6 @@ template void copyMakeBorder_gpu<float, 1>(const DevMem2Db& src, const DevMem2Db
template void copyMakeBorder_gpu<float, 3>(const DevMem2Db& src, const DevMem2Db& dst, int top, int left, int borderMode, const float* borderValue, cudaStream_t stream);
template void copyMakeBorder_gpu<float, 4>(const DevMem2Db& src, const DevMem2Db& dst, int top, int left, int borderMode, const float* borderValue, cudaStream_t stream);
} // namespace copy_make_border
} // namespace imgproc
END_OPENCV_DEVICE_NAMESPACE

@ -48,7 +48,7 @@
BEGIN_OPENCV_DEVICE_NAMESPACE
namespace pyr_down {
namespace imgproc {
template <typename T, typename B> __global__ void pyrDown(const PtrStep<T> src, PtrStep<T> dst, const B b, int dst_cols)
{
@ -182,6 +182,6 @@ template void pyrDown_gpu<float, 2>(const DevMem2Db& src, const DevMem2Db& dst,
template void pyrDown_gpu<float, 3>(const DevMem2Db& src, const DevMem2Db& dst, int borderType, cudaStream_t stream);
template void pyrDown_gpu<float, 4>(const DevMem2Db& src, const DevMem2Db& dst, int borderType, cudaStream_t stream);
} // namespace pyr_down
} // namespace imgproc
END_OPENCV_DEVICE_NAMESPACE

@ -48,7 +48,7 @@
BEGIN_OPENCV_DEVICE_NAMESPACE
namespace pyr_up {
namespace imgproc {
template <typename T, typename B> __global__ void pyrUp(const PtrStep<T> src, DevMem2D_<T> dst, const B b)
{
@ -177,6 +177,6 @@ template void pyrUp_gpu<float, 2>(const DevMem2Db& src, const DevMem2Db& dst, in
template void pyrUp_gpu<float, 3>(const DevMem2Db& src, const DevMem2Db& dst, int borderType, cudaStream_t stream);
template void pyrUp_gpu<float, 4>(const DevMem2Db& src, const DevMem2Db& dst, int borderType, cudaStream_t stream);
} // namespace pyr_up
} // namespace imgproc
END_OPENCV_DEVICE_NAMESPACE

@ -49,7 +49,7 @@
BEGIN_OPENCV_DEVICE_NAMESPACE
namespace remap {
namespace imgproc {
template <typename Ptr2D, typename T> __global__ void remap(const Ptr2D src, const PtrStepf mapx, const PtrStepf mapy, DevMem2D_<T> dst)
{
@ -249,6 +249,6 @@ template void remap_gpu<float >(const DevMem2Db& src, const DevMem2Df& xmap, con
template void remap_gpu<float3>(const DevMem2Db& src, const DevMem2Df& xmap, const DevMem2Df& ymap, const DevMem2Db& dst, int interpolation, int borderMode, const float* borderValue, cudaStream_t stream, int cc);
template void remap_gpu<float4>(const DevMem2Db& src, const DevMem2Df& xmap, const DevMem2Df& ymap, const DevMem2Db& dst, int interpolation, int borderMode, const float* borderValue, cudaStream_t stream, int cc);
} // namespace remap
} // namespace imgproc
END_OPENCV_DEVICE_NAMESPACE

@ -49,7 +49,7 @@
BEGIN_OPENCV_DEVICE_NAMESPACE
namespace resize {
namespace imgproc {
template <typename Ptr2D, typename T> __global__ void resize(const Ptr2D src, float fx, float fy, DevMem2D_<T> dst)
{
@ -260,6 +260,6 @@ template void resize_gpu<float >(const DevMem2Db& src, float fx, float fy, const
template void resize_gpu<float3>(const DevMem2Db& src, float fx, float fy, const DevMem2Db& dst, int interpolation, cudaStream_t stream);
template void resize_gpu<float4>(const DevMem2Db& src, float fx, float fy, const DevMem2Db& dst, int interpolation, cudaStream_t stream);
} // namespace resize
} // namespace imgproc
END_OPENCV_DEVICE_NAMESPACE

@ -109,7 +109,7 @@ void cv::gpu::CannyBuf::release() { throw_nogpu(); }
BEGIN_OPENCV_DEVICE_NAMESPACE
namespace remap
namespace imgproc
{
template <typename T>
void remap_gpu(const DevMem2Db& src, const DevMem2Df& xmap, const DevMem2Df& ymap, const DevMem2Db& dst,
@ -120,7 +120,7 @@ END_OPENCV_DEVICE_NAMESPACE
void cv::gpu::remap(const GpuMat& src, GpuMat& dst, const GpuMat& xmap, const GpuMat& ymap, int interpolation, int borderMode, const Scalar& borderValue, Stream& stream)
{
using namespace OPENCV_DEVICE_NAMESPACE_ remap;
using namespace OPENCV_DEVICE_NAMESPACE_ imgproc;
typedef void (*caller_t)(const DevMem2Db& src, const DevMem2Df& xmap, const DevMem2Df& ymap, const DevMem2Db& dst, int interpolation,
int borderMode, const float* borderValue, cudaStream_t stream, int cc);
@ -311,7 +311,7 @@ void cv::gpu::reprojectImageTo3D(const GpuMat& disp, GpuMat& xyzw, const Mat& Q,
BEGIN_OPENCV_DEVICE_NAMESPACE
namespace resize
namespace imgproc
{
template <typename T> void resize_gpu(const DevMem2Db& src, float fx, float fy, const DevMem2Db& dst, int interpolation, cudaStream_t stream);
}
@ -380,7 +380,7 @@ void cv::gpu::resize(const GpuMat& src, GpuMat& dst, Size dsize, double fx, doub
}
else
{
using namespace OPENCV_DEVICE_NAMESPACE_ resize;
using namespace OPENCV_DEVICE_NAMESPACE_ imgproc;
typedef void (*caller_t)(const DevMem2Db& src, float fx, float fy, const DevMem2Db& dst, int interpolation, cudaStream_t stream);
static const caller_t callers[6][4] =
@ -402,7 +402,7 @@ void cv::gpu::resize(const GpuMat& src, GpuMat& dst, Size dsize, double fx, doub
BEGIN_OPENCV_DEVICE_NAMESPACE
namespace copy_make_border
namespace imgproc
{
template <typename T, int cn> void copyMakeBorder_gpu(const DevMem2Db& src, const DevMem2Db& dst, int top, int left, int borderMode, const T* borderValue, cudaStream_t stream);
}
@ -413,7 +413,7 @@ namespace
{
template <typename T, int cn> void copyMakeBorder_caller(const DevMem2Db& src, const DevMem2Db& dst, int top, int left, int borderType, const Scalar& value, cudaStream_t stream)
{
using namespace OPENCV_DEVICE_NAMESPACE_ copy_make_border;
using namespace OPENCV_DEVICE_NAMESPACE_ imgproc;
Scalar_<T> val(saturate_cast<T>(value[0]), saturate_cast<T>(value[1]), saturate_cast<T>(value[2]), saturate_cast<T>(value[3]));
@ -1813,7 +1813,7 @@ void cv::gpu::convolve(const GpuMat& image, const GpuMat& templ, GpuMat& result,
BEGIN_OPENCV_DEVICE_NAMESPACE
namespace pyr_down
namespace imgproc
{
template <typename T, int cn> void pyrDown_gpu(const DevMem2Db& src, const DevMem2Db& dst, int borderType, cudaStream_t stream);
}
@ -1822,7 +1822,7 @@ END_OPENCV_DEVICE_NAMESPACE
void cv::gpu::pyrDown(const GpuMat& src, GpuMat& dst, int borderType, Stream& stream)
{
using namespace OPENCV_DEVICE_NAMESPACE_ pyr_down;
using namespace OPENCV_DEVICE_NAMESPACE_ imgproc;
typedef void (*func_t)(const DevMem2Db& src, const DevMem2Db& dst, int borderType, cudaStream_t stream);
@ -1853,7 +1853,7 @@ void cv::gpu::pyrDown(const GpuMat& src, GpuMat& dst, int borderType, Stream& st
BEGIN_OPENCV_DEVICE_NAMESPACE
namespace pyr_up
namespace imgproc
{
template <typename T, int cn> void pyrUp_gpu(const DevMem2Db& src, const DevMem2Db& dst, int borderType, cudaStream_t stream);
}
@ -1862,7 +1862,7 @@ END_OPENCV_DEVICE_NAMESPACE
void cv::gpu::pyrUp(const GpuMat& src, GpuMat& dst, int borderType, Stream& stream)
{
using namespace OPENCV_DEVICE_NAMESPACE_ pyr_up;
using namespace OPENCV_DEVICE_NAMESPACE_ imgproc;
typedef void (*func_t)(const DevMem2Db& src, const DevMem2Db& dst, int borderType, cudaStream_t stream);

@ -272,61 +272,61 @@ OPENCV_GPU_IMPLEMENT_MINMAX(minimum, double, ::fmin)
// Math functions
#define OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(func) \
template <typename T> struct func ## _func : unary_function<T, float> \
#define OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(name, func) \
template <typename T> struct name ## _func : unary_function<T, float> \
{ \
__device__ __forceinline__ float operator ()(typename TypeTraits<T>::ParameterType v) const \
{ \
return :: ## func ## f(v); \
return func ## f(v); \
} \
}; \
template <> struct func ## _func<double> : unary_function<double, double> \
template <> struct name ## _func<double> : unary_function<double, double> \
{ \
__device__ __forceinline__ double operator ()(double v) const \
{ \
return :: ## func(v); \
return func(v); \
} \
};
#define OPENCV_GPU_IMPLEMENT_BIN_FUNCTOR(func) \
template <typename T> struct func ## _func : binary_function<T, T, float> \
#define OPENCV_GPU_IMPLEMENT_BIN_FUNCTOR(name, func) \
template <typename T> struct name ## _func : binary_function<T, T, float> \
{ \
__device__ __forceinline__ float operator ()(typename TypeTraits<T>::ParameterType v1, typename TypeTraits<T>::ParameterType v2) const \
{ \
return :: ## func ## f(v1, v2); \
return func ## f(v1, v2); \
} \
}; \
template <> struct func ## _func<double> : binary_function<double, double, double> \
template <> struct name ## _func<double> : binary_function<double, double, double> \
{ \
__device__ __forceinline__ double operator ()(double v1, double v2) const \
{ \
return :: ## func(v1, v2); \
return func(v1, v2); \
} \
};
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(fabs)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(sqrt)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(exp)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(exp2)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(exp10)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(log)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(log2)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(log10)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(sin)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(cos)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(tan)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(asin)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(acos)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(atan)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(sinh)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(cosh)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(tanh)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(asinh)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(acosh)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(atanh)
OPENCV_GPU_IMPLEMENT_BIN_FUNCTOR(hypot)
OPENCV_GPU_IMPLEMENT_BIN_FUNCTOR(atan2)
OPENCV_GPU_IMPLEMENT_BIN_FUNCTOR(pow)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(fabs, ::fabs)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(sqrt, ::sqrt)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(exp, ::exp)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(exp2, ::exp2)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(exp10, ::exp10)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(log, ::log)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(log2, ::log2)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(log10, ::log10)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(sin, ::sin)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(cos, ::cos)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(tan, ::tan)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(asin, ::asin)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(acos, ::acos)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(atan, ::atan)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(sinh, ::sinh)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(cosh, ::cosh)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(tanh, ::tanh)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(asinh, ::asinh)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(acosh, ::acosh)
OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(atanh, ::atanh)
OPENCV_GPU_IMPLEMENT_BIN_FUNCTOR(hypot, ::hypot)
OPENCV_GPU_IMPLEMENT_BIN_FUNCTOR(atan2, ::atan2)
OPENCV_GPU_IMPLEMENT_BIN_FUNCTOR(pow, ::pow)
#undef OPENCV_GPU_IMPLEMENT_UN_FUNCTOR
#undef OPENCV_GPU_IMPLEMENT_BIN_FUNCTOR

Loading…
Cancel
Save