diff --git a/components/renderers/cpp/include/ftl/render/tri_render.hpp b/components/renderers/cpp/include/ftl/render/tri_render.hpp index 60d3dcbf6a68576981c1d72194dae731e8592fec..3d9183643e9f2fe183499cf8bbcc97699328ef60 100644 --- a/components/renderers/cpp/include/ftl/render/tri_render.hpp +++ b/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> diff --git a/components/renderers/cpp/src/reprojection.cu b/components/renderers/cpp/src/reprojection.cu index f1a9022a52a3ed4efda457efc257fbf393c9014b..8a075b65382f7b574c6c4496557364a7a3aef1ee 100644 --- a/components/renderers/cpp/src/reprojection.cu +++ b/components/renderers/cpp/src/reprojection.cu @@ -274,3 +274,51 @@ template void ftl::cuda::reproject( const ftl::render::SplatParams ¶ms, 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() ); +} diff --git a/components/renderers/cpp/src/splatter_cuda.hpp b/components/renderers/cpp/src/splatter_cuda.hpp index 838eda409761c7420a35944a45900ace0912f124..010a637188523030ea2a40ac3c6cdd720c3a0148 100644 --- a/components/renderers/cpp/src/splatter_cuda.hpp +++ b/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, diff --git a/components/renderers/cpp/src/tri_render.cpp b/components/renderers/cpp/src/tri_render.cpp index f9e89ffb3adbb1b212ea8798441761b43ec5e5bc..b3fba82a7de676f2aa963e80467a45dbf3f0bd13 100644 --- a/components/renderers/cpp/src/tri_render.cpp +++ b/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); - out.get<GpuMat>(Channel::Colour).setTo(background_, cvstream); + + if (env_image_.empty() || !value("environment_enabled", false)) { + 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;