Implements #181 back facing 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() {
@@ -76,6 +81,7 @@ void Splatter::renderChannel(
 	bool is_4chan = scene_->frames[0].get<GpuMat>(channel).type() == CV_32FC4;
 	
 	// Render each camera into virtual view
+	// TODO: Move out of renderChannel, this is a common step to all channels
 	for (size_t i=0; i < scene_->frames.size(); ++i) {
 		auto &f = scene_->frames[i];
 		auto *s = scene_->sources[i];
@@ -87,8 +93,9 @@ void Splatter::renderChannel(
 
 		ftl::cuda::dibr_merge(
 			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,13 +20,31 @@ 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, SplatParams params) {
+ template <bool CULLING>
+ __global__ void dibr_merge_kernel(TextureObject<float4> points,
+		TextureObject<float4> normals,
+		TextureObject<int> depth, SplatParams params) {
 	const int x = blockIdx.x*blockDim.x + threadIdx.x;
 	const int y = blockIdx.y*blockDim.y + threadIdx.y;
 
 	const float4 worldPos = points.tex2D(x, y);
 	if (worldPos.x == MINF || (!(params.m_flags & ftl::render::kShowDisconMask) && worldPos.w < 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);
 	if (camPos.z < params.camera.minDepth) return;
@@ -43,11 +61,12 @@ using ftl::render::SplatParams;
 	}
 }
 
-void ftl::cuda::dibr_merge(TextureObject<float4> &points, 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, 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

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