From 303181924c560ab97c5e276397ce236d0891029e Mon Sep 17 00:00:00 2001 From: Tim Shen Date: Thu, 19 Nov 2020 14:41:02 -0800 Subject: [PATCH] [XLA/GPU] Migrate nested reduce emitter to take LMHLO. PiperOrigin-RevId: 343373088 --- include/mlir-hlo/Dialect/mhlo/IR/lhlo_ops.td | 9 +++++---- 1 file changed, 5 insertions(+), 4 deletions(-) diff --git a/include/mlir-hlo/Dialect/mhlo/IR/lhlo_ops.td b/include/mlir-hlo/Dialect/mhlo/IR/lhlo_ops.td index 0ee44f3..ff052b5 100644 --- a/include/mlir-hlo/Dialect/mhlo/IR/lhlo_ops.td +++ b/include/mlir-hlo/Dialect/mhlo/IR/lhlo_ops.td @@ -197,10 +197,11 @@ def LHLO_XorOp : LHLO_BinaryElementwiseOp<"xor", LHLO_PredOrIntBuffer>, BASE_HLO //===----------------------------------------------------------------------===// // TODO(b/139813999): specify required function signature in a type-safe way. -def LHLO_ReduceOp: LHLO_Op<"reduce", [ - SameVariadicOperandSize, - SingleBlockImplicitTerminator<"TerminatorOp"> - ]>, BASE_HLO_ReduceOp { +// +// 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 { let arguments = (ins Arg, "", [MemRead]>:$operands, Arg, "", [MemRead]>:$init_values,