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

Allow for point cloud mode

parent 5509f29f
No related branches found
No related tags found
1 merge request!151Implements #216 triangle renderer
...@@ -171,3 +171,104 @@ template void ftl::cuda::reproject( ...@@ -171,3 +171,104 @@ template void ftl::cuda::reproject(
const ftl::render::SplatParams &params, const ftl::render::SplatParams &params,
const ftl::rgbd::Camera &camera, const ftl::rgbd::Camera &camera,
const float4x4 &poseInv, cudaStream_t stream); const float4x4 &poseInv, cudaStream_t stream);
//==============================================================================
// Without normals
//==============================================================================
/*
* Pass 2: Accumulate attribute contributions if the points pass a visibility test.
*/
template <typename A, typename B>
__global__ void reprojection_kernel(
TextureObject<A> in, // Attribute input
TextureObject<float> depth_src,
TextureObject<int> depth_in, // Virtual depth map
TextureObject<B> out, // Accumulated output
TextureObject<float> contrib,
SplatParams params,
Camera camera, float4x4 poseInv) {
const int x = (blockIdx.x*blockDim.x + threadIdx.x);
const int y = blockIdx.y*blockDim.y + threadIdx.y;
const float d = (float)depth_in.tex2D((int)x, (int)y) / 1000.0f;
if (d < params.camera.minDepth || d > params.camera.maxDepth) return;
const float3 worldPos = params.m_viewMatrixInverse * params.camera.screenToCam(x, y, d);
//if (worldPos.x == MINF || (!(params.m_flags & ftl::render::kShowDisconMask) && worldPos.w < 0.0f)) return;
const float3 camPos = poseInv * worldPos;
if (camPos.z < camera.minDepth) return;
if (camPos.z > camera.maxDepth) return;
const uint2 screenPos = camera.camToScreen<uint2>(camPos);
// Not on screen so stop now...
if (screenPos.x >= depth_src.width() || screenPos.y >= depth_src.height()) return;
const float d2 = depth_src.tex2D((int)screenPos.x, (int)screenPos.y);
const A input = in.tex2D((int)screenPos.x, (int)screenPos.y); //generateInput(in.tex2D((int)screenPos.x, (int)screenPos.y), params, worldPos);
float weight = ftl::cuda::weighting(fabs(camPos.z - d2), 0.02f);
const B weighted = make<B>(input) * weight;
if (weight > 0.0f) {
accumulateOutput(out, contrib, make_uint2(x,y), weighted, weight);
//out(screenPos.x, screenPos.y) = input;
}
}
template <typename A, typename B>
void ftl::cuda::reproject(
TextureObject<A> &in,
TextureObject<float> &depth_src, // Original 3D points
TextureObject<int> &depth_in, // Virtual depth map
TextureObject<B> &out, // Accumulated output
TextureObject<float> &contrib,
const SplatParams &params,
const Camera &camera, const float4x4 &poseInv, cudaStream_t stream) {
const dim3 gridSize((out.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (out.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
reprojection_kernel<<<gridSize, blockSize, 0, stream>>>(
in,
depth_src,
depth_in,
out,
contrib,
params,
camera,
poseInv
);
cudaSafeCall( cudaGetLastError() );
}
template void ftl::cuda::reproject(
ftl::cuda::TextureObject<uchar4> &in, // Original colour image
ftl::cuda::TextureObject<float> &depth_src, // Original 3D points
ftl::cuda::TextureObject<int> &depth_in, // Virtual depth map
ftl::cuda::TextureObject<float4> &out, // Accumulated output
ftl::cuda::TextureObject<float> &contrib,
const ftl::render::SplatParams &params,
const ftl::rgbd::Camera &camera,
const float4x4 &poseInv, cudaStream_t stream);
template void ftl::cuda::reproject(
ftl::cuda::TextureObject<float> &in, // Original colour image
ftl::cuda::TextureObject<float> &depth_src, // Original 3D points
ftl::cuda::TextureObject<int> &depth_in, // Virtual depth map
ftl::cuda::TextureObject<float> &out, // Accumulated output
ftl::cuda::TextureObject<float> &contrib,
const ftl::render::SplatParams &params,
const ftl::rgbd::Camera &camera,
const float4x4 &poseInv, cudaStream_t stream);
template void ftl::cuda::reproject(
ftl::cuda::TextureObject<float4> &in, // Original colour image
ftl::cuda::TextureObject<float> &depth_src, // Original 3D points
ftl::cuda::TextureObject<int> &depth_in, // Virtual depth map
ftl::cuda::TextureObject<float4> &out, // Accumulated output
ftl::cuda::TextureObject<float> &contrib,
const ftl::render::SplatParams &params,
const ftl::rgbd::Camera &camera,
const float4x4 &poseInv, cudaStream_t stream);
...@@ -67,6 +67,17 @@ namespace cuda { ...@@ -67,6 +67,17 @@ namespace cuda {
const ftl::rgbd::Camera &camera, const ftl::rgbd::Camera &camera,
const float4x4 &poseInv, cudaStream_t stream); const float4x4 &poseInv, cudaStream_t stream);
template <typename A, typename B>
void reproject(
ftl::cuda::TextureObject<A> &in, // Original colour image
ftl::cuda::TextureObject<float> &depth_src, // Original 3D points
ftl::cuda::TextureObject<int> &depth_in, // Virtual depth map
ftl::cuda::TextureObject<B> &out, // Accumulated output
ftl::cuda::TextureObject<float> &contrib,
const ftl::render::SplatParams &params,
const ftl::rgbd::Camera &camera,
const float4x4 &poseInv, cudaStream_t stream);
template <typename A, typename B> template <typename A, typename B>
void dibr_normalise( void dibr_normalise(
ftl::cuda::TextureObject<A> &in, ftl::cuda::TextureObject<A> &in,
......
...@@ -212,6 +212,7 @@ void Triangular::__reprojectChannel(ftl::rgbd::Frame &output, ftl::codecs::Chann ...@@ -212,6 +212,7 @@ void Triangular::__reprojectChannel(ftl::rgbd::Frame &output, ftl::codecs::Chann
auto poseInv = MatrixConversion::toCUDA(s->getPose().cast<float>().inverse()); auto poseInv = MatrixConversion::toCUDA(s->getPose().cast<float>().inverse());
if (splat_) {
ftl::cuda::reproject( ftl::cuda::reproject(
f.createTexture<T>(in), f.createTexture<T>(in),
f.createTexture<float>(Channel::Depth), // TODO: Use depth? f.createTexture<float>(Channel::Depth), // TODO: Use depth?
...@@ -223,6 +224,19 @@ void Triangular::__reprojectChannel(ftl::rgbd::Frame &output, ftl::codecs::Chann ...@@ -223,6 +224,19 @@ void Triangular::__reprojectChannel(ftl::rgbd::Frame &output, ftl::codecs::Chann
s->parameters(), s->parameters(),
poseInv, stream poseInv, stream
); );
} else {
// Can't use normals with point cloud version
ftl::cuda::reproject(
f.createTexture<T>(in),
f.createTexture<float>(Channel::Depth), // TODO: Use depth?
temp_.getTexture<int>(Channel::Depth2),
temp_.createTexture<typename AccumSelector<T>::type>(AccumSelector<T>::channel),
temp_.getTexture<float>(Channel::Contribution),
params_,
s->parameters(),
poseInv, stream
);
}
} }
ftl::cuda::dibr_normalise( ftl::cuda::dibr_normalise(
......
0% Loading or .
You are about to add 0 people to the discussion. Proceed with caution.
Please register or to comment