diff --git a/bench/00_transform/einsum.cu b/bench/00_transform/einsum.cu index 4c7cf7a2..713c912a 100644 --- a/bench/00_transform/einsum.cu +++ b/bench/00_transform/einsum.cu @@ -1,6 +1,5 @@ #include #include "matx.h" -#include "matx_einsum.h" #if ENABLE_CUTENSOR diff --git a/examples/convolution.cu b/examples/convolution.cu index 6abe61d8..d458ad8d 100644 --- a/examples/convolution.cu +++ b/examples/convolution.cu @@ -30,9 +30,7 @@ // OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. ///////////////////////////////////////////////////////////////////////////////// -#include "matx_type_utils.h" -#include "matx_tensor.h" -#include "matx_conv.h" +#include "matx.h" #include #include diff --git a/examples/fft_conv.cu b/examples/fft_conv.cu index 210f6c81..c085e3b6 100644 --- a/examples/fft_conv.cu +++ b/examples/fft_conv.cu @@ -31,7 +31,6 @@ ///////////////////////////////////////////////////////////////////////////////// #include "matx.h" -#include "matx_filter.h" #include #include #include diff --git a/examples/recursive_filter.cu b/examples/recursive_filter.cu index 040e57b0..69fed5ee 100644 --- a/examples/recursive_filter.cu +++ b/examples/recursive_filter.cu @@ -31,7 +31,6 @@ ///////////////////////////////////////////////////////////////////////////////// #include "matx.h" -#include "matx_filter.h" #include #include #include diff --git a/examples/spectrogram.cu b/examples/spectrogram.cu index 35be79a0..23c30841 100644 --- a/examples/spectrogram.cu +++ b/examples/spectrogram.cu @@ -30,14 +30,13 @@ // OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. ///////////////////////////////////////////////////////////////////////////////// -#include "matx.h" -#include "matx_viz.h" -#include "matx_make.h" #include #include #include #include +#include "matx.h" + using namespace matx; #define FFT_TYPE CUFFT_C2C @@ -156,4 +155,4 @@ int main([[maybe_unused]] int argc, [[maybe_unused]] char **argv) cudaStreamDestroy(stream); CUDA_CHECK_LAST_ERROR(); MATX_EXIT_HANDLER(); -} \ No newline at end of file +} diff --git a/examples/spectrogram_graph.cu b/examples/spectrogram_graph.cu index 18bdf8a4..86ca7a44 100644 --- a/examples/spectrogram_graph.cu +++ b/examples/spectrogram_graph.cu @@ -31,7 +31,6 @@ ///////////////////////////////////////////////////////////////////////////////// #include "matx.h" -#include "matx_viz.h" #include #include #include @@ -165,4 +164,4 @@ int main([[maybe_unused]] int argc, [[maybe_unused]] char **argv) cudaStreamDestroy(stream); CUDA_CHECK_LAST_ERROR(); MATX_EXIT_HANDLER(); -} \ No newline at end of file +} diff --git a/include/matx.h b/include/matx.h index d5dcc1e1..9aa6263a 100644 --- a/include/matx.h +++ b/include/matx.h @@ -32,28 +32,19 @@ #pragma once #include -#include "matx_defines.h" -#include "matx_half_complex.h" -#include "matx_half.h" +#include "matx/core/defines.h" +#include "matx/core/error.h" +#include "matx/core/file_io.h" +#include "matx/core/half_complex.h" +#include "matx/core/half.h" +#include "matx/core/pybind.h" +#include "matx/core/tensor.h" +#include "matx/core/utils.h" -#include "matx_utils.h" -#include "matx_error.h" -#include "matx_tensor.h" -#include "matx_random.h" -#include "matx_tensor_generators.h" +#include "matx/executors/executors.h" +#include "matx/generators/generators.h" #include "matx/operators/operators.h" #include "matx/transforms/transforms.h" -#include "matx_exec_kernel.h" -#include "matx_fft.h" -#include "matx_conv.h" -#include "matx_corr.h" -#include "matx_matmul.h" -#include "matx_reduce.h" -#include "matx_inverse.h" -#include "matx_solver.h" -#include "matx_cov.h" -#include "matx_cub.h" - using fcomplex = cuda::std::complex; using dcomplex = cuda::std::complex; diff --git a/include/matx_allocator.h b/include/matx/core/allocator.h similarity index 99% rename from include/matx_allocator.h rename to include/matx/core/allocator.h index 1afa2cc4..fb693e8a 100644 --- a/include/matx_allocator.h +++ b/include/matx/core/allocator.h @@ -41,7 +41,7 @@ #include #endif -#include "matx_error.h" +#include "matx/core/error.h" #pragma once diff --git a/include/matx_cache.h b/include/matx/core/cache.h similarity index 99% rename from include/matx_cache.h rename to include/matx/core/cache.h index 30726f9d..10514ff9 100644 --- a/include/matx_cache.h +++ b/include/matx/core/cache.h @@ -33,11 +33,12 @@ #pragma once -#include "matx_error.h" #include #include #include +#include "matx/core/error.h" + namespace matx { namespace detail { diff --git a/include/matx_defines.h b/include/matx/core/defines.h similarity index 100% rename from include/matx_defines.h rename to include/matx/core/defines.h diff --git a/include/matx_error.h b/include/matx/core/error.h similarity index 99% rename from include/matx_error.h rename to include/matx/core/error.h index 4e5688f4..f875ad24 100644 --- a/include/matx_error.h +++ b/include/matx/core/error.h @@ -37,7 +37,7 @@ #include #include -#include "matx_stacktrace.h" +#include "matx/core/stacktrace.h" namespace matx { diff --git a/include/matx_file_io.h b/include/matx/core/file_io.h similarity index 98% rename from include/matx_file_io.h rename to include/matx/core/file_io.h index 459904d5..a86b8fff 100644 --- a/include/matx_file_io.h +++ b/include/matx/core/file_io.h @@ -30,8 +30,6 @@ // OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. ///////////////////////////////////////////////////////////////////////////////// -#include "matx_pybind.h" -#include "matx_tensor.h" #include #include #include @@ -41,7 +39,9 @@ // #include // #include -#include "matx_error.h" +#include "matx/core/pybind.h" +#include "matx/core/tensor.h" +#include "matx/core/error.h" #pragma once @@ -231,4 +231,4 @@ void WriteMAT(const TensorType &t, const std::string fname, } }; // namespace io -}; // namespace matx \ No newline at end of file +}; // namespace matx diff --git a/include/matx_get_grid_dims.h b/include/matx/core/get_grid_dims.h similarity index 100% rename from include/matx_get_grid_dims.h rename to include/matx/core/get_grid_dims.h diff --git a/include/matx_half.h b/include/matx/core/half.h similarity index 99% rename from include/matx_half.h rename to include/matx/core/half.h index 64341616..14f3a65b 100644 --- a/include/matx_half.h +++ b/include/matx/core/half.h @@ -32,12 +32,13 @@ #pragma once -#include "cuda_bf16.h" -#include "cuda_fp16.h" #include -#include "matx_defines.h" #include +#include "cuda_bf16.h" +#include "cuda_fp16.h" +#include "matx/core/defines.h" + namespace matx { /** @@ -1209,4 +1210,4 @@ __MATX_HOST__ __MATX_DEVICE__ __MATX_INLINE__ matxHalf tanh(const matxHalf using matxFp16 = matxHalf<__half>; ///< Alias for fp16 using matxBf16 = matxHalf<__nv_bfloat16>; ///< Alias for bf16 -}; // namespace matx \ No newline at end of file +}; // namespace matx diff --git a/include/matx_half_complex.h b/include/matx/core/half_complex.h similarity index 99% rename from include/matx_half_complex.h rename to include/matx/core/half_complex.h index f9b91209..710de3bc 100644 --- a/include/matx_half_complex.h +++ b/include/matx/core/half_complex.h @@ -33,13 +33,14 @@ #pragma once // Next line is a workaround for libcuda++ 11.6.0 on CTK 11.4.100+. Remove once libcuda++ is fixed #define _LIBCUDACXX_HAS_NO_INT128 -#include "cuComplex.h" -#include "matx_half.h" #include #include #include #include +#include "cuComplex.h" +#include "matx/core/half.h" + namespace matx { /** @@ -1043,4 +1044,4 @@ tanh(const matxHalfComplex &x) using matxFp16Complex = matxHalfComplex; ///< Alias for a MatX fp16 complex wrapper using matxBf16Complex = matxHalfComplex; ///< Alias for a MatXbf16 complex wrapper -}; // namespace matx \ No newline at end of file +}; // namespace matx diff --git a/include/matx_make.h b/include/matx/core/make_tensor.h similarity index 99% rename from include/matx_make.h rename to include/matx/core/make_tensor.h index 995dc314..d9446ac1 100644 --- a/include/matx_make.h +++ b/include/matx/core/make_tensor.h @@ -33,9 +33,9 @@ ///////////////////////////////////////////////////////////////////////////////// #pragma once -#include "matx_storage.h" -#include "matx_tensor_desc.h" -#include "matx_shape.h" +#include "matx/core/storage.h" +#include "matx/core/tensor_desc.h" +#include "matx/core/shape.h" namespace matx { @@ -353,4 +353,4 @@ auto make_static_tensor() { return tensor_t{std::move(s), std::move(desc)}; } -} // namespace matx \ No newline at end of file +} // namespace matx diff --git a/include/matx_pybind.h b/include/matx/core/pybind.h similarity index 99% rename from include/matx_pybind.h rename to include/matx/core/pybind.h index 7f1ee707..3ac12042 100644 --- a/include/matx_pybind.h +++ b/include/matx/core/pybind.h @@ -31,7 +31,7 @@ ///////////////////////////////////////////////////////////////////////////////// #pragma once -#include "matx.h" +#include "matx/core/type_utils.h" #include diff --git a/include/matx_shape.h b/include/matx/core/shape.h similarity index 98% rename from include/matx_shape.h rename to include/matx/core/shape.h index d599ad58..56f2c2c2 100644 --- a/include/matx_shape.h +++ b/include/matx/core/shape.h @@ -36,9 +36,9 @@ #include #include -#include "matx_allocator.h" -#include "matx_error.h" -#include "matx_type_utils.h" +#include "matx/core/allocator.h" +#include "matx/core/error.h" +#include "matx/core/type_utils.h" namespace matx { diff --git a/include/matx_stacktrace.h b/include/matx/core/stacktrace.h similarity index 100% rename from include/matx_stacktrace.h rename to include/matx/core/stacktrace.h diff --git a/include/matx_storage.h b/include/matx/core/storage.h similarity index 99% rename from include/matx_storage.h rename to include/matx/core/storage.h index e951fba1..0f2513f6 100644 --- a/include/matx_storage.h +++ b/include/matx/core/storage.h @@ -32,9 +32,9 @@ #pragma once -#include "matx_type_utils.h" -#include "matx_allocator.h" -#include "matx_error.h" +#include "matx/core/type_utils.h" +#include "matx/core/allocator.h" +#include "matx/core/error.h" namespace matx { @@ -538,4 +538,4 @@ namespace matx */ template using DefaultStorage = basic_storage>>; -}; \ No newline at end of file +}; diff --git a/include/matx_tensor.h b/include/matx/core/tensor.h similarity index 99% rename from include/matx_tensor.h rename to include/matx/core/tensor.h index 4d42a235..631d176b 100644 --- a/include/matx_tensor.h +++ b/include/matx/core/tensor.h @@ -40,12 +40,12 @@ #include #include -#include "matx_allocator.h" -#include "matx_error.h" -#include "matx_storage.h" -#include "matx_tensor_impl.h" -#include "matx_utility_kernels.cuh" -#include "matx_tensor_utils.h" +#include "matx/core/allocator.h" +#include "matx/core/error.h" +#include "matx/core/storage.h" +#include "matx/core/tensor_impl.h" +#include "matx/core/tensor_utils.h" +#include "matx/kernels/utility.cuh" static constexpr int MAX_TENSOR_DIM = 4; diff --git a/include/matx_tensor_desc.h b/include/matx/core/tensor_desc.h similarity index 99% rename from include/matx_tensor_desc.h rename to include/matx/core/tensor_desc.h index 7f10b654..86c6c44a 100644 --- a/include/matx_tensor_desc.h +++ b/include/matx/core/tensor_desc.h @@ -36,7 +36,7 @@ #include #include -#include "matx_error.h" +#include "matx/core/error.h" namespace matx { @@ -445,4 +445,4 @@ template using DefaultDescriptor = tensor_desc_cr_ds_64_64_t; -}; // namespace matx \ No newline at end of file +}; // namespace matx diff --git a/include/matx_tensor_impl.h b/include/matx/core/tensor_impl.h similarity index 98% rename from include/matx_tensor_impl.h rename to include/matx/core/tensor_impl.h index fbfcf7a0..a43fb3b0 100644 --- a/include/matx_tensor_impl.h +++ b/include/matx/core/tensor_impl.h @@ -34,15 +34,15 @@ #include #include -#include "matx_error.h" -#include "matx_set.h" -#include "matx_defines.h" -#include "matx_tensor_desc.h" -#include "matx_type_utils.h" -#include "matx_tensor_utils.h" -#include "matx_exec_kernel.h" -#include "matx_executor.h" -#include "matx_make.h" +#include "matx/core/error.h" +#include "matx/core/defines.h" +#include "matx/core/tensor_desc.h" +#include "matx/core/type_utils.h" +#include "matx/core/tensor_utils.h" +#include "matx/operators/set.h" +//#include "matx_exec_kernel.h" +//#include "matx_executor.h" +#include "matx/core/make_tensor.h" namespace matx { diff --git a/include/matx_tensor_utils.h b/include/matx/core/tensor_utils.h similarity index 99% rename from include/matx_tensor_utils.h rename to include/matx/core/tensor_utils.h index 30649a77..4c19057d 100644 --- a/include/matx_tensor_utils.h +++ b/include/matx/core/tensor_utils.h @@ -34,7 +34,8 @@ #include #include -#include "matx_make.h" + +#include "matx/core/make_tensor.h" namespace matx { @@ -618,4 +619,4 @@ void Print(const Op &op, Args... dims) { // }, tup); // } -} \ No newline at end of file +} diff --git a/include/matx_type_utils.h b/include/matx/core/type_utils.h similarity index 99% rename from include/matx_type_utils.h rename to include/matx/core/type_utils.h index 697a6d3f..98fbe446 100644 --- a/include/matx_type_utils.h +++ b/include/matx/core/type_utils.h @@ -32,9 +32,6 @@ #pragma once -#include "cuda_fp16.h" -#include "matx_half.h" -#include "matx_half_complex.h" #include #include #include @@ -43,6 +40,10 @@ #include #include +#include "cuda_fp16.h" +#include "matx/core/half.h" +#include "matx/core/half_complex.h" + /** * Defines type traits for host and device compilers. This file should be includable by * the host compiler, so no device code should be present diff --git a/include/matx_utils.h b/include/matx/core/utils.h similarity index 97% rename from include/matx_utils.h rename to include/matx/core/utils.h index 3d9f795b..d433f730 100644 --- a/include/matx_utils.h +++ b/include/matx/core/utils.h @@ -31,8 +31,8 @@ ///////////////////////////////////////////////////////////////////////////////// #pragma once -#include "matx_defines.h" -#include "matx_error.h" +#include "matx/core/defines.h" +#include "matx/core/error.h" #define AMPERE_CC 8 #define VOLTA_CC 7 diff --git a/include/matx_viz.h b/include/matx/core/viz.h similarity index 99% rename from include/matx_viz.h rename to include/matx/core/viz.h index 39af6de2..adb1395f 100644 --- a/include/matx_viz.h +++ b/include/matx/core/viz.h @@ -32,12 +32,13 @@ #pragma once -#include "matx_error.h" -#include "matx_tensor.h" -#include "matx_pybind.h" #include #include +#include "matx/core/error.h" +#include "matx/core/tensor.h" +#include "matx/core/pybind.h" + #if MATX_ENABLE_VIZ namespace matx { @@ -378,4 +379,4 @@ void surf( }; // viz }; // matx -#endif \ No newline at end of file +#endif diff --git a/include/matx_executor.h b/include/matx/executors/device.h similarity index 96% rename from include/matx_executor.h rename to include/matx/executors/device.h index 9e9ee143..7fd84b61 100644 --- a/include/matx_executor.h +++ b/include/matx/executors/device.h @@ -32,9 +32,9 @@ #pragma once -#include "matx_defines.h" -#include "matx_exec_host.h" -#include "matx_exec_kernel.h" +#include "matx/core/defines.h" +#include "matx/executors/host.h" +#include "matx/executors/kernel.h" namespace matx { @@ -63,4 +63,4 @@ namespace detail { ex.Exec(op); } } -}; \ No newline at end of file +}; diff --git a/include/matx/executors/executors.h b/include/matx/executors/executors.h new file mode 100644 index 00000000..1592058b --- /dev/null +++ b/include/matx/executors/executors.h @@ -0,0 +1,36 @@ +//////////////////////////////////////////////////////////////////////////////// +// 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/executors/device.h" +#include "matx/executors/host.h" diff --git a/include/matx_exec_host.h b/include/matx/executors/host.h similarity index 98% rename from include/matx_exec_host.h rename to include/matx/executors/host.h index 2aa05956..82df2220 100644 --- a/include/matx_exec_host.h +++ b/include/matx/executors/host.h @@ -33,8 +33,8 @@ #pragma once #include -#include "matx_error.h" -#include "matx_get_grid_dims.h" +#include "matx/core/error.h" +#include "matx/core/get_grid_dims.h" namespace matx { @@ -105,4 +105,4 @@ class SingleThreadHostExecutor { } }; -} \ No newline at end of file +} diff --git a/include/matx_exec_kernel.h b/include/matx/executors/kernel.h similarity index 99% rename from include/matx_exec_kernel.h rename to include/matx/executors/kernel.h index 89bbc5f0..e94e24ff 100644 --- a/include/matx_exec_kernel.h +++ b/include/matx/executors/kernel.h @@ -33,8 +33,8 @@ #pragma once #include -#include "matx_error.h" -#include "matx_get_grid_dims.h" +#include "matx/core/error.h" +#include "matx/core/get_grid_dims.h" namespace matx { namespace detail { diff --git a/include/matx/generators/alternate.h b/include/matx/generators/alternate.h new file mode 100644 index 00000000..4fe7e3cd --- /dev/null +++ b/include/matx/generators/alternate.h @@ -0,0 +1,100 @@ +//////////////////////////////////////////////////////////////////////////////// +// 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/generators/generator1d.h" + +namespace matx +{ + namespace detail { + template class Alternating { + private: + index_t size_; + + public: + using scalar_type = T; + + inline __MATX_HOST__ __MATX_DEVICE__ Alternating(index_t size) : size_(size) {}; + inline __MATX_HOST__ __MATX_DEVICE__ T operator()(index_t i) const + { + return (-2 * (i & 1)) + 1; + } + }; + } + + + /** + * Creates an alternating +1/-1 sequence + * + * @tparam T + * Data type + * @tparam Dim + * Dimension to create window over + * @tparam RANK + * The RANK of the shape, can be deduced from shape + * + * @param s + * The shape of the tensor + * + * Returns values for alternating sequence + */ + template ::type>, bool> = true> + inline auto alternate(ShapeType &&s) + { + constexpr int RANK = std::tuple_size>::value; + static_assert(RANK > Dim); + detail::Alternating h( *(s.begin() + Dim)); + return detail::matxGenerator1D_t, Dim, ShapeType>(std::forward(s), h); + } + + /** + * Creates an alternating +1/-1 sequence + * + * @tparam T + * Data type + * @tparam Dim + * Dimension to create window over + * @tparam RANK + * The RANK of the shape, can be deduced from shape + * + * @param s + * C array representing shape of the tensor + * + */ + template + inline auto alternate(const index_t (&s)[RANK]) + { + return alternate(detail::to_array(s)); + } +} // end namespace matx diff --git a/include/matx/generators/bartlett.h b/include/matx/generators/bartlett.h new file mode 100644 index 00000000..80c67d3d --- /dev/null +++ b/include/matx/generators/bartlett.h @@ -0,0 +1,105 @@ +//////////////////////////////////////////////////////////////////////////////// +// 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/generators/generator1d.h" + +namespace matx +{ + namespace detail { + template class Bartlett { + private: + index_t size_; + + public: + using scalar_type = T; + inline __MATX_HOST__ __MATX_DEVICE__ Bartlett(index_t size) : size_(size){}; + + inline __MATX_HOST__ __MATX_DEVICE__ T operator()(index_t i) const + { + return (T(2) / (T(size_) - 1)) * + (((T(size_) - 1) / T(2)) - + cuda::std::abs(T(i) - ((T(size_) - 1) / T(2)))); + } + }; + } + + /** + * Creates a Bartlett window operator of shape s with the + * window applies along the x, y, z, or w dimension + * + * @tparam T + * Data type + * @tparam Dim + * Dimension to create window over + * @tparam RANK + * The RANK of the shape + * + * @param s + * The shape of the tensor + * + * Returns values for a Bartlett window across the selected dimension. + */ + template ::type>, bool> = true> + inline auto bartlett(ShapeType &&s) + { + constexpr int RANK = std::tuple_size>::value; + static_assert(RANK > Dim); + detail::Bartlett h( *(s.begin() + Dim)); + return detail::matxGenerator1D_t, Dim, ShapeType>(std::forward(s), h); + } + + /** + * Creates a Bartlett window operator of shape s with the + * window applies along the x, y, z, or w dimension + * + * @tparam T + * Data type + * @tparam Dim + * Dimension to create window over + * @tparam RANK + * The RANK of the shape + * + * @param s + * The C array shape of the tensor + * + * Returns values for a Bartlett window across the selected dimension. + */ + template + inline auto bartlett(const index_t (&s)[RANK]) + { + return bartlett(detail::to_array(s)); + } + +} // end namespace matx diff --git a/include/matx/generators/blackman.h b/include/matx/generators/blackman.h new file mode 100644 index 00000000..0a5ba47e --- /dev/null +++ b/include/matx/generators/blackman.h @@ -0,0 +1,106 @@ +//////////////////////////////////////////////////////////////////////////////// +// 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/generators/generator1d.h" + +namespace matx +{ + namespace detail { + template class Blackman { + private: + index_t size_; + + public: + using scalar_type = T; + inline __MATX_HOST__ __MATX_DEVICE__ Blackman(index_t size) : size_(size){}; + + inline __MATX_HOST__ __MATX_DEVICE__ T operator()(index_t i) const + { + return T(0.42) + + ((T)0.5 * + cuda::std::cos(T(M_PI) * (1 - size_ + 2 * T(i)) / T(size_ - 1))) + + ((T)0.08 * cuda::std::cos(T(2 * M_PI) * (1 - size_ + 2 * T(i)) / + T(size_ - 1))); + } + }; + } + + /** + * Creates a Blackman window operator of shape s with the + * window applies along the x, y, z, or w dimension + * + * @tparam T + * Data type + * @tparam Dim + * Dimension to create window over + * @tparam RANK + * The RANK of the shape + * + * @param s + * The shape of the tensor + * + * Returns values for a Blackman window across the selected dimension. + */ + template ::type>, bool> = true> + inline auto blackman(ShapeType &&s) + { + constexpr int RANK = std::tuple_size>::value; + static_assert(RANK > Dim); + detail::Blackman h( *(s.begin() + Dim)); + return detail::matxGenerator1D_t, Dim, ShapeType>(std::forward(s), h); + } + + /** + * Creates a Blackman window operator of shape s with the + * window applies along the x, y, z, or w dimension + * + * @tparam T + * Data type + * @tparam Dim + * Dimension to create window over + * @tparam RANK + * The RANK of the shape + * + * @param s + * The shape of the tensor + * + * Returns values for a Blackman window across the selected dimension. + */ + template + inline auto blackman(const index_t (&s)[RANK]) + { + return blackman(detail::to_array(s)); + } +} // end namespace matx diff --git a/include/matx/generators/chirp.h b/include/matx/generators/chirp.h new file mode 100644 index 00000000..5b239556 --- /dev/null +++ b/include/matx/generators/chirp.h @@ -0,0 +1,238 @@ +//////////////////////////////////////////////////////////////////////////////// +// 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/generators/linspace.h" + +namespace matx +{ + enum class ChirpMethod { + CHIRP_METHOD_LINEAR + }; + + enum class ChirpType { + CHIRP_TYPE_REAL, + CHIRP_TYPE_COMPLEX + }; + + namespace detail { + template + class Chirp { + using space_type = typename SpaceOp::scalar_type; + + + private: + SpaceOp sop_; + FreqType f0_; + FreqType f1_; + space_type t1_; + ChirpMethod method_; + + public: + using scalar_type = FreqType; + using matxop = bool; + inline __MATX_HOST__ __MATX_DEVICE__ Chirp(SpaceOp sop, FreqType f0, space_type t1, FreqType f1, ChirpMethod method) : + sop_(sop), + f0_(f0), + t1_(t1), + f1_(f1), + method_(method) + {} + + inline __MATX_HOST__ __MATX_DEVICE__ auto operator()(index_t i) const + { + if (method_ == ChirpMethod::CHIRP_METHOD_LINEAR) { + return cuda::std::cos(2.0f * M_PI * (f0_ * sop_(i) + 0.5f * ((f1_ - f0_) / t1_) * sop_(i) * sop_(i))); + } + + return 0.0; + } + + constexpr inline __MATX_HOST__ __MATX_DEVICE__ index_t Size([[maybe_unused]] int dim) const + { + return sop_.Size(0); + } + static inline constexpr __MATX_HOST__ __MATX_DEVICE__ int32_t Rank() { return 1; } + }; + + template + class ComplexChirp { + using space_type = typename SpaceOp::scalar_type; + + + private: + SpaceOp sop_; + FreqType f0_; + FreqType f1_; + space_type t1_; + ChirpMethod method_; + + public: + using scalar_type = cuda::std::complex; + using matxop = bool; + inline __MATX_HOST__ __MATX_DEVICE__ ComplexChirp(SpaceOp sop, FreqType f0, space_type t1, FreqType f1, ChirpMethod method) : + sop_(sop), + f0_(f0), + t1_(t1), + f1_(f1), + method_(method) + {} + + inline __MATX_HOST__ __MATX_DEVICE__ auto operator()(index_t i) const + { + if (method_ == ChirpMethod::CHIRP_METHOD_LINEAR) { + FreqType real = cuda::std::cos(2.0f * M_PI * (f0_ * sop_(i) + 0.5f * ((f1_ - f0_) / t1_) * sop_(i) * sop_(i))); + FreqType imag = -cuda::std::cos(2.0f * M_PI * (f0_ * sop_(i) + 0.5f * ((f1_ - f0_) / t1_) * sop_(i) * sop_(i) + 90.0/360.0)); + return cuda::std::complex{real, imag}; + } + + return cuda::std::complex{0, 0}; + } + + constexpr inline __MATX_HOST__ __MATX_DEVICE__ index_t Size([[maybe_unused]] int dim) const + { + return sop_.Size(0); + } + static inline constexpr __MATX_HOST__ __MATX_DEVICE__ int32_t Rank() { return 1; } + }; + } + + /** + * Creates a real chirp signal (swept-frequency cosine) + * + * SpaceOp provides the time vector with custom spacing. + * + * @tparam FreqType + * Frequency data type + * @tparam SpaceOp + * Operator type of spacer + * + * @param t + * Vector representing values in time + * @param f0 + * Instantenous frequency at time 0 + * @param t1 + * Time for f1 + * @param f1 + * Frequency (Hz) at time t1 + * @param method + * Chirp method (CHIRP_METHOD_LINEAR) + * + * @returns The chirp operator + */ + template + inline auto chirp(SpaceOp t, FreqType f0, typename SpaceOp::scalar_type t1, FreqType f1, ChirpMethod method = ChirpMethod::CHIRP_METHOD_LINEAR) + { + MATX_ASSERT_STR(method == ChirpMethod::CHIRP_METHOD_LINEAR, matxInvalidType, "Only linear chirps are supported") + + return detail::Chirp(t, f0, t1, f1, method); + } + + /** + * Creates a complex chirp signal (swept-frequency cosine) + * + * SpaceOp provides the time vector with custom spacing. + * + * @tparam FreqType + * Frequency data type + * @tparam SpaceOp + * Operator type of spacer + * + * @param t + * Vector representing values in time + * @param f0 + * Instantenous frequency at time 0 + * @param t1 + * Time for f1 + * @param f1 + * Frequency (Hz) at time t1 + * @param method + * Chirp method (CHIRP_METHOD_LINEAR) + * + * @returns The chirp operator + */ + template + inline auto cchirp(SpaceOp t, FreqType f0, typename SpaceOp::scalar_type t1, FreqType f1, ChirpMethod method = ChirpMethod::CHIRP_METHOD_LINEAR) + { + MATX_ASSERT_STR(method == ChirpMethod::CHIRP_METHOD_LINEAR, matxInvalidType, "Only linear chirps are supported") + + return detail::ComplexChirp(t, f0, t1, f1, method); + } + + /** + * Creates a chirp signal (swept-frequency cosine) + * + * Creates a linearly-spaced sequence from 0 to "last" with "num" elements in between. Each step is + * of size 1/num. + * + * @tparam FreqType + * Frequency data type + * @tparam TimeType + * Type of time vector + * @tparam Method + * Chirp method (CHIRP_METHOD_LINEAR) + * + * @param num + * Number of time samples + * @param last + * Last time sample value + * @param f0 + * Instantenous frequency at time 0 + * @param t1 + * Time for f1 + * @param f1 + * Frequency (Hz) at time t1 + * @param method + * Method to use to generate the chirp + * + * @returns The chirp operator + */ + template + inline auto chirp(index_t num, TimeType last, FreqType f0, TimeType t1, FreqType f1, ChirpMethod method = ChirpMethod::CHIRP_METHOD_LINEAR) + { + std::array shape = {num}; + auto space = linspace<0>(std::move(shape), (TimeType)0, last); + return chirp(space, f0, t1, f1, method); + } + + template + inline auto cchirp(index_t num, TimeType last, FreqType f0, TimeType t1, FreqType f1, ChirpMethod method = ChirpMethod::CHIRP_METHOD_LINEAR) + { + std::array shape = {num}; + auto space = linspace<0>(std::move(shape), (TimeType)0, last); + return cchirp(space, f0, t1, f1, method); + } + + + +} // end namespace matx diff --git a/include/matx/generators/diag.h b/include/matx/generators/diag.h new file mode 100644 index 00000000..6a464409 --- /dev/null +++ b/include/matx/generators/diag.h @@ -0,0 +1,142 @@ +//////////////////////////////////////////////////////////////////////////////// +// 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 + + +namespace matx +{ + namespace detail { + template class Diag { + static constexpr int RANK = std::tuple_size>::value; + + private: + ShapeType s_; + T val_; + + public: + // dummy type to signal this is a matxop + using matxop = bool; + using scalar_type = T; + + Diag(ShapeType &&s, T val) : s_(std::forward(s)), val_(val) + { + static_assert(Rank() > 1, "Diagonal generator must be used with an operator of rank 1 or higher"); + }; + + template + __MATX_INLINE__ __MATX_DEVICE__ __MATX_HOST__ auto operator()(Is... indices) const { + if (((pp_get<0>(indices...) == indices) && ...)) { + return T(val_); + } + else { + return T(0.0f); + } + } + + constexpr inline __MATX_HOST__ __MATX_DEVICE__ auto Size(int dim) const + { + return *(s_.begin() + dim); + } + static inline constexpr __MATX_HOST__ __MATX_DEVICE__ int32_t Rank() { return RANK; } + }; + } + + /** + * Creates a diagonal tensor with a given value on the diagonals + * + * diag returns a given value on all elements on the diagonals of a tensor, and + * 0 otherwise. In other words, if the index of every dimension is the same, the + * value is returned, otherwise a zero is returned. + * + * @tparam T + * Data type + * + */ + template ::type>, bool> = true> + inline auto diag(ShapeType &&s, T val) + { + return detail::Diag(std::forward(s), val); + } + + /** + * Creates a diagonal tensor with a given value on the diagonals + * + * diag returns a given value on all elements on the diagonals of a tensor, and + * 0 otherwise. In other words, if the index of every dimension is the same, the + * value is returned, otherwise a zero is returned. + * + * @tparam T + * Data type + * + */ + template + inline auto diag(const index_t (&s)[RANK], T val) + { + return diag(detail::to_array(s), val); + } + + /** + * Creates an identity patterns on the tensor + * + * eye returns 1 on all elements on the diagonals of a tensor, and 0 otherwise. + * In other words, if the index of every dimension is the same, a 1 is returned, + * otherwise a zero is returned. + * + * @tparam T + * Data type + * + */ + template ::type>, bool> = true> + inline auto eye(ShapeType &&s) + { + return detail::Diag(std::forward(s), T(1)); + } + + /** + * Creates an identity patterns on the tensor + * + * eye returns 1 on all elements on the diagonals of a tensor, and 0 otherwise. + * In other words, if the index of every dimension is the same, a 1 is returned, + * otherwise a zero is returned. + * + * @tparam T + * Data type + * + */ + template inline auto eye(const index_t (&s)[RANK]) + { + return eye(detail::to_array(s)); + } +} // end namespace matx diff --git a/include/matx/generators/flattop.h b/include/matx/generators/flattop.h new file mode 100644 index 00000000..91759fa8 --- /dev/null +++ b/include/matx/generators/flattop.h @@ -0,0 +1,115 @@ +//////////////////////////////////////////////////////////////////////////////// +// 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/generators/generator1d.h" + +namespace matx +{ + namespace detail { + template class FlatTop { + private: + index_t size_; + + static constexpr T a0 = 0.21557895; + static constexpr T a1 = 0.41663158; + static constexpr T a2 = 0.277263158; + static constexpr T a3 = 0.083578947; + static constexpr T a4 = 0.006947368; + + public: + + using scalar_type = T; + inline __MATX_HOST__ __MATX_DEVICE__ FlatTop(index_t size) : size_(size){}; + + inline __MATX_HOST__ __MATX_DEVICE__ T operator()(index_t i) const + { + return a0 + - a1 * cuda::std::cos(2*M_PI*i / (size_ - 1)) + + a2 * cuda::std::cos(4*M_PI*i / (size_ - 1)) + - a3 * cuda::std::cos(6*M_PI*i / (size_ - 1)) + + a4 * cuda::std::cos(8*M_PI*i / (size_ - 1)); + } + }; + } + + /** + * Creates a Flattop window operator of shape s with the + * window applies along the x, y, z, or w dimension + * + * @tparam T + * Data type + * @tparam Dim + * Dimension to create window over + * @tparam RANK + * The RANK of the shape + * + * @param s + * The shape of the tensor + * + * Returns values for a Flattop window across the selected dimension. + */ + template ::type>, bool> = true> + inline auto flattop(ShapeType &&s) + { + constexpr int RANK = std::tuple_size>::value; + static_assert(RANK > Dim); + detail::FlatTop h( *(s.begin() + Dim)); + return detail::matxGenerator1D_t, Dim, ShapeType>(std::forward(s), h); + } + + /** + * Creates a Flattop window operator of shape s with the + * window applies along the x, y, z, or w dimension + * + * @tparam T + * Data type + * @tparam Dim + * Dimension to create window over + * @tparam RANK + * The RANK of the shape + * + * @param s + * The shape of the tensor + * + * Returns values for a Flattop window across the selected dimension. + */ + template + inline auto flattop(const index_t (&s)[RANK]) + { + return flattop(detail::to_array(s)); + } + + +} // end namespace matx diff --git a/include/matx/generators/generator1d.h b/include/matx/generators/generator1d.h new file mode 100644 index 00000000..911ef030 --- /dev/null +++ b/include/matx/generators/generator1d.h @@ -0,0 +1,65 @@ +//////////////////////////////////////////////////////////////////////////////// +// 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 + + +namespace matx +{ + namespace detail { + template class matxGenerator1D_t { + public: + // dummy type to signal this is a matxop + using matxop = bool; + using scalar_type = typename Generator1D::scalar_type; + static constexpr int RANK = std::tuple_size>::value; + + template + matxGenerator1D_t(S &&s, Generator1D f) : f_(f), s_(std::forward(s)) {} + + template + __MATX_INLINE__ __MATX_DEVICE__ __MATX_HOST__ auto operator()(Is... indices) const { + return f_(pp_get(indices...)); + } + + constexpr inline __MATX_HOST__ __MATX_DEVICE__ auto Size(int dim) const + { + return *(s_.begin() + dim); + } + static inline constexpr __MATX_HOST__ __MATX_DEVICE__ int32_t Rank() { return RANK; } + + private: + Generator1D f_; + ShapeType s_; + }; + } +} // end namespace matx diff --git a/include/matx/generators/generators.h b/include/matx/generators/generators.h new file mode 100644 index 00000000..3c212468 --- /dev/null +++ b/include/matx/generators/generators.h @@ -0,0 +1,49 @@ +//////////////////////////////////////////////////////////////////////////////// +// 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/generators/alternate.h" +#include "matx/generators/bartlett.h" +#include "matx/generators/blackman.h" +#include "matx/generators/chirp.h" +#include "matx/generators/diag.h" +#include "matx/generators/flattop.h" +#include "matx/generators/hamming.h" +#include "matx/generators/hanning.h" +#include "matx/generators/linspace.h" +#include "matx/generators/logspace.h" +#include "matx/generators/meshgrid.h" +#include "matx/generators/ones.h" +#include "matx/generators/random.h" +#include "matx/generators/range.h" +#include "matx/generators/zeros.h" diff --git a/include/matx/generators/hamming.h b/include/matx/generators/hamming.h new file mode 100644 index 00000000..bca1887e --- /dev/null +++ b/include/matx/generators/hamming.h @@ -0,0 +1,104 @@ +//////////////////////////////////////////////////////////////////////////////// +// 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/generators/generator1d.h" + +namespace matx +{ + namespace detail { + template class Hamming { + private: + index_t size_; + + public: + using scalar_type = T; + + inline __MATX_HOST__ __MATX_DEVICE__ Hamming(index_t size) : size_(size){}; + + inline __MATX_HOST__ __MATX_DEVICE__ T operator()(index_t i) const + { + return T(.54) - T(.46) * cuda::std::cos(T(2 * M_PI) * T(i) / T(size_ - 1)); + } + }; + } + + + /** + * Creates a Hamming window operator of shape s with the + * window applies along the specified dimension + * + * @tparam T + * Data type + * @tparam Dim + * Dimension to create window over + * @tparam RANK + * The RANK of the shape, can be deduced from shape + * + * @param s + * The shape of the tensor + * + * Returns values for a Hamming window across the selected dimension. + */ + template ::type>, bool> = true> + inline auto hamming(ShapeType &&s) + { + constexpr int RANK = std::tuple_size>::value; + static_assert(RANK > Dim); + detail::Hamming h( *(s.begin() + Dim)); + return detail::matxGenerator1D_t, Dim, ShapeType>(std::forward(s), h); + } + + /** + * Creates a Hamming window operator of C-array shape s with the + * window applies along the specified dimension + * + * @tparam T + * Data type + * @tparam Dim + * Dimension to create window over + * @tparam RANK + * The RANK of the shape, can be deduced from shape + * + * @param s + * C array representing shape of the tensor + * + * Returns values for a Hamming window across the selected dimension. + */ + template + inline auto hamming(const index_t (&s)[RANK]) + { + return hamming(detail::to_array(s)); + } +} // end namespace matx diff --git a/include/matx/generators/hanning.h b/include/matx/generators/hanning.h new file mode 100644 index 00000000..ecb4ba25 --- /dev/null +++ b/include/matx/generators/hanning.h @@ -0,0 +1,102 @@ +//////////////////////////////////////////////////////////////////////////////// +// 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/generators/generator1d.h" + +namespace matx +{ + namespace detail { + template class Hanning { + private: + index_t size_; + + public: + using scalar_type = T; + inline __MATX_HOST__ __MATX_DEVICE__ Hanning(index_t size) : size_(size){}; + + inline __MATX_HOST__ __MATX_DEVICE__ T operator()(index_t i) const + { + return T(0.5) * (1 - cuda::std::cos(T(2 * M_PI) * T(i) / T(size_ - 1))); + } + }; + } + + /** + * Creates a Hanning window operator of shape s with the + * window applies along the x, y, z, or w dimension + * + * @tparam T + * Data type + * @tparam Dim + * Dimension to create window over + * @tparam RANK + * The RANK of the shape + * + * @param s + * The shape of the tensor + * + * Returns values for a Hanning window across the selected dimension. + */ + template ::type>, bool> = true> + inline auto hanning(ShapeType &&s) + { + constexpr int RANK = std::tuple_size>::value; + static_assert(RANK > Dim); + detail::Hanning h( *(s.begin() + Dim)); + return detail::matxGenerator1D_t, Dim, ShapeType>(std::forward(s), h); + } + + /** + * Creates a Hanning window operator of shape s with the + * window applies along the x, y, z, or w dimension + * + * @tparam T + * Data type + * @tparam Dim + * Dimension to create window over + * @tparam RANK + * The RANK of the shape + * + * @param s + * The C-array shape of the tensor + * + * Returns values for a Hanning window across the selected dimension. + */ + template + inline auto hanning(const index_t (&s)[RANK]) + { + return hanning(detail::to_array(s)); + } +} // end namespace matx diff --git a/include/matx/generators/linspace.h b/include/matx/generators/linspace.h new file mode 100644 index 00000000..9c2f7b3e --- /dev/null +++ b/include/matx/generators/linspace.h @@ -0,0 +1,115 @@ +//////////////////////////////////////////////////////////////////////////////// +// 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/generators/range.h" + +namespace matx +{ + namespace detail { + template class Linspace { + private: + Range range_; + + public: + using scalar_type = T; + + inline Linspace(T first, T last, index_t count) + { +#ifdef __CUDA_ARCH__ + range_ = Range{first, (last - first) / static_cast(count - 1)}; +#else + // Host has no support for most half precision operators/intrinsics + if constexpr (is_matx_half_v) { + range_ = Range{static_cast(first), + (static_cast(last) - static_cast(first)) / + static_cast(count - 1)}; + } + else { + range_ = Range{first, (last - first) / static_cast(count - 1)}; + } +#endif + } + + __MATX_DEVICE__ __MATX_HOST__ __MATX_INLINE__ T operator()(index_t idx) const { return range_(idx); } + }; + } + + + /** + * @brief Create a linearly-spaced range of values + * + * Creates a set of values using a start and end that are linearly- + * spaced apart over the set of values. Distance is determined + * by the shape and selected dimension. + * + * @tparam T Operator type + * @tparam Dim Dimension to operate over + * @tparam ShapeType Shape type + * @param s Shape object + * @param first First value + * @param last Last value + * @return Operator with linearly-spaced values + */ + template ::type>, bool> = true> + inline auto linspace(ShapeType &&s, T first, T last) + { + constexpr int RANK = std::tuple_size>::value; + static_assert(RANK > Dim); + auto count = *(s.begin() + Dim); + detail::Linspace l(first, last, count); + return detail::matxGenerator1D_t, Dim, ShapeType>(std::forward(s), l); + } + + /** + * @brief Create a linearly-spaced range of values + * + * Creates a set of values using a start and end that are linearly- + * spaced apart over the set of values. Distance is determined + * by the shape and selected dimension. + * + * @tparam T Operator type + * @tparam Dim Dimension to operate over + * @tparam ShapeType Shape type + * @param s Shape object + * @param first First value + * @param last Last value + * @return Operator with linearly-spaced values + */ + template + inline auto linspace(const index_t (&s)[RANK], T first, T last) + { + return linspace(detail::to_array(s), first, last); + } +} // end namespace matx diff --git a/include/matx/generators/logspace.h b/include/matx/generators/logspace.h new file mode 100644 index 00000000..ac099d39 --- /dev/null +++ b/include/matx/generators/logspace.h @@ -0,0 +1,139 @@ +//////////////////////////////////////////////////////////////////////////////// +// 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/generators/generator1d.h" + +namespace matx +{ + namespace detail { + template class Logspace { + private: + Range range_; + + public: + using scalar_type = T; + + inline Logspace(T first, T last, index_t count) + { +#ifdef __CUDA_ARCH__ + if constexpr (is_matx_half_v) { + range_ = Range{first, (last - first) / static_cast(count - 1.0f)}; + } + else { + range_ = Range{first, (last - first) / static_cast(count - 1)}; + } +#else + // Host has no support for most half precision operators/intrinsics + if constexpr (is_matx_half_v) { + range_ = Range{static_cast(first), + (static_cast(last) - static_cast(first)) / + static_cast(count - 1)}; + } + else { + range_ = Range{first, (last - first) / static_cast(count - 1)}; + } +#endif + } + + __MATX_DEVICE__ __MATX_HOST__ __MATX_INLINE__ T operator()(index_t idx) const + { + if constexpr (is_matx_half_v) { + return static_cast( + cuda::std::pow(10, static_cast(range_(idx)))); + } + else { + return cuda::std::pow(10, range_(idx)); + } + + // WAR for compiler bug. + if constexpr (!is_matx_half_v) { + return cuda::std::pow(10, range_(idx)); + } + else { + return static_cast( + cuda::std::pow(10, static_cast(range_(idx)))); + } + } + }; + } + + + /** + * @brief Create a log10-spaced range of values + * + * Creates a set of values using a start and end that are log10- + * spaced apart over the set of values. Distance is determined + * by the shape and selected dimension. + * + * @tparam T Operator type + * @tparam Dim Dimension to operate over + * @tparam ShapeType Shape type + * @param s Shape object + * @param first First value + * @param last Last value + * @return Operator with log10-spaced values + */ + template ::type>, bool> = true> + inline auto logspace(ShapeType &&s, T first, T last) + { + constexpr int RANK = std::tuple_size>::value; + static_assert(RANK > Dim); + auto count = *(s.begin() + Dim); + detail::Logspace l(first, last, count); + return detail::matxGenerator1D_t, Dim, ShapeType>(std::forward(s), l); + } + + /** + * @brief Create a log10-spaced range of values + * + * Creates a set of values using a start and end that are log10- + * spaced apart over the set of values. Distance is determined + * by the shape and selected dimension. + * + * @tparam T Operator type + * @tparam Dim Dimension to operate over + * @tparam ShapeType Shape type + * @param s Shape object + * @param first First value + * @param last Last value + * @return Operator with log10-spaced values + */ + template + inline auto logspace(const index_t (&s)[RANK], T first, T last) + { + return logspace(detail::to_array(s), first, last); + } + +} // end namespace matx diff --git a/include/matx/generators/meshgrid.h b/include/matx/generators/meshgrid.h new file mode 100644 index 00000000..2728eba3 --- /dev/null +++ b/include/matx/generators/meshgrid.h @@ -0,0 +1,115 @@ +//////////////////////////////////////////////////////////////////////////////// +// 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 + + +namespace matx +{ + namespace detail { + template class Meshgrid_X { + private: + std::array x_; + std::array y_; + + public: + // dummy type to signal this is a matxop + using matxop = bool; + using scalar_type = T; + + Meshgrid_X(std::array x, std::array y) : x_(x), y_(y) {} + + inline __MATX_DEVICE__ T operator()(index_t i, index_t j) const + { + return x_[0] + j * (x_[1] - x_[0]) / (x_[2] - 1); + } + + constexpr inline __MATX_HOST__ __MATX_DEVICE__ auto Size(int dim) const + { + return (dim == 0) ? y_[2] : x_[2]; + } + static inline constexpr __MATX_HOST__ __MATX_DEVICE__ int32_t Rank() { return 2; } + }; + + template class Meshgrid_Y { + private: + std::array x_; + std::array y_; + + public: + // dummy type to signal this is a matxop + using matxop = bool; + using scalar_type = T; + + Meshgrid_Y(std::array x, std::array y) : x_(x), y_(y) {} + + inline __MATX_DEVICE__ T operator()(index_t i, index_t j) const + { + return y_[0] + i * (y_[1] - y_[0]) / (y_[2] - 1); + }; + + constexpr inline __MATX_HOST__ __MATX_DEVICE__ index_t Size(uint32_t dim) const + { + return (dim == 0) ? y_[2] : x_[2]; + } + static inline constexpr __MATX_HOST__ __MATX_DEVICE__ int32_t Rank() { return 2; } + }; + } + + /** + * Creates an mesh grid X matrix + * + * + * @tparam T + * Data type + * + */ + template + inline auto meshgrid_x(const std::array &x, const std::array &y) + { + return detail::Meshgrid_X(x, y); + } + + /** + * Creates an mesh grid Y matrix + * + * + * @tparam T + * Data type + * + */ + template + inline auto meshgrid_y(const std::array &x, const std::array &y) + { + return detail::Meshgrid_Y(x, y); + } +} // end namespace matx diff --git a/include/matx/generators/ones.h b/include/matx/generators/ones.h new file mode 100644 index 00000000..431735b6 --- /dev/null +++ b/include/matx/generators/ones.h @@ -0,0 +1,76 @@ +//////////////////////////////////////////////////////////////////////////////// +// 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/operators/constval.h" + +namespace matx +{ + /** + * Return one for all elements + * + * Ones is used as an operator that always returns a 1 type for all + * elements. It can be used in place of memset to set all values to 1. + * + * @tparam T + * Data type + * + * @param s + * Shape of tensor + */ + template ::type>, bool> = true> + inline auto ones(ShapeType &&s) + { + return detail::ConstVal(std::forward(s), T(1)); + } + + /** + * Return one for all elements + * + * Ones is used as an operator that always returns a 1 type for all + * elements. It can be used in place of memset to set all values to 1. + * + * @tparam T + * Data type + * + * @param s + * Shape of tensor + */ + template + inline auto ones(const index_t (&s)[RANK]) + { + return ones(detail::to_array(s)); + } + +} // end namespace matx diff --git a/include/matx_random.h b/include/matx/generators/random.h similarity index 99% rename from include/matx_random.h rename to include/matx/generators/random.h index 211979b0..0bfb9dc6 100644 --- a/include/matx_random.h +++ b/include/matx/generators/random.h @@ -32,8 +32,8 @@ #pragma once -#include "matx_error.h" -#include "matx_shape.h" +#include "matx/core/error.h" +#include "matx/core/shape.h" #include #include #include diff --git a/include/matx/generators/range.h b/include/matx/generators/range.h new file mode 100644 index 00000000..97dd8065 --- /dev/null +++ b/include/matx/generators/range.h @@ -0,0 +1,115 @@ +//////////////////////////////////////////////////////////////////////////////// +// 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/generators/generator1d.h" + +namespace matx +{ + namespace detail { + template class Range { + private: + T first_; + T step_; + + public: + using scalar_type = T; + + Range() = default; + + Range(T first, T step) : first_(first), step_(step) {} + + __MATX_DEVICE__ __MATX_HOST__ __MATX_INLINE__ T operator()(index_t idx) const + { + if constexpr (is_matx_half_v) { + return first_ + T(static_cast((float)idx) * step_); + } + else { + return first_ + T(static_cast(idx) * step_); + } + + if constexpr (!is_matx_half_v) { + return first_ + T(static_cast(idx) * step_); + } + else { + return first_ + T(static_cast((float)idx) * step_); + } + } + }; + } + + /** + * Create a range of values along the x dimension + * + * Creates a range of values of type T with a start and step size. + * Value is determined by the index in operator() + * + * @param s + * Tensor shape + * @param first + * Starting value + * @param step + * Step size + * + */ + template ::type>, bool> = true> + inline auto range(ShapeType &&s, T first, T step) + { + constexpr int RANK = std::tuple_size>::value; + static_assert(RANK > Dim); + detail::Range r(first, step); + return detail::matxGenerator1D_t, Dim, ShapeType>(std::forward(s), r); + } + + /** + * Create a range of values along the x dimension + * + * Creates a range of values of type T with a start and step size. + * Value is determined by the index in operator() + * + * @param s + * Tensor shape + * @param first + * Starting value + * @param step + * Step size + * + */ + template + inline auto range(const index_t (&s)[RANK], T first, T step) + { + return range(detail::to_array(s), first, step); + } + +} // end namespace matx diff --git a/include/matx/generators/zeros.h b/include/matx/generators/zeros.h new file mode 100644 index 00000000..59d46b86 --- /dev/null +++ b/include/matx/generators/zeros.h @@ -0,0 +1,74 @@ +//////////////////////////////////////////////////////////////////////////////// +// 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 + + +namespace matx +{ + /** + * Return zero for all elements + * + * Zeros is used as an operator that always returns a 0 type for all + * elements. It can be used in place of memset to zero a block of memory. + * + * @tparam T + * Data type + * + * @param s + * Shape of tensor + */ + template ::type>, bool> = true> + inline auto zeros(ShapeType &&s) + { + return detail::ConstVal(std::forward(s), T(0)); + } + + /** + * Return zero for all elements + * + * Zeros is used as an operator that always returns a 0 type for all + * elements. It can be used in place of memset to zero a block of memory. + * + * @tparam T + * Data type + * + * @param s + * Shape of tensor + */ + template + inline auto zeros(const index_t (&s)[RANK]) + { + return zeros(detail::to_array(s)); + } +} // end namespace matx diff --git a/include/kernels/matx_conv_kernels.cuh b/include/matx/kernels/conv.cuh old mode 100755 new mode 100644 similarity index 99% rename from include/kernels/matx_conv_kernels.cuh rename to include/matx/kernels/conv.cuh index ff38f313..8b30209b --- a/include/kernels/matx_conv_kernels.cuh +++ b/include/matx/kernels/conv.cuh @@ -1,9 +1,6 @@ #pragma once -#include "cuComplex.h" -#include "matx_type_utils.h" -#include "matx_tensor_utils.h" #include #include #include @@ -11,6 +8,10 @@ #include #include +#include "cuComplex.h" +#include "matx/core/type_utils.h" +#include "matx/core/tensor_utils.h" + #define BLOCK_SIZE_NON_RECURSIVE 1024 namespace matx { diff --git a/include/kernels/matx_filter_kernels.cuh b/include/matx/kernels/filter.cuh old mode 100755 new mode 100644 similarity index 100% rename from include/kernels/matx_filter_kernels.cuh rename to include/matx/kernels/filter.cuh diff --git a/include/kernels/matx_transpose.cuh b/include/matx/kernels/transpose.cuh similarity index 98% rename from include/kernels/matx_transpose.cuh rename to include/matx/kernels/transpose.cuh index 5bfd7e7c..7a37c460 100644 --- a/include/kernels/matx_transpose.cuh +++ b/include/matx/kernels/transpose.cuh @@ -1,8 +1,6 @@ #pragma once -#include "cuComplex.h" -#include "matx_type_utils.h" #include #include #include @@ -10,6 +8,9 @@ #include #include +#include "cuComplex.h" +#include "matx/core/type_utils.h" + namespace matx { #ifdef __CUDACC__ diff --git a/include/kernels/matx_utility_kernels.cuh b/include/matx/kernels/utility.cuh similarity index 82% rename from include/kernels/matx_utility_kernels.cuh rename to include/matx/kernels/utility.cuh index bba3b016..b0776d21 100644 --- a/include/kernels/matx_utility_kernels.cuh +++ b/include/matx/kernels/utility.cuh @@ -1,7 +1,5 @@ #pragma once -#include "matx_tensor.h" -#include "matx_type_utils.h" #include #include #include @@ -9,6 +7,9 @@ #include #include +#include "matx/core/tensor.h" +#include "matx/core/type_utils.h" + namespace matx { diff --git a/include/matx/operators/base_operator.h b/include/matx/operators/base_operator.h index b7c0e9a9..a82907e5 100644 --- a/include/matx/operators/base_operator.h +++ b/include/matx/operators/base_operator.h @@ -33,7 +33,7 @@ #pragma once -#include "matx_type_utils.h" +#include "matx/core/type_utils.h" namespace matx { diff --git a/include/matx/operators/binary_operators.h b/include/matx/operators/binary_operators.h index 94bac9ab..ae8c9ac4 100644 --- a/include/matx/operators/binary_operators.h +++ b/include/matx/operators/binary_operators.h @@ -33,7 +33,7 @@ #pragma once -#include "matx_type_utils.h" +#include "matx/core/type_utils.h" #include "matx/operators/base_operator.h" #include "matx/operators/scalar_ops.h" diff --git a/include/matx/operators/cast.h b/include/matx/operators/cast.h index a575790c..39e53b45 100644 --- a/include/matx/operators/cast.h +++ b/include/matx/operators/cast.h @@ -33,7 +33,7 @@ #pragma once -#include "matx_type_utils.h" +#include "matx/core/type_utils.h" #include "matx/operators/base_operator.h" namespace matx diff --git a/include/matx/operators/clone.h b/include/matx/operators/clone.h index 2fc59105..82538cbf 100644 --- a/include/matx/operators/clone.h +++ b/include/matx/operators/clone.h @@ -33,7 +33,7 @@ #pragma once -#include "matx_type_utils.h" +#include "matx/core/type_utils.h" #include "matx/operators/base_operator.h" namespace matx diff --git a/include/matx/operators/collapse.h b/include/matx/operators/collapse.h index 4663635a..b5834482 100644 --- a/include/matx/operators/collapse.h +++ b/include/matx/operators/collapse.h @@ -33,7 +33,7 @@ #pragma once -#include "matx_type_utils.h" +#include "matx/core/type_utils.h" #include "matx/operators/base_operator.h" namespace matx diff --git a/include/matx/operators/comma.h b/include/matx/operators/comma.h index f08a17c2..919570bc 100644 --- a/include/matx/operators/comma.h +++ b/include/matx/operators/comma.h @@ -33,7 +33,7 @@ #pragma once -#include "matx_type_utils.h" +#include "matx/core/type_utils.h" #include "matx/operators/base_operator.h" namespace matx diff --git a/include/matx/operators/concat.h b/include/matx/operators/concat.h index dd2a6484..e6573113 100644 --- a/include/matx/operators/concat.h +++ b/include/matx/operators/concat.h @@ -33,7 +33,7 @@ #pragma once -#include "matx_type_utils.h" +#include "matx/core/type_utils.h" #include "matx/operators/base_operator.h" namespace matx diff --git a/include/matx/operators/constval.h b/include/matx/operators/constval.h new file mode 100644 index 00000000..5001eeda --- /dev/null +++ b/include/matx/operators/constval.h @@ -0,0 +1,64 @@ +//////////////////////////////////////////////////////////////////////////////// +// 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 + + +namespace matx +{ + namespace detail { + template class ConstVal { + static constexpr int RANK = std::tuple_size::type>::value; + + private: + ShapeType s_; + T v_; + + public: + // dummy type to signal this is a matxop + using matxop = bool; + using scalar_type = T; + + ConstVal(ShapeType &&s, T val) : s_(std::forward(s)), v_(val){}; + + template + __MATX_INLINE__ __MATX_DEVICE__ __MATX_HOST__ auto operator()(Is...) const { + return v_; }; + + constexpr inline __MATX_HOST__ __MATX_DEVICE__ auto Size(int dim) const + { + return *(s_.begin() + dim); + } + static inline constexpr __MATX_HOST__ __MATX_DEVICE__ int32_t Rank() { return RANK; } + }; + } +} // end namespace matx diff --git a/include/matx_signal.h b/include/matx/operators/dct.h similarity index 94% rename from include/matx_signal.h rename to include/matx/operators/dct.h index 3308cf20..f736fa93 100644 --- a/include/matx_signal.h +++ b/include/matx/operators/dct.h @@ -35,14 +35,12 @@ #include #include -#include "matx_allocator.h" -#include "matx_error.h" -#include "matx_shape.h" -#include "matx_tensor.h" -#include "matx_type_utils.h" +#include "matx/core/allocator.h" +#include "matx/core/error.h" +#include "matx/core/tensor.h" +#include "matx/core/type_utils.h" namespace matx { -namespace signal { /* Operator for perfoming the 2*exp(-j*pi*k/(2N)) part of the DCT */ namespace detail { @@ -109,5 +107,4 @@ void dct(OutputTensor &out, const InputTensor &in, detail::dctOp(out, s, N).run(stream); } -}; // namespace signal -}; // namespace matx \ No newline at end of file +}; // namespace matx diff --git a/include/matx/operators/diag.h b/include/matx/operators/diag.h index 2549c5e7..fca440c8 100644 --- a/include/matx/operators/diag.h +++ b/include/matx/operators/diag.h @@ -33,7 +33,7 @@ #pragma once -#include "matx_type_utils.h" +#include "matx/core/type_utils.h" #include "matx/operators/base_operator.h" namespace matx diff --git a/include/matx/operators/fftshift.h b/include/matx/operators/fftshift.h index 4ea87f96..be065284 100644 --- a/include/matx/operators/fftshift.h +++ b/include/matx/operators/fftshift.h @@ -33,7 +33,7 @@ #pragma once -#include "matx_type_utils.h" +#include "matx/core/type_utils.h" #include "matx/operators/base_operator.h" namespace matx diff --git a/include/matx/operators/flatten.h b/include/matx/operators/flatten.h index b0934e32..b9856be0 100644 --- a/include/matx/operators/flatten.h +++ b/include/matx/operators/flatten.h @@ -33,7 +33,7 @@ #pragma once -#include "matx_type_utils.h" +#include "matx/core/type_utils.h" #include "matx/operators/base_operator.h" namespace matx diff --git a/include/matx/operators/hermitian.h b/include/matx/operators/hermitian.h index f3f2e25c..5dc94e8f 100644 --- a/include/matx/operators/hermitian.h +++ b/include/matx/operators/hermitian.h @@ -33,7 +33,7 @@ #pragma once -#include "matx_type_utils.h" +#include "matx/core/type_utils.h" #include "matx/operators/base_operator.h" namespace matx diff --git a/include/matx/operators/if.h b/include/matx/operators/if.h index 72faddb4..66aafba7 100644 --- a/include/matx/operators/if.h +++ b/include/matx/operators/if.h @@ -33,7 +33,7 @@ #pragma once -#include "matx_type_utils.h" +#include "matx/core/type_utils.h" #include "matx/operators/base_operator.h" namespace matx diff --git a/include/matx/operators/ifelse.h b/include/matx/operators/ifelse.h index 5d72786a..9b34d732 100644 --- a/include/matx/operators/ifelse.h +++ b/include/matx/operators/ifelse.h @@ -33,7 +33,7 @@ #pragma once -#include "matx_type_utils.h" +#include "matx/core/type_utils.h" #include "matx/operators/base_operator.h" namespace matx diff --git a/include/matx/operators/interleaved.h b/include/matx/operators/interleaved.h index be8c8b5e..0d6d000f 100644 --- a/include/matx/operators/interleaved.h +++ b/include/matx/operators/interleaved.h @@ -33,7 +33,7 @@ #pragma once -#include "matx_type_utils.h" +#include "matx/core/type_utils.h" #include "matx/operators/base_operator.h" namespace matx diff --git a/include/matx/operators/kronecker.h b/include/matx/operators/kronecker.h index d7fcdaf0..31b41a97 100644 --- a/include/matx/operators/kronecker.h +++ b/include/matx/operators/kronecker.h @@ -33,7 +33,7 @@ #pragma once -#include "matx_type_utils.h" +#include "matx/core/type_utils.h" #include "matx/operators/base_operator.h" namespace matx diff --git a/include/matx/operators/operators.h b/include/matx/operators/operators.h index 1c33f10b..28a99c71 100644 --- a/include/matx/operators/operators.h +++ b/include/matx/operators/operators.h @@ -37,10 +37,12 @@ #include "matx/operators/collapse.h" #include "matx/operators/concat.h" +#include "matx/operators/constval.h" #include "matx/operators/cast.h" #include "matx/operators/clone.h" #include "matx/operators/comma.h" #include "matx/operators/diag.h" +#include "matx/operators/dct.h" #include "matx/operators/fftshift.h" #include "matx/operators/flatten.h" #include "matx/operators/hermitian.h" @@ -56,5 +58,6 @@ #include "matx/operators/reverse.h" #include "matx/operators/select.h" #include "matx/operators/self.h" +#include "matx/operators/set.h" #include "matx/operators/shift.h" #include "matx/operators/slice.h" diff --git a/include/matx/operators/permute.h b/include/matx/operators/permute.h index c5cb7f76..560212b2 100644 --- a/include/matx/operators/permute.h +++ b/include/matx/operators/permute.h @@ -33,7 +33,7 @@ #pragma once -#include "matx_type_utils.h" +#include "matx/core/type_utils.h" #include "matx/operators/base_operator.h" namespace matx diff --git a/include/matx/operators/planar.h b/include/matx/operators/planar.h index bd283ffd..92c6b2a7 100644 --- a/include/matx/operators/planar.h +++ b/include/matx/operators/planar.h @@ -33,7 +33,7 @@ #pragma once -#include "matx_type_utils.h" +#include "matx/core/type_utils.h" #include "matx/operators/base_operator.h" namespace matx diff --git a/include/matx/operators/r2c.h b/include/matx/operators/r2c.h index 9d937bad..c6017f9b 100644 --- a/include/matx/operators/r2c.h +++ b/include/matx/operators/r2c.h @@ -33,7 +33,7 @@ #pragma once -#include "matx_type_utils.h" +#include "matx/core/type_utils.h" #include "matx/operators/base_operator.h" namespace matx diff --git a/include/matx/operators/remap.h b/include/matx/operators/remap.h index 92748396..095f0d54 100644 --- a/include/matx/operators/remap.h +++ b/include/matx/operators/remap.h @@ -33,7 +33,7 @@ #pragma once -#include "matx_type_utils.h" +#include "matx/core/type_utils.h" #include "matx/operators/base_operator.h" namespace matx diff --git a/include/matx/operators/repmat.h b/include/matx/operators/repmat.h index 4e95fd07..c6210482 100644 --- a/include/matx/operators/repmat.h +++ b/include/matx/operators/repmat.h @@ -33,7 +33,7 @@ #pragma once -#include "matx_type_utils.h" +#include "matx/core/type_utils.h" #include "matx/operators/base_operator.h" namespace matx diff --git a/include/matx/operators/reverse.h b/include/matx/operators/reverse.h index 63063eb8..501b1df6 100644 --- a/include/matx/operators/reverse.h +++ b/include/matx/operators/reverse.h @@ -33,7 +33,7 @@ #pragma once -#include "matx_type_utils.h" +#include "matx/core/type_utils.h" #include "matx/operators/base_operator.h" namespace matx diff --git a/include/matx/operators/select.h b/include/matx/operators/select.h index d2313a79..d8b83c54 100644 --- a/include/matx/operators/select.h +++ b/include/matx/operators/select.h @@ -33,7 +33,7 @@ #pragma once -#include "matx_type_utils.h" +#include "matx/core/type_utils.h" #include "matx/operators/base_operator.h" namespace matx diff --git a/include/matx/operators/self.h b/include/matx/operators/self.h index 7556833c..3fd1bc57 100644 --- a/include/matx/operators/self.h +++ b/include/matx/operators/self.h @@ -33,7 +33,7 @@ #pragma once -#include "matx_type_utils.h" +#include "matx/core/type_utils.h" #include "matx/operators/base_operator.h" namespace matx diff --git a/include/matx_set.h b/include/matx/operators/set.h similarity index 98% rename from include/matx_set.h rename to include/matx/operators/set.h index f213e45e..61b84ee7 100644 --- a/include/matx_set.h +++ b/include/matx/operators/set.h @@ -34,9 +34,9 @@ #include -#include "matx_error.h" -#include "matx_type_utils.h" -#include "matx_tensor_utils.h" +#include "matx/core/error.h" +#include "matx/core/type_utils.h" +#include "matx/core/tensor_utils.h" namespace matx { template class tensor_t; ///< Tensor detail type diff --git a/include/matx/operators/shift.h b/include/matx/operators/shift.h index c051b8ec..f5adb39d 100644 --- a/include/matx/operators/shift.h +++ b/include/matx/operators/shift.h @@ -33,7 +33,7 @@ #pragma once -#include "matx_type_utils.h" +#include "matx/core/type_utils.h" #include "matx/operators/base_operator.h" namespace matx diff --git a/include/matx/operators/slice.h b/include/matx/operators/slice.h index 817d99b5..c22c5804 100644 --- a/include/matx/operators/slice.h +++ b/include/matx/operators/slice.h @@ -33,7 +33,7 @@ #pragma once -#include "matx_type_utils.h" +#include "matx/core/type_utils.h" #include "matx/operators/base_operator.h" namespace matx diff --git a/include/matx/operators/unary_operators.h b/include/matx/operators/unary_operators.h index 332ce92b..1a23f681 100644 --- a/include/matx/operators/unary_operators.h +++ b/include/matx/operators/unary_operators.h @@ -33,7 +33,7 @@ #pragma once -#include "matx_type_utils.h" +#include "matx/core/type_utils.h" #include "matx/operators/scalar_ops.h" #include "matx/operators/base_operator.h" diff --git a/include/matx_radar.h b/include/matx/transforms/ambgfun.h similarity index 96% rename from include/matx_radar.h rename to include/matx/transforms/ambgfun.h index 57b4cdca..1ce5b428 100644 --- a/include/matx_radar.h +++ b/include/matx/transforms/ambgfun.h @@ -35,21 +35,20 @@ #include #include -#include "matx_allocator.h" -#include "matx_error.h" -#include "matx_shape.h" -#include "matx_tensor.h" -#include "matx_type_utils.h" +#include "matx/core/allocator.h" +#include "matx/core/error.h" +#include "matx/core/tensor.h" +#include "matx/core/type_utils.h" + +#include "matx/transforms/copy.h" namespace matx { -namespace signal{ typedef enum { AMGBFUN_CUT_TYPE_2D, AMGBFUN_CUT_TYPE_DELAY, AMGBFUN_CUT_TYPE_DOPPLER, } AMBGFunCutType_t; }; -}; namespace matx { namespace detail { @@ -160,7 +159,7 @@ template class AmbgDoppX : public BaseOp> { template void InternalAmbgFun(AMFTensor &amf, XTensor &x, std::optional &y, - [[maybe_unused]] double fs, ::matx::signal::AMBGFunCutType_t cut, + [[maybe_unused]] double fs, ::matx::AMBGFunCutType_t cut, [[maybe_unused]] float cut_val, cudaStream_t stream = 0) { constexpr int RANK = XTensor::Rank(); @@ -207,7 +206,7 @@ void InternalAmbgFun(AMFTensor &amf, XTensor &x, powf(2.0, static_cast(std::ceil(std::log2(len_seq - 1))))); index_t xlen = x_normdiv_v.Size(RANK - 1); - if (cut == ::matx::signal::AMGBFUN_CUT_TYPE_2D) { + if (cut == ::matx::AMGBFUN_CUT_TYPE_2D) { T1 *new_ynorm; matxAlloc(reinterpret_cast(&new_ynorm), sizeof(T1) * (len_seq - 1) * xlen, MATX_ASYNC_DEVICE_MEMORY, @@ -240,7 +239,7 @@ void InternalAmbgFun(AMFTensor &amf, XTensor &x, (amf_tmp_v = (float)nfreq * abs(fftshift1D(fullfft))).run(stream); matx::copy(amf, amf_tmp_v.RealView(), stream); } - else if (cut == ::matx::signal::AMGBFUN_CUT_TYPE_DELAY) { + else if (cut == ::matx::AMGBFUN_CUT_TYPE_DELAY) { T1 *fft_data_x, *fft_data_y, *amf_tmp; matxAlloc(reinterpret_cast(&fft_data_x), sizeof(*fft_data_x) * nfreq, MATX_ASYNC_DEVICE_MEMORY, stream); @@ -276,7 +275,7 @@ void InternalAmbgFun(AMFTensor &amf, XTensor &x, auto amfv = make_tensor(amf_tmp_v.GetStorage(), amfv_size); matx::copy(amf, amfv.RealView(), stream); } - else if (cut == ::matx::signal::AMGBFUN_CUT_TYPE_DOPPLER) { + else if (cut == ::matx::AMGBFUN_CUT_TYPE_DOPPLER) { T1 *fft_data_x, *fft_data_y, *amf_tmp; matxAlloc(reinterpret_cast(&fft_data_x), sizeof(*fft_data_x) * (len_seq - 1), MATX_ASYNC_DEVICE_MEMORY, @@ -321,7 +320,6 @@ void InternalAmbgFun(AMFTensor &amf, XTensor &x, } namespace matx { -namespace signal { /** * Cross-ambiguity function @@ -408,5 +406,4 @@ inline void ambgfun(AMFTensor &amf, XTensor &x, } -}; // namespace signal }; // namespace matx diff --git a/include/matx_conv.h b/include/matx/transforms/conv.h similarity index 99% rename from include/matx_conv.h rename to include/matx/transforms/conv.h index c9b5810c..d4d004a2 100644 --- a/include/matx_conv.h +++ b/include/matx/transforms/conv.h @@ -37,11 +37,10 @@ #include #include -#include "kernels/matx_conv_kernels.cuh" -#include "matx_error.h" -#include "matx_tensor.h" - +#include "matx/core/error.h" +#include "matx/core/tensor.h" #include "matx/operators/clone.h" +#include "matx/kernels/conv.cuh" namespace matx { namespace detail { diff --git a/include/matx/transforms/copy.h b/include/matx/transforms/copy.h index 537ff8f2..f2af2b3d 100644 --- a/include/matx/transforms/copy.h +++ b/include/matx/transforms/copy.h @@ -33,7 +33,7 @@ #pragma once -#include "matx_type_utils.h" +#include "matx/core/type_utils.h" namespace matx { diff --git a/include/matx_corr.h b/include/matx/transforms/corr.h similarity index 96% rename from include/matx_corr.h rename to include/matx/transforms/corr.h index 54c4ef3c..48c0e99b 100644 --- a/include/matx_corr.h +++ b/include/matx/transforms/corr.h @@ -37,9 +37,9 @@ #include #include -#include "matx_conv.h" -#include "matx_error.h" -#include "matx_tensor.h" +#include "matx/transforms/conv.h" +#include "matx/core/error.h" +#include "matx/core/tensor.h" namespace matx { diff --git a/include/matx_cov.h b/include/matx/transforms/cov.h similarity index 98% rename from include/matx_cov.h rename to include/matx/transforms/cov.h index 3f6d4ee7..866496aa 100644 --- a/include/matx_cov.h +++ b/include/matx/transforms/cov.h @@ -33,13 +33,13 @@ #pragma once -#include "matx_error.h" -#include "matx_matmul.h" -#include "matx_tensor.h" - #include #include +#include "matx/core/error.h" +#include "matx/core/tensor.h" +#include "matx/transforms/matmul.h" + namespace matx { namespace detail { /** diff --git a/include/matx_cub.h b/include/matx/transforms/cub.h similarity index 99% rename from include/matx_cub.h rename to include/matx/transforms/cub.h index ced6688d..cf038512 100644 --- a/include/matx_cub.h +++ b/include/matx/transforms/cub.h @@ -33,8 +33,6 @@ #pragma once -#include "matx_error.h" -#include "matx_tensor.h" #include #include #ifdef __CUDACC__ @@ -42,6 +40,9 @@ #endif #include +#include "matx/core/error.h" +#include "matx/core/tensor.h" + namespace matx { using namespace std::placeholders; @@ -1762,4 +1763,4 @@ void unique(OutputTensor &a_out, CountTensor &num_found, const InputOperator &a, } #endif } -}; // namespace matx \ No newline at end of file +}; // namespace matx diff --git a/include/matx_einsum.h b/include/matx/transforms/einsum.h similarity index 100% rename from include/matx_einsum.h rename to include/matx/transforms/einsum.h diff --git a/include/matx_fft.h b/include/matx/transforms/fft.h similarity index 99% rename from include/matx_fft.h rename to include/matx/transforms/fft.h index 5b474ad5..eddfc992 100644 --- a/include/matx_fft.h +++ b/include/matx/transforms/fft.h @@ -35,10 +35,10 @@ #include #include -#include "matx_cache.h" -#include "matx_error.h" -#include "matx_make.h" -#include "matx_tensor.h" +#include "matx/core/cache.h" +#include "matx/core/error.h" +#include "matx/core/make_tensor.h" +#include "matx/core/tensor.h" #include #include diff --git a/include/matx_filter.h b/include/matx/transforms/filter.h similarity index 99% rename from include/matx_filter.h rename to include/matx/transforms/filter.h index dafcf7aa..f7a883ee 100644 --- a/include/matx_filter.h +++ b/include/matx/transforms/filter.h @@ -32,16 +32,17 @@ #pragma once -#include "matx_conv.h" -#include "matx_error.h" -#include "matx_filter_kernels.cuh" -#include "matx_tensor.h" #include #include #include #include #include +#include "matx/core/error.h" +#include "matx/core/tensor.h" +#include "matx/transforms/conv.h" +#include "matx/kernels/filter.cuh" + namespace matx { namespace detail { template #include diff --git a/include/matx_matmul.h b/include/matx/transforms/matmul.h similarity index 99% rename from include/matx_matmul.h rename to include/matx/transforms/matmul.h index 60f0c2df..d3a60b6f 100644 --- a/include/matx_matmul.h +++ b/include/matx/transforms/matmul.h @@ -32,9 +32,6 @@ #pragma once -#include "cublas_v2.h" -#include "matx_error.h" -#include "matx_tensor.h" #include #if MATX_ENABLE_CUTLASS == 1 @@ -45,6 +42,11 @@ #include #include +#include "cublas_v2.h" +#include "matx/core/error.h" +#include "matx/core/tensor.h" +#include "matx/core/cache.h" + namespace matx { /** diff --git a/include/matx/transforms/permute.h b/include/matx/transforms/permute.h index 84a15ed4..fe4471ab 100644 --- a/include/matx/transforms/permute.h +++ b/include/matx/transforms/permute.h @@ -33,7 +33,7 @@ #pragma once -#include "matx_type_utils.h" +#include "matx/core/type_utils.h" namespace matx { diff --git a/include/matx_reduce.h b/include/matx/transforms/reduce.h similarity index 99% rename from include/matx_reduce.h rename to include/matx/transforms/reduce.h index 66d4bc2c..567185a4 100644 --- a/include/matx_reduce.h +++ b/include/matx/transforms/reduce.h @@ -32,11 +32,11 @@ #pragma once -#include "matx_cub.h" -#include "matx_error.h" -#include "matx_get_grid_dims.h" -#include "matx_tensor.h" -#include "matx_type_utils.h" +#include "matx/core/error.h" +#include "matx/core/get_grid_dims.h" +#include "matx/core/tensor.h" +#include "matx/core/type_utils.h" +#include "matx/transforms/cub.h" #include #ifdef __CUDACC__ diff --git a/include/matx_solver.h b/include/matx/transforms/solver.h similarity index 99% rename from include/matx_solver.h rename to include/matx/transforms/solver.h index 391b115a..abde65cd 100644 --- a/include/matx_solver.h +++ b/include/matx/transforms/solver.h @@ -34,8 +34,8 @@ #include "cublas_v2.h" #include "cusolverDn.h" -#include "matx_error.h" -#include "matx_tensor.h" +#include "matx/core/error.h" +#include "matx/core/tensor.h" #include #include diff --git a/include/matx/transforms/transforms.h b/include/matx/transforms/transforms.h index e9c2625a..51c69e8a 100644 --- a/include/matx/transforms/transforms.h +++ b/include/matx/transforms/transforms.h @@ -32,6 +32,18 @@ #pragma once +#include "matx/transforms/ambgfun.h" +#include "matx/transforms/conv.h" #include "matx/transforms/copy.h" -#include "matx/transforms/transpose.h" +#include "matx/transforms/corr.h" +#include "matx/transforms/cov.h" +#include "matx/transforms/cub.h" +#include "matx/transforms/einsum.h" +#include "matx/transforms/fft.h" +#include "matx/transforms/filter.h" +#include "matx/transforms/inverse.h" +#include "matx/transforms/matmul.h" #include "matx/transforms/permute.h" +#include "matx/transforms/reduce.h" +#include "matx/transforms/solver.h" +#include "matx/transforms/transpose.h" diff --git a/include/matx/transforms/transpose.h b/include/matx/transforms/transpose.h index 828c8353..1030b71d 100644 --- a/include/matx/transforms/transpose.h +++ b/include/matx/transforms/transpose.h @@ -33,8 +33,8 @@ #pragma once -#include "matx_type_utils.h" -#include "kernels/matx_transpose.cuh" +#include "matx/core/type_utils.h" +#include "matx/kernels/transpose.cuh" namespace matx { diff --git a/include/matx_tensor_generators.h b/include/matx_tensor_generators.h deleted file mode 100644 index 4049f5b8..00000000 --- a/include/matx_tensor_generators.h +++ /dev/null @@ -1,1216 +0,0 @@ -//////////////////////////////////////////////////////////////////////////////// -// 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 -#include - -#include "matx_tensor.h" - -namespace matx { - -namespace detail { -template class ConstVal { - static constexpr int RANK = std::tuple_size::type>::value; - -private: - ShapeType s_; - T v_; - -public: - // dummy type to signal this is a matxop - using matxop = bool; - using scalar_type = T; - - ConstVal(ShapeType &&s, T val) : s_(std::forward(s)), v_(val){}; - - template - __MATX_INLINE__ __MATX_DEVICE__ __MATX_HOST__ auto operator()(Is...) const { - return v_; }; - - constexpr inline __MATX_HOST__ __MATX_DEVICE__ auto Size(int dim) const - { - return *(s_.begin() + dim); - } - static inline constexpr __MATX_HOST__ __MATX_DEVICE__ int32_t Rank() { return RANK; } -}; -} - -/** - * Return zero for all elements - * - * Zeros is used as an operator that always returns a 0 type for all - * elements. It can be used in place of memset to zero a block of memory. - * - * @tparam T - * Data type - * - * @param s - * Shape of tensor - */ -template ::type>, bool> = true> -inline auto zeros(ShapeType &&s) -{ - return detail::ConstVal(std::forward(s), T(0)); -} - -/** - * Return zero for all elements - * - * Zeros is used as an operator that always returns a 0 type for all - * elements. It can be used in place of memset to zero a block of memory. - * - * @tparam T - * Data type - * - * @param s - * Shape of tensor - */ -template -inline auto zeros(const index_t (&s)[RANK]) -{ - return zeros(detail::to_array(s)); -} - -/** - * Return one for all elements - * - * Ones is used as an operator that always returns a 1 type for all - * elements. It can be used in place of memset to set all values to 1. - * - * @tparam T - * Data type - * - * @param s - * Shape of tensor - */ -template ::type>, bool> = true> -inline auto ones(ShapeType &&s) -{ - return detail::ConstVal(std::forward(s), T(1)); -} - -/** - * Return one for all elements - * - * Ones is used as an operator that always returns a 1 type for all - * elements. It can be used in place of memset to set all values to 1. - * - * @tparam T - * Data type - * - * @param s - * Shape of tensor - */ -template -inline auto ones(const index_t (&s)[RANK]) -{ - return ones(detail::to_array(s)); -} - -namespace detail { -template class Diag { - static constexpr int RANK = std::tuple_size>::value; - -private: - ShapeType s_; - T val_; - -public: - // dummy type to signal this is a matxop - using matxop = bool; - using scalar_type = T; - - Diag(ShapeType &&s, T val) : s_(std::forward(s)), val_(val) - { - static_assert(Rank() > 1, "Diagonal generator must be used with an operator of rank 1 or higher"); - }; - - template - __MATX_INLINE__ __MATX_DEVICE__ __MATX_HOST__ auto operator()(Is... indices) const { - if (((pp_get<0>(indices...) == indices) && ...)) { - return T(val_); - } - else { - return T(0.0f); - } - } - - constexpr inline __MATX_HOST__ __MATX_DEVICE__ auto Size(int dim) const - { - return *(s_.begin() + dim); - } - static inline constexpr __MATX_HOST__ __MATX_DEVICE__ int32_t Rank() { return RANK; } -}; -} - -/** - * Creates a diagonal tensor with a given value on the diagonals - * - * diag returns a given value on all elements on the diagonals of a tensor, and - * 0 otherwise. In other words, if the index of every dimension is the same, the - * value is returned, otherwise a zero is returned. - * - * @tparam T - * Data type - * - */ -template ::type>, bool> = true> -inline auto diag(ShapeType &&s, T val) -{ - return detail::Diag(std::forward(s), val); -} - -/** - * Creates a diagonal tensor with a given value on the diagonals - * - * diag returns a given value on all elements on the diagonals of a tensor, and - * 0 otherwise. In other words, if the index of every dimension is the same, the - * value is returned, otherwise a zero is returned. - * - * @tparam T - * Data type - * - */ -template -inline auto diag(const index_t (&s)[RANK], T val) -{ - return diag(detail::to_array(s), val); -} - -/** - * Creates an identity patterns on the tensor - * - * eye returns 1 on all elements on the diagonals of a tensor, and 0 otherwise. - * In other words, if the index of every dimension is the same, a 1 is returned, - * otherwise a zero is returned. - * - * @tparam T - * Data type - * - */ -template ::type>, bool> = true> -inline auto eye(ShapeType &&s) -{ - return detail::Diag(std::forward(s), T(1)); -} - -/** - * Creates an identity patterns on the tensor - * - * eye returns 1 on all elements on the diagonals of a tensor, and 0 otherwise. - * In other words, if the index of every dimension is the same, a 1 is returned, - * otherwise a zero is returned. - * - * @tparam T - * Data type - * - */ -template inline auto eye(const index_t (&s)[RANK]) -{ - return eye(detail::to_array(s)); -} - -namespace detail { -template class matxGenerator1D_t { -public: - // dummy type to signal this is a matxop - using matxop = bool; - using scalar_type = typename Generator1D::scalar_type; - static constexpr int RANK = std::tuple_size>::value; - - template - matxGenerator1D_t(S &&s, Generator1D f) : f_(f), s_(std::forward(s)) {} - - template - __MATX_INLINE__ __MATX_DEVICE__ __MATX_HOST__ auto operator()(Is... indices) const { - return f_(pp_get(indices...)); - } - - constexpr inline __MATX_HOST__ __MATX_DEVICE__ auto Size(int dim) const - { - return *(s_.begin() + dim); - } - static inline constexpr __MATX_HOST__ __MATX_DEVICE__ int32_t Rank() { return RANK; } - -private: - Generator1D f_; - ShapeType s_; -}; -} - - -namespace detail { -template class Alternating { -private: - index_t size_; - -public: - using scalar_type = T; - - inline __MATX_HOST__ __MATX_DEVICE__ Alternating(index_t size) : size_(size) {}; - inline __MATX_HOST__ __MATX_DEVICE__ T operator()(index_t i) const - { - return (-2 * (i & 1)) + 1; - } -}; -} - - -/** - * Creates an alternating +1/-1 sequence - * - * @tparam T - * Data type - * @tparam Dim - * Dimension to create window over - * @tparam RANK - * The RANK of the shape, can be deduced from shape - * - * @param s - * The shape of the tensor - * - * Returns values for alternating sequence - */ -template ::type>, bool> = true> -inline auto alternate(ShapeType &&s) -{ - constexpr int RANK = std::tuple_size>::value; - static_assert(RANK > Dim); - detail::Alternating h( *(s.begin() + Dim)); - return detail::matxGenerator1D_t, Dim, ShapeType>(std::forward(s), h); -} - -/** - * Creates an alternating +1/-1 sequence - * - * @tparam T - * Data type - * @tparam Dim - * Dimension to create window over - * @tparam RANK - * The RANK of the shape, can be deduced from shape - * - * @param s - * C array representing shape of the tensor - * - */ -template -inline auto alternate(const index_t (&s)[RANK]) -{ - return alternate(detail::to_array(s)); -} - -namespace detail { -template class Hamming { -private: - index_t size_; - -public: - using scalar_type = T; - - inline __MATX_HOST__ __MATX_DEVICE__ Hamming(index_t size) : size_(size){}; - - inline __MATX_HOST__ __MATX_DEVICE__ T operator()(index_t i) const - { - return T(.54) - T(.46) * cuda::std::cos(T(2 * M_PI) * T(i) / T(size_ - 1)); - } -}; -} - - -/** - * Creates a Hamming window operator of shape s with the - * window applies along the specified dimension - * - * @tparam T - * Data type - * @tparam Dim - * Dimension to create window over - * @tparam RANK - * The RANK of the shape, can be deduced from shape - * - * @param s - * The shape of the tensor - * - * Returns values for a Hamming window across the selected dimension. - */ -template ::type>, bool> = true> -inline auto hamming(ShapeType &&s) -{ - constexpr int RANK = std::tuple_size>::value; - static_assert(RANK > Dim); - detail::Hamming h( *(s.begin() + Dim)); - return detail::matxGenerator1D_t, Dim, ShapeType>(std::forward(s), h); -} - -/** - * Creates a Hamming window operator of C-array shape s with the - * window applies along the specified dimension - * - * @tparam T - * Data type - * @tparam Dim - * Dimension to create window over - * @tparam RANK - * The RANK of the shape, can be deduced from shape - * - * @param s - * C array representing shape of the tensor - * - * Returns values for a Hamming window across the selected dimension. - */ -template -inline auto hamming(const index_t (&s)[RANK]) -{ - return hamming(detail::to_array(s)); -} - - -namespace detail { -template class Hanning { -private: - index_t size_; - -public: - using scalar_type = T; - inline __MATX_HOST__ __MATX_DEVICE__ Hanning(index_t size) : size_(size){}; - - inline __MATX_HOST__ __MATX_DEVICE__ T operator()(index_t i) const - { - return T(0.5) * (1 - cuda::std::cos(T(2 * M_PI) * T(i) / T(size_ - 1))); - } -}; -} - -/** - * Creates a Hanning window operator of shape s with the - * window applies along the x, y, z, or w dimension - * - * @tparam T - * Data type - * @tparam Dim - * Dimension to create window over - * @tparam RANK - * The RANK of the shape - * - * @param s - * The shape of the tensor - * - * Returns values for a Hanning window across the selected dimension. - */ -template ::type>, bool> = true> -inline auto hanning(ShapeType &&s) -{ - constexpr int RANK = std::tuple_size>::value; - static_assert(RANK > Dim); - detail::Hanning h( *(s.begin() + Dim)); - return detail::matxGenerator1D_t, Dim, ShapeType>(std::forward(s), h); -} - -/** - * Creates a Hanning window operator of shape s with the - * window applies along the x, y, z, or w dimension - * - * @tparam T - * Data type - * @tparam Dim - * Dimension to create window over - * @tparam RANK - * The RANK of the shape - * - * @param s - * The C-array shape of the tensor - * - * Returns values for a Hanning window across the selected dimension. - */ -template -inline auto hanning(const index_t (&s)[RANK]) -{ - return hanning(detail::to_array(s)); -} - - -namespace detail { -template class FlatTop { -private: - index_t size_; - - static constexpr T a0 = 0.21557895; - static constexpr T a1 = 0.41663158; - static constexpr T a2 = 0.277263158; - static constexpr T a3 = 0.083578947; - static constexpr T a4 = 0.006947368; - -public: - - using scalar_type = T; - inline __MATX_HOST__ __MATX_DEVICE__ FlatTop(index_t size) : size_(size){}; - - inline __MATX_HOST__ __MATX_DEVICE__ T operator()(index_t i) const - { - return a0 - - a1 * cuda::std::cos(2*M_PI*i / (size_ - 1)) - + a2 * cuda::std::cos(4*M_PI*i / (size_ - 1)) - - a3 * cuda::std::cos(6*M_PI*i / (size_ - 1)) - + a4 * cuda::std::cos(8*M_PI*i / (size_ - 1)); - } -}; -} - -/** - * Creates a Flattop window operator of shape s with the - * window applies along the x, y, z, or w dimension - * - * @tparam T - * Data type - * @tparam Dim - * Dimension to create window over - * @tparam RANK - * The RANK of the shape - * - * @param s - * The shape of the tensor - * - * Returns values for a Flattop window across the selected dimension. - */ -template ::type>, bool> = true> -inline auto flattop(ShapeType &&s) -{ - constexpr int RANK = std::tuple_size>::value; - static_assert(RANK > Dim); - detail::FlatTop h( *(s.begin() + Dim)); - return detail::matxGenerator1D_t, Dim, ShapeType>(std::forward(s), h); -} - -/** - * Creates a Flattop window operator of shape s with the - * window applies along the x, y, z, or w dimension - * - * @tparam T - * Data type - * @tparam Dim - * Dimension to create window over - * @tparam RANK - * The RANK of the shape - * - * @param s - * The shape of the tensor - * - * Returns values for a Flattop window across the selected dimension. - */ -template -inline auto flattop(const index_t (&s)[RANK]) -{ - return flattop(detail::to_array(s)); -} - - -namespace detail { -template class Blackman { -private: - index_t size_; - -public: - using scalar_type = T; - inline __MATX_HOST__ __MATX_DEVICE__ Blackman(index_t size) : size_(size){}; - - inline __MATX_HOST__ __MATX_DEVICE__ T operator()(index_t i) const - { - return T(0.42) + - ((T)0.5 * - cuda::std::cos(T(M_PI) * (1 - size_ + 2 * T(i)) / T(size_ - 1))) + - ((T)0.08 * cuda::std::cos(T(2 * M_PI) * (1 - size_ + 2 * T(i)) / - T(size_ - 1))); - } -}; -} - -/** - * Creates a Blackman window operator of shape s with the - * window applies along the x, y, z, or w dimension - * - * @tparam T - * Data type - * @tparam Dim - * Dimension to create window over - * @tparam RANK - * The RANK of the shape - * - * @param s - * The shape of the tensor - * - * Returns values for a Blackman window across the selected dimension. - */ -template ::type>, bool> = true> -inline auto blackman(ShapeType &&s) -{ - constexpr int RANK = std::tuple_size>::value; - static_assert(RANK > Dim); - detail::Blackman h( *(s.begin() + Dim)); - return detail::matxGenerator1D_t, Dim, ShapeType>(std::forward(s), h); -} - -/** - * Creates a Blackman window operator of shape s with the - * window applies along the x, y, z, or w dimension - * - * @tparam T - * Data type - * @tparam Dim - * Dimension to create window over - * @tparam RANK - * The RANK of the shape - * - * @param s - * The shape of the tensor - * - * Returns values for a Blackman window across the selected dimension. - */ -template -inline auto blackman(const index_t (&s)[RANK]) -{ - return blackman(detail::to_array(s)); -} - -namespace detail { -template class Bartlett { -private: - index_t size_; - -public: - using scalar_type = T; - inline __MATX_HOST__ __MATX_DEVICE__ Bartlett(index_t size) : size_(size){}; - - inline __MATX_HOST__ __MATX_DEVICE__ T operator()(index_t i) const - { - return (T(2) / (T(size_) - 1)) * - (((T(size_) - 1) / T(2)) - - cuda::std::abs(T(i) - ((T(size_) - 1) / T(2)))); - } -}; -} - -/** - * Creates a Bartlett window operator of shape s with the - * window applies along the x, y, z, or w dimension - * - * @tparam T - * Data type - * @tparam Dim - * Dimension to create window over - * @tparam RANK - * The RANK of the shape - * - * @param s - * The shape of the tensor - * - * Returns values for a Bartlett window across the selected dimension. - */ -template ::type>, bool> = true> -inline auto bartlett(ShapeType &&s) -{ - constexpr int RANK = std::tuple_size>::value; - static_assert(RANK > Dim); - detail::Bartlett h( *(s.begin() + Dim)); - return detail::matxGenerator1D_t, Dim, ShapeType>(std::forward(s), h); -} - -/** - * Creates a Bartlett window operator of shape s with the - * window applies along the x, y, z, or w dimension - * - * @tparam T - * Data type - * @tparam Dim - * Dimension to create window over - * @tparam RANK - * The RANK of the shape - * - * @param s - * The C array shape of the tensor - * - * Returns values for a Bartlett window across the selected dimension. - */ -template -inline auto bartlett(const index_t (&s)[RANK]) -{ - return bartlett(detail::to_array(s)); -} - - - -namespace detail { -template class Range { -private: - T first_; - T step_; - -public: - using scalar_type = T; - - Range() = default; - - Range(T first, T step) : first_(first), step_(step) {} - - __MATX_DEVICE__ __MATX_HOST__ __MATX_INLINE__ T operator()(index_t idx) const - { - if constexpr (is_matx_half_v) { - return first_ + T(static_cast((float)idx) * step_); - } - else { - return first_ + T(static_cast(idx) * step_); - } - - if constexpr (!is_matx_half_v) { - return first_ + T(static_cast(idx) * step_); - } - else { - return first_ + T(static_cast((float)idx) * step_); - } - } -}; -} - -/** - * Create a range of values along the x dimension - * - * Creates a range of values of type T with a start and step size. - * Value is determined by the index in operator() - * - * @param s - * Tensor shape - * @param first - * Starting value - * @param step - * Step size - * - */ -template ::type>, bool> = true> -inline auto range(ShapeType &&s, T first, T step) -{ - constexpr int RANK = std::tuple_size>::value; - static_assert(RANK > Dim); - detail::Range r(first, step); - return detail::matxGenerator1D_t, Dim, ShapeType>(std::forward(s), r); -} - -/** - * Create a range of values along the x dimension - * - * Creates a range of values of type T with a start and step size. - * Value is determined by the index in operator() - * - * @param s - * Tensor shape - * @param first - * Starting value - * @param step - * Step size - * - */ -template -inline auto range(const index_t (&s)[RANK], T first, T step) -{ - return range(detail::to_array(s), first, step); -} - - -namespace detail { -template class Linspace { -private: - Range range_; - -public: - using scalar_type = T; - - inline Linspace(T first, T last, index_t count) - { -#ifdef __CUDA_ARCH__ - range_ = Range{first, (last - first) / static_cast(count - 1)}; -#else - // Host has no support for most half precision operators/intrinsics - if constexpr (is_matx_half_v) { - range_ = Range{static_cast(first), - (static_cast(last) - static_cast(first)) / - static_cast(count - 1)}; - } - else { - range_ = Range{first, (last - first) / static_cast(count - 1)}; - } -#endif - } - - __MATX_DEVICE__ __MATX_HOST__ __MATX_INLINE__ T operator()(index_t idx) const { return range_(idx); } -}; -} - - -/** - * @brief Create a linearly-spaced range of values - * - * Creates a set of values using a start and end that are linearly- - * spaced apart over the set of values. Distance is determined - * by the shape and selected dimension. - * - * @tparam T Operator type - * @tparam Dim Dimension to operate over - * @tparam ShapeType Shape type - * @param s Shape object - * @param first First value - * @param last Last value - * @return Operator with linearly-spaced values - */ -template ::type>, bool> = true> -inline auto linspace(ShapeType &&s, T first, T last) -{ - constexpr int RANK = std::tuple_size>::value; - static_assert(RANK > Dim); - auto count = *(s.begin() + Dim); - detail::Linspace l(first, last, count); - return detail::matxGenerator1D_t, Dim, ShapeType>(std::forward(s), l); -} - -/** - * @brief Create a linearly-spaced range of values - * - * Creates a set of values using a start and end that are linearly- - * spaced apart over the set of values. Distance is determined - * by the shape and selected dimension. - * - * @tparam T Operator type - * @tparam Dim Dimension to operate over - * @tparam ShapeType Shape type - * @param s Shape object - * @param first First value - * @param last Last value - * @return Operator with linearly-spaced values - */ -template -inline auto linspace(const index_t (&s)[RANK], T first, T last) -{ - return linspace(detail::to_array(s), first, last); -} - -namespace detail { -template class Logspace { -private: - Range range_; - -public: - using scalar_type = T; - - inline Logspace(T first, T last, index_t count) - { -#ifdef __CUDA_ARCH__ - if constexpr (is_matx_half_v) { - range_ = Range{first, (last - first) / static_cast(count - 1.0f)}; - } - else { - range_ = Range{first, (last - first) / static_cast(count - 1)}; - } -#else - // Host has no support for most half precision operators/intrinsics - if constexpr (is_matx_half_v) { - range_ = Range{static_cast(first), - (static_cast(last) - static_cast(first)) / - static_cast(count - 1)}; - } - else { - range_ = Range{first, (last - first) / static_cast(count - 1)}; - } -#endif - } - - __MATX_DEVICE__ __MATX_HOST__ __MATX_INLINE__ T operator()(index_t idx) const - { - if constexpr (is_matx_half_v) { - return static_cast( - cuda::std::pow(10, static_cast(range_(idx)))); - } - else { - return cuda::std::pow(10, range_(idx)); - } - - // WAR for compiler bug. - if constexpr (!is_matx_half_v) { - return cuda::std::pow(10, range_(idx)); - } - else { - return static_cast( - cuda::std::pow(10, static_cast(range_(idx)))); - } - } -}; -} - - -/** - * @brief Create a log10-spaced range of values - * - * Creates a set of values using a start and end that are log10- - * spaced apart over the set of values. Distance is determined - * by the shape and selected dimension. - * - * @tparam T Operator type - * @tparam Dim Dimension to operate over - * @tparam ShapeType Shape type - * @param s Shape object - * @param first First value - * @param last Last value - * @return Operator with log10-spaced values - */ -template ::type>, bool> = true> -inline auto logspace(ShapeType &&s, T first, T last) -{ - constexpr int RANK = std::tuple_size>::value; - static_assert(RANK > Dim); - auto count = *(s.begin() + Dim); - detail::Logspace l(first, last, count); - return detail::matxGenerator1D_t, Dim, ShapeType>(std::forward(s), l); -} - -/** - * @brief Create a log10-spaced range of values - * - * Creates a set of values using a start and end that are log10- - * spaced apart over the set of values. Distance is determined - * by the shape and selected dimension. - * - * @tparam T Operator type - * @tparam Dim Dimension to operate over - * @tparam ShapeType Shape type - * @param s Shape object - * @param first First value - * @param last Last value - * @return Operator with log10-spaced values - */ -template -inline auto logspace(const index_t (&s)[RANK], T first, T last) -{ - return logspace(detail::to_array(s), first, last); -} - - -enum class ChirpMethod { - CHIRP_METHOD_LINEAR -}; - -enum class ChirpType { - CHIRP_TYPE_REAL, - CHIRP_TYPE_COMPLEX -}; - -namespace detail { -template -class Chirp { - using space_type = typename SpaceOp::scalar_type; - - -private: - SpaceOp sop_; - FreqType f0_; - FreqType f1_; - space_type t1_; - ChirpMethod method_; - -public: - using scalar_type = FreqType; - using matxop = bool; - inline __MATX_HOST__ __MATX_DEVICE__ Chirp(SpaceOp sop, FreqType f0, space_type t1, FreqType f1, ChirpMethod method) : - sop_(sop), - f0_(f0), - t1_(t1), - f1_(f1), - method_(method) - {} - - inline __MATX_HOST__ __MATX_DEVICE__ auto operator()(index_t i) const - { - if (method_ == ChirpMethod::CHIRP_METHOD_LINEAR) { - return cuda::std::cos(2.0f * M_PI * (f0_ * sop_(i) + 0.5f * ((f1_ - f0_) / t1_) * sop_(i) * sop_(i))); - } - - return 0.0; - } - - constexpr inline __MATX_HOST__ __MATX_DEVICE__ index_t Size([[maybe_unused]] int dim) const - { - return sop_.Size(0); - } - static inline constexpr __MATX_HOST__ __MATX_DEVICE__ int32_t Rank() { return 1; } -}; - -template -class ComplexChirp { - using space_type = typename SpaceOp::scalar_type; - - -private: - SpaceOp sop_; - FreqType f0_; - FreqType f1_; - space_type t1_; - ChirpMethod method_; - -public: - using scalar_type = cuda::std::complex; - using matxop = bool; - inline __MATX_HOST__ __MATX_DEVICE__ ComplexChirp(SpaceOp sop, FreqType f0, space_type t1, FreqType f1, ChirpMethod method) : - sop_(sop), - f0_(f0), - t1_(t1), - f1_(f1), - method_(method) - {} - - inline __MATX_HOST__ __MATX_DEVICE__ auto operator()(index_t i) const - { - if (method_ == ChirpMethod::CHIRP_METHOD_LINEAR) { - FreqType real = cuda::std::cos(2.0f * M_PI * (f0_ * sop_(i) + 0.5f * ((f1_ - f0_) / t1_) * sop_(i) * sop_(i))); - FreqType imag = -cuda::std::cos(2.0f * M_PI * (f0_ * sop_(i) + 0.5f * ((f1_ - f0_) / t1_) * sop_(i) * sop_(i) + 90.0/360.0)); - return cuda::std::complex{real, imag}; - } - - return cuda::std::complex{0, 0}; - } - - constexpr inline __MATX_HOST__ __MATX_DEVICE__ index_t Size([[maybe_unused]] int dim) const - { - return sop_.Size(0); - } - static inline constexpr __MATX_HOST__ __MATX_DEVICE__ int32_t Rank() { return 1; } -}; -} - -namespace signal { -/** - * Creates a real chirp signal (swept-frequency cosine) - * - * SpaceOp provides the time vector with custom spacing. - * - * @tparam FreqType - * Frequency data type - * @tparam SpaceOp - * Operator type of spacer - * - * @param t - * Vector representing values in time - * @param f0 - * Instantenous frequency at time 0 - * @param t1 - * Time for f1 - * @param f1 - * Frequency (Hz) at time t1 - * @param method - * Chirp method (CHIRP_METHOD_LINEAR) - * - * @returns The chirp operator - */ -template -inline auto chirp(SpaceOp t, FreqType f0, typename SpaceOp::scalar_type t1, FreqType f1, ChirpMethod method = ChirpMethod::CHIRP_METHOD_LINEAR) -{ - MATX_ASSERT_STR(method == ChirpMethod::CHIRP_METHOD_LINEAR, matxInvalidType, "Only linear chirps are supported") - - return detail::Chirp(t, f0, t1, f1, method); -} - -/** - * Creates a complex chirp signal (swept-frequency cosine) - * - * SpaceOp provides the time vector with custom spacing. - * - * @tparam FreqType - * Frequency data type - * @tparam SpaceOp - * Operator type of spacer - * - * @param t - * Vector representing values in time - * @param f0 - * Instantenous frequency at time 0 - * @param t1 - * Time for f1 - * @param f1 - * Frequency (Hz) at time t1 - * @param method - * Chirp method (CHIRP_METHOD_LINEAR) - * - * @returns The chirp operator - */ -template -inline auto cchirp(SpaceOp t, FreqType f0, typename SpaceOp::scalar_type t1, FreqType f1, ChirpMethod method = ChirpMethod::CHIRP_METHOD_LINEAR) -{ - MATX_ASSERT_STR(method == ChirpMethod::CHIRP_METHOD_LINEAR, matxInvalidType, "Only linear chirps are supported") - - return detail::ComplexChirp(t, f0, t1, f1, method); -} - -/** - * Creates a chirp signal (swept-frequency cosine) - * - * Creates a linearly-spaced sequence from 0 to "last" with "num" elements in between. Each step is - * of size 1/num. - * - * @tparam FreqType - * Frequency data type - * @tparam TimeType - * Type of time vector - * @tparam Method - * Chirp method (CHIRP_METHOD_LINEAR) - * - * @param num - * Number of time samples - * @param last - * Last time sample value - * @param f0 - * Instantenous frequency at time 0 - * @param t1 - * Time for f1 - * @param f1 - * Frequency (Hz) at time t1 - * @param method - * Method to use to generate the chirp - * - * @returns The chirp operator - */ -template -inline auto chirp(index_t num, TimeType last, FreqType f0, TimeType t1, FreqType f1, ChirpMethod method = ChirpMethod::CHIRP_METHOD_LINEAR) -{ - std::array shape = {num}; - auto space = linspace<0>(std::move(shape), (TimeType)0, last); - return chirp(space, f0, t1, f1, method); -} - -template -inline auto cchirp(index_t num, TimeType last, FreqType f0, TimeType t1, FreqType f1, ChirpMethod method = ChirpMethod::CHIRP_METHOD_LINEAR) -{ - std::array shape = {num}; - auto space = linspace<0>(std::move(shape), (TimeType)0, last); - return cchirp(space, f0, t1, f1, method); -} -} - - - -namespace detail { -template class Meshgrid_X { -private: - std::array x_; - std::array y_; - -public: - // dummy type to signal this is a matxop - using matxop = bool; - using scalar_type = T; - - Meshgrid_X(std::array x, std::array y) : x_(x), y_(y) {} - - inline __MATX_DEVICE__ T operator()(index_t i, index_t j) const - { - return x_[0] + j * (x_[1] - x_[0]) / (x_[2] - 1); - } - - constexpr inline __MATX_HOST__ __MATX_DEVICE__ auto Size(int dim) const - { - return (dim == 0) ? y_[2] : x_[2]; - } - static inline constexpr __MATX_HOST__ __MATX_DEVICE__ int32_t Rank() { return 2; } -}; - -template class Meshgrid_Y { -private: - std::array x_; - std::array y_; - -public: - // dummy type to signal this is a matxop - using matxop = bool; - using scalar_type = T; - - Meshgrid_Y(std::array x, std::array y) : x_(x), y_(y) {} - - inline __MATX_DEVICE__ T operator()(index_t i, index_t j) const - { - return y_[0] + i * (y_[1] - y_[0]) / (y_[2] - 1); - }; - - constexpr inline __MATX_HOST__ __MATX_DEVICE__ index_t Size(uint32_t dim) const - { - return (dim == 0) ? y_[2] : x_[2]; - } - static inline constexpr __MATX_HOST__ __MATX_DEVICE__ int32_t Rank() { return 2; } -}; -} - -/** - * Creates an mesh grid X matrix - * - * - * @tparam T - * Data type - * - */ -template -inline auto meshgrid_x(const std::array &x, const std::array &y) -{ - return detail::Meshgrid_X(x, y); -} - -/** - * Creates an mesh grid Y matrix - * - * - * @tparam T - * Data type - * - */ -template -inline auto meshgrid_y(const std::array &x, const std::array &y) -{ - return detail::Meshgrid_Y(x, y); -} -} // end namespace matx diff --git a/test/00_io/FileIOTests.cu b/test/00_io/FileIOTests.cu index 9d995bfd..9acd5f3e 100644 --- a/test/00_io/FileIOTests.cu +++ b/test/00_io/FileIOTests.cu @@ -32,8 +32,6 @@ #include "assert.h" #include "matx.h" -#include "matx_file_io.h" -#include "matx_pybind.h" #include "test_types.h" #include "utilities.h" #include "gtest/gtest.h" diff --git a/test/00_operators/GeneratorTests.cu b/test/00_operators/GeneratorTests.cu index 221a7285..94af1162 100644 --- a/test/00_operators/GeneratorTests.cu +++ b/test/00_operators/GeneratorTests.cu @@ -32,7 +32,6 @@ #include "assert.h" #include "matx.h" -#include "matx_pybind.h" #include "test_types.h" #include "utilities.h" #include "gtest/gtest.h" @@ -540,11 +539,11 @@ TYPED_TEST(BasicGeneratorTestsFloatNonComplexNonHalf, Chirp) "01_signal", "chirp", "run", {count, static_cast(end), static_cast(f0), static_cast(f1)}); auto t1 = make_tensor({count}); - (t1 = signal::chirp(count, end, f0, end, f1)).run(); + (t1 = chirp(count, end, f0, end, f1)).run(); MATX_TEST_ASSERT_COMPARE(pb, t1, "Y", 0.01); auto t1c = make_tensor>({count}); - (t1c = signal::cchirp(count, end, f0, end, f1, ChirpMethod::CHIRP_METHOD_LINEAR)).run(); + (t1c = cchirp(count, end, f0, end, f1, ChirpMethod::CHIRP_METHOD_LINEAR)).run(); pb.reset(); MATX_EXIT_HANDLER(); diff --git a/test/00_operators/OperatorTests.cu b/test/00_operators/OperatorTests.cu index 1fe082c1..b90417d7 100644 --- a/test/00_operators/OperatorTests.cu +++ b/test/00_operators/OperatorTests.cu @@ -32,7 +32,6 @@ #include "assert.h" #include "matx.h" -#include "matx_pybind.h" #include "test_types.h" #include "utilities.h" #include "gtest/gtest.h" diff --git a/test/00_operators/PythonEmbed.cu b/test/00_operators/PythonEmbed.cu index 8d20d565..cd1eade1 100644 --- a/test/00_operators/PythonEmbed.cu +++ b/test/00_operators/PythonEmbed.cu @@ -30,7 +30,6 @@ // OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. ///////////////////////////////////////////////////////////////////////////////// #include "matx.h" -#include "matx_pybind.h" #include "utilities.h" #include "gtest/gtest.h" @@ -41,4 +40,4 @@ TEST(BasicPythonTest, Print) auto pb = detail::MatXPybind{}; pb.InitAndRunTVGenerator("00_python_tests", "matx_python_tests", "run", {}); -} \ No newline at end of file +} diff --git a/test/00_operators/ReductionTests.cu b/test/00_operators/ReductionTests.cu index 521ae38e..a68bc21d 100644 --- a/test/00_operators/ReductionTests.cu +++ b/test/00_operators/ReductionTests.cu @@ -32,7 +32,6 @@ #include "assert.h" #include "matx.h" -#include "matx_pybind.h" #include "test_types.h" #include "utilities.h" #include "gtest/gtest.h" diff --git a/test/00_solver/Cholesky.cu b/test/00_solver/Cholesky.cu index 8b2eb4eb..ba2e6974 100644 --- a/test/00_solver/Cholesky.cu +++ b/test/00_solver/Cholesky.cu @@ -32,7 +32,6 @@ #include "assert.h" #include "matx.h" -#include "matx_pybind.h" #include "test_types.h" #include "utilities.h" #include "gtest/gtest.h" diff --git a/test/00_solver/Det.cu b/test/00_solver/Det.cu index 33bd2fc7..961ec569 100644 --- a/test/00_solver/Det.cu +++ b/test/00_solver/Det.cu @@ -32,7 +32,6 @@ #include "assert.h" #include "matx.h" -#include "matx_pybind.h" #include "test_types.h" #include "utilities.h" #include "gtest/gtest.h" diff --git a/test/00_solver/Eigen.cu b/test/00_solver/Eigen.cu index 1eef9c94..9f80fa7f 100644 --- a/test/00_solver/Eigen.cu +++ b/test/00_solver/Eigen.cu @@ -32,7 +32,6 @@ #include "assert.h" #include "matx.h" -#include "matx_pybind.h" #include "test_types.h" #include "utilities.h" #include "gtest/gtest.h" diff --git a/test/00_solver/LU.cu b/test/00_solver/LU.cu index fcd0cfd2..758f3cbb 100644 --- a/test/00_solver/LU.cu +++ b/test/00_solver/LU.cu @@ -32,7 +32,6 @@ #include "assert.h" #include "matx.h" -#include "matx_pybind.h" #include "test_types.h" #include "utilities.h" #include "gtest/gtest.h" diff --git a/test/00_solver/QR.cu b/test/00_solver/QR.cu index f73e3311..40a9913e 100644 --- a/test/00_solver/QR.cu +++ b/test/00_solver/QR.cu @@ -32,7 +32,6 @@ #include "assert.h" #include "matx.h" -#include "matx_pybind.h" #include "test_types.h" #include "utilities.h" #include "gtest/gtest.h" diff --git a/test/00_solver/SVD.cu b/test/00_solver/SVD.cu index 6e45bc86..7e7889d6 100644 --- a/test/00_solver/SVD.cu +++ b/test/00_solver/SVD.cu @@ -32,7 +32,6 @@ #include "assert.h" #include "matx.h" -#include "matx_pybind.h" #include "test_types.h" #include "utilities.h" #include "gtest/gtest.h" diff --git a/test/00_tensor/CUBTests.cu b/test/00_tensor/CUBTests.cu index 5c3494d5..aa63062c 100644 --- a/test/00_tensor/CUBTests.cu +++ b/test/00_tensor/CUBTests.cu @@ -32,7 +32,6 @@ #include "assert.h" #include "matx.h" -#include "matx_cub.h" #include "test_types.h" #include "utilities.h" #include "gtest/gtest.h" diff --git a/test/00_tensor/EinsumTests.cu b/test/00_tensor/EinsumTests.cu index 8d572cf9..9584b4a5 100644 --- a/test/00_tensor/EinsumTests.cu +++ b/test/00_tensor/EinsumTests.cu @@ -35,8 +35,6 @@ #include "test_types.h" #include "utilities.h" #include "gtest/gtest.h" -#include "matx_pybind.h" -#include "matx_einsum.h" using namespace matx; diff --git a/test/00_transform/ConvCorr.cu b/test/00_transform/ConvCorr.cu index d6eec078..3a72238a 100644 --- a/test/00_transform/ConvCorr.cu +++ b/test/00_transform/ConvCorr.cu @@ -32,7 +32,6 @@ #include "assert.h" #include "matx.h" -#include "matx_pybind.h" #include "test_types.h" #include "utilities.h" #include "gtest/gtest.h" diff --git a/test/00_transform/Cov.cu b/test/00_transform/Cov.cu index 066ba0b7..544dd8c1 100644 --- a/test/00_transform/Cov.cu +++ b/test/00_transform/Cov.cu @@ -32,7 +32,6 @@ #include "assert.h" #include "matx.h" -#include "matx_pybind.h" #include "test_types.h" #include "utilities.h" #include "gtest/gtest.h" diff --git a/test/00_transform/FFT.cu b/test/00_transform/FFT.cu index 7b4bbf4f..88414eea 100644 --- a/test/00_transform/FFT.cu +++ b/test/00_transform/FFT.cu @@ -32,7 +32,6 @@ #include "assert.h" #include "matx.h" -#include "matx_pybind.h" #include "test_types.h" #include "utilities.h" #include "gtest/gtest.h" diff --git a/test/00_transform/MatMul.cu b/test/00_transform/MatMul.cu index b0708ff7..5f9606a5 100644 --- a/test/00_transform/MatMul.cu +++ b/test/00_transform/MatMul.cu @@ -32,8 +32,6 @@ #include "assert.h" #include "matx.h" -#include "matx_matmul.h" -#include "matx_pybind.h" #include "test_types.h" #include "utilities.h" #include "gtest/gtest.h" @@ -267,4 +265,4 @@ TYPED_TEST(MatMulTestFloatTypes, MediumRectBatched) MATX_TEST_ASSERT_COMPARE(this->pb, c, "c", this->thresh); MATX_EXIT_HANDLER(); -} \ No newline at end of file +} diff --git a/test/01_radar/MVDRBeamformer.cu b/test/01_radar/MVDRBeamformer.cu index 07e82516..c2ce47fc 100644 --- a/test/01_radar/MVDRBeamformer.cu +++ b/test/01_radar/MVDRBeamformer.cu @@ -32,7 +32,6 @@ #include "assert.h" #include "matx.h" -#include "matx_pybind.h" #include "mvdr_beamformer.h" #include "utilities.h" #include "gtest/gtest.h" @@ -75,4 +74,4 @@ TEST(Radar, MVDRBeamformer) MATX_TEST_ASSERT_COMPARE(pb, cov_inv, "cov_inv", 0.01); MATX_EXIT_HANDLER(); -} \ No newline at end of file +} diff --git a/test/01_radar/MultiChannelRadarPipeline.cu b/test/01_radar/MultiChannelRadarPipeline.cu index 58f132e7..2d8c1af0 100644 --- a/test/01_radar/MultiChannelRadarPipeline.cu +++ b/test/01_radar/MultiChannelRadarPipeline.cu @@ -32,7 +32,6 @@ #include "assert.h" #include "matx.h" -#include "matx_pybind.h" #include "simple_pipeline.h" #include "test_types.h" #include "utilities.h" diff --git a/test/01_radar/ambgfun.cu b/test/01_radar/ambgfun.cu index ede2f164..92251ba2 100644 --- a/test/01_radar/ambgfun.cu +++ b/test/01_radar/ambgfun.cu @@ -32,8 +32,6 @@ #include "assert.h" #include "matx.h" -#include "matx_pybind.h" -#include "matx_radar.h" #include "utilities.h" #include "gtest/gtest.h" @@ -67,7 +65,7 @@ TEST_F(RadarAmbiguityFunction, Cut2D) {2 * sig_size - 1, (index_t)pow(2, std::ceil(std::log2(2 * sig_size - 1)))}); - signal::ambgfun(amf2dv, xv, 1e3, signal::AMGBFUN_CUT_TYPE_2D, 1.0); + ambgfun(amf2dv, xv, 1e3, AMGBFUN_CUT_TYPE_2D, 1.0); MATX_TEST_ASSERT_COMPARE(pb, amf2dv, "amf_2d", 0.01); MATX_EXIT_HANDLER(); @@ -80,7 +78,7 @@ TEST_F(RadarAmbiguityFunction, CutDelay) tensor_t amf_delay_v( {1, (index_t)pow(2, std::ceil(std::log2(2 * sig_size - 1)))}); - signal::ambgfun(amf_delay_v, xv, 1e3, signal::AMGBFUN_CUT_TYPE_DELAY, 1.0); + ambgfun(amf_delay_v, xv, 1e3, AMGBFUN_CUT_TYPE_DELAY, 1.0); auto delay1d = amf_delay_v.Slice<1>({0, 0}, {matxDropDim, matxEnd}); MATX_TEST_ASSERT_COMPARE(pb, delay1d, "amf_delay", 0.01); @@ -94,11 +92,11 @@ TEST_F(RadarAmbiguityFunction, CutDoppler) tensor_t amf_doppler_v({1, xv.Size(0) * 2 - 1}); - signal::ambgfun(amf_doppler_v, xv, 1e3, signal::AMGBFUN_CUT_TYPE_DOPPLER, + ambgfun(amf_doppler_v, xv, 1e3, AMGBFUN_CUT_TYPE_DOPPLER, 1.0); auto doppler1d = amf_doppler_v.Slice<1>({0, 0}, {matxDropDim, matxEnd}); MATX_TEST_ASSERT_COMPARE(pb, doppler1d, "amf_doppler", 0.01); MATX_EXIT_HANDLER(); -} \ No newline at end of file +} diff --git a/test/01_radar/dct.cu b/test/01_radar/dct.cu index 254e25b6..12c8fa4d 100644 --- a/test/01_radar/dct.cu +++ b/test/01_radar/dct.cu @@ -32,9 +32,6 @@ #include "assert.h" #include "matx.h" -#include "matx_pybind.h" -#include "matx_radar.h" -#include "matx_signal.h" #include "utilities.h" #include "gtest/gtest.h" @@ -65,7 +62,7 @@ TEST_F(DctTests, Real1DN100) MATX_ENTER_HANDLER(); tensor_t out{{sig_size}}; - signal::dct(out, xv); + dct(out, xv); MATX_TEST_ASSERT_COMPARE(pb, out, "Y", 0.01); MATX_EXIT_HANDLER(); diff --git a/test/include/test_types.h b/test/include/test_types.h index 4fbbc608..256e8410 100644 --- a/test/include/test_types.h +++ b/test/include/test_types.h @@ -31,11 +31,9 @@ ///////////////////////////////////////////////////////////////////////////////// #pragma once -#include "matx_half.h" -#include "matx_half_complex.h" -#include "matx_type_utils.h" -#include "gtest/gtest.h" #include +#include "gtest/gtest.h" +#include "matx.h" using testing::Types;