diff --git a/components/renderers/cpp/include/ftl/render/CUDARender.hpp b/components/renderers/cpp/include/ftl/render/CUDARender.hpp index 3b933821450aed45c2c91aebf3ee8f84ed8755ac..7369d5c9cc9762159b6aaf241f1768ac908fda3c 100644 --- a/components/renderers/cpp/include/ftl/render/CUDARender.hpp +++ b/components/renderers/cpp/include/ftl/render/CUDARender.hpp @@ -58,6 +58,7 @@ class CUDARender : public ftl::render::FSRenderer { ftl::cuda::TextureObject<float4> accum_; // 2 is number of channels can render together ftl::cuda::TextureObject<int> contrib_; //ftl::cuda::TextureObject<half4> normals_; + cv::cuda::GpuMat colour_scale_; std::list<cv::cuda::GpuMat*> screen_buffers_; std::list<cv::cuda::GpuMat*> depth_buffers_; diff --git a/components/renderers/cpp/src/CUDARender.cpp b/components/renderers/cpp/src/CUDARender.cpp index 0e16d901377019275954e146b82d7e6bd77a0268..5d45b28b1ccf61313aba36823970238116ed1994 100644 --- a/components/renderers/cpp/src/CUDARender.cpp +++ b/components/renderers/cpp/src/CUDARender.cpp @@ -8,6 +8,7 @@ #include <ftl/cuda/transform.hpp> #include <ftl/operators/cuda/smoothing_cuda.hpp> #include "carver.hpp" +#include <ftl/utility/image_debug.hpp> #include <ftl/cuda/colour_cuda.hpp> @@ -297,6 +298,9 @@ void CUDARender::_mesh(ftl::rgbd::Frame &out, const Eigen::Matrix4d &t, cudaStre ftl::cuda::depth_carve( f.create<cv::cuda::GpuMat>(_getDepthChannel()), ref.get<cv::cuda::GpuMat>(Channel::Depth), + f.get<cv::cuda::GpuMat>(Channel::Colour), + ref.get<cv::cuda::GpuMat>(Channel::Colour), + colour_scale_, transformR, f.getLeft(), ref.getLeft(), @@ -306,6 +310,8 @@ void CUDARender::_mesh(ftl::rgbd::Frame &out, const Eigen::Matrix4d &t, cudaStre } } + if (!colour_scale_.empty()) ftl::utility::show_image(colour_scale_, "CScale", 1.0f, ftl::utility::ImageVisualisation::HEAT_MAPPED); + // For each source depth map for (size_t i=0; i < scene_->frames.size(); ++i) { //if (!scene_->hasFrame(i)) continue; @@ -417,7 +423,7 @@ void CUDARender::_mesh(ftl::rgbd::Frame &out, const Eigen::Matrix4d &t, cudaStre ftl::cuda::touch_merge(depth_out_, out.createTexture<float>(_getDepthChannel()), collisions_, 1024, touch_dist_, stream_); // For each source depth map, verify results - if (value("post_carve_result", false)) { + /*if (value("post_carve_result", false)) { for (size_t i=0; i < scene_->frames.size(); ++i) { //if (!scene_->hasFrame(i)) continue; auto &f = scene_->frames[i].cast<ftl::rgbd::Frame>(); @@ -446,7 +452,7 @@ void CUDARender::_mesh(ftl::rgbd::Frame &out, const Eigen::Matrix4d &t, cudaStre stream_ ); } - } + }*/ //filters_->filter(out, src, stream); diff --git a/components/renderers/cpp/src/carver.cu b/components/renderers/cpp/src/carver.cu index 96e5b16fd11157589f7e07f14e0748f9e73b447e..02fccd23a977d5119fd7702c62e2b1454ce79573 100644 --- a/components/renderers/cpp/src/carver.cu +++ b/components/renderers/cpp/src/carver.cu @@ -1,4 +1,5 @@ #include "carver.hpp" +#include <cudatl/fixed.hpp> __device__ inline float depthErrorCoef(const ftl::rgbd::Camera &cam, float disps=1.0f) { return disps / (cam.baseline*cam.fx); @@ -6,6 +7,7 @@ __device__ inline float depthErrorCoef(const ftl::rgbd::Camera &cam, float disps // ==== Reverse Verify Result ================================================== +// No colour scale calculations __global__ void reverse_check_kernel( float* __restrict__ depth_in, const float* __restrict__ depth_original, @@ -38,21 +40,127 @@ __global__ void reverse_check_kernel( // If the value is significantly further then carve. Depth error // is not always easy to calculate, depends on source. if (!(d2 < ointrin.maxDepth && d2 - campos.z > d2*d2*err_coef)) break; + + d += 0.002f; + } else break; + } + + // Too much carving means just outright remove the point. + depth_in[y*pitch4+x] = (count < 0) ? 0.0f : d; +} + +__global__ void reverse_check_kernel( + float* __restrict__ depth_in, + const float* __restrict__ depth_original, + const uchar4* __restrict__ in_colour, + const uchar4* __restrict__ ref_colour, + int8_t* __restrict__ colour_scale, + int pitch4, + int pitch, + int opitch4, + int in_col_pitch4, + int o_col_pitch4, + int cwidth, + int cheight, + float4x4 transformR, + 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 >= vintrin.width || y < 0 || y >= vintrin.height) return; + + float d = depth_in[y*pitch4+x]; + + const float err_coef = 0.001f; //depthErrorCoef(ointrin); + + int ox = 0; + int oy = 0; + + bool match = false; + + int count = 10; // Allow max 2cm of carving. + while (--count >= 0) { + float3 campos = transformR * vintrin.screenToCam(x,y,d); + int2 spos = ointrin.camToScreen<int2>(campos); + ox = spos.x; + oy = spos.y; + + if (campos.z > 0.0f && ox >= 0 && ox < ointrin.width && oy >= 0 && oy < ointrin.height) { + float d2 = depth_original[oy*opitch4+ox]; + + // TODO: Threshold comes from depth error characteristics + // If the value is significantly further then carve. Depth error + // is not always easy to calculate, depends on source. + if (!(d2 < ointrin.maxDepth && d2 - campos.z > d2*d2*err_coef)) { + match = fabsf(campos.z - d2) < d2*d2*err_coef; break; + } + d += 0.002f; } else break; } + // We found a match, so do a colour check + float scale = 1.0f; + if (match) { + // Generate colour scaling + const float ximgscale = float(cwidth) / float(ointrin.width); + ox = float(ox) * ximgscale; + const float yimgscale = float(cheight) / float(ointrin.height); + oy = float(oy) * yimgscale; + + int cy = float(y) * yimgscale; + int cx = float(x) * ximgscale; + + const uchar4 ocol = ref_colour[oy*o_col_pitch4+ox]; + const uchar4 vcol = in_colour[cy*in_col_pitch4+cx]; + + const float scaleX = (vcol.x == 0) ? 1.0f : float(ocol.x) / float(vcol.x); + const float scaleY = (vcol.y == 0) ? 1.0f : float(ocol.y) / float(vcol.y); + const float scaleZ = (vcol.z == 0) ? 1.0f : float(ocol.z) / float(vcol.z); + scale = (scaleX+scaleY+scaleZ) / 3.0f; + } + colour_scale[x+pitch*y] = cudatl::float2fixed8<6>(scale); + // Too much carving means just outright remove the point. depth_in[y*pitch4+x] = (count < 0) ? 0.0f : d; } -void ftl::cuda::depth_carve(cv::cuda::GpuMat &depth_in, const cv::cuda::GpuMat &depth_original, const float4x4 &transformR, const ftl::rgbd::Camera &vintrin, const ftl::rgbd::Camera &ointrin, cudaStream_t stream) { +void ftl::cuda::depth_carve( + cv::cuda::GpuMat &depth_in, + const cv::cuda::GpuMat &depth_original, + const cv::cuda::GpuMat &in_colour, + const cv::cuda::GpuMat &ref_colour, + cv::cuda::GpuMat &colour_scale, + const float4x4 &transformR, + const ftl::rgbd::Camera &vintrin, + const ftl::rgbd::Camera &ointrin, + cudaStream_t stream) +{ static constexpr int THREADS_X = 16; static constexpr int THREADS_Y = 8; const dim3 gridSize((depth_in.cols + THREADS_X - 1)/THREADS_X, (depth_in.rows + THREADS_Y - 1)/THREADS_Y); const dim3 blockSize(THREADS_X, THREADS_Y); - reverse_check_kernel<<<gridSize, blockSize, 0, stream>>>(depth_in.ptr<float>(), depth_original.ptr<float>(), depth_in.step1(), depth_original.step1(), transformR, vintrin, ointrin); + colour_scale.create(depth_in.size(), CV_8U); + + reverse_check_kernel<<<gridSize, blockSize, 0, stream>>>( + depth_in.ptr<float>(), + depth_original.ptr<float>(), + in_colour.ptr<uchar4>(), + ref_colour.ptr<uchar4>(), + colour_scale.ptr<int8_t>(), + depth_in.step1(), + colour_scale.step1(), + depth_original.step1(), + in_colour.step1()/4, + ref_colour.step1()/4, + in_colour.cols, + in_colour.rows, + transformR, + vintrin, ointrin); + cudaSafeCall( cudaGetLastError() ); } diff --git a/components/renderers/cpp/src/carver.hpp b/components/renderers/cpp/src/carver.hpp index f3e6f671a531aa921c0c61ae81d3f407b5a7519d..8574685f1371bd0c90e51c53e88f6684620ccc9e 100644 --- a/components/renderers/cpp/src/carver.hpp +++ b/components/renderers/cpp/src/carver.hpp @@ -14,6 +14,9 @@ namespace cuda { void depth_carve( cv::cuda::GpuMat &in, const cv::cuda::GpuMat &ref, + const cv::cuda::GpuMat &in_colour, + const cv::cuda::GpuMat &ref_colour, + cv::cuda::GpuMat &colour_scale, const float4x4 &transform, const ftl::rgbd::Camera &incam, const ftl::rgbd::Camera &refcam, diff --git a/lib/cudatl/include/cudatl/fixed.hpp b/lib/cudatl/include/cudatl/fixed.hpp index 3009f180c3bfb674a7a06cb56a05b61d7cfb3ecb..12afa70213cbd9326703ada220fdab5e723595aa 100644 --- a/lib/cudatl/include/cudatl/fixed.hpp +++ b/lib/cudatl/include/cudatl/fixed.hpp @@ -10,6 +10,12 @@ __device__ inline float fixed2float(short v); template <int FRAC> __device__ inline short float2fixed(float v); +template <int FRAC> +__device__ inline float fixed2float8(int8_t v); + +template <int FRAC> +__device__ inline int8_t float2fixed8(float v); + } #include <cudatl/impl/fixed.hpp> diff --git a/lib/cudatl/include/cudatl/impl/fixed.hpp b/lib/cudatl/include/cudatl/impl/fixed.hpp index 289eafbdf72e4fe057e3bc4d2934bb157706f0a4..083641732acb6735212e899a84e1e530dc37e1d7 100644 --- a/lib/cudatl/include/cudatl/impl/fixed.hpp +++ b/lib/cudatl/include/cudatl/impl/fixed.hpp @@ -8,4 +8,16 @@ template <int FRAC> __device__ inline short cudatl::float2fixed(float v) { return short(v * float(1 << FRAC)); -} \ No newline at end of file +} + +template <int FRAC> +__device__ inline float cudatl::fixed2float8(int8_t v) +{ + return float(v) / float(1 << FRAC); +} + +template <int FRAC> +__device__ inline int8_t cudatl::float2fixed8(float v) +{ + return int8_t(v * float(1 << FRAC)); +}