From c6eb5cb214c1260a974038b7ceccfc49f9748f1a Mon Sep 17 00:00:00 2001
From: Nicolas Pope <nicolas.pope@utu.fi>
Date: Thu, 12 Dec 2019 15:19:36 +0200
Subject: [PATCH] High resolution colour
---
.../reconstruct/src/ilw/discontinuity.cu | 30 +++--
applications/reconstruct/src/ilw/ilw_cuda.hpp | 54 ++++----
applications/reconstruct/src/main.cpp | 5 +-
.../reconstruct/src/reconstruction.cpp | 26 ++++
.../reconstruct/src/reconstruction.hpp | 3 +
.../codecs/include/ftl/codecs/bitrates.hpp | 8 +-
.../codecs/include/ftl/codecs/channels.hpp | 123 +++++++++---------
.../codecs/include/ftl/codecs/encoder.hpp | 20 +--
.../include/ftl/codecs/nvpipe_encoder.hpp | 22 ++--
components/codecs/src/bitrates.cpp | 44 ++++---
components/codecs/src/encoder.cpp | 13 +-
components/codecs/src/opencv_decoder.cpp | 3 +-
components/codecs/src/opencv_encoder.cpp | 13 +-
components/codecs/test/nvpipe_codec_unit.cpp | 14 +-
components/codecs/test/opencv_codec_unit.cpp | 12 +-
components/operators/src/colours.cpp | 13 +-
components/operators/src/mask.cpp | 4 +-
components/operators/src/mask.cu | 31 +++--
components/operators/src/mask_cuda.hpp | 8 +-
components/renderers/cpp/src/reprojection.cu | 11 +-
components/renderers/cpp/src/tri_render.cpp | 10 +-
components/rgbd-sources/src/source.cpp | 21 ---
.../rgbd-sources/src/sources/net/net.cpp | 17 ++-
.../src/sources/stereovideo/calibrate.cpp | 20 +++
.../src/sources/stereovideo/calibrate.hpp | 10 +-
.../src/sources/stereovideo/stereovideo.cpp | 108 +++++++++------
.../src/sources/stereovideo/stereovideo.hpp | 12 +-
components/rgbd-sources/src/streamer.cpp | 2 +-
28 files changed, 386 insertions(+), 271 deletions(-)
diff --git a/applications/reconstruct/src/ilw/discontinuity.cu b/applications/reconstruct/src/ilw/discontinuity.cu
index fe78d4715..fcadde03e 100644
--- a/applications/reconstruct/src/ilw/discontinuity.cu
+++ b/applications/reconstruct/src/ilw/discontinuity.cu
@@ -5,11 +5,12 @@
using ftl::cuda::Mask;
template <int RADIUS>
-__global__ void discontinuity_kernel(ftl::cuda::TextureObject<int> mask_out, ftl::cuda::TextureObject<float> depth, ftl::rgbd::Camera params) {
+__global__ void discontinuity_kernel(ftl::cuda::TextureObject<int> mask_out, ftl::cuda::TextureObject<float> depth,
+ const cv::Size size, const double minDepth, const double maxDepth) {
const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
- if (x < params.width && y < params.height) {
+ if (x < size.width && y < size.height) {
Mask mask(0);
const float d = depth.tex2D((int)x, (int)y);
@@ -17,7 +18,7 @@ __global__ void discontinuity_kernel(ftl::cuda::TextureObject<int> mask_out, ftl
// Calculate depth between 0.0 and 1.0
//float p = (d - params.minDepth) / (params.maxDepth - params.minDepth);
- if (d >= params.minDepth && d <= params.maxDepth) {
+ if (d >= minDepth && d <= maxDepth) {
/* Orts-Escolano S. et al. 2016. Holoportation: Virtual 3D teleportation in real-time. */
// Is there a discontinuity nearby?
for (int u=-RADIUS; u<=RADIUS; ++u) {
@@ -26,22 +27,25 @@ __global__ void discontinuity_kernel(ftl::cuda::TextureObject<int> mask_out, ftl
if (fabs(depth.tex2D((int)x+u, (int)y+v) - d) > 0.1f) mask.isDiscontinuity(true);
}
}
- }
-
- mask_out(x,y) = (int)mask;
+ }
+
+ mask_out(x,y) = (int)mask;
}
}
-void ftl::cuda::discontinuity(ftl::cuda::TextureObject<int> &mask_out, ftl::cuda::TextureObject<float> &depth, const ftl::rgbd::Camera ¶ms, uint discon, cudaStream_t stream) {
- const dim3 gridSize((params.width + T_PER_BLOCK - 1)/T_PER_BLOCK, (params.height + T_PER_BLOCK - 1)/T_PER_BLOCK);
+void ftl::cuda::discontinuity(ftl::cuda::TextureObject<int> &mask_out, ftl::cuda::TextureObject<float> &depth,
+ const cv::Size size, const double minDepth, const double maxDepth,
+ uint discon, cudaStream_t stream) {
+
+ const dim3 gridSize((size.width + T_PER_BLOCK - 1)/T_PER_BLOCK, (size.height + T_PER_BLOCK - 1)/T_PER_BLOCK);
const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
switch (discon) {
- case 5 : discontinuity_kernel<5><<<gridSize, blockSize, 0, stream>>>(mask_out, depth, params); break;
- case 4 : discontinuity_kernel<4><<<gridSize, blockSize, 0, stream>>>(mask_out, depth, params); break;
- case 3 : discontinuity_kernel<3><<<gridSize, blockSize, 0, stream>>>(mask_out, depth, params); break;
- case 2 : discontinuity_kernel<2><<<gridSize, blockSize, 0, stream>>>(mask_out, depth, params); break;
- case 1 : discontinuity_kernel<1><<<gridSize, blockSize, 0, stream>>>(mask_out, depth, params); break;
+ case 5 : discontinuity_kernel<5><<<gridSize, blockSize, 0, stream>>>(mask_out, depth, size, minDepth, maxDepth); break;
+ case 4 : discontinuity_kernel<4><<<gridSize, blockSize, 0, stream>>>(mask_out, depth, size, minDepth, maxDepth); break;
+ case 3 : discontinuity_kernel<3><<<gridSize, blockSize, 0, stream>>>(mask_out, depth, size, minDepth, maxDepth); break;
+ case 2 : discontinuity_kernel<2><<<gridSize, blockSize, 0, stream>>>(mask_out, depth, size, minDepth, maxDepth); break;
+ case 1 : discontinuity_kernel<1><<<gridSize, blockSize, 0, stream>>>(mask_out, depth, size, minDepth, maxDepth); break;
default: break;
}
cudaSafeCall( cudaGetLastError() );
diff --git a/applications/reconstruct/src/ilw/ilw_cuda.hpp b/applications/reconstruct/src/ilw/ilw_cuda.hpp
index fad97afbd..94e522347 100644
--- a/applications/reconstruct/src/ilw/ilw_cuda.hpp
+++ b/applications/reconstruct/src/ilw/ilw_cuda.hpp
@@ -10,15 +10,15 @@ namespace ftl {
namespace cuda {
struct ILWParams {
- float spatial_smooth;
- float colour_smooth;
+ float spatial_smooth;
+ float colour_smooth;
float fill_match;
float fill_threshold;
float match_threshold;
- float cost_ratio;
- float cost_threshold;
+ float cost_ratio;
+ float cost_threshold;
float range;
- uint flags;
+ uint flags;
};
static const uint kILWFlag_IgnoreBad = 0x0001;
@@ -29,7 +29,9 @@ static const uint kILWFlag_ColourConfidenceOnly = 0x0008;
void discontinuity(
ftl::cuda::TextureObject<int> &mask_out,
ftl::cuda::TextureObject<float> &depth,
- const ftl::rgbd::Camera ¶ms,
+ const cv::Size size,
+ const double minDepth,
+ const double maxDepth,
uint discon, cudaStream_t stream
);
@@ -49,32 +51,32 @@ void preprocess_depth(
);
void correspondence(
- ftl::cuda::TextureObject<float> &d1,
- ftl::cuda::TextureObject<float> &d2,
- ftl::cuda::TextureObject<uchar4> &c1,
- ftl::cuda::TextureObject<uchar4> &c2,
- ftl::cuda::TextureObject<float> &dout,
- ftl::cuda::TextureObject<float> &conf,
+ ftl::cuda::TextureObject<float> &d1,
+ ftl::cuda::TextureObject<float> &d2,
+ ftl::cuda::TextureObject<uchar4> &c1,
+ ftl::cuda::TextureObject<uchar4> &c2,
+ ftl::cuda::TextureObject<float> &dout,
+ ftl::cuda::TextureObject<float> &conf,
ftl::cuda::TextureObject<int> &mask,
- float4x4 &pose1,
- float4x4 &pose1_inv,
- float4x4 &pose2,
- const ftl::rgbd::Camera &cam1,
- const ftl::rgbd::Camera &cam2,
- const ILWParams ¶ms, int win,
- cudaStream_t stream
+ float4x4 &pose1,
+ float4x4 &pose1_inv,
+ float4x4 &pose2,
+ const ftl::rgbd::Camera &cam1,
+ const ftl::rgbd::Camera &cam2,
+ const ILWParams ¶ms, int win,
+ cudaStream_t stream
);
void move_points(
- ftl::cuda::TextureObject<float> &d_old,
- ftl::cuda::TextureObject<float> &d_new,
+ ftl::cuda::TextureObject<float> &d_old,
+ ftl::cuda::TextureObject<float> &d_new,
ftl::cuda::TextureObject<float> &conf,
- const ftl::rgbd::Camera &camera,
- const float4x4 &pose,
+ const ftl::rgbd::Camera &camera,
+ const float4x4 &pose,
const ILWParams ¶ms,
- float rate,
- int radius,
- cudaStream_t stream
+ float rate,
+ int radius,
+ cudaStream_t stream
);
}
diff --git a/applications/reconstruct/src/main.cpp b/applications/reconstruct/src/main.cpp
index 33db96ae0..dac32881b 100644
--- a/applications/reconstruct/src/main.cpp
+++ b/applications/reconstruct/src/main.cpp
@@ -82,8 +82,8 @@ static Eigen::Affine3d create_rotation_matrix(float ax, float ay, float az) {
return rz * rx * ry;
}
-// TODO: * Remove this class (requires more general solution). Also does not
-// process disconnections/reconnections/types etc. correctly.
+// TODO: * Remove this class (requires more general solution). Also does
+// not process disconnections/reconnections/types etc. correctly.
// * Update when new options become available.
class ConfigProxy {
@@ -216,6 +216,7 @@ static void run(ftl::Configurable *root) {
for (auto &input : sources) {
string uri = input->getURI();
+
auto T = transformations.find(uri);
if (T == transformations.end()) {
LOG(WARNING) << "Camera pose for " + uri + " not found in transformations";
diff --git a/applications/reconstruct/src/reconstruction.cpp b/applications/reconstruct/src/reconstruction.cpp
index be99d1581..8a868e56e 100644
--- a/applications/reconstruct/src/reconstruction.cpp
+++ b/applications/reconstruct/src/reconstruction.cpp
@@ -60,10 +60,36 @@ Reconstruction::Reconstruction(nlohmann::json &config, const std::string name) :
ftl::pool.push([this](int id) {
UNIQUE_LOCK(fs_align_.mtx, lk);
+
+ /*rgb_.resize(fs_align_.frames.size());
+ for (size_t i = 0; i < rgb_.size(); i++) {
+ auto &depth = fs_align_.frames[i].get<cv::cuda::GpuMat>(ftl::codecs::Channel::Depth);
+ auto &color = fs_align_.frames[i].get<cv::cuda::GpuMat>(ftl::codecs::Channel::Colour);
+
+ if (depth.size() != color.size()) {
+ std::swap(rgb_[i], color);
+ cv::cuda::resize(rgb_[i], color, depth.size(), 0.0, 0.0, cv::INTER_LINEAR);
+ }
+ }*/
+
pipeline_->apply(fs_align_, fs_align_, 0);
// TODO: To use second GPU, could do a download, swap, device change,
// then upload to other device. Or some direct device-2-device copy.
+ /*
+ for (size_t i = 0; i < rgb_.size(); i++) {
+ auto &depth = fs_align_.frames[i].get<cv::cuda::GpuMat>(ftl::codecs::Channel::Depth);
+ auto &color = fs_align_.frames[i].get<cv::cuda::GpuMat>(ftl::codecs::Channel::Colour);
+ auto &tmp = rgb_[i];
+
+ // TODO doesn't always work correctly if resolution changes
+ if (!tmp.empty() && (depth.size() != tmp.size())) {
+ std::swap(tmp, color);
+ fs_align_.frames[i].resetTexture(ftl::codecs::Channel::Colour);
+ fs_align_.frames[i].createTexture<uchar4>(ftl::codecs::Channel::Colour, true);
+ }
+ }*/
+
fs_align_.swapTo(fs_render_);
LOG(INFO) << "Align complete... " << fs_align_.timestamp;
diff --git a/applications/reconstruct/src/reconstruction.hpp b/applications/reconstruct/src/reconstruction.hpp
index 6546f85c1..50441bedc 100644
--- a/applications/reconstruct/src/reconstruction.hpp
+++ b/applications/reconstruct/src/reconstruction.hpp
@@ -27,11 +27,14 @@ class Reconstruction : public ftl::Configurable {
private:
bool busy_;
+
ftl::rgbd::FrameSet fs_render_;
ftl::rgbd::FrameSet fs_align_;
ftl::rgbd::Group *group_;
ftl::operators::Graph *pipeline_;
ftl::render::Triangular *renderer_;
+
+ std::vector<cv::cuda::GpuMat> rgb_;
};
}
diff --git a/components/codecs/include/ftl/codecs/bitrates.hpp b/components/codecs/include/ftl/codecs/bitrates.hpp
index d34ede8ad..fbacb4979 100644
--- a/components/codecs/include/ftl/codecs/bitrates.hpp
+++ b/components/codecs/include/ftl/codecs/bitrates.hpp
@@ -49,6 +49,8 @@ enum struct definition_t : uint8_t {
Invalid
};
+definition_t findDefinition(int width, int height);
+
/**
* Get width in pixels of definition.
*/
@@ -97,10 +99,8 @@ static const preset_t kPresetMinimum = -1;
* Represents the details of each preset codec configuration.
*/
struct CodecPreset {
- definition_t colour_res;
- definition_t depth_res;
- bitrate_t colour_qual;
- bitrate_t depth_qual;
+ definition_t res;
+ bitrate_t qual;
};
/**
diff --git a/components/codecs/include/ftl/codecs/channels.hpp b/components/codecs/include/ftl/codecs/channels.hpp
index 6673275fe..9aa214302 100644
--- a/components/codecs/include/ftl/codecs/channels.hpp
+++ b/components/codecs/include/ftl/codecs/channels.hpp
@@ -8,30 +8,31 @@ namespace ftl {
namespace codecs {
enum struct Channel : int {
- None = -1,
- Colour = 0, // 8UC3 or 8UC4
- Left = 0,
- Depth = 1, // 32S or 32F
- Right = 2, // 8UC3 or 8UC4
- Colour2 = 2,
- Disparity = 3,
- Depth2 = 3,
- Deviation = 4,
- Screen = 4,
- Normals = 5, // 32FC4
- Points = 6, // 32FC4 (should be deprecated)
- Confidence = 7, // 32F
- Contribution = 7, // 32F
- EnergyVector = 8, // 32FC4
- Flow = 9, // 32F
- Smoothing = 9, // 32F
- Energy = 10, // 32F
+ None = -1,
+ Colour = 0, // 8UC3 or 8UC4
+ Left = 0,
+ Depth = 1, // 32S or 32F
+ Right = 2, // 8UC3 or 8UC4
+ Colour2 = 2,
+ Disparity = 3,
+ Depth2 = 3,
+ Deviation = 4,
+ Screen = 4,
+ Normals = 5, // 32FC4
+ Points = 6, // 32FC4 (should be deprecated)
+ Confidence = 7, // 32F
+ Contribution = 7, // 32F
+ EnergyVector = 8, // 32FC4
+ Flow = 9, // 32F
+ Smoothing = 9, // 32F
+ Energy = 10, // 32F
Mask = 11, // 32U
Density = 12, // 32F
Support1 = 13, // 8UC4 (currently)
Support2 = 14, // 8UC4 (currently)
- Segmentation = 15, // 32S?
- ColourNormals = 16, // 8UC4
+ Segmentation = 15, // 32S?
+ ColourNormals = 16, // 8UC4
+ ColourHighRes = 20, // 8UC3 or 8UC4
AudioLeft = 32,
AudioRight = 33,
@@ -39,7 +40,7 @@ enum struct Channel : int {
Configuration = 64, // JSON Data
Calibration = 65, // Camera Parameters Object
Pose = 66, // Eigen::Matrix4d
- Index = 67,
+ Index = 67,
Data = 2048 // Custom data, any codec.
};
@@ -51,7 +52,7 @@ std::string name(Channel c);
int type(Channel c);
class Channels {
- public:
+ public:
class iterator {
public:
@@ -67,48 +68,48 @@ class Channels {
unsigned int ix_;
};
- inline Channels() { mask = 0; }
- inline explicit Channels(unsigned int m) { mask = m; }
- inline explicit Channels(Channel c) { mask = (c == Channel::None) ? 0 : 0x1 << static_cast<unsigned int>(c); }
- inline Channels &operator=(Channel c) { mask = (c == Channel::None) ? 0 : 0x1 << static_cast<unsigned int>(c); return *this; }
- inline Channels operator|(Channel c) const { return (c == Channel::None) ? Channels(mask) : Channels(mask | (0x1 << static_cast<unsigned int>(c))); }
- inline Channels operator+(Channel c) const { return (c == Channel::None) ? Channels(mask) : Channels(mask | (0x1 << static_cast<unsigned int>(c))); }
- inline Channels &operator|=(Channel c) { mask |= (c == Channel::None) ? 0 : (0x1 << static_cast<unsigned int>(c)); return *this; }
- inline Channels &operator+=(Channel c) { mask |= (c == Channel::None) ? 0 : (0x1 << static_cast<unsigned int>(c)); return *this; }
- inline Channels &operator-=(Channel c) { mask &= ~((c == Channel::None) ? 0 : (0x1 << static_cast<unsigned int>(c))); return *this; }
- inline Channels &operator+=(unsigned int c) { mask |= (0x1 << c); return *this; }
- inline Channels &operator-=(unsigned int c) { mask &= ~(0x1 << c); return *this; }
-
- inline bool has(Channel c) const {
- return (c == Channel::None) ? true : mask & (0x1 << static_cast<unsigned int>(c));
- }
-
- inline bool has(unsigned int c) const {
- return mask & (0x1 << c);
- }
+ inline Channels() { mask = 0; }
+ inline explicit Channels(unsigned int m) { mask = m; }
+ inline explicit Channels(Channel c) { mask = (c == Channel::None) ? 0 : 0x1 << static_cast<unsigned int>(c); }
+ inline Channels &operator=(Channel c) { mask = (c == Channel::None) ? 0 : 0x1 << static_cast<unsigned int>(c); return *this; }
+ inline Channels operator|(Channel c) const { return (c == Channel::None) ? Channels(mask) : Channels(mask | (0x1 << static_cast<unsigned int>(c))); }
+ inline Channels operator+(Channel c) const { return (c == Channel::None) ? Channels(mask) : Channels(mask | (0x1 << static_cast<unsigned int>(c))); }
+ inline Channels &operator|=(Channel c) { mask |= (c == Channel::None) ? 0 : (0x1 << static_cast<unsigned int>(c)); return *this; }
+ inline Channels &operator+=(Channel c) { mask |= (c == Channel::None) ? 0 : (0x1 << static_cast<unsigned int>(c)); return *this; }
+ inline Channels &operator-=(Channel c) { mask &= ~((c == Channel::None) ? 0 : (0x1 << static_cast<unsigned int>(c))); return *this; }
+ inline Channels &operator+=(unsigned int c) { mask |= (0x1 << c); return *this; }
+ inline Channels &operator-=(unsigned int c) { mask &= ~(0x1 << c); return *this; }
+
+ inline bool has(Channel c) const {
+ return (c == Channel::None) ? true : mask & (0x1 << static_cast<unsigned int>(c));
+ }
+
+ inline bool has(unsigned int c) const {
+ return mask & (0x1 << c);
+ }
inline iterator begin() { return iterator(*this, 0); }
inline iterator end() { return iterator(*this, 32); }
- inline operator unsigned int() { return mask; }
- inline operator bool() { return mask > 0; }
- inline operator Channel() {
- if (mask == 0) return Channel::None;
- int ix = 0;
- int tmask = mask;
- while (!(tmask & 0x1) && ++ix < 32) tmask >>= 1;
- return static_cast<Channel>(ix);
- }
-
- inline size_t count() { return std::bitset<32>(mask).count(); }
- inline void clear() { mask = 0; }
-
- static const size_t kMax = 32;
+ inline operator unsigned int() { return mask; }
+ inline operator bool() { return mask > 0; }
+ inline operator Channel() {
+ if (mask == 0) return Channel::None;
+ int ix = 0;
+ int tmask = mask;
+ while (!(tmask & 0x1) && ++ix < 32) tmask >>= 1;
+ return static_cast<Channel>(ix);
+ }
+
+ inline size_t count() { return std::bitset<32>(mask).count(); }
+ inline void clear() { mask = 0; }
+
+ static const size_t kMax = 32;
static Channels All();
- private:
- unsigned int mask;
+ private:
+ unsigned int mask;
};
inline Channels::iterator Channels::iterator::operator++() { Channels::iterator i = *this; while (++ix_ < 32 && !channels_.has(ix_)); return i; }
@@ -124,9 +125,9 @@ static const Channels kAllChannels(0xFFFFFFFFu);
inline bool isFloatChannel(ftl::codecs::Channel chan) {
switch (chan) {
case Channel::Depth :
- //case Channel::Normals :
+ //case Channel::Normals :
case Channel::Confidence:
- case Channel::Flow :
+ case Channel::Flow :
case Channel::Density:
case Channel::Energy : return true;
default : return false;
@@ -139,11 +140,11 @@ inline bool isFloatChannel(ftl::codecs::Channel chan) {
MSGPACK_ADD_ENUM(ftl::codecs::Channel);
inline ftl::codecs::Channels operator|(ftl::codecs::Channel a, ftl::codecs::Channel b) {
- return ftl::codecs::Channels(a) | b;
+ return ftl::codecs::Channels(a) | b;
}
inline ftl::codecs::Channels operator+(ftl::codecs::Channel a, ftl::codecs::Channel b) {
- return ftl::codecs::Channels(a) | b;
+ return ftl::codecs::Channels(a) | b;
}
#endif // _FTL_RGBD_CHANNELS_HPP_
diff --git a/components/codecs/include/ftl/codecs/encoder.hpp b/components/codecs/include/ftl/codecs/encoder.hpp
index 9c3aa8fef..ed817f7b1 100644
--- a/components/codecs/include/ftl/codecs/encoder.hpp
+++ b/components/codecs/include/ftl/codecs/encoder.hpp
@@ -46,16 +46,16 @@ void free(Encoder *&e);
* convert an OpenCV Mat or GpuMat into a compressed byte array of some form.
*/
class Encoder {
- public:
- friend Encoder *allocateEncoder(ftl::codecs::definition_t,
+ public:
+ friend Encoder *allocateEncoder(ftl::codecs::definition_t,
ftl::codecs::device_t, ftl::codecs::codec_t);
- friend void free(Encoder *&);
+ friend void free(Encoder *&);
- public:
- Encoder(ftl::codecs::definition_t maxdef,
+ public:
+ Encoder(ftl::codecs::definition_t maxdef,
ftl::codecs::definition_t mindef,
ftl::codecs::device_t dev);
- virtual ~Encoder();
+ virtual ~Encoder();
/**
* Wrapper encode to allow use of presets.
@@ -76,21 +76,21 @@ class Encoder {
* @param cb Callback containing compressed data
* @return True if succeeded with encoding.
*/
- virtual bool encode(
+ virtual bool encode(
const cv::cuda::GpuMat &in,
ftl::codecs::definition_t definition,
ftl::codecs::bitrate_t bitrate,
const std::function<void(const ftl::codecs::Packet&)> &cb)=0;
// TODO: Eventually, use GPU memory directly since some encoders can support this
- //virtual bool encode(const cv::cuda::GpuMat &in, std::vector<uint8_t> &out, bitrate_t bix, bool)=0;
+ //virtual bool encode(const cv::cuda::GpuMat &in, std::vector<uint8_t> &out, bitrate_t bix, bool)=0;
virtual void reset() {}
virtual bool supports(ftl::codecs::codec_t codec)=0;
- protected:
- bool available;
+ protected:
+ bool available;
const ftl::codecs::definition_t max_definition;
const ftl::codecs::definition_t min_definition;
const ftl::codecs::device_t device;
diff --git a/components/codecs/include/ftl/codecs/nvpipe_encoder.hpp b/components/codecs/include/ftl/codecs/nvpipe_encoder.hpp
index 5d04068c5..07c874d12 100644
--- a/components/codecs/include/ftl/codecs/nvpipe_encoder.hpp
+++ b/components/codecs/include/ftl/codecs/nvpipe_encoder.hpp
@@ -8,20 +8,20 @@ namespace ftl {
namespace codecs {
class NvPipeEncoder : public ftl::codecs::Encoder {
- public:
- NvPipeEncoder(ftl::codecs::definition_t maxdef,
+ public:
+ NvPipeEncoder(ftl::codecs::definition_t maxdef,
ftl::codecs::definition_t mindef);
- ~NvPipeEncoder();
+ ~NvPipeEncoder();
bool encode(const cv::cuda::GpuMat &in, ftl::codecs::preset_t preset,
const std::function<void(const ftl::codecs::Packet&)> &cb) {
return Encoder::encode(in, preset, cb);
}
- bool encode(const cv::cuda::GpuMat &in, ftl::codecs::definition_t definition, ftl::codecs::bitrate_t bitrate,
+ bool encode(const cv::cuda::GpuMat &in, ftl::codecs::definition_t definition, ftl::codecs::bitrate_t bitrate,
const std::function<void(const ftl::codecs::Packet&)>&) override;
- //bool encode(const cv::cuda::GpuMat &in, std::vector<uint8_t> &out, bitrate_t bix, bool);
+ //bool encode(const cv::cuda::GpuMat &in, std::vector<uint8_t> &out, bitrate_t bix, bool);
void reset();
@@ -29,18 +29,18 @@ class NvPipeEncoder : public ftl::codecs::Encoder {
static constexpr int kFlagRGB = 0x00000001;
- private:
- NvPipe *nvenc_;
- definition_t current_definition_;
- bool is_float_channel_;
+ private:
+ NvPipe *nvenc_;
+ definition_t current_definition_;
+ bool is_float_channel_;
bool was_reset_;
ftl::codecs::codec_t preference_;
cv::cuda::GpuMat tmp_;
cv::cuda::GpuMat tmp2_;
cv::cuda::Stream stream_;
- bool _encoderMatch(const cv::cuda::GpuMat &in, definition_t def);
- bool _createEncoder(const cv::cuda::GpuMat &in, definition_t def, bitrate_t rate);
+ bool _encoderMatch(const cv::cuda::GpuMat &in, definition_t def);
+ bool _createEncoder(const cv::cuda::GpuMat &in, definition_t def, bitrate_t rate);
ftl::codecs::definition_t _verifiedDefinition(ftl::codecs::definition_t def, const cv::cuda::GpuMat &in);
};
diff --git a/components/codecs/src/bitrates.cpp b/components/codecs/src/bitrates.cpp
index 45a505768..37889f5a5 100644
--- a/components/codecs/src/bitrates.cpp
+++ b/components/codecs/src/bitrates.cpp
@@ -8,21 +8,18 @@ using ftl::codecs::preset_t;
using ftl::codecs::definition_t;
using ftl::codecs::codec_t;
+
static const CodecPreset special_presets[] = {
- definition_t::HTC_VIVE, definition_t::HTC_VIVE, bitrate_t::High, bitrate_t::High
+ definition_t::HTC_VIVE, bitrate_t::High
};
static const CodecPreset presets[] = {
- definition_t::HD1080, definition_t::HD1080, bitrate_t::High, bitrate_t::High,
- definition_t::HD1080, definition_t::HD720, bitrate_t::Standard, bitrate_t::Standard,
- definition_t::HD720, definition_t::HD720, bitrate_t::High, bitrate_t::High,
- definition_t::HD720, definition_t::SD576, bitrate_t::Standard, bitrate_t::Standard,
- definition_t::SD576, definition_t::SD576, bitrate_t::High, bitrate_t::High,
- definition_t::SD576, definition_t::SD480, bitrate_t::Standard, bitrate_t::Standard,
- definition_t::SD480, definition_t::SD480, bitrate_t::High, bitrate_t::High,
- definition_t::SD480, definition_t::LD360, bitrate_t::Standard, bitrate_t::Standard,
- definition_t::LD360, definition_t::LD360, bitrate_t::Standard, bitrate_t::Standard,
- definition_t::LD360, definition_t::LD360, bitrate_t::Low, bitrate_t::Low
+ definition_t::HD1080, bitrate_t::High,
+ definition_t::HD720, bitrate_t::High,
+ definition_t::SD576, bitrate_t::High,
+ definition_t::SD480, bitrate_t::High,
+ definition_t::LD360, bitrate_t::Standard,
+ definition_t::LD360, bitrate_t::Low
};
static const float kAspectRatio = 1.777778f;
@@ -53,11 +50,27 @@ int ftl::codecs::getHeight(definition_t d) {
return resolutions[static_cast<int>(d)].height;
}
+definition_t ftl::codecs::findDefinition(int width, int height) {
+ int best = 0;
+ bool smaller = true;
+
+ for(const Resolution res : resolutions) {
+ if ((res.width == width) && (res.height == height)) {
+ return static_cast<definition_t>(best);
+ }
+ best++;
+ }
+
+ // TODO error!
+ return definition_t::Any;
+}
+
+/*
const CodecPreset &ftl::codecs::getPreset(preset_t p) {
if (p < 0 && p >= -1) return special_presets[std::abs(p+1)];
- if (p > kPresetWorst) return presets[kPresetWorst];
- if (p < kPresetBest) return presets[kPresetBest];
- return presets[p];
+ if (p > kPresetWorst) return presets[kPresetWorst];
+ if (p < kPresetBest) return presets[kPresetBest];
+ return presets[p];
}
preset_t ftl::codecs::findPreset(size_t width, size_t height) {
@@ -80,10 +93,11 @@ preset_t ftl::codecs::findPreset(size_t width, size_t height) {
for (preset_t i=kPresetMinimum; i<=kPresetWorst; ++i) {
const auto &preset = getPreset(i);
- if ((int)preset.colour_res == best_def && (int)preset.depth_res == best_def) {
+ if ((int)preset.res == best_def) {
return i;
}
}
return kPresetWorst;
}
+*/
diff --git a/components/codecs/src/encoder.cpp b/components/codecs/src/encoder.cpp
index 9a7eac72d..7c7f9a358 100644
--- a/components/codecs/src/encoder.cpp
+++ b/components/codecs/src/encoder.cpp
@@ -36,7 +36,7 @@ static MUTEX mutex;
Encoder *ftl::codecs::allocateEncoder(ftl::codecs::definition_t maxdef,
ftl::codecs::device_t dev, ftl::codecs::codec_t codec) {
- UNIQUE_LOCK(mutex, lk);
+ UNIQUE_LOCK(mutex, lk);
if (!has_been_init) init_encoders();
for (auto i=encoders.begin(); i!=encoders.end(); ++i) {
@@ -55,10 +55,10 @@ Encoder *ftl::codecs::allocateEncoder(ftl::codecs::definition_t maxdef,
}
void ftl::codecs::free(Encoder *&enc) {
- UNIQUE_LOCK(mutex, lk);
- enc->reset();
+ UNIQUE_LOCK(mutex, lk);
+ enc->reset();
enc->available = true;
- enc = nullptr;
+ enc = nullptr;
}
Encoder::Encoder(definition_t maxdef, definition_t mindef, device_t dev) :
@@ -72,9 +72,8 @@ Encoder::~Encoder() {
bool Encoder::encode(const cv::cuda::GpuMat &in, preset_t preset,
const std::function<void(const ftl::codecs::Packet&)> &cb) {
- const auto &settings = ftl::codecs::getPreset(preset);
- const definition_t definition = (in.type() == CV_32F) ? settings.depth_res : settings.colour_res;
- const bitrate_t bitrate = (in.type() == CV_32F) ? settings.depth_qual : settings.colour_qual;
+ const definition_t definition = ftl::codecs::findDefinition(in.size().width, in.size().height);
+ const bitrate_t bitrate = bitrate_t::High;
return encode(in, definition, bitrate, cb);
}
diff --git a/components/codecs/src/opencv_decoder.cpp b/components/codecs/src/opencv_decoder.cpp
index 0b9feea46..c3c5e9567 100644
--- a/components/codecs/src/opencv_decoder.cpp
+++ b/components/codecs/src/opencv_decoder.cpp
@@ -18,7 +18,7 @@ bool OpenCVDecoder::accepts(const ftl::codecs::Packet &pkt) {
}
bool OpenCVDecoder::decode(const ftl::codecs::Packet &pkt, cv::cuda::GpuMat &out) {
-
+ //CHECK(cv::Size(ftl::codecs::getWidth(pkt.definition), ftl::codecs::getHeight(pkt.definition)) == out.size());
int chunk_dim = std::sqrt(pkt.block_total);
int chunk_width = out.cols / chunk_dim;
int chunk_height = out.rows / chunk_dim;
@@ -37,7 +37,6 @@ bool OpenCVDecoder::decode(const ftl::codecs::Packet &pkt, cv::cuda::GpuMat &out
// Apply colour correction to chunk
//ftl::rgbd::colourCorrection(tmp_rgb, gamma_, temperature_);
-
// TODO:(Nick) Decode directly into double buffer if no scaling
// Can either check JPG/PNG headers or just use pkt definition.
diff --git a/components/codecs/src/opencv_encoder.cpp b/components/codecs/src/opencv_encoder.cpp
index 5dc1995a8..772922e7b 100644
--- a/components/codecs/src/opencv_encoder.cpp
+++ b/components/codecs/src/opencv_encoder.cpp
@@ -17,7 +17,7 @@ OpenCVEncoder::OpenCVEncoder(ftl::codecs::definition_t maxdef,
}
OpenCVEncoder::~OpenCVEncoder() {
-
+
}
bool OpenCVEncoder::supports(ftl::codecs::codec_t codec) {
@@ -30,9 +30,12 @@ bool OpenCVEncoder::supports(ftl::codecs::codec_t codec) {
bool OpenCVEncoder::encode(const cv::cuda::GpuMat &in, definition_t definition, bitrate_t bitrate, const std::function<void(const ftl::codecs::Packet&)> &cb) {
bool is_colour = in.type() != CV_32F;
- current_definition_ = definition;
+
+ // Ensure definition does not exceed max
+ current_definition_ = ((int)definition < (int)max_definition) ? max_definition : definition;
in.download(tmp_);
+ //CHECK(cv::Size(ftl::codecs::getWidth(definition), ftl::codecs::getHeight(definition)) == in.size());
// Scale down image to match requested definition...
if (ftl::codecs::getHeight(current_definition_) < in.rows) {
@@ -42,11 +45,12 @@ bool OpenCVEncoder::encode(const cv::cuda::GpuMat &in, definition_t definition,
}
// Represent float at 16bit int
- if (!is_colour) {
+ if (!is_colour) {
tmp_.convertTo(tmp_, CV_16UC1, 1000);
}
- chunk_dim_ = (definition == definition_t::LD360) ? 1 : 4;
+ // FIXME: Chunking is broken so forced to single chunk
+ chunk_dim_ = 1; //(definition == definition_t::LD360) ? 1 : 4;
chunk_count_ = chunk_dim_ * chunk_dim_;
jobs_ = chunk_count_;
@@ -94,6 +98,7 @@ bool OpenCVEncoder::_encodeBlock(const cv::Mat &in, ftl::codecs::Packet &pkt, bi
int cx = (pkt.block_number % chunk_dim_) * chunk_width;
int cy = (pkt.block_number / chunk_dim_) * chunk_height;
cv::Rect roi(cx,cy,chunk_width,chunk_height);
+
cv::Mat chunkHead = in(roi);
if (pkt.codec == codec_t::PNG) {
diff --git a/components/codecs/test/nvpipe_codec_unit.cpp b/components/codecs/test/nvpipe_codec_unit.cpp
index 609ce56a5..dc63131f7 100644
--- a/components/codecs/test/nvpipe_codec_unit.cpp
+++ b/components/codecs/test/nvpipe_codec_unit.cpp
@@ -22,19 +22,18 @@ namespace ftl {
}
}
+/*
TEST_CASE( "NvPipeEncoder::encode() - A colour test image at preset 0" ) {
ftl::codecs::NvPipeEncoder encoder(definition_t::HD1080, definition_t::SD480);
cv::cuda::GpuMat m(cv::Size(1920,1080), CV_8UC3, cv::Scalar(0,0,0));
int block_total = 0;
std::atomic<int> block_count = 0;
-
- const CodecPreset &preset = ftl::codecs::getPreset(ftl::codecs::kPreset0);
-
- bool r = encoder.encode(m, ftl::codecs::kPreset0, [&block_total, &block_count, preset, m](const ftl::codecs::Packet &pkt) {
+ encoder.encode()
+ bool r = encoder.encode(m, definition::H, [&block_total, &block_count, preset, m](const ftl::codecs::Packet &pkt) {
REQUIRE( pkt.codec == codec_t::HEVC );
REQUIRE( pkt.data.size() > 0 );
- REQUIRE( pkt.definition == preset.colour_res );
+ REQUIRE( pkt.definition == definition_t::HD1080 );
block_total = pkt.block_total;
block_count++;
@@ -51,12 +50,10 @@ TEST_CASE( "NvPipeEncoder::encode() - A depth test image at preset 0" ) {
int block_total = 0;
std::atomic<int> block_count = 0;
- const CodecPreset &preset = ftl::codecs::getPreset(ftl::codecs::kPreset0);
-
bool r = encoder.encode(m, ftl::codecs::kPreset0, [&block_total, &block_count, preset](const ftl::codecs::Packet &pkt) {
REQUIRE( pkt.codec == codec_t::HEVC );
REQUIRE( pkt.data.size() > 0 );
- REQUIRE( pkt.definition == preset.depth_res );
+ REQUIRE( pkt.definition == definition_t::HD1080 );
block_total = pkt.block_total;
block_count++;
@@ -65,6 +62,7 @@ TEST_CASE( "NvPipeEncoder::encode() - A depth test image at preset 0" ) {
REQUIRE( r );
REQUIRE( block_count == block_total );
}
+*/
TEST_CASE( "NvPipeDecoder::decode() - A colour test image" ) {
ftl::codecs::NvPipeEncoder encoder(definition_t::HD1080, definition_t::SD480);
diff --git a/components/codecs/test/opencv_codec_unit.cpp b/components/codecs/test/opencv_codec_unit.cpp
index 2505eeb89..961658db5 100644
--- a/components/codecs/test/opencv_codec_unit.cpp
+++ b/components/codecs/test/opencv_codec_unit.cpp
@@ -21,15 +21,17 @@ namespace ftl {
}
}
}
-
+/*
TEST_CASE( "OpenCVEncoder::encode() - A colour test image at preset 0" ) {
ftl::codecs::OpenCVEncoder encoder(definition_t::HD1080, definition_t::SD480);
- cv::cuda::GpuMat m(cv::Size(1024,576), CV_8UC3, cv::Scalar(0,0,0));
int block_total = 0;
std::atomic<int> block_count = 0;
const CodecPreset &preset = ftl::codecs::getPreset(ftl::codecs::kPreset4);
+ cv::cuda::GpuMat m(cv::Size(ftl::codecs::getWidth(preset.res),
+ ftl::codecs::getHeight(preset.res)),
+ CV_8UC3, cv::Scalar(0,0,0));
std::mutex mtx;
@@ -37,7 +39,7 @@ TEST_CASE( "OpenCVEncoder::encode() - A colour test image at preset 0" ) {
std::unique_lock<std::mutex> lk(mtx);
REQUIRE( pkt.codec == codec_t::JPG );
REQUIRE( pkt.data.size() > 0 );
- REQUIRE( pkt.definition == preset.colour_res );
+ REQUIRE( pkt.definition == preset.res );
block_total = pkt.block_total;
block_count++;
@@ -66,7 +68,7 @@ TEST_CASE( "OpenCVEncoder::encode() - A depth test image at preset 0" ) {
std::unique_lock<std::mutex> lk(mtx);
REQUIRE( pkt.codec == codec_t::PNG );
REQUIRE( pkt.data.size() > 0 );
- REQUIRE( pkt.definition == preset.depth_res );
+ REQUIRE( pkt.definition == preset.res );
block_total = pkt.block_total;
block_count++;
@@ -78,7 +80,7 @@ TEST_CASE( "OpenCVEncoder::encode() - A depth test image at preset 0" ) {
REQUIRE( r );
REQUIRE( block_count == block_total );
}
-
+*/
TEST_CASE( "OpenCVDecoder::decode() - A colour test image no resolution change" ) {
ftl::codecs::OpenCVEncoder encoder(definition_t::HD1080, definition_t::SD480);
ftl::codecs::OpenCVDecoder decoder;
diff --git a/components/operators/src/colours.cpp b/components/operators/src/colours.cpp
index 9c6fff8b8..6a49f6ede 100644
--- a/components/operators/src/colours.cpp
+++ b/components/operators/src/colours.cpp
@@ -14,11 +14,12 @@ ColourChannels::~ColourChannels() {
bool ColourChannels::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd::Source *s, cudaStream_t stream) {
auto cvstream = cv::cuda::StreamAccessor::wrapStream(stream);
+ auto &col = in.get<cv::cuda::GpuMat>(Channel::Colour);
+
// Convert colour from BGR to BGRA if needed
- if (in.get<cv::cuda::GpuMat>(Channel::Colour).type() == CV_8UC3) {
+ if (col.type() == CV_8UC3) {
//cv::cuda::Stream cvstream = cv::cuda::StreamAccessor::wrapStream(stream);
// Convert to 4 channel colour
- auto &col = in.get<cv::cuda::GpuMat>(Channel::Colour);
temp_.create(col.size(), CV_8UC4);
cv::cuda::swap(col, temp_);
cv::cuda::cvtColor(temp_,col, cv::COLOR_BGR2BGRA, 0, cvstream);
@@ -27,5 +28,13 @@ bool ColourChannels::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgb
//in.resetTexture(Channel::Colour);
in.createTexture<uchar4>(Channel::Colour, true);
+ auto &depth = in.get<cv::cuda::GpuMat>(Channel::Depth);
+ if (depth.size() != col.size()) {
+ auto &col2 = in.create<cv::cuda::GpuMat>(Channel::ColourHighRes);
+ cv::cuda::resize(col, col2, depth.size(), 0.0, 0.0, cv::INTER_LINEAR, cvstream);
+ in.createTexture<uchar4>(Channel::ColourHighRes, true);
+ in.swapChannels(Channel::Colour, Channel::ColourHighRes);
+ }
+
return true;
}
diff --git a/components/operators/src/mask.cpp b/components/operators/src/mask.cpp
index f923f11d0..c7dcbb2ac 100644
--- a/components/operators/src/mask.cpp
+++ b/components/operators/src/mask.cpp
@@ -22,7 +22,9 @@ bool DiscontinuityMask::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::
out.createTexture<int>(Channel::Mask, ftl::rgbd::Format<int>(in.get<cv::cuda::GpuMat>(Channel::Depth).size())),
in.createTexture<uchar4>(Channel::Support1),
in.createTexture<float>(Channel::Depth),
- s->parameters(), radius, threshold, stream
+ in.get<cv::cuda::GpuMat>(Channel::Depth).size(),
+ s->parameters().minDepth, s->parameters().maxDepth,
+ radius, threshold, stream
);
return true;
diff --git a/components/operators/src/mask.cu b/components/operators/src/mask.cu
index e385f41b1..91ddf19dd 100644
--- a/components/operators/src/mask.cu
+++ b/components/operators/src/mask.cu
@@ -4,16 +4,21 @@
using ftl::cuda::Mask;
-__global__ void discontinuity_kernel(ftl::cuda::TextureObject<int> mask_out, ftl::cuda::TextureObject<uchar4> support, ftl::cuda::TextureObject<float> depth, ftl::rgbd::Camera params, float threshold, int radius) {
+__global__ void discontinuity_kernel(ftl::cuda::TextureObject<int> mask_out,
+ ftl::cuda::TextureObject<uchar4> support,
+ ftl::cuda::TextureObject<float> depth,
+ const cv::Size size, const double minDepth, const double maxDepth,
+ float threshold, int radius) {
+
const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
- if (x < params.width && y < params.height) {
+ if (x < size.width && y < size.height) {
Mask mask(0);
const float d = depth.tex2D((int)x, (int)y);
- if (d >= params.minDepth && d <= params.maxDepth) {
+ if (d >= minDepth && d <= maxDepth) {
/* Orts-Escolano S. et al. 2016. Holoportation: Virtual 3D teleportation in real-time. */
// If colour cross support region terminates within the requested
@@ -37,17 +42,21 @@ __global__ void discontinuity_kernel(ftl::cuda::TextureObject<int> mask_out, ftl
float dS = depth.tex2D((int)x, (int)y + sup.w + radius);
if (fabs(dS - d) > threshold) mask.isDiscontinuity(true);
}
- }
-
- mask_out(x,y) = (int)mask;
+ }
+
+ mask_out(x,y) = (int)mask;
}
}
-void ftl::cuda::discontinuity(ftl::cuda::TextureObject<int> &mask_out, ftl::cuda::TextureObject<uchar4> &support, ftl::cuda::TextureObject<float> &depth, const ftl::rgbd::Camera ¶ms, int discon, float thresh, cudaStream_t stream) {
- const dim3 gridSize((params.width + T_PER_BLOCK - 1)/T_PER_BLOCK, (params.height + T_PER_BLOCK - 1)/T_PER_BLOCK);
+void ftl::cuda::discontinuity( ftl::cuda::TextureObject<int> &mask_out, ftl::cuda::TextureObject<uchar4> &support,
+ ftl::cuda::TextureObject<float> &depth,
+ const cv::Size size, const double minDepth, const double maxDepth,
+ int discon, float thresh, cudaStream_t stream) {
+
+ const dim3 gridSize((size.width + T_PER_BLOCK - 1)/T_PER_BLOCK, (size.height + T_PER_BLOCK - 1)/T_PER_BLOCK);
const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
- discontinuity_kernel<<<gridSize, blockSize, 0, stream>>>(mask_out, support, depth, params, thresh, discon);
+ discontinuity_kernel<<<gridSize, blockSize, 0, stream>>>(mask_out, support, depth, size, minDepth, maxDepth, thresh, discon);
cudaSafeCall( cudaGetLastError() );
#ifdef _DEBUG
@@ -55,8 +64,6 @@ void ftl::cuda::discontinuity(ftl::cuda::TextureObject<int> &mask_out, ftl::cuda
#endif
}
-
-
__global__ void cull_discontinuity_kernel(ftl::cuda::TextureObject<int> mask, ftl::cuda::TextureObject<float> depth) {
const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
@@ -71,7 +78,7 @@ void ftl::cuda::cull_discontinuity(ftl::cuda::TextureObject<int> &mask, ftl::cud
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);
- cull_discontinuity_kernel<<<gridSize, blockSize, 0, stream>>>(mask, depth);
+ cull_discontinuity_kernel<<<gridSize, blockSize, 0, stream>>>(mask, depth);
cudaSafeCall( cudaGetLastError() );
#ifdef _DEBUG
diff --git a/components/operators/src/mask_cuda.hpp b/components/operators/src/mask_cuda.hpp
index 6a02aafdb..20c266290 100644
--- a/components/operators/src/mask_cuda.hpp
+++ b/components/operators/src/mask_cuda.hpp
@@ -19,7 +19,7 @@ class Mask {
#endif
__device__ inline operator int() const { return v_; }
- __device__ inline bool is(int m) const { return v_ & m; }
+ __device__ inline bool is(int m) const { return v_ & m; }
__device__ inline bool isFilled() const { return v_ & kMask_Filled; }
__device__ inline bool isDiscontinuity() const { return v_ & kMask_Discontinuity; }
@@ -31,7 +31,7 @@ class Mask {
__device__ inline void hasCorrespondence(bool v) { v_ = (v) ? v_ | kMask_Correspondence : v_ & (~kMask_Correspondence); }
__device__ inline void isBad(bool v) { v_ = (v) ? v_ | kMask_Bad : v_ & (~kMask_Bad); }
- static constexpr int kMask_Filled = 0x0001;
+ static constexpr int kMask_Filled = 0x0001;
static constexpr int kMask_Discontinuity = 0x0002;
static constexpr int kMask_Correspondence = 0x0004;
static constexpr int kMask_Bad = 0x0008;
@@ -44,7 +44,9 @@ void discontinuity(
ftl::cuda::TextureObject<int> &mask,
ftl::cuda::TextureObject<uchar4> &support,
ftl::cuda::TextureObject<float> &depth,
- const ftl::rgbd::Camera ¶ms,
+ const cv::Size size,
+ const double minDepth,
+ const double maxDepth,
int radius, float threshold,
cudaStream_t stream);
diff --git a/components/renderers/cpp/src/reprojection.cu b/components/renderers/cpp/src/reprojection.cu
index 9c414f892..72b7cd072 100644
--- a/components/renderers/cpp/src/reprojection.cu
+++ b/components/renderers/cpp/src/reprojection.cu
@@ -94,7 +94,10 @@ __global__ void reprojection_kernel(
const float dotproduct = (max(dot(ray,n),-0.1f)+0.1) / 1.1f;
const float d2 = depth_src.tex2D(int(screenPos.x+0.5f), int(screenPos.y+0.5f));
- const auto input = in.tex2D(screenPos.x, screenPos.y); //generateInput(in.tex2D((int)screenPos.x, (int)screenPos.y), params, worldPos);
+
+ const float inSX = float(in.width()) / float(depth_src.width());
+ const float inSY = float(in.height()) / float(depth_src.height());
+ const auto input = in.tex2D(screenPos.x*inSX, screenPos.y*inSY); //generateInput(in.tex2D((int)screenPos.x, (int)screenPos.y), params, worldPos);
// TODO: Z checks need to interpolate between neighbors if large triangles are used
//float weight = ftl::cuda::weighting(fabs(camPos.z - d2), params.depthThreshold);
@@ -213,7 +216,11 @@ __global__ void reprojection_kernel(
if (screenPos.x >= depth_src.width() || screenPos.y >= depth_src.height()) return;
const float d2 = depth_src.tex2D((int)(screenPos.x+0.5f), (int)(screenPos.y+0.5f));
- const auto input = in.tex2D(screenPos.x, screenPos.y); //generateInput(in.tex2D((int)screenPos.x, (int)screenPos.y), params, worldPos);
+
+ const float inSX = float(in.width()) / float(depth_src.width());
+ const float inSY = float(in.height()) / float(depth_src.height());
+ const auto input = in.tex2D(screenPos.x*inSX, screenPos.y*inSY); //generateInput(in.tex2D((int)screenPos.x, (int)screenPos.y), params, worldPos);
+
float weight = ftl::cuda::weighting(fabs(camPos.z - d2), 0.02f);
const B weighted = make<B>(input) * weight;
diff --git a/components/renderers/cpp/src/tri_render.cpp b/components/renderers/cpp/src/tri_render.cpp
index 06d4fe262..d1d0894f4 100644
--- a/components/renderers/cpp/src/tri_render.cpp
+++ b/components/renderers/cpp/src/tri_render.cpp
@@ -220,13 +220,13 @@ void Triangular::__reprojectChannel(ftl::rgbd::Frame &output, ftl::codecs::Chann
auto &f = scene_->frames[i];
auto *s = scene_->sources[i];
- if (f.get<GpuMat>(in).type() == CV_8UC3) {
+ /*if (f.get<GpuMat>(in).type() == CV_8UC3) {
// Convert to 4 channel colour
auto &col = f.get<GpuMat>(in);
GpuMat tmp(col.size(), CV_8UC4);
cv::cuda::swap(col, tmp);
cv::cuda::cvtColor(tmp,col, cv::COLOR_BGR2BGRA);
- }
+ }*/
auto transform = MatrixConversion::toCUDA(s->getPose().cast<float>().inverse() * t.cast<float>().inverse()) * params_.m_viewMatrixInverse;
auto transformR = MatrixConversion::toCUDA(s->getPose().cast<float>().inverse()).getFloat3x3();
@@ -607,7 +607,11 @@ bool Triangular::render(ftl::rgbd::VirtualSource *src, ftl::rgbd::Frame &out, co
}
// Reprojection of colours onto surface
- _renderChannel(out, Channel::Colour, Channel::Colour, t, stream_);
+ auto main_channel = (scene_->frames[0].hasChannel(Channel::ColourHighRes)) ? Channel::ColourHighRes : Channel::Colour;
+ //if (scene_->frames[0].hasChannel(Channel::ColourHighRes)) {
+ // LOG(INFO) << "HAVE HIGH RES: " << scene_->frames[0].get<GpuMat>(Channel::ColourHighRes).rows;
+ //}
+ _renderChannel(out, main_channel, Channel::Colour, t, stream_);
if (value("cool_effect", false)) {
auto pose = params.m_viewMatrixInverse.getFloat3x3();
diff --git a/components/rgbd-sources/src/source.cpp b/components/rgbd-sources/src/source.cpp
index 13cdd5487..4a7873515 100644
--- a/components/rgbd-sources/src/source.cpp
+++ b/components/rgbd-sources/src/source.cpp
@@ -303,27 +303,6 @@ void Source::notify(int64_t ts, cv::cuda::GpuMat &c1, cv::cuda::GpuMat &c2) {
int max_width = max(impl_->params_.width, max(c1.cols, c2.cols));
int max_height = max(impl_->params_.height, max(c1.rows, c2.rows));
- // Do we need to scale camera parameters
- if (impl_->params_.width < max_width || impl_->params_.height < max_height) {
- impl_->params_ = impl_->params_.scaled(max_width, max_height);
- }
-
- // Should channel 1 be scaled?
- if (c1.cols < max_width || c1.rows < max_height) {
- LOG(WARNING) << "Resizing on GPU";
- cv::cuda::resize(c1, c1, cv::Size(max_width, max_height));
- }
-
- // Should channel 2 be scaled?
- if (!c2.empty() && (c2.cols < max_width || c2.rows < max_height)) {
- LOG(WARNING) << "Resizing on GPU";
- if (c2.type() == CV_32F) {
- cv::cuda::resize(c2, c2, cv::Size(max_width, max_height), 0.0, 0.0, cv::INTER_NEAREST);
- } else {
- cv::cuda::resize(c2, c2, cv::Size(max_width, max_height));
- }
- }
-
if (callback_) callback_(ts, c1, c2);
}
diff --git a/components/rgbd-sources/src/sources/net/net.cpp b/components/rgbd-sources/src/sources/net/net.cpp
index e4073536a..694aa50f8 100644
--- a/components/rgbd-sources/src/sources/net/net.cpp
+++ b/components/rgbd-sources/src/sources/net/net.cpp
@@ -52,8 +52,8 @@ NetFrame &NetFrameQueue::getFrame(int64_t ts, const cv::Size &s, int c1type, int
f.chunk_total[1] = 0;
f.channel_count = 0;
f.tx_size = 0;
- f.channel[0].create(s, c1type);
- f.channel[1].create(s, c2type);
+ //f.channel[0].create(s, c1type);
+ //f.channel[1].create(s, c2type);
return f;
}
oldest = (f.timestamp < oldest) ? f.timestamp : oldest;
@@ -72,8 +72,8 @@ NetFrame &NetFrameQueue::getFrame(int64_t ts, const cv::Size &s, int c1type, int
f.chunk_total[1] = 0;
f.channel_count = 0;
f.tx_size = 0;
- f.channel[0].create(s, c1type);
- f.channel[1].create(s, c2type);
+ //f.channel[0].create(s, c1type);
+ //f.channel[1].create(s, c2type);
return f;
}
}
@@ -276,12 +276,15 @@ void NetSource::_recvPacket(short ttimeoff, const ftl::codecs::StreamPacket &spk
LOG(WARNING) << "Missing calibration, skipping frame";
return;
}
-
- NetFrame &frame = queue_.getFrame(spkt.timestamp, cv::Size(params_.width, params_.height), CV_8UC3, (isFloatChannel(chan) ? CV_32FC1 : CV_8UC3));
+
+ const cv::Size size = cv::Size(ftl::codecs::getWidth(pkt.definition), ftl::codecs::getHeight(pkt.definition));
+ NetFrame &frame = queue_.getFrame(spkt.timestamp, size, CV_8UC3, (isFloatChannel(chan) ? CV_32FC1 : CV_8UC3));
// Update frame statistics
frame.tx_size += pkt.data.size();
+ frame.channel[channum].create(size, (isFloatChannel(rchan) ? CV_32FC1 : CV_8UC3));
+
// Only decode if this channel is wanted.
if (rchan == Channel::Colour || rchan == chan) {
_createDecoder(channum, pkt);
@@ -290,7 +293,7 @@ void NetSource::_recvPacket(short ttimeoff, const ftl::codecs::StreamPacket &spk
LOG(ERROR) << "No frame decoder available";
return;
}
-
+
decoder->decode(pkt, frame.channel[channum]);
} else if (chan != Channel::None && rchan != Channel::Colour) {
// Didn't receive correct second channel so just clear the images
diff --git a/components/rgbd-sources/src/sources/stereovideo/calibrate.cpp b/components/rgbd-sources/src/sources/stereovideo/calibrate.cpp
index fc99d701f..88d69ab9c 100644
--- a/components/rgbd-sources/src/sources/stereovideo/calibrate.cpp
+++ b/components/rgbd-sources/src/sources/stereovideo/calibrate.cpp
@@ -211,6 +211,26 @@ void Calibrate::_updateIntrinsics() {
map2_gpu_.second.upload(map2_.second);
}
+cv::Mat Calibrate::getCameraMatrixLeft(const cv::Size res) {
+ double scale_x = ((double) res.width) / ((double) img_size_.width);
+ double scale_y = ((double) res.height) / ((double) img_size_.height);
+ Mat scale(cv::Size(3, 3), CV_64F, 0.0);
+ scale.at<double>(0, 0) = scale_x;
+ scale.at<double>(1, 1) = scale_y;
+ scale.at<double>(2, 2) = 1.0;
+ return scale * Kl_;
+}
+
+cv::Mat Calibrate::getCameraMatrixRight(const cv::Size res) {
+ double scale_x = ((double) res.width) / ((double) img_size_.width);
+ double scale_y = ((double) res.height) / ((double) img_size_.height);
+ Mat scale(cv::Size(3, 3), CV_64F, 0.0);
+ scale.at<double>(0, 0) = scale_x;
+ scale.at<double>(1, 1) = scale_y;
+ scale.at<double>(2, 2) = 1.0;
+ return scale * Kr_;
+}
+
void Calibrate::rectifyStereo(GpuMat &l, GpuMat &r, Stream &stream) {
// cv::cuda::remap() can not use same Mat for input and output
diff --git a/components/rgbd-sources/src/sources/stereovideo/calibrate.hpp b/components/rgbd-sources/src/sources/stereovideo/calibrate.hpp
index 4561b90a7..39ec301a9 100644
--- a/components/rgbd-sources/src/sources/stereovideo/calibrate.hpp
+++ b/components/rgbd-sources/src/sources/stereovideo/calibrate.hpp
@@ -49,15 +49,11 @@ class Calibrate : public ftl::Configurable {
void updateCalibration(const ftl::rgbd::Camera &p);
- /**
- * Get the camera matrix. Used to convert disparity map back to depth and
- * a 3D point cloud.
- */
+ // Get disparity to depth matrix.
const cv::Mat &getQ() const { return Q_; }
- const cv::Mat &getCameraMatrixLeft() { return Kl_; }
- const cv::Mat &getCameraMatrixRight() { return Kr_; }
- const cv::Mat &getCameraMatrix() { return getCameraMatrixLeft(); }
+ cv::Mat getCameraMatrixLeft(const cv::Size res);
+ cv::Mat getCameraMatrixRight(const cv::Size res);
private:
void _updateIntrinsics();
diff --git a/components/rgbd-sources/src/sources/stereovideo/stereovideo.cpp b/components/rgbd-sources/src/sources/stereovideo/stereovideo.cpp
index 84d1e574b..e8eff732b 100644
--- a/components/rgbd-sources/src/sources/stereovideo/stereovideo.cpp
+++ b/components/rgbd-sources/src/sources/stereovideo/stereovideo.cpp
@@ -8,7 +8,6 @@
#include "ftl/operators/opticalflow.hpp"
#endif
-
#include "ftl/operators/smoothing.hpp"
#include "ftl/operators/colours.hpp"
#include "ftl/operators/normals.hpp"
@@ -69,14 +68,35 @@ void StereoVideoSource::init(const string &file) {
lsrc_ = ftl::create<LocalSource>(host_, "feed");
}
- cv::Size size = cv::Size(lsrc_->width(), lsrc_->height());
+ color_size_ = cv::Size(lsrc_->width(), lsrc_->height());
frames_ = std::vector<Frame>(2);
- calib_ = ftl::create<Calibrate>(host_, "calibration", size, stream_);
+ pipeline_input_ = ftl::config::create<ftl::operators::Graph>(host_, "input");
+ #ifdef HAVE_OPTFLOW
+ pipeline_input_->append<ftl::operators::NVOpticalFlow>("optflow");
+ #endif
+
+ pipeline_depth_ = ftl::config::create<ftl::operators::Graph>(host_, "disparity");
+ depth_size_ = cv::Size( pipeline_depth_->value("width", color_size_.width),
+ pipeline_depth_->value("height", color_size_.height));
+
+ pipeline_depth_->append<ftl::operators::FixstarsSGM>("algorithm");
+ #ifdef HAVE_OPTFLOW
+ pipeline_depth_->append<ftl::operators::OpticalFlowTemporalSmoothing>("optflow_filter");
+ #endif
+ pipeline_depth_->append<ftl::operators::DisparityBilateralFilter>("bilateral_filter");
+ pipeline_depth_->append<ftl::operators::DisparityToDepth>("calculate_depth");
+ pipeline_depth_->append<ftl::operators::ColourChannels>("colour"); // Convert BGR to BGRA
+ pipeline_depth_->append<ftl::operators::Normals>("normals"); // Estimate surface normals
+ pipeline_depth_->append<ftl::operators::CrossSupport>("cross");
+ pipeline_depth_->append<ftl::operators::DiscontinuityMask>("discontinuity_mask");
+ pipeline_depth_->append<ftl::operators::AggreMLS>("mls"); // Perform MLS (using smoothing channel)
+
+ calib_ = ftl::create<Calibrate>(host_, "calibration", color_size_, stream_);
if (!calib_->isCalibrated()) LOG(WARNING) << "Cameras are not calibrated!";
// Generate camera parameters from camera matrix
- cv::Mat K = calib_->getCameraMatrix();
+ cv::Mat K = calib_->getCameraMatrixLeft(depth_size_);
params_ = {
K.at<double>(0,0), // Fx
K.at<double>(1,1), // Fy
@@ -126,49 +146,34 @@ void StereoVideoSource::init(const string &file) {
mask_l_gpu.download(mask_l);
mask_l_ = (mask_l == 0);
- pipeline_input_ = ftl::config::create<ftl::operators::Graph>(host_, "input");
- #ifdef HAVE_OPTFLOW
- pipeline_input_->append<ftl::operators::NVOpticalFlow>("optflow");
- #endif
-
- pipeline_depth_ = ftl::config::create<ftl::operators::Graph>(host_, "disparity");
- pipeline_depth_->append<ftl::operators::FixstarsSGM>("algorithm");
-
- #ifdef HAVE_OPTFLOW
- pipeline_depth_->append<ftl::operators::OpticalFlowTemporalSmoothing>("optflow_filter");
- #endif
- pipeline_depth_->append<ftl::operators::DisparityBilateralFilter>("bilateral_filter");
- pipeline_depth_->append<ftl::operators::DisparityToDepth>("calculate_depth");
- pipeline_depth_->append<ftl::operators::ColourChannels>("colour"); // Convert BGR to BGRA
- pipeline_depth_->append<ftl::operators::Normals>("normals"); // Estimate surface normals
- pipeline_depth_->append<ftl::operators::CrossSupport>("cross");
- pipeline_depth_->append<ftl::operators::DiscontinuityMask>("discontinuity_mask");
- pipeline_depth_->append<ftl::operators::AggreMLS>("mls"); // Perform MLS (using smoothing channel)
-
LOG(INFO) << "StereoVideo source ready...";
ready_ = true;
}
ftl::rgbd::Camera StereoVideoSource::parameters(Channel chan) {
+ cv::Mat K;
+
if (chan == Channel::Right) {
- cv::Mat q = calib_->getCameraMatrixRight();
- ftl::rgbd::Camera params = {
- q.at<double>(0,0), // Fx
- q.at<double>(1,1), // Fy
- -q.at<double>(0,2), // Cx
- -q.at<double>(1,2), // Cy
- (unsigned int)lsrc_->width(),
- (unsigned int)lsrc_->height(),
- 0.0f, // 0m min
- 15.0f, // 15m max
- 1.0 / calib_->getQ().at<double>(3,2), // Baseline
- 0.0f // doffs
- };
- return params;
- //params_.doffs = -calib_->getQ().at<double>(3,3) * params_.baseline;
+ K = calib_->getCameraMatrixRight(depth_size_);
} else {
- return params_;
+ K = calib_->getCameraMatrixLeft(depth_size_);
}
+
+ // TODO: remove hardcoded values (min/max)
+ ftl::rgbd::Camera params = {
+ K.at<double>(0,0), // Fx
+ K.at<double>(1,1), // Fy
+ -K.at<double>(0,2), // Cx
+ -K.at<double>(1,2), // Cy
+ (unsigned int) depth_size_.width,
+ (unsigned int) depth_size_.height,
+ 0.0f, // 0m min
+ 15.0f, // 15m max
+ 1.0 / calib_->getQ().at<double>(3,2), // Baseline
+ 0.0f // doffs
+ };
+
+ return params;
}
bool StereoVideoSource::capture(int64_t ts) {
@@ -205,8 +210,31 @@ bool StereoVideoSource::compute(int n, int b) {
}
if (chan == Channel::Depth) {
- pipeline_depth_->apply(frame, frame, host_, cv::cuda::StreamAccessor::getStream(stream_));
+ // stereo algorithms assume input same size as output
+ bool resize = (depth_size_ != color_size_);
+
+ cv::cuda::GpuMat& left = frame.get<cv::cuda::GpuMat>(Channel::Left);
+ cv::cuda::GpuMat& right = frame.get<cv::cuda::GpuMat>(Channel::Right);
+
+ if (left.empty() || right.empty()) {
+ return false;
+ }
+
+ if (resize) {
+ cv::cuda::swap(fullres_left_, left);
+ cv::cuda::swap(fullres_right_, right);
+ cv::cuda::resize(fullres_left_, left, depth_size_, 0, 0, cv::INTER_CUBIC, stream_);
+ cv::cuda::resize(fullres_right_, right, depth_size_, 0, 0, cv::INTER_CUBIC, stream_);
+ }
+
+ pipeline_depth_->apply(frame, frame, host_, cv::cuda::StreamAccessor::getStream(stream_));
stream_.waitForCompletion();
+
+ if (resize) {
+ cv::cuda::swap(fullres_left_, left);
+ cv::cuda::swap(fullres_right_, right);
+ }
+
host_->notify(timestamp_,
frame.get<cv::cuda::GpuMat>(Channel::Left),
frame.get<cv::cuda::GpuMat>(Channel::Depth));
diff --git a/components/rgbd-sources/src/sources/stereovideo/stereovideo.hpp b/components/rgbd-sources/src/sources/stereovideo/stereovideo.hpp
index 78fcdbcf8..9532e6188 100644
--- a/components/rgbd-sources/src/sources/stereovideo/stereovideo.hpp
+++ b/components/rgbd-sources/src/sources/stereovideo/stereovideo.hpp
@@ -17,9 +17,7 @@ class Disparity;
/**
* RGBD source from either a stereo video file with left + right images, or
- * direct from two camera devices. A variety of algorithms are included for
- * calculating disparity, before converting to depth. Calibration of the images
- * is also performed.
+ * direct from two camera devices.
*/
class StereoVideoSource : public detail::Source {
public:
@@ -32,15 +30,21 @@ class StereoVideoSource : public detail::Source {
bool retrieve();
bool compute(int n, int b);
bool isReady();
- Camera parameters(ftl::codecs::Channel chan);
+ Camera parameters(ftl::codecs::Channel chan) override;
private:
LocalSource *lsrc_;
Calibrate *calib_;
+ cv::Size color_size_;
+ cv::Size depth_size_;
+
ftl::operators::Graph *pipeline_input_;
ftl::operators::Graph *pipeline_depth_;
+ cv::cuda::GpuMat fullres_left_;
+ cv::cuda::GpuMat fullres_right_;
+
bool ready_;
cv::cuda::Stream stream_;
diff --git a/components/rgbd-sources/src/streamer.cpp b/components/rgbd-sources/src/streamer.cpp
index e05e7faf9..96626022c 100644
--- a/components/rgbd-sources/src/streamer.cpp
+++ b/components/rgbd-sources/src/streamer.cpp
@@ -193,7 +193,7 @@ void Streamer::add(Source *src) {
if (spkt.channel == Channel::Calibration) {
// Calibration changed, so lets re-check the bitrate presets
const auto ¶ms = src->parameters();
- s->hq_bitrate = ftl::codecs::findPreset(params.width, params.height);
+ s->hq_bitrate = ftl::codecs::kPresetBest;
}
//LOG(INFO) << "RAW CALLBACK";
--
GitLab