From 4487d1ea6886e2fb624eb1570941ed439065673e Mon Sep 17 00:00:00 2001 From: cliffburdick Date: Wed, 24 Jul 2024 15:29:38 -0700 Subject: [PATCH] Added toeplitz operator --- include/matx/core/file_io.h | 2 +- include/matx/core/iterator.h | 6 +- include/matx/core/make_tensor.h | 30 ++++---- include/matx/core/operator_utils.h | 12 +-- include/matx/core/pybind.h | 12 +-- include/matx/core/tensor.h | 2 +- include/matx/core/tensor_impl.h | 1 - include/matx/core/tensor_utils.h | 10 +-- include/matx/core/tie.h | 5 +- include/matx/core/type_utils.h | 20 ++--- include/matx/generators/alternate.h | 2 +- include/matx/generators/bartlett.h | 2 +- include/matx/generators/blackman.h | 2 +- include/matx/generators/chirp.h | 12 +-- include/matx/generators/diag.h | 6 +- include/matx/generators/fftfreq.h | 2 +- include/matx/generators/flattop.h | 2 +- include/matx/generators/generator1d.h | 2 +- include/matx/generators/hamming.h | 2 +- include/matx/generators/hanning.h | 2 +- include/matx/generators/linspace.h | 2 +- include/matx/generators/logspace.h | 2 +- include/matx/generators/meshgrid.h | 4 +- include/matx/generators/ones.h | 2 +- include/matx/generators/random.h | 4 +- include/matx/generators/range.h | 2 +- include/matx/generators/zeros.h | 2 +- include/matx/kernels/channelize_poly.cuh | 20 ++--- include/matx/kernels/conv.cuh | 12 +-- include/matx/kernels/resample_poly.cuh | 22 +++--- include/matx/kernels/transpose.cuh | 2 +- include/matx/operators/all.h | 6 +- include/matx/operators/ambgfun.h | 6 +- include/matx/operators/any.h | 6 +- include/matx/operators/argmax.h | 2 +- include/matx/operators/argmin.h | 2 +- include/matx/operators/at.h | 2 +- include/matx/operators/binary_operators.h | 6 +- include/matx/operators/cart2sph.h | 2 +- include/matx/operators/cast.h | 6 +- include/matx/operators/cgsolve.h | 6 +- include/matx/operators/channelize_poly.h | 4 +- include/matx/operators/chol.h | 4 +- include/matx/operators/clone.h | 2 +- include/matx/operators/collapse.h | 4 +- include/matx/operators/concat.h | 6 +- include/matx/operators/constval.h | 2 +- include/matx/operators/conv.h | 14 ++-- include/matx/operators/corr.h | 6 +- include/matx/operators/cov.h | 6 +- include/matx/operators/cumsum.h | 6 +- include/matx/operators/dct.h | 2 +- include/matx/operators/det.h | 6 +- include/matx/operators/diag.h | 2 +- include/matx/operators/eig.h | 2 +- include/matx/operators/einsum.h | 2 +- include/matx/operators/fft.h | 26 +++---- include/matx/operators/fftshift.h | 8 +- include/matx/operators/filter.h | 6 +- include/matx/operators/find.h | 2 +- include/matx/operators/find_idx.h | 2 +- include/matx/operators/flatten.h | 2 +- include/matx/operators/frexp.h | 12 +-- include/matx/operators/hermitian.h | 2 +- include/matx/operators/hist.h | 12 +-- include/matx/operators/if.h | 2 +- include/matx/operators/ifelse.h | 2 +- include/matx/operators/index.h | 2 +- include/matx/operators/interleaved.h | 12 +-- include/matx/operators/inverse.h | 6 +- include/matx/operators/isclose.h | 4 +- include/matx/operators/kronecker.h | 2 +- include/matx/operators/legendre.h | 14 ++-- include/matx/operators/lu.h | 6 +- include/matx/operators/matmul.h | 6 +- include/matx/operators/matvec.h | 6 +- include/matx/operators/max.h | 6 +- include/matx/operators/mean.h | 6 +- include/matx/operators/median.h | 6 +- include/matx/operators/min.h | 6 +- include/matx/operators/norm.h | 8 +- include/matx/operators/operators.h | 1 + include/matx/operators/outer.h | 6 +- include/matx/operators/overlap.h | 2 +- include/matx/operators/percentile.h | 6 +- include/matx/operators/permute.h | 2 +- include/matx/operators/planar.h | 6 +- include/matx/operators/polyval.h | 6 +- include/matx/operators/prod.h | 6 +- include/matx/operators/pwelch.h | 6 +- include/matx/operators/qr.h | 6 +- include/matx/operators/r2c.h | 2 +- include/matx/operators/reduce.h | 6 +- include/matx/operators/remap.h | 4 +- include/matx/operators/repmat.h | 2 +- include/matx/operators/resample_poly.h | 6 +- include/matx/operators/reshape.h | 2 +- include/matx/operators/reverse.h | 2 +- include/matx/operators/scalar_ops.h | 8 +- include/matx/operators/select.h | 2 +- include/matx/operators/self.h | 2 +- include/matx/operators/set.h | 2 +- include/matx/operators/shift.h | 2 +- include/matx/operators/sign.h | 16 ++-- include/matx/operators/slice.h | 2 +- include/matx/operators/softmax.h | 6 +- include/matx/operators/sort.h | 6 +- include/matx/operators/sph2cart.h | 2 +- include/matx/operators/stack.h | 6 +- include/matx/operators/stdd.h | 6 +- include/matx/operators/sum.h | 6 +- include/matx/operators/svd.h | 6 +- include/matx/operators/trace.h | 6 +- include/matx/operators/transpose.h | 4 +- include/matx/operators/unary_operators.h | 8 +- include/matx/operators/unique.h | 2 +- include/matx/operators/updownsample.h | 4 +- include/matx/operators/var.h | 6 +- include/matx/transforms/ambgfun.h | 2 +- include/matx/transforms/cgsolve.h | 22 +++--- include/matx/transforms/channelize_poly.h | 22 +++--- include/matx/transforms/conv.h | 26 +++---- include/matx/transforms/cov.h | 2 +- include/matx/transforms/cub.h | 44 +++++------ include/matx/transforms/einsum.h | 2 +- include/matx/transforms/fft/fft_common.h | 8 +- include/matx/transforms/fft/fft_cuda.h | 26 +++---- include/matx/transforms/fft/fft_fftw.h | 44 +++++------ include/matx/transforms/filter.h | 2 +- include/matx/transforms/inverse.h | 2 +- include/matx/transforms/matmul/matmul_cblas.h | 66 ++++++++--------- include/matx/transforms/matmul/matmul_cuda.h | 70 +++++++++--------- include/matx/transforms/percentile.h | 8 +- include/matx/transforms/pwelch.h | 4 +- include/matx/transforms/qr.h | 4 +- include/matx/transforms/reduce.h | 74 +++++++++---------- include/matx/transforms/resample_poly.h | 6 +- include/matx/transforms/solver.h | 34 ++++----- include/matx/transforms/svd.h | 8 +- include/matx/transforms/transpose.h | 2 +- test/00_operators/OperatorTests.cu | 40 ++++++++++ test/00_solver/SVD.cu | 12 +-- test/test_vectors/generators/00_operators.py | 26 +++++++ 143 files changed, 640 insertions(+), 575 deletions(-) diff --git a/include/matx/core/file_io.h b/include/matx/core/file_io.h index 3f09ed61a..61c38a9e5 100644 --- a/include/matx/core/file_io.h +++ b/include/matx/core/file_io.h @@ -158,7 +158,7 @@ void read_csv(TensorType &t, const std::string fname, auto np = pybind11::module_::import("numpy"); auto obj = np.attr("genfromtxt")("fname"_a = fname.c_str(), "delimiter"_a = delimiter, "skip_header"_a = skip_header, - "dtype"_a = detail::MatXPybind::GetNumpyDtype()); + "dtype"_a = detail::MatXPybind::GetNumpyDtype()); pb->NumpyToTensorView(t, obj); } diff --git a/include/matx/core/iterator.h b/include/matx/core/iterator.h index 33b119448..8cae6f759 100644 --- a/include/matx/core/iterator.h +++ b/include/matx/core/iterator.h @@ -47,8 +47,7 @@ namespace matx { template struct RandomOperatorIterator { using self_type = RandomOperatorIterator; - using value_type = typename std::conditional_t, typename OperatorType::scalar_type>; - using scalar_type = value_type; + using value_type = typename std::conditional_t, typename OperatorType::value_type>; // using stride_type = std::conditional_t, typename OperatorType::desc_type::stride_type, // index_t>; using stride_type = index_t; @@ -174,8 +173,7 @@ __MATX_INLINE__ __MATX_HOST__ __MATX_DEVICE__ index_t operator-(const RandomOper template struct RandomOperatorOutputIterator { using self_type = RandomOperatorOutputIterator; - using value_type = typename std::conditional_t, typename OperatorType::scalar_type>; - using scalar_type = value_type; + using value_type = typename std::conditional_t, typename OperatorType::value_type>; // using stride_type = std::conditional_t, typename OperatorType::desc_type::stride_type, // index_t>; using stride_type = index_t; diff --git a/include/matx/core/make_tensor.h b/include/matx/core/make_tensor.h index abec01be0..fde6c1e5a 100644 --- a/include/matx/core/make_tensor.h +++ b/include/matx/core/make_tensor.h @@ -78,7 +78,7 @@ void make_tensor( TensorType &tensor, cudaStream_t stream = 0) { MATX_NVTX_START("", matx::MATX_NVTX_LOG_API) - auto tmp = make_tensor(shape, space, stream); + auto tmp = make_tensor(shape, space, stream); tensor.Shallow(tmp); } @@ -166,7 +166,7 @@ auto make_tensor( TensorType &tensor, cudaStream_t stream = 0) { MATX_NVTX_START("", matx::MATX_NVTX_LOG_API) - auto tmp = make_tensor(std::forward(shape), space, stream); + auto tmp = make_tensor(std::forward(shape), space, stream); tensor.Shallow(tmp); } @@ -225,7 +225,7 @@ auto make_tensor_p( TensorType &tensor, cudaStream_t stream = 0) { MATX_NVTX_START("", matx::MATX_NVTX_LOG_API) - auto tmp = make_tensor(std::forward(shape), space, stream); + auto tmp = make_tensor(std::forward(shape), space, stream); tensor.Shallow(tmp); } @@ -261,7 +261,7 @@ template ({}, space, stream); + auto tmp = make_tensor({}, space, stream); tensor.Shallow(tmp); } @@ -322,12 +322,12 @@ auto make_tensor( T *data, template , bool> = true> auto make_tensor( TensorType &tensor, - typename TensorType::scalar_type *data, + typename TensorType::value_type *data, const index_t (&shape)[TensorType::Rank()], bool owning = false) { MATX_NVTX_START("", matx::MATX_NVTX_LOG_API) - auto tmp = make_tensor(data, shape, owning); + auto tmp = make_tensor(data, shape, owning); tensor.Shallow(tmp); } @@ -373,12 +373,12 @@ auto make_tensor( T *data, template , bool> = true> auto make_tensor( TensorType &tensor, - typename TensorType::scalar_type *data, + typename TensorType::value_type *data, typename TensorType::shape_container &&shape, bool owning = false) { MATX_NVTX_START("", matx::MATX_NVTX_LOG_API) - auto tmp = make_tensor(data, std::forward(shape), owning); + auto tmp = make_tensor(data, std::forward(shape), owning); tensor.Shallow(tmp); } @@ -414,9 +414,9 @@ auto make_tensor( T *ptr, template , bool> = true> auto make_tensor( TensorType &tensor, - typename TensorType::scalar_type *ptr, + typename TensorType::value_type *ptr, bool owning = false) { - auto tmp = make_tensor(ptr, owning); + auto tmp = make_tensor(ptr, owning); tensor.Shallow(tmp); } @@ -534,12 +534,12 @@ auto make_tensor( T* const data, template , bool> = true> auto make_tensor( TensorType &tensor, - typename TensorType::scalar_type* const data, + typename TensorType::value_type* const data, typename TensorType::desc_type &&desc, bool owning = false) { MATX_NVTX_START("", matx::MATX_NVTX_LOG_API) - auto tmp = make_tensor(data, std::forward(desc), owning); + auto tmp = make_tensor(data, std::forward(desc), owning); tensor.Shallow(tmp); } @@ -585,7 +585,7 @@ auto make_tensor( TensorType &&tensor, cudaStream_t stream = 0) { MATX_NVTX_START("", matx::MATX_NVTX_LOG_API) - auto tmp = make_tensor(std::forward(desc), space, stream); + auto tmp = make_tensor(std::forward(desc), space, stream); tensor.Shallow(tmp); } @@ -633,13 +633,13 @@ auto make_tensor( T *const data, template , bool> = true> auto make_tensor( TensorType &tensor, - typename TensorType::scalar_type *const data, + typename TensorType::value_type *const data, const index_t (&shape)[TensorType::Rank()], const index_t (&strides)[TensorType::Rank()], bool owning = false) { MATX_NVTX_START("", matx::MATX_NVTX_LOG_API) - auto tmp = make_tensor(data, shape, strides, owning); + auto tmp = make_tensor(data, shape, strides, owning); tensor.Shallow(tmp); } diff --git a/include/matx/core/operator_utils.h b/include/matx/core/operator_utils.h index 2fa7aee46..47eb3d9d3 100644 --- a/include/matx/core/operator_utils.h +++ b/include/matx/core/operator_utils.h @@ -44,13 +44,13 @@ namespace matx { if (out.IsContiguous()) { if constexpr(ConvertType) { return func( in, - reinterpret_cast::scalar_type> *>(out.Data()), + reinterpret_cast::value_type> *>(out.Data()), bi, ei); } else { return func( in, - reinterpret_cast::scalar_type *>(out.Data()), + reinterpret_cast::value_type *>(out.Data()), bi, ei); } @@ -70,14 +70,14 @@ namespace matx { if constexpr (ConvertType) { return ReduceOutput( std::forward(func), std::forward(out), - reinterpret_cast::scalar_type> *>(in_base.Data()), + reinterpret_cast::value_type> *>(in_base.Data()), BeginOffset{in_base}, EndOffset{in_base}); } else { return ReduceOutput( std::forward(func), std::forward(out), - reinterpret_cast::scalar_type *>(in_base.Data()), + reinterpret_cast::value_type *>(in_base.Data()), BeginOffset{in_base}, EndOffset{in_base}); } @@ -118,9 +118,9 @@ namespace matx { namespace detail { // Used inside of transforms to allocate temporary output template - __MATX_HOST__ __MATX_INLINE__ void AllocateTempTensor(TensorType &tensor, Executor &&ex, ShapeType &&shape, typename TensorType::scalar_type **ptr) { + __MATX_HOST__ __MATX_INLINE__ void AllocateTempTensor(TensorType &tensor, Executor &&ex, ShapeType &&shape, typename TensorType::value_type **ptr) { const auto ttl_size = std::accumulate(shape.begin(), shape.end(), static_cast(1), - std::multiplies()) * sizeof(typename TensorType::scalar_type); + std::multiplies()) * sizeof(typename TensorType::value_type); if constexpr (is_cuda_executor_v) { matxAlloc((void**)ptr, ttl_size, MATX_ASYNC_DEVICE_MEMORY, ex.getStream()); make_tensor(tensor, *ptr, shape); diff --git a/include/matx/core/pybind.h b/include/matx/core/pybind.h index e562712e8..a4cc6c529 100644 --- a/include/matx/core/pybind.h +++ b/include/matx/core/pybind.h @@ -176,7 +176,7 @@ class MatXPybind { template static pybind11::object GetEmptyNumpy(const TensorType &ten) { - using T = typename TensorType::scalar_type; + using T = typename TensorType::value_type; auto np = pybind11::module_::import("numpy"); pybind11::list dims; @@ -329,7 +329,7 @@ class MatXPybind { void NumpyToTensorView(TensorType ten, const pybind11::object &np_ten) { - using T = typename TensorType::scalar_type; + using T = typename TensorType::value_type; constexpr int RANK = TensorType::Rank(); static_assert(RANK <=5, "NumpyToTensorView only supports max(RANK) = 5 at the moment."); @@ -377,7 +377,7 @@ class MatXPybind { template auto NumpyToTensorView(const pybind11::object &np_ten) { - using T = typename TensorType::scalar_type; + using T = typename TensorType::value_type; constexpr int RANK = TensorType::Rank(); using ntype = matx_convert_complex_type; auto ften = pybind11::array_t(np_ten); @@ -398,7 +398,7 @@ class MatXPybind { template auto TensorViewToNumpy(const TensorType &ten) { - using tensor_type = typename TensorType::scalar_type; + using tensor_type = typename TensorType::value_type; using ntype = matx_convert_complex_type; constexpr int RANK = TensorType::Rank(); @@ -466,12 +466,12 @@ class MatXPybind { template > + typename CT = matx_convert_cuda_complex_type> std::optional> CompareOutput(const TensorType &ten, const std::string fname, double thresh, bool debug = false) { - using raw_type = typename TensorType::scalar_type; + using raw_type = typename TensorType::value_type; using ntype = matx_convert_complex_type; using ctype = matx_convert_cuda_complex_type; auto resobj = res_dict[fname.c_str()]; diff --git a/include/matx/core/tensor.h b/include/matx/core/tensor.h index 6330b2b67..cc8514092 100644 --- a/include/matx/core/tensor.h +++ b/include/matx/core/tensor.h @@ -87,7 +87,7 @@ class tensor_t : public detail::tensor_impl_t { public: // Type specifier for reflection on class using type = T; ///< Type of traits - using scalar_type = T; ///< Type of traits + using value_type = T; ///< Type of traits // Type specifier for signaling this is a matx operation or tensor view using matxop = bool; ///< Indicate this is a MatX operator using matxoplvalue = bool; ///< Indicate this is a MatX operator that can be on the lhs of an equation diff --git a/include/matx/core/tensor_impl.h b/include/matx/core/tensor_impl.h index c6e53f5d4..8103a16b3 100644 --- a/include/matx/core/tensor_impl.h +++ b/include/matx/core/tensor_impl.h @@ -71,7 +71,6 @@ class tensor_impl_t { public: // Type specifier for reflection on class using type = T; // TODO is this necessary - using scalar_type = T; using value_type = T; using tensor_view = bool; using desc_type = Desc; diff --git a/include/matx/core/tensor_utils.h b/include/matx/core/tensor_utils.h index 97e81728d..e5177df39 100644 --- a/include/matx/core/tensor_utils.h +++ b/include/matx/core/tensor_utils.h @@ -842,7 +842,7 @@ namespace detail { PrintKernel<<<1, 1>>>(op, dims...); } else { - auto tmpv = make_tensor(op.Shape()); + auto tmpv = make_tensor(op.Shape()); (tmpv = op).run(); PrintData(fp, tmpv, dims...); } @@ -911,7 +911,7 @@ void PrintData(FILE* fp, const Op &op, Args... dims) { } } else { - auto tmpv = make_tensor(op.Shape()); + auto tmpv = make_tensor(op.Shape()); (tmpv = op).run(); cudaStreamSynchronize(0); InternalPrint(fp, tmpv, dims...); @@ -962,7 +962,7 @@ void PrintData(FILE* fp, const Op &op, Args... dims) { // PrintKernel<<<1, 1>>>(op, dims...); // } // else { -// auto tmpv = make_tensor(op.Shape()); +// auto tmpv = make_tensor(op.Shape()); // (tmpv = op).run(); // PrintData(tmpv, dims...); // } @@ -1004,7 +1004,7 @@ void fprint(FILE* fp, const Op &op, Args... dims) // print tensor size info first std::string type = (is_tensor_view_v) ? "Tensor" : "Operator"; - fprintf(fp, "%s{%s} Rank: %d, Sizes:[", type.c_str(), detail::GetTensorType().c_str(), op.Rank()); + fprintf(fp, "%s{%s} Rank: %d, Sizes:[", type.c_str(), detail::GetTensorType().c_str(), op.Rank()); for (index_t dimIdx = 0; dimIdx < (op.Rank() ); dimIdx++ ) { @@ -1110,7 +1110,7 @@ void print(const Op &op) template auto OpToTensor(Op &&op, [[maybe_unused]] cudaStream_t stream) { if constexpr (!is_tensor_view_v) { - return make_tensor::scalar_type>(op.Shape(), MATX_ASYNC_DEVICE_MEMORY, stream); + return make_tensor::value_type>(op.Shape(), MATX_ASYNC_DEVICE_MEMORY, stream); } else { return op; } diff --git a/include/matx/core/tie.h b/include/matx/core/tie.h index fdb2d0d90..21f8ad428 100644 --- a/include/matx/core/tie.h +++ b/include/matx/core/tie.h @@ -43,8 +43,7 @@ namespace matx { template struct mtie : public BaseOp>{ using mtie_type = bool; - using scalar_type = void; // Doesn't matter since it's not used - using value_type = void; // Doesn't matter since it's not used + using value_type = void; // Doesn't matter since it's not used using shape_type = index_t; using matxoplvalue = bool; @@ -97,7 +96,7 @@ struct mtie : public BaseOp>{ // Run the PreRun on the inner type to avoid allocation but allow transforms using MatX operators // to do any setup needed if constexpr (sizeof...(Ts) == 2) { - cuda::std::get(ts_).InnerPreRun(NoShape{}, std::forward(ex)); + cuda::std::get(ts_).InnerPreRun(detail::NoShape{}, std::forward(ex)); } cuda::std::get(ts_).Exec(ts_, std::forward(ex)); } diff --git a/include/matx/core/type_utils.h b/include/matx/core/type_utils.h index f9e91e8fb..945ba369b 100644 --- a/include/matx/core/type_utils.h +++ b/include/matx/core/type_utils.h @@ -62,9 +62,11 @@ enum class MemoryLayout { MEMORY_LAYOUT_COL_MAJOR, }; -struct NoShape{}; namespace detail { +struct NoShape{}; +struct EmptyOp{}; + template struct is_noshape : std::integral_constant> {}; }; @@ -688,23 +690,23 @@ template inline constexpr bool is_matx_type_v = detail::is_matx_type::value; namespace detail { -template struct extract_scalar_type_impl { - using scalar_type = T; +template struct extract_value_type_impl { + using value_type = T; }; template -struct extract_scalar_type_impl> { - using scalar_type = typename T::scalar_type; +struct extract_value_type_impl()>> { + using value_type = typename T::value_type; }; } /** - * @brief Extract the scalar_type type + * @brief Extract the value_type type * * @tparam T Type to extract from */ template -using extract_scalar_type_t = typename detail::extract_scalar_type_impl::scalar_type; +using extract_value_type_t = typename detail::extract_value_type_impl::value_type; /** * @brief Promote half precision floating point value to fp32, or leave untouched if not half @@ -779,7 +781,7 @@ struct base_type { template struct base_type>> { - using type = tensor_impl_t; + using type = tensor_impl_t; }; template using base_type_t = typename base_type::type>::type; @@ -903,7 +905,7 @@ __MATX_INLINE__ __MATX_DEVICE__ __MATX_HOST__ auto select_tuple(Tuple&& tuple, s template ) && ...), bool> = true> constexpr bool TensorTypesMatch() { using first_type = cuda::std::tuple_element_t<0, cuda::std::tuple>; - return ((std::is_same_v) && ...); + return ((std::is_same_v) && ...); } struct no_permute_t{}; diff --git a/include/matx/generators/alternate.h b/include/matx/generators/alternate.h index e2db8aa17..2dfde91c4 100644 --- a/include/matx/generators/alternate.h +++ b/include/matx/generators/alternate.h @@ -43,7 +43,7 @@ namespace matx index_t size_; public: - using scalar_type = T; + using value_type = T; using matxop = bool; __MATX_INLINE__ std::string str() const { return "alternate"; } diff --git a/include/matx/generators/bartlett.h b/include/matx/generators/bartlett.h index 23eaef40a..d4ef14da6 100644 --- a/include/matx/generators/bartlett.h +++ b/include/matx/generators/bartlett.h @@ -42,7 +42,7 @@ namespace matx index_t size_; public: - using scalar_type = T; + using value_type = T; __MATX_INLINE__ std::string str() const { return "bartlett"; } inline __MATX_HOST__ __MATX_DEVICE__ Bartlett(index_t size) : size_(size){}; diff --git a/include/matx/generators/blackman.h b/include/matx/generators/blackman.h index bd74e9abe..5c53e6f37 100644 --- a/include/matx/generators/blackman.h +++ b/include/matx/generators/blackman.h @@ -42,7 +42,7 @@ namespace matx index_t size_; public: - using scalar_type = T; + using value_type = T; __MATX_INLINE__ std::string str() const { return "blackman"; } diff --git a/include/matx/generators/chirp.h b/include/matx/generators/chirp.h index 027f5400d..c329f1aa6 100644 --- a/include/matx/generators/chirp.h +++ b/include/matx/generators/chirp.h @@ -48,7 +48,7 @@ namespace matx namespace detail { template class Chirp : public BaseOp> { - using space_type = typename SpaceOp::scalar_type; + using space_type = typename SpaceOp::value_type; private: @@ -59,7 +59,7 @@ namespace matx ChirpMethod method_; public: - using scalar_type = FreqType; + using value_type = FreqType; using matxop = bool; __MATX_INLINE__ std::string str() const { return "chirp"; } @@ -90,7 +90,7 @@ namespace matx template class ComplexChirp : public BaseOp> { - using space_type = typename SpaceOp::scalar_type; + using space_type = typename SpaceOp::value_type; private: @@ -101,7 +101,7 @@ namespace matx ChirpMethod method_; public: - using scalar_type = cuda::std::complex; + using value_type = cuda::std::complex; using matxop = bool; __MATX_INLINE__ std::string str() const { return "cchirp"; } @@ -157,7 +157,7 @@ namespace matx * @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) + inline auto chirp(SpaceOp t, FreqType f0, typename SpaceOp::value_type t1, FreqType f1, ChirpMethod method = ChirpMethod::CHIRP_METHOD_LINEAR) { MATX_ASSERT_STR(method == ChirpMethod::CHIRP_METHOD_LINEAR, matxInvalidType, "Only linear chirps are supported") @@ -188,7 +188,7 @@ namespace matx * @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) + inline auto cchirp(SpaceOp t, FreqType f0, typename SpaceOp::value_type t1, FreqType f1, ChirpMethod method = ChirpMethod::CHIRP_METHOD_LINEAR) { MATX_ASSERT_STR(method == ChirpMethod::CHIRP_METHOD_LINEAR, matxInvalidType, "Only linear chirps are supported") diff --git a/include/matx/generators/diag.h b/include/matx/generators/diag.h index 793dcd90b..20d4bf109 100644 --- a/include/matx/generators/diag.h +++ b/include/matx/generators/diag.h @@ -47,7 +47,7 @@ namespace matx public: // dummy type to signal this is a matxop using matxop = bool; - using scalar_type = T; + using value_type = T; __MATX_INLINE__ std::string str() const { return "diag"; } @@ -106,7 +106,7 @@ namespace matx template , bool> = true> inline auto diag(T val) { - return detail::Diag(NoShape{}, T(val)); + return detail::Diag(detail::NoShape{}, T(val)); } /** @@ -197,6 +197,6 @@ namespace matx template inline auto eye() { - return detail::Diag(NoShape{}, T(1)); + return detail::Diag(detail::NoShape{}, T(1)); } } // end namespace matx diff --git a/include/matx/generators/fftfreq.h b/include/matx/generators/fftfreq.h index 9fb604a89..5e4f711d2 100644 --- a/include/matx/generators/fftfreq.h +++ b/include/matx/generators/fftfreq.h @@ -43,7 +43,7 @@ namespace matx float d_; public: - using scalar_type = T; + using value_type = T; using matxop = bool; __MATX_INLINE__ std::string str() const { return "fftfreq"; } diff --git a/include/matx/generators/flattop.h b/include/matx/generators/flattop.h index c401ec6cd..f5a1ec3dc 100644 --- a/include/matx/generators/flattop.h +++ b/include/matx/generators/flattop.h @@ -49,7 +49,7 @@ namespace matx public: - using scalar_type = T; + using value_type = T; __MATX_INLINE__ std::string str() const { return "flattop"; } diff --git a/include/matx/generators/generator1d.h b/include/matx/generators/generator1d.h index 9dc149678..cdeca6282 100644 --- a/include/matx/generators/generator1d.h +++ b/include/matx/generators/generator1d.h @@ -41,7 +41,7 @@ namespace matx public: // dummy type to signal this is a matxop using matxop = bool; - using scalar_type = typename Generator1D::scalar_type; + using value_type = typename Generator1D::value_type; static constexpr int RANK = cuda::std::tuple_size>::value; __MATX_INLINE__ std::string str() const { return "gen1d"; } diff --git a/include/matx/generators/hamming.h b/include/matx/generators/hamming.h index 867c5a689..73dd19979 100644 --- a/include/matx/generators/hamming.h +++ b/include/matx/generators/hamming.h @@ -42,7 +42,7 @@ namespace matx index_t size_; public: - using scalar_type = T; + using value_type = T; __MATX_INLINE__ std::string str() const { return "hamming"; } diff --git a/include/matx/generators/hanning.h b/include/matx/generators/hanning.h index 62ab94aff..3230cccb8 100644 --- a/include/matx/generators/hanning.h +++ b/include/matx/generators/hanning.h @@ -42,7 +42,7 @@ namespace matx index_t size_; public: - using scalar_type = T; + using value_type = T; __MATX_INLINE__ std::string str() const { return "hanning"; } diff --git a/include/matx/generators/linspace.h b/include/matx/generators/linspace.h index 1c8d2e114..0127e0bcf 100644 --- a/include/matx/generators/linspace.h +++ b/include/matx/generators/linspace.h @@ -42,7 +42,7 @@ namespace matx Range range_; public: - using scalar_type = T; + using value_type = T; using matxop = bool; __MATX_INLINE__ std::string str() const { return "linspace"; } diff --git a/include/matx/generators/logspace.h b/include/matx/generators/logspace.h index 965c3ee89..65ff0f348 100644 --- a/include/matx/generators/logspace.h +++ b/include/matx/generators/logspace.h @@ -42,7 +42,7 @@ namespace matx Range range_; public: - using scalar_type = T; + using value_type = T; __MATX_INLINE__ std::string str() const { return "logspace"; } diff --git a/include/matx/generators/meshgrid.h b/include/matx/generators/meshgrid.h index 153976298..882ac2bd5 100644 --- a/include/matx/generators/meshgrid.h +++ b/include/matx/generators/meshgrid.h @@ -47,8 +47,8 @@ namespace matx public: using matxop = bool; - using scalar_type = typename T1::scalar_type; - //typedef typename T1::scalar_type scalar_type; + using value_type = typename T1::value_type; + //typedef typename T1::value_type value_type; __MATX_INLINE__ std::string str() const { return "meshgrid"; } diff --git a/include/matx/generators/ones.h b/include/matx/generators/ones.h index e4ac43f40..2ac6b109a 100644 --- a/include/matx/generators/ones.h +++ b/include/matx/generators/ones.h @@ -88,7 +88,7 @@ namespace matx template inline auto ones() { - return ones(NoShape{}); + return ones(detail::NoShape{}); } } // end namespace matx diff --git a/include/matx/generators/random.h b/include/matx/generators/random.h index c19f6cfea..4aaf7df35 100644 --- a/include/matx/generators/random.h +++ b/include/matx/generators/random.h @@ -255,7 +255,7 @@ template class randomTensorView_t { public: using type = T; ///< Type trait to get type - using scalar_type = T; ///< Type trait to get type + using value_type = T; ///< Type trait to get type // dummy type to signal this is a matxop using matxop = bool; ///< Type trait to indicate this is an operator @@ -398,7 +398,7 @@ template class randomTensorView_t { public: - using scalar_type = T; + using value_type = T; using matxop = bool; static_assert(std::is_same_v || diff --git a/include/matx/generators/range.h b/include/matx/generators/range.h index d189d53ef..ceeafb5bc 100644 --- a/include/matx/generators/range.h +++ b/include/matx/generators/range.h @@ -43,7 +43,7 @@ namespace matx T step_; public: - using scalar_type = T; + using value_type = T; Range() = default; diff --git a/include/matx/generators/zeros.h b/include/matx/generators/zeros.h index 76752c3af..7a2dd9a60 100644 --- a/include/matx/generators/zeros.h +++ b/include/matx/generators/zeros.h @@ -87,6 +87,6 @@ namespace matx template inline auto zeros() { - return zeros(NoShape{}); + return zeros(detail::NoShape{}); } } // end namespace matx diff --git a/include/matx/kernels/channelize_poly.cuh b/include/matx/kernels/channelize_poly.cuh index bc787a8c7..455205ee6 100644 --- a/include/matx/kernels/channelize_poly.cuh +++ b/include/matx/kernels/channelize_poly.cuh @@ -54,9 +54,9 @@ template __launch_bounds__(THREADS) __global__ void ChannelizePoly1D(OutType output, InType input, FilterType filter) { - using output_t = typename OutType::scalar_type; - using input_t = typename InType::scalar_type; - using filter_t = typename FilterType::scalar_type; + using output_t = typename OutType::value_type; + using input_t = typename InType::value_type; + using filter_t = typename FilterType::value_type; // Opportunistically store the filter taps in shared memory if the static shared memory // size is sufficient. Otherwise, we will read directly from global memory on use. @@ -183,9 +183,9 @@ __global__ void ChannelizePoly1D(OutType output, InType input, FilterType filter template __global__ void ChannelizePoly1D_Smem(OutType output, InType input, FilterType filter, index_t elems_per_channel_per_cta) { - using output_t = typename OutType::scalar_type; - using input_t = typename InType::scalar_type; - using filter_t = typename FilterType::scalar_type; + using output_t = typename OutType::value_type; + using input_t = typename InType::value_type; + using filter_t = typename FilterType::value_type; extern __shared__ uint8_t __attribute((aligned(16))) smem_dyn_align16[]; @@ -324,9 +324,9 @@ template (smem_filter); @@ -232,7 +232,7 @@ __global__ void ResamplePoly1D_PhaseBlock(OutType output, InType input, FilterTy } template -__device__ inline void ResamplePoly1D_LoadFilter(typename FilterType::scalar_type *s_filter, const FilterType &filter) +__device__ inline void ResamplePoly1D_LoadFilter(typename FilterType::value_type *s_filter, const FilterType &filter) { const index_t filter_len = filter.Size(0); const int tid = threadIdx.x; @@ -241,7 +241,7 @@ __device__ inline void ResamplePoly1D_LoadFilter(typename FilterType::scalar_typ s_filter[t+1] = filter.operator()(t); } if (tid == 0) { - s_filter[0] = static_cast(0); + s_filter[0] = static_cast(0); } } else { for (int t = tid; t < filter_len; t += THREADS) { @@ -256,9 +256,9 @@ __launch_bounds__(MATX_RESAMPLE_POLY_MAX_NUM_THREADS) __global__ void ResamplePoly1D_ElemBlock(OutType output, InType input, FilterType filter, index_t up, index_t down, index_t elems_per_thread) { - using output_t = typename OutType::scalar_type; - using input_t = typename InType::scalar_type; - using filter_t = typename FilterType::scalar_type; + using output_t = typename OutType::value_type; + using input_t = typename InType::value_type; + using filter_t = typename FilterType::value_type; extern __shared__ uint8_t smem_filter[]; filter_t *s_filter = reinterpret_cast(smem_filter); @@ -368,9 +368,9 @@ __launch_bounds__(MATX_RESAMPLE_POLY_MAX_NUM_THREADS) __global__ void ResamplePoly1D_WarpCentric(OutType output, InType input, FilterType filter, index_t up, index_t down, index_t elems_per_warp) { - using output_t = typename OutType::scalar_type; - using input_t = typename InType::scalar_type; - using filter_t = typename FilterType::scalar_type; + using output_t = typename OutType::value_type; + using input_t = typename InType::value_type; + using filter_t = typename FilterType::value_type; auto block = cg::this_thread_block(); auto tile = cg::tiled_partition(block); diff --git a/include/matx/kernels/transpose.cuh b/include/matx/kernels/transpose.cuh index 7a37c460b..f1a5adb55 100644 --- a/include/matx/kernels/transpose.cuh +++ b/include/matx/kernels/transpose.cuh @@ -24,7 +24,7 @@ template __global__ void transpose_kernel_oop(OutputTensor out, const InputTensor in) { - using T = typename OutputTensor::scalar_type; + using T = typename OutputTensor::value_type; constexpr int RANK = OutputTensor::Rank(); extern __shared__ float diff --git a/include/matx/operators/all.h b/include/matx/operators/all.h index 0407f53f1..9eed9c6b7 100644 --- a/include/matx/operators/all.h +++ b/include/matx/operators/all.h @@ -49,12 +49,12 @@ namespace detail { private: OpA a_; cuda::std::array out_dims_; - mutable detail::tensor_impl_t::scalar_type, ORank> tmp_out_; - mutable typename remove_cvref_t::scalar_type *ptr; + mutable detail::tensor_impl_t::value_type, ORank> tmp_out_; + mutable typename remove_cvref_t::value_type *ptr; public: using matxop = bool; - using scalar_type = typename remove_cvref_t::scalar_type; + using value_type = typename remove_cvref_t::value_type; using matx_transform_op = bool; using all_xform_op = bool; diff --git a/include/matx/operators/ambgfun.h b/include/matx/operators/ambgfun.h index 551ab3298..d3543b109 100644 --- a/include/matx/operators/ambgfun.h +++ b/include/matx/operators/ambgfun.h @@ -50,12 +50,12 @@ namespace matx AMBGFunCutType_t cut_; float cut_val_; cuda::std::array out_dims_; - mutable detail::tensor_impl_t::scalar_type, 2> tmp_out_; - mutable typename remove_cvref_t::scalar_type *ptr; + mutable detail::tensor_impl_t::value_type, 2> tmp_out_; + mutable typename remove_cvref_t::value_type *ptr; public: using matxop = bool; - using scalar_type = typename OpX::scalar_type; + using value_type = typename OpX::value_type; using matx_transform_op = bool; using ambgfun_xform_op = bool; diff --git a/include/matx/operators/any.h b/include/matx/operators/any.h index a2600346d..2d34f50c6 100644 --- a/include/matx/operators/any.h +++ b/include/matx/operators/any.h @@ -49,12 +49,12 @@ namespace detail { private: OpA a_; cuda::std::array out_dims_; - mutable detail::tensor_impl_t::scalar_type, ORank> tmp_out_; - mutable typename remove_cvref_t::scalar_type *ptr; + mutable detail::tensor_impl_t::value_type, ORank> tmp_out_; + mutable typename remove_cvref_t::value_type *ptr; public: using matxop = bool; - using scalar_type = typename remove_cvref_t::scalar_type; + using value_type = typename remove_cvref_t::value_type; using matx_transform_op = bool; using any_xform_op = bool; diff --git a/include/matx/operators/argmax.h b/include/matx/operators/argmax.h index 2e2b6acc9..5e57ae8a5 100644 --- a/include/matx/operators/argmax.h +++ b/include/matx/operators/argmax.h @@ -51,7 +51,7 @@ namespace detail { public: using matxop = bool; - using scalar_type = typename remove_cvref_t::scalar_type; + using value_type = typename remove_cvref_t::value_type; using matx_transform_op = bool; using argmax_xform_op = bool; diff --git a/include/matx/operators/argmin.h b/include/matx/operators/argmin.h index a96cde932..160f040fe 100644 --- a/include/matx/operators/argmin.h +++ b/include/matx/operators/argmin.h @@ -51,7 +51,7 @@ namespace detail { public: using matxop = bool; - using scalar_type = typename remove_cvref_t::scalar_type; + using value_type = typename remove_cvref_t::value_type; using matx_transform_op = bool; using argmin_xform_op = bool; diff --git a/include/matx/operators/at.h b/include/matx/operators/at.h index 1305c4718..7f34512e6 100644 --- a/include/matx/operators/at.h +++ b/include/matx/operators/at.h @@ -50,7 +50,7 @@ namespace matx public: using matxop = bool; - using scalar_type = typename Op::scalar_type; + using value_type = typename Op::value_type; __MATX_INLINE__ std::string str() const { return "at()"; } __MATX_INLINE__ AtOp(Op op, Is... is) : op_(op), idx_{is...} {}; diff --git a/include/matx/operators/binary_operators.h b/include/matx/operators/binary_operators.h index 79f7dcde6..87fd66daa 100644 --- a/include/matx/operators/binary_operators.h +++ b/include/matx/operators/binary_operators.h @@ -42,8 +42,8 @@ is_matx_op()>> \ [[nodiscard]] __MATX_INLINE__ auto FUNCTION(I1 i1, I2 i2) \ { \ - using I1Type = extract_scalar_type_t; \ - using I2Type = extract_scalar_type_t; \ + using I1Type = extract_value_type_t; \ + using I2Type = extract_value_type_t; \ using Op = TENSOR_OP; \ const typename detail::base_type::type &base1 = i1; \ const typename detail::base_type::type &base2 = i2; \ @@ -101,7 +101,7 @@ namespace matx public: // dummy type to signal this is a matxop using matxop = bool; - using scalar_type = typename Op::scalar_type; + using value_type = typename Op::value_type; using self_type = matxBinaryOp; __MATX_INLINE__ const std::string str() const { diff --git a/include/matx/operators/cart2sph.h b/include/matx/operators/cart2sph.h index ad2f443f3..c898a53f9 100644 --- a/include/matx/operators/cart2sph.h +++ b/include/matx/operators/cart2sph.h @@ -52,7 +52,7 @@ namespace matx public: using matxop = bool; - using scalar_type = typename T1::scalar_type; + using value_type = typename T1::value_type; __MATX_INLINE__ std::string str() const { return "cart2sph(" + get_type_str(x_) + "," + get_type_str(y_) + "," + get_type_str(z_) + ")"; } diff --git a/include/matx/operators/cast.h b/include/matx/operators/cast.h index f0c5c2c31..ee6ea3d8d 100644 --- a/include/matx/operators/cast.h +++ b/include/matx/operators/cast.h @@ -66,7 +66,7 @@ namespace matx public: using matxop = bool; - using scalar_type = NewType; + using value_type = NewType; __MATX_INLINE__ std::string str() const { return as_type_str() + "(" + op_.str() + ")"; } __MATX_INLINE__ CastOp(T op) : op_(op){}; @@ -118,7 +118,7 @@ namespace matx public: using matxop = bool; - using scalar_type = NewType; + using value_type = NewType; static_assert(!is_complex_v && !is_complex_half_v, "T1 input operator cannot be complex"); static_assert(!is_complex_v && !is_complex_half_v, "T2 input operator cannot be complex"); static_assert(is_complex_v || is_complex_half_v, "ComplexCastOp output type should be complex"); @@ -192,7 +192,7 @@ namespace matx template auto __MATX_INLINE__ as_type(T t) { - if constexpr (std::is_same_v) { + if constexpr (std::is_same_v) { // optimized path when type is the same to avoid creating unecessary operators return t; } else { diff --git a/include/matx/operators/cgsolve.h b/include/matx/operators/cgsolve.h index fd7ba3cc7..63cb53156 100644 --- a/include/matx/operators/cgsolve.h +++ b/include/matx/operators/cgsolve.h @@ -49,12 +49,12 @@ namespace matx double tol_; int max_iters_; cuda::std::array out_dims_; - mutable detail::tensor_impl_t tmp_out_; - mutable typename OpA::scalar_type *ptr; + mutable detail::tensor_impl_t tmp_out_; + mutable typename OpA::value_type *ptr; public: using matxop = bool; - using scalar_type = typename OpA::scalar_type; + using value_type = typename OpA::value_type; using matx_transform_op = bool; using cgsolve_xform_op = bool; diff --git a/include/matx/operators/channelize_poly.h b/include/matx/operators/channelize_poly.h index d9e81301d..d27b033d5 100644 --- a/include/matx/operators/channelize_poly.h +++ b/include/matx/operators/channelize_poly.h @@ -49,7 +49,7 @@ namespace detail { // type cuda::std::complex and the filter has type double, out_t // will be cuda::std::complex. using out_t = cuda::std::common_type_t< - complex_from_scalar_t, complex_from_scalar_t>; + complex_from_scalar_t, complex_from_scalar_t>; OpA a_; FilterType f_; index_t num_channels_; @@ -62,7 +62,7 @@ namespace detail { using matxop = bool; using matx_transform_op = bool; using channelize_poly_xform_op = bool; - using scalar_type = out_t; + using value_type = out_t; __MATX_INLINE__ std::string str() const { return "channelize_poly(" + get_type_str(a_) + "," + get_type_str(f_) + ")";} __MATX_INLINE__ ChannelizePolyOp(OpA a, const FilterType &f, index_t num_channels, index_t decimation_factor) : diff --git a/include/matx/operators/chol.h b/include/matx/operators/chol.h index 82efde6c2..331a03cec 100644 --- a/include/matx/operators/chol.h +++ b/include/matx/operators/chol.h @@ -45,11 +45,11 @@ namespace detail { private: OpA a_; cublasFillMode_t uplo_; - mutable matx::tensor_t tmp_out_; + mutable matx::tensor_t tmp_out_; public: using matxop = bool; - using scalar_type = typename OpA::scalar_type; + using value_type = typename OpA::value_type; using matx_transform_op = bool; using chol_xform_op = bool; diff --git a/include/matx/operators/clone.h b/include/matx/operators/clone.h index a8f46ff04..b3c832d05 100644 --- a/include/matx/operators/clone.h +++ b/include/matx/operators/clone.h @@ -49,7 +49,7 @@ namespace matx public: using matxop = bool; - using scalar_type = typename T::scalar_type; + using value_type = typename T::value_type; __MATX_INLINE__ std::string str() const { return "clone(" + op_.str() + ")"; } diff --git a/include/matx/operators/collapse.h b/include/matx/operators/collapse.h index a25c39fe9..a84768136 100644 --- a/include/matx/operators/collapse.h +++ b/include/matx/operators/collapse.h @@ -48,7 +48,7 @@ namespace matx public: using matxop = bool; - using scalar_type = typename T1::scalar_type; + using value_type = typename T1::value_type; using shape_type = index_t; using matxoplvalue = bool; using self_type = LCollapseOp; @@ -203,7 +203,7 @@ namespace matx public: using matxop = bool; - using scalar_type = typename T1::scalar_type; + using value_type = typename T1::value_type; using shape_type = index_t; using matxlvalue = bool; using self_type = RCollapseOp; diff --git a/include/matx/operators/concat.h b/include/matx/operators/concat.h index 76e442d3f..4fc33fd76 100644 --- a/include/matx/operators/concat.h +++ b/include/matx/operators/concat.h @@ -50,7 +50,7 @@ namespace matx class ConcatOp : public BaseOp> { using first_type = cuda::std::tuple_element_t<0, cuda::std::tuple>; - using first_value_type = typename first_type::scalar_type; + using first_value_type = typename first_type::value_type; using self_type = ConcatOp; static constexpr int RANK = first_type::Rank(); @@ -61,7 +61,7 @@ namespace matx using shape_type = index_t; // Scalar type of operation - using scalar_type = first_value_type; + using value_type = first_value_type; template __MATX_INLINE__ std::string get_str() const { @@ -96,7 +96,7 @@ namespace matx if constexpr ( I == N ) { // This should never happen - return scalar_type(-9999); + return value_type(-9999); // returning this to satisfy lvalue requirements } else { const auto &op = cuda::std::get(ops_); diff --git a/include/matx/operators/constval.h b/include/matx/operators/constval.h index ca033f5f3..3df28affc 100644 --- a/include/matx/operators/constval.h +++ b/include/matx/operators/constval.h @@ -46,7 +46,7 @@ namespace matx public: // dummy type to signal this is a matxop using matxop = bool; - using scalar_type = T; + using value_type = T; __MATX_INLINE__ std::string str() const { return "constval"; } ConstVal(ShapeType &&s, T val) : s_(std::forward(s)), v_(val){}; diff --git a/include/matx/operators/conv.h b/include/matx/operators/conv.h index 0536556ac..b4de3befc 100644 --- a/include/matx/operators/conv.h +++ b/include/matx/operators/conv.h @@ -44,8 +44,8 @@ namespace matx class Conv1DOp : public BaseOp> { private: - using out_t = std::conditional_t, - typename OpA::scalar_type, typename OpB::scalar_type>; + using out_t = std::conditional_t, + typename OpA::value_type, typename OpB::value_type>; constexpr static int max_rank = cuda::std::max(OpA::Rank(), OpB::Rank()); OpA a_; OpB b_; @@ -60,7 +60,7 @@ namespace matx public: using matxop = bool; - using scalar_type = out_t; + using value_type = out_t; using matx_transform_op = bool; using conv_xform_op = bool; @@ -71,7 +71,7 @@ namespace matx __MATX_INLINE__ Conv1DOp(const OpA &A, const OpB &B, matxConvCorrMode_t mode, matxConvCorrMethod_t method, PermDims perm) : a_(A), b_(B), mode_(mode), method_(method), perm_(perm) { - MATX_ASSERT_STR((!is_matx_type_v && !is_matx_type_v) || + MATX_ASSERT_STR((!is_matx_type_v && !is_matx_type_v) || method == MATX_C_METHOD_DIRECT, matxInvalidType, "FFT convolutions do not support half precision float currently"); @@ -234,8 +234,8 @@ namespace detail { class Conv2DOp : public BaseOp> { private: - using out_t = std::conditional_t, - typename OpA::scalar_type, typename OpB::scalar_type>; + using out_t = std::conditional_t, + typename OpA::value_type, typename OpB::value_type>; constexpr static int max_rank = cuda::std::max(OpA::Rank(), OpB::Rank()); OpA a_; OpB b_; @@ -247,7 +247,7 @@ namespace detail { public: using matxop = bool; - using scalar_type = out_t; + using value_type = out_t; using matx_transform_op = bool; using conv_xform_op = bool; diff --git a/include/matx/operators/corr.h b/include/matx/operators/corr.h index 1847fc759..52762ba7e 100644 --- a/include/matx/operators/corr.h +++ b/include/matx/operators/corr.h @@ -44,8 +44,8 @@ namespace matx class CorrOp : public BaseOp> { private: - using out_t = std::conditional_t, - typename OpA::scalar_type, typename OpB::scalar_type>; + using out_t = std::conditional_t, + typename OpA::value_type, typename OpB::value_type>; constexpr static int max_rank = cuda::std::max(OpA::Rank(), OpB::Rank()); OpA a_; OpB b_; @@ -58,7 +58,7 @@ namespace matx public: using matxop = bool; - using scalar_type = out_t; + using value_type = out_t; using matx_transform_op = bool; using conv_xform_op = bool; diff --git a/include/matx/operators/cov.h b/include/matx/operators/cov.h index d6ccfac26..48b502081 100644 --- a/include/matx/operators/cov.h +++ b/include/matx/operators/cov.h @@ -46,12 +46,12 @@ namespace matx private: OpA a_; cuda::std::array out_dims_; - mutable detail::tensor_impl_t::scalar_type, OpA::Rank()> tmp_out_; - mutable typename remove_cvref_t::scalar_type *ptr; + mutable detail::tensor_impl_t::value_type, OpA::Rank()> tmp_out_; + mutable typename remove_cvref_t::value_type *ptr; public: using matxop = bool; - using scalar_type = typename OpA::scalar_type; + using value_type = typename OpA::value_type; using matx_transform_op = bool; using cov_xform_op = bool; diff --git a/include/matx/operators/cumsum.h b/include/matx/operators/cumsum.h index a872e7630..c050b7e83 100644 --- a/include/matx/operators/cumsum.h +++ b/include/matx/operators/cumsum.h @@ -48,12 +48,12 @@ namespace detail { private: OpA a_; cuda::std::array out_dims_; - mutable detail::tensor_impl_t::scalar_type, OpA::Rank()> tmp_out_; - mutable typename remove_cvref_t::scalar_type *ptr; + mutable detail::tensor_impl_t::value_type, OpA::Rank()> tmp_out_; + mutable typename remove_cvref_t::value_type *ptr; public: using matxop = bool; - using scalar_type = typename OpA::scalar_type; + using value_type = typename OpA::value_type; using matx_transform_op = bool; using cumsum_xform_op = bool; diff --git a/include/matx/operators/dct.h b/include/matx/operators/dct.h index 033100404..5e7f263dc 100644 --- a/include/matx/operators/dct.h +++ b/include/matx/operators/dct.h @@ -101,7 +101,7 @@ void dct(OutputTensor &out, const InputTensor &in, MATX_STATIC_ASSERT(OutputTensor::Rank() == 1, matxInvalidDim); index_t N = in.Size(OutputTensor::Rank() - 1); - tensor_t, 1> tmp{{N + 1}}; + tensor_t, 1> tmp{{N + 1}}; fft_impl(tmp, in, 0, FFTNorm::BACKWARD, stream); auto s = tmp.Slice({0}, {N}); diff --git a/include/matx/operators/det.h b/include/matx/operators/det.h index a74411746..78a466858 100644 --- a/include/matx/operators/det.h +++ b/include/matx/operators/det.h @@ -44,12 +44,12 @@ namespace detail { { private: OpA a_; - mutable detail::tensor_impl_t::scalar_type, OpA::Rank()> tmp_out_; - mutable typename remove_cvref_t::scalar_type *ptr; + mutable detail::tensor_impl_t::value_type, OpA::Rank()> tmp_out_; + mutable typename remove_cvref_t::value_type *ptr; public: using matxop = bool; - using scalar_type = typename OpA::scalar_type; + using value_type = typename OpA::value_type; using matx_transform_op = bool; using det_xform_op = bool; diff --git a/include/matx/operators/diag.h b/include/matx/operators/diag.h index 74b9e53c0..4667b7d8a 100644 --- a/include/matx/operators/diag.h +++ b/include/matx/operators/diag.h @@ -55,7 +55,7 @@ namespace matx public: using matxop = bool; - using scalar_type = typename T1::scalar_type; + using value_type = typename T1::value_type; __MATX_INLINE__ std::string str() const { return "diag(" + op_.str() + ")"; } diff --git a/include/matx/operators/eig.h b/include/matx/operators/eig.h index 7cf9cf0f4..676275aa7 100644 --- a/include/matx/operators/eig.h +++ b/include/matx/operators/eig.h @@ -52,7 +52,7 @@ namespace detail { public: using matxop = bool; - using scalar_type = typename OpA::scalar_type; + using value_type = typename OpA::value_type; using matx_transform_op = bool; using eig_xform_op = bool; diff --git a/include/matx/operators/einsum.h b/include/matx/operators/einsum.h index 982fcfdb4..3b35f5779 100644 --- a/include/matx/operators/einsum.h +++ b/include/matx/operators/einsum.h @@ -53,7 +53,7 @@ namespace detail { public: using matxop = bool; - using scalar_type = void; + using value_type = void; using matx_transform_op = bool; using einsum_xform_op = bool; diff --git a/include/matx/operators/fft.h b/include/matx/operators/fft.h index 89cd0cbc7..8eaeef78c 100644 --- a/include/matx/operators/fft.h +++ b/include/matx/operators/fft.h @@ -53,18 +53,18 @@ namespace matx FFTType type_; FFTNorm norm_; cuda::std::array out_dims_; - using ttype = std::conditional_t, - typename OpA::scalar_type, - typename scalar_to_complex::ctype>; + using ttype = std::conditional_t, + typename OpA::value_type, + typename scalar_to_complex::ctype>; // This should be tensor_impl_t, but need to work around issues with temp types returned in fft mutable matx::tensor_t tmp_out_; mutable ttype *ptr; public: using matxop = bool; - using scalar_type = std::conditional_t, - typename OpA::scalar_type, - typename scalar_to_complex::ctype>; + using value_type = std::conditional_t, + typename OpA::value_type, + typename scalar_to_complex::ctype>; using matx_transform_op = bool; using fft_xform_op = bool; @@ -84,7 +84,7 @@ namespace matx } if (fft_size_ != 0) { - if constexpr (is_complex_v) { + if constexpr (is_complex_v) { if constexpr (std::is_same_v) { out_dims_[Rank() - 1] = fft_size_; } @@ -102,7 +102,7 @@ namespace matx } } else { - if constexpr (!is_complex_v) { // C2C uses the same input/output size. R2C is N/2+1 + if constexpr (!is_complex_v) { // C2C uses the same input/output size. R2C is N/2+1 if constexpr (!std::is_same_v) { out_dims_[perm_[Rank()-1]] = out_dims_[perm_[Rank()-1]] / 2 + 1; } @@ -303,16 +303,16 @@ namespace matx FFTType type_; FFTNorm norm_; cuda::std::array out_dims_; - using ttype = std::conditional_t, - typename OpA::scalar_type, - typename scalar_to_complex::ctype>; + using ttype = std::conditional_t, + typename OpA::value_type, + typename scalar_to_complex::ctype>; // This should be tensor_impl_t, but need to work around issues with temp types returned in fft mutable matx::tensor_t tmp_out_; mutable ttype *ptr; public: using matxop = bool; - using scalar_type = typename OpA::scalar_type; + using value_type = typename OpA::value_type; using matx_transform_op = bool; using fft2_xform_op = bool; @@ -330,7 +330,7 @@ namespace matx out_dims_[r] = a_.Size(r); } - if constexpr (!is_complex_v) { // C2C uses the same input/output size. R2C is N/2+1 + if constexpr (!is_complex_v) { // C2C uses the same input/output size. R2C is N/2+1 if constexpr (!std::is_same_v) { out_dims_[perm_[0]] = out_dims_[perm_[0]] / 2 + 1; out_dims_[perm_[1]] = out_dims_[perm_[1]] / 2 + 1; diff --git a/include/matx/operators/fftshift.h b/include/matx/operators/fftshift.h index b111ebffc..91c0592b1 100644 --- a/include/matx/operators/fftshift.h +++ b/include/matx/operators/fftshift.h @@ -47,7 +47,7 @@ namespace matx public: using matxop = bool; - using scalar_type = typename T1::scalar_type; + using value_type = typename T1::value_type; __MATX_INLINE__ std::string str() const { return "fftshift(" + op_.str() + ")"; } @@ -128,7 +128,7 @@ namespace matx public: using matxop = bool; - using scalar_type = typename T1::scalar_type; + using value_type = typename T1::value_type; __MATX_INLINE__ FFTShift2DOp(T1 op) : op_(op){ static_assert(Rank() >= 2, "2D FFT shift must have a rank 2 operator or higher"); @@ -199,7 +199,7 @@ namespace matx public: using matxop = bool; - using scalar_type = typename T1::scalar_type; + using value_type = typename T1::value_type; __MATX_INLINE__ IFFTShift1DOp(T1 op) : op_(op) { static_assert(Rank() >= 1, "1D IFFT shift must have a rank 1 operator or higher"); @@ -269,7 +269,7 @@ namespace matx public: using matxop = bool; - using scalar_type = typename T1::scalar_type; + using value_type = typename T1::value_type; __MATX_INLINE__ IFFTShift2DOp(T1 op) : op_(op) { static_assert(Rank() >= 2, "2D IFFT shift must have a rank 2 operator or higher"); diff --git a/include/matx/operators/filter.h b/include/matx/operators/filter.h index e8c43728f..accf37e26 100644 --- a/include/matx/operators/filter.h +++ b/include/matx/operators/filter.h @@ -50,12 +50,12 @@ namespace detail { cuda::std::array h_rec_; cuda::std::array h_nonrec_; cuda::std::array out_dims_; - mutable detail::tensor_impl_t::scalar_type, OpA::Rank()> tmp_out_; - mutable typename remove_cvref_t::scalar_type *ptr; + mutable detail::tensor_impl_t::value_type, OpA::Rank()> tmp_out_; + mutable typename remove_cvref_t::value_type *ptr; public: using matxop = bool; - using scalar_type = void; + using value_type = void; using matx_transform_op = bool; using filter_xform_op = bool; diff --git a/include/matx/operators/find.h b/include/matx/operators/find.h index ac8147112..3c92c848d 100644 --- a/include/matx/operators/find.h +++ b/include/matx/operators/find.h @@ -51,7 +51,7 @@ namespace detail { public: using matxop = bool; - using scalar_type = typename OpA::scalar_type; + using value_type = typename OpA::value_type; using matx_transform_op = bool; using find_xform_op = bool; diff --git a/include/matx/operators/find_idx.h b/include/matx/operators/find_idx.h index f690b1042..ac0365e78 100644 --- a/include/matx/operators/find_idx.h +++ b/include/matx/operators/find_idx.h @@ -51,7 +51,7 @@ namespace detail { public: using matxop = bool; - using scalar_type = typename OpA::scalar_type; + using value_type = typename OpA::value_type; using matx_transform_op = bool; using find_idx_xform_op = bool; diff --git a/include/matx/operators/flatten.h b/include/matx/operators/flatten.h index 646bf0d64..b16a2a15b 100644 --- a/include/matx/operators/flatten.h +++ b/include/matx/operators/flatten.h @@ -48,7 +48,7 @@ namespace matx public: using matxop = bool; - using scalar_type = typename T1::scalar_type; + using value_type = typename T1::value_type; __MATX_INLINE__ std::string str() const { return "flatten(" + op1_.str() + ")"; } diff --git a/include/matx/operators/frexp.h b/include/matx/operators/frexp.h index 6de5b94fb..7e3f5ab44 100644 --- a/include/matx/operators/frexp.h +++ b/include/matx/operators/frexp.h @@ -48,12 +48,12 @@ namespace detail { public: using matxop = bool; - using scalar_type = typename OpA::scalar_type; + using value_type = typename OpA::value_type; __MATX_INLINE__ std::string str() const { return "frexp()"; } __MATX_INLINE__ FrexpOp(OpA a) : a_(a) { - static_assert(std::is_floating_point_v || - is_cuda_complex_v, "frexp() must take a floating point input"); + static_assert(std::is_floating_point_v || + is_cuda_complex_v, "frexp() must take a floating point input"); }; @@ -61,8 +61,8 @@ namespace detail { __MATX_INLINE__ __MATX_DEVICE__ __MATX_HOST__ auto operator()(Is... indices) const { [[maybe_unused]] int rexp; - if constexpr (is_cuda_complex_v) { - if constexpr (std::is_same_v) { + if constexpr (is_cuda_complex_v) { + if constexpr (std::is_same_v) { if constexpr (WHICH == 0) { // real fractional const auto frac = cuda::std::frexpf(a_(indices...).real(), &rexp); return frac; @@ -94,7 +94,7 @@ namespace detail { } } else { - if constexpr (std::is_same_v) { + if constexpr (std::is_same_v) { [[maybe_unused]] const float frac = cuda::std::frexpf(a_(indices...), &rexp); if constexpr (WHICH == 0) { // fractional return frac; diff --git a/include/matx/operators/hermitian.h b/include/matx/operators/hermitian.h index 2317db9ed..4a1090d8f 100644 --- a/include/matx/operators/hermitian.h +++ b/include/matx/operators/hermitian.h @@ -53,7 +53,7 @@ namespace matx public: using matxop = bool; - using scalar_type = typename T1::scalar_type; + using value_type = typename T1::value_type; __MATX_INLINE__ std::string str() const { return "hermitian(" + op_.str() + ")"; } __MATX_INLINE__ HermitianTransOp(T1 op) : op_(op) { diff --git a/include/matx/operators/hist.h b/include/matx/operators/hist.h index 6060b61dc..396da339a 100644 --- a/include/matx/operators/hist.h +++ b/include/matx/operators/hist.h @@ -47,20 +47,20 @@ namespace detail { { private: OpA a_; - typename OpA::scalar_type lower_; - typename OpA::scalar_type upper_; + typename OpA::value_type lower_; + typename OpA::value_type upper_; cuda::std::array out_dims_; mutable detail::tensor_impl_t tmp_out_; mutable int *ptr; public: using matxop = bool; - using scalar_type = typename OpA::scalar_type; + using value_type = typename OpA::value_type; using matx_transform_op = bool; using hist_xform_op = bool; __MATX_INLINE__ std::string str() const { return "hist()"; } - __MATX_INLINE__ HistOp(OpA a, typename OpA::scalar_type lower, typename OpA::scalar_type upper) : a_(a), lower_(lower), upper_(upper) { + __MATX_INLINE__ HistOp(OpA a, typename OpA::value_type lower, typename OpA::value_type upper) : a_(a), lower_(lower), upper_(upper) { for (int r = 0; r < Rank(); r++) { out_dims_[r] = a_.Size(r); } @@ -138,8 +138,8 @@ namespace detail { */ template __MATX_INLINE__ auto hist(const InputOperator &a, - const typename InputOperator::scalar_type lower, - const typename InputOperator::scalar_type upper) { + const typename InputOperator::value_type lower, + const typename InputOperator::value_type upper) { return detail::HistOp(a, lower, upper); } diff --git a/include/matx/operators/if.h b/include/matx/operators/if.h index 531a55868..ac3719158 100644 --- a/include/matx/operators/if.h +++ b/include/matx/operators/if.h @@ -58,7 +58,7 @@ namespace matx cuda::std::array(), detail::get_rank())> size_; public: - using scalar_type = void; ///< Scalar type for type extraction + using value_type = void; ///< Scalar type for type extraction __MATX_INLINE__ std::string str() const { return "if(" + cond_.str() + ") then {" + op_.str() + "}"; } /** diff --git a/include/matx/operators/ifelse.h b/include/matx/operators/ifelse.h index 1aeba1097..c2682e3a6 100644 --- a/include/matx/operators/ifelse.h +++ b/include/matx/operators/ifelse.h @@ -60,7 +60,7 @@ namespace matx cuda::std::array(), detail::get_rank(), detail::get_rank())> size_; public: - using scalar_type = void; ///< Scalar type for type extraction + using value_type = void; ///< Scalar type for type extraction __MATX_INLINE__ std::string str() const { return "if(" + detail::get_type_str(cond_) + ") then {" + detail::get_type_str(op1_) + "} else {" + detail::get_type_str(op2_) + "}"; } diff --git a/include/matx/operators/index.h b/include/matx/operators/index.h index 3863d6954..d6320ebe7 100644 --- a/include/matx/operators/index.h +++ b/include/matx/operators/index.h @@ -50,7 +50,7 @@ namespace matx public: using matxop = bool; - using scalar_type = index_t; + using value_type = index_t; __MATX_INLINE__ std::string str() const { return "index()"; } __MATX_INLINE__ IndexOp(int dim) : dim_(dim){}; diff --git a/include/matx/operators/interleaved.h b/include/matx/operators/interleaved.h index dc0515426..19c964bbc 100644 --- a/include/matx/operators/interleaved.h +++ b/include/matx/operators/interleaved.h @@ -47,15 +47,15 @@ namespace matx public: using matxop = bool; - using scalar_type = typename T1::scalar_type; + using value_type = typename T1::value_type; - using complex_type = std::conditional_t, - matxHalfComplex, - cuda::std::complex>; + using complex_type = std::conditional_t, + matxHalfComplex, + cuda::std::complex>; __MATX_INLINE__ std::string str() const { return "interleaved(" + op_.str() + ")"; } __MATX_INLINE__ ComplexInterleavedOp(T1 op) : op_(op) { - static_assert(!is_complex_v>, "Complex interleaved op only works on scalar input types"); + static_assert(!is_complex_v>, "Complex interleaved op only works on scalar input types"); static_assert(Rank() > 0); }; @@ -139,7 +139,7 @@ namespace matx template auto interleaved(T1 t) { - static_assert(!is_complex_v>, "Input to interleaved operator must be real-valued"); + static_assert(!is_complex_v>, "Input to interleaved operator must be real-valued"); return detail::ComplexInterleavedOp(t); } } // end namespace matx diff --git a/include/matx/operators/inverse.h b/include/matx/operators/inverse.h index 21c3fb56b..69684f6a2 100644 --- a/include/matx/operators/inverse.h +++ b/include/matx/operators/inverse.h @@ -46,12 +46,12 @@ namespace detail { { private: OpA a_; - mutable detail::tensor_impl_t::scalar_type, OpA::Rank()> tmp_out_; - mutable typename remove_cvref_t::scalar_type *ptr; + mutable detail::tensor_impl_t::value_type, OpA::Rank()> tmp_out_; + mutable typename remove_cvref_t::value_type *ptr; public: using matxop = bool; - using scalar_type = typename OpA::scalar_type; + using value_type = typename OpA::value_type; using matx_transform_op = bool; using inv_xform_op = bool; diff --git a/include/matx/operators/isclose.h b/include/matx/operators/isclose.h index 278160f41..9d9df0bc6 100644 --- a/include/matx/operators/isclose.h +++ b/include/matx/operators/isclose.h @@ -46,8 +46,8 @@ namespace matx { public: using matxop = bool; - using scalar_type = typename remove_cvref_t::scalar_type; - using inner_type = typename inner_op_type_t::type; + using value_type = typename remove_cvref_t::value_type; + using inner_type = typename inner_op_type_t::type; __MATX_INLINE__ std::string str() const { return "isclose()"; } diff --git a/include/matx/operators/kronecker.h b/include/matx/operators/kronecker.h index 70ff0dc85..1b5e2f18f 100644 --- a/include/matx/operators/kronecker.h +++ b/include/matx/operators/kronecker.h @@ -55,7 +55,7 @@ namespace matx public: using matxop = bool; - using scalar_type = typename T1::scalar_type; + using value_type = typename T1::value_type; __MATX_INLINE__ std::string str() const { return "kron(" + op1_.str() + "," + op2_.str() + ")"; } diff --git a/include/matx/operators/legendre.h b/include/matx/operators/legendre.h index c1606053a..2d38e53e3 100644 --- a/include/matx/operators/legendre.h +++ b/include/matx/operators/legendre.h @@ -90,7 +90,7 @@ namespace matx public: using matxop = bool; - using scalar_type = typename T3::scalar_type; + using value_type = typename T3::value_type; __MATX_INLINE__ std::string str() const { return "legendre(" + get_type_str(n_) + "," + get_type_str(m_) + "," + get_type_str(in_) + ")"; } @@ -100,7 +100,7 @@ namespace matx } template - __MATX_INLINE__ __MATX_DEVICE__ __MATX_HOST__ scalar_type operator()(Is... indices) const + __MATX_INLINE__ __MATX_DEVICE__ __MATX_HOST__ value_type operator()(Is... indices) const { cuda::std::array inds{indices...}; cuda::std::array xinds{}; @@ -130,13 +130,13 @@ namespace matx auto x = cuda::std::apply(in_, xinds); - scalar_type ret; + value_type ret; // if we are half precision up cast to float - if constexpr (is_complex_half_v) { - ret = static_cast(legendre(n, m, cuda::std::complex(x))); - } else if constexpr (is_matx_half_v) { - ret = static_cast(legendre(n, m, float(x))); + if constexpr (is_complex_half_v) { + ret = static_cast(legendre(n, m, cuda::std::complex(x))); + } else if constexpr (is_matx_half_v) { + ret = static_cast(legendre(n, m, float(x))); } else { ret = legendre(n, m, x); } diff --git a/include/matx/operators/lu.h b/include/matx/operators/lu.h index 16dd5146d..37978aef2 100644 --- a/include/matx/operators/lu.h +++ b/include/matx/operators/lu.h @@ -44,12 +44,12 @@ namespace detail { { private: OpA a_; - mutable detail::tensor_impl_t::scalar_type, OpA::Rank()> tmp_out_; - mutable typename remove_cvref_t::scalar_type *ptr; + mutable detail::tensor_impl_t::value_type, OpA::Rank()> tmp_out_; + mutable typename remove_cvref_t::value_type *ptr; public: using matxop = bool; - using scalar_type = typename OpA::scalar_type; + using value_type = typename OpA::value_type; using matx_transform_op = bool; using lu_xform_op = bool; diff --git a/include/matx/operators/matmul.h b/include/matx/operators/matmul.h index 17dfbda78..7533a056f 100644 --- a/include/matx/operators/matmul.h +++ b/include/matx/operators/matmul.h @@ -55,12 +55,12 @@ namespace matx static constexpr int out_rank = cuda::std::max(OpA::Rank(), OpB::Rank()); cuda::std::array out_dims_; // This should be tensor_impl_t, but need to work around issues with temp types returned in matmul - mutable matx::tensor_t::scalar_type, out_rank> tmp_out_; - mutable typename remove_cvref_t::scalar_type *ptr; + mutable matx::tensor_t::value_type, out_rank> tmp_out_; + mutable typename remove_cvref_t::value_type *ptr; public: using matxop = bool; - using scalar_type = typename OpA::scalar_type; + using value_type = typename OpA::value_type; using matx_transform_op = bool; using matmul_xform_op = bool; diff --git a/include/matx/operators/matvec.h b/include/matx/operators/matvec.h index 1ae0344a6..1a18cef91 100644 --- a/include/matx/operators/matvec.h +++ b/include/matx/operators/matvec.h @@ -50,12 +50,12 @@ namespace matx float beta_; static constexpr int RANK = remove_cvref_t::Rank(); cuda::std::array out_dims_; - mutable detail::tensor_impl_t::scalar_type, RANK> tmp_out_; - mutable typename remove_cvref_t::scalar_type *ptr; + mutable detail::tensor_impl_t::value_type, RANK> tmp_out_; + mutable typename remove_cvref_t::value_type *ptr; public: using matxop = bool; - using scalar_type = typename OpA::scalar_type; + using value_type = typename OpA::value_type; using matx_transform_op = bool; using matvec_xform_op = bool; diff --git a/include/matx/operators/max.h b/include/matx/operators/max.h index 69aee65c1..62f665124 100644 --- a/include/matx/operators/max.h +++ b/include/matx/operators/max.h @@ -49,12 +49,12 @@ namespace detail { private: OpA a_; cuda::std::array out_dims_; - mutable detail::tensor_impl_t::scalar_type, ORank> tmp_out_; - mutable typename remove_cvref_t::scalar_type *ptr; + mutable detail::tensor_impl_t::value_type, ORank> tmp_out_; + mutable typename remove_cvref_t::value_type *ptr; public: using matxop = bool; - using scalar_type = typename remove_cvref_t::scalar_type; + using value_type = typename remove_cvref_t::value_type; using matx_transform_op = bool; using max_xform_op = bool; diff --git a/include/matx/operators/mean.h b/include/matx/operators/mean.h index 2cf22955a..0274874e5 100644 --- a/include/matx/operators/mean.h +++ b/include/matx/operators/mean.h @@ -49,12 +49,12 @@ namespace detail { private: OpA a_; cuda::std::array out_dims_; - mutable detail::tensor_impl_t::scalar_type, ORank> tmp_out_; - mutable typename remove_cvref_t::scalar_type *ptr; + mutable detail::tensor_impl_t::value_type, ORank> tmp_out_; + mutable typename remove_cvref_t::value_type *ptr; public: using matxop = bool; - using scalar_type = typename remove_cvref_t::scalar_type; + using value_type = typename remove_cvref_t::value_type; using matx_transform_op = bool; using mean_xform_op = bool; diff --git a/include/matx/operators/median.h b/include/matx/operators/median.h index cce0a4672..80e79cfdd 100644 --- a/include/matx/operators/median.h +++ b/include/matx/operators/median.h @@ -49,12 +49,12 @@ namespace detail { private: OpA a_; cuda::std::array out_dims_; - mutable detail::tensor_impl_t::scalar_type, ORank> tmp_out_; - mutable typename remove_cvref_t::scalar_type *ptr; + mutable detail::tensor_impl_t::value_type, ORank> tmp_out_; + mutable typename remove_cvref_t::value_type *ptr; public: using matxop = bool; - using scalar_type = typename remove_cvref_t::scalar_type; + using value_type = typename remove_cvref_t::value_type; using matx_transform_op = bool; using median_xform_op = bool; diff --git a/include/matx/operators/min.h b/include/matx/operators/min.h index 71049914e..e05aeda58 100644 --- a/include/matx/operators/min.h +++ b/include/matx/operators/min.h @@ -49,12 +49,12 @@ namespace detail { private: OpA a_; cuda::std::array out_dims_; - mutable detail::tensor_impl_t::scalar_type, ORank> tmp_out_; - mutable typename remove_cvref_t::scalar_type *ptr; + mutable detail::tensor_impl_t::value_type, ORank> tmp_out_; + mutable typename remove_cvref_t::value_type *ptr; public: using matxop = bool; - using scalar_type = typename remove_cvref_t::scalar_type; + using value_type = typename remove_cvref_t::value_type; using matx_transform_op = bool; using min_xform_op = bool; diff --git a/include/matx/operators/norm.h b/include/matx/operators/norm.h index 546584baa..30cf5c3af 100644 --- a/include/matx/operators/norm.h +++ b/include/matx/operators/norm.h @@ -44,17 +44,17 @@ namespace matx class NormOp : public BaseOp> { private: - using out_type = typename inner_op_type_t::scalar_type>::type; + using out_type = typename inner_op_type_t::value_type>::type; OpA a_; NormOrder order_; static constexpr int ORank = std::is_same_v ? OpA::Rank() - 1 : OpA::Rank() - 2; cuda::std::array out_dims_; - mutable detail::tensor_impl_t::scalar_type, ORank> tmp_out_; - mutable typename remove_cvref_t::scalar_type *ptr; + mutable detail::tensor_impl_t::value_type, ORank> tmp_out_; + mutable typename remove_cvref_t::value_type *ptr; public: using matxop = bool; - using scalar_type = out_type; + using value_type = out_type; using matx_transform_op = bool; using norm_xform_op = bool; using matx_inner_op_impl = bool; // Indicates this operator uses matx operators for its implementation diff --git a/include/matx/operators/operators.h b/include/matx/operators/operators.h index 52482ed9c..e3972ed56 100644 --- a/include/matx/operators/operators.h +++ b/include/matx/operators/operators.h @@ -102,6 +102,7 @@ #include "matx/operators/stack.h" #include "matx/operators/stdd.h" #include "matx/operators/svd.h" +#include "matx/operators/toeplitz.h" #include "matx/operators/trace.h" #include "matx/operators/transpose.h" #include "matx/operators/unique.h" diff --git a/include/matx/operators/outer.h b/include/matx/operators/outer.h index f27f860cb..71c43bcc6 100644 --- a/include/matx/operators/outer.h +++ b/include/matx/operators/outer.h @@ -50,12 +50,12 @@ namespace matx float beta_; static constexpr int RANK = cuda::std::max(remove_cvref_t::Rank(), remove_cvref_t::Rank()) + 1; cuda::std::array out_dims_; - mutable matx::tensor_t::scalar_type, RANK> tmp_out_; - mutable typename remove_cvref_t::scalar_type *ptr; + mutable matx::tensor_t::value_type, RANK> tmp_out_; + mutable typename remove_cvref_t::value_type *ptr; public: using matxop = bool; - using scalar_type = typename OpA::scalar_type; + using value_type = typename OpA::value_type; using matx_transform_op = bool; using outer_xform_op = bool; diff --git a/include/matx/operators/overlap.h b/include/matx/operators/overlap.h index 41f718a58..a99418d9a 100644 --- a/include/matx/operators/overlap.h +++ b/include/matx/operators/overlap.h @@ -46,7 +46,7 @@ namespace matx class OverlapOp : public BaseOp> { public: - using scalar_type = typename T::scalar_type; + using value_type = typename T::value_type; using shape_type = index_t; using self_type = OverlapOp; diff --git a/include/matx/operators/percentile.h b/include/matx/operators/percentile.h index c404850b0..c9b95a4a2 100644 --- a/include/matx/operators/percentile.h +++ b/include/matx/operators/percentile.h @@ -49,12 +49,12 @@ namespace detail { uint32_t q_; PercentileMethod method_; cuda::std::array out_dims_; - mutable detail::tensor_impl_t::scalar_type, ORank> tmp_out_; - mutable typename remove_cvref_t::scalar_type *ptr; + mutable detail::tensor_impl_t::value_type, ORank> tmp_out_; + mutable typename remove_cvref_t::value_type *ptr; public: using matxop = bool; - using scalar_type = typename remove_cvref_t::scalar_type; + using value_type = typename remove_cvref_t::value_type; using matx_transform_op = bool; using prod_xform_op = bool; diff --git a/include/matx/operators/permute.h b/include/matx/operators/permute.h index cfde597a4..58a0391e2 100644 --- a/include/matx/operators/permute.h +++ b/include/matx/operators/permute.h @@ -46,7 +46,7 @@ namespace matx class PermuteOp : public BaseOp> { public: - using scalar_type = typename T::scalar_type; + using value_type = typename T::value_type; using self_type = PermuteOp; private: diff --git a/include/matx/operators/planar.h b/include/matx/operators/planar.h index b368de239..ed5c3af94 100644 --- a/include/matx/operators/planar.h +++ b/include/matx/operators/planar.h @@ -47,12 +47,12 @@ namespace matx public: using matxop = bool; - using scalar_type = typename T1::scalar_type; + using value_type = typename T1::value_type; __MATX_INLINE__ std::string str() const { return "planar(" + op_.str() + ")"; } __MATX_INLINE__ ComplexPlanarOp(T1 op) : op_(op) { - static_assert(is_complex_v>, "Complex planar op only works on complex types"); + static_assert(is_complex_v>, "Complex planar op only works on complex types"); static_assert(Rank() > 0); }; @@ -135,7 +135,7 @@ namespace matx template auto planar(T1 t) { - static_assert(is_complex_v>, "Input to interleaved operator must be complex-valued"); + static_assert(is_complex_v>, "Input to interleaved operator must be complex-valued"); return detail::ComplexPlanarOp(t); } } // end namespace matx diff --git a/include/matx/operators/polyval.h b/include/matx/operators/polyval.h index a3155f674..ae3671bee 100644 --- a/include/matx/operators/polyval.h +++ b/include/matx/operators/polyval.h @@ -52,7 +52,7 @@ namespace matx public: using matxop = bool; - using scalar_type = typename Op::scalar_type; + using value_type = typename Op::value_type; __MATX_INLINE__ std::string str() const { return "polyval()"; } __MATX_INLINE__ PolyvalOp(const Op &op, const Coeffs &coeffs) : op_(op), coeffs_(coeffs) { @@ -60,10 +60,10 @@ namespace matx MATX_STATIC_ASSERT_STR(Op::Rank() == 1, matxInvalidDim, "Input operator must be rank 1"); }; - __MATX_INLINE__ __MATX_DEVICE__ __MATX_HOST__ scalar_type operator()(index_t idx) const + __MATX_INLINE__ __MATX_DEVICE__ __MATX_HOST__ value_type operator()(index_t idx) const { // Horner's method for computing polynomial - scalar_type ttl{coeffs_(0)}; + value_type ttl{coeffs_(0)}; for(int i = 1; i < coeffs_.Size(0); i++) { ttl = ttl * op_(idx) + coeffs_(i); } diff --git a/include/matx/operators/prod.h b/include/matx/operators/prod.h index df08bc113..fee3f9bdd 100644 --- a/include/matx/operators/prod.h +++ b/include/matx/operators/prod.h @@ -49,12 +49,12 @@ namespace detail { private: OpA a_; cuda::std::array out_dims_; - mutable detail::tensor_impl_t::scalar_type, ORank> tmp_out_; - mutable typename remove_cvref_t::scalar_type *ptr; + mutable detail::tensor_impl_t::value_type, ORank> tmp_out_; + mutable typename remove_cvref_t::value_type *ptr; public: using matxop = bool; - using scalar_type = typename remove_cvref_t::scalar_type; + using value_type = typename remove_cvref_t::value_type; using matx_transform_op = bool; using prod_xform_op = bool; diff --git a/include/matx/operators/pwelch.h b/include/matx/operators/pwelch.h index a255fc8b7..84f9275e2 100644 --- a/include/matx/operators/pwelch.h +++ b/include/matx/operators/pwelch.h @@ -51,12 +51,12 @@ namespace matx index_t noverlap_; index_t nfft_; cuda::std::array out_dims_; - mutable detail::tensor_impl_t::scalar_type, 1> tmp_out_; - mutable typename remove_cvref_t::scalar_type *ptr; + mutable detail::tensor_impl_t::value_type, 1> tmp_out_; + mutable typename remove_cvref_t::value_type *ptr; public: using matxop = bool; - using scalar_type = typename OpX::scalar_type::value_type; + using value_type = typename OpX::value_type::value_type; using matx_transform_op = bool; using pwelch_xform_op = bool; diff --git a/include/matx/operators/qr.h b/include/matx/operators/qr.h index 28950345e..3c9dffc05 100644 --- a/include/matx/operators/qr.h +++ b/include/matx/operators/qr.h @@ -48,7 +48,7 @@ namespace detail { public: using matxop = bool; - using scalar_type = typename OpA::scalar_type; + using value_type = typename OpA::value_type; using matx_transform_op = bool; using qr_xform_op = bool; @@ -111,11 +111,11 @@ namespace detail { { private: OpA a_; - matx::tensor_t tmp_out_; + matx::tensor_t tmp_out_; public: using matxop = bool; - using scalar_type = typename OpA::scalar_type; + using value_type = typename OpA::value_type; using matx_transform_op = bool; using cusolver_qr_xform_op = bool; diff --git a/include/matx/operators/r2c.h b/include/matx/operators/r2c.h index dade6c38d..8d76a814e 100644 --- a/include/matx/operators/r2c.h +++ b/include/matx/operators/r2c.h @@ -48,7 +48,7 @@ namespace matx public: using matxop = bool; - using scalar_type = typename T1::scalar_type; + using value_type = typename T1::value_type; __MATX_INLINE__ std::string str() const { return "r2c(" + op_.str() + ")"; } diff --git a/include/matx/operators/reduce.h b/include/matx/operators/reduce.h index b3c0b8dcb..576f91c5f 100644 --- a/include/matx/operators/reduce.h +++ b/include/matx/operators/reduce.h @@ -50,12 +50,12 @@ namespace matx ReductionOp reduction_op_; bool init_; cuda::std::array out_dims_; - mutable detail::tensor_impl_t::scalar_type, ORank> tmp_out_; - mutable typename remove_cvref_t::scalar_type *ptr; + mutable detail::tensor_impl_t::value_type, ORank> tmp_out_; + mutable typename remove_cvref_t::value_type *ptr; public: using matxop = bool; - using scalar_type = typename OpA::scalar_type; + using value_type = typename OpA::value_type; using matx_transform_op = bool; using reduce_xform_op = bool; diff --git a/include/matx/operators/remap.h b/include/matx/operators/remap.h index ba3cda1a6..19538cc89 100644 --- a/include/matx/operators/remap.h +++ b/include/matx/operators/remap.h @@ -54,9 +54,9 @@ namespace matx using matxop = bool; using matxoplvalue = bool; - using scalar_type = typename T::scalar_type; + using value_type = typename T::value_type; using shape_type = std::conditional_t, typename T::shape_type, index_t>; - using index_type = typename IdxType::scalar_type; + using index_type = typename IdxType::value_type; using self_type = RemapOp; static_assert(std::is_integral::value, "RemapOp: Type for index operator must be integral"); static_assert(IdxType::Rank() <= 1, "RemapOp: Rank of index operator must be 0 or 1"); diff --git a/include/matx/operators/repmat.h b/include/matx/operators/repmat.h index b8fd38831..9d22d358f 100644 --- a/include/matx/operators/repmat.h +++ b/include/matx/operators/repmat.h @@ -57,7 +57,7 @@ namespace matx public: using matxop = bool; - using scalar_type = typename T1::scalar_type; + using value_type = typename T1::value_type; __MATX_INLINE__ std::string str() const { return "repmat(" + op_.str() + ")"; } diff --git a/include/matx/operators/resample_poly.h b/include/matx/operators/resample_poly.h index 7fe63ef67..bf8d32191 100644 --- a/include/matx/operators/resample_poly.h +++ b/include/matx/operators/resample_poly.h @@ -46,8 +46,8 @@ namespace detail { class ResamplePolyOp : public BaseOp> { private: - using out_t = std::conditional_t, - typename FilterType::scalar_type, typename FilterType::scalar_type>; + using out_t = std::conditional_t, + typename FilterType::value_type, typename FilterType::value_type>; OpA a_; FilterType f_; index_t up_; @@ -60,7 +60,7 @@ namespace detail { using matxop = bool; using matx_transform_op = bool; using resample_poly_xform_op = bool; - using scalar_type = out_t; + using value_type = out_t; __MATX_INLINE__ std::string str() const { return "resample_poly(" + get_type_str(a_) + "," + get_type_str(f_) + ")";} __MATX_INLINE__ ResamplePolyOp(OpA a, const FilterType &f, index_t up, index_t down) : diff --git a/include/matx/operators/reshape.h b/include/matx/operators/reshape.h index de3a06302..c5fe61e01 100644 --- a/include/matx/operators/reshape.h +++ b/include/matx/operators/reshape.h @@ -47,7 +47,7 @@ namespace matx class ReshapeOp : public BaseOp> { public: - using scalar_type = typename T::scalar_type; + using value_type = typename T::value_type; private: T op_; diff --git a/include/matx/operators/reverse.h b/include/matx/operators/reverse.h index 39701d8fa..f5dd3943d 100644 --- a/include/matx/operators/reverse.h +++ b/include/matx/operators/reverse.h @@ -56,7 +56,7 @@ namespace matx public: using matxop = bool; using matxoplvalue = bool; - using scalar_type = typename T1::scalar_type; + using value_type = typename T1::value_type; using self_type = ReverseOp; __MATX_INLINE__ std::string str() const { return "reverse(" + op_.str() + ")"; } diff --git a/include/matx/operators/scalar_ops.h b/include/matx/operators/scalar_ops.h index 941a96278..fa7e27168 100644 --- a/include/matx/operators/scalar_ops.h +++ b/include/matx/operators/scalar_ops.h @@ -92,7 +92,7 @@ template class UnOp { __MATX_INLINE__ __MATX_DEVICE__ __MATX_HOST__ decltype(auto) operator()(const T1 &v1) const { return op(v1); } - using scalar_type = std::invoke_result_t; + using value_type = std::invoke_result_t; }; template class BinOp { @@ -111,7 +111,7 @@ template class BinOp { return op(v1, v2); } - using scalar_type = std::invoke_result_t; + using value_type = std::invoke_result_t; }; template class TerOp { @@ -127,7 +127,7 @@ template class TerOp { return op(v1, v2, v3); } - using scalar_type = std::invoke_result_t; + using value_type = std::invoke_result_t; }; MATX_UNARY_OP_GEN(ceil, Ceil); @@ -200,7 +200,7 @@ static __MATX_INLINE__ __MATX_HOST__ __MATX_DEVICE__ auto _internal_conj(T v1) return cuda::std::conj(v1); } else { - return conj(v1); + return matx::conj(v1); } } template struct ConjF { diff --git a/include/matx/operators/select.h b/include/matx/operators/select.h index 0bb9580b5..52801f982 100644 --- a/include/matx/operators/select.h +++ b/include/matx/operators/select.h @@ -52,7 +52,7 @@ namespace matx public: using matxop = bool; - using scalar_type = typename T::scalar_type; + using value_type = typename T::value_type; static_assert(IdxType::Rank() == 1, "Rank of index operator must be 1"); __MATX_INLINE__ std::string str() const { return "select(" + op_.str() + ")"; } diff --git a/include/matx/operators/self.h b/include/matx/operators/self.h index 9238b6522..6a3172596 100644 --- a/include/matx/operators/self.h +++ b/include/matx/operators/self.h @@ -53,7 +53,7 @@ namespace matx public: using matxop = bool; - using scalar_type = typename T1::scalar_type; + using value_type = typename T1::value_type; __MATX_INLINE__ std::string str() const { return "self(" + op_.str() + ")"; } diff --git a/include/matx/operators/set.h b/include/matx/operators/set.h index 8bdc8e411..a73ffbfd5 100644 --- a/include/matx/operators/set.h +++ b/include/matx/operators/set.h @@ -67,7 +67,7 @@ class set : public BaseOp> { public: // Type specifier for reflection on class - using scalar_type = typename T::scalar_type; + using value_type = typename T::value_type; using tensor_type = T; using op_type = Op; using matx_setop = bool; diff --git a/include/matx/operators/shift.h b/include/matx/operators/shift.h index e8a02663f..e8da0f591 100644 --- a/include/matx/operators/shift.h +++ b/include/matx/operators/shift.h @@ -60,7 +60,7 @@ namespace matx public: using matxop = bool; using matxoplvalue = bool; - using scalar_type = typename T1::scalar_type; + using value_type = typename T1::value_type; using self_type = ShiftOp; __MATX_INLINE__ std::string str() const { return "shift(" + op_.str() + ")"; } diff --git a/include/matx/operators/sign.h b/include/matx/operators/sign.h index f056923f2..734ee54ef 100644 --- a/include/matx/operators/sign.h +++ b/include/matx/operators/sign.h @@ -51,28 +51,28 @@ namespace matx public: using matxop = bool; - using scalar_type = typename T::scalar_type; + using value_type = typename T::value_type; - scalar_type zval_; + value_type zval_; __MATX_INLINE__ std::string str() const { return "sign(" + get_type_str(op_) + ")"; } - __MATX_INLINE__ SignOp(T op, scalar_type zval) : op_(op), zval_(zval) {}; + __MATX_INLINE__ SignOp(T op, value_type zval) : op_(op), zval_(zval) {}; template __MATX_INLINE__ __MATX_DEVICE__ __MATX_HOST__ auto operator()(Is... indices) const { auto v = get_value(op_,indices...); - if constexpr (is_complex_v ) { - if ( v == scalar_type(0)) { + if constexpr (is_complex_v ) { + if ( v == value_type(0)) { return zval_; } else { return v / abs(v); // sign defintion for complex values } } else { // real branch if( v < 0) - return scalar_type(-1); + return value_type(-1); else if ( v > 0 ) - return scalar_type(1); + return value_type(1); else return zval_; } @@ -106,7 +106,7 @@ namespace matx } // end namespace detail template - __MATX_INLINE__ auto sign(T op, typename T::scalar_type zval=0) { + __MATX_INLINE__ auto sign(T op, typename T::value_type zval=0) { return detail::SignOp(op,zval); } diff --git a/include/matx/operators/slice.h b/include/matx/operators/slice.h index e8a518614..ff84d7f6a 100644 --- a/include/matx/operators/slice.h +++ b/include/matx/operators/slice.h @@ -46,7 +46,7 @@ namespace matx class SliceOp : public BaseOp> { public: - using scalar_type = typename T::scalar_type; + using value_type = typename T::value_type; using shape_type = index_t; using self_type = SliceOp; diff --git a/include/matx/operators/softmax.h b/include/matx/operators/softmax.h index f8617233a..c97c51227 100644 --- a/include/matx/operators/softmax.h +++ b/include/matx/operators/softmax.h @@ -47,12 +47,12 @@ namespace matx OpA a_; PermDims perm_; cuda::std::array out_dims_; - mutable detail::tensor_impl_t::scalar_type, OpA::Rank()> tmp_out_; - mutable typename remove_cvref_t::scalar_type *ptr; + mutable detail::tensor_impl_t::value_type, OpA::Rank()> tmp_out_; + mutable typename remove_cvref_t::value_type *ptr; public: using matxop = bool; - using scalar_type = typename OpA::scalar_type; + using value_type = typename OpA::value_type; using matx_transform_op = bool; using softmax_xform_op = bool; diff --git a/include/matx/operators/sort.h b/include/matx/operators/sort.h index 5817cab4f..b16410d6a 100644 --- a/include/matx/operators/sort.h +++ b/include/matx/operators/sort.h @@ -49,12 +49,12 @@ namespace detail { OpA a_; SortDirection_t dir_; cuda::std::array out_dims_; - mutable detail::tensor_impl_t::scalar_type, OpA::Rank()> tmp_out_; - mutable typename remove_cvref_t::scalar_type *ptr; + mutable detail::tensor_impl_t::value_type, OpA::Rank()> tmp_out_; + mutable typename remove_cvref_t::value_type *ptr; public: using matxop = bool; - using scalar_type = typename OpA::scalar_type; + using value_type = typename OpA::value_type; using matx_transform_op = bool; using sort_xform_op = bool; diff --git a/include/matx/operators/sph2cart.h b/include/matx/operators/sph2cart.h index 8ca29e54c..64de3e044 100644 --- a/include/matx/operators/sph2cart.h +++ b/include/matx/operators/sph2cart.h @@ -52,7 +52,7 @@ namespace matx public: using matxop = bool; - using scalar_type = typename T1::scalar_type; + using value_type = typename T1::value_type; __MATX_INLINE__ std::string str() const { return "sph2cart(" + get_type_str(theta_) + "," + get_type_str(phi_) + "," + get_type_str(r_) + ")"; } diff --git a/include/matx/operators/stack.h b/include/matx/operators/stack.h index c2de85245..aece2ade2 100644 --- a/include/matx/operators/stack.h +++ b/include/matx/operators/stack.h @@ -49,7 +49,7 @@ namespace matx class StackOp : public BaseOp> { using first_type = cuda::std::tuple_element_t<0, cuda::std::tuple>; - using first_value_type = typename first_type::scalar_type; + using first_value_type = typename first_type::value_type; using self_type = StackOp; static constexpr int RANK = first_type::Rank(); @@ -60,7 +60,7 @@ namespace matx using shape_type = index_t; // Scalar type of operation - using scalar_type = first_value_type; + using value_type = first_value_type; template __MATX_INLINE__ std::string get_str() const { @@ -90,7 +90,7 @@ namespace matx if constexpr ( I == N ) { // This should never happen - return scalar_type(-9999); + return value_type(-9999); } else { if ( I < oidx ) { // this is not the correct operator, recurse diff --git a/include/matx/operators/stdd.h b/include/matx/operators/stdd.h index 94c6faff4..82b904b14 100644 --- a/include/matx/operators/stdd.h +++ b/include/matx/operators/stdd.h @@ -49,12 +49,12 @@ namespace detail { private: OpA a_; cuda::std::array out_dims_; - mutable detail::tensor_impl_t::scalar_type, ORank> tmp_out_; - mutable typename remove_cvref_t::scalar_type *ptr; + mutable detail::tensor_impl_t::value_type, ORank> tmp_out_; + mutable typename remove_cvref_t::value_type *ptr; public: using matxop = bool; - using scalar_type = typename remove_cvref_t::scalar_type; + using value_type = typename remove_cvref_t::value_type; using matx_transform_op = bool; using stdd_xform_op = bool; diff --git a/include/matx/operators/sum.h b/include/matx/operators/sum.h index 36e32ad27..03545ea5d 100644 --- a/include/matx/operators/sum.h +++ b/include/matx/operators/sum.h @@ -49,12 +49,12 @@ namespace detail { private: OpA a_; cuda::std::array out_dims_; - mutable detail::tensor_impl_t::scalar_type, ORank> tmp_out_; - mutable typename remove_cvref_t::scalar_type *ptr; + mutable detail::tensor_impl_t::value_type, ORank> tmp_out_; + mutable typename remove_cvref_t::value_type *ptr; public: using matxop = bool; - using scalar_type = typename remove_cvref_t::scalar_type; + using value_type = typename remove_cvref_t::value_type; using matx_transform_op = bool; using sum_xform_op = bool; diff --git a/include/matx/operators/svd.h b/include/matx/operators/svd.h index ba4dad557..33c36202c 100644 --- a/include/matx/operators/svd.h +++ b/include/matx/operators/svd.h @@ -52,7 +52,7 @@ namespace detail { public: using matxop = bool; - using scalar_type = typename OpA::scalar_type; + using value_type = typename OpA::value_type; using matx_transform_op = bool; using svd_xform_op = bool; @@ -110,7 +110,7 @@ namespace detail { public: using matxop = bool; - using scalar_type = typename OpA::scalar_type; + using value_type = typename OpA::value_type; using matx_transform_op = bool; using svd_xform_op = bool; @@ -187,7 +187,7 @@ namespace detail { public: using matxop = bool; - using scalar_type = typename OpA::scalar_type; + using value_type = typename OpA::value_type; using matx_transform_op = bool; using svd_xform_op = bool; diff --git a/include/matx/operators/trace.h b/include/matx/operators/trace.h index 53da7669d..8cbb9f022 100644 --- a/include/matx/operators/trace.h +++ b/include/matx/operators/trace.h @@ -47,12 +47,12 @@ namespace detail { { private: OpA a_; - mutable detail::tensor_impl_t::scalar_type, 0> tmp_out_; - mutable typename remove_cvref_t::scalar_type *ptr; + mutable detail::tensor_impl_t::value_type, 0> tmp_out_; + mutable typename remove_cvref_t::value_type *ptr; public: using matxop = bool; - using scalar_type = typename OpA::scalar_type; + using value_type = typename OpA::value_type; using matx_transform_op = bool; using trace_xform_op = bool; diff --git a/include/matx/operators/transpose.h b/include/matx/operators/transpose.h index c8d7a61ab..7cd73752e 100644 --- a/include/matx/operators/transpose.h +++ b/include/matx/operators/transpose.h @@ -48,11 +48,11 @@ namespace detail { private: OpA a_; cuda::std::array out_dims_; - mutable matx::tensor_t tmp_out_; + mutable matx::tensor_t tmp_out_; public: using matxop = bool; - using scalar_type = typename OpA::scalar_type; + using value_type = typename OpA::value_type; using shape_type = std::conditional_t, typename OpA::shape_type, index_t>; using matx_transform_op = bool; using matxoplvalue = bool; diff --git a/include/matx/operators/unary_operators.h b/include/matx/operators/unary_operators.h index 563345578..111545d58 100644 --- a/include/matx/operators/unary_operators.h +++ b/include/matx/operators/unary_operators.h @@ -42,7 +42,7 @@ typename = typename std::enable_if_t()>> \ [[nodiscard]] __MATX_INLINE__ auto FUNCTION(I1 i1) \ { \ - using I1Type = extract_scalar_type_t; \ + using I1Type = extract_value_type_t; \ using Op = TENSOR_OP; \ const typename detail::base_type::type &base = i1; \ return detail::matxUnaryOp(base, Op()); \ @@ -64,7 +64,7 @@ namespace matx public: // dummy type to signal this is a matxop using matxop = bool; - using scalar_type = typename Op::scalar_type; + using value_type = typename Op::value_type; using self_type = matxUnaryOp; __MATX_INLINE__ const std::string str() const { @@ -377,7 +377,7 @@ namespace matx template ()>> [[nodiscard]] __MATX_INLINE__ auto conj(I1 i1) { - using I1Type = extract_scalar_type_t; + using I1Type = extract_value_type_t; if constexpr (is_complex_v) { using Op = detail::ConjOp; const typename detail::base_type::type &base = i1; @@ -414,7 +414,7 @@ namespace matx template ()>> [[nodiscard]] __MATX_INLINE__ auto real(I1 i1) { - using I1Type = extract_scalar_type_t; + using I1Type = extract_value_type_t; if constexpr (is_complex_v) { using Op = detail::RealOp; const typename detail::base_type::type &base = i1; diff --git a/include/matx/operators/unique.h b/include/matx/operators/unique.h index 289c1e45e..1bd21f1c7 100644 --- a/include/matx/operators/unique.h +++ b/include/matx/operators/unique.h @@ -50,7 +50,7 @@ namespace detail { public: using matxop = bool; - using scalar_type = typename OpA::scalar_type; + using value_type = typename OpA::value_type; using matx_transform_op = bool; using unique_xform_op = bool; diff --git a/include/matx/operators/updownsample.h b/include/matx/operators/updownsample.h index d211c60ab..8a7138ce8 100644 --- a/include/matx/operators/updownsample.h +++ b/include/matx/operators/updownsample.h @@ -53,7 +53,7 @@ namespace matx public: using matxop = bool; using matxoplvalue = bool; - using scalar_type = typename T::scalar_type; + using value_type = typename T::value_type; using self_type = UpsampleOp; static __MATX_INLINE__ constexpr __MATX_HOST__ __MATX_DEVICE__ int32_t Rank() @@ -81,7 +81,7 @@ namespace matx return cuda::std::apply(op_, ind); } - return static_cast(0); + return static_cast(0); } constexpr __MATX_INLINE__ __MATX_HOST__ __MATX_DEVICE__ index_t Size(int32_t dim) const diff --git a/include/matx/operators/var.h b/include/matx/operators/var.h index e1771ddd6..c1f04e8e9 100644 --- a/include/matx/operators/var.h +++ b/include/matx/operators/var.h @@ -50,12 +50,12 @@ namespace detail { OpA a_; int ddof_; cuda::std::array out_dims_; - mutable detail::tensor_impl_t::scalar_type, ORank> tmp_out_; - mutable typename remove_cvref_t::scalar_type *ptr; + mutable detail::tensor_impl_t::value_type, ORank> tmp_out_; + mutable typename remove_cvref_t::value_type *ptr; public: using matxop = bool; - using scalar_type = typename remove_cvref_t::scalar_type; + using value_type = typename remove_cvref_t::value_type; using matx_transform_op = bool; using var_xform_op = bool; diff --git a/include/matx/transforms/ambgfun.h b/include/matx/transforms/ambgfun.h index 6c45e0401..8ff75b55d 100644 --- a/include/matx/transforms/ambgfun.h +++ b/include/matx/transforms/ambgfun.h @@ -168,7 +168,7 @@ void ambgfun_impl(AMFTensor &amf, XTensor &x, MATX_NVTX_START("", matx::MATX_NVTX_LOG_INTERNAL) constexpr int RANK = XTensor::Rank(); - using T1 = typename XTensor::scalar_type; + using T1 = typename XTensor::value_type; MATX_STATIC_ASSERT(is_cuda_complex_v, matxInvalidType); auto ry = x.View(); diff --git a/include/matx/transforms/cgsolve.h b/include/matx/transforms/cgsolve.h index 5327a0207..0f389ca9e 100644 --- a/include/matx/transforms/cgsolve.h +++ b/include/matx/transforms/cgsolve.h @@ -61,7 +61,7 @@ namespace matx template __MATX_INLINE__ void cgsolve_impl(XType X, AType A, BType B, double tol=1e-6, int max_iters=4, cudaStream_t stream=0) { - using scalar_type = typename XType::scalar_type; + using value_type = typename XType::value_type; const int VRANK = XType::Rank(); const int SRANK = XType::Rank() - 1; MATX_NVTX_START("", matx::MATX_NVTX_LOG_API) @@ -82,9 +82,9 @@ namespace matx int converged_host = false; // Construct 3 temporary vectors - auto r0 = make_tensor(X.Shape(), MATX_ASYNC_DEVICE_MEMORY, stream); - auto p = make_tensor(X.Shape(), MATX_ASYNC_DEVICE_MEMORY, stream); - auto Ap = make_tensor(X.Shape(), MATX_ASYNC_DEVICE_MEMORY, stream); + auto r0 = make_tensor(X.Shape(), MATX_ASYNC_DEVICE_MEMORY, stream); + auto p = make_tensor(X.Shape(), MATX_ASYNC_DEVICE_MEMORY, stream); + auto Ap = make_tensor(X.Shape(), MATX_ASYNC_DEVICE_MEMORY, stream); // create aliases to reuse memory auto r1 = r0; @@ -94,13 +94,13 @@ namespace matx for(int i = 0 ; i < SRANK; i++) { scalar_shape[i] = X.Size(i); } - scalar_type N = scalar_type(X.Size(SRANK)); + value_type N = value_type(X.Size(SRANK)); // Construct temporary scalars - auto r0r0 = make_tensor(scalar_shape, MATX_ASYNC_DEVICE_MEMORY, stream); - auto r1r1 = make_tensor(scalar_shape, MATX_ASYNC_DEVICE_MEMORY, stream); - auto pAp = make_tensor(scalar_shape, MATX_ASYNC_DEVICE_MEMORY, stream); - auto norm = make_tensor(scalar_shape, MATX_ASYNC_DEVICE_MEMORY, stream); + auto r0r0 = make_tensor(scalar_shape, MATX_ASYNC_DEVICE_MEMORY, stream); + auto r1r1 = make_tensor(scalar_shape, MATX_ASYNC_DEVICE_MEMORY, stream); + auto pAp = make_tensor(scalar_shape, MATX_ASYNC_DEVICE_MEMORY, stream); + auto norm = make_tensor(scalar_shape, MATX_ASYNC_DEVICE_MEMORY, stream); auto converged = make_tensor({}, MATX_ASYNC_DEVICE_MEMORY, stream); @@ -152,7 +152,7 @@ namespace matx auto updateOp = ( r1 = r0 - (r0r0c/pApc) * Ap, X = X + (r0r0c/pApc) * p); - (IF( pApc != scalar_type(0), updateOp)).run(stream); + (IF( pApc != value_type(0), updateOp)).run(stream); // r1r1 = dot(r1, r1) (r1r1 = sum(r1*r1)).run(stream); @@ -176,7 +176,7 @@ namespace matx // p = r1 + b * p auto updateP = ( p = r1 + (r1r1c/r0r0c) * p); - (IF( pApc != scalar_type(0), updateP)).run(stream); + (IF( pApc != value_type(0), updateP)).run(stream); // Advance residual swap(r0r0, r1r1); diff --git a/include/matx/transforms/channelize_poly.h b/include/matx/transforms/channelize_poly.h index 27490cf5e..2e164ba3a 100644 --- a/include/matx/transforms/channelize_poly.h +++ b/include/matx/transforms/channelize_poly.h @@ -64,8 +64,8 @@ inline void matxChannelizePoly1DInternal(OutType o, const InType &i, #ifdef __CUDACC__ MATX_NVTX_START("", matx::MATX_NVTX_LOG_INTERNAL) - using input_t = typename InType::scalar_type; - using filter_t = typename FilterType::scalar_type; + using input_t = typename InType::value_type; + using filter_t = typename FilterType::value_type; index_t filter_len = filter.Size(FilterType::Rank()-1); @@ -86,8 +86,8 @@ inline void matxChannelizePoly1DInternal(OutType o, const InType &i, template inline size_t matxChannelizePoly1DInternal_SmemSizeBytes(const OutType &o, const InType &, const FilterType &filter) { - using input_t = typename InType::scalar_type; - using filter_t = typename FilterType::scalar_type; + using input_t = typename InType::value_type; + using filter_t = typename FilterType::value_type; index_t filter_len = filter.Size(FilterType::Rank()-1); @@ -126,8 +126,8 @@ inline void matxChannelizePoly1DInternal_Smem(OutType o, const InType &i, const #ifdef __CUDACC__ MATX_NVTX_START("", matx::MATX_NVTX_LOG_INTERNAL) - using input_t = typename InType::scalar_type; - using filter_t = typename FilterType::scalar_type; + using input_t = typename InType::value_type; + using filter_t = typename FilterType::value_type; index_t filter_len = filter.Size(FilterType::Rank()-1); @@ -154,8 +154,8 @@ inline void matxChannelizePoly1DInternal_FusedChan(OutType o, const InType &i, #ifdef __CUDACC__ MATX_NVTX_START("", matx::MATX_NVTX_LOG_INTERNAL) - using input_t = typename InType::scalar_type; - using filter_t = typename FilterType::scalar_type; + using input_t = typename InType::value_type; + using filter_t = typename FilterType::value_type; index_t filter_len = filter.Size(FilterType::Rank()-1); @@ -232,9 +232,9 @@ template inline void channelize_poly_impl(OutType out, const InType &in, const FilterType &f, index_t num_channels, [[maybe_unused]] index_t decimation_factor, cudaStream_t stream = 0) { MATX_NVTX_START("", matx::MATX_NVTX_LOG_API) - using input_t = typename InType::scalar_type; - using filter_t = typename FilterType::scalar_type; - using output_t = typename OutType::scalar_type; + using input_t = typename InType::value_type; + using filter_t = typename FilterType::value_type; + using output_t = typename OutType::value_type; constexpr int IN_RANK = InType::Rank(); constexpr int OUT_RANK = OutType::Rank(); diff --git a/include/matx/transforms/conv.h b/include/matx/transforms/conv.h index c13d522d6..b7e2d127d 100644 --- a/include/matx/transforms/conv.h +++ b/include/matx/transforms/conv.h @@ -64,9 +64,9 @@ inline void matxFFTConv1DInternal(OutputType &o, const InType &i, auto allocate_tensor = [&](auto shape) { if constexpr (is_cuda_executor_v) { - return make_tensor>(shape, MATX_ASYNC_DEVICE_MEMORY, exec.getStream()); + return make_tensor>(shape, MATX_ASYNC_DEVICE_MEMORY, exec.getStream()); } else { - return make_tensor>(shape, MATX_HOST_MALLOC_MEMORY); + return make_tensor>(shape, MATX_HOST_MALLOC_MEMORY); } }; @@ -74,11 +74,11 @@ inline void matxFFTConv1DInternal(OutputType &o, const InType &i, auto s2 = allocate_tensor(in_shape_padded); auto sifft = allocate_tensor(in_shape_padded); - if constexpr (! is_complex_v) { + if constexpr (! is_complex_v) { slice_end[InType::Rank() - 1] = padded_size/2 + 1; } auto s1s = slice(s1, slice_start, slice_end); - if constexpr (! is_complex_v) { + if constexpr (! is_complex_v) { slice_end[FilterType::Rank() - 1] = padded_size/2 + 1; } auto s2s = slice(s2, slice_start, slice_end); @@ -87,9 +87,9 @@ inline void matxFFTConv1DInternal(OutputType &o, const InType &i, // If this is real-valued input we need to accomodate cuFFT's output of N/2+1 complex // samples and use r2c to convert back to N. - if constexpr (!is_complex_v) { + if constexpr (!is_complex_v) { slice_end[InType::Rank() - 1] = padded_size / 2 + 1; - if constexpr (!is_complex_v) { + if constexpr (!is_complex_v) { (sifft = r2c(slice(s1, slice_start, slice_end) * slice(s2, slice_start, slice_end), padded_size)). run(exec); } @@ -98,7 +98,7 @@ inline void matxFFTConv1DInternal(OutputType &o, const InType &i, } } else { - if constexpr (!is_complex_v) { + if constexpr (!is_complex_v) { (sifft = s1 * r2c(slice(s2, slice_start, slice_end), padded_size)).run(exec); } else { @@ -112,7 +112,7 @@ inline void matxFFTConv1DInternal(OutputType &o, const InType &i, // Write directly to output in FULL mode. if (mode == MATX_C_MODE_FULL) { - if constexpr (is_complex_v || is_complex_v) { + if constexpr (is_complex_v || is_complex_v) { (o = ifft(sifft)).run(exec); } else { @@ -130,7 +130,7 @@ inline void matxFFTConv1DInternal(OutputType &o, const InType &i, slice_end[InType::Rank() - 1] = padded_size - filter_size / 2; - if constexpr (is_complex_v || is_complex_v) { + if constexpr (is_complex_v || is_complex_v) { (o = slice(sifft, slice_start, slice_end)).run(exec); } else { @@ -142,7 +142,7 @@ inline void matxFFTConv1DInternal(OutputType &o, const InType &i, slice_start[InType::Rank() - 1] = filter_size - 1; slice_end[InType::Rank() - 1] = padded_size - filter_size + 1; - if constexpr (is_complex_v || is_complex_v) { + if constexpr (is_complex_v || is_complex_v) { (o = slice(sifft, slice_start, slice_end)).run(exec); } else { @@ -169,8 +169,8 @@ inline void matxDirectConv1DInternal(OutputType &o, const InType &i, const auto stream = exec.getStream(); - using strip_input_t = typename InType::scalar_type; - using strip_filter_t = typename FilterType::scalar_type; + using strip_input_t = typename InType::value_type; + using strip_filter_t = typename FilterType::value_type; using shape_type = std::conditional_t, typename OutputType::shape_type, index_t>; size_t filter_len = filter.Size(filter.Rank()-1); @@ -259,7 +259,7 @@ inline void conv1d_impl_internal(OutputType &o, const In1Type &i1, const In2Type } const int Rank = In1Type::Rank(); - //detail::tensor_impl_t &o_base = o; + //detail::tensor_impl_t &o_base = o; typename detail::base_type::type &o_base = o; const typename detail::base_type::type &in1_base = i1; const typename detail::base_type::type &in2_base = i2; diff --git a/include/matx/transforms/cov.h b/include/matx/transforms/cov.h index ba036cea6..3695a970b 100644 --- a/include/matx/transforms/cov.h +++ b/include/matx/transforms/cov.h @@ -57,7 +57,7 @@ struct CovParams_t { template class matxCovHandle_t { public: static constexpr int RANK = TensorTypeA::Rank(); - using T1 = typename TensorTypeA::scalar_type; + using T1 = typename TensorTypeA::value_type; /** * Construct a handle for computing a covariance matrix * diff --git a/include/matx/transforms/cub.h b/include/matx/transforms/cub.h index b8fae9882..2d36aef00 100644 --- a/include/matx/transforms/cub.h +++ b/include/matx/transforms/cub.h @@ -116,8 +116,8 @@ struct EmptyParams_t {}; template class matxCubPlan_t { static constexpr int RANK = OutputTensor::Rank(); - using T1 = typename InputOperator::scalar_type; - using T2 = typename OutputTensor::scalar_type; + using T1 = typename InputOperator::value_type; + using T2 = typename OutputTensor::value_type; public: /** @@ -270,7 +270,7 @@ class matxCubPlan_t { } } else { - const tensor_impl_t base = a; + const tensor_impl_t base = a; for (size_t iter = 0; iter < total_iter; iter++) { auto aop = cuda::std::apply([&a_out](auto... param) { return a_out.GetPointer(param...); }, idx); @@ -319,7 +319,7 @@ class matxCubPlan_t { #ifdef __CUDACC__ MATX_NVTX_START("", matx::MATX_NVTX_LOG_INTERNAL) - const tensor_impl_t base = a; + const tensor_impl_t base = a; if (RANK == 1 || d_temp == nullptr) { if constexpr (is_tensor_view_v) { if (a.IsContiguous()) { @@ -375,7 +375,7 @@ class matxCubPlan_t { if (RANK == 1 || d_temp == nullptr) { if constexpr (is_tensor_view_v) { - const tensor_impl_t base = a; + const tensor_impl_t base = a; if (a.IsContiguous()) { cub::DeviceScan::InclusiveSum(d_temp, temp_storage_bytes, a.Data(), a_out.Data(), static_cast(a.Size(a.Rank()-1)), @@ -904,7 +904,7 @@ inline void ExecSort(OutputTensor &a_out, MATX_NVTX_START("", matx::MATX_NVTX_LOG_INTERNAL) if constexpr (is_tensor_view_v) { - const tensor_impl_t base = a; + const tensor_impl_t base = a; if (a.IsContiguous()) { cub::DeviceSelect::If(d_temp, temp_storage_bytes, @@ -993,7 +993,7 @@ inline void ExecSort(OutputTensor &a_out, stream); } else { - tensor_impl_t base = a; + tensor_impl_t base = a; cub::DeviceSelect::If(d_temp, temp_storage_bytes, cub::CountingInputIterator{0}, @@ -1043,7 +1043,7 @@ inline void ExecSort(OutputTensor &a_out, MATX_NVTX_START("", matx::MATX_NVTX_LOG_INTERNAL) if constexpr (is_tensor_view_v) { - const tensor_impl_t base = a; + const tensor_impl_t base = a; if (a.IsContiguous()) { cub::DeviceSelect::Unique(d_temp, temp_storage_bytes, @@ -1159,13 +1159,13 @@ using cub_cache_t = std::unordered_map -void cub_reduce(OutputTensor &a_out, const InputOperator &a, typename InputOperator::scalar_type init, +void cub_reduce(OutputTensor &a_out, const InputOperator &a, typename InputOperator::value_type init, const cudaStream_t stream = 0) { #ifdef __CUDACC__ MATX_NVTX_START("", matx::MATX_NVTX_LOG_API) // Get parameters required by these tensors - using param_type = typename detail::ReduceParams_t; + using param_type = typename detail::ReduceParams_t; auto reduce_params = param_type{ReduceOp{}, init}; #ifndef MATX_DISABLE_CUB_CACHE @@ -1433,7 +1433,7 @@ void sort_impl(OutputTensor &a_out, const InputOperator &a, lin + a.Size(0), lout, lout + a_out.Size(0), - std::greater()); + std::greater()); } } else { @@ -1449,7 +1449,7 @@ void sort_impl(OutputTensor &a_out, const InputOperator &a, lin + (b+1)*a.Size(1), lout + b*a.Size(1), lout + (b+1)*a.Size(1), - std::greater()); + std::greater()); } } } @@ -1565,14 +1565,14 @@ void cumsum_impl(OutputTensor &a_out, const InputOperator &a, */ template void hist_impl(OutputTensor &a_out, const InputOperator &a, - const typename InputOperator::scalar_type lower, - const typename InputOperator::scalar_type upper, const cudaStream_t stream = 0) + const typename InputOperator::value_type lower, + const typename InputOperator::value_type upper, const cudaStream_t stream = 0) { - static_assert(std::is_same_v, "Output histogram operator must use int type"); + static_assert(std::is_same_v, "Output histogram operator must use int type"); #ifdef __CUDACC__ MATX_NVTX_START("", matx::MATX_NVTX_LOG_API) - detail::HistEvenParams_t hp{lower, upper}; + detail::HistEvenParams_t hp{lower, upper}; #ifndef MATX_DISABLE_CUB_CACHE // Get parameters required by these tensors auto params = @@ -1586,8 +1586,8 @@ void hist_impl(OutputTensor &a_out, const InputOperator &a, auto tmp = new detail::matxCubPlan_t< OutputTensor, InputOperator, detail::CUB_OP_HIST_EVEN, - detail::HistEvenParams_t>{ - a_out, a, detail::HistEvenParams_t{hp}, stream}; + detail::HistEvenParams_t>{ + a_out, a, detail::HistEvenParams_t{hp}, stream}; tmp->ExecHistEven(a_out, a, lower, upper, stream); detail::cub_cache.Insert(params, static_cast(tmp)); @@ -1595,7 +1595,7 @@ void hist_impl(OutputTensor &a_out, const InputOperator &a, else { auto sort_type = static_cast> *>( + detail::CUB_OP_HIST_EVEN, detail::HistEvenParams_t> *>( ret.value()); sort_type->ExecHistEven(a_out, a, lower, upper, stream); } @@ -1603,8 +1603,8 @@ void hist_impl(OutputTensor &a_out, const InputOperator &a, auto tmp = detail::matxCubPlan_t< OutputTensor, InputOperator, detail::CUB_OP_HIST_EVEN, - detail::HistEvenParams_t>{ - a_out, a, detail::HistEvenParams_t{hp}, stream}; + detail::HistEvenParams_t>{ + a_out, a, detail::HistEvenParams_t{hp}, stream}; tmp.ExecHistEven(a_out, a, lower, upper, stream); #endif @@ -1957,7 +1957,7 @@ void unique_impl(OutputTensor &a_out, CountTensor &num_found, const InputOperato cudaStream_t stream = exec.getStream(); // Allocate space for sorted input since CUB doesn't do unique over unsorted inputs - auto sort_tensor = make_tensor(a.Shape(), MATX_ASYNC_DEVICE_MEMORY, stream); + auto sort_tensor = make_tensor(a.Shape(), MATX_ASYNC_DEVICE_MEMORY, stream); matx::sort_impl(sort_tensor, a, SORT_DIR_ASC, stream); diff --git a/include/matx/transforms/einsum.h b/include/matx/transforms/einsum.h index 7592c26f6..f914a0d1c 100644 --- a/include/matx/transforms/einsum.h +++ b/include/matx/transforms/einsum.h @@ -119,7 +119,7 @@ class matxEinsumHandle_t { extents[sizeof...(InT)], strides[sizeof...(InT)], modes[sizeof...(InT)], - MatXTypeToCudaType(), + MatXTypeToCudaType(), CUTENSORNET_COMPUTE_32F, &descNet_); MATX_ASSERT_STR(status == CUTENSORNET_STATUS_SUCCESS, diff --git a/include/matx/transforms/fft/fft_common.h b/include/matx/transforms/fft/fft_common.h index 227f87de1..423f433cd 100644 --- a/include/matx/transforms/fft/fft_common.h +++ b/include/matx/transforms/fft/fft_common.h @@ -66,8 +66,8 @@ namespace detail { MATX_NVTX_START("", matx::MATX_NVTX_LOG_INTERNAL) using index_type = typename OutputTensor::shape_type; - using T1 = typename OutputTensor::scalar_type; - using T2 = typename InputTensor::scalar_type; + using T1 = typename OutputTensor::value_type; + using T2 = typename InputTensor::value_type; constexpr int RANK = OutputTensor::Rank(); index_type starts[RANK] = {0}; @@ -164,8 +164,8 @@ namespace detail { template constexpr __MATX_INLINE__ FFTType DeduceFFTTransformType() { - using T1 = typename T1T::scalar_type; - using T2 = typename T2T::scalar_type; + using T1 = typename T1T::value_type; + using T2 = typename T2T::value_type; // Deduce plan type from view types if constexpr (std::is_same_v>) { diff --git a/include/matx/transforms/fft/fft_cuda.h b/include/matx/transforms/fft/fft_cuda.h index 79e51a259..cbc584e8c 100644 --- a/include/matx/transforms/fft/fft_cuda.h +++ b/include/matx/transforms/fft/fft_cuda.h @@ -75,8 +75,8 @@ struct FftCUDAParams_t { template class matxCUDAFFTPlan_t { public: using index_type = typename OutTensorType::shape_type; - using T1 = typename OutTensorType::scalar_type; - using T2 = typename InTensorType::scalar_type; + using T1 = typename OutTensorType::value_type; + using T2 = typename InTensorType::value_type; static constexpr auto RANK = OutTensorType::Rank(); static_assert(OutTensorType::Rank() == InTensorType::Rank(), "Input and output FFT ranks must match"); @@ -101,7 +101,7 @@ template class matxCUDAFFTPlan_t cufftSetStream(this->plan_, stream); // Normalize input if necessary - using s_type = typename detail::value_promote_t; + using s_type = typename detail::value_promote_t; s_type factor; if (params_.fft_rank == 1) { factor = static_cast(params_.n[0]); @@ -142,7 +142,7 @@ template class matxCUDAFFTPlan_t // cuFFT doesn't scale IFFT the same as MATLAB/Python. Scale it here to // match - using s_type = typename detail::value_promote_t; + using s_type = typename detail::value_promote_t; s_type factor; if (params_.fft_rank == 1) { factor = static_cast(params_.n[0]); @@ -193,7 +193,7 @@ template class matxCUDAFFTPlan_t params.batch = 1; for (int dim = i.Rank() - 2; dim >= 0; dim--) { - if (static_cast(params.batch * i.Size(dim) * sizeof(typename InTensorType::scalar_type)) > max_for_fft_workspace) { + if (static_cast(params.batch * i.Size(dim) * sizeof(typename InTensorType::value_type)) > max_for_fft_workspace) { break; } @@ -369,8 +369,8 @@ template class matxCUDAFFTPlan_t template class matxCUDAFFTPlan1D_t : public matxCUDAFFTPlan_t { public: - using T1 = typename OutTensorType::scalar_type; - using T2 = typename InTensorType::scalar_type; + using T1 = typename OutTensorType::value_type; + using T2 = typename InTensorType::value_type; static constexpr int RANK = OutTensorType::Rank(); /** @@ -506,8 +506,8 @@ template class matxCUDAFFTPlan2D_t : public matxCUDAFFTPlan_t { public: static constexpr int RANK = OutTensorType::Rank(); - using T1 = typename OutTensorType::scalar_type; - using T2 = typename InTensorType::scalar_type; + using T1 = typename OutTensorType::value_type; + using T2 = typename InTensorType::value_type; /** * Construct a 2D FFT plan @@ -667,7 +667,7 @@ using fft_cuda_cache_t = std::unordered_map __MATX_INLINE__ auto getCufft1DSupportedTensor( const TensorOp &in, cudaStream_t stream) { if constexpr ( !(is_tensor_view_v)) { - return make_tensor(in.Shape(), MATX_ASYNC_DEVICE_MEMORY, stream); + return make_tensor(in.Shape(), MATX_ASYNC_DEVICE_MEMORY, stream); } else { bool supported = true; @@ -676,7 +676,7 @@ __MATX_INLINE__ auto getCufft1DSupportedTensor( const TensorOp &in, cudaStream_t if (supported) { return in; } else { - return make_tensor(in.Shape(), MATX_ASYNC_DEVICE_MEMORY, stream); + return make_tensor(in.Shape(), MATX_ASYNC_DEVICE_MEMORY, stream); } } } @@ -687,7 +687,7 @@ __MATX_INLINE__ auto getCufft2DSupportedTensor( const TensorOp &in, cudaStream_t constexpr int IRANK = TensorOp::Rank(); if constexpr ( !is_tensor_view_v) { - return make_tensor(in.Shape(), MATX_ASYNC_DEVICE_MEMORY, stream); + return make_tensor(in.Shape(), MATX_ASYNC_DEVICE_MEMORY, stream); } else { bool supported = true; @@ -703,7 +703,7 @@ __MATX_INLINE__ auto getCufft2DSupportedTensor( const TensorOp &in, cudaStream_t if (supported) { return in; } else { - return make_tensor(in.Shape(), MATX_ASYNC_DEVICE_MEMORY, stream); + return make_tensor(in.Shape(), MATX_ASYNC_DEVICE_MEMORY, stream); } } } diff --git a/include/matx/transforms/fft/fft_fftw.h b/include/matx/transforms/fft/fft_fftw.h index 8aaf1f2b5..d7d490ed4 100644 --- a/include/matx/transforms/fft/fft_fftw.h +++ b/include/matx/transforms/fft/fft_fftw.h @@ -86,8 +86,8 @@ struct FftFFTWParams_t { MATX_NVTX_START("", matx::MATX_NVTX_LOG_INTERNAL) FftFFTWParams_t params; constexpr auto RANK = OutTensorType::Rank(); - using T1 = typename OutTensorType::scalar_type; - using T2 = typename InTensorType::scalar_type; + using T1 = typename OutTensorType::value_type; + using T2 = typename InTensorType::value_type; params.irank = i.Rank(); params.orank = o.Rank(); @@ -150,9 +150,9 @@ struct FftFFTWParams_t { params.odist = (RANK<=2) ? 1 : (int) static_cast(o.Stride(RANK-3)); } - params.is_fp32 = is_fp32_inner_type_v; - if constexpr (std::is_same_v) { + params.is_fp32 = is_fp32_inner_type_v; + if constexpr (std::is_same_v) { params.in_place = o.Data() == i.Data(); } else { params.in_place = false; @@ -318,8 +318,8 @@ class FFTWPlanManager { */ template class matxFFTWPlan_t { public: - using out_scalar_type = typename OutTensorType::scalar_type; - using plan_type = std::conditional_t, fftwf_plan, fftw_plan>; + using out_value_type = typename OutTensorType::value_type; + using plan_type = std::conditional_t, fftwf_plan, fftw_plan>; template matxFFTWPlan_t(OutTensorType &o, @@ -490,7 +490,7 @@ template class matxFFTWPlan_t { } private: - static constexpr bool is_fp32_ = is_fp32_inner_type_v; + static constexpr bool is_fp32_ = is_fp32_inner_type_v; FftFFTWParams_t params_; plan_type plan_; @@ -503,7 +503,7 @@ template class matxFFTWPlan_t { constexpr int RANK=TensorOp::Rank(); if constexpr ( !(is_tensor_view_v)) { - return make_tensor(in.Shape(), MATX_HOST_MALLOC_MEMORY); + return make_tensor(in.Shape(), MATX_HOST_MALLOC_MEMORY); } else { bool supported = true; @@ -523,7 +523,7 @@ template class matxFFTWPlan_t { if (supported) { return in; } else { - return make_tensor(in.Shape(), MATX_HOST_MALLOC_MEMORY); + return make_tensor(in.Shape(), MATX_HOST_MALLOC_MEMORY); } } } @@ -534,7 +534,7 @@ template class matxFFTWPlan_t { constexpr int RANK=TensorOp::Rank(); if constexpr ( !is_tensor_view_v) { - return make_tensor(in.Shape(), MATX_HOST_MALLOC_MEMORY); + return make_tensor(in.Shape(), MATX_HOST_MALLOC_MEMORY); } else { bool supported = true; @@ -550,7 +550,7 @@ template class matxFFTWPlan_t { if (supported) { return in; } else { - return make_tensor(in.Shape(), MATX_HOST_MALLOC_MEMORY); + return make_tensor(in.Shape(), MATX_HOST_MALLOC_MEMORY); } } } @@ -606,7 +606,7 @@ template class matxFFTWPlan_t { (o = out).run(exec); } - using s_type = typename detail::value_promote_t::type>; + using s_type = typename detail::value_promote_t::type>; s_type factor; constexpr s_type s_one = static_cast(1.0); @@ -658,7 +658,7 @@ template class matxFFTWPlan_t { (o = out).run(exec); } - using s_type = typename detail::value_promote_t::type>; + using s_type = typename detail::value_promote_t::type>; s_type factor; constexpr s_type s_one = static_cast(1.0); @@ -689,8 +689,8 @@ template class matxFFTWPlan_t { { MATX_STATIC_ASSERT_STR(OutputTensor::Rank() == InputTensor::Rank(), matxInvalidDim, "Input and output tensor ranks must match"); - MATX_STATIC_ASSERT_STR( (is_fp32_inner_type_v || - is_fp64_inner_type_v), matxInvalidType, + MATX_STATIC_ASSERT_STR( (is_fp32_inner_type_v || + is_fp64_inner_type_v), matxInvalidType, "Host FFTs only support single or double precision floats"); MATX_NVTX_START("", matx::MATX_NVTX_LOG_INTERNAL) @@ -707,8 +707,8 @@ template class matxFFTWPlan_t { { MATX_STATIC_ASSERT_STR(OutputTensor::Rank() == InputTensor::Rank(), matxInvalidDim, "Input and output tensor ranks must match"); - MATX_STATIC_ASSERT_STR( (is_fp32_inner_type_v || - is_fp64_inner_type_v), matxInvalidType, + MATX_STATIC_ASSERT_STR( (is_fp32_inner_type_v || + is_fp64_inner_type_v), matxInvalidType, "Host FFTs only support single or double precision floats"); MATX_NVTX_START("", matx::MATX_NVTX_LOG_INTERNAL) @@ -727,8 +727,8 @@ template class matxFFTWPlan_t { MATX_STATIC_ASSERT_STR(OutputTensor::Rank() == InputTensor::Rank(), matxInvalidDim, "Input and output tensor ranks must match"); MATX_STATIC_ASSERT_STR(InputTensor::Rank() >= 2, matxInvalidSize, "2D FFT must be rank 2 tensor or higher"); - MATX_STATIC_ASSERT_STR( (is_fp32_inner_type_v || - is_fp64_inner_type_v), matxInvalidType, + MATX_STATIC_ASSERT_STR( (is_fp32_inner_type_v || + is_fp64_inner_type_v), matxInvalidType, "Host FFTs only support single or double precision floats"); MATX_NVTX_START("", matx::MATX_NVTX_LOG_INTERNAL) @@ -745,8 +745,8 @@ template class matxFFTWPlan_t { MATX_STATIC_ASSERT_STR(OutputTensor::Rank() == InputTensor::Rank(), matxInvalidDim, "Input and output tensor ranks must match"); MATX_STATIC_ASSERT_STR(InputTensor::Rank() >= 2, matxInvalidSize, "2D FFT must be rank 2 tensor or higher"); - MATX_STATIC_ASSERT_STR( (is_fp32_inner_type_v || - is_fp64_inner_type_v), matxInvalidType, + MATX_STATIC_ASSERT_STR( (is_fp32_inner_type_v || + is_fp64_inner_type_v), matxInvalidType, "Host FFTs only support single or double precision floats"); MATX_NVTX_START("", matx::MATX_NVTX_LOG_INTERNAL) diff --git a/include/matx/transforms/filter.h b/include/matx/transforms/filter.h index adbea437b..5b73472b1 100644 --- a/include/matx/transforms/filter.h +++ b/include/matx/transforms/filter.h @@ -472,7 +472,7 @@ void filter_impl([[maybe_unused]] OutType &o, [[maybe_unused]] const InType &i, params.nonrec.push_back(h_nonrec[j]); } - params.dtype = detail::TypeToInt(); + params.dtype = detail::TypeToInt(); params.ftype = detail::TypeToInt(); // Update when we support different types params.hash = rhash + nrhash; diff --git a/include/matx/transforms/inverse.h b/include/matx/transforms/inverse.h index 39974b856..b794d0d77 100644 --- a/include/matx/transforms/inverse.h +++ b/include/matx/transforms/inverse.h @@ -73,7 +73,7 @@ template constexpr bool CompatibleGemmCBLASTypes() { // All 3 should be the same type - if constexpr (!std::is_same_v || - !std::is_same_v) { + if constexpr (!std::is_same_v || + !std::is_same_v) { return false; } // List of accepted types when A/B/C match - return std::is_same_v || - std::is_same_v || - std::is_same_v> || - std::is_same_v>; + return std::is_same_v || + std::is_same_v || + std::is_same_v> || + std::is_same_v>; } #if MATX_EN_CPU_MATMUL @@ -114,7 +114,7 @@ static MatMulCBLASParams_t GetGemmParams(TensorTypeC &c, that the tensors are in column-major ordering. */ MatMulCBLASParams_t params; - params.dtype = TypeToInt(); + params.dtype = TypeToInt(); params.rank = c.Rank(); // Batches @@ -240,7 +240,7 @@ __MATX_INLINE__ void matmul_exec(TensorTypeC &c, MATX_NVTX_START("", matx::MATX_NVTX_LOG_INTERNAL) static constexpr int RANK = TensorTypeC::Rank(); - using scalar_type = typename TensorTypeC::scalar_type; + using value_type = typename TensorTypeC::value_type; // Prep for batch looping using shape_type = typename TensorTypeA::desc_type::shape_type; @@ -257,15 +257,15 @@ __MATX_INLINE__ void matmul_exec(TensorTypeC &c, std::multiplies()); } - scalar_type salpha; - scalar_type sbeta; - if constexpr (std::is_same_v> || - std::is_same_v>) { + value_type salpha; + value_type sbeta; + if constexpr (std::is_same_v> || + std::is_same_v>) { salpha = {alpha, 0}; sbeta = {beta, 0}; } - else if constexpr (std::is_same_v || - std::is_same_v) { + else if constexpr (std::is_same_v || + std::is_same_v) { salpha = alpha;; sbeta = beta; } @@ -276,25 +276,25 @@ __MATX_INLINE__ void matmul_exec(TensorTypeC &c, auto b_ptr = b.Data(); auto c_ptr = c.Data(); - if constexpr (std::is_same_v) { + if constexpr (std::is_same_v) { cblas_sgemm_batch_strided(CblasRowMajor, params.opA, params.opB, params.m, params.n, params.k, salpha, a_ptr, params.lda, params.astride, b_ptr, params.ldb, params.bstride, sbeta, c_ptr, params.ldc, params.cstride, params.batch); - } else if constexpr (std::is_same_v) { + } else if constexpr (std::is_same_v) { cblas_dgemm_batch_strided(CblasRowMajor, params.opA, params.opB, params.m, params.n, params.k, salpha, a_ptr, params.lda, params.astride, b_ptr, params.ldb, params.bstride, sbeta, c_ptr, params.ldc, params.cstride, params.batch); - } else if constexpr (std::is_same_v>) { + } else if constexpr (std::is_same_v>) { cblas_cgemm_batch_strided(CblasRowMajor, params.opA, params.opB, params.m, params.n, params.k, (void *)&salpha, (void *)a_ptr, params.lda, params.astride, (void *)b_ptr, params.ldb, params.bstride, (void *)&sbeta, (void *)c_ptr, params.ldc, params.cstride, params.batch); - } else if constexpr (std::is_same_v>) { + } else if constexpr (std::is_same_v>) { cblas_zgemm_batch_strided(CblasRowMajor, params.opA, params.opB, params.m, params.n, params.k, (void *)&salpha, (void *)a_ptr, params.lda, params.astride, @@ -309,25 +309,25 @@ __MATX_INLINE__ void matmul_exec(TensorTypeC &c, auto bp = cuda::std::apply([&b](auto... param) { return b.GetPointer(param...); }, b_idx); auto cp = cuda::std::apply([&c](auto... param) { return c.GetPointer(param...); }, c_idx); - if constexpr (std::is_same_v) { + if constexpr (std::is_same_v) { cblas_sgemm_batch_strided(CblasRowMajor, params.opA, params.opB, params.m, params.n, params.k, salpha, ap, params.lda, params.astride, bp, params.ldb, params.bstride, sbeta, cp, params.ldc, params.cstride, params.batch); - } else if constexpr (std::is_same_v) { + } else if constexpr (std::is_same_v) { cblas_dgemm_batch_strided(CblasRowMajor, params.opA, params.opB, params.m, params.n, params.k, salpha, ap, params.lda, params.astride, bp, params.ldb, params.bstride, sbeta, cp, params.ldc, params.cstride, params.batch); - } else if constexpr (std::is_same_v>) { + } else if constexpr (std::is_same_v>) { cblas_cgemm_batch_strided(CblasRowMajor, params.opA, params.opB, params.m, params.n, params.k, (void *)&salpha, (void *)ap, params.lda, params.astride, (void *)bp, params.ldb, params.bstride, (void *)&sbeta, (void *)cp, params.ldc, params.cstride, params.batch); - } else if constexpr (std::is_same_v>) { + } else if constexpr (std::is_same_v>) { cblas_zgemm_batch_strided(CblasRowMajor, params.opA, params.opB, params.m, params.n, params.k, (void *)&salpha, (void *)ap, params.lda, params.astride, @@ -358,25 +358,25 @@ __MATX_INLINE__ void matmul_exec(TensorTypeC &c, auto bp = cuda::std::apply([&b](auto... param) { return b.GetPointer(param...); }, b_idx); auto cp = cuda::std::apply([&c](auto... param) { return c.GetPointer(param...); }, c_idx); - if constexpr (std::is_same_v) { + if constexpr (std::is_same_v) { cblas_sgemm(CblasRowMajor, params.opA, params.opB, params.m, params.n, params.k, salpha, ap, params.lda, bp, params.ldb, sbeta, cp, params.ldc); - } else if constexpr (std::is_same_v) { + } else if constexpr (std::is_same_v) { cblas_dgemm(CblasRowMajor, params.opA, params.opB, params.m, params.n, params.k, salpha, ap, params.lda, bp, params.ldb, sbeta, cp, params.ldc); - } else if constexpr (std::is_same_v>) { + } else if constexpr (std::is_same_v>) { cblas_cgemm(CblasRowMajor, params.opA, params.opB, params.m, params.n, params.k, (void *)&salpha, (const void **)ap, params.lda, (const void **)bp, params.ldb, (void *)&sbeta, (void **)cp, params.ldc); - } else if constexpr (std::is_same_v>) { + } else if constexpr (std::is_same_v>) { cblas_zgemm(CblasRowMajor, params.opA, params.opB, params.m, params.n, params.k, (void *)&salpha, (const void **)ap, params.lda, @@ -436,7 +436,7 @@ __MATX_INLINE__ auto getCBLASSupportedTensor(const Op &in) { constexpr int RANK = Op::Rank(); if constexpr (!(is_tensor_view_v)) { - return make_tensor(in.Shape(), MATX_HOST_MALLOC_MEMORY); + return make_tensor(in.Shape(), MATX_HOST_MALLOC_MEMORY); } else { bool supported = true; @@ -453,7 +453,7 @@ __MATX_INLINE__ auto getCBLASSupportedTensor(const Op &in) { if (supported) { return in; } else { - return make_tensor(in.Shape(), MATX_HOST_MALLOC_MEMORY); + return make_tensor(in.Shape(), MATX_HOST_MALLOC_MEMORY); } } } @@ -495,17 +495,17 @@ void matmul_impl([[maybe_unused]] TensorTypeC C, MATX_ASSERT_STR(MATX_EN_CPU_MATMUL, matxInvalidExecutor, "Trying to run MatMul on host executor but host MatMul support is not configured"); #if MATX_EN_CPU_MATMUL - constexpr auto is_c_complex = is_complex_v; + constexpr auto is_c_complex = is_complex_v; if constexpr (is_c_complex) { - constexpr auto is_a_complex = is_complex_v; - constexpr auto is_b_complex = is_complex_v; + constexpr auto is_a_complex = is_complex_v; + constexpr auto is_b_complex = is_complex_v; MATX_STATIC_ASSERT_STR((is_a_complex || is_b_complex), matxInvalidType, "If C is complex, then either A or B must be complex"); } // promote A and B to the type of C - auto A_ = as_type(A); - auto B_ = as_type(B); + auto A_ = as_type(A); + auto B_ = as_type(B); MATX_STATIC_ASSERT_STR((detail::CompatibleGemmCBLASTypes()), matxInvalidType, "Combination of A/B/C types are not supported for host MatMul"); diff --git a/include/matx/transforms/matmul/matmul_cuda.h b/include/matx/transforms/matmul/matmul_cuda.h index aca63dd78..abb77b7d6 100644 --- a/include/matx/transforms/matmul/matmul_cuda.h +++ b/include/matx/transforms/matmul/matmul_cuda.h @@ -73,35 +73,35 @@ typedef enum { template constexpr bool CompatibleGemmCUDATypes() { - if constexpr (!std::is_same_v && - !std::is_same_v && - !std::is_same_v) { + if constexpr (!std::is_same_v && + !std::is_same_v && + !std::is_same_v) { return false; } if constexpr (PROV == PROVIDER_TYPE_CUBLASLT) { - if constexpr (std::is_same_v && - std::is_same_v) { + if constexpr (std::is_same_v && + std::is_same_v) { // List of accepted types when A/B/C match - return std::is_same_v || - std::is_same_v || - std::is_same_v || - std::is_same_v || - std::is_same_v> || - std::is_same_v> || - std::is_same_v || - std::is_same_v || - std::is_same_v; + return std::is_same_v || + std::is_same_v || + std::is_same_v || + std::is_same_v || + std::is_same_v> || + std::is_same_v> || + std::is_same_v || + std::is_same_v || + std::is_same_v; } // Accumulator type different from A/B - else if constexpr ( std::is_same_v && - !std::is_same_v) { - return (std::is_same_v && std::is_same_v) || - (std::is_same_v && std::is_same_v) || - (std::is_same_v && std::is_same_v) || - (std::is_same_v && std::is_same_v) || - (std::is_same_v && std::is_same_v); + else if constexpr ( std::is_same_v && + !std::is_same_v) { + return (std::is_same_v && std::is_same_v) || + (std::is_same_v && std::is_same_v) || + (std::is_same_v && std::is_same_v) || + (std::is_same_v && std::is_same_v) || + (std::is_same_v && std::is_same_v); } } @@ -141,9 +141,9 @@ template class MatMulCUDAHandle_t { public: - using T1 = typename TensorTypeC::scalar_type; - using T2 = typename TensorTypeA::scalar_type; - using T3 = typename TensorTypeB::scalar_type; + using T1 = typename TensorTypeC::value_type; + using T2 = typename TensorTypeA::value_type; + using T3 = typename TensorTypeB::value_type; static constexpr int RANK = TensorTypeC::Rank(); // We allow a batch stride of 0 on one of the tensors, so only make sure C's rank matches one of them @@ -355,7 +355,7 @@ class MatMulCUDAHandle_t { } else { if constexpr (PROV == PROVIDER_TYPE_CUBLASLT) { - if constexpr (is_complex_half_v) { + if constexpr (is_complex_half_v) { // For half complex we always copy to a new tensor so it is always cublas op N params.opA = CUBLAS_OP_N; } else if ( a.Stride(TensorTypeA::Rank()-1) > 1 // last stride > 1 @@ -365,7 +365,7 @@ class MatMulCUDAHandle_t { params.opA = CUBLAS_OP_N; } - if constexpr (is_complex_half_v) { + if constexpr (is_complex_half_v) { // For half complex we always copy to a new tensor so it is always cublas op N params.opB = CUBLAS_OP_N; } else if ( b.Stride(TensorTypeB::Rank()-1) > 1 // last stride > 1 @@ -400,12 +400,12 @@ class MatMulCUDAHandle_t { } // for complex half we have copied to planar row-major - if (is_complex_half_v) { + if (is_complex_half_v) { params.ldb = b.Size(TensorTypeB::Rank()-1); } // for complex half we have copied to planar row-major - if constexpr (is_complex_half_v) { + if constexpr (is_complex_half_v) { params.lda = a.Size(TensorTypeA::Rank()-1); } @@ -1104,7 +1104,7 @@ __MATX_INLINE__ auto getCublasSupportedTensor( const Op &in, cudaStream_t stream constexpr int RANK=Op::Rank(); if constexpr ( !(is_tensor_view_v)) { - return make_tensor(in.Shape(), MATX_ASYNC_DEVICE_MEMORY, stream); + return make_tensor(in.Shape(), MATX_ASYNC_DEVICE_MEMORY, stream); } else { bool supported = true; @@ -1122,7 +1122,7 @@ __MATX_INLINE__ auto getCublasSupportedTensor( const Op &in, cudaStream_t stream if(supported) { return in; } else { - return make_tensor(in.Shape(), MATX_ASYNC_DEVICE_MEMORY, stream); + return make_tensor(in.Shape(), MATX_ASYNC_DEVICE_MEMORY, stream); } } } @@ -1167,17 +1167,17 @@ void matmul_impl(TensorTypeC C, const TensorTypeA A, MATX_NVTX_START("", matx::MATX_NVTX_LOG_API) const auto stream = exec.getStream(); - constexpr auto is_c_complex = is_complex_v; + constexpr auto is_c_complex = is_complex_v; if constexpr (is_c_complex) { - constexpr auto is_a_complex = is_complex_v; - constexpr auto is_b_complex = is_complex_v; + constexpr auto is_a_complex = is_complex_v; + constexpr auto is_b_complex = is_complex_v; static_assert(is_a_complex || is_b_complex, "If C is complex then either A or B should be complex "); } // promote A and B to the type of C - auto A_ = as_type(A); - auto B_ = as_type(B); + auto A_ = as_type(A); + auto B_ = as_type(B); static_assert(detail::CompatibleGemmCUDATypes(), "Combination of A/B/C types are not supported"); diff --git a/include/matx/transforms/percentile.h b/include/matx/transforms/percentile.h index 0f5d40806..d500e0bbf 100644 --- a/include/matx/transforms/percentile.h +++ b/include/matx/transforms/percentile.h @@ -102,7 +102,7 @@ void __MATX_INLINE__ percentile_impl(OutType dest, const InType &in, uint32_t q, if constexpr (OutType::Rank() == 0) { auto insize = TotalSize(in); - matx::tensor_t sort_out; + matx::tensor_t sort_out; if constexpr (is_cuda_executor_v) { make_tensor(sort_out, {insize}, MATX_ASYNC_DEVICE_MEMORY, exec.getStream()); } @@ -133,7 +133,7 @@ void __MATX_INLINE__ percentile_impl(OutType dest, const InType &in, uint32_t q, subidx = q/100. * (static_cast(insize) - alpha - beta + 1) + alpha - 1; auto int_val = at(sort_out, static_cast(subidx)); (dest = at(sort_out, static_cast(subidx)) + - as_type( + as_type( (subidx - std::floor(subidx)) * as_type(at(sort_out, static_cast(subidx + 1)) - int_val) ) @@ -149,10 +149,10 @@ void __MATX_INLINE__ percentile_impl(OutType dest, const InType &in, uint32_t q, break; } case PercentileMethod::MIDPOINT: { - (dest = as_type( + (dest = as_type( (at(sort_out, static_cast(cuda::std::ceil(base_index))) + at(sort_out, static_cast(cuda::std::ceil(base_index + 1)))) / - static_cast(2)) + static_cast(2)) ).run(exec); break; } diff --git a/include/matx/transforms/pwelch.h b/include/matx/transforms/pwelch.h index 7479bc488..33dfbeed4 100644 --- a/include/matx/transforms/pwelch.h +++ b/include/matx/transforms/pwelch.h @@ -48,7 +48,7 @@ namespace matx // Create temporary space for fft outputs index_t batches = x_with_overlaps.Shape()[0]; - auto X_with_overlaps = make_tensor>({batches,static_cast(nfft)},MATX_ASYNC_DEVICE_MEMORY,stream); + auto X_with_overlaps = make_tensor>({batches,static_cast(nfft)},MATX_ASYNC_DEVICE_MEMORY,stream); if constexpr (std::is_same_v) { @@ -64,7 +64,7 @@ namespace matx auto mag_sq_X_with_overlaps = X_with_overlaps.RealView(); // Perform the reduction across 'batches' rows and normalize - auto norm_factor = static_cast(1.) / static_cast(batches); + auto norm_factor = static_cast(1.) / static_cast(batches); (Pxx = sum(mag_sq_X_with_overlaps, {0}) * norm_factor).run(stream); } diff --git a/include/matx/transforms/qr.h b/include/matx/transforms/qr.h index e30085fc4..7fce7bd92 100644 --- a/include/matx/transforms/qr.h +++ b/include/matx/transforms/qr.h @@ -45,7 +45,7 @@ namespace detail { template inline auto qr_internal_workspace(const AType &A, cudaStream_t stream) { - using ATypeS = typename AType::scalar_type; + using ATypeS = typename AType::value_type; const int RANK = AType::Rank(); index_t m = A.Size(RANK-2); @@ -79,7 +79,7 @@ namespace detail { MATX_ASSERT_STR(AType::Rank() == QType::Rank(), matxInvalidDim, "qr: A and Q must have the same rank"); MATX_ASSERT_STR(AType::Rank() == RType::Rank(), matxInvalidDim, "qr: A and R must have the same rank"); - using ATypeS = typename AType::scalar_type; + using ATypeS = typename AType::value_type; using NTypeS = typename inner_op_type_t::type; const int RANK = AType::Rank(); diff --git a/include/matx/transforms/reduce.h b/include/matx/transforms/reduce.h index 708d6ac00..431a25612 100644 --- a/include/matx/transforms/reduce.h +++ b/include/matx/transforms/reduce.h @@ -1042,8 +1042,8 @@ __global__ void matxReduceKernel(OutType dest, const InType in, constexpr uint32_t RANK = OutType::Rank(); constexpr uint32_t DRANK = InType::Rank() - RANK; cuda::std::array indices; - using scalar_type = typename InType::scalar_type; - using T = typename OutType::scalar_type; + using value_type = typename InType::value_type; + using T = typename OutType::value_type; [[maybe_unused]] bool valid; // This is for 2 stage reduction @@ -1051,10 +1051,10 @@ __global__ void matxReduceKernel(OutType dest, const InType in, // nvcc limitation here. we have to declare shared memory with the same type // across all template functions then recast to the type we want extern __shared__ char smemc_[]; - scalar_type *smem_ = reinterpret_cast(smemc_); + value_type *smem_ = reinterpret_cast(smemc_); int s2_size, soff; - scalar_type *smem; + value_type *smem; // if blockDim.x > 32 we need a 2 stage reduction if (blockDim.x > 32) { @@ -1194,8 +1194,8 @@ __global__ void matxReduceKernel(OutType dest, const InType in, template __global__ void matxIndexKernel(OutType dest, TensorIndexType idest, InType in, [[maybe_unused]] index_t mult) { - using index_type = typename TensorIndexType::scalar_type; - using T = typename OutType::scalar_type; + using index_type = typename TensorIndexType::value_type; + using T = typename OutType::value_type; T in_val; constexpr uint32_t RANK = TensorIndexType::Rank(); constexpr uint32_t DRANK = InType::Rank() - RANK; @@ -1366,8 +1366,8 @@ void __MATX_INLINE__ reduce(OutType dest, [[maybe_unused]] TensorIndexType idest #ifdef __CUDACC__ MATX_NVTX_START("", matx::MATX_NVTX_LOG_API) - using scalar_type = typename InType::scalar_type; - using T = typename OutType::scalar_type; + using value_type = typename InType::value_type; + using T = typename OutType::value_type; static_assert(OutType::Rank() < InType::Rank()); static_assert(is_matx_reduction_v, "Must use a reduction operator for reducing"); @@ -1391,7 +1391,7 @@ void __MATX_INLINE__ reduce(OutType dest, [[maybe_unused]] TensorIndexType idest auto mult = std::accumulate(sizes.begin() + 1, sizes.end(), 1, std::multiplies()); - detail::matxReduceKernel<<>>( + detail::matxReduceKernel<<>>( dest, in, ReduceOp(), mult); // If we need the indices too, launch that kernel @@ -1488,7 +1488,7 @@ void __MATX_INLINE__ mean_impl(OutType dest, const InType &in, MATX_NVTX_START("mean_impl(" + get_type_str(in) + ")", matx::MATX_NVTX_LOG_API) static_assert(OutType::Rank() < InType::Rank(), "reduction dimensions must be <= Rank of input"); - using inner_type = typename inner_op_type_t::type; + using inner_type = typename inner_op_type_t::type; inner_type scale = 1; cudaStream_t stream = exec.getStream(); @@ -1535,17 +1535,17 @@ void __MATX_INLINE__ mean_impl(OutType dest, const InType &in, [[maybe_unused]] MATX_NVTX_START("mean_impl(" + get_type_str(in) + ")", matx::MATX_NVTX_LOG_API) static_assert(OutType::Rank() < InType::Rank(), "reduction dimensions must be <= Rank of input"); - using inner_type = typename inner_op_type_t::type; + using inner_type = typename inner_op_type_t::type; inner_type scale = 1; auto ft = [&](auto &&lin, auto &&lout, [[maybe_unused]] auto &&lbegin, [[maybe_unused]] auto &&lend) { if constexpr (OutType::Rank() == 0) { auto ts = TotalSize(in); - *lout = std::accumulate(lin, lin + ts, static_cast(0)) / static_cast(ts); + *lout = std::accumulate(lin, lin + ts, static_cast(0)) / static_cast(ts); } else { for (index_t b = 0; b < lin.Size(0); b++) { - *(lout + b) = std::accumulate(lin + lbegin[b], lin + lend[b], static_cast(0)) / static_cast(lin.Size(1)); + *(lout + b) = std::accumulate(lin + lbegin[b], lin + lend[b], static_cast(0)) / static_cast(lin.Size(1)); } } }; @@ -1584,8 +1584,8 @@ void __MATX_INLINE__ softmax_impl(OutType dest, const InType &in, #ifdef __CUDACC__ MATX_NVTX_START("softmax_impl(" + get_type_str(in) + ")", matx::MATX_NVTX_LOG_API) - auto tmp_sum = make_tensor({}, MATX_ASYNC_DEVICE_MEMORY, stream); - auto tmp_max = make_tensor({}, MATX_ASYNC_DEVICE_MEMORY, stream); + auto tmp_sum = make_tensor({}, MATX_ASYNC_DEVICE_MEMORY, stream); + auto tmp_max = make_tensor({}, MATX_ASYNC_DEVICE_MEMORY, stream); max_impl(tmp_max, in, cudaExecutor{stream}); sum_impl(tmp_sum, exp(in - tmp_max), stream); (dest = exp(in - tmp_max) / tmp_sum).run(stream); @@ -1653,8 +1653,8 @@ void __MATX_INLINE__ softmax_impl(OutType dest, const InType &in, PermDims dims, } } - auto tmp_sum = make_tensor(red_shape, MATX_ASYNC_DEVICE_MEMORY, stream); - auto tmp_max = make_tensor(red_shape, MATX_ASYNC_DEVICE_MEMORY, stream); + auto tmp_sum = make_tensor(red_shape, MATX_ASYNC_DEVICE_MEMORY, stream); + auto tmp_max = make_tensor(red_shape, MATX_ASYNC_DEVICE_MEMORY, stream); max_impl(tmp_max, permute(in, perm), stream); sum_impl(tmp_sum, exp(permute(in, perm) - clone(tmp_max, clone_dims)), stream); @@ -1693,7 +1693,7 @@ void __MATX_INLINE__ median_impl(OutType dest, #ifdef __CUDACC__ if constexpr ( OutType::Rank() <= 1 && InType::Rank() <=2 ) { MATX_NVTX_START("median_impl(" + get_type_str(in) + ")", matx::MATX_NVTX_LOG_API) - using T = typename OutType::scalar_type; + using T = typename OutType::value_type; constexpr int RANK_IN = InType::Rank(); static_assert(RANK_IN <= 2 && (RANK_IN == OutType::Rank() + 1)); @@ -1802,7 +1802,7 @@ void __MATX_INLINE__ median_impl(OutType dest, const InType &in, [[maybe_unused] auto ft = [&](auto &&lin, auto &&lout, [[maybe_unused]] auto &&lbegin, [[maybe_unused]] auto &&lend) { if constexpr (OutType::Rank() == 0) { auto insize = TotalSize(in); - auto tin = new typename InType::scalar_type[insize]; + auto tin = new typename InType::value_type[insize]; std::partial_sort_copy( lin, lin + insize, tin, @@ -1818,7 +1818,7 @@ void __MATX_INLINE__ median_impl(OutType dest, const InType &in, [[maybe_unused] } else { auto insize = lin.Size(1); - auto tin = new typename InType::scalar_type[insize]; + auto tin = new typename InType::value_type[insize]; for (index_t b = 0; b < lin.Size(0); b++) { std::partial_sort_copy( lin + lbegin[b], lin + lend[b], @@ -1893,11 +1893,11 @@ void __MATX_INLINE__ sum_impl(OutType dest, const InType &in, [[maybe_unused]] H MATX_NVTX_START("sum_impl(" + get_type_str(in) + ")", matx::MATX_NVTX_LOG_API) auto ft = [&](auto &&lin, auto &&lout, [[maybe_unused]] auto &&lbegin, [[maybe_unused]] auto &&lend) { if constexpr (OutType::Rank() == 0) { - *lout = std::accumulate(lin, lin + lin.Size(0), static_cast(0)); + *lout = std::accumulate(lin, lin + lin.Size(0), static_cast(0)); } else { for (index_t b = 0; b < lin.Size(0); b++) { - auto f = std::accumulate(lin + lbegin[b], lin + lend[b], static_cast(0)); + auto f = std::accumulate(lin + lbegin[b], lin + lend[b], static_cast(0)); *(lout + b) = f; } } @@ -1934,7 +1934,7 @@ void __MATX_INLINE__ prod_impl(OutType dest, const InType &in, cudaExecutor exec cudaStream_t stream = exec.getStream(); // example-begin reduce-1 // Reduce "in" into "dest" using a product operation as the reduction type - reduce(dest, in, detail::reduceOpProd(), stream, true); + reduce(dest, in, detail::reduceOpProd(), stream, true); // example-end reduce-1 #endif } @@ -1964,15 +1964,15 @@ void __MATX_INLINE__ prod_impl(OutType dest, const InType &in, [[maybe_unused]] if constexpr (OutType::Rank() == 0) { *lout = std::accumulate(lin, lin + TotalSize(in), - static_cast(1), - std::multiplies()); + static_cast(1), + std::multiplies()); } else { for (index_t b = 0; b < lin.Size(0); b++) { *(lout + b) = std::accumulate(lin + lbegin[b], lin + lend[b], - static_cast(1), - std::multiplies()); + static_cast(1), + std::multiplies()); } } }; @@ -2085,7 +2085,7 @@ void __MATX_INLINE__ argmax_impl(OutType dest, TensorIndexType &idest, const InT // example-begin reduce-2 // Reduce "in" into both "dest" and "idest" using a max operation for the reduction. "dest" will contain // the reduced values while "idest" will include the indices of the reduced values - reduce(dest, idest, in, detail::reduceOpMax(), stream, true); + reduce(dest, idest, in, detail::reduceOpMax(), stream, true); // example-end reduce-2 #endif } @@ -2233,7 +2233,7 @@ void __MATX_INLINE__ argmin_impl(OutType dest, TensorIndexType &idest, const InT MATX_NVTX_START("argmin_impl(" + get_type_str(in) + ")", matx::MATX_NVTX_LOG_API) cudaStream_t stream = exec.getStream(); - reduce(dest, idest, in, detail::reduceOpMin(), stream, true); + reduce(dest, idest, in, detail::reduceOpMin(), stream, true); #endif } @@ -2308,7 +2308,7 @@ void __MATX_INLINE__ any_impl(OutType dest, const InType &in, cudaExecutor exec #ifdef __CUDACC__ MATX_NVTX_START("any_impl(" + get_type_str(in) + ")", matx::MATX_NVTX_LOG_API) cudaStream_t stream = exec.getStream(); - reduce(dest, in, detail::reduceOpAny(), stream, true); + reduce(dest, in, detail::reduceOpAny(), stream, true); #endif } @@ -2338,13 +2338,13 @@ void __MATX_INLINE__ any_impl(OutType dest, const InType &in, [[maybe_unused]] H auto ft = [&](auto &&lin, auto &&lout, [[maybe_unused]] auto &&lbegin, [[maybe_unused]] auto &&lend) { if constexpr (OutType::Rank() == 0) { - *lout = std::any_of(lin, lin + TotalSize(in), [](typename InType::scalar_type vin) { + *lout = std::any_of(lin, lin + TotalSize(in), [](typename InType::value_type vin) { return vin != 0; }); } else { for (index_t b = 0; b < lin.Size(0); b++) { - lout[b] = std::any_of(lin + lbegin[b], lin + lend[b], [](typename InType::scalar_type vin) { + lout[b] = std::any_of(lin + lbegin[b], lin + lend[b], [](typename InType::value_type vin) { return vin != 0; }); } @@ -2381,7 +2381,7 @@ void __MATX_INLINE__ all_impl(OutType dest, const InType &in, cudaExecutor exec #ifdef __CUDACC__ MATX_NVTX_START("all_impl(" + get_type_str(in) + ")", matx::MATX_NVTX_LOG_API) cudaStream_t stream = exec.getStream(); - reduce(dest, in, detail::reduceOpAll(), stream, true); + reduce(dest, in, detail::reduceOpAll(), stream, true); #endif } @@ -2411,13 +2411,13 @@ void __MATX_INLINE__ all_impl(OutType dest, const InType &in, [[maybe_unused]] H auto ft = [&](auto &&lin, auto &&lout, [[maybe_unused]] auto &&lbegin, [[maybe_unused]] auto &&lend) { if constexpr (OutType::Rank() == 0) { - *lout = std::all_of(lin, lin + TotalSize(in), [](typename InType::scalar_type vin) { + *lout = std::all_of(lin, lin + TotalSize(in), [](typename InType::value_type vin) { return vin != 0; }); } else { for (index_t b = 0; b < lin.Size(0); b++) { - lout[b] = std::all_of(lin + lbegin[b], lin + lend[b], [](typename InType::scalar_type vin) { + lout[b] = std::all_of(lin + lbegin[b], lin + lend[b], [](typename InType::value_type vin) { return vin != 0; }); } @@ -2541,7 +2541,7 @@ void __MATX_INLINE__ var_impl(OutType dest, const InType &in, Executor &&exec, i { MATX_NVTX_START("var_impl(" + get_type_str(in) + ")", matx::MATX_NVTX_LOG_API) matxMemorySpace_t space; - using inner_type = typename inner_op_type_t::type; + using inner_type = typename inner_op_type_t::type; if constexpr (is_cuda_executor_v) { space = MATX_ASYNC_DEVICE_MEMORY; @@ -2550,7 +2550,7 @@ void __MATX_INLINE__ var_impl(OutType dest, const InType &in, Executor &&exec, i space = MATX_HOST_MALLOC_MEMORY; } - auto mean_tns = make_tensor(dest.Descriptor(), space); + auto mean_tns = make_tensor(dest.Descriptor(), space); mean_impl(mean_tns, in, exec); // need to clone along right most dims diff --git a/include/matx/transforms/resample_poly.h b/include/matx/transforms/resample_poly.h index 2e2b3a90e..a3443a8b3 100644 --- a/include/matx/transforms/resample_poly.h +++ b/include/matx/transforms/resample_poly.h @@ -53,9 +53,9 @@ inline void matxResamplePoly1DInternal(OutType &o, const InType &i, #ifdef __CUDACC__ MATX_NVTX_START("", matx::MATX_NVTX_LOG_INTERNAL) - using input_t = typename InType::scalar_type; - using filter_t = typename FilterType::scalar_type; - using output_t = typename OutType::scalar_type; + using input_t = typename InType::value_type; + using filter_t = typename FilterType::value_type; + using output_t = typename OutType::value_type; using shape_type = typename OutType::shape_type; // Even-length filters will be prepended with a single 0 to make them odd-length diff --git a/include/matx/transforms/solver.h b/include/matx/transforms/solver.h index ae69df79d..06cfcd079 100644 --- a/include/matx/transforms/solver.h +++ b/include/matx/transforms/solver.h @@ -120,7 +120,7 @@ class matxDnSolver_t { */ template static inline auto - TransposeCopy(typename TensorType::scalar_type *tp, const TensorType &a, cudaStream_t stream = 0) + TransposeCopy(typename TensorType::value_type *tp, const TensorType &a, cudaStream_t stream = 0) { MATX_NVTX_START("", matx::MATX_NVTX_LOG_INTERNAL) @@ -183,7 +183,7 @@ template class matxDnCholSolverPlan_t : public matxDnSolver_t { using OutTensor_t = remove_cvref_t; static_assert(OutTensor_t::Rank() == remove_cvref_t::Rank(), "Cholesky input/output tensor ranks must match"); - using T1 = typename OutTensor_t::scalar_type; + using T1 = typename OutTensor_t::value_type; static constexpr int RANK = OutTensor_t::Rank(); public: @@ -337,9 +337,9 @@ template class matxDnLUSolverPlan_t : public matxDnSolver_t { using OutTensor_t = remove_cvref_t; static constexpr int RANK = OutTensor_t::Rank(); - using T1 = typename OutTensor_t::scalar_type; + using T1 = typename OutTensor_t::value_type; static_assert(RANK-1 == PivotTensor::Rank(), "Pivot tensor rank must be one less than output"); - static_assert(std::is_same_v, "Pivot tensor type must be int64_t"); + static_assert(std::is_same_v, "Pivot tensor type must be int64_t"); public: /** @@ -508,7 +508,7 @@ struct DnQRParams_t { template class matxDnQRSolverPlan_t : public matxDnSolver_t { using out_type_t = remove_cvref_t; - using T1 = typename out_type_t::scalar_type; + using T1 = typename out_type_t::value_type; static constexpr int RANK = out_type_t::Rank(); static_assert(out_type_t::Rank()-1 == TauTensor::Rank(), "Tau tensor must be one rank less than output tensor"); static_assert(out_type_t::Rank() == ATensor::Rank(), "Output tensor must match A tensor rank in SVD"); @@ -690,10 +690,10 @@ struct DnSVDParams_t { template class matxDnSVDSolverPlan_t : public matxDnSolver_t { - using T1 = typename ATensor::scalar_type; - using T2 = typename UTensor::scalar_type; - using T3 = typename STensor::scalar_type; - using T4 = typename VTensor::scalar_type; + using T1 = typename ATensor::value_type; + using T2 = typename UTensor::value_type; + using T3 = typename STensor::value_type; + using T4 = typename VTensor::value_type; static constexpr int RANK = UTensor::Rank(); static_assert(UTensor::Rank()-1 == STensor::Rank(), "S tensor must be 1 rank lower than U tensor in SVD"); static_assert(UTensor::Rank() == ATensor::Rank(), "U tensor must match A tensor rank in SVD"); @@ -907,8 +907,8 @@ struct DnEigParams_t { template class matxDnEigSolverPlan_t : public matxDnSolver_t { public: - using T2 = typename WTensor::scalar_type; - using T1 = typename ATensor::scalar_type; + using T2 = typename WTensor::value_type; + using T1 = typename ATensor::value_type; static constexpr int RANK = remove_cvref_t::Rank(); static_assert(RANK == ATensor::Rank(), "Output and A tensor ranks must match for eigen solver"); static_assert(RANK-1 == WTensor::Rank(), "W tensor must be one rank lower than output for eigen solver"); @@ -1115,7 +1115,7 @@ void chol_impl(OutputTensor &&out, const ATensor &a, MATX_NVTX_START("", matx::MATX_NVTX_LOG_API) using OutputTensor_t = remove_cvref_t; - using T1 = typename OutputTensor_t::scalar_type; + using T1 = typename OutputTensor_t::value_type; auto a_new = OpToTensor(a, stream); @@ -1197,7 +1197,7 @@ void lu_impl(OutputTensor &&out, PivotTensor &&piv, { MATX_NVTX_START("", matx::MATX_NVTX_LOG_API) - using T1 = typename remove_cvref_t::scalar_type; + using T1 = typename remove_cvref_t::value_type; auto piv_new = OpToTensor(piv, stream); auto a_new = OpToTensor(a, stream); @@ -1291,7 +1291,7 @@ void det_impl(OutputTensor &out, const InputTensor &a, s[RANK - 2] = cuda::std::min(a_new.Size(RANK - 1), a_new.Size(RANK - 2)); auto piv = make_tensor(s, MATX_ASYNC_DEVICE_MEMORY, stream); - auto ac = make_tensor(a_new.Shape(), MATX_ASYNC_DEVICE_MEMORY, stream); + auto ac = make_tensor(a_new.Shape(), MATX_ASYNC_DEVICE_MEMORY, stream); lu_impl(ac, piv, a_new, stream); (out = prod(diag(ac))).run(stream); @@ -1327,7 +1327,7 @@ void cusolver_qr_impl(OutTensor &&out, TauTensor &&tau, { MATX_NVTX_START("", matx::MATX_NVTX_LOG_API) - using T1 = typename remove_cvref_t::scalar_type; + using T1 = typename remove_cvref_t::value_type; auto tau_new = OpToTensor(tau, stream); auto a_new = OpToTensor(a, stream); @@ -1413,7 +1413,7 @@ void svd_impl(UTensor &&u, STensor &&s, { MATX_NVTX_START("", matx::MATX_NVTX_LOG_API) - using T1 = typename ATensor::scalar_type; + using T1 = typename ATensor::value_type; auto u_new = OpToTensor(u, stream); auto s_new = OpToTensor(s, stream); @@ -1508,7 +1508,7 @@ void eig_impl(OutputTensor &&out, WTensor &&w, to transpose in and out of the function. Eventually this may be fixed in cuSolver. */ - using T1 = typename remove_cvref_t::scalar_type; + using T1 = typename remove_cvref_t::value_type; auto w_new = OpToTensor(w, stream); auto a_new = OpToTensor(a, stream); diff --git a/include/matx/transforms/svd.h b/include/matx/transforms/svd.h index 2168e8891..f8fc05f53 100644 --- a/include/matx/transforms/svd.h +++ b/include/matx/transforms/svd.h @@ -83,8 +83,8 @@ void svdpi_impl(UType &U, SType &S, VTType &VT, AType &A, X0Type &x0, int iterat static_assert(VTType::Rank() == AType::Rank()); static_assert(SType::Rank() == AType::Rank()-1); - using ATypeS = typename AType::scalar_type; - using STypeS = typename SType::scalar_type; + using ATypeS = typename AType::value_type; + using STypeS = typename SType::value_type; const int RANK = AType::Rank(); auto m = A.Size(RANK-2); // rows @@ -305,7 +305,7 @@ void svdpi_impl(UType &U, SType &S, VTType &VT, AType &A, X0Type &x0, int iterat template inline auto svdbpi_impl_workspace(const AType &A, cudaStream_t stream) { - using ATypeS = typename AType::scalar_type; + using ATypeS = typename AType::value_type; const int RANK = AType::Rank(); auto m = A.Size(RANK-2); // rows @@ -380,7 +380,7 @@ inline void svdbpi_impl(UType &U, SType &S, VTType &VT, const AType &A, int max_ static_assert(VTType::Rank() == AType::Rank()); static_assert(SType::Rank() == AType::Rank()-1); - using STypeS = typename SType::scalar_type; + using STypeS = typename SType::value_type; const int RANK = AType::Rank(); auto m = A.Size(RANK-2); // rows diff --git a/include/matx/transforms/transpose.h b/include/matx/transforms/transpose.h index a75ac89de..7940bcba3 100644 --- a/include/matx/transforms/transpose.h +++ b/include/matx/transforms/transpose.h @@ -78,7 +78,7 @@ namespace matx } #ifdef __CUDACC__ - size_t shm = sizeof(typename OutputTensor::scalar_type) * TILE_DIM * (TILE_DIM + 1); + size_t shm = sizeof(typename OutputTensor::value_type) * TILE_DIM * (TILE_DIM + 1); if constexpr (RANK == 2) { dim3 block(TILE_DIM, TILE_DIM); diff --git a/test/00_operators/OperatorTests.cu b/test/00_operators/OperatorTests.cu index f1e3d0aa2..28053ba82 100644 --- a/test/00_operators/OperatorTests.cu +++ b/test/00_operators/OperatorTests.cu @@ -1945,6 +1945,46 @@ TYPED_TEST(OperatorTestsFloatNonComplexAllExecs, NDOperatorFuncs) MATX_EXIT_HANDLER(); } +TYPED_TEST(OperatorTestsFloatAllExecs, Toeplitz) +{ + MATX_ENTER_HANDLER(); + using TestType = cuda::std::tuple_element_t<0, TypeParam>; + using ExecType = cuda::std::tuple_element_t<1, TypeParam>; + + auto pb = std::make_unique(); + pb->InitAndRunTVGenerator("00_operators", "toeplitz", "run", {5, 5, 6}); + + ExecType exec{}; + auto out1 = make_tensor({5, 5}); + auto out2 = make_tensor({5, 5}); + auto out3 = make_tensor({5, 6}); + auto c = make_tensor({5}); + auto r = make_tensor({5}); + auto r2 = make_tensor({6}); + + pb->NumpyToTensorView(c, "c"); + pb->NumpyToTensorView(r, "r"); + pb->NumpyToTensorView(r2, "r2"); + + // example-begin toeplitz-test-1 + (out1 = toeplitz(c)).run(exec); + // example-end toeplitz-test-1 + exec.sync(); + MATX_TEST_ASSERT_COMPARE(pb, out1, "out1", 0.01); + + // example-begin toeplitz-test-2 + (out2 = toeplitz(c, r)).run(exec); + // example-end toeplitz-test-2 + exec.sync(); + MATX_TEST_ASSERT_COMPARE(pb, out2, "out2", 0.01); + + (out3 = toeplitz(c, r2)).run(exec); + exec.sync(); + MATX_TEST_ASSERT_COMPARE(pb, out3, "out3", 0.01); + + MATX_EXIT_HANDLER(); +} + TYPED_TEST(OperatorTestsNumericNonComplexAllExecs, OperatorFuncs) { MATX_ENTER_HANDLER(); diff --git a/test/00_solver/SVD.cu b/test/00_solver/SVD.cu index 3720b16d4..1f82c5bd1 100644 --- a/test/00_solver/SVD.cu +++ b/test/00_solver/SVD.cu @@ -68,14 +68,14 @@ TYPED_TEST(SVDSolverTestNonHalfTypes, SVDBasic) using TestType = cuda::std::tuple_element_t<0, TypeParam>; using ExecType = cuda::std::tuple_element_t<1, TypeParam>; - using scalar_type = typename inner_op_type_t::type; + using value_type = typename inner_op_type_t::type; tensor_t Av{{m, n}}; tensor_t Atv{{n, m}}; - tensor_t Sv{{std::min(m, n)}}; + tensor_t Sv{{std::min(m, n)}}; tensor_t Uv{{m, m}}; tensor_t Vv{{n, n}}; - tensor_t Sav{{m, n}}; + tensor_t Sav{{m, n}}; tensor_t SSolav{{m, n}}; tensor_t Uav{{m, m}}; tensor_t Vav{{n, n}}; @@ -149,18 +149,18 @@ TYPED_TEST(SVDSolverTestNonHalfTypes, SVDBasicBatched) constexpr index_t batches = 10; - using scalar_type = typename inner_op_type_t::type; + using value_type = typename inner_op_type_t::type; auto Av1 = make_tensor({m, n}); this->pb->NumpyToTensorView(Av1, "A"); auto Av = make_tensor({batches, m, n}); auto Atv = make_tensor({batches, n, m}); (Av = Av1).run(this->exec); - auto Sv = make_tensor({batches, std::min(m, n)}); + auto Sv = make_tensor({batches, std::min(m, n)}); auto Uv = make_tensor({batches, m, m}); auto Vv = make_tensor({batches, n, n}); - auto Sav = make_tensor({batches, m, n}); + auto Sav = make_tensor({batches, m, n}); auto SSolav = make_tensor({batches, m, n}); auto Uav = make_tensor({batches, m, m}); auto Vav = make_tensor({batches, n, n}); diff --git a/test/test_vectors/generators/00_operators.py b/test/test_vectors/generators/00_operators.py index e5e38cff8..897c65812 100644 --- a/test/test_vectors/generators/00_operators.py +++ b/test/test_vectors/generators/00_operators.py @@ -2,7 +2,9 @@ import numpy as np import scipy.signal as ss +import scipy.linalg as sl from typing import Dict, List +import matx_common class polyval_operator: def __init__(self, dtype: str, size: List[int]): @@ -104,6 +106,30 @@ def run(self) -> Dict[str, np.array]: 'c_float3d': c1 } +class toeplitz: + def __init__(self, dtype: str, size: List[int]): + self.size = size + self.dtype = dtype + pass + + def run(self) -> Dict[str, np.array]: + c = matx_common.randn_ndarray((self.size[0],), self.dtype) + r = matx_common.randn_ndarray((self.size[1],), self.dtype) + r2 = matx_common.randn_ndarray((self.size[2],), self.dtype) + + # The first element in each array must match + r[0] = c[0] + r2[0] = c[0] + + return { + 'c': c, + 'r': r, + 'r2': r2, + 'out1': sl.toeplitz(c), + 'out2': sl.toeplitz(c, r), + 'out3': sl.toeplitz(c, r2) + } + class pwelch_operators: def __init__(self, dtype: str, cfg: Dict): #PWelchGeneratorCfg): self.dtype = dtype