diff --git a/CMakeLists.txt b/CMakeLists.txt
index 1991d5345a5c0b567b5d69897fdf57b3c6689d7d..c4192b464b45274d3bccf77f5719df45e7b677dd 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -225,6 +225,23 @@ find_library(UUID_LIBRARIES NAMES uuid libuuid)
else()
endif()
+# For screen capture
+if (WIN32)
+ message(STATUS "Screen capture not supported")
+else()
+ find_package(X11)
+ if (X11_FOUND)
+ if (X11_XShm_FOUND)
+ message(STATUS "Using X11 for screen capture")
+ set(HAVE_X11 TRUE)
+ else()
+ message(STATUS "No X11 Shared memory extension")
+ endif()
+ else()
+ message(STATUS "No X11, screen capture disabled")
+ endif()
+endif()
+
# For ftl2mkv
check_include_file("libavformat/avformat.h" HAVE_AVFORMAT)
diff --git a/applications/gui/src/camera.cpp b/applications/gui/src/camera.cpp
index 615cdd86f06390bc405a1d0c7d51ddc2b82de939..a6f578179789ec079e2f58722257b645e30dccbf 100644
--- a/applications/gui/src/camera.cpp
+++ b/applications/gui/src/camera.cpp
@@ -141,7 +141,9 @@ static Eigen::Affine3d create_rotation_matrix(float ax, float ay, float az) {
return rz * rx * ry;
}
-ftl::gui::Camera::Camera(ftl::gui::Screen *screen, int fsid, int fid, ftl::codecs::Channel c) : screen_(screen), fsid_(fsid), fid_(fid), channel_(c),channels_(0u) {
+static int vcamcount = 0;
+
+ftl::gui::Camera::Camera(ftl::gui::Screen *screen, int fsmask, int fid, ftl::codecs::Channel c) : screen_(screen), fsmask_(fsmask), fid_(fid), channel_(c),channels_(0u) {
eye_ = Eigen::Vector3d(0.0f, 0.0f, 0.0f);
neye_ = Eigen::Vector4d(0.0f, 0.0f, 0.0f, 0.0f);
rotmat_.setIdentity();
@@ -170,6 +172,7 @@ ftl::gui::Camera::Camera(ftl::gui::Screen *screen, int fsid, int fid, ftl::codec
renderer_ = nullptr;
post_pipe_ = nullptr;
record_stream_ = nullptr;
+ transform_ix_ = -1;
/*src->setCallback([this](int64_t ts, ftl::rgbd::Frame &frame) {
UNIQUE_LOCK(mutex_, lk);
@@ -189,35 +192,29 @@ ftl::gui::Camera::Camera(ftl::gui::Screen *screen, int fsid, int fid, ftl::codec
}
});*/
- auto *host = screen->root();
+ //auto *host = screen->root();
// Is virtual camera?
if (fid == 255) {
- state_.getLeft().width = host->value("width", 1280);
- state_.getLeft().height = host->value("height", 720);
- state_.getLeft().fx = host->value("focal", 700.0f);
- state_.getLeft().fy = state_.getLeft().fx;
- state_.getLeft().cx = -(double)state_.getLeft().width / 2.0;
- state_.getLeft().cy = -(double)state_.getLeft().height / 2.0;
- state_.getLeft().minDepth = host->value("minDepth", 0.1f);
- state_.getLeft().maxDepth = host->value("maxDepth", 15.0f);
- state_.getLeft().doffs = 0;
- state_.getLeft().baseline = host->value("baseline", 0.05f);
-
- state_.getRight().width = host->value("width", 1280);
- state_.getRight().height = host->value("height", 720);
- state_.getRight().fx = host->value("focal_right", 700.0f);
- state_.getRight().fy = state_.getRight().fx;
- state_.getRight().cx = host->value("centre_x_right", -(double)state_.getLeft().width / 2.0);
- state_.getRight().cy = host->value("centre_y_right", -(double)state_.getLeft().height / 2.0);
- state_.getRight().minDepth = host->value("minDepth", 0.1f);
- state_.getRight().maxDepth = host->value("maxDepth", 15.0f);
- state_.getRight().doffs = 0;
- state_.getRight().baseline = host->value("baseline", 0.05f);
+ renderer_ = ftl::create<ftl::render::CUDARender>(screen_->root(), std::string("vcam")+std::to_string(vcamcount++));
+ // Allow mask to be changed
+ fsmask_ = renderer_->value("fsmask", fsmask_);
+ renderer_->on("fsmask", [this](const ftl::config::Event &e) {
+ fsmask_ = renderer_->value("fsmask", fsmask_);
+ });
+
+ intrinsics_ = ftl::create<ftl::Configurable>(renderer_, "intrinsics");
+
+ state_.getLeft() = ftl::rgbd::Camera::from(intrinsics_);
+ state_.getRight() = state_.getLeft();
Eigen::Matrix4d pose;
pose.setIdentity();
state_.setPose(pose);
+
+ for (auto &t : transforms_) {
+ t.setIdentity();
+ }
}
}
@@ -228,13 +225,14 @@ ftl::gui::Camera::~Camera() {
void ftl::gui::Camera::draw(std::vector<ftl::rgbd::FrameSet*> &fss) {
if (fid_ != 255) return;
- if (fsid_ >= fss.size()) return;
+ //if (fsid_ >= fss.size()) return;
+
+ //auto &fs = *fss[fsid_];
- auto &fs = *fss[fsid_];
-
- UNIQUE_LOCK(fs.mtx,lk);
UNIQUE_LOCK(mutex_, lk2);
- _draw(fs);
+ //state_.getLeft().fx = intrinsics_->value("focal", 700.0f);
+ //state_.getLeft().fy = state_.getLeft().fx;
+ _draw(fss);
for (auto *fset : fss) {
for (const auto &f : fset->frames) {
@@ -264,16 +262,19 @@ void ftl::gui::Camera::draw(std::vector<ftl::rgbd::FrameSet*> &fss) {
}
}
-void ftl::gui::Camera::_draw(ftl::rgbd::FrameSet &fs) {
+void ftl::gui::Camera::_draw(std::vector<ftl::rgbd::FrameSet*> &fss) {
frame_.reset();
frame_.setOrigin(&state_);
- if (!renderer_) renderer_ = ftl::create<ftl::render::Triangular>(screen_->root(), "vcam1");
- Eigen::Matrix4d t;
- t.setIdentity();
- renderer_->render(fs, frame_, channel_, t);
- // TODO: Insert post-render pipeline.
- // FXAA + Bad colour removal
+ renderer_->begin(frame_);
+ for (auto *fs : fss) {
+ if (!usesFrameset(fs->id)) continue;
+
+ // FIXME: Should perhaps remain locked until after end is called?
+ UNIQUE_LOCK(fs->mtx,lk);
+ renderer_->submit(fs, Channels<0>(channel_), transforms_[fs->id]);
+ }
+ renderer_->end();
if (!post_pipe_) {
post_pipe_ = ftl::config::create<ftl::operators::Graph>(screen_->root(), "post_filters");
@@ -289,14 +290,17 @@ void ftl::gui::Camera::_draw(ftl::rgbd::FrameSet &fs) {
over_col.create(im1_.size(), CV_8UC4);
over_depth.create(im1_.size(), CV_32F);
- for (int i=0; i<fs.frames.size(); ++i) {
- auto pose = fs.frames[i].getPose().inverse() * state_.getPose();
- Eigen::Vector4d pos = pose.inverse() * Eigen::Vector4d(0,0,0,1);
- pos /= pos[3];
+ for (auto *fs : fss) {
+ if (!usesFrameset(fs->id)) continue;
+ for (int i=0; i<fs->frames.size(); ++i) {
+ auto pose = fs->frames[i].getPose().inverse() * state_.getPose();
+ Eigen::Vector4d pos = pose.inverse() * Eigen::Vector4d(0,0,0,1);
+ pos /= pos[3];
- auto name = fs.frames[i].get<std::string>("name");
- ftl::overlay::drawCamera(state_.getLeft(), im1_, over_depth, fs.frames[i].getLeftCamera(), pose, cv::Scalar(0,0,255,255), 0.2,screen_->root()->value("show_frustrum", false));
- if (name) ftl::overlay::drawText(state_.getLeft(), im1_, over_depth, *name, pos, 0.5, cv::Scalar(0,0,255,255));
+ auto name = fs->frames[i].get<std::string>("name");
+ ftl::overlay::drawCamera(state_.getLeft(), im1_, over_depth, fs->frames[i].getLeftCamera(), pose, cv::Scalar(0,0,255,255), 0.2,screen_->root()->value("show_frustrum", false));
+ if (name) ftl::overlay::drawText(state_.getLeft(), im1_, over_depth, *name, pos, 0.5, cv::Scalar(0,0,255,255));
+ }
}
}
@@ -308,7 +312,7 @@ void ftl::gui::Camera::_draw(ftl::rgbd::FrameSet &fs) {
fs2.mask = 1;
fs2.stale = false;
frame_.swapTo(Channels<0>(Channel::Colour), f); // Channel::Colour + Channel::Depth
- fs2.timestamp = fs.timestamp;
+ fs2.timestamp = ftl::timer::get_time();
fs2.id = 0;
record_sender_->post(fs2);
record_stream_->select(0, Channels<0>(Channel::Colour));
@@ -341,27 +345,30 @@ void ftl::gui::Camera::_downloadFrames(ftl::rgbd::Frame *frame) {
void ftl::gui::Camera::update(std::vector<ftl::rgbd::FrameSet*> &fss) {
UNIQUE_LOCK(mutex_, lk);
- if (fss.size() <= fsid_) return;
- auto &fs = *fss[fsid_];
-
- ftl::rgbd::Frame *frame = nullptr;
-
+ //if (fss.size() <= fsid_) return;
if (fid_ == 255) {
name_ = "Virtual Camera";
// Do a draw if not active. If active the draw function will be called
// directly.
if (screen_->activeCamera() != this) {
- _draw(fs);
+ _draw(fss);
}
} else {
- if (fid_ >= fs.frames.size()) return;
- frame = &fs.frames[fid_];
- _downloadFrames(frame);
- auto n = frame->get<std::string>("name");
- if (n) {
- name_ = *n;
- } else {
- name_ = "No name";
+ for (auto *fs : fss) {
+ if (!usesFrameset(fs->id)) continue;
+
+ ftl::rgbd::Frame *frame = nullptr;
+
+ if (fid_ >= fs->frames.size()) return;
+ frame = &fs->frames[fid_];
+ _downloadFrames(frame);
+ auto n = frame->get<std::string>("name");
+ if (n) {
+ name_ = *n;
+ } else {
+ name_ = "No name";
+ }
+ return;
}
}
}
@@ -424,6 +431,10 @@ void ftl::gui::Camera::keyMovement(int key, int modifiers) {
float scalar = (key == 266) ? -mag : mag;
neye_ += rotmat_*Eigen::Vector4d(0.0,scalar,0.0,1.0);
return;
+ } else if (key >= '0' && key <= '5') {
+ int ix = key - (int)('0');
+ transform_ix_ = ix-1;
+ return;
}
}
@@ -623,7 +634,13 @@ const GLTexture &ftl::gui::Camera::captureFrame() {
Eigen::Affine3d t(trans);
Eigen::Matrix4d viewPose = t.matrix() * rotmat_;
- if (isVirtual()) state_.setPose(viewPose);
+ if (isVirtual()) {
+ if (transform_ix_ < 0) {
+ state_.setPose(viewPose);
+ } else {
+ transforms_[transform_ix_] = viewPose;
+ }
+ }
//src_->grab();
diff --git a/applications/gui/src/camera.hpp b/applications/gui/src/camera.hpp
index 15f373e9bc9834c9f1f16856c167c40e4eb8a7ee..b92f241541c7af9e8a2f284eab6cf60cdf55bf56 100644
--- a/applications/gui/src/camera.hpp
+++ b/applications/gui/src/camera.hpp
@@ -2,7 +2,7 @@
#define _FTL_GUI_CAMERA_HPP_
#include <ftl/rgbd/frameset.hpp>
-#include <ftl/render/tri_render.hpp>
+#include <ftl/render/CUDARender.hpp>
#include <ftl/codecs/writer.hpp>
#include "gltexture.hpp"
@@ -10,6 +10,7 @@
#include <ftl/streams/sender.hpp>
#include <string>
+#include <array>
#ifdef HAVE_OPENVR
#include <openvr/openvr.h>
@@ -25,7 +26,7 @@ class PoseWindow;
class Camera {
public:
- Camera(ftl::gui::Screen *screen, int fsid, int fid, ftl::codecs::Channel chan=ftl::codecs::Channel::Colour);
+ Camera(ftl::gui::Screen *screen, int fsmask, int fid, ftl::codecs::Channel chan=ftl::codecs::Channel::Colour);
~Camera();
Camera(const Camera &)=delete;
@@ -33,7 +34,9 @@ class Camera {
int width() const { return width_; }
int height() const { return height_; }
- int getFramesetId() const { return fsid_; }
+ int getFramesetMask() const { return fsmask_; }
+
+ bool usesFrameset(int id) const { return fsmask_ & (1 << id); }
void setPose(const Eigen::Matrix4d &p);
@@ -93,7 +96,7 @@ class Camera {
cv::Mat visualizeActiveChannel();
Screen *screen_;
- int fsid_;
+ unsigned int fsmask_; // Frameset Mask
int fid_;
int width_;
@@ -119,7 +122,8 @@ class Camera {
cv::Mat im1_; // first channel (left)
cv::Mat im2_; // second channel ("right")
- ftl::render::Triangular *renderer_;
+ ftl::render::CUDARender *renderer_;
+ ftl::Configurable *intrinsics_;
ftl::operators::Graph *post_pipe_;
ftl::rgbd::Frame frame_;
ftl::rgbd::FrameState state_;
@@ -128,6 +132,9 @@ class Camera {
std::string name_;
+ int transform_ix_;
+ std::array<Eigen::Matrix4d,ftl::stream::kMaxStreams> transforms_; // Frameset transforms for virtual cam
+
MUTEX mutex_;
#ifdef HAVE_OPENVR
@@ -137,7 +144,7 @@ class Camera {
#endif
void _downloadFrames(ftl::rgbd::Frame *frame);
- void _draw(ftl::rgbd::FrameSet &fs);
+ void _draw(std::vector<ftl::rgbd::FrameSet*> &fss);
};
}
diff --git a/applications/gui/src/screen.cpp b/applications/gui/src/screen.cpp
index 9a997d8403702bfab5fc2a26cf41bea4d1726cf5..70a79163b4a8ab7716df5bfbcf1846805301e962 100644
--- a/applications/gui/src/screen.cpp
+++ b/applications/gui/src/screen.cpp
@@ -463,7 +463,7 @@ bool ftl::gui::Screen::keyboardEvent(int key, int scancode, int action, int modi
} else {
//LOG(INFO) << "Key press " << key << " - " << action << " - " << modifiers;
- if (key >= 262 && key <= 267) {
+ if ((key >= 262 && key <= 267) || (key >= '0' && key <= '9')) {
if (camera_) camera_->keyMovement(key, modifiers);
return true;
} else if (action == 1 && key == 'H') {
diff --git a/applications/gui/src/src_window.cpp b/applications/gui/src/src_window.cpp
index 2ec91f8e41387f57fc3d4a8f163fb33863ba1e3f..c8adf88195352bd3e1ab543a9855048eed0e8501 100644
--- a/applications/gui/src/src_window.cpp
+++ b/applications/gui/src/src_window.cpp
@@ -109,7 +109,7 @@ SourceWindow::SourceWindow(ftl::gui::Screen *screen)
}*/
const auto *cstream = interceptor_;
- _createDefaultCameras(*framesets_[fs.id], cstream->available(fs.id).has(Channel::Depth));
+ _createDefaultCameras(*framesets_[fs.id], true); // cstream->available(fs.id).has(Channel::Depth)
//LOG(INFO) << "Channels = " << (unsigned int)cstream->available(fs.id);
@@ -151,7 +151,7 @@ SourceWindow::SourceWindow(ftl::gui::Screen *screen)
// Check paths for FTL files to load.
auto paths = (*screen->root()->get<nlohmann::json>("paths"));
- int ftl_count = 0;
+ int ftl_count = available_.size();
for (auto &x : paths.items()) {
std::string path = x.value().get<std::string>();
auto eix = path.find_last_of('.');
@@ -205,7 +205,7 @@ void SourceWindow::stopRecordingVideo() {
ftl::codecs::Channels<0> SourceWindow::_aggregateChannels(int id) {
ftl::codecs::Channels<0> cs = ftl::codecs::Channels<0>(Channel::Colour);
for (auto cam : cameras_) {
- if (cam.second.camera->getFramesetId() == id) {
+ if (cam.second.camera->usesFrameset(id)) {
if (cam.second.camera->isVirtual()) {
cs += Channel::Depth;
} else {
@@ -215,6 +215,7 @@ ftl::codecs::Channels<0> SourceWindow::_aggregateChannels(int id) {
}
}
}
+
return cs;
}
@@ -222,7 +223,7 @@ void SourceWindow::_createDefaultCameras(ftl::rgbd::FrameSet &fs, bool makevirtu
for (int i=0; i<fs.frames.size(); ++i) {
int id = (fs.id << 8) + i;
if (cameras_.find(id) == cameras_.end()) {
- auto *cam = new ftl::gui::Camera(screen_, fs.id, i);
+ auto *cam = new ftl::gui::Camera(screen_, 1 << fs.id, i);
cameras_[id] = {
cam,
nullptr
@@ -231,7 +232,7 @@ void SourceWindow::_createDefaultCameras(ftl::rgbd::FrameSet &fs, bool makevirtu
}
if (makevirtual && cameras_.find((fs.id << 8) + 255) == cameras_.end()) {
- auto *cam = new ftl::gui::Camera(screen_, fs.id, 255);
+ auto *cam = new ftl::gui::Camera(screen_, 1 << fs.id, 255);
cameras_[(fs.id << 8) + 255] = {
cam,
nullptr
diff --git a/applications/reconstruct/src/main.cpp b/applications/reconstruct/src/main.cpp
index 0b5a71d7d8a15ee025e5135d4f703df528f2963d..3e35b5ad96eec52c81a26dd50d6e66e8ac52405c 100644
--- a/applications/reconstruct/src/main.cpp
+++ b/applications/reconstruct/src/main.cpp
@@ -19,9 +19,6 @@
#include "reconstruction.hpp"
-#include "ilw/ilw.hpp"
-#include <ftl/render/tri_render.hpp>
-
#include <fstream>
#include <string>
#include <vector>
diff --git a/applications/reconstruct/src/reconstruction.cpp b/applications/reconstruct/src/reconstruction.cpp
index 524128dd5af15784609c89fd1f3c5681a566bb54..9ae89fc0d15fb791cb7c7b35ba9fe979b0757e4b 100644
--- a/applications/reconstruct/src/reconstruction.cpp
+++ b/applications/reconstruct/src/reconstruction.cpp
@@ -128,7 +128,7 @@ void Reconstruction::addRawCallback(const std::function<void(ftl::rgbd::Source *
group_->addRawCallback(cb);
}*/
-bool Reconstruction::render(ftl::rgbd::VirtualSource *vs, ftl::rgbd::Frame &out) {
+/*bool Reconstruction::render(ftl::rgbd::VirtualSource *vs, ftl::rgbd::Frame &out) {
{
UNIQUE_LOCK(exchange_mtx_, lk);
if (new_frame_) {
@@ -136,11 +136,7 @@ bool Reconstruction::render(ftl::rgbd::VirtualSource *vs, ftl::rgbd::Frame &out)
fs_align_.swapTo(fs_render_);
}
}
- /*if (fs_render_.stale || fs_render_.timestamp <= 0) {
- LOG(ERROR) << "STALE FRAME TO RENDER";
- return false;
- }
- fs_render_.stale = true;*/
+
// Create scene transform, intended for axis aligning the walls and floor
Eigen::Matrix4d transform;
@@ -166,4 +162,4 @@ bool Reconstruction::render(ftl::rgbd::VirtualSource *vs, ftl::rgbd::Frame &out)
bool res = false; //renderer_->render(vs, out, sm.matrix() * transform);
//fs_render_.resetFull();
return res;
-}
\ No newline at end of file
+}*/
\ No newline at end of file
diff --git a/applications/reconstruct/src/reconstruction.hpp b/applications/reconstruct/src/reconstruction.hpp
index 82cecd9ed8fa3466c15f1608c039dd8f2a9f8a8a..4e545cf8520158e23891d6d37819f354768c95b9 100644
--- a/applications/reconstruct/src/reconstruction.hpp
+++ b/applications/reconstruct/src/reconstruction.hpp
@@ -7,7 +7,6 @@
#include "ftl/rgbd/group.hpp"
#include "ftl/rgbd/frameset.hpp"
#include "ftl/operators/operator.hpp"
-#include "ftl/render/tri_render.hpp"
namespace ftl {
@@ -22,11 +21,6 @@ class Reconstruction : public ftl::Configurable, public ftl::rgbd::Generator {
void setGenerator(ftl::rgbd::Generator *);
- /**
- * Do the render for a specified virtual camera.
- */
- bool render(ftl::rgbd::VirtualSource *vs, ftl::rgbd::Frame &out);
-
/** Number of frames in last frameset. This can change over time. */
size_t size() override;
@@ -51,7 +45,6 @@ class Reconstruction : public ftl::Configurable, public ftl::rgbd::Generator {
ftl::rgbd::FrameSet fs_align_;
ftl::rgbd::Generator *gen_;
ftl::operators::Graph *pipeline_;
- ftl::render::Triangular *renderer_;
ftl::rgbd::VideoCallback cb_;
diff --git a/components/common/cpp/include/ftl/config.h.in b/components/common/cpp/include/ftl/config.h.in
index f36fe3b73b61d3d32b5a2ee564430cf56d7eec89..bcbebd958a9b6f1c6aff58f9c0907a882905ae15 100644
--- a/components/common/cpp/include/ftl/config.h.in
+++ b/components/common/cpp/include/ftl/config.h.in
@@ -25,6 +25,7 @@
#cmakedefine HAVE_OPENVR
#cmakedefine HAVE_NVPIPE
#cmakedefine HAVE_PORTAUDIO
+#cmakedefine HAVE_X11
extern const char *FTL_BRANCH;
extern const char *FTL_VERSION_LONG;
diff --git a/components/renderers/cpp/CMakeLists.txt b/components/renderers/cpp/CMakeLists.txt
index 0d292b397c07d5406ab68e9d8084fb27cc9129b7..66a930c2c515c818010ddae5557ad976a9ca2d70 100644
--- a/components/renderers/cpp/CMakeLists.txt
+++ b/components/renderers/cpp/CMakeLists.txt
@@ -6,7 +6,7 @@ add_library(ftlrender
src/screen.cu
src/triangle_render.cu
src/reprojection.cu
- src/tri_render.cpp
+ src/CUDARender.cpp
)
target_include_directories(ftlrender PUBLIC
diff --git a/components/renderers/cpp/include/ftl/render/tri_render.hpp b/components/renderers/cpp/include/ftl/render/CUDARender.hpp
similarity index 74%
rename from components/renderers/cpp/include/ftl/render/tri_render.hpp
rename to components/renderers/cpp/include/ftl/render/CUDARender.hpp
index 8e47d80e0d2e0a74f6ca9f9a1b9c7be4997f4f24..616f78a199a4fa0e36ced654e7c18b8a3452c450 100644
--- a/components/renderers/cpp/include/ftl/render/tri_render.hpp
+++ b/components/renderers/cpp/include/ftl/render/CUDARender.hpp
@@ -1,10 +1,11 @@
-#ifndef _FTL_RECONSTRUCTION_TRI_HPP_
-#define _FTL_RECONSTRUCTION_TRI_HPP_
+#ifndef _FTL_RENDER_CUDA_HPP_
+#define _FTL_RENDER_CUDA_HPP_
#include <ftl/render/renderer.hpp>
#include <ftl/rgbd/frameset.hpp>
#include <ftl/render/splat_params.hpp>
#include <ftl/cuda/points.hpp>
+#include <ftl/codecs/channels.hpp>
//#include <ftl/filters/filter.hpp>
namespace ftl {
@@ -14,12 +15,15 @@ namespace render {
* Generate triangles between connected points and render those. Colour is done
* by weighted reprojection to the original source images.
*/
-class Triangular : public ftl::render::Renderer {
+class CUDARender : public ftl::render::Renderer {
public:
- explicit Triangular(nlohmann::json &config);
- ~Triangular();
+ explicit CUDARender(nlohmann::json &config);
+ ~CUDARender();
- bool render(ftl::rgbd::FrameSet &in, ftl::rgbd::Frame &out, ftl::codecs::Channel, const Eigen::Matrix4d &t) override;
+ void begin(ftl::rgbd::Frame &) override;
+ void end() override;
+
+ bool submit(ftl::rgbd::FrameSet *in, ftl::codecs::Channels<0>, const Eigen::Matrix4d &t) override;
//void setOutputDevice(int);
protected:
@@ -29,6 +33,7 @@ class Triangular : public ftl::render::Renderer {
int device_;
ftl::rgbd::Frame temp_;
ftl::rgbd::Frame accum_;
+ ftl::rgbd::Frame *out_;
ftl::rgbd::FrameSet *scene_;
ftl::cuda::ClipSpace clip_;
bool clipping_;
@@ -51,6 +56,14 @@ class Triangular : public ftl::render::Renderer {
//ftl::Filters *filters_;
+ struct SubmitState {
+ ftl::rgbd::FrameSet *fs;
+ ftl::codecs::Channels<0> channels;
+ Eigen::Matrix4d transform;
+ };
+
+ std::vector<SubmitState> sets_;
+
template <typename T>
void __reprojectChannel(ftl::rgbd::Frame &, ftl::codecs::Channel in, ftl::codecs::Channel out, const Eigen::Matrix4d &t, cudaStream_t);
void _reprojectChannel(ftl::rgbd::Frame &, ftl::codecs::Channel in, ftl::codecs::Channel out, const Eigen::Matrix4d &t, cudaStream_t);
@@ -65,7 +78,8 @@ class Triangular : public ftl::render::Renderer {
void _renderDensity(ftl::rgbd::Frame &out, const Eigen::Matrix4d &t);
void _renderRight(ftl::rgbd::Frame &out, const Eigen::Matrix4d &t);
void _renderSecond(ftl::rgbd::Frame &out, ftl::codecs::Channel chan, const Eigen::Matrix4d &t);
- void _render(ftl::rgbd::FrameSet &in, ftl::rgbd::Frame &out, ftl::codecs::Channel, const Eigen::Matrix4d &t);
+ void _renderPass1(const Eigen::Matrix4d &t);
+ void _renderPass2(ftl::codecs::Channels<0>, const Eigen::Matrix4d &t);
bool _alreadySeen() const { return last_frame_ == scene_->timestamp; }
};
@@ -73,4 +87,4 @@ class Triangular : public ftl::render::Renderer {
}
}
-#endif // _FTL_RECONSTRUCTION_TRI_HPP_
+#endif // _FTL_RENDER_CUDA_HPP_
diff --git a/components/renderers/cpp/include/ftl/render/renderer.hpp b/components/renderers/cpp/include/ftl/render/renderer.hpp
index 3a36655147fc58a34c2b5605f28f949d93238fab..ec479ce51e3b494b4df6818303c1aeb5db480bad 100644
--- a/components/renderers/cpp/include/ftl/render/renderer.hpp
+++ b/components/renderers/cpp/include/ftl/render/renderer.hpp
@@ -21,11 +21,26 @@ class Renderer : public ftl::Configurable {
explicit Renderer(nlohmann::json &config) : Configurable(config) {};
virtual ~Renderer() {};
+ /**
+ * Begin a new render. This clears memory, allocates buffers etc. The RGBD
+ * frame given as parameter is where the output channels are rendered to.
+ */
+ virtual void begin(ftl::rgbd::Frame &)=0;
+
+ /**
+ * Finish a render. Post process the output as required, or finish
+ * generating it from internal buffers. The output frame is only valid
+ * after this is called.
+ */
+ virtual void end()=0;
+
/**
- * Generate a single virtual frame. The frame takes its pose and calibration
- * from the output frame pose and calibration channels.
+ * Render all frames of a frameset into the output frame. This can be called
+ * multiple times between `begin` and `end` to combine multiple framesets.
+ * Note that the frameset pointer must remain valid until `end` is called,
+ * and ideally should not be swapped between
*/
- virtual bool render(ftl::rgbd::FrameSet &, ftl::rgbd::Frame &, ftl::codecs::Channel, const Eigen::Matrix4d &)=0;
+ virtual bool submit(ftl::rgbd::FrameSet *, ftl::codecs::Channels<0>, const Eigen::Matrix4d &)=0;
};
}
diff --git a/components/renderers/cpp/include/ftl/render/splat_render.hpp b/components/renderers/cpp/include/ftl/render/splat_render.hpp
deleted file mode 100644
index 8e51aadf15b8e32d05b3253e72cb2adc7e69f98b..0000000000000000000000000000000000000000
--- a/components/renderers/cpp/include/ftl/render/splat_render.hpp
+++ /dev/null
@@ -1,67 +0,0 @@
-#ifndef _FTL_RECONSTRUCTION_SPLAT_HPP_
-#define _FTL_RECONSTRUCTION_SPLAT_HPP_
-
-#include <ftl/render/renderer.hpp>
-#include <ftl/rgbd/frameset.hpp>
-#include <ftl/render/splat_params.hpp>
-#include <ftl/cuda/points.hpp>
-
-namespace ftl {
-namespace render {
-
-/**
- * Render the voxel hash structure by generating image points for surface
- * voxels and expanding those into interpolated splats. This is a two pass
- * algorithm with the option of completing the second pass on a separate GPU.
- * It also possible to only complete the first pass and perform the second step
- * on a separate machine or at a later time, the advantage being to save local
- * processing resources and that the first pass result may compress better.
- */
-class Splatter : public ftl::render::Renderer {
- public:
- explicit Splatter(nlohmann::json &config, ftl::rgbd::FrameSet *fs);
- ~Splatter();
-
- bool render(ftl::rgbd::VirtualSource *src, ftl::rgbd::Frame &out, const Eigen::Matrix4d &t) override;
- //void setOutputDevice(int);
-
- protected:
- void _renderChannel(ftl::rgbd::Frame &out, ftl::codecs::Channel channel_in, ftl::codecs::Channel channel_out, cudaStream_t stream);
-
- private:
- int device_;
- /*ftl::cuda::TextureObject<int> depth1_;
- ftl::cuda::TextureObject<int> depth3_;
- ftl::cuda::TextureObject<uchar4> colour1_;
- ftl::cuda::TextureObject<float4> colour_tmp_;
- ftl::cuda::TextureObject<float> depth2_;
- ftl::cuda::TextureObject<uchar4> colour2_;
- ftl::cuda::TextureObject<float4> normal1_;*/
- //SplatParams params_;
-
- ftl::rgbd::Frame temp_;
- ftl::rgbd::Frame accum_;
- ftl::rgbd::FrameSet *scene_;
- ftl::cuda::ClipSpace clip_;
- bool clipping_;
- float norm_filter_;
- bool backcull_;
- cv::Scalar background_;
- bool splat_;
- float3 light_dir_;
- uchar4 light_diffuse_;
- uchar4 light_ambient_;
- ftl::render::SplatParams params_;
- cudaStream_t stream_;
- float3 light_pos_;
-
- template <typename T>
- void __blendChannel(ftl::rgbd::Frame &, ftl::codecs::Channel in, ftl::codecs::Channel out, cudaStream_t);
- void _blendChannel(ftl::rgbd::Frame &, ftl::codecs::Channel in, ftl::codecs::Channel out, cudaStream_t);
- void _dibr(cudaStream_t);
-};
-
-}
-}
-
-#endif // _FTL_RECONSTRUCTION_SPLAT_HPP_
diff --git a/components/renderers/cpp/src/tri_render.cpp b/components/renderers/cpp/src/CUDARender.cpp
similarity index 78%
rename from components/renderers/cpp/src/tri_render.cpp
rename to components/renderers/cpp/src/CUDARender.cpp
index 5120746fc12d32ed60fd2fed28bc4420d966e548..a40c792251dfce60f88ee0f6883307b7fa727cc7 100644
--- a/components/renderers/cpp/src/tri_render.cpp
+++ b/components/renderers/cpp/src/CUDARender.cpp
@@ -1,4 +1,4 @@
-#include <ftl/render/tri_render.hpp>
+#include <ftl/render/CUDARender.hpp>
#include <ftl/utility/matrix_conversion.hpp>
#include "splatter_cuda.hpp"
#include <ftl/cuda/points.hpp>
@@ -16,7 +16,7 @@
#include <string>
-using ftl::render::Triangular;
+using ftl::render::CUDARender;
using ftl::codecs::Channel;
using ftl::codecs::Channels;
using ftl::rgbd::Format;
@@ -72,7 +72,7 @@ static uchar4 parseCUDAColour(const std::string &colour) {
return make_uchar4(0,0,0,0);
}
-Triangular::Triangular(nlohmann::json &config) : ftl::render::Renderer(config), scene_(nullptr) {
+CUDARender::CUDARender(nlohmann::json &config) : ftl::render::Renderer(config), scene_(nullptr) {
/*if (config["clipping"].is_object()) {
auto &c = config["clipping"];
float rx = c.value("pitch", 0.0f);
@@ -155,7 +155,7 @@ Triangular::Triangular(nlohmann::json &config) : ftl::render::Renderer(config),
last_frame_ = -1;
}
-Triangular::~Triangular() {
+CUDARender::~CUDARender() {
}
@@ -173,66 +173,12 @@ struct AccumSelector<float> {
//static constexpr cv::Scalar value = cv::Scalar(0.0f);
};
-/*template <typename T>
-void Triangular::__blendChannel(ftl::rgbd::Frame &output, ftl::codecs::Channel in, ftl::codecs::Channel out, cudaStream_t stream) {
- cv::cuda::Stream cvstream = cv::cuda::StreamAccessor::wrapStream(stream);
- temp_.create<GpuMat>(
- AccumSelector<T>::channel,
- Format<typename AccumSelector<T>::type>(params_.camera.width, params_.camera.height)
- ).setTo(cv::Scalar(0.0f), cvstream);
- temp_.get<GpuMat>(Channel::Contribution).setTo(cv::Scalar(0.0f), cvstream);
-
- temp_.createTexture<float>(Channel::Contribution);
-
- for (auto &f : scene_->frames) {
- if (f.get<GpuMat>(in).type() == CV_8UC3) {
- // Convert to 4 channel colour
- auto &col = f.get<GpuMat>(in);
- GpuMat tmp(col.size(), CV_8UC4);
- cv::cuda::swap(col, tmp);
- cv::cuda::cvtColor(tmp,col, cv::COLOR_BGR2BGRA);
- }
-
- ftl::cuda::dibr_attribute(
- f.createTexture<T>(in),
- f.createTexture<float4>(Channel::Points),
- temp_.getTexture<int>(Channel::Depth2),
- temp_.createTexture<typename AccumSelector<T>::type>(AccumSelector<T>::channel),
- temp_.getTexture<float>(Channel::Contribution),
- params_, stream
- );
- }
-
- ftl::cuda::dibr_normalise(
- temp_.getTexture<typename AccumSelector<T>::type>(AccumSelector<T>::channel),
- output.createTexture<T>(out),
- temp_.getTexture<float>(Channel::Contribution),
- stream
- );
-}*/
-
template <typename T>
-void Triangular::__reprojectChannel(ftl::rgbd::Frame &output, ftl::codecs::Channel in, ftl::codecs::Channel out, const Eigen::Matrix4d &t, cudaStream_t stream) {
+void CUDARender::__reprojectChannel(ftl::rgbd::Frame &output, ftl::codecs::Channel in, ftl::codecs::Channel out, const Eigen::Matrix4d &t, cudaStream_t stream) {
cv::cuda::Stream cvstream = cv::cuda::StreamAccessor::wrapStream(stream);
- temp_.create<GpuMat>(
- AccumSelector<T>::channel,
- Format<typename AccumSelector<T>::type>(params_.camera.width, params_.camera.height)
- ).setTo(cv::Scalar(0.0f), cvstream);
- temp_.get<GpuMat>(Channel::Contribution).setTo(cv::Scalar(0.0f), cvstream);
-
- temp_.createTexture<float>(Channel::Contribution);
for (size_t i=0; i < scene_->frames.size(); ++i) {
auto &f = scene_->frames[i];
- //auto *s = scene_->sources[i];
-
- /*if (f.get<GpuMat>(in).type() == CV_8UC3) {
- // Convert to 4 channel colour
- auto &col = f.get<GpuMat>(in);
- GpuMat tmp(col.size(), CV_8UC4);
- cv::cuda::swap(col, tmp);
- cv::cuda::cvtColor(tmp,col, cv::COLOR_BGR2BGRA);
- }*/
if (!f.hasChannel(in)) {
LOG(ERROR) << "Reprojecting unavailable channel";
@@ -243,52 +189,59 @@ void Triangular::__reprojectChannel(ftl::rgbd::Frame &output, ftl::codecs::Chann
auto transformR = MatrixConversion::toCUDA(f.getPose().cast<float>().inverse()).getFloat3x3();
if (mesh_) {
- ftl::cuda::reproject(
- f.createTexture<T>(in),
- f.createTexture<float>(Channel::Depth),
- output.getTexture<float>(Channel::Depth),
- output.getTexture<float4>(Channel::Normals),
- temp_.createTexture<typename AccumSelector<T>::type>(AccumSelector<T>::channel),
- temp_.getTexture<float>(Channel::Contribution),
- params_,
- f.getLeftCamera(),
- transform, transformR, stream
- );
+ if (f.hasChannel(Channel::Depth)) {
+ ftl::cuda::reproject(
+ f.createTexture<T>(in),
+ f.createTexture<float>(Channel::Depth),
+ output.getTexture<float>(Channel::Depth),
+ output.getTexture<float4>(Channel::Normals),
+ temp_.createTexture<typename AccumSelector<T>::type>(AccumSelector<T>::channel),
+ temp_.getTexture<float>(Channel::Contribution),
+ params_,
+ f.getLeftCamera(),
+ transform, transformR, stream
+ );
+ } else {
+ // Reproject without depth channel or normals
+ ftl::cuda::reproject(
+ f.createTexture<T>(in),
+ output.getTexture<float>(Channel::Depth),
+ temp_.createTexture<typename AccumSelector<T>::type>(AccumSelector<T>::channel),
+ temp_.getTexture<float>(Channel::Contribution),
+ params_,
+ f.getLeftCamera(),
+ transform, stream
+ );
+ }
} else {
// Can't use normals with point cloud version
- ftl::cuda::reproject(
- f.createTexture<T>(in),
- f.createTexture<float>(Channel::Depth),
- output.getTexture<float>(Channel::Depth),
- temp_.createTexture<typename AccumSelector<T>::type>(AccumSelector<T>::channel),
- temp_.getTexture<float>(Channel::Contribution),
- params_,
- f.getLeftCamera(),
- transform, stream
- );
+ if (f.hasChannel(Channel::Depth)) {
+ ftl::cuda::reproject(
+ f.createTexture<T>(in),
+ f.createTexture<float>(Channel::Depth),
+ output.getTexture<float>(Channel::Depth),
+ temp_.createTexture<typename AccumSelector<T>::type>(AccumSelector<T>::channel),
+ temp_.getTexture<float>(Channel::Contribution),
+ params_,
+ f.getLeftCamera(),
+ transform, stream
+ );
+ } else {
+ ftl::cuda::reproject(
+ f.createTexture<T>(in),
+ output.getTexture<float>(Channel::Depth),
+ temp_.createTexture<typename AccumSelector<T>::type>(AccumSelector<T>::channel),
+ temp_.getTexture<float>(Channel::Contribution),
+ params_,
+ f.getLeftCamera(),
+ transform, stream
+ );
+ }
}
}
-
- ftl::cuda::dibr_normalise(
- temp_.getTexture<typename AccumSelector<T>::type>(AccumSelector<T>::channel),
- output.createTexture<T>(out),
- temp_.getTexture<float>(Channel::Contribution),
- stream
- );
}
-/*void Triangular::_blendChannel(ftl::rgbd::Frame &output, ftl::codecs::Channel in, ftl::codecs::Channel out, cudaStream_t stream) {
- int type = output.get<GpuMat>(out).type(); // == CV_32F; //ftl::rgbd::isFloatChannel(channel);
-
- switch (type) {
- case CV_32F : __blendChannel<float>(output, in, out, stream); break;
- case CV_32FC4 : __blendChannel<float4>(output, in, out, stream); break;
- case CV_8UC4 : __blendChannel<uchar4>(output, in, out, stream); break;
- default : LOG(ERROR) << "Invalid output channel format";
- }
-}*/
-
-void Triangular::_reprojectChannel(ftl::rgbd::Frame &output, ftl::codecs::Channel in, ftl::codecs::Channel out, const Eigen::Matrix4d &t, cudaStream_t stream) {
+void CUDARender::_reprojectChannel(ftl::rgbd::Frame &output, ftl::codecs::Channel in, ftl::codecs::Channel out, const Eigen::Matrix4d &t, cudaStream_t stream) {
int type = output.get<GpuMat>(out).type(); // == CV_32F; //ftl::rgbd::isFloatChannel(channel);
switch (type) {
@@ -299,7 +252,7 @@ void Triangular::_reprojectChannel(ftl::rgbd::Frame &output, ftl::codecs::Channe
}
}
-void Triangular::_dibr(ftl::rgbd::Frame &out, const Eigen::Matrix4d &t, cudaStream_t stream) {
+void CUDARender::_dibr(ftl::rgbd::Frame &out, const Eigen::Matrix4d &t, cudaStream_t stream) {
cv::cuda::Stream cvstream = cv::cuda::StreamAccessor::wrapStream(stream);
temp_.get<GpuMat>(Channel::Depth2).setTo(cv::Scalar(0x7FFFFFFF), cvstream);
@@ -307,27 +260,36 @@ void Triangular::_dibr(ftl::rgbd::Frame &out, const Eigen::Matrix4d &t, cudaStre
auto &f = scene_->frames[i];
//auto *s = scene_->sources[i];
- if (f.empty(Channel::Depth + Channel::Colour)) {
+ if (f.empty(Channel::Colour)) {
LOG(ERROR) << "Missing required channel";
continue;
}
auto transform = params_.m_viewMatrix * MatrixConversion::toCUDA(t.cast<float>() * f.getPose().cast<float>());
- ftl::cuda::dibr_merge(
- f.createTexture<float>(Channel::Depth),
- temp_.createTexture<int>(Channel::Depth2),
- transform,
- f.getLeftCamera(),
- params_, stream
- );
+ if (f.hasChannel(Channel::Depth)) {
+ ftl::cuda::dibr_merge(
+ f.createTexture<float>(Channel::Depth),
+ temp_.createTexture<int>(Channel::Depth2),
+ transform,
+ f.getLeftCamera(),
+ params_, stream
+ );
+ } else {
+ ftl::cuda::dibr_merge(
+ temp_.createTexture<int>(Channel::Depth2),
+ transform,
+ f.getLeftCamera(),
+ params_, stream
+ );
+ }
}
// Convert from int depth to float depth
temp_.get<GpuMat>(Channel::Depth2).convertTo(out.get<GpuMat>(Channel::Depth), CV_32F, 1.0f / 100000.0f, cvstream);
}
-void Triangular::_mesh(ftl::rgbd::Frame &out, const Eigen::Matrix4d &t, cudaStream_t stream) {
+void CUDARender::_mesh(ftl::rgbd::Frame &out, const Eigen::Matrix4d &t, cudaStream_t stream) {
cv::cuda::Stream cvstream = cv::cuda::StreamAccessor::wrapStream(stream);
bool do_blend = value("mesh_blend", true);
@@ -344,7 +306,7 @@ void Triangular::_mesh(ftl::rgbd::Frame &out, const Eigen::Matrix4d &t, cudaStre
auto &f = scene_->frames[i];
//auto *s = scene_->sources[i];
- if (f.empty(Channel::Depth + Channel::Colour)) {
+ if (f.empty(Channel::Colour)) {
LOG(ERROR) << "Missing required channel";
continue;
}
@@ -353,12 +315,21 @@ void Triangular::_mesh(ftl::rgbd::Frame &out, const Eigen::Matrix4d &t, cudaStre
auto transform = params_.m_viewMatrix * MatrixConversion::toCUDA(t.cast<float>() * f.getPose().cast<float>());
// Calculate and save virtual view screen position of each source pixel
- ftl::cuda::screen_coord(
- f.createTexture<float>(Channel::Depth),
- f.createTexture<float>(Channel::Depth2, Format<float>(f.get<GpuMat>(Channel::Depth).size())),
- f.createTexture<short2>(Channel::Screen, Format<short2>(f.get<GpuMat>(Channel::Depth).size())),
- params_, transform, f.getLeftCamera(), stream
- );
+ if (f.hasChannel(Channel::Depth)) {
+ ftl::cuda::screen_coord(
+ f.createTexture<float>(Channel::Depth),
+ f.createTexture<float>(Channel::Depth2, Format<float>(f.get<GpuMat>(Channel::Depth).size())),
+ f.createTexture<short2>(Channel::Screen, Format<short2>(f.get<GpuMat>(Channel::Depth).size())),
+ params_, transform, f.getLeftCamera(), stream
+ );
+ } else {
+ // Constant depth version
+ ftl::cuda::screen_coord(
+ f.createTexture<float>(Channel::Depth2, Format<float>(f.get<GpuMat>(Channel::Colour).size())),
+ f.createTexture<short2>(Channel::Screen, Format<short2>(f.get<GpuMat>(Channel::Colour).size())),
+ params_, transform, f.getLeftCamera(), stream
+ );
+ }
// Must reset depth channel if blending
if (do_blend) {
@@ -399,7 +370,7 @@ void Triangular::_mesh(ftl::rgbd::Frame &out, const Eigen::Matrix4d &t, cudaStre
params_.camera, params_.m_viewMatrix.getFloat3x3(), params_.m_viewMatrixInverse.getFloat3x3(), stream_);
}
-void Triangular::_renderChannel(
+void CUDARender::_renderChannel(
ftl::rgbd::Frame &out,
Channel channel_in, Channel channel_out, const Eigen::Matrix4d &t, cudaStream_t stream)
{
@@ -476,7 +447,7 @@ static cv::Scalar HSVtoRGB(int H, double S, double V) {
return cv::Scalar((Bs + m) * 255, (Gs + m) * 255, (Rs + m) * 255, 0);
}
-void Triangular::_allocateChannels(ftl::rgbd::Frame &out) {
+void CUDARender::_allocateChannels(ftl::rgbd::Frame &out) {
const auto &camera = out.getLeftCamera();
cv::cuda::Stream cvstream = cv::cuda::StreamAccessor::wrapStream(stream_);
@@ -497,7 +468,7 @@ void Triangular::_allocateChannels(ftl::rgbd::Frame &out) {
temp_.createTexture<int>(Channel::Depth);
}
-void Triangular::_updateParameters(ftl::rgbd::Frame &out) {
+void CUDARender::_updateParameters(ftl::rgbd::Frame &out) {
const auto &camera = out.getLeftCamera();
// Parameters object to pass to CUDA describing the camera
@@ -510,7 +481,7 @@ void Triangular::_updateParameters(ftl::rgbd::Frame &out) {
params_.camera = camera;
}
-void Triangular::_preprocessColours() {
+void CUDARender::_preprocessColours() {
bool show_discon = value("show_discontinuity_mask", false);
bool show_fill = value("show_filled", false);
bool colour_sources = value("colour_sources", false);
@@ -540,7 +511,7 @@ void Triangular::_preprocessColours() {
}
}
-void Triangular::_postprocessColours(ftl::rgbd::Frame &out) {
+void CUDARender::_postprocessColours(ftl::rgbd::Frame &out) {
if (value("cool_effect", false)) {
auto pose = params_.m_viewMatrixInverse.getFloat3x3();
auto col = parseCUDAColour(value("cool_effect_colour", std::string("#2222ff")));
@@ -574,7 +545,7 @@ void Triangular::_postprocessColours(ftl::rgbd::Frame &out) {
}
}
-void Triangular::_renderNormals(ftl::rgbd::Frame &out) {
+void CUDARender::_renderNormals(ftl::rgbd::Frame &out) {
cv::cuda::Stream cvstream = cv::cuda::StreamAccessor::wrapStream(stream_);
// Visualise normals to RGBA
@@ -586,14 +557,14 @@ void Triangular::_renderNormals(ftl::rgbd::Frame &out) {
light_ambient_, stream_);
}
-void Triangular::_renderDensity(ftl::rgbd::Frame &out, const Eigen::Matrix4d &t) {
+void CUDARender::_renderDensity(ftl::rgbd::Frame &out, const Eigen::Matrix4d &t) {
cv::cuda::Stream cvstream = cv::cuda::StreamAccessor::wrapStream(stream_);
out.create<GpuMat>(Channel::Density, Format<float>(params_.camera.width, params_.camera.height));
out.get<GpuMat>(Channel::Density).setTo(cv::Scalar(0.0f), cvstream);
_renderChannel(out, Channel::Depth, Channel::Density, t, stream_);
}
-void Triangular::_renderRight(ftl::rgbd::Frame &out, const Eigen::Matrix4d &t) {
+void CUDARender::_renderRight(ftl::rgbd::Frame &out, const Eigen::Matrix4d &t) {
cv::cuda::Stream cvstream = cv::cuda::StreamAccessor::wrapStream(stream_);
float baseline = params_.camera.baseline;
@@ -619,7 +590,7 @@ void Triangular::_renderRight(ftl::rgbd::Frame &out, const Eigen::Matrix4d &t) {
_renderChannel(out, Channel::Left, Channel::Right, t, stream_);
}
-void Triangular::_renderSecond(ftl::rgbd::Frame &out, ftl::codecs::Channel chan, const Eigen::Matrix4d &t) {
+void CUDARender::_renderSecond(ftl::rgbd::Frame &out, ftl::codecs::Channel chan, const Eigen::Matrix4d &t) {
cv::cuda::Stream cvstream = cv::cuda::StreamAccessor::wrapStream(stream_);
if (ftl::codecs::isFloatChannel(chan)) {
@@ -632,10 +603,49 @@ void Triangular::_renderSecond(ftl::rgbd::Frame &out, ftl::codecs::Channel chan,
_renderChannel(out, chan, chan, t, stream_);
}
-void Triangular::_render(ftl::rgbd::FrameSet &in, ftl::rgbd::Frame &out, Channel chan, const Eigen::Matrix4d &t) {
+void CUDARender::_renderPass1(const Eigen::Matrix4d &t) {
+ const auto &camera = out_->getLeftCamera();
+ cv::cuda::Stream cvstream = cv::cuda::StreamAccessor::wrapStream(stream_);
+
+ // Render source specific debug info into colour channels
+ _preprocessColours();
+
+ if (mesh_) {
+ // Render depth channel using triangles
+ _mesh(*out_, t, stream_);
+ } else {
+ // Render depth channel as a point cloud
+ _dibr(*out_, t, stream_);
+ }
+}
+
+void CUDARender::_renderPass2(Channels<0> chans, const Eigen::Matrix4d &t) {
+ const auto &camera = out_->getLeftCamera();
+ cv::cuda::Stream cvstream = cv::cuda::StreamAccessor::wrapStream(stream_);
+
+ // Reprojection of colours onto surface
+ auto main_channel = (scene_->frames[0].hasChannel(Channel::ColourHighRes)) ? Channel::ColourHighRes : Channel::Colour;
+ _renderChannel(*out_, main_channel, Channel::Colour, t, stream_);
+
+
+ // Support rendering of a second channel without redoing all the work
+ for (auto chan : chans) {
+ switch(chan) {
+ case Channel::None :
+ case Channel::Left :
+ case Channel::Depth : break;
+ case Channel::ColourNormals : _renderNormals(*out_); break;
+ case Channel::Density : _renderDensity(*out_, t); break;
+ case Channel::Right : _renderRight(*out_, t); break;
+ default : _renderSecond(*out_, chan, t);
+ }
+ }
+}
+
+void CUDARender::begin(ftl::rgbd::Frame &out) {
+ out_ = &out;
const auto &camera = out.getLeftCamera();
cv::cuda::Stream cvstream = cv::cuda::StreamAccessor::wrapStream(stream_);
- //cudaSafeCall(cudaSetDevice(scene_->getCUDADevice()));
_updateParameters(out);
@@ -653,52 +663,54 @@ void Triangular::_render(ftl::rgbd::FrameSet &in, ftl::rgbd::Frame &out, Channel
camera, pose, stream_);
}
- //cudaSafeCall(cudaStreamSynchronize(stream_));
+ temp_.create<GpuMat>(
+ AccumSelector<uchar4>::channel,
+ Format<typename AccumSelector<uchar4>::type>(params_.camera.width, params_.camera.height)
+ ).setTo(cv::Scalar(0.0f), cvstream);
+ temp_.get<GpuMat>(Channel::Contribution).setTo(cv::Scalar(0.0f), cvstream);
- // Render source specific debug info into colour channels
- _preprocessColours();
+ temp_.createTexture<float>(Channel::Contribution);
- //cudaSafeCall(cudaStreamSynchronize(stream_));
+ sets_.clear();
+}
- if (mesh_) {
- // Render depth channel using triangles
- _mesh(out, t, stream_);
- } else {
- // Render depth channel as a point cloud
- _dibr(out, t, stream_);
+void CUDARender::end() {
+ /*ftl::cuda::dibr_normalise(
+ temp_.getTexture<typename AccumSelector<T>::type>(AccumSelector<T>::channel),
+ output.createTexture<T>(out),
+ temp_.getTexture<float>(Channel::Contribution),
+ stream
+ );*/
+
+ for (auto &s : sets_) {
+ scene_ = s.fs;
+ try {
+ _renderPass2(s.channels, s.transform);
+ } catch(std::exception &e) {
+ LOG(ERROR) << "Exception in render: " << e.what();
+ }
}
+ scene_ = nullptr;
- //cudaSafeCall(cudaStreamSynchronize(stream_));
-
- // Reprojection of colours onto surface
- auto main_channel = (scene_->frames[0].hasChannel(Channel::ColourHighRes)) ? Channel::ColourHighRes : Channel::Colour;
- _renderChannel(out, main_channel, Channel::Colour, t, stream_);
-
- //cudaSafeCall(cudaStreamSynchronize(stream_));
-
- // Debug colour info relating to the rendering process
- _postprocessColours(out);
+ // FIXME: Allow for other channel accumulations
+ ftl::cuda::dibr_normalise(
+ temp_.getTexture<typename AccumSelector<uchar4>::type>(AccumSelector<uchar4>::channel),
+ out_->createTexture<uchar4>(Channel::Colour),
+ temp_.getTexture<float>(Channel::Contribution),
+ stream_
+ );
- //cudaSafeCall(cudaStreamSynchronize(stream_));
+ _postprocessColours(*out_);
- // Support rendering of a second channel without redoing all the work
- switch(chan) {
- case Channel::None :
- case Channel::Left :
- case Channel::Depth : break;
- case Channel::ColourNormals : _renderNormals(out); break;
- case Channel::Density : _renderDensity(out, t); break;
- case Channel::Right : _renderRight(out, t); break;
- default : _renderSecond(out, chan, t);
- }
+ cudaSafeCall(cudaStreamSynchronize(stream_));
}
-bool Triangular::render(ftl::rgbd::FrameSet &in, ftl::rgbd::Frame &out, Channel chan, const Eigen::Matrix4d &t) {
+bool CUDARender::submit(ftl::rgbd::FrameSet *in, Channels<0> chans, const Eigen::Matrix4d &t) {
if (scene_) {
LOG(WARNING) << "Already rendering...";
return false;
}
- scene_ = ∈
+ scene_ = in;
if (scene_->frames.size() == 0) {
scene_ = nullptr;
return false;
@@ -707,13 +719,18 @@ bool Triangular::render(ftl::rgbd::FrameSet &in, ftl::rgbd::Frame &out, Channel
bool success = true;
try {
- _render(in, out, chan, t);
- cudaSafeCall(cudaStreamSynchronize(stream_));
+ _renderPass1(t);
+ //cudaSafeCall(cudaStreamSynchronize(stream_));
} catch (std::exception &e) {
LOG(ERROR) << "Exception in render: " << e.what();
success = false;
}
+ auto &s = sets_.emplace_back();
+ s.fs = in;
+ s.channels = chans;
+ s.transform = t;
+
last_frame_ = scene_->timestamp;
scene_ = nullptr;
return success;
diff --git a/components/renderers/cpp/src/reprojection.cu b/components/renderers/cpp/src/reprojection.cu
index 2e342d1e2ceb903416fd7e22f668a0819127cd27..79b01ff534cf36d26c4d58c3e15d3b32fbd95020 100644
--- a/components/renderers/cpp/src/reprojection.cu
+++ b/components/renderers/cpp/src/reprojection.cu
@@ -286,6 +286,103 @@ template void ftl::cuda::reproject(
const ftl::rgbd::Camera &camera,
const float4x4 &poseInv, cudaStream_t stream);
+//==============================================================================
+// Without normals or depth
+//==============================================================================
+
+/*
+ * Pass 2: Accumulate attribute contributions if the points pass a visibility test.
+ */
+ template <typename A, typename B>
+__global__ void reprojection_kernel(
+ TextureObject<A> in, // Attribute input
+ TextureObject<float> depth_in, // Virtual depth map
+ TextureObject<B> out, // Accumulated output
+ TextureObject<float> contrib,
+ SplatParams params,
+ Camera camera, float4x4 poseInv) {
+
+ const int x = (blockIdx.x*blockDim.x + threadIdx.x);
+ const int y = blockIdx.y*blockDim.y + threadIdx.y;
+
+ const float d = depth_in.tex2D((int)x, (int)y);
+ if (d < params.camera.minDepth || d > params.camera.maxDepth) return;
+
+ //const float3 worldPos = params.m_viewMatrixInverse * params.camera.screenToCam(x, y, d);
+ //if (worldPos.x == MINF || (!(params.m_flags & ftl::render::kShowDisconMask) && worldPos.w < 0.0f)) return;
+
+ const float3 camPos = poseInv * params.camera.screenToCam(x, y, d);
+ //if (camPos.z < camera.minDepth) return;
+ //if (camPos.z > camera.maxDepth) return;
+ const float2 screenPos = camera.camToScreen<float2>(camPos);
+
+ // Not on screen so stop now...
+ if (screenPos.x >= in.width() || screenPos.y >= in.height()) return;
+
+ const float d2 = camera.maxDepth;
+
+ const auto input = in.tex2D(screenPos.x, screenPos.y); //generateInput(in.tex2D((int)screenPos.x, (int)screenPos.y), params, worldPos);
+
+ float weight = ftl::cuda::weighting(fabs(camPos.z - d2), 0.02f);
+ const B weighted = make<B>(input) * weight;
+
+ if (weight > 0.0f) {
+ accumulateOutput(out, contrib, make_uint2(x,y), weighted, weight);
+ //out(screenPos.x, screenPos.y) = input;
+ }
+}
+
+
+template <typename A, typename B>
+void ftl::cuda::reproject(
+ TextureObject<A> &in,
+ TextureObject<float> &depth_in, // Virtual depth map
+ TextureObject<B> &out, // Accumulated output
+ TextureObject<float> &contrib,
+ const SplatParams ¶ms,
+ const Camera &camera, const float4x4 &poseInv, cudaStream_t stream) {
+ const dim3 gridSize((out.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (out.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
+ const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
+
+ reprojection_kernel<<<gridSize, blockSize, 0, stream>>>(
+ in,
+ depth_in,
+ out,
+ contrib,
+ params,
+ camera,
+ poseInv
+ );
+ cudaSafeCall( cudaGetLastError() );
+}
+
+template void ftl::cuda::reproject(
+ ftl::cuda::TextureObject<uchar4> &in, // Original colour image
+ ftl::cuda::TextureObject<float> &depth_in, // Virtual depth map
+ ftl::cuda::TextureObject<float4> &out, // Accumulated output
+ ftl::cuda::TextureObject<float> &contrib,
+ const ftl::render::SplatParams ¶ms,
+ const ftl::rgbd::Camera &camera,
+ const float4x4 &poseInv, cudaStream_t stream);
+
+template void ftl::cuda::reproject(
+ ftl::cuda::TextureObject<float> &in, // Original colour image
+ ftl::cuda::TextureObject<float> &depth_in, // Virtual depth map
+ ftl::cuda::TextureObject<float> &out, // Accumulated output
+ ftl::cuda::TextureObject<float> &contrib,
+ const ftl::render::SplatParams ¶ms,
+ const ftl::rgbd::Camera &camera,
+ const float4x4 &poseInv, cudaStream_t stream);
+
+template void ftl::cuda::reproject(
+ ftl::cuda::TextureObject<float4> &in, // Original colour image
+ ftl::cuda::TextureObject<float> &depth_in, // Virtual depth map
+ ftl::cuda::TextureObject<float4> &out, // Accumulated output
+ ftl::cuda::TextureObject<float> &contrib,
+ const ftl::render::SplatParams ¶ms,
+ const ftl::rgbd::Camera &camera,
+ const float4x4 &poseInv, cudaStream_t stream);
+
// ===== Equirectangular Reprojection ==========================================
diff --git a/components/renderers/cpp/src/screen.cu b/components/renderers/cpp/src/screen.cu
index c4991ee6e65dd6801acd847a67a4713be8edfd13..f12f255f0b9a3fbdd5f418864567f5688e489ba2 100644
--- a/components/renderers/cpp/src/screen.cu
+++ b/components/renderers/cpp/src/screen.cu
@@ -47,3 +47,42 @@ void ftl::cuda::screen_coord(TextureObject<float> &depth, TextureObject<float> &
screen_coord_kernel<<<gridSize, blockSize, 0, stream>>>(depth, depth_out, screen_out, params.camera, pose, camera);
cudaSafeCall( cudaGetLastError() );
}
+
+
+// ==== Constant depth version =================================================
+
+/*
+ * Convert source screen position to output screen coordinates. Assumes a
+ * constant depth of 1m instead of using a depth channel input.
+ */
+ __global__ void screen_coord_kernel(TextureObject<float> depth_out,
+ TextureObject<short2> screen_out, Camera vcamera, float4x4 pose, Camera camera) {
+ const int x = blockIdx.x*blockDim.x + threadIdx.x;
+ const int y = blockIdx.y*blockDim.y + threadIdx.y;
+
+ if (x >= 0 && y >= 0 && x < depth_out.width() && y < depth_out.height()) {
+ uint2 screenPos = make_uint2(30000,30000);
+ const float d = camera.maxDepth;
+
+ // Find the virtual screen position of current point
+ const float3 camPos = pose * camera.screenToCam(x,y,d);
+ screenPos = vcamera.camToScreen<uint2>(camPos);
+
+ if ( camPos.z < vcamera.minDepth ||
+ camPos.z > vcamera.maxDepth ||
+ screenPos.x >= vcamera.width ||
+ screenPos.y >= vcamera.height)
+ screenPos = make_uint2(30000,30000);
+
+ screen_out(x,y) = make_short2(screenPos.x, screenPos.y);
+ depth_out(x,y) = camPos.z;
+ }
+}
+
+void ftl::cuda::screen_coord(TextureObject<float> &depth_out, TextureObject<short2> &screen_out, const SplatParams ¶ms, const float4x4 &pose, const Camera &camera, cudaStream_t stream) {
+ const dim3 gridSize((screen_out.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (screen_out.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
+ const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
+
+ screen_coord_kernel<<<gridSize, blockSize, 0, stream>>>(depth_out, screen_out, params.camera, pose, camera);
+ cudaSafeCall( cudaGetLastError() );
+}
diff --git a/components/renderers/cpp/src/splatter.cpp b/components/renderers/cpp/src/splatter.cpp
deleted file mode 100644
index 731efec404dc24a972f290c61d8fc0c6d2ba34c7..0000000000000000000000000000000000000000
--- a/components/renderers/cpp/src/splatter.cpp
+++ /dev/null
@@ -1,498 +0,0 @@
-#include <ftl/render/splat_render.hpp>
-#include <ftl/utility/matrix_conversion.hpp>
-#include "splatter_cuda.hpp"
-#include <ftl/cuda/points.hpp>
-#include <ftl/cuda/normals.hpp>
-#include <ftl/cuda/mask.hpp>
-
-#include <opencv2/core/cuda_stream_accessor.hpp>
-
-#include <string>
-
-using ftl::render::Splatter;
-using ftl::codecs::Channel;
-using ftl::codecs::Channels;
-using ftl::rgbd::Format;
-using cv::cuda::GpuMat;
-using std::stoul;
-using ftl::cuda::Mask;
-
-static Eigen::Affine3d create_rotation_matrix(float ax, float ay, float az) {
- Eigen::Affine3d rx =
- Eigen::Affine3d(Eigen::AngleAxisd(ax, Eigen::Vector3d(1, 0, 0)));
- Eigen::Affine3d ry =
- Eigen::Affine3d(Eigen::AngleAxisd(ay, Eigen::Vector3d(0, 1, 0)));
- Eigen::Affine3d rz =
- Eigen::Affine3d(Eigen::AngleAxisd(az, Eigen::Vector3d(0, 0, 1)));
- return rz * rx * ry;
-}
-
-/*
- * Parse a CSS style colour string into a scalar.
- */
-static cv::Scalar parseCVColour(const std::string &colour) {
- std::string c = colour;
- if (c[0] == '#') {
- c.erase(0, 1);
- unsigned long value = stoul(c.c_str(), nullptr, 16);
- return cv::Scalar(
- (value >> 0) & 0xff,
- (value >> 8) & 0xff,
- (value >> 16) & 0xff,
- (value >> 24) & 0xff
- );
- }
-
- return cv::Scalar(0,0,0,0);
-}
-
-/*
- * Parse a CSS style colour string into a scalar.
- */
-static uchar4 parseCUDAColour(const std::string &colour) {
- std::string c = colour;
- if (c[0] == '#') {
- c.erase(0, 1);
- unsigned long value = stoul(c.c_str(), nullptr, 16);
- return make_uchar4(
- (value >> 0) & 0xff,
- (value >> 8) & 0xff,
- (value >> 16) & 0xff,
- (value >> 24) & 0xff
- );
- }
-
- return make_uchar4(0,0,0,0);
-}
-
-Splatter::Splatter(nlohmann::json &config, ftl::rgbd::FrameSet *fs) : ftl::render::Renderer(config), scene_(fs) {
- if (config["clipping"].is_object()) {
- auto &c = config["clipping"];
- float rx = c.value("pitch", 0.0f);
- float ry = c.value("yaw", 0.0f);
- float rz = c.value("roll", 0.0f);
- float x = c.value("x", 0.0f);
- float y = c.value("y", 0.0f);
- float z = c.value("z", 0.0f);
- float width = c.value("width", 1.0f);
- float height = c.value("height", 1.0f);
- float depth = c.value("depth", 1.0f);
-
- Eigen::Affine3f r = create_rotation_matrix(rx, ry, rz).cast<float>();
- Eigen::Translation3f trans(Eigen::Vector3f(x,y,z));
- Eigen::Affine3f t(trans);
-
- clip_.origin = MatrixConversion::toCUDA(r.matrix() * t.matrix());
- clip_.size = make_float3(width, height, depth);
- clipping_ = value("clipping_enabled", true);
- } else {
- clipping_ = false;
- }
-
- on("clipping_enabled", [this](const ftl::config::Event &e) {
- clipping_ = value("clipping_enabled", true);
- });
-
- norm_filter_ = value("normal_filter", -1.0f);
- on("normal_filter", [this](const ftl::config::Event &e) {
- norm_filter_ = value("normal_filter", -1.0f);
- });
-
- backcull_ = value("back_cull", true);
- on("back_cull", [this](const ftl::config::Event &e) {
- backcull_ = value("back_cull", true);
- });
-
- splat_ = value("splatting", true);
- on("splatting", [this](const ftl::config::Event &e) {
- splat_ = value("splatting", true);
- });
-
- background_ = parseCVColour(value("background", std::string("#4c4c4c")));
- on("background", [this](const ftl::config::Event &e) {
- background_ = parseCVColour(value("background", std::string("#4c4c4c")));
- });
-
- light_diffuse_ = parseCUDAColour(value("diffuse", std::string("#e0e0e0")));
- on("diffuse", [this](const ftl::config::Event &e) {
- light_diffuse_ = parseCUDAColour(value("diffuse", std::string("#e0e0e0")));
- });
-
- light_ambient_ = parseCUDAColour(value("ambient", std::string("#0e0e0e")));
- on("ambient", [this](const ftl::config::Event &e) {
- light_ambient_ = parseCUDAColour(value("ambient", std::string("#0e0e0e")));
- });
-
- light_pos_ = make_float3(value("light_x", 0.3f), value("light_y", 0.2f), value("light_z", 1.0f));
- on("light_x", [this](const ftl::config::Event &e) { light_pos_.x = value("light_x", 0.3f); });
- on("light_y", [this](const ftl::config::Event &e) { light_pos_.y = value("light_y", 0.3f); });
- on("light_z", [this](const ftl::config::Event &e) { light_pos_.z = value("light_z", 0.3f); });
-
- cudaSafeCall(cudaStreamCreate(&stream_));
-}
-
-Splatter::~Splatter() {
-
-}
-
-template <typename T>
-struct AccumSelector {
- typedef float4 type;
- static constexpr Channel channel = Channel::Colour;
- //static constexpr cv::Scalar value = cv::Scalar(0.0f,0.0f,0.0f,0.0f);
-};
-
-template <>
-struct AccumSelector<float> {
- typedef float type;
- static constexpr Channel channel = Channel::Colour2;
- //static constexpr cv::Scalar value = cv::Scalar(0.0f);
-};
-
-template <typename T>
-void Splatter::__blendChannel(ftl::rgbd::Frame &output, ftl::codecs::Channel in, ftl::codecs::Channel out, cudaStream_t stream) {
- cv::cuda::Stream cvstream = cv::cuda::StreamAccessor::wrapStream(stream);
- temp_.create<GpuMat>(
- AccumSelector<T>::channel,
- Format<typename AccumSelector<T>::type>(params_.camera.width, params_.camera.height)
- ).setTo(cv::Scalar(0.0f), cvstream);
- temp_.get<GpuMat>(Channel::Contribution).setTo(cv::Scalar(0.0f), cvstream);
-
- temp_.createTexture<float>(Channel::Contribution);
-
- for (auto &f : scene_->frames) {
- if (f.get<GpuMat>(in).type() == CV_8UC3) {
- // Convert to 4 channel colour
- auto &col = f.get<GpuMat>(in);
- GpuMat tmp(col.size(), CV_8UC4);
- cv::cuda::swap(col, tmp);
- cv::cuda::cvtColor(tmp,col, cv::COLOR_BGR2BGRA);
- }
-
- ftl::cuda::dibr_attribute(
- f.createTexture<T>(in),
- f.createTexture<float4>(Channel::Points),
- temp_.getTexture<int>(Channel::Depth2),
- temp_.createTexture<typename AccumSelector<T>::type>(AccumSelector<T>::channel),
- temp_.getTexture<float>(Channel::Contribution),
- params_, stream
- );
- }
-
- ftl::cuda::dibr_normalise(
- temp_.getTexture<typename AccumSelector<T>::type>(AccumSelector<T>::channel),
- output.createTexture<T>(out),
- temp_.getTexture<float>(Channel::Contribution),
- stream
- );
-}
-
-void Splatter::_blendChannel(ftl::rgbd::Frame &output, ftl::codecs::Channel in, ftl::codecs::Channel out, cudaStream_t stream) {
- int type = output.get<GpuMat>(out).type(); // == CV_32F; //ftl::rgbd::isFloatChannel(channel);
-
- switch (type) {
- case CV_32F : __blendChannel<float>(output, in, out, stream); break;
- case CV_32FC4 : __blendChannel<float4>(output, in, out, stream); break;
- case CV_8UC4 : __blendChannel<uchar4>(output, in, out, stream); break;
- default : LOG(ERROR) << "Invalid output channel format";
- }
-}
-
-void Splatter::_dibr(cudaStream_t stream) {
- cv::cuda::Stream cvstream = cv::cuda::StreamAccessor::wrapStream(stream);
- temp_.get<GpuMat>(Channel::Depth2).setTo(cv::Scalar(0x7FFFFFFF), cvstream);
-
- for (size_t i=0; i < scene_->frames.size(); ++i) {
- auto &f = scene_->frames[i];
- //auto *s = scene_->sources[i];
-
- if (f.empty(Channel::Depth + Channel::Colour)) {
- LOG(ERROR) << "Missing required channel";
- continue;
- }
-
- ftl::cuda::dibr_merge(
- f.createTexture<float4>(Channel::Points),
- f.createTexture<float4>(Channel::Normals),
- temp_.createTexture<int>(Channel::Depth2),
- params_, backcull_, stream
- );
-
- //LOG(INFO) << "DIBR DONE";
- }
-}
-
-void Splatter::_renderChannel(
- ftl::rgbd::Frame &out,
- Channel channel_in, Channel channel_out, cudaStream_t stream)
-{
- if (channel_out == Channel::None || channel_in == Channel::None) return;
- cv::cuda::Stream cvstream = cv::cuda::StreamAccessor::wrapStream(stream);
-
- if (scene_->frames.size() < 1) return;
- bool is_float = out.get<GpuMat>(channel_out).type() == CV_32F; //ftl::rgbd::isFloatChannel(channel);
- bool is_4chan = out.get<GpuMat>(channel_out).type() == CV_32FC4;
-
-
- temp_.createTexture<float4>(Channel::Colour);
- temp_.createTexture<float>(Channel::Contribution);
-
- // Generate initial normals for the splats
- accum_.create<GpuMat>(Channel::Normals, Format<float4>(params_.camera.width, params_.camera.height));
- _blendChannel(accum_, Channel::Normals, Channel::Normals, stream);
- // Put normals in camera space here...
- ftl::cuda::transform_normals(accum_.getTexture<float4>(Channel::Normals), params_.m_viewMatrix.getFloat3x3(), stream);
-
- // Estimate point density
- accum_.create<GpuMat>(Channel::Density, Format<float>(params_.camera.width, params_.camera.height));
- _blendChannel(accum_, Channel::Depth, Channel::Density, stream);
-
- // FIXME: Using colour 2 in this way seems broken since it is already used
- if (is_4chan) {
- accum_.create<GpuMat>(channel_out, Format<float4>(params_.camera.width, params_.camera.height));
- accum_.get<GpuMat>(channel_out).setTo(cv::Scalar(0.0f,0.0f,0.0f,0.0f), cvstream);
- } else if (is_float) {
- accum_.create<GpuMat>(channel_out, Format<float>(params_.camera.width, params_.camera.height));
- accum_.get<GpuMat>(channel_out).setTo(cv::Scalar(0.0f), cvstream);
- } else {
- accum_.create<GpuMat>(channel_out, Format<uchar4>(params_.camera.width, params_.camera.height));
- accum_.get<GpuMat>(channel_out).setTo(cv::Scalar(0,0,0,0), cvstream);
- }
-
- //if (splat_) {
- _blendChannel(accum_, channel_in, channel_out, stream);
- //} else {
- // _blendChannel(out, channel, channel, stream);
- //}
-
- // Now splat the points
- if (splat_) {
- if (is_4chan) {
- ftl::cuda::splat(
- accum_.getTexture<float4>(Channel::Normals),
- accum_.getTexture<float>(Channel::Density),
- accum_.getTexture<float4>(channel_out),
- temp_.getTexture<int>(Channel::Depth2),
- out.createTexture<float>(Channel::Depth),
- out.createTexture<float4>(channel_out),
- params_, stream
- );
- } else if (is_float) {
- ftl::cuda::splat(
- accum_.getTexture<float4>(Channel::Normals),
- accum_.getTexture<float>(Channel::Density),
- accum_.getTexture<float>(channel_out),
- temp_.getTexture<int>(Channel::Depth2),
- out.createTexture<float>(Channel::Depth),
- out.createTexture<float>(channel_out),
- params_, stream
- );
- } else {
- ftl::cuda::splat(
- accum_.getTexture<float4>(Channel::Normals),
- accum_.getTexture<float>(Channel::Density),
- accum_.getTexture<uchar4>(channel_out),
- temp_.getTexture<int>(Channel::Depth2),
- out.createTexture<float>(Channel::Depth),
- out.createTexture<uchar4>(channel_out),
- params_, stream
- );
- }
- } else {
- // Swap accum frames directly to output.
- accum_.swapTo(Channels(channel_out), out);
- }
-}
-
-bool Splatter::render(ftl::rgbd::VirtualSource *src, ftl::rgbd::Frame &out, const Eigen::Matrix4d &t) {
- SHARED_LOCK(scene_->mtx, lk);
- if (!src->isReady()) return false;
-
- scene_->upload(Channel::Colour + Channel::Depth, stream_);
-
- const auto &camera = src->parameters();
- //cudaSafeCall(cudaSetDevice(scene_->getCUDADevice()));
-
- // 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));
-
-
- if (scene_->frames.size() == 0) return false;
- auto &g = scene_->frames[0].get<GpuMat>(Channel::Colour);
-
- temp_.create<GpuMat>(Channel::Colour, Format<float4>(camera.width, camera.height));
- temp_.create<GpuMat>(Channel::Contribution, Format<float>(camera.width, camera.height));
- temp_.create<GpuMat>(Channel::Depth, Format<int>(camera.width, camera.height));
- temp_.create<GpuMat>(Channel::Depth2, Format<int>(camera.width, camera.height));
- temp_.create<GpuMat>(Channel::Normals, Format<float4>(g.cols, g.rows));
-
- cv::cuda::Stream cvstream = cv::cuda::StreamAccessor::wrapStream(stream_);
-
- // Parameters object to pass to CUDA describing the camera
- SplatParams ¶ms = params_;
- params.m_flags = 0;
- //if () params.m_flags |= ftl::render::kShowDisconMask;
- if (value("normal_weight_colours", true)) params.m_flags |= ftl::render::kNormalWeightColours;
- params.m_viewMatrix = MatrixConversion::toCUDA(src->getPose().cast<float>().inverse());
- params.m_viewMatrixInverse = MatrixConversion::toCUDA(src->getPose().cast<float>());
- params.camera = camera;
- // Clear all channels to 0 or max depth
-
- out.get<GpuMat>(Channel::Depth).setTo(cv::Scalar(1000.0f), cvstream);
- out.get<GpuMat>(Channel::Colour).setTo(background_, cvstream);
-
- //LOG(INFO) << "Render ready: " << camera.width << "," << camera.height;
-
- bool show_discon = value("show_discontinuity_mask", false);
- bool show_fill = value("show_filled", false);
-
- temp_.createTexture<int>(Channel::Depth);
- //temp_.get<GpuMat>(Channel::Normals).setTo(cv::Scalar(0.0f,0.0f,0.0f,0.0f), cvstream);
-
- // First make sure each input has normals
- temp_.createTexture<float4>(Channel::Normals);
- for (int i=0; i<scene_->frames.size(); ++i) {
- auto &f = scene_->frames[i];
- //auto s = scene_->sources[i];
-
- if (f.hasChannel(Channel::Mask)) {
- if (show_discon) {
- ftl::cuda::show_mask(f.getTexture<uchar4>(Channel::Colour), f.getTexture<int>(Channel::Mask), Mask::kMask_Discontinuity, make_uchar4(0,0,255,255), stream_);
- }
- if (show_fill) {
- ftl::cuda::show_mask(f.getTexture<uchar4>(Channel::Colour), f.getTexture<int>(Channel::Mask), Mask::kMask_Filled, make_uchar4(0,255,0,255), stream_);
- }
- }
-
- // Needs to create points channel first?
- if (!f.hasChannel(Channel::Points)) {
- //LOG(INFO) << "Creating points... " << s->parameters().width;
-
- auto &t = f.createTexture<float4>(Channel::Points, Format<float4>(f.get<GpuMat>(Channel::Colour).size()));
- auto pose = MatrixConversion::toCUDA(f.getPose().cast<float>()); //.inverse());
- ftl::cuda::point_cloud(t, f.createTexture<float>(Channel::Depth), f.getLeftCamera(), pose, 0, stream_);
-
- //LOG(INFO) << "POINTS Added";
- }
-
- // Clip first?
- if (clipping_) {
- ftl::cuda::clipping(f.createTexture<float4>(Channel::Points), clip_, stream_);
- }
-
- if (!f.hasChannel(Channel::Normals)) {
- Eigen::Matrix4f matrix = f.getPose().cast<float>().transpose();
- auto pose = MatrixConversion::toCUDA(matrix);
-
- auto &g = f.get<GpuMat>(Channel::Colour);
- ftl::cuda::normals(f.createTexture<float4>(Channel::Normals, Format<float4>(g.cols, g.rows)),
- temp_.getTexture<float4>(Channel::Normals),
- f.getTexture<float4>(Channel::Points),
- 1, 0.02f,
- f.getLeftCamera(), pose.getFloat3x3(), stream_);
-
- if (norm_filter_ > -0.1f) {
- ftl::cuda::normal_filter(f.getTexture<float4>(Channel::Normals), f.getTexture<float4>(Channel::Points), f.getLeftCamera(), pose, norm_filter_, stream_);
- }
- }
- }
-
- Channel chan = src->getChannel();
-
- int aligned_source = value("aligned_source",-1);
- if (aligned_source >= 0 && aligned_source < scene_->frames.size()) {
- // FIXME: Output may not be same resolution as source!
- cudaSafeCall(cudaStreamSynchronize(stream_));
- scene_->frames[aligned_source].copyTo(Channel::Depth + Channel::Colour, out);
-
- if (chan == Channel::Normals) {
- // Convert normal to single float value
- temp_.create<GpuMat>(Channel::Colour, Format<uchar4>(camera.width, camera.height)).setTo(cv::Scalar(0,0,0,0), cvstream);
- ftl::cuda::normal_visualise(scene_->frames[aligned_source].getTexture<float4>(Channel::Normals), temp_.createTexture<uchar4>(Channel::Colour),
- light_pos_,
- light_diffuse_,
- light_ambient_, stream_);
-
- // Put in output as single float
- cv::cuda::swap(temp_.get<GpuMat>(Channel::Colour), out.create<GpuMat>(Channel::Normals));
- out.resetTexture(Channel::Normals);
- }
-
- return true;
- }
-
- _dibr(stream_);
- _renderChannel(out, Channel::Colour, Channel::Colour, stream_);
-
- if (chan == Channel::Depth)
- {
- //temp_.get<GpuMat>(Channel::Depth).convertTo(out.get<GpuMat>(Channel::Depth), CV_32F, 1.0f / 1000.0f, cvstream);
- } else if (chan == Channel::Normals) {
- out.create<GpuMat>(Channel::Normals, Format<float4>(camera.width, camera.height));
-
- // Render normal attribute
- _renderChannel(out, Channel::Normals, Channel::Normals, stream_);
-
- // Convert normal to single float value
- temp_.create<GpuMat>(Channel::Colour, Format<uchar4>(camera.width, camera.height)).setTo(cv::Scalar(0,0,0,0), cvstream);
- ftl::cuda::normal_visualise(out.getTexture<float4>(Channel::Normals), temp_.createTexture<uchar4>(Channel::Colour),
- light_pos_,
- light_diffuse_,
- light_ambient_, stream_);
-
- // Put in output as single float
- cv::cuda::swap(temp_.get<GpuMat>(Channel::Colour), out.create<GpuMat>(Channel::Normals));
- out.resetTexture(Channel::Normals);
- }
- //else if (chan == Channel::Contribution)
- //{
- // cv::cuda::swap(temp_.get<GpuMat>(Channel::Contribution), out.create<GpuMat>(Channel::Contribution));
- //}
- else if (chan == Channel::Density) {
- out.create<GpuMat>(chan, Format<float>(camera.width, camera.height));
- out.get<GpuMat>(chan).setTo(cv::Scalar(0.0f), cvstream);
- _renderChannel(out, Channel::Depth, Channel::Density, stream_);
- }
- else if (chan == Channel::Right)
- {
- float baseline = camera.baseline;
-
- //Eigen::Translation3f translation(baseline, 0.0f, 0.0f);
- //Eigen::Affine3f transform(translation);
- //Eigen::Matrix4f matrix = transform.matrix() * src->getPose().cast<float>();
-
- Eigen::Matrix4f transform = Eigen::Matrix4f::Identity();
- transform(0, 3) = baseline;
- Eigen::Matrix4f matrix = transform.inverse() * src->getPose().cast<float>();
-
- params.m_viewMatrix = MatrixConversion::toCUDA(matrix.inverse());
- params.m_viewMatrixInverse = MatrixConversion::toCUDA(matrix);
-
- params.camera = src->parameters(Channel::Right);
-
- out.create<GpuMat>(Channel::Right, Format<uchar4>(camera.width, camera.height));
- out.get<GpuMat>(Channel::Right).setTo(background_, cvstream);
-
- _dibr(stream_); // Need to re-dibr due to pose change
- _renderChannel(out, Channel::Left, Channel::Right, stream_);
-
- } else if (chan != Channel::None) {
- if (ftl::codecs::isFloatChannel(chan)) {
- out.create<GpuMat>(chan, Format<float>(camera.width, camera.height));
- out.get<GpuMat>(chan).setTo(cv::Scalar(0.0f), cvstream);
- } else {
- out.create<GpuMat>(chan, Format<uchar4>(camera.width, camera.height));
- out.get<GpuMat>(chan).setTo(background_, cvstream);
- }
- _renderChannel(out, chan, chan, stream_);
- }
-
- cudaSafeCall(cudaStreamSynchronize(stream_));
- return true;
-}
-
-//void Splatter::setOutputDevice(int device) {
-// device_ = device;
-//}
diff --git a/components/renderers/cpp/src/splatter.cu b/components/renderers/cpp/src/splatter.cu
index 55706b0856750738134a3417dfdb017205ab2bdf..7696433bd78f6187132a0bf2547cc653cc5c9b2e 100644
--- a/components/renderers/cpp/src/splatter.cu
+++ b/components/renderers/cpp/src/splatter.cu
@@ -137,6 +137,39 @@ using ftl::cuda::warpSum;
}
}
+/*
+ * Pass 1: Directly render each camera into virtual view but with no upsampling
+ * for sparse points.
+ */
+ __global__ void dibr_merge_kernel(
+ TextureObject<int> depth_out,
+ float4x4 transform,
+ ftl::rgbd::Camera cam,
+ SplatParams params) {
+ const int x = blockIdx.x*blockDim.x + threadIdx.x;
+ const int y = blockIdx.y*blockDim.y + threadIdx.y;
+
+ const float d0 = 1.0f;
+
+ const float3 camPos = transform * cam.screenToCam(x,y,d0);
+ //if (worldPos.x == MINF || (!(params.m_flags & ftl::render::kShowDisconMask) && worldPos.w < 0.0f)) return;
+
+ // Find the virtual screen position of current point
+ //const float3 camPos = params.m_viewMatrix * make_float3(worldPos);
+ //if (camPos.z < params.camera.minDepth) return;
+ //if (camPos.z > params.camera.maxDepth) return;
+
+ const float d = camPos.z;
+
+ const uint2 screenPos = params.camera.camToScreen<uint2>(camPos);
+ const unsigned int cx = screenPos.x;
+ const unsigned int cy = screenPos.y;
+ if (d > params.camera.minDepth && d < params.camera.maxDepth && cx < depth_out.width() && cy < depth_out.height()) {
+ // Transform estimated point to virtual cam space and output z
+ atomicMin(&depth_out(cx,cy), d * 100000.0f);
+ }
+}
+
void ftl::cuda::dibr_merge(TextureObject<float4> &points, TextureObject<float4> &normals, TextureObject<int> &depth, SplatParams params, bool culling, 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);
@@ -155,13 +188,21 @@ void ftl::cuda::dibr_merge(TextureObject<float4> &points, TextureObject<int> &de
}
void ftl::cuda::dibr_merge(TextureObject<float> &depth, TextureObject<int> &depth_out, const float4x4 &transform, const ftl::rgbd::Camera &cam, SplatParams params, cudaStream_t stream) {
- const dim3 gridSize((depth_out.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (depth_out.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
+ 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>>>(depth, depth_out, transform, cam, params);
cudaSafeCall( cudaGetLastError() );
}
+void ftl::cuda::dibr_merge(TextureObject<int> &depth_out, const float4x4 &transform, const ftl::rgbd::Camera &cam, SplatParams params, cudaStream_t stream) {
+ const dim3 gridSize((cam.width + T_PER_BLOCK - 1)/T_PER_BLOCK, (cam.height + T_PER_BLOCK - 1)/T_PER_BLOCK);
+ const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
+
+ dibr_merge_kernel<<<gridSize, blockSize, 0, stream>>>(depth_out, transform, cam, params);
+ cudaSafeCall( cudaGetLastError() );
+}
+
//==============================================================================
diff --git a/components/renderers/cpp/src/splatter_cuda.hpp b/components/renderers/cpp/src/splatter_cuda.hpp
index 5d027b2c6ff0c12efa28ca277d0404c11c5eb49c..33f0340f2a94007a87981dbf400e953869624392 100644
--- a/components/renderers/cpp/src/splatter_cuda.hpp
+++ b/components/renderers/cpp/src/splatter_cuda.hpp
@@ -15,6 +15,14 @@ namespace cuda {
const ftl::rgbd::Camera &camera,
cudaStream_t stream);
+ void screen_coord(
+ ftl::cuda::TextureObject<float> &depth_out,
+ ftl::cuda::TextureObject<short2> &screen_out,
+ const ftl::render::SplatParams ¶ms,
+ const float4x4 &pose,
+ const ftl::rgbd::Camera &camera,
+ cudaStream_t stream);
+
void triangle_render1(
ftl::cuda::TextureObject<float> &depth_in,
ftl::cuda::TextureObject<int> &depth_out,
@@ -51,6 +59,13 @@ namespace cuda {
ftl::render::SplatParams params,
cudaStream_t stream);
+ void dibr_merge(
+ ftl::cuda::TextureObject<int> &depth_out,
+ const float4x4 &transform,
+ const ftl::rgbd::Camera &cam,
+ ftl::render::SplatParams params,
+ cudaStream_t stream);
+
template <typename T>
void splat(
ftl::cuda::TextureObject<float4> &normals,
@@ -93,6 +108,16 @@ namespace cuda {
const ftl::rgbd::Camera &camera,
const float4x4 &poseInv, cudaStream_t stream);
+ template <typename A, typename B>
+ void reproject(
+ ftl::cuda::TextureObject<A> &in, // Original colour image
+ ftl::cuda::TextureObject<float> &depth_in, // Virtual depth map
+ ftl::cuda::TextureObject<B> &out, // Accumulated output
+ ftl::cuda::TextureObject<float> &contrib,
+ const ftl::render::SplatParams ¶ms,
+ const ftl::rgbd::Camera &camera,
+ const float4x4 &poseInv, cudaStream_t stream);
+
void equirectangular_reproject(
ftl::cuda::TextureObject<uchar4> &image_in,
ftl::cuda::TextureObject<uchar4> &image_out,
diff --git a/components/renderers/cpp/src/triangle_render.cu b/components/renderers/cpp/src/triangle_render.cu
index 61f44f320e0b889649d1f82c5f2c090d326e1221..fbbeee9fe3636aaa801366287f4f0f6361f9b5b8 100644
--- a/components/renderers/cpp/src/triangle_render.cu
+++ b/components/renderers/cpp/src/triangle_render.cu
@@ -145,8 +145,9 @@ float getZAtCoordinate(const float3 &barycentricCoord, const float (&tri)[3]) {
// Remove really large triangles
if ((maxX - minX) * (maxY - minY) > params.triangle_limit) return;
- for (int sy=minY; sy <= maxY; ++sy) {
- for (int sx=minX; sx <= maxX; ++sx) {
+ // TODO: Verify that < is correct, was <= before but < is faster.
+ for (int sy=minY; sy < maxY; ++sy) {
+ for (int sx=minX; sx < maxX; ++sx) {
if (sx >= params.camera.width || sx < 0 || sy >= params.camera.height || sy < 0) continue;
float3 baryCentricCoordinate = calculateBarycentricCoordinate(v, make_short2(sx, sy));
@@ -154,6 +155,7 @@ float getZAtCoordinate(const float3 &barycentricCoord, const float (&tri)[3]) {
if (isBarycentricCoordInBounds(baryCentricCoordinate)) {
float new_depth = getZAtCoordinate(baryCentricCoordinate, d);
atomicMin(&depth_out(sx,sy), int(new_depth*100000.0f));
+ //depth_out(sx,sy) = int(new_depth*100000.0f);
}
}
}
diff --git a/components/rgbd-sources/CMakeLists.txt b/components/rgbd-sources/CMakeLists.txt
index e611e97d55a89fef5a83f0bd1d5520eee5b62ef3..c45597f8b6777c972b8754fc5207daef0018d0eb 100644
--- a/components/rgbd-sources/CMakeLists.txt
+++ b/components/rgbd-sources/CMakeLists.txt
@@ -14,6 +14,8 @@ set(RGBDSRC
#src/sources/virtual/virtual.cpp
#src/sources/ftlfile/file_source.cpp
#src/sources/ftlfile/player.cpp
+ src/sources/screencapture/screencapture.cpp
+ src/camera.cpp
)
if (HAVE_REALSENSE)
@@ -41,7 +43,7 @@ if (CUDA_FOUND)
set_property(TARGET ftlrgbd PROPERTY CUDA_SEPARABLE_COMPILATION OFF)
endif()
-target_link_libraries(ftlrgbd ftlcommon ${OpenCV_LIBS} ${LIBSGM_LIBRARIES} ${CUDA_LIBRARIES} Eigen3::Eigen realsense ftlnet ${LibArchive_LIBRARIES} ftlcodecs ftloperators ftldata)
+target_link_libraries(ftlrgbd ftlcommon ${OpenCV_LIBS} ${LIBSGM_LIBRARIES} ${CUDA_LIBRARIES} Eigen3::Eigen realsense ftlnet ${LibArchive_LIBRARIES} ftlcodecs ftloperators ftldata ${X11_X11_LIB} ${X11_Xext_LIB})
if (BUILD_TESTS)
add_subdirectory(test)
diff --git a/components/rgbd-sources/include/ftl/rgbd/camera.hpp b/components/rgbd-sources/include/ftl/rgbd/camera.hpp
index 82a981d736e97ed42f403e62ef67f2f3a1dd2909..772a884b0b47b963d3ba8db0dcc8a6532cd83412 100644
--- a/components/rgbd-sources/include/ftl/rgbd/camera.hpp
+++ b/components/rgbd-sources/include/ftl/rgbd/camera.hpp
@@ -8,6 +8,7 @@
#ifndef __CUDACC__
#include <ftl/utility/msgpack.hpp>
+#include <ftl/configurable.hpp>
#endif
namespace ftl{
@@ -18,16 +19,16 @@ namespace rgbd {
* operate on CPU and GPU.
*/
struct __align__(16) Camera {
- double fx; // Focal length X
- double fy; // Focal length Y (usually same as fx)
- double cx; // Principle point Y
- double cy; // Principle point Y
+ float fx; // Focal length X
+ float fy; // Focal length Y (usually same as fx)
+ float cx; // Principle point Y
+ float cy; // Principle point Y
unsigned int width; // Pixel width
unsigned int height; // Pixel height
- double minDepth; // Near clip in meters
- double maxDepth; // Far clip in meters
- double baseline; // For stereo pair
- double doffs; // Disparity offset
+ float minDepth; // Near clip in meters
+ float maxDepth; // Far clip in meters
+ float baseline; // For stereo pair
+ float doffs; // Disparity offset
Camera scaled(int width, int height) const;
@@ -43,6 +44,11 @@ struct __align__(16) Camera {
#ifndef __CUDACC__
MSGPACK_DEFINE(fx,fy,cx,cy,width,height,minDepth,maxDepth,baseline,doffs);
+
+ /**
+ * Make a camera struct from a configurable.
+ */
+ static Camera from(ftl::Configurable*);
#endif
};
diff --git a/components/rgbd-sources/src/camera.cpp b/components/rgbd-sources/src/camera.cpp
new file mode 100644
index 0000000000000000000000000000000000000000..90d77f34ba1226405e5cb1c0a5b98459033294d6
--- /dev/null
+++ b/components/rgbd-sources/src/camera.cpp
@@ -0,0 +1,18 @@
+#include <ftl/rgbd/camera.hpp>
+
+using ftl::rgbd::Camera;
+
+Camera Camera::from(ftl::Configurable *cfg) {
+ Camera r;
+ r.width = cfg->value("width", 1280);
+ r.height = cfg->value("height", 720);
+ r.fx = cfg->value("focal", 700.0f);
+ r.fy = r.fx;
+ r.cx = -(float)r.width / 2.0f;
+ r.cy = -(float)r.height / 2.0f;
+ r.minDepth = cfg->value("minDepth", 0.1f);
+ r.maxDepth = cfg->value("maxDepth", 15.0f);
+ r.doffs = 0;
+ r.baseline = cfg->value("baseline", 0.05f);
+ return r;
+}
diff --git a/components/rgbd-sources/src/source.cpp b/components/rgbd-sources/src/source.cpp
index a71ea24e9ed6106425bf88f3726f226416a39e76..84a20c710b61d2282b658ea341005c789e0910a6 100644
--- a/components/rgbd-sources/src/source.cpp
+++ b/components/rgbd-sources/src/source.cpp
@@ -6,6 +6,7 @@
#include "sources/stereovideo/stereovideo.hpp"
#include "sources/image/image.hpp"
#include "sources/middlebury/middlebury_source.hpp"
+#include "sources/screencapture/screencapture.hpp"
#ifdef HAVE_LIBARCHIVE
#include <ftl/rgbd/snapshot.hpp>
@@ -28,6 +29,7 @@ using ftl::rgbd::detail::StereoVideoSource;
//using ftl::rgbd::detail::NetSource;
using ftl::rgbd::detail::ImageSource;
using ftl::rgbd::detail::MiddleburySource;
+using ftl::rgbd::detail::ScreenCapture;
using ftl::rgbd::capability_t;
using ftl::codecs::Channel;
//using ftl::rgbd::detail::FileSource;
@@ -177,6 +179,8 @@ ftl::rgbd::detail::Source *Source::_createDeviceImpl(const ftl::URI &uri) {
#else
LOG(ERROR) << "You do not have 'librealsense2' installed";
#endif
+ } else if (uri.getPathSegment(0) == "screen") {
+ return new ScreenCapture(this);
} else {
/*params_.width = value("width", 1280);
params_.height = value("height", 720);
diff --git a/components/rgbd-sources/src/sources/screencapture/screencapture.cpp b/components/rgbd-sources/src/sources/screencapture/screencapture.cpp
new file mode 100644
index 0000000000000000000000000000000000000000..baf7091bbe1c04be7b4328d91b9d902ff58a646e
--- /dev/null
+++ b/components/rgbd-sources/src/sources/screencapture/screencapture.cpp
@@ -0,0 +1,140 @@
+#include "screencapture.hpp"
+
+#define LOGURU_REPLACE_GLOG 1
+#include <loguru.hpp>
+#include <ftl/threads.hpp>
+#include <ftl/rgbd/source.hpp>
+
+using ftl::rgbd::detail::ScreenCapture;
+using ftl::codecs::Channel;
+using cv::cuda::GpuMat;
+
+#ifdef HAVE_X11
+#include <X11/Xlib.h>
+#include <X11/Xutil.h>
+
+#include <X11/extensions/XShm.h>
+#include <sys/ipc.h>
+#include <sys/shm.h>
+
+namespace ftl {
+namespace rgbd {
+namespace detail {
+struct X11State {
+ Display* display;
+ Window root;
+ XWindowAttributes window_attributes;
+ Screen* screen = window_attributes.screen;
+ XShmSegmentInfo shminfo;
+ XImage* ximg;
+};
+}
+}
+}
+#endif
+
+ScreenCapture::ScreenCapture(ftl::rgbd::Source *host)
+ : ftl::rgbd::detail::Source(host) {
+ capabilities_ = kCapVideo;
+
+ const uint WIDTH = 1280;
+ const uint HEIGHT = 720;
+
+ ready_ = false;
+
+ #ifdef HAVE_X11
+
+ impl_state_ = new X11State;
+ auto &s = *impl_state_;
+
+ s.display = XOpenDisplay(NULL);
+ if (!s.display) {
+ LOG(ERROR) << "Could not open X11 display";
+ return;
+ }
+
+ s.root = DefaultRootWindow(s.display); // TODO: Could choose windows?
+
+ if (!XGetWindowAttributes(s.display, s.root, &s.window_attributes)) {
+ LOG(ERROR) << "Could not get X11 window attributes";
+ return;
+ }
+
+ s.screen = s.window_attributes.screen;
+ params_.width = s.window_attributes.width;
+ params_.height = s.window_attributes.height;
+
+ s.ximg = XShmCreateImage(s.display, DefaultVisualOfScreen(s.screen),
+ DefaultDepthOfScreen(s.screen), ZPixmap, NULL, &s.shminfo,
+ params_.width, params_.height);
+
+ // TODO: Can this happen?
+ if (!s.ximg) {
+ LOG(ERROR) << "Didn't get shared memory image from X11 screen";
+ return;
+ }
+
+ s.shminfo.shmid = shmget(IPC_PRIVATE, s.ximg->bytes_per_line * s.ximg->height, IPC_CREAT|0777);
+ s.shminfo.shmaddr = s.ximg->data = (char*)shmat(s.shminfo.shmid, 0, 0);
+ s.shminfo.readOnly = False;
+ if(s.shminfo.shmid < 0) {
+ LOG(ERROR) << "Fatal shminfo error!";
+ return;
+ }
+
+ if (!XShmAttach(impl_state_->display, &impl_state_->shminfo)) {
+ LOG(ERROR) << "X11 Shared Memory attach failure";
+ return;
+ }
+
+ ready_ = true;
+
+ #endif
+
+ params_.cx = -(params_.width / 2.0);
+ params_.cy = -(params_.height / 2.0);
+ params_.fx = 700.0;
+ params_.fy = 700.0;
+ params_.maxDepth = host_->value("depth", 1.0f);
+ params_.minDepth = 0.0f;
+ params_.doffs = 0.0;
+
+ state_.getLeft() = params_;
+
+}
+
+ScreenCapture::~ScreenCapture() {
+ #ifdef HAVE_X11
+ delete impl_state_;
+ #endif
+}
+
+bool ScreenCapture::retrieve() {
+ if (!ready_) return false;
+ cv::Mat img;
+
+ #ifdef HAVE_X11
+ XShmGetImage(impl_state_->display, impl_state_->root, impl_state_->ximg, 0, 0, 0x00ffffff);
+ img = cv::Mat(params_.height, params_.width, CV_8UC4, impl_state_->ximg->data);
+ #endif
+
+ frame_.reset();
+ frame_.setOrigin(&state_);
+
+ if (!img.empty()) {
+ frame_.create<cv::cuda::GpuMat>(Channel::Colour).upload(img);
+ }
+
+ return true;
+}
+
+bool ScreenCapture::compute(int n, int b) {
+ if (!ready_) return false;
+ host_->notify(timestamp_, frame_);
+ return true;
+}
+
+bool ScreenCapture::isReady() {
+ return ready_;
+}
+
diff --git a/components/rgbd-sources/src/sources/screencapture/screencapture.hpp b/components/rgbd-sources/src/sources/screencapture/screencapture.hpp
new file mode 100644
index 0000000000000000000000000000000000000000..e1511cb9abf39b47fe14d84a55697e0199761ab8
--- /dev/null
+++ b/components/rgbd-sources/src/sources/screencapture/screencapture.hpp
@@ -0,0 +1,40 @@
+#ifndef _FTL_RGBD_SCREENCAPTURE_HPP_
+#define _FTL_RGBD_SCREENCAPTURE_HPP_
+
+#include <ftl/rgbd/detail/source.hpp>
+#include <ftl/config.h>
+
+namespace ftl {
+
+namespace rgbd {
+
+namespace detail {
+
+#ifdef HAVE_X11
+struct X11State;
+typedef X11State ImplState;
+#else
+typedef int ImplState;
+#endif
+
+class ScreenCapture : public ftl::rgbd::detail::Source {
+ public:
+ explicit ScreenCapture(ftl::rgbd::Source *host);
+ ~ScreenCapture();
+
+ bool capture(int64_t ts) { timestamp_ = ts; return true; };
+ bool retrieve();
+ bool compute(int n=-1, int b=-1);
+ bool isReady();
+
+ private:
+ bool ready_;
+
+ ImplState *impl_state_;
+};
+
+}
+}
+}
+
+#endif // _FTL_RGBD_SCREENCAPTURE_HPP_
diff --git a/components/rgbd-sources/src/sources/stereovideo/stereovideo.cpp b/components/rgbd-sources/src/sources/stereovideo/stereovideo.cpp
index fd97848563536de22cde51796c4350795f5eaac9..7d5fb662fe467a07912fd6a4f208af5d5aadc11d 100644
--- a/components/rgbd-sources/src/sources/stereovideo/stereovideo.cpp
+++ b/components/rgbd-sources/src/sources/stereovideo/stereovideo.cpp
@@ -192,19 +192,19 @@ void StereoVideoSource::updateParameters() {
cv::Mat K;
// same for left and right
- double baseline = 1.0 / calib_->getQ().at<double>(3,2);
- double doff = -calib_->getQ().at<double>(3,3) * baseline;
- double min_depth = this->host_->getConfig().value<double>("min_depth", 0.0);
- double max_depth = this->host_->getConfig().value<double>("max_depth", 15.0);
+ float baseline = static_cast<float>(1.0 / calib_->getQ().at<double>(3,2));
+ float doff = static_cast<float>(-calib_->getQ().at<double>(3,3) * baseline);
+ float min_depth = this->host_->getConfig().value<double>("min_depth", 0.0);
+ float max_depth = this->host_->getConfig().value<double>("max_depth", 15.0);
// left
K = calib_->getCameraMatrixLeft(color_size_);
state_.getLeft() = {
- K.at<double>(0,0), // Fx
- K.at<double>(1,1), // Fy
- -K.at<double>(0,2), // Cx
- -K.at<double>(1,2), // Cy
+ static_cast<float>(K.at<double>(0,0)), // Fx
+ static_cast<float>(K.at<double>(1,1)), // Fy
+ static_cast<float>(-K.at<double>(0,2)), // Cx
+ static_cast<float>(-K.at<double>(1,2)), // Cy
(unsigned int) color_size_.width,
(unsigned int) color_size_.height,
min_depth,
@@ -223,10 +223,10 @@ void StereoVideoSource::updateParameters() {
K = calib_->getCameraMatrixRight(color_size_);
state_.getRight() = {
- K.at<double>(0,0), // Fx
- K.at<double>(1,1), // Fy
- -K.at<double>(0,2), // Cx
- -K.at<double>(1,2), // Cy
+ static_cast<float>(K.at<double>(0,0)), // Fx
+ static_cast<float>(K.at<double>(1,1)), // Fy
+ static_cast<float>(-K.at<double>(0,2)), // Cx
+ static_cast<float>(-K.at<double>(1,2)), // Cy
(unsigned int) color_size_.width,
(unsigned int) color_size_.height,
min_depth,
diff --git a/components/rgbd-sources/test/source_unit.cpp b/components/rgbd-sources/test/source_unit.cpp
index 4ff46d7649694055bade34278feeb74042f00044..c1dbd76f5daf060f20fb8f09a7e1357dc783e4c7 100644
--- a/components/rgbd-sources/test/source_unit.cpp
+++ b/components/rgbd-sources/test/source_unit.cpp
@@ -45,6 +45,18 @@ class ImageSource : public ftl::rgbd::detail::Source {
bool isReady() { return true; };
};
+class ScreenCapture : public ftl::rgbd::detail::Source {
+ public:
+ explicit ScreenCapture(ftl::rgbd::Source *host) : ftl::rgbd::detail::Source(host) {
+ last_type = "screen";
+ }
+
+ bool capture(int64_t ts) { return true; }
+ bool retrieve() { return true; }
+ bool compute(int n, int b) { return true; };
+ bool isReady() { return true; };
+};
+
class StereoVideoSource : public ftl::rgbd::detail::Source {
public:
explicit StereoVideoSource(ftl::rgbd::Source *host) : ftl::rgbd::detail::Source(host) {
@@ -133,6 +145,7 @@ class MiddleburySource : public ftl::rgbd::detail::Source {
#define _FTL_RGBD_SNAPSHOT_SOURCE_HPP_
#define _FTL_RGBD_IMAGE_HPP_
#define _FTL_RGBD_REALSENSE_HPP_
+#define _FTL_RGBD_SCREENCAPTURE_HPP_
#define _FTL_RGBD_MIDDLEBURY_SOURCE_HPP_
#define _FTL_RGBD_FILE_SOURCE_HPP_
diff --git a/components/streams/src/injectors.cpp b/components/streams/src/injectors.cpp
index c9ccce40121ca3789bbf030ccc6a037ac760f4d7..2a58e5eac0fe174233b39a4a7ba501f3a54e3795 100644
--- a/components/streams/src/injectors.cpp
+++ b/components/streams/src/injectors.cpp
@@ -62,7 +62,7 @@ void ftl::stream::injectCalibration(ftl::stream::Stream *stream, const ftl::rgbd
ftl::codecs::StreamPacket spkt = {
4,
ts,
- fsid,
+ static_cast<uint8_t>(fsid),
static_cast<uint8_t>(ix),
(right) ? Channel::Calibration2 : Channel::Calibration
};
diff --git a/components/streams/src/netstream.cpp b/components/streams/src/netstream.cpp
index d9533de4958646edc3e0b03bbfbb04f1023cdfb3..bd16f8c0a6becce1bcc3e6a1ba4dcfb047199065 100644
--- a/components/streams/src/netstream.cpp
+++ b/components/streams/src/netstream.cpp
@@ -81,7 +81,7 @@ bool Net::post(const ftl::codecs::StreamPacket &spkt, const ftl::codecs::Packet
// Quality filter the packets
if (pkt.bitrate > 0 && pkt.bitrate != client.quality) {
//++c;
- LOG(INFO) << "Incorrect quality: " << (int)pkt.bitrate << " but requested " << (int)client.quality;
+ //LOG(INFO) << "Incorrect quality: " << (int)pkt.bitrate << " but requested " << (int)client.quality;
//continue;
}
diff --git a/components/streams/src/receiver.cpp b/components/streams/src/receiver.cpp
index 6ef07b742476ab3fd86a8d63cfcc911f30e534df..6338b1943520ef992be559f36ec2f22ac1987fdf 100644
--- a/components/streams/src/receiver.cpp
+++ b/components/streams/src/receiver.cpp
@@ -259,7 +259,8 @@ void Receiver::_processVideo(const StreamPacket &spkt, const Packet &pkt) {
InternalVideoStates &vidstate = _getVideoFrame(spkt,i);
auto &frame = builder_[spkt.streamID].get(spkt.timestamp, spkt.frame_number+i);
- auto sel = stream_->selected(spkt.frameSetID());
+ const auto *cs = stream_;
+ auto sel = stream_->selected(spkt.frameSetID()) & cs->available(spkt.frameSetID());
frame.create<cv::cuda::GpuMat>(spkt.channel);
diff --git a/components/streams/test/receiver_unit.cpp b/components/streams/test/receiver_unit.cpp
index d683d4abddfc5f3e764150aa50710fa242718829..062dce960ab7daab5fec0ec6a354c9247f1a5869 100644
--- a/components/streams/test/receiver_unit.cpp
+++ b/components/streams/test/receiver_unit.cpp
@@ -306,21 +306,21 @@ TEST_CASE( "Receiver sync bugs" ) {
return true;
});
- stream.post(spkt, pkt);
+ try { stream.post(spkt, pkt); } catch(...) {}
spkt.timestamp = 10;
spkt.channel = Channel::ColourHighRes;
- stream.post(spkt, pkt);
+ try { stream.post(spkt, pkt); } catch(...) {}
spkt.timestamp = 20;
spkt.channel = Channel::Colour2;
- stream.post(spkt, pkt);
+ try { stream.post(spkt, pkt); } catch(...) {}
spkt.timestamp = 20;
spkt.channel = Channel::Colour;
- stream.post(spkt, pkt);
+ try { stream.post(spkt, pkt); } catch(...) {}
int i=10;
- while (i-- > 0 && count < 1) std::this_thread::sleep_for(std::chrono::milliseconds(10));
+ while (i-- > 0 && count < 2) std::this_thread::sleep_for(std::chrono::milliseconds(10));
- REQUIRE( count == 1 );
+ REQUIRE( count == 2 );
REQUIRE( ts == 20 );
REQUIRE( !haswrongchan );
}