Add overlap operator

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<index_t, N> &windows, const std::array<index_t, N> &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:

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

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

include/matx/core/tensor.h

@@ -1071,9 +1071,10 @@ class tensor_t : public detail::tensor_impl_t<T,RANK,Desc> {
    * @returns Overlapping view of data
    *
    */
+  template <int N>  
   __MATX_INLINE__ auto
-  OverlapView(std::initializer_list<typename Desc::shape_type> const &windows,
-              std::initializer_list<typename Desc::stride_type> const &strides) const
+  OverlapView(const std::array<typename Desc::shape_type, N> &windows,
+              const std::array<typename Desc::stride_type, N> &strides) const
   {
     static_assert(RANK == 1, "Overlapped views only supported on 1D tensors.");
 

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"

include/matx/operators/overlap.h

@@ -0,0 +1,171 @@
+////////////////////////////////////////////////////////////////////////////////
+// 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 <int DIM, typename T>
+      class OverlapOp : public BaseOp<OverlapOp<DIM, T>>
+    {
+      public: 
+        using scalar_type = typename T::scalar_type;
+        using shape_type = index_t; 
+
+      private:
+        typename base_type<T>::type op_;
+        std::array<int32_t, DIM> dims_;
+        std::array<shape_type, DIM+1> n_;
+        std::array<shape_type, DIM+1> 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<shape_type, DIM> &windows,
+                                      const std::array<shape_type, DIM> &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;
+          adj_el -= adj_el % stride_size;
+
+          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<typename R>
+        __MATX_INLINE__ auto operator=(const R &rhs) {
+          if constexpr (is_matx_transform_op<R>()) {
+            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 <typename OpType, int N>
+  __MATX_INLINE__ auto overlap( const OpType &op, 
+      const std::array<index_t, N> &windows,
+      const std::array<index_t, N> &strides)
+  {
+    if constexpr (is_tensor_view_v<OpType>) {
+      return op.template OverlapView<N>(windows, strides);
+    } else {
+      return detail::OverlapOp<N, OpType>(op, windows, strides);
+    }
+  }
+
+  template <typename OpType, int N>
+  __MATX_INLINE__ auto overlap( const OpType &op, 
+      const index_t (&windows)[N],
+      const index_t (&strides)[N]) 
+  {
+    return overlap<OpType, N>(op, 
+        detail::to_array(windows), 
+        detail::to_array(strides));
+  }
+
+} // end namespace matx

test/00_operators/OperatorTests.cu

@@ -2378,6 +2378,87 @@ 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<TestType>::type;
+
+  ExecType exec{}; 
+
+
+  tensor_t<TestType, 1> a{{10}};
+  a.SetVals({0, 1, 2, 3, 4, 5, 6, 7, 8, 9});
+  auto ao = overlap(a, {4}, {2});
+
+  tensor_t<TestType, 2> 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<TestType, 2> 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<TestType, 2> 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<TestType, 2> 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<TestType, 2> 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  
+
+  ASSERT_EQ(b4out.Size(0), 4);
+  ASSERT_EQ(b4out.Size(1), 3);
+
+  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)
 {

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<TypeParam, 1> a{{10}};
-  a.SetVals({0, 1, 2, 3, 4, 5, 6, 7, 8, 9});
-  auto ao = a.OverlapView({4}, {2});
-
-  tensor_t<TypeParam, 2> 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<TypeParam, 2> 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<TypeParam, 2> 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<TypeParam, 2> 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();