diff --git a/CMakeLists.txt b/CMakeLists.txt index fb73e5cf9d35ad7de7e1a2cbf66826d802c0b749..6640f5d7ee46c3511254604ae0840e49cb9e91d6 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -230,6 +230,7 @@ add_subdirectory(applications/calibration) add_subdirectory(applications/groupview) add_subdirectory(applications/player) add_subdirectory(applications/recorder) +add_subdirectory(applications/merger) if (HAVE_AVFORMAT) add_subdirectory(applications/ftl2mkv) diff --git a/applications/gui/src/main.cpp b/applications/gui/src/main.cpp index 7875e52bb04afc59953ac1a728dafe64008b6d6c..5ebc552fa0ee0664c7a0b7898a723b90bf87f29c 100644 --- a/applications/gui/src/main.cpp +++ b/applications/gui/src/main.cpp @@ -25,10 +25,12 @@ int main(int argc, char **argv) { std::map<ftl::UUID, std::vector<ftl::NetConfigurable*>> peerConfigurables; + // FIXME: Move this elsewhere, it is not just for GUI net->onConnect([&controller, &peerConfigurables](ftl::net::Peer *p) { ftl::UUID peer = p->id(); auto cs = controller->getConfigurables(peer); for (auto c : cs) { + //LOG(INFO) << "NET CONFIG: " << c; ftl::config::json_t *configuration = new ftl::config::json_t; *configuration = controller->get(peer, c); if (!configuration->empty()) { diff --git a/applications/gui/src/media_panel.cpp b/applications/gui/src/media_panel.cpp index 142eb9c93a42c4b9403596c78fd31eedbd289bfe..693024bd3c6a89ccb0ac6e893b01641dbd77c574 100644 --- a/applications/gui/src/media_panel.cpp +++ b/applications/gui/src/media_panel.cpp @@ -98,8 +98,9 @@ MediaPanel::MediaPanel(ftl::gui::Screen *screen) : nanogui::Window(screen, ""), button = new Button(this, "", ENTYPO_ICON_CONTROLLER_PAUS); button->setCallback([this,button]() { - paused_ = !paused_; - screen_->control()->pause(); + //paused_ = !paused_; + paused_ = !(bool)ftl::config::get("[reconstruction]/controls/paused"); + ftl::config::update("[reconstruction]/controls/paused", paused_); if (paused_) { button->setIcon(ENTYPO_ICON_CONTROLLER_PLAY); } else { diff --git a/applications/gui/src/screen.cpp b/applications/gui/src/screen.cpp index 811d96cf851781add465e4ec58ee97e93e5c98f2..76665281fc1b9fcbb1b0b45609c509d02f88a69d 100644 --- a/applications/gui/src/screen.cpp +++ b/applications/gui/src/screen.cpp @@ -84,6 +84,8 @@ ftl::gui::Screen::Screen(ftl::Configurable *proot, ftl::net::Universe *pnet, ftl pos_y_ = root_->value("position_y", 0.0f); }); + shortcuts_ = ftl::create<ftl::Configurable>(root_, "shortcuts"); + setSize(Vector2i(1280,720)); toolbuttheme = new Theme(*theme()); @@ -434,12 +436,31 @@ bool ftl::gui::Screen::mouseButtonEvent(const nanogui::Vector2i &p, int button, } } +static std::string generateKeyComboStr(int key, int modifiers) { + std::string res = ""; + + switch(modifiers) { + case 1: res += "Shift+"; break; + case 2: res += "Ctrl+"; break; + case 3: res += "Ctrl+Shift+"; break; + case 4: res += "Alt+"; break; + default: break; + } + + if (key < 127 && key >= 32) { + char buf[2] = { (char)key, 0 }; + return res + std::string(buf); + } else { + return ""; + } +} + bool ftl::gui::Screen::keyboardEvent(int key, int scancode, int action, int modifiers) { using namespace Eigen; if (nanogui::Screen::keyboardEvent(key, scancode, action, modifiers)) { return true; } else { - LOG(INFO) << "Key press " << key << " - " << action << " - " << modifiers; + //LOG(INFO) << "Key press " << key << " - " << action << " - " << modifiers; if (key >= 262 && key <= 267) { if (camera_) camera_->keyMovement(key, modifiers); @@ -447,9 +468,42 @@ bool ftl::gui::Screen::keyboardEvent(int key, int scancode, int action, int modi } else if (action == 1 && key == 'H') { swindow_->setVisible(false); cwindow_->setVisible(false); - } else if (action == 1 && key == 32) { - ctrl_->pause(); - return true; + } else if (action == 1) { + std::string combo = generateKeyComboStr(key, modifiers); + + if (combo.size() > 0) { + LOG(INFO) << "Key combo = " << combo; + + auto s = shortcuts_->get<nlohmann::json>(combo); + if (s) { + //LOG(INFO) << "FOUND KEYBOARD SHORTCUT"; + std::string op = (*s).value("op",std::string("=")); + std::string uri = (*s).value("uri",std::string("")); + + if (op == "toggle") { + auto v = ftl::config::get(uri); + if (v.is_boolean()) { + ftl::config::update(uri, !v.get<bool>()); + } + } else if (op == "+=") { + auto v = ftl::config::get(uri); + if (v.is_number_float()) { + ftl::config::update(uri, v.get<float>() + (*s).value("value",0.0f)); + } else if (v.is_number_integer()) { + ftl::config::update(uri, v.get<int>() + (*s).value("value",0)); + } + } else if (op == "-=") { + auto v = ftl::config::get(uri); + if (v.is_number_float()) { + ftl::config::update(uri, v.get<float>() - (*s).value("value",0.0f)); + } else if (v.is_number_integer()) { + ftl::config::update(uri, v.get<int>() - (*s).value("value",0)); + } + } else if (op == "=") { + ftl::config::update(uri, (*s)["value"]); + } + } + } } return false; } diff --git a/applications/gui/src/screen.hpp b/applications/gui/src/screen.hpp index 9dbed7ead2c064394e1eb877a6a9043ed5b364b4..90cd519bb3681126f4dc4f0d258e342953562433 100644 --- a/applications/gui/src/screen.hpp +++ b/applications/gui/src/screen.hpp @@ -101,6 +101,8 @@ class Screen : public nanogui::Screen { bool show_two_images_ = false; + ftl::Configurable *shortcuts_; + #ifdef HAVE_OPENVR vr::IVRSystem *HMD_; #endif diff --git a/applications/merger/CMakeLists.txt b/applications/merger/CMakeLists.txt new file mode 100644 index 0000000000000000000000000000000000000000..ccc625afb0241966628bf8bffa790cf8a041afa7 --- /dev/null +++ b/applications/merger/CMakeLists.txt @@ -0,0 +1,11 @@ +set(FTLMERGER + src/main.cpp +) + +add_executable(ftl-merge ${FTLMERGER}) + +target_include_directories(ftl-merge PRIVATE src) + +target_link_libraries(ftl-merge ftlcommon ftlcodecs ftlrgbd Threads::Threads ${OpenCV_LIBS}) + + diff --git a/applications/merger/src/main.cpp b/applications/merger/src/main.cpp new file mode 100644 index 0000000000000000000000000000000000000000..d7b3f2b92d5976082f1995420cc6486061b714b2 --- /dev/null +++ b/applications/merger/src/main.cpp @@ -0,0 +1,104 @@ +#include <loguru.hpp> +#include <ftl/configuration.hpp> +#include <ftl/codecs/reader.hpp> +#include <ftl/codecs/writer.hpp> +#include <ftl/codecs/packet.hpp> +#include <ftl/rgbd/camera.hpp> +#include <ftl/codecs/hevc.hpp> + +#include <fstream> + +int main(int argc, char **argv) { + auto root = ftl::configure(argc, argv, "merger_default"); + + std::string outputfile = root->value("out", std::string("output.ftl")); + std::vector<std::string> paths = *root->get<std::vector<std::string>>("paths"); + int timeoff = int(root->value("offset", 0.0f) * 1000.0f); + int stream_mask1 = root->value("mask1",0xFF); + int stream_mask2 = root->value("mask2",0xFF); + + if (paths.size() == 0) { + LOG(ERROR) << "Missing input ftl file(s)."; + return -1; + } + + // Generate the output writer... + std::ofstream of; + of.open(outputfile); + if (!of.is_open()) { + LOG(ERROR) << "Could not open output file: " << outputfile; + return -1; + } + + ftl::codecs::Writer out(of); + out.begin(); + + std::vector<std::ifstream> fs; + std::vector<ftl::codecs::Reader*> rs; + fs.resize(paths.size()); + rs.resize(paths.size()); + + for (size_t i=0; i<paths.size(); ++i) { + fs[i].open(paths[i]); + if (!fs[i].is_open()) { + LOG(ERROR) << "Could not open file: " << paths[i]; + return -1; + } + + LOG(INFO) << "Opening("<< i <<"): " << paths[i]; + + rs[i] = new ftl::codecs::Reader(fs[i]); + if (!rs[i]->begin()) { + LOG(ERROR) << "Bad ftl file format"; + return -1; + } + } + + std::map<int,int> idmap; + int lastid = 0; + + bool res = rs[0]->read(90000000000000, [&rs,&out,&idmap,&lastid,stream_mask1,stream_mask2,timeoff](const ftl::codecs::StreamPacket &spkt, const ftl::codecs::Packet &pkt) { + if (((0x1 << spkt.streamID) & stream_mask1) == 0) return; + + ftl::codecs::StreamPacket spkt2 = spkt; + if (idmap.find(spkt.streamID) == idmap.end()) { + idmap[spkt.streamID] = lastid++; + } + spkt2.streamID = idmap[spkt.streamID]; + + // Now read all other sources up to the same packet timestamp. + out.write(spkt2, pkt); + + for (size_t j=1; j<rs.size(); ++j) { + ftl::codecs::Reader *r = rs[j]; + + // FIXME: Need to truncate other stream if the following returns + // no frames, meaning the timeshift causes this stream to run out + // before the main stream. + rs[j]->read(spkt.timestamp+timeoff+1, [&out,&idmap,&lastid,j,r,stream_mask2,timeoff](const ftl::codecs::StreamPacket &spkt, const ftl::codecs::Packet &pkt) { + if (((0x1 << spkt.streamID) & stream_mask2) == 0) return; + if (int(spkt.channel) < 32 && spkt.timestamp < r->getStartTime()+timeoff) return; + + ftl::codecs::StreamPacket spkt2 = spkt; + if (idmap.find(spkt.streamID + (j << 16)) == idmap.end()) { + idmap[spkt.streamID+(j << 16)] = lastid++; + } + spkt2.streamID = idmap[spkt.streamID + (j << 16)]; + spkt2.timestamp -= timeoff; + + out.write(spkt2, pkt); + }); + } + }); + + out.end(); + of.close(); + + for (size_t i=0; i<rs.size(); ++i) { + rs[i]->end(); + delete rs[i]; + fs[i].close(); + } + + return 0; +} \ No newline at end of file diff --git a/applications/reconstruct/src/main.cpp b/applications/reconstruct/src/main.cpp index 551a2dfe309bf8a5a29c17cc4856acb48f70f09c..f3a92c93340de67bbee302a89dae93f01b550bcd 100644 --- a/applications/reconstruct/src/main.cpp +++ b/applications/reconstruct/src/main.cpp @@ -241,7 +241,7 @@ 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"); - root->set("tags", nlohmann::json::array({ root->getID()+"/virtual" })); + //root->set("tags", nlohmann::json::array({ root->getID()+"/virtual" })); ftl::render::Triangular *splat = ftl::create<ftl::render::Triangular>(root, "renderer", &scene_B); ftl::rgbd::Group *group = new ftl::rgbd::Group; ftl::ILW *align = ftl::create<ftl::ILW>(root, "merge"); @@ -331,6 +331,7 @@ static void run(ftl::Configurable *root) { //pipeline1->append<ftl::operators::ScanFieldFill>("filling"); // Generate a smoothing channel pipeline1->append<ftl::operators::CrossSupport>("cross"); pipeline1->append<ftl::operators::DiscontinuityMask>("discontinuity"); + pipeline1->append<ftl::operators::CrossSupport>("cross2")->set("discon_support", true); pipeline1->append<ftl::operators::CullDiscontinuity>("remove_discontinuity"); //pipeline1->append<ftl::operators::AggreMLS>("mls"); // Perform MLS (using smoothing channel) pipeline1->append<ftl::operators::VisCrossSupport>("viscross")->set("enabled", false); diff --git a/components/common/cpp/include/ftl/configuration.hpp b/components/common/cpp/include/ftl/configuration.hpp index 9e4ba8aed4bcf8b5f2c8cec2c1079edb945359c3..18aaf89f9433911dd71ec3d3519ff2125317e78e 100644 --- a/components/common/cpp/include/ftl/configuration.hpp +++ b/components/common/cpp/include/ftl/configuration.hpp @@ -47,6 +47,8 @@ void removeConfigurable(Configurable *cfg); */ bool update(const std::string &puri, const json_t &value); +json_t &get(const std::string &puri); + /** * Resolve a JSON schema reference, but do not wait for a remote reference * if it is not available. A null entity is returned if not resolved. diff --git a/components/common/cpp/src/configuration.cpp b/components/common/cpp/src/configuration.cpp index aa236c8a06772e71d59edd3886ce7a9cb14e2eec..efeea93e6789c8b30fca1f1a67838e26e08546ac 100644 --- a/components/common/cpp/src/configuration.cpp +++ b/components/common/cpp/src/configuration.cpp @@ -194,6 +194,7 @@ ftl::Configurable *ftl::config::find(const std::string &uri) { actual_uri = rootCFG->getID() + uri; } } + auto ix = config_instance.find(actual_uri); if (ix == config_instance.end()) return nullptr; else return (*ix).second; @@ -228,6 +229,7 @@ void ftl::config::registerConfigurable(ftl::Configurable *cfg) { auto tags = cfg->get<vector<string>>("tags"); if (tags) { for (auto &t : *tags) { + //LOG(INFO) << "REGISTER TAG: " << t; tag_index[t].push_back(cfg); } } @@ -236,22 +238,47 @@ void ftl::config::registerConfigurable(ftl::Configurable *cfg) { json_t null_json; +/* To allow for custom tag format */ +static std::string preprocessURI(const std::string &uri) { + if (uri[0] == '[') { + size_t closing = uri.find_last_of(']'); + string tags = uri.substr(1, closing-1); + + // TODO: Allow for multiple tags + + const auto &cfgs = ftl::config::findByTag(tags); + + // FIXME: Check for more than one tag result + if (cfgs.size() > 0) { + //LOG(INFO) << "PREPROC URI " << cfgs[0]->getID() + uri.substr(closing+1); + return cfgs[0]->getID() + uri.substr(closing+1); + } else { + return uri; + } + } else if (uri[0] == '/') { + return rootCFG->getID() + uri; + } else { + return uri; + } +} + bool ftl::config::update(const std::string &puri, const json_t &value) { // Remove last component of URI string tail = ""; string head = ""; - size_t last_hash = puri.find_last_of('#'); + string uri = preprocessURI(puri); + size_t last_hash = uri.find_last_of('#'); if (last_hash != string::npos) { - size_t last = puri.find_last_of('/'); + size_t last = uri.find_last_of('/'); if (last != string::npos && last > last_hash) { - tail = puri.substr(last+1); - head = puri.substr(0, last); + tail = uri.substr(last+1); + head = uri.substr(0, last); } else { - tail = puri.substr(last_hash+1); - head = puri.substr(0, last_hash); + tail = uri.substr(last_hash+1); + head = uri.substr(0, last_hash); } } else { - LOG(WARNING) << "Expected a # in an update URI: " << puri; + LOG(WARNING) << "Expected a # in an update URI: " << uri; return false; } @@ -284,6 +311,35 @@ bool ftl::config::update(const std::string &puri, const json_t &value) { } } +json_t &ftl::config::get(const std::string &puri) { + // Remove last component of URI + string tail = ""; + string head = ""; + string uri = preprocessURI(puri); + size_t last_hash = uri.find_last_of('#'); + if (last_hash != string::npos) { + size_t last = uri.find_last_of('/'); + if (last != string::npos && last > last_hash) { + tail = uri.substr(last+1); + head = uri.substr(0, last); + } else { + tail = uri.substr(last_hash+1); + head = uri.substr(0, last_hash); + } + } else { + LOG(WARNING) << "Expected a # in an update URI: " << uri; + return null_json; + } + + Configurable *cfg = find(head); + + if (cfg) { + return cfg->getConfig()[tail]; + } else { + return null_json; + } +} + json_t &ftl::config::resolve(const std::string &puri, bool eager) { string uri_str = puri; @@ -528,6 +584,7 @@ Configurable *ftl::config::configure(int argc, char **argv, const std::string &r vector<string> paths; while (argc-- > 0) { paths.push_back(argv[0]); + argv++; } if (!findConfiguration(options["config"], paths)) { diff --git a/components/operators/src/mvmls.cpp b/components/operators/src/mvmls.cpp index 4c02a7607f1d6a506c1f60cbd1f25eabd2cb5732..e85f8271149537c920b838e9885c3a406bbb5b5b 100644 --- a/components/operators/src/mvmls.cpp +++ b/components/operators/src/mvmls.cpp @@ -61,7 +61,7 @@ bool MultiViewMLS::apply(ftl::rgbd::FrameSet &in, ftl::rgbd::FrameSet &out, cuda LOG(ERROR) << "Required normals channel missing for MLS"; return false; } - if (!f.hasChannel(Channel::Support1)) { + if (!f.hasChannel(Channel::Support2)) { LOG(ERROR) << "Required cross support channel missing for MLS"; return false; } @@ -214,7 +214,7 @@ bool MultiViewMLS::apply(ftl::rgbd::FrameSet &in, ftl::rgbd::FrameSet &out, cuda } ftl::cuda::mls_aggr_horiz( - f.createTexture<uchar4>(Channel::Support1), + f.createTexture<uchar4>(Channel::Support2), f.createTexture<float4>(Channel::Normals), normals_horiz_[i], f.createTexture<float>(Channel::Depth), @@ -228,7 +228,7 @@ bool MultiViewMLS::apply(ftl::rgbd::FrameSet &in, ftl::rgbd::FrameSet &out, cuda ); ftl::cuda::mls_aggr_vert( - f.getTexture<uchar4>(Channel::Support1), + f.getTexture<uchar4>(Channel::Support2), normals_horiz_[i], f.getTexture<float4>(Channel::Normals), centroid_horiz_[i], diff --git a/components/operators/src/segmentation.cpp b/components/operators/src/segmentation.cpp index fb015b6b16e1aa7e1dcac9735103c8c6af41567e..9bf7605e1ea9ef6b5dd73a4c3078ba70bb61abb9 100644 --- a/components/operators/src/segmentation.cpp +++ b/components/operators/src/segmentation.cpp @@ -14,18 +14,17 @@ CrossSupport::~CrossSupport() { } bool CrossSupport::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd::Source *s, cudaStream_t stream) { - bool use_depth = config()->value("depth_region", false); + bool use_mask = config()->value("discon_support", false); - if (use_depth) { + if (use_mask) { ftl::cuda::support_region( - in.createTexture<float>(Channel::Depth), + in.createTexture<int>(Channel::Mask), out.createTexture<uchar4>(Channel::Support2, ftl::rgbd::Format<uchar4>(in.get<cv::cuda::GpuMat>(Channel::Colour).size())), - config()->value("depth_tau", 0.04f), config()->value("v_max", 5), config()->value("h_max", 5), - config()->value("symmetric", true), stream + config()->value("symmetric", false), stream ); - } //else { + } else { ftl::cuda::support_region( in.createTexture<uchar4>(Channel::Colour), out.createTexture<uchar4>(Channel::Support1, ftl::rgbd::Format<uchar4>(in.get<cv::cuda::GpuMat>(Channel::Colour).size())), @@ -34,7 +33,7 @@ bool CrossSupport::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd: config()->value("h_max", 5), config()->value("symmetric", true), stream ); - //} + } return true; } diff --git a/components/operators/src/segmentation.cu b/components/operators/src/segmentation.cu index c16e647931f7d4c4d92551f44f6242fc415bd91d..aaa81e2ad98b5171572b6768754b93f32abcd1eb 100644 --- a/components/operators/src/segmentation.cu +++ b/components/operators/src/segmentation.cu @@ -1,8 +1,10 @@ #include "segmentation_cuda.hpp" +#include "mask_cuda.hpp" #define T_PER_BLOCK 8 using ftl::cuda::TextureObject; +using ftl::cuda::Mask; template <typename T> __device__ inline float cross(T p1, T p2); @@ -89,6 +91,64 @@ __device__ uchar4 calculate_support_region(const TextureObject<T> &img, int x, i return result; } +__device__ uchar4 calculate_support_region(const TextureObject<int> &img, int x, int y, int v_max, int h_max) { + int x_min = max(0, x - h_max); + int x_max = min(img.width()-1, x + h_max); + int y_min = max(0, y - v_max); + int y_max = min(img.height()-1, y + v_max); + + uchar4 result = make_uchar4(0, 0, 0, 0); + + Mask m1(img.tex2D(x,y)); + + int u; + for (u=x-1; u >= x_min; --u) { + Mask m2(img.tex2D(u,y)); + if (m2.isDiscontinuity()) { + result.x = x - u - 1; + break; + } + } + if (u < x_min) result.x = x - x_min; + + for (u=x+1; u <= x_max; ++u) { + Mask m2(img.tex2D(u,y)); + if (m2.isDiscontinuity()) { + result.y = u - x - 1; + break; + } + } + if (u > x_max) result.y = x_max - x; + + int v; + for (v=y-1; v >= y_min; --v) { + Mask m2(img.tex2D(x,v)); + if (m2.isDiscontinuity()) { + result.z = y - v - 1; + break; + } + } + if (v < y_min) result.z = y - y_min; + + for (v=y+1; v <= y_max; ++v) { + Mask m2(img.tex2D(x,v)); + if (m2.isDiscontinuity()) { + result.w = v - y - 1; + break; + } + } + if (v > y_max) result.w = y_max - y; + + // Make symetric left/right and up/down + if (false) { + result.x = min(result.x, result.y); + result.y = result.x; + result.z = min(result.z, result.w); + result.w = result.z; + } + return result; +} + template <typename T, bool SYM> __global__ void support_region_kernel(TextureObject<T> img, TextureObject<uchar4> region, float tau, int v_max, int h_max) { const int x = blockIdx.x*blockDim.x + threadIdx.x; @@ -99,6 +159,15 @@ __global__ void support_region_kernel(TextureObject<T> img, TextureObject<uchar4 region(x,y) = calculate_support_region<T,SYM>(img, x, y, tau, v_max, h_max); } +__global__ void support_region_kernel(TextureObject<int> img, TextureObject<uchar4> region, int v_max, int h_max) { + const int x = blockIdx.x*blockDim.x + threadIdx.x; + const int y = blockIdx.y*blockDim.y + threadIdx.y; + + if (x < 0 || y < 0 || x >= img.width() || y >= img.height()) return; + + region(x,y) = calculate_support_region(img, x, y, v_max, h_max); +} + void ftl::cuda::support_region( ftl::cuda::TextureObject<uchar4> &colour, ftl::cuda::TextureObject<uchar4> ®ion, @@ -142,6 +211,26 @@ void ftl::cuda::support_region( #endif } +void ftl::cuda::support_region( + ftl::cuda::TextureObject<int> &mask, + ftl::cuda::TextureObject<uchar4> ®ion, + int v_max, + int h_max, + bool sym, + cudaStream_t stream) { + + const dim3 gridSize((region.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (region.height() + T_PER_BLOCK - 1)/T_PER_BLOCK); + const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK); + + support_region_kernel<<<gridSize, blockSize, 0, stream>>>(mask, region, v_max, h_max); + cudaSafeCall( cudaGetLastError() ); + + + #ifdef _DEBUG + cudaSafeCall(cudaDeviceSynchronize()); + #endif +} + __global__ void vis_support_region_kernel(TextureObject<uchar4> colour, TextureObject<uchar4> region, uchar4 bcolour, uchar4 acolour, int ox, int oy, int dx, int dy) { const int x = blockIdx.x*blockDim.x + threadIdx.x; diff --git a/components/operators/src/segmentation_cuda.hpp b/components/operators/src/segmentation_cuda.hpp index c2cb390d9c0ee62a33127eec1720cf0d6fee8cae..1383489337dc968a33ec630facb597920518e2bf 100644 --- a/components/operators/src/segmentation_cuda.hpp +++ b/components/operators/src/segmentation_cuda.hpp @@ -18,6 +18,12 @@ void support_region( float tau, int v_max, int h_max, bool sym, cudaStream_t stream); +void support_region( + ftl::cuda::TextureObject<int> &mask, + ftl::cuda::TextureObject<uchar4> ®ion, + int v_max, int h_max, bool sym, + cudaStream_t stream); + void vis_support_region( ftl::cuda::TextureObject<uchar4> &colour, ftl::cuda::TextureObject<uchar4> ®ion, diff --git a/components/renderers/cpp/include/ftl/cuda/normals.hpp b/components/renderers/cpp/include/ftl/cuda/normals.hpp index dc3d0265ce4c142861bb0ca0b2e841dc81887449..bbf690f4f66178297158dea35528ba01629445a1 100644 --- a/components/renderers/cpp/include/ftl/cuda/normals.hpp +++ b/components/renderers/cpp/include/ftl/cuda/normals.hpp @@ -42,6 +42,11 @@ void normal_visualise(ftl::cuda::TextureObject<float4> &norm, const float3 &light, const uchar4 &diffuse, const uchar4 &ambient, cudaStream_t stream); +void cool_blue(ftl::cuda::TextureObject<float4> &norm, + ftl::cuda::TextureObject<uchar4> &output, + const uchar4 &colouring, const float3x3 &pose, + cudaStream_t stream); + void normal_filter(ftl::cuda::TextureObject<float4> &norm, ftl::cuda::TextureObject<float4> &points, const ftl::rgbd::Camera &camera, const float4x4 &pose, diff --git a/components/renderers/cpp/src/normals.cu b/components/renderers/cpp/src/normals.cu index 97452555dfb61fd5015aa6a1745b39e2ddcca409..31034c0af6059f2b8014d7af780bd082bb34868a 100644 --- a/components/renderers/cpp/src/normals.cu +++ b/components/renderers/cpp/src/normals.cu @@ -399,6 +399,50 @@ void ftl::cuda::normal_visualise(ftl::cuda::TextureObject<float4> &norm, //============================================================================== +__global__ void cool_blue_kernel(ftl::cuda::TextureObject<float4> norm, + ftl::cuda::TextureObject<uchar4> output, + uchar4 colouring, float3x3 pose) { + const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x; + const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y; + + if(x >= norm.width() || y >= norm.height()) return; + + //output(x,y) = make_uchar4(0,0,0,0); + float3 ray = pose * make_float3(0.0f,0.0f,1.0f); + ray = ray / length(ray); + float3 n = make_float3(norm.tex2D((int)x,(int)y)); + float l = length(n); + if (l == 0) return; + n /= l; + + const float d = 1.0f - max(dot(ray, n), 0.0f); + uchar4 original = output(x,y); //.tex2D(x,y); + + output(x,y) = make_uchar4( + min(255.0f, colouring.x*d + original.x), + min(255.0f, colouring.y*d + original.y), + min(255.0f, colouring.z*d + original.z), 255); +} + +void ftl::cuda::cool_blue(ftl::cuda::TextureObject<float4> &norm, + ftl::cuda::TextureObject<uchar4> &output, + const uchar4 &colouring, const float3x3 &pose, + cudaStream_t stream) { + + const dim3 gridSize((norm.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (norm.height() + T_PER_BLOCK - 1)/T_PER_BLOCK); + const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK); + + cool_blue_kernel<<<gridSize, blockSize, 0, stream>>>(norm, output, colouring, pose); + + cudaSafeCall( cudaGetLastError() ); + #ifdef _DEBUG + cudaSafeCall(cudaDeviceSynchronize()); + //cutilCheckMsg(__FUNCTION__); + #endif +} + +//============================================================================== + __global__ void filter_normals_kernel(ftl::cuda::TextureObject<float4> norm, ftl::cuda::TextureObject<float4> output, ftl::rgbd::Camera camera, float4x4 pose, float thresh) { diff --git a/components/renderers/cpp/src/tri_render.cpp b/components/renderers/cpp/src/tri_render.cpp index 5912b6480b109db2c3cce960ae5bca6abe6531f8..5d38742cf979575fd24388689b6cb5899821af65 100644 --- a/components/renderers/cpp/src/tri_render.cpp +++ b/components/renderers/cpp/src/tri_render.cpp @@ -609,6 +609,18 @@ bool Triangular::render(ftl::rgbd::VirtualSource *src, ftl::rgbd::Frame &out) { // Reprojection of colours onto surface _renderChannel(out, Channel::Colour, Channel::Colour, stream_); + if (value("cool_effect", false)) { + auto pose = params.m_viewMatrixInverse.getFloat3x3(); + auto col = parseCUDAColour(value("cool_effect_colour", std::string("#2222ff"))); + + ftl::cuda::cool_blue( + out.getTexture<float4>(Channel::Normals), + out.getTexture<uchar4>(Channel::Colour), + col, pose, + stream_ + ); + } + if (value("show_bad_colour", false)) { ftl::cuda::show_missing_colour( out.getTexture<float>(Channel::Depth),