From bb06ab78240bdcaafa9681a05b7b117a139cc8fb Mon Sep 17 00:00:00 2001 From: cliffburdick Date: Mon, 28 Aug 2023 15:20:51 -0700 Subject: [PATCH] Add overlap operator --- .../api/manipulation/rearranging/overlap.rst | 33 ++++ examples/spectrogram.cu | 2 +- examples/spectrogram_graph.cu | 2 +- include/matx/core/tensor.h | 5 +- include/matx/operators/operators.h | 1 + include/matx/operators/overlap.h | 173 ++++++++++++++++++ test/00_operators/OperatorTests.cu | 78 ++++++++ test/00_tensor/ViewTests.cu | 53 ------ 8 files changed, 290 insertions(+), 57 deletions(-) create mode 100644 docs_input/api/manipulation/rearranging/overlap.rst create mode 100644 include/matx/operators/overlap.h diff --git a/docs_input/api/manipulation/rearranging/overlap.rst b/docs_input/api/manipulation/rearranging/overlap.rst new file mode 100644 index 000000000..3ddabe3d1 --- /dev/null +++ b/docs_input/api/manipulation/rearranging/overlap.rst @@ -0,0 +1,33 @@ +.. _overlap_func: + +overlap +####### + +Create an overlapping view an of input operator giving a higher-rank view of the input + +For example, the following 1D tensor [1 2 3 4 5] could be cloned into a 2d tensor with a +window size of 2 and overlap of 1, resulting in: + [1 2 + 2 3 + 3 4 + 4 5] + +Currently this only works on 1D tensors going to 2D, but may be expanded +for higher dimensions in the future. Note that if the window size does not +divide evenly into the existing column dimension, the view may chop off the +end of the data to make the tensor rectangular. + +.. note:: + Only 1D input operators are accepted at this time + +.. doxygenfunction:: overlap( const OpType &op, const index_t (&windows)[N], const index_t (&strides)[N]) +.. doxygenfunction:: overlap( const OpType &op, const std::array &windows, const std::array &strides) + +Examples +~~~~~~~~ + +.. literalinclude:: ../../../../test/00_operators/OperatorTests.cu + :language: cpp + :start-after: example-begin overlap-test-1 + :end-before: example-end overlap-test-1 + :dedent: diff --git a/examples/spectrogram.cu b/examples/spectrogram.cu index 9e194b370..ad6b24c4f 100644 --- a/examples/spectrogram.cu +++ b/examples/spectrogram.cu @@ -115,7 +115,7 @@ int main([[maybe_unused]] int argc, [[maybe_unused]] char **argv) .run(stream); // Create overlapping matrix of segments. - auto stackedMatrix = x.OverlapView({nperseg}, {nstep}); + auto stackedMatrix = overlap(x, {nperseg}, {nstep}); // FFT along rows (fftStackedMatrix = fft(stackedMatrix)).run(stream); // Absolute value diff --git a/examples/spectrogram_graph.cu b/examples/spectrogram_graph.cu index 7a61d3431..fc1b7013a 100644 --- a/examples/spectrogram_graph.cu +++ b/examples/spectrogram_graph.cu @@ -117,7 +117,7 @@ int main([[maybe_unused]] int argc, [[maybe_unused]] char **argv) .run(stream); // Create overlapping matrix of segments. - auto stackedMatrix = x.OverlapView({nperseg}, {nstep}); + auto stackedMatrix = overlap(x, {nperseg}, {nstep}); // FFT along rows (fftStackedMatrix = fft(stackedMatrix)).run(stream); // Absolute value diff --git a/include/matx/core/tensor.h b/include/matx/core/tensor.h index 7d6658650..9f350aabd 100644 --- a/include/matx/core/tensor.h +++ b/include/matx/core/tensor.h @@ -1071,9 +1071,10 @@ class tensor_t : public detail::tensor_impl_t { * @returns Overlapping view of data * */ + template __MATX_INLINE__ auto - OverlapView(std::initializer_list const &windows, - std::initializer_list const &strides) const + OverlapView(const std::array &windows, + const std::array &strides) const { static_assert(RANK == 1, "Overlapped views only supported on 1D tensors."); diff --git a/include/matx/operators/operators.h b/include/matx/operators/operators.h index bbc735bed..ad4778997 100644 --- a/include/matx/operators/operators.h +++ b/include/matx/operators/operators.h @@ -74,6 +74,7 @@ #include "matx/operators/legendre.h" #include "matx/operators/lu.h" #include "matx/operators/matmul.h" +#include "matx/operators/overlap.h" #include "matx/operators/percentile.h" #include "matx/operators/permute.h" #include "matx/operators/planar.h" diff --git a/include/matx/operators/overlap.h b/include/matx/operators/overlap.h new file mode 100644 index 000000000..519fe46c4 --- /dev/null +++ b/include/matx/operators/overlap.h @@ -0,0 +1,173 @@ +//////////////////////////////////////////////////////////////////////////////// +// BSD 3-Clause License +// +// Copyright (c) 2021, NVIDIA Corporation +// All rights reserved. +// +// Redistribution and use in source and binary forms, with or without +// modification, are permitted provided that the following conditions are met: +// +// 1. Redistributions of source code must retain the above copyright notice, this +// list of conditions and the following disclaimer. +// +// 2. Redistributions in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// 3. Neither the name of the copyright holder nor the names of its +// contributors may be used to endorse or promote products derived from +// this software without specific prior written permission. +// +// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" +// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE +// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE +// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE +// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL +// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR +// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER +// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, +// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +///////////////////////////////////////////////////////////////////////////////// + +#pragma once + + +#include "matx/core/type_utils.h" +#include "matx/operators/base_operator.h" + +namespace matx +{ + /** + * Slices elements from an operator/tensor. + */ + namespace detail { + template + class OverlapOp : public BaseOp> + { + public: + using scalar_type = typename T::scalar_type; + using shape_type = index_t; + + private: + typename base_type::type op_; + std::array dims_; + std::array n_; + std::array s_; + + public: + using matxop = bool; + using matxoplvalue = bool; + + static_assert(DIM == 1, "overlap() only supports input rank 1 currently"); + + __MATX_INLINE__ std::string str() const { return "overlap(" + op_.str() + ")"; } + __MATX_INLINE__ OverlapOp(T op, const std::array &windows, + const std::array &strides) : op_(op) { + + // This only works for 1D tensors going to 2D at the moment. Generalize to + // higher dims later + index_t window_size = windows[0]; + index_t stride_size = strides[0]; + + MATX_ASSERT(stride_size < window_size, matxInvalidSize); + MATX_ASSERT(stride_size > 0, matxInvalidSize); + + // Figure out the actual length of the sequence we can use. It might be + // shorter than the original operator if the window/stride doesn't line up + // properly to make a rectangular matrix. + shape_type adj_el = op_.Size(0) - window_size; + while ((adj_el % stride_size) != 0) { + adj_el--; + } + + n_[1] = window_size; + s_[1] = 1; + n_[0] = adj_el / stride_size + 1; + s_[0] = stride_size; + }; + + __MATX_INLINE__ __MATX_DEVICE__ __MATX_HOST__ auto operator()(index_t i0, index_t i1) const + { + return op_(i0*s_[0] + i1); + } + + __MATX_INLINE__ __MATX_DEVICE__ __MATX_HOST__ auto& operator()(index_t i0, index_t i1) + { + return op_(i0*s_[0] + i1); + } + + static __MATX_INLINE__ constexpr __MATX_HOST__ __MATX_DEVICE__ int32_t Rank() + { + return DIM + 1; + } + constexpr __MATX_INLINE__ __MATX_HOST__ __MATX_DEVICE__ shape_type Size(int32_t dim) const + { + return n_[dim]; + } + + template + __MATX_INLINE__ auto operator=(const R &rhs) { + if constexpr (is_matx_transform_op()) { + return mtie(*this, rhs); + } + else { + return set(*this, rhs); + } + } + }; + } + + /** + * @brief Create an overlapping tensor view + * + * Creates and overlapping tensor view where an existing tensor can be + * repeated into a higher rank with overlapping elements. For example, the + * following 1D tensor [1 2 3 4 5] could be cloned into a 2d tensor with a + * window size of 2 and overlap of 1, resulting in: + * + * [1 2 + * 2 3 + * 3 4 + * 4 5] + * + * Currently this only works on 1D tensors going to 2D, but may be expanded + * for higher dimensions in the future. Note that if the window size does not + * divide evenly into the existing column dimension, the view may chop off the + * end of the data to make the tensor rectangular. + * + * @tparam OpType + * Type of operator input + * @tparam N + * Rank of overlapped window + * @param windows + * Window size (columns in output) + * @param strides + * Strides between data elements + * + * @returns Overlapping view of data + * + */ + template + __MATX_INLINE__ auto overlap( const OpType &op, + const std::array &windows, + const std::array &strides) + { + if constexpr (is_tensor_view_v) { + return op.template OverlapView(windows, strides); + } else { + return detail::OverlapOp(op, windows, strides); + } + } + + template + __MATX_INLINE__ auto overlap( const OpType &op, + const index_t (&windows)[N], + const index_t (&strides)[N]) + { + return overlap(op, + detail::to_array(windows), + detail::to_array(strides)); + } + +} // end namespace matx diff --git a/test/00_operators/OperatorTests.cu b/test/00_operators/OperatorTests.cu index 20f8f3316..ca078552a 100644 --- a/test/00_operators/OperatorTests.cu +++ b/test/00_operators/OperatorTests.cu @@ -2378,6 +2378,84 @@ TYPED_TEST(OperatorTestsNumericAllExecs, Reshape) MATX_EXIT_HANDLER(); } +TYPED_TEST(OperatorTestsNumericNonComplexAllExecs, Overlap) +{ + MATX_ENTER_HANDLER(); + + using TestType = std::tuple_element_t<0, TypeParam>; + using ExecType = std::tuple_element_t<1, TypeParam>; + using inner_type = typename inner_op_type_t::type; + + ExecType exec{}; + + + tensor_t a{{10}}; + a.SetVals({0, 1, 2, 3, 4, 5, 6, 7, 8, 9}); + auto ao = overlap(a, {4}, {2}); + + tensor_t b{{4, 4}}; + b.SetVals({{0, 1, 2, 3}, {2, 3, 4, 5}, {4, 5, 6, 7}, {6, 7, 8, 9}}); + for (index_t i = 0; i < b.Size(0); i++) { + for (index_t j = 0; j < b.Size(1); j++) { + ASSERT_EQ(ao(i, j), b(i, j)); + } + } + + auto ao2 = overlap(a, {4}, {1}); + + tensor_t b2{{7, 4}}; + b2.SetVals({{0, 1, 2, 3}, + {1, 2, 3, 4}, + {2, 3, 4, 5}, + {3, 4, 5, 6}, + {4, 5, 6, 7}, + {5, 6, 7, 8}, + {6, 7, 8, 9}}); + for (index_t i = 0; i < b2.Size(0); i++) { + for (index_t j = 0; j < b2.Size(1); j++) { + ASSERT_EQ(ao2(i, j), b2(i, j)); + } + } + + auto ao3 = overlap(a, {4}, {3}); + tensor_t b3{{3, 4}}; + b3.SetVals({{0, 1, 2, 3}, {3, 4, 5, 6}, {6, 7, 8, 9}}); + for (index_t i = 0; i < b3.Size(0); i++) { + for (index_t j = 0; j < b3.Size(1); j++) { + ASSERT_EQ(ao3(i, j), b3(i, j)); + } + } + + auto ao4 = overlap(a, {3}, {2}); + tensor_t b4{{4, 3}}; + b4.SetVals({{0, 1, 2}, {2, 3, 4}, {4, 5, 6}, {6, 7, 8}}); + for (index_t i = 0; i < b4.Size(0); i++) { + for (index_t j = 0; j < b4.Size(1); j++) { + ASSERT_EQ(ao4(i, j), b4(i, j)); + } + } + + // Test with an operator input + // example-begin overlap-test-1 + auto aop = linspace<0>(a.Shape(), (TestType)0, (TestType)9); + tensor_t b4out{{4, 3}}; + + // Input is {0, 1, 2, 3, 4, 5, 6, 7, 8, 9} + // Output is: {{0, 1, 2}, {2, 3, 4}, {4, 5, 6}, {6, 7, 8}} + (b4out = overlap(aop, {3}, {2})).run(exec); + // example-end overlap-test-1 + + cudaStreamSynchronize(0); + for (index_t i = 0; i < b4.Size(0); i++) { + for (index_t j = 0; j < b4.Size(1); j++) { + ASSERT_EQ(b4out(i, j), b4(i, j)); + } + } + + MATX_EXIT_HANDLER(); +} + + TYPED_TEST(OperatorTestsNumericAllExecs, Broadcast) { diff --git a/test/00_tensor/ViewTests.cu b/test/00_tensor/ViewTests.cu index 17d4c06da..a513642cd 100644 --- a/test/00_tensor/ViewTests.cu +++ b/test/00_tensor/ViewTests.cu @@ -124,59 +124,6 @@ TYPED_TEST_SUITE(ViewTestsIntegralAllExecs, MatXTypesIntegralAllExecs); TYPED_TEST_SUITE(ViewTestsBooleanAllExecs, MatXTypesBooleanAllExecs); -TYPED_TEST(ViewTestsNumericNonComplex, OverlapView) -{ - MATX_ENTER_HANDLER(); - - tensor_t a{{10}}; - a.SetVals({0, 1, 2, 3, 4, 5, 6, 7, 8, 9}); - auto ao = a.OverlapView({4}, {2}); - - tensor_t b{{4, 4}}; - b.SetVals({{0, 1, 2, 3}, {2, 3, 4, 5}, {4, 5, 6, 7}, {6, 7, 8, 9}}); - for (index_t i = 0; i < b.Size(0); i++) { - for (index_t j = 0; j < b.Size(1); j++) { - ASSERT_EQ(ao(i, j), b(i, j)); - } - } - - auto ao2 = a.OverlapView({4}, {1}); - - tensor_t b2{{7, 4}}; - b2.SetVals({{0, 1, 2, 3}, - {1, 2, 3, 4}, - {2, 3, 4, 5}, - {3, 4, 5, 6}, - {4, 5, 6, 7}, - {5, 6, 7, 8}, - {6, 7, 8, 9}}); - for (index_t i = 0; i < b2.Size(0); i++) { - for (index_t j = 0; j < b2.Size(1); j++) { - ASSERT_EQ(ao2(i, j), b2(i, j)); - } - } - - auto ao3 = a.OverlapView({4}, {3}); - tensor_t b3{{3, 4}}; - b3.SetVals({{0, 1, 2, 3}, {3, 4, 5, 6}, {6, 7, 8, 9}}); - for (index_t i = 0; i < b3.Size(0); i++) { - for (index_t j = 0; j < b3.Size(1); j++) { - ASSERT_EQ(ao3(i, j), b3(i, j)); - } - } - - auto ao4 = a.OverlapView({3}, {2}); - tensor_t b4{{4, 3}}; - b4.SetVals({{0, 1, 2}, {2, 3, 4}, {4, 5, 6}, {6, 7, 8}}); - for (index_t i = 0; i < b4.Size(0); i++) { - for (index_t j = 0; j < b4.Size(1); j++) { - ASSERT_EQ(ao4(i, j), b4(i, j)); - } - } - - MATX_EXIT_HANDLER(); -} - TYPED_TEST(ViewTestsAll, Stride) { MATX_ENTER_HANDLER();