diff --git a/applications/reconstruct/src/ilw.cpp b/applications/reconstruct/src/ilw.cpp
index ae6b6ae4e2ee53a287697b3b5830a2b4ba347e0b..12d1b32f1febb11b6e7a12d365e13627820a977c 100644
--- a/applications/reconstruct/src/ilw.cpp
+++ b/applications/reconstruct/src/ilw.cpp
@@ -3,7 +3,7 @@
using ftl::ILW;
using ftl::detail::ILWData;
-ILW::ILW() {
+ILW::ILW(nlohmann::json &config) : ftl::Configurable(config) {
}
@@ -32,14 +32,19 @@ bool ILW::_phase0(ftl::rgbd::FrameSet &fs) {
// Clear points channel...
// Upload camera data?
+ return true;
}
bool ILW::_phase1(ftl::rgbd::FrameSet &fs) {
// Run correspondence kernel to find points
// For each camera combination
+
+ return true;
}
bool ILW::_phase2(ftl::rgbd::FrameSet &fs) {
// Run energies and motion kernel
+
+ return true;
}
diff --git a/applications/reconstruct/src/main.cpp b/applications/reconstruct/src/main.cpp
index 2af70b290c0849dae87445f36977299f10dbb6ba..d7ce5f9d4f84ee172c23edc9656f1c1d30d7c0b8 100644
--- a/applications/reconstruct/src/main.cpp
+++ b/applications/reconstruct/src/main.cpp
@@ -10,7 +10,7 @@
#include <ftl/configuration.hpp>
#include <ftl/depth_camera.hpp>
#include <ftl/rgbd.hpp>
-#include <ftl/virtual_source.hpp>
+#include <ftl/rgbd/virtual.hpp>
#include <ftl/rgbd/streamer.hpp>
#include <ftl/slave.hpp>
#include <ftl/rgbd/group.hpp>
@@ -97,14 +97,14 @@ static void run(ftl::Configurable *root) {
//ftl::voxhash::SceneRep *scene = ftl::create<ftl::voxhash::SceneRep>(root, "voxelhash");
ftl::rgbd::Streamer *stream = ftl::create<ftl::rgbd::Streamer>(root, "stream", net);
- ftl::rgbd::VirtualSource *virt = ftl::create<ftl::rgbd::VirtualSource>(root, "virtual", net);
- ftl::render::Splatter *splat = ftl::create<ftl::render::Splatter>(root, "renderer", scene);
+ ftl::rgbd::VirtualSource *virt = ftl::create<ftl::rgbd::VirtualSource>(root, "virtual");
+ ftl::render::Splatter *splat = ftl::create<ftl::render::Splatter>(root, "renderer", &scene_B);
ftl::rgbd::Group group;
ftl::ILW *align = ftl::create<ftl::ILW>(root, "merge");
// Generate virtual camera render when requested by streamer
- virt->onRender([splat,&scene_B](ftl::rgbd::Frame &out) {
- splat->render(scene_B, out);
+ virt->onRender([splat,virt](ftl::rgbd::Frame &out) {
+ splat->render(virt, out);
});
stream->add(virt);
@@ -119,7 +119,7 @@ static void run(ftl::Configurable *root) {
bool busy = false;
group.setName("ReconGroup");
- group.sync([splat,virt,&busy,&slave](ftl::rgbd::FrameSet &fs) -> bool {
+ group.sync([splat,virt,&busy,&slave,&scene_A,&scene_B,&align](ftl::rgbd::FrameSet &fs) -> bool {
//cudaSetDevice(scene->getCUDADevice());
if (slave.isPaused()) return true;
@@ -133,14 +133,14 @@ static void run(ftl::Configurable *root) {
// Swap the entire frameset to allow rapid return
fs.swapTo(scene_A);
- ftl::pool.push([&scene_B,&scene_A,&busy,&slave](int id) {
+ ftl::pool.push([&scene_B,&scene_A,&busy,&slave,&align](int id) {
//cudaSetDevice(scene->getCUDADevice());
// TODO: Release frameset here...
//cudaSafeCall(cudaStreamSynchronize(scene->getIntegrationStream()));
// Send all frames to GPU, block until done?
scene_A.upload(Channel::Colour + Channel::Depth); // TODO: (Nick) Add scene stream.
- align.process(scene_A);
+ align->process(scene_A);
// TODO: To use second GPU, could do a download, swap, device change,
// then upload to other device. Or some direct device-2-device copy.
diff --git a/components/renderers/cpp/CMakeLists.txt b/components/renderers/cpp/CMakeLists.txt
index 89cfe508deac5307fd92f1446d87eaeac202071b..b575721587262e2e468d6cb48cf8c44c6771e6fc 100644
--- a/components/renderers/cpp/CMakeLists.txt
+++ b/components/renderers/cpp/CMakeLists.txt
@@ -1,6 +1,7 @@
add_library(ftlrender
src/splat_render.cpp
src/splatter.cu
+ src/points.cu
)
# These cause errors in CI build and are being removed from PCL in newer versions
diff --git a/components/renderers/cpp/include/ftl/cuda/intersections.hpp b/components/renderers/cpp/include/ftl/cuda/intersections.hpp
new file mode 100644
index 0000000000000000000000000000000000000000..9cfdbc2544d9c1bd32c9f5e12a0a161f45c50d54
--- /dev/null
+++ b/components/renderers/cpp/include/ftl/cuda/intersections.hpp
@@ -0,0 +1,88 @@
+#ifndef _FTL_CUDA_INTERSECTIONS_HPP_
+#define _FTL_CUDA_INTERSECTIONS_HPP_
+
+#ifndef PINF
+#define PINF __int_as_float(0x7f800000)
+#endif
+
+namespace ftl {
+namespace cuda {
+
+__device__ inline bool intersectPlane(const float3 &n, const float3 &p0, const float3 &l0, const float3 &l, float &t) {
+ // assuming vectors are all normalized
+ float denom = dot(n, l);
+ if (denom > 1e-6) {
+ t = dot(p0 - l0, n) / denom;
+ return (t >= 0);
+ }
+
+ return false;
+}
+
+__device__ inline bool intersectPlane(const float3 &n, const float3 &p0, const float3 &l, float &t) {
+ // assuming vectors are all normalized
+ float denom = dot(n, l);
+ if (denom > 1e-6) {
+ t = dot(p0, n) / denom;
+ return (t >= 0);
+ }
+ return false;
+}
+
+__device__ inline bool intersectDisk(const float3 &n, const float3 &p0, float radius, const float3 &l0, const float3 &l) {
+ float t = 0;
+ if (intersectPlane(n, p0, l0, l, t)) {
+ float3 p = l0 + l * t;
+ float3 v = p - p0;
+ float d2 = dot(v, v);
+ return (sqrt(d2) <= radius);
+ // or you can use the following optimisation (and precompute radius^2)
+ // return d2 <= radius2; // where radius2 = radius * radius
+ }
+ return false;
+}
+
+/**
+ * Get the radius of a ray intersection with a disk.
+ * @param n Normalised normal of disk.
+ * @param p0 Centre of disk in camera space
+ * @param l Normalised ray direction in camera space
+ * @return Radius from centre of disk where intersection occurred.
+ */
+__device__ inline float intersectDistance(const float3 &n, const float3 &p0, const float3 &l0, const float3 &l) {
+ float t = 0;
+ if (intersectPlane(n, p0, l0, l, t)) {
+ const float3 p = l0 + l * t;
+ const float3 v = p - p0;
+ const float d2 = dot(v, v);
+ return sqrt(d2);
+ // or you can use the following optimisation (and precompute radius^2)
+ // return d2 <= radius2; // where radius2 = radius * radius
+ }
+ return PINF;
+}
+
+/**
+ * Get the radius of a ray intersection with a disk.
+ * @param n Normalised normal of disk.
+ * @param p0 Centre of disk in camera space
+ * @param l Normalised ray direction in camera space
+ * @return Radius from centre of disk where intersection occurred.
+ */
+__device__ inline float intersectDistance(const float3 &n, const float3 &p0, const float3 &l) {
+ float t = 0;
+ if (intersectPlane(n, p0, l, t)) {
+ const float3 p = l * t;
+ const float3 v = p - p0;
+ const float d2 = dot(v, v);
+ return sqrt(d2);
+ // or you can use the following optimisation (and precompute radius^2)
+ // return d2 <= radius2; // where radius2 = radius * radius
+ }
+ return PINF;
+}
+
+}
+}
+
+#endif // _FTL_CUDA_INTERSECTIONS_HPP_
diff --git a/components/renderers/cpp/include/ftl/cuda/points.hpp b/components/renderers/cpp/include/ftl/cuda/points.hpp
new file mode 100644
index 0000000000000000000000000000000000000000..deffe32777789e2b58a96aef2106975ad37e0cdd
--- /dev/null
+++ b/components/renderers/cpp/include/ftl/cuda/points.hpp
@@ -0,0 +1,16 @@
+#ifndef _FTL_CUDA_POINTS_HPP_
+#define _FTL_CUDA_POINTS_HPP_
+
+#include <ftl/cuda_common.hpp>
+#include <ftl/rgbd/camera.hpp>
+#include <ftl/cuda_matrix_util.hpp>
+
+namespace ftl {
+namespace cuda {
+
+void point_cloud(ftl::cuda::TextureObject<float4> &output, ftl::cuda::TextureObject<float> &depth, const ftl::rgbd::Camera ¶ms, const float4x4 &pose, cudaStream_t stream);
+
+}
+}
+
+#endif // _FTL_CUDA_POINTS_HPP_
diff --git a/components/renderers/cpp/include/ftl/render/renderer.hpp b/components/renderers/cpp/include/ftl/render/renderer.hpp
index f564ac7a564c5ec7134f40d22444c983b245601e..1871b9f9f2a8e1fda0766e1c2e74d2169f47f3fa 100644
--- a/components/renderers/cpp/include/ftl/render/renderer.hpp
+++ b/components/renderers/cpp/include/ftl/render/renderer.hpp
@@ -26,7 +26,7 @@ class Renderer : public ftl::Configurable {
* the virtual camera object passed, and writes the result into the
* virtual camera.
*/
- virtual bool render(ftl::rgbd::VirtualSource *, cudaStream_t)=0;
+ virtual bool render(ftl::rgbd::VirtualSource *, ftl::rgbd::Frame &, cudaStream_t)=0;
};
}
diff --git a/components/renderers/cpp/src/splat_params.hpp b/components/renderers/cpp/include/ftl/render/splat_params.hpp
similarity index 100%
rename from components/renderers/cpp/src/splat_params.hpp
rename to components/renderers/cpp/include/ftl/render/splat_params.hpp
diff --git a/components/renderers/cpp/include/ftl/render/splat_render.hpp b/components/renderers/cpp/include/ftl/render/splat_render.hpp
index 89934aa5b8f745cd5d147b7c08dd97d77d98a577..55522c483c25f58d843543ee8d5ba42aae9c32c8 100644
--- a/components/renderers/cpp/include/ftl/render/splat_render.hpp
+++ b/components/renderers/cpp/include/ftl/render/splat_render.hpp
@@ -3,7 +3,7 @@
#include <ftl/render/renderer.hpp>
#include <ftl/rgbd/frameset.hpp>
-#include "splat_params.hpp"
+#include <ftl/render/splat_params.hpp>
namespace ftl {
namespace render {
@@ -18,10 +18,10 @@ namespace render {
*/
class Splatter : public ftl::render::Renderer {
public:
- explicit Splatter(nlohmann::json &config, const ftl::rgbd::FrameSet &fs);
+ explicit Splatter(nlohmann::json &config, ftl::rgbd::FrameSet *fs);
~Splatter();
- bool render(ftl::rgbd::VirtualSource *src, cudaStream_t stream=0) override;
+ bool render(ftl::rgbd::VirtualSource *src, ftl::rgbd::Frame &out, cudaStream_t stream=0) override;
//void setOutputDevice(int);
@@ -36,9 +36,8 @@ class Splatter : public ftl::render::Renderer {
ftl::cuda::TextureObject<float4> normal1_;*/
//SplatParams params_;
- ftl::rgbd::Frame output_;
ftl::rgbd::Frame temp_;
- const ftl::rgbd::FrameSet &scene_;
+ ftl::rgbd::FrameSet *scene_;
};
}
diff --git a/applications/reconstruct/include/ftl/matrix_conversion.hpp b/components/renderers/cpp/include/ftl/utility/matrix_conversion.hpp
similarity index 100%
rename from applications/reconstruct/include/ftl/matrix_conversion.hpp
rename to components/renderers/cpp/include/ftl/utility/matrix_conversion.hpp
diff --git a/components/renderers/cpp/src/points.cu b/components/renderers/cpp/src/points.cu
new file mode 100644
index 0000000000000000000000000000000000000000..39764e4c8aba523caf2758262d9f41f8782ac9dc
--- /dev/null
+++ b/components/renderers/cpp/src/points.cu
@@ -0,0 +1,28 @@
+#include <ftl/cuda/points.hpp>
+
+#define T_PER_BLOCK 8
+
+__global__ void point_cloud_kernel(ftl::cuda::TextureObject<float4> output, ftl::cuda::TextureObject<float> depth, ftl::rgbd::Camera params, float4x4 pose)
+{
+ const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
+ const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
+
+ if (x < params.width && y < params.height) {
+ float d = depth.tex2D((int)x, (int)y);
+
+ output(x,y) = (d >= params.minDepth && d <= params.maxDepth) ?
+ make_float4(pose * params.screenToCam(x, y, d), 0.0f) :
+ make_float4(MINF, MINF, MINF, MINF);
+ }
+}
+
+void ftl::cuda::point_cloud(ftl::cuda::TextureObject<float4> &output, ftl::cuda::TextureObject<float> &depth, const ftl::rgbd::Camera ¶ms, const float4x4 &pose, cudaStream_t stream) {
+ const dim3 gridSize((params.width + T_PER_BLOCK - 1)/T_PER_BLOCK, (params.height + T_PER_BLOCK - 1)/T_PER_BLOCK);
+ const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
+
+ point_cloud_kernel<<<gridSize, blockSize, 0, stream>>>(output, depth, params, pose);
+
+#ifdef _DEBUG
+ cudaSafeCall(cudaDeviceSynchronize());
+#endif
+}
diff --git a/components/renderers/cpp/src/splat_render.cpp b/components/renderers/cpp/src/splat_render.cpp
index dc07c126e036d7e892cccb2f84eb31fa82c39e0d..2097ae284655b6e56155d641d49252f744994bfc 100644
--- a/components/renderers/cpp/src/splat_render.cpp
+++ b/components/renderers/cpp/src/splat_render.cpp
@@ -1,12 +1,16 @@
#include <ftl/render/splat_render.hpp>
+#include <ftl/utility/matrix_conversion.hpp>
#include "splatter_cuda.hpp"
+#include <ftl/cuda/points.hpp>
+
+#include <opencv2/core/cuda_stream_accessor.hpp>
using ftl::render::Splatter;
using ftl::rgbd::Channel;
using ftl::rgbd::Format;
using cv::cuda::GpuMat;
-Splatter::Splatter(nlohmann::json &config, const ftl::rgbd::FrameSet &fs) : ftl::render::Renderer(config), scene_(fs) {
+Splatter::Splatter(nlohmann::json &config, ftl::rgbd::FrameSet *fs) : ftl::render::Renderer(config), scene_(fs) {
}
@@ -14,15 +18,16 @@ Splatter::~Splatter() {
}
-bool Splatter::render(ftl::rgbd::VirtualSource *src, cudaStream_t stream) {
- if (!src->isReady()) return;
+bool Splatter::render(ftl::rgbd::VirtualSource *src, ftl::rgbd::Frame &out, cudaStream_t stream) {
+ if (!src->isReady()) return false;
const auto &camera = src->parameters();
//cudaSafeCall(cudaSetDevice(scene_->getCUDADevice()));
- output_.create<GpuMat>(Channel::Depth, Format<float>(camera.width, camera.height));
- output_.create<GpuMat>(Channel::Colour, Format<uchar4>(camera.width, camera.height));
+ // Create all the required channels
+ out.create<GpuMat>(Channel::Depth, Format<float>(camera.width, camera.height));
+ out.create<GpuMat>(Channel::Colour, Format<uchar4>(camera.width, camera.height));
temp_.create<GpuMat>(Channel::Colour, Format<float4>(camera.width, camera.height));
temp_.create<GpuMat>(Channel::Colour2, Format<uchar4>(camera.width, camera.height));
@@ -62,40 +67,33 @@ bool Splatter::render(ftl::rgbd::VirtualSource *src, cudaStream_t stream) {
if (src->value("splatting", true) == false) params.m_flags |= ftl::render::kNoSplatting;
if (src->value("upsampling", true) == false) params.m_flags |= ftl::render::kNoUpsampling;
if (src->value("texturing", true) == false) params.m_flags |= ftl::render::kNoTexturing;
-
params.m_viewMatrix = MatrixConversion::toCUDA(src->getPose().cast<float>().inverse());
params.m_viewMatrixInverse = MatrixConversion::toCUDA(src->getPose().cast<float>());
- //params.voxelSize = scene_->getHashParams().m_virtualVoxelSize;
params.camera = camera;
- /*params.camera.fx = camera.fx;
- params.camera.fy = camera.fy;
- params.camera.mx = -camera.cx;
- params.camera.my = -camera.cy;
- params.camera.m_imageWidth = camera.width;
- params.camera.m_imageHeight = camera.height;
- params.camera.m_sensorDepthWorldMax = camera.maxDepth;
- params.camera.m_sensorDepthWorldMin = camera.minDepth;*/
-
- //ftl::cuda::compactifyAllocated(scene_->getHashData(), scene_->getHashParams(), stream);
- //LOG(INFO) << "Occupied: " << scene_->getOccupiedCount();
-
- //if (scene_->value("voxels", false)) {
- // TODO:(Nick) Stereo for voxel version
- //ftl::cuda::isosurface_point_image(scene_->getHashData(), depth1_, params, stream);
- //ftl::cuda::splat_points(depth1_, depth2_, params, stream);
- //ftl::cuda::dibr(depth2_, colour1_, scene_->cameraCount(), params, stream);
- //src->writeFrames(ts, colour1_, depth2_, stream);
- //} else {
-
- //ftl::cuda::clear_depth(depth1_, stream);
- //ftl::cuda::clear_depth(depth3_, stream);
- //ftl::cuda::clear_depth(depth2_, stream);
- //ftl::cuda::clear_colour(colour2_, stream);
-
- temp_.get<GpuMat>(Channel::Depth).setTo(cv::Scalar(0), cvstream);
- temp_.get<GpuMat>(Channel::Depth2).setTo(cv::Scalar(0), cvstream);
- output_.get<GpuMat>(Channel::Depth).setTo(cv::Scalar(0.0f), cvstream);
- output_.get<GpuMat>(Channel::Colour).setTo(cv::Scalar(0,0,0), cvstream);
+
+ // Clear all channels to 0 or max depth
+ temp_.get<GpuMat>(Channel::Depth).setTo(cv::Scalar(0x7FFFFFFF), cvstream);
+ temp_.get<GpuMat>(Channel::Depth2).setTo(cv::Scalar(0x7FFFFFFF), cvstream);
+ out.get<GpuMat>(Channel::Depth).setTo(cv::Scalar(1000.0f), cvstream);
+ out.get<GpuMat>(Channel::Colour).setTo(cv::Scalar(0,0,0), cvstream);
+
+ // Render each camera into virtual view
+ for (auto &f : scene_->frames) {
+ // Needs to create points channel first?
+ if (!f.hasChannel(Channel::Points)) {
+ auto &t = f.createTexture<float4>(Channel::Points, Format<float4>(f.get<GpuMat>(Channel::Colour).size()));
+ auto pose = MatrixConversion::toCUDA(f.source()->getPose().cast<float>().inverse());
+ ftl::cuda::point_cloud(t, f.createTexture<float>(Channel::Depth), f.source()->parameters(), pose, stream);
+ }
+
+ ftl::cuda::dibr_merge(
+ f.createTexture<float4>(Channel::Points),
+ temp_.createTexture<int>(Channel::Depth),
+ params, stream
+ );
+ }
+
+ //ftl::cuda::dibr(depth1_, colour1_, normal1_, depth2_, colour_tmp_, depth3_, scene_->cameraCount(), params, stream);
// Step 1: Put all points into virtual view to gather them
//ftl::cuda::dibr_raw(depth1_, scene_->cameraCount(), params, stream);
@@ -109,14 +107,14 @@ bool Splatter::render(ftl::rgbd::VirtualSource *src, cudaStream_t stream) {
//ftl::cuda::splat_points(depth1_, colour1_, normal1_, depth2_, colour2_, params, stream);
//ftl::cuda::int_to_float(depth1_, depth2_, 1.0f / 1000.0f, stream);
- temp_.get<GpuMat>(Channel::Depth).convertTo(output_.get<GpuMat>(Channel::Depth), CV_32F, 1.0f / 1000.0f, cvstream);
+ temp_.get<GpuMat>(Channel::Depth).convertTo(out.get<GpuMat>(Channel::Depth), CV_32F, 1.0f / 1000.0f, cvstream);
//src->writeFrames(ts, colour1_, depth2_, stream);
- src->write(scene_.timestamp, output_, stream);
+ //src->write(scene_.timestamp, output_, stream);
} else {
- temp_.get<GpuMat>(Channel::Depth).convertTo(output_.get<GpuMat>(Channel::Depth), CV_32F, 1.0f / 1000.0f, cvstream);
+ temp_.get<GpuMat>(Channel::Depth).convertTo(out.get<GpuMat>(Channel::Depth), CV_32F, 1.0f / 1000.0f, cvstream);
//ftl::cuda::int_to_float(depth1_, depth2_, 1.0f / 1000.0f, stream);
//src->writeFrames(ts, colour1_, depth2_, stream);
- src->write(scene_.timestamp, output_, stream);
+ //src->write(scene_.timestamp, output_, stream);
}
} else if (src->getChannel() == Channel::Energy) {
//ftl::cuda::int_to_float(depth1_, depth2_, 1.0f / 1000.0f, stream);
@@ -124,7 +122,7 @@ bool Splatter::render(ftl::rgbd::VirtualSource *src, cudaStream_t stream) {
//ftl::cuda::splat_points(depth1_, colour1_, normal1_, depth2_, colour2_, params, stream);
//ftl::cuda::int_to_float(depth1_, depth2_, 1.0f / 1000.0f, stream);
//src->writeFrames(ts, colour1_, depth2_, stream);
- src->write(scene_.timestamp, output_, stream);
+ //src->write(scene_.timestamp, output_, stream);
//} else {
//ftl::cuda::int_to_float(depth1_, depth2_, 1.0f / 1000.0f, stream);
// src->writeFrames(colour1_, depth2_, stream);
@@ -137,19 +135,19 @@ bool Splatter::render(ftl::rgbd::VirtualSource *src, cudaStream_t stream) {
params.m_viewMatrixInverse = MatrixConversion::toCUDA(matrix);
//ftl::cuda::clear_depth(depth1_, stream);
- ftl::cuda::dibr(depth1_, colour1_, normal1_, depth2_, colour_tmp_, depth3_, scene_->cameraCount(), params, stream);
+ //ftl::cuda::dibr(depth1_, colour1_, normal1_, depth2_, colour_tmp_, depth3_, scene_->cameraCount(), params, stream);
//src->writeFrames(ts, colour1_, colour2_, stream);
- src->write(scene_.timestamp, output_, stream);
+ //src->write(scene_.timestamp, output_, stream);
} else {
if (value("splatting", false)) {
//ftl::cuda::splat_points(depth1_, colour1_, normal1_, depth2_, colour2_, params, stream);
//src->writeFrames(ts, colour1_, depth2_, stream);
- src->write(scene_.timestamp, output_, stream);
+ //src->write(scene_.timestamp, out, stream);
} else {
//ftl::cuda::int_to_float(depth1_, depth2_, 1.0f / 1000.0f, stream);
- temp_.get<GpuMat>(Channel::Depth).convertTo(output_.get<GpuMat>(Channel::Depth), CV_32F, 1.0f / 1000.0f, cvstream);
+ temp_.get<GpuMat>(Channel::Depth).convertTo(out.get<GpuMat>(Channel::Depth), CV_32F, 1.0f / 1000.0f, cvstream);
//src->writeFrames(ts, colour1_, depth2_, stream);
- src->write(scene_.timestamp, output_, stream);
+ //src->write(scene_.timestamp, output_, stream);
}
}
//}
@@ -162,6 +160,6 @@ bool Splatter::render(ftl::rgbd::VirtualSource *src, cudaStream_t stream) {
return true;
}
-void Splatter::setOutputDevice(int device) {
- device_ = device;
-}
+//void Splatter::setOutputDevice(int device) {
+// device_ = device;
+//}
diff --git a/components/renderers/cpp/src/splatter.cu b/components/renderers/cpp/src/splatter.cu
index 3c7d94f3918d6b6f36433c6beb28462c5c9ed23a..ad3a6e8a9d1860ea9735cc105ab45a4061c04afb 100644
--- a/components/renderers/cpp/src/splatter.cu
+++ b/components/renderers/cpp/src/splatter.cu
@@ -1,4 +1,5 @@
-#include "splat_params.hpp"
+#include <ftl/render/splat_params.hpp>
+#include "splatter_cuda.hpp"
#include <ftl/rgbd/camera.hpp>
#include <ftl/cuda_common.hpp>
@@ -19,7 +20,7 @@ using ftl::render::SplatParams;
* Pass 1: Directly render each camera into virtual view but with no upsampling
* for sparse points.
*/
- __global__ void dibr_merge_kernel(TextureObject<float4> points, TextureObject<int> depth, int cam, SplatParams params) {
+ __global__ void dibr_merge_kernel(TextureObject<float4> points, TextureObject<int> depth, SplatParams params) {
const int x = blockIdx.x*blockDim.x + threadIdx.x;
const int y = blockIdx.y*blockDim.y + threadIdx.y;
@@ -42,6 +43,14 @@ using ftl::render::SplatParams;
}
}
+void ftl::cuda::dibr_merge(TextureObject<float4> &points, TextureObject<int> &depth, SplatParams params, cudaStream_t stream) {
+ const dim3 gridSize((depth.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (depth.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
+ const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
+
+ dibr_merge_kernel<<<gridSize, blockSize, 0, stream>>>(points, depth, params);
+ cudaSafeCall( cudaGetLastError() );
+}
+
__device__ inline float4 make_float4(const uchar4 &c) {
return make_float4(c.x,c.y,c.z,c.w);
}
diff --git a/components/renderers/cpp/src/splatter_cuda.hpp b/components/renderers/cpp/src/splatter_cuda.hpp
index 125d43390be2514a0e8f301b37d426d3fccb4350..b1b9bbd3db89b19ece3a9379b2fc28f3120a69d4 100644
--- a/components/renderers/cpp/src/splatter_cuda.hpp
+++ b/components/renderers/cpp/src/splatter_cuda.hpp
@@ -2,108 +2,11 @@
#define _FTL_RECONSTRUCTION_SPLAT_CUDA_HPP_
#include <ftl/cuda_common.hpp>
-#include "splat_params.hpp"
+#include <ftl/render/splat_params.hpp>
namespace ftl {
namespace cuda {
-
-__device__ inline bool intersectPlane(const float3 &n, const float3 &p0, const float3 &l0, const float3 &l, float &t) {
- // assuming vectors are all normalized
- float denom = dot(n, l);
- if (denom > 1e-6) {
- t = dot(p0 - l0, n) / denom;
- return (t >= 0);
- }
-
- return false;
-}
-
-__device__ inline bool intersectPlane(const float3 &n, const float3 &p0, const float3 &l, float &t) {
- // assuming vectors are all normalized
- float denom = dot(n, l);
- if (denom > 1e-6) {
- t = dot(p0, n) / denom;
- return (t >= 0);
- }
- return false;
-}
-
-__device__ inline bool intersectDisk(const float3 &n, const float3 &p0, float radius, const float3 &l0, const float3 &l) {
- float t = 0;
- if (intersectPlane(n, p0, l0, l, t)) {
- float3 p = l0 + l * t;
- float3 v = p - p0;
- float d2 = dot(v, v);
- return (sqrt(d2) <= radius);
- // or you can use the following optimisation (and precompute radius^2)
- // return d2 <= radius2; // where radius2 = radius * radius
- }
- return false;
-}
-
-/**
- * Get the radius of a ray intersection with a disk.
- * @param n Normalised normal of disk.
- * @param p0 Centre of disk in camera space
- * @param l Normalised ray direction in camera space
- * @return Radius from centre of disk where intersection occurred.
- */
-__device__ inline float intersectDistance(const float3 &n, const float3 &p0, const float3 &l0, const float3 &l) {
- float t = 0;
- if (intersectPlane(n, p0, l0, l, t)) {
- const float3 p = l0 + l * t;
- const float3 v = p - p0;
- const float d2 = dot(v, v);
- return sqrt(d2);
- // or you can use the following optimisation (and precompute radius^2)
- // return d2 <= radius2; // where radius2 = radius * radius
- }
- return PINF;
-}
-
-/**
- * Get the radius of a ray intersection with a disk.
- * @param n Normalised normal of disk.
- * @param p0 Centre of disk in camera space
- * @param l Normalised ray direction in camera space
- * @return Radius from centre of disk where intersection occurred.
- */
-__device__ inline float intersectDistance(const float3 &n, const float3 &p0, const float3 &l) {
- float t = 0;
- if (intersectPlane(n, p0, l, t)) {
- const float3 p = l * t;
- const float3 v = p - p0;
- const float d2 = dot(v, v);
- return sqrt(d2);
- // or you can use the following optimisation (and precompute radius^2)
- // return d2 <= radius2; // where radius2 = radius * radius
- }
- return PINF;
-}
-
-void splat_points(const ftl::cuda::TextureObject<int> &depth_in,
- const ftl::cuda::TextureObject<uchar4> &colour_in,
- const ftl::cuda::TextureObject<float4> &normal_in,
- const ftl::cuda::TextureObject<float> &depth_out,
- const ftl::cuda::TextureObject<uchar4> &colour_out, const ftl::render::SplatParams ¶ms, cudaStream_t stream);
-
-void dibr(const ftl::cuda::TextureObject<int> &depth_out,
- const ftl::cuda::TextureObject<uchar4> &colour_out,
- const ftl::cuda::TextureObject<float4> &normal_out,
- const ftl::cuda::TextureObject<float> &confidence_out,
- const ftl::cuda::TextureObject<float4> &tmp_colour,
- const ftl::cuda::TextureObject<int> &tmp_depth, int numcams,
- const ftl::render::SplatParams ¶ms, cudaStream_t stream);
-
-/**
- * Directly render input depth maps to virtual view with clipping.
- */
-void dibr_raw(const ftl::cuda::TextureObject<int> &depth_out, int numcams,
- const ftl::render::SplatParams ¶ms, cudaStream_t stream);
-
-void dibr(const ftl::cuda::TextureObject<float> &depth_out,
- const ftl::cuda::TextureObject<uchar4> &colour_out, int numcams, const ftl::render::SplatParams ¶ms, cudaStream_t stream);
-
+ void dibr_merge(ftl::cuda::TextureObject<float4> &points, ftl::cuda::TextureObject<int> &depth, ftl::render::SplatParams params, cudaStream_t stream);
}
}
diff --git a/components/rgbd-sources/CMakeLists.txt b/components/rgbd-sources/CMakeLists.txt
index 3ac2e2de5f173a12bd7107d8d40b5d2eca3b4b5d..2b056d009a73dd7ee82adaedbf03bb98194d636f 100644
--- a/components/rgbd-sources/CMakeLists.txt
+++ b/components/rgbd-sources/CMakeLists.txt
@@ -4,6 +4,7 @@ set(RGBDSRC
src/disparity.cpp
src/source.cpp
src/frame.cpp
+ src/frameset.cpp
src/stereovideo.cpp
src/middlebury_source.cpp
src/net.cpp
@@ -17,6 +18,7 @@ set(RGBDSRC
src/cb_segmentation.cpp
src/abr.cpp
src/offilter.cpp
+ src/virtual.cpp
)
if (HAVE_REALSENSE)
diff --git a/components/rgbd-sources/include/ftl/rgbd/channels.hpp b/components/rgbd-sources/include/ftl/rgbd/channels.hpp
index 7c014e682b5b3a53ef45911088d3a6951e0e9a3e..00d26b93763042cf3ef2db6e32e2db22c0150156 100644
--- a/components/rgbd-sources/include/ftl/rgbd/channels.hpp
+++ b/components/rgbd-sources/include/ftl/rgbd/channels.hpp
@@ -9,18 +9,19 @@ namespace rgbd {
enum struct Channel : int {
None = -1,
- Colour = 0,
+ Colour = 0, // 8UC3 or 8UC4
Left = 0,
- Depth = 1,
- Right = 2,
+ Depth = 1, // 32S or 32F
+ Right = 2, // 8UC3 or 8UC4
Colour2 = 2,
Disparity = 3,
Depth2 = 3,
Deviation = 4,
- Normals,
- Confidence,
- Flow,
- Energy,
+ Normals, // 32FC4
+ Points, // 32FC4
+ Confidence, // 32F
+ Flow, // 32F
+ Energy, // 32F
LeftGray,
RightGray,
Overlay1
diff --git a/components/rgbd-sources/include/ftl/rgbd/frame.hpp b/components/rgbd-sources/include/ftl/rgbd/frame.hpp
index fff3550352da10b2dc04744b58c1c0746ab55803..89523fc9bb1d09f1ed3ca003f234ffb5df2534ec 100644
--- a/components/rgbd-sources/include/ftl/rgbd/frame.hpp
+++ b/components/rgbd-sources/include/ftl/rgbd/frame.hpp
@@ -6,6 +6,7 @@
#include <ftl/exception.hpp>
#include <opencv2/core.hpp>
#include <opencv2/core/cuda.hpp>
+#include <opencv2/core/cuda_stream_accessor.hpp>
#include <ftl/rgbd/channels.hpp>
#include <ftl/rgbd/format.hpp>
@@ -23,22 +24,31 @@ namespace rgbd {
// NN for depth/disparity/optflow, linear/cubic/etc. for RGB
class Frame;
+class Source;
/**
* Manage a set of image channels corresponding to a single camera frame.
*/
class Frame {
public:
- Frame() {}
+ Frame() : src_(nullptr) {}
+ explicit Frame(ftl::rgbd::Source *src) : src_(src) {}
+
+ inline ftl::rgbd::Source *source() const { return src_; }
// Prevent frame copy, instead use a move.
//Frame(const Frame &)=delete;
//Frame &operator=(const Frame &)=delete;
- void download(ftl::rgbd::Channel c, cv::cuda::Stream& stream=cv::cuda::Stream::Null());
- void upload(ftl::rgbd::Channel c, cv::cuda::Stream& stream=cv::cuda::Stream::Null());
- void download(ftl::rgbd::Channels c, cv::cuda::Stream& stream=cv::cuda::Stream::Null());
- void upload(ftl::rgbd::Channels c, cv::cuda::Stream& stream=cv::cuda::Stream::Null());
+ void download(ftl::rgbd::Channel c, cv::cuda::Stream stream);
+ void upload(ftl::rgbd::Channel c, cv::cuda::Stream stream);
+ void download(ftl::rgbd::Channels c, cv::cuda::Stream stream);
+ void upload(ftl::rgbd::Channels c, cv::cuda::Stream stream);
+
+ inline void download(ftl::rgbd::Channel c, cudaStream_t stream=0) { download(c, cv::cuda::StreamAccessor::wrapStream(stream)); };
+ inline void upload(ftl::rgbd::Channel c, cudaStream_t stream=0) { upload(c, cv::cuda::StreamAccessor::wrapStream(stream)); };
+ inline void download(ftl::rgbd::Channels c, cudaStream_t stream=0) { download(c, cv::cuda::StreamAccessor::wrapStream(stream)); };
+ inline void upload(ftl::rgbd::Channels c, cudaStream_t stream=0) { upload(c, cv::cuda::StreamAccessor::wrapStream(stream)); };
/**
* Perform a buffer swap of the selected channels. This is intended to be
@@ -138,6 +148,8 @@ private:
ftl::rgbd::Channels channels_; // Does it have a channel
ftl::rgbd::Channels gpu_; // Is the channel on a GPU
+ ftl::rgbd::Source *src_;
+
inline ChannelData &_get(ftl::rgbd::Channel c) { return data_[static_cast<unsigned int>(c)]; }
inline const ChannelData &_get(ftl::rgbd::Channel c) const { return data_[static_cast<unsigned int>(c)]; }
};
diff --git a/components/rgbd-sources/include/ftl/rgbd/frameset.hpp b/components/rgbd-sources/include/ftl/rgbd/frameset.hpp
index 4831d1cb852f492e77b515f3ea7a25286480f621..2fa39e2eacf19339860e98fa98df44f687ac64c7 100644
--- a/components/rgbd-sources/include/ftl/rgbd/frameset.hpp
+++ b/components/rgbd-sources/include/ftl/rgbd/frameset.hpp
@@ -1,6 +1,7 @@
#ifndef _FTL_RGBD_FRAMESET_HPP_
#define _FTL_RGBD_FRAMESET_HPP_
+#include <ftl/threads.hpp>
#include <ftl/rgbd/frame.hpp>
#include <opencv2/opencv.hpp>
diff --git a/components/rgbd-sources/include/ftl/rgbd/source.hpp b/components/rgbd-sources/include/ftl/rgbd/source.hpp
index 5d296e699bf897c243ccfe9371635ff45e2010cd..0ee163add0009023ec24e6df6bd18a1da927af1e 100644
--- a/components/rgbd-sources/include/ftl/rgbd/source.hpp
+++ b/components/rgbd-sources/include/ftl/rgbd/source.hpp
@@ -26,6 +26,7 @@ namespace rgbd {
static inline bool isValidDepth(float d) { return (d > 0.01f) && (d < 39.99f); }
class SnapshotReader;
+class VirtualSource;
/**
* RGBD Generic data source configurable entity. This class hides the
@@ -40,6 +41,7 @@ class Source : public ftl::Configurable {
public:
template <typename T, typename... ARGS>
friend T *ftl::config::create(ftl::config::json_t &, ARGS ...);
+ friend class VirtualSource;
//template <typename T, typename... ARGS>
//friend T *ftl::config::create(ftl::Configurable *, const std::string &, ARGS ...);
@@ -51,7 +53,7 @@ class Source : public ftl::Configurable {
Source(const Source&)=delete;
Source &operator=(const Source&) =delete;
- private:
+ protected:
explicit Source(ftl::config::json_t &cfg);
Source(ftl::config::json_t &cfg, ftl::rgbd::SnapshotReader *);
Source(ftl::config::json_t &cfg, ftl::net::Universe *net);
diff --git a/components/rgbd-sources/include/ftl/rgbd/virtual.hpp b/components/rgbd-sources/include/ftl/rgbd/virtual.hpp
index 55dfb9e8950be9bfe86917fb79f92c7d7f8e730a..f0ab3a93bdb0c7bbf81fe61b4724000f3b020b31 100644
--- a/components/rgbd-sources/include/ftl/rgbd/virtual.hpp
+++ b/components/rgbd-sources/include/ftl/rgbd/virtual.hpp
@@ -11,13 +11,15 @@ class VirtualSource : public ftl::rgbd::Source {
explicit VirtualSource(ftl::config::json_t &cfg);
~VirtualSource();
+ void onRender(const std::function<void(ftl::rgbd::Frame &)> &);
+
/**
* Write frames into source buffers from an external renderer. Virtual
* sources do not have an internal generator of frames but instead have
* their data provided from an external rendering class. This function only
* works when there is no internal generator.
*/
- void write(int64_t ts, ftl::rgbd::Frame &frame, cudaStream_t stream=0);
+ //void write(int64_t ts, ftl::rgbd::Frame &frame, cudaStream_t stream=0);
};
}
diff --git a/components/rgbd-sources/src/frame.cpp b/components/rgbd-sources/src/frame.cpp
index f8a4b348a167d6b4f372f0c8fa5563676513e48d..ea207404b7213d147062bee328e5a48b3be24cff 100644
--- a/components/rgbd-sources/src/frame.cpp
+++ b/components/rgbd-sources/src/frame.cpp
@@ -13,15 +13,15 @@ void Frame::reset() {
gpu_.clear();
}
-void Frame::download(Channel c, cv::cuda::Stream& stream) {
+void Frame::download(Channel c, cv::cuda::Stream stream) {
download(Channels(c), stream);
}
-void Frame::upload(Channel c, cv::cuda::Stream& stream) {
+void Frame::upload(Channel c, cv::cuda::Stream stream) {
upload(Channels(c), stream);
}
-void Frame::download(Channels c, cv::cuda::Stream& stream) {
+void Frame::download(Channels c, cv::cuda::Stream stream) {
for (size_t i=0u; i<Channels::kMax; ++i) {
if (c.has(i) && channels_.has(i) && gpu_.has(i)) {
data_[i].gpu.download(data_[i].host, stream);
@@ -30,7 +30,7 @@ void Frame::download(Channels c, cv::cuda::Stream& stream) {
}
}
-void Frame::upload(Channels c, cv::cuda::Stream& stream) {
+void Frame::upload(Channels c, cv::cuda::Stream stream) {
for (size_t i=0u; i<Channels::kMax; ++i) {
if (c.has(i) && channels_.has(i) && !gpu_.has(i)) {
data_[i].gpu.upload(data_[i].host, stream);
diff --git a/components/rgbd-sources/src/frameset.cpp b/components/rgbd-sources/src/frameset.cpp
index 4b830cf7c24a59b5ec89bcabec287a1c446c42db..72cb5cb50af72dc035d3b113269ed4af38c2d5b9 100644
--- a/components/rgbd-sources/src/frameset.cpp
+++ b/components/rgbd-sources/src/frameset.cpp
@@ -27,12 +27,12 @@ void FrameSet::swapTo(ftl::rgbd::FrameSet &fs) {
}
fs.timestamp = timestamp;
- fs.count = count;
+ fs.count = static_cast<int>(count);
fs.stale = stale;
- fs.mask = mask;
+ fs.mask = static_cast<unsigned int>(mask);
for (size_t i=0; i<frames.size(); ++i) {
- frames[i].swap(Channels::All(), fs.frames[i]);
+ frames[i].swapTo(Channels::All(), fs.frames[i]);
}
stale = true;
diff --git a/components/rgbd-sources/src/virtual.cpp b/components/rgbd-sources/src/virtual.cpp
index 9d64f53c73b2902d8fe074d84a8192657c36d1fc..63ee5ddbb712db1d9306f27a14e1165edb6576ef 100644
--- a/components/rgbd-sources/src/virtual.cpp
+++ b/components/rgbd-sources/src/virtual.cpp
@@ -1,3 +1,20 @@
+#include <ftl/rgbd/virtual.hpp>
+
+using ftl::rgbd::VirtualSource;
+using ftl::rgbd::Source;
+
+VirtualSource::VirtualSource(ftl::config::json_t &cfg) : Source(cfg) {
+
+}
+
+VirtualSource::~VirtualSource() {
+
+}
+
+void VirtualSource::onRender(const std::function<void(ftl::rgbd::Frame &)> &f) {
+
+}
+
/*
void Source::writeFrames(int64_t ts, const cv::Mat &rgb, const cv::Mat &depth) {
if (!impl_) {