Merge remote-tracking branch 'upstream/2.4' into 2.4
This commit is contained in:
commit
34c7162104
@ -85,10 +85,10 @@ namespace
|
|||||||
};
|
};
|
||||||
size_t colors_mum = sizeof(colors)/sizeof(colors[0]);
|
size_t colors_mum = sizeof(colors)/sizeof(colors[0]);
|
||||||
|
|
||||||
#if (defined __cplusplus && __cplusplus > 199711L) || defined _STLPORT_MAJOR
|
template<class FwIt, class T> inline void _iota(FwIt first, FwIt last, T value)
|
||||||
#else
|
{
|
||||||
template<class FwIt, class T> void iota(FwIt first, FwIt last, T value) { while(first != last) *first++ = value++; }
|
while(first != last) *first++ = value++;
|
||||||
#endif
|
}
|
||||||
|
|
||||||
void computeNormals( const Octree& Octree, const vector<Point3f>& centers, vector<Point3f>& normals,
|
void computeNormals( const Octree& Octree, const vector<Point3f>& centers, vector<Point3f>& normals,
|
||||||
vector<uchar>& mask, float normalRadius, int minNeighbors = 20)
|
vector<uchar>& mask, float normalRadius, int minNeighbors = 20)
|
||||||
@ -799,14 +799,14 @@ void cv::SpinImageModel::selectRandomSubset(float ratio)
|
|||||||
else if (setSize == vtxSize)
|
else if (setSize == vtxSize)
|
||||||
{
|
{
|
||||||
subset.resize(vtxSize);
|
subset.resize(vtxSize);
|
||||||
iota(subset.begin(), subset.end(), 0);
|
_iota(subset.begin(), subset.end(), 0);
|
||||||
}
|
}
|
||||||
else
|
else
|
||||||
{
|
{
|
||||||
RNG& rnd = theRNG();
|
RNG& rnd = theRNG();
|
||||||
|
|
||||||
vector<size_t> left(vtxSize);
|
vector<size_t> left(vtxSize);
|
||||||
iota(left.begin(), left.end(), (size_t)0);
|
_iota(left.begin(), left.end(), (size_t)0);
|
||||||
|
|
||||||
subset.resize(setSize);
|
subset.resize(setSize);
|
||||||
for(size_t i = 0; i < setSize; ++i)
|
for(size_t i = 0; i < setSize; ++i)
|
||||||
@ -879,7 +879,7 @@ void cv::SpinImageModel::compute()
|
|||||||
{
|
{
|
||||||
mesh.computeNormals(normalRadius, minNeighbors);
|
mesh.computeNormals(normalRadius, minNeighbors);
|
||||||
subset.resize(mesh.vtx.size());
|
subset.resize(mesh.vtx.size());
|
||||||
iota(subset.begin(), subset.end(), 0);
|
_iota(subset.begin(), subset.end(), 0);
|
||||||
}
|
}
|
||||||
else
|
else
|
||||||
mesh.computeNormals(subset, normalRadius, minNeighbors);
|
mesh.computeNormals(subset, normalRadius, minNeighbors);
|
||||||
|
@ -1341,7 +1341,12 @@ PERF_TEST_P(Sz_Depth_Code, ImgProc_CvtColorBayer,
|
|||||||
Values(CvtColorInfo(1, 3, cv::COLOR_BayerBG2BGR),
|
Values(CvtColorInfo(1, 3, cv::COLOR_BayerBG2BGR),
|
||||||
CvtColorInfo(1, 3, cv::COLOR_BayerGB2BGR),
|
CvtColorInfo(1, 3, cv::COLOR_BayerGB2BGR),
|
||||||
CvtColorInfo(1, 3, cv::COLOR_BayerRG2BGR),
|
CvtColorInfo(1, 3, cv::COLOR_BayerRG2BGR),
|
||||||
CvtColorInfo(1, 3, cv::COLOR_BayerGR2BGR))))
|
CvtColorInfo(1, 3, cv::COLOR_BayerGR2BGR),
|
||||||
|
|
||||||
|
CvtColorInfo(1, 1, cv::COLOR_BayerBG2GRAY),
|
||||||
|
CvtColorInfo(1, 1, cv::COLOR_BayerGB2GRAY),
|
||||||
|
CvtColorInfo(1, 1, cv::COLOR_BayerRG2GRAY),
|
||||||
|
CvtColorInfo(1, 1, cv::COLOR_BayerGR2GRAY))))
|
||||||
{
|
{
|
||||||
const cv::Size size = GET_PARAM(0);
|
const cv::Size size = GET_PARAM(0);
|
||||||
const int depth = GET_PARAM(1);
|
const int depth = GET_PARAM(1);
|
||||||
|
@ -1640,6 +1640,43 @@ namespace
|
|||||||
{
|
{
|
||||||
bayer_to_bgr(src, dst, dcn, true, true, stream);
|
bayer_to_bgr(src, dst, dcn, true, true, stream);
|
||||||
}
|
}
|
||||||
|
void bayer_to_gray(const GpuMat& src, GpuMat& dst, bool blue_last, bool start_with_green, Stream& stream)
|
||||||
|
{
|
||||||
|
typedef void (*func_t)(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream);
|
||||||
|
static const func_t funcs[3] =
|
||||||
|
{
|
||||||
|
Bayer2BGR_8u_gpu<1>,
|
||||||
|
0,
|
||||||
|
Bayer2BGR_16u_gpu<1>,
|
||||||
|
};
|
||||||
|
|
||||||
|
CV_Assert(src.type() == CV_8UC1 || src.type() == CV_16UC1);
|
||||||
|
CV_Assert(src.rows > 2 && src.cols > 2);
|
||||||
|
|
||||||
|
dst.create(src.size(), CV_MAKETYPE(src.depth(), 1));
|
||||||
|
|
||||||
|
funcs[src.depth()](src, dst, blue_last, start_with_green, StreamAccessor::getStream(stream));
|
||||||
|
}
|
||||||
|
|
||||||
|
void bayerBG_to_gray(const GpuMat& src, GpuMat& dst, int /*dcn*/, Stream& stream)
|
||||||
|
{
|
||||||
|
bayer_to_gray(src, dst, false, false, stream);
|
||||||
|
}
|
||||||
|
|
||||||
|
void bayerGB_to_gray(const GpuMat& src, GpuMat& dst, int /*dcn*/, Stream& stream)
|
||||||
|
{
|
||||||
|
bayer_to_gray(src, dst, false, true, stream);
|
||||||
|
}
|
||||||
|
|
||||||
|
void bayerRG_to_gray(const GpuMat& src, GpuMat& dst, int /*dcn*/, Stream& stream)
|
||||||
|
{
|
||||||
|
bayer_to_gray(src, dst, true, false, stream);
|
||||||
|
}
|
||||||
|
|
||||||
|
void bayerGR_to_gray(const GpuMat& src, GpuMat& dst, int /*dcn*/, Stream& stream)
|
||||||
|
{
|
||||||
|
bayer_to_gray(src, dst, true, true, stream);
|
||||||
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
void cv::gpu::cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn, Stream& stream)
|
void cv::gpu::cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn, Stream& stream)
|
||||||
@ -1756,10 +1793,10 @@ void cv::gpu::cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn, Stream
|
|||||||
yuv_to_bgr, // CV_YUV2BGR = 84
|
yuv_to_bgr, // CV_YUV2BGR = 84
|
||||||
yuv_to_rgb, // CV_YUV2RGB = 85
|
yuv_to_rgb, // CV_YUV2RGB = 85
|
||||||
|
|
||||||
0, // CV_BayerBG2GRAY = 86
|
bayerBG_to_gray, // CV_BayerBG2GRAY = 86
|
||||||
0, // CV_BayerGB2GRAY = 87
|
bayerGB_to_gray, // CV_BayerGB2GRAY = 87
|
||||||
0, // CV_BayerRG2GRAY = 88
|
bayerRG_to_gray, // CV_BayerRG2GRAY = 88
|
||||||
0, // CV_BayerGR2GRAY = 89
|
bayerGR_to_gray, // CV_BayerGR2GRAY = 89
|
||||||
|
|
||||||
//YUV 4:2:0 formats family
|
//YUV 4:2:0 formats family
|
||||||
0, // CV_YUV2RGB_NV12 = 90,
|
0, // CV_YUV2RGB_NV12 = 90,
|
||||||
|
@ -42,42 +42,37 @@
|
|||||||
|
|
||||||
#if !defined CUDA_DISABLER
|
#if !defined CUDA_DISABLER
|
||||||
|
|
||||||
#include <opencv2/gpu/device/common.hpp>
|
#include "opencv2/gpu/device/common.hpp"
|
||||||
#include <opencv2/gpu/device/vec_traits.hpp>
|
#include "opencv2/gpu/device/vec_traits.hpp"
|
||||||
#include <opencv2/gpu/device/vec_math.hpp>
|
#include "opencv2/gpu/device/vec_math.hpp"
|
||||||
#include <opencv2/gpu/device/limits.hpp>
|
#include "opencv2/gpu/device/limits.hpp"
|
||||||
|
#include "opencv2/gpu/device/color.hpp"
|
||||||
|
|
||||||
namespace cv { namespace gpu {
|
namespace cv { namespace gpu { namespace device
|
||||||
namespace device
|
{
|
||||||
|
template <typename T> struct Bayer2BGR;
|
||||||
|
|
||||||
|
template <> struct Bayer2BGR<uchar>
|
||||||
{
|
{
|
||||||
template <typename D>
|
uchar3 res0;
|
||||||
__global__ void Bayer2BGR_8u(const PtrStepb src, PtrStepSz<D> dst, const bool blue_last, const bool start_with_green)
|
uchar3 res1;
|
||||||
|
uchar3 res2;
|
||||||
|
uchar3 res3;
|
||||||
|
|
||||||
|
__device__ void apply(const PtrStepSzb& src, int s_x, int s_y, bool blue_last, bool start_with_green)
|
||||||
{
|
{
|
||||||
const int s_x = blockIdx.x * blockDim.x + threadIdx.x;
|
|
||||||
int s_y = blockIdx.y * blockDim.y + threadIdx.y;
|
|
||||||
|
|
||||||
if (s_y >= dst.rows || (s_x << 2) >= dst.cols)
|
|
||||||
return;
|
|
||||||
|
|
||||||
s_y = ::min(::max(s_y, 1), dst.rows - 2);
|
|
||||||
|
|
||||||
uchar4 patch[3][3];
|
uchar4 patch[3][3];
|
||||||
patch[0][1] = ((const uchar4*) src.ptr(s_y - 1))[s_x];
|
patch[0][1] = ((const uchar4*) src.ptr(s_y - 1))[s_x];
|
||||||
patch[0][0] = ((const uchar4*) src.ptr(s_y - 1))[::max(s_x - 1, 0)];
|
patch[0][0] = ((const uchar4*) src.ptr(s_y - 1))[::max(s_x - 1, 0)];
|
||||||
patch[0][2] = ((const uchar4*) src.ptr(s_y - 1))[::min(s_x + 1, ((dst.cols + 3) >> 2) - 1)];
|
patch[0][2] = ((const uchar4*) src.ptr(s_y - 1))[::min(s_x + 1, ((src.cols + 3) >> 2) - 1)];
|
||||||
|
|
||||||
patch[1][1] = ((const uchar4*) src.ptr(s_y))[s_x];
|
patch[1][1] = ((const uchar4*) src.ptr(s_y))[s_x];
|
||||||
patch[1][0] = ((const uchar4*) src.ptr(s_y))[::max(s_x - 1, 0)];
|
patch[1][0] = ((const uchar4*) src.ptr(s_y))[::max(s_x - 1, 0)];
|
||||||
patch[1][2] = ((const uchar4*) src.ptr(s_y))[::min(s_x + 1, ((dst.cols + 3) >> 2) - 1)];
|
patch[1][2] = ((const uchar4*) src.ptr(s_y))[::min(s_x + 1, ((src.cols + 3) >> 2) - 1)];
|
||||||
|
|
||||||
patch[2][1] = ((const uchar4*) src.ptr(s_y + 1))[s_x];
|
patch[2][1] = ((const uchar4*) src.ptr(s_y + 1))[s_x];
|
||||||
patch[2][0] = ((const uchar4*) src.ptr(s_y + 1))[::max(s_x - 1, 0)];
|
patch[2][0] = ((const uchar4*) src.ptr(s_y + 1))[::max(s_x - 1, 0)];
|
||||||
patch[2][2] = ((const uchar4*) src.ptr(s_y + 1))[::min(s_x + 1, ((dst.cols + 3) >> 2) - 1)];
|
patch[2][2] = ((const uchar4*) src.ptr(s_y + 1))[::min(s_x + 1, ((src.cols + 3) >> 2) - 1)];
|
||||||
|
|
||||||
D res0 = VecTraits<D>::all(numeric_limits<uchar>::max());
|
|
||||||
D res1 = VecTraits<D>::all(numeric_limits<uchar>::max());
|
|
||||||
D res2 = VecTraits<D>::all(numeric_limits<uchar>::max());
|
|
||||||
D res3 = VecTraits<D>::all(numeric_limits<uchar>::max());
|
|
||||||
|
|
||||||
if ((s_y & 1) ^ start_with_green)
|
if ((s_y & 1) ^ start_with_green)
|
||||||
{
|
{
|
||||||
@ -181,45 +176,69 @@ namespace cv { namespace gpu {
|
|||||||
res3.z = t7;
|
res3.z = t7;
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
const int d_x = (blockIdx.x * blockDim.x + threadIdx.x) << 2;
|
|
||||||
const int d_y = blockIdx.y * blockDim.y + threadIdx.y;
|
|
||||||
|
|
||||||
dst(d_y, d_x) = res0;
|
|
||||||
if (d_x + 1 < dst.cols)
|
|
||||||
dst(d_y, d_x + 1) = res1;
|
|
||||||
if (d_x + 2 < dst.cols)
|
|
||||||
dst(d_y, d_x + 2) = res2;
|
|
||||||
if (d_x + 3 < dst.cols)
|
|
||||||
dst(d_y, d_x + 3) = res3;
|
|
||||||
}
|
}
|
||||||
|
};
|
||||||
|
|
||||||
template <typename D>
|
template <typename D> __device__ __forceinline__ D toDst(const uchar3& pix);
|
||||||
__global__ void Bayer2BGR_16u(const PtrStepb src, PtrStepSz<D> dst, const bool blue_last, const bool start_with_green)
|
template <> __device__ __forceinline__ uchar toDst<uchar>(const uchar3& pix)
|
||||||
|
{
|
||||||
|
typename bgr_to_gray_traits<uchar>::functor_type f = bgr_to_gray_traits<uchar>::create_functor();
|
||||||
|
return f(pix);
|
||||||
|
}
|
||||||
|
template <> __device__ __forceinline__ uchar3 toDst<uchar3>(const uchar3& pix)
|
||||||
|
{
|
||||||
|
return pix;
|
||||||
|
}
|
||||||
|
template <> __device__ __forceinline__ uchar4 toDst<uchar4>(const uchar3& pix)
|
||||||
|
{
|
||||||
|
return make_uchar4(pix.x, pix.y, pix.z, 255);
|
||||||
|
}
|
||||||
|
|
||||||
|
template <typename D>
|
||||||
|
__global__ void Bayer2BGR_8u(const PtrStepSzb src, PtrStep<D> dst, const bool blue_last, const bool start_with_green)
|
||||||
|
{
|
||||||
|
const int s_x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||||
|
int s_y = blockIdx.y * blockDim.y + threadIdx.y;
|
||||||
|
|
||||||
|
if (s_y >= src.rows || (s_x << 2) >= src.cols)
|
||||||
|
return;
|
||||||
|
|
||||||
|
s_y = ::min(::max(s_y, 1), src.rows - 2);
|
||||||
|
|
||||||
|
Bayer2BGR<uchar> bayer;
|
||||||
|
bayer.apply(src, s_x, s_y, blue_last, start_with_green);
|
||||||
|
|
||||||
|
const int d_x = (blockIdx.x * blockDim.x + threadIdx.x) << 2;
|
||||||
|
const int d_y = blockIdx.y * blockDim.y + threadIdx.y;
|
||||||
|
|
||||||
|
dst(d_y, d_x) = toDst<D>(bayer.res0);
|
||||||
|
if (d_x + 1 < src.cols)
|
||||||
|
dst(d_y, d_x + 1) = toDst<D>(bayer.res1);
|
||||||
|
if (d_x + 2 < src.cols)
|
||||||
|
dst(d_y, d_x + 2) = toDst<D>(bayer.res2);
|
||||||
|
if (d_x + 3 < src.cols)
|
||||||
|
dst(d_y, d_x + 3) = toDst<D>(bayer.res3);
|
||||||
|
}
|
||||||
|
|
||||||
|
template <> struct Bayer2BGR<ushort>
|
||||||
|
{
|
||||||
|
ushort3 res0;
|
||||||
|
ushort3 res1;
|
||||||
|
|
||||||
|
__device__ void apply(const PtrStepSzb& src, int s_x, int s_y, bool blue_last, bool start_with_green)
|
||||||
{
|
{
|
||||||
const int s_x = blockIdx.x * blockDim.x + threadIdx.x;
|
|
||||||
int s_y = blockIdx.y * blockDim.y + threadIdx.y;
|
|
||||||
|
|
||||||
if (s_y >= dst.rows || (s_x << 1) >= dst.cols)
|
|
||||||
return;
|
|
||||||
|
|
||||||
s_y = ::min(::max(s_y, 1), dst.rows - 2);
|
|
||||||
|
|
||||||
ushort2 patch[3][3];
|
ushort2 patch[3][3];
|
||||||
patch[0][1] = ((const ushort2*) src.ptr(s_y - 1))[s_x];
|
patch[0][1] = ((const ushort2*) src.ptr(s_y - 1))[s_x];
|
||||||
patch[0][0] = ((const ushort2*) src.ptr(s_y - 1))[::max(s_x - 1, 0)];
|
patch[0][0] = ((const ushort2*) src.ptr(s_y - 1))[::max(s_x - 1, 0)];
|
||||||
patch[0][2] = ((const ushort2*) src.ptr(s_y - 1))[::min(s_x + 1, ((dst.cols + 1) >> 1) - 1)];
|
patch[0][2] = ((const ushort2*) src.ptr(s_y - 1))[::min(s_x + 1, ((src.cols + 1) >> 1) - 1)];
|
||||||
|
|
||||||
patch[1][1] = ((const ushort2*) src.ptr(s_y))[s_x];
|
patch[1][1] = ((const ushort2*) src.ptr(s_y))[s_x];
|
||||||
patch[1][0] = ((const ushort2*) src.ptr(s_y))[::max(s_x - 1, 0)];
|
patch[1][0] = ((const ushort2*) src.ptr(s_y))[::max(s_x - 1, 0)];
|
||||||
patch[1][2] = ((const ushort2*) src.ptr(s_y))[::min(s_x + 1, ((dst.cols + 1) >> 1) - 1)];
|
patch[1][2] = ((const ushort2*) src.ptr(s_y))[::min(s_x + 1, ((src.cols + 1) >> 1) - 1)];
|
||||||
|
|
||||||
patch[2][1] = ((const ushort2*) src.ptr(s_y + 1))[s_x];
|
patch[2][1] = ((const ushort2*) src.ptr(s_y + 1))[s_x];
|
||||||
patch[2][0] = ((const ushort2*) src.ptr(s_y + 1))[::max(s_x - 1, 0)];
|
patch[2][0] = ((const ushort2*) src.ptr(s_y + 1))[::max(s_x - 1, 0)];
|
||||||
patch[2][2] = ((const ushort2*) src.ptr(s_y + 1))[::min(s_x + 1, ((dst.cols + 1) >> 1) - 1)];
|
patch[2][2] = ((const ushort2*) src.ptr(s_y + 1))[::min(s_x + 1, ((src.cols + 1) >> 1) - 1)];
|
||||||
|
|
||||||
D res0 = VecTraits<D>::all(numeric_limits<ushort>::max());
|
|
||||||
D res1 = VecTraits<D>::all(numeric_limits<ushort>::max());
|
|
||||||
|
|
||||||
if ((s_y & 1) ^ start_with_green)
|
if ((s_y & 1) ^ start_with_green)
|
||||||
{
|
{
|
||||||
@ -279,53 +298,87 @@ namespace cv { namespace gpu {
|
|||||||
res1.z = t3;
|
res1.z = t3;
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
const int d_x = (blockIdx.x * blockDim.x + threadIdx.x) << 1;
|
|
||||||
const int d_y = blockIdx.y * blockDim.y + threadIdx.y;
|
|
||||||
|
|
||||||
dst(d_y, d_x) = res0;
|
|
||||||
if (d_x + 1 < dst.cols)
|
|
||||||
dst(d_y, d_x + 1) = res1;
|
|
||||||
}
|
}
|
||||||
|
};
|
||||||
|
|
||||||
template <int cn>
|
template <typename D> __device__ __forceinline__ D toDst(const ushort3& pix);
|
||||||
void Bayer2BGR_8u_gpu(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream)
|
template <> __device__ __forceinline__ ushort toDst<ushort>(const ushort3& pix)
|
||||||
{
|
{
|
||||||
typedef typename TypeVec<uchar, cn>::vec_type dst_t;
|
typename bgr_to_gray_traits<ushort>::functor_type f = bgr_to_gray_traits<ushort>::create_functor();
|
||||||
|
return f(pix);
|
||||||
const dim3 block(32, 8);
|
|
||||||
const dim3 grid(divUp(dst.cols, 4 * block.x), divUp(dst.rows, block.y));
|
|
||||||
|
|
||||||
cudaSafeCall( cudaFuncSetCacheConfig(Bayer2BGR_8u<dst_t>, cudaFuncCachePreferL1) );
|
|
||||||
|
|
||||||
Bayer2BGR_8u<dst_t><<<grid, block, 0, stream>>>(src, (PtrStepSz<dst_t>)dst, blue_last, start_with_green);
|
|
||||||
cudaSafeCall( cudaGetLastError() );
|
|
||||||
|
|
||||||
if (stream == 0)
|
|
||||||
cudaSafeCall( cudaDeviceSynchronize() );
|
|
||||||
}
|
|
||||||
template <int cn>
|
|
||||||
void Bayer2BGR_16u_gpu(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream)
|
|
||||||
{
|
|
||||||
typedef typename TypeVec<ushort, cn>::vec_type dst_t;
|
|
||||||
|
|
||||||
const dim3 block(32, 8);
|
|
||||||
const dim3 grid(divUp(dst.cols, 2 * block.x), divUp(dst.rows, block.y));
|
|
||||||
|
|
||||||
cudaSafeCall( cudaFuncSetCacheConfig(Bayer2BGR_16u<dst_t>, cudaFuncCachePreferL1) );
|
|
||||||
|
|
||||||
Bayer2BGR_16u<dst_t><<<grid, block, 0, stream>>>(src, (PtrStepSz<dst_t>)dst, blue_last, start_with_green);
|
|
||||||
cudaSafeCall( cudaGetLastError() );
|
|
||||||
|
|
||||||
if (stream == 0)
|
|
||||||
cudaSafeCall( cudaDeviceSynchronize() );
|
|
||||||
}
|
|
||||||
|
|
||||||
template void Bayer2BGR_8u_gpu<3>(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream);
|
|
||||||
template void Bayer2BGR_8u_gpu<4>(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream);
|
|
||||||
template void Bayer2BGR_16u_gpu<3>(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream);
|
|
||||||
template void Bayer2BGR_16u_gpu<4>(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream);
|
|
||||||
}
|
}
|
||||||
}}
|
template <> __device__ __forceinline__ ushort3 toDst<ushort3>(const ushort3& pix)
|
||||||
|
{
|
||||||
|
return pix;
|
||||||
|
}
|
||||||
|
template <> __device__ __forceinline__ ushort4 toDst<ushort4>(const ushort3& pix)
|
||||||
|
{
|
||||||
|
return make_ushort4(pix.x, pix.y, pix.z, numeric_limits<ushort>::max());
|
||||||
|
}
|
||||||
|
|
||||||
|
template <typename D>
|
||||||
|
__global__ void Bayer2BGR_16u(const PtrStepSzb src, PtrStep<D> dst, const bool blue_last, const bool start_with_green)
|
||||||
|
{
|
||||||
|
const int s_x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||||
|
int s_y = blockIdx.y * blockDim.y + threadIdx.y;
|
||||||
|
|
||||||
|
if (s_y >= src.rows || (s_x << 1) >= src.cols)
|
||||||
|
return;
|
||||||
|
|
||||||
|
s_y = ::min(::max(s_y, 1), src.rows - 2);
|
||||||
|
|
||||||
|
Bayer2BGR<ushort> bayer;
|
||||||
|
bayer.apply(src, s_x, s_y, blue_last, start_with_green);
|
||||||
|
|
||||||
|
const int d_x = (blockIdx.x * blockDim.x + threadIdx.x) << 1;
|
||||||
|
const int d_y = blockIdx.y * blockDim.y + threadIdx.y;
|
||||||
|
|
||||||
|
dst(d_y, d_x) = toDst<D>(bayer.res0);
|
||||||
|
if (d_x + 1 < src.cols)
|
||||||
|
dst(d_y, d_x + 1) = toDst<D>(bayer.res1);
|
||||||
|
}
|
||||||
|
|
||||||
|
template <int cn>
|
||||||
|
void Bayer2BGR_8u_gpu(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream)
|
||||||
|
{
|
||||||
|
typedef typename TypeVec<uchar, cn>::vec_type dst_t;
|
||||||
|
|
||||||
|
const dim3 block(32, 8);
|
||||||
|
const dim3 grid(divUp(src.cols, 4 * block.x), divUp(src.rows, block.y));
|
||||||
|
|
||||||
|
cudaSafeCall( cudaFuncSetCacheConfig(Bayer2BGR_8u<dst_t>, cudaFuncCachePreferL1) );
|
||||||
|
|
||||||
|
Bayer2BGR_8u<dst_t><<<grid, block, 0, stream>>>(src, (PtrStepSz<dst_t>)dst, blue_last, start_with_green);
|
||||||
|
cudaSafeCall( cudaGetLastError() );
|
||||||
|
|
||||||
|
if (stream == 0)
|
||||||
|
cudaSafeCall( cudaDeviceSynchronize() );
|
||||||
|
}
|
||||||
|
|
||||||
|
template <int cn>
|
||||||
|
void Bayer2BGR_16u_gpu(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream)
|
||||||
|
{
|
||||||
|
typedef typename TypeVec<ushort, cn>::vec_type dst_t;
|
||||||
|
|
||||||
|
const dim3 block(32, 8);
|
||||||
|
const dim3 grid(divUp(src.cols, 2 * block.x), divUp(src.rows, block.y));
|
||||||
|
|
||||||
|
cudaSafeCall( cudaFuncSetCacheConfig(Bayer2BGR_16u<dst_t>, cudaFuncCachePreferL1) );
|
||||||
|
|
||||||
|
Bayer2BGR_16u<dst_t><<<grid, block, 0, stream>>>(src, (PtrStepSz<dst_t>)dst, blue_last, start_with_green);
|
||||||
|
cudaSafeCall( cudaGetLastError() );
|
||||||
|
|
||||||
|
if (stream == 0)
|
||||||
|
cudaSafeCall( cudaDeviceSynchronize() );
|
||||||
|
}
|
||||||
|
|
||||||
|
template void Bayer2BGR_8u_gpu<1>(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream);
|
||||||
|
template void Bayer2BGR_8u_gpu<3>(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream);
|
||||||
|
template void Bayer2BGR_8u_gpu<4>(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream);
|
||||||
|
|
||||||
|
template void Bayer2BGR_16u_gpu<1>(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream);
|
||||||
|
template void Bayer2BGR_16u_gpu<3>(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream);
|
||||||
|
template void Bayer2BGR_16u_gpu<4>(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream);
|
||||||
|
}}}
|
||||||
|
|
||||||
#endif /* CUDA_DISABLER */
|
#endif /* CUDA_DISABLER */
|
@ -2218,6 +2218,70 @@ GPU_TEST_P(CvtColor, BayerGR2BGR4)
|
|||||||
EXPECT_MAT_NEAR(dst_gold(cv::Rect(1, 1, dst.cols - 2, dst.rows - 2)), dst3(cv::Rect(1, 1, dst.cols - 2, dst.rows - 2)), 0);
|
EXPECT_MAT_NEAR(dst_gold(cv::Rect(1, 1, dst.cols - 2, dst.rows - 2)), dst3(cv::Rect(1, 1, dst.cols - 2, dst.rows - 2)), 0);
|
||||||
}
|
}
|
||||||
|
|
||||||
|
GPU_TEST_P(CvtColor, BayerBG2Gray)
|
||||||
|
{
|
||||||
|
if ((depth != CV_8U && depth != CV_16U) || useRoi)
|
||||||
|
return;
|
||||||
|
|
||||||
|
cv::Mat src = randomMat(size, depth);
|
||||||
|
|
||||||
|
cv::gpu::GpuMat dst;
|
||||||
|
cv::gpu::cvtColor(loadMat(src, useRoi), dst, cv::COLOR_BayerBG2GRAY);
|
||||||
|
|
||||||
|
cv::Mat dst_gold;
|
||||||
|
cv::cvtColor(src, dst_gold, cv::COLOR_BayerBG2GRAY);
|
||||||
|
|
||||||
|
EXPECT_MAT_NEAR(dst_gold(cv::Rect(1, 1, dst.cols - 2, dst.rows - 2)), dst(cv::Rect(1, 1, dst.cols - 2, dst.rows - 2)), 2);
|
||||||
|
}
|
||||||
|
|
||||||
|
GPU_TEST_P(CvtColor, BayerGB2Gray)
|
||||||
|
{
|
||||||
|
if ((depth != CV_8U && depth != CV_16U) || useRoi)
|
||||||
|
return;
|
||||||
|
|
||||||
|
cv::Mat src = randomMat(size, depth);
|
||||||
|
|
||||||
|
cv::gpu::GpuMat dst;
|
||||||
|
cv::gpu::cvtColor(loadMat(src, useRoi), dst, cv::COLOR_BayerGB2GRAY);
|
||||||
|
|
||||||
|
cv::Mat dst_gold;
|
||||||
|
cv::cvtColor(src, dst_gold, cv::COLOR_BayerGB2GRAY);
|
||||||
|
|
||||||
|
EXPECT_MAT_NEAR(dst_gold(cv::Rect(1, 1, dst.cols - 2, dst.rows - 2)), dst(cv::Rect(1, 1, dst.cols - 2, dst.rows - 2)), 2);
|
||||||
|
}
|
||||||
|
|
||||||
|
GPU_TEST_P(CvtColor, BayerRG2Gray)
|
||||||
|
{
|
||||||
|
if ((depth != CV_8U && depth != CV_16U) || useRoi)
|
||||||
|
return;
|
||||||
|
|
||||||
|
cv::Mat src = randomMat(size, depth);
|
||||||
|
|
||||||
|
cv::gpu::GpuMat dst;
|
||||||
|
cv::gpu::cvtColor(loadMat(src, useRoi), dst, cv::COLOR_BayerRG2GRAY);
|
||||||
|
|
||||||
|
cv::Mat dst_gold;
|
||||||
|
cv::cvtColor(src, dst_gold, cv::COLOR_BayerRG2GRAY);
|
||||||
|
|
||||||
|
EXPECT_MAT_NEAR(dst_gold(cv::Rect(1, 1, dst.cols - 2, dst.rows - 2)), dst(cv::Rect(1, 1, dst.cols - 2, dst.rows - 2)), 2);
|
||||||
|
}
|
||||||
|
|
||||||
|
GPU_TEST_P(CvtColor, BayerGR2Gray)
|
||||||
|
{
|
||||||
|
if ((depth != CV_8U && depth != CV_16U) || useRoi)
|
||||||
|
return;
|
||||||
|
|
||||||
|
cv::Mat src = randomMat(size, depth);
|
||||||
|
|
||||||
|
cv::gpu::GpuMat dst;
|
||||||
|
cv::gpu::cvtColor(loadMat(src, useRoi), dst, cv::COLOR_BayerGR2GRAY);
|
||||||
|
|
||||||
|
cv::Mat dst_gold;
|
||||||
|
cv::cvtColor(src, dst_gold, cv::COLOR_BayerGR2GRAY);
|
||||||
|
|
||||||
|
EXPECT_MAT_NEAR(dst_gold(cv::Rect(1, 1, dst.cols - 2, dst.rows - 2)), dst(cv::Rect(1, 1, dst.cols - 2, dst.rows - 2)), 2);
|
||||||
|
}
|
||||||
|
|
||||||
INSTANTIATE_TEST_CASE_P(GPU_ImgProc, CvtColor, testing::Combine(
|
INSTANTIATE_TEST_CASE_P(GPU_ImgProc, CvtColor, testing::Combine(
|
||||||
ALL_DEVICES,
|
ALL_DEVICES,
|
||||||
DIFFERENT_SIZES,
|
DIFFERENT_SIZES,
|
||||||
|
@ -15,6 +15,7 @@ import android.content.DialogInterface;
|
|||||||
import android.content.res.TypedArray;
|
import android.content.res.TypedArray;
|
||||||
import android.graphics.Bitmap;
|
import android.graphics.Bitmap;
|
||||||
import android.graphics.Canvas;
|
import android.graphics.Canvas;
|
||||||
|
import android.graphics.Rect;
|
||||||
import android.util.AttributeSet;
|
import android.util.AttributeSet;
|
||||||
import android.util.Log;
|
import android.util.Log;
|
||||||
import android.view.SurfaceHolder;
|
import android.view.SurfaceHolder;
|
||||||
@ -44,6 +45,7 @@ public abstract class CameraBridgeViewBase extends SurfaceView implements Surfac
|
|||||||
protected int mFrameHeight;
|
protected int mFrameHeight;
|
||||||
protected int mMaxHeight;
|
protected int mMaxHeight;
|
||||||
protected int mMaxWidth;
|
protected int mMaxWidth;
|
||||||
|
protected float mScale = 0;
|
||||||
protected int mPreviewFormat = Highgui.CV_CAP_ANDROID_COLOR_FRAME_RGBA;
|
protected int mPreviewFormat = Highgui.CV_CAP_ANDROID_COLOR_FRAME_RGBA;
|
||||||
protected int mCameraIndex = -1;
|
protected int mCameraIndex = -1;
|
||||||
protected boolean mEnabled;
|
protected boolean mEnabled;
|
||||||
@ -389,7 +391,22 @@ public abstract class CameraBridgeViewBase extends SurfaceView implements Surfac
|
|||||||
Canvas canvas = getHolder().lockCanvas();
|
Canvas canvas = getHolder().lockCanvas();
|
||||||
if (canvas != null) {
|
if (canvas != null) {
|
||||||
canvas.drawColor(0, android.graphics.PorterDuff.Mode.CLEAR);
|
canvas.drawColor(0, android.graphics.PorterDuff.Mode.CLEAR);
|
||||||
canvas.drawBitmap(mCacheBitmap, (canvas.getWidth() - mCacheBitmap.getWidth()) / 2, (canvas.getHeight() - mCacheBitmap.getHeight()) / 2, null);
|
Log.d(TAG, "mStretch value: " + mScale);
|
||||||
|
|
||||||
|
if (mScale != 0) {
|
||||||
|
canvas.drawBitmap(mCacheBitmap, new Rect(0,0,mCacheBitmap.getWidth(), mCacheBitmap.getHeight()),
|
||||||
|
new Rect((int)((canvas.getWidth() - mScale*mCacheBitmap.getWidth()) / 2),
|
||||||
|
(int)((canvas.getHeight() - mScale*mCacheBitmap.getHeight()) / 2),
|
||||||
|
(int)((canvas.getWidth() - mScale*mCacheBitmap.getWidth()) / 2 + mScale*mCacheBitmap.getWidth()),
|
||||||
|
(int)((canvas.getHeight() - mScale*mCacheBitmap.getHeight()) / 2 + mScale*mCacheBitmap.getHeight())), null);
|
||||||
|
} else {
|
||||||
|
canvas.drawBitmap(mCacheBitmap, new Rect(0,0,mCacheBitmap.getWidth(), mCacheBitmap.getHeight()),
|
||||||
|
new Rect((canvas.getWidth() - mCacheBitmap.getWidth()) / 2,
|
||||||
|
(canvas.getHeight() - mCacheBitmap.getHeight()) / 2,
|
||||||
|
(canvas.getWidth() - mCacheBitmap.getWidth()) / 2 + mCacheBitmap.getWidth(),
|
||||||
|
(canvas.getHeight() - mCacheBitmap.getHeight()) / 2 + mCacheBitmap.getHeight()), null);
|
||||||
|
}
|
||||||
|
|
||||||
if (mFpsMeter != null) {
|
if (mFpsMeter != null) {
|
||||||
mFpsMeter.measure();
|
mFpsMeter.measure();
|
||||||
mFpsMeter.draw(canvas, 20, 30);
|
mFpsMeter.draw(canvas, 20, 30);
|
||||||
|
@ -10,6 +10,7 @@ import android.hardware.Camera.PreviewCallback;
|
|||||||
import android.os.Build;
|
import android.os.Build;
|
||||||
import android.util.AttributeSet;
|
import android.util.AttributeSet;
|
||||||
import android.util.Log;
|
import android.util.Log;
|
||||||
|
import android.view.ViewGroup.LayoutParams;
|
||||||
|
|
||||||
import org.opencv.core.CvType;
|
import org.opencv.core.CvType;
|
||||||
import org.opencv.core.Mat;
|
import org.opencv.core.Mat;
|
||||||
@ -130,6 +131,11 @@ public class JavaCameraView extends CameraBridgeViewBase implements PreviewCallb
|
|||||||
mFrameWidth = params.getPreviewSize().width;
|
mFrameWidth = params.getPreviewSize().width;
|
||||||
mFrameHeight = params.getPreviewSize().height;
|
mFrameHeight = params.getPreviewSize().height;
|
||||||
|
|
||||||
|
if ((getLayoutParams().width == LayoutParams.MATCH_PARENT) && (getLayoutParams().height == LayoutParams.MATCH_PARENT))
|
||||||
|
mScale = Math.min(((float)height)/mFrameHeight, ((float)width)/mFrameWidth);
|
||||||
|
else
|
||||||
|
mScale = 0;
|
||||||
|
|
||||||
if (mFpsMeter != null) {
|
if (mFpsMeter != null) {
|
||||||
mFpsMeter.setResolution(mFrameWidth, mFrameHeight);
|
mFpsMeter.setResolution(mFrameWidth, mFrameHeight);
|
||||||
}
|
}
|
||||||
|
@ -8,6 +8,7 @@ import org.opencv.highgui.VideoCapture;
|
|||||||
import android.content.Context;
|
import android.content.Context;
|
||||||
import android.util.AttributeSet;
|
import android.util.AttributeSet;
|
||||||
import android.util.Log;
|
import android.util.Log;
|
||||||
|
import android.view.ViewGroup.LayoutParams;
|
||||||
|
|
||||||
/**
|
/**
|
||||||
* This class is an implementation of a bridge between SurfaceView and native OpenCV camera.
|
* This class is an implementation of a bridge between SurfaceView and native OpenCV camera.
|
||||||
@ -102,6 +103,11 @@ public class NativeCameraView extends CameraBridgeViewBase {
|
|||||||
mFrameWidth = (int)frameSize.width;
|
mFrameWidth = (int)frameSize.width;
|
||||||
mFrameHeight = (int)frameSize.height;
|
mFrameHeight = (int)frameSize.height;
|
||||||
|
|
||||||
|
if ((getLayoutParams().width == LayoutParams.MATCH_PARENT) && (getLayoutParams().height == LayoutParams.MATCH_PARENT))
|
||||||
|
mScale = Math.min(((float)height)/mFrameHeight, ((float)width)/mFrameWidth);
|
||||||
|
else
|
||||||
|
mScale = 0;
|
||||||
|
|
||||||
if (mFpsMeter != null) {
|
if (mFpsMeter != null) {
|
||||||
mFpsMeter.setResolution(mFrameWidth, mFrameHeight);
|
mFpsMeter.setResolution(mFrameWidth, mFrameHeight);
|
||||||
}
|
}
|
||||||
|
@ -510,7 +510,7 @@ static bool pyopencv_to(PyObject* obj, double& value, const char* name = "<unkno
|
|||||||
(void)name;
|
(void)name;
|
||||||
if(!obj || obj == Py_None)
|
if(!obj || obj == Py_None)
|
||||||
return true;
|
return true;
|
||||||
if(PyInt_CheckExact(obj))
|
if(!!PyInt_CheckExact(obj))
|
||||||
value = (double)PyInt_AS_LONG(obj);
|
value = (double)PyInt_AS_LONG(obj);
|
||||||
else
|
else
|
||||||
value = PyFloat_AsDouble(obj);
|
value = PyFloat_AsDouble(obj);
|
||||||
@ -527,7 +527,7 @@ static bool pyopencv_to(PyObject* obj, float& value, const char* name = "<unknow
|
|||||||
(void)name;
|
(void)name;
|
||||||
if(!obj || obj == Py_None)
|
if(!obj || obj == Py_None)
|
||||||
return true;
|
return true;
|
||||||
if(PyInt_CheckExact(obj))
|
if(!!PyInt_CheckExact(obj))
|
||||||
value = (float)PyInt_AS_LONG(obj);
|
value = (float)PyInt_AS_LONG(obj);
|
||||||
else
|
else
|
||||||
value = (float)PyFloat_AsDouble(obj);
|
value = (float)PyFloat_AsDouble(obj);
|
||||||
@ -623,7 +623,7 @@ static inline bool pyopencv_to(PyObject* obj, Point& p, const char* name = "<unk
|
|||||||
(void)name;
|
(void)name;
|
||||||
if(!obj || obj == Py_None)
|
if(!obj || obj == Py_None)
|
||||||
return true;
|
return true;
|
||||||
if(PyComplex_CheckExact(obj))
|
if(!!PyComplex_CheckExact(obj))
|
||||||
{
|
{
|
||||||
Py_complex c = PyComplex_AsCComplex(obj);
|
Py_complex c = PyComplex_AsCComplex(obj);
|
||||||
p.x = saturate_cast<int>(c.real);
|
p.x = saturate_cast<int>(c.real);
|
||||||
@ -638,7 +638,7 @@ static inline bool pyopencv_to(PyObject* obj, Point2f& p, const char* name = "<u
|
|||||||
(void)name;
|
(void)name;
|
||||||
if(!obj || obj == Py_None)
|
if(!obj || obj == Py_None)
|
||||||
return true;
|
return true;
|
||||||
if(PyComplex_CheckExact(obj))
|
if(!!PyComplex_CheckExact(obj))
|
||||||
{
|
{
|
||||||
Py_complex c = PyComplex_AsCComplex(obj);
|
Py_complex c = PyComplex_AsCComplex(obj);
|
||||||
p.x = saturate_cast<float>(c.real);
|
p.x = saturate_cast<float>(c.real);
|
||||||
@ -989,7 +989,7 @@ static bool pyopencv_to(PyObject *o, cv::flann::IndexParams& p, const char *name
|
|||||||
const char* value = PyString_AsString(item);
|
const char* value = PyString_AsString(item);
|
||||||
p.setString(k, value);
|
p.setString(k, value);
|
||||||
}
|
}
|
||||||
else if( PyBool_Check(item) )
|
else if( !!PyBool_Check(item) )
|
||||||
p.setBool(k, item == Py_True);
|
p.setBool(k, item == Py_True);
|
||||||
else if( PyInt_Check(item) )
|
else if( PyInt_Check(item) )
|
||||||
{
|
{
|
||||||
|
@ -1158,7 +1158,7 @@ static PyObject* cvseq_map_getitem(PyObject *o, PyObject *item)
|
|||||||
if (i < 0)
|
if (i < 0)
|
||||||
i += (int)cvseq_seq_length(o);
|
i += (int)cvseq_seq_length(o);
|
||||||
return cvseq_seq_getitem(o, i);
|
return cvseq_seq_getitem(o, i);
|
||||||
} else if (PySlice_Check(item)) {
|
} else if (!!PySlice_Check(item)) {
|
||||||
Py_ssize_t start, stop, step, slicelength, cur, i;
|
Py_ssize_t start, stop, step, slicelength, cur, i;
|
||||||
PyObject* result;
|
PyObject* result;
|
||||||
|
|
||||||
@ -1975,7 +1975,7 @@ struct dims
|
|||||||
|
|
||||||
static int convert_to_dim(PyObject *item, int i, dims *dst, CvArr *cva, const char *name = "no_name")
|
static int convert_to_dim(PyObject *item, int i, dims *dst, CvArr *cva, const char *name = "no_name")
|
||||||
{
|
{
|
||||||
if (PySlice_Check(item)) {
|
if (!!PySlice_Check(item)) {
|
||||||
Py_ssize_t start, stop, step, slicelength;
|
Py_ssize_t start, stop, step, slicelength;
|
||||||
PySlice_GetIndicesEx((PySliceObject*)item, cvGetDimSize(cva, i), &start, &stop, &step, &slicelength);
|
PySlice_GetIndicesEx((PySliceObject*)item, cvGetDimSize(cva, i), &start, &stop, &step, &slicelength);
|
||||||
dst->i[i] = (int)start;
|
dst->i[i] = (int)start;
|
||||||
|
@ -4,8 +4,8 @@
|
|||||||
android:layout_height="match_parent" >
|
android:layout_height="match_parent" >
|
||||||
|
|
||||||
<org.opencv.samples.tutorial3.Tutorial3View
|
<org.opencv.samples.tutorial3.Tutorial3View
|
||||||
android:layout_width="fill_parent"
|
android:layout_width="match_parent"
|
||||||
android:layout_height="fill_parent"
|
android:layout_height="match_parent"
|
||||||
android:visibility="gone"
|
android:visibility="gone"
|
||||||
android:id="@+id/tutorial3_activity_java_surface_view" />
|
android:id="@+id/tutorial3_activity_java_surface_view" />
|
||||||
|
|
||||||
|
Loading…
x
Reference in New Issue
Block a user