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

Initial working back point cull

parent 50d8d9a6
No related branches found
No related tags found
1 merge request!117Implements #181 back facing cull
This commit is part of merge request !117. Comments created here will be created in the context of that merge request.
......@@ -76,6 +76,7 @@ void Splatter::renderChannel(
bool is_4chan = scene_->frames[0].get<GpuMat>(channel).type() == CV_32FC4;
// Render each camera into virtual view
// TODO: Move out of renderChannel, this is a common step to all channels
for (size_t i=0; i < scene_->frames.size(); ++i) {
auto &f = scene_->frames[i];
auto *s = scene_->sources[i];
......@@ -87,6 +88,7 @@ void Splatter::renderChannel(
ftl::cuda::dibr_merge(
f.createTexture<float4>(Channel::Points),
f.createTexture<float4>(Channel::Normals),
temp_.getTexture<int>(Channel::Depth),
params, stream
);
......
......@@ -20,13 +20,27 @@ using ftl::render::SplatParams;
* Pass 1: Directly render each camera into virtual view but with no upsampling
* for sparse points.
*/
__global__ void dibr_merge_kernel(TextureObject<float4> points, TextureObject<int> depth, SplatParams params) {
__global__ void dibr_merge_kernel(TextureObject<float4> points,
TextureObject<float4> normals,
TextureObject<int> depth, SplatParams params) {
const int x = blockIdx.x*blockDim.x + threadIdx.x;
const int y = blockIdx.y*blockDim.y + threadIdx.y;
const float4 worldPos = points.tex2D(x, y);
if (worldPos.x == MINF || (!(params.m_flags & ftl::render::kShowDisconMask) && worldPos.w < 0.0f)) return;
float3 ray = params.m_viewMatrixInverse.getFloat3x3() * params.camera.screenToCam(x,y,1.0f);
ray = ray / length(ray);
float3 n = make_float3(normals.tex2D((int)x,(int)y));
float l = length(n);
if (l == 0) {
return;
}
n /= l;
const float facing = dot(ray, n);
if (facing <= 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;
......@@ -43,11 +57,11 @@ using ftl::render::SplatParams;
}
}
void ftl::cuda::dibr_merge(TextureObject<float4> &points, TextureObject<int> &depth, SplatParams params, cudaStream_t stream) {
void ftl::cuda::dibr_merge(TextureObject<float4> &points, TextureObject<float4> &normals, TextureObject<int> &depth, 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>>>(points, depth, params);
dibr_merge_kernel<<<gridSize, blockSize, 0, stream>>>(points, normals, depth, params);
cudaSafeCall( cudaGetLastError() );
}
......
......@@ -8,6 +8,7 @@ namespace ftl {
namespace cuda {
void dibr_merge(
ftl::cuda::TextureObject<float4> &points,
ftl::cuda::TextureObject<float4> &normals,
ftl::cuda::TextureObject<int> &depth,
ftl::render::SplatParams params,
cudaStream_t stream);
......
0% Loading or .
You are about to add 0 people to the discussion. Proceed with caution.
Please register or to comment