diff --git a/components/renderers/cpp/include/ftl/cuda/normals.hpp b/components/renderers/cpp/include/ftl/cuda/normals.hpp index b4d2ced197c383a6d2dc185bf1c8883d6dde356a..b9801743a5e505cf38a7b632127755f248e618c9 100644 --- a/components/renderers/cpp/include/ftl/cuda/normals.hpp +++ b/components/renderers/cpp/include/ftl/cuda/normals.hpp @@ -10,7 +10,9 @@ namespace cuda { void normals(ftl::cuda::TextureObject<float4> &output, ftl::cuda::TextureObject<float4> &temp, - ftl::cuda::TextureObject<float4> &input, cudaStream_t stream); + ftl::cuda::TextureObject<float4> &input, + const ftl::rgbd::Camera &camera, + const float3x3 &pose, cudaStream_t stream); void normal_visualise(ftl::cuda::TextureObject<float4> &norm, ftl::cuda::TextureObject<uchar4> &output, diff --git a/components/renderers/cpp/include/ftl/render/splat_params.hpp b/components/renderers/cpp/include/ftl/render/splat_params.hpp index 5bd7d2edbc0b30863e5ecfe04f2e7bc0ce662f20..0509ee37f0d85163fde2b462ad0871f8a824070b 100644 --- a/components/renderers/cpp/include/ftl/render/splat_params.hpp +++ b/components/renderers/cpp/include/ftl/render/splat_params.hpp @@ -9,9 +9,7 @@ namespace ftl { namespace render { static const uint kShowDisconMask = 0x00000001; -static const uint kNoSplatting = 0x00000002; -static const uint kNoUpsampling = 0x00000004; -static const uint kNoTexturing = 0x00000008; +static const uint kNormalWeightColours = 0x00000002; struct __align__(16) SplatParams { float4x4 m_viewMatrix; diff --git a/components/renderers/cpp/src/normals.cu b/components/renderers/cpp/src/normals.cu index cbbd0e96797ee11fffd43de0df399a03efcf8db8..4d66869316bb3cc4f7f9bd56f889bf92ecc4635c 100644 --- a/components/renderers/cpp/src/normals.cu +++ b/components/renderers/cpp/src/normals.cu @@ -34,7 +34,8 @@ __global__ void computeNormals_kernel(ftl::cuda::TextureObject<float4> output, template <int RADIUS> __global__ void smooth_normals_kernel(ftl::cuda::TextureObject<float4> norms, ftl::cuda::TextureObject<float4> output, - ftl::cuda::TextureObject<float4> points, float smoothing) { + ftl::cuda::TextureObject<float4> points, + ftl::rgbd::Camera camera, float3x3 pose, float smoothing) { const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x; const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y; @@ -63,21 +64,28 @@ __global__ void smooth_normals_kernel(ftl::cuda::TextureObject<float4> norms, } } - // FIXME: USE A DIFFERENT OUTPUT BUFFER - //__syncthreads(); - output(x,y) = (contrib > 0.0f) ? make_float4(nsum / contrib, 1.0f) : make_float4(0.0f); + // Compute dot product of normal with camera to obtain measure of how + // well this point faces the source camera, a measure of confidence + float3 ray = pose * camera.screenToCam(x, y, 1.0f); + ray = ray / length(ray); + nsum /= contrib; + nsum /= length(nsum); + + output(x,y) = (contrib > 0.0f) ? make_float4(nsum, dot(nsum, ray)) : make_float4(0.0f); } void ftl::cuda::normals(ftl::cuda::TextureObject<float4> &output, ftl::cuda::TextureObject<float4> &temp, - ftl::cuda::TextureObject<float4> &input, cudaStream_t stream) { + ftl::cuda::TextureObject<float4> &input, + const ftl::rgbd::Camera &camera, + const float3x3 &pose,cudaStream_t stream) { const dim3 gridSize((input.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (input.height() + T_PER_BLOCK - 1)/T_PER_BLOCK); const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK); computeNormals_kernel<<<gridSize, blockSize, 0, stream>>>(temp, input); cudaSafeCall( cudaGetLastError() ); - smooth_normals_kernel<3><<<gridSize, blockSize, 0, stream>>>(temp, output, input, 0.04f); + smooth_normals_kernel<3><<<gridSize, blockSize, 0, stream>>>(temp, output, input, camera, pose, 0.04f); cudaSafeCall( cudaGetLastError() ); #ifdef _DEBUG diff --git a/components/renderers/cpp/src/points.cu b/components/renderers/cpp/src/points.cu index f73eaa2085e901991f9cc40580a3273864ceda4f..a188adaeffaad12be42762a3e1b529926dee225e 100644 --- a/components/renderers/cpp/src/points.cu +++ b/components/renderers/cpp/src/points.cu @@ -12,7 +12,9 @@ __global__ void point_cloud_kernel(ftl::cuda::TextureObject<float4> output, ftl: output(x,y) = make_float4(MINF, MINF, MINF, MINF); const float d = depth.tex2D((int)x, (int)y); - float p = d; + + // Calculate depth between 0.0 and 1.0 + float p = (d - params.minDepth) / (params.maxDepth - params.minDepth); if (d >= params.minDepth && d <= params.maxDepth) { /* Orts-Escolano S. et al. 2016. Holoportation: Virtual 3D teleportation in real-time. */ diff --git a/components/renderers/cpp/src/splat_render.cpp b/components/renderers/cpp/src/splat_render.cpp index d993e145856840e7af545de7365e1e77af7f0866..b22b377e158faff8143d49c279955a195ce52dcc 100644 --- a/components/renderers/cpp/src/splat_render.cpp +++ b/components/renderers/cpp/src/splat_render.cpp @@ -288,10 +288,8 @@ bool Splatter::render(ftl::rgbd::VirtualSource *src, ftl::rgbd::Frame &out, cuda // Parameters object to pass to CUDA describing the camera SplatParams params; params.m_flags = 0; - if (src->value("show_discontinuity_mask", false)) params.m_flags |= ftl::render::kShowDisconMask; - //if (src->value("splatting", true) == false) params.m_flags |= ftl::render::kNoSplatting; - //if (src->value("upsampling", true) == false) params.m_flags |= ftl::render::kNoUpsampling; - //if (src->value("texturing", true) == false) params.m_flags |= ftl::render::kNoTexturing; + if (value("show_discontinuity_mask", false)) params.m_flags |= ftl::render::kShowDisconMask; + if (value("normal_weight_colours", true)) params.m_flags |= ftl::render::kNormalWeightColours; params.m_viewMatrix = MatrixConversion::toCUDA(src->getPose().cast<float>().inverse()); params.m_viewMatrixInverse = MatrixConversion::toCUDA(src->getPose().cast<float>()); params.camera = camera; @@ -328,14 +326,15 @@ bool Splatter::render(ftl::rgbd::VirtualSource *src, ftl::rgbd::Frame &out, cuda } if (!f.hasChannel(Channel::Normals)) { + Eigen::Matrix4f matrix = s->getPose().cast<float>(); + auto pose = MatrixConversion::toCUDA(matrix); + auto &g = f.get<GpuMat>(Channel::Colour); ftl::cuda::normals(f.createTexture<float4>(Channel::Normals, Format<float4>(g.cols, g.rows)), temp_.getTexture<float4>(Channel::Normals), // FIXME: Uses assumption of vcam res same as input res - f.getTexture<float4>(Channel::Points), stream); + f.getTexture<float4>(Channel::Points), s->parameters(), pose.getFloat3x3(), stream); if (norm_filter_ > -0.1f) { - Eigen::Matrix4f matrix = s->getPose().cast<float>(); - auto pose = MatrixConversion::toCUDA(matrix); ftl::cuda::normal_filter(f.getTexture<float4>(Channel::Normals), f.getTexture<float4>(Channel::Points), s->parameters(), pose, norm_filter_, stream); } } diff --git a/components/renderers/cpp/src/splatter.cu b/components/renderers/cpp/src/splatter.cu index a4f5d97ff91b62561cfa4fda990c44fc7790c72f..3567d5c07dcdc819b93989992069b3244149855f 100644 --- a/components/renderers/cpp/src/splatter.cu +++ b/components/renderers/cpp/src/splatter.cu @@ -184,39 +184,34 @@ __device__ inline float make(const float4 &v) { const float3 camPos = params.camera.screenToCam((int)(x+u),(int)(y+v),d); const float3 camPos2 = params.camera.screenToCam((int)(x),(int)(y),d); - //if (length(camPos - camPos2) > 2.0f*(camPos.z/params.camera.fx)) continue; const float3 worldPos = params.m_viewMatrixInverse * camPos; - //const float3 camPos2 = pose_inv * worldPos; - //const uint2 screenPos = camera.camToScreen<uint2>(camPos2); - - //if (screenPos.x < points.width() && screenPos.y < points.height()) { - // Can now read points, normals and colours from source cam - - // What is contribution of our current point at this pixel? - //const float3 p = make_float3(points.tex2D((int)screenPos.x, (int)screenPos.y)); - //const float weight = ftl::cuda::spatialWeighting(worldPos, p, (camPos.z/params.camera.fx)); //*(camPos2.z/camera.fx)); - //if (weight <= 0.0f) continue; - - float3 n = make_float3(normals.tex2D((int)(x+u), (int)(y+v))); - const float l = length(n); - if (l == 0.0f) continue; - n /= l; - - // Does the ray intersect plane of splat? - float t = 1000.0f; - if (ftl::cuda::intersectPlane(n, worldPos, origin, ray, t)) { //} && fabs(t-camPos.z) < 0.01f) { - //t *= (params.m_viewMatrix.getFloat3x3() * ray).z; - t *= scale; - const float3 camPos3 = params.camera.screenToCam((int)(x),(int)(y),t); - const float weight = ftl::cuda::spatialWeighting(camPos, camPos3, 2.0f*(camPos3.z/params.camera.fx)); //*(camPos2.z/camera.fx)); - if (weight == 0.0f) continue; - //depth += t * weight; - //contrib += weight; - depth = min(depth, t); - results[i/WARP_SIZE] = {weight, t, in.tex2D((int)x+u, (int)y+v)}; //make_float2(t, weight); - //atomicMin(&depth_out(x,y), (int)(depth * 1000.0f)); - } - //} + + + // Assumed to be normalised + float4 n = normals.tex2D((int)(x+u), (int)(y+v)); + + // Does the ray intersect plane of splat? + float t = 1000.0f; + if (ftl::cuda::intersectPlane(make_float3(n), worldPos, origin, ray, t)) { //} && fabs(t-camPos.z) < 0.01f) { + // Adjust from normalised ray back to original meters units + t *= scale; + const float3 camPos3 = params.camera.screenToCam((int)(x),(int)(y),t); + float weight = ftl::cuda::spatialWeighting(camPos, camPos3, 2.0f*(camPos3.z/params.camera.fx)); + + /* Buehler C. et al. 2001. Unstructured Lumigraph Rendering. */ + /* Orts-Escolano S. et al. 2016. Holoportation: Virtual 3D teleportation in real-time. */ + // This is the simple naive colour weighting. It might be good + // enough for our purposes if the alignment step prevents ghosting + // TODO: Use depth and perhaps the neighbourhood consistency in: + // Kuster C. et al. 2011. FreeCam: A hybrid camera system for interactive free-viewpoint video + if (params.m_flags & ftl::render::kNormalWeightColours) weight *= n.w * n.w; + //if (params.m_flags & ftl::render::kDepthWeightColours) weight *= ??? + + if (weight <= 0.0f) continue; + + depth = min(depth, t); + results[i/WARP_SIZE] = {weight, t, in.tex2D((int)x+u, (int)y+v)}; + } } depth = warpMin(depth);