Skip to content
Snippets Groups Projects
Commit 275a9dea authored by Sebastian Hahta's avatar Sebastian Hahta
Browse files

Merge branch 'master' into feature/python

parents b1c56157 c9629aed
No related branches found
No related tags found
1 merge request!155Feature/python
Pipeline #16072 failed
Showing
with 458 additions and 348 deletions
......@@ -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);
});
}
......
......@@ -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
......@@ -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);
};
}
......
......@@ -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){
......
......@@ -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;
......
......@@ -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;
......
......@@ -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();
}
......
......@@ -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})
......
#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
}
#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_
#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;
}
#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() );
}
......@@ -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);
......
......@@ -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);
......
......@@ -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);
......
......@@ -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);
}
}
......
......@@ -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; };
......
......@@ -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;
};
......
......@@ -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;
......
#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_
0% Loading or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment