WIP massive renderer refactor

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

components/renderers/cpp/CMakeLists.txt

 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)

components/renderers/cpp/include/ftl/render/renderer.hpp

 #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;
 };
 
 }

components/renderers/cpp/include/ftl/render/splat_render.hpp

 #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_;
 };
 
 }

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;
 };
 
 }

components/renderers/cpp/src/splat_render.cpp

-#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);

components/renderers/cpp/src/splatter.cu

+#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;
+    }
+}
+}

components/renderers/cpp/src/splat_render_cuda.hpp → components/renderers/cpp/src/splatter_cuda.hpp

 #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 &params, cudaStream_t stream);
-
-//void isosurface_point_image_stereo(const ftl::voxhash::HashData& hashData,
-//		const ftl::voxhash::HashParams& hashParams,
-//		const RayCastData &rayCastData, const RayCastParams &params,
-//		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,

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,

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_;

components/rgbd-sources/include/ftl/rgbd/virtual.hpp

+#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_

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_));

components/rgbd-sources/src/virtual.cpp

+/*
+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