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

Function complement

include/cute/layout.hpp:1178–1227  ·  view source on GitHub ↗

Source from the content-addressed store, hash-verified

1176// @pre @a layout has been filtered (flattened and no stride-0 or size-1 modes).
1177template <class Shape, class Stride, class CoTarget>
1178CUTE_HOST_DEVICE constexpr
1179auto
1180complement(Shape const& shape, Stride const& stride, CoTarget const& cotarget)
1181{
1182 if constexpr (is_constant<0, Stride>::value) {
1183 // Special case for irreducible rank-1 stride-0 layout
1184 return make_layout(coalesce(cotarget));
1185 } else {
1186 // General case
1187 constexpr int R = rank_v<Shape>;
1188 static_assert(R == 1 || is_static<Stride>::value,
1189 "Dynamic-stride complement only for rank-1 layouts");
1190
1191 // Should just be a sort and a fold...
1192 // Then we could even handle dynamic strides (but they would destroy all static strides)
1193 auto [shape_, stride_, result_shape_, result_stride] =
1194 fold(make_seq<R-1>{},
1195 cute::make_tuple(shape, stride, cute::make_tuple(), cute::make_tuple(Int<1>{})),
1196 [](auto const& init, auto i)
1197 {
1198 auto [shape, stride, result_shape, result_stride] = init;
1199 auto min_stride = cute::min(stride);
1200 auto min_idx = cute::find(stride, min_stride);
1201 auto new_shape = min_stride / get<i>(result_stride);
1202 auto new_stride = min_stride * get<min_idx>(shape);
1203 static_assert(not is_constant<0, decltype(new_shape)>::value, "Non-injective Layout detected in complement.");
1204
1205 return cute::make_tuple(remove<min_idx>(shape), // Remove the min_idx from shape
1206 remove<min_idx>(stride), // Remove the min_idx from stride
1207 append(result_shape , new_shape ), // new shape = min_stride / last_stride
1208 append(result_stride, new_stride)); // new stride = min_stride * curr_shape
1209 });
1210
1211 // Append the last shape mode
1212 auto new_shape = get<0>(stride_) / get<R-1>(result_stride); // new shape = min_stride / last_stride
1213 static_assert(not is_constant<0, decltype(new_shape)>::value, "Non-injective Layout detected in complement.");
1214 auto result_shape = append(result_shape_, new_shape);
1215
1216 // Compute the rest_shape and rest_stride
1217 auto new_stride = get<0>(stride_) * get<0>(shape_); // new stride = min_stride * curr_shape
1218 auto rest_shape = coalesce(ceil_div(cotarget, new_stride));
1219 auto rest_stride = compact_major<LayoutLeft>(rest_shape, new_stride);
1220
1221 // Coalesce and append (rest_shape, rest_stride)
1222 return coalesce(make_layout(make_shape (result_shape , rest_shape ),
1223 make_stride(result_stride, rest_stride)));
1224 }
1225
1226 CUTE_GCC_UNREACHABLE;
1227}
1228
1229} // end namespace detail
1230

Callers 6

logical_divideFunction · 0.70
ceil_divFunction · 0.70
logical_productFunction · 0.70
get_layoutC_TVMethod · 0.50
get_layoutA_TVMethod · 0.50
get_layoutB_TVMethod · 0.50

Calls 12

foldFunction · 0.85
make_layoutFunction · 0.70
coalesceFunction · 0.70
minFunction · 0.70
appendFunction · 0.70
ceil_divFunction · 0.70
filterFunction · 0.70
shapeFunction · 0.70
cosizeFunction · 0.70
findFunction · 0.50
shapeMethod · 0.45
strideMethod · 0.45

Tested by

no test coverage detected