diff --git a/include/mlir-hlo/Dialect/mhlo/IR/hlo_ops.td b/include/mlir-hlo/Dialect/mhlo/IR/hlo_ops.td index 2f08797..77b5159 100644 --- a/include/mlir-hlo/Dialect/mhlo/IR/hlo_ops.td +++ b/include/mlir-hlo/Dialect/mhlo/IR/hlo_ops.td @@ -401,12 +401,18 @@ def HLO_InfeedOp : HLO_Op<"infeed", []> { of the data. Multiple Infeed operations are allowed in a computation, but there must be a total order among the Infeed operations. + Attributes: + layout: Array attribute. Same shape as the output of the infeed, except + that every tensor is replaced by a minor_to_major array for the + tensor's layout. + See https://www.tensorflow.org/xla/operation_semantics#infeed. }]; let arguments = (ins HLO_Token:$token, - DefaultValuedAttr:$infeed_config + DefaultValuedAttr:$infeed_config, + OptionalAttr:$layout ); let results = (outs HLO_Tuple); let hasCustomHLOConverter = 1; diff --git a/tests/ops.mlir b/tests/ops.mlir index 76d8896..5651f04 100644 --- a/tests/ops.mlir +++ b/tests/ops.mlir @@ -442,7 +442,7 @@ func @dot_bad_precision_config(%arg0: tensor<2x2xi32>, %arg1: tensor<2x2xi32>) - func @infeed_invalid_number_of_results(%token: !mhlo.token) -> tuple>, !mhlo.token, tensor> { // expected-error@+1 {{result is expected to be a tuple of size 2, but got 3}} - %0 = "mhlo.infeed"(%token) {infeed_config = "foobar"} : (!mhlo.token) -> tuple>, !mhlo.token, tensor> + %0 = "mhlo.infeed"(%token) {infeed_config = "foobar", layout = [[[0]], unit, [0]]} : (!mhlo.token) -> tuple>, !mhlo.token, tensor> return %0 : tuple>, !mhlo.token, tensor> } @@ -450,7 +450,7 @@ func @infeed_invalid_number_of_results(%token: !mhlo.token) -> tuple tuple>, tensor> { // expected-error@+1 {{second element of result tuple is expected to be of token type, but got 'tensor'}} - %0 = "mhlo.infeed"(%token) {infeed_config = "foobar"} : (!mhlo.token) -> tuple>, tensor> + %0 = "mhlo.infeed"(%token) {infeed_config = "foobar", layout = [[[0]], [0]]} : (!mhlo.token) -> tuple>, tensor> return %0 : tuple>, tensor> }