| 1176 | // @pre @a layout has been filtered (flattened and no stride-0 or size-1 modes). |
| 1177 | template <class Shape, class Stride, class CoTarget> |
| 1178 | CUTE_HOST_DEVICE constexpr |
| 1179 | auto |
| 1180 | complement(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 |
no test coverage detected