Skip to content
Snippets Groups Projects
Commit 9cb56ada authored by Nicolas Pope's avatar Nicolas Pope
Browse files

Move carving to fusion step

parent 185aed3d
No related branches found
No related tags found
1 merge request!354Add full 3D MLS and carving
Pipeline #33593 passed
Stage: all
Hide whitespace changes
Inline Side-by-side
Showing
with 128 additions and 369 deletions
components/operators/CMakeLists.txt
+ 1
− 0
View file @ 9cb56ada
......@@ -27,6 +27,7 @@ set(OPERSRC
src/fusion/correspondence_util.cu
src/fusion/mls_aggr.cu
src/fusion/smoothing/mls_multi_weighted.cu
src/fusion/carving/carver.cu
src/fusion/fusion.cpp
src/misc/clipping.cpp
src/disparity/depth.cpp
......
components/renderers/cpp/src/carver.hpp → components/operators/include/ftl/operators/cuda/carver.hpp
+ 34
− 0
View file @ 9cb56ada
......@@ -14,9 +14,9 @@ namespace cuda {
void depth_carve(
cv::cuda::GpuMat &in,
const cv::cuda::GpuMat &ref,
const cv::cuda::GpuMat &in_colour,
const cv::cuda::GpuMat &ref_colour,
cv::cuda::GpuMat &colour_scale,
//const cv::cuda::GpuMat &in_colour,
//const cv::cuda::GpuMat &ref_colour,
//cv::cuda::GpuMat &colour_scale,
const float4x4 &transform,
const ftl::rgbd::Camera &incam,
const ftl::rgbd::Camera &refcam,
......@@ -28,32 +28,6 @@ void apply_colour_scaling(
int radius,
cudaStream_t stream);
void mls_reduce(
const cv::cuda::GpuMat &centroid,
const cv::cuda::GpuMat &normals,
const cv::cuda::GpuMat &contrib,
cv::cuda::GpuMat &normals_out,
cv::cuda::GpuMat &depth,
const ftl::rgbd::Camera &camera,
cudaStream_t stream
);
void mls_gather(
const cv::cuda::GpuMat &normals_in, // Source frame
cv::cuda::GpuMat &normals_out,
const cv::cuda::GpuMat &depth_origin, // Rendered image
const cv::cuda::GpuMat &depth_in,
cv::cuda::GpuMat &centroid_out,
cv::cuda::GpuMat &contrib_out,
float smoothing,
const float4x4 &o_2_in,
const float4x4 &in_2_o,
const float3x3 &in_2_o33,
const ftl::rgbd::Camera &camera_origin, // Virtual camera
const ftl::rgbd::Camera &camera_in,
cudaStream_t stream
);
}
}
......
components/renderers/cpp/src/carver.cu → components/operators/src/fusion/carving/carver.cu
+ 23
− 214
View file @ 9cb56ada
#include "carver.hpp"
#include <ftl/operators/cuda/carver.hpp>
#include <cudatl/fixed.hpp>
#include <ftl/cuda/weighting.hpp>
......@@ -9,7 +9,7 @@ __device__ inline float depthErrorCoef(const ftl::rgbd::Camera &cam, float disps
// ==== Reverse Verify Result ==================================================
// No colour scale calculations
__global__ void reverse_check_kernel(
/*__global__ void reverse_check_kernel(
float* __restrict__ depth_in,
const float* __restrict__ depth_original,
int pitch4,
......@@ -48,21 +48,21 @@ __global__ void reverse_check_kernel(
// Too much carving means just outright remove the point.
depth_in[y*pitch4+x] = (count < 0) ? 0.0f : d;
}
}*/
__global__ void reverse_check_kernel(
float* __restrict__ depth_in,
const float* __restrict__ depth_original,
const uchar4* __restrict__ in_colour,
const uchar4* __restrict__ ref_colour,
int8_t* __restrict__ colour_scale,
//const uchar4* __restrict__ in_colour,
//const uchar4* __restrict__ ref_colour,
//int8_t* __restrict__ colour_scale,
int pitch4,
int pitch,
//int pitch,
int opitch4,
int in_col_pitch4,
int o_col_pitch4,
int cwidth,
int cheight,
//int in_col_pitch4,
//int o_col_pitch4,
//int cwidth,
//int cheight,
float4x4 transformR,
ftl::rgbd::Camera vintrin,
ftl::rgbd::Camera ointrin
......@@ -136,9 +136,9 @@ __global__ void reverse_check_kernel(
void ftl::cuda::depth_carve(
cv::cuda::GpuMat &depth_in,
const cv::cuda::GpuMat &depth_original,
const cv::cuda::GpuMat &in_colour,
const cv::cuda::GpuMat &ref_colour,
cv::cuda::GpuMat &colour_scale,
//const cv::cuda::GpuMat &in_colour,
//const cv::cuda::GpuMat &ref_colour,
//cv::cuda::GpuMat &colour_scale,
const float4x4 &transformR,
const ftl::rgbd::Camera &vintrin,
const ftl::rgbd::Camera &ointrin,
......@@ -150,218 +150,27 @@ void ftl::cuda::depth_carve(
const dim3 gridSize((depth_in.cols + THREADS_X - 1)/THREADS_X, (depth_in.rows + THREADS_Y - 1)/THREADS_Y);
const dim3 blockSize(THREADS_X, THREADS_Y);
colour_scale.create(depth_in.size(), CV_8U);
//colour_scale.create(depth_in.size(), CV_8U);
reverse_check_kernel<<<gridSize, blockSize, 0, stream>>>(
depth_in.ptr<float>(),
depth_original.ptr<float>(),
in_colour.ptr<uchar4>(),
ref_colour.ptr<uchar4>(),
colour_scale.ptr<int8_t>(),
//in_colour.ptr<uchar4>(),
//ref_colour.ptr<uchar4>(),
//colour_scale.ptr<int8_t>(),
depth_in.step1(),
colour_scale.step1(),
//colour_scale.step1(),
depth_original.step1(),
in_colour.step1()/4,
ref_colour.step1()/4,
in_colour.cols,
in_colour.rows,
//in_colour.step1()/4,
//ref_colour.step1()/4,
//in_colour.cols,
//in_colour.rows,
transformR,
vintrin, ointrin);
cudaSafeCall( cudaGetLastError() );
}
// ==== Multi image MLS ========================================================
/*
* Gather points for Moving Least Squares, from each source image
*/
template <int SEARCH_RADIUS>
__global__ void mls_gather_kernel(
const half4* __restrict__ normals_in,
half4* __restrict__ normals_out,
const float* __restrict__ depth_origin,
const float* __restrict__ depth_in,
float4* __restrict__ centroid_out,
float* __restrict__ contrib_out,
float smoothing,
float4x4 o_2_in,
float4x4 in_2_o,
float3x3 in_2_o33,
ftl::rgbd::Camera camera_origin,
ftl::rgbd::Camera camera_in,
int npitch_out,
int cpitch_out,
int wpitch_out,
int dpitch_o,
int dpitch_i,
int npitch_in
) {
const int x = blockIdx.x*blockDim.x + threadIdx.x;
const int y = blockIdx.y*blockDim.y + threadIdx.y;
if (x < 0 || y < 0 || x >= camera_origin.width || y >= camera_origin.height) return;
float3 nX = make_float3(normals_out[y*npitch_out+x]);
float3 aX = make_float3(centroid_out[y*cpitch_out+x]);
float contrib = contrib_out[y*wpitch_out+x];
float d0 = depth_origin[x+y*dpitch_o];
if (d0 <= camera_origin.minDepth || d0 >= camera_origin.maxDepth) return;
float3 X = camera_origin.screenToCam((int)(x),(int)(y),d0);
int2 s = camera_in.camToScreen<int2>(o_2_in * X);
// Neighbourhood
for (int v=-SEARCH_RADIUS; v<=SEARCH_RADIUS; ++v) {
for (int u=-SEARCH_RADIUS; u<=SEARCH_RADIUS; ++u) {
const float d = (s.x+u >= 0 && s.x+u < camera_in.width && s.y+v >= 0 && s.y+v < camera_in.height) ? depth_in[s.x+u+(s.y+v)*dpitch_i] : 0.0f;
if (d <= camera_in.minDepth || d >= camera_in.maxDepth) continue;
// Point and normal of neighbour
const float3 Xi = in_2_o * camera_in.screenToCam(s.x+u, s.y+v, d);
const float3 Ni = make_float3(normals_in[s.x+u+(s.y+v)*npitch_in]);
// Gauss approx weighting function using point distance
const float w = (length(Ni) > 0.0f) ? ftl::cuda::spatialWeighting(X,Xi,smoothing) : 0.0f;
aX += Xi*w;
nX += (in_2_o33 * Ni)*w;
contrib += w;
}
}
normals_out[y*npitch_out+x] = make_half4(nX, 0.0f);
centroid_out[y*cpitch_out+x] = make_float4(aX, 0.0f);
contrib_out[y*wpitch_out+x] = contrib;
}
/**
* Convert accumulated values into estimate of depth and normals at pixel.
*/
__global__ void mls_reduce_kernel(
const float4* __restrict__ centroid,
const half4* __restrict__ normals,
const float* __restrict__ contrib_out,
half4* __restrict__ normals_out,
float* __restrict__ depth,
ftl::rgbd::Camera camera,
int npitch_in,
int cpitch_in,
int wpitch,
int npitch,
int dpitch
) {
const int x = blockIdx.x*blockDim.x + threadIdx.x;
const int y = blockIdx.y*blockDim.y + threadIdx.y;
if (x >= 0 && y >= 0 && x < camera.width && y < camera.height) {
float3 nX = make_float3(normals[y*npitch_in+x]);
float3 aX = make_float3(centroid[y*cpitch_in+x]);
float contrib = contrib_out[y*wpitch+x];
//depth[x+y*dpitch] = X.z;
normals_out[x+y*npitch] = make_half4(0.0f, 0.0f, 0.0f, 0.0f);
float d0 = depth[x+y*dpitch];
//depth[x+y*dpitch] = 0.0f;
if (d0 <= camera.minDepth || d0 >= camera.maxDepth || contrib == 0.0f) return;
float3 X = camera.screenToCam((int)(x),(int)(y),d0);
nX /= contrib; // Weighted average normal
aX /= contrib; // Weighted average point (centroid)
// Signed-Distance Field function
float fX = nX.x * (X.x - aX.x) + nX.y * (X.y - aX.y) + nX.z * (X.z - aX.z);
// Calculate new point using SDF function to adjust depth (and position)
X = X - nX * fX;
depth[x+y*dpitch] = X.z;
normals_out[x+y*npitch] = make_half4(nX / length(nX), 0.0f);
}
}
#define T_PER_BLOCK 8
void ftl::cuda::mls_gather(
const cv::cuda::GpuMat &normals_in, // Source frame
cv::cuda::GpuMat &normals_out,
const cv::cuda::GpuMat &depth_origin, // Rendered image
const cv::cuda::GpuMat &depth_in,
cv::cuda::GpuMat &centroid_out,
cv::cuda::GpuMat &contrib_out,
float smoothing,
const float4x4 &o_2_in,
const float4x4 &in_2_o,
const float3x3 &in_2_o33,
const ftl::rgbd::Camera &camera_origin, // Virtual camera
const ftl::rgbd::Camera &camera_in,
cudaStream_t stream
) {
const dim3 gridSize((depth_origin.cols + T_PER_BLOCK - 1)/T_PER_BLOCK, (depth_origin.rows + T_PER_BLOCK - 1)/T_PER_BLOCK);
const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
normals_out.create(depth_origin.size(), CV_16FC4);
centroid_out.create(depth_origin.size(), CV_32FC4);
contrib_out.create(depth_origin.size(), CV_32F);
mls_gather_kernel<3><<<gridSize, blockSize, 0, stream>>>(
normals_in.ptr<half4>(),
normals_out.ptr<half4>(),
depth_origin.ptr<float>(),
depth_in.ptr<float>(),
centroid_out.ptr<float4>(),
contrib_out.ptr<float>(),
smoothing,
o_2_in,
in_2_o,
in_2_o33,
camera_origin,
camera_in,
normals_out.step1()/4,
centroid_out.step1()/4,
contrib_out.step1(),
depth_origin.step1(),
depth_in.step1(),
normals_in.step1()/4
);
cudaSafeCall( cudaGetLastError() );
}
void ftl::cuda::mls_reduce(
const cv::cuda::GpuMat &centroid,
const cv::cuda::GpuMat &normals,
const cv::cuda::GpuMat &contrib,
cv::cuda::GpuMat &normals_out,
cv::cuda::GpuMat &depth,
const ftl::rgbd::Camera &camera,
cudaStream_t stream
) {
const dim3 gridSize((depth.cols + T_PER_BLOCK - 1)/T_PER_BLOCK, (depth.rows + T_PER_BLOCK - 1)/T_PER_BLOCK);
const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
normals_out.create(depth.size(), CV_16FC4);
mls_reduce_kernel<<<gridSize, blockSize, 0, stream>>>(
centroid.ptr<float4>(),
normals.ptr<half4>(),
contrib.ptr<float>(),
normals_out.ptr<half4>(),
depth.ptr<float>(),
camera,
normals.step1()/4,
centroid.step1()/4,
contrib.step1(),
normals_out.step1()/4,
depth.step1()
);
cudaSafeCall( cudaGetLastError() );
}
// ==== Apply colour scale =====================================================
template <int RADIUS>
......
components/operators/src/fusion/fusion.cpp
+ 33
− 1
View file @ 9cb56ada
#include <ftl/operators/fusion.hpp>
#include <ftl/operators/cuda/carver.hpp>
#include <ftl/utility/matrix_conversion.hpp>
#include <opencv2/core/cuda_stream_accessor.hpp>
......@@ -32,7 +33,38 @@ bool Fusion::apply(ftl::rgbd::FrameSet &in, ftl::rgbd::FrameSet &out, cudaStream
cv::cuda::cvtColor(col, temp_, cv::COLOR_BGRA2GRAY, 0, cvstream);
cv::cuda::resize(temp_, weights_[i], d.size(), 0, 0, cv::INTER_LINEAR, cvstream);
}
}
if (config()->value("visibility_carving", true)) {
for (size_t i=0; i < in.frames.size(); ++i) {
if (!in.hasFrame(i)) continue;
auto &f = in.frames[i].cast<ftl::rgbd::Frame>();
for (size_t j=0; j < in.frames.size(); ++j) {
if (i == j) continue;
if (!in.hasFrame(j)) continue;
auto &ref = in.frames[j].cast<ftl::rgbd::Frame>();
auto transformR = MatrixConversion::toCUDA(ref.getPose().cast<float>().inverse() * f.getPose().cast<float>());
ftl::cuda::depth_carve(
f.create<cv::cuda::GpuMat>(Channel::Depth),
ref.get<cv::cuda::GpuMat>(Channel::Depth),
//f.get<cv::cuda::GpuMat>(Channel::Colour),
//ref.get<cv::cuda::GpuMat>(Channel::Colour),
//colour_scale_,
transformR,
f.getLeft(),
ref.getLeft(),
stream
);
}
//ftl::cuda::apply_colour_scaling(colour_scale_, f.create<cv::cuda::GpuMat>(Channel::Colour), 3, stream_);
}
}
for (int iters=0; iters < mls_iters; ++iters) {
for (size_t i=0; i<in.frames.size(); ++i) {
......
components/renderers/cpp/CMakeLists.txt
+ 0
− 1
View file @ 9cb56ada
......@@ -13,7 +13,6 @@ add_library(ftlrender
src/gltexture.cpp
src/touch.cu
src/GLRender.cpp
src/carver.cu
#src/assimp_render.cpp
#src/assimp_scene.cpp
)
......
components/renderers/cpp/src/CUDARender.cpp
+ 37
− 153
View file @ 9cb56ada
......@@ -7,7 +7,6 @@
#include <ftl/render/colouriser.hpp>
#include <ftl/cuda/transform.hpp>
#include <ftl/operators/cuda/mls_cuda.hpp>
#include "carver.hpp"
#include <ftl/utility/image_debug.hpp>
#include <ftl/cuda/colour_cuda.hpp>
......@@ -256,62 +255,6 @@ void CUDARender::_mesh(ftl::rgbd::Frame &out, const Eigen::Matrix4d &t, cudaStre
// FIXME: Is it possible to remember previously if there should be depth?
bool use_depth = scene_->anyHasChannel(Channel::Depth) || scene_->anyHasChannel(Channel::GroundTruth);
// For each source depth map, verify results
if (value("pre_carve_inputs", false) && !_alreadySeen()) {
for (size_t i=0; i < scene_->frames.size(); ++i) {
//if (!scene_->hasFrame(i)) continue;
auto &f = scene_->frames[i].cast<ftl::rgbd::Frame>();
//auto *s = scene_->sources[i];
if (!f.has(Channel::Colour)) {
//LOG(ERROR) << "Missing required channel";
continue;
}
// We have the needed depth data?
if (use_depth && !f.hasOwn(Channel::Depth) && !f.hasOwn(Channel::GroundTruth)) {
continue;
}
for (size_t j=0; j < scene_->frames.size(); ++j) {
if (i == j) continue;
//if (!scene_->hasFrame(i)) continue;
auto &ref = scene_->frames[j].cast<ftl::rgbd::Frame>();
//auto *s = scene_->sources[i];
if (!ref.has(Channel::Colour)) {
//LOG(ERROR) << "Missing required channel";
continue;
}
// We have the needed depth data?
if (use_depth && !ref.hasOwn(Channel::Depth) && !ref.hasOwn(Channel::GroundTruth)) {
continue;
}
//auto pose = MatrixConversion::toCUDA(t.cast<float>() * f.getPose().cast<float>());
//Eigen::Matrix4f refpose = ref.getPose().cast<float>();
auto transformR = MatrixConversion::toCUDA(ref.getPose().cast<float>().inverse() * f.getPose().cast<float>()); //poseInverse_;
//auto transform = pose_ * MatrixConversion::toCUDA(t.cast<float>() * f.getPose().cast<float>());
ftl::cuda::depth_carve(
f.create<cv::cuda::GpuMat>(_getDepthChannel()),
ref.get<cv::cuda::GpuMat>(Channel::Depth),
f.get<cv::cuda::GpuMat>(Channel::Colour),
ref.get<cv::cuda::GpuMat>(Channel::Colour),
colour_scale_,
transformR,
f.getLeft(),
ref.getLeft(),
stream_
);
}
//ftl::cuda::apply_colour_scaling(colour_scale_, f.create<cv::cuda::GpuMat>(Channel::Colour), 3, stream_);
}
}
//if (!colour_scale_.empty()) ftl::utility::show_image(colour_scale_, "CScale", 1.0f, ftl::utility::ImageVisualisation::HEAT_MAPPED);
// For each source depth map
......@@ -427,110 +370,51 @@ void CUDARender::_mesh(ftl::rgbd::Frame &out, const Eigen::Matrix4d &t, cudaStre
// Generate actual depth map using MLS with mesh as estimate
float mls_smoothing = value("mls_smooth", 0.05f);
int mls_iter = value("mls_iter", 3);
if (value("mls_full", false)) {
// Clear buffers
mls_centroid_.create(params_.camera.height, params_.camera.width, CV_32FC4);
mls_contrib_.create(params_.camera.height, params_.camera.width, CV_32F);
mls_normals_.create(params_.camera.height, params_.camera.width, CV_16FC4);
for (int iter=0; iter<mls_iter; ++iter) {
mls_centroid_.setTo(cv::Scalar(0,0,0,0), cvstream);
mls_contrib_.setTo(cv::Scalar(0), cvstream);
mls_normals_.setTo(cv::Scalar(0,0,0,0), cvstream);
for (size_t i=0; i < scene_->frames.size(); ++i) {
auto &f = scene_->frames[i].cast<ftl::rgbd::Frame>();
if (!f.has(Channel::Colour)) continue;
// We have the needed depth data?
if (use_depth && !f.hasOwn(Channel::Depth) && !f.hasOwn(Channel::GroundTruth)) {
continue;
}
int mls_radius = value("mls_radius", 0);
//auto pose = MatrixConversion::toCUDA(t.cast<float>() * f.getPose().cast<float>());
// VCAM to Original
auto transformR = MatrixConversion::toCUDA(f.getPose().cast<float>().inverse() * t.cast<float>().inverse()) * poseInverse_;
// Original to VCAM
auto transform = pose_ * MatrixConversion::toCUDA(t.cast<float>() * f.getPose().cast<float>());
auto transform33 = pose_.getFloat3x3() * MatrixConversion::toCUDA(t.cast<float>() * f.getPose().cast<float>()).getFloat3x3();
ftl::cuda::mls_gather(
f.get<cv::cuda::GpuMat>(Channel::Normals),
mls_normals_,
out.get<cv::cuda::GpuMat>(_getDepthChannel()),
f.get<cv::cuda::GpuMat>(Channel::Depth),
mls_centroid_,
mls_contrib_,
mls_smoothing,
transformR,
transform,
transform33,
params_.camera,
f.getLeft(),
stream_
);
}
if (mls_radius == 0) {
ftl::cuda::normals(
out.createTexture<half4>(_getNormalsChannel()),
out.getTexture<float>(_getDepthChannel()),
params_.camera, stream_);
} else {
ftl::cuda::normals(
temp_.createTexture<half4>(Channel::Normals),
out.getTexture<float>(_getDepthChannel()),
params_.camera, stream_);
// Now reduce MLS results to new depth+normals
ftl::cuda::mls_reduce(
mls_centroid_,
mls_normals_,
mls_contrib_,
ftl::cuda::mls_smooth(
temp_.get<cv::cuda::GpuMat>(Channel::Normals),
out.create<cv::cuda::GpuMat>(_getNormalsChannel()),
out.create<cv::cuda::GpuMat>(_getDepthChannel()),
out.get<cv::cuda::GpuMat>(_getDepthChannel()),
//out.getTexture<float>(_getDepthChannel()),
value("mls_smooth", 0.01f),
value("mls_radius", 2),
params_.camera,
stream_
);
}
} else {
int mls_radius = value("mls_radius", 0);
if (mls_radius == 0) {
ftl::cuda::normals(
out.createTexture<half4>(_getNormalsChannel()),
out.getTexture<float>(_getDepthChannel()),
params_.camera, stream_);
} else {
ftl::cuda::normals(
temp_.createTexture<half4>(Channel::Normals),
out.getTexture<float>(_getDepthChannel()),
params_.camera, stream_);
ftl::cuda::mls_smooth(
temp_.get<cv::cuda::GpuMat>(Channel::Normals),
out.create<cv::cuda::GpuMat>(_getNormalsChannel()),
out.get<cv::cuda::GpuMat>(_getDepthChannel()),
//out.getTexture<float>(_getDepthChannel()),
value("mls_smooth", 0.01f),
value("mls_radius", 2),
params_.camera,
stream_
);
/*ftl::cuda::mls_smooth(
out.createTexture<half4>(_getNormalsChannel()),
temp_.createTexture<half4>(Channel::Normals),
out.getTexture<float>(_getDepthChannel()),
//out.getTexture<float>(_getDepthChannel()),
value("mls_smooth", 0.01f),
value("mls_radius", 2),
params_.camera,
stream_
);
/*ftl::cuda::mls_smooth(
out.createTexture<half4>(_getNormalsChannel()),
temp_.createTexture<half4>(Channel::Normals),
out.getTexture<float>(_getDepthChannel()),
//out.getTexture<float>(_getDepthChannel()),
value("mls_smooth", 0.01f),
value("mls_radius", 2),
params_.camera,
stream_
);
ftl::cuda::mls_smooth(
temp_.createTexture<half4>(Channel::Normals),
out.createTexture<half4>(_getNormalsChannel()),
out.getTexture<float>(_getDepthChannel()),
//out.getTexture<float>(_getDepthChannel()),
value("mls_smooth", 0.01f),
value("mls_radius", 2),
params_.camera,
stream_
);*/
}
ftl::cuda::mls_smooth(
temp_.createTexture<half4>(Channel::Normals),
out.createTexture<half4>(_getNormalsChannel()),
out.getTexture<float>(_getDepthChannel()),
//out.getTexture<float>(_getDepthChannel()),
value("mls_smooth", 0.01f),
value("mls_radius", 2),
params_.camera,
stream_
);*/
}
ftl::cuda::transform_normals(
......
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment