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

Fix for point render

parent f44a0d4c
No related branches found
No related tags found
1 merge request!175Resolves #251 rendering of point clouds
...@@ -103,6 +103,40 @@ using ftl::cuda::warpSum; ...@@ -103,6 +103,40 @@ using ftl::cuda::warpSum;
} }
} }
/*
* Pass 1: Directly render each camera into virtual view but with no upsampling
* for sparse points.
*/
__global__ void dibr_merge_kernel(TextureObject<float> depth,
TextureObject<int> depth_out,
float4x4 transform,
ftl::rgbd::Camera cam,
SplatParams params) {
const int x = blockIdx.x*blockDim.x + threadIdx.x;
const int y = blockIdx.y*blockDim.y + threadIdx.y;
const float d0 = depth.tex2D(x, y);
if (d0 <= cam.minDepth || d0 >= cam.maxDepth) return;
const float3 camPos = transform * cam.screenToCam(x,y,d0);
//if (worldPos.x == MINF || (!(params.m_flags & ftl::render::kShowDisconMask) && worldPos.w < 0.0f)) return;
// Find the virtual screen position of current point
//const float3 camPos = params.m_viewMatrix * make_float3(worldPos);
//if (camPos.z < params.camera.minDepth) return;
//if (camPos.z > params.camera.maxDepth) return;
const float d = camPos.z;
const uint2 screenPos = params.camera.camToScreen<uint2>(camPos);
const unsigned int cx = screenPos.x;
const unsigned int cy = screenPos.y;
if (d > params.camera.minDepth && d < params.camera.maxDepth && cx < depth.width() && cy < depth.height()) {
// Transform estimated point to virtual cam space and output z
atomicMin(&depth_out(cx,cy), d * 100000.0f);
}
}
void ftl::cuda::dibr_merge(TextureObject<float4> &points, TextureObject<float4> &normals, TextureObject<int> &depth, SplatParams params, bool culling, cudaStream_t stream) { void ftl::cuda::dibr_merge(TextureObject<float4> &points, TextureObject<float4> &normals, TextureObject<int> &depth, SplatParams params, bool culling, cudaStream_t stream) {
const dim3 gridSize((depth.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (depth.height() + T_PER_BLOCK - 1)/T_PER_BLOCK); const dim3 gridSize((depth.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (depth.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK); const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
...@@ -120,6 +154,14 @@ void ftl::cuda::dibr_merge(TextureObject<float4> &points, TextureObject<int> &de ...@@ -120,6 +154,14 @@ void ftl::cuda::dibr_merge(TextureObject<float4> &points, TextureObject<int> &de
cudaSafeCall( cudaGetLastError() ); cudaSafeCall( cudaGetLastError() );
} }
void ftl::cuda::dibr_merge(TextureObject<float> &depth, TextureObject<int> &depth_out, const float4x4 &transform, const ftl::rgbd::Camera &cam, SplatParams params, cudaStream_t stream) {
const dim3 gridSize((depth.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (depth.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
dibr_merge_kernel<<<gridSize, blockSize, 0, stream>>>(depth, depth_out, transform, cam, params);
cudaSafeCall( cudaGetLastError() );
}
//============================================================================== //==============================================================================
......
...@@ -43,6 +43,14 @@ namespace cuda { ...@@ -43,6 +43,14 @@ namespace cuda {
ftl::render::SplatParams params, ftl::render::SplatParams params,
cudaStream_t stream); cudaStream_t stream);
void dibr_merge(
ftl::cuda::TextureObject<float> &depth,
ftl::cuda::TextureObject<int> &depth_out,
const float4x4 &transform,
const ftl::rgbd::Camera &cam,
ftl::render::SplatParams params,
cudaStream_t stream);
template <typename T> template <typename T>
void splat( void splat(
ftl::cuda::TextureObject<float4> &normals, ftl::cuda::TextureObject<float4> &normals,
......
...@@ -300,9 +300,13 @@ void Triangular::_dibr(ftl::rgbd::Frame &out, cudaStream_t stream) { ...@@ -300,9 +300,13 @@ void Triangular::_dibr(ftl::rgbd::Frame &out, cudaStream_t stream) {
continue; continue;
} }
auto transform = params_.m_viewMatrix * MatrixConversion::toCUDA(s->getPose().cast<float>());
ftl::cuda::dibr_merge( ftl::cuda::dibr_merge(
f.createTexture<float4>(Channel::Points), f.createTexture<float>(Channel::Depth),
temp_.createTexture<int>(Channel::Depth2), temp_.createTexture<int>(Channel::Depth2),
transform,
s->parameters(),
params_, stream params_, stream
); );
} }
......
0% Loading or .
You are about to add 0 people to the discussion. Proceed with caution.
Please register or to comment