diff --git a/SDK/C/include/ftl/ftl.h b/SDK/C/include/ftl/ftl.h
index 2820e2436ce1be05cbad2668cd44e8e9d5738838..6a9afabeb5f359e8eb6b3b1ee266f517793592fb 100644
--- a/SDK/C/include/ftl/ftl.h
+++ b/SDK/C/include/ftl/ftl.h
@@ -32,7 +32,11 @@ enum ftlError_t {
FTLERROR_STREAM_BAD_IMAGE_SIZE,
FTLERROR_STREAM_NO_INTRINSICS,
FTLERROR_STREAM_NO_DATA,
- FTLERROR_STREAM_DUPLICATE
+ FTLERROR_STREAM_BAD_SOURCE,
+ FTLERROR_STREAM_DUPLICATE,
+ FTLERROR_OPERATOR_MISSING_CHANNEL,
+ FTLERROR_OPERATOR_INVALID_FRAME,
+ FTLERROR_OPERATOR_BAD_FORMAT
};
enum ftlChannel_t {
@@ -97,6 +101,11 @@ enum ftlImageFormat_t {
FTLIMAGE_RGB_FLOAT // Used by Blender
};
+enum ftlPipeline_t {
+ FTLPIPE_DEPTH,
+ FTLPIPE_RECONSTRUCT
+};
+
#ifdef __cplusplus
extern "C" {
#endif
@@ -107,20 +116,66 @@ extern "C" {
*/
ftlError_t ftlGetLastStreamError(ftlStream_t stream);
+const char *ftlGetErrorString(ftlError_t err);
+
// ==== FTL Stream API =========================================================
/**
* Create a new file or net stream from a URI. This is for writing or sending
- * data and makes a write only stream.
+ * data and makes a write only stream. Note that a NULL uri string is valid
+ * and will result in an in-memory stream to allow for temporary framesets
+ * that might then be transformed or rendered into an actual output stream.
*/
ftlStream_t ftlCreateWriteStream(const char *uri);
/**
* Open an existing file or network stream from the URI. These streams are
- * read only.
+ * read only. The URI can also be a device URI.
*/
ftlStream_t ftlCreateReadStream(const char *uri);
+/**
+ * Set a configurable property for this stream. The `name` parameter is the
+ * path part of a configurable URI, this is prepended with the streams internal
+ * URI so that this path is relative to the stream object.
+ */
+ftlError_t ftlSetConfigString(ftlStream_t stream, const char *name, const char *value);
+
+/**
+ * Set a configurable property for this stream. The `name` parameter is the
+ * path part of a configurable URI, this is prepended with the streams internal
+ * URI so that this path is relative to the stream object. An example:
+ * `ftlSetConfigInt(stream, "sender/codec", 2);`
+ *
+ * The example makes the streams encoder use HEVC codec.
+ */
+ftlError_t ftlSetConfigInt(ftlStream_t stream, const char *name, int value);
+
+ftlError_t ftlSetConfigFloat(ftlStream_t stream, const char *name, float value);
+
+ftlError_t ftlSetConfigJSON(ftlStream_t stream, const char *name, const char *json);
+
+/**
+ * Set a configurable property for this stream. The `name` parameter is the
+ * path part of a configurable URI, this is prepended with the streams internal
+ * URI so that this path is relative to the stream object. An example:
+ * `ftlSetConfigBool(stream, "sender/lossless", false);`
+ *
+ * The example makes the streams encoder use lossy encoding.
+ */
+ftlError_t ftlSetConfigBool(ftlStream_t stream, const char *name, bool value);
+
+/**
+ * General write function to put any data into a frames channel. This could
+ * be used instead of `ftlImageWrite` or `ftlPoseWrite`, for example.
+ */
+//ftlError_t ftlWrite(ftlStream_t stream, int32_t sourceId, ftlChannel_t channel, ftlDataType_t type, int32_t size, const void *data);
+
+/**
+ * Read data from stream source channel.
+ */
+//ftlError_t ftlRead(ftlStream_t stream, int32_t sourceId, ftlChannel_t channel, ftlDataType_t type, int32_t *size, void *data);
+
/**
* Write raw image data to a frame. The width and height of
* the image contained in `data` are contained within the intrinsics data which
@@ -189,6 +244,86 @@ ftlError_t ftlIntrinsicsWriteRight(ftlStream_t stream, int32_t sourceId, int32_t
*/
ftlError_t ftlDestroyStream(ftlStream_t stream);
+// ==== Operator functions =====================================================
+
+/**
+ * Enable a pipeline to be applied before the current frameset is encoded into
+ * the stream. A frameset is only encoded and streamed after a call to
+ * `ftlNextFrame` or `ftlDestroyStream`, therefore this can be enabled any time
+ * before the first such call. Multiple pipelines can be enabled and will be
+ * applied in the order enabled. Default settings for pipelines can be
+ * changed using `ftlSetConfig...` functions with a URI of "pipes/depth/...",
+ * for example.
+ */
+ftlError_t ftlEnablePipeline(ftlStream_t stream, ftlPipeline_t pipeline);
+
+ftlError_t ftlDisablePipeline(ftlStream_t stream, ftlPipeline_t pipeline);
+
+/**
+ * Mark an additional channel as needing saving to the stream. The results of
+ * various pipelines produce channels in addition to the `ftlImageWrite`
+ * function, these need to be selected here.
+ */
+ftlError_t ftlSelect(ftlStream_t stream, ftlChannel_t channel);
+
+/**
+ * Apply the depth operator to a frameset to generate a depth channel. This
+ * requires that the frames already have a left and right channel.
+ */
+//ftlError_t ftlDepth(ftlStream_t stream);
+
+/**
+ * Remove occlusions from a depth map. The `data` passed is a
+ * right depth image with same resolution as in the intrinsics. The `channel`
+ * parameter says which channel in the frame contains the depth map to compare
+ * against, allowing ground truth channel to be used.
+ */
+ftlError_t ftlRemoveOcclusion(ftlStream_t stream, int32_t sourceId, ftlChannel_t channel, uint32_t pitch, const float *data);
+
+/**
+ * Instead of removing occluded pixels, this will add a bit to the mask channel
+ * to indicate an occlusion. It also implicitely selects the mask channel for
+ * encoding into the stream.
+ */
+ftlError_t ftlMaskOcclusion(ftlStream_t stream, int32_t sourceId, ftlChannel_t channel, uint32_t pitch, const float *data);
+
+/**
+ * Apply a discontinuity detection to the specified channel. The given channel
+ * must be a depth map. The operator adds a discontinuity flag to the mask
+ * channel.
+ */
+//ftlError_t ftlDiscontinuity(ftlStream_t stream, ftlChannel_t channel);
+
+/**
+ * Performs a frameset reconstruction on the current stream frameset. This
+ * updates and modifies depths maps. The depth channel must exist for this to
+ * work, if not then it will run `ftlDepth` automatically.
+ */
+//ftlError_t ftlReconstruction(ftlStream_t stream);
+
+// ==== Rendering ==============================================================
+
+/*ftlError_t ftlRender(
+ ftlStream_t input,
+ ftlStream_t output,
+ const float *pose,
+ ...
+ )*/
+
+// ==== Misc frame operations ==================================================
+
+/**
+ * Set config properties within a frame. For example, set the `name` property
+ * to give this source frame a label.
+ */
+ftlError_t ftlSetPropertyString(ftlStream_t stream, int32_t sourceId, const char *prop, const char *value);
+
+ftlError_t ftlSetPropertyInt(ftlStream_t stream, int32_t sourceId, const char *prop, int value);
+
+ftlError_t ftlSetPropertyFloat(ftlStream_t stream, int32_t sourceId, const char *prop, float value);
+
+// TODO: Possibly add shapes and other kinds of data to the frame?
+
#ifdef __cplusplus
};
#endif
diff --git a/SDK/C/src/streams.cpp b/SDK/C/src/streams.cpp
index 94b1fb01413dfda8fd9725bb04617c9d097f5862..01350a45944b44b970b98451aa99ce2ba2de5f4c 100644
--- a/SDK/C/src/streams.cpp
+++ b/SDK/C/src/streams.cpp
@@ -4,6 +4,10 @@
#include <ftl/streams/sender.hpp>
#include <ftl/streams/filestream.hpp>
#include <ftl/streams/netstream.hpp>
+#include <ftl/operators/cuda/disparity.hpp>
+#include <ftl/operators/operator.hpp>
+#include <ftl/operators/disparity.hpp>
+#include <ftl/operators/mvmls.hpp>
#include <opencv2/imgproc.hpp>
@@ -21,6 +25,8 @@ struct FTLStream {
int64_t interval;
bool has_fresh_data;
+ ftl::operators::Graph *pipelines;
+
std::vector<ftl::rgbd::FrameState> video_states;
ftl::rgbd::FrameSet video_fs;
};
@@ -55,6 +61,7 @@ ftlStream_t ftlCreateWriteStream(const char *uri) {
s->last_error = FTLERROR_OK;
s->stream = nullptr;
s->sender = nullptr;
+ s->pipelines = nullptr;
s->video_fs.id = 0;
s->video_fs.count = 0;
s->video_fs.mask = 0;
@@ -194,8 +201,6 @@ ftlError_t ftlIntrinsicsWriteLeft(ftlStream_t stream, int32_t sourceId, int32_t
stream->video_fs.frames[sourceId].setLeft(cam);
stream->has_fresh_data = true;
- LOG(INFO) << "INTRIN " << f << "," << cx << "," << cy << "," << baseline;
-
return FTLERROR_OK;
}
@@ -221,8 +226,6 @@ ftlError_t ftlIntrinsicsWriteRight(ftlStream_t stream, int32_t sourceId, int32_t
stream->video_fs.frames[sourceId].setRight(cam);
stream->has_fresh_data = true;
- LOG(INFO) << "INTRINR " << f << "," << cx << "," << cy << "," << baseline;
-
return FTLERROR_OK;
}
@@ -244,6 +247,100 @@ ftlError_t ftlPoseWrite(ftlStream_t stream, int32_t sourceId, const float *data)
return FTLERROR_OK;
}
+ftlError_t ftlRemoveOcclusion(ftlStream_t stream, int32_t sourceId, ftlChannel_t channel, uint32_t pitch, const float *data) {
+ if (!stream) return FTLERROR_STREAM_INVALID_STREAM;
+ if (!stream->stream) return FTLERROR_STREAM_INVALID_STREAM;
+ if (sourceId < 0 || sourceId >= 32)
+ return FTLERROR_STREAM_INVALID_PARAMETER;
+ if (!stream->video_fs.hasFrame(sourceId))
+ return FTLERROR_STREAM_NO_INTRINSICS;
+ if (static_cast<int>(channel) < 0 || static_cast<int>(channel) > 32)
+ return FTLERROR_STREAM_BAD_CHANNEL;
+ if (!stream->video_fs.frames[sourceId].hasChannel(static_cast<ftl::codecs::Channel>(channel)))
+ return FTLERROR_STREAM_BAD_CHANNEL;
+
+ auto &frame = stream->video_fs.frames[sourceId];
+ //auto &mask = frame.create<cv::cuda::GpuMat>(ftl::codecs::Channel::Mask);
+ auto &depth = frame.get<cv::cuda::GpuMat>(static_cast<ftl::codecs::Channel>(channel));
+ auto &intrin = frame.getLeft();
+
+ cv::Mat depthR(intrin.height, intrin.width, CV_32F, const_cast<float*>(data), pitch);
+ cv::Mat tmp;
+ cv::flip(depthR, tmp, 0);
+ cv::cuda::GpuMat depthRGPU;
+ depthRGPU.upload(tmp);
+ ftl::cuda::remove_occlusions(depth, depthRGPU, intrin, 0);
+
+ return FTLERROR_OK;
+}
+
+ftlError_t ftlMaskOcclusion(ftlStream_t stream, int32_t sourceId, ftlChannel_t channel, uint32_t pitch, const float *data) {
+ if (!stream) return FTLERROR_STREAM_INVALID_STREAM;
+ if (!stream->stream) return FTLERROR_STREAM_INVALID_STREAM;
+ if (sourceId < 0 || sourceId >= 32)
+ return FTLERROR_STREAM_INVALID_PARAMETER;
+ if (!stream->video_fs.hasFrame(sourceId))
+ return FTLERROR_STREAM_NO_INTRINSICS;
+ if (static_cast<int>(channel) < 0 || static_cast<int>(channel) > 32)
+ return FTLERROR_STREAM_BAD_CHANNEL;
+ if (!stream->video_fs.frames[sourceId].hasChannel(static_cast<ftl::codecs::Channel>(channel)))
+ return FTLERROR_STREAM_BAD_CHANNEL;
+
+ auto &frame = stream->video_fs.frames[sourceId];
+ auto &mask = frame.create<cv::cuda::GpuMat>(ftl::codecs::Channel::Mask);
+ auto &depth = frame.get<cv::cuda::GpuMat>(static_cast<ftl::codecs::Channel>(channel));
+ auto &intrin = frame.getLeft();
+
+ mask.create(depth.size(), CV_8UC1);
+
+ cv::Mat depthR(intrin.height, intrin.width, CV_32F, const_cast<float*>(data), pitch);
+ cv::Mat tmp;
+ cv::flip(depthR, tmp, 0);
+ cv::cuda::GpuMat depthRGPU;
+ depthRGPU.upload(tmp);
+ ftl::cuda::mask_occlusions(depth, depthRGPU, mask, intrin, 0);
+
+ ftlSelect(stream, FTLCHANNEL_Mask);
+
+ return FTLERROR_OK;
+}
+
+ftlError_t ftlEnablePipeline(ftlStream_t stream, ftlPipeline_t pipeline) {
+ if (!stream) return FTLERROR_STREAM_INVALID_STREAM;
+ if (!stream->stream) return FTLERROR_STREAM_INVALID_STREAM;
+
+ if (stream->pipelines == nullptr) {
+ stream->pipelines = ftl::create<ftl::operators::Graph>(root, "pipes");
+ }
+
+ switch (pipeline) {
+ case FTLPIPE_DEPTH : stream->pipelines->append<ftl::operators::DepthChannel>("depth");
+ ftlSelect(stream, FTLCHANNEL_Depth); // Enable saving of depth automatrically
+ break;
+ case FTLPIPE_RECONSTRUCT : stream->pipelines->append<ftl::operators::MultiViewMLS>("reconstruct"); break;
+ default: return FTLERROR_STREAM_INVALID_PARAMETER;
+ }
+
+ return FTLERROR_OK;
+}
+
+ftlError_t ftlDisablePipeline(ftlStream_t stream, ftlPipeline_t pipeline) {
+ return FTLERROR_OK;
+}
+
+ftlError_t ftlSelect(ftlStream_t stream, ftlChannel_t channel) {
+ if (!stream) return FTLERROR_STREAM_INVALID_STREAM;
+ if (!stream->stream) return FTLERROR_STREAM_INVALID_STREAM;
+ if (static_cast<int>(channel) < 0 || static_cast<int>(channel) > 32)
+ return FTLERROR_STREAM_BAD_CHANNEL;
+
+ ftl::codecs::Channels<0> channels;
+ if (stream->stream->size() > static_cast<unsigned int>(stream->video_fs.id)) channels = stream->stream->selected(stream->video_fs.id);
+ channels += static_cast<ftl::codecs::Channel>(channel);
+ stream->stream->select(stream->video_fs.id, channels, true);
+ return FTLERROR_OK;
+}
+
ftlError_t ftlSetFrameRate(ftlStream_t stream, float fps) {
if (!stream) return FTLERROR_STREAM_INVALID_STREAM;
if (!stream->stream) return FTLERROR_STREAM_INVALID_STREAM;
@@ -260,6 +357,9 @@ ftlError_t ftlNextFrame(ftlStream_t stream) {
try {
cudaSetDevice(0);
+ if (stream->pipelines) {
+ stream->pipelines->apply(stream->video_fs, stream->video_fs, 0);
+ }
stream->sender->post(stream->video_fs);
} catch (const std::exception &e) {
return FTLERROR_STREAM_ENCODE_FAILED;
@@ -292,9 +392,13 @@ ftlError_t ftlDestroyStream(ftlStream_t stream) {
if (stream->has_fresh_data) {
try {
cudaSetDevice(0);
+ if (stream->pipelines) {
+ stream->pipelines->apply(stream->video_fs, stream->video_fs, 0);
+ }
stream->sender->post(stream->video_fs);
} catch (const std::exception &e) {
err = FTLERROR_STREAM_ENCODE_FAILED;
+ LOG(ERROR) << "Sender exception: " << e.what();
}
}
@@ -302,11 +406,55 @@ ftlError_t ftlDestroyStream(ftlStream_t stream) {
err = FTLERROR_STREAM_FILE_CREATE_FAILED;
}
if (stream->sender) delete stream->sender;
+ if (stream->pipelines) delete stream->pipelines;
delete stream->stream;
stream->sender = nullptr;
stream->stream = nullptr;
+ stream->pipelines = nullptr;
delete stream;
//});
//return FTLERROR_OK;
return err;
}
+
+ftlError_t ftlSetPropertyString(ftlStream_t stream, int32_t sourceId, const char *prop, const char *value) {
+ if (!stream) return FTLERROR_STREAM_INVALID_STREAM;
+ if (!stream->stream) return FTLERROR_STREAM_INVALID_STREAM;
+ if (sourceId < 0 || sourceId >= 32)
+ return FTLERROR_STREAM_INVALID_PARAMETER;
+ if (!stream->video_fs.hasFrame(sourceId))
+ return FTLERROR_STREAM_NO_INTRINSICS;
+ if (!value) return FTLERROR_STREAM_INVALID_PARAMETER;
+ if (!prop) return FTLERROR_STREAM_INVALID_PARAMETER;
+
+ stream->video_fs.frames[sourceId].set(std::string(prop), std::string(value));
+ return FTLERROR_OK;
+}
+
+ftlError_t ftlSetPropertyInt(ftlStream_t stream, int32_t sourceId, const char *prop, int value) {
+ if (!stream) return FTLERROR_STREAM_INVALID_STREAM;
+ if (!stream->stream) return FTLERROR_STREAM_INVALID_STREAM;
+ if (sourceId < 0 || sourceId >= 32)
+ return FTLERROR_STREAM_INVALID_PARAMETER;
+ if (!stream->video_fs.hasFrame(sourceId))
+ return FTLERROR_STREAM_NO_INTRINSICS;
+ if (!value) return FTLERROR_STREAM_INVALID_PARAMETER;
+ if (!prop) return FTLERROR_STREAM_INVALID_PARAMETER;
+
+ stream->video_fs.frames[sourceId].set(std::string(prop), value);
+ return FTLERROR_OK;
+}
+
+ftlError_t ftlSetPropertyFloat(ftlStream_t stream, int32_t sourceId, const char *prop, float value) {
+ if (!stream) return FTLERROR_STREAM_INVALID_STREAM;
+ if (!stream->stream) return FTLERROR_STREAM_INVALID_STREAM;
+ if (sourceId < 0 || sourceId >= 32)
+ return FTLERROR_STREAM_INVALID_PARAMETER;
+ if (!stream->video_fs.hasFrame(sourceId))
+ return FTLERROR_STREAM_NO_INTRINSICS;
+ if (!value) return FTLERROR_STREAM_INVALID_PARAMETER;
+ if (!prop) return FTLERROR_STREAM_INVALID_PARAMETER;
+
+ stream->video_fs.frames[sourceId].set(std::string(prop), value);
+ return FTLERROR_OK;
+}
diff --git a/SDK/Python/blender_script.py b/SDK/Python/blender_script.py
index 1a567e429af02eec04e2746398ebe6c5ab8830f8..96dea4cb4616416ef123a93d6bcf46c1e21166a6 100644
--- a/SDK/Python/blender_script.py
+++ b/SDK/Python/blender_script.py
@@ -251,6 +251,22 @@ ftlImageWrite = ftl.ftlImageWrite
ftlImageWrite.restype = c_int
ftlImageWrite.argtypes = [c_void_p, c_int, c_int, c_int, c_int, c_void_p]
+ftlRemoveOcclusion = ftl.ftlRemoveOcclusion
+ftlRemoveOcclusion.restype = c_int
+ftlRemoveOcclusion.argtypes = [c_void_p, c_int, c_int, c_int, c_void_p]
+
+ftlMaskOcclusion = ftl.ftlMaskOcclusion
+ftlMaskOcclusion.restype = c_int
+ftlMaskOcclusion.argtypes = [c_void_p, c_int, c_int, c_int, c_void_p]
+
+ftlEnablePipeline = ftl.ftlEnablePipeline
+ftlEnablePipeline.restype = c_int
+ftlEnablePipeline.argtypes = [c_void_p, c_int]
+
+ftlSelect = ftl.ftlSelect
+ftlSelect.restype = c_int
+ftlSelect.argtypes = [c_void_p, c_int]
+
ftlPoseWrite = ftl.ftlPoseWrite
ftlPoseWrite.restype = c_int
ftlPoseWrite.argtypes = [c_void_p, c_int, c_void_p]
@@ -273,6 +289,8 @@ def render_and_save(filename, cameras):
print("Could not create FTL stream")
return
+ ftlCheck(ftlEnablePipeline(stream, 0))
+
for camname in cameras:
obj = bpy.context.scene.objects[camname]
if obj.type == 'CAMERA':
@@ -281,7 +299,6 @@ def render_and_save(filename, cameras):
ftlCheck(ftlIntrinsicsWriteLeft(c_void_p(stream), c_int(i), c_int(int(image.intrinsics.width)), c_int(int(image.intrinsics.height)), c_float(image.intrinsics.fx), c_float(image.intrinsics.cx), c_float(image.intrinsics.cy), c_float(image.intrinsics.baseline), c_float(image.intrinsics.min_depth), c_float(image.intrinsics.max_depth)))
ftlCheck(ftlIntrinsicsWriteRight(c_void_p(stream), c_int(i), c_int(int(image.intrinsics.width)), c_int(int(image.intrinsics.height)), c_float(image.intrinsics.fx), c_float(image.intrinsics.cx), c_float(image.intrinsics.cy), c_float(image.intrinsics.baseline), c_float(image.intrinsics.min_depth), c_float(image.intrinsics.max_depth)))
- # This line needs fixing
M = np.identity(4,dtype=np.float32)
M[0:3,0:4] = image.pose
M = np.transpose(M)
@@ -291,9 +308,10 @@ def render_and_save(filename, cameras):
ftlCheck(ftlImageWrite(stream, c_int(i), 0, 5, 0, image.imL.ctypes.data_as(c_void_p)))
ftlCheck(ftlImageWrite(stream, c_int(i), 2, 5, 0, image.imR.ctypes.data_as(c_void_p)))
ftlCheck(ftlImageWrite(stream, c_int(i), 22, 0, 0, image.depthL.ctypes.data_as(c_void_p)))
+ ftlCheck(ftlMaskOcclusion(stream, c_int(i), 22, 0, image.depthR.ctypes.data_as(c_void_p)))
i = i + 1
ftlCheck(ftlDestroyStream(stream))
-render_and_save(b'./blender.ftl', ['Camera', 'Camera.002'])
+render_and_save(b'./blender.ftl', ['Camera'])
diff --git a/components/codecs/src/channels.cpp b/components/codecs/src/channels.cpp
index 75dc784fb09a6f9882f1a0156c09314641fdaef4..1e9d8d5260e1c969d4e04325fd4084498cc5d390 100644
--- a/components/codecs/src/channels.cpp
+++ b/components/codecs/src/channels.cpp
@@ -13,23 +13,23 @@ static ChannelInfo info[] = {
"Right", CV_8UC4, // 2
"DepthRight", CV_32F, // 3
"Deviation", CV_32F, // 4
- "Normals", CV_32FC4, // 5
+ "Normals", CV_16FC4, // 5
"Weights", CV_16SC1, // 6
"Confidence", CV_32F, // 7
"EnergyVector", CV_32FC4, // 8
"Flow", CV_32F, // 9
"Energy", CV_32F, // 10
- "Mask", CV_32S, // 11
+ "Mask", CV_8U, // 11
"Density", CV_32F, // 12
"Support1", CV_8UC4, // 13
"Support2", CV_8UC4, // 14
"Segmentation", CV_32S, // 15
"ColourNormals", 0, // 16
- "ColourHighRes", 0, // 17
- "Disparity", 0, // 18
+ "ColourHighRes", CV_8UC4, // 17
+ "Disparity", CV_32F, // 18
"Smoothing", 0, // 19
- "Colour2HighRes", 0, // 20
+ "Colour2HighRes", CV_8UC4, // 20
"Overlay", 0, // 21
"GroundTruth", CV_32F, // 22
"NoName", 0,
diff --git a/components/codecs/src/generate.cpp b/components/codecs/src/generate.cpp
index f39167dfd6e1ac1e953be8813ca9725c600e47de..64dd2a042a42b6e5dbf272a01f4a922ed00168fe 100644
--- a/components/codecs/src/generate.cpp
+++ b/components/codecs/src/generate.cpp
@@ -31,6 +31,8 @@ void init_encoders() {
encoders.push_back(new ftl::codecs::OpenCVEncoder(definition_t::HD1080, definition_t::HD720));
encoders.push_back(new ftl::codecs::OpenCVEncoder(definition_t::HD1080, definition_t::HD720));
encoders.push_back(new ftl::codecs::OpenCVEncoder(definition_t::HD1080, definition_t::HD720));
+ encoders.push_back(new ftl::codecs::OpenCVEncoder(definition_t::HD1080, definition_t::HD720));
+ encoders.push_back(new ftl::codecs::OpenCVEncoder(definition_t::HD1080, definition_t::HD720));
encoders.push_back(new ftl::codecs::OpenCVEncoder(definition_t::SD576, definition_t::LD360));
encoders.push_back(new ftl::codecs::OpenCVEncoder(definition_t::SD576, definition_t::LD360));
encoders.push_back(new ftl::codecs::OpenCVEncoder(definition_t::SD576, definition_t::LD360));
diff --git a/components/codecs/src/opencv_decoder.cpp b/components/codecs/src/opencv_decoder.cpp
index be3c800d63611b940b0c1f2f4032b3681bdf6585..981bac10e6155a05fab467407483e9d47656d2c9 100644
--- a/components/codecs/src/opencv_decoder.cpp
+++ b/components/codecs/src/opencv_decoder.cpp
@@ -59,6 +59,8 @@ bool OpenCVDecoder::decode(const ftl::codecs::Packet &pkt, cv::cuda::GpuMat &out
chunkHead.upload(tmp_);
} else if (!tmp_.empty() && tmp_.type() == CV_16U && chunkHead.type() == CV_16U) {
chunkHead.upload(tmp_);
+ } else if (!tmp_.empty() && tmp_.type() == CV_8UC1 && chunkHead.type() == CV_8UC1) {
+ chunkHead.upload(tmp_);
} else {
// Silent ignore?
}
diff --git a/components/codecs/src/opencv_encoder.cpp b/components/codecs/src/opencv_encoder.cpp
index 1bb96812802eb59b0ac97bd07f9700e53596a57e..6cf2a3183ded4e3090dbcd9cb89f4b3574b4fe90 100644
--- a/components/codecs/src/opencv_encoder.cpp
+++ b/components/codecs/src/opencv_encoder.cpp
@@ -30,7 +30,7 @@ bool OpenCVEncoder::supports(ftl::codecs::codec_t codec) {
bool OpenCVEncoder::encode(const cv::cuda::GpuMat &in, ftl::codecs::Packet &pkt) {
bool is_colour = !(pkt.flags & ftl::codecs::kFlagFloat);
- if (is_colour && in.type() != CV_8UC4) return false;
+ if (is_colour && in.type() != CV_8UC4 && in.type() != CV_8UC1) return false;
if (!is_colour && in.type() == CV_8UC4) {
LOG(ERROR) << "OpenCV Encoder doesn't support lossy depth";
return false;
@@ -74,7 +74,7 @@ bool OpenCVEncoder::encode(const cv::cuda::GpuMat &in, ftl::codecs::Packet &pkt)
return false;
}
- if (pkt.codec == codec_t::Any) pkt.codec = (is_colour) ? codec_t::JPG : codec_t::PNG;
+ if (pkt.codec == codec_t::Any) pkt.codec = (is_colour && in.type() != CV_8UC1) ? codec_t::JPG : codec_t::PNG;
//for (int i=0; i<chunk_count_; ++i) {
// Add chunk job to thread pool
diff --git a/components/operators/src/disparity/cuda.hpp b/components/operators/include/ftl/operators/cuda/disparity.hpp
similarity index 65%
rename from components/operators/src/disparity/cuda.hpp
rename to components/operators/include/ftl/operators/cuda/disparity.hpp
index 74102111e31b6f44a22ae083fa8cf698723ec904..e18fcd4fe972b82706fdefecb285546e345c2778 100644
--- a/components/operators/src/disparity/cuda.hpp
+++ b/components/operators/include/ftl/operators/cuda/disparity.hpp
@@ -9,6 +9,14 @@ void disparity_to_depth(const cv::cuda::GpuMat &disparity, cv::cuda::GpuMat &dep
void depth_to_disparity(cv::cuda::GpuMat &disparity, const cv::cuda::GpuMat &depth,
const ftl::rgbd::Camera &c, cudaStream_t &stream);
+void remove_occlusions(cv::cuda::GpuMat &depth, const cv::cuda::GpuMat &depthR,
+ const ftl::rgbd::Camera &c, cudaStream_t stream);
+
+void mask_occlusions(const cv::cuda::GpuMat &depth,
+ const cv::cuda::GpuMat &depthR,
+ cv::cuda::GpuMat &mask,
+ const ftl::rgbd::Camera &c, cudaStream_t stream);
+
void optflow_filter(cv::cuda::GpuMat &disp, const cv::cuda::GpuMat &optflow,
cv::cuda::GpuMat &history, cv::cuda::GpuMat &support, int n_max, float threshold, bool fill,
diff --git a/components/operators/include/ftl/operators/gt_cuda.hpp b/components/operators/include/ftl/operators/cuda/gt.hpp
similarity index 100%
rename from components/operators/include/ftl/operators/gt_cuda.hpp
rename to components/operators/include/ftl/operators/cuda/gt.hpp
diff --git a/components/operators/include/ftl/operators/mask_cuda.hpp b/components/operators/include/ftl/operators/cuda/mask.hpp
similarity index 98%
rename from components/operators/include/ftl/operators/mask_cuda.hpp
rename to components/operators/include/ftl/operators/cuda/mask.hpp
index aadab5e5a640c7200bc87f76b4b80d496b6dbe59..8e136371879c451a9b03997998c46e7b50a8c943 100644
--- a/components/operators/include/ftl/operators/mask_cuda.hpp
+++ b/components/operators/include/ftl/operators/cuda/mask.hpp
@@ -41,6 +41,7 @@ class Mask {
static constexpr type kMask_Correspondence = 0x04;
static constexpr type kMask_Bad = 0x08;
static constexpr type kMask_Noise = 0x10;
+ static constexpr type kMask_Occlusion = 0x20;
private:
type v_;
diff --git a/components/operators/include/ftl/operators/gt_analysis.hpp b/components/operators/include/ftl/operators/gt_analysis.hpp
index 9397a96da2ff80985d08d90d7f68697935fc895f..a89230caa12e42fc05027506c6484be1bf348d91 100644
--- a/components/operators/include/ftl/operators/gt_analysis.hpp
+++ b/components/operators/include/ftl/operators/gt_analysis.hpp
@@ -3,7 +3,7 @@
#include <ftl/operators/operator.hpp>
#include <ftl/cuda_common.hpp>
-#include <ftl/operators/gt_cuda.hpp>
+#include <ftl/operators/cuda/gt.hpp>
namespace ftl {
namespace operators {
diff --git a/components/operators/src/disparity/bilateral_filter.cpp b/components/operators/src/disparity/bilateral_filter.cpp
index 694a23a9a0e9bab1b144ed9b279f4952dce022ff..0f515ed41b842df11099927542498ae861dc1df9 100644
--- a/components/operators/src/disparity/bilateral_filter.cpp
+++ b/components/operators/src/disparity/bilateral_filter.cpp
@@ -1,7 +1,7 @@
#include <ftl/operators/disparity.hpp>
#include "opencv/joint_bilateral.hpp"
-#include "cuda.hpp"
+#include <ftl/operators/cuda/disparity.hpp>
#include <opencv2/cudaimgproc.hpp>
diff --git a/components/operators/src/disparity/disp2depth.cu b/components/operators/src/disparity/disp2depth.cu
index e39afe6f60e236530985874dba9c0f598fe6a5e3..cc023c6c2b551043042ae93cc13830aec307bc98 100644
--- a/components/operators/src/disparity/disp2depth.cu
+++ b/components/operators/src/disparity/disp2depth.cu
@@ -1,6 +1,8 @@
#include <ftl/cuda_common.hpp>
#include <ftl/rgbd/camera.hpp>
#include <opencv2/core/cuda_stream_accessor.hpp>
+#include <ftl/operators/cuda/disparity.hpp>
+#include <ftl/operators/cuda/mask.hpp>
#ifndef PINF
#define PINF __int_as_float(0x7f800000)
@@ -60,3 +62,66 @@ void depth_to_disparity(cv::cuda::GpuMat &disparity, const cv::cuda::GpuMat &dep
}
}
}
+
+// =============================================================================
+
+__global__ void remove_occ_kernel(cv::cuda::PtrStepSz<float> depth, cv::cuda::PtrStepSz<float> depthR,
+ ftl::rgbd::Camera cam)
+{
+ for (STRIDE_Y(v,depth.rows)) {
+ for (STRIDE_X(u,depth.cols)) {
+ float d = depth(v,u);
+ int disparity = int((d > cam.maxDepth || d < cam.minDepth) ? 0.0f : ((cam.baseline*cam.fx) / d) + cam.doffs);
+
+ if (disparity > 0 && u-disparity > 0) {
+ float dR = depthR(v,u-disparity);
+ if (fabsf(d-dR) > 0.01f*d) {
+ depth(v,u) = 0.0f;
+ }
+ }
+ }
+ }
+}
+
+void ftl::cuda::remove_occlusions(cv::cuda::GpuMat &depth, const cv::cuda::GpuMat &depthR,
+ const ftl::rgbd::Camera &c, cudaStream_t stream) {
+ dim3 grid(1,1,1);
+ dim3 threads(128, 4, 1);
+ grid.x = cv::cuda::device::divUp(depth.cols, 128);
+ grid.y = cv::cuda::device::divUp(depth.rows, 4);
+ remove_occ_kernel<<<grid, threads, 0, stream>>>(
+ depth, depthR, c);
+ cudaSafeCall( cudaGetLastError() );
+}
+
+__global__ void mask_occ_kernel(cv::cuda::PtrStepSz<float> depth,
+ cv::cuda::PtrStepSz<float> depthR,
+ cv::cuda::PtrStepSz<uchar> mask,
+ ftl::rgbd::Camera cam)
+{
+ for (STRIDE_Y(v,depth.rows)) {
+ for (STRIDE_X(u,depth.cols)) {
+ float d = depth(v,u);
+ int disparity = int((d > cam.maxDepth || d < cam.minDepth) ? 0.0f : ((cam.baseline*cam.fx) / d) + cam.doffs);
+
+ if (disparity > 0 && u-disparity > 0) {
+ float dR = depthR(v,u-disparity);
+ if (fabsf(d-dR) > 0.01f*d) {
+ mask(v,u) = mask(v,u) | ftl::cuda::Mask::kMask_Occlusion;
+ }
+ }
+ }
+ }
+}
+
+void ftl::cuda::mask_occlusions(const cv::cuda::GpuMat &depth, const cv::cuda::GpuMat &depthR,
+ cv::cuda::GpuMat &mask,
+ const ftl::rgbd::Camera &c, cudaStream_t stream) {
+ dim3 grid(1,1,1);
+ dim3 threads(128, 4, 1);
+ grid.x = cv::cuda::device::divUp(depth.cols, 128);
+ grid.y = cv::cuda::device::divUp(depth.rows, 4);
+ mask_occ_kernel<<<grid, threads, 0, stream>>>(
+ depth, depthR, mask, c);
+ cudaSafeCall( cudaGetLastError() );
+}
diff --git a/components/operators/src/disparity/disparity_to_depth.cpp b/components/operators/src/disparity/disparity_to_depth.cpp
index d00cced27bb6e5e8fc5870e6ac5cde432f964917..5dbe97b335e4d40d06cd61f63bac37f1d2a2d215 100644
--- a/components/operators/src/disparity/disparity_to_depth.cpp
+++ b/components/operators/src/disparity/disparity_to_depth.cpp
@@ -1,5 +1,5 @@
#include "ftl/operators/disparity.hpp"
-#include "disparity/cuda.hpp"
+#include <ftl/operators/cuda/disparity.hpp>
using ftl::operators::DisparityToDepth;
using ftl::codecs::Channel;
diff --git a/components/operators/src/disparity/optflow_smoothing.cpp b/components/operators/src/disparity/optflow_smoothing.cpp
index 82476395a7ed7a5ec90b79f3aebbd1b7e9b3fb16..e2f1e6f936fec0236ce0be92ba05335e5150e8f3 100644
--- a/components/operators/src/disparity/optflow_smoothing.cpp
+++ b/components/operators/src/disparity/optflow_smoothing.cpp
@@ -1,7 +1,7 @@
#include <loguru.hpp>
#include "ftl/operators/disparity.hpp"
-#include "disparity/cuda.hpp"
+#include <ftl/operators/cuda/disparity.hpp>
#ifdef HAVE_OPTFLOW
diff --git a/components/operators/src/disparity/optflow_smoothing.cu b/components/operators/src/disparity/optflow_smoothing.cu
index 8e3293cd99c8136925acbd999d1c9be90b09e2f4..c30302b85ac218d464624e6fbb6c740b95a304ee 100644
--- a/components/operators/src/disparity/optflow_smoothing.cu
+++ b/components/operators/src/disparity/optflow_smoothing.cu
@@ -3,7 +3,7 @@
#include <opencv2/core/cuda_stream_accessor.hpp>
#include "disparity/qsort.h"
-#include "disparity/cuda.hpp"
+#include <ftl/operators/cuda/disparity.hpp>
__device__ void quicksort(float A[], size_t n)
{
diff --git a/components/operators/src/fusion/correspondence.cu b/components/operators/src/fusion/correspondence.cu
index b19a84af903af05b3908a3d2f782c16eeb6559d8..be72902428eb2e2e0d985755b8d55cf4e66391e6 100644
--- a/components/operators/src/fusion/correspondence.cu
+++ b/components/operators/src/fusion/correspondence.cu
@@ -1,6 +1,6 @@
#include "mvmls_cuda.hpp"
#include <ftl/cuda/weighting.hpp>
-#include <ftl/operators/mask_cuda.hpp>
+#include <ftl/operators/cuda/mask.hpp>
#include <ftl/cuda/warp.hpp>
using ftl::cuda::TextureObject;
diff --git a/components/operators/src/fusion/correspondence_depth.cu b/components/operators/src/fusion/correspondence_depth.cu
index 9f83fa79f33da8e9c275b13e7951f74f418cafed..f7303cfc413f03b49bdc75510906a5cbd409b626 100644
--- a/components/operators/src/fusion/correspondence_depth.cu
+++ b/components/operators/src/fusion/correspondence_depth.cu
@@ -1,6 +1,6 @@
#include "mvmls_cuda.hpp"
#include <ftl/cuda/weighting.hpp>
-#include <ftl/operators/mask_cuda.hpp>
+#include <ftl/operators/cuda/mask.hpp>
#include <ftl/cuda/warp.hpp>
using ftl::cuda::TextureObject;
diff --git a/components/operators/src/fusion/correspondence_util.cu b/components/operators/src/fusion/correspondence_util.cu
index de91eabe6b2c4f8a52d6ee4abb993ee8d0ecebcc..5145f69767866e66ed01518b6e492328f022f595 100644
--- a/components/operators/src/fusion/correspondence_util.cu
+++ b/components/operators/src/fusion/correspondence_util.cu
@@ -1,6 +1,6 @@
#include "mvmls_cuda.hpp"
#include <ftl/cuda/weighting.hpp>
-#include <ftl/operators/mask_cuda.hpp>
+#include <ftl/operators/cuda/mask.hpp>
#include <ftl/cuda/warp.hpp>
using ftl::cuda::TextureObject;
diff --git a/components/operators/src/fusion/mls_aggr.cu b/components/operators/src/fusion/mls_aggr.cu
index 9332117bc11f382f418151c5305f2d694cf3e7e0..d597654480909974c1b2605ff643fc5516a90fc1 100644
--- a/components/operators/src/fusion/mls_aggr.cu
+++ b/components/operators/src/fusion/mls_aggr.cu
@@ -1,6 +1,6 @@
#include "mvmls_cuda.hpp"
#include <ftl/cuda/weighting.hpp>
-#include <ftl/operators/mask_cuda.hpp>
+#include <ftl/operators/cuda/mask.hpp>
#include <ftl/cuda/warp.hpp>
using ftl::cuda::TextureObject;
diff --git a/components/operators/src/gt_analysis.cpp b/components/operators/src/gt_analysis.cpp
index 91d59ca3d362a0f59580773b11fa0a9d7e3ac405..d4a29a5894580bdaf6bb0839ea9b11fb058c776c 100644
--- a/components/operators/src/gt_analysis.cpp
+++ b/components/operators/src/gt_analysis.cpp
@@ -1,5 +1,5 @@
#include <ftl/operators/gt_analysis.hpp>
-#include <ftl/operators/gt_cuda.hpp>
+#include <ftl/operators/cuda/gt.hpp>
using ftl::operators::GTAnalysis;
using ftl::codecs::Channel;
diff --git a/components/operators/src/gt_analysis.cu b/components/operators/src/gt_analysis.cu
index 230f2fd20c44dce8e6db125434d6801b69183cef..53b76cd6332f6c00878be70c187d5cefdaa930a8 100644
--- a/components/operators/src/gt_analysis.cu
+++ b/components/operators/src/gt_analysis.cu
@@ -1,4 +1,4 @@
-#include <ftl/operators/gt_cuda.hpp>
+#include <ftl/operators/cuda/gt.hpp>
#ifndef WARP_SIZE
#define WARP_SIZE 32
diff --git a/components/operators/src/mask.cpp b/components/operators/src/mask.cpp
index 1c53630cf7b43112d2fe7ceba19110e9b89024b8..e39fcabea0fa322ef2c6d8333a6366aef664e0e4 100644
--- a/components/operators/src/mask.cpp
+++ b/components/operators/src/mask.cpp
@@ -1,5 +1,5 @@
#include <ftl/operators/mask.hpp>
-#include <ftl/operators/mask_cuda.hpp>
+#include <ftl/operators/cuda/mask.hpp>
using ftl::operators::DiscontinuityMask;
using ftl::operators::BorderMask;
@@ -27,6 +27,12 @@ bool DiscontinuityMask::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, cudaS
if (!in.hasChannel(Channel::Depth) || !in.hasChannel(Channel::Support1)) return false;
+ if (!out.hasChannel(Channel::Mask)) {
+ auto &m = out.create<cv::cuda::GpuMat>(Channel::Mask);
+ m.create(in.get<cv::cuda::GpuMat>(Channel::Depth).size(), CV_8UC1);
+ m.setTo(cv::Scalar(0));
+ }
+
/*ftl::cuda::discontinuity(
out.createTexture<uint8_t>(Channel::Mask, ftl::rgbd::Format<uint8_t>(in.get<cv::cuda::GpuMat>(Channel::Depth).size())),
in.createTexture<uchar4>(Channel::Support1),
@@ -37,7 +43,7 @@ bool DiscontinuityMask::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, cudaS
);*/
ftl::cuda::discontinuity(
- out.createTexture<uint8_t>(Channel::Mask, ftl::rgbd::Format<uint8_t>(in.get<cv::cuda::GpuMat>(Channel::Depth).size())),
+ out.createTexture<uint8_t>(Channel::Mask),
in.createTexture<uchar4>(Channel::Support1),
in.createTexture<float>(Channel::Depth),
in.get<cv::cuda::GpuMat>(Channel::Depth).size(),
diff --git a/components/operators/src/mask.cu b/components/operators/src/mask.cu
index 85fde5190a7c1e4fd69432b61b515a0c868501de..8632420d7c2510c1c547f2617c7c61a96a4e2d55 100644
--- a/components/operators/src/mask.cu
+++ b/components/operators/src/mask.cu
@@ -1,4 +1,4 @@
-#include <ftl/operators/mask_cuda.hpp>
+#include <ftl/operators/cuda/mask.hpp>
#define T_PER_BLOCK 8
@@ -14,7 +14,7 @@ __global__ void discontinuity_kernel(ftl::cuda::TextureObject<uint8_t> mask_out,
const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
if (x < size.width && y < size.height) {
- Mask mask(0);
+ Mask mask(mask_out(x,y));
const float d = depth.tex2D((int)x, (int)y);
// Multiples of pixel size at given depth
diff --git a/components/operators/src/nvopticalflow.cpp b/components/operators/src/nvopticalflow.cpp
index deee3ecdbc27d9d99228cf607d02703beca8c866..422fc3040d27edd0ee49e0e37f62e1185a1d872e 100644
--- a/components/operators/src/nvopticalflow.cpp
+++ b/components/operators/src/nvopticalflow.cpp
@@ -3,6 +3,9 @@
#include <opencv2/cudaimgproc.hpp>
+#define LOGURU_REPLACE_GLOG 1
+#include <loguru.hpp>
+
using ftl::rgbd::Frame;
using ftl::rgbd::Source;
using ftl::codecs::Channel;
@@ -29,7 +32,9 @@ NVOpticalFlow::~NVOpticalFlow() {
bool NVOpticalFlow::init() {
if (!ftl::cuda::hasCompute(7,5)) {
config()->set("enabled", false);
- throw FTL_Error("GPU does not support optical flow");
+ //throw FTL_Error("GPU does not support optical flow");
+ LOG(ERROR) << "GPU does not support optical flow";
+ return false;
}
nvof_ = cv::cuda::NvidiaOpticalFlow_1_0::create(
size_.width, size_.height,
@@ -47,7 +52,7 @@ bool NVOpticalFlow::apply(Frame &in, Frame &out, cudaStream_t stream) {
if (in.get<GpuMat>(channel_in_).size() != size_) {
size_ = in.get<GpuMat>(channel_in_).size();
- init();
+ if (!init()) return false;
}
auto cvstream = cv::cuda::StreamAccessor::wrapStream(stream);
diff --git a/components/operators/src/segmentation.cu b/components/operators/src/segmentation.cu
index 8704cd275939f9350ead23538ef419b5da6ff0a2..de1268931a7b91bb09f0e3dc6385a8d9c7a4a9d8 100644
--- a/components/operators/src/segmentation.cu
+++ b/components/operators/src/segmentation.cu
@@ -1,5 +1,5 @@
#include "segmentation_cuda.hpp"
-#include <ftl/operators/mask_cuda.hpp>
+#include <ftl/operators/cuda/mask.hpp>
#define T_PER_BLOCK 8
diff --git a/components/operators/src/weighting.cpp b/components/operators/src/weighting.cpp
index 5445d6786de3bc04bd515315910740a49d35290a..549a22f144d7627277e4e0396dd24331feaaacf3 100644
--- a/components/operators/src/weighting.cpp
+++ b/components/operators/src/weighting.cpp
@@ -37,10 +37,16 @@ bool PixelWeights::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, cudaStream
Channel dchan = (in.hasChannel(Channel::Depth)) ? Channel::Depth : Channel::GroundTruth;
+ if (!out.hasChannel(Channel::Mask)) {
+ auto &m = out.create<cv::cuda::GpuMat>(Channel::Mask);
+ m.create(in.get<cv::cuda::GpuMat>(dchan).size(), CV_8UC1);
+ m.setTo(cv::Scalar(0));
+ }
+
if (output_normals) {
ftl::cuda::pixel_weighting(
out.createTexture<short>(Channel::Weights, ftl::rgbd::Format<short>(in.get<cv::cuda::GpuMat>(dchan).size())),
- out.createTexture<uint8_t>(Channel::Mask, ftl::rgbd::Format<uint8_t>(in.get<cv::cuda::GpuMat>(dchan).size())),
+ out.createTexture<uint8_t>(Channel::Mask),
out.createTexture<half4>(Channel::Normals, ftl::rgbd::Format<half4>(in.get<cv::cuda::GpuMat>(dchan).size())),
in.createTexture<uchar4>(Channel::Support1),
in.createTexture<float>(dchan),
diff --git a/components/operators/src/weighting.cu b/components/operators/src/weighting.cu
index ccb5d7f4e31389edd17a22a139b9391d40b16068..90aa5fd494f19e20a28eb7f6dc3ff822bd5c05ce 100644
--- a/components/operators/src/weighting.cu
+++ b/components/operators/src/weighting.cu
@@ -24,7 +24,7 @@ __global__ void pixel_weight_kernel(
const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
if (x < size.width && y < size.height) {
- Mask mask(0);
+ Mask mask(mask_out(x,y));
const float d = depth.tex2D((int)x, (int)y);
// Multiples of pixel size at given depth
diff --git a/components/operators/src/weighting_cuda.hpp b/components/operators/src/weighting_cuda.hpp
index 782aae3df36b214f2962093890c0b8a72b3f1ef9..69b679196e3880991f41e64cc1c2e8308b31f44b 100644
--- a/components/operators/src/weighting_cuda.hpp
+++ b/components/operators/src/weighting_cuda.hpp
@@ -1,7 +1,7 @@
#ifndef _FTL_OPERATORS_WEIGHTING_CUDA_HPP_
#define _FTL_OPERATORS_WEIGHTING_CUDA_HPP_
-#include <ftl/operators/mask_cuda.hpp>
+#include <ftl/operators/cuda/mask.hpp>
namespace ftl {
namespace cuda {
diff --git a/components/renderers/cpp/src/CUDARender.cpp b/components/renderers/cpp/src/CUDARender.cpp
index c2673073bf35d959a6c79441be8a67155b94a8e6..b00efb036525426c8ec7bc61a01f247b83d8436d 100644
--- a/components/renderers/cpp/src/CUDARender.cpp
+++ b/components/renderers/cpp/src/CUDARender.cpp
@@ -3,7 +3,7 @@
#include "splatter_cuda.hpp"
#include <ftl/cuda/points.hpp>
#include <ftl/cuda/normals.hpp>
-#include <ftl/operators/mask_cuda.hpp>
+#include <ftl/operators/cuda/mask.hpp>
#include <ftl/render/colouriser.hpp>
#include <ftl/cuda/transform.hpp>
diff --git a/components/renderers/cpp/src/mask.cu b/components/renderers/cpp/src/mask.cu
index 7e36d9b32e5c024e2fc3094daf66d55248304b8d..9952cf80ce04f785862c2c6a958ab6b267e4697c 100644
--- a/components/renderers/cpp/src/mask.cu
+++ b/components/renderers/cpp/src/mask.cu
@@ -1,5 +1,5 @@
#include "splatter_cuda.hpp"
-#include <ftl/operators/mask_cuda.hpp>
+#include <ftl/operators/cuda/mask.hpp>
using ftl::cuda::TextureObject;
using ftl::cuda::Mask;
diff --git a/components/streams/src/receiver.cpp b/components/streams/src/receiver.cpp
index dfd8f0003bda646ae4588003c6339d219623192c..6a274d56c3efc52007b996b20d0f8bdf7b6b67ef 100644
--- a/components/streams/src/receiver.cpp
+++ b/components/streams/src/receiver.cpp
@@ -208,7 +208,11 @@ void Receiver::_processVideo(const StreamPacket &spkt, const Packet &pkt) {
auto &surface = ividstate.surface[static_cast<int>(spkt.channel)];
// Allocate a decode surface, this is a tiled image to be split later
- surface.create(height*ty, width*tx, ((isFloatChannel(spkt.channel)) ? ((pkt.flags & 0x2) ? CV_16UC4 : CV_16U) : CV_8UC4));
+ int cvtype = ftl::codecs::type(spkt.channel);
+ if (cvtype == CV_32F) cvtype = (pkt.flags & 0x2) ? CV_16UC4 : CV_16U;
+
+ //surface.create(height*ty, width*tx, ((isFloatChannel(spkt.channel)) ? ((pkt.flags & 0x2) ? CV_16UC4 : CV_16U) : CV_8UC4));
+ surface.create(height*ty, width*tx, cvtype);
bool is_static = ividstate.decoders[channum] && (spkt.hint_capability & ftl::codecs::kStreamCap_Static);
@@ -279,7 +283,7 @@ void Receiver::_processVideo(const StreamPacket &spkt, const Packet &pkt) {
// Add channel to frame and allocate memory if required
const cv::Size size = cv::Size(width, height);
- frame.getBuffer<cv::cuda::GpuMat>(spkt.channel).create(size, (isFloatChannel(rchan) ? CV_32FC1 : CV_8UC4));
+ frame.getBuffer<cv::cuda::GpuMat>(spkt.channel).create(size, ftl::codecs::type(spkt.channel)); //(isFloatChannel(rchan) ? CV_32FC1 : CV_8UC4));
cv::Rect roi((i % tx)*width, (i / tx)*height, width, height);
cv::cuda::GpuMat sroi = surface(roi);
@@ -300,6 +304,8 @@ void Receiver::_processVideo(const StreamPacket &spkt, const Packet &pkt) {
ftl::cuda::vuya_to_depth(frame.getBuffer<cv::cuda::GpuMat>(spkt.channel), sroi, 16.0f, cvstream);
} else if (isFloatChannel(rchan)) {
sroi.convertTo(frame.getBuffer<cv::cuda::GpuMat>(spkt.channel), CV_32FC1, 1.0f/1000.0f, cvstream);
+ } else if (sroi.type() == CV_8UC1) {
+ sroi.copyTo(frame.getBuffer<cv::cuda::GpuMat>(spkt.channel), cvstream);
} else {
cv::cuda::cvtColor(sroi, frame.getBuffer<cv::cuda::GpuMat>(spkt.channel), cv::COLOR_RGBA2BGRA, 0, cvstream);
}
diff --git a/components/streams/src/sender.cpp b/components/streams/src/sender.cpp
index 623baf76de7f541d487614467851ac82327db2bb..236f96f5ec959f124dc04fcea909b362591760fb 100644
--- a/components/streams/src/sender.cpp
+++ b/components/streams/src/sender.cpp
@@ -409,7 +409,14 @@ int Sender::_generateTiles(const ftl::rgbd::FrameSet &fs, int offset, Channel c,
int tilecount = tx*ty;
uint32_t count = 0;
- surface.surface.create(height, width, (lossless && m.type() == CV_32F) ? CV_16U : CV_8UC4);
+ int cvtype = CV_8UC4;
+ switch (m.type()) {
+ case CV_32F : cvtype = (lossless && m.type() == CV_32F) ? CV_16U : CV_8UC4; break;
+ case CV_8UC1 : cvtype = CV_8UC1; break;
+ default : cvtype = CV_8UC4;
+ }
+
+ surface.surface.create(height, width, cvtype);
// Loop over tiles with ROI mats and do colour conversions.
while (tilecount > 0 && count+offset < fs.frames.size()) {
@@ -429,6 +436,8 @@ int Sender::_generateTiles(const ftl::rgbd::FrameSet &fs, int offset, Channel c,
cv::cuda::cvtColor(m, sroi, cv::COLOR_BGRA2RGBA, 0, stream);
} else if (m.type() == CV_8UC3) {
cv::cuda::cvtColor(m, sroi, cv::COLOR_BGR2RGBA, 0, stream);
+ } else if (m.type() == CV_8UC1) {
+ m.copyTo(sroi, stream);
} else {
LOG(ERROR) << "Unsupported colour format: " << m.type();
return 0;