diff --git a/components/codecs/include/ftl/codecs/channels.hpp b/components/codecs/include/ftl/codecs/channels.hpp
index 7cd5827f33bc8d2506024e92873754e99e75c007..43ef166a203e427add9835a732c41e7c39018b30 100644
--- a/components/codecs/include/ftl/codecs/channels.hpp
+++ b/components/codecs/include/ftl/codecs/channels.hpp
@@ -22,7 +22,8 @@ enum struct Channel : int {
 	Confidence		= 7,	// 32F
 	Contribution	= 7,	// 32F
 	EnergyVector	= 8,	// 32FC4
-	Flow			= 9,	// 32F
+	Flow			= 9,	// 16SC2
+	Flow2			= 10,	// 16SC2
 	Energy			= 10,	// 32F
 	Mask			= 11,	// 32U
 	Density			= 12,	// 32F
diff --git a/components/operators/CMakeLists.txt b/components/operators/CMakeLists.txt
index 0fd5466aed057c3d1034501fee00518de5a44bf9..becb9f60b3acd370bde53a6b54617860ce624bb8 100644
--- a/components/operators/CMakeLists.txt
+++ b/components/operators/CMakeLists.txt
@@ -42,6 +42,7 @@ endif (HAVE_LIBSGM)
 if (HAVE_OPTFLOW)
 	list(APPEND OPERSRC
 		src/nvopticalflow.cpp
+		src/opticalflow.cu
 		src/disparity/optflow_smoothing.cu
 		src/disparity/optflow_smoothing.cpp)
 endif()
diff --git a/components/operators/include/ftl/cuda/fixed.hpp b/components/operators/include/ftl/cuda/fixed.hpp
new file mode 100644
index 0000000000000000000000000000000000000000..efb0c09caabd5e41f024679e8e80a2c6ea2b6df3
--- /dev/null
+++ b/components/operators/include/ftl/cuda/fixed.hpp
@@ -0,0 +1,14 @@
+#ifndef _FTL_CUDA_FIXED_HPP_
+#define _FTL_CUDA_FIXED_HPP_
+
+template <int FRAC>
+__device__ inline float fixed2float(short v) {
+    return v / (1 << FRAC);
+}
+
+template <int FRAC>
+__device__ inline short float2fixed(float v) {
+    return short(v * float(1 << FRAC));
+}
+
+#endif
\ No newline at end of file
diff --git a/components/operators/include/ftl/operators/cuda/disparity.hpp b/components/operators/include/ftl/operators/cuda/disparity.hpp
index 75034d9e1808e6023712988e1a66ec2d8f118a0d..4f421969cd8609972e8895e2d82ca6c3f30741fb 100644
--- a/components/operators/include/ftl/operators/cuda/disparity.hpp
+++ b/components/operators/include/ftl/operators/cuda/disparity.hpp
@@ -29,6 +29,8 @@ void show_rpe(const cv::cuda::GpuMat &disp, cv::cuda::GpuMat &left, const cv::cu
 void show_disp_density(const cv::cuda::GpuMat &disp, cv::cuda::GpuMat &left,
 			float scale, cudaStream_t stream);
 
+void merge_disparities(cv::cuda::GpuMat &disp, const cv::cuda::GpuMat &estimate, 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/opticalflow.hpp b/components/operators/include/ftl/operators/opticalflow.hpp
index 43edf6b8440d311c0cfde052e6a9c1a6bf3b78b8..8ee77736dcba7d95becf381759266dfeb543cba3 100644
--- a/components/operators/include/ftl/operators/opticalflow.hpp
+++ b/components/operators/include/ftl/operators/opticalflow.hpp
@@ -13,7 +13,7 @@ namespace operators {
 class NVOpticalFlow : public ftl::operators::Operator {
 	public:
 	explicit NVOpticalFlow(ftl::Configurable*);
-	NVOpticalFlow(ftl::Configurable*, const std::tuple<ftl::codecs::Channel,ftl::codecs::Channel> &channels);
+	NVOpticalFlow(ftl::Configurable*, const std::tuple<ftl::codecs::Channel,ftl::codecs::Channel,ftl::codecs::Channel,ftl::codecs::Channel> &channels);
 	~NVOpticalFlow();
 
 	inline Operator::Type type() const override { return Operator::Type::OneToOne; }
@@ -27,12 +27,14 @@ class NVOpticalFlow : public ftl::operators::Operator {
 	cv::Size size_;
 	
 	// TODO: Colour to Flow always assumed, could also calculate something else?
-	ftl::codecs::Channel channel_in_;
-	ftl::codecs::Channel channel_out_;
+	ftl::codecs::Channel channel_in_[2];
+	ftl::codecs::Channel channel_out_[2];
 
 	cv::Ptr<cv::cuda::NvidiaOpticalFlow_1_0> nvof_;
 	cv::cuda::GpuMat left_gray_;
 	cv::cuda::GpuMat left_gray_prev_;
+	cv::cuda::GpuMat right_gray_;
+	cv::cuda::GpuMat right_gray_prev_;
 };
 
 }
diff --git a/components/operators/src/colours.cpp b/components/operators/src/colours.cpp
index 2c102482ec3738a649a57a618b867388d209f0e4..8ca2c62bf8613e77df8100f6781ec5c36f47427e 100644
--- a/components/operators/src/colours.cpp
+++ b/components/operators/src/colours.cpp
@@ -45,6 +45,9 @@ bool ColourChannels::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, cudaStre
 
 	//in.resetTexture(Channel::Colour);
 	in.createTexture<uchar4>(Channel::Colour, true);
+	if (in.hasChannel(Channel::Right)) {
+		in.createTexture<uchar4>(Channel::Right, true);
+	}
 
 	if (in.hasChannel(Channel::Depth)) {
 		auto &depth = in.get<cv::cuda::GpuMat>(Channel::Depth);
diff --git a/components/operators/src/depth.cpp b/components/operators/src/depth.cpp
index 5ffb6bbaf296d7b1f65ae850cda9ee7b6aae146d..c3749419b98da34f3885166248a620c2a68a789f 100644
--- a/components/operators/src/depth.cpp
+++ b/components/operators/src/depth.cpp
@@ -140,13 +140,13 @@ void DepthChannel::_createPipeline(size_t size) {
 
 	pipe_->append<ftl::operators::ColourChannels>("colour");  // Convert BGR to BGRA
 	pipe_->append<ftl::operators::CrossSupport>("cross");
+	#ifdef HAVE_OPTFLOW
+	pipe_->append<ftl::operators::NVOpticalFlow>("optflow", Channel::Colour, Channel::Flow, Channel::Colour2, Channel::Flow2);
+	//if (size == 1) pipe_->append<ftl::operators::OpticalFlowTemporalSmoothing>("optflow_filter", Channel::Disparity);
+	#endif
 	#ifdef HAVE_LIBSGM
 	pipe_->append<ftl::operators::FixstarsSGM>("algorithm");
 	#endif
-	#ifdef HAVE_OPTFLOW
-	pipe_->append<ftl::operators::NVOpticalFlow>("optflow", Channel::Colour, Channel::Flow);
-	if (size == 1) pipe_->append<ftl::operators::OpticalFlowTemporalSmoothing>("optflow_filter", Channel::Disparity);
-	#endif
 	pipe_->append<ftl::operators::DisparityBilateralFilter>("bilateral_filter");
 	//pipe_->append<ftl::operators::OpticalFlowTemporalSmoothing>("optflow_filter", Channel::Disparity);
 	pipe_->append<ftl::operators::DisparityToDepth>("calculate_depth");
diff --git a/components/operators/src/disparity/disp2depth.cu b/components/operators/src/disparity/disp2depth.cu
index bed3bc3e88dbeb83fdbcd35b78ec31abad76c019..b8f878ac08eee93116aeea0d7d3283a118adb480 100644
--- a/components/operators/src/disparity/disp2depth.cu
+++ b/components/operators/src/disparity/disp2depth.cu
@@ -3,6 +3,7 @@
 #include <opencv2/core/cuda_stream_accessor.hpp>
 #include <ftl/operators/cuda/disparity.hpp>
 #include <ftl/operators/cuda/mask.hpp>
+#include <ftl/cuda/fixed.hpp>
 
 #ifndef PINF
 #define PINF __int_as_float(0x7f800000)
@@ -208,6 +209,34 @@ void ftl::cuda::show_rpe(const cv::cuda::GpuMat &disp, cv::cuda::GpuMat &left, c
 // =============================================================================
 
 
+__global__ void merge_disp_kernel(cv::cuda::PtrStepSz<short> disp,
+	cv::cuda::PtrStepSz<short> estimate)
+{
+	for (STRIDE_Y(v,disp.rows)) {
+	for (STRIDE_X(u,disp.cols)) {
+		short cd = disp(v,u);
+		float d = fixed2float<4>((cd >= 4096) ? 0 : cd);
+		float e = fixed2float<4>(estimate(v,u));
+
+		if (e == 0.0f) d = 0.0f;
+		if (fabsf(d-e) > 4.0f) d = 0.0f;
+		disp(v,u) = float2fixed<4>(d);
+	}
+	}
+}
+
+void ftl::cuda::merge_disparities(cv::cuda::GpuMat &disp, const cv::cuda::GpuMat &estimate, cudaStream_t stream) {
+	dim3 grid(1,1,1);
+	dim3 threads(128, 4, 1);
+	grid.x = cv::cuda::device::divUp(disp.cols, 128);
+	grid.y = cv::cuda::device::divUp(disp.rows, 4);
+	merge_disp_kernel<<<grid, threads, 0, stream>>>(disp, estimate);
+	cudaSafeCall( cudaGetLastError() );
+}
+
+// =============================================================================
+
+
 template <int MAX_DISP>
 __global__ void show_disp_density_kernel(cv::cuda::PtrStepSz<short> disp,
 	cv::cuda::PtrStepSz<uchar4> left,
diff --git a/components/operators/src/disparity/fixstars_sgm.cpp b/components/operators/src/disparity/fixstars_sgm.cpp
index b55a7af6d7f7c4b006f1e8988aa4369be7db82e9..0a8bdf0dd016c2137286b06bff78dd6e72a19560 100644
--- a/components/operators/src/disparity/fixstars_sgm.cpp
+++ b/components/operators/src/disparity/fixstars_sgm.cpp
@@ -155,7 +155,8 @@ bool FixstarsSGM::apply(Frame &in, Frame &out, cudaStream_t stream) {
 		if (!init()) { return false; }
 	}
 
-	auto &disp = out.create<GpuMat>(Channel::Disparity, Format<short>(l.size()));
+	bool has_estimate = in.hasChannel(Channel::Disparity);
+	auto &disp = (!has_estimate) ? out.create<GpuMat>(Channel::Disparity, Format<short>(l.size())) : in.get<GpuMat>(Channel::Disparity);
 
 	auto cvstream = cv::cuda::StreamAccessor::wrapStream(stream);
 	cv::cuda::cvtColor(l, lbw_, cv::COLOR_BGRA2GRAY, 0, cvstream);
@@ -169,6 +170,10 @@ bool FixstarsSGM::apply(Frame &in, Frame &out, cudaStream_t stream) {
 	// GpuMat left_pixels(dispt_, cv::Rect(0, 0, max_disp_, dispt_.rows));
 	// left_pixels.setTo(0);
 
+	if (config()->value("merge_estimates", false) && has_estimate) {
+		ftl::cuda::merge_disparities(disp_int_, disp, stream);
+	}
+
 	cv::cuda::threshold(disp_int_, disp, 4096.0f, 0.0f, cv::THRESH_TOZERO_INV, cvstream);
 
 	if (config()->value("check_reprojection", false)) {
diff --git a/components/operators/src/disparity/optflow_smoothing.cu b/components/operators/src/disparity/optflow_smoothing.cu
index c30302b85ac218d464624e6fbb6c740b95a304ee..c85c22cd2fdf9027092a59cfb4617e9e222b552c 100644
--- a/components/operators/src/disparity/optflow_smoothing.cu
+++ b/components/operators/src/disparity/optflow_smoothing.cu
@@ -50,7 +50,7 @@ __global__ void temporal_median_filter_kernel(
 
 		float t = area * threshold + 0.25f;  // 0.25 is the 1/4 pixel accuracy NVIDIA claim
 
-		if (max(abs(flow.x),abs(flow.y)) > makeFixed<5>(t))
+		if (sqrt(float(flow.x*flow.x) + float(flow.y*flow.y)) > makeFixed<5>(t))  // max(abs(flow.x),abs(flow.y))
 		{
 			// TODO: Perhaps rather than discard it could follow the optical flow
 			// This would require the generation of a depth flow also.
@@ -73,7 +73,7 @@ __global__ void temporal_median_filter_kernel(
 
 		const float disparity = disp(y, x);
 
-		if (isValidDisparity(disparity))
+		//if (isValidDisparity(disparity))
 		{
 			history(y, (x + 1) * n_max - 1) += 1.0;
 			count++;
diff --git a/components/operators/src/nvopticalflow.cpp b/components/operators/src/nvopticalflow.cpp
index 422fc3040d27edd0ee49e0e37f62e1185a1d872e..2426c10197a165b02929f444544cdc338de98cf7 100644
--- a/components/operators/src/nvopticalflow.cpp
+++ b/components/operators/src/nvopticalflow.cpp
@@ -1,5 +1,8 @@
 #include <ftl/operators/opticalflow.hpp>
 #include <ftl/exception.hpp>
+#include <ftl/operators/cuda/disparity.hpp>
+
+#include "opticalflow_cuda.hpp"
 
 #include <opencv2/cudaimgproc.hpp>
 
@@ -16,14 +19,17 @@ using cv::Size;
 using cv::cuda::GpuMat;
 
 NVOpticalFlow::NVOpticalFlow(ftl::Configurable* cfg) :
-		ftl::operators::Operator(cfg), channel_in_(ftl::codecs::Channel::Colour), channel_out_(ftl::codecs::Channel::Flow) {
+		ftl::operators::Operator(cfg), channel_in_{ftl::codecs::Channel::Colour,ftl::codecs::Channel::Colour2}, channel_out_{ftl::codecs::Channel::Flow,ftl::codecs::Channel::Flow2} {
 	size_ = Size(0, 0);
 
 }
 
-NVOpticalFlow::NVOpticalFlow(ftl::Configurable*cfg, const std::tuple<ftl::codecs::Channel,ftl::codecs::Channel> &channels) : ftl::operators::Operator(cfg) {
-	channel_in_ = std::get<0>(channels);
-	channel_out_ = std::get<1>(channels);
+NVOpticalFlow::NVOpticalFlow(ftl::Configurable*cfg, const std::tuple<ftl::codecs::Channel,ftl::codecs::Channel,ftl::codecs::Channel,ftl::codecs::Channel> &channels) : ftl::operators::Operator(cfg) {
+	channel_in_[0] = std::get<0>(channels);
+	channel_out_[0] = std::get<1>(channels);
+	channel_in_[1] = std::get<2>(channels);
+	channel_out_[1] = std::get<3>(channels);
+	size_ = Size(0, 0);
 }
 
 NVOpticalFlow::~NVOpticalFlow() {
@@ -43,27 +49,65 @@ bool NVOpticalFlow::init() {
 	
 	left_gray_.create(size_, CV_8UC1);
 	left_gray_prev_.create(size_, CV_8UC1);
+	right_gray_.create(size_, CV_8UC1);
+	right_gray_prev_.create(size_, CV_8UC1);
 	return true;
 }
 
 bool NVOpticalFlow::apply(Frame &in, Frame &out, cudaStream_t stream) {
-	if (!in.hasChannel(channel_in_)) return false;
-	if (in.hasChannel(channel_out_)) return true;
+	bool both_channels = config()->value("both_channels", false);
+	float scale = config()->value("viz_scale", 200.0f);
+
+	if (!in.hasChannel(channel_in_[0])) return false;
+	if (in.hasChannel(channel_out_[0])) return true;
+	if (both_channels) {
+		if (!in.hasChannel(channel_in_[1])) return false;
+		if (in.hasChannel(channel_out_[1])) return true;
+	}
 
-	if (in.get<GpuMat>(channel_in_).size() != size_) {
-		size_ = in.get<GpuMat>(channel_in_).size();
+	if (in.get<GpuMat>(channel_in_[0]).size() != size_) {
+		size_ = in.get<GpuMat>(channel_in_[0]).size();
 		if (!init()) return false;
 	}
 	
 	auto cvstream = cv::cuda::StreamAccessor::wrapStream(stream);
-	auto &flow = out.create<GpuMat>(channel_out_);
+	auto &flow1 = out.create<GpuMat>(channel_out_[0]);
 
-	cv::cuda::cvtColor(in.get<GpuMat>(channel_in_), left_gray_, cv::COLOR_BGRA2GRAY, 0, cvstream);
+	cv::cuda::cvtColor(in.get<GpuMat>(channel_in_[0]), left_gray_, cv::COLOR_BGRA2GRAY, 0, cvstream);
+	cv::cuda::cvtColor(in.get<GpuMat>(channel_in_[1]), right_gray_, cv::COLOR_BGRA2GRAY, 0, cvstream);
 
 	// TODO: Use optical flow confidence output, perhaps combined with a
 	// sensitivity adjustment
-	nvof_->calc(left_gray_, left_gray_prev_, flow, cvstream);
-	std::swap(left_gray_, left_gray_prev_);
+	nvof_->calc(left_gray_, right_gray_, flow1, cvstream);
+	//std::swap(left_gray_, left_gray_prev_);
+
+	if (both_channels) {
+		auto &flow2 = out.create<GpuMat>(channel_out_[1]);
+		nvof_->calc(right_gray_, left_gray_, flow2, cvstream);
+		//std::swap(right_gray_, right_gray_prev_);
+	}
+
+	if (config()->value("show_flow", false)) {
+		ftl::cuda::show_optical_flow(flow1, in.getTexture<uchar4>(channel_in_[0]), scale, stream);
+		if (both_channels)
+			ftl::cuda::show_optical_flow(out.get<GpuMat>(channel_out_[1]), in.getTexture<uchar4>(channel_in_[1]), scale, stream);
+	}
+
+	if (config()->value("generate_disparity", false)) {
+		if (!in.hasChannel(Channel::Disparity)) {
+			in.create<GpuMat>(Channel::Disparity).create(size_, CV_16SC1);
+		}
+		ftl::cuda::disparity_from_flow(
+			flow1,
+			out.get<GpuMat>(channel_out_[1]),
+			in.createTexture<short>(Channel::Disparity), stream);
+	}
+
+	if (config()->value("check_reprojection", false)) {
+		ftl::cuda::check_reprojection(in.get<cv::cuda::GpuMat>(Channel::Disparity), in.getTexture<uchar4>(Channel::Colour),
+			in.createTexture<uchar4>(Channel::Colour2, true),
+			stream);
+	}
 
 	return true;
 }
\ No newline at end of file
diff --git a/components/operators/src/opticalflow.cu b/components/operators/src/opticalflow.cu
new file mode 100644
index 0000000000000000000000000000000000000000..0876fe1fa627cebc6c2e2d7842351b9558972af7
--- /dev/null
+++ b/components/operators/src/opticalflow.cu
@@ -0,0 +1,102 @@
+#include "opticalflow_cuda.hpp"
+#include <ftl/cuda/fixed.hpp>
+
+#define T_PER_BLOCK 8
+
+using ftl::cuda::TextureObject;
+
+__global__ void show_optflow_kernel(
+        cv::cuda::PtrStepSz<short2> optflow,
+        TextureObject<uchar4> colour, float scale) {
+
+    const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
+    const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
+
+    if (x < colour.width() && y < colour.height()) {
+        short2 flow = optflow(y/4, x/4);  // 4 is granularity
+
+        float fx = max(-1.0f, min(1.0f, fixed2float<5>(flow.x) / scale));
+        float fy = max(-1.0f, min(1.0f, fixed2float<5>(flow.y) / scale));
+        float f = sqrt(fx*fx+fy*fy);
+
+        float4 c = colour.tex2D(float(x)+0.5f, float(y)+0.5f);
+        c += make_float4(f*255.0f, 0.0f, f*255.0f, 0.0f);
+        colour(x,y) = make_uchar4(min(255.0f, c.x), min(255.0f, c.y), min(255.0f, c.z), 0.0f);
+    }
+}
+
+void ftl::cuda::show_optical_flow(const cv::cuda::GpuMat &optflow, const TextureObject<uchar4> &colour, float scale, cudaStream_t stream) {
+    const dim3 gridSize((colour.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (colour.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
+    const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
+
+    show_optflow_kernel<<<gridSize, blockSize, 0, stream>>>(optflow, colour, scale);
+    cudaSafeCall( cudaGetLastError() );
+}
+
+// ==== Flow difference ========================================================
+
+__global__ void gen_disparity_kernel(
+        cv::cuda::PtrStepSz<short2> optflow1,
+        cv::cuda::PtrStepSz<short2> optflow2,
+        TextureObject<short> disparity) {
+
+    const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
+    const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
+
+    if (x < disparity.width() && y < disparity.height()) {
+        short2 flow1 = optflow1(y/4, x/4);  // 4 is granularity
+        float disp = -fixed2float<5>(flow1.x);
+
+        // Do a consistency check
+        if (disp > 0.0f && x-disp-0.5f > 0) { //< colour.width()) {
+            short2 flow2 = optflow2(y/4, (x-round(disp))/4);  // 4 is granularity
+            if (fabsf(disp - fixed2float<5>(flow2.x)) > 1.0f) disp = 0.0f;
+        }
+
+        disparity(x,y) = float2fixed<4>(disp);
+    }
+}
+
+void ftl::cuda::disparity_from_flow(const cv::cuda::GpuMat &optflow1, const cv::cuda::GpuMat &optflow2, const TextureObject<short> &disparity, cudaStream_t stream) {
+    const dim3 gridSize((disparity.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (disparity.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
+    const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
+
+    gen_disparity_kernel<<<gridSize, blockSize, 0, stream>>>(optflow1, optflow2, disparity);
+    cudaSafeCall( cudaGetLastError() );
+}
+
+__global__ void show_optflow_diff_kernel(
+        cv::cuda::PtrStepSz<short2> optflow1,
+        cv::cuda::PtrStepSz<short2> optflow2,
+        TextureObject<uchar4> colour, float scale,
+        ftl::rgbd::Camera cam) {
+
+    const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
+    const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
+
+    if (x < colour.width() && y < colour.height()) {
+        short2 flow1 = optflow1(y/4, x/4);  // 4 is granularity
+        float disp = fixed2float<5>(flow1.x);
+
+        if (disp > 0.0f && x-disp-0.5f > 0) { //< colour.width()) {
+            short2 flow2 = optflow2(y/4, (x-round(disp))/4);  // 4 is granularity
+
+            float dx = (fixed2float<5>(flow1.x) + fixed2float<5>(flow2.x)) / disp;
+            float fR = max(0.0f, min(1.0f, dx / scale));
+            float fB = -max(-1.0f, min(0.0f, dx / scale));
+            float fG = (fR == 1.0f || fB == 1.0f) ? 0.0f : 1.0f;
+
+            float4 c = colour.tex2D(float(x)+0.5f, float(y)+0.5f);
+            c += make_float4(fG*fB*255.0f, (1.0f-fG)*255.0f, fG*fR*255.0f, 0.0f);
+            colour(x,y) = make_uchar4(min(255.0f, c.x), min(255.0f, c.y), min(255.0f, c.z), 0.0f);
+        }
+    }
+}
+
+void ftl::cuda::show_optical_flow_diff(const cv::cuda::GpuMat &optflow1, const cv::cuda::GpuMat &optflow2, const TextureObject<uchar4> &colour, float scale, const ftl::rgbd::Camera &cam, cudaStream_t stream) {
+    const dim3 gridSize((colour.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (colour.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
+    const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
+
+    show_optflow_diff_kernel<<<gridSize, blockSize, 0, stream>>>(optflow1, optflow2, colour, scale, cam);
+    cudaSafeCall( cudaGetLastError() );
+}
diff --git a/components/operators/src/opticalflow_cuda.hpp b/components/operators/src/opticalflow_cuda.hpp
new file mode 100644
index 0000000000000000000000000000000000000000..f5b22b75d780cebbfe4af49e64ab138b2de00f63
--- /dev/null
+++ b/components/operators/src/opticalflow_cuda.hpp
@@ -0,0 +1,19 @@
+#ifndef _FTL_OPTFLOW_CUDA_HPP_
+#define _FTL_OPTFLOW_CUDA_HPP_
+
+#include <ftl/cuda_common.hpp>
+#include <ftl/rgbd/camera.hpp>
+
+namespace ftl {
+namespace cuda {
+
+void show_optical_flow(const cv::cuda::GpuMat &optflow, const ftl::cuda::TextureObject<uchar4> &colour, float scale, cudaStream_t);
+
+void disparity_from_flow(const cv::cuda::GpuMat &optflow1, const cv::cuda::GpuMat &optflow2, const ftl::cuda::TextureObject<short> &disparity, cudaStream_t stream);
+
+void show_optical_flow_diff(const cv::cuda::GpuMat &optflow1, const cv::cuda::GpuMat &optflow2, const ftl::cuda::TextureObject<uchar4> &colour, float scale, const ftl::rgbd::Camera &cam, cudaStream_t stream);
+
+}
+}
+
+#endif
diff --git a/components/rgbd-sources/src/sources/stereovideo/stereovideo.cpp b/components/rgbd-sources/src/sources/stereovideo/stereovideo.cpp
index 4ddc992b63f6ac30e1aeb34bbfdb7f04d85385ab..8f8297c588f59bef62ab3d5b711be531e972beb4 100644
--- a/components/rgbd-sources/src/sources/stereovideo/stereovideo.cpp
+++ b/components/rgbd-sources/src/sources/stereovideo/stereovideo.cpp
@@ -79,7 +79,7 @@ void StereoVideoSource::init(const string &file) {
 
 	pipeline_input_ = ftl::config::create<ftl::operators::Graph>(host_, "input");
 	#ifdef HAVE_OPTFLOW
-	pipeline_input_->append<ftl::operators::NVOpticalFlow>("optflow", Channel::Colour, Channel::Flow);
+	pipeline_input_->append<ftl::operators::NVOpticalFlow>("optflow");
 	#endif
 	pipeline_input_->append<ftl::operators::ColourChannels>("colour");
 
diff --git a/lib/libsgm/src/horizontal_path_aggregation.cu b/lib/libsgm/src/horizontal_path_aggregation.cu
index b3772b6f64ddf16c7dfe70b3c2c2ea9aaf0ecf18..860bb3d69a77924bf6953b42a7be8f297a6e71a9 100644
--- a/lib/libsgm/src/horizontal_path_aggregation.cu
+++ b/lib/libsgm/src/horizontal_path_aggregation.cu
@@ -21,7 +21,7 @@ limitations under the License.
 namespace sgm {
 namespace path_aggregation {
 
-static constexpr unsigned int DP_BLOCK_SIZE = 8u;
+static constexpr unsigned int DP_BLOCK_SIZE = 16u;
 static constexpr unsigned int DP_BLOCKS_PER_THREAD = 1u;
 
 static constexpr unsigned int WARPS_PER_BLOCK = 4u;