MCPcopy Create free account
hub / github.com/NVIDIA/cutlass / recast

Function recast

include/cute/tensor_impl.hpp:757–783  ·  view source on GitHub ↗

Source from the content-addressed store, hash-verified

755
756template <class NewType_, class Tensor>
757CUTE_HOST_DEVICE constexpr
758auto
759recast(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

Callers

nothing calls this directly

Calls 9

foldFunction · 0.85
make_tensorFunction · 0.70
flattenFunction · 0.70
minFunction · 0.70
transformFunction · 0.50
dataMethod · 0.45
layoutMethod · 0.45
shapeMethod · 0.45
strideMethod · 0.45

Tested by

no test coverage detected