From c3cb8131680041c6fbc9859dbe954f4c930c8acc Mon Sep 17 00:00:00 2001
From: Nicolas Pope <nwpope@utu.fi>
Date: Mon, 2 Nov 2020 18:21:06 +0200
Subject: [PATCH] Attempt to luma adjust pixels
---
components/renderers/cpp/src/CUDARender.cpp | 4 +-
components/renderers/cpp/src/carver.cu | 98 ++++++++++++++++++---
components/renderers/cpp/src/carver.hpp | 6 ++
3 files changed, 95 insertions(+), 13 deletions(-)
diff --git a/components/renderers/cpp/src/CUDARender.cpp b/components/renderers/cpp/src/CUDARender.cpp
index 5d45b28b1..9ee7b0457 100644
--- a/components/renderers/cpp/src/CUDARender.cpp
+++ b/components/renderers/cpp/src/CUDARender.cpp
@@ -257,7 +257,7 @@ void CUDARender::_mesh(ftl::rgbd::Frame &out, const Eigen::Matrix4d &t, cudaStre
bool use_depth = scene_->anyHasChannel(Channel::Depth) || scene_->anyHasChannel(Channel::GroundTruth);
// For each source depth map, verify results
- if (value("pre_carve_inputs", true)) {
+ if (value("pre_carve_inputs", true) && !_alreadySeen()) {
for (size_t i=0; i < scene_->frames.size(); ++i) {
//if (!scene_->hasFrame(i)) continue;
auto &f = scene_->frames[i].cast<ftl::rgbd::Frame>();
@@ -307,6 +307,8 @@ void CUDARender::_mesh(ftl::rgbd::Frame &out, const Eigen::Matrix4d &t, cudaStre
stream_
);
}
+
+ ftl::cuda::apply_colour_scaling(colour_scale_, f.create<cv::cuda::GpuMat>(Channel::Colour), 3, stream_);
}
}
diff --git a/components/renderers/cpp/src/carver.cu b/components/renderers/cpp/src/carver.cu
index 02fccd23a..79c78fc36 100644
--- a/components/renderers/cpp/src/carver.cu
+++ b/components/renderers/cpp/src/carver.cu
@@ -73,7 +73,8 @@ __global__ void reverse_check_kernel(
float d = depth_in[y*pitch4+x];
- const float err_coef = 0.001f; //depthErrorCoef(ointrin);
+ // TODO: Externally provide the error coefficient
+ const float err_coef = 0.0005f; //depthErrorCoef(ointrin);
int ox = 0;
int oy = 0;
@@ -97,13 +98,13 @@ __global__ void reverse_check_kernel(
match = fabsf(campos.z - d2) < d2*d2*err_coef; break;
}
- d += 0.002f;
+ d += 0.002f; // TODO: Should this be += error or what?
} else break;
}
// We found a match, so do a colour check
- float scale = 1.0f;
- if (match) {
+ float idiff = 0.0f;
+ //if (match) {
// Generate colour scaling
const float ximgscale = float(cwidth) / float(ointrin.width);
ox = float(ox) * ximgscale;
@@ -113,18 +114,22 @@ __global__ void reverse_check_kernel(
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 uchar4 ocol = (match) ? ref_colour[oy*o_col_pitch4+ox] : vcol;
- 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);
+ float i1 = (0.2126f*float(vcol.z) + 0.7152f*float(vcol.y) + 0.0722f*float(vcol.x));
+ float i2 = (0.2126f*float(ocol.z) + 0.7152f*float(ocol.y) + 0.0722f*float(ocol.x));
+ idiff = i2-i1;
+
+ //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 = (0.2126f*scaleZ + 0.7152f*scaleY + 0.0722f*scaleX);
+ //}
+ colour_scale[x+pitch*y] = int8_t(max(-127.0f,min(127.0f,idiff)));
// Too much carving means just outright remove the point.
- depth_in[y*pitch4+x] = (count < 0) ? 0.0f : d;
+ depth_in[y*pitch4+x] = (count < 0 || fabsf(idiff) > 50.0f) ? 0.0f : d;
}
void ftl::cuda::depth_carve(
@@ -164,3 +169,72 @@ void ftl::cuda::depth_carve(
cudaSafeCall( cudaGetLastError() );
}
+
+// ==== Apply colour scale =====================================================
+
+template <int RADIUS>
+__global__ void apply_colour_scaling_kernel(
+ const int8_t* __restrict__ scale,
+ uchar4* __restrict__ colour,
+ int spitch,
+ int cpitch,
+ int swidth,
+ int sheight,
+ int cwidth,
+ int cheight
+) {
+ const int x = blockIdx.x*blockDim.x + threadIdx.x;
+ const int y = blockIdx.y*blockDim.y + threadIdx.y;
+
+ if (x >= 0 && x < cwidth && y >= 0 && y < cheight) {
+ int sx = (float(swidth) / float(cwidth)) * float(x);
+ int sy = (float(sheight) / float(cheight)) * float(y);
+
+ float s = 0.0f;
+ //float mindiff = 100.0f;
+
+ for (int v=-RADIUS; v<=RADIUS; ++v) {
+ #pragma unroll
+ for (int u=-RADIUS; u<=RADIUS; ++u) {
+ float ns = (sx >= RADIUS && sy >= RADIUS && sx < swidth-RADIUS && sy < sheight-RADIUS) ? scale[sx+u+(sy+v)*spitch] : 0.0f;
+ s += ns;
+ }
+ }
+
+ s /= float((2*RADIUS+1)*(2*RADIUS+1));
+
+ uchar4 c = colour[x+y*cpitch];
+ colour[x+y*cpitch] = make_uchar4(
+ max(0.0f, min(255.0f, float(c.x) + 0.0722f*s)),
+ max(0.0f, min(255.0f, float(c.y) + 0.7152f*s)),
+ max(0.0f, min(255.0f, float(c.z) + 0.2126f*s)),
+ 255.0f
+ );
+ }
+}
+
+void ftl::cuda::apply_colour_scaling(
+ const cv::cuda::GpuMat &scale,
+ cv::cuda::GpuMat &colour,
+ int radius,
+ cudaStream_t stream)
+{
+ static constexpr int THREADS_X = 16;
+ static constexpr int THREADS_Y = 8;
+
+ const dim3 gridSize((colour.cols + THREADS_X - 1)/THREADS_X, (colour.rows + THREADS_Y - 1)/THREADS_Y);
+ const dim3 blockSize(THREADS_X, THREADS_Y);
+
+ apply_colour_scaling_kernel<2><<<gridSize, blockSize, 0, stream>>>(
+ scale.ptr<int8_t>(),
+ colour.ptr<uchar4>(),
+ scale.step1(),
+ colour.step1()/4,
+ scale.cols,
+ scale.rows,
+ colour.cols,
+ colour.rows
+ );
+
+ cudaSafeCall( cudaGetLastError() );
+}
diff --git a/components/renderers/cpp/src/carver.hpp b/components/renderers/cpp/src/carver.hpp
index 8574685f1..a50a8d396 100644
--- a/components/renderers/cpp/src/carver.hpp
+++ b/components/renderers/cpp/src/carver.hpp
@@ -22,6 +22,12 @@ void depth_carve(
const ftl::rgbd::Camera &refcam,
cudaStream_t stream);
+void apply_colour_scaling(
+ const cv::cuda::GpuMat &scale,
+ cv::cuda::GpuMat &colour,
+ int radius,
+ cudaStream_t stream);
+
}
}
--
GitLab