diff --git a/applications/reconstruct/src/ilw/ilw.cu b/applications/reconstruct/src/ilw/ilw.cu
index 831ff641b50dd7396a7282b99119a72164f6a34e..97ea74eebdc8f3b3d8aeec75f5153ce468808a33 100644
--- a/applications/reconstruct/src/ilw/ilw.cu
+++ b/applications/reconstruct/src/ilw/ilw.cu
@@ -195,14 +195,14 @@ __global__ void correspondence_energy_vector_kernel(
const float avgcolour = totalcolour/(float)count;
const float confidence = bestcolour / totalcolour; //bestcolour - avgcolour;
- if (bestweight > 0.0f) {
+ //if (bestweight > 0.0f) {
float old = conf.tex2D(x,y);
- if (bestweight * confidence > old) {
+ if (bestweight * confidence >= old) {
dout(x,y) = bestdepth;
conf(x,y) = bestweight * confidence;
}
- }
+ //}
}
void ftl::cuda::correspondence(
diff --git a/components/renderers/cpp/src/splat_render.cpp b/components/renderers/cpp/src/splat_render.cpp
index 1ccaf156cc31d7ba2097364e627d6d93500d293a..895ea79ff5bdcf1dd476ecb441fad2af7fcc6f79 100644
--- a/components/renderers/cpp/src/splat_render.cpp
+++ b/components/renderers/cpp/src/splat_render.cpp
@@ -134,8 +134,8 @@ void Splatter::renderChannel(
cv::cuda::Stream cvstream = cv::cuda::StreamAccessor::wrapStream(stream);
temp_.get<GpuMat>(Channel::Depth).setTo(cv::Scalar(0x7FFFFFFF), cvstream);
temp_.get<GpuMat>(Channel::Depth2).setTo(cv::Scalar(0x7FFFFFFF), cvstream);
- //temp_.get<GpuMat>(Channel::Colour).setTo(cv::Scalar(0.0f,0.0f,0.0f,0.0f), cvstream);
- //temp_.get<GpuMat>(Channel::Contribution).setTo(cv::Scalar(0.0f), cvstream);
+ temp_.get<GpuMat>(Channel::Colour).setTo(cv::Scalar(0.0f,0.0f,0.0f,0.0f), cvstream);
+ temp_.get<GpuMat>(Channel::Contribution).setTo(cv::Scalar(0.0f), cvstream);
if (scene_->frames.size() < 1) return;
bool is_float = out.get<GpuMat>(channel).type() == CV_32F; //ftl::rgbd::isFloatChannel(channel);
@@ -162,23 +162,33 @@ void Splatter::renderChannel(
//LOG(INFO) << "DIBR DONE";
}
- //temp_.createTexture<float4>(Channel::Colour);
- //temp_.createTexture<float>(Channel::Contribution);
+ temp_.createTexture<float4>(Channel::Colour);
+ temp_.createTexture<float>(Channel::Contribution);
out.create<GpuMat>(Channel::Normals, Format<float4>(params.camera.width, params.camera.height));
+ out.get<GpuMat>(Channel::Normals).setTo(cv::Scalar(0.0f,0.0f,0.0f,0.0f), cvstream);
// Create normals first
for (auto &f : scene_->frames) {
+
ftl::cuda::dibr_attribute(
f.createTexture<float4>(Channel::Normals),
f.createTexture<float4>(Channel::Points),
temp_.getTexture<int>(Channel::Depth2),
out.createTexture<float4>(Channel::Normals),
+ temp_.getTexture<float>(Channel::Contribution),
params, stream
);
}
- //temp_.get<GpuMat>(Channel::Colour).setTo(cv::Scalar(0.0f,0.0f,0.0f,0.0f), cvstream);
- //temp_.get<GpuMat>(Channel::Contribution).setTo(cv::Scalar(0.0f), cvstream);
+ ftl::cuda::dibr_normalise(
+ out.getTexture<float4>(Channel::Normals),
+ out.getTexture<float4>(Channel::Normals),
+ temp_.getTexture<float>(Channel::Contribution),
+ stream
+ );
+
+ temp_.get<GpuMat>(Channel::Colour).setTo(cv::Scalar(0.0f,0.0f,0.0f,0.0f), cvstream);
+ temp_.get<GpuMat>(Channel::Contribution).setTo(cv::Scalar(0.0f), cvstream);
if (is_4chan) {
temp_.create<GpuMat>(Channel::Colour2, Format<float4>(params.camera.width, params.camera.height));
@@ -207,7 +217,8 @@ void Splatter::renderChannel(
f.createTexture<float4>(channel),
f.createTexture<float4>(Channel::Points),
temp_.getTexture<int>(Channel::Depth2),
- (splat_) ? temp_.createTexture<float4>(Channel::Colour2) : out.createTexture<float4>(channel),
+ temp_.getTexture<float4>(Channel::Colour),
+ temp_.getTexture<float>(Channel::Contribution),
params, stream
);
} else if (is_float) {
@@ -215,7 +226,8 @@ void Splatter::renderChannel(
f.createTexture<float>(channel),
f.createTexture<float4>(Channel::Points),
temp_.getTexture<int>(Channel::Depth2),
- (splat_) ? temp_.createTexture<float>(Channel::Colour2) : out.createTexture<float>(channel),
+ temp_.createTexture<float>(Channel::Colour2),
+ temp_.getTexture<float>(Channel::Contribution),
params, stream
);
} else {
@@ -223,12 +235,36 @@ void Splatter::renderChannel(
f.createTexture<uchar4>(channel),
f.createTexture<float4>(Channel::Points),
temp_.getTexture<int>(Channel::Depth2),
- (splat_) ? temp_.createTexture<uchar4>(Channel::Colour2) : out.createTexture<uchar4>(channel),
+ temp_.createTexture<float4>(Channel::Colour),
+ temp_.getTexture<float>(Channel::Contribution),
params, stream
);
}
}
+ if (is_4chan) {
+ ftl::cuda::dibr_normalise(
+ temp_.getTexture<float4>(Channel::Colour),
+ (splat_) ? temp_.createTexture<float4>(Channel::Colour2) : out.createTexture<float4>(channel),
+ temp_.getTexture<float>(Channel::Contribution),
+ stream
+ );
+ } else if (is_float) {
+ ftl::cuda::dibr_normalise(
+ temp_.createTexture<float>(Channel::Colour2),
+ (splat_) ? temp_.createTexture<float>(Channel::Colour2) : out.createTexture<float>(channel),
+ temp_.getTexture<float>(Channel::Contribution),
+ stream
+ );
+ } else {
+ ftl::cuda::dibr_normalise(
+ temp_.getTexture<float4>(Channel::Colour),
+ (splat_) ? temp_.createTexture<uchar4>(Channel::Colour2) : out.createTexture<uchar4>(channel),
+ temp_.getTexture<float>(Channel::Contribution),
+ stream
+ );
+ }
+
//out.get<GpuMat>(Channel::Left).setTo(cv::Scalar(0,0,0,0), cvstream);
// Now splat the points
@@ -277,8 +313,8 @@ bool Splatter::render(ftl::rgbd::VirtualSource *src, ftl::rgbd::Frame &out, cuda
out.create<GpuMat>(Channel::Colour, Format<uchar4>(camera.width, camera.height));
// FIXME: Use source resolutions, not virtual resolution
- //temp_.create<GpuMat>(Channel::Colour, Format<float4>(camera.width, camera.height));
- //temp_.create<GpuMat>(Channel::Contribution, Format<float>(camera.width, camera.height));
+ temp_.create<GpuMat>(Channel::Colour, Format<float4>(camera.width, camera.height));
+ temp_.create<GpuMat>(Channel::Contribution, Format<float>(camera.width, camera.height));
temp_.create<GpuMat>(Channel::Depth, Format<int>(camera.width, camera.height));
temp_.create<GpuMat>(Channel::Depth2, Format<int>(camera.width, camera.height));
temp_.create<GpuMat>(Channel::Normals, Format<float4>(camera.width, camera.height));
@@ -302,6 +338,7 @@ bool Splatter::render(ftl::rgbd::VirtualSource *src, ftl::rgbd::Frame &out, cuda
//LOG(INFO) << "Render ready: " << camera.width << "," << camera.height;
temp_.createTexture<int>(Channel::Depth);
+ //temp_.get<GpuMat>(Channel::Normals).setTo(cv::Scalar(0.0f,0.0f,0.0f,0.0f), cvstream);
// First make sure each input has normals
temp_.createTexture<float4>(Channel::Normals);
diff --git a/components/renderers/cpp/src/splatter.cu b/components/renderers/cpp/src/splatter.cu
index 3567d5c07dcdc819b93989992069b3244149855f..36c4fa6d5ca01a4c5bbbc57bce3f376c08925d36 100644
--- a/components/renderers/cpp/src/splatter.cu
+++ b/components/renderers/cpp/src/splatter.cu
@@ -127,6 +127,22 @@ __device__ inline float make(const float4 &v) {
return v.x;
}
+template <typename T>
+__device__ inline T make(const uchar4 &v);
+
+template <>
+__device__ inline float4 make(const uchar4 &v) {
+ return make_float4((float)v.x, (float)v.y, (float)v.z, (float)v.w);
+}
+
+template <typename T>
+__device__ inline T make(float v);
+
+template <>
+__device__ inline float make(float v) {
+ return v;
+}
+
/*
* Pass 1b: Expand splats to full size and merge
*/
@@ -302,15 +318,45 @@ __device__ inline uchar4 generateInput(const uchar4 &in, const SplatParams ¶
in;
}
+template <typename A, typename B>
+__device__ inline B weightInput(const A &in, float weight) {
+ return in * weight;
+}
+
+template <>
+__device__ inline float4 weightInput(const uchar4 &in, float weight) {
+ return make_float4(
+ (float)in.x * weight,
+ (float)in.y * weight,
+ (float)in.z * weight,
+ (float)in.w * weight);
+}
+
+template <typename T>
+__device__ inline void accumulateOutput(TextureObject<T> &out, TextureObject<float> &contrib, const uint2 &pos, const T &in, float w) {
+ atomicAdd(&out(pos.x, pos.y), in);
+ atomicAdd(&contrib(pos.x, pos.y), w);
+}
+
+template <>
+__device__ inline void accumulateOutput(TextureObject<float4> &out, TextureObject<float> &contrib, const uint2 &pos, const float4 &in, float w) {
+ atomicAdd((float*)&out(pos.x, pos.y), in.x);
+ atomicAdd(((float*)&out(pos.x, pos.y))+1, in.y);
+ atomicAdd(((float*)&out(pos.x, pos.y))+2, in.z);
+ atomicAdd(((float*)&out(pos.x, pos.y))+3, in.w);
+ atomicAdd(&contrib(pos.x, pos.y), w);
+}
+
/*
* Pass 2: Accumulate attribute contributions if the points pass a visibility test.
*/
- template <typename T>
+ template <typename A, typename B>
__global__ void dibr_attribute_contrib_kernel(
- TextureObject<T> in, // Attribute input
+ TextureObject<A> in, // Attribute input
TextureObject<float4> points, // Original 3D points
TextureObject<int> depth_in, // Virtual depth map
- TextureObject<T> out, // Accumulated output
+ TextureObject<B> out, // Accumulated output
+ TextureObject<float> contrib,
SplatParams params) {
const int x = (blockIdx.x*blockDim.x + threadIdx.x);
@@ -328,25 +374,26 @@ __global__ void dibr_attribute_contrib_kernel(
if (screenPos.x >= depth_in.width() || screenPos.y >= depth_in.height()) return;
// Is this point near the actual surface and therefore a contributor?
- const int d = depth_in.tex2D((int)screenPos.x, (int)screenPos.y);
-
- const T input = generateInput(in.tex2D(x, y), params, worldPos);
+ const float d = (float)depth_in.tex2D((int)screenPos.x, (int)screenPos.y) / 1000.0f;
- //const float3 nearest = params.camera.screenToCam((int)(screenPos.x),(int)(screenPos.y),d);
+ const A input = generateInput(in.tex2D(x, y), params, worldPos);
+ const float weight = ftl::cuda::weighting(fabs(camPos.z - d), 0.01f);
+ const B weighted = make<B>(input) * weight; //weightInput(input, weight);
- //const float l = length(nearest - camPos);
- if (d == (int)(camPos.z*1000.0f)) {
- out(screenPos.x, screenPos.y) = input;
+ if (weight > 0.0f) {
+ accumulateOutput(out, contrib, screenPos, weighted, weight);
+ //out(screenPos.x, screenPos.y) = input;
}
}
-template <typename T>
+template <typename A, typename B>
void ftl::cuda::dibr_attribute(
- TextureObject<T> &in,
+ TextureObject<A> &in,
TextureObject<float4> &points, // Original 3D points
TextureObject<int> &depth_in, // Virtual depth map
- TextureObject<T> &out, // Accumulated output
+ TextureObject<B> &out, // Accumulated output
+ TextureObject<float> &contrib,
SplatParams ¶ms, cudaStream_t stream) {
const dim3 gridSize((in.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (in.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
@@ -355,115 +402,68 @@ void ftl::cuda::dibr_attribute(
in,
points,
depth_in,
- out,
+ out,
+ contrib,
params
);
cudaSafeCall( cudaGetLastError() );
}
-template void ftl::cuda::dibr_attribute<uchar4>(
+template void ftl::cuda::dibr_attribute(
ftl::cuda::TextureObject<uchar4> &in, // Original colour image
ftl::cuda::TextureObject<float4> &points, // Original 3D points
ftl::cuda::TextureObject<int> &depth_in, // Virtual depth map
- ftl::cuda::TextureObject<uchar4> &out, // Accumulated output
+ ftl::cuda::TextureObject<float4> &out, // Accumulated output
+ ftl::cuda::TextureObject<float> &contrib,
ftl::render::SplatParams ¶ms, cudaStream_t stream);
-template void ftl::cuda::dibr_attribute<float>(
+template void ftl::cuda::dibr_attribute(
ftl::cuda::TextureObject<float> &in, // Original colour image
ftl::cuda::TextureObject<float4> &points, // Original 3D points
ftl::cuda::TextureObject<int> &depth_in, // Virtual depth map
ftl::cuda::TextureObject<float> &out, // Accumulated output
+ ftl::cuda::TextureObject<float> &contrib,
ftl::render::SplatParams ¶ms, cudaStream_t stream);
-template void ftl::cuda::dibr_attribute<float4>(
+template void ftl::cuda::dibr_attribute(
ftl::cuda::TextureObject<float4> &in, // Original colour image
ftl::cuda::TextureObject<float4> &points, // Original 3D points
ftl::cuda::TextureObject<int> &depth_in, // Virtual depth map
ftl::cuda::TextureObject<float4> &out, // Accumulated output
+ ftl::cuda::TextureObject<float> &contrib,
ftl::render::SplatParams ¶ms, cudaStream_t stream);
//==============================================================================
-/*__global__ void dibr_normalise_kernel(
- TextureObject<float4> colour_in,
- TextureObject<uchar4> colour_out,
- //TextureObject<float4> normals,
- TextureObject<float> contribs) {
- const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
- const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
-
- if (x < colour_in.width() && y < colour_in.height()) {
- const float4 colour = colour_in.tex2D((int)x,(int)y);
- //const float4 normal = normals.tex2D((int)x,(int)y);
- const float contrib = contribs.tex2D((int)x,(int)y);
-
- if (contrib > 0.0f) {
- colour_out(x,y) = make_uchar4(colour.x / contrib, colour.y / contrib, colour.z / contrib, 0);
- //normals(x,y) = normal / contrib;
- }
- }
-}
-
+template <typename A, typename B>
__global__ void dibr_normalise_kernel(
- TextureObject<float4> colour_in,
- TextureObject<float> colour_out,
- //TextureObject<float4> normals,
+ TextureObject<A> in,
+ TextureObject<B> out,
TextureObject<float> contribs) {
const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
- if (x < colour_in.width() && y < colour_in.height()) {
- const float4 colour = colour_in.tex2D((int)x,(int)y);
+ if (x < in.width() && y < in.height()) {
+ const A a = in.tex2D((int)x,(int)y);
//const float4 normal = normals.tex2D((int)x,(int)y);
const float contrib = contribs.tex2D((int)x,(int)y);
if (contrib > 0.0f) {
- colour_out(x,y) = colour.x / contrib;
+ out(x,y) = make<B>(a / contrib);
//normals(x,y) = normal / contrib;
}
}
}
-__global__ void dibr_normalise_kernel(
- TextureObject<float4> colour_in,
- TextureObject<float4> colour_out,
- //TextureObject<float4> normals,
- TextureObject<float> contribs) {
- const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
- const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
-
- if (x < colour_in.width() && y < colour_in.height()) {
- const float4 colour = colour_in.tex2D((int)x,(int)y);
- //const float4 normal = normals.tex2D((int)x,(int)y);
- const float contrib = contribs.tex2D((int)x,(int)y);
-
- if (contrib > 0.0f) {
- colour_out(x,y) = make_float4(colour.x / contrib, colour.y / contrib, colour.z / contrib, 0);
- //normals(x,y) = normal / contrib;
- }
- }
-}
-
-void ftl::cuda::dibr_normalise(TextureObject<float4> &colour_in, TextureObject<uchar4> &colour_out, TextureObject<float> &contribs, cudaStream_t stream) {
- const dim3 gridSize((colour_in.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (colour_in.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
- const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
-
- dibr_normalise_kernel<<<gridSize, blockSize, 0, stream>>>(colour_in, colour_out, contribs);
- cudaSafeCall( cudaGetLastError() );
-}
-
-void ftl::cuda::dibr_normalise(TextureObject<float4> &colour_in, TextureObject<float> &colour_out, TextureObject<float> &contribs, cudaStream_t stream) {
- const dim3 gridSize((colour_in.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (colour_in.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
+template <typename A, typename B>
+void ftl::cuda::dibr_normalise(TextureObject<A> &in, TextureObject<B> &out, TextureObject<float> &contribs, cudaStream_t stream) {
+ const dim3 gridSize((in.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (in.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
- dibr_normalise_kernel<<<gridSize, blockSize, 0, stream>>>(colour_in, colour_out, contribs);
+ dibr_normalise_kernel<<<gridSize, blockSize, 0, stream>>>(in, out, contribs);
cudaSafeCall( cudaGetLastError() );
}
-void ftl::cuda::dibr_normalise(TextureObject<float4> &colour_in, TextureObject<float4> &colour_out, TextureObject<float> &contribs, cudaStream_t stream) {
- const dim3 gridSize((colour_in.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (colour_in.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
- const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
-
- dibr_normalise_kernel<<<gridSize, blockSize, 0, stream>>>(colour_in, colour_out, contribs);
- cudaSafeCall( cudaGetLastError() );
-}*/
+template void ftl::cuda::dibr_normalise<float4,uchar4>(TextureObject<float4> &in, TextureObject<uchar4> &out, TextureObject<float> &contribs, cudaStream_t stream);
+template void ftl::cuda::dibr_normalise<float,float>(TextureObject<float> &in, TextureObject<float> &out, TextureObject<float> &contribs, cudaStream_t stream);
+template void ftl::cuda::dibr_normalise<float4,float4>(TextureObject<float4> &in, TextureObject<float4> &out, TextureObject<float> &contribs, cudaStream_t stream);
diff --git a/components/renderers/cpp/src/splatter_cuda.hpp b/components/renderers/cpp/src/splatter_cuda.hpp
index 1888a586720d388567a6c4d48f4a29b6b50848ac..463beefb18b54879a401580083cf21ca58b5d38f 100644
--- a/components/renderers/cpp/src/splatter_cuda.hpp
+++ b/components/renderers/cpp/src/splatter_cuda.hpp
@@ -23,13 +23,21 @@ namespace cuda {
ftl::cuda::TextureObject<T> &colour_out,
const ftl::render::SplatParams ¶ms, cudaStream_t stream);
- template <typename T>
+ template <typename A, typename B>
void dibr_attribute(
- ftl::cuda::TextureObject<T> &in, // Original colour image
+ ftl::cuda::TextureObject<A> &in, // Original colour image
ftl::cuda::TextureObject<float4> &points, // Original 3D points
ftl::cuda::TextureObject<int> &depth_in, // Virtual depth map
- ftl::cuda::TextureObject<T> &out, // Accumulated output
+ ftl::cuda::TextureObject<B> &out, // Accumulated output
+ ftl::cuda::TextureObject<float> &contrib,
ftl::render::SplatParams ¶ms, cudaStream_t stream);
+
+ template <typename A, typename B>
+ void dibr_normalise(
+ ftl::cuda::TextureObject<A> &in,
+ ftl::cuda::TextureObject<B> &out,
+ ftl::cuda::TextureObject<float> &contribs,
+ cudaStream_t stream);
}
}