| 54 | } |
| 55 | |
| 56 | XlaOp GetDiagonalMask(XlaOp x, int diagonal) { |
| 57 | XlaBuilder* builder = x.builder(); |
| 58 | return builder->ReportErrorOrReturn([&]() -> StatusOr<XlaOp> { |
| 59 | TF_ASSIGN_OR_RETURN(Shape shape, builder->GetShape(x)); |
| 60 | auto n_dims = static_cast<int32>(shape.rank()); |
| 61 | TF_RET_CHECK(n_dims >= 2); |
| 62 | auto m = shape.dimensions(n_dims - 2); |
| 63 | auto n = shape.dimensions(n_dims - 1); |
| 64 | absl::Span<const int64> major_dims = |
| 65 | AsInt64Slice(shape.dimensions()).subspan(/*pos=*/0, /*len=*/n_dims - 2); |
| 66 | auto a = Iota(builder, S32, n); |
| 67 | auto b = Iota(builder, S32, m) + ConstantR0WithType(builder, S32, diagonal); |
| 68 | auto indicator = Eq(b, Broadcast(a, {m}), /*broadcast_dimensions=*/{0}); |
| 69 | auto mask = Broadcast(indicator, major_dims); |
| 70 | return mask; |
| 71 | }); |
| 72 | } |
| 73 | |
| 74 | XlaOp GetMatrixDiagonal(XlaOp x, int k) { |
| 75 | XlaBuilder* builder = x.builder(); |
no test coverage detected