diff --git a/components/renderers/cpp/CMakeLists.txt b/components/renderers/cpp/CMakeLists.txt
index fe93b78bb84000a5a499eda8dd136a5f7c153fae..4b3c7e2d72ed9c50ed347b2801b6992e18ed08c4 100644
--- a/components/renderers/cpp/CMakeLists.txt
+++ b/components/renderers/cpp/CMakeLists.txt
@@ -13,6 +13,7 @@ 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 d44a8cfd102b4a1bf90d250ded75a543f09e01b2..0e16d901377019275954e146b82d7e6bd77a0268 100644
--- a/components/renderers/cpp/src/CUDARender.cpp
+++ b/components/renderers/cpp/src/CUDARender.cpp
@@ -7,6 +7,7 @@
#include <ftl/render/colouriser.hpp>
#include <ftl/cuda/transform.hpp>
#include <ftl/operators/cuda/smoothing_cuda.hpp>
+#include "carver.hpp"
#include <ftl/cuda/colour_cuda.hpp>
@@ -254,6 +255,57 @@ 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", true)) {
+ 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),
+ transformR,
+ f.getLeft(),
+ ref.getLeft(),
+ stream_
+ );
+ }
+ }
+ }
+
// For each source depth map
for (size_t i=0; i < scene_->frames.size(); ++i) {
//if (!scene_->hasFrame(i)) continue;
@@ -365,7 +417,7 @@ void CUDARender::_mesh(ftl::rgbd::Frame &out, const Eigen::Matrix4d &t, cudaStre
ftl::cuda::touch_merge(depth_out_, out.createTexture<float>(_getDepthChannel()), collisions_, 1024, touch_dist_, stream_);
// For each source depth map, verify results
- if (value("carve_result", true)) {
+ if (value("post_carve_result", false)) {
for (size_t i=0; i < scene_->frames.size(); ++i) {
//if (!scene_->hasFrame(i)) continue;
auto &f = scene_->frames[i].cast<ftl::rgbd::Frame>();
@@ -383,13 +435,12 @@ void CUDARender::_mesh(ftl::rgbd::Frame &out, const Eigen::Matrix4d &t, cudaStre
//auto pose = MatrixConversion::toCUDA(t.cast<float>() * f.getPose().cast<float>());
auto transformR = MatrixConversion::toCUDA(f.getPose().cast<float>().inverse() * t.cast<float>().inverse()) * poseInverse_;
- auto transform = pose_ * MatrixConversion::toCUDA(t.cast<float>() * f.getPose().cast<float>());
+ //auto transform = pose_ * MatrixConversion::toCUDA(t.cast<float>() * f.getPose().cast<float>());
- ftl::cuda::reverse_verify(
+ ftl::cuda::depth_carve(
out.create<cv::cuda::GpuMat>(_getDepthChannel()),
f.get<cv::cuda::GpuMat>(Channel::Depth),
transformR,
- transform,
params_.camera,
f.getLeft(),
stream_
diff --git a/components/renderers/cpp/src/carver.cu b/components/renderers/cpp/src/carver.cu
new file mode 100644
index 0000000000000000000000000000000000000000..19435598cee8fdf5c70cec4e3b1344f2752ad36c
--- /dev/null
+++ b/components/renderers/cpp/src/carver.cu
@@ -0,0 +1,47 @@
+#include "carver.hpp"
+
+// ==== Reverse Verify Result ==================================================
+
+__global__ void reverse_check_kernel(
+ float* __restrict__ depth_in,
+ const float* __restrict__ depth_original,
+ int pitch4,
+ int opitch4,
+ float4x4 transformR,
+ ftl::rgbd::Camera vintrin,
+ ftl::rgbd::Camera ointrin
+) {
+ const int x = blockIdx.x*blockDim.x + threadIdx.x;
+ const int y = blockIdx.y*blockDim.y + threadIdx.y;
+
+ if (x < 0 || x >= vintrin.width || y < 0 || y >= vintrin.height) return;
+
+ float d = depth_in[y*pitch4+x];
+
+ // FIXME: This is dangerous, need to check through alternates instead
+ while (true) {
+ float3 campos = transformR * vintrin.screenToCam(x,y,d);
+ int2 spos = ointrin.camToScreen<int2>(campos);
+ int ox = spos.x;
+ int oy = spos.y;
+
+ if (campos.z > 0.0f && ox >= 0 && ox < ointrin.width && oy >= 0 && oy < ointrin.height) {
+ float d2 = depth_original[oy*opitch4+ox];
+ if (!(d2 < ointrin.maxDepth && d2 - campos.z > d2*0.001f)) break;
+ d += 0.002f;
+ } else break;
+ }
+
+ depth_in[y*pitch4+x] = d;
+}
+
+void ftl::cuda::depth_carve(cv::cuda::GpuMat &depth_in, const cv::cuda::GpuMat &depth_original, const float4x4 &transformR, const ftl::rgbd::Camera &vintrin, const ftl::rgbd::Camera &ointrin, cudaStream_t stream) {
+ static constexpr int THREADS_X = 8;
+ static constexpr int THREADS_Y = 8;
+
+ 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);
+
+ reverse_check_kernel<<<gridSize, blockSize, 0, stream>>>(depth_in.ptr<float>(), depth_original.ptr<float>(), depth_in.step1(), depth_original.step1(), transformR, vintrin, ointrin);
+ cudaSafeCall( cudaGetLastError() );
+}
diff --git a/components/renderers/cpp/src/carver.hpp b/components/renderers/cpp/src/carver.hpp
new file mode 100644
index 0000000000000000000000000000000000000000..f3e6f671a531aa921c0c61ae81d3f407b5a7519d
--- /dev/null
+++ b/components/renderers/cpp/src/carver.hpp
@@ -0,0 +1,25 @@
+#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 float4x4 &transform,
+ const ftl::rgbd::Camera &incam,
+ const ftl::rgbd::Camera &refcam,
+ cudaStream_t stream);
+
+}
+}
+
+#endif
\ No newline at end of file
diff --git a/components/renderers/cpp/src/triangle_render.cu b/components/renderers/cpp/src/triangle_render.cu
index 4c8770bdbaeabbfa8e2288c1d15a3f7a3aa0e27c..301cf2a226d0cff47c9a1672fe7d78b124f3fc99 100644
--- a/components/renderers/cpp/src/triangle_render.cu
+++ b/components/renderers/cpp/src/triangle_render.cu
@@ -224,53 +224,6 @@ void ftl::cuda::merge_convert_depth(TextureObject<int> &depth_in, TextureObject<
cudaSafeCall( cudaGetLastError() );
}
-// ==== Reverse Verify Result ==================================================
-
-__global__ void reverse_check_kernel(
- float* __restrict__ depth_in,
- const float* __restrict__ depth_original,
- int pitch4,
- int opitch4,
- float4x4 transformR,
- float4x4 transform,
- ftl::rgbd::Camera vintrin,
- ftl::rgbd::Camera ointrin
-) {
- const int x = blockIdx.x*blockDim.x + threadIdx.x;
- const int y = blockIdx.y*blockDim.y + threadIdx.y;
-
- if (x < 0 || x >= vintrin.width || y < 0 || y >= vintrin.height) return;
-
- float d = depth_in[y*pitch4+x];
-
- // FIXME: This is dangerous, need to check through alternates instead
- while (true) {
- float3 campos = transformR * vintrin.screenToCam(x,y,d);
- int2 spos = ointrin.camToScreen<int2>(campos);
- int ox = spos.x;
- int oy = spos.y;
-
- if (campos.z > 0.0f && ox >= 0 && ox < ointrin.width && oy >= 0 && oy < ointrin.height) {
- float d2 = depth_original[oy*opitch4+ox];
- if (!(d2 < ointrin.maxDepth && d2 - campos.z > d2*0.001f)) break;
- d += 0.002f;
- } else break;
- }
-
- depth_in[y*pitch4+x] = d;
-}
-
-void ftl::cuda::reverse_verify(cv::cuda::GpuMat &depth_in, const cv::cuda::GpuMat &depth_original, const float4x4 &transformR, const float4x4 &transform, const ftl::rgbd::Camera &vintrin, const ftl::rgbd::Camera &ointrin, cudaStream_t stream) {
- static constexpr int THREADS_X = 8;
- static constexpr int THREADS_Y = 8;
-
- 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);
-
- reverse_check_kernel<<<gridSize, blockSize, 0, stream>>>(depth_in.ptr<float>(), depth_original.ptr<float>(), depth_in.step1(), depth_original.step1(), transformR, transform, vintrin, ointrin);
- cudaSafeCall( cudaGetLastError() );
-}
-
// ==== BLENDER ========
/*