| 755 | |
| 756 | template <class NewType_, class Tensor> |
| 757 | CUTE_HOST_DEVICE constexpr |
| 758 | auto |
| 759 | recast(Tensor&& tensor) |
| 760 | { |
| 761 | using OldType = typename remove_cvref_t<Tensor>::element_type; |
| 762 | using NewType = copy_cv_t<OldType, NewType_>; |
| 763 | |
| 764 | if constexpr (is_same<NewType, OldType>::value) { |
| 765 | return make_tensor(static_cast<Tensor&&>(tensor).data(), tensor.layout()); |
| 766 | } else { |
| 767 | auto old_layout = tensor.layout(); |
| 768 | auto new_layout = recast_layout<OldType,NewType>(old_layout); |
| 769 | |
| 770 | // If this is an upcast of a normal Layout with static negative strides, then offset as well |
| 771 | if constexpr (sizeof(OldType) < sizeof(NewType) && not is_composed_layout<decltype(old_layout)>::value) { |
| 772 | auto shape_diff = transform(flatten(old_layout.shape()), flatten(new_layout.shape()), minus{}); |
| 773 | auto extent_diff = transform(shape_diff, flatten(old_layout.stride()), multiplies{}); |
| 774 | auto offset = fold(extent_diff, Int<0>{}, [](auto const& i, auto const& a) { return i + cute::min(a,Int<0>{}); }); |
| 775 | |
| 776 | return make_tensor(recast_ptr<NewType>(static_cast<Tensor&&>(tensor).data() + offset), new_layout); |
| 777 | } else { |
| 778 | return make_tensor(recast_ptr<NewType>(static_cast<Tensor&&>(tensor).data() ), new_layout); |
| 779 | } |
| 780 | } |
| 781 | |
| 782 | CUTE_GCC_UNREACHABLE; |
| 783 | } |
| 784 | |
| 785 | // |
| 786 | // max_common_vector |