diff --git a/applications/gui/src/src_window.cpp b/applications/gui/src/src_window.cpp index 328e29ad6b278fdbc7e676f18cb5b29538427673..a63861b24705a06866804e3bddd472dcd79ef317 100644 --- a/applications/gui/src/src_window.cpp +++ b/applications/gui/src/src_window.cpp @@ -220,7 +220,7 @@ bool SourceWindow::_processFrameset(ftl::rgbd::FrameSet &fs, bool fromstream) { void SourceWindow::_checkFrameSets(int id) { while (framesets_.size() <= id) { auto *p = ftl::config::create<ftl::operators::Graph>(screen_->root(), "pre_filters"); - //pre_pipeline_->append<ftl::operators::ColourChannels>("colour"); // Convert BGR to BGRA + //p->append<ftl::operators::ColourChannels>("colour"); // Convert BGR to BGRA p->append<ftl::operators::DetectAndTrack>("facedetection")->value("enabled", false); p->append<ftl::operators::ArUco>("aruco")->value("enabled", false); //p->append<ftl::operators::HFSmoother>("hfnoise"); diff --git a/components/common/cpp/include/ftl/cuda_common.hpp b/components/common/cpp/include/ftl/cuda_common.hpp index 4741d99d0af4ab20fe73878cb7b5787ef1c7c28f..8210fdb92a4bcd3cc820e92679ec3141aa861551 100644 --- a/components/common/cpp/include/ftl/cuda_common.hpp +++ b/components/common/cpp/include/ftl/cuda_common.hpp @@ -380,36 +380,17 @@ TextureObject<T>::~TextureObject() { free(); } -/*template <> -void TextureObject<uchar4>::upload(const cv::Mat &m, cudaStream_t stream); - -template <> -void TextureObject<float>::upload(const cv::Mat &m, cudaStream_t stream); - -template <> -void TextureObject<float2>::upload(const cv::Mat &m, cudaStream_t stream); - -template <> -void TextureObject<float4>::upload(const cv::Mat &m, cudaStream_t stream); - -template <> -void TextureObject<uchar>::upload(const cv::Mat &m, cudaStream_t stream); - - -template <> -void TextureObject<uchar4>::download(cv::Mat &m, cudaStream_t stream) const; - -template <> -void TextureObject<float>::download(cv::Mat &m, cudaStream_t stream) const; - -template <> -void TextureObject<float2>::download(cv::Mat &m, cudaStream_t stream) const; - -template <> -void TextureObject<float4>::download(cv::Mat &m, cudaStream_t stream) const; - -template <> -void TextureObject<uchar>::download(cv::Mat &m, cudaStream_t stream) const;*/ +/** + * Read a texture value using coordinates in the range of `b`, but from the + * texture `a` which may have a different resolution. + */ +template <typename A, typename B> +__device__ inline A getScaledTex2D(int x, int y, ftl::cuda::TextureObject<A> &a, ftl::cuda::TextureObject<B> &b) { + return a.tex2D( + (int)((float)x * ((float)a.width() / (float)b.width())), + (int)((float)y * ((float)a.height() / (float)b.height())) + ); +} } } diff --git a/components/operators/src/mask.cu b/components/operators/src/mask.cu index 128ae7d1778a0b79c3f743b79ae925e60a57d312..9f25b50ffc73c9e7881a1ce623a7e33d67b7d7f2 100644 --- a/components/operators/src/mask.cu +++ b/components/operators/src/mask.cu @@ -84,7 +84,8 @@ __global__ void discontinuity_kernel(ftl::cuda::TextureObject<uint8_t> mask_out, const float g = max(g1,max(g2,(max(g3,g4)))); // Calculate support window area - const uchar4 sup = support.tex2D((int)x, (int)y); + //const uchar4 sup = support.tex2D((int)x, (int)y); + const uchar4 sup = getScaledTex2D(x, y, support, depth); const float supx = min(sup.x,sup.y); const float supy = min(sup.z,sup.w); const float area = supx * supy; diff --git a/components/renderers/cpp/src/mask.cu b/components/renderers/cpp/src/mask.cu index e23012c7be7988bdf5d599c296db2df6ebce442d..7e36d9b32e5c024e2fc3094daf66d55248304b8d 100644 --- a/components/renderers/cpp/src/mask.cu +++ b/components/renderers/cpp/src/mask.cu @@ -3,6 +3,7 @@ using ftl::cuda::TextureObject; using ftl::cuda::Mask; +using ftl::cuda::getScaledTex2D; #define T_PER_BLOCK 16 @@ -15,10 +16,11 @@ __global__ void show_mask_kernel( const int y = blockIdx.y*blockDim.y + threadIdx.y; if (x >= 0 && x < colour.width() && y >=0 && y < colour.height()) { - float xscale = (float)x * ((float)mask.width() / (float)colour.width()); - float yscale = (float)y * ((float)mask.height() / (float)colour.height()); + //float xscale = (float)x * ((float)mask.width() / (float)colour.width()); + //float yscale = (float)y * ((float)mask.height() / (float)colour.height()); - Mask m(mask.tex2D((int)xscale,(int)yscale)); + //Mask m(mask.tex2D((int)xscale,(int)yscale)); + Mask m(getScaledTex2D(x, y, mask, colour)); if (m.is(id)) { colour(x,y) = style;