Skip to content
Snippets Groups Projects
Commit 3719d74e authored by Nicolas Pope's avatar Nicolas Pope
Browse files

Accum the raw data instead of exact match

parent 6327e630
No related branches found
No related tags found
1 merge request!123Implements #189 using density to estimate radius
Show whitespace changes
Inline Side-by-side
applications/reconstruct/src/ilw/ilw.cu
+ 3
− 3
View file @ 3719d74e
......@@ -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(
......
components/renderers/cpp/src/splat_render.cpp
+ 48
− 11
View file @ 3719d74e
......@@ -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);
......
components/renderers/cpp/src/splatter.cu
+ 82
− 82
View file @ 3719d74e
......@@ -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 &para
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 &params, 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);
......@@ -356,114 +403,67 @@ void ftl::cuda::dibr_attribute(
points,
depth_in,
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 &params, 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 &params, 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 &params, 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);
components/renderers/cpp/src/splatter_cuda.hpp
+ 11
− 3
View file @ 3719d74e
......@@ -23,13 +23,21 @@ namespace cuda {
ftl::cuda::TextureObject<T> &colour_out,
const ftl::render::SplatParams &params, 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 &params, 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);
}
}
......
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Please register or to comment