From bae7c7bd2b4b59e49aa357362b4c37c568c6ef8b Mon Sep 17 00:00:00 2001 From: "A. Unique TensorFlower" Date: Thu, 19 Nov 2020 15:58:53 -0800 Subject: [PATCH] [XLA/GPU] Migrate nested reduce emitter to take LMHLO. PiperOrigin-RevId: 343388053 --- include/mlir-hlo/Dialect/mhlo/IR/lhlo_ops.td | 9 ++++----- 1 file changed, 4 insertions(+), 5 deletions(-) diff --git a/include/mlir-hlo/Dialect/mhlo/IR/lhlo_ops.td b/include/mlir-hlo/Dialect/mhlo/IR/lhlo_ops.td index ff052b5..0ee44f3 100644 --- a/include/mlir-hlo/Dialect/mhlo/IR/lhlo_ops.td +++ b/include/mlir-hlo/Dialect/mhlo/IR/lhlo_ops.td @@ -197,11 +197,10 @@ def LHLO_XorOp : LHLO_BinaryElementwiseOp<"xor", LHLO_PredOrIntBuffer>, BASE_HLO //===----------------------------------------------------------------------===// // TODO(b/139813999): specify required function signature in a type-safe way. -// -// The region `body` may return lmhlo.TerminatorOp or mhlo.ReturnOp. We are -// moving towards mhlo.ReturnOp, but some code that needs cleanup still assumes lmhlo.TerminatorOp. -// TODO(timshen): cleanup lmhlo.TerminatorOp. -def LHLO_ReduceOp: LHLO_Op<"reduce", [SameVariadicOperandSize]>, BASE_HLO_ReduceOp { +def LHLO_ReduceOp: LHLO_Op<"reduce", [ + SameVariadicOperandSize, + SingleBlockImplicitTerminator<"TerminatorOp"> + ]>, BASE_HLO_ReduceOp { let arguments = (ins Arg, "", [MemRead]>:$operands, Arg, "", [MemRead]>:$init_values,