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
Select Git revision

Target

Select target project
  • nicolaspope/ftl
1 result
Select Git revision
Show changes
Show whitespace changes
Inline Side-by-side
Showing
with 1165 additions and 901 deletions
components/renderers/cpp/src/CUDARender.cpp
View file @ 70063e8d
...@@ -25,13 +25,14 @@ using ftl::render::CUDARender; ...@@ -25,13 +25,14 @@ using ftl::render::CUDARender;
using ftl::codecs::Channel; using ftl::codecs::Channel;
using ftl::codecs::Channels; using ftl::codecs::Channels;
using ftl::rgbd::Format; using ftl::rgbd::Format;
using ftl::rgbd::VideoFrame;
using cv::cuda::GpuMat; using cv::cuda::GpuMat;
using std::stoul; using std::stoul;
using ftl::cuda::Mask; using ftl::cuda::Mask;
using ftl::render::parseCUDAColour; using ftl::render::parseCUDAColour;
using ftl::render::parseCVColour; using ftl::render::parseCVColour;
CUDARender::CUDARender(nlohmann::json &config) : ftl::render::FSRenderer(config), scene_(nullptr) { CUDARender::CUDARender(nlohmann::json &config) : ftl::render::FSRenderer(config), temp_d_(ftl::data::Frame::make_standalone()), temp_(temp_d_.cast<ftl::rgbd::Frame>()), scene_(nullptr) {
/*if (config["clipping"].is_object()) { /*if (config["clipping"].is_object()) {
auto &c = config["clipping"]; auto &c = config["clipping"];
float rx = c.value("pitch", 0.0f); float rx = c.value("pitch", 0.0f);
...@@ -59,27 +60,29 @@ CUDARender::CUDARender(nlohmann::json &config) : ftl::render::FSRenderer(config) ...@@ -59,27 +60,29 @@ CUDARender::CUDARender(nlohmann::json &config) : ftl::render::FSRenderer(config)
colouriser_ = ftl::create<ftl::render::Colouriser>(this, "colouriser"); colouriser_ = ftl::create<ftl::render::Colouriser>(this, "colouriser");
on("clipping_enabled", [this](const ftl::config::Event &e) { on("touch_sensitivity", touch_dist_, 0.04f);
on("clipping_enabled", [this]() {
clipping_ = value("clipping_enabled", true); clipping_ = value("clipping_enabled", true);
}); });
norm_filter_ = value("normal_filter", -1.0f); norm_filter_ = value("normal_filter", -1.0f);
on("normal_filter", [this](const ftl::config::Event &e) { on("normal_filter", [this]() {
norm_filter_ = value("normal_filter", -1.0f); norm_filter_ = value("normal_filter", -1.0f);
}); });
backcull_ = value("back_cull", true); backcull_ = value("back_cull", true);
on("back_cull", [this](const ftl::config::Event &e) { on("back_cull", [this]() {
backcull_ = value("back_cull", true); backcull_ = value("back_cull", true);
}); });
mesh_ = value("meshing", true); mesh_ = value("meshing", true);
on("meshing", [this](const ftl::config::Event &e) { on("meshing", [this]() {
mesh_ = value("meshing", true); mesh_ = value("meshing", true);
}); });
background_ = parseCVColour(value("background", std::string("#4c4c4c"))); background_ = parseCVColour(value("background", std::string("#4c4c4c")));
on("background", [this](const ftl::config::Event &e) { on("background", [this]() {
background_ = parseCVColour(value("background", std::string("#4c4c4c"))); background_ = parseCVColour(value("background", std::string("#4c4c4c")));
}); });
...@@ -96,12 +99,18 @@ CUDARender::CUDARender(nlohmann::json &config) : ftl::render::FSRenderer(config) ...@@ -96,12 +99,18 @@ CUDARender::CUDARender(nlohmann::json &config) : ftl::render::FSRenderer(config)
} }
} }
cudaSafeCall(cudaStreamCreate(&stream_)); //cudaSafeCall(cudaStreamCreate(&stream_));
stream_ = 0;
last_frame_ = -1; last_frame_ = -1;
temp_.store();
// Allocate collisions buffer
cudaSafeCall(cudaMalloc(&collisions_, 1024*sizeof(ftl::cuda::Collision)));
} }
CUDARender::~CUDARender() { CUDARender::~CUDARender() {
delete colouriser_;
cudaFree(collisions_);
} }
void CUDARender::_renderChannel(ftl::rgbd::Frame &output, ftl::codecs::Channel in, const Eigen::Matrix4d &t, cudaStream_t stream) { void CUDARender::_renderChannel(ftl::rgbd::Frame &output, ftl::codecs::Channel in, const Eigen::Matrix4d &t, cudaStream_t stream) {
...@@ -110,12 +119,12 @@ void CUDARender::_renderChannel(ftl::rgbd::Frame &output, ftl::codecs::Channel i ...@@ -110,12 +119,12 @@ void CUDARender::_renderChannel(ftl::rgbd::Frame &output, ftl::codecs::Channel i
if (in == Channel::None) return; if (in == Channel::None) return;
for (size_t i=0; i < scene_->frames.size(); ++i) { for (size_t i=0; i < scene_->frames.size(); ++i) {
if (!scene_->hasFrame(i)) continue; //if (!scene_->hasFrame(i)) continue;
auto &f = scene_->frames[i]; auto &f = scene_->frames[i].cast<ftl::rgbd::Frame>();
if (!f.hasChannel(in)) { if (!f.hasChannel(in)) {
LOG(ERROR) << "Reprojecting unavailable channel"; //LOG(ERROR) << "Reprojecting unavailable channel";
return; continue;
} }
_adjustDepthThresholds(f.getLeftCamera()); _adjustDepthThresholds(f.getLeftCamera());
...@@ -169,14 +178,14 @@ void CUDARender::_renderChannel(ftl::rgbd::Frame &output, ftl::codecs::Channel i ...@@ -169,14 +178,14 @@ void CUDARender::_renderChannel(ftl::rgbd::Frame &output, ftl::codecs::Channel i
void CUDARender::_dibr(ftl::rgbd::Frame &out, const Eigen::Matrix4d &t, cudaStream_t stream) { void CUDARender::_dibr(ftl::rgbd::Frame &out, const Eigen::Matrix4d &t, cudaStream_t stream) {
cv::cuda::Stream cvstream = cv::cuda::StreamAccessor::wrapStream(stream); cv::cuda::Stream cvstream = cv::cuda::StreamAccessor::wrapStream(stream);
temp_.get<GpuMat>(Channel::Depth2).setTo(cv::Scalar(0x7FFFFFFF), cvstream); temp_.set<GpuMat>(Channel::Depth2).setTo(cv::Scalar(0x7FFFFFFF), cvstream);
for (size_t i=0; i < scene_->frames.size(); ++i) { for (size_t i=0; i < scene_->frames.size(); ++i) {
if (!scene_->hasFrame(i)) continue; if (!scene_->hasFrame(i)) continue;
auto &f = scene_->frames[i]; auto &f = scene_->frames[i].cast<ftl::rgbd::Frame>();
//auto *s = scene_->sources[i]; //auto *s = scene_->sources[i];
if (f.empty(Channel::Colour)) { if (!f.has(Channel::Colour)) {
LOG(ERROR) << "Missing required channel"; LOG(ERROR) << "Missing required channel";
continue; continue;
} }
...@@ -233,23 +242,35 @@ void CUDARender::_mesh(ftl::rgbd::Frame &out, const Eigen::Matrix4d &t, cudaStre ...@@ -233,23 +242,35 @@ void CUDARender::_mesh(ftl::rgbd::Frame &out, const Eigen::Matrix4d &t, cudaStre
bool do_blend = value("mesh_blend", false); bool do_blend = value("mesh_blend", false);
float blend_alpha = value("blend_alpha", 0.02f); float blend_alpha = value("blend_alpha", 0.02f);
if (do_blend) { if (do_blend) {
temp_.get<GpuMat>(Channel::Depth).setTo(cv::Scalar(0x7FFFFFFF), cvstream); temp_.set<GpuMat>(Channel::Depth).setTo(cv::Scalar(0x7FFFFFFF), cvstream);
temp_.get<GpuMat>(Channel::Weights).setTo(cv::Scalar(0.0f), cvstream); temp_.set<GpuMat>(Channel::Weights).setTo(cv::Scalar(0.0f), cvstream);
} else { } else {
temp_.get<GpuMat>(Channel::Depth2).setTo(cv::Scalar(0x7FFFFFFF), cvstream); temp_.set<GpuMat>(Channel::Depth2).setTo(cv::Scalar(0x7FFFFFFF), cvstream);
} }
int valid_count = 0;
// FIXME: Is it possible to remember previously if there should be depth?
bool use_depth = scene_->anyHasChannel(Channel::Depth) || scene_->anyHasChannel(Channel::GroundTruth);
// For each source depth map // For each source depth map
for (size_t i=0; i < scene_->frames.size(); ++i) { for (size_t i=0; i < scene_->frames.size(); ++i) {
if (!scene_->hasFrame(i)) continue; //if (!scene_->hasFrame(i)) continue;
auto &f = scene_->frames[i]; auto &f = scene_->frames[i].cast<ftl::rgbd::Frame>();
//auto *s = scene_->sources[i]; //auto *s = scene_->sources[i];
if (f.empty(Channel::Colour)) { if (!f.has(Channel::Colour)) {
LOG(ERROR) << "Missing required channel"; //LOG(ERROR) << "Missing required channel";
continue; continue;
} }
// We have the needed depth data?
if (use_depth && !f.hasOwn(Channel::Depth) && !f.hasOwn(Channel::GroundTruth)) {
continue;
}
++valid_count;
//auto pose = MatrixConversion::toCUDA(t.cast<float>() * f.getPose().cast<float>()); //auto pose = MatrixConversion::toCUDA(t.cast<float>() * f.getPose().cast<float>());
auto transform = pose_ * MatrixConversion::toCUDA(t.cast<float>() * f.getPose().cast<float>()); auto transform = pose_ * MatrixConversion::toCUDA(t.cast<float>() * f.getPose().cast<float>());
...@@ -260,6 +281,7 @@ void CUDARender::_mesh(ftl::rgbd::Frame &out, const Eigen::Matrix4d &t, cudaStre ...@@ -260,6 +281,7 @@ void CUDARender::_mesh(ftl::rgbd::Frame &out, const Eigen::Matrix4d &t, cudaStre
auto &screenbuffer = _getScreenBuffer(bufsize); auto &screenbuffer = _getScreenBuffer(bufsize);
// Calculate and save virtual view screen position of each source pixel // Calculate and save virtual view screen position of each source pixel
if (use_depth) {
if (f.hasChannel(Channel::Depth)) { if (f.hasChannel(Channel::Depth)) {
ftl::cuda::screen_coord( ftl::cuda::screen_coord(
f.createTexture<float>(Channel::Depth), f.createTexture<float>(Channel::Depth),
...@@ -274,6 +296,7 @@ void CUDARender::_mesh(ftl::rgbd::Frame &out, const Eigen::Matrix4d &t, cudaStre ...@@ -274,6 +296,7 @@ void CUDARender::_mesh(ftl::rgbd::Frame &out, const Eigen::Matrix4d &t, cudaStre
screenbuffer, screenbuffer,
params_, transform, f.getLeftCamera(), stream params_, transform, f.getLeftCamera(), stream
); );
}
} else { } else {
// Constant depth version // Constant depth version
ftl::cuda::screen_coord( ftl::cuda::screen_coord(
...@@ -285,9 +308,11 @@ void CUDARender::_mesh(ftl::rgbd::Frame &out, const Eigen::Matrix4d &t, cudaStre ...@@ -285,9 +308,11 @@ void CUDARender::_mesh(ftl::rgbd::Frame &out, const Eigen::Matrix4d &t, cudaStre
// Must reset depth channel if blending // Must reset depth channel if blending
if (do_blend) { if (do_blend) {
temp_.get<GpuMat>(Channel::Depth).setTo(cv::Scalar(0x7FFFFFFF), cvstream); temp_.set<GpuMat>(Channel::Depth).setTo(cv::Scalar(0x7FFFFFFF), cvstream);
} }
depth_out_.to_gpumat().setTo(cv::Scalar(1000.0f), cvstream);
// Decide on and render triangles around each point // Decide on and render triangles around each point
ftl::cuda::triangle_render1( ftl::cuda::triangle_render1(
depthbuffer, depthbuffer,
...@@ -303,7 +328,8 @@ void CUDARender::_mesh(ftl::rgbd::Frame &out, const Eigen::Matrix4d &t, cudaStre ...@@ -303,7 +328,8 @@ void CUDARender::_mesh(ftl::rgbd::Frame &out, const Eigen::Matrix4d &t, cudaStre
// Blend this sources mesh with previous meshes // Blend this sources mesh with previous meshes
ftl::cuda::mesh_blender( ftl::cuda::mesh_blender(
temp_.getTexture<int>(Channel::Depth), temp_.getTexture<int>(Channel::Depth),
out.createTexture<float>(_getDepthChannel()), //out.createTexture<float>(_getDepthChannel()),
depth_out_,
f.createTexture<short>(Channel::Weights), f.createTexture<short>(Channel::Weights),
temp_.createTexture<float>(Channel::Weights), temp_.createTexture<float>(Channel::Weights),
params_, params_,
...@@ -315,20 +341,28 @@ void CUDARender::_mesh(ftl::rgbd::Frame &out, const Eigen::Matrix4d &t, cudaStre ...@@ -315,20 +341,28 @@ void CUDARender::_mesh(ftl::rgbd::Frame &out, const Eigen::Matrix4d &t, cudaStre
} }
} }
if (valid_count == 0) return;
// Convert from int depth to float depth // Convert from int depth to float depth
//temp_.get<GpuMat>(Channel::Depth2).convertTo(out.get<GpuMat>(Channel::Depth), CV_32F, 1.0f / 100000.0f, cvstream); //temp_.get<GpuMat>(Channel::Depth2).convertTo(out.get<GpuMat>(Channel::Depth), CV_32F, 1.0f / 100000.0f, cvstream);
if (do_blend) { if (do_blend) {
ftl::cuda::dibr_normalise( ftl::cuda::dibr_normalise(
out.getTexture<float>(_getDepthChannel()), //out.getTexture<float>(_getDepthChannel()),
out.getTexture<float>(_getDepthChannel()), //out.getTexture<float>(_getDepthChannel()),
depth_out_,
depth_out_,
temp_.getTexture<float>(Channel::Weights), temp_.getTexture<float>(Channel::Weights),
stream_ stream_
); );
} else { } else {
ftl::cuda::merge_convert_depth(temp_.getTexture<int>(Channel::Depth2), out.createTexture<float>(_getDepthChannel()), 1.0f / 100000.0f, stream_); //ftl::cuda::merge_convert_depth(temp_.getTexture<int>(Channel::Depth2), out.createTexture<float>(_getDepthChannel()), 1.0f / 100000.0f, stream_);
ftl::cuda::merge_convert_depth(temp_.getTexture<int>(Channel::Depth2), depth_out_, 1.0f / 100000.0f, stream_);
} }
// Now merge new render to any existing frameset render, detecting collisions
ftl::cuda::touch_merge(depth_out_, out.createTexture<float>(_getDepthChannel()), collisions_, 1024, touch_dist_, stream_);
//filters_->filter(out, src, stream); //filters_->filter(out, src, stream);
// Generate normals for final virtual image // Generate normals for final virtual image
...@@ -347,29 +381,30 @@ void CUDARender::_allocateChannels(ftl::rgbd::Frame &out, ftl::codecs::Channel c ...@@ -347,29 +381,30 @@ void CUDARender::_allocateChannels(ftl::rgbd::Frame &out, ftl::codecs::Channel c
// Allocate left channel buffers and clear them // Allocate left channel buffers and clear them
if (chan == Channel::Colour) { if (chan == Channel::Colour) {
//if (!out.hasChannel(Channel::Depth)) { //if (!out.hasChannel(Channel::Depth)) {
out.create<GpuMat>(Channel::Depth, Format<float>(camera.width, camera.height)); out.create<VideoFrame>(Channel::Depth).createGPU(Format<float>(camera.width, camera.height));
out.create<GpuMat>(Channel::Colour, Format<uchar4>(camera.width, camera.height)); out.create<VideoFrame>(Channel::Colour).createGPU(Format<uchar4>(camera.width, camera.height));
out.create<GpuMat>(Channel::Normals, Format<half4>(camera.width, camera.height)); out.create<VideoFrame>(Channel::Normals).createGPU(Format<half4>(camera.width, camera.height));
out.createTexture<uchar4>(Channel::Colour, true); // Force interpolated colour out.createTexture<uchar4>(Channel::Colour, ftl::rgbd::Format<uchar4>(camera.width, camera.height), true); // Force interpolated colour
out.get<GpuMat>(Channel::Depth).setTo(cv::Scalar(1000.0f), cvstream); out.set<GpuMat>(Channel::Depth).setTo(cv::Scalar(1000.0f), cvstream);
//} //}
// Allocate right channel buffers and clear them // Allocate right channel buffers and clear them
} else { } else {
if (!out.hasChannel(Channel::Depth2)) { //if (!out.hasChannel(Channel::Depth2)) {
out.create<GpuMat>(Channel::Depth2, Format<float>(camera.width, camera.height)); out.create<VideoFrame>(Channel::Depth2).createGPU(Format<float>(camera.width, camera.height));
out.create<GpuMat>(Channel::Colour2, Format<uchar4>(camera.width, camera.height)); out.create<VideoFrame>(Channel::Colour2).createGPU(Format<uchar4>(camera.width, camera.height));
out.create<GpuMat>(Channel::Normals2, Format<half4>(camera.width, camera.height)); out.create<VideoFrame>(Channel::Normals2).createGPU(Format<half4>(camera.width, camera.height));
out.createTexture<uchar4>(Channel::Colour2, true); // Force interpolated colour out.createTexture<uchar4>(Channel::Colour2, ftl::rgbd::Format<uchar4>(camera.width, camera.height), true); // Force interpolated colour
out.get<GpuMat>(Channel::Depth2).setTo(cv::Scalar(1000.0f), cvstream); out.set<GpuMat>(Channel::Depth2).setTo(cv::Scalar(1000.0f), cvstream);
} //}
} }
temp_.create<GpuMat>(Channel::Depth, Format<int>(camera.width, camera.height)); temp_.create<VideoFrame>(Channel::Depth).createGPU(Format<int>(camera.width, camera.height));
temp_.create<GpuMat>(Channel::Depth2, Format<int>(camera.width, camera.height)); temp_.create<VideoFrame>(Channel::Depth2).createGPU(Format<int>(camera.width, camera.height));
temp_.create<GpuMat>(Channel::Normals, Format<half4>(camera.width, camera.height)); temp_.create<VideoFrame>(Channel::Normals).createGPU(Format<half4>(camera.width, camera.height));
temp_.create<GpuMat>(Channel::Weights, Format<float>(camera.width, camera.height)); temp_.create<VideoFrame>(Channel::Weights).createGPU(Format<float>(camera.width, camera.height));
temp_.createTexture<int>(Channel::Depth); temp_.createTexture<int>(Channel::Depth);
depth_out_.create(camera.width, camera.height);
accum_.create(camera.width, camera.height); accum_.create(camera.width, camera.height);
contrib_.create(camera.width, camera.height); contrib_.create(camera.width, camera.height);
...@@ -404,6 +439,7 @@ void CUDARender::_updateParameters(ftl::rgbd::Frame &out, ftl::codecs::Channel c ...@@ -404,6 +439,7 @@ void CUDARender::_updateParameters(ftl::rgbd::Frame &out, ftl::codecs::Channel c
params_.disconDisparities = value("discon_disparities", 2.0f); params_.disconDisparities = value("discon_disparities", 2.0f);
params_.accumulationMode = static_cast<ftl::render::AccumulationFunction>(value("accumulation_func", 0)); params_.accumulationMode = static_cast<ftl::render::AccumulationFunction>(value("accumulation_func", 0));
params_.m_flags = 0; params_.m_flags = 0;
params_.projection = static_cast<ftl::rgbd::Projection>(value("projection",0));
if (value("normal_weight_colours", true)) params_.m_flags |= ftl::render::kNormalWeightColours; if (value("normal_weight_colours", true)) params_.m_flags |= ftl::render::kNormalWeightColours;
if (value("channel_weights", false)) params_.m_flags |= ftl::render::kUseWeightsChannel; if (value("channel_weights", false)) params_.m_flags |= ftl::render::kUseWeightsChannel;
} }
...@@ -437,7 +473,7 @@ void CUDARender::_postprocessColours(ftl::rgbd::Frame &out) { ...@@ -437,7 +473,7 @@ void CUDARender::_postprocessColours(ftl::rgbd::Frame &out) {
params_.camera, params_.camera,
stream_ stream_
); );
} else if (out.hasChannel(_getDepthChannel()) && out.hasChannel(out_chan_)) { } else if (mesh_ && out.hasChannel(_getDepthChannel()) && out.hasChannel(out_chan_)) {
ftl::cuda::fix_bad_colour( ftl::cuda::fix_bad_colour(
out.getTexture<float>(_getDepthChannel()), out.getTexture<float>(_getDepthChannel()),
out.getTexture<uchar4>(out_chan_), out.getTexture<uchar4>(out_chan_),
...@@ -466,17 +502,27 @@ void CUDARender::_renderPass2(Channels<0> chans, const Eigen::Matrix4d &t) { ...@@ -466,17 +502,27 @@ void CUDARender::_renderPass2(Channels<0> chans, const Eigen::Matrix4d &t) {
for (auto chan : chans) { for (auto chan : chans) {
ftl::codecs::Channel mapped = chan; ftl::codecs::Channel mapped = chan;
if (chan == Channel::Colour && scene_->firstFrame().hasChannel(Channel::ColourHighRes)) mapped = Channel::ColourHighRes; // FIXME: Doesn't seem to work
//if (chan == Channel::Colour && scene_->firstFrame().hasChannel(Channel::ColourHighRes)) mapped = Channel::ColourHighRes;
_renderChannel(*out_, mapped, t, stream_); _renderChannel(*out_, mapped, t, stream_);
} }
} }
void CUDARender::cancel() {
out_ = nullptr;
scene_ = nullptr;
stage_ = Stage::Finished;
cudaSafeCall(cudaStreamSynchronize(stream_));
}
void CUDARender::begin(ftl::rgbd::Frame &out, ftl::codecs::Channel chan) { void CUDARender::begin(ftl::rgbd::Frame &out, ftl::codecs::Channel chan) {
if (stage_ != Stage::Finished) { if (stage_ != Stage::Finished) {
throw FTL_Error("Already rendering"); throw FTL_Error("Already rendering");
} }
stream_ = out.stream();
out_ = &out; out_ = &out;
const auto &camera = out.getLeftCamera(); const auto &camera = out.getLeftCamera();
cv::cuda::Stream cvstream = cv::cuda::StreamAccessor::wrapStream(stream_); cv::cuda::Stream cvstream = cv::cuda::StreamAccessor::wrapStream(stream_);
...@@ -491,7 +537,7 @@ void CUDARender::begin(ftl::rgbd::Frame &out, ftl::codecs::Channel chan) { ...@@ -491,7 +537,7 @@ void CUDARender::begin(ftl::rgbd::Frame &out, ftl::codecs::Channel chan) {
// Apply a colour background // Apply a colour background
if (env_image_.empty() || !value("environment_enabled", false)) { if (env_image_.empty() || !value("environment_enabled", false)) {
out.get<GpuMat>(chan).setTo(background_, cvstream); out.set<GpuMat>(chan).setTo(background_, cvstream);
} else { } else {
auto pose = poseInverse_.getFloat3x3(); auto pose = poseInverse_.getFloat3x3();
ftl::cuda::equirectangular_reproject( ftl::cuda::equirectangular_reproject(
...@@ -502,6 +548,9 @@ void CUDARender::begin(ftl::rgbd::Frame &out, ftl::codecs::Channel chan) { ...@@ -502,6 +548,9 @@ void CUDARender::begin(ftl::rgbd::Frame &out, ftl::codecs::Channel chan) {
sets_.clear(); sets_.clear();
stage_ = Stage::ReadySubmit; stage_ = Stage::ReadySubmit;
// Reset collision data.
cudaSafeCall(cudaMemsetAsync(collisions_, 0, sizeof(int), stream_));
} }
void CUDARender::render() { void CUDARender::render() {
...@@ -566,14 +615,45 @@ void CUDARender::_endSubmit() { ...@@ -566,14 +615,45 @@ void CUDARender::_endSubmit() {
void CUDARender::_end() { void CUDARender::_end() {
_postprocessColours(*out_); _postprocessColours(*out_);
// Final OpenGL flip // Final OpenGL flip (easier to do in shader?)
ftl::cuda::flip(out_->getTexture<uchar4>(out_chan_), stream_); /*ftl::cuda::flip(out_->getTexture<uchar4>(out_chan_), stream_);*/
ftl::cuda::flip(out_->getTexture<float>(_getDepthChannel()), stream_); /*ftl::cuda::flip(out_->getTexture<float>(_getDepthChannel()), stream_);*/
cudaSafeCall(cudaMemcpyAsync(collisions_host_, collisions_, sizeof(ftl::cuda::Collision)*1024, cudaMemcpyDeviceToHost, stream_));
cudaSafeCall(cudaStreamSynchronize(stream_)); cudaSafeCall(cudaStreamSynchronize(stream_));
// Convert collisions into camera coordinates.
collision_points_.resize(collisions_host_[0].screen);
for (uint i=1; i<collisions_host_[0].screen+1; ++i) {
collision_points_[i-1] = make_float4(collisions_host_[i].x(), collisions_host_[i].y(), collisions_host_[i].depth, collisions_host_[i].strength());
}
// Do something with the collisions
/*if (collisions_host_[0].screen > 0) {
float x = 0.0f;
float y = 0.0f;
float d = 0.0f;
float w = 0.0f;
for (uint i=1; i<collisions_host_[0].screen+1; ++i) {
float inum = collisions_host_[i].strength();
int ix = collisions_host_[i].x();
int iy = collisions_host_[i].y();
x += float(ix)*inum;
y += float(iy)*inum;
d += collisions_host_[i].depth*inum;
w += inum;
} }
bool CUDARender::submit(ftl::rgbd::FrameSet *in, Channels<0> chans, const Eigen::Matrix4d &t) { x /= w;
y /= w;
d /= w;
LOG(INFO) << "Collision at: " << x << "," << y << ", " << d;
}*/
}
bool CUDARender::submit(ftl::data::FrameSet *in, Channels<0> chans, const Eigen::Matrix4d &t) {
if (stage_ != Stage::ReadySubmit) { if (stage_ != Stage::ReadySubmit) {
throw FTL_Error("Renderer not ready for submits"); throw FTL_Error("Renderer not ready for submits");
} }
...@@ -587,9 +667,8 @@ bool CUDARender::submit(ftl::rgbd::FrameSet *in, Channels<0> chans, const Eigen: ...@@ -587,9 +667,8 @@ bool CUDARender::submit(ftl::rgbd::FrameSet *in, Channels<0> chans, const Eigen:
bool success = true; bool success = true;
try { try {
_renderPass1(in->pose); _renderPass1(t);
//cudaSafeCall(cudaStreamSynchronize(stream_)); } catch (const ftl::exception &e) {
} catch (std::exception &e) {
LOG(ERROR) << "Exception in render: " << e.what(); LOG(ERROR) << "Exception in render: " << e.what();
success = false; success = false;
} }
...@@ -597,9 +676,9 @@ bool CUDARender::submit(ftl::rgbd::FrameSet *in, Channels<0> chans, const Eigen: ...@@ -597,9 +676,9 @@ bool CUDARender::submit(ftl::rgbd::FrameSet *in, Channels<0> chans, const Eigen:
auto &s = sets_.emplace_back(); auto &s = sets_.emplace_back();
s.fs = in; s.fs = in;
s.channels = chans; s.channels = chans;
s.transform = in->pose; s.transform = t;
last_frame_ = scene_->timestamp; last_frame_ = scene_->timestamp();
scene_ = nullptr; scene_ = nullptr;
return success; return success;
} }
components/renderers/cpp/src/clipping.cu
View file @ 70063e8d
...@@ -27,12 +27,15 @@ __global__ void clipping_kernel(ftl::cuda::TextureObject<float> depth, ftl::cuda ...@@ -27,12 +27,15 @@ __global__ void clipping_kernel(ftl::cuda::TextureObject<float> depth, ftl::cuda
const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x; const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y; const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
if (x < depth.width() && y < depth.height()) { const float sx = float(x) / float(colour.width()) * float(depth.width());
float d = depth(x,y); const float sy = float(y) / float(colour.height()) * float(depth.height());
float4 p = make_float4(camera.screenToCam(x,y,d), 0.0f);
if (sx >= 0.0f && sx < depth.width() && sy < depth.height() && sy >= 0.0f) {
float d = depth(sx,sy);
float4 p = make_float4(camera.screenToCam(sx,sy,d), 0.0f);
if (d <= camera.minDepth || d >= camera.maxDepth || isClipped(p, clip)) { if (d <= camera.minDepth || d >= camera.maxDepth || isClipped(p, clip)) {
depth(x,y) = 0.0f; depth(sx,sy) = 0.0f;
colour(x,y) = make_uchar4(0,0,0,0); colour(x,y) = make_uchar4(0,0,0,0);
} }
} }
...@@ -54,7 +57,7 @@ void ftl::cuda::clipping(ftl::cuda::TextureObject<float> &depth, ...@@ -54,7 +57,7 @@ void ftl::cuda::clipping(ftl::cuda::TextureObject<float> &depth,
const ftl::rgbd::Camera &camera, const ftl::rgbd::Camera &camera,
const ClipSpace &clip, cudaStream_t stream) { const ClipSpace &clip, cudaStream_t stream) {
const dim3 gridSize((depth.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (depth.height() + T_PER_BLOCK - 1)/T_PER_BLOCK); const dim3 gridSize((colour.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (colour.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK); const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
clipping_kernel<<<gridSize, blockSize, 0, stream>>>(depth, colour, camera, clip); clipping_kernel<<<gridSize, blockSize, 0, stream>>>(depth, colour, camera, clip);
......
components/renderers/cpp/src/colouriser.cpp
View file @ 70063e8d
...@@ -2,6 +2,7 @@ ...@@ -2,6 +2,7 @@
#include "splatter_cuda.hpp" #include "splatter_cuda.hpp"
#include <ftl/cuda/colour_cuda.hpp> #include <ftl/cuda/colour_cuda.hpp>
#include <ftl/cuda/normals.hpp> #include <ftl/cuda/normals.hpp>
#include <ftl/operators/cuda/mask.hpp>
#include <opencv2/cudaarithm.hpp> #include <opencv2/cudaarithm.hpp>
#include <opencv2/cudaimgproc.hpp> #include <opencv2/cudaimgproc.hpp>
...@@ -113,9 +114,13 @@ Colouriser::~Colouriser() { ...@@ -113,9 +114,13 @@ Colouriser::~Colouriser() {
} }
TextureObject<uchar4> &Colouriser::colourise(ftl::rgbd::Frame &f, Channel c, cudaStream_t stream) { TextureObject<uchar4> &Colouriser::colourise(ftl::rgbd::Frame &f, Channel c, cudaStream_t stream) {
const auto &vf = f.get<ftl::rgbd::VideoFrame>(c);
if (!vf.isGPU()) {
f.upload(c);
}
switch (c) { switch (c) {
case Channel::Overlay : return f.createTexture<uchar4>(c); case Channel::Overlay :
case Channel::ColourHighRes :
case Channel::Colour : case Channel::Colour :
case Channel::Colour2 : return _processColour(f,c,stream); case Channel::Colour2 : return _processColour(f,c,stream);
case Channel::GroundTruth : case Channel::GroundTruth :
...@@ -183,7 +188,7 @@ TextureObject<uchar4> &Colouriser::_processColour(ftl::rgbd::Frame &f, Channel c ...@@ -183,7 +188,7 @@ TextureObject<uchar4> &Colouriser::_processColour(ftl::rgbd::Frame &f, Channel c
bool colour_sources = value("colour_sources", false); bool colour_sources = value("colour_sources", false);
if (!colour_sources && show_mask == 0) { if (!colour_sources && show_mask == 0) {
return f.createTexture<uchar4>(c); return f.createTexture<uchar4>(c, true);
} }
cv::cuda::Stream cvstream = cv::cuda::StreamAccessor::wrapStream(stream); cv::cuda::Stream cvstream = cv::cuda::StreamAccessor::wrapStream(stream);
...@@ -192,7 +197,7 @@ TextureObject<uchar4> &Colouriser::_processColour(ftl::rgbd::Frame &f, Channel c ...@@ -192,7 +197,7 @@ TextureObject<uchar4> &Colouriser::_processColour(ftl::rgbd::Frame &f, Channel c
auto &buf = _getBuffer(size.width, size.height); auto &buf = _getBuffer(size.width, size.height);
if (colour_sources) { if (colour_sources) {
auto colour = HSVtoRGB(360 / 8 * f.id, 0.6, 0.85); auto colour = HSVtoRGB(360 / 8 * f.source(), 0.6, 0.85);
buf.to_gpumat().setTo(colour, cvstream); buf.to_gpumat().setTo(colour, cvstream);
} }
......
components/renderers/cpp/src/dibr.cu
View file @ 70063e8d
...@@ -9,10 +9,12 @@ ...@@ -9,10 +9,12 @@
using ftl::cuda::TextureObject; using ftl::cuda::TextureObject;
using ftl::render::Parameters; using ftl::render::Parameters;
using ftl::rgbd::Projection;
/* /*
* DIBR point cloud with a depth check * DIBR point cloud with a depth check
*/ */
template <Projection PROJECT>
__global__ void dibr_merge_kernel(TextureObject<float> depth, __global__ void dibr_merge_kernel(TextureObject<float> depth,
TextureObject<int> depth_out, TextureObject<int> depth_out,
float4x4 transform, float4x4 transform,
...@@ -26,11 +28,13 @@ using ftl::render::Parameters; ...@@ -26,11 +28,13 @@ using ftl::render::Parameters;
const float3 camPos = transform * cam.screenToCam(x,y,d0); const float3 camPos = transform * cam.screenToCam(x,y,d0);
const float d = camPos.z; //const float d = camPos.z;
const uint2 screenPos = params.camera.camToScreen<uint2>(camPos); //const uint2 screenPos = params.camera.camToScreen<uint2>(camPos);
const unsigned int cx = screenPos.x; const float3 screenPos = params.camera.project<PROJECT>(camPos);
const unsigned int cy = screenPos.y; const unsigned int cx = (unsigned int)(screenPos.x+0.5f);
const unsigned int cy = (unsigned int)(screenPos.y+0.5f);
const float d = screenPos.z;
if (d > params.camera.minDepth && d < params.camera.maxDepth && cx < depth_out.width() && cy < depth_out.height()) { if (d > params.camera.minDepth && d < params.camera.maxDepth && cx < depth_out.width() && cy < depth_out.height()) {
// Transform estimated point to virtual cam space and output z // Transform estimated point to virtual cam space and output z
atomicMin(&depth_out(cx,cy), d * 100000.0f); atomicMin(&depth_out(cx,cy), d * 100000.0f);
...@@ -67,7 +71,11 @@ void ftl::cuda::dibr_merge(TextureObject<float> &depth, TextureObject<int> &dept ...@@ -67,7 +71,11 @@ void ftl::cuda::dibr_merge(TextureObject<float> &depth, TextureObject<int> &dept
const dim3 gridSize((depth.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (depth.height() + T_PER_BLOCK - 1)/T_PER_BLOCK); const dim3 gridSize((depth.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (depth.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK); const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
dibr_merge_kernel<<<gridSize, blockSize, 0, stream>>>(depth, depth_out, transform, cam, params); if (params.projection == Projection::PERSPECTIVE) {
dibr_merge_kernel<Projection::PERSPECTIVE><<<gridSize, blockSize, 0, stream>>>(depth, depth_out, transform, cam, params);
} else {
dibr_merge_kernel<Projection::ORTHOGRAPHIC><<<gridSize, blockSize, 0, stream>>>(depth, depth_out, transform, cam, params);
}
cudaSafeCall( cudaGetLastError() ); cudaSafeCall( cudaGetLastError() );
} }
......
applications/gui/src/gltexture.cpp → components/renderers/cpp/src/gltexture.cpp
View file @ 70063e8d
#include "gltexture.hpp" #include <ftl/utility/gltexture.hpp>
#include <nanogui/opengl.h> #include <nanogui/opengl.h>
#include <loguru.hpp> #include <loguru.hpp>
#include <ftl/cuda_common.hpp> #include <ftl/cuda_common.hpp>
#include <cuda_gl_interop.h> #include <cuda_gl_interop.h>
#include <opencv2/core/cuda_stream_accessor.hpp>
#include <ftl/exception.hpp> #include <ftl/exception.hpp>
using ftl::gui::GLTexture; void log_error() {
auto err = glGetError();
if (err != 0) LOG(ERROR) << "OpenGL Texture error: " << err;
}
GLTexture::GLTexture(GLTexture::Type type) { using ftl::utility::GLTexture;
GLTexture::GLTexture() {
glid_ = std::numeric_limits<unsigned int>::max(); glid_ = std::numeric_limits<unsigned int>::max();
glbuf_ = std::numeric_limits<unsigned int>::max(); glbuf_ = std::numeric_limits<unsigned int>::max();
cuda_res_ = nullptr; cuda_res_ = nullptr;
width_ = 0; width_ = 0;
height_ = 0; height_ = 0;
changed_ = true; type_ = Type::RGBA;
type_ = type;
} }
GLTexture::~GLTexture() { GLTexture::~GLTexture() {
//glDeleteTextures(1, &glid_); free(); // Note: Do not simply remove this...
} }
void GLTexture::update(cv::Mat &m) { void GLTexture::make(int width, int height, Type type) {
LOG(INFO) << "DEPRECATED"; if (width != width_ || height != height_ || type_ != type) {
if (m.rows == 0) return;
if (glid_ == std::numeric_limits<unsigned int>::max()) {
glGenTextures(1, &glid_);
glBindTexture(GL_TEXTURE_2D, glid_);
//cv::Mat m(cv::Size(100,100), CV_8UC3);
if (type_ == Type::BGRA) {
glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA8, m.cols, m.rows, 0, GL_BGRA, GL_UNSIGNED_BYTE, m.data);
} else if (type_ == Type::Float) {
glTexImage2D(GL_TEXTURE_2D, 0, GL_R32F, m.cols, m.rows, 0, GL_RED, GL_FLOAT, m.data);
}
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_CLAMP_TO_EDGE);
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_CLAMP_TO_EDGE);
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST);
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_LINEAR);
} else {
//glBindTexture(GL_TEXTURE_2D, glid_);
// TODO Allow for other formats
//glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA8, m.cols, m.rows, 0, GL_BGRA, GL_UNSIGNED_BYTE, m.data);
}
auto err = glGetError();
if (err != 0) LOG(ERROR) << "OpenGL Texture error: " << err;
}
void GLTexture::make(int width, int height) {
if (width != width_ || height != height_) {
free(); free();
} }
...@@ -58,31 +39,33 @@ void GLTexture::make(int width, int height) { ...@@ -58,31 +39,33 @@ void GLTexture::make(int width, int height) {
width_ = width; width_ = width;
height_ = height; height_ = height;
stride_ = ((width*4) % ALIGNMENT != 0) ? ((width*4) + (ALIGNMENT - ((width*4) % ALIGNMENT))) / 4 : width; stride_ = ((width*4) % ALIGNMENT != 0) ?
((width*4) + (ALIGNMENT - ((width*4) % ALIGNMENT))) / 4:
width;
type_ = type;
if (width == 0 || height == 0) { if (width == 0 || height == 0) {
throw FTL_Error("Invalid texture size"); throw FTL_Error("Invalid texture size");
} }
if (glid_ == std::numeric_limits<unsigned int>::max()) { if (glid_ == std::numeric_limits<unsigned int>::max()) {
glGenTextures(1, &glid_); glGenTextures(1, &glid_); log_error();
glBindTexture(GL_TEXTURE_2D, glid_); glBindTexture(GL_TEXTURE_2D, glid_); log_error();
glPixelStorei(GL_UNPACK_ROW_LENGTH, stride_); log_error();
glPixelStorei(GL_UNPACK_ROW_LENGTH, stride_);
//cv::Mat m(cv::Size(100,100), CV_8UC3);
//glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA8, width, height, 0, GL_BGRA, GL_UNSIGNED_BYTE, nullptr);
if (type_ == Type::BGRA) { if (type_ == Type::BGRA) {
glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA8, width, height, 0, GL_BGRA, GL_UNSIGNED_BYTE, nullptr); glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA8, width, height, 0, GL_BGRA, GL_UNSIGNED_BYTE, nullptr);
} else if (type_ == Type::Float) { } else if (type_ == Type::Float) {
glTexImage2D(GL_TEXTURE_2D, 0, GL_R32F, width, height, 0, GL_RED, GL_FLOAT, nullptr); glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA8, width, height, 0, GL_BGRA, GL_UNSIGNED_BYTE, nullptr);
} }
log_error();
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_CLAMP_TO_EDGE); glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_CLAMP_TO_EDGE);
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_CLAMP_TO_EDGE); glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_CLAMP_TO_EDGE);
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST); glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST);
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_LINEAR); glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_LINEAR);
auto err = glGetError(); log_error();
if (err != 0) LOG(ERROR) << "OpenGL Texture error: " << err;
glPixelStorei(GL_UNPACK_ROW_LENGTH, 0); glPixelStorei(GL_UNPACK_ROW_LENGTH, 0);
...@@ -91,11 +74,12 @@ void GLTexture::make(int width, int height) { ...@@ -91,11 +74,12 @@ void GLTexture::make(int width, int height) {
glGenBuffers(1, &glbuf_); glGenBuffers(1, &glbuf_);
// Make this the current UNPACK buffer (OpenGL is state-based) // Make this the current UNPACK buffer (OpenGL is state-based)
glBindBuffer(GL_PIXEL_UNPACK_BUFFER, glbuf_); glBindBuffer(GL_PIXEL_UNPACK_BUFFER, glbuf_);
// Allocate data for the buffer. 4-channel 8-bit image // Allocate data for the buffer. 4-channel 8-bit image or 1-channel float
glBufferData(GL_PIXEL_UNPACK_BUFFER, stride_ * height * 4, NULL, GL_DYNAMIC_COPY); glBufferData(GL_PIXEL_UNPACK_BUFFER, stride_ * height * 4, NULL, GL_DYNAMIC_COPY);
cudaSafeCall(cudaGraphicsGLRegisterBuffer(&cuda_res_, glbuf_, cudaGraphicsRegisterFlagsWriteDiscard)); cudaSafeCall(cudaGraphicsGLRegisterBuffer(&cuda_res_, glbuf_, cudaGraphicsRegisterFlagsWriteDiscard));
glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0); glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0);
log_error();
} }
} }
...@@ -114,6 +98,7 @@ void GLTexture::free() { ...@@ -114,6 +98,7 @@ void GLTexture::free() {
} }
cv::cuda::GpuMat GLTexture::map(cudaStream_t stream) { cv::cuda::GpuMat GLTexture::map(cudaStream_t stream) {
mtx_.lock();
void *devptr; void *devptr;
size_t size; size_t size;
cudaSafeCall(cudaGraphicsMapResources(1, &cuda_res_, stream)); cudaSafeCall(cudaGraphicsMapResources(1, &cuda_res_, stream));
...@@ -122,13 +107,15 @@ cv::cuda::GpuMat GLTexture::map(cudaStream_t stream) { ...@@ -122,13 +107,15 @@ cv::cuda::GpuMat GLTexture::map(cudaStream_t stream) {
} }
void GLTexture::unmap(cudaStream_t stream) { void GLTexture::unmap(cudaStream_t stream) {
// note: code must not throw, otherwise mtx_.unlock() does not happen
cudaSafeCall(cudaGraphicsUnmapResources(1, &cuda_res_, stream)); cudaSafeCall(cudaGraphicsUnmapResources(1, &cuda_res_, stream));
changed_ = true;
//glActiveTexture(GL_TEXTURE0); //glActiveTexture(GL_TEXTURE0);
glBindBuffer(GL_PIXEL_UNPACK_BUFFER, glbuf_); glBindBuffer(GL_PIXEL_UNPACK_BUFFER, glbuf_);
// Select the appropriate texture // Select the appropriate texture
glBindTexture(GL_TEXTURE_2D, glid_); glBindTexture(GL_TEXTURE_2D, glid_);
glPixelStorei(GL_UNPACK_ROW_LENGTH, stride_); glPixelStorei(GL_UNPACK_ROW_LENGTH, stride_);
// Make a texture from the buffer // Make a texture from the buffer
if (type_ == Type::BGRA) { if (type_ == Type::BGRA) {
...@@ -139,6 +126,8 @@ void GLTexture::unmap(cudaStream_t stream) { ...@@ -139,6 +126,8 @@ void GLTexture::unmap(cudaStream_t stream) {
glPixelStorei(GL_UNPACK_ROW_LENGTH, 0); glPixelStorei(GL_UNPACK_ROW_LENGTH, 0);
glBindTexture(GL_TEXTURE_2D, 0); glBindTexture(GL_TEXTURE_2D, 0);
glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0); glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0);
mtx_.unlock();
} }
unsigned int GLTexture::texture() const { unsigned int GLTexture::texture() const {
...@@ -153,6 +142,47 @@ unsigned int GLTexture::texture() const { ...@@ -153,6 +142,47 @@ unsigned int GLTexture::texture() const {
return glid_; return glid_;
} else { } else {
return glid_; throw FTL_Error("No OpenGL texture; use make() first");
}
}
void GLTexture::copyFrom(const ftl::cuda::TextureObject<uchar4> &buffer, cudaStream_t stream) {
if (buffer.width() == 0 || buffer.height() == 0) {
return;
} }
make(buffer.width(), buffer.height(), ftl::utility::GLTexture::Type::BGRA);
auto dst = map(stream);
cudaSafeCall(cudaMemcpy2D( dst.data, dst.step, buffer.devicePtr(), buffer.pitch(),
buffer.width()*4, buffer.height(), cudaMemcpyDeviceToDevice));
unmap(stream);
}
void GLTexture::copyFrom(const cv::Mat &im, cudaStream_t stream) {
if (im.rows == 0 || im.cols == 0 || im.channels() != 4 || im.type() != CV_8UC4) {
LOG(ERROR) << __FILE__ << ":" << __LINE__ << ": " << "bad OpenCV format";
return;
}
auto cvstream = cv::cuda::StreamAccessor::wrapStream(stream);
make(im.cols, im.rows, ftl::utility::GLTexture::Type::BGRA);
auto dst = map(stream);
dst.upload(im);
unmap(stream);
}
void GLTexture::copyFrom(const cv::cuda::GpuMat &im, cudaStream_t stream) {
if (im.rows == 0 || im.cols == 0 || im.channels() != 4 || im.type() != CV_8UC4) {
LOG(ERROR) << __FILE__ << ":" << __LINE__ << ": " << "bad OpenCV format";
return;
}
auto cvstream = cv::cuda::StreamAccessor::wrapStream(stream);
make(im.cols, im.rows, ftl::utility::GLTexture::Type::BGRA);
auto dst = map(stream);
im.copyTo(dst, cvstream);
unmap(stream);
} }
components/renderers/cpp/src/overlay.cpp
View file @ 70063e8d
#include <ftl/render/overlay.hpp> #include <ftl/render/overlay.hpp>
#include <ftl/utility/matrix_conversion.hpp> #include <ftl/utility/matrix_conversion.hpp>
#include <ftl/cuda_common.hpp>
#include <opencv2/imgproc.hpp> #include <opencv2/imgproc.hpp>
#include <ftl/codecs/shapes.hpp> #include <ftl/codecs/shapes.hpp>
#include <ftl/operators/poser.hpp>
#define LOGURU_REPLACE_GLOG 1 #define LOGURU_REPLACE_GLOG 1
#include <loguru.hpp> #include <loguru.hpp>
...@@ -21,6 +23,8 @@ namespace { ...@@ -21,6 +23,8 @@ namespace {
uniform float height; uniform float height;
uniform float far; uniform float far;
uniform float near; uniform float near;
//uniform float offset_x;
//uniform float offset_y;
uniform mat4 pose; uniform mat4 pose;
uniform vec3 scale; uniform vec3 scale;
...@@ -32,8 +36,8 @@ namespace { ...@@ -32,8 +36,8 @@ namespace {
// (vert.z-near) / (far-near) * 2.0 - 1.0, 1.0); // (vert.z-near) / (far-near) * 2.0 - 1.0, 1.0);
vec4 pos = vec4( vec4 pos = vec4(
vert.x*focal / (width/2.0), (vert.x*focal) / (width/2.0),
-vert.y*focal / (height/2.0), (-vert.y*focal) / (height/2.0),
-vert.z * ((far+near) / (far-near)) + (2.0 * near * far / (far-near)), -vert.z * ((far+near) / (far-near)) + (2.0 * near * far / (far-near)),
//((vert.z - near) / (far - near) * 2.0 - 1.0) * vert.z, //((vert.z - near) / (far - near) * 2.0 - 1.0) * vert.z,
vert.z vert.z
...@@ -165,8 +169,8 @@ void Overlay::_createShapes() { ...@@ -165,8 +169,8 @@ void Overlay::_createShapes() {
shapes_[Shape::XZPLANE] = {88, 2*3, 94, 40*2}; shapes_[Shape::XZPLANE] = {88, 2*3, 94, 40*2};
shapes_[Shape::AXIS] = {0, 0, 82, 2*3}; shapes_[Shape::AXIS] = {0, 0, 82, 2*3};
oShader.uploadAttrib("vertex", sizeof(float3)*shape_verts_.size(), 3, sizeof(float), GL_FLOAT, false, shape_verts_.data()); oShader.uploadAttrib("vertex", 3*shape_verts_.size(), 3, sizeof(float), GL_FLOAT, false, shape_verts_.data());
oShader.uploadAttrib ("indices", sizeof(int)*shape_tri_indices_.size(), 1, sizeof(int), GL_UNSIGNED_INT, true, shape_tri_indices_.data()); oShader.uploadAttrib ("indices", 1*shape_tri_indices_.size(), 1, sizeof(int), GL_UNSIGNED_INT, true, shape_tri_indices_.data());
} }
void Overlay::_drawFilledShape(Shape shape, const Eigen::Matrix4d &pose, float scale, uchar4 c) { void Overlay::_drawFilledShape(Shape shape, const Eigen::Matrix4d &pose, float scale, uchar4 c) {
...@@ -241,12 +245,12 @@ void Overlay::_drawAxis(const Eigen::Matrix4d &pose, const Eigen::Vector3f &scal ...@@ -241,12 +245,12 @@ void Overlay::_drawAxis(const Eigen::Matrix4d &pose, const Eigen::Vector3f &scal
(const void *)(loffset * sizeof(uint32_t))); (const void *)(loffset * sizeof(uint32_t)));
} }
void Overlay::draw(ftl::rgbd::FrameSet &fs, ftl::rgbd::FrameState &state, const Eigen::Vector2f &screenSize) { void Overlay::draw(NVGcontext *ctx, ftl::data::FrameSet &fs, ftl::rgbd::Frame &frame, const Eigen::Vector2f &screenSize, const Eigen::Vector2f &imageSize, const Eigen::Vector2f &offset, const Eigen::Matrix4d &cursor) {
if (!value("enabled", false)) return; if (!value("enabled", false)) return;
double zfar = 8.0f; double zfar = 8.0f;
auto intrin = state.getLeft(); auto intrin = frame.getLeft();
intrin = intrin.scaled(screenSize[0], screenSize[1]); intrin = intrin.scaled(imageSize[0], imageSize[1]);
if (!init_) { if (!init_) {
oShader.init("OverlayShader", overlayVertexShader, overlayFragmentShader); oShader.init("OverlayShader", overlayVertexShader, overlayFragmentShader);
...@@ -263,7 +267,7 @@ void Overlay::draw(ftl::rgbd::FrameSet &fs, ftl::rgbd::FrameState &state, const ...@@ -263,7 +267,7 @@ void Overlay::draw(ftl::rgbd::FrameSet &fs, ftl::rgbd::FrameState &state, const
{0.8f, -0.4f, 2.0f} {0.8f, -0.4f, 2.0f}
}; };
auto pose = MatrixConversion::toCUDA(state.getPose().cast<float>().inverse()); auto pose = MatrixConversion::toCUDA(frame.getPose().cast<float>().inverse());
tris[0] = pose * tris[0]; tris[0] = pose * tris[0];
tris[1] = pose * tris[1]; tris[1] = pose * tris[1];
...@@ -278,10 +282,12 @@ void Overlay::draw(ftl::rgbd::FrameSet &fs, ftl::rgbd::FrameState &state, const ...@@ -278,10 +282,12 @@ void Overlay::draw(ftl::rgbd::FrameSet &fs, ftl::rgbd::FrameState &state, const
glEnable(GL_LINE_SMOOTH); glEnable(GL_LINE_SMOOTH);
oShader.setUniform("focal", intrin.fx); oShader.setUniform("focal", intrin.fx);
oShader.setUniform("width", float(intrin.width)); oShader.setUniform("width", screenSize[0]);
oShader.setUniform("height", float(intrin.height)); oShader.setUniform("height", screenSize[1]);
oShader.setUniform("far", zfar); oShader.setUniform("far", zfar);
oShader.setUniform("near", 0.1f); // TODO: but make sure CUDA depth is also normalised like this oShader.setUniform("near", 0.1f); // TODO: but make sure CUDA depth is also normalised like this
//oShader.setUniform("offset_x", offset[0]);
//oShader.setUniform("offset_y", offset[1]);
/*oShader.setUniform("blockColour", Eigen::Vector4f(1.0f,1.0f,0.0f,0.5f)); /*oShader.setUniform("blockColour", Eigen::Vector4f(1.0f,1.0f,0.0f,0.5f));
oShader.uploadAttrib("vertex", sizeof(tris), 3, sizeof(float), GL_FLOAT, false, tris); oShader.uploadAttrib("vertex", sizeof(tris), 3, sizeof(float), GL_FLOAT, false, tris);
...@@ -293,13 +299,26 @@ void Overlay::draw(ftl::rgbd::FrameSet &fs, ftl::rgbd::FrameState &state, const ...@@ -293,13 +299,26 @@ void Overlay::draw(ftl::rgbd::FrameSet &fs, ftl::rgbd::FrameState &state, const
//glFinish(); //glFinish();
if (value("show_poses", false)) { if (value("show_poses", true)) {
for (size_t i=0; i<fs.frames.size(); ++i) { for (size_t i=0; i<fs.frames.size(); ++i) {
auto pose = fs.frames[i].getPose(); //.inverse() * state.getPose(); auto &f = fs.frames[i].cast<ftl::rgbd::Frame>();
if (f.id().id == frame.id().id) continue;
auto name = fs.frames[i].get<std::string>("name"); auto pose = f.getPose(); //.inverse() * state.getPose();
_drawOutlinedShape(Shape::CAMERA, state.getPose().inverse() * pose, Eigen::Vector3f(0.2f,0.2f,0.2f), make_uchar4(255,0,0,80), make_uchar4(255,0,0,255));
_drawAxis(state.getPose().inverse() * pose, Eigen::Vector3f(0.2f, 0.2f, 0.2f)); std::string name = fs.frames[0].name();
auto tpose = frame.getPose().inverse() * pose;
_drawOutlinedShape(Shape::CAMERA, tpose, Eigen::Vector3f(0.2f,0.2f,0.2f), make_uchar4(255,0,0,80), make_uchar4(255,0,0,255));
_drawAxis(tpose, Eigen::Vector3f(0.2f, 0.2f, 0.2f));
float3 textpos;
textpos.x = tpose(0,3);
textpos.y = tpose(1,3);
textpos.z = tpose(2,3);
float2 textscreen = f.getLeft().camToScreen<float2>(textpos);
if (textpos.z > 0.1f) nvgText(ctx, textscreen.x, textscreen.y, name.c_str(), nullptr);
//ftl::overlay::drawCamera(state.getLeft(), out, over_depth_, fs.frames[i].getLeftCamera(), pose, cv::Scalar(0,0,255,255), 0.2,value("show_frustrum", false)); //ftl::overlay::drawCamera(state.getLeft(), out, over_depth_, fs.frames[i].getLeftCamera(), pose, cv::Scalar(0,0,255,255), 0.2,value("show_frustrum", false));
//if (name) ftl::overlay::drawText(state.getLeft(), out, over_depth_, *name, pos, 0.5, cv::Scalar(0,0,255,255)); //if (name) ftl::overlay::drawText(state.getLeft(), out, over_depth_, *name, pos, 0.5, cv::Scalar(0,0,255,255));
...@@ -308,15 +327,20 @@ void Overlay::draw(ftl::rgbd::FrameSet &fs, ftl::rgbd::FrameState &state, const ...@@ -308,15 +327,20 @@ void Overlay::draw(ftl::rgbd::FrameSet &fs, ftl::rgbd::FrameState &state, const
if (value("show_xz_plane", false)) { if (value("show_xz_plane", false)) {
float gscale = value("grid_scale",0.5f); float gscale = value("grid_scale",0.5f);
_drawOutlinedShape(Shape::XZPLANE, state.getPose().inverse(), Eigen::Vector3f(gscale,gscale,gscale), make_uchar4(200,200,200,50), make_uchar4(255,255,255,100)); _drawOutlinedShape(Shape::XZPLANE, frame.getPose().inverse(), Eigen::Vector3f(gscale,gscale,gscale), make_uchar4(200,200,200,50), make_uchar4(255,255,255,100));
} }
if (value("show_axis", true)) { if (value("show_axis", true)) {
_drawAxis(state.getPose().inverse(), Eigen::Vector3f(0.5f, 0.5f, 0.5f)); _drawAxis(frame.getPose().inverse(), Eigen::Vector3f(0.5f, 0.5f, 0.5f));
}
if (value("show_cursor", true)) {
_drawAxis(frame.getPose().inverse() * cursor.inverse(), Eigen::Vector3f(0.2f, 0.2f, 0.2f));
_drawOutlinedShape(Shape::XZPLANE, frame.getPose().inverse() * cursor.inverse(), Eigen::Vector3f(0.05f, 0.05f, 0.05f), make_uchar4(200,200,200,50), make_uchar4(255,255,255,100));
} }
if (value("show_shapes", false)) { if (value("show_shapes", true)) {
if (fs.hasChannel(Channel::Shapes3D)) { /*if (fs.hasChannel(Channel::Shapes3D)) {
std::vector<ftl::codecs::Shape3D> shapes; std::vector<ftl::codecs::Shape3D> shapes;
fs.get(Channel::Shapes3D, shapes); fs.get(Channel::Shapes3D, shapes);
...@@ -340,32 +364,43 @@ void Overlay::draw(ftl::rgbd::FrameSet &fs, ftl::rgbd::FrameState &state, const ...@@ -340,32 +364,43 @@ void Overlay::draw(ftl::rgbd::FrameSet &fs, ftl::rgbd::FrameState &state, const
//ftl::overlay::drawBox(state.getLeft(), out, over_depth_, pose, cv::Scalar(0,0,255,255), s.size.cast<double>()); //ftl::overlay::drawBox(state.getLeft(), out, over_depth_, pose, cv::Scalar(0,0,255,255), s.size.cast<double>());
//ftl::overlay::drawText(state.getLeft(), out, over_depth_, s.label, pos, 0.5, cv::Scalar(0,0,255,100)); //ftl::overlay::drawText(state.getLeft(), out, over_depth_, s.label, pos, 0.5, cv::Scalar(0,0,255,100));
} }
} }*/
for (size_t i=0; i<fs.frames.size(); ++i) { auto shapes = ftl::operators::Poser::getAll(fs.frameset());
if (fs.frames[i].hasChannel(Channel::Shapes3D)) {
std::vector<ftl::codecs::Shape3D> shapes; for (auto *ps : shapes) {
fs.frames[i].get(Channel::Shapes3D, shapes); auto &s = *ps;
for (auto &s : shapes) {
auto pose = s.pose.cast<double>(); //.inverse() * state.getPose(); auto pose = s.pose.cast<double>(); //.inverse() * state.getPose();
//Eigen::Vector4d pos = pose.inverse() * Eigen::Vector4d(0,0,0,1); //Eigen::Vector4d pos = pose.inverse() * Eigen::Vector4d(0,0,0,1);
//pos /= pos[3]; //pos /= pos[3];
Eigen::Vector3f scale(s.size[0]/2.0f, s.size[1]/2.0f, s.size[2]/2.0f); Eigen::Vector3f scale(s.size[0]/2.0f, s.size[1]/2.0f, s.size[2]/2.0f);
auto tpose = frame.getPose().inverse() * pose;
switch (s.type) { switch (s.type) {
case ftl::codecs::Shape3DType::CLIPPING: _drawOutlinedShape(Shape::BOX, state.getPose().inverse() * pose, scale, make_uchar4(255,0,255,80), make_uchar4(255,0,255,255)); break; case ftl::codecs::Shape3DType::CAMERA: _drawOutlinedShape(Shape::CAMERA, tpose, scale, make_uchar4(255,0,0,80), make_uchar4(255,0,0,255)); break;
case ftl::codecs::Shape3DType::ARUCO: _drawAxis(state.getPose().inverse() * pose, Eigen::Vector3f(0.2f, 0.2f, 0.2f)); break; case ftl::codecs::Shape3DType::CLIPPING: _drawOutlinedShape(Shape::BOX, tpose, scale, make_uchar4(255,0,255,80), make_uchar4(255,0,255,255)); break;
case ftl::codecs::Shape3DType::ARUCO: _drawAxis(tpose, Eigen::Vector3f(0.2f, 0.2f, 0.2f)); break;
case ftl::codecs::Shape3DType::CURSOR: _drawAxis(tpose, Eigen::Vector3f(0.2f, 0.2f, 0.2f)); break;
default: break; default: break;
} }
if (s.label.size() > 0) {
float3 textpos;
textpos.x = tpose(0,3);
textpos.y = tpose(1,3);
textpos.z = tpose(2,3);
float2 textscreen = frame.getLeft().camToScreen<float2>(textpos);
if (textpos.z > 0.1f) nvgText(ctx, textscreen.x, textscreen.y, s.label.c_str(), nullptr);
}
//ftl::overlay::drawBox(state.getLeft(), out, over_depth_, pose, cv::Scalar(0,0,255,100), s.size.cast<double>()); //ftl::overlay::drawBox(state.getLeft(), out, over_depth_, pose, cv::Scalar(0,0,255,100), s.size.cast<double>());
//ftl::overlay::drawText(state.getLeft(), out, over_depth_, s.label, pos, 0.5, cv::Scalar(0,0,255,100)); //ftl::overlay::drawText(state.getLeft(), out, over_depth_, s.label, pos, 0.5, cv::Scalar(0,0,255,100));
} }
} }
}
}
glDisable(GL_LINE_SMOOTH); glDisable(GL_LINE_SMOOTH);
glDisable(GL_BLEND); glDisable(GL_BLEND);
......
components/renderers/cpp/src/reprojection.cu
View file @ 70063e8d
...@@ -14,6 +14,7 @@ using ftl::render::Parameters; ...@@ -14,6 +14,7 @@ using ftl::render::Parameters;
using ftl::rgbd::Camera; using ftl::rgbd::Camera;
using ftl::render::ViewPortMode; using ftl::render::ViewPortMode;
using ftl::render::AccumulationFunction; using ftl::render::AccumulationFunction;
using ftl::rgbd::Projection;
/*template <typename T> /*template <typename T>
__device__ inline T generateInput(const T &in, const SplatParams &params, const float4 &worldPos) { __device__ inline T generateInput(const T &in, const SplatParams &params, const float4 &worldPos) {
...@@ -133,13 +134,13 @@ __device__ inline float2 convertScreen<ViewPortMode::Stretch>(const Parameters & ...@@ -133,13 +134,13 @@ __device__ inline float2 convertScreen<ViewPortMode::Stretch>(const Parameters &
} }
template <typename A> template <typename A>
__device__ inline auto getInput(TextureObject<A> &in, const float2 &screen, float width, float height) { __device__ inline auto getInput(TextureObject<A> &in, const float3 &screen, float width, float height) {
const float inSX = float(in.width()) / width; const float inSX = float(in.width()) / width;
const float inSY = float(in.height()) / height; const float inSY = float(in.height()) / height;
return in.tex2D(screen.x*inSX, screen.y*inSY); return in.tex2D(screen.x*inSX, screen.y*inSY);
} }
__device__ float weightByNormal(TextureObject<half4> &normals, int x, int y, const float3x3 &transformR, const float2 &screenPos, const ftl::rgbd::Camera &camera) { __device__ float weightByNormal(TextureObject<half4> &normals, int x, int y, const float3x3 &transformR, const float3 &screenPos, const ftl::rgbd::Camera &camera) {
// Calculate the dot product of surface normal and camera ray // Calculate the dot product of surface normal and camera ray
const float3 n = transformR * make_float3(normals.tex2D(x, y)); const float3 n = transformR * make_float3(normals.tex2D(x, y));
float3 ray = camera.screenToCam(screenPos.x, screenPos.y, 1.0f); float3 ray = camera.screenToCam(screenPos.x, screenPos.y, 1.0f);
...@@ -159,7 +160,7 @@ __device__ float depthMatching(const Parameters &params, float d1, float d2) { ...@@ -159,7 +160,7 @@ __device__ float depthMatching(const Parameters &params, float d1, float d2) {
/* /*
* Full reprojection with normals and depth * Full reprojection with normals and depth
*/ */
template <typename A, typename B, ViewPortMode VPMODE, AccumulationFunction ACCUM> template <typename A, typename B, AccumulationFunction ACCUM, Projection PROJECT>
__global__ void reprojection_kernel( __global__ void reprojection_kernel(
TextureObject<A> in, // Attribute input TextureObject<A> in, // Attribute input
TextureObject<float> depth_src, TextureObject<float> depth_src,
...@@ -176,10 +177,11 @@ __global__ void reprojection_kernel( ...@@ -176,10 +177,11 @@ __global__ void reprojection_kernel(
const float d = depth_in.tex2D((int)x, (int)y); const float d = depth_in.tex2D((int)x, (int)y);
if (d > params.camera.minDepth && d < params.camera.maxDepth) { if (d > params.camera.minDepth && d < params.camera.maxDepth) {
const float2 rpt = convertScreen<VPMODE>(params, x, y); //const float2 rpt = convertScreen<VPMODE>(params, x, y);
const float3 camPos = transform * params.camera.screenToCam(rpt.x, rpt.y, d); //const float3 camPos = transform * params.camera.screenToCam(rpt.x, rpt.y, d);
const float3 camPos = transform * params.camera.unproject<PROJECT>(make_float3(x, y, d));
if (camPos.z > camera.minDepth && camPos.z < camera.maxDepth) { if (camPos.z > camera.minDepth && camPos.z < camera.maxDepth) {
const float2 screenPos = camera.camToScreen<float2>(camPos); const float3 screenPos = camera.project<Projection::PERSPECTIVE>(camPos);
// Not on screen so stop now... // Not on screen so stop now...
if (screenPos.x < depth_src.width() && screenPos.y < depth_src.height()) { if (screenPos.x < depth_src.width() && screenPos.y < depth_src.height()) {
...@@ -220,7 +222,7 @@ __global__ void reprojection_kernel( ...@@ -220,7 +222,7 @@ __global__ void reprojection_kernel(
/* /*
* Full reprojection without normals * Full reprojection without normals
*/ */
template <typename A, typename B, ViewPortMode VPMODE, AccumulationFunction ACCUM> template <typename A, typename B, AccumulationFunction ACCUM, Projection PROJECT>
__global__ void reprojection_kernel( __global__ void reprojection_kernel(
TextureObject<A> in, // Attribute input TextureObject<A> in, // Attribute input
TextureObject<float> depth_src, TextureObject<float> depth_src,
...@@ -236,10 +238,10 @@ __global__ void reprojection_kernel( ...@@ -236,10 +238,10 @@ __global__ void reprojection_kernel(
const float d = depth_in.tex2D((int)x, (int)y); const float d = depth_in.tex2D((int)x, (int)y);
if (d > params.camera.minDepth && d < params.camera.maxDepth) { if (d > params.camera.minDepth && d < params.camera.maxDepth) {
const float2 rpt = convertScreen<VPMODE>(params, x, y); //const float2 rpt = convertScreen<VPMODE>(params, x, y);
const float3 camPos = transform * params.camera.screenToCam(rpt.x, rpt.y, d); const float3 camPos = transform * params.camera.unproject<PROJECT>(make_float3(x, y, d));
if (camPos.z > camera.minDepth && camPos.z < camera.maxDepth) { if (camPos.z > camera.minDepth && camPos.z < camera.maxDepth) {
const float2 screenPos = camera.camToScreen<float2>(camPos); const float3 screenPos = camera.project<Projection::PERSPECTIVE>(camPos);
// Not on screen so stop now... // Not on screen so stop now...
if (screenPos.x < depth_src.width() && screenPos.y < depth_src.height()) { if (screenPos.x < depth_src.width() && screenPos.y < depth_src.height()) {
...@@ -248,6 +250,7 @@ __global__ void reprojection_kernel( ...@@ -248,6 +250,7 @@ __global__ void reprojection_kernel(
// Boolean match (0 or 1 weight). 1.0 if depths are sufficiently close // Boolean match (0 or 1 weight). 1.0 if depths are sufficiently close
float weight = depthMatching(params, camPos.z, d2); float weight = depthMatching(params, camPos.z, d2);
if (params.m_flags & ftl::render::kUseWeightsChannel)
weight *= float(weights.tex2D(int(screenPos.x+0.5f), int(screenPos.y+0.5f))) / 32767.0f; weight *= float(weights.tex2D(int(screenPos.x+0.5f), int(screenPos.y+0.5f))) / 32767.0f;
const B output = make<B>(input); // * weight; //weightInput(input, weight); const B output = make<B>(input); // * weight; //weightInput(input, weight);
...@@ -278,66 +281,76 @@ void ftl::cuda::reproject( ...@@ -278,66 +281,76 @@ void ftl::cuda::reproject(
if (normals) { if (normals) {
if (params.accumulationMode == AccumulationFunction::CloseWeights) { if (params.accumulationMode == AccumulationFunction::CloseWeights) {
switch (params.viewPortMode) { switch (params.projection) {
case ViewPortMode::Disabled: reprojection_kernel<A,B,ViewPortMode::Disabled,AccumulationFunction::CloseWeights><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, *normals, out, contrib, params, camera, transform, transformR); break; case Projection::PERSPECTIVE: reprojection_kernel<A,B,AccumulationFunction::CloseWeights, Projection::PERSPECTIVE><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, *normals, out, contrib, params, camera, transform, transformR); break;
case ViewPortMode::Clipping: reprojection_kernel<A,B,ViewPortMode::Clipping,AccumulationFunction::CloseWeights><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, *normals, out, contrib, params, camera, transform, transformR); break; case Projection::ORTHOGRAPHIC: reprojection_kernel<A,B,AccumulationFunction::CloseWeights, Projection::ORTHOGRAPHIC><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, *normals, out, contrib, params, camera, transform, transformR); break;
case ViewPortMode::Stretch: reprojection_kernel<A,B,ViewPortMode::Stretch,AccumulationFunction::CloseWeights><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, *normals, out, contrib, params, camera, transform, transformR); break; case Projection::EQUIRECTANGULAR: reprojection_kernel<A,B,AccumulationFunction::CloseWeights, Projection::EQUIRECTANGULAR><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, *normals, out, contrib, params, camera, transform, transformR); break;
//case ViewPortMode::Stretch: reprojection_kernel<A,B,ViewPortMode::Stretch,AccumulationFunction::CloseWeights><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, *normals, out, contrib, params, camera, transform, transformR); break;
} }
} else if (params.accumulationMode == AccumulationFunction::BestWeight) { } else if (params.accumulationMode == AccumulationFunction::BestWeight) {
switch (params.viewPortMode) { switch (params.projection) {
case ViewPortMode::Disabled: reprojection_kernel<A,B,ViewPortMode::Disabled,AccumulationFunction::BestWeight><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, *normals, out, contrib, params, camera, transform, transformR); break; case Projection::PERSPECTIVE: reprojection_kernel<A,B,AccumulationFunction::BestWeight, Projection::PERSPECTIVE><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, *normals, out, contrib, params, camera, transform, transformR); break;
case ViewPortMode::Clipping: reprojection_kernel<A,B,ViewPortMode::Clipping,AccumulationFunction::BestWeight><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, *normals, out, contrib, params, camera, transform, transformR); break; case Projection::ORTHOGRAPHIC: reprojection_kernel<A,B,AccumulationFunction::BestWeight, Projection::ORTHOGRAPHIC><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, *normals, out, contrib, params, camera, transform, transformR); break;
case ViewPortMode::Stretch: reprojection_kernel<A,B,ViewPortMode::Stretch,AccumulationFunction::BestWeight><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, *normals, out, contrib, params, camera, transform, transformR); break; case Projection::EQUIRECTANGULAR: reprojection_kernel<A,B,AccumulationFunction::BestWeight, Projection::EQUIRECTANGULAR><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, *normals, out, contrib, params, camera, transform, transformR); break;
//case ViewPortMode::Stretch: reprojection_kernel<A,B,ViewPortMode::Stretch,AccumulationFunction::BestWeight><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, *normals, out, contrib, params, camera, transform, transformR); break;
} }
} else if (params.accumulationMode == AccumulationFunction::Simple) { } else if (params.accumulationMode == AccumulationFunction::Simple) {
switch (params.viewPortMode) { switch (params.projection) {
case ViewPortMode::Disabled: reprojection_kernel<A,B,ViewPortMode::Disabled,AccumulationFunction::Simple><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, *normals, out, contrib, params, camera, transform, transformR); break; case Projection::PERSPECTIVE: reprojection_kernel<A,B,AccumulationFunction::Simple, Projection::PERSPECTIVE><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, *normals, out, contrib, params, camera, transform, transformR); break;
case ViewPortMode::Clipping: reprojection_kernel<A,B,ViewPortMode::Clipping,AccumulationFunction::Simple><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, *normals, out, contrib, params, camera, transform, transformR); break; case Projection::ORTHOGRAPHIC: reprojection_kernel<A,B,AccumulationFunction::Simple, Projection::ORTHOGRAPHIC><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, *normals, out, contrib, params, camera, transform, transformR); break;
case ViewPortMode::Stretch: reprojection_kernel<A,B,ViewPortMode::Stretch,AccumulationFunction::Simple><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, *normals, out, contrib, params, camera, transform, transformR); break; case Projection::EQUIRECTANGULAR: reprojection_kernel<A,B,AccumulationFunction::Simple, Projection::EQUIRECTANGULAR><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, *normals, out, contrib, params, camera, transform, transformR); break;
//case ViewPortMode::Stretch: reprojection_kernel<A,B,ViewPortMode::Stretch,AccumulationFunction::Simple><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, *normals, out, contrib, params, camera, transform, transformR); break;
} }
} else if (params.accumulationMode == AccumulationFunction::ColourDiscard) { } else if (params.accumulationMode == AccumulationFunction::ColourDiscard) {
switch (params.viewPortMode) { switch (params.projection) {
case ViewPortMode::Disabled: reprojection_kernel<A,B,ViewPortMode::Disabled,AccumulationFunction::ColourDiscard><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, *normals, out, contrib, params, camera, transform, transformR); break; case Projection::PERSPECTIVE: reprojection_kernel<A,B,AccumulationFunction::ColourDiscard, Projection::PERSPECTIVE><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, *normals, out, contrib, params, camera, transform, transformR); break;
case ViewPortMode::Clipping: reprojection_kernel<A,B,ViewPortMode::Clipping,AccumulationFunction::ColourDiscard><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, *normals, out, contrib, params, camera, transform, transformR); break; case Projection::ORTHOGRAPHIC: reprojection_kernel<A,B,AccumulationFunction::ColourDiscard, Projection::ORTHOGRAPHIC><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, *normals, out, contrib, params, camera, transform, transformR); break;
case ViewPortMode::Stretch: reprojection_kernel<A,B,ViewPortMode::Stretch,AccumulationFunction::ColourDiscard><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, *normals, out, contrib, params, camera, transform, transformR); break; case Projection::EQUIRECTANGULAR: reprojection_kernel<A,B,AccumulationFunction::ColourDiscard, Projection::EQUIRECTANGULAR><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, *normals, out, contrib, params, camera, transform, transformR); break;
//case ViewPortMode::Stretch: reprojection_kernel<A,B,ViewPortMode::Stretch,AccumulationFunction::ColourDiscard><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, *normals, out, contrib, params, camera, transform, transformR); break;
} }
} else if (params.accumulationMode == AccumulationFunction::ColourDiscardSmooth) { } else if (params.accumulationMode == AccumulationFunction::ColourDiscardSmooth) {
switch (params.viewPortMode) { switch (params.projection) {
case ViewPortMode::Disabled: reprojection_kernel<A,B,ViewPortMode::Disabled,AccumulationFunction::ColourDiscardSmooth><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, *normals, out, contrib, params, camera, transform, transformR); break; case Projection::PERSPECTIVE: reprojection_kernel<A,B,AccumulationFunction::ColourDiscardSmooth, Projection::PERSPECTIVE><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, *normals, out, contrib, params, camera, transform, transformR); break;
case ViewPortMode::Clipping: reprojection_kernel<A,B,ViewPortMode::Clipping,AccumulationFunction::ColourDiscardSmooth><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, *normals, out, contrib, params, camera, transform, transformR); break; case Projection::ORTHOGRAPHIC: reprojection_kernel<A,B,AccumulationFunction::ColourDiscardSmooth, Projection::ORTHOGRAPHIC><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, *normals, out, contrib, params, camera, transform, transformR); break;
case ViewPortMode::Stretch: reprojection_kernel<A,B,ViewPortMode::Stretch,AccumulationFunction::ColourDiscardSmooth><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, *normals, out, contrib, params, camera, transform, transformR); break; case Projection::EQUIRECTANGULAR: reprojection_kernel<A,B,AccumulationFunction::ColourDiscardSmooth, Projection::EQUIRECTANGULAR><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, *normals, out, contrib, params, camera, transform, transformR); break;
//case ViewPortMode::Stretch: reprojection_kernel<A,B,ViewPortMode::Stretch,AccumulationFunction::ColourDiscardSmooth><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, *normals, out, contrib, params, camera, transform, transformR); break;
} }
} }
} else { } else {
if (params.accumulationMode == AccumulationFunction::CloseWeights) { if (params.accumulationMode == AccumulationFunction::CloseWeights) {
switch (params.viewPortMode) { switch (params.projection) {
case ViewPortMode::Disabled: reprojection_kernel<A,B,ViewPortMode::Disabled,AccumulationFunction::CloseWeights><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, out, contrib, params, camera, transform, transformR); break; case Projection::PERSPECTIVE: reprojection_kernel<A,B,AccumulationFunction::CloseWeights, Projection::PERSPECTIVE><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, out, contrib, params, camera, transform, transformR); break;
case ViewPortMode::Clipping: reprojection_kernel<A,B,ViewPortMode::Clipping,AccumulationFunction::CloseWeights><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, out, contrib, params, camera, transform, transformR); break; case Projection::ORTHOGRAPHIC: reprojection_kernel<A,B,AccumulationFunction::CloseWeights, Projection::ORTHOGRAPHIC><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, out, contrib, params, camera, transform, transformR); break;
case ViewPortMode::Stretch: reprojection_kernel<A,B,ViewPortMode::Stretch,AccumulationFunction::CloseWeights><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, out, contrib, params, camera, transform, transformR); break; case Projection::EQUIRECTANGULAR: reprojection_kernel<A,B,AccumulationFunction::CloseWeights, Projection::EQUIRECTANGULAR><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, out, contrib, params, camera, transform, transformR); break;
//case ViewPortMode::Stretch: reprojection_kernel<A,B,ViewPortMode::Stretch,AccumulationFunction::CloseWeights><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, out, contrib, params, camera, transform, transformR); break;
} }
} else if (params.accumulationMode == AccumulationFunction::BestWeight) { } else if (params.accumulationMode == AccumulationFunction::BestWeight) {
switch (params.viewPortMode) { switch (params.projection) {
case ViewPortMode::Disabled: reprojection_kernel<A,B,ViewPortMode::Disabled,AccumulationFunction::BestWeight><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, out, contrib, params, camera, transform, transformR); break; case Projection::PERSPECTIVE: reprojection_kernel<A,B,AccumulationFunction::BestWeight, Projection::PERSPECTIVE><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, out, contrib, params, camera, transform, transformR); break;
case ViewPortMode::Clipping: reprojection_kernel<A,B,ViewPortMode::Clipping,AccumulationFunction::BestWeight><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, out, contrib, params, camera, transform, transformR); break; case Projection::ORTHOGRAPHIC: reprojection_kernel<A,B,AccumulationFunction::BestWeight, Projection::ORTHOGRAPHIC><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, out, contrib, params, camera, transform, transformR); break;
case ViewPortMode::Stretch: reprojection_kernel<A,B,ViewPortMode::Stretch,AccumulationFunction::BestWeight><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, out, contrib, params, camera, transform, transformR); break; case Projection::EQUIRECTANGULAR: reprojection_kernel<A,B,AccumulationFunction::BestWeight, Projection::EQUIRECTANGULAR><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, out, contrib, params, camera, transform, transformR); break;
//case ViewPortMode::Stretch: reprojection_kernel<A,B,ViewPortMode::Stretch,AccumulationFunction::BestWeight><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, out, contrib, params, camera, transform, transformR); break;
} }
} else if (params.accumulationMode == AccumulationFunction::Simple) { } else if (params.accumulationMode == AccumulationFunction::Simple) {
switch (params.viewPortMode) { switch (params.projection) {
case ViewPortMode::Disabled: reprojection_kernel<A,B,ViewPortMode::Disabled,AccumulationFunction::Simple><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, out, contrib, params, camera, transform, transformR); break; case Projection::PERSPECTIVE: reprojection_kernel<A,B,AccumulationFunction::Simple, Projection::PERSPECTIVE><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, out, contrib, params, camera, transform, transformR); break;
case ViewPortMode::Clipping: reprojection_kernel<A,B,ViewPortMode::Clipping,AccumulationFunction::Simple><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, out, contrib, params, camera, transform, transformR); break; case Projection::ORTHOGRAPHIC: reprojection_kernel<A,B,AccumulationFunction::Simple, Projection::ORTHOGRAPHIC><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, out, contrib, params, camera, transform, transformR); break;
case ViewPortMode::Stretch: reprojection_kernel<A,B,ViewPortMode::Stretch,AccumulationFunction::Simple><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, out, contrib, params, camera, transform, transformR); break; case Projection::EQUIRECTANGULAR: reprojection_kernel<A,B,AccumulationFunction::Simple, Projection::EQUIRECTANGULAR><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, out, contrib, params, camera, transform, transformR); break;
//case ViewPortMode::Stretch: reprojection_kernel<A,B,ViewPortMode::Stretch,AccumulationFunction::Simple><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, out, contrib, params, camera, transform, transformR); break;
} }
} else if (params.accumulationMode == AccumulationFunction::ColourDiscard) { } else if (params.accumulationMode == AccumulationFunction::ColourDiscard) {
switch (params.viewPortMode) { switch (params.projection) {
case ViewPortMode::Disabled: reprojection_kernel<A,B,ViewPortMode::Disabled,AccumulationFunction::ColourDiscard><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, out, contrib, params, camera, transform, transformR); break; case Projection::PERSPECTIVE: reprojection_kernel<A,B,AccumulationFunction::ColourDiscard, Projection::PERSPECTIVE><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, out, contrib, params, camera, transform, transformR); break;
case ViewPortMode::Clipping: reprojection_kernel<A,B,ViewPortMode::Clipping,AccumulationFunction::ColourDiscard><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, out, contrib, params, camera, transform, transformR); break; case Projection::ORTHOGRAPHIC: reprojection_kernel<A,B,AccumulationFunction::ColourDiscard, Projection::ORTHOGRAPHIC><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, out, contrib, params, camera, transform, transformR); break;
case ViewPortMode::Stretch: reprojection_kernel<A,B,ViewPortMode::Stretch,AccumulationFunction::ColourDiscard><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, out, contrib, params, camera, transform, transformR); break; case Projection::EQUIRECTANGULAR: reprojection_kernel<A,B,AccumulationFunction::ColourDiscard, Projection::EQUIRECTANGULAR><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, out, contrib, params, camera, transform, transformR); break;
//case ViewPortMode::Stretch: reprojection_kernel<A,B,ViewPortMode::Stretch,AccumulationFunction::ColourDiscard><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, out, contrib, params, camera, transform, transformR); break;
} }
} else if (params.accumulationMode == AccumulationFunction::ColourDiscardSmooth) { } else if (params.accumulationMode == AccumulationFunction::ColourDiscardSmooth) {
switch (params.viewPortMode) { switch (params.projection) {
case ViewPortMode::Disabled: reprojection_kernel<A,B,ViewPortMode::Disabled,AccumulationFunction::ColourDiscardSmooth><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, out, contrib, params, camera, transform, transformR); break; case Projection::PERSPECTIVE: reprojection_kernel<A,B,AccumulationFunction::ColourDiscardSmooth, Projection::PERSPECTIVE><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, out, contrib, params, camera, transform, transformR); break;
case ViewPortMode::Clipping: reprojection_kernel<A,B,ViewPortMode::Clipping,AccumulationFunction::ColourDiscardSmooth><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, out, contrib, params, camera, transform, transformR); break; case Projection::ORTHOGRAPHIC: reprojection_kernel<A,B,AccumulationFunction::ColourDiscardSmooth, Projection::ORTHOGRAPHIC><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, out, contrib, params, camera, transform, transformR); break;
case ViewPortMode::Stretch: reprojection_kernel<A,B,ViewPortMode::Stretch,AccumulationFunction::ColourDiscardSmooth><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, out, contrib, params, camera, transform, transformR); break; case Projection::EQUIRECTANGULAR: reprojection_kernel<A,B,AccumulationFunction::ColourDiscardSmooth, Projection::EQUIRECTANGULAR><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, out, contrib, params, camera, transform, transformR); break;
//case ViewPortMode::Stretch: reprojection_kernel<A,B,ViewPortMode::Stretch,AccumulationFunction::ColourDiscardSmooth><<<gridSize, blockSize, 0, stream>>>(in, depth_src, depth_in, weights, out, contrib, params, camera, transform, transformR); break;
} }
} }
} }
...@@ -405,8 +418,9 @@ __global__ void reprojection_kernel( ...@@ -405,8 +418,9 @@ __global__ void reprojection_kernel(
const float d = depth_in.tex2D((int)x, (int)y); const float d = depth_in.tex2D((int)x, (int)y);
if (d > params.camera.minDepth && d < params.camera.maxDepth) { if (d > params.camera.minDepth && d < params.camera.maxDepth) {
const float3 camPos = poseInv * params.camera.screenToCam(x, y, d); //const float3 camPos = poseInv * params.camera.screenToCam(x, y, d);
const float2 screenPos = camera.camToScreen<float2>(camPos); const float3 camPos = poseInv * params.camera.unproject<Projection::PERSPECTIVE>(make_float3(x, y, d));
const float3 screenPos = camera.project<Projection::PERSPECTIVE>(camPos);
if (screenPos.x < in.width() && screenPos.y < in.height()) { if (screenPos.x < in.width() && screenPos.y < in.height()) {
const auto input = in.tex2D(screenPos.x, screenPos.y); const auto input = in.tex2D(screenPos.x, screenPos.y);
......
components/renderers/cpp/src/screen.cu
View file @ 70063e8d
...@@ -7,6 +7,7 @@ using ftl::rgbd::Camera; ...@@ -7,6 +7,7 @@ using ftl::rgbd::Camera;
using ftl::cuda::TextureObject; using ftl::cuda::TextureObject;
using ftl::render::Parameters; using ftl::render::Parameters;
using ftl::render::ViewPortMode; using ftl::render::ViewPortMode;
using ftl::rgbd::Projection;
#define T_PER_BLOCK 8 #define T_PER_BLOCK 8
...@@ -45,7 +46,7 @@ __device__ inline uint2 convertToScreen<ViewPortMode::Warping>(const Parameters ...@@ -45,7 +46,7 @@ __device__ inline uint2 convertToScreen<ViewPortMode::Warping>(const Parameters
/* /*
* Convert source screen position to output screen coordinates. * Convert source screen position to output screen coordinates.
*/ */
template <ftl::render::ViewPortMode VPMODE> template <ftl::render::ViewPortMode VPMODE, Projection PROJECT>
__global__ void screen_coord_kernel(TextureObject<float> depth, __global__ void screen_coord_kernel(TextureObject<float> depth,
TextureObject<float> depth_out, TextureObject<float> depth_out,
TextureObject<short2> screen_out, Parameters params, float4x4 pose, Camera camera) { TextureObject<short2> screen_out, Parameters params, float4x4 pose, Camera camera) {
...@@ -53,22 +54,23 @@ __device__ inline uint2 convertToScreen<ViewPortMode::Warping>(const Parameters ...@@ -53,22 +54,23 @@ __device__ inline uint2 convertToScreen<ViewPortMode::Warping>(const Parameters
const int y = blockIdx.y*blockDim.y + threadIdx.y; const int y = blockIdx.y*blockDim.y + threadIdx.y;
if (x >= 0 && y >= 0 && x < depth.width() && y < depth.height()) { if (x >= 0 && y >= 0 && x < depth.width() && y < depth.height()) {
uint2 screenPos = make_uint2(30000,30000); //uint2 screenPos = make_uint2(30000,30000);
const float d = depth.tex2D(x, y); const float d = depth.tex2D(x, y);
// Find the virtual screen position of current point // Find the virtual screen position of current point
const float3 camPos = (d > camera.minDepth && d < camera.maxDepth) ? pose * camera.screenToCam(x,y,d) : make_float3(0.0f,0.0f,0.0f); const float3 camPos = (d > camera.minDepth && d < camera.maxDepth) ? pose * camera.screenToCam(x,y,d) : make_float3(0.0f,0.0f,0.0f);
screenPos = convertToScreen<VPMODE>(params, camPos); float3 screenPos = params.camera.project<PROJECT>(camPos); //convertToScreen<VPMODE>(params, camPos);
if ( camPos.z < params.camera.minDepth || if ( screenPos.z < params.camera.minDepth ||
camPos.z > params.camera.maxDepth || screenPos.z > params.camera.maxDepth ||
//!vp.inside(screenPos.x, screenPos.y)) //!vp.inside(screenPos.x, screenPos.y))
screenPos.x < 0.0f || screenPos.y < 0.0f ||
screenPos.x >= params.camera.width || screenPos.x >= params.camera.width ||
screenPos.y >= params.camera.height) screenPos.y >= params.camera.height)
screenPos = make_uint2(30000,30000); screenPos = make_float3(30000,30000,0);
screen_out(x,y) = make_short2(screenPos.x, screenPos.y); screen_out(x,y) = make_short2(screenPos.x, screenPos.y);
depth_out(x,y) = camPos.z; depth_out(x,y) = screenPos.z;
} }
} }
...@@ -78,10 +80,24 @@ void ftl::cuda::screen_coord(TextureObject<float> &depth, TextureObject<float> & ...@@ -78,10 +80,24 @@ void ftl::cuda::screen_coord(TextureObject<float> &depth, TextureObject<float> &
const dim3 gridSize((depth.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (depth.height() + T_PER_BLOCK - 1)/T_PER_BLOCK); const dim3 gridSize((depth.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (depth.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK); const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
if (params.projection == Projection::PERSPECTIVE) {
switch (params.viewPortMode) {
case ViewPortMode::Disabled: screen_coord_kernel<ViewPortMode::Disabled, Projection::PERSPECTIVE><<<gridSize, blockSize, 0, stream>>>(depth, depth_out, screen_out, params, pose, camera); break;
case ViewPortMode::Clipping: screen_coord_kernel<ViewPortMode::Clipping, Projection::PERSPECTIVE><<<gridSize, blockSize, 0, stream>>>(depth, depth_out, screen_out, params, pose, camera); break;
case ViewPortMode::Stretch: screen_coord_kernel<ViewPortMode::Stretch, Projection::PERSPECTIVE><<<gridSize, blockSize, 0, stream>>>(depth, depth_out, screen_out, params, pose, camera); break;
}
} else if (params.projection == Projection::EQUIRECTANGULAR) {
switch (params.viewPortMode) {
case ViewPortMode::Disabled: screen_coord_kernel<ViewPortMode::Disabled, Projection::EQUIRECTANGULAR><<<gridSize, blockSize, 0, stream>>>(depth, depth_out, screen_out, params, pose, camera); break;
case ViewPortMode::Clipping: screen_coord_kernel<ViewPortMode::Clipping, Projection::EQUIRECTANGULAR><<<gridSize, blockSize, 0, stream>>>(depth, depth_out, screen_out, params, pose, camera); break;
case ViewPortMode::Stretch: screen_coord_kernel<ViewPortMode::Stretch, Projection::EQUIRECTANGULAR><<<gridSize, blockSize, 0, stream>>>(depth, depth_out, screen_out, params, pose, camera); break;
}
} else if (params.projection == Projection::ORTHOGRAPHIC) {
switch (params.viewPortMode) { switch (params.viewPortMode) {
case ViewPortMode::Disabled: screen_coord_kernel<ViewPortMode::Disabled><<<gridSize, blockSize, 0, stream>>>(depth, depth_out, screen_out, params, pose, camera); break; case ViewPortMode::Disabled: screen_coord_kernel<ViewPortMode::Disabled, Projection::ORTHOGRAPHIC><<<gridSize, blockSize, 0, stream>>>(depth, depth_out, screen_out, params, pose, camera); break;
case ViewPortMode::Clipping: screen_coord_kernel<ViewPortMode::Clipping><<<gridSize, blockSize, 0, stream>>>(depth, depth_out, screen_out, params, pose, camera); break; case ViewPortMode::Clipping: screen_coord_kernel<ViewPortMode::Clipping, Projection::ORTHOGRAPHIC><<<gridSize, blockSize, 0, stream>>>(depth, depth_out, screen_out, params, pose, camera); break;
case ViewPortMode::Stretch: screen_coord_kernel<ViewPortMode::Stretch><<<gridSize, blockSize, 0, stream>>>(depth, depth_out, screen_out, params, pose, camera); break; case ViewPortMode::Stretch: screen_coord_kernel<ViewPortMode::Stretch, Projection::ORTHOGRAPHIC><<<gridSize, blockSize, 0, stream>>>(depth, depth_out, screen_out, params, pose, camera); break;
}
} }
cudaSafeCall( cudaGetLastError() ); cudaSafeCall( cudaGetLastError() );
} }
......
components/renderers/cpp/src/splatter_cuda.hpp
View file @ 70063e8d
...@@ -165,11 +165,6 @@ namespace cuda { ...@@ -165,11 +165,6 @@ namespace cuda {
uchar4 bad_colour, uchar4 bad_colour,
cudaStream_t stream); cudaStream_t stream);
void show_mask(
ftl::cuda::TextureObject<uchar4> &colour,
ftl::cuda::TextureObject<uint8_t> &mask,
int id, uchar4 style, cudaStream_t stream);
void merge_convert_depth( void merge_convert_depth(
ftl::cuda::TextureObject<int> &d1, ftl::cuda::TextureObject<int> &d1,
ftl::cuda::TextureObject<float> &d2, ftl::cuda::TextureObject<float> &d2,
......
components/renderers/cpp/src/touch.cu 0 → 100644
View file @ 70063e8d
#include <ftl/cuda/touch.hpp>
#include <ftl/cuda/warp.hpp>
using ftl::cuda::TextureObject;
using ftl::cuda::warpSum;
__device__ inline ftl::cuda::Collision pack_collision(int cx, int cy, int num, float cd) {
return ftl::cuda::Collision{(num << 24) | (cx << 12) | (cy), cd};
}
__global__ void touch_kernel(TextureObject<float> depth_in, TextureObject<float> depth_out, ftl::cuda::Collision *collisions, int max_collisions, float dist) {
const int x = blockIdx.x*blockDim.x + threadIdx.x;
const int y = blockIdx.y*blockDim.y + threadIdx.y;
bool collision = false;
float cd = 0.0f;
if (x >= 0 && y >= 0 && x < depth_in.width() && y < depth_in.height()) {
//uint2 screenPos = make_uint2(30000,30000);
const float din = depth_in.tex2D(x, y);
const float dout = depth_out.tex2D(x, y);
collision = (din < 1000.0f && fabsf(din-dout) < dist);
cd = fminf(din,dout);
depth_out(x,y) = cd;
}
int num_collisions = __popc(__ballot_sync(0xFFFFFFFF, collision));
float cx = warpSum((collision) ? float(x) : 0.0f) / float(num_collisions);
float cy = warpSum((collision) ? float(y) : 0.0f) / float(num_collisions);
cd = warpSum((collision) ? float(cd) : 0.0f) / float(num_collisions);
if ((threadIdx.x+threadIdx.y*blockDim.x) % 32 == 0) {
if (num_collisions > 0) {
//printf("Collision: %f,%f [%d]\n", cx, cy, num_collisions);
int ix = atomicInc(&collisions[0].screen, max_collisions-1);
collisions[ix+1] = pack_collision(cx, cy, num_collisions, cd);
}
}
}
#define T_PER_BLOCK 8
void ftl::cuda::touch_merge(TextureObject<float> &depth_in, TextureObject<float> &depth_out, ftl::cuda::Collision *collisions, int max_collisions, float dist, cudaStream_t stream) {
const dim3 gridSize((depth_in.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (depth_in.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
touch_kernel<<<gridSize, blockSize, 0, stream>>>(depth_in, depth_out, collisions, max_collisions, dist);
cudaSafeCall( cudaGetLastError() );
}
components/renderers/cpp/test/CMakeLists.txt 0 → 100644
View file @ 70063e8d
### Renderer Unit ##############################################################
add_executable(render_unit
$<TARGET_OBJECTS:CatchTest>
./render_unit.cpp
)
target_include_directories(render_unit PUBLIC "${CMAKE_CURRENT_SOURCE_DIR}/../include")
target_link_libraries(render_unit
ftlcommon ftlcodecs ftldata ftlrgbd)
target_precompile_headers(render_unit REUSE_FROM ftlcommon)
add_test(RenderUnitTest render_unit)
\ No newline at end of file
components/renderers/cpp/test/render_unit.cpp 0 → 100644
View file @ 70063e8d
#include "catch.hpp"
#include <ftl/data/new_frameset.hpp>
#include <ftl/data/framepool.hpp>
#include <ftl/render/CUDARender.hpp>
#include <nlohmann/json.hpp>
using ftl::data::Frame;
using ftl::data::FrameSet;
using ftl::config::json_t;
using ftl::codecs::Channel;
TEST_CASE("Renderer Single Frame", "") {
json_t global = json_t{{"$id","ftl://test"}};
auto *root = ftl::config::configure(global);
ftl::data::Pool pool(5,7);
Frame f = pool.allocate(ftl::data::FrameID(0,0), 1000);
f.store();
auto fsptr = FrameSet::fromFrame(f);
auto renderer = std::unique_ptr<ftl::render::CUDARender>(
ftl::create<ftl::render::CUDARender>(root, "renderer")
);
Frame out = pool.allocate(ftl::data::FrameID(1,0), 1000);
out.store();
ftl::rgbd::Frame &rgbdframe = out.cast<ftl::rgbd::Frame>();
auto &calib = rgbdframe.setLeft();
calib.width = 640;
calib.height = 480;
calib.fx = 700;
calib.fy = 700;
calib.cx = -250;
calib.cy = -200;
calib.minDepth = 0.1f;
calib.maxDepth = 10.0f;
rgbdframe.setPose() = Eigen::Matrix4d::Identity();
int width = rgbdframe.getLeft().width;
int height = rgbdframe.getLeft().height;
auto &colour = rgbdframe.create<cv::cuda::GpuMat>(Channel::Colour);
colour.create(height, width, CV_8UC4);
rgbdframe.create<cv::cuda::GpuMat>(Channel::Depth).create(height, width, CV_32F);
rgbdframe.createTexture<float>(Channel::Depth);
SECTION("copes with single frame missing colour") {
for (int i=0; i<20; ++i) {
renderer->begin(out.cast<ftl::rgbd::Frame>(), Channel::Colour);
Eigen::Matrix4d pose;
pose.setIdentity();
renderer->submit(fsptr.get(), ftl::codecs::Channels<0>(Channel::Colour), pose);
renderer->render();
renderer->end();
}
}
/*SECTION("single colour empty mat") {
fsptr->frames[0].create<cv::cuda::GpuMat>(Channel::Colour);
fsptr->frames[0].cast<ftl::rgbd::Frame>().setLeft() = calib;
fsptr->frames[0].cast<ftl::rgbd::Frame>().setPose() = Eigen::Matrix4d::Identity();
for (int i=0; i<20; ++i) {
renderer->begin(out.cast<ftl::rgbd::Frame>(), Channel::Colour);
Eigen::Matrix4d pose;
pose.setIdentity();
renderer->submit(fsptr.get(), ftl::codecs::Channels<0>(Channel::Colour), pose);
renderer->render();
renderer->end();
}
}*/
SECTION("single colour only frame") {
fsptr->frames[0].create<cv::cuda::GpuMat>(Channel::Colour).create(640,480,CV_8UC4);
fsptr->frames[0].cast<ftl::rgbd::Frame>().setLeft() = calib;
fsptr->frames[0].cast<ftl::rgbd::Frame>().setPose() = Eigen::Matrix4d::Identity();
for (int i=0; i<20; ++i) {
renderer->begin(out.cast<ftl::rgbd::Frame>(), Channel::Colour);
Eigen::Matrix4d pose;
pose.setIdentity();
renderer->submit(fsptr.get(), ftl::codecs::Channels<0>(Channel::Colour), pose);
renderer->render();
renderer->end();
}
}
SECTION("single full only frame") {
fsptr->frames[0].create<cv::cuda::GpuMat>(Channel::Colour).create(640,480,CV_8UC4);
fsptr->frames[0].cast<ftl::rgbd::Frame>().setLeft() = calib;
fsptr->frames[0].cast<ftl::rgbd::Frame>().setPose() = Eigen::Matrix4d::Identity();
auto &depth = fsptr->frames[0].create<cv::cuda::GpuMat>(Channel::Colour);
depth.create(640,480,CV_8UC4);
depth.setTo(cv::Scalar(5.0f));
for (int i=0; i<20; ++i) {
renderer->begin(out.cast<ftl::rgbd::Frame>(), Channel::Colour);
Eigen::Matrix4d pose;
pose.setIdentity();
renderer->submit(fsptr.get(), ftl::codecs::Channels<0>(Channel::Colour), pose);
renderer->render();
renderer->end();
}
}
SECTION("single frame empty depth") {
fsptr->frames[0].create<cv::cuda::GpuMat>(Channel::Colour).create(640,480,CV_8UC4);
fsptr->frames[0].cast<ftl::rgbd::Frame>().setLeft() = calib;
fsptr->frames[0].cast<ftl::rgbd::Frame>().setPose() = Eigen::Matrix4d::Identity();
auto &depth = fsptr->frames[0].create<cv::cuda::GpuMat>(Channel::Colour);
//depth.create(640,480,CV_8UC4);
//depth.setTo(cv::Scalar(5.0f));
for (int i=0; i<20; ++i) {
renderer->begin(out.cast<ftl::rgbd::Frame>(), Channel::Colour);
Eigen::Matrix4d pose;
pose.setIdentity();
renderer->submit(fsptr.get(), ftl::codecs::Channels<0>(Channel::Colour), pose);
renderer->render();
renderer->end();
}
}
}
components/rgbd-sources/CMakeLists.txt
View file @ 70063e8d
set(RGBDSRC set(RGBDSRC
src/sources/stereovideo/calibrate.cpp src/sources/stereovideo/rectification.cpp
src/sources/stereovideo/local.cpp src/sources/stereovideo/opencv.cpp
src/source.cpp src/source.cpp
src/frame.cpp src/frame.cpp
src/frameset.cpp #src/frameset.cpp
src/sources/stereovideo/stereovideo.cpp src/sources/stereovideo/stereovideo.cpp
#src/colour.cpp #src/colour.cpp
src/group.cpp #src/group.cpp
src/cb_segmentation.cpp #src/cb_segmentation.cpp
#src/abr.cpp #src/abr.cpp
src/sources/screencapture/screencapture.cpp src/sources/screencapture/screencapture.cpp
src/camera.cpp src/camera.cpp
#src/init.cpp
) )
if (HAVE_REALSENSE) if (HAVE_REALSENSE)
list(APPEND RGBDSRC "src/sources/realsense/realsense_source.cpp") list(APPEND RGBDSRC "src/sources/realsense/realsense_source.cpp")
endif() endif()
if (LibArchive_FOUND) if (HAVE_PYLON)
list(APPEND RGBDSRC list(APPEND RGBDSRC "src/sources/stereovideo/pylon.cpp")
src/sources/snapshot/snapshot.cpp endif()
src/sources/snapshot/snapshot_source.cpp
)
endif (LibArchive_FOUND)
add_library(ftlrgbd ${RGBDSRC}) add_library(ftlrgbd ${RGBDSRC})
...@@ -38,7 +36,9 @@ if (CUDA_FOUND) ...@@ -38,7 +36,9 @@ if (CUDA_FOUND)
set_property(TARGET ftlrgbd PROPERTY CUDA_SEPARABLE_COMPILATION OFF) set_property(TARGET ftlrgbd PROPERTY CUDA_SEPARABLE_COMPILATION OFF)
endif() endif()
target_link_libraries(ftlrgbd ftlcalibration ftlcommon ${OpenCV_LIBS} ${LIBSGM_LIBRARIES} ${CUDA_LIBRARIES} Eigen3::Eigen realsense ftlnet ${LibArchive_LIBRARIES} ftlcodecs ftloperators ftldata ${X11_X11_LIB} ${X11_Xext_LIB}) target_link_libraries(ftlrgbd ftlcalibration ftlcommon ${OpenCV_LIBS} ${CUDA_LIBRARIES} Eigen3::Eigen realsense ftlnet ${LibArchive_LIBRARIES} ftlcodecs ftloperators ftldata ${X11_X11_LIB} ${X11_Xext_LIB} ${X11_Xtst_LIB} ${X11_XTest_LIB} Pylon)
target_precompile_headers(ftlrgbd REUSE_FROM ftldata)
if (BUILD_TESTS) if (BUILD_TESTS)
add_subdirectory(test) add_subdirectory(test)
......
components/rgbd-sources/include/ftl/cb_segmentation.hpp deleted 100644 → 0
View file @ e2e259ab
#pragma once
#include <opencv2/core.hpp>
namespace ftl {
/**
* @brief Codebook segmentation and depthmap filling.
* @param Input image width
* @param Input image height
*
* Codebook segmentation based on
*
* Kim, K., Chalidabhongse, T. H., Harwood, D., & Davis, L. (2005).
* Real-time foreground-background segmentation using codebook model.
* Real-Time Imaging. https://doi.org/10.1016/j.rti.2004.12.004
*
* and fixed size codebook optimization in
*
* Rodriguez-Gomez, R., Fernandez-Sanchez, E. J., Diaz, J., & Ros, E.
* (2015). Codebook hardware implementation on FPGA for background
* subtraction. Journal of Real-Time Image Processing.
* https://doi.org/10.1007/s11554-012-0249-6
*
* Additional modifications to include depth maps as part of the
* background model.
*/
class CBSegmentation {
public:
CBSegmentation(char codebook_size, size_t width, size_t height, float alpha, float beta, float epsilon, float sigma, int T_add, int T_del, int T_h);
/**
* @brief Segment image.
* @param Input image (3-channels)
* @param Output Mat. Background pixels set to 0, foreground pixels > 0.
*
* @todo Template method on OpenCV type
*/
void apply(cv::Mat &in, cv::Mat &out, cv::Mat &depth, bool fill=false);
void apply(cv::Mat &in, cv::Mat &out);
protected:
class Pixel {
public:
int idx;
float r;
float g;
float b;
float i;
int d;
long t;
Pixel(const int &index, const uchar *bgr, const int &depth, const long &time);
};
class Codeword {
public:
float r;
float g;
float b;
float i_min, i_max;
long f, lambda, p, q;
float d_m;
float d_f;
float d_S;
void set(CBSegmentation::Pixel &pixel);
void update(CBSegmentation::Pixel &pixel);
bool colordiff(CBSegmentation::Pixel &pixel, float epsilon);
bool brightness(CBSegmentation::Pixel &pixel, float alpha, float beta);
bool depthdiff(CBSegmentation::Pixel &pixel, float sigma);
inline int freq() { return f; }
inline long getLambda() { return lambda; }
inline long ctime() { return p; }
inline long atime() { return q; }
};
enum EntryType { H, M };
union Entry {
char size;
struct Data {
EntryType type;
CBSegmentation::Codeword cw;
} data ;
};
struct CompareEntry{
bool operator()(const Entry &a,const Entry &b) const{
return !((a.data.type == M && b.data.type == H) ||
(a.data.cw.f < b.data.cw.f));
}
};
bool processPixel(Pixel &px, Codeword *codeword=nullptr);
size_t size_;
size_t width_;
size_t height_;
float alpha_;
float beta_;
float epsilon_;
float sigma_;
int T_add_;
int T_del_;
int T_h_;
private:
long t_ = 1;
std::vector<Entry> cb_;
};
}
\ No newline at end of file