Allow both raycast and hash based renderer

applications/reconstruct/include/ftl/ray_cast_sdf.hpp

 #pragma once
 
+#include <ftl/configurable.hpp>
 #include <ftl/matrix_conversion.hpp>
 #include <ftl/cuda_matrix_util.hpp>
 #include <ftl/depth_camera.hpp>
@@ -8,23 +9,24 @@
 
 // #include "DX11RayIntervalSplatting.h"
 
-class CUDARayCastSDF
+class CUDARayCastSDF : public ftl::Configurable
 {
 public:
-	CUDARayCastSDF(const RayCastParams& params) {
-		create(params);
+	CUDARayCastSDF(nlohmann::json& config) : ftl::Configurable(config) {
+		create(parametersFromConfig(config));
+		hash_render_ = config.value("hash_renderer", false);
 	}
 
 	~CUDARayCastSDF(void) {
 		destroy();
 	}
 
-	static RayCastParams parametersFromConfig(const nlohmann::json& gas, const Eigen::Matrix4f& intrinsics, const Eigen::Matrix4f& intrinsicsInv) {
+	static RayCastParams parametersFromConfig(const nlohmann::json& gas) {
 		RayCastParams params;
 		params.m_width = gas["adapterWidth"].get<unsigned int>();
 		params.m_height = gas["adapterHeight"].get<unsigned int>();
-		params.m_intrinsics = MatrixConversion::toCUDA(intrinsics);
-		params.m_intrinsicsInverse = MatrixConversion::toCUDA(intrinsicsInv);
+		params.m_intrinsics = MatrixConversion::toCUDA(Eigen::Matrix4f());
+		params.m_intrinsicsInverse = MatrixConversion::toCUDA(Eigen::Matrix4f());
 		params.m_minDepth = gas["sensorDepthMin"].get<float>();
 		params.m_maxDepth = gas["sensorDepthMax"].get<float>();
 		params.m_rayIncrement = gas["SDFRayIncrementFactor"].get<float>() * gas["SDFTruncation"].get<float>();
@@ -59,6 +61,7 @@ private:
 
 	RayCastParams m_params;
 	RayCastData m_data;
+	bool hash_render_;
 
 	// DX11RayIntervalSplatting m_rayIntervalSplatting;
 };

applications/reconstruct/src/main.cpp

@@ -388,16 +388,15 @@ static void run() {
 	vector<PointCloud<PointXYZRGB>::Ptr> clouds(inputs.size());
 	Display display_merged(config["display"], "Merged"); // todo
 
-	auto rayparams = CUDARayCastSDF::parametersFromConfig(config["voxelhash"], Eigen::Matrix4f(), Eigen::Matrix4f());
-	CUDARayCastSDF rays(rayparams);
+	CUDARayCastSDF rays(config["voxelhash"]);
 
 	LOG(INFO) << "About to create scene";
 	ftl::voxhash::SceneRep scene(config["voxelhash"]);
 	LOG(INFO) << "Scene created";
 
-	float3 *point_array = new float3[rayparams.m_width*rayparams.m_height];
+	float3 *point_array = new float3[rays.getRayCastParams().m_width*rays.getRayCastParams().m_height];
 	//uchar3 *colour_array = new uchar3[rayparams.m_width*rayparams.m_height];
-	cv::Mat colour_array(cv::Size(rayparams.m_width,rayparams.m_height), CV_8UC3);
+	cv::Mat colour_array(cv::Size(rays.getRayCastParams().m_width,rays.getRayCastParams().m_height), CV_8UC3);
 
 	float bounce = 0.0;
 	int bounce_dir = 1;
@@ -478,7 +477,7 @@ static void run() {
 		rays.render(scene.getHashData(), scene.getHashParams(), inputs[0].gpu, viewPose);
 		//LOG(INFO) << "Download points";
 		
-		rays.getRayCastData().download(point_array, (uchar3*)colour_array.data, rayparams);
+		rays.getRayCastData().download(point_array, (uchar3*)colour_array.data, rays.getRayCastParams());
 
 		int pc = 0;
 

applications/reconstruct/src/ray_cast_sdf.cpp

@@ -43,8 +43,8 @@ void CUDARayCastSDF::render(const ftl::voxhash::HashData& hashData, const ftl::v
 	//m_data.d_rayIntervalSplatMinArray = m_rayIntervalSplatting.mapMinToCuda();
 	//m_data.d_rayIntervalSplatMaxArray = m_rayIntervalSplatting.mapMaxToCuda();
 
-	//renderCS(hashData, m_data, cameraData, m_params);
-	nickRenderCUDA(hashData, hashParams, m_data, cameraData, m_params);
+	if (hash_render_) nickRenderCUDA(hashData, hashParams, m_data, cameraData, m_params);
+	else renderCS(hashData, m_data, cameraData, m_params);
 
 	//convertToCameraSpace(cameraData);
 	if (!m_params.m_useGradients)

applications/reconstruct/src/ray_cast_sdf.cu

@@ -153,6 +153,22 @@ __device__ inline uint blockLinear(int x, int y, int z) {
 	return x + (y << 1) + (z << 2);
 }
 
+__device__ inline void trilinearInterp(const ftl::voxhash::HashData &hashData, const ftl::voxhash::Voxel *voxels, const uint *ix, const float3 &pos, float &depth, uchar3 &colour) {
+	float3 colorFloat = make_float3(0.0f, 0.0f, 0.0f);
+	const float3 weight = frac(hashData.worldToVirtualVoxelPosFloat(pos));  // Should be world position of ray, not voxel??
+	float dist = 0.0f;
+	dist+= (1.0f-weight.x)*(1.0f-weight.y)*(1.0f-weight.z)*voxels[ix[0]].sdf; colorFloat+= (1.0f-weight.x)*(1.0f-weight.y)*(1.0f-weight.z)*make_float3(voxels[ix[0]].color);
+	dist+=	   weight.x *(1.0f-weight.y)*(1.0f-weight.z)*voxels[ix[1]].sdf; colorFloat+=	   weight.x *(1.0f-weight.y)*(1.0f-weight.z)*make_float3(voxels[ix[1]].color);
+	dist+= (1.0f-weight.x)*	   weight.y *(1.0f-weight.z)*voxels[ix[2]].sdf; colorFloat+= (1.0f-weight.x)*	   weight.y *(1.0f-weight.z)*make_float3(voxels[ix[2]].color);
+	dist+= (1.0f-weight.x)*(1.0f-weight.y)*	   weight.z *voxels[ix[3]].sdf; colorFloat+= (1.0f-weight.x)*(1.0f-weight.y)*	   weight.z *make_float3(voxels[ix[3]].color);
+	dist+=	   weight.x *	   weight.y *(1.0f-weight.z)*voxels[ix[4]].sdf; colorFloat+=	   weight.x *	   weight.y *(1.0f-weight.z)*make_float3(voxels[ix[4]].color);
+	dist+= (1.0f-weight.x)*	   weight.y *	   weight.z *voxels[ix[5]].sdf; colorFloat+= (1.0f-weight.x)*	   weight.y *	   weight.z *make_float3(voxels[ix[5]].color);
+	dist+=	   weight.x *(1.0f-weight.y)*	   weight.z *voxels[ix[6]].sdf; colorFloat+=	   weight.x *(1.0f-weight.y)*	   weight.z *make_float3(voxels[ix[6]].color);
+	dist+=	   weight.x *	   weight.y *	   weight.z *voxels[ix[7]].sdf; colorFloat+=	   weight.x *	   weight.y *	   weight.z *make_float3(voxels[ix[7]].color);
+	depth = dist;
+	colour = make_uchar3(colorFloat);
+}
+
 __global__ void nickRenderKernel(ftl::voxhash::HashData hashData, RayCastData rayCastData, DepthCameraData cameraData, RayCastParams params) {
 	// TODO(Nick) Reduce bank conflicts by aligning these
 	__shared__ ftl::voxhash::Voxel voxels[SDF_BLOCK_BUFFER];
@@ -213,8 +229,8 @@ __global__ void nickRenderKernel(ftl::voxhash::HashData hashData, RayCastData ra
 	if (missweight) return;
 
 	// Trilinear Interpolation (simple and fast)
-	float3 colorFloat = make_float3(0.0f, 0.0f, 0.0f);
-	const float3 weight = frac(hashData.worldToVirtualVoxelPosFloat(worldPos));
+	/*float3 colorFloat = make_float3(0.0f, 0.0f, 0.0f);
+	const float3 weight = frac(hashData.worldToVirtualVoxelPosFloat(worldPos));  // Should be world position of ray, not voxel??
 	float dist = 0.0f;
 	dist+= (1.0f-weight.x)*(1.0f-weight.y)*(1.0f-weight.z)*voxels[ix[0]].sdf; colorFloat+= (1.0f-weight.x)*(1.0f-weight.y)*(1.0f-weight.z)*make_float3(voxels[ix[0]].color);
 	dist+=	   weight.x *(1.0f-weight.y)*(1.0f-weight.z)*voxels[ix[1]].sdf; colorFloat+=	   weight.x *(1.0f-weight.y)*(1.0f-weight.z)*make_float3(voxels[ix[1]].color);
@@ -228,11 +244,13 @@ __global__ void nickRenderKernel(ftl::voxhash::HashData hashData, RayCastData ra
 	// Must finish using colours before updating colours
 	__syncthreads();
 
-	voxels[j].color = make_uchar3(colorFloat);
-	voxels[j].sdf = dist;
+	//voxels[j].color = make_uchar3(colorFloat);
+	//voxels[j].sdf = dist;
 
 	// What happens if fitlered voxel is put back?
-	hashData.d_SDFBlocks[blocks[0].ptr + i] = voxels[j];
+	//hashData.d_SDFBlocks[blocks[0].ptr + i] = voxels[j];
+
+	//return;*/
 
 	__syncthreads();
 
@@ -260,21 +278,48 @@ __global__ void nickRenderKernel(ftl::voxhash::HashData hashData, RayCastData ra
 	// TODO Could adjust weights, strengthen on surface, weaken otherwise??
 	if (!is_surface) return;
 
-	const float3 camPos = params.m_viewMatrixInverse * worldPos;
-	const uint2 screenPos = make_uint2(cameraData.cameraToKinectScreenInt(camPos));
+	const float3 camPos = params.m_viewMatrix * worldPos;
+	const float2 screenPosf = cameraData.cameraToKinectScreenFloat(camPos);
+	const uint2 screenPos = make_uint2(make_int2(screenPosf)); //  + make_float2(0.5f, 0.5f)
+
+	/*if (screenPos.x < params.m_width && screenPos.y < params.m_height && 
+			rayCastData.d_depth[(screenPos.y)*params.m_width+screenPos.x] > camPos.z) {
+		rayCastData.d_depth[(screenPos.y)*params.m_width+screenPos.x] = camPos.z;
+		rayCastData.d_colors[(screenPos.y)*params.m_width+screenPos.x] = voxels[j].color;
+	}*/
+
+	//return;
 
 	// For this voxel in hash, get its screen position and check it is on screen
-	// TODO (Nick) MAKE THIS ATOMIC!!!! (seems to not matter but...)
 	// Convert depth map to int by x1000 and use atomicMin
-	const int pixsize = static_cast<int>(c_hashParams.m_virtualVoxelSize*c_depthCameraParams.fx/camPos.z)+1;
-	for (int x=0; x<pixsize; x++) {
-		for (int y=0; y<pixsize; y++) {
+	//const int pixsize = static_cast<int>(c_hashParams.m_virtualVoxelSize*c_depthCameraParams.fx/camPos.z)+1;
+	int pixsizeX = 10;  // Max voxel pixels
+	int pixsizeY = 10;
+
+	for (int y=0; y<pixsizeY; y++) {
+		for (int x=0; x<pixsizeX; x++) {
+			// TODO(Nick) Within a window, check pixels that have same voxel id
+			// Then trilinear interpolate between current voxel and neighbors.
+			const float3 pixelWorldPos = params.m_viewMatrixInverse * cameraData.kinectDepthToSkeleton(screenPos.x+x,screenPos.y+y, camPos.z);
+			const float3 posInVoxel = (pixelWorldPos - worldPos) / make_float3(c_hashParams.m_virtualVoxelSize,c_hashParams.m_virtualVoxelSize,c_hashParams.m_virtualVoxelSize);
+
+			if (posInVoxel.x >= 1.0f || posInVoxel.y >= 1.0f || posInVoxel.z >= 1.0f) {
+				pixsizeX = x;
+				continue;
+			}
+
+			/*float depth;
+			uchar3 col;
+			trilinearInterp(hashData, voxels, ix, pixelWorldPos, depth, col);*/
+
+			// TODO (Nick) MAKE THIS ATOMIC!!!!
 			if (screenPos.x+x < params.m_width && screenPos.y+y < params.m_height && 
-					rayCastData.d_depth[(screenPos.y+y)*params.m_width+screenPos.x+x] > camPos.z-0.02f) {  // The 0.02 is a hack to reduce pixalation effects, should be done another way.
+					rayCastData.d_depth[(screenPos.y+y)*params.m_width+screenPos.x+x] > camPos.z) {
 				rayCastData.d_depth[(screenPos.y+y)*params.m_width+screenPos.x+x] = camPos.z;
 				rayCastData.d_colors[(screenPos.y+y)*params.m_width+screenPos.x+x] = voxels[j].color;
 			}
 		}
+		if (pixsizeX == 0) break;
 	}
 }
 

applications/reconstruct/src/scene_rep_hash_sdf.cu

@@ -540,7 +540,7 @@ __global__ void integrateDepthMapKernel(HashData hashData, DepthCameraData camer
 				uint idx = entry.ptr + i;
 				hashData.d_SDFBlocks[idx].weight = 0;
 				//hashData.d_SDFBlocks[idx].sdf = PINF;
-				hashData.d_SDFBlocks[idx].color = make_uchar3(255,0,0);
+				hashData.d_SDFBlocks[idx].color = make_uchar3(0,0,0);
 			}
 			else if (sdf > -truncation) // && depthZeroOne >= 0.0f && depthZeroOne <= 1.0f) //check if in truncation range should already be made in depth map computation
 			{