diff --git a/components/renderers/cpp/src/splat_render.cpp b/components/renderers/cpp/src/splat_render.cpp
index 0619192370d2bca13883c0dfa93c0f9bf8dab4cd..bfb627f6b2f0ca352d7635a9f418c3bec5af4166 100644
--- a/components/renderers/cpp/src/splat_render.cpp
+++ b/components/renderers/cpp/src/splat_render.cpp
@@ -73,12 +73,12 @@ void Splatter::renderChannel(
cv::cuda::Stream cvstream = cv::cuda::StreamAccessor::wrapStream(stream);
temp_.get<GpuMat>(Channel::Depth).setTo(cv::Scalar(0x7FFFFFFF), cvstream);
temp_.get<GpuMat>(Channel::Depth2).setTo(cv::Scalar(0x7FFFFFFF), cvstream);
- temp_.get<GpuMat>(Channel::Colour).setTo(cv::Scalar(0.0f,0.0f,0.0f,0.0f), cvstream);
- temp_.get<GpuMat>(Channel::Contribution).setTo(cv::Scalar(0.0f), cvstream);
+ //temp_.get<GpuMat>(Channel::Colour).setTo(cv::Scalar(0.0f,0.0f,0.0f,0.0f), cvstream);
+ //temp_.get<GpuMat>(Channel::Contribution).setTo(cv::Scalar(0.0f), cvstream);
if (scene_->frames.size() < 1) return;
- bool is_float = scene_->frames[0].get<GpuMat>(channel).type() == CV_32F; //ftl::rgbd::isFloatChannel(channel);
- bool is_4chan = scene_->frames[0].get<GpuMat>(channel).type() == CV_32FC4;
+ bool is_float = out.get<GpuMat>(channel).type() == CV_32F; //ftl::rgbd::isFloatChannel(channel);
+ bool is_4chan = out.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
@@ -101,8 +101,8 @@ void Splatter::renderChannel(
//LOG(INFO) << "DIBR DONE";
}
- temp_.createTexture<float4>(Channel::Colour);
- temp_.createTexture<float>(Channel::Contribution);
+ //temp_.createTexture<float4>(Channel::Colour);
+ //temp_.createTexture<float>(Channel::Contribution);
out.create<GpuMat>(Channel::Normals, Format<float4>(params.camera.width, params.camera.height));
// Create normals first
@@ -116,11 +116,21 @@ void Splatter::renderChannel(
);
}
- temp_.get<GpuMat>(Channel::Colour).setTo(cv::Scalar(0.0f,0.0f,0.0f,0.0f), cvstream);
- temp_.get<GpuMat>(Channel::Colour2).setTo(cv::Scalar(0,0,0,0), cvstream);
- temp_.get<GpuMat>(Channel::Contribution).setTo(cv::Scalar(0.0f), cvstream);
+ //temp_.get<GpuMat>(Channel::Colour).setTo(cv::Scalar(0.0f,0.0f,0.0f,0.0f), cvstream);
+ //temp_.get<GpuMat>(Channel::Contribution).setTo(cv::Scalar(0.0f), cvstream);
- // Create colours first
+ if (is_4chan) {
+ temp_.create<GpuMat>(Channel::Colour2, Format<float4>(params.camera.width, params.camera.height));
+ temp_.get<GpuMat>(Channel::Colour2).setTo(cv::Scalar(0.0f,0.0f,0.0f,0.0f), cvstream);
+ } else if (is_float) {
+ temp_.create<GpuMat>(Channel::Colour2, Format<float>(params.camera.width, params.camera.height));
+ temp_.get<GpuMat>(Channel::Colour2).setTo(cv::Scalar(0.0f), cvstream);
+ } else {
+ temp_.create<GpuMat>(Channel::Colour2, Format<uchar4>(params.camera.width, params.camera.height));
+ temp_.get<GpuMat>(Channel::Colour2).setTo(cv::Scalar(0,0,0,0), cvstream);
+ }
+
+ // Create attribute first
for (auto &f : scene_->frames) {
// Convert colour from BGR to BGRA if needed
if (f.get<GpuMat>(channel).type() == CV_8UC3) {
@@ -131,26 +141,64 @@ void Splatter::renderChannel(
cv::cuda::cvtColor(tmp,col, cv::COLOR_BGR2BGRA);
}
- ftl::cuda::dibr_attribute(
- f.createTexture<uchar4>(Channel::Left),
- f.createTexture<float4>(Channel::Points),
- temp_.getTexture<int>(Channel::Depth2),
- temp_.createTexture<uchar4>(Channel::Colour2),
- params, stream
- );
+ if (is_4chan) {
+ ftl::cuda::dibr_attribute(
+ f.createTexture<float4>(channel),
+ f.createTexture<float4>(Channel::Points),
+ temp_.getTexture<int>(Channel::Depth2),
+ temp_.createTexture<float4>(Channel::Colour2),
+ params, stream
+ );
+ } else if (is_float) {
+ ftl::cuda::dibr_attribute(
+ f.createTexture<float>(channel),
+ f.createTexture<float4>(Channel::Points),
+ temp_.getTexture<int>(Channel::Depth2),
+ temp_.createTexture<float>(Channel::Colour2),
+ params, stream
+ );
+ } else {
+ ftl::cuda::dibr_attribute(
+ f.createTexture<uchar4>(channel),
+ f.createTexture<float4>(Channel::Points),
+ temp_.getTexture<int>(Channel::Depth2),
+ temp_.createTexture<uchar4>(Channel::Colour2),
+ params, stream
+ );
+ }
}
//out.get<GpuMat>(Channel::Left).setTo(cv::Scalar(0,0,0,0), cvstream);
// Now splat the points
- ftl:cuda::splat(
- out.getTexture<float4>(Channel::Normals),
- temp_.getTexture<uchar4>(Channel::Colour2),
- temp_.getTexture<int>(Channel::Depth2),
- out.createTexture<float>(Channel::Depth),
- out.createTexture<uchar4>(Channel::Left),
- params, stream
- );
+ if (is_4chan) {
+ ftl::cuda::splat(
+ out.getTexture<float4>(Channel::Normals),
+ temp_.getTexture<float4>(Channel::Colour2),
+ temp_.getTexture<int>(Channel::Depth2),
+ out.createTexture<float>(Channel::Depth),
+ out.createTexture<float4>(channel),
+ params, stream
+ );
+ } else if (is_float) {
+ ftl::cuda::splat(
+ out.getTexture<float4>(Channel::Normals),
+ temp_.getTexture<float>(Channel::Colour2),
+ temp_.getTexture<int>(Channel::Depth2),
+ out.createTexture<float>(Channel::Depth),
+ out.createTexture<float>(channel),
+ params, stream
+ );
+ } else {
+ ftl::cuda::splat(
+ out.getTexture<float4>(Channel::Normals),
+ temp_.getTexture<uchar4>(Channel::Colour2),
+ temp_.getTexture<int>(Channel::Depth2),
+ out.createTexture<float>(Channel::Depth),
+ out.createTexture<uchar4>(channel),
+ params, stream
+ );
+ }
}
bool Splatter::render(ftl::rgbd::VirtualSource *src, ftl::rgbd::Frame &out, cudaStream_t stream) {
@@ -166,38 +214,14 @@ bool Splatter::render(ftl::rgbd::VirtualSource *src, ftl::rgbd::Frame &out, cuda
out.create<GpuMat>(Channel::Colour, Format<uchar4>(camera.width, camera.height));
// FIXME: Use source resolutions, not virtual resolution
- temp_.create<GpuMat>(Channel::Colour, Format<float4>(camera.width, camera.height));
- temp_.create<GpuMat>(Channel::Colour2, Format<uchar4>(camera.width, camera.height));
- temp_.create<GpuMat>(Channel::Contribution, Format<float>(camera.width, camera.height));
+ //temp_.create<GpuMat>(Channel::Colour, Format<float4>(camera.width, camera.height));
+ //temp_.create<GpuMat>(Channel::Contribution, Format<float>(camera.width, camera.height));
temp_.create<GpuMat>(Channel::Depth, Format<int>(camera.width, camera.height));
temp_.create<GpuMat>(Channel::Depth2, Format<int>(camera.width, camera.height));
temp_.create<GpuMat>(Channel::Normals, Format<float4>(camera.width, camera.height));
cv::cuda::Stream cvstream = cv::cuda::StreamAccessor::wrapStream(stream);
- // Create buffers if they don't exist
- /*if ((unsigned int)depth1_.width() != camera.width || (unsigned int)depth1_.height() != camera.height) {
- depth1_ = ftl::cuda::TextureObject<int>(camera.width, camera.height);
- }
- if ((unsigned int)depth3_.width() != camera.width || (unsigned int)depth3_.height() != camera.height) {
- depth3_ = ftl::cuda::TextureObject<int>(camera.width, camera.height);
- }
- if ((unsigned int)colour1_.width() != camera.width || (unsigned int)colour1_.height() != camera.height) {
- colour1_ = ftl::cuda::TextureObject<uchar4>(camera.width, camera.height);
- }
- if ((unsigned int)colour_tmp_.width() != camera.width || (unsigned int)colour_tmp_.height() != camera.height) {
- colour_tmp_ = ftl::cuda::TextureObject<float4>(camera.width, camera.height);
- }
- if ((unsigned int)normal1_.width() != camera.width || (unsigned int)normal1_.height() != camera.height) {
- normal1_ = ftl::cuda::TextureObject<float4>(camera.width, camera.height);
- }
- if ((unsigned int)depth2_.width() != camera.width || (unsigned int)depth2_.height() != camera.height) {
- depth2_ = ftl::cuda::TextureObject<float>(camera.width, camera.height);
- }
- if ((unsigned int)colour2_.width() != camera.width || (unsigned int)colour2_.height() != camera.height) {
- colour2_ = ftl::cuda::TextureObject<uchar4>(camera.width, camera.height);
- }*/
-
// Parameters object to pass to CUDA describing the camera
SplatParams params;
params.m_flags = 0;
@@ -267,7 +291,8 @@ bool Splatter::render(ftl::rgbd::VirtualSource *src, ftl::rgbd::Frame &out, cuda
renderChannel(params, out, Channel::Normals, stream);
// Convert normal to single float value
- ftl::cuda::normal_visualise(out.getTexture<float4>(Channel::Normals), temp_.getTexture<float>(Channel::Contribution), camera, params.m_viewMatrixInverse, stream);
+ temp_.create<GpuMat>(Channel::Contribution, Format<float>(camera.width, camera.height));
+ ftl::cuda::normal_visualise(out.getTexture<float4>(Channel::Normals), temp_.createTexture<float>(Channel::Contribution), camera, params.m_viewMatrixInverse, stream);
// Put in output as single float
cv::cuda::swap(temp_.get<GpuMat>(Channel::Contribution), out.create<GpuMat>(Channel::Normals));
diff --git a/components/renderers/cpp/src/splatter.cu b/components/renderers/cpp/src/splatter.cu
index 16ff3e7d48a8315412d5df7dc9c91c784a7080f4..a4f5d97ff91b62561cfa4fda990c44fc7790c72f 100644
--- a/components/renderers/cpp/src/splatter.cu
+++ b/components/renderers/cpp/src/splatter.cu
@@ -87,17 +87,57 @@ __device__ inline float4 make_float4(const uchar4 &c) {
return make_float4(c.x,c.y,c.z,c.w);
}
+__device__ inline float4 make_float4(const float4 &v) {
+ return v;
+}
+
+template <typename T>
+__device__ inline T make();
+
+template <>
+__device__ inline uchar4 make() {
+ return make_uchar4(0,0,0,0);
+}
+
+template <>
+__device__ inline float4 make() {
+ return make_float4(0.0f,0.0f,0.0f,0.0f);
+}
+
+template <>
+__device__ inline float make() {
+ return 0.0f;
+}
+
+template <typename T>
+__device__ inline T make(const float4 &);
+
+template <>
+__device__ inline uchar4 make(const float4 &v) {
+ return make_uchar4((int)v.x, (int)v.y, (int)v.z, (int)v.w);
+}
+
+template <>
+__device__ inline float4 make(const float4 &v) {
+ return v;
+}
+
+template <>
+__device__ inline float make(const float4 &v) {
+ return v.x;
+}
+
/*
* Pass 1b: Expand splats to full size and merge
*/
- template <int SEARCH_DIAMETER>
+ template <int SEARCH_DIAMETER, typename T>
__global__ void splat_kernel(
//TextureObject<float4> points, // Original 3D points
TextureObject<float4> normals,
- TextureObject<uchar4> colour_in,
+ TextureObject<T> in,
TextureObject<int> depth_in, // Virtual depth map
TextureObject<float> depth_out, // Accumulated output
- TextureObject<uchar4> colour_out,
+ TextureObject<T> out,
//ftl::rgbd::Camera camera,
//float4x4 pose_inv,
SplatParams params) {
@@ -124,7 +164,7 @@ __device__ inline float4 make_float4(const uchar4 &c) {
struct Result {
float weight;
float depth;
- uchar4 colour;
+ T attr;
};
Result results[(SEARCH_DIAMETER*SEARCH_DIAMETER) / WARP_SIZE];
@@ -135,7 +175,7 @@ __device__ inline float4 make_float4(const uchar4 &c) {
const float u = (i % SEARCH_DIAMETER) - (SEARCH_DIAMETER / 2);
const float v = (i / SEARCH_DIAMETER) - (SEARCH_DIAMETER / 2);
- results[i/WARP_SIZE] = {0.0f, 0.0f, make_uchar4(0,0,0,0)};
+ results[i/WARP_SIZE] = {0.0f, 0.0f, make<T>()};
// Use the depth buffer to determine this pixels 3D position in camera space
const float d = ((float)depth_in.tex2D(x+u, y+v)/1000.0f);
@@ -173,7 +213,7 @@ __device__ inline float4 make_float4(const uchar4 &c) {
//depth += t * weight;
//contrib += weight;
depth = min(depth, t);
- results[i/WARP_SIZE] = {weight, t, colour_in.tex2D((int)x+u, (int)y+v)}; //make_float2(t, weight);
+ results[i/WARP_SIZE] = {weight, t, in.tex2D((int)x+u, (int)y+v)}; //make_float2(t, weight);
//atomicMin(&depth_out(x,y), (int)(depth * 1000.0f));
}
//}
@@ -183,41 +223,42 @@ __device__ inline float4 make_float4(const uchar4 &c) {
float adepth = 0.0f;
float contrib = 0.0f;
- float4 colour = make_float4(0.0f);
+ float4 attr = make_float4(0.0f);
// Loop over results array
for (int i=0; i<(SEARCH_DIAMETER*SEARCH_DIAMETER) / WARP_SIZE; ++i) {
if (results[i].depth - depth < 0.04f) {
adepth += results[i].depth * results[i].weight;
- colour += make_float4(results[i].colour) * results[i].weight;
+ attr += make_float4(results[i].attr) * results[i].weight;
contrib += results[i].weight;
}
}
// Sum all attributes and contributions
adepth = warpSum(adepth);
- colour.x = warpSum(colour.x);
- colour.y = warpSum(colour.y);
- colour.z = warpSum(colour.z);
+ attr.x = warpSum(attr.x);
+ attr.y = warpSum(attr.y);
+ attr.z = warpSum(attr.z);
contrib = warpSum(contrib);
if (lane == 0 && contrib > 0.0f) {
depth_out(x,y) = adepth / contrib;
- colour_out(x,y) = make_uchar4(colour.x / contrib, colour.y / contrib, colour.z / contrib, 255.0f);
+ out(x,y) = make<T>(attr / contrib);
}
}
+template <typename T>
void ftl::cuda::splat(
TextureObject<float4> &normals,
- TextureObject<uchar4> &colour_in,
+ TextureObject<T> &colour_in,
TextureObject<int> &depth_in, // Virtual depth map
TextureObject<float> &depth_out,
- TextureObject<uchar4> &colour_out,
+ TextureObject<T> &colour_out,
const SplatParams ¶ms, cudaStream_t stream) {
const dim3 gridSize((depth_in.width() + 2 - 1)/2, (depth_in.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
const dim3 blockSize(2*WARP_SIZE, T_PER_BLOCK);
- splat_kernel<8><<<gridSize, blockSize, 0, stream>>>(
+ splat_kernel<8,T><<<gridSize, blockSize, 0, stream>>>(
normals,
colour_in,
depth_in,
@@ -228,6 +269,30 @@ void ftl::cuda::splat(
cudaSafeCall( cudaGetLastError() );
}
+template void ftl::cuda::splat<uchar4>(
+ TextureObject<float4> &normals,
+ TextureObject<uchar4> &colour_in,
+ TextureObject<int> &depth_in, // Virtual depth map
+ TextureObject<float> &depth_out,
+ TextureObject<uchar4> &colour_out,
+ const SplatParams ¶ms, cudaStream_t stream);
+
+template void ftl::cuda::splat<float4>(
+ TextureObject<float4> &normals,
+ TextureObject<float4> &colour_in,
+ TextureObject<int> &depth_in, // Virtual depth map
+ TextureObject<float> &depth_out,
+ TextureObject<float4> &colour_out,
+ const SplatParams ¶ms, cudaStream_t stream);
+
+template void ftl::cuda::splat<float>(
+ TextureObject<float4> &normals,
+ TextureObject<float> &colour_in,
+ TextureObject<int> &depth_in, // Virtual depth map
+ TextureObject<float> &depth_out,
+ TextureObject<float> &colour_out,
+ const SplatParams ¶ms, cudaStream_t stream);
+
//==============================================================================
template <typename T>
diff --git a/components/renderers/cpp/src/splatter_cuda.hpp b/components/renderers/cpp/src/splatter_cuda.hpp
index 56cf4f58334102ab0b46029da2599a43d9bbd1ad..1888a586720d388567a6c4d48f4a29b6b50848ac 100644
--- a/components/renderers/cpp/src/splatter_cuda.hpp
+++ b/components/renderers/cpp/src/splatter_cuda.hpp
@@ -14,12 +14,13 @@ namespace cuda {
bool culling,
cudaStream_t stream);
+ template <typename T>
void splat(
ftl::cuda::TextureObject<float4> &normals,
- ftl::cuda::TextureObject<uchar4> &colour_in,
+ ftl::cuda::TextureObject<T> &colour_in,
ftl::cuda::TextureObject<int> &depth_in, // Virtual depth map
ftl::cuda::TextureObject<float> &depth_out,
- ftl::cuda::TextureObject<uchar4> &colour_out,
+ ftl::cuda::TextureObject<T> &colour_out,
const ftl::render::SplatParams ¶ms, cudaStream_t stream);
template <typename T>