diff --git a/components/operators/CMakeLists.txt b/components/operators/CMakeLists.txt
index 9d6793817d5d20c5312e4c94a01fca1fe4a1d374..23a80cb7c0fd20bfd9d63842a467051d737f415c 100644
--- a/components/operators/CMakeLists.txt
+++ b/components/operators/CMakeLists.txt
@@ -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
diff --git a/components/operators/include/ftl/operators/cuda/carver.hpp b/components/operators/include/ftl/operators/cuda/carver.hpp
new file mode 100644
index 0000000000000000000000000000000000000000..1dd0c04a884610211ca7e6ffa8c19702ff1f25cb
--- /dev/null
+++ b/components/operators/include/ftl/operators/cuda/carver.hpp
@@ -0,0 +1,34 @@
+#ifndef _FTL_CUDA_CARVER_HPP_
+#define _FTL_CUDA_CARVER_HPP_
+
+#include <ftl/cuda_common.hpp>
+#include <ftl/cuda_matrix_util.hpp>
+#include <ftl/rgbd/camera.hpp>
+
+namespace ftl {
+namespace cuda {
+
+/**
+ * Carve `in` using `ref` as visibility reference.
+ */
+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 float4x4 &transform,
+ const ftl::rgbd::Camera &incam,
+ const ftl::rgbd::Camera &refcam,
+ cudaStream_t stream);
+
+void apply_colour_scaling(
+ const cv::cuda::GpuMat &scale,
+ cv::cuda::GpuMat &colour,
+ int radius,
+ cudaStream_t stream);
+
+}
+}
+
+#endif
\ No newline at end of file
diff --git a/components/renderers/cpp/src/carver.cu b/components/operators/src/fusion/carving/carver.cu
similarity index 50%
rename from components/renderers/cpp/src/carver.cu
rename to components/operators/src/fusion/carving/carver.cu
index f804c1e271eed66889511856569fd391bfb3deae..d1dd480ef0cbbf2ea4a75595fd6af4fee712297d 100644
--- a/components/renderers/cpp/src/carver.cu
+++ b/components/operators/src/fusion/carving/carver.cu
@@ -1,4 +1,4 @@
-#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 ¢roid_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 ¢roid,
- 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>
diff --git a/components/operators/src/fusion/fusion.cpp b/components/operators/src/fusion/fusion.cpp
index 489edb4c2e704bd1f7a38ae86d036dbd4358fbd2..334820a3a12d590bbfcdee51390b2fef1f6b41a7 100644
--- a/components/operators/src/fusion/fusion.cpp
+++ b/components/operators/src/fusion/fusion.cpp
@@ -1,4 +1,5 @@
#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) {
diff --git a/components/renderers/cpp/CMakeLists.txt b/components/renderers/cpp/CMakeLists.txt
index 4b3c7e2d72ed9c50ed347b2801b6992e18ed08c4..fe93b78bb84000a5a499eda8dd136a5f7c153fae 100644
--- a/components/renderers/cpp/CMakeLists.txt
+++ b/components/renderers/cpp/CMakeLists.txt
@@ -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
)
diff --git a/components/renderers/cpp/src/CUDARender.cpp b/components/renderers/cpp/src/CUDARender.cpp
index da271fa404b7e85d72edacd6d408325cf776ddb7..0f888dec2b9b75b1de77c791f9295d1563c113b5 100644
--- a/components/renderers/cpp/src/CUDARender.cpp
+++ b/components/renderers/cpp/src/CUDARender.cpp
@@ -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(
diff --git a/components/renderers/cpp/src/carver.hpp b/components/renderers/cpp/src/carver.hpp
deleted file mode 100644
index 3f9a66ec567b20ca6ae71869e6c2d6086644ce93..0000000000000000000000000000000000000000
--- a/components/renderers/cpp/src/carver.hpp
+++ /dev/null
@@ -1,60 +0,0 @@
-#ifndef _FTL_CUDA_CARVER_HPP_
-#define _FTL_CUDA_CARVER_HPP_
-
-#include <ftl/cuda_common.hpp>
-#include <ftl/cuda_matrix_util.hpp>
-#include <ftl/rgbd/camera.hpp>
-
-namespace ftl {
-namespace cuda {
-
-/**
- * Carve `in` using `ref` as visibility reference.
- */
-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 float4x4 &transform,
- const ftl::rgbd::Camera &incam,
- const ftl::rgbd::Camera &refcam,
- cudaStream_t stream);
-
-void apply_colour_scaling(
- const cv::cuda::GpuMat &scale,
- cv::cuda::GpuMat &colour,
- int radius,
- cudaStream_t stream);
-
-void mls_reduce(
- const cv::cuda::GpuMat ¢roid,
- 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 ¢roid_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
-);
-
-}
-}
-
-#endif
\ No newline at end of file