experiments/STT-tensorflow
1
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

Integrate LLVM at llvm/llvm-project@0cf7e4b252

Browse Source 
Updates LLVM usage to match
[0cf7e4b252fe](https://github.com/llvm/llvm-project/commit/0cf7e4b252fe)

PiperOrigin-RevId: 347948887
Change-Id: I87697e08bcfc29dbf259e523eb481366b4795537
This commit is contained in:
A. Unique TensorFlower 2020-12-16 20:29:15 -08:00 committed by TensorFlower Gardener
parent a87bf6e6fe
commit 173dd7b6d5
17 changed files with 435 additions and 54 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

5
tensorflow/compiler/mlir/hlo/BUILD
View File

@ -570,7 +570,7 @@ cc_library(
"@llvm-project//mlir:SCFDialect",
"@llvm-project//mlir:StandardOps",
"@llvm-project//mlir:Support",
"@llvm-project//mlir:Transforms",
"@llvm-project//mlir:TensorDialect",
],
)
@ -740,6 +740,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,
@ -809,11 +810,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,
)

7
tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/hlo_legalize_to_lhlo.cc
View File

@ -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 "mlir/Dialect/Tensor/IR/Tensor.h" // from @llvm-project
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<ConstantIndexOp>(loc, shape_element.index());
Value alloc_operand =
rewriter->create<ExtractElementOp>(loc, shape_operand, index);
rewriter->create<tensor::ExtractOp>(loc, shape_operand, index);
if (!alloc_operand.getType().isIndex()) {
alloc_operand = rewriter->create<IndexCastOp>(loc, alloc_operand,
rewriter->getIndexType());
@ -292,7 +294,7 @@ class HloToLhloDynamicBroadcastInDimOpConverter
for (int i = 0; i < result_rank; ++i) {
Value i_val = b->create<ConstantIndexOp>(loc, i);
Value result_dim_size =
b->create<ExtractElementOp>(loc, op.output_dimensions(), i_val);
b->create<tensor::ExtractOp>(loc, op.output_dimensions(), i_val);
if (!result_dim_size.getType().isIndex()) {
result_dim_size =
b->create<IndexCastOp>(loc, result_dim_size, b->getIndexType());
@ -567,6 +569,7 @@ struct HloLegalizeToLhlo
ConversionTarget target(context);
target.addLegalDialect<lmhlo::LmhloDialect>();
target.addLegalDialect<StandardOpsDialect>();
target.addLegalDialect<tensor::TensorDialect>();
target.addIllegalOp<mlir::TensorLoadOp>();
target.addIllegalOp<mlir::TensorStoreOp>();
target.addIllegalDialect<mhlo::MhloDialect>();

10
tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/legalize_control_flow.cc
View File

@ -31,6 +31,7 @@ limitations under the License.
#include "mlir/Pass/Pass.h"
#include "mlir/Pass/PassRegistry.h"
#include "mlir/Support/LogicalResult.h"
#include "mlir/Dialect/Tensor/IR/Tensor.h" // from @llvm-project
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<mlir::ExtractElementOp>(loc, if_op.pred());
auto cond_value = builder.create<mlir::tensor::ExtractOp>(loc, if_op.pred());
builder.create<mlir::CondBranchOp>(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<mlir::BranchOp>(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):
// <inlined conditional region>
// "mhlo".return(%1)
@ -150,7 +151,7 @@ LogicalResult LowerWhileOp(mlir::mhlo::WhileOp while_op) {
// Into:
// ^cond(%0):
// <inlined conditional region>
// %2 = extract_element %1[] : tensor<i1> // Extract the condition value.
// %2 = tensor.extract %1[] : tensor<i1> // 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<mlir::ExtractElementOp>(loc, return_value);
auto cond_value =
builder.create<mlir::tensor::ExtractOp>(loc, return_value);
// Get the body block arguments.
llvm::SmallVector<Value, 4> successor_args(cond_block->args_begin(),

3
tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/mhlo_control_flow_to_scf.cc
View File

@ -24,6 +24,7 @@ limitations under the License.
#include "mlir/IR/Value.h"
#include "mlir/Pass/Pass.h"
#include "mlir/Support/LLVM.h"
#include "mlir/Dialect/Tensor/IR/Tensor.h" // from @llvm-project
#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<ExtractElementOp>(
return b.create<tensor::ExtractOp>(
loc, b.create<IndexCastOp>(loc, tensorIndexType, val), ValueRange());
};

14
tensorflow/compiler/mlir/hlo/tests/hlo-legalize-to-lhlo.mlir
View File

@ -177,16 +177,16 @@ func @dyn_broadcast(%operand: memref<?x?xf32>) -> index {
// CHECK: %[[OP_STRIDE_0:.*]] = muli %[[C1]], %[[OPER_DIM_1]] : index
// CHECK: %[[OPER_DIM_0:.*]] = dim %[[OPERAND]], %[[C0]] : memref<?x?xf32>
// 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<?x?xf32>, %rhs: tensor<?x?xf32>) {
// CHECK: %[[DIM1:.*]] = dim %arg0, %[[C1]] : memref<?x?xf32>
// 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<?x?xf32>, memref<?x?xf32>, memref<?x?xf32>) -> ()
@ -577,9 +577,9 @@ func @tanh_dyn(%arg0: tensor<?x?xf32>) {
// CHECK: %[[DIM1:.*]] = dim %arg0, %[[C1]] : memref<?x?xf32>
// 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<?x?xf32>, memref<?x?xf32>) -> ()

10
tensorflow/compiler/mlir/hlo/tests/legalize-control-flow.mlir
View File

@ -5,7 +5,7 @@ func @while(%arg0: tensor<i64>) -> tensor<i64> {
//CHECK: br ^bb1(%arg0 : tensor<i64>)
//CHECK: ^bb1([[VAL0:%.+]]: tensor<i64>):
//CHECK: [[VAL1:%.+]] = "mhlo.compare"([[VAL0]], [[VAL0]])
//CHECK: [[VAL2:%.+]] = extract_element [[VAL1]][] : tensor<i1>
//CHECK: [[VAL2:%.+]] = tensor.extract [[VAL1]][] : tensor<i1>
//CHECK: cond_br [[VAL2]], ^bb2([[VAL0]] : tensor<i64>), ^bb3([[VAL0]] : tensor<i64>)
//CHECK: ^bb2([[VAL3:%.+]]: tensor<i64>):
//CHECK: [[VAL4:%.+]] = mhlo.add [[VAL3]], [[VAL3]]
@ -33,7 +33,7 @@ func @conditional(%arg0: tensor<f32>) -> tensor<f32> {
// CHECK: [[VAL0:%.+]] = "mhlo.compare"(%arg0, [[C0]]) {comparison_direction = "LT"} : (tensor<f32>, tensor<f32>) -> tensor<i1>
%0 = "mhlo.compare"(%arg0, %cst) {comparison_direction = "LT"} : (tensor<f32>, tensor<f32>) -> tensor<i1>
// CHECK: [[VAL1:%.+]] = extract_element [[VAL0]][] : tensor<i1>
// CHECK: [[VAL1:%.+]] = tensor.extract [[VAL0]][] : tensor<i1>
// CHECK: cond_br [[VAL1]], ^bb1(%arg0 : tensor<f32>), ^bb2(%arg0 : tensor<f32>)
%1 = "mhlo.if"(%0, %arg0, %arg0) ( {
@ -63,7 +63,7 @@ func @while_with_multiple_blocks_in_body(%arg0: tensor<i64>) -> tensor<i64> {
// CHECK: br ^[[COND_ENTRY:.+]](%arg0 : tensor<i64>)
// CHECK: ^[[COND_ENTRY]](%0: tensor<i64>):
// CHECK: %1 = "mhlo.compare"(%0, %0) {comparison_direction = "LT"} : (tensor<i64>, tensor<i64>) -> tensor<i1>
// CHECK: %2 = extract_element %1[] : tensor<i1>
// CHECK: %2 = tensor.extract %1[] : tensor<i1>
// CHECK: cond_br %2, ^[[BODY_ENTRY:.+]](%0 : tensor<i64>), ^[[EXIT:.+]](%0 : tensor<i64>)
// CHECK: ^[[BODY_ENTRY]](%3: tensor<i64>):
// CHECK: br ^[[BODY_SUCC:.+]](%3 : tensor<i64>)
@ -95,7 +95,7 @@ func @while_with_multiple_blocks_in_cond(%arg0: tensor<i64>) -> tensor<i64> {
// CHECK: br ^[[COND_SUCC:.+]](%0 : tensor<i64>)
// CHECK: ^[[COND_SUCC]](%1: tensor<i64>):
// CHECK: %2 = "mhlo.compare"(%1, %1) {comparison_direction = "LT"} : (tensor<i64>, tensor<i64>) -> tensor<i1>
// CHECK: %3 = extract_element %2[] : tensor<i1>
// CHECK: %3 = tensor.extract %2[] : tensor<i1>
// CHECK: cond_br %3, ^[[BODY_ENTRY:.+]](%0 : tensor<i64>), ^[[EXIT:.+]](%0 : tensor<i64>)
// CHECK: ^[[BODY_ENTRY]](%4: tensor<i64>):
// CHECK: br ^[[COND_ENTRY]](%4 : tensor<i64>)
@ -118,7 +118,7 @@ func @while_with_multiple_blocks_in_cond(%arg0: tensor<i64>) -> tensor<i64> {
// CHECK-LABEL: func @conditional_with_multiple_blocks(%arg0: tensor<f32>, %arg1: tensor<f32>, %arg2: tensor<i1>) -> tensor<f32> {
func @conditional_with_multiple_blocks(%arg0: tensor<f32>, %arg1: tensor<f32>, %pred: tensor<i1>) -> tensor<f32> {
// CHECK: %0 = extract_element %arg2[] : tensor<i1>
// CHECK: %0 = tensor.extract %arg2[] : tensor<i1>
// CHECK: cond_br %0, ^[[THEN_ENTRY:.+]](%arg0 : tensor<f32>), ^[[ELSE_ENTRY:.+]](%arg1 : tensor<f32>)
// CHECK: ^[[THEN_ENTRY]](%1: tensor<f32>):
// CHECK: br ^[[THEN_SUCC:.+]](%1 : tensor<f32>)

6
tensorflow/compiler/mlir/hlo/tests/legalize_to_scf.mlir
View File

@ -30,9 +30,9 @@ func @lt_loop(%arg0: tensor<4xf32>, %arg1: tensor<f32>, %arg2: tensor<f32>, %arg
// CHECK: %[[VAL_11:.*]] = constant dense<0> : tensor<i32>
// CHECK: %[[VAL_12:.*]] = constant dense<1000> : tensor<i32>
// CHECK: %[[VAL_14:.*]] = index_cast %[[VAL_11]] : tensor<i32> to tensor<index>
// CHECK: %[[VAL_15:.*]] = extract_element %[[VAL_14]][] : tensor<index>
// CHECK: %[[VAL_15:.*]] = tensor.extract %[[VAL_14]][] : tensor<index>
// CHECK: %[[VAL_16:.*]] = index_cast %[[VAL_12]] : tensor<i32> to tensor<index>
// CHECK: %[[VAL_17:.*]] = extract_element %[[VAL_16]][] : tensor<index>
// CHECK: %[[VAL_17:.*]] = tensor.extract %[[VAL_16]][] : tensor<index>
// CHECK: %[[VAL_18:.*]] = index_cast %[[VAL_10]] : tensor<i32> to tensor<index>
// CHECK: %[[VAL_19:.*]] = extract_element %[[VAL_18]][] : tensor<index>
// CHECK: %[[VAL_19:.*]] = tensor.extract %[[VAL_18]][] : tensor<index>
// CHECK: scf.for %[[VAL_21:.*]] = %[[VAL_15]] to %[[VAL_17]] step %[[VAL_19]] iter_args(%[[VAL_22:.*]] = %[[VAL_9]], %[[VAL_23:.*]] = %[[VAL_12]])

1
tensorflow/compiler/mlir/tensorflow/BUILD
View File

@ -1011,6 +1011,7 @@ cc_library(
"@llvm-project//mlir:Rewrite",
"@llvm-project//mlir:StandardOps",
"@llvm-project//mlir:Support",
"@llvm-project//mlir:TensorDialect",
"@llvm-project//mlir:TransformUtils",
"@llvm-project//mlir:Transforms",
],

16
tensorflow/compiler/mlir/tensorflow/tests/functional-control-flow-to-cfg.mlir
View File

@ -11,7 +11,7 @@ func @testIf1Result(tensor<i1>, tensor<*xf32>, tensor<*xf32>) -> tensor<*xf32> {
} : (tensor<i1>, tensor<*xf32>, tensor<*xf32>) -> tensor<*xf32>
// CHECK: [[TOBOOL:%.+]] = "tf.ToBool"(%arg0) : (tensor<i1>) -> tensor<i1>
// CHECK: [[PRED:%.+]] = extract_element [[TOBOOL]][] : tensor<i1>
// CHECK: [[PRED:%.+]] = tensor.extract [[TOBOOL]][] : tensor<i1>
// CHECK: cond_br [[PRED]], ^bb1, ^bb2
// CHECK: ^bb1:
// CHECK: [[THEN:%.+]] = call @testIf1Then(%arg1, %arg2)
@ -36,7 +36,7 @@ func @testIf3Result(tensor<i1>, tensor<*xf32>) -> (tensor<*xf32>, tensor<*xi8>,
} : (tensor<i1>, tensor<*xf32>) -> (tensor<*xf32>, tensor<*xi8>, tensor<*xbf16>)
// CHECK: [[TOBOOL:%.+]] = "tf.ToBool"(%arg0) : (tensor<i1>) -> tensor<i1>
// CHECK: [[PRED:%.+]] = extract_element [[TOBOOL]][] : tensor<i1>
// CHECK: [[PRED:%.+]] = tensor.extract [[TOBOOL]][] : tensor<i1>
// CHECK: cond_br [[PRED]], ^bb1, ^bb2
// CHECK: ^bb1:
// CHECK: [[THEN:%.+]]:3 = call @testIf3Then(%arg1)
@ -65,7 +65,7 @@ func @testIfCasts(%arg0: tensor<i1>, %arg1: tensor<!tf.variant<tensor<f32>>>) ->
} : (tensor<i1>, tensor<!tf.variant<tensor<f32>>>) -> tensor<!tf.variant<tensor<f32>>>
return %0: tensor<!tf.variant<tensor<f32>>>
// CHECK: [[TOBOOL:%.+]] = "tf.ToBool"(%arg0) : (tensor<i1>) -> tensor<i1>
// CHECK: [[PRED:%.+]] = extract_element [[TOBOOL]][] : tensor<i1>
// CHECK: [[PRED:%.+]] = tensor.extract [[TOBOOL]][] : tensor<i1>
// CHECK: cond_br [[PRED]], ^bb1, ^bb2
// CHECK: ^bb1:
// CHECK: [[CAST0:%.+]] = "tf.Cast"(%arg1) {Truncate = false} : (tensor<!tf.variant<tensor<f32>>>) -> tensor<!tf.variant>
@ -93,7 +93,7 @@ func @testIf1x4(tensor<4xi1>, tensor<*xf32>, tensor<*xf32>) -> tensor<*xf32> {
^bb0(%arg0: tensor<4xi1>, %arg1: tensor<*xf32>, %arg2: tensor<*xf32>):
// CHECK: [[TOBOOL:%.+]] = "tf.ToBool"(%arg0) : (tensor<4xi1>) -> tensor<i1>
// CHECK: [[PRED:%.+]] = extract_element [[TOBOOL]][] : tensor<i1>
// CHECK: [[PRED:%.+]] = tensor.extract [[TOBOOL]][] : tensor<i1>
%1 = "tf.If"(%arg0, %arg1, %arg2) {
then_branch = @testIf1Then, else_branch = @testIf1Else, is_stateless = false
} : (tensor<4xi1>, tensor<*xf32>, tensor<*xf32>) -> tensor<*xf32>
@ -118,7 +118,7 @@ func @testWhile2Result(tensor<*xf32>, tensor<*xf32>) -> (tensor<*xf32>, tensor<*
// CHECK: ^bb1([[CONDARG0:%.+]]: tensor<*xf32>, [[CONDARG1:%.+]]: tensor<*xf32>):
// CHECK: [[CONTINUE:%.+]] = call @testWhile2Cond(%0, %1) : (tensor<*xf32>, tensor<*xf32>) -> tensor<i1>
// CHECK: [[TOBOOL:%.+]] = "tf.ToBool"([[CONTINUE]]) : (tensor<i1>) -> tensor<i1>
// CHECK: [[PRED:%.+]] = extract_element [[TOBOOL]][] : tensor<i1>
// CHECK: [[PRED:%.+]] = tensor.extract [[TOBOOL]][] : tensor<i1>
// CHECK: cond_br [[PRED]], ^bb2([[CONDARG0]], [[CONDARG1]] : tensor<*xf32>, tensor<*xf32>), ^bb3([[CONDARG0]], [[CONDARG1]] : tensor<*xf32>, tensor<*xf32>)
// CHECK: ^bb2([[BODYARG0:%.+]]: tensor<*xf32>, [[BODYARG1:%.+]]: tensor<*xf32>):
// CHECK: [[BODYRETS:%.+]]:2 = call @testWhile2Body([[BODYARG0]], [[BODYARG1]]) : (tensor<*xf32>, tensor<*xf32>) -> (tensor<*xf32>, tensor<*xf32>)
@ -142,7 +142,7 @@ func @testWhile0Result() {
// CHECK: ^bb1:
// CHECK: [[CONTINUE:%.+]] = call @testWhile0Cond() : () -> tensor<i1>
// CHECK: [[TOBOOL:%.+]] = "tf.ToBool"([[CONTINUE]]) : (tensor<i1>) -> tensor<i1>
// CHECK: [[PRED:%.+]] = extract_element [[TOBOOL]][] : tensor<i1>
// CHECK: [[PRED:%.+]] = tensor.extract [[TOBOOL]][] : tensor<i1>
// CHECK: cond_br [[PRED]], ^bb2, ^bb3
// CHECK: ^bb2:
// CHECK: call @testWhile0Body() : () -> ()
@ -166,7 +166,7 @@ func @testComplexWhile1Result(%arg0: tensor<*xf32>, %arg1: tensor<*xf32>) -> (te
// CHECK: ^bb1([[CONDARG0:%.+]]: tensor<*xf32>, [[CONDARG1:%.+]]: tensor<*xf32>):
// CHECK: [[CONTINUE:%.+]] = call @testWhile2Cond([[CONDARG0]], [[CONDARG1]]) : (tensor<*xf32>, tensor<*xf32>) -> tensor<i1>
// CHECK: [[TOBOOL:%.+]] = "tf.ToBool"([[CONTINUE]]) : (tensor<i1>) -> tensor<i1>
// CHECK: [[PRED:%.+]] = extract_element [[TOBOOL]][] : tensor<i1>
// CHECK: [[PRED:%.+]] = tensor.extract [[TOBOOL]][] : tensor<i1>
// CHECK: cond_br [[PRED]], ^bb2([[CONDARG0]], [[CONDARG1]] : tensor<*xf32>, tensor<*xf32>), ^bb3([[CONDARG0]], [[CONDARG1]] : tensor<*xf32>, tensor<*xf32>)
// CHECK: ^bb2([[BODYARG0:%.+]]: tensor<*xf32>, [[BODYARG1:%.+]]: tensor<*xf32>):
// CHECK: [[BODYRETS:%.+]]:2 = call @testWhile2Body([[BODYARG0]], [[BODYARG1]]) : (tensor<*xf32>, tensor<*xf32>) -> (tensor<*xf32>, tensor<*xf32>)
@ -206,7 +206,7 @@ func @testWhileCasts(%arg0: tensor<!tf.variant<tensor<1x3xf32>>>) -> (tensor<!tf
// CHECK: ^bb1([[CONDARG0:%.+]]: tensor<!tf.variant>): // 2 preds: ^bb0, ^bb2
// CHECK: [[CONTINUE:%.+]] = call @testWhileCond([[CONDARG0]]) : (tensor<!tf.variant>) -> tensor<i1>
// CHECK: [[TOBOOL:%.+]] = "tf.ToBool"([[CONTINUE]]) : (tensor<i1>) -> tensor<i1>
// CHECK: [[PRED:%.+]] = extract_element [[TOBOOL]][] : tensor<i1>
// CHECK: [[PRED:%.+]] = tensor.extract [[TOBOOL]][] : tensor<i1>
// CHECK: [[CASTCONDARG0:%.+]] = "tf.Cast"([[CONDARG0]]) {Truncate = false} : (tensor<!tf.variant>) -> tensor<!tf.variant<tensor<1x?xf32>>>
// CHECK: cond_br [[PRED]], ^bb2([[CASTCONDARG0]] : tensor<!tf.variant<tensor<1x?xf32>>>), ^bb3([[CASTCONDARG0]] : tensor<!tf.variant<tensor<1x?xf32>>>)
// CHECK: ^bb2([[BODYARG0:%.+]]: tensor<!tf.variant<tensor<1x?xf32>>>): // pred: ^bb1

3
tensorflow/compiler/mlir/tensorflow/transforms/functional_control_flow_to_cfg.cc
View File

@ -17,6 +17,7 @@ limitations under the License.
// TensorFlow dialect to MLIR Control Flow Graph (CFG) form.
#include "mlir/Dialect/StandardOps/IR/Ops.h" // from @llvm-project
#include "mlir/Dialect/Tensor/IR/Tensor.h" // from @llvm-project
#include "mlir/IR/Attributes.h" // from @llvm-project
#include "mlir/IR/Builders.h" // from @llvm-project
#include "mlir/IR/Operation.h" // from @llvm-project
@ -42,7 +43,7 @@ struct FunctionalControlFlowToCFG
// control flow op into an i1 value.
static Value LowerCondition(Location loc, Value value, OpBuilder* builder) {
auto zero_d = builder->create<ToBoolOp>(loc, value);
auto scalar = builder->create<ExtractElementOp>(loc, zero_d);
auto scalar = builder->create<tensor::ExtractOp>(loc, zero_d);
return scalar.getResult();
}

10
tensorflow/compiler/mlir/tools/kernel_gen/tests/bufferize.mlir
View File

@ -2,14 +2,14 @@
// RUN: kernel-gen-opt %s --func-bufferize --final-bufferize --promote-buffers-to-stack | FileCheck %s --check-prefixes=CHECK,ALLOCA
// CHECK-LABEL: @extract_element
// CHECK-LABEL: @tensor.extract
// CHECK-SAME: (%[[ARG:.*]]: memref<?xf32>) -> f32
func @extract_element(%arg : tensor<?xf32>) -> f32 {
func @tensor.extract(%arg : tensor<?xf32>) -> f32 {
// CHECK: %[[C0:.*]] = constant 0 : index
// CHECK: %[[RESULT:.*]] = load %[[ARG]][%[[C0]]]
// CHECK: return %[[RESULT]]
%c0 = constant 0 : index
%result = extract_element %arg[%c0] : tensor<?xf32>
%result = tensor.extract %arg[%c0] : tensor<?xf32>
return %result : f32
}
@ -30,7 +30,7 @@ func @tensor_from_elements(%a : f32) -> f32 {
%c = constant 2.3 : f32
%tfe = tensor_from_elements %a, %b, %c : tensor<3xf32>
%c0 = constant 0 : index
%result = extract_element %tfe[%c0] : tensor<3xf32>
%result = tensor.extract %tfe[%c0] : tensor<3xf32>
return %result : f32
}
@ -54,7 +54,7 @@ func @dynamic_tensor_from_elements(%arg : tensor<*xf32>) -> index {
yield %elem : index
} : tensor<?xindex>
%c0 = constant 0 : index
%result = extract_element %tfe[%c0] : tensor<?xindex>
%result = tensor.extract %tfe[%c0] : tensor<?xindex>
return %result : index
}

2
tensorflow/compiler/mlir/tools/kernel_gen/transforms/BUILD
View File

@ -87,6 +87,8 @@ cc_library(
copts = if_cuda_is_configured(["-DGOOGLE_CUDA=1"]) + if_rocm_is_configured(["-DTENSORFLOW_USE_ROCM=1"]),
deps = [
"@llvm-project//mlir:Affine",
"@llvm-project//mlir:TensorDialect",
"@llvm-project//mlir:TensorTransforms",
"//tensorflow/compiler/xla/service:hlo_module_config",
"//tensorflow/compiler/xla:debug_options_flags",
"//tensorflow/compiler/xla:statusor",

12
tensorflow/compiler/mlir/tools/kernel_gen/transforms/bufferize_pass.cc
View File

@ -30,6 +30,8 @@ limitations under the License.
#include "mlir/Dialect/StandardOps/IR/Ops.h" // from @llvm-project
#include "mlir/Dialect/StandardOps/Transforms/FuncConversions.h" // from @llvm-project
#include "mlir/Dialect/StandardOps/Transforms/Passes.h" // from @llvm-project
#include "mlir/Dialect/Tensor/IR/Tensor.h" // from @llvm-project
#include "mlir/Dialect/Tensor/Transforms/Passes.h" // from @llvm-project
#include "mlir/IR/BuiltinOps.h" // from @llvm-project
#include "mlir/IR/BuiltinTypes.h" // from @llvm-project
#include "mlir/IR/MLIRContext.h" // from @llvm-project
@ -109,7 +111,10 @@ struct HloBufferizePass : public HloBufferizePassBase<HloBufferizePass> {
OwningRewritePatternList patterns;
auto& context = getContext();
ConversionTarget target(context);
target.addLegalDialect<lmhlo::LmhloDialect>();
target.addLegalDialect<StandardOpsDialect>();
target.addLegalDialect<lmhlo::LmhloDialect, StandardOpsDialect>();
target.addLegalDialect<tensor::TensorDialect>();
target.addIllegalDialect<mhlo::MhloDialect>();
CustomBufferizeTypeConverter converter;
@ -149,7 +154,8 @@ struct FinalBufferizePass : public FinalBufferizePassBase<FinalBufferizePass> {
// TODO(b/173201243): Move to tablegen.
void getDependentDialects(DialectRegistry& registry) const override {
registry.insert<AffineDialect, scf::SCFDialect, shape::ShapeDialect,
tf_framework::TFFrameworkDialect, lmhlo::LmhloDialect>();
tensor::TensorDialect, tf_framework::TFFrameworkDialect,
lmhlo::LmhloDialect>();
}
public:
@ -157,13 +163,14 @@ struct FinalBufferizePass : public FinalBufferizePassBase<FinalBufferizePass> {
auto& context = getContext();
ConversionTarget target(context);
target.addLegalDialect<scf::SCFDialect, StandardOpsDialect,
tensor::TensorDialect,
tf_framework::TFFrameworkDialect, AffineDialect,
shape::ShapeDialect, lmhlo::LmhloDialect,
linalg::LinalgDialect>();
target.addLegalOp<FuncOp, ModuleOp, ModuleTerminatorOp>();
target.addIllegalDialect<mhlo::MhloDialect>();
target.addIllegalOp<DynamicTensorFromElementsOp, ExtractElementOp,
target.addIllegalOp<DynamicTensorFromElementsOp, tensor::ExtractOp,
TensorFromElementsOp, TensorCastOp, TensorLoadOp,
TensorToMemrefOp>();
BufferizeTypeConverter converter;
@ -175,6 +182,7 @@ struct FinalBufferizePass : public FinalBufferizePassBase<FinalBufferizePass> {
typesAreLegal);
OwningRewritePatternList patterns;
populateTensorBufferizePatterns(&context, converter, patterns);
populateStdBufferizePatterns(&context, converter, patterns);
populateEliminateBufferizeMaterializationsPatterns(&context, converter,
patterns);

2
tensorflow/compiler/mlir/tools/kernel_gen/transforms/shape_to_descriptors_pass.cc
View File

@ -21,6 +21,7 @@ limitations under the License.
#include "mlir/Dialect/Shape/IR/Shape.h" // from @llvm-project
#include "mlir/Dialect/Shape/Transforms/Passes.h" // from @llvm-project
#include "mlir/Dialect/StandardOps/IR/Ops.h" // from @llvm-project
#include "mlir/Dialect/Tensor/IR/Tensor.h" // from @llvm-project
#include "mlir/IR/MLIRContext.h" // from @llvm-project
#include "mlir/IR/PatternMatch.h" // from @llvm-project
#include "mlir/Transforms/DialectConversion.h" // from @llvm-project
@ -49,6 +50,7 @@ struct ShapeToDescriptorsPass
target.addIllegalDialect<shape::ShapeDialect>();
target.addLegalDialect<scf::SCFDialect>();
target.addLegalDialect<StandardOpsDialect>();
target.addLegalDialect<tensor::TensorDialect>();
// Don't mark the primary Cstr/Assuming ops as illegal, so they can be
// lowered at a later time to assertions.
target.addLegalOp<shape::AssumingOp, shape::AssumingYieldOp,

4
tensorflow/workspace.bzl
View File

@ -685,8 +685,8 @@ def tf_repositories(path_prefix = "", tf_repo_name = ""):
)
# Check out LLVM and MLIR from llvm-project.
LLVM_COMMIT = "7aeb3804c46cc6c8f291415ca09ae34021301eb8"
LLVM_SHA256 = "945718051c37e55fc3aa79f4752dbcb22240ed19bc8037d41b4d461acd34c000"
LLVM_COMMIT = "0cf7e4b252fe1458fddb8e3dbfcae43450e9c04c"
LLVM_SHA256 = "6f663272226c5d36a213a308865af67ea58bcf7986c1f5551d48bf7eed83edfc"
LLVM_URLS = [
"https://storage.googleapis.com/mirror.tensorflow.org/github.com/llvm/llvm-project/archive/{commit}.tar.gz".format(commit = LLVM_COMMIT),
"https://github.com/llvm/llvm-project/archive/{commit}.tar.gz".format(commit = LLVM_COMMIT),

370
third_party/mlir/BUILD vendored
View File

@ -550,6 +550,185 @@ cc_library(
],
)
##---------------------------------------------------------------------------##
# ArmSVE dialect.
##---------------------------------------------------------------------------##
filegroup(
name = "ArmSVETdFiles",
srcs = [
"include/mlir/Dialect/ArmSVE/ArmSVE.td",
"include/mlir/Dialect/LLVMIR/LLVMOpBase.td",
"include/mlir/IR/OpBase.td",
":SideEffectTdFiles",
],
)
gentbl(
name = "ArmSVEIncGen",
strip_include_prefix = "include",
tbl_outs = [
(
"-gen-op-decls",
"include/mlir/Dialect/ArmSVE/ArmSVE.h.inc",
),
(
"-gen-op-defs",
"include/mlir/Dialect/ArmSVE/ArmSVE.cpp.inc",
),
(
"-gen-typedef-decls",
"include/mlir/Dialect/ArmSVE/ArmSVETypes.h.inc",
),
(
"-gen-typedef-defs",
"include/mlir/Dialect/ArmSVE/ArmSVETypes.cpp.inc",
),
(
"-gen-dialect-decls -dialect arm_sve",
"include/mlir/Dialect/ArmSVE/ArmSVEDialect.h.inc",
),
],
tblgen = ":mlir-tblgen",
td_file = "include/mlir/Dialect/ArmSVE/ArmSVE.td",
td_srcs = [
":ArmSVETdFiles",
],
)
cc_library(
name = "ArmSVE",
srcs = [
"lib/Dialect/ArmSVE/IR/ArmSVEDialect.cpp",
],
hdrs = [
"include/mlir/Dialect/ArmSVE/ArmSVEDialect.h",
],
includes = ["include"],
deps = [
":ArmSVEIncGen",
":IR",
":SideEffectInterfaces",
":VectorOps",
"@llvm-project//llvm:Core",
"@llvm-project//llvm:Support",
],
)
cc_library(
name = "ArmSVEToLLVM",
srcs = glob([
"lib/Conversion/ArmSVEToLLVM/*.cpp",
]) + ["lib/Conversion/PassDetail.h"],
hdrs = glob([
"include/mlir/Conversion/ArmSVEToLLVM/*.h",
]),
includes = ["include"],
deps = [
":ArmSVE",
":ConversionPassIncGen",
":EDSC",
":IR",
":LLVMArmSVE",
":LLVMDialect",
":Pass",
":StandardOps",
":StandardToLLVM",
":Support",
":Transforms",
":VectorOps",
"@llvm-project//llvm:Core",
"@llvm-project//llvm:Support",
],
)
filegroup(
name = "LLVMArmSVETdFiles",
srcs = [
"include/mlir/Dialect/LLVMIR/LLVMArmSVE.td",
":LLVMOpsTdFiles",
],
)
gentbl(
name = "LLVMArmSVEIncGen",
strip_include_prefix = "include",
tbl_outs = [
(
"-gen-dialect-decls -dialect=llvm_arm_sve",
"include/mlir/Dialect/LLVMIR/LLVMArmSVEDialect.h.inc",
),
(
"-gen-op-decls",
"include/mlir/Dialect/LLVMIR/LLVMArmSVE.h.inc",
),
(
"-gen-op-defs",
"include/mlir/Dialect/LLVMIR/LLVMArmSVE.cpp.inc",
),
(
"-gen-op-doc",
"g3doc/Dialects/LLVMIR/LLVMArmSVE.md",
),
],
tblgen = ":mlir-tblgen",
td_file = "include/mlir/Dialect/LLVMIR/LLVMArmSVE.td",
td_srcs = [
":LLVMArmSVETdFiles",
],
)
cc_library(
name = "LLVMArmSVE",
srcs = [
"lib/Dialect/LLVMIR/IR/LLVMArmSVEDialect.cpp",
],
hdrs = [
"include/mlir/Dialect/LLVMIR/LLVMArmSVEDialect.h",
],
includes = ["include"],
deps = [
":IR",
":LLVMArmSVEIncGen",
":LLVMDialect",
"@llvm-project//llvm:Core",
"@llvm-project//llvm:Support",
],
)
gentbl(
name = "LLVMArmSVEConversionIncGen",
strip_include_prefix = "include",
tbl_outs = [
(
"-gen-llvmir-conversions",
"include/mlir/Dialect/LLVMIR/LLVMArmSVEConversions.inc",
),
],
tblgen = ":mlir-tblgen",
td_file = "include/mlir/Dialect/LLVMIR/LLVMArmSVE.td",
td_srcs = [
":LLVMArmSVETdFiles",
],
)
cc_library(
name = "TargetLLVMArmSVEIntr",
srcs = [
"lib/Target/LLVMIR/LLVMArmSVEIntr.cpp",
],
includes = ["include"],
deps = [
":IR",
":LLVMArmSVE",
":LLVMArmSVEConversionIncGen",
":LLVMIRModuleTranslation",
":Translation",
"@llvm-project//llvm:Core",
"@llvm-project//llvm:Support",
],
)
##---------------------------------------------------------------------------##
# AVX512 dialect.
##---------------------------------------------------------------------------##
@ -1314,6 +1493,7 @@ cc_library(
":ShapeToStandardGen",
":StandardOps",
":Support",
":TensorDialect",
":Transforms",
],
)
@ -1371,6 +1551,7 @@ cc_library(
":SideEffectInterfaces",
":StandardOpsIncGen",
":Support",
":TensorDialect",
":VectorInterfaces",
":ViewLikeInterface",
"@llvm-project//llvm:Support",
@ -1406,6 +1587,7 @@ cc_library(
":StandardOps",
":StandardOpsTransformsPassIncGen",
":Support",
":TensorDialect",
":Transforms",
"@llvm-project//llvm:Support",
],
@ -1508,6 +1690,8 @@ cc_library(
"lib/Dialect/LLVMIR/IR/*AVX512*.h",
"lib/Dialect/LLVMIR/IR/*ArmNeon*.cpp",
"lib/Dialect/LLVMIR/IR/*ArmNeon*.h",
"lib/Dialect/LLVMIR/IR/*ArmSVE*.cpp",
"lib/Dialect/LLVMIR/IR/*ArmSVE*.h",
"lib/Dialect/LLVMIR/IR/NVVM*.cpp",
"lib/Dialect/LLVMIR/IR/NVVM*.h",
"lib/Dialect/LLVMIR/IR/ROCDL*.cpp",
@ -1521,6 +1705,7 @@ cc_library(
exclude = [
"include/mlir/Dialect/LLVMIR/*AVX512*.h",
"include/mlir/Dialect/LLVMIR/*ArmNeon*.h",
"include/mlir/Dialect/LLVMIR/*ArmSVE*.h",
"include/mlir/Dialect/LLVMIR/NVVM*.h",
"include/mlir/Dialect/LLVMIR/ROCDL*.h",
],
@ -2623,18 +2808,12 @@ alias(
)
cc_library(
name = "SPIRVSerialization",
srcs = glob(
[
"lib/Dialect/SPIRV/Serialization/*.cpp",
name = "SPIRVBinaryUtils",
srcs = [
"lib/Target/SPIRV/SPIRVBinaryUtils.cpp",
],
exclude = [
"lib/Dialect/SPIRV/Serialization/TranslateRegistration.cpp",
],
),
hdrs = [
"include/mlir/Dialect/SPIRV/SPIRVBinaryUtils.h",
"include/mlir/Dialect/SPIRV/Serialization.h",
"include/mlir/Target/SPIRV/SPIRVBinaryUtils.h",
],
includes = ["include"],
deps = [
@ -2642,6 +2821,48 @@ cc_library(
":SPIRVDialect",
":SPIRVOpUtilsIncGen",
":SPIRVOpsIncGen",
":Support",
"@llvm-project//llvm:Support",
],
)
cc_library(
name = "SPIRVSerialization",
srcs = [
"lib/Target/SPIRV/Serialization.cpp",
],
hdrs = [
"include/mlir/Target/SPIRV/Serialization.h",
],
includes = ["include"],
deps = [
":IR",
":SPIRVBinaryUtils",
":SPIRVDialect",
":SPIRVOpUtilsIncGen",
":SPIRVOpsIncGen",
":SPIRVSerializationGen",
":Support",
":Transforms",
"@llvm-project//llvm:Support",
],
)
cc_library(
name = "SPIRVDeserialization",
srcs = [
"lib/Target/SPIRV/Deserialization.cpp",
],
hdrs = [
"include/mlir/Target/SPIRV/Deserialization.h",
],
includes = ["include"],
deps = [
":IR",
":SPIRVBinaryUtils",
":SPIRVDialect",
":SPIRVOpUtilsIncGen",
":SPIRVOpsIncGen",
":SPIRVSerializationGen",
":Support",
":Transforms",
@ -2671,12 +2892,13 @@ cc_library(
cc_library(
name = "SPIRVTranslateRegistration",
srcs = [
"lib/Dialect/SPIRV/Serialization/TranslateRegistration.cpp",
"lib/Target/SPIRV/TranslateRegistration.cpp",
],
includes = ["include"],
deps = [
":IR",
":Parser",
":SPIRVDeserialization",
":SPIRVDialect",
":SPIRVSerialization",
":Support",
@ -2685,6 +2907,120 @@ cc_library(
],
)
filegroup(
name = "TensorOpsTdFiles",
srcs = [
"include/mlir/Dialect/Tensor/IR/TensorBase.td",
"include/mlir/Dialect/Tensor/IR/TensorOps.td",
":OpBaseTdFiles",
":SideEffectTdFiles",
],
)
gentbl(
name = "TensorBaseIncGen",
strip_include_prefix = "include",
tbl_outs = [
(
"-gen-dialect-decls -dialect=tensor",
"include/mlir/Dialect/Tensor/IR/TensorOpsDialect.h.inc",
),
],
tblgen = ":mlir-tblgen",
td_file = "include/mlir/Dialect/Tensor/IR/TensorBase.td",
td_srcs = [
":TensorOpsTdFiles",
],
)
gentbl(
name = "TensorOpsIncGen",
strip_include_prefix = "include",
tbl_outs = [
(
"-gen-op-decls",
"include/mlir/Dialect/Tensor/IR/TensorOps.h.inc",
),
(
"-gen-op-defs",
"include/mlir/Dialect/Tensor/IR/TensorOps.cpp.inc",
),
],
tblgen = ":mlir-tblgen",
td_file = "include/mlir/Dialect/Tensor/IR/TensorOps.td",
td_srcs = [
":TensorOpsTdFiles",
],
)
cc_library(
name = "TensorDialect",
srcs = glob(
[
"lib/Dialect/Tensor/IR/*.cpp",
"lib/Dialect/Tensor/IR/*.h",
],
) + [
"include/mlir/Transforms/InliningUtils.h",
],
hdrs = [
"include/mlir/Dialect/Tensor/IR/Tensor.h",
],
includes = ["include"],
deps = [
":IR",
":SideEffectInterfaces",
":Support",
":TensorBaseIncGen",
":TensorOpsIncGen",
"@llvm-project//llvm:Support",
],
)
gentbl(
name = "TensorPassIncGen",
strip_include_prefix = "include",
tbl_outs = [
(
"-gen-pass-decls -name Tensor",
"include/mlir/Dialect/Tensor/Transforms/Passes.h.inc",
),
],
tblgen = ":mlir-tblgen",
td_file = "include/mlir/Dialect/Tensor/Transforms/Passes.td",
td_srcs = [
":PassBaseTdFiles",
],
)
cc_library(
name = "TensorTransforms",
srcs = glob(
[
"lib/Dialect/Tensor/Transforms/*.cpp",
"lib/Dialect/Tensor/Transforms/*.h",
],
),
hdrs = [
"include/mlir/Dialect/Tensor/Transforms/Passes.h",
],
includes = ["include"],
deps = [
":Async",
":EDSC",
":IR",
":ParallelLoopMapperAttrGen",
":Pass",
":SCFDialect",
":StandardOps",
":Support",
":TensorDialect",
":TensorPassIncGen",
":Transforms",
"@llvm-project//llvm:Support",
],
)
cc_library(
name = "Rewrite",
srcs = glob([
@ -3320,6 +3656,7 @@ cc_library(
":Support",
":TargetLLVMAVX512Intr",
":TargetLLVMArmNeonIntr",
":TargetLLVMArmSVEIntr",
":Translation",
"@llvm-project//llvm:Core",
"@llvm-project//llvm:IRReader",
@ -3468,6 +3805,7 @@ cc_library(
"@llvm-project//mlir/test:TestSPIRV",
"@llvm-project//mlir/test:TestShapeDialect",
"@llvm-project//mlir/test:TestTransforms",
"@llvm-project//mlir/test:TestTypeDialect",
],
)
@ -3519,6 +3857,8 @@ cc_library(
":AffineTransforms",
":ArmNeon",
":ArmNeonToLLVM",
":ArmSVE",
":ArmSVEToLLVM",
":Async",
":AsyncPassIncGen",
":AsyncToLLVM",
@ -3535,6 +3875,7 @@ cc_library(
":IR",
":LLVMAVX512",
":LLVMArmNeon",
":LLVMArmSVE",
":LLVMDialect",
":LLVMIRTransforms",
":LLVMPassIncGen",
@ -3573,6 +3914,8 @@ cc_library(
":StandardOpsTransformsPassIncGen",
":StandardToLLVM",
":StandardToSPIRVTransforms",
":TensorDialect",
":TensorTransforms",
":TosaDialect",
":Transforms",
":TransformsPassIncGen",
@ -3618,6 +3961,7 @@ cc_binary(
"@llvm-project//mlir/test:TestSPIRV",
"@llvm-project//mlir/test:TestTosaDialect",
"@llvm-project//mlir/test:TestTransforms",
"@llvm-project//mlir/test:TestTypeDialect",
],
)
@ -4417,6 +4761,7 @@ cc_library(
":StandardOpsTransforms",
":StandardToLLVM",
":Support",
":TensorDialect",
":TransformUtils",
":Transforms",
":TransformsPassIncGen",
@ -4481,12 +4826,15 @@ cc_library(
":AVX512ToLLVM",
":ArmNeon",
":ArmNeonToLLVM",
":ArmSVE",
":ArmSVEToLLVM",
":ConversionPassIncGen",
":DialectUtils",
":EDSC",
":IR",
":LLVMAVX512",
":LLVMArmNeon",
":LLVMArmSVE",
":LLVMDialect",
":LLVMIRModuleTranslation",
":Pass",

12
third_party/mlir/test.BUILD vendored
View File

@ -314,6 +314,18 @@ cc_library(
],
)
cc_library(
name = "TestTypeDialect",
srcs = glob([
"lib/Dialect/LLVMIR/*.cpp",
]),
deps = [
":TestDialect",
"@llvm-project//mlir:IR",
"@llvm-project//mlir:LLVMDialect",
],
)
cc_library(
name = "TestTosaDialect",
srcs = glob([