mirror of
https://github.com/opencv/opencv.git
synced 2024-12-28 11:58:13 +08:00
220 lines
8.2 KiB
Plaintext
220 lines
8.2 KiB
Plaintext
/*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*/
|
|
|
|
#include "opencv2/opencv_modules.hpp"
|
|
|
|
#ifndef HAVE_OPENCV_CUDEV
|
|
|
|
#error "opencv_cudev is required"
|
|
|
|
#else
|
|
|
|
#include "opencv2/cudev.hpp"
|
|
|
|
using namespace cv::cudev;
|
|
|
|
void cmpMat(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat&, double, Stream& stream, int cmpop);
|
|
|
|
namespace
|
|
{
|
|
template <class Op, typename T> struct CmpOp : binary_function<T, T, uchar>
|
|
{
|
|
__device__ __forceinline__ uchar operator()(T a, T b) const
|
|
{
|
|
Op op;
|
|
return -op(a, b);
|
|
}
|
|
};
|
|
|
|
template <typename ScalarDepth> struct TransformPolicy : DefaultTransformPolicy
|
|
{
|
|
};
|
|
template <> struct TransformPolicy<double> : DefaultTransformPolicy
|
|
{
|
|
enum {
|
|
shift = 1
|
|
};
|
|
};
|
|
|
|
template <template <typename> class Op, typename T>
|
|
void cmpMat_v1(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream)
|
|
{
|
|
CmpOp<Op<T>, T> op;
|
|
gridTransformBinary_< TransformPolicy<T> >(globPtr<T>(src1), globPtr<T>(src2), globPtr<uchar>(dst), op, stream);
|
|
}
|
|
|
|
struct VCmpEq4 : binary_function<uint, uint, uint>
|
|
{
|
|
__device__ __forceinline__ uint operator ()(uint a, uint b) const
|
|
{
|
|
return vcmpeq4(a, b);
|
|
}
|
|
};
|
|
struct VCmpNe4 : binary_function<uint, uint, uint>
|
|
{
|
|
__device__ __forceinline__ uint operator ()(uint a, uint b) const
|
|
{
|
|
return vcmpne4(a, b);
|
|
}
|
|
};
|
|
struct VCmpLt4 : binary_function<uint, uint, uint>
|
|
{
|
|
__device__ __forceinline__ uint operator ()(uint a, uint b) const
|
|
{
|
|
return vcmplt4(a, b);
|
|
}
|
|
};
|
|
struct VCmpLe4 : binary_function<uint, uint, uint>
|
|
{
|
|
__device__ __forceinline__ uint operator ()(uint a, uint b) const
|
|
{
|
|
return vcmple4(a, b);
|
|
}
|
|
};
|
|
|
|
void cmpMatEq_v4(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream)
|
|
{
|
|
const int vcols = src1.cols >> 2;
|
|
|
|
GlobPtrSz<uint> src1_ = globPtr((uint*) src1.data, src1.step, src1.rows, vcols);
|
|
GlobPtrSz<uint> src2_ = globPtr((uint*) src2.data, src2.step, src1.rows, vcols);
|
|
GlobPtrSz<uint> dst_ = globPtr((uint*) dst.data, dst.step, src1.rows, vcols);
|
|
|
|
gridTransformBinary(src1_, src2_, dst_, VCmpEq4(), stream);
|
|
}
|
|
void cmpMatNe_v4(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream)
|
|
{
|
|
const int vcols = src1.cols >> 2;
|
|
|
|
GlobPtrSz<uint> src1_ = globPtr((uint*) src1.data, src1.step, src1.rows, vcols);
|
|
GlobPtrSz<uint> src2_ = globPtr((uint*) src2.data, src2.step, src1.rows, vcols);
|
|
GlobPtrSz<uint> dst_ = globPtr((uint*) dst.data, dst.step, src1.rows, vcols);
|
|
|
|
gridTransformBinary(src1_, src2_, dst_, VCmpNe4(), stream);
|
|
}
|
|
void cmpMatLt_v4(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream)
|
|
{
|
|
const int vcols = src1.cols >> 2;
|
|
|
|
GlobPtrSz<uint> src1_ = globPtr((uint*) src1.data, src1.step, src1.rows, vcols);
|
|
GlobPtrSz<uint> src2_ = globPtr((uint*) src2.data, src2.step, src1.rows, vcols);
|
|
GlobPtrSz<uint> dst_ = globPtr((uint*) dst.data, dst.step, src1.rows, vcols);
|
|
|
|
gridTransformBinary(src1_, src2_, dst_, VCmpLt4(), stream);
|
|
}
|
|
void cmpMatLe_v4(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream)
|
|
{
|
|
const int vcols = src1.cols >> 2;
|
|
|
|
GlobPtrSz<uint> src1_ = globPtr((uint*) src1.data, src1.step, src1.rows, vcols);
|
|
GlobPtrSz<uint> src2_ = globPtr((uint*) src2.data, src2.step, src1.rows, vcols);
|
|
GlobPtrSz<uint> dst_ = globPtr((uint*) dst.data, dst.step, src1.rows, vcols);
|
|
|
|
gridTransformBinary(src1_, src2_, dst_, VCmpLe4(), stream);
|
|
}
|
|
}
|
|
|
|
void cmpMat(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat&, double, Stream& stream, int cmpop)
|
|
{
|
|
typedef void (*func_t)(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream);
|
|
static const func_t funcs[7][4] =
|
|
{
|
|
{cmpMat_v1<equal_to, uchar> , cmpMat_v1<not_equal_to, uchar> , cmpMat_v1<less, uchar> , cmpMat_v1<less_equal, uchar> },
|
|
{cmpMat_v1<equal_to, schar> , cmpMat_v1<not_equal_to, schar> , cmpMat_v1<less, schar> , cmpMat_v1<less_equal, schar> },
|
|
{cmpMat_v1<equal_to, ushort>, cmpMat_v1<not_equal_to, ushort>, cmpMat_v1<less, ushort>, cmpMat_v1<less_equal, ushort>},
|
|
{cmpMat_v1<equal_to, short> , cmpMat_v1<not_equal_to, short> , cmpMat_v1<less, short> , cmpMat_v1<less_equal, short> },
|
|
{cmpMat_v1<equal_to, int> , cmpMat_v1<not_equal_to, int> , cmpMat_v1<less, int> , cmpMat_v1<less_equal, int> },
|
|
{cmpMat_v1<equal_to, float> , cmpMat_v1<not_equal_to, float> , cmpMat_v1<less, float> , cmpMat_v1<less_equal, float> },
|
|
{cmpMat_v1<equal_to, double>, cmpMat_v1<not_equal_to, double>, cmpMat_v1<less, double>, cmpMat_v1<less_equal, double>}
|
|
};
|
|
|
|
typedef void (*func_v4_t)(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream);
|
|
static const func_v4_t funcs_v4[] =
|
|
{
|
|
cmpMatEq_v4, cmpMatNe_v4, cmpMatLt_v4, cmpMatLe_v4
|
|
};
|
|
|
|
const int depth = src1.depth();
|
|
|
|
CV_DbgAssert( depth <= CV_64F );
|
|
|
|
static const int codes[] =
|
|
{
|
|
0, 2, 3, 2, 3, 1
|
|
};
|
|
const GpuMat* psrc1[] =
|
|
{
|
|
&src1, &src2, &src2, &src1, &src1, &src1
|
|
};
|
|
const GpuMat* psrc2[] =
|
|
{
|
|
&src2, &src1, &src1, &src2, &src2, &src2
|
|
};
|
|
|
|
const int code = codes[cmpop];
|
|
|
|
GpuMat src1_ = psrc1[cmpop]->reshape(1);
|
|
GpuMat src2_ = psrc2[cmpop]->reshape(1);
|
|
GpuMat dst_ = dst.reshape(1);
|
|
|
|
if (depth == CV_8U && (src1_.cols & 3) == 0)
|
|
{
|
|
const intptr_t src1ptr = reinterpret_cast<intptr_t>(src1_.data);
|
|
const intptr_t src2ptr = reinterpret_cast<intptr_t>(src2_.data);
|
|
const intptr_t dstptr = reinterpret_cast<intptr_t>(dst_.data);
|
|
|
|
const bool isAllAligned = (src1ptr & 31) == 0 && (src2ptr & 31) == 0 && (dstptr & 31) == 0;
|
|
|
|
if (isAllAligned)
|
|
{
|
|
funcs_v4[code](src1_, src2_, dst_, stream);
|
|
return;
|
|
}
|
|
}
|
|
|
|
const func_t func = funcs[depth][code];
|
|
|
|
func(src1_, src2_, dst_, stream);
|
|
}
|
|
|
|
#endif
|