Chiebot-Mirror
/
opencv
8
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

fixed some bugs in gpu module under Windows:

Browse Source 
* convertTo alignment error
* compare with scalars error
pull/2/head
Vladislav Vinogradov 13 years ago
parent c492dc6a83
commit 67b718f1cb
4 changed files with 116 additions and 49 deletions
Whitespace
Show all changes Ignore whitespace when comparing lines Ignore changes in amount of whitespace Ignore changes in whitespace at EOL
Split View
Diff Options
Show Stats Download Patch File Download Diff File
  1. 12
      modules/core/src/gpumat.cpp
  2. 141
      modules/gpu/src/cuda/element_operations.cu
  3. 10
      modules/gpu/test/test_core.cpp
  4. 2
      modules/gpu/test/test_filters.cpp

12
modules/core/src/gpumat.cpp
Unescape View File

@ -1037,6 +1037,11 @@ namespace
}
};
template <typename T> static inline bool isAligned(const T* ptr, size_t size)
{
return reinterpret_cast<size_t>(ptr) % size == 0;
}
//////////////////////////////////////////////////////////////////////////
// CudaFuncTable
@ -1165,6 +1170,13 @@ namespace
CV_Error(CV_StsUnsupportedFormat, "The device doesn't support double");
}
bool aligned = isAligned(src.data, 16) && isAligned(dst.data, 16);
if (!aligned)
{
cv::gpu::convertTo(src, dst);
return;
}
const func_t func = funcs[src.depth()][dst.depth()][src.channels() - 1];
CV_DbgAssert(func != 0);

141
modules/gpu/src/cuda/element_operations.cu
Unescape View File

@ -1158,11 +1158,8 @@ namespace cv { namespace gpu { namespace device
//////////////////////////////////////////////////////////////////////////////////////
// Compare
#define TYPE_VEC(type, cn) typename TypeVec<type, cn>::vec_type
template <template <typename> class Op, typename T, int cn> struct Compare;
template <template <typename> class Op, typename T>
struct Compare<Op, T, 1>: binary_function<T, T, uchar>
struct Compare: binary_function<T, T, uchar>
{
__device__ __forceinline__ uchar operator()(T src1, T src2) const
{
@ -1170,47 +1167,9 @@ namespace cv { namespace gpu { namespace device
return static_cast<uchar>(static_cast<int>(op(src1, src2)) * 255);
}
};
template <template <typename> class Op, typename T>
struct Compare<Op, T, 2>: binary_function<TYPE_VEC(T, 2), TYPE_VEC(T, 2), TYPE_VEC(uchar, 2)>
{
__device__ __forceinline__ TYPE_VEC(uchar, 2) operator()(const TYPE_VEC(T, 2) & src1, const TYPE_VEC(T, 2) & src2) const
{
Op<T> op;
return VecTraits<TYPE_VEC(uchar, 2)>::make(
static_cast<uchar>(static_cast<int>(op(src1.x, src2.x)) * 255),
static_cast<uchar>(static_cast<int>(op(src1.y, src2.y)) * 255));
}
};
template <template <typename> class Op, typename T>
struct Compare<Op, T, 3>: binary_function<TYPE_VEC(T, 3), TYPE_VEC(T, 3), TYPE_VEC(uchar, 3)>
{
__device__ __forceinline__ TYPE_VEC(uchar, 3) operator()(const TYPE_VEC(T, 3) & src1, const TYPE_VEC(T, 3) & src2) const
{
Op<T> op;
return VecTraits<TYPE_VEC(uchar, 3)>::make(
static_cast<uchar>(static_cast<int>(op(src1.x, src2.x)) * 255),
static_cast<uchar>(static_cast<int>(op(src1.y, src2.y)) * 255),
static_cast<uchar>(static_cast<int>(op(src1.z, src2.z)) * 255));
}
};
template <template <typename> class Op, typename T>
struct Compare<Op, T, 4>: binary_function<TYPE_VEC(T, 4), TYPE_VEC(T, 4), TYPE_VEC(uchar, 4)>
{
__device__ __forceinline__ TYPE_VEC(uchar, 4) operator()(const TYPE_VEC(T, 4) & src1, const TYPE_VEC(T, 4) & src2) const
{
Op<T> op;
return VecTraits<TYPE_VEC(uchar, 4)>::make(
static_cast<uchar>(static_cast<int>(op(src1.x, src2.x)) * 255),
static_cast<uchar>(static_cast<int>(op(src1.y, src2.y)) * 255),
static_cast<uchar>(static_cast<int>(op(src1.z, src2.z)) * 255),
static_cast<uchar>(static_cast<int>(op(src1.w, src2.w)) * 255));
}
};
#undef TYPE_VEC
#define IMPLEMENT_COMPARE_TRANSFORM_FUNCTOR_TRAITS(op, type, block_dim_y, shift) \
template <> struct TransformFunctorTraits< Compare<op, type, 1> > : DefaultTransformFunctorTraits< Compare<op, type, 1> > \
template <> struct TransformFunctorTraits< Compare<op, type> > : DefaultTransformFunctorTraits< Compare<op, type> > \
{ \
enum { smart_block_dim_y = block_dim_y }; \
enum { smart_shift = shift }; \
@ -1233,7 +1192,7 @@ namespace cv { namespace gpu { namespace device
template <template <typename> class Op, typename T> void compare(DevMem2Db src1, DevMem2Db src2, DevMem2Db dst, cudaStream_t stream)
{
Compare<Op, T, 1> op;
Compare<Op, T> op;
cv::gpu::device::transform(static_cast< DevMem2D_<T> >(src1), static_cast< DevMem2D_<T> >(src2), dst, op, WithOutMask(), stream);
}
@ -1286,6 +1245,95 @@ namespace cv { namespace gpu { namespace device
template void compare_le<float >(DevMem2Db src1, DevMem2Db src2, DevMem2Db dst, cudaStream_t stream);
template void compare_le<double>(DevMem2Db src1, DevMem2Db src2, DevMem2Db dst, cudaStream_t stream);
#define TYPE_VEC(type, cn) typename TypeVec<type, cn>::vec_type
template <template <typename> class Op, typename T, int cn> struct CompareScalar;
template <template <typename> class Op, typename T>
struct CompareScalar<Op, T, 1>: unary_function<T, uchar>
{
const T val;
__host__ explicit CompareScalar(T val) : val(val) {}
__device__ __forceinline__ uchar operator()(T src) const
{
Op<T> op;
return static_cast<uchar>(static_cast<int>(op(src, val)) * 255);
}
};
template <template <typename> class Op, typename T>
struct CompareScalar<Op, T, 2>: unary_function<TYPE_VEC(T, 2), TYPE_VEC(uchar, 2)>
{
const TYPE_VEC(T, 2) val;
__host__ explicit CompareScalar(TYPE_VEC(T, 2) val) : val(val) {}
__device__ __forceinline__ TYPE_VEC(uchar, 2) operator()(const TYPE_VEC(T, 2) & src) const
{
Op<T> op;
return VecTraits<TYPE_VEC(uchar, 2)>::make(
static_cast<uchar>(static_cast<int>(op(src.x, val.x)) * 255),
static_cast<uchar>(static_cast<int>(op(src.y, val.y)) * 255));
}
};
template <template <typename> class Op, typename T>
struct CompareScalar<Op, T, 3>: unary_function<TYPE_VEC(T, 3), TYPE_VEC(uchar, 3)>
{
const TYPE_VEC(T, 3) val;
__host__ explicit CompareScalar(TYPE_VEC(T, 3) val) : val(val) {}
__device__ __forceinline__ TYPE_VEC(uchar, 3) operator()(const TYPE_VEC(T, 3) & src) const
{
Op<T> op;
return VecTraits<TYPE_VEC(uchar, 3)>::make(
static_cast<uchar>(static_cast<int>(op(src.x, val.x)) * 255),
static_cast<uchar>(static_cast<int>(op(src.y, val.y)) * 255),
static_cast<uchar>(static_cast<int>(op(src.z, val.z)) * 255));
}
};
template <template <typename> class Op, typename T>
struct CompareScalar<Op, T, 4>: unary_function<TYPE_VEC(T, 4), TYPE_VEC(uchar, 4)>
{
const TYPE_VEC(T, 4) val;
__host__ explicit CompareScalar(TYPE_VEC(T, 4) val) : val(val) {}
__device__ __forceinline__ TYPE_VEC(uchar, 4) operator()(const TYPE_VEC(T, 4) & src) const
{
Op<T> op;
return VecTraits<TYPE_VEC(uchar, 4)>::make(
static_cast<uchar>(static_cast<int>(op(src.x, val.x)) * 255),
static_cast<uchar>(static_cast<int>(op(src.y, val.y)) * 255),
static_cast<uchar>(static_cast<int>(op(src.z, val.z)) * 255),
static_cast<uchar>(static_cast<int>(op(src.w, val.w)) * 255));
}
};
#undef TYPE_VEC
#define IMPLEMENT_COMPARE_TRANSFORM_FUNCTOR_TRAITS(op, type, block_dim_y, shift) \
template <> struct TransformFunctorTraits< CompareScalar<op, type, 1> > : DefaultTransformFunctorTraits< CompareScalar<op, type, 1> > \
{ \
enum { smart_block_dim_y = block_dim_y }; \
enum { smart_shift = shift }; \
};
IMPLEMENT_COMPARE_TRANSFORM_FUNCTOR_TRAITS(equal_to, int, 8, 4)
IMPLEMENT_COMPARE_TRANSFORM_FUNCTOR_TRAITS(equal_to, float, 8, 4)
IMPLEMENT_COMPARE_TRANSFORM_FUNCTOR_TRAITS(not_equal_to, int, 8, 4)
IMPLEMENT_COMPARE_TRANSFORM_FUNCTOR_TRAITS(not_equal_to, float, 8, 4)
IMPLEMENT_COMPARE_TRANSFORM_FUNCTOR_TRAITS(greater, int, 8, 4)
IMPLEMENT_COMPARE_TRANSFORM_FUNCTOR_TRAITS(greater, float, 8, 4)
IMPLEMENT_COMPARE_TRANSFORM_FUNCTOR_TRAITS(less, int, 8, 4)
IMPLEMENT_COMPARE_TRANSFORM_FUNCTOR_TRAITS(less, float, 8, 4)
IMPLEMENT_COMPARE_TRANSFORM_FUNCTOR_TRAITS(greater_equal, int, 8, 4)
IMPLEMENT_COMPARE_TRANSFORM_FUNCTOR_TRAITS(greater_equal, float, 8, 4)
IMPLEMENT_COMPARE_TRANSFORM_FUNCTOR_TRAITS(less_equal, int, 8, 4)
IMPLEMENT_COMPARE_TRANSFORM_FUNCTOR_TRAITS(less_equal, float, 8, 4)
#undef IMPLEMENT_COMPARE_TRANSFORM_FUNCTOR_TRAITS
template <template <typename> class Op, typename T, int cn> void compare(DevMem2Db src, double val[4], DevMem2Db dst, cudaStream_t stream)
{
typedef typename TypeVec<T, cn>::vec_type src_t;
@ -1294,9 +1342,9 @@ namespace cv { namespace gpu { namespace device
T sval[] = {static_cast<T>(val[0]), static_cast<T>(val[1]), static_cast<T>(val[2]), static_cast<T>(val[3])};
src_t val1 = VecTraits<src_t>::make(sval);
Compare<Op, T, cn> op;
CompareScalar<Op, T, cn> op(val1);
cv::gpu::device::transform(static_cast< DevMem2D_<src_t> >(src), static_cast< DevMem2D_<dst_t> >(dst), cv::gpu::device::bind2nd(op, val1), WithOutMask(), stream);
cv::gpu::device::transform(static_cast< DevMem2D_<src_t> >(src), static_cast< DevMem2D_<dst_t> >(dst), op, WithOutMask(), stream);
}
template <typename T> void compare_eq(DevMem2Db src, int cn, double val[4], DevMem2Db dst, cudaStream_t stream)
@ -1432,7 +1480,6 @@ namespace cv { namespace gpu { namespace device
template void compare_ge<float >(DevMem2Db src, int cn, double val[4], DevMem2Db dst, cudaStream_t stream);
template void compare_ge<double>(DevMem2Db src, int cn, double val[4], DevMem2Db dst, cudaStream_t stream);
//////////////////////////////////////////////////////////////////////////
// Unary bitwise logical matrix operations

10
modules/gpu/test/test_core.cpp
Unescape View File

@ -850,7 +850,7 @@ TEST_P(Divide_Array, WithScale)
cv::Mat dst_gold;
cv::divide(mat1, mat2, dst_gold, scale, depth.second);
EXPECT_MAT_NEAR(dst_gold, dst, depth.first >= CV_32F || depth.second >= CV_32F ? 1e-4 : 1.0);
EXPECT_MAT_NEAR(dst_gold, dst, depth.first >= CV_32F || depth.second >= CV_32F ? 1e-2 : 1.0);
}
}
@ -1612,6 +1612,14 @@ TEST_P(Compare_Scalar, Accuracy)
cv::Mat src = randomMat(size, type);
cv::Scalar sc = randomScalar(0.0, 255.0);
if (src.depth() < CV_32F)
{
sc.val[0] = cvRound(sc.val[0]);
sc.val[1] = cvRound(sc.val[1]);
sc.val[2] = cvRound(sc.val[2]);
sc.val[3] = cvRound(sc.val[3]);
}
if (src.depth() == CV_64F && !supportFeature(devInfo, cv::gpu::NATIVE_DOUBLE))
{
try

2
modules/gpu/test/test_filters.cpp
Unescape View File

@ -334,7 +334,7 @@ TEST_P(Laplacian, Accuracy)
cv::Mat dst_gold;
cv::Laplacian(src, dst_gold, -1, ksize.width);
EXPECT_MAT_NEAR(dst_gold, dst, 0.0);
EXPECT_MAT_NEAR(dst_gold, dst, src.depth() < CV_32F ? 0.0 : 1e-3);
}
INSTANTIATE_TEST_CASE_P(GPU_Filter, Laplacian, testing::Combine(

Write Preview
Loading…
Cancel
Save