Integrate llvm/llvm-project@ca23b7ca476fb (#11429)
* Reset third_party/llvm-project:
ca23b7ca476fb66771d1c6d02c6837938fde6c7e (2022-12-02 19:05:43 +0000):
[AsmPrinter] .addrsig_sym: remove isTransitiveUsedByMetadataOnly
* Updated to tensorflow/tensorflow@a4ed969
* Updated to tensorflow/mlir-hlo@9e3d164
* Called `expand-strided-metadata` pass for LLVMGPU
Co-authored-by: Thomas Raoux <thomasraoux@google.com>
diff --git a/compiler/src/iree/compiler/Codegen/Dialect/test/lowering_config_attr.mlir b/compiler/src/iree/compiler/Codegen/Dialect/test/lowering_config_attr.mlir
index ff9a552..29e2315 100644
--- a/compiler/src/iree/compiler/Codegen/Dialect/test/lowering_config_attr.mlir
+++ b/compiler/src/iree/compiler/Codegen/Dialect/test/lowering_config_attr.mlir
@@ -39,7 +39,10 @@
return
}
}
-// CHECK: #compilation = #iree_codegen.compilation_info<lowering_config = <tile_sizes = []>, translation_info = <CPUDefault>>
+// CHECK: #config = #iree_codegen.lowering_config<tile_sizes = []>
+// CHECK: #translation = #iree_codegen.translation_info<CPUDefault>
+// CHECK: #compilation = #iree_codegen.compilation_info<lowering_config = #config, translation_info = #translation>
+
// -----
@@ -53,4 +56,7 @@
return
}
}
-// CHECK: #compilation = #iree_codegen.compilation_info<lowering_config = <tile_sizes = []>, translation_info = <CPUDefault>, workgroup_size = [16, 4, 1], subgroup_size = 32>
+// CHECK: #config = #iree_codegen.lowering_config<tile_sizes = []>
+// CHECK: #translation = #iree_codegen.translation_info<CPUDefault>
+// CHECK: #compilation = #iree_codegen.compilation_info<lowering_config = #config, translation_info = #translation, workgroup_size = [16, 4, 1], subgroup_size = 32>
+
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/Passes.cpp b/compiler/src/iree/compiler/Codegen/LLVMGPU/Passes.cpp
index 5294c62..b99c356 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMGPU/Passes.cpp
+++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/Passes.cpp
@@ -394,6 +394,7 @@
pm.addNestedPass<func::FuncOp>(arith::createArithExpandOpsPass());
pm.addNestedPass<func::FuncOp>(memref::createExpandOpsPass());
+ pm.addPass(memref::createExpandStridedMetadataPass());
pm.addPass(createLowerAffinePass());
// Strip out the debug info for the kernel as CUDA driver doesn't diggest PTX
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/workgroup_specialization_pipeline_test.mlir b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/workgroup_specialization_pipeline_test.mlir
index aece8ac..15250eb 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/workgroup_specialization_pipeline_test.mlir
+++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/workgroup_specialization_pipeline_test.mlir
@@ -96,10 +96,10 @@
// CHECK: %[[ARR:.*]] = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer) offset(%[[c0]]) alignment(64) : memref<102401xf32>
// CHECK: %[[ARR2:.*]] = hal.interface.binding.subspan set(0) binding(1) type(storage_buffer) offset(%[[c0]]) alignment(64) : memref<102401xf32>
// CHECK: %[[BLKX:.*]] = hal.interface.workgroup.id[0] : index
-// CHECK: %[[BLKX2:.*]] = affine.min #map2()[%[[BLKX]]]
+// CHECK: %[[BLKX2:.*]] = affine.min #{{.+}}()[%[[BLKX]]]
// CHECK: %[[CMP:.*]] = arith.cmpi eq, %[[BLKX2]], %[[c256]] : index
// CHECK: scf.if %[[CMP]]
// CHECK: %[[TIDX:.*]] = gpu.thread_id x
-// CHECK: %[[AFF:.*]] = affine.apply #map3(%[[TIDX]])[%[[BLKX]]]
+// CHECK: %[[AFF:.*]] = affine.apply #{{.+}}(%[[TIDX]])[%[[BLKX]]]
// CHECK: vector.transfer_read %[[ARR]][%[[AFF]]], %[[cst]] {in_bounds = [true]} : memref<102401xf32>, vector<4xf32>
-// CHECK: vector.transfer_read %[[ARR2]][%[[AFF]]], %[[cst]] {in_bounds = [true]} : memref<102401xf32>, vector<4xf32>
\ No newline at end of file
+// CHECK: vector.transfer_read %[[ARR2]][%[[AFF]]], %[[cst]] {in_bounds = [true]} : memref<102401xf32>, vector<4xf32>
diff --git a/compiler/src/iree/compiler/Tools/BUILD b/compiler/src/iree/compiler/Tools/BUILD
index ea46a1c..c1268a7 100644
--- a/compiler/src/iree/compiler/Tools/BUILD
+++ b/compiler/src/iree/compiler/Tools/BUILD
@@ -104,7 +104,6 @@
"@llvm-project//mlir:LinalgDialect",
"@llvm-project//mlir:LinalgPassIncGen",
"@llvm-project//mlir:LinalgToLLVM",
- "@llvm-project//mlir:LinalgToSPIRV",
"@llvm-project//mlir:LinalgTransforms",
"@llvm-project//mlir:MLProgramDialect",
"@llvm-project//mlir:MathDialect",
diff --git a/compiler/src/iree/compiler/Tools/CMakeLists.txt b/compiler/src/iree/compiler/Tools/CMakeLists.txt
index 6a2960a..66091ec 100644
--- a/compiler/src/iree/compiler/Tools/CMakeLists.txt
+++ b/compiler/src/iree/compiler/Tools/CMakeLists.txt
@@ -133,7 +133,6 @@
MLIRLLVMDialect
MLIRLinalgDialect
MLIRLinalgToLLVM
- MLIRLinalgToSPIRV
MLIRLinalgTransforms
MLIRMLProgramDialect
MLIRQuantDialect
diff --git a/compiler/src/iree/compiler/Tools/init_mlir_passes.h b/compiler/src/iree/compiler/Tools/init_mlir_passes.h
index 8c36168..fc9e784 100644
--- a/compiler/src/iree/compiler/Tools/init_mlir_passes.h
+++ b/compiler/src/iree/compiler/Tools/init_mlir_passes.h
@@ -79,7 +79,6 @@
registerConvertGPUToSPIRVPass();
registerConvertControlFlowToSPIRVPass();
registerConvertFuncToSPIRVPass();
- registerConvertLinalgToSPIRVPass();
}
} // namespace mlir
diff --git a/integrations/tensorflow/WORKSPACE b/integrations/tensorflow/WORKSPACE
index 5e9d4dd..e1b8fc6 100644
--- a/integrations/tensorflow/WORKSPACE
+++ b/integrations/tensorflow/WORKSPACE
@@ -7,7 +7,7 @@
load("@bazel_tools//tools/build_defs/repo:git.bzl", "git_repository")
-TENSORFLOW_COMMIT = "cfe950a3f9b48555816aeece894ed6a7a2271749"
+TENSORFLOW_COMMIT = "a4ed969207533e0e82cdda414d1eb78cb1072f6d"
git_repository(
name = "org_tensorflow",
diff --git a/integrations/tensorflow/iree-dialects/lib/Dialect/LinalgExt/Passes/TileAndDecomposeWinogradPass.cpp b/integrations/tensorflow/iree-dialects/lib/Dialect/LinalgExt/Passes/TileAndDecomposeWinogradPass.cpp
index 7214789..67b54fd 100644
--- a/integrations/tensorflow/iree-dialects/lib/Dialect/LinalgExt/Passes/TileAndDecomposeWinogradPass.cpp
+++ b/integrations/tensorflow/iree-dialects/lib/Dialect/LinalgExt/Passes/TileAndDecomposeWinogradPass.cpp
@@ -191,7 +191,7 @@
loopNest.loops.back().getBody()->getTerminator())) {
rewriter.replaceOpWithNewOp<scf::YieldOp>(yieldOp, updatedOutput);
}
- inputOp.getResults()[0].replaceAllUsesWith(loopNest.getResults()[0]);
+ inputOp.getResults()[0].replaceAllUsesWith(loopNest.results[0]);
return success();
}
};
@@ -338,7 +338,7 @@
loopNest.loops.back().getBody()->getTerminator())) {
rewriter.replaceOpWithNewOp<scf::YieldOp>(yieldOp, updatedOutput);
}
- outputOp.getResults()[0].replaceAllUsesWith(loopNest.getResults()[0]);
+ outputOp.getResults()[0].replaceAllUsesWith(loopNest.results[0]);
return success();
}
};
diff --git a/integrations/tensorflow/iree-dialects/lib/Dialect/LinalgExt/Transforms/Tiling.cpp b/integrations/tensorflow/iree-dialects/lib/Dialect/LinalgExt/Transforms/Tiling.cpp
index e8ff056..de4bff6 100644
--- a/integrations/tensorflow/iree-dialects/lib/Dialect/LinalgExt/Transforms/Tiling.cpp
+++ b/integrations/tensorflow/iree-dialects/lib/Dialect/LinalgExt/Transforms/Tiling.cpp
@@ -90,7 +90,7 @@
createMatchingSubsetInsertOp(b, loc, sliceOp, sliceOp, iterArgs[0]);
return scf::ValueVector({yieldValue});
});
- return loopNest.getResults();
+ return loopNest.results;
}
namespace {
diff --git a/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/Passes/TileAndDecomposeWinogradPass.cpp b/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/Passes/TileAndDecomposeWinogradPass.cpp
index 7214789..67b54fd 100644
--- a/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/Passes/TileAndDecomposeWinogradPass.cpp
+++ b/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/Passes/TileAndDecomposeWinogradPass.cpp
@@ -191,7 +191,7 @@
loopNest.loops.back().getBody()->getTerminator())) {
rewriter.replaceOpWithNewOp<scf::YieldOp>(yieldOp, updatedOutput);
}
- inputOp.getResults()[0].replaceAllUsesWith(loopNest.getResults()[0]);
+ inputOp.getResults()[0].replaceAllUsesWith(loopNest.results[0]);
return success();
}
};
@@ -338,7 +338,7 @@
loopNest.loops.back().getBody()->getTerminator())) {
rewriter.replaceOpWithNewOp<scf::YieldOp>(yieldOp, updatedOutput);
}
- outputOp.getResults()[0].replaceAllUsesWith(loopNest.getResults()[0]);
+ outputOp.getResults()[0].replaceAllUsesWith(loopNest.results[0]);
return success();
}
};
diff --git a/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/Transforms/Tiling.cpp b/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/Transforms/Tiling.cpp
index e8ff056..de4bff6 100644
--- a/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/Transforms/Tiling.cpp
+++ b/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/Transforms/Tiling.cpp
@@ -90,7 +90,7 @@
createMatchingSubsetInsertOp(b, loc, sliceOp, sliceOp, iterArgs[0]);
return scf::ValueVector({yieldValue});
});
- return loopNest.getResults();
+ return loopNest.results;
}
namespace {
diff --git a/llvm-external-projects/iree-dialects/test/Dialect/iree_linalg_ext/tiling.mlir b/llvm-external-projects/iree-dialects/test/Dialect/iree_linalg_ext/tiling.mlir
index 44cab77..d89d4ed 100644
--- a/llvm-external-projects/iree-dialects/test/Dialect/iree_linalg_ext/tiling.mlir
+++ b/llvm-external-projects/iree-dialects/test/Dialect/iree_linalg_ext/tiling.mlir
@@ -1322,8 +1322,6 @@
ins(%arg0 : memref<1x10x10x1280xf32>) outs(%arg1 : memref<8x8x1x2x2x1280xf32>)
return
}
-// CHECK-DAG: #[[MAP:.+]] = affine_map<(d0, d1, d2, d3) -> (d0, d1, d2, d3)>
-// CHECK-DAG: #[[MAP1:.+]] = affine_map<(d0, d1, d2, d3, d4, d5) -> (d0, d1, d2, d3, d4, d5)>
// CHECK-DAG: #[[MAP2:.+]] = affine_map<(d0)[s0, s1] -> (1, -d0 + s1)>
// CHECK-DAG: #[[MAP3:.+]] = affine_map<(d0)[s0, s1] -> (32, -d0 + s1)>
// CHECK: func.func @winograd_input_transform_memref(%[[ARG0:[a-zA-Z0-9_]+]]: memref<1x10x10x1280xf32>,
@@ -1398,8 +1396,6 @@
ins(%arg0 : memref<8x8x1x2x2x32xf32>) outs(%arg1 : memref<1x12x12x32xf32>)
return
}
-// CHECK-DAG: #[[MAP:.+]] = affine_map<(d0, d1, d2, d3, d4, d5) -> (d0, d1, d2, d3, d4, d5)>
-// CHECK-DAG: #[[MAP1:.+]] = affine_map<(d0, d1, d2, d3) -> (d0, d1, d2, d3)>
// CHECK-DAG: #[[MAP2:.+]] = affine_map<(d0)[s0, s1] -> (1, -d0 + s1)>
// CHECK-DAG: #[[MAP3:.+]] = affine_map<(d0)[s0, s1] -> (32, -d0 + s1)>
// CHECK: func.func @winograd_output_transform_memref(%[[ARG0:[a-zA-Z0-9_]+]]: memref<8x8x1x2x2x32xf32>,
diff --git a/third_party/llvm-project b/third_party/llvm-project
index 3d22047..ca23b7c 160000
--- a/third_party/llvm-project
+++ b/third_party/llvm-project
@@ -1 +1 @@
-Subproject commit 3d2204760a89def9b6412e0737905d93babeb112
+Subproject commit ca23b7ca476fb66771d1c6d02c6837938fde6c7e
diff --git a/third_party/mlir-hlo b/third_party/mlir-hlo
index 8d929fc..9e3d164 160000
--- a/third_party/mlir-hlo
+++ b/third_party/mlir-hlo
@@ -1 +1 @@
-Subproject commit 8d929fcdd31e269371f00192958754c02f8bc7df
+Subproject commit 9e3d164c0e7f4bd4eee80ce95cef0941755d5a8e
