diff --git a/components/renderers/cpp/src/CUDARender.cpp b/components/renderers/cpp/src/CUDARender.cpp index dab897d5696f78ebffbec097485778bebd57c03b..c274e136fd5348ce7c982a26e51e151e99679aec 100644 --- a/components/renderers/cpp/src/CUDARender.cpp +++ b/components/renderers/cpp/src/CUDARender.cpp @@ -363,6 +363,37 @@ void CUDARender::_mesh(ftl::rgbd::Frame &out, const Eigen::Matrix4d &t, cudaStre // Now merge new render to any existing frameset render, detecting collisions ftl::cuda::touch_merge(depth_out_, out.createTexture<float>(_getDepthChannel()), collisions_, 1024, touch_dist_, stream_); + // For each source depth map, verify results + for (size_t i=0; i < scene_->frames.size(); ++i) { + //if (!scene_->hasFrame(i)) continue; + auto &f = scene_->frames[i].cast<ftl::rgbd::Frame>(); + //auto *s = scene_->sources[i]; + + if (!f.has(Channel::Colour)) { + //LOG(ERROR) << "Missing required channel"; + continue; + } + + // We have the needed depth data? + if (use_depth && !f.hasOwn(Channel::Depth) && !f.hasOwn(Channel::GroundTruth)) { + continue; + } + + //auto pose = MatrixConversion::toCUDA(t.cast<float>() * f.getPose().cast<float>()); + auto transformR = MatrixConversion::toCUDA(f.getPose().cast<float>().inverse() * t.cast<float>().inverse()) * poseInverse_; + auto transform = pose_ * MatrixConversion::toCUDA(t.cast<float>() * f.getPose().cast<float>()); + + ftl::cuda::reverse_verify( + out.getTexture<float>(_getDepthChannel()), + f.getTexture<float>(Channel::Depth), + transformR, + transform, + params_.camera, + f.getLeft(), + stream_ + ); + } + //filters_->filter(out, src, stream); // Generate normals for final virtual image diff --git a/components/renderers/cpp/src/splatter_cuda.hpp b/components/renderers/cpp/src/splatter_cuda.hpp index 41088cc226e34f4916c6919d8edc465906dbd499..147ccf99479674f26ab6b7b5d86ee493c18f9123 100644 --- a/components/renderers/cpp/src/splatter_cuda.hpp +++ b/components/renderers/cpp/src/splatter_cuda.hpp @@ -170,6 +170,15 @@ namespace cuda { ftl::cuda::TextureObject<float> &d2, float factor, cudaStream_t stream); + void reverse_verify( + ftl::cuda::TextureObject<float> &depth_in, + ftl::cuda::TextureObject<float> &depth_original, + const float4x4 &transformR, + const float4x4 &transform, + const ftl::rgbd::Camera &vintrin, + const ftl::rgbd::Camera &ointrin, + cudaStream_t stream); + void fix_bad_colour( ftl::cuda::TextureObject<float> &depth, ftl::cuda::TextureObject<uchar4> &out, diff --git a/components/renderers/cpp/src/triangle_render.cu b/components/renderers/cpp/src/triangle_render.cu index dccdd276ad1ab8a6520a19a5b908c51db6b54823..3b45d1315d2ef3bc486fe97acc7f7a3a6f163e0c 100644 --- a/components/renderers/cpp/src/triangle_render.cu +++ b/components/renderers/cpp/src/triangle_render.cu @@ -202,7 +202,8 @@ __global__ void merge_convert_kernel( if (x < 0 || x >= depth_in.width() || y < 0 || y >= depth_in.height()) return; - float a = float(depth_in.tex2D(x,y))*alpha; + int d = depth_in.tex2D(x,y); + float a = float(d)*alpha; float b = depth_out.tex2D(x,y); depth_out(x,y) = min(a,b); } @@ -215,6 +216,47 @@ void ftl::cuda::merge_convert_depth(TextureObject<int> &depth_in, TextureObject< cudaSafeCall( cudaGetLastError() ); } +// ==== Reverse Verify Result ================================================== + +// ==== Merge convert =========== + +__global__ void reverse_check_kernel( + TextureObject<float> depth_in, + TextureObject<float> depth_original, + float4x4 transformR, + float4x4 transform, + ftl::rgbd::Camera vintrin, + ftl::rgbd::Camera ointrin +) { + const int x = blockIdx.x*blockDim.x + threadIdx.x; + const int y = blockIdx.y*blockDim.y + threadIdx.y; + + if (x < 0 || x >= depth_in.width() || y < 0 || y >= depth_in.height()) return; + + float d = depth_in.tex2D(x,y); + float3 campos = transformR * vintrin.screenToCam(x,y,d); + int2 spos = ointrin.camToScreen<int2>(campos); + int ox = spos.x; + int oy = spos.y; + + if (campos.z > 0.0f && ox >= 0 && ox < ointrin.width && oy >= 0 && oy < ointrin.height) { + float d2 = depth_original.tex2D(ox,oy); + + if (d2 < ointrin.maxDepth && d2 - campos.z > 0.002f) { + //printf("Original %f, %f\n", d2, campos.z); + depth_in(x,y) = 0.0f; //(transform * ointrin.screenToCam(ox,oy,d2)).z; + } + } +} + +void ftl::cuda::reverse_verify(TextureObject<float> &depth_in, TextureObject<float> &depth_original, const float4x4 &transformR, const float4x4 &transform, const ftl::rgbd::Camera &vintrin, const ftl::rgbd::Camera &ointrin, cudaStream_t stream) { + const dim3 gridSize((depth_in.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (depth_in.height() + T_PER_BLOCK - 1)/T_PER_BLOCK); + const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK); + + reverse_check_kernel<<<gridSize, blockSize, 0, stream>>>(depth_in, depth_original, transformR, transform, vintrin, ointrin); + cudaSafeCall( cudaGetLastError() ); +} + // ==== BLENDER ======== /*