diff --git a/applications/reconstruct/CMakeLists.txt b/applications/reconstruct/CMakeLists.txt
index be1fe443ef1c8cf3fb14d207cc75ee81bf7d8437..9d9065d3972a9a915059d2e7c88f17f33c676c42 100644
--- a/applications/reconstruct/CMakeLists.txt
+++ b/applications/reconstruct/CMakeLists.txt
@@ -4,14 +4,14 @@
set(REPSRC
src/main.cpp
- src/voxel_scene.cpp
+ #src/voxel_scene.cpp
#src/ray_cast_sdf.cu
src/camera_util.cu
#src/ray_cast_sdf.cpp
src/registration.cpp
#src/virtual_source.cpp
- src/splat_render.cpp
- src/dibr.cu
+ #src/splat_render.cpp
+ #src/dibr.cu
src/mls.cu
src/depth_camera.cu
src/depth_camera.cpp
diff --git a/applications/reconstruct/include/ftl/cuda_matrix_util.hpp b/components/common/cpp/include/ftl/cuda_matrix_util.hpp
similarity index 100%
rename from applications/reconstruct/include/ftl/cuda_matrix_util.hpp
rename to components/common/cpp/include/ftl/cuda_matrix_util.hpp
diff --git a/components/renderers/cpp/CMakeLists.txt b/components/renderers/cpp/CMakeLists.txt
index 33f910ca0342096bb551b430374776a86de89b2f..89cfe508deac5307fd92f1446d87eaeac202071b 100644
--- a/components/renderers/cpp/CMakeLists.txt
+++ b/components/renderers/cpp/CMakeLists.txt
@@ -1,6 +1,6 @@
add_library(ftlrender
- src/display.cpp
- src/rgbd_display.cpp
+ src/splat_render.cpp
+ src/splatter.cu
)
# These cause errors in CI build and are being removed from PCL in newer versions
@@ -11,6 +11,6 @@ target_include_directories(ftlrender PUBLIC
$<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}/include>
$<INSTALL_INTERFACE:include>
PRIVATE src)
-target_link_libraries(ftlrender ftlrgbd ftlcommon ftlnet Eigen3::Eigen Threads::Threads glog::glog ${OpenCV_LIBS} ${PCL_LIBRARIES})
+target_link_libraries(ftlrender ftlrgbd ftlcommon Eigen3::Eigen Threads::Threads ${OpenCV_LIBS})
#ADD_SUBDIRECTORY(test)
diff --git a/components/renderers/cpp/include/ftl/render/renderer.hpp b/components/renderers/cpp/include/ftl/render/renderer.hpp
index d8817ad9192de3e31a83108e531dfb009490a066..f564ac7a564c5ec7134f40d22444c983b245601e 100644
--- a/components/renderers/cpp/include/ftl/render/renderer.hpp
+++ b/components/renderers/cpp/include/ftl/render/renderer.hpp
@@ -1,13 +1,32 @@
#ifndef _FTL_RENDER_RENDERER_HPP_
#define _FTL_RENDER_RENDERER_HPP_
+#include <ftl/configurable.hpp>
+#include <ftl/rgbd/virtual.hpp>
+#include <ftl/cuda_common.hpp>
+
namespace ftl {
namespace render {
+/**
+ * Abstract class for all renderers. A renderer takes some 3D scene and
+ * generates a virtual camera perspective of that scene. The scene might be
+ * based upon a point cloud, or an entirely virtual mesh or procedural scene.
+ * It is intended that multiple scenes can be rendered into a single virtual
+ * view using a compositing renderer, such a renderer accepting any kind of
+ * renderer for compositing and hence relying on this base class.
+ */
class Renderer : public ftl::Configurable {
public:
- Renderer();
- virtual ~Renderer();
+ explicit Renderer(nlohmann::json &config) : Configurable(config) {};
+ virtual ~Renderer() {};
+
+ /**
+ * Generate a single virtual camera frame. The frame takes its pose from
+ * the virtual camera object passed, and writes the result into the
+ * virtual camera.
+ */
+ virtual bool render(ftl::rgbd::VirtualSource *, cudaStream_t)=0;
};
}
diff --git a/components/renderers/cpp/include/ftl/render/splat_render.hpp b/components/renderers/cpp/include/ftl/render/splat_render.hpp
index 828db11fad6bbb5e3aeeae0899bc15549fab1708..5357bf5c30c600b7a544f5aee87abc6eb9947c69 100644
--- a/components/renderers/cpp/include/ftl/render/splat_render.hpp
+++ b/components/renderers/cpp/include/ftl/render/splat_render.hpp
@@ -1,13 +1,8 @@
#ifndef _FTL_RECONSTRUCTION_SPLAT_HPP_
#define _FTL_RECONSTRUCTION_SPLAT_HPP_
-#include <ftl/configurable.hpp>
-#include <ftl/rgbd/source.hpp>
-#include <ftl/depth_camera.hpp>
-#include <ftl/voxel_scene.hpp>
-//#include <ftl/ray_cast_util.hpp>
-#include <ftl/cuda_common.hpp>
-
+#include <ftl/render/renderer.hpp>
+#include <ftl/rgbd/frameset.hpp>
#include "splat_params.hpp"
namespace ftl {
@@ -21,26 +16,29 @@ namespace render {
* on a separate machine or at a later time, the advantage being to save local
* processing resources and that the first pass result may compress better.
*/
-class Splatter {
+class Splatter : public ftl::render::Renderer {
public:
- explicit Splatter(ftl::voxhash::SceneRep *scene);
+ explicit Splatter(nlohmann::json &config, const ftl::rgbd::FrameSet &fs);
~Splatter();
- void render(int64_t ts, ftl::rgbd::Source *src, cudaStream_t stream=0);
+ void render(ftl::rgbd::VirtualSource *src, cudaStream_t stream=0);
- void setOutputDevice(int);
+ //void setOutputDevice(int);
private:
int device_;
- ftl::cuda::TextureObject<int> depth1_;
+ /*ftl::cuda::TextureObject<int> depth1_;
ftl::cuda::TextureObject<int> depth3_;
ftl::cuda::TextureObject<uchar4> colour1_;
ftl::cuda::TextureObject<float4> colour_tmp_;
ftl::cuda::TextureObject<float> depth2_;
ftl::cuda::TextureObject<uchar4> colour2_;
- ftl::cuda::TextureObject<float4> normal1_;
- SplatParams params_;
- ftl::voxhash::SceneRep *scene_;
+ ftl::cuda::TextureObject<float4> normal1_;*/
+ //SplatParams params_;
+
+ ftl::rgbd::Frame output_;
+ ftl::rgbd::Frame temp_;
+ const ftl::rgbd::FrameSet &scene_;
};
}
diff --git a/components/renderers/cpp/src/splat_params.hpp b/components/renderers/cpp/src/splat_params.hpp
index 8ae5bf345e4e0d348c414cc1ce7bb52190d5ffff..4f9c8882b161d7774388e8d9fff7337cb1d6e685 100644
--- a/components/renderers/cpp/src/splat_params.hpp
+++ b/components/renderers/cpp/src/splat_params.hpp
@@ -3,7 +3,7 @@
#include <ftl/cuda_util.hpp>
#include <ftl/cuda_matrix_util.hpp>
-#include <ftl/depth_camera_params.hpp>
+#include <ftl/rgbd/camera.hpp>
namespace ftl {
namespace render {
@@ -18,10 +18,10 @@ struct __align__(16) SplatParams {
float4x4 m_viewMatrixInverse;
uint m_flags;
- float voxelSize;
+ //float voxelSize;
float depthThreshold;
- DepthCameraParams camera;
+ ftl::rgbd::Camera camera;
};
}
diff --git a/components/renderers/cpp/src/splat_render.cpp b/components/renderers/cpp/src/splat_render.cpp
index 21e1b1db3952422e82be757fec56d3c75de0b4c8..fc92d541092b79f7eb744fedd3f73445f187b8ee 100644
--- a/components/renderers/cpp/src/splat_render.cpp
+++ b/components/renderers/cpp/src/splat_render.cpp
@@ -1,12 +1,11 @@
-#include "splat_render.hpp"
-#include "splat_render_cuda.hpp"
-#include "compactors.hpp"
-#include "depth_camera_cuda.hpp"
+#include <ftl/render/splat_render.hpp>
+#include "splatter_cuda.hpp"
using ftl::render::Splatter;
using ftl::rgbd::Channel;
+using ftl::rgbd::Format;
-Splatter::Splatter(ftl::voxhash::SceneRep *scene) : scene_(scene) {
+Splatter::Splatter(nlohmann::json &config, const ftl::rgbd::FrameSet &fs) : ftl::render::Renderer(config), scene_(fs) {
}
@@ -14,15 +13,25 @@ Splatter::~Splatter() {
}
-void Splatter::render(int64_t ts, ftl::rgbd::Source *src, cudaStream_t stream) {
+void Splatter::render(ftl::rgbd::VirtualSource *src, cudaStream_t stream) {
if (!src->isReady()) return;
const auto &camera = src->parameters();
- cudaSafeCall(cudaSetDevice(scene_->getCUDADevice()));
+ //cudaSafeCall(cudaSetDevice(scene_->getCUDADevice()));
- // Create buffers if they don't exists
- if ((unsigned int)depth1_.width() != camera.width || (unsigned int)depth1_.height() != camera.height) {
+ output_.create<cv::cuda::GpuMat>(Channel::Depth, Format<float>(camera.width, camera.height));
+ output_.create<cv::cuda::GpuMat>(Channel::Colour, Format<uchar4>(camera.width, camera.height));
+
+ temp_.create<cv::cuda::GpuMat>(Channel::Colour, Format<float4>(camera.width, camera.height));
+ temp_.create<cv::cuda::GpuMat>(Channel::Colour2, Format<uchar4>(camera.width, camera.height));
+ temp_.create<cv::cuda::GpuMat>(Channel::Confidence, Format<float>(camera.width, camera.height));
+ temp_.create<cv::cuda::GpuMat>(Channel::Depth, Format<int>(camera.width, camera.height));
+ temp_.create<cv::cuda::GpuMat>(Channel::Depth2, Format<int>(camera.width, camera.height));
+ temp_.create<cv::cuda::GpuMat>(Channel::Normals, Format<float4>(camera.width, camera.height));
+
+ // Create buffers if they don't exist
+ /*if ((unsigned int)depth1_.width() != camera.width || (unsigned int)depth1_.height() != camera.height) {
depth1_ = ftl::cuda::TextureObject<int>(camera.width, camera.height);
}
if ((unsigned int)depth3_.width() != camera.width || (unsigned int)depth3_.height() != camera.height) {
@@ -42,7 +51,7 @@ void Splatter::render(int64_t ts, ftl::rgbd::Source *src, cudaStream_t stream) {
}
if ((unsigned int)colour2_.width() != camera.width || (unsigned int)colour2_.height() != camera.height) {
colour2_ = ftl::cuda::TextureObject<uchar4>(camera.width, camera.height);
- }
+ }*/
// Parameters object to pass to CUDA describing the camera
SplatParams params;
@@ -53,27 +62,27 @@ void Splatter::render(int64_t ts, ftl::rgbd::Source *src, cudaStream_t stream) {
params.m_viewMatrix = MatrixConversion::toCUDA(src->getPose().cast<float>().inverse());
params.m_viewMatrixInverse = MatrixConversion::toCUDA(src->getPose().cast<float>());
- params.voxelSize = scene_->getHashParams().m_virtualVoxelSize;
- params.camera.flags = 0;
- params.camera.fx = camera.fx;
+ //params.voxelSize = scene_->getHashParams().m_virtualVoxelSize;
+ params.camera = camera;
+ /*params.camera.fx = camera.fx;
params.camera.fy = camera.fy;
params.camera.mx = -camera.cx;
params.camera.my = -camera.cy;
params.camera.m_imageWidth = camera.width;
params.camera.m_imageHeight = camera.height;
params.camera.m_sensorDepthWorldMax = camera.maxDepth;
- params.camera.m_sensorDepthWorldMin = camera.minDepth;
+ params.camera.m_sensorDepthWorldMin = camera.minDepth;*/
//ftl::cuda::compactifyAllocated(scene_->getHashData(), scene_->getHashParams(), stream);
//LOG(INFO) << "Occupied: " << scene_->getOccupiedCount();
- if (scene_->value("voxels", false)) {
+ //if (scene_->value("voxels", false)) {
// TODO:(Nick) Stereo for voxel version
- ftl::cuda::isosurface_point_image(scene_->getHashData(), depth1_, params, stream);
+ //ftl::cuda::isosurface_point_image(scene_->getHashData(), depth1_, params, stream);
//ftl::cuda::splat_points(depth1_, depth2_, params, stream);
//ftl::cuda::dibr(depth2_, colour1_, scene_->cameraCount(), params, stream);
- src->writeFrames(ts, colour1_, depth2_, stream);
- } else {
+ //src->writeFrames(ts, colour1_, depth2_, stream);
+ //} else {
ftl::cuda::clear_depth(depth1_, stream);
ftl::cuda::clear_depth(depth3_, stream);
ftl::cuda::clear_depth(depth2_, stream);
@@ -88,7 +97,7 @@ void Splatter::render(int64_t ts, ftl::rgbd::Source *src, cudaStream_t stream) {
if (src->getChannel() == Channel::Depth) {
//ftl::cuda::int_to_float(depth1_, depth2_, 1.0f / 1000.0f, stream);
- if (src->value("splatting", false)) {
+ if (value("splatting", false)) {
//ftl::cuda::splat_points(depth1_, colour1_, normal1_, depth2_, colour2_, params, stream);
ftl::cuda::int_to_float(depth1_, depth2_, 1.0f / 1000.0f, stream);
src->writeFrames(ts, colour1_, depth2_, stream);
@@ -117,7 +126,7 @@ void Splatter::render(int64_t ts, ftl::rgbd::Source *src, cudaStream_t stream) {
ftl::cuda::dibr(depth1_, colour1_, normal1_, depth2_, colour_tmp_, depth3_, scene_->cameraCount(), params, stream);
src->writeFrames(ts, colour1_, colour2_, stream);
} else {
- if (src->value("splatting", false)) {
+ if (value("splatting", false)) {
//ftl::cuda::splat_points(depth1_, colour1_, normal1_, depth2_, colour2_, params, stream);
src->writeFrames(ts, colour1_, depth2_, stream);
} else {
@@ -125,7 +134,7 @@ void Splatter::render(int64_t ts, ftl::rgbd::Source *src, cudaStream_t stream) {
src->writeFrames(ts, colour1_, depth2_, stream);
}
}
- }
+ //}
//ftl::cuda::median_filter(depth1_, depth2_, stream);
//ftl::cuda::splat_points(depth1_, depth2_, params, stream);
diff --git a/components/renderers/cpp/src/splatter.cu b/components/renderers/cpp/src/splatter.cu
new file mode 100644
index 0000000000000000000000000000000000000000..586eb40fd97bf56d27c9d87c5ab4d7d68c11da5b
--- /dev/null
+++ b/components/renderers/cpp/src/splatter.cu
@@ -0,0 +1,134 @@
+#include "splat_params.hpp"
+#include <ftl/rgbd/camera.hpp>
+#include <ftl/cuda_common.hpp>
+
+using ftl::cuda::TextureObject;
+using ftl::render::SplatParams;
+
+/*
+ * Pass 1: Directly render each camera into virtual view but with no upsampling
+ * for sparse points.
+ */
+ __global__ void dibr_merge_kernel(TextureObject<float4> points, TextureObject<int> depth, int cam, SplatParams params) {
+ const int x = blockIdx.x*blockDim.x + threadIdx.x;
+ const int y = blockIdx.y*blockDim.y + threadIdx.y;
+
+ const float3 worldPos = make_float3(tex2D<float4>(points, x, y));
+ if (worldPos.x == MINF) return;
+
+ // Find the virtual screen position of current point
+ const float3 camPos = params.m_viewMatrix * worldPos;
+ if (camPos.z < params.camera.minDepth) return;
+ if (camPos.z > params.camera.maxDepth) return;
+
+ const float d = camPos.z;
+
+ const uint2 screenPos = params.camera.camToScreen<uint2>(camPos);
+ const unsigned int cx = screenPos.x;
+ const unsigned int cy = screenPos.y;
+ if (d > params.camera.minDepth && d < params.camera.maxDepth && cx < depth.width() && cy < depth.height()) {
+ // Transform estimated point to virtual cam space and output z
+ atomicMin(&depth(cx,cy), d * 1000.0f);
+ }
+}
+
+__device__ inline float4 make_float4(const uchar4 &c) {
+ return make_float4(c.x,c.y,c.z,c.w);
+}
+
+/*
+ * Pass 2: Accumulate attribute contributions if the points pass a visibility test.
+ */
+__global__ void dibr_attribute_contrib_kernel(
+ TextureObject<uchar4> colour_in, // Original colour image
+ TextureObject<float4> points, // Original 3D points
+ TextureObject<int> depth_in, // Virtual depth map
+ TextureObject<float4> colour_out, // Accumulated output
+ //TextureObject<float4> normal_out,
+ TextureObject<float> contrib_out,
+ SplatParams params) {
+
+ //const ftl::voxhash::DepthCameraCUDA &camera = c_cameras[cam];
+
+ const int tid = (threadIdx.x + threadIdx.y * blockDim.x);
+ //const int warp = tid / WARP_SIZE;
+ const int x = (blockIdx.x*blockDim.x + threadIdx.x) / WARP_SIZE;
+ const int y = blockIdx.y*blockDim.y + threadIdx.y;
+
+ const float3 worldPos = make_float3(tex2D<float4>(camera.points, x, y));
+ //const float3 normal = make_float3(tex2D<float4>(camera.normal, x, y));
+ if (worldPos.x == MINF) return;
+ //const float r = (camera.poseInverse * worldPos).z / camera.params.fx;
+
+ const float3 camPos = params.m_viewMatrix * worldPos;
+ if (camPos.z < params.camera.minDepth) return;
+ if (camPos.z > params.camera.maxDepth) return;
+ const uint2 screenPos = params.camera.camToScreen<uint2>(camPos);
+
+ const int upsample = 8; //min(UPSAMPLE_MAX, int((5.0f*r) * params.camera.fx / camPos.z));
+
+ // Not on screen so stop now...
+ if (screenPos.x >= depth_in.width() || screenPos.y >= depth_in.height()) return;
+
+ // Is this point near the actual surface and therefore a contributor?
+ const float d = ((float)depth_in.tex2D((int)screenPos.x, (int)screenPos.y)/1000.0f);
+ //if (abs(d - camPos.z) > DEPTH_THRESHOLD) return;
+
+ // TODO:(Nick) Should just one thread load these to shared mem?
+ const float4 colour = make_float4(colour_in.tex2D<uchar4>(x, y));
+ //const float4 normal = tex2D<float4>(camera.normal, x, y);
+
+ // Each thread in warp takes an upsample point and updates corresponding depth buffer.
+ const int lane = tid % WARP_SIZE;
+ for (int i=lane; i<upsample*upsample; i+=WARP_SIZE) {
+ const float u = (i % upsample) - (upsample / 2);
+ const float v = (i / upsample) - (upsample / 2);
+
+ // Use the depth buffer to determine this pixels 3D position in camera space
+ const float d = ((float)depth_in.tex2D(screenPos.x+u, screenPos.y+v)/1000.0f);
+ const float3 nearest = params.camera.screenToCam((int)(screenPos.x+u),(int)(screenPos.y+v),d);
+
+ // What is contribution of our current point at this pixel?
+ const float weight = ftl::cuda::spatialWeighting(length(nearest - camPos), SMOOTHING_MULTIPLIER_C*(nearest.z/params.camera.fx));
+ if (screenPos.x+u < colour_out.width() && screenPos.y+v < colour_out.height() && weight > 0.0f) { // TODO: Use confidence threshold here
+ const float4 wcolour = colour * weight;
+ const float4 wnormal = normal * weight;
+
+ //printf("Z %f\n", d);
+
+ // Add this points contribution to the pixel buffer
+ atomicAdd((float*)&colour_out(screenPos.x+u, screenPos.y+v), wcolour.x);
+ atomicAdd((float*)&colour_out(screenPos.x+u, screenPos.y+v)+1, wcolour.y);
+ atomicAdd((float*)&colour_out(screenPos.x+u, screenPos.y+v)+2, wcolour.z);
+ atomicAdd((float*)&colour_out(screenPos.x+u, screenPos.y+v)+3, wcolour.w);
+ //atomicAdd((float*)&normal_out(screenPos.x+u, screenPos.y+v), wnormal.x);
+ //atomicAdd((float*)&normal_out(screenPos.x+u, screenPos.y+v)+1, wnormal.y);
+ //atomicAdd((float*)&normal_out(screenPos.x+u, screenPos.y+v)+2, wnormal.z);
+ //atomicAdd((float*)&normal_out(screenPos.x+u, screenPos.y+v)+3, wnormal.w);
+ atomicAdd(&contrib_out(screenPos.x+u, screenPos.y+v), weight);
+ }
+ }
+}
+
+
+
+
+__global__ void dibr_normalise_kernel(
+ TextureObject<float4> colour_in,
+ TextureObject<uchar4> colour_out,
+ //TextureObject<float4> normals,
+ TextureObject<float> contribs) {
+const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
+const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
+
+if (x < colour_in.width() && y < colour_in.height()) {
+ const float4 colour = colour_in.tex2D((int)x,(int)y);
+ //const float4 normal = normals.tex2D((int)x,(int)y);
+ const float contrib = contribs.tex2D((int)x,(int)y);
+
+ if (contrib > 0.0f) {
+ colour_out(x,y) = make_uchar4(colour.x / contrib, colour.y / contrib, colour.z / contrib, 0);
+ //normals(x,y) = normal / contrib;
+ }
+}
+}
diff --git a/components/renderers/cpp/src/splat_render_cuda.hpp b/components/renderers/cpp/src/splatter_cuda.hpp
similarity index 85%
rename from components/renderers/cpp/src/splat_render_cuda.hpp
rename to components/renderers/cpp/src/splatter_cuda.hpp
index e60fc8c27c0d39ef8798803a526891c5da2fca62..125d43390be2514a0e8f301b37d426d3fccb4350 100644
--- a/components/renderers/cpp/src/splat_render_cuda.hpp
+++ b/components/renderers/cpp/src/splatter_cuda.hpp
@@ -1,10 +1,7 @@
#ifndef _FTL_RECONSTRUCTION_SPLAT_CUDA_HPP_
#define _FTL_RECONSTRUCTION_SPLAT_CUDA_HPP_
-#include <ftl/depth_camera.hpp>
-#include <ftl/voxel_hash.hpp>
-//#include <ftl/ray_cast_util.hpp>
-
+#include <ftl/cuda_common.hpp>
#include "splat_params.hpp"
namespace ftl {
@@ -84,21 +81,6 @@ __device__ inline float intersectDistance(const float3 &n, const float3 &p0, con
return PINF;
}
-/**
- * NOTE: Not strictly isosurface currently since it includes the internals
- * of objects up to at most truncation depth.
- */
-void isosurface_point_image(const ftl::voxhash::HashData& hashData,
- const ftl::cuda::TextureObject<int> &depth,
- const ftl::render::SplatParams ¶ms, cudaStream_t stream);
-
-//void isosurface_point_image_stereo(const ftl::voxhash::HashData& hashData,
-// const ftl::voxhash::HashParams& hashParams,
-// const RayCastData &rayCastData, const RayCastParams ¶ms,
-// cudaStream_t stream);
-
-// TODO: isosurface_point_cloud
-
void splat_points(const ftl::cuda::TextureObject<int> &depth_in,
const ftl::cuda::TextureObject<uchar4> &colour_in,
const ftl::cuda::TextureObject<float4> &normal_in,
diff --git a/components/rgbd-sources/include/ftl/rgbd/channels.hpp b/components/rgbd-sources/include/ftl/rgbd/channels.hpp
index 591e0e312b0c5a161dc40ed0d4966852d8781971..9c7f24b036c95271f990780d61395c1b595fbb1e 100644
--- a/components/rgbd-sources/include/ftl/rgbd/channels.hpp
+++ b/components/rgbd-sources/include/ftl/rgbd/channels.hpp
@@ -12,9 +12,11 @@ enum struct Channel : int {
Colour = 0,
Left = 0,
Depth = 1,
- Right,
- Disparity,
- Deviation,
+ Right = 2,
+ Colour2 = 2,
+ Disparity = 3,
+ Depth2 = 3,
+ Deviation = 4,
Normals,
Confidence,
Flow,
diff --git a/components/rgbd-sources/include/ftl/rgbd/source.hpp b/components/rgbd-sources/include/ftl/rgbd/source.hpp
index 07391d76f10f4eec69f4a7df97123dfc59270a0c..e14950e63e45df89ea144f599e3750b1daa02444 100644
--- a/components/rgbd-sources/include/ftl/rgbd/source.hpp
+++ b/components/rgbd-sources/include/ftl/rgbd/source.hpp
@@ -121,17 +121,6 @@ class Source : public ftl::Configurable {
*/
void getDepth(cv::Mat &d);
- /**
- * Write frames into source buffers from an external renderer. Virtual
- * sources do not have an internal generator of frames but instead have
- * their data provided from an external rendering class. This function only
- * works when there is no internal generator.
- */
- void writeFrames(int64_t ts, const cv::Mat &rgb, const cv::Mat &depth);
- void writeFrames(int64_t ts, const ftl::cuda::TextureObject<uchar4> &rgb, const ftl::cuda::TextureObject<uint> &depth, cudaStream_t stream);
- void writeFrames(int64_t ts, const ftl::cuda::TextureObject<uchar4> &rgb, const ftl::cuda::TextureObject<float> &depth, cudaStream_t stream);
- void writeFrames(int64_t ts, const ftl::cuda::TextureObject<uchar4> &rgb, const ftl::cuda::TextureObject<uchar4> &rgb2, cudaStream_t stream);
-
int64_t timestamp() const { return timestamp_; }
/**
@@ -214,7 +203,7 @@ class Source : public ftl::Configurable {
void removeCallback() { callback_ = nullptr; }
- private:
+ protected:
detail::Source *impl_;
cv::Mat rgb_;
cv::Mat depth_;
diff --git a/components/rgbd-sources/include/ftl/rgbd/virtual.hpp b/components/rgbd-sources/include/ftl/rgbd/virtual.hpp
new file mode 100644
index 0000000000000000000000000000000000000000..e8e8edd406779191234b76bd36d6893dd4cd8d9b
--- /dev/null
+++ b/components/rgbd-sources/include/ftl/rgbd/virtual.hpp
@@ -0,0 +1,25 @@
+#ifndef _FTL_RGBD_VIRTUAL_HPP_
+#define _FTL_RGBD_VIRTUAL_HPP_
+
+#include <ftl/rgbd/source.hpp>
+
+namespace ftl {
+namespace rgbd {
+
+class VirtualSource : public ftl::rgbd::Source {
+ public:
+ explicit VirtualSource(ftl::config::json_t &cfg);
+
+ /**
+ * Write frames into source buffers from an external renderer. Virtual
+ * sources do not have an internal generator of frames but instead have
+ * their data provided from an external rendering class. This function only
+ * works when there is no internal generator.
+ */
+ void write(int64_t ts, ftl::rgbd::Frame &frame, cudaStream_t stream=0);
+};
+
+}
+}
+
+#endif // _FTL_RGBD_VIRTUAL_HPP_
diff --git a/components/rgbd-sources/src/source.cpp b/components/rgbd-sources/src/source.cpp
index d33473fc38fbf55670e8b8844ffa5a03487e606d..35d23f27ad7edac18d3e3e02247296f1382be5e2 100644
--- a/components/rgbd-sources/src/source.cpp
+++ b/components/rgbd-sources/src/source.cpp
@@ -273,71 +273,6 @@ bool Source::compute(int N, int B) {
return false;
}
-void Source::writeFrames(int64_t ts, const cv::Mat &rgb, const cv::Mat &depth) {
- if (!impl_) {
- UNIQUE_LOCK(mutex_,lk);
- rgb.copyTo(rgb_);
- depth.copyTo(depth_);
- timestamp_ = ts;
- }
-}
-
-void Source::writeFrames(int64_t ts, const ftl::cuda::TextureObject<uchar4> &rgb, const ftl::cuda::TextureObject<uint> &depth, cudaStream_t stream) {
- if (!impl_) {
- UNIQUE_LOCK(mutex_,lk);
- timestamp_ = ts;
- rgb_.create(rgb.height(), rgb.width(), CV_8UC4);
- cudaSafeCall(cudaMemcpy2DAsync(rgb_.data, rgb_.step, rgb.devicePtr(), rgb.pitch(), rgb_.cols * sizeof(uchar4), rgb_.rows, cudaMemcpyDeviceToHost, stream));
- depth_.create(depth.height(), depth.width(), CV_32SC1);
- cudaSafeCall(cudaMemcpy2DAsync(depth_.data, depth_.step, depth.devicePtr(), depth.pitch(), depth_.cols * sizeof(uint), depth_.rows, cudaMemcpyDeviceToHost, stream));
- //cudaSafeCall(cudaStreamSynchronize(stream)); // TODO:(Nick) Don't wait here.
- stream_ = stream;
- //depth_.convertTo(depth_, CV_32F, 1.0f / 1000.0f);
- } else {
- LOG(ERROR) << "writeFrames cannot be done on this source: " << getURI();
- }
-}
-
-void Source::writeFrames(int64_t ts, const ftl::cuda::TextureObject<uchar4> &rgb, const ftl::cuda::TextureObject<float> &depth, cudaStream_t stream) {
- if (!impl_) {
- UNIQUE_LOCK(mutex_,lk);
- timestamp_ = ts;
- rgb.download(rgb_, stream);
- //rgb_.create(rgb.height(), rgb.width(), CV_8UC4);
- //cudaSafeCall(cudaMemcpy2DAsync(rgb_.data, rgb_.step, rgb.devicePtr(), rgb.pitch(), rgb_.cols * sizeof(uchar4), rgb_.rows, cudaMemcpyDeviceToHost, stream));
- depth.download(depth_, stream);
- //depth_.create(depth.height(), depth.width(), CV_32FC1);
- //cudaSafeCall(cudaMemcpy2DAsync(depth_.data, depth_.step, depth.devicePtr(), depth.pitch(), depth_.cols * sizeof(float), depth_.rows, cudaMemcpyDeviceToHost, stream));
-
- stream_ = stream;
- cudaSafeCall(cudaStreamSynchronize(stream_));
- cv::cvtColor(rgb_,rgb_, cv::COLOR_BGRA2BGR);
- cv::cvtColor(rgb_,rgb_, cv::COLOR_Lab2BGR);
-
- if (callback_) callback_(timestamp_, rgb_, depth_);
- }
-}
-
-void Source::writeFrames(int64_t ts, const ftl::cuda::TextureObject<uchar4> &rgb, const ftl::cuda::TextureObject<uchar4> &rgb2, cudaStream_t stream) {
- if (!impl_) {
- UNIQUE_LOCK(mutex_,lk);
- timestamp_ = ts;
- rgb.download(rgb_, stream);
- //rgb_.create(rgb.height(), rgb.width(), CV_8UC4);
- //cudaSafeCall(cudaMemcpy2DAsync(rgb_.data, rgb_.step, rgb.devicePtr(), rgb.pitch(), rgb_.cols * sizeof(uchar4), rgb_.rows, cudaMemcpyDeviceToHost, stream));
- rgb2.download(depth_, stream);
- //depth_.create(depth.height(), depth.width(), CV_32FC1);
- //cudaSafeCall(cudaMemcpy2DAsync(depth_.data, depth_.step, depth.devicePtr(), depth.pitch(), depth_.cols * sizeof(float), depth_.rows, cudaMemcpyDeviceToHost, stream));
-
- stream_ = stream;
- cudaSafeCall(cudaStreamSynchronize(stream_));
- cv::cvtColor(rgb_,rgb_, cv::COLOR_BGRA2BGR);
- cv::cvtColor(rgb_,rgb_, cv::COLOR_Lab2BGR);
- cv::cvtColor(depth_,depth_, cv::COLOR_BGRA2BGR);
- cv::cvtColor(depth_,depth_, cv::COLOR_Lab2BGR);
- }
-}
-
bool Source::thumbnail(cv::Mat &t) {
if (!impl_ && stream_ != 0) {
cudaSafeCall(cudaStreamSynchronize(stream_));
diff --git a/components/rgbd-sources/src/virtual.cpp b/components/rgbd-sources/src/virtual.cpp
new file mode 100644
index 0000000000000000000000000000000000000000..9d64f53c73b2902d8fe074d84a8192657c36d1fc
--- /dev/null
+++ b/components/rgbd-sources/src/virtual.cpp
@@ -0,0 +1,66 @@
+/*
+void Source::writeFrames(int64_t ts, const cv::Mat &rgb, const cv::Mat &depth) {
+ if (!impl_) {
+ UNIQUE_LOCK(mutex_,lk);
+ rgb.copyTo(rgb_);
+ depth.copyTo(depth_);
+ timestamp_ = ts;
+ }
+}
+
+void Source::writeFrames(int64_t ts, const ftl::cuda::TextureObject<uchar4> &rgb, const ftl::cuda::TextureObject<uint> &depth, cudaStream_t stream) {
+ if (!impl_) {
+ UNIQUE_LOCK(mutex_,lk);
+ timestamp_ = ts;
+ rgb_.create(rgb.height(), rgb.width(), CV_8UC4);
+ cudaSafeCall(cudaMemcpy2DAsync(rgb_.data, rgb_.step, rgb.devicePtr(), rgb.pitch(), rgb_.cols * sizeof(uchar4), rgb_.rows, cudaMemcpyDeviceToHost, stream));
+ depth_.create(depth.height(), depth.width(), CV_32SC1);
+ cudaSafeCall(cudaMemcpy2DAsync(depth_.data, depth_.step, depth.devicePtr(), depth.pitch(), depth_.cols * sizeof(uint), depth_.rows, cudaMemcpyDeviceToHost, stream));
+ //cudaSafeCall(cudaStreamSynchronize(stream)); // TODO:(Nick) Don't wait here.
+ stream_ = stream;
+ //depth_.convertTo(depth_, CV_32F, 1.0f / 1000.0f);
+ } else {
+ LOG(ERROR) << "writeFrames cannot be done on this source: " << getURI();
+ }
+}
+
+void Source::writeFrames(int64_t ts, const ftl::cuda::TextureObject<uchar4> &rgb, const ftl::cuda::TextureObject<float> &depth, cudaStream_t stream) {
+ if (!impl_) {
+ UNIQUE_LOCK(mutex_,lk);
+ timestamp_ = ts;
+ rgb.download(rgb_, stream);
+ //rgb_.create(rgb.height(), rgb.width(), CV_8UC4);
+ //cudaSafeCall(cudaMemcpy2DAsync(rgb_.data, rgb_.step, rgb.devicePtr(), rgb.pitch(), rgb_.cols * sizeof(uchar4), rgb_.rows, cudaMemcpyDeviceToHost, stream));
+ depth.download(depth_, stream);
+ //depth_.create(depth.height(), depth.width(), CV_32FC1);
+ //cudaSafeCall(cudaMemcpy2DAsync(depth_.data, depth_.step, depth.devicePtr(), depth.pitch(), depth_.cols * sizeof(float), depth_.rows, cudaMemcpyDeviceToHost, stream));
+
+ stream_ = stream;
+ cudaSafeCall(cudaStreamSynchronize(stream_));
+ cv::cvtColor(rgb_,rgb_, cv::COLOR_BGRA2BGR);
+ cv::cvtColor(rgb_,rgb_, cv::COLOR_Lab2BGR);
+
+ if (callback_) callback_(timestamp_, rgb_, depth_);
+ }
+}
+
+void Source::writeFrames(int64_t ts, const ftl::cuda::TextureObject<uchar4> &rgb, const ftl::cuda::TextureObject<uchar4> &rgb2, cudaStream_t stream) {
+ if (!impl_) {
+ UNIQUE_LOCK(mutex_,lk);
+ timestamp_ = ts;
+ rgb.download(rgb_, stream);
+ //rgb_.create(rgb.height(), rgb.width(), CV_8UC4);
+ //cudaSafeCall(cudaMemcpy2DAsync(rgb_.data, rgb_.step, rgb.devicePtr(), rgb.pitch(), rgb_.cols * sizeof(uchar4), rgb_.rows, cudaMemcpyDeviceToHost, stream));
+ rgb2.download(depth_, stream);
+ //depth_.create(depth.height(), depth.width(), CV_32FC1);
+ //cudaSafeCall(cudaMemcpy2DAsync(depth_.data, depth_.step, depth.devicePtr(), depth.pitch(), depth_.cols * sizeof(float), depth_.rows, cudaMemcpyDeviceToHost, stream));
+
+ stream_ = stream;
+ cudaSafeCall(cudaStreamSynchronize(stream_));
+ cv::cvtColor(rgb_,rgb_, cv::COLOR_BGRA2BGR);
+ cv::cvtColor(rgb_,rgb_, cv::COLOR_Lab2BGR);
+ cv::cvtColor(depth_,depth_, cv::COLOR_BGRA2BGR);
+ cv::cvtColor(depth_,depth_, cv::COLOR_Lab2BGR);
+ }
+}
+*/
\ No newline at end of file