Initial working 360 image background

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

@@ -43,6 +43,9 @@ class Triangular : public ftl::render::Renderer {
 	cudaStream_t stream_;
 	float3 light_pos_;
 
+	cv::cuda::GpuMat env_image_;
+	ftl::cuda::TextureObject<uchar4> env_tex_;
+
 	//ftl::Filters *filters_;
 
 	template <typename T>

components/renderers/cpp/src/reprojection.cu

@@ -274,3 +274,51 @@ template void ftl::cuda::reproject(
 		const ftl::render::SplatParams &params,
 		const ftl::rgbd::Camera &camera,
 		const float4x4 &poseInv, cudaStream_t stream);
+
+
+// ===== Equirectangular Reprojection ==========================================
+
+__device__ inline float2 equirect_reprojection(int x_img, int y_img, double f, const float3x3 &rot, int w1, int h1, const ftl::rgbd::Camera &cam) {
+	float3 ray3d = cam.screenToCam(x_img, y_img, 1.0f);
+	ray3d /= length(ray3d);
+	ray3d = rot * ray3d;
+
+    //inverse formula for spherical projection, reference Szeliski book "Computer Vision: Algorithms and Applications" p439.
+    float theta = atan2(ray3d.y,sqrt(ray3d.x*ray3d.x+ray3d.z*ray3d.z));
+	float phi = atan2(ray3d.x, ray3d.z);
+	
+	const float pi = 3.14f;
+
+    //get 2D point on equirectangular map
+    float x_sphere = (((phi*w1)/pi+w1)/2); 
+    float y_sphere = (theta+ pi/2)*h1/pi;
+
+    return make_float2(x_sphere,y_sphere);
+};
+
+__global__ void equirectangular_kernel(
+		TextureObject<uchar4> image_in,
+		TextureObject<uchar4> image_out,
+		Camera camera, float3x3 pose) {
+		
+	const int x = (blockIdx.x*blockDim.x + threadIdx.x);
+	const int y = blockIdx.y*blockDim.y + threadIdx.y;
+
+	if (x >= 0 && y >= 0 && x < image_out.width() && y < image_out.height()) {
+		const float2 p = equirect_reprojection(x,y, camera.fx, pose, image_in.width(), image_in.height(), camera);
+		const float4 colour = image_in.tex2D(p.x, p.y);
+		image_out(x,y) = make_uchar4(colour.x, colour.y, colour.z, 0);
+	}
+}
+
+void ftl::cuda::equirectangular_reproject(
+		ftl::cuda::TextureObject<uchar4> &image_in,
+		ftl::cuda::TextureObject<uchar4> &image_out,
+		const ftl::rgbd::Camera &camera, const float3x3 &pose, cudaStream_t stream) {
+
+	const dim3 gridSize((image_out.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (image_out.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
+	const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
+
+	equirectangular_kernel<<<gridSize, blockSize, 0, stream>>>(image_in, image_out, camera, pose);
+	cudaSafeCall( cudaGetLastError() );
+}

components/renderers/cpp/src/splatter_cuda.hpp

@@ -85,6 +85,11 @@ namespace cuda {
 		const ftl::rgbd::Camera &camera,
 		const float4x4 &poseInv, cudaStream_t stream);
 
+	void equirectangular_reproject(
+		ftl::cuda::TextureObject<uchar4> &image_in,
+		ftl::cuda::TextureObject<uchar4> &image_out,
+		const ftl::rgbd::Camera &camera, const float3x3 &pose, cudaStream_t stream);
+
 	template <typename A, typename B>
 	void dibr_normalise(
 		ftl::cuda::TextureObject<A> &in,

components/renderers/cpp/src/tri_render.cpp

@@ -130,6 +130,19 @@ Triangular::Triangular(nlohmann::json &config, ftl::rgbd::FrameSet *fs) : ftl::r
 	on("light_y", [this](const ftl::config::Event &e) { light_pos_.y = value("light_y", 0.3f); });
 	on("light_z", [this](const ftl::config::Event &e) { light_pos_.z = value("light_z", 0.3f); });
 
+	// Load any environment texture
+	std::string envimage = value("environment", std::string(""));
+	if (envimage.size() > 0) {
+		cv::Mat envim = cv::imread(envimage);
+		if (!envim.empty()) {
+			if (envim.type() == CV_8UC3) {
+				cv::cvtColor(envim,envim, cv::COLOR_BGR2BGRA);
+			}
+			env_image_.upload(envim);
+			env_tex_ = std::move(ftl::cuda::TextureObject<uchar4>(env_image_, true));
+		}
+	}
+
 	cudaSafeCall(cudaStreamCreate(&stream_));
 
 	//filters_ = ftl::create<ftl::Filters>(this, "filters");
@@ -485,7 +498,16 @@ bool Triangular::render(ftl::rgbd::VirtualSource *src, ftl::rgbd::Frame &out) {
 	// Clear all channels to 0 or max depth
 
 	out.get<GpuMat>(Channel::Depth).setTo(cv::Scalar(1000.0f), cvstream);
+
+	if (env_image_.empty()) {
 		out.get<GpuMat>(Channel::Colour).setTo(background_, cvstream);
+	} else {
+		auto pose = params.m_viewMatrixInverse.getFloat3x3();
+		ftl::cuda::equirectangular_reproject(
+			env_tex_,
+			out.createTexture<uchar4>(Channel::Colour, true),
+			camera, pose, stream_);
+	}
 
 	//LOG(INFO) << "Render ready: " << camera.width << "," << camera.height;