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

Fix for use of wrong depth in render

parent 3533f450
No related branches found
No related tags found
1 merge request!151Implements #216 triangle renderer
Pipeline #15996 passed
......@@ -13,6 +13,7 @@ using ftl::render::SplatParams;
* Convert source screen position to output screen coordinates.
*/
__global__ void screen_coord_kernel(TextureObject<float> depth,
TextureObject<float> depth_out,
TextureObject<short2> screen_out, SplatParams params, float4x4 pose, Camera camera) {
const int x = blockIdx.x*blockDim.x + threadIdx.x;
const int y = blockIdx.y*blockDim.y + threadIdx.y;
......@@ -31,12 +32,13 @@ using ftl::render::SplatParams;
if (camPos.z < params.camera.minDepth || camPos.z > params.camera.maxDepth || screenPos.x >= params.camera.width || screenPos.y >= params.camera.height)
screenPos = make_uint2(30000,30000);
screen_out(x,y) = make_short2(screenPos.x, screenPos.y);
depth_out(x,y) = camPos.z;
}
void ftl::cuda::screen_coord(TextureObject<float> &depth, TextureObject<short2> &screen_out, const SplatParams &params, const float4x4 &pose, const Camera &camera, cudaStream_t stream) {
void ftl::cuda::screen_coord(TextureObject<float> &depth, TextureObject<float> &depth_out, TextureObject<short2> &screen_out, const SplatParams &params, const float4x4 &pose, const Camera &camera, 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);
screen_coord_kernel<<<gridSize, blockSize, 0, stream>>>(depth, screen_out, params, pose, camera);
screen_coord_kernel<<<gridSize, blockSize, 0, stream>>>(depth, depth_out, screen_out, params, pose, camera);
cudaSafeCall( cudaGetLastError() );
}
......@@ -278,12 +278,13 @@ void Splatter::_dibr(cudaStream_t stream) {
ftl::cuda::screen_coord(
f.createTexture<float>(Channel::Depth),
f.createTexture<float>(Channel::Depth2, Format<float>(f.get<GpuMat>(Channel::Depth).size())),
f.createTexture<short2>(Channel::Screen, Format<short2>(f.get<GpuMat>(Channel::Depth).size())),
params_, pose, s->parameters(), stream
);
ftl::cuda::triangle_render1(
f.getTexture<float>(Channel::Depth),
f.getTexture<float>(Channel::Depth2),
temp_.createTexture<int>(Channel::Depth2),
f.getTexture<short2>(Channel::Screen),
params_, stream
......
......@@ -8,6 +8,7 @@ namespace ftl {
namespace cuda {
void screen_coord(
ftl::cuda::TextureObject<float> &depth,
ftl::cuda::TextureObject<float> &depth_out,
ftl::cuda::TextureObject<short2> &screen_out,
const ftl::render::SplatParams &params,
const float4x4 &pose,
......
0% Loading or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment