Allow config of back point cull

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

@@ -44,6 +44,7 @@ class Splatter : public ftl::render::Renderer {
 	ftl::cuda::ClipSpace clip_;
 	bool clipping_;
 	float norm_filter_;
+	bool backcull_;
 };
 
 }

components/renderers/cpp/src/splat_render.cpp

@@ -54,6 +54,11 @@ Splatter::Splatter(nlohmann::json &config, ftl::rgbd::FrameSet *fs) : ftl::rende
 	on("normal_filter", [this](const ftl::config::Event &e) {
 		norm_filter_ = value("normal_filter", -1.0f);
 	});
+
+	backcull_ = value("back_cull", true);
+	on("back_cull", [this](const ftl::config::Event &e) {
+		backcull_ = value("back_cull", true);
+	});
 }
 
 Splatter::~Splatter() {
@@ -90,7 +95,7 @@ void Splatter::renderChannel(
 			f.createTexture<float4>(Channel::Points),
 			f.createTexture<float4>(Channel::Normals),
 			temp_.getTexture<int>(Channel::Depth),
-			params, stream
+			params, backcull_, stream
 		);
 
 		//LOG(INFO) << "DIBR DONE";

components/renderers/cpp/src/splatter.cu

@@ -20,6 +20,7 @@ using ftl::render::SplatParams;
  * Pass 1: Directly render each camera into virtual view but with no upsampling
  * for sparse points.
  */
+ template <bool CULLING>
  __global__ void dibr_merge_kernel(TextureObject<float4> points,
 		TextureObject<float4> normals,
 		TextureObject<int> depth, SplatParams params) {
@@ -29,17 +30,20 @@ using ftl::render::SplatParams;
 	const float4 worldPos = points.tex2D(x, y);
 	if (worldPos.x == MINF || (!(params.m_flags & ftl::render::kShowDisconMask) && worldPos.w < 0.0f)) return;
 
-	float3 ray = params.m_viewMatrixInverse.getFloat3x3() * params.camera.screenToCam(x,y,1.0f);
-    ray = ray / length(ray);
-    float3 n = make_float3(normals.tex2D((int)x,(int)y));
-    float l = length(n);
-    if (l == 0) {
-        return;
-    }
-    n /= l;
-
-	const float facing = dot(ray, n);
-	if (facing <= 0.0f) return;
+	// Compile time enable/disable of culling back facing points
+	if (CULLING) {
+		float3 ray = params.m_viewMatrixInverse.getFloat3x3() * params.camera.screenToCam(x,y,1.0f);
+		ray = ray / length(ray);
+		float3 n = make_float3(normals.tex2D((int)x,(int)y));
+		float l = length(n);
+		if (l == 0) {
+			return;
+		}
+		n /= l;
+
+		const float facing = dot(ray, n);
+		if (facing <= 0.0f) return;
+	}
 
     // Find the virtual screen position of current point
 	const float3 camPos = params.m_viewMatrix * make_float3(worldPos);
@@ -57,11 +61,12 @@ using ftl::render::SplatParams;
 	}
 }
 
-void ftl::cuda::dibr_merge(TextureObject<float4> &points, TextureObject<float4> &normals, TextureObject<int> &depth, SplatParams params, cudaStream_t stream) {
+void ftl::cuda::dibr_merge(TextureObject<float4> &points, TextureObject<float4> &normals, TextureObject<int> &depth, SplatParams params, bool culling, cudaStream_t stream) {
     const dim3 gridSize((depth.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (depth.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
     const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
 
-    dibr_merge_kernel<<<gridSize, blockSize, 0, stream>>>(points, normals, depth, params);
+	if (culling) dibr_merge_kernel<true><<<gridSize, blockSize, 0, stream>>>(points, normals, depth, params);
+	else dibr_merge_kernel<false><<<gridSize, blockSize, 0, stream>>>(points, normals, depth, params);
     cudaSafeCall( cudaGetLastError() );
 }
 

components/renderers/cpp/src/splatter_cuda.hpp

@@ -11,6 +11,7 @@ namespace cuda {
 		ftl::cuda::TextureObject<float4> &normals,
 		ftl::cuda::TextureObject<int> &depth,
 		ftl::render::SplatParams params,
+		bool culling,
 		cudaStream_t stream);
 
 	void dibr_attribute(