diff --git a/components/renderers/cpp/include/ftl/cuda/normals.hpp b/components/renderers/cpp/include/ftl/cuda/normals.hpp
index 1b4dbc86c46a3fddb3d42b3cc0f8023ab3c62483..b9801743a5e505cf38a7b632127755f248e618c9 100644
--- a/components/renderers/cpp/include/ftl/cuda/normals.hpp
+++ b/components/renderers/cpp/include/ftl/cuda/normals.hpp
@@ -11,7 +11,8 @@ namespace cuda {
void normals(ftl::cuda::TextureObject<float4> &output,
ftl::cuda::TextureObject<float4> &temp,
ftl::cuda::TextureObject<float4> &input,
- const ftl::rgbd::Camera &camera, cudaStream_t stream);
+ const ftl::rgbd::Camera &camera,
+ const float3x3 &pose, cudaStream_t stream);
void normal_visualise(ftl::cuda::TextureObject<float4> &norm,
ftl::cuda::TextureObject<uchar4> &output,
diff --git a/components/renderers/cpp/src/normals.cu b/components/renderers/cpp/src/normals.cu
index de31404a62b8cf0f9d248f7b6f6a6a53cec8b01b..4d66869316bb3cc4f7f9bd56f889bf92ecc4635c 100644
--- a/components/renderers/cpp/src/normals.cu
+++ b/components/renderers/cpp/src/normals.cu
@@ -35,7 +35,7 @@ template <int RADIUS>
__global__ void smooth_normals_kernel(ftl::cuda::TextureObject<float4> norms,
ftl::cuda::TextureObject<float4> output,
ftl::cuda::TextureObject<float4> points,
- ftl::rgbd::Camera camera, float smoothing) {
+ ftl::rgbd::Camera camera, float3x3 pose, float smoothing) {
const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
@@ -66,7 +66,7 @@ __global__ void smooth_normals_kernel(ftl::cuda::TextureObject<float4> norms,
// Compute dot product of normal with camera to obtain measure of how
// well this point faces the source camera, a measure of confidence
- float3 ray = camera.screenToCam(x, y, 1.0f);
+ float3 ray = pose * camera.screenToCam(x, y, 1.0f);
ray = ray / length(ray);
nsum /= contrib;
nsum /= length(nsum);
@@ -77,14 +77,15 @@ __global__ void smooth_normals_kernel(ftl::cuda::TextureObject<float4> norms,
void ftl::cuda::normals(ftl::cuda::TextureObject<float4> &output,
ftl::cuda::TextureObject<float4> &temp,
ftl::cuda::TextureObject<float4> &input,
- const ftl::rgbd::Camera &camera, cudaStream_t stream) {
+ const ftl::rgbd::Camera &camera,
+ const float3x3 &pose,cudaStream_t stream) {
const dim3 gridSize((input.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (input.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
computeNormals_kernel<<<gridSize, blockSize, 0, stream>>>(temp, input);
cudaSafeCall( cudaGetLastError() );
- smooth_normals_kernel<3><<<gridSize, blockSize, 0, stream>>>(temp, output, input, camera, 0.04f);
+ smooth_normals_kernel<3><<<gridSize, blockSize, 0, stream>>>(temp, output, input, camera, pose, 0.04f);
cudaSafeCall( cudaGetLastError() );
#ifdef _DEBUG
diff --git a/components/renderers/cpp/src/splat_render.cpp b/components/renderers/cpp/src/splat_render.cpp
index c801c27ce573e988daf7a0087b305e882c267230..6b8736effb30895b9ce0f5e503acec78f44c7224 100644
--- a/components/renderers/cpp/src/splat_render.cpp
+++ b/components/renderers/cpp/src/splat_render.cpp
@@ -328,14 +328,15 @@ bool Splatter::render(ftl::rgbd::VirtualSource *src, ftl::rgbd::Frame &out, cuda
}
if (!f.hasChannel(Channel::Normals)) {
+ Eigen::Matrix4f matrix = s->getPose().cast<float>();
+ auto pose = MatrixConversion::toCUDA(matrix);
+
auto &g = f.get<GpuMat>(Channel::Colour);
ftl::cuda::normals(f.createTexture<float4>(Channel::Normals, Format<float4>(g.cols, g.rows)),
temp_.getTexture<float4>(Channel::Normals), // FIXME: Uses assumption of vcam res same as input res
- f.getTexture<float4>(Channel::Points), s->parameters(), stream);
+ f.getTexture<float4>(Channel::Points), s->parameters(), pose.getFloat3x3(), stream);
if (norm_filter_ > -0.1f) {
- Eigen::Matrix4f matrix = s->getPose().cast<float>();
- auto pose = MatrixConversion::toCUDA(matrix);
ftl::cuda::normal_filter(f.getTexture<float4>(Channel::Normals), f.getTexture<float4>(Channel::Points), s->parameters(), pose, norm_filter_, stream);
}
}
diff --git a/components/renderers/cpp/src/splatter.cu b/components/renderers/cpp/src/splatter.cu
index a4f5d97ff91b62561cfa4fda990c44fc7790c72f..3c55263d812229fe7701504af013deadfe04bada 100644
--- a/components/renderers/cpp/src/splatter.cu
+++ b/components/renderers/cpp/src/splatter.cu
@@ -197,19 +197,19 @@ __device__ inline float make(const float4 &v) {
//const float weight = ftl::cuda::spatialWeighting(worldPos, p, (camPos.z/params.camera.fx)); //*(camPos2.z/camera.fx));
//if (weight <= 0.0f) continue;
- float3 n = make_float3(normals.tex2D((int)(x+u), (int)(y+v)));
- const float l = length(n);
- if (l == 0.0f) continue;
- n /= l;
+ float4 n = normals.tex2D((int)(x+u), (int)(y+v));
+ //const float l = length(n);
+ //if (l == 0.0f) continue;
+ //n /= l;
// Does the ray intersect plane of splat?
float t = 1000.0f;
- if (ftl::cuda::intersectPlane(n, worldPos, origin, ray, t)) { //} && fabs(t-camPos.z) < 0.01f) {
+ if (ftl::cuda::intersectPlane(make_float3(n), worldPos, origin, ray, t)) { //} && fabs(t-camPos.z) < 0.01f) {
//t *= (params.m_viewMatrix.getFloat3x3() * ray).z;
t *= scale;
const float3 camPos3 = params.camera.screenToCam((int)(x),(int)(y),t);
- const float weight = ftl::cuda::spatialWeighting(camPos, camPos3, 2.0f*(camPos3.z/params.camera.fx)); //*(camPos2.z/camera.fx));
- if (weight == 0.0f) continue;
+ const float weight = ftl::cuda::spatialWeighting(camPos, camPos3, 2.0f*(camPos3.z/params.camera.fx)) * n.w * n.w;
+ if (weight <= 0.0f) continue;
//depth += t * weight;
//contrib += weight;
depth = min(depth, t);