Skip to content

Commit

Permalink
Added toeplitz operator
Browse files Browse the repository at this point in the history
  • Loading branch information
@cliffburdick
cliffburdick committed Jul 24, 2024
1 parent 4fd3986 commit 4487d1e
Show file tree
Hide file tree
Showing 143 changed files with 640 additions and 575 deletions.
2 changes: 1 addition & 1 deletion include/matx/core/file_io.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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<typename TensorType::scalar_type>());
"dtype"_a = detail::MatXPybind::GetNumpyDtype<typename TensorType::value_type>());
pb->NumpyToTensorView(t, obj);
}

Expand Down
6 changes: 2 additions & 4 deletions include/matx/core/iterator.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -47,8 +47,7 @@ namespace matx {
template <typename OperatorType, bool ConvertType = true>
struct RandomOperatorIterator {
using self_type = RandomOperatorIterator<OperatorType, ConvertType>;
using value_type = typename std::conditional_t<ConvertType, detail::convert_matx_type_t<typename OperatorType::scalar_type>, typename OperatorType::scalar_type>;
using scalar_type = value_type;
using value_type = typename std::conditional_t<ConvertType, detail::convert_matx_type_t<typename OperatorType::value_type>, typename OperatorType::value_type>;
// using stride_type = std::conditional_t<is_tensor_view_v<OperatorType>, typename OperatorType::desc_type::stride_type,
// index_t>;
using stride_type = index_t;
Expand Down Expand Up @@ -174,8 +173,7 @@ __MATX_INLINE__ __MATX_HOST__ __MATX_DEVICE__ index_t operator-(const RandomOper
template <typename OperatorType, bool ConvertType = true>
struct RandomOperatorOutputIterator {
using self_type = RandomOperatorOutputIterator<OperatorType, ConvertType>;
using value_type = typename std::conditional_t<ConvertType, detail::convert_matx_type_t<typename OperatorType::scalar_type>, typename OperatorType::scalar_type>;
using scalar_type = value_type;
using value_type = typename std::conditional_t<ConvertType, detail::convert_matx_type_t<typename OperatorType::value_type>, typename OperatorType::value_type>;
// using stride_type = std::conditional_t<is_tensor_view_v<OperatorType>, typename OperatorType::desc_type::stride_type,
// index_t>;
using stride_type = index_t;
Expand Down
30 changes: 15 additions & 15 deletions include/matx/core/make_tensor.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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<typename TensorType::scalar_type, TensorType::Rank()>(shape, space, stream);
auto tmp = make_tensor<typename TensorType::value_type, TensorType::Rank()>(shape, space, stream);
tensor.Shallow(tmp);
}

Expand Down Expand Up @@ -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<typename TensorType::scalar_type, ShapeType>(std::forward<ShapeType>(shape), space, stream);
auto tmp = make_tensor<typename TensorType::value_type, ShapeType>(std::forward<ShapeType>(shape), space, stream);
tensor.Shallow(tmp);
}

Expand Down Expand Up @@ -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<typename TensorType::scalar_type, typename TensorType::shape_container>(std::forward<typename TensorType::shape_container>(shape), space, stream);
auto tmp = make_tensor<typename TensorType::value_type, typename TensorType::shape_container>(std::forward<typename TensorType::shape_container>(shape), space, stream);
tensor.Shallow(tmp);
}

Expand Down Expand Up @@ -261,7 +261,7 @@ template <typename TensorType,
auto make_tensor( TensorType &tensor,
matxMemorySpace_t space = MATX_MANAGED_MEMORY,
cudaStream_t stream = 0) {
auto tmp = make_tensor<typename TensorType::scalar_type>({}, space, stream);
auto tmp = make_tensor<typename TensorType::value_type>({}, space, stream);
tensor.Shallow(tmp);
}

Expand Down Expand Up @@ -322,12 +322,12 @@ auto make_tensor( T *data,
template <typename TensorType,
std::enable_if_t<is_tensor_view_v<TensorType>, 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<typename TensorType::scalar_type, TensorType::Rank()>(data, shape, owning);
auto tmp = make_tensor<typename TensorType::value_type, TensorType::Rank()>(data, shape, owning);
tensor.Shallow(tmp);
}

Expand Down Expand Up @@ -373,12 +373,12 @@ auto make_tensor( T *data,
template <typename TensorType,
std::enable_if_t<is_tensor_view_v<TensorType>, 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<typename TensorType::scalar_type, typename TensorType::shape_container>(data, std::forward<typename TensorType::shape_container>(shape), owning);
auto tmp = make_tensor<typename TensorType::value_type, typename TensorType::shape_container>(data, std::forward<typename TensorType::shape_container>(shape), owning);
tensor.Shallow(tmp);
}

Expand Down Expand Up @@ -414,9 +414,9 @@ auto make_tensor( T *ptr,
template <typename TensorType,
std::enable_if_t<is_tensor_view_v<TensorType>, bool> = true>
auto make_tensor( TensorType &tensor,
typename TensorType::scalar_type *ptr,
typename TensorType::value_type *ptr,
bool owning = false) {
auto tmp = make_tensor<typename TensorType::scalar_type>(ptr, owning);
auto tmp = make_tensor<typename TensorType::value_type>(ptr, owning);
tensor.Shallow(tmp);
}

Expand Down Expand Up @@ -534,12 +534,12 @@ auto make_tensor( T* const data,
template <typename TensorType,
std::enable_if_t<is_tensor_view_v<TensorType>, 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<typename TensorType::scalar_type, typename TensorType::desc_type>(data, std::forward<typename TensorType::desc_type>(desc), owning);
auto tmp = make_tensor<typename TensorType::value_type, typename TensorType::desc_type>(data, std::forward<typename TensorType::desc_type>(desc), owning);
tensor.Shallow(tmp);
}

Expand Down Expand Up @@ -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<typename TensorType::scalar_type, typename TensorType::desc_type>(std::forward<typename TensorType::desc_type>(desc), space, stream);
auto tmp = make_tensor<typename TensorType::value_type, typename TensorType::desc_type>(std::forward<typename TensorType::desc_type>(desc), space, stream);
tensor.Shallow(tmp);
}

Expand Down Expand Up @@ -633,13 +633,13 @@ auto make_tensor( T *const data,
template <typename TensorType,
std::enable_if_t<is_tensor_view_v<TensorType>, 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<typename TensorType::scalar_type, TensorType::Rank()>(data, shape, strides, owning);
auto tmp = make_tensor<typename TensorType::value_type, TensorType::Rank()>(data, shape, strides, owning);
tensor.Shallow(tmp);
}

Expand Down
12 changes: 6 additions & 6 deletions include/matx/core/operator_utils.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -44,13 +44,13 @@ namespace matx {
if (out.IsContiguous()) {
if constexpr(ConvertType) {
return func( in,
reinterpret_cast<detail::convert_matx_type_t<typename remove_cvref_t<OutputOp>::scalar_type> *>(out.Data()),
reinterpret_cast<detail::convert_matx_type_t<typename remove_cvref_t<OutputOp>::value_type> *>(out.Data()),
bi,
ei);
}
else {
return func( in,
reinterpret_cast<typename remove_cvref_t<OutputOp>::scalar_type *>(out.Data()),
reinterpret_cast<typename remove_cvref_t<OutputOp>::value_type *>(out.Data()),
bi,
ei);
}
Expand All @@ -70,14 +70,14 @@ namespace matx {
if constexpr (ConvertType) {
return ReduceOutput<ConvertType>( std::forward<Func>(func),
std::forward<OutputOp>(out),
reinterpret_cast<detail::convert_matx_type_t<typename remove_cvref_t<InputOp>::scalar_type> *>(in_base.Data()),
reinterpret_cast<detail::convert_matx_type_t<typename remove_cvref_t<InputOp>::value_type> *>(in_base.Data()),
BeginOffset{in_base},
EndOffset{in_base});
}
else {
return ReduceOutput<ConvertType>( std::forward<Func>(func),
std::forward<OutputOp>(out),
reinterpret_cast<typename remove_cvref_t<InputOp>::scalar_type *>(in_base.Data()),
reinterpret_cast<typename remove_cvref_t<InputOp>::value_type *>(in_base.Data()),
BeginOffset{in_base},
EndOffset{in_base});
}
Expand Down Expand Up @@ -118,9 +118,9 @@ namespace matx {
namespace detail {
// Used inside of transforms to allocate temporary output
template <typename TensorType, typename Executor, typename ShapeType>
__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<index_t>(1),
std::multiplies<index_t>()) * sizeof(typename TensorType::scalar_type);
std::multiplies<index_t>()) * sizeof(typename TensorType::value_type);
if constexpr (is_cuda_executor_v<Executor>) {
matxAlloc((void**)ptr, ttl_size, MATX_ASYNC_DEVICE_MEMORY, ex.getStream());
make_tensor(tensor, *ptr, shape);
Expand Down
12 changes: 6 additions & 6 deletions include/matx/core/pybind.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -176,7 +176,7 @@ class MatXPybind {
template <typename TensorType>
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;

Expand Down Expand Up @@ -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.");

Expand Down Expand Up @@ -377,7 +377,7 @@ class MatXPybind {
template <typename TensorType>
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<T>;
auto ften = pybind11::array_t<ntype, pybind11::array::c_style | pybind11::array::forcecast>(np_ten);
Expand All @@ -398,7 +398,7 @@ class MatXPybind {

template <typename TensorType>
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<tensor_type>;
constexpr int RANK = TensorType::Rank();

Expand Down Expand Up @@ -466,12 +466,12 @@ class MatXPybind {


template <typename TensorType,
typename CT = matx_convert_cuda_complex_type<typename TensorType::scalar_type>>
typename CT = matx_convert_cuda_complex_type<typename TensorType::value_type>>
std::optional<TestFailResult<CT>>
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<raw_type>;
using ctype = matx_convert_cuda_complex_type<raw_type>;
auto resobj = res_dict[fname.c_str()];
Expand Down
2 changes: 1 addition & 1 deletion include/matx/core/tensor.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -87,7 +87,7 @@ class tensor_t : public detail::tensor_impl_t<T,RANK,Desc> {
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
Expand Down
1 change: 0 additions & 1 deletion include/matx/core/tensor_impl.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down
10 changes: 5 additions & 5 deletions include/matx/core/tensor_utils.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -842,7 +842,7 @@ namespace detail {
PrintKernel<<<1, 1>>>(op, dims...);
}
else {
auto tmpv = make_tensor<typename Op::scalar_type>(op.Shape());
auto tmpv = make_tensor<typename Op::value_type>(op.Shape());
(tmpv = op).run();
PrintData(fp, tmpv, dims...);
}
Expand Down Expand Up @@ -911,7 +911,7 @@ void PrintData(FILE* fp, const Op &op, Args... dims) {
}
}
else {
auto tmpv = make_tensor<typename Op::scalar_type>(op.Shape());
auto tmpv = make_tensor<typename Op::value_type>(op.Shape());
(tmpv = op).run();
cudaStreamSynchronize(0);
InternalPrint(fp, tmpv, dims...);
Expand Down Expand Up @@ -962,7 +962,7 @@ void PrintData(FILE* fp, const Op &op, Args... dims) {
// PrintKernel<<<1, 1>>>(op, dims...);
// }
// else {
// auto tmpv = make_tensor<typename Op::scalar_type>(op.Shape());
// auto tmpv = make_tensor<typename Op::value_type>(op.Shape());
// (tmpv = op).run();
// PrintData(tmpv, dims...);
// }
Expand Down Expand Up @@ -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<Op>) ? "Tensor" : "Operator";

fprintf(fp, "%s{%s} Rank: %d, Sizes:[", type.c_str(), detail::GetTensorType<typename Op::scalar_type>().c_str(), op.Rank());
fprintf(fp, "%s{%s} Rank: %d, Sizes:[", type.c_str(), detail::GetTensorType<typename Op::value_type>().c_str(), op.Rank());

for (index_t dimIdx = 0; dimIdx < (op.Rank() ); dimIdx++ )
{
Expand Down Expand Up @@ -1110,7 +1110,7 @@ void print(const Op &op)
template <typename Op>
auto OpToTensor(Op &&op, [[maybe_unused]] cudaStream_t stream) {
if constexpr (!is_tensor_view_v<Op>) {
return make_tensor<typename remove_cvref<Op>::scalar_type>(op.Shape(), MATX_ASYNC_DEVICE_MEMORY, stream);
return make_tensor<typename remove_cvref<Op>::value_type>(op.Shape(), MATX_ASYNC_DEVICE_MEMORY, stream);
} else {
return op;
}
Expand Down
5 changes: 2 additions & 3 deletions include/matx/core/tie.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -43,8 +43,7 @@ namespace matx {
template <typename... Ts>
struct mtie : public BaseOp<mtie<Ts...>>{
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;

Expand Down Expand Up @@ -97,7 +96,7 @@ struct mtie : public BaseOp<mtie<Ts...>>{
// 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<sizeof...(Ts) - 1>(ts_).InnerPreRun(NoShape{}, std::forward<Executor>(ex));
cuda::std::get<sizeof...(Ts) - 1>(ts_).InnerPreRun(detail::NoShape{}, std::forward<Executor>(ex));
}
cuda::std::get<sizeof...(Ts) - 1>(ts_).Exec(ts_, std::forward<Executor>(ex));
}
Expand Down
20 changes: 11 additions & 9 deletions include/matx/core/type_utils.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -62,9 +62,11 @@ enum class MemoryLayout {
MEMORY_LAYOUT_COL_MAJOR,
};

struct NoShape{};

namespace detail {
struct NoShape{};
struct EmptyOp{};

template <typename T>
struct is_noshape : std::integral_constant<bool, std::is_same_v<NoShape, T>> {};
};
Expand Down Expand Up @@ -688,23 +690,23 @@ template <class T>
inline constexpr bool is_matx_type_v = detail::is_matx_type<T>::value;

namespace detail {
template <typename T, typename = void> struct extract_scalar_type_impl {
using scalar_type = T;
template <typename T, typename = void> struct extract_value_type_impl {
using value_type = T;
};

template <typename T>
struct extract_scalar_type_impl<T, std::void_t<typename T::scalar_type>> {
using scalar_type = typename T::scalar_type;
struct extract_value_type_impl<T, typename std::enable_if_t<is_matx_op<T>()>> {
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 <typename T>
using extract_scalar_type_t = typename detail::extract_scalar_type_impl<T>::scalar_type;
using extract_value_type_t = typename detail::extract_value_type_impl<T>::value_type;

/**
* @brief Promote half precision floating point value to fp32, or leave untouched if not half
Expand Down Expand Up @@ -779,7 +781,7 @@ struct base_type {

template <typename T>
struct base_type<T, typename std::enable_if_t<is_tensor_view_v<T>>> {
using type = tensor_impl_t<typename T::scalar_type, T::Rank(), typename T::desc_type>;
using type = tensor_impl_t<typename T::value_type, T::Rank(), typename T::desc_type>;
};

template <typename T> using base_type_t = typename base_type<typename remove_cvref<T>::type>::type;
Expand Down Expand Up @@ -903,7 +905,7 @@ __MATX_INLINE__ __MATX_DEVICE__ __MATX_HOST__ auto select_tuple(Tuple&& tuple, s
template <typename... T, std::enable_if_t<((is_tensor_view_v<T>) && ...), bool> = true>
constexpr bool TensorTypesMatch() {
using first_type = cuda::std::tuple_element_t<0, cuda::std::tuple<T...>>;
return ((std::is_same_v<typename first_type::scalar_type, typename T::scalar_type>) && ...);
return ((std::is_same_v<typename first_type::value_type, typename T::value_type>) && ...);
}

struct no_permute_t{};
Expand Down
2 changes: 1 addition & 1 deletion include/matx/generators/alternate.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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"; }
Expand Down
Loading

0 comments on commit 4487d1e

Please sign in to comment.