Skip to content
Snippets Groups Projects

Compare revisions

Changes are shown as if the source revision was being merged into the target revision. Learn more about comparing revisions.

Source

Select target project
No results found

Target

Select target project
  • nicolaspope/ftl
1 result
Show changes
Commits on Source (31)
Hide whitespace changes
Inline Side-by-side
Showing
with 458 additions and 348 deletions
applications/gui/src/camera.cpp
View file @ c9629aed
......@@ -148,17 +148,20 @@ ftl::gui::Camera::Camera(ftl::gui::Screen *screen, ftl::rgbd::Source *src) : scr
posewin_->setTheme(screen->windowtheme);
posewin_->setVisible(false);
src->setCallback([this](int64_t ts, cv::Mat &channel1, cv::Mat &channel2) {
src->setCallback([this](int64_t ts, cv::cuda::GpuMat &channel1, cv::cuda::GpuMat &channel2) {
UNIQUE_LOCK(mutex_, lk);
im1_.create(channel1.size(), channel1.type());
im2_.create(channel2.size(), channel2.type());
//cv::swap(channel1, im1_);
//cv::swap(channel2, im2_);
channel1.download(im1_);
channel2.download(im2_);
// OpenGL (0,0) bottom left
cv::flip(channel1, im1_, 0);
cv::flip(channel2, im2_, 0);
cv::flip(im1_, im1_, 0);
cv::flip(im2_, im2_, 0);
});
}
......
applications/gui/src/config_window.cpp
View file @ c9629aed
......@@ -15,8 +15,11 @@ using std::string;
using std::vector;
using ftl::config::json_t;
ConfigWindow::ConfigWindow(nanogui::Widget *parent, ftl::ctrl::Master *ctrl, const ftl::UUID &peer) : ConfigWindow(parent, ctrl, std::optional<ftl::UUID>(peer)) {
ConfigWindow::ConfigWindow(nanogui::Widget *parent, ftl::ctrl::Master *ctrl, const ftl::UUID &peer)
}
ConfigWindow::ConfigWindow(nanogui::Widget *parent, ftl::ctrl::Master *ctrl, const std::optional<ftl::UUID> &peer)
: nanogui::Window(parent, "Settings"), ctrl_(ctrl), peer_(peer) {
using namespace nanogui;
......@@ -24,78 +27,115 @@ ConfigWindow::ConfigWindow(nanogui::Widget *parent, ftl::ctrl::Master *ctrl, con
setPosition(Vector2i(parent->width()/2.0f - 100.0f, parent->height()/2.0f - 100.0f));
//setModal(true);
configurables_ = ctrl->getConfigurables(peer);
if (peer) {
configurables_ = ctrl->getConfigurables(peer.value());
} else {
configurables_ = ftl::config::list();
}
new Label(this, "Select Configurable","sans-bold");
auto select = new ComboBox(this, configurables_);
select->setCallback([this](int ix) {
LOG(INFO) << "Change configurable: " << ix;
_buildForm(configurables_[ix], ctrl_->get(peer_, configurables_[ix]));
setVisible(false);
//this->parent()->removeChild(this);
//delete this;
//screen()->removeChild(this);
});
for (auto c : configurables_) {
auto itembutton = new Button(this, c);
itembutton->setCallback([this,c]() {
LOG(INFO) << "Change configurable: " << c;
_buildForm(c);
setVisible(false);
//this->parent()->removeChild(this);
//delete this;
//screen()->removeChild(this);
});
}
}
ConfigWindow::~ConfigWindow() {
}
void ConfigWindow::_addElements(nanogui::FormHelper *form, const std::string &suri, const ftl::config::json_t &data) {
class ConfigWindow::References {
public:
References(ftl::NetConfigurable* nc, ftl::config::json_t* config, const std::string* suri) : nc(nc), config(config), suri(suri) {
}
~References() {
delete nc;
delete config;
delete suri;
}
private:
ftl::NetConfigurable* nc;
ftl::config::json_t* config;
const std::string* suri;
};
std::vector<ftl::gui::ConfigWindow::References *> ConfigWindow::_addElements(nanogui::FormHelper *form, ftl::Configurable &nc, const std::string &suri, std::function<ftl::Configurable*(const std::string*, std::vector<References *>&)> construct) {
using namespace nanogui;
std::vector<References *> references;
auto data = nc.getConfig();
for (auto i=data.begin(); i!=data.end(); ++i) {
if (i.key() == "$id") continue;
if (i.key() == "$ref" && i.value().is_string()) {
LOG(INFO) << "Follow $ref: " << i.value();
_addElements(form, suri, ctrl_->get(peer_, i.value().get<string>()));
const std::string* suri = new std::string(i.value().get<string>());
ftl::Configurable* rc = construct(suri, references);
auto new_references = _addElements(form, *rc, *suri, construct);
references.insert(references.end(), new_references.begin(), new_references.end());
continue;
}
if (i.value().is_boolean()) {
string key = i.key();
form->addVariable<bool>(i.key(), [this,data,key,suri](const bool &b){
ctrl_->set(peer_, suri + string("/") + key, json_t(b));
form->addVariable<bool>(i.key(), [this,data,key,&nc](const bool &b){
nc.set(key, b);
}, [data,key]() -> bool {
return data[key].get<bool>();
});
} else if (i.value().is_number_integer()) {
string key = i.key();
form->addVariable<int>(i.key(), [this,data,key,suri](const int &f){
ctrl_->set(peer_, suri + string("/") + key, json_t(f));
form->addVariable<int>(i.key(), [this,data,key,&nc](const int &f){
nc.set(key, f);
}, [data,key]() -> int {
return data[key].get<int>();
});
} else if (i.value().is_number_float()) {
string key = i.key();
form->addVariable<float>(i.key(), [this,data,key,suri](const float &f){
ctrl_->set(peer_, suri + string("/") + key, json_t(f));
form->addVariable<float>(i.key(), [this,data,key,&nc](const float &f){
nc.set(key, f);
}, [data,key]() -> float {
return data[key].get<float>();
});
} else if (i.value().is_string()) {
string key = i.key();
form->addVariable<string>(i.key(), [this,data,key,suri](const string &f){
ctrl_->set(peer_, suri + string("/") + key, json_t(f));
form->addVariable<string>(i.key(), [this,data,key,&nc](const string &f){
nc.set(key, f);
}, [data,key]() -> string {
return data[key].get<string>();
});
} else if (i.value().is_object()) {
string key = i.key();
const ftl::config::json_t &v = i.value();
form->addButton(i.key(), [this,form,suri,key,v]() {
_buildForm(suri+string("/")+key, v);
})->setIcon(ENTYPO_ICON_FOLDER);
// Checking the URI with exists() prevents unloaded local configurations from being shown.
if (suri.find('#') != string::npos && exists(suri+string("/")+key)) {
form->addButton(key, [this,suri,key]() {
_buildForm(suri+string("/")+key);
})->setIcon(ENTYPO_ICON_FOLDER);
} else if (exists(suri+string("#")+key)) {
form->addButton(key, [this,suri,key]() {
_buildForm(suri+string("#")+key);
})->setIcon(ENTYPO_ICON_FOLDER);
}
}
}
return references;
}
void ConfigWindow::_buildForm(const std::string &suri, ftl::config::json_t data) {
void ConfigWindow::_buildForm(const std::string &suri) {
using namespace nanogui;
ftl::URI uri(suri);
......@@ -105,11 +145,50 @@ void ConfigWindow::_buildForm(const std::string &suri, ftl::config::json_t data)
form->addWindow(Vector2i(100,50), uri.getFragment());
form->window()->setTheme(theme());
_addElements(form, suri, data);
ftl::config::json_t* config;
config = new ftl::config::json_t;
const std::string* allocated_suri = new std::string(suri);
std::vector<ftl::gui::ConfigWindow::References *> references;
ftl::Configurable* nc;
if (peer_) {
*config = ctrl_->get(peer_.value(), suri);
nc = new ftl::NetConfigurable(peer_.value(), *allocated_suri, *ctrl_, *config);
references = _addElements(form, *nc, *allocated_suri, [this](auto suri, auto &references) {
ftl::config::json_t* config = new ftl::config::json_t;
*config = ctrl_->get(peer_.value(), *suri);
auto nc = new ftl::NetConfigurable(peer_.value(), *suri, *ctrl_, *config);
auto r = new References(nc, config, suri);
references.push_back(r);
return nc;
});
} else {
nc = ftl::config::find(suri);
if (nc) {
references = _addElements(form, *nc, *allocated_suri, [this](auto suri, auto &references) {
return ftl::config::find(*suri);
});
}
}
auto closebutton = form->addButton("Close", [form]() {
auto closebutton = form->addButton("Close", [this,form,config,allocated_suri,nc,references]() {
form->window()->setVisible(false);
for(auto r : references) {
delete r;
}
if (peer_) {
delete nc;
}
delete config;
delete allocated_suri;
delete form;
});
closebutton->setIcon(ENTYPO_ICON_CROSS);
}
bool ConfigWindow::exists(const std::string &uri) {
// If the Configurable is a NetConfigurable, the URI is not checked.
return peer_ || ftl::config::find(uri);
}
\ No newline at end of file
applications/gui/src/config_window.hpp
View file @ c9629aed
......@@ -6,6 +6,7 @@
#include <ftl/uuid.hpp>
#include <nanogui/formhelper.h>
#include <ftl/net_configurable.hpp>
namespace ftl {
namespace gui {
......@@ -16,15 +17,23 @@ namespace gui {
class ConfigWindow : public nanogui::Window {
public:
ConfigWindow(nanogui::Widget *parent, ftl::ctrl::Master *ctrl, const ftl::UUID &peer);
ConfigWindow(nanogui::Widget *parent, ftl::ctrl::Master *ctrl, const std::optional<ftl::UUID> &peer = std::nullopt);
~ConfigWindow();
private:
/*
References holds the pointers to a NetConfigurable and all its members so that
they can all be returned from _addElements() and then simultaneously deleted
as the form is closed.
*/
class References;
ftl::ctrl::Master *ctrl_;
ftl::UUID peer_;
std::optional<ftl::UUID> peer_;
std::vector<std::string> configurables_;
void _buildForm(const std::string &uri, ftl::config::json_t data);
void _addElements(nanogui::FormHelper *form, const std::string &suri, const ftl::config::json_t &data);
void _buildForm(const std::string &uri);
std::vector<References *> _addElements(nanogui::FormHelper *form, ftl::Configurable &nc, const std::string &suri, std::function<ftl::Configurable*(const std::string*, std::vector<References *>&)> construct);
bool exists(const std::string &uri);
};
}
......
applications/gui/src/main.cpp
View file @ c9629aed
......@@ -13,7 +13,7 @@ int main(int argc, char **argv) {
ftl::net::Universe *net = ftl::create<ftl::net::Universe>(root, "net");
net->start();
//net->waitConnections();
net->waitConnections();
ftl::ctrl::Master *controller = new ftl::ctrl::Master(root, net);
controller->onLog([](const ftl::ctrl::LogEvent &e){
......
applications/gui/src/screen.cpp
View file @ c9629aed
......@@ -17,6 +17,7 @@
#include "ctrl_window.hpp"
#include "src_window.hpp"
#include "config_window.hpp"
#include "camera.hpp"
#include "media_panel.hpp"
......@@ -212,12 +213,41 @@ ftl::gui::Screen::Screen(ftl::Configurable *proot, ftl::net::Universe *pnet, ftl
popup->setVisible(false);
});
button = new ToolButton(toolbar, ENTYPO_ICON_COG);
button->setIconExtraScale(1.5f);
button->setTheme(toolbuttheme);
button->setTooltip("Settings");
button->setFixedSize(Vector2i(40,40));
button->setPosition(Vector2i(5,height()-50));
popbutton = new PopupButton(innertool, "", ENTYPO_ICON_COG);
popbutton->setIconExtraScale(1.5f);
popbutton->setTheme(toolbuttheme);
popbutton->setTooltip("Settings");
popbutton->setFixedSize(Vector2i(40,40));
popbutton->setSide(Popup::Side::Right);
popbutton->setChevronIcon(0);
// popbutton->setPosition(Vector2i(5,height()-50));
popup = popbutton->popup();
popup->setLayout(new GroupLayout());
popup->setTheme(toolbuttheme);
//net_->onConnect([this,popup](ftl::net::Peer *p) {
{
LOG(INFO) << "NET CONNECT";
auto node_details = ctrl_->getSlaves();
std::vector<std::string> node_titles;
for (auto &d : node_details) {
LOG(INFO) << "ADDING TITLE: " << d.dump();
auto peer = ftl::UUID(d["id"].get<std::string>());
auto itembutton = new Button(popup, d["title"].get<std::string>());
itembutton->setCallback([this,popup,peer]() {
auto config_window = new ConfigWindow(this, ctrl_, peer);
config_window->setTheme(windowtheme);
});
}
}
//});
itembutton = new Button(popup, "Local");
itembutton->setCallback([this,popup]() {
auto config_window = new ConfigWindow(this, ctrl_);
config_window->setTheme(windowtheme);
});
//configwindow_ = new ConfigWindow(parent, ctrl_);
cwindow_ = new ftl::gui::ControlWindow(this, controller);
......@@ -268,12 +298,9 @@ bool ftl::gui::Screen::initVR() {
{
HMD_ = nullptr;
LOG(ERROR) << "Unable to init VR runtime: " << vr::VR_GetVRInitErrorAsEnglishDescription(eError);
return false;
}
uint32_t size_x, size_y;
HMD_->GetRecommendedRenderTargetSize(&size_x, &size_y);
LOG(INFO) << size_x << ", " << size_y;
LOG(INFO) << "\n" << getCameraMatrix(HMD_, vr::Eye_Left);
return true;
}
......@@ -349,18 +376,18 @@ bool ftl::gui::Screen::mouseButtonEvent(const nanogui::Vector2i &p, int button,
float sx = ((float)p[0] - positionAfterOffset[0]) / mScale;
float sy = ((float)p[1] - positionAfterOffset[1]) / mScale;
Eigen::Vector4f camPos;
//Eigen::Vector4f camPos;
try {
camPos = camera_->source()->point(sx,sy).cast<float>();
} catch(...) {
return true;
}
//try {
//camPos = camera_->source()->point(sx,sy).cast<float>();
//} catch(...) {
// return true;
//}
camPos *= -1.0f;
//camPos *= -1.0f;
//Eigen::Vector4f worldPos = camera_->source()->getPose().cast<float>() * camPos;
//lookPoint_ = Eigen::Vector3f(worldPos[0],worldPos[1],worldPos[2]);
LOG(INFO) << "Depth at click = " << -camPos[2];
//LOG(INFO) << "Depth at click = " << -camPos[2];
return true;
}
return false;
......
applications/gui/src/screen.hpp
View file @ c9629aed
......@@ -62,9 +62,6 @@ class Screen : public nanogui::Screen {
bool hasVR() const { return false; }
#endif
void setDualView(bool v) { show_two_images_ = v; LOG(INFO) << "CLICK"; }
bool getDualView() { return show_two_images_; }
nanogui::Theme *windowtheme;
nanogui::Theme *specialtheme;
nanogui::Theme *mediatheme;
......
applications/player/src/main.cpp
View file @ c9629aed
......@@ -77,11 +77,13 @@ int main(int argc, char **argv) {
//LOG(INFO) << "Reading packet: (" << (int)spkt.streamID << "," << (int)spkt.channel << ") " << (int)pkt.codec << ", " << (int)pkt.definition;
cv::Mat frame(cv::Size(ftl::codecs::getWidth(pkt.definition),ftl::codecs::getHeight(pkt.definition)), (spkt.channel == Channel::Depth) ? CV_32F : CV_8UC3);
cv::cuda::GpuMat gframe(cv::Size(ftl::codecs::getWidth(pkt.definition),ftl::codecs::getHeight(pkt.definition)), (spkt.channel == Channel::Depth) ? CV_32F : CV_8UC3);
cv::Mat frame;
createDecoder(pkt);
try {
decoder->decode(pkt, frame);
decoder->decode(pkt, gframe);
gframe.download(frame);
} catch (std::exception &e) {
LOG(INFO) << "Decoder exception: " << e.what();
}
......
applications/reconstruct/CMakeLists.txt
View file @ c9629aed
......@@ -20,6 +20,7 @@ set(REPSRC
src/ilw/fill.cu
src/ilw/discontinuity.cu
src/ilw/correspondence.cu
src/filters/smoothing.cu
)
add_executable(ftl-reconstruct ${REPSRC})
......
applications/reconstruct/src/filters/smoothing.cu 0 → 100644
View file @ c9629aed
#include "smoothing.hpp"
#include <ftl/cuda/weighting.hpp>
#define T_PER_BLOCK 8
template <int RADIUS>
__global__ void depth_smooth_kernel(
ftl::cuda::TextureObject<float> depth_in,
ftl::cuda::TextureObject<uchar4> colour_in,
ftl::cuda::TextureObject<float> depth_out,
ftl::rgbd::Camera camera,
float factor, float thresh) {
const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
if (x < depth_in.width() && y < depth_in.height()) {
float d = depth_in.tex2D((int)x,(int)y);
depth_out(x,y) = 0.0f;
if (d < camera.minDepth || d > camera.maxDepth) return;
uchar4 c = colour_in.tex2D((int)x, (int)y);
float3 pos = camera.screenToCam(x,y,d);
float contrib = 0.0f;
float new_depth = 0.0f;
for (int v=-RADIUS; v<=RADIUS; ++v) {
for (int u=-RADIUS; u<=RADIUS; ++u) {
// Get colour difference to center
const uchar4 cN = colour_in.tex2D((int)x+u, (int)y+v);
const float colourWeight = ftl::cuda::colourWeighting(c, cN, thresh);
const float dN = depth_in.tex2D((int)x + u, (int)y + v);
const float3 posN = camera.screenToCam(x+u, y+v, dN);
const float weight = ftl::cuda::spatialWeighting(posN, pos, factor * colourWeight);
contrib += weight;
new_depth += dN * weight;
}
}
if (contrib > 0.0f) {
depth_out(x,y) = new_depth / contrib;
}
}
}
void ftl::cuda::depth_smooth(
ftl::cuda::TextureObject<float> &depth_in,
ftl::cuda::TextureObject<uchar4> &colour_in,
ftl::cuda::TextureObject<float> &depth_out,
const ftl::rgbd::Camera &camera,
int radius, float factor, float thresh, int iters, 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 blockSize(T_PER_BLOCK, T_PER_BLOCK);
for (int n=0; n<iters; ++n) {
switch (radius) {
case 5 : depth_smooth_kernel<5><<<gridSize, blockSize, 0, stream>>>(depth_in, colour_in, depth_out, camera, factor, thresh); break;
case 4 : depth_smooth_kernel<4><<<gridSize, blockSize, 0, stream>>>(depth_in, colour_in, depth_out, camera, factor, thresh); break;
case 3 : depth_smooth_kernel<3><<<gridSize, blockSize, 0, stream>>>(depth_in, colour_in, depth_out, camera, factor, thresh); break;
case 2 : depth_smooth_kernel<2><<<gridSize, blockSize, 0, stream>>>(depth_in, colour_in, depth_out, camera, factor, thresh); break;
case 1 : depth_smooth_kernel<1><<<gridSize, blockSize, 0, stream>>>(depth_in, colour_in, depth_out, camera, factor, thresh); break;
default: break;
}
cudaSafeCall( cudaGetLastError() );
switch (radius) {
case 5 : depth_smooth_kernel<5><<<gridSize, blockSize, 0, stream>>>(depth_out, colour_in, depth_in, camera, factor, thresh); break;
case 4 : depth_smooth_kernel<4><<<gridSize, blockSize, 0, stream>>>(depth_out, colour_in, depth_in, camera, factor, thresh); break;
case 3 : depth_smooth_kernel<3><<<gridSize, blockSize, 0, stream>>>(depth_out, colour_in, depth_in, camera, factor, thresh); break;
case 2 : depth_smooth_kernel<2><<<gridSize, blockSize, 0, stream>>>(depth_out, colour_in, depth_in, camera, factor, thresh); break;
case 1 : depth_smooth_kernel<1><<<gridSize, blockSize, 0, stream>>>(depth_out, colour_in, depth_in, camera, factor, thresh); break;
default: break;
}
cudaSafeCall( cudaGetLastError() );
}
#ifdef _DEBUG
cudaSafeCall(cudaDeviceSynchronize());
#endif
}
applications/reconstruct/src/filters/smoothing.hpp 0 → 100644
View file @ c9629aed
#ifndef _FTL_CUDA_SMOOTHING_HPP_
#define _FTL_CUDA_SMOOTHING_HPP_
#include <ftl/rgbd/camera.hpp>
#include <ftl/cuda_common.hpp>
namespace ftl {
namespace cuda {
void depth_smooth(
ftl::cuda::TextureObject<float> &depth_in,
ftl::cuda::TextureObject<uchar4> &colour_in,
ftl::cuda::TextureObject<float> &depth_out,
const ftl::rgbd::Camera &camera,
int radius, float factor, float thresh, int iters,
cudaStream_t stream);
}
}
#endif // _FTL_CUDA_SMOOTHING_HPP_
applications/reconstruct/src/ilw.cpp deleted 100644 → 0
View file @ 345449ce
#include "ilw.hpp"
#include <ftl/utility/matrix_conversion.hpp>
#include <ftl/rgbd/source.hpp>
#include <ftl/cuda/points.hpp>
#include <loguru.hpp>
#include "ilw_cuda.hpp"
using ftl::ILW;
using ftl::detail::ILWData;
using ftl::codecs::Channel;
using ftl::codecs::Channels;
using ftl::rgbd::Format;
using cv::cuda::GpuMat;
ILW::ILW(nlohmann::json &config) : ftl::Configurable(config) {
}
ILW::~ILW() {
}
bool ILW::process(ftl::rgbd::FrameSet &fs, cudaStream_t stream) {
_phase0(fs, stream);
//for (int i=0; i<2; ++i) {
_phase1(fs, stream);
//for (int j=0; j<3; ++j) {
// _phase2(fs);
//}
// TODO: Break if no time left
//}
return true;
}
bool ILW::_phase0(ftl::rgbd::FrameSet &fs, cudaStream_t stream) {
// Make points channel...
for (size_t i=0; i<fs.frames.size(); ++i) {
auto &f = fs.frames[i];
auto *s = fs.sources[i];
if (f.empty(Channel::Depth + Channel::Colour)) {
LOG(ERROR) << "Missing required channel";
continue;
}
auto &t = f.createTexture<float4>(Channel::Points, Format<float4>(f.get<GpuMat>(Channel::Colour).size()));
auto pose = MatrixConversion::toCUDA(s->getPose().cast<float>()); //.inverse());
ftl::cuda::point_cloud(t, f.createTexture<float>(Channel::Depth), s->parameters(), pose, stream);
// TODO: Create energy vector texture and clear it
// Create energy and clear it
// Convert colour from BGR to BGRA if needed
if (f.get<GpuMat>(Channel::Colour).type() == CV_8UC3) {
// Convert to 4 channel colour
auto &col = f.get<GpuMat>(Channel::Colour);
GpuMat tmp(col.size(), CV_8UC4);
cv::cuda::swap(col, tmp);
cv::cuda::cvtColor(tmp,col, cv::COLOR_BGR2BGRA);
}
f.createTexture<float4>(Channel::EnergyVector, Format<float4>(f.get<GpuMat>(Channel::Colour).size()));
f.createTexture<float>(Channel::Energy, Format<float>(f.get<GpuMat>(Channel::Colour).size()));
f.createTexture<uchar4>(Channel::Colour);
cv::cuda::Stream cvstream = cv::cuda::StreamAccessor::wrapStream(stream);
f.get<GpuMat>(Channel::EnergyVector).setTo(cv::Scalar(0.0f,0.0f,0.0f,0.0f), cvstream);
f.get<GpuMat>(Channel::Energy).setTo(cv::Scalar(0.0f), cvstream);
}
return true;
}
bool ILW::_phase1(ftl::rgbd::FrameSet &fs, cudaStream_t stream) {
// Run correspondence kernel to create an energy vector
// For each camera combination
for (size_t i=0; i<fs.frames.size(); ++i) {
for (size_t j=0; j<fs.frames.size(); ++j) {
if (i == j) continue;
LOG(INFO) << "Running phase1";
auto &f1 = fs.frames[i];
auto &f2 = fs.frames[j];
//auto s1 = fs.frames[i];
auto s2 = fs.sources[j];
auto pose = MatrixConversion::toCUDA(s2->getPose().cast<float>().inverse());
try {
//Calculate energy vector to best correspondence
ftl::cuda::correspondence_energy_vector(
f1.getTexture<float4>(Channel::Points),
f2.getTexture<float4>(Channel::Points),
f1.getTexture<uchar4>(Channel::Colour),
f2.getTexture<uchar4>(Channel::Colour),
// TODO: Add normals and other things...
f1.getTexture<float4>(Channel::EnergyVector),
f1.getTexture<float>(Channel::Energy),
pose,
s2->parameters(),
stream
);
} catch (ftl::exception &e) {
LOG(ERROR) << "Exception in correspondence: " << e.what();
}
LOG(INFO) << "Correspondences done... " << i;
}
}
return true;
}
bool ILW::_phase2(ftl::rgbd::FrameSet &fs) {
// Run energies and motion kernel
return true;
}
applications/reconstruct/src/ilw.cu deleted 100644 → 0
View file @ 345449ce
#include "ilw_cuda.hpp"
using ftl::cuda::TextureObject;
using ftl::rgbd::Camera;
#define WARP_SIZE 32
#define T_PER_BLOCK 8
#define FULL_MASK 0xffffffff
__device__ inline float warpMax(float e) {
for (int i = WARP_SIZE/2; i > 0; i /= 2) {
const float other = __shfl_xor_sync(FULL_MASK, e, i, WARP_SIZE);
e = max(e, other);
}
return e;
}
__global__ void correspondence_energy_vector_kernel(
TextureObject<float4> p1,
TextureObject<float4> p2,
TextureObject<uchar4> c1,
TextureObject<uchar4> c2,
TextureObject<float4> vout,
TextureObject<float> eout,
float4x4 pose2, // Inverse
Camera cam2) {
// Each warp picks point in p1
const int tid = (threadIdx.x + threadIdx.y * blockDim.x);
const int x = (blockIdx.x*blockDim.x + threadIdx.x) / WARP_SIZE;
const int y = blockIdx.y*blockDim.y + threadIdx.y;
const float3 world1 = make_float3(p1.tex2D(x, y));
if (world1.x == MINF) {
vout(x,y) = make_float4(0.0f);
eout(x,y) = 0.0f;
return;
}
const float3 camPos2 = pose2 * world1;
const uint2 screen2 = cam2.camToScreen<uint2>(camPos2);
const int upsample = 8;
// Project to p2 using cam2
// Each thread takes a possible correspondence and calculates a weighting
const int lane = tid % WARP_SIZE;
for (int i=lane; i<upsample*upsample; i+=WARP_SIZE) {
const float u = (i % upsample) - (upsample / 2);
const float v = (i / upsample) - (upsample / 2);
const float3 world2 = make_float3(p2.tex2D(screen2.x+u, screen2.y+v));
if (world2.x == MINF) continue;
// Determine degree of correspondence
const float confidence = 1.0f / length(world1 - world2);
const float maxconf = warpMax(confidence);
// This thread has best confidence value
if (maxconf == confidence) {
vout(x,y) = vout.tex2D(x, y) + make_float4(
(world1.x - world2.x) * maxconf,
(world1.y - world2.y) * maxconf,
(world1.z - world2.z) * maxconf,
maxconf);
eout(x,y) = eout.tex2D(x,y) + length(world1 - world2)*maxconf;
}
}
}
void ftl::cuda::correspondence_energy_vector(
TextureObject<float4> &p1,
TextureObject<float4> &p2,
TextureObject<uchar4> &c1,
TextureObject<uchar4> &c2,
TextureObject<float4> &vout,
TextureObject<float> &eout,
float4x4 &pose2,
const Camera &cam2,
cudaStream_t stream) {
const dim3 gridSize((p1.width() + 2 - 1)/2, (p1.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
const dim3 blockSize(2*WARP_SIZE, T_PER_BLOCK);
printf("COR SIZE %d,%d\n", p1.width(), p1.height());
correspondence_energy_vector_kernel<<<gridSize, blockSize, 0, stream>>>(
p1, p2, c1, c2, vout, eout, pose2, cam2
);
cudaSafeCall( cudaGetLastError() );
}
applications/reconstruct/src/ilw/ilw.cpp
View file @ c9629aed
......@@ -144,7 +144,7 @@ ILW::~ILW() {
bool ILW::process(ftl::rgbd::FrameSet &fs) {
if (!enabled_) return false;
fs.upload(Channel::Colour + Channel::Depth, stream_);
//fs.upload(Channel::Colour + Channel::Depth, stream_);
_phase0(fs, stream_);
params_.range = value("search_range", 0.05f);
......
applications/reconstruct/src/main.cpp
View file @ c9629aed
......@@ -16,9 +16,10 @@
#include <ftl/rgbd/group.hpp>
#include <ftl/threads.hpp>
#include <ftl/codecs/writer.hpp>
#include <ftl/codecs/reader.hpp>
#include "ilw/ilw.hpp"
#include <ftl/render/splat_render.hpp>
#include <ftl/render/tri_render.hpp>
#include <fstream>
#include <string>
......@@ -29,6 +30,7 @@
#include <opencv2/opencv.hpp>
#include <ftl/net/universe.hpp>
#include "filters/smoothing.hpp"
#include <ftl/registration.hpp>
#include <cuda_profiler_api.h>
......@@ -65,34 +67,6 @@ static Eigen::Affine3d create_rotation_matrix(float ax, float ay, float az) {
return rz * rx * ry;
}
static void writeSourceProperties(ftl::codecs::Writer &writer, int id, ftl::rgbd::Source *src) {
ftl::codecs::StreamPacket spkt;
ftl::codecs::Packet pkt;
spkt.timestamp = 0;
spkt.streamID = id;
spkt.channel = Channel::Calibration;
spkt.channel_count = 1;
pkt.codec = ftl::codecs::codec_t::CALIBRATION;
pkt.definition = ftl::codecs::definition_t::Any;
pkt.block_number = 0;
pkt.block_total = 1;
pkt.flags = 0;
pkt.data = std::move(std::vector<uint8_t>((uint8_t*)&src->parameters(), (uint8_t*)&src->parameters() + sizeof(ftl::rgbd::Camera)));
writer.write(spkt, pkt);
spkt.channel = Channel::Pose;
pkt.codec = ftl::codecs::codec_t::POSE;
pkt.definition = ftl::codecs::definition_t::Any;
pkt.block_number = 0;
pkt.block_total = 1;
pkt.flags = 0;
pkt.data = std::move(std::vector<uint8_t>((uint8_t*)src->getPose().data(), (uint8_t*)src->getPose().data() + 4*4*sizeof(double)));
writer.write(spkt, pkt);
}
static void run(ftl::Configurable *root) {
Universe *net = ftl::create<Universe>(root, "net");
ftl::ctrl::Slave slave(net, root);
......@@ -102,6 +76,38 @@ static void run(ftl::Configurable *root) {
net->start();
net->waitConnections();
// Check paths for an FTL file to load...
auto paths = (*root->get<nlohmann::json>("paths"));
for (auto &x : paths.items()) {
std::string path = x.value().get<std::string>();
auto eix = path.find_last_of('.');
auto ext = path.substr(eix+1);
// Command line path is ftl file
if (ext == "ftl") {
// Create temp reader to count number of sources found in file
std::ifstream file;
file.open(path);
ftl::codecs::Reader reader(file);
reader.begin();
int max_stream = 0;
reader.read(reader.getStartTime()+100, [&max_stream](const ftl::codecs::StreamPacket &spkt, const ftl::codecs::Packet &pkt) {
max_stream = max(max_stream, spkt.streamID);
});
reader.end();
LOG(INFO) << "Found " << (max_stream+1) << " sources in " << path;
int N = root->value("N", 100);
// For each stream found, add a source object
for (int i=0; i<=min(max_stream,N-1); ++i) {
root->getConfig()["sources"].push_back(nlohmann::json{{"uri",std::string("file://") + path + std::string("#") + std::to_string(i)}});
}
}
}
// Create a vector of all input RGB-Depth sources
auto sources = ftl::createArray<Source>(root, "sources", net);
......@@ -165,7 +171,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");
ftl::render::Splatter *splat = ftl::create<ftl::render::Splatter>(root, "renderer", &scene_B);
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");
......@@ -174,7 +180,7 @@ static void run(ftl::Configurable *root) {
// Generate virtual camera render when requested by streamer
virt->onRender([splat,virt,&scene_B,align](ftl::rgbd::Frame &out) {
virt->setTimestamp(scene_B.timestamp);
//virt->setTimestamp(scene_B.timestamp);
// Do we need to convert Lab to BGR?
if (align->isLabColour()) {
for (auto &f : scene_B.frames) {
......@@ -215,14 +221,15 @@ static void run(ftl::Configurable *root) {
fileout.open(std::string(timestamp) + ".ftl");
writer.begin();
group->addRawCallback(std::function(recorder));
// TODO: Write pose+calibration+config packets
auto sources = group->sources();
for (int i=0; i<sources.size(); ++i) {
writeSourceProperties(writer, i, sources[i]);
//writeSourceProperties(writer, i, sources[i]);
sources[i]->inject(Channel::Calibration, sources[i]->parameters(), Channel::Left, sources[i]->getCapabilities());
sources[i]->inject(sources[i]->getPose());
}
group->addRawCallback(std::function(recorder));
} else {
group->removeRawCallback(recorder);
writer.end();
......@@ -238,9 +245,11 @@ static void run(ftl::Configurable *root) {
bool busy = false;
auto *filter = ftl::config::create<ftl::Configurable>(root, "filters");
group->setLatency(4);
group->setName("ReconGroup");
group->sync([splat,virt,&busy,&slave,&scene_A,&scene_B,&align,controls](ftl::rgbd::FrameSet &fs) -> bool {
group->sync([splat,virt,&busy,&slave,&scene_A,&scene_B,&align,controls,filter](ftl::rgbd::FrameSet &fs) -> bool {
//cudaSetDevice(scene->getCUDADevice());
//if (slave.isPaused()) return true;
......@@ -255,13 +264,46 @@ static void run(ftl::Configurable *root) {
// Swap the entire frameset to allow rapid return
fs.swapTo(scene_A);
ftl::pool.push([&scene_B,&scene_A,&busy,&slave,&align](int id) {
ftl::pool.push([&scene_B,&scene_A,&busy,&slave,&align, filter](int id) {
//cudaSetDevice(scene->getCUDADevice());
// TODO: Release frameset here...
//cudaSafeCall(cudaStreamSynchronize(scene->getIntegrationStream()));
UNIQUE_LOCK(scene_A.mtx, lk);
cv::cuda::GpuMat tmp;
float factor = filter->value("smooth_factor", 0.4f);
float colour_limit = filter->value("colour_limit", 30.0f);
bool do_smooth = filter->value("pre_smooth", false);
int iters = filter->value("iterations", 3);
int radius = filter->value("radius", 5);
if (do_smooth) {
// Presmooth...
for (int i=0; i<scene_A.frames.size(); ++i) {
auto &f = scene_A.frames[i];
auto s = scene_A.sources[i];
// Convert colour from BGR to BGRA if needed
if (f.get<cv::cuda::GpuMat>(Channel::Colour).type() == CV_8UC3) {
//cv::cuda::Stream cvstream = cv::cuda::StreamAccessor::wrapStream(stream);
// Convert to 4 channel colour
auto &col = f.get<cv::cuda::GpuMat>(Channel::Colour);
tmp.create(col.size(), CV_8UC4);
cv::cuda::swap(col, tmp);
cv::cuda::cvtColor(tmp,col, cv::COLOR_BGR2BGRA, 0);
}
ftl::cuda::depth_smooth(
f.createTexture<float>(Channel::Depth),
f.createTexture<uchar4>(Channel::Colour),
f.createTexture<float>(Channel::Depth2, ftl::rgbd::Format<float>(f.get<cv::cuda::GpuMat>(Channel::Depth).size())),
s->parameters(),
radius, factor, colour_limit, iters, 0
);
}
}
// Send all frames to GPU, block until done?
//scene_A.upload(Channel::Colour + Channel::Depth); // TODO: (Nick) Add scene stream.
align->process(scene_A);
......
applications/recorder/src/main.cpp
View file @ c9629aed
......@@ -62,34 +62,6 @@ static Eigen::Affine3d create_rotation_matrix(float ax, float ay, float az) {
return rz * rx * ry;
}
static void writeSourceProperties(ftl::codecs::Writer &writer, int id, ftl::rgbd::Source *src) {
ftl::codecs::StreamPacket spkt;
ftl::codecs::Packet pkt;
spkt.timestamp = 0;
spkt.streamID = id;
spkt.channel = Channel::Calibration;
spkt.channel_count = 1;
pkt.codec = ftl::codecs::codec_t::CALIBRATION;
pkt.definition = ftl::codecs::definition_t::Any;
pkt.block_number = 0;
pkt.block_total = 1;
pkt.flags = 0;
pkt.data = std::move(std::vector<uint8_t>((uint8_t*)&src->parameters(), (uint8_t*)&src->parameters() + sizeof(ftl::rgbd::Camera)));
writer.write(spkt, pkt);
spkt.channel = Channel::Pose;
pkt.codec = ftl::codecs::codec_t::POSE;
pkt.definition = ftl::codecs::definition_t::Any;
pkt.block_number = 0;
pkt.block_total = 1;
pkt.flags = 0;
pkt.data = std::move(std::vector<uint8_t>((uint8_t*)src->getPose().data(), (uint8_t*)src->getPose().data() + 4*4*sizeof(double)));
writer.write(spkt, pkt);
}
static void run(ftl::Configurable *root) {
Universe *net = ftl::create<Universe>(root, "net");
ftl::ctrl::Slave slave(net, root);
......@@ -199,7 +171,9 @@ static void run(ftl::Configurable *root) {
// TODO: Write pose+calibration+config packets
auto sources = group->sources();
for (int i=0; i<sources.size(); ++i) {
writeSourceProperties(writer, i, sources[i]);
//writeSourceProperties(writer, i, sources[i]);
sources[i]->inject(Channel::Calibration, sources[i]->parameters(), Channel::Left, sources[i]->getCapabilities());
sources[i]->inject(sources[i]->getPose());
}
} else {
//group->removeRawCallback(recorder);
......
components/codecs/include/ftl/codecs/bitrates.hpp
View file @ c9629aed
......@@ -22,7 +22,11 @@ enum struct codec_t : uint8_t {
JSON = 100, // A JSON string
CALIBRATION, // Camera parameters object
POSE, // 4x4 eigen matrix
RAW // Some unknown binary format (msgpack?)
MSGPACK,
STRING, // Null terminated string
RAW, // Some unknown binary format
Any = 255
};
/**
......@@ -38,9 +42,11 @@ enum struct definition_t : uint8_t {
LD360 = 6,
Any = 7,
HTC_VIVE = 8
HTC_VIVE = 8,
// TODO: Add audio definitions
Invalid
};
/**
......@@ -85,6 +91,8 @@ static const preset_t kPresetLQThreshold = 4;
static const preset_t kPresetHTCVive = -1;
static const preset_t kPresetMinimum = -1;
/**
* Represents the details of each preset codec configuration.
*/
......@@ -123,6 +131,8 @@ preset_t findPreset(definition_t, definition_t);
*/
preset_t findPreset(definition_t);
preset_t findPreset(size_t width, size_t height);
}
}
......
components/codecs/include/ftl/codecs/channels.hpp
View file @ c9629aed
......@@ -17,6 +17,7 @@ enum struct Channel : int {
Disparity = 3,
Depth2 = 3,
Deviation = 4,
Screen = 4,
Normals = 5, // 32FC4
Points = 6, // 32FC4
Confidence = 7, // 32F
......@@ -36,7 +37,8 @@ enum struct Channel : int {
Configuration = 64, // JSON Data
Calibration = 65, // Camera Parameters Object
Pose = 66, // Eigen::Matrix4d
Data = 67 // Custom data, any codec.
Index = 67,
Data = 2048 // Custom data, any codec.
};
inline bool isVideo(Channel c) { return (int)c < 32; };
......
components/codecs/include/ftl/codecs/decoder.hpp
View file @ c9629aed
......@@ -35,7 +35,7 @@ class Decoder {
Decoder() {};
virtual ~Decoder() {};
virtual bool decode(const ftl::codecs::Packet &pkt, cv::Mat &out)=0;
virtual bool decode(const ftl::codecs::Packet &pkt, cv::cuda::GpuMat &out)=0;
virtual bool accepts(const ftl::codecs::Packet &)=0;
};
......
components/codecs/include/ftl/codecs/encoder.hpp
View file @ c9629aed
......@@ -33,7 +33,8 @@ enum class device_t {
*/
Encoder *allocateEncoder(
ftl::codecs::definition_t maxdef=ftl::codecs::definition_t::HD1080,
ftl::codecs::device_t dev=ftl::codecs::device_t::Any);
ftl::codecs::device_t dev=ftl::codecs::device_t::Any,
ftl::codecs::codec_t codec=ftl::codecs::codec_t::Any);
/**
* Release an encoder to be reused by some other stream.
......@@ -47,7 +48,7 @@ void free(Encoder *&e);
class Encoder {
public:
friend Encoder *allocateEncoder(ftl::codecs::definition_t,
ftl::codecs::device_t);
ftl::codecs::device_t, ftl::codecs::codec_t);
friend void free(Encoder *&);
public:
......@@ -59,7 +60,7 @@ class Encoder {
/**
* Wrapper encode to allow use of presets.
*/
virtual bool encode(const cv::Mat &in, ftl::codecs::preset_t preset,
virtual bool encode(const cv::cuda::GpuMat &in, ftl::codecs::preset_t preset,
const std::function<void(const ftl::codecs::Packet&)> &cb);
/**
......@@ -76,7 +77,7 @@ class Encoder {
* @return True if succeeded with encoding.
*/
virtual bool encode(
const cv::Mat &in,
const cv::cuda::GpuMat &in,
ftl::codecs::definition_t definition,
ftl::codecs::bitrate_t bitrate,
const std::function<void(const ftl::codecs::Packet&)> &cb)=0;
......@@ -86,6 +87,8 @@ class Encoder {
virtual void reset() {}
virtual bool supports(ftl::codecs::codec_t codec)=0;
protected:
bool available;
const ftl::codecs::definition_t max_definition;
......
components/codecs/include/ftl/codecs/h264.hpp 0 → 100644
View file @ c9629aed
#ifndef _FTL_CODECS_H264_HPP_
#define _FTL_CODECS_H264_HPP_
namespace ftl {
namespace codecs {
/**
* H.264 codec utility functions.
*/
namespace h264 {
/**
* H264 Network Abstraction Layer Unit types.
*/
enum class NALType : int {
UNSPECIFIED_0 = 0,
CODED_SLICE_NON_IDR = 1,
CODED_SLICE_PART_A = 2,
CODED_SLICE_PART_B = 3,
CODED_SLICE_PART_C = 4,
CODED_SLICE_IDR = 5,
SEI = 6,
SPS = 7,
PPS = 8,
ACCESS_DELIMITER = 9,
EO_SEQ = 10,
EO_STREAM = 11,
FILTER_DATA = 12,
SPS_EXT = 13,
PREFIX_NAL_UNIT = 14,
SUBSET_SPS = 15,
RESERVED_16 = 16,
RESERVED_17 = 17,
RESERVED_18 = 18,
CODED_SLICE_AUX = 19,
CODED_SLICE_EXT = 20,
CODED_SLICE_DEPTH = 21,
RESERVED_22 = 22,
RESERVED_23 = 23,
UNSPECIFIED_24 = 24,
UNSPECIFIED_25,
UNSPECIFIED_26,
UNSPECIFIED_27,
UNSPECIFIED_28,
UNSPECIFIED_29,
UNSPECIFIED_30,
UNSPECIFIED_31
};
/**
* Extract the NAL unit type from the first NAL header.
* With NvPipe, the 5th byte contains the NAL Unit header.
*/
inline NALType getNALType(const std::vector<uint8_t> &data) {
return static_cast<NALType>(data[4] & 0x1F);
}
/**
* Check the H264 bitstream for an I-Frame. With NvPipe, all I-Frames start
* with a SPS NAL unit so just check for this.
*/
inline bool isIFrame(const std::vector<uint8_t> &data) {
return getNALType(data) == NALType::SPS;
}
}
}
}
#endif // _FTL_CODECS_H264_HPP_