From 5da9190dd92020d74c2ce133b07a5c8e7d899a36 Mon Sep 17 00:00:00 2001 From: "A. Unique TensorFlower" Date: Wed, 16 Dec 2020 20:29:15 -0800 Subject: [PATCH] Integrate LLVM at llvm/llvm-project@0cf7e4b252fe Updates LLVM usage to match [0cf7e4b252fe](https://github.com/llvm/llvm-project/commit/0cf7e4b252fe) PiperOrigin-RevId: 347948887 --- BUILD | 5 +++-- WORKSPACE | 4 ++-- build_tools/llvm_version.txt | 2 +- .../mhlo/transforms/hlo_legalize_to_lhlo.cc | 7 +++++-- .../mhlo/transforms/legalize_control_flow.cc | 10 ++++++---- .../mhlo/transforms/mhlo_control_flow_to_scf.cc | 3 ++- tests/end2end/broadcast.mlir | 3 ++- tests/hlo-legalize-to-lhlo.mlir | 14 +++++++------- tests/legalize-control-flow.mlir | 10 +++++----- tests/legalize_to_scf.mlir | 6 +++--- 10 files changed, 36 insertions(+), 28 deletions(-) diff --git a/BUILD b/BUILD index fb3b89e..a6317ea 100644 --- a/BUILD +++ b/BUILD @@ -580,7 +580,7 @@ cc_library( "@llvm-project//mlir:SCFDialect", "@llvm-project//mlir:StandardOps", "@llvm-project//mlir:Support", - "@llvm-project//mlir:Transforms", + "@llvm-project//mlir:TensorDialect", ], ) @@ -750,6 +750,7 @@ cc_library( "@llvm-project//mlir:StandardOps", "@llvm-project//mlir:StandardOpsTransforms", "@llvm-project//mlir:Support", + "@llvm-project//mlir:TensorDialect", "@llvm-project//mlir:Transforms", ], alwayslink = 1, @@ -811,11 +812,11 @@ cc_library( deps = [ ":hlo", "@llvm-project//llvm:Support", - "@llvm-project//mlir:Analysis", "@llvm-project//mlir:IR", "@llvm-project//mlir:Pass", "@llvm-project//mlir:StandardOps", "@llvm-project//mlir:Support", + "@llvm-project//mlir:TensorDialect", ], alwayslink = 1, ) diff --git a/WORKSPACE b/WORKSPACE index 7ab34f1..067fcb9 100644 --- a/WORKSPACE +++ b/WORKSPACE @@ -15,9 +15,9 @@ load("@bazel_tools//tools/build_defs/repo:http.bzl", "http_archive") -LLVM_COMMIT = "7aeb3804c46cc6c8f291415ca09ae34021301eb8" +LLVM_COMMIT = "0cf7e4b252fe1458fddb8e3dbfcae43450e9c04c" -LLVM_SHA256 = "945718051c37e55fc3aa79f4752dbcb22240ed19bc8037d41b4d461acd34c000" +LLVM_SHA256 = "6f663272226c5d36a213a308865af67ea58bcf7986c1f5551d48bf7eed83edfc" LLVM_BAZEL_TAG = "llvm-project-{commit}".format(commit = LLVM_COMMIT) diff --git a/build_tools/llvm_version.txt b/build_tools/llvm_version.txt index f0fde1f..12cf86b 100644 --- a/build_tools/llvm_version.txt +++ b/build_tools/llvm_version.txt @@ -1,2 +1,2 @@ -7aeb3804c46cc6c8f291415ca09ae34021301eb8 +0cf7e4b252fe1458fddb8e3dbfcae43450e9c04c diff --git a/lib/Dialect/mhlo/transforms/hlo_legalize_to_lhlo.cc b/lib/Dialect/mhlo/transforms/hlo_legalize_to_lhlo.cc index 11f9159..0a864cc 100644 --- a/lib/Dialect/mhlo/transforms/hlo_legalize_to_lhlo.cc +++ b/lib/Dialect/mhlo/transforms/hlo_legalize_to_lhlo.cc @@ -24,6 +24,7 @@ limitations under the License. #include "mlir/Dialect/Shape/Transforms/Passes.h" #include "mlir/Dialect/StandardOps/IR/Ops.h" #include "mlir/Dialect/StandardOps/Transforms/FuncConversions.h" +#include "mlir/Dialect/Tensor/IR/Tensor.h" #include "mlir/IR/AffineMap.h" #include "mlir/IR/Attributes.h" #include "mlir/IR/BlockAndValueMapping.h" @@ -37,6 +38,7 @@ limitations under the License. #include "mlir/Pass/Pass.h" #include "mlir/Transforms/Bufferize.h" #include "mlir/Transforms/DialectConversion.h" +#include "third_party/llvm/llvm-project/mlir/include/mlir/Dialect/Tensor/IR/Tensor.h" namespace mlir { namespace mhlo { @@ -62,7 +64,7 @@ Value InsertDynamicAllocAndDealloc(Location loc, Value result, if (shape_element.value() != ShapedType::kDynamicSize) continue; Value index = rewriter->create(loc, shape_element.index()); Value alloc_operand = - rewriter->create(loc, shape_operand, index); + rewriter->create(loc, shape_operand, index); if (!alloc_operand.getType().isIndex()) { alloc_operand = rewriter->create(loc, alloc_operand, rewriter->getIndexType()); @@ -292,7 +294,7 @@ class HloToLhloDynamicBroadcastInDimOpConverter for (int i = 0; i < result_rank; ++i) { Value i_val = b->create(loc, i); Value result_dim_size = - b->create(loc, op.output_dimensions(), i_val); + b->create(loc, op.output_dimensions(), i_val); if (!result_dim_size.getType().isIndex()) { result_dim_size = b->create(loc, result_dim_size, b->getIndexType()); @@ -567,6 +569,7 @@ struct HloLegalizeToLhlo ConversionTarget target(context); target.addLegalDialect(); target.addLegalDialect(); + target.addLegalDialect(); target.addIllegalOp(); target.addIllegalOp(); target.addIllegalDialect(); diff --git a/lib/Dialect/mhlo/transforms/legalize_control_flow.cc b/lib/Dialect/mhlo/transforms/legalize_control_flow.cc index 53472cb..253bd96 100644 --- a/lib/Dialect/mhlo/transforms/legalize_control_flow.cc +++ b/lib/Dialect/mhlo/transforms/legalize_control_flow.cc @@ -31,6 +31,7 @@ limitations under the License. #include "mlir/Pass/Pass.h" #include "mlir/Pass/PassRegistry.h" #include "mlir/Support/LogicalResult.h" +#include "third_party/llvm/llvm-project/mlir/include/mlir/Dialect/Tensor/IR/Tensor.h" namespace mlir { namespace mhlo { @@ -83,7 +84,7 @@ LogicalResult LowerIfOp(mlir::mhlo::IfOp if_op) { // Extract the predicate for checking branching, then branch to the true and // false regions appropriately. - auto cond_value = builder.create(loc, if_op.pred()); + auto cond_value = builder.create(loc, if_op.pred()); builder.create(loc, cond_value, true_block, if_op.true_arg(), false_block, if_op.false_arg()); @@ -142,7 +143,7 @@ LogicalResult LowerWhileOp(mlir::mhlo::WhileOp while_op) { builder.create(loc, cond_block, while_op.getOperand()); // Updates the inlined condition blocks by replacing the return op with an - // extract_element and conditional branch. This changes the block below: + // tensor.extract and conditional branch. This changes the block below: // ^cond(%0): // // "mhlo".return(%1) @@ -150,7 +151,7 @@ LogicalResult LowerWhileOp(mlir::mhlo::WhileOp while_op) { // Into: // ^cond(%0): // - // %2 = extract_element %1[] : tensor // Extract the condition value. + // %2 = tensor.extract %1[] : tensor // Extract the condition value. // cond_br %2, ^body(%0), ^tail(%0) // Branch. builder.setInsertionPointToStart(cond_block); @@ -166,7 +167,8 @@ LogicalResult LowerWhileOp(mlir::mhlo::WhileOp while_op) { builder.setInsertionPointToEnd(new_block); auto return_value = return_op.getOperand(0); - auto cond_value = builder.create(loc, return_value); + auto cond_value = + builder.create(loc, return_value); // Get the body block arguments. llvm::SmallVector successor_args(cond_block->args_begin(), diff --git a/lib/Dialect/mhlo/transforms/mhlo_control_flow_to_scf.cc b/lib/Dialect/mhlo/transforms/mhlo_control_flow_to_scf.cc index b79624d..517b7ce 100644 --- a/lib/Dialect/mhlo/transforms/mhlo_control_flow_to_scf.cc +++ b/lib/Dialect/mhlo/transforms/mhlo_control_flow_to_scf.cc @@ -24,6 +24,7 @@ limitations under the License. #include "mlir/IR/Value.h" #include "mlir/Pass/Pass.h" #include "mlir/Support/LLVM.h" +#include "third_party/llvm/llvm-project/mlir/include/mlir/Dialect/Tensor/IR/Tensor.h" #define DEBUG_TYPE "mhlo-control-flow-to-scf" @@ -119,7 +120,7 @@ void MatchAndRewrite(WhileOp whileOp) { auto tensorIndexType = RankedTensorType::get({}, b.getIndexType()); auto getAsIndex = [&](Value val) { auto loc = whileOp.getLoc(); - return b.create( + return b.create( loc, b.create(loc, tensorIndexType, val), ValueRange()); }; diff --git a/tests/end2end/broadcast.mlir b/tests/end2end/broadcast.mlir index 544efb8..bd8dfc4 100644 --- a/tests/end2end/broadcast.mlir +++ b/tests/end2end/broadcast.mlir @@ -1,5 +1,6 @@ // RUN: mlir-hlo-opt %s -chlo-legalize-to-hlo -hlo-legalize-to-lhlo \ -// RUN: -std-bufferize --canonicalize -buffer-hoisting -buffer-deallocation \ +// RUN: -std-bufferize -tensor-bufferize -finalizing-bufferize \ +// RUN: --canonicalize -buffer-hoisting -buffer-deallocation \ // RUN: -copy-removal -canonicalize -cse -lhlo-legalize-to-linalg \ // RUN: -lhlo-fuse-linalg -convert-linalg-to-loops -canonicalize -cse \ // RUN: -convert-linalg-to-llvm -convert-std-to-llvm \ diff --git a/tests/hlo-legalize-to-lhlo.mlir b/tests/hlo-legalize-to-lhlo.mlir index 5c05d5e..8cfffb3 100644 --- a/tests/hlo-legalize-to-lhlo.mlir +++ b/tests/hlo-legalize-to-lhlo.mlir @@ -177,16 +177,16 @@ func @dyn_broadcast(%operand: memref) -> index { // CHECK: %[[OP_STRIDE_0:.*]] = muli %[[C1]], %[[OPER_DIM_1]] : index // CHECK: %[[OPER_DIM_0:.*]] = dim %[[OPERAND]], %[[C0]] : memref -// CHECK: %[[EL0:.*]] = extract_element %[[SHAPE]]{{\[}}%[[C0]]] : tensor<3xi64> +// CHECK: %[[EL0:.*]] = tensor.extract %[[SHAPE]]{{\[}}%[[C0]]] : tensor<3xi64> // CHECK: %[[SIZE_0:.*]] = index_cast %[[EL0]] : i64 to index -// CHECK: %[[EL1:.*]] = extract_element %[[SHAPE]]{{\[}}%[[C1]]] : tensor<3xi64> +// CHECK: %[[EL1:.*]] = tensor.extract %[[SHAPE]]{{\[}}%[[C1]]] : tensor<3xi64> // CHECK: %[[SIZE_1:.*]] = index_cast %[[EL1]] : i64 to index // CHECK: %[[EXPAND_1:.*]] = cmpi "slt", %[[OPER_DIM_0]], %[[SIZE_1]] : index // CHECK: %[[STRIDE_1:.*]] = select %[[EXPAND_1]], %[[C0]], %[[OP_STRIDE_0]] : index // CHECK: %[[C2:.*]] = constant 2 : index -// CHECK: %[[EL2:.*]] = extract_element %[[SHAPE]]{{\[}}%[[C2]]] : tensor<3xi64> +// CHECK: %[[EL2:.*]] = tensor.extract %[[SHAPE]]{{\[}}%[[C2]]] : tensor<3xi64> // CHECK: %[[SIZE_2:.*]] = index_cast %[[EL2]] : i64 to index // CHECK: %[[EXPAND_2:.*]] = cmpi "slt", %[[OPER_DIM_1]], %[[SIZE_2]] : index // CHECK: %[[STRIDE_2:.*]] = select %[[EXPAND_2]], %[[C0]], %[[C1]] : index @@ -554,9 +554,9 @@ func @add_dyn(%lhs: tensor, %rhs: tensor) { // CHECK: %[[DIM1:.*]] = dim %arg0, %[[C1]] : memref // CHECK: %[[IC1:.*]] = index_cast %[[DIM1]] : index to i64 // CHECK: %[[SHAPE:.*]] = tensor_from_elements %[[IC0]], %[[IC1]] : tensor<2xi64> - // CHECK: %[[EE0:.*]] = extract_element %[[SHAPE]][%[[C0]]] : tensor<2xi64> + // CHECK: %[[EE0:.*]] = tensor.extract %[[SHAPE]][%[[C0]]] : tensor<2xi64> // CHECK: %[[ICS0:.*]] = index_cast %[[EE0]] : i64 to index - // CHECK: %[[EE1:.*]] = extract_element %[[SHAPE]][%[[C1]]] : tensor<2xi64> + // CHECK: %[[EE1:.*]] = tensor.extract %[[SHAPE]][%[[C1]]] : tensor<2xi64> // CHECK: %[[ICS1:.*]] = index_cast %[[EE1]] : i64 to index // CHECK: %[[RESULT:.*]] = alloc(%[[ICS0]], %[[ICS1]]) // CHECK: "lmhlo.add"(%arg0, %arg1, %[[RESULT]]) : (memref, memref, memref) -> () @@ -577,9 +577,9 @@ func @tanh_dyn(%arg0: tensor) { // CHECK: %[[DIM1:.*]] = dim %arg0, %[[C1]] : memref // CHECK: %[[IC1:.*]] = index_cast %[[DIM1]] : index to i64 // CHECK: %[[SHAPE:.*]] = tensor_from_elements %[[IC0]], %[[IC1]] : tensor<2xi64> - // CHECK: %[[EE0:.*]] = extract_element %[[SHAPE]][%[[C0]]] : tensor<2xi64> + // CHECK: %[[EE0:.*]] = tensor.extract %[[SHAPE]][%[[C0]]] : tensor<2xi64> // CHECK: %[[ICS0:.*]] = index_cast %[[EE0]] : i64 to index - // CHECK: %[[EE1:.*]] = extract_element %[[SHAPE]][%[[C1]]] : tensor<2xi64> + // CHECK: %[[EE1:.*]] = tensor.extract %[[SHAPE]][%[[C1]]] : tensor<2xi64> // CHECK: %[[ICS1:.*]] = index_cast %[[EE1]] : i64 to index // CHECK: %[[RESULT:.*]] = alloc(%[[ICS0]], %[[ICS1]]) // CHECK: "lmhlo.tanh"(%arg0, %[[RESULT]]) : (memref, memref) -> () diff --git a/tests/legalize-control-flow.mlir b/tests/legalize-control-flow.mlir index 274792e..8e5e18a 100644 --- a/tests/legalize-control-flow.mlir +++ b/tests/legalize-control-flow.mlir @@ -5,7 +5,7 @@ func @while(%arg0: tensor) -> tensor { //CHECK: br ^bb1(%arg0 : tensor) //CHECK: ^bb1([[VAL0:%.+]]: tensor): //CHECK: [[VAL1:%.+]] = "mhlo.compare"([[VAL0]], [[VAL0]]) - //CHECK: [[VAL2:%.+]] = extract_element [[VAL1]][] : tensor + //CHECK: [[VAL2:%.+]] = tensor.extract [[VAL1]][] : tensor //CHECK: cond_br [[VAL2]], ^bb2([[VAL0]] : tensor), ^bb3([[VAL0]] : tensor) //CHECK: ^bb2([[VAL3:%.+]]: tensor): //CHECK: [[VAL4:%.+]] = mhlo.add [[VAL3]], [[VAL3]] @@ -33,7 +33,7 @@ func @conditional(%arg0: tensor) -> tensor { // CHECK: [[VAL0:%.+]] = "mhlo.compare"(%arg0, [[C0]]) {comparison_direction = "LT"} : (tensor, tensor) -> tensor %0 = "mhlo.compare"(%arg0, %cst) {comparison_direction = "LT"} : (tensor, tensor) -> tensor - // CHECK: [[VAL1:%.+]] = extract_element [[VAL0]][] : tensor + // CHECK: [[VAL1:%.+]] = tensor.extract [[VAL0]][] : tensor // CHECK: cond_br [[VAL1]], ^bb1(%arg0 : tensor), ^bb2(%arg0 : tensor) %1 = "mhlo.if"(%0, %arg0, %arg0) ( { @@ -63,7 +63,7 @@ func @while_with_multiple_blocks_in_body(%arg0: tensor) -> tensor { // CHECK: br ^[[COND_ENTRY:.+]](%arg0 : tensor) // CHECK: ^[[COND_ENTRY]](%0: tensor): // CHECK: %1 = "mhlo.compare"(%0, %0) {comparison_direction = "LT"} : (tensor, tensor) -> tensor - // CHECK: %2 = extract_element %1[] : tensor + // CHECK: %2 = tensor.extract %1[] : tensor // CHECK: cond_br %2, ^[[BODY_ENTRY:.+]](%0 : tensor), ^[[EXIT:.+]](%0 : tensor) // CHECK: ^[[BODY_ENTRY]](%3: tensor): // CHECK: br ^[[BODY_SUCC:.+]](%3 : tensor) @@ -95,7 +95,7 @@ func @while_with_multiple_blocks_in_cond(%arg0: tensor) -> tensor { // CHECK: br ^[[COND_SUCC:.+]](%0 : tensor) // CHECK: ^[[COND_SUCC]](%1: tensor): // CHECK: %2 = "mhlo.compare"(%1, %1) {comparison_direction = "LT"} : (tensor, tensor) -> tensor - // CHECK: %3 = extract_element %2[] : tensor + // CHECK: %3 = tensor.extract %2[] : tensor // CHECK: cond_br %3, ^[[BODY_ENTRY:.+]](%0 : tensor), ^[[EXIT:.+]](%0 : tensor) // CHECK: ^[[BODY_ENTRY]](%4: tensor): // CHECK: br ^[[COND_ENTRY]](%4 : tensor) @@ -118,7 +118,7 @@ func @while_with_multiple_blocks_in_cond(%arg0: tensor) -> tensor { // CHECK-LABEL: func @conditional_with_multiple_blocks(%arg0: tensor, %arg1: tensor, %arg2: tensor) -> tensor { func @conditional_with_multiple_blocks(%arg0: tensor, %arg1: tensor, %pred: tensor) -> tensor { - // CHECK: %0 = extract_element %arg2[] : tensor + // CHECK: %0 = tensor.extract %arg2[] : tensor // CHECK: cond_br %0, ^[[THEN_ENTRY:.+]](%arg0 : tensor), ^[[ELSE_ENTRY:.+]](%arg1 : tensor) // CHECK: ^[[THEN_ENTRY]](%1: tensor): // CHECK: br ^[[THEN_SUCC:.+]](%1 : tensor) diff --git a/tests/legalize_to_scf.mlir b/tests/legalize_to_scf.mlir index 9c887a7..101800d 100644 --- a/tests/legalize_to_scf.mlir +++ b/tests/legalize_to_scf.mlir @@ -30,9 +30,9 @@ func @lt_loop(%arg0: tensor<4xf32>, %arg1: tensor, %arg2: tensor, %arg // CHECK: %[[VAL_11:.*]] = constant dense<0> : tensor // CHECK: %[[VAL_12:.*]] = constant dense<1000> : tensor // CHECK: %[[VAL_14:.*]] = index_cast %[[VAL_11]] : tensor to tensor -// CHECK: %[[VAL_15:.*]] = extract_element %[[VAL_14]][] : tensor +// CHECK: %[[VAL_15:.*]] = tensor.extract %[[VAL_14]][] : tensor // CHECK: %[[VAL_16:.*]] = index_cast %[[VAL_12]] : tensor to tensor -// CHECK: %[[VAL_17:.*]] = extract_element %[[VAL_16]][] : tensor +// CHECK: %[[VAL_17:.*]] = tensor.extract %[[VAL_16]][] : tensor // CHECK: %[[VAL_18:.*]] = index_cast %[[VAL_10]] : tensor to tensor -// CHECK: %[[VAL_19:.*]] = extract_element %[[VAL_18]][] : tensor +// CHECK: %[[VAL_19:.*]] = tensor.extract %[[VAL_18]][] : tensor // CHECK: scf.for %[[VAL_21:.*]] = %[[VAL_15]] to %[[VAL_17]] step %[[VAL_19]] iter_args(%[[VAL_22:.*]] = %[[VAL_9]], %[[VAL_23:.*]] = %[[VAL_12]])