Fix bug in high res colour scaling

applications/gui/src/src_window.cpp

@@ -220,7 +220,7 @@ bool SourceWindow::_processFrameset(ftl::rgbd::FrameSet &fs, bool fromstream) {
 void SourceWindow::_checkFrameSets(int id) {
 	while (framesets_.size() <= id) {
 		auto *p = ftl::config::create<ftl::operators::Graph>(screen_->root(), "pre_filters");
-		//pre_pipeline_->append<ftl::operators::ColourChannels>("colour");  // Convert BGR to BGRA
+		//p->append<ftl::operators::ColourChannels>("colour");  // Convert BGR to BGRA
 		p->append<ftl::operators::DetectAndTrack>("facedetection")->value("enabled", false);
 		p->append<ftl::operators::ArUco>("aruco")->value("enabled", false);
 		//p->append<ftl::operators::HFSmoother>("hfnoise");

components/common/cpp/include/ftl/cuda_common.hpp

@@ -380,36 +380,17 @@ TextureObject<T>::~TextureObject() {
 	free();
 }
 
-/*template <>
-void TextureObject<uchar4>::upload(const cv::Mat &m, cudaStream_t stream);
-
-template <>
-void TextureObject<float>::upload(const cv::Mat &m, cudaStream_t stream);
-
-template <>
-void TextureObject<float2>::upload(const cv::Mat &m, cudaStream_t stream);
-
-template <>
-void TextureObject<float4>::upload(const cv::Mat &m, cudaStream_t stream);
-
-template <>
-void TextureObject<uchar>::upload(const cv::Mat &m, cudaStream_t stream);
-
-
-template <>
-void TextureObject<uchar4>::download(cv::Mat &m, cudaStream_t stream) const;
-
-template <>
-void TextureObject<float>::download(cv::Mat &m, cudaStream_t stream) const;
-
-template <>
-void TextureObject<float2>::download(cv::Mat &m, cudaStream_t stream) const;
-
-template <>
-void TextureObject<float4>::download(cv::Mat &m, cudaStream_t stream) const;
-
-template <>
-void TextureObject<uchar>::download(cv::Mat &m, cudaStream_t stream) const;*/
+/**
+ * Read a texture value using coordinates in the range of `b`, but from the
+ * texture `a` which may have a different resolution.
+ */
+template <typename A, typename B>
+__device__ inline A getScaledTex2D(int x, int y, ftl::cuda::TextureObject<A> &a, ftl::cuda::TextureObject<B> &b) {
+	return a.tex2D(
+		(int)((float)x * ((float)a.width() / (float)b.width())),
+		(int)((float)y * ((float)a.height() / (float)b.height()))
+	);
+}
 
 }
 }

components/operators/src/mask.cu

@@ -84,7 +84,8 @@ __global__ void discontinuity_kernel(ftl::cuda::TextureObject<uint8_t> mask_out,
 			const float g = max(g1,max(g2,(max(g3,g4))));
 
 			// Calculate support window area
-			const uchar4 sup = support.tex2D((int)x, (int)y);
+			//const uchar4 sup = support.tex2D((int)x, (int)y);
+			const uchar4 sup = getScaledTex2D(x, y, support, depth);
 			const float supx = min(sup.x,sup.y);
 			const float supy = min(sup.z,sup.w);
 			const float area = supx * supy;

components/renderers/cpp/src/mask.cu

@@ -3,6 +3,7 @@
 
 using ftl::cuda::TextureObject;
 using ftl::cuda::Mask;
+using ftl::cuda::getScaledTex2D;
 
 #define T_PER_BLOCK 16
 
@@ -15,10 +16,11 @@ __global__ void show_mask_kernel(
 	const int y = blockIdx.y*blockDim.y + threadIdx.y;
 
 	if (x >= 0 && x < colour.width() && y >=0 && y < colour.height()) {
-        float xscale = (float)x * ((float)mask.width() / (float)colour.width());
-        float yscale = (float)y * ((float)mask.height() / (float)colour.height());
+        //float xscale = (float)x * ((float)mask.width() / (float)colour.width());
+        //float yscale = (float)y * ((float)mask.height() / (float)colour.height());
 
-        Mask m(mask.tex2D((int)xscale,(int)yscale));
+        //Mask m(mask.tex2D((int)xscale,(int)yscale));
+        Mask m(getScaledTex2D(x, y, mask, colour));
 
         if (m.is(id)) {
             colour(x,y) = style;