mirror of
https://github.com/opencv/opencv.git
synced 2026-07-31 00:03:03 +04:00
Merge branch 4.x
This commit is contained in:
@@ -54,23 +54,23 @@ namespace cv { namespace cuda { namespace device
|
||||
{
|
||||
// TODO use intrinsics like __sinf and so on
|
||||
|
||||
namespace build_warp_maps
|
||||
struct WarpParams
|
||||
{
|
||||
|
||||
__constant__ float ck_rinv[9];
|
||||
__constant__ float cr_kinv[9];
|
||||
__constant__ float ct[3];
|
||||
__constant__ float cscale;
|
||||
}
|
||||
|
||||
float k_rinv[9];
|
||||
float r_kinv[9];
|
||||
float t[3];
|
||||
float scale;
|
||||
};
|
||||
|
||||
class PlaneMapper
|
||||
{
|
||||
public:
|
||||
static __device__ __forceinline__ void mapBackward(float u, float v, float &x, float &y)
|
||||
static __device__ __forceinline__ void mapBackward(float u, float v, float &x, float &y,
|
||||
const WarpParams& params)
|
||||
{
|
||||
using namespace build_warp_maps;
|
||||
|
||||
const float *ck_rinv = params.k_rinv;
|
||||
const float *ct = params.t;
|
||||
const float cscale = params.scale;
|
||||
float x_ = u / cscale - ct[0];
|
||||
float y_ = v / cscale - ct[1];
|
||||
|
||||
@@ -88,10 +88,11 @@ namespace cv { namespace cuda { namespace device
|
||||
class CylindricalMapper
|
||||
{
|
||||
public:
|
||||
static __device__ __forceinline__ void mapBackward(float u, float v, float &x, float &y)
|
||||
static __device__ __forceinline__ void mapBackward(float u, float v, float &x, float &y,
|
||||
const WarpParams& params)
|
||||
{
|
||||
using namespace build_warp_maps;
|
||||
|
||||
const float *ck_rinv = params.k_rinv;
|
||||
const float cscale = params.scale;
|
||||
u /= cscale;
|
||||
float x_ = ::sinf(u);
|
||||
float y_ = v / cscale;
|
||||
@@ -111,10 +112,11 @@ namespace cv { namespace cuda { namespace device
|
||||
class SphericalMapper
|
||||
{
|
||||
public:
|
||||
static __device__ __forceinline__ void mapBackward(float u, float v, float &x, float &y)
|
||||
static __device__ __forceinline__ void mapBackward(float u, float v, float &x, float &y,
|
||||
const WarpParams& params)
|
||||
{
|
||||
using namespace build_warp_maps;
|
||||
|
||||
const float *ck_rinv = params.k_rinv;
|
||||
const float cscale = params.scale;
|
||||
v /= cscale;
|
||||
u /= cscale;
|
||||
|
||||
@@ -136,7 +138,8 @@ namespace cv { namespace cuda { namespace device
|
||||
|
||||
template <typename Mapper>
|
||||
__global__ void buildWarpMapsKernel(int tl_u, int tl_v, int cols, int rows,
|
||||
PtrStepf map_x, PtrStepf map_y)
|
||||
PtrStepf map_x, PtrStepf map_y,
|
||||
const WarpParams params)
|
||||
{
|
||||
int du = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
int dv = blockIdx.y * blockDim.y + threadIdx.y;
|
||||
@@ -145,7 +148,7 @@ namespace cv { namespace cuda { namespace device
|
||||
float u = tl_u + du;
|
||||
float v = tl_v + dv;
|
||||
float x, y;
|
||||
Mapper::mapBackward(u, v, x, y);
|
||||
Mapper::mapBackward(u, v, x, y, params);
|
||||
map_x.ptr(dv)[du] = x;
|
||||
map_y.ptr(dv)[du] = y;
|
||||
}
|
||||
@@ -156,10 +159,16 @@ namespace cv { namespace cuda { namespace device
|
||||
const float k_rinv[9], const float r_kinv[9], const float t[3],
|
||||
float scale, cudaStream_t stream)
|
||||
{
|
||||
cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::ck_rinv, k_rinv, 9*sizeof(float)));
|
||||
cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cr_kinv, r_kinv, 9*sizeof(float)));
|
||||
cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::ct, t, 3*sizeof(float)));
|
||||
cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cscale, &scale, sizeof(float)));
|
||||
WarpParams params;
|
||||
for (int i = 0; i < 9; ++i)
|
||||
{
|
||||
params.k_rinv[i] = k_rinv[i];
|
||||
params.r_kinv[i] = r_kinv[i];
|
||||
}
|
||||
params.t[0] = t[0];
|
||||
params.t[1] = t[1];
|
||||
params.t[2] = t[2];
|
||||
params.scale = scale;
|
||||
|
||||
int cols = map_x.cols;
|
||||
int rows = map_x.rows;
|
||||
@@ -167,7 +176,8 @@ namespace cv { namespace cuda { namespace device
|
||||
dim3 threads(32, 8);
|
||||
dim3 grid(divUp(cols, threads.x), divUp(rows, threads.y));
|
||||
|
||||
buildWarpMapsKernel<PlaneMapper><<<grid,threads>>>(tl_u, tl_v, cols, rows, map_x, map_y);
|
||||
buildWarpMapsKernel<PlaneMapper><<<grid,threads>>>(tl_u, tl_v, cols, rows,
|
||||
map_x, map_y, params);
|
||||
cudaSafeCall(cudaGetLastError());
|
||||
if (stream == 0)
|
||||
cudaSafeCall(cudaDeviceSynchronize());
|
||||
@@ -178,9 +188,16 @@ namespace cv { namespace cuda { namespace device
|
||||
const float k_rinv[9], const float r_kinv[9], float scale,
|
||||
cudaStream_t stream)
|
||||
{
|
||||
cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::ck_rinv, k_rinv, 9*sizeof(float)));
|
||||
cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cr_kinv, r_kinv, 9*sizeof(float)));
|
||||
cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cscale, &scale, sizeof(float)));
|
||||
WarpParams params;
|
||||
for (int i = 0; i < 9; ++i)
|
||||
{
|
||||
params.k_rinv[i] = k_rinv[i];
|
||||
params.r_kinv[i] = r_kinv[i];
|
||||
}
|
||||
params.t[0] = 0.f;
|
||||
params.t[1] = 0.f;
|
||||
params.t[2] = 0.f;
|
||||
params.scale = scale;
|
||||
|
||||
int cols = map_x.cols;
|
||||
int rows = map_x.rows;
|
||||
@@ -188,7 +205,8 @@ namespace cv { namespace cuda { namespace device
|
||||
dim3 threads(32, 8);
|
||||
dim3 grid(divUp(cols, threads.x), divUp(rows, threads.y));
|
||||
|
||||
buildWarpMapsKernel<CylindricalMapper><<<grid,threads>>>(tl_u, tl_v, cols, rows, map_x, map_y);
|
||||
buildWarpMapsKernel<CylindricalMapper><<<grid,threads>>>(tl_u, tl_v, cols, rows,
|
||||
map_x, map_y, params);
|
||||
cudaSafeCall(cudaGetLastError());
|
||||
if (stream == 0)
|
||||
cudaSafeCall(cudaDeviceSynchronize());
|
||||
@@ -199,9 +217,16 @@ namespace cv { namespace cuda { namespace device
|
||||
const float k_rinv[9], const float r_kinv[9], float scale,
|
||||
cudaStream_t stream)
|
||||
{
|
||||
cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::ck_rinv, k_rinv, 9*sizeof(float)));
|
||||
cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cr_kinv, r_kinv, 9*sizeof(float)));
|
||||
cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cscale, &scale, sizeof(float)));
|
||||
WarpParams params;
|
||||
for (int i = 0; i < 9; ++i)
|
||||
{
|
||||
params.k_rinv[i] = k_rinv[i];
|
||||
params.r_kinv[i] = r_kinv[i];
|
||||
}
|
||||
params.t[0] = 0.f;
|
||||
params.t[1] = 0.f;
|
||||
params.t[2] = 0.f;
|
||||
params.scale = scale;
|
||||
|
||||
int cols = map_x.cols;
|
||||
int rows = map_x.rows;
|
||||
@@ -209,7 +234,8 @@ namespace cv { namespace cuda { namespace device
|
||||
dim3 threads(32, 8);
|
||||
dim3 grid(divUp(cols, threads.x), divUp(rows, threads.y));
|
||||
|
||||
buildWarpMapsKernel<SphericalMapper><<<grid,threads>>>(tl_u, tl_v, cols, rows, map_x, map_y);
|
||||
buildWarpMapsKernel<SphericalMapper><<<grid,threads>>>(tl_u, tl_v, cols, rows,
|
||||
map_x, map_y, params);
|
||||
cudaSafeCall(cudaGetLastError());
|
||||
if (stream == 0)
|
||||
cudaSafeCall(cudaDeviceSynchronize());
|
||||
|
||||
@@ -338,8 +338,8 @@ void DpSeamFinder::findComponents()
|
||||
states_.push_back(SECOND);
|
||||
|
||||
floodFill(labels_, Point(x, y), ++ncomps_);
|
||||
tls_.push_back(Point(x, y));
|
||||
brs_.push_back(Point(x+1, y+1));
|
||||
tls_.emplace_back(x,y);
|
||||
brs_.emplace_back(x+1, y+1);
|
||||
contours_.push_back(std::vector<Point>());
|
||||
}
|
||||
|
||||
@@ -356,7 +356,7 @@ void DpSeamFinder::findComponents()
|
||||
if ((x == 0 || labels_(y, x-1) != l) || (x == unionSize_.width-1 || labels_(y, x+1) != l) ||
|
||||
(y == 0 || labels_(y-1, x) != l) || (y == unionSize_.height-1 || labels_(y+1, x) != l))
|
||||
{
|
||||
contours_[ci].push_back(Point(x, y));
|
||||
contours_[ci].emplace_back(x,y);
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -517,7 +517,7 @@ void DpSeamFinder::resolveConflicts(
|
||||
if ((x == 0 || labels_(y, x-1) != l[i]) || (x == unionSize_.width-1 || labels_(y, x+1) != l[i]) ||
|
||||
(y == 0 || labels_(y-1, x) != l[i]) || (y == unionSize_.height-1 || labels_(y+1, x) != l[i]))
|
||||
{
|
||||
contours_[c[i]].push_back(Point(x, y));
|
||||
contours_[c[i]].emplace_back(x,y);
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -637,7 +637,7 @@ bool DpSeamFinder::getSeamTips(int comp1, int comp2, Point &p1, Point &p2)
|
||||
(x < unionSize_.width-1 && labels_(y, x+1) == l2) ||
|
||||
(y < unionSize_.height-1 && labels_(y+1, x) == l2)))
|
||||
{
|
||||
specialPoints.push_back(Point(x, y));
|
||||
specialPoints.emplace_back(x,y);
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
Reference in New Issue
Block a user