Merge branch 'main' into google-to-main
diff --git a/SUBMODULE_VERSIONS.txt b/SUBMODULE_VERSIONS.txt index 8815076..813f325 100644 --- a/SUBMODULE_VERSIONS.txt +++ b/SUBMODULE_VERSIONS.txt
@@ -4,14 +4,14 @@ aa533abfd4232b01f9e57041d70114d5a77e6de0 third_party/googletest 88b845dee001723c4a0db1fe5477de735b6d3bb0 third_party/liburing acd6f6f014c25e46363e718381e0b35205df2d83 third_party/libyaml -57b40b5f34383634949d1639e64a5c2acd0dc5f6 third_party/llvm-project -a04c82f16281fc6288edb8e34d9b2f2fad767ab2 third_party/mlir-hlo +6fbc812883ab949d2ce8ee8768f986f2cb39ac35 third_party/llvm-project +d343113a490691716eac87db61d77b07a7304436 third_party/mlir-hlo 3f701faace7addc75d16dea8a6cd769fa5b3f260 third_party/musl 4c7697dbe973ed01ae6fbec37d186ebd05982e1f third_party/pybind11 2e1b5fb39ebc2ef4cb77005f8267e4f3a6241ba1 third_party/spirv_cross f5417a4b6633c3217c9a1bc2f0c70b1454975ba7 third_party/spirv_headers b42009b3b9d4ca35bc703f5310eedc74f584be58 third_party/stblib -0a1636bb1667586c9109934625e15814aadadba1 third_party/tensorflow +55e85b4c921995b6130026a0988e2dec69ad244b third_party/tensorflow 058e89011fceca912d43638ebb6b85992147fcfe third_party/tracy 9d10a96f2d57c3c37e167f2e73c9a31ac2e51fa5 third_party/vulkan_headers 8d4a9e9174a9c6ad6a3a3ae981b915ef13fc12c4 third_party/vulkan_memory_allocator
diff --git a/benchmarks/TensorFlow/CMakeLists.txt b/benchmarks/TensorFlow/CMakeLists.txt index 87b466e..6f95f86 100644 --- a/benchmarks/TensorFlow/CMakeLists.txt +++ b/benchmarks/TensorFlow/CMakeLists.txt
@@ -21,7 +21,7 @@ "fp16" # MODULE_TAGS # This uses the same input MLIR source as fp32 to save download time. # It requires users to have "--iree-flow-demote-f32-to-f16". - "https://storage.googleapis.com/iree-model-artifacts/MobileBertSquad-9c0042d19.tar.gz" # MLIR_SOURCE + "https://storage.googleapis.com/iree-model-artifacts/MobileBertSquad-89edfa50d.tar.gz" # MLIR_SOURCE "serving_default" # ENTRY_FUNCTION # The conversion done by "--iree-flow-demote-f32-to-f16" won't change the # original input signature. @@ -31,7 +31,7 @@ set(MOBILEBERT_FP32_MODULE "MobileBertSquad" # MODULE_NAME "fp32" # MODULE_TAGS - "https://storage.googleapis.com/iree-model-artifacts/MobileBertSquad-9c0042d19.tar.gz" # MLIR_SOURCE + "https://storage.googleapis.com/iree-model-artifacts/MobileBertSquad-89edfa50d.tar.gz" # MLIR_SOURCE "serving_default" # ENTRY_FUNCTION "1x384xi32,1x384xi32,1x384xi32" # FUNCTION_INPUTS ) @@ -39,7 +39,7 @@ set(MOBILENET_V2_MODULE "MobileNetV2" # MODULE_NAME "fp32,imagenet" # MODULE_TAGS - "https://storage.googleapis.com/iree-model-artifacts/MobileNetV2-9c0042d19.tar.gz" # MLIR_SOURCE + "https://storage.googleapis.com/iree-model-artifacts/MobileNetV2-89edfa50d.tar.gz" # MLIR_SOURCE "call" # ENTRY_FUNCTION "1x224x224x3xf32" # FUNCTION_INPUTS ) @@ -47,7 +47,7 @@ set(MOBILENET_V3SMALL_MODULE "MobileNetV3Small" # MODULE_NAME "fp32,imagenet" # MODULE_TAGS - "https://storage.googleapis.com/iree-model-artifacts/MobileNetV3Small-9c0042d19.tar.gz" # MLIR_SOURCE + "https://storage.googleapis.com/iree-model-artifacts/MobileNetV3Small-89edfa50d.tar.gz" # MLIR_SOURCE "call" # ENTRY_FUNCTION "1x224x224x3xf32" # FUNCTION_INPUTS )
diff --git a/bindings/python/iree/runtime/vm_test.py b/bindings/python/iree/runtime/vm_test.py index e903952..41fac47 100644 --- a/bindings/python/iree/runtime/vm_test.py +++ b/bindings/python/iree/runtime/vm_test.py
@@ -18,7 +18,7 @@ binary = iree.compiler.compile_str( """ func @add_scalar(%arg0: i32, %arg1: i32) -> i32 { - %0 = addi %arg0, %arg1 : i32 + %0 = arith.addi %arg0, %arg1 : i32 return %0 : i32 } """,
diff --git a/build_tools/bazel_to_cmake/bazel_to_cmake_targets.py b/build_tools/bazel_to_cmake/bazel_to_cmake_targets.py index 32dfb56..bbfc8e3 100644 --- a/build_tools/bazel_to_cmake/bazel_to_cmake_targets.py +++ b/build_tools/bazel_to_cmake/bazel_to_cmake_targets.py
@@ -33,6 +33,7 @@ "@llvm-project//mlir:LLVMDialect": ["MLIRLLVMIR"], "@llvm-project//mlir:LLVMTransforms": ["MLIRStandardToLLVM"], "@llvm-project//mlir:MathDialect": ["MLIRMath"], + "@llvm-project//mlir:ArithmeticDialect": ["MLIRArithmetic"], "@llvm-project//mlir:MemRefDialect": ["MLIRMemRef"], "@llvm-project//mlir:SCFToGPUPass": ["MLIRSCFToGPU"], "@llvm-project//mlir:SCFDialect": ["MLIRSCF"],
diff --git a/colab/edge_detection.ipynb b/colab/edge_detection.ipynb index a04bd5f..95526cb 100644 --- a/colab/edge_detection.ipynb +++ b/colab/edge_detection.ipynb
@@ -359,11 +359,11 @@ "#map = affine_map<(d0, d1, d2, d3) -> (d0, d1, d2, d3)>\n", "module {\n", " func @edge_detect_sobel_operator(%arg0: tensor<1x128x128x1xf32> {tf._user_specified_name = \"image\"} loc(unknown)) -> tensor<1x128x128x1xf32> attributes {iree.module.export, iree.reflection = {abi = \"sip\", abiv = 1 : i32, sip = \"I8!S5!k0_0R3!_0\"}, tf._construction_context = \"kEagerRuntime\", tf._input_shapes = [#tf.shape<1x128x128x1>]} {\n", - " %cst = constant 0.000000e+00 : f32 loc(#loc1)\n", - " %c0 = constant 0 : index loc(#loc1)\n", - " %c1 = constant 1 : index loc(#loc1)\n", - " %cst_0 = constant dense<[[[-1.000000e+00], [0.000000e+00], [1.000000e+00]], [[-2.000000e+00], [0.000000e+00], [2.000000e+00]], [[-1.000000e+00], [0.000000e+00], [1.000000e+00]]]> : tensor<3x3x1xf32> loc(#loc1)\n", - " %cst_1 = constant dense<[[[1.000000e+00], [2.000000e+00], [1.000000e+00]], [[0.000000e+00], [0.000000e+00], [0.000000e+00]], [[-1.000000e+00], [-2.000000e+00], [-1.000000e+00]]]> : tensor<3x3x1xf32> loc(#loc2)\n", + " %cst = arith.constant 0.000000e+00 : f32 loc(#loc1)\n", + " %c0 = arith.constant 0 : index loc(#loc1)\n", + " %c1 = arith.constant 1 : index loc(#loc1)\n", + " %cst_0 = arith.constant dense<[[[-1.000000e+00], [0.000000e+00], [1.000000e+00]], [[-2.000000e+00], [0.000000e+00], [2.000000e+00]], [[-1.000000e+00], [0.000000e+00], [1.000000e+00]]]> : tensor<3x3x1xf32> loc(#loc1)\n", + " %cst_1 = arith.constant dense<[[[1.000000e+00], [2.000000e+00], [1.000000e+00]], [[0.000000e+00], [0.000000e+00], [0.000000e+00]], [[-1.000000e+00], [-2.000000e+00], [-1.000000e+00]]]> : tensor<3x3x1xf32> loc(#loc2)\n", " %0 = linalg.pad_tensor %arg0 low[%c0, %c1, %c1, %c0] high[%c0, %c1, %c1, %c0] {\n", " ^bb0(%arg1: index loc(unknown), %arg2: index loc(unknown), %arg3: index loc(unknown), %arg4: index loc(unknown)): // no predecessors\n", " linalg.yield %cst : f32 loc(#loc1)\n", @@ -374,7 +374,7 @@ " %4 = linalg.init_tensor [1, 128, 128, 1] : tensor<1x128x128x1xf32> loc(#loc3)\n", " %5 = linalg.generic {indexing_maps = [#map, #map], iterator_types = [\"parallel\", \"parallel\", \"parallel\", \"parallel\"]} ins(%3 : tensor<1x128x128x1xf32>) outs(%4 : tensor<1x128x128x1xf32>) {\n", " ^bb0(%arg1: f32 loc(unknown), %arg2: f32 loc(unknown)): // no predecessors\n", - " %15 = mulf %arg1, %arg1 : f32 loc(#loc3)\n", + " %15 = arith.mulf %arg1, %arg1 : f32 loc(#loc3)\n", " linalg.yield %15 : f32 loc(#loc3)\n", " } -> tensor<1x128x128x1xf32> loc(#loc3)\n", " %6 = linalg.init_tensor [1, 128, 128, 1] : tensor<1x128x128x1xf32> loc(#loc2)\n", @@ -383,13 +383,13 @@ " %9 = linalg.init_tensor [1, 128, 128, 1] : tensor<1x128x128x1xf32> loc(#loc4)\n", " %10 = linalg.generic {indexing_maps = [#map, #map], iterator_types = [\"parallel\", \"parallel\", \"parallel\", \"parallel\"]} ins(%8 : tensor<1x128x128x1xf32>) outs(%9 : tensor<1x128x128x1xf32>) {\n", " ^bb0(%arg1: f32 loc(unknown), %arg2: f32 loc(unknown)): // no predecessors\n", - " %15 = mulf %arg1, %arg1 : f32 loc(#loc4)\n", + " %15 = arith.mulf %arg1, %arg1 : f32 loc(#loc4)\n", " linalg.yield %15 : f32 loc(#loc4)\n", " } -> tensor<1x128x128x1xf32> loc(#loc4)\n", " %11 = linalg.init_tensor [1, 128, 128, 1] : tensor<1x128x128x1xf32> loc(#loc5)\n", " %12 = linalg.generic {indexing_maps = [#map, #map, #map], iterator_types = [\"parallel\", \"parallel\", \"parallel\", \"parallel\"]} ins(%5, %10 : tensor<1x128x128x1xf32>, tensor<1x128x128x1xf32>) outs(%11 : tensor<1x128x128x1xf32>) {\n", " ^bb0(%arg1: f32 loc(unknown), %arg2: f32 loc(unknown), %arg3: f32 loc(unknown)): // no predecessors\n", - " %15 = addf %arg1, %arg2 : f32 loc(#loc5)\n", + " %15 = arith.addf %arg1, %arg2 : f32 loc(#loc5)\n", " linalg.yield %15 : f32 loc(#loc5)\n", " } -> tensor<1x128x128x1xf32> loc(#loc5)\n", " %13 = linalg.init_tensor [1, 128, 128, 1] : tensor<1x128x128x1xf32> loc(#loc6)\n",
diff --git a/colab/low_level_invoke_function.ipynb b/colab/low_level_invoke_function.ipynb index b06c22f..937d774 100644 --- a/colab/low_level_invoke_function.ipynb +++ b/colab/low_level_invoke_function.ipynb
@@ -112,7 +112,7 @@ " module @arithmetic {\n", " func @simple_mul(%arg0: tensor<4xf32>, %arg1: tensor<4xf32>) -> tensor<4xf32>\n", " {\n", - " %0 = mulf %arg0, %arg1 : tensor<4xf32>\n", + " %0 = arith.mulf %arg0, %arg1 : tensor<4xf32>\n", " return %0 : tensor<4xf32>\n", " } \n", " }\n",
diff --git a/docs/developers/design_docs/simple_ir_walkthrough.md b/docs/developers/design_docs/simple_ir_walkthrough.md index 7268d00..9e17d98 100644 --- a/docs/developers/design_docs/simple_ir_walkthrough.md +++ b/docs/developers/design_docs/simple_ir_walkthrough.md
@@ -127,9 +127,9 @@ ```mlir func @simple_mul(%arg0: tensor<4xf32>, %arg1: tensor<4xf32>) -> tensor<4xf32> { - %cst = constant dense<[4, 1, 1]> : tensor<3xi32> + %cst = arith.constant dense<[4, 1, 1]> : tensor<3xi32> %0 = iree.dispatch_region[%cst : tensor<3xi32>](%arg2 = %arg0 : tensor<4xf32>, %arg3 = %arg1 : tensor<4xf32>) : tensor<4xf32> { - %1 = mulf %arg2, %arg3 : tensor<4xf32> + %1 = arith.mulf %arg2, %arg3 : tensor<4xf32> iree.return %1 : tensor<4xf32> } return %0 : tensor<4xf32> @@ -169,7 +169,7 @@ attributes {iree.executable.export} { %0 = iree.load_input(%arg0 : memref<4xf32>) : tensor<4xf32> %1 = iree.load_input(%arg1 : memref<4xf32>) : tensor<4xf32> - %2 = mulf %0, %1 : tensor<4xf32> + %2 = arith.mulf %0, %1 : tensor<4xf32> iree.store_output(%2 : tensor<4xf32>, %arg2 : memref<4xf32>) iree.return } @@ -215,7 +215,7 @@ attributes {iree.executable.export, iree.executable.workload = dense<[4, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} { %0 = iree.load_input(%arg0 : memref<4xf32>) : tensor<4xf32> %1 = iree.load_input(%arg1 : memref<4xf32>) : tensor<4xf32> - %2 = mulf %0, %1 : tensor<4xf32> + %2 = arith.mulf %0, %1 : tensor<4xf32> iree.store_output(%2 : tensor<4xf32>, %arg2 : memref<4xf32>) iree.return }
diff --git a/docs/developers/design_roadmap.md b/docs/developers/design_roadmap.md index 39312d1..1baed85 100644 --- a/docs/developers/design_roadmap.md +++ b/docs/developers/design_roadmap.md
@@ -117,8 +117,8 @@ loop iteration counter and comparison. ```mlir -%start = constant dense<1> : tensor<i32> -%bound = constant dense<3> : tensor<i32> +%start = arith.constant dense<1> : tensor<i32> +%bound = arith.constant dense<3> : tensor<i32> %res = "mhlo.while"(%start) ( { ^bb0(%count: tensor<i32>): %1 = "mhlo.compare"(%count, %bound) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1> @@ -136,9 +136,9 @@ ```mlir func @main() -> tensor<i32> attributes {iree.reflection = {f = "I1!R6!B3!t6", fv = "1"}} { - %cst = constant dense<1> : tensor<i32> - %cst_0 = constant dense<3> : tensor<i32> - %cst_1 = constant dense<1> : vector<3xi32> + %cst = arith.constant dense<1> : tensor<i32> + %cst_0 = arith.constant dense<3> : tensor<i32> + %cst_1 = arith.constant dense<1> : vector<3xi32> br ^bb1(%cst : tensor<i32>) ^bb1(%2: tensor<i32>): // 2 preds: ^bb0, ^bb2 %3 = flow.ex.stream.fragment(%arg0 = %cst_1 : vector<3xi32>, %arg1 = %2 : tensor<i32>, %arg2 = %cst_0 : tensor<i32>) -> tensor<i1> { @@ -267,7 +267,7 @@ flow.stream.append[%s0](...) { flow.tensor.update ... } - %b = cmpi ne %some_flag, ... + %b = arith.cmpi ne %some_flag, ... cond_br %b, ^a(%s0), ^b(%s0) ^a(%s1): flow.stream.append[%s1](...) { @@ -640,14 +640,14 @@ Allocating tensors from the ringbuffer does not require sophisticated runtime packing as we can emit IR to calculate required sizes for dynamically shaped -tensors. Whether a basic block reserves `%sz = constant 42 : index` bytes or -`%sz = std.muli %cst, %dyn_dim : index` bytes doesn't materially change how the -allocations are performed. Since almost all usage involves simple write head +tensors. Whether a basic block reserves `%sz = arith.constant 42 : index` bytes +or `%sz = std.muli %cst, %dyn_dim : index` bytes doesn't materially change how +the allocations are performed. Since almost all usage involves simple write head bumps there is no need for ahead-of-time memory planning or large fixed allocations, and since no buffer within the ringbuffer can alias we can have -coarse (_read: low overhead_) guarantees about the availability of certain -regions of the ringbuffer (_"when this event is signaled all prior ringbuffer -writes have completed"_). +coarse (*read: low overhead*) guarantees about the availability of certain +regions of the ringbuffer (*"when this event is signaled all prior ringbuffer +writes have completed"*). Usually any planning we may want to perform can be done in IR via code motion. For example applying traditional algorithms used to reduce register pressure
diff --git a/docs/developers/developing_iree/testing_guide.md b/docs/developers/developing_iree/testing_guide.md index 926630c..0fcc23c 100644 --- a/docs/developers/developing_iree/testing_guide.md +++ b/docs/developers/developing_iree/testing_guide.md
@@ -271,7 +271,7 @@ around ```mlir -%expected = constant dense<101.0> : tensor<f32> +%expected = arith.constant dense<101.0> : tensor<f32> check.expect_almost_eq(%result, %expected) : tensor<f32> ```
diff --git a/integrations/tensorflow/iree_tf_compiler/BUILD b/integrations/tensorflow/iree_tf_compiler/BUILD index 96d7dbb..bf987f5 100644 --- a/integrations/tensorflow/iree_tf_compiler/BUILD +++ b/integrations/tensorflow/iree_tf_compiler/BUILD
@@ -113,6 +113,7 @@ "//iree_tf_compiler/MHLO", "@llvm-project//llvm:Support", "@llvm-project//mlir:IR", + "@llvm-project//mlir:MathDialect", "@llvm-project//mlir:Parser", "@llvm-project//mlir:Pass", "@llvm-project//mlir:StandardOps",
diff --git a/integrations/tensorflow/iree_tf_compiler/TF/ConvertToMHLO.cpp b/integrations/tensorflow/iree_tf_compiler/TF/ConvertToMHLO.cpp index b2b0a03..daedf83 100644 --- a/integrations/tensorflow/iree_tf_compiler/TF/ConvertToMHLO.cpp +++ b/integrations/tensorflow/iree_tf_compiler/TF/ConvertToMHLO.cpp
@@ -33,12 +33,13 @@ // the IREE requires. class ConvertToMHLOPass : public PassWrapper<ConvertToMHLOPass, FunctionPass> { void getDependentDialects(DialectRegistry ®istry) const override { - registry.insert<mlir::linalg::LinalgDialect, mlir::TF::TensorFlowDialect, - mlir::tf_executor::TensorFlowExecutorDialect, - mlir::tf_device::TensorFlowDeviceDialect, - mlir::tf_saved_model::TensorFlowSavedModelDialect, - chlo::HloClientDialect, mhlo::MhloDialect, - shape::ShapeDialect, StandardOpsDialect>(); + registry + .insert<mlir::linalg::LinalgDialect, mlir::TF::TensorFlowDialect, + mlir::tf_executor::TensorFlowExecutorDialect, + mlir::tf_device::TensorFlowDeviceDialect, + mlir::tf_saved_model::TensorFlowSavedModelDialect, + chlo::HloClientDialect, mhlo::MhloDialect, shape::ShapeDialect, + mlir::arith::ArithmeticDialect, StandardOpsDialect>(); } StringRef getArgument() const override { return "iree-tf-convert-to-mhlo"; } @@ -89,7 +90,8 @@ target.addLegalDialect<chlo::HloClientDialect>(); target.addLegalDialect<linalg::LinalgDialect>(); target.addLegalDialect<mhlo::MhloDialect>(); - target.addLegalDialect<mlir::StandardOpsDialect>(); + target.addLegalDialect<mlir::StandardOpsDialect, + mlir::arith::ArithmeticDialect>(); target.addLegalDialect<shape::ShapeDialect>(); target.addLegalDialect<tensor::TensorDialect>(); target.addLegalOp<mlir::CallOp>();
diff --git a/integrations/tensorflow/iree_tf_compiler/TF/SavedModelToIreeABI.cpp b/integrations/tensorflow/iree_tf_compiler/TF/SavedModelToIreeABI.cpp index 0c41825..287bc38 100644 --- a/integrations/tensorflow/iree_tf_compiler/TF/SavedModelToIreeABI.cpp +++ b/integrations/tensorflow/iree_tf_compiler/TF/SavedModelToIreeABI.cpp
@@ -259,15 +259,15 @@ } // Recurse into sequence (index can be sparse on child ikey). if (type == LevelType::List || type == LevelType::Tuple) { - Value listSizeValue = - builder.create<ConstantOp>(loc, builder.getIndexType(), - builder.getIndexAttr(getNeededListSize())); + Value listSizeValue = builder.create<arith::ConstantOp>( + loc, builder.getIndexType(), + builder.getIndexAttr(getNeededListSize())); Value listValue = builder.create<IREE::Util::ListCreateOp>( loc, getIrType(builder), listSizeValue); builder.create<IREE::Util::ListResizeOp>(loc, listValue, listSizeValue); for (StructureLevel &child : children) { Value childValue = child.emitCreateReturns(loc, builder, callReturns); - Value indexValue = builder.create<ConstantOp>( + Value indexValue = builder.create<arith::ConstantOp>( loc, builder.getIndexType(), builder.getIndexAttr(child.ikey)); builder.create<IREE::Util::ListSetOp>(loc, listValue, indexValue, childValue); @@ -277,16 +277,16 @@ // Recurse into dict (modeled as a dense tuple of children). if (type == LevelType::Dict) { - Value listSizeValue = - builder.create<ConstantOp>(loc, builder.getIndexType(), - builder.getIndexAttr(getNeededListSize())); + Value listSizeValue = builder.create<arith::ConstantOp>( + loc, builder.getIndexType(), + builder.getIndexAttr(getNeededListSize())); Value listValue = builder.create<IREE::Util::ListCreateOp>( loc, getIrType(builder), listSizeValue); builder.create<IREE::Util::ListResizeOp>(loc, listValue, listSizeValue); for (auto it : llvm::enumerate(children)) { StructureLevel &child = it.value(); Value childValue = child.emitCreateReturns(loc, builder, callReturns); - Value indexValue = builder.create<ConstantOp>( + Value indexValue = builder.create<arith::ConstantOp>( loc, builder.getIndexType(), builder.getIndexAttr(it.index())); builder.create<IREE::Util::ListSetOp>(loc, listValue, indexValue, childValue); @@ -300,8 +300,8 @@ // given index. Value emitGetFromList(Location loc, OpBuilder &builder, Value parentList, int index) { - Value indexValue = builder.create<ConstantOp>(loc, builder.getIndexType(), - builder.getIndexAttr(index)); + Value indexValue = builder.create<arith::ConstantOp>( + loc, builder.getIndexType(), builder.getIndexAttr(index)); Value itemValue = builder.create<IREE::Util::ListGetOp>( loc, getIrType(builder), parentList, indexValue); // TODO: Null check, etc. How does that work if returning a tensor? Need
diff --git a/integrations/tensorflow/iree_tf_compiler/TF/test/lower_global_tensors_complex.mlir b/integrations/tensorflow/iree_tf_compiler/TF/test/lower_global_tensors_complex.mlir index 4758edd..e5c05b0 100644 --- a/integrations/tensorflow/iree_tf_compiler/TF/test/lower_global_tensors_complex.mlir +++ b/integrations/tensorflow/iree_tf_compiler/TF/test/lower_global_tensors_complex.mlir
@@ -33,7 +33,7 @@ // CHECK: func @f(%arg0: tensor<?xf32> {tf_saved_model.index_path = [0]}) -> (tensor<?xf32> {tf_saved_model.index_path = [0]}) attributes {tf_saved_model.exported_names = ["f"]} { // CHECK-NEXT: [[PTR0:%.+]] = util.global.address [[V]] : !util.ptr<tensor<?xf32>> // CHECK-NEXT: [[PTR1:%.+]] = util.global.address [[V1]] : !util.ptr<tensor<?xf32>> -// CHECK-NEXT: %[[FALSE:.+]] = constant false +// CHECK-NEXT: %[[FALSE:.+]] = arith.constant false // CHECK-NEXT: cond_br %[[FALSE]], ^bb1([[PTR0]] : !util.ptr<tensor<?xf32>>), ^bb1([[PTR1]] : !util.ptr<tensor<?xf32>>) // CHECK-NEXT: ^bb1([[PTR:%.+]]: !util.ptr<tensor<?xf32>>): // 2 preds: ^bb0, ^bb0 // CHECK-NEXT: [[T:%.+]] = util.global.load.indirect [[PTR]] : !util.ptr<tensor<?xf32>> -> tensor<?xf32> @@ -42,7 +42,7 @@ "tf_saved_model.global_tensor"() { is_mutable, sym_name = "v", type = tensor<?xf32>, value = dense<1.> : tensor<1xf32> } : () -> () "tf_saved_model.global_tensor"() { is_mutable, sym_name = "v1", type = tensor<?xf32>, value = dense<1.> : tensor<1xf32> } : () -> () func @f(%arg0: tensor<?xf32> {tf_saved_model.index_path = [0]}, %v: tensor<!tf_type.resource<tensor<?xf32>>> {tf_saved_model.bound_input = @v}, %v1: tensor<!tf_type.resource<tensor<?xf32>>> {tf_saved_model.bound_input = @v1}) -> (tensor<?xf32> {tf_saved_model.index_path = [0]}) attributes {tf_saved_model.exported_names = ["f"]} { - %pred = constant false + %pred = arith.constant false cond_br %pred, ^bb1(%v : tensor<!tf_type.resource<tensor<?xf32>>>), ^bb1(%v1 : tensor<!tf_type.resource<tensor<?xf32>>>) ^bb1(%either: tensor<!tf_type.resource<tensor<?xf32>>>): %ret = "tf.ReadVariableOp"(%either) : (tensor<!tf_type.resource<tensor<?xf32>>>) -> tensor<?xf32>
diff --git a/integrations/tensorflow/iree_tf_compiler/TF/test/saved_model_to_iree_abi.mlir b/integrations/tensorflow/iree_tf_compiler/TF/test/saved_model_to_iree_abi.mlir index 8f909c0..9195f59 100644 --- a/integrations/tensorflow/iree_tf_compiler/TF/test/saved_model_to_iree_abi.mlir +++ b/integrations/tensorflow/iree_tf_compiler/TF/test/saved_model_to_iree_abi.mlir
@@ -61,41 +61,41 @@ // CHECK-LABEL: module @dict_nest // CHECK: func @dict_nest(%arg0: !util.list<?>, %arg1: !hal.buffer_view) -> (!util.list<?>, !util.list<?>) // CHECK-SAME{LITERAL}: iree.abi = "{\22a\22:[[\22sdict\22,[\22dict\22,[\22sdict\22,[\22a\22,[\22ndarray\22,\22f32\22,1,16]],[\22b\22,[\22ndarray\22,\22f32\22,1,16]]]],[\22list\22,[\22slist\22,[\22ndarray\22,\22f32\22,1,16],[\22ndarray\22,\22f32\22,1,16]]]],[\22ndarray\22,\22f32\22,0]],\22r\22:[[\22sdict\22,[\22dict\22,[\22sdict\22,[\22a\22,[\22ndarray\22,\22f32\22,1,16]],[\22b\22,[\22ndarray\22,\22f32\22,1,16]]]],[\22list\22,[\22stuple\22,[\22ndarray\22,\22f32\22,1,16],[\22ndarray\22,\22f32\22,1,16]]]]],\22v\22:1}" -// CHECK: %[[c0:.+]] = constant 0 : index +// CHECK: %[[c0:.+]] = arith.constant 0 : index // CHECK: %[[L0:.+]] = util.list.get %arg0[%[[c0]]] : !util.list<?> -> !util.list<?> -// CHECK: %[[c0_0:.+]] = constant 0 : index +// CHECK: %[[c0_0:.+]] = arith.constant 0 : index // CHECK: %[[L1:.+]] = util.list.get %[[L0]][%[[c0_0]]] : !util.list<?> -> !hal.buffer_view // CHECK: %[[L1_TENSOR:.+]] = hal.tensor.cast %[[L1]] : !hal.buffer_view -> tensor<16xf32> -// CHECK: %[[c1:.+]] = constant 1 : index +// CHECK: %[[c1:.+]] = arith.constant 1 : index // CHECK: %[[L2:.+]] = util.list.get %[[L0]][%[[c1]]] : !util.list<?> -> !hal.buffer_view // CHECK: %[[L2_TENSOR:.+]] = hal.tensor.cast %[[L2]] : !hal.buffer_view -> tensor<16xf32> -// CHECK: %[[c1_1:.+]] = constant 1 : index +// CHECK: %[[c1_1:.+]] = arith.constant 1 : index // CHECK: %[[L3:.+]] = util.list.get %arg0[%[[c1_1]]] : !util.list<?> -> !util.list<?> -// CHECK: %[[c0_2:.+]] = constant 0 : index +// CHECK: %[[c0_2:.+]] = arith.constant 0 : index // CHECK: %[[L4:.+]] = util.list.get %[[L3]][%[[c0_2]]] : !util.list<?> -> !hal.buffer_view // CHECK: %[[L4_TENSOR:.+]] = hal.tensor.cast %[[L4]] : !hal.buffer_view -> tensor<16xf32> -// CHECK: %[[c1_3:.+]] = constant 1 : index +// CHECK: %[[c1_3:.+]] = arith.constant 1 : index // CHECK: %[[L5:.+]] = util.list.get %[[L3]][%[[c1_3]]] : !util.list<?> -> !hal.buffer_view // CHECK: %[[L5_TENSOR:.+]] = hal.tensor.cast %[[L5]] : !hal.buffer_view -> tensor<16xf32> // CHECK: %[[ARG1_TENSOR:.+]] = hal.tensor.cast %arg1 : !hal.buffer_view -> tensor<f32> // CHECK: %[[RESULT:.+]]:4 = call @__inference_dict_nest_190(%[[L1_TENSOR]], %[[L2_TENSOR]], %[[L4_TENSOR]], %[[L5_TENSOR]], %[[ARG1_TENSOR]]) : (tensor<16xf32>, tensor<16xf32>, tensor<16xf32>, tensor<16xf32>, tensor<f32>) -> (tensor<16xf32>, tensor<16xf32>, tensor<16xf32>, tensor<16xf32>) -// CHECK: %[[c2:.+]] = constant 2 : index +// CHECK: %[[c2:.+]] = arith.constant 2 : index // CHECK: %[[R7:.+]] = util.list.create %[[c2]] : !util.list<?> // CHECK: util.list.resize %[[R7]], %[[c2]] // CHECK: %[[R0_BV:.+]] = hal.tensor.cast %[[RESULT]]#0 : tensor<16xf32> -> !hal.buffer_view -// CHECK: %[[c0_4:.+]] = constant 0 : index +// CHECK: %[[c0_4:.+]] = arith.constant 0 : index // CHECK: util.list.set %[[R7]][%[[c0_4]]], %[[R0_BV]] : !hal.buffer_view -> !util.list<?> // CHECK: %[[R1_BV:.+]] = hal.tensor.cast %[[RESULT]]#1 : tensor<16xf32> -> !hal.buffer_view -// CHECK: %[[c1_5:.+]] = constant 1 : index +// CHECK: %[[c1_5:.+]] = arith.constant 1 : index // CHECK: util.list.set %[[R7]][%[[c1_5]]], %[[R1_BV]] : !hal.buffer_view -> !util.list<?> -// CHECK: %[[c2_8:.+]] = constant 2 : index +// CHECK: %[[c2_8:.+]] = arith.constant 2 : index // CHECK: %[[R9:.+]] = util.list.create %[[c2_8]] : !util.list<?> // CHECK: util.list.resize %[[R9]], %[[c2_8]] // CHECK: %[[R2_BV:.+]] = hal.tensor.cast %[[RESULT]]#2 : tensor<16xf32> -> !hal.buffer_view -// CHECK: %[[c0_9:.+]] = constant 0 : index +// CHECK: %[[c0_9:.+]] = arith.constant 0 : index // CHECK: util.list.set %[[R9]][%[[c0_9]]], %[[R2_BV]] : !hal.buffer_view -> !util.list<?> // CHECK: %[[R3_BV:.+]] = hal.tensor.cast %[[RESULT]]#3 : tensor<16xf32> -> !hal.buffer_view -// CHECK: %[[c1_10:.+]] = constant 1 : index +// CHECK: %[[c1_10:.+]] = arith.constant 1 : index // CHECK: util.list.set %[[R9]][%[[c1_10]]], %[[R3_BV]] : !hal.buffer_view -> !util.list<?> // return %[[R7]], %[[R8]] : !util.list<?>, !util.list<?> // CHECK: func private @__inference_dict_nest_190
diff --git a/integrations/tensorflow/iree_tf_compiler/iree-import-tflite-main.cpp b/integrations/tensorflow/iree_tf_compiler/iree-import-tflite-main.cpp index 6043499..d7cb458 100644 --- a/integrations/tensorflow/iree_tf_compiler/iree-import-tflite-main.cpp +++ b/integrations/tensorflow/iree_tf_compiler/iree-import-tflite-main.cpp
@@ -64,7 +64,7 @@ registry.insert<mlir::TFL::TensorFlowLiteDialect>(); registry.insert<mlir::tosa::TosaDialect>(); registry.insert<quant::QuantizationDialect>(); - registry.insert<StandardOpsDialect>(); + registry.insert<StandardOpsDialect, mlir::arith::ArithmeticDialect>(); // Convert the Module proto into MLIR. MLIRContext context(registry);
diff --git a/integrations/tensorflow/iree_tf_compiler/iree-import-xla-main.cpp b/integrations/tensorflow/iree_tf_compiler/iree-import-xla-main.cpp index 08f270a..dcbf2e7 100644 --- a/integrations/tensorflow/iree_tf_compiler/iree-import-xla-main.cpp +++ b/integrations/tensorflow/iree_tf_compiler/iree-import-xla-main.cpp
@@ -16,6 +16,7 @@ #include "llvm/Support/InitLLVM.h" #include "llvm/Support/ToolOutputFile.h" #include "mlir-hlo/Dialect/mhlo/IR/register.h" +#include "mlir/Dialect/Math/IR/Math.h" #include "mlir/Dialect/StandardOps/IR/Ops.h" #include "mlir/IR/AsmState.h" #include "mlir/IR/BuiltinOps.h" @@ -144,7 +145,8 @@ DialectRegistry registry; mlir::mhlo::registerAllMhloDialects(registry); - registry.insert<mlir::StandardOpsDialect>(); + registry.insert<mlir::StandardOpsDialect, mlir::arith::ArithmeticDialect, + mlir::math::MathDialect>(); MLIRContext context; OwningModuleRef module = ModuleOp::create(mlir::UnknownLoc::get(&context)); context.appendDialectRegistry(registry);
diff --git a/iree/base/internal/BUILD b/iree/base/internal/BUILD index d16aec5..27db3b0 100644 --- a/iree/base/internal/BUILD +++ b/iree/base/internal/BUILD
@@ -287,11 +287,6 @@ name = "synchronization_benchmark", testonly = True, srcs = ["synchronization_benchmark.cc"], - tags = [ - "manual", - "nobuilder", - "notap", - ], deps = [ ":synchronization", "//iree/testing:benchmark_main", @@ -302,11 +297,6 @@ run_binary_test( name = "synchronization_benchmark_test", args = ["--benchmark_min_time=0"], - tags = [ - "manual", - "nobuilder", - "notap", - ], test_binary = ":synchronization_benchmark", )
diff --git a/iree/base/internal/CMakeLists.txt b/iree/base/internal/CMakeLists.txt index bc2d6f4..97bc9fc 100644 --- a/iree/base/internal/CMakeLists.txt +++ b/iree/base/internal/CMakeLists.txt
@@ -315,10 +315,6 @@ "--benchmark_min_time=0" TEST_BINARY ::synchronization_benchmark - LABELS - "manual" - "nobuilder" - "notap" ) iree_cc_test(
diff --git a/iree/compiler/Bindings/Native/Transforms/WrapEntryPoints.cpp b/iree/compiler/Bindings/Native/Transforms/WrapEntryPoints.cpp index 3f17b8e..03ff4d7 100644 --- a/iree/compiler/Bindings/Native/Transforms/WrapEntryPoints.cpp +++ b/iree/compiler/Bindings/Native/Transforms/WrapEntryPoints.cpp
@@ -29,7 +29,8 @@ : public PassWrapper<WrapEntryPointsPass, OperationPass<ModuleOp>> { public: void getDependentDialects(DialectRegistry ®istry) const override { - registry.insert<StandardOpsDialect, IREE::HAL::HALDialect, + registry.insert<StandardOpsDialect, mlir::arith::ArithmeticDialect, + IREE::HAL::HALDialect, // TODO: memref is here because the **tensor** dim op was // moved there for some reason. When that goes away we can // drop this dependency.
diff --git a/iree/compiler/Bindings/TFLite/Transforms/MaterializeShapeSupport.cpp b/iree/compiler/Bindings/TFLite/Transforms/MaterializeShapeSupport.cpp index 8327bf6..5f14a58 100644 --- a/iree/compiler/Bindings/TFLite/Transforms/MaterializeShapeSupport.cpp +++ b/iree/compiler/Bindings/TFLite/Transforms/MaterializeShapeSupport.cpp
@@ -42,7 +42,7 @@ registry.insert<iree_compiler::IREE::Flow::FlowDialect>(); registry.insert<iree_compiler::IREE::Util::UtilDialect>(); registry.insert<iree_compiler::ShapeDialect>(); - registry.insert<StandardOpsDialect>(); + registry.insert<StandardOpsDialect, mlir::arith::ArithmeticDialect>(); } StringRef getArgument() const override { @@ -249,7 +249,8 @@ } // Clear the dirty flag now that the shapes have been updated. - auto falseValue = exitBuilder.createOrFold<ConstantIntOp>(exitLoc, 0, 1); + auto falseValue = + exitBuilder.createOrFold<arith::ConstantIntOp>(exitLoc, 0, 1); exitBuilder.create<IREE::Util::GlobalStoreOp>(exitLoc, falseValue, dirtyGlobalOp.getName()); exitBuilder.create<ReturnOp>(exitLoc); @@ -285,9 +286,10 @@ builder.create<BranchOp>(loc, compareBlocks[0]); for (size_t i = 0; i < caseCount; ++i) { auto compareBuilder = OpBuilder::atBlockBegin(compareBlocks[i]); - auto caseValue = compareBuilder.createOrFold<ConstantIndexOp>(loc, i); - auto eqValue = compareBuilder.createOrFold<CmpIOp>(loc, CmpIPredicate::eq, - indexValue, caseValue); + auto caseValue = + compareBuilder.createOrFold<arith::ConstantIndexOp>(loc, i); + auto eqValue = compareBuilder.createOrFold<arith::CmpIOp>( + loc, arith::CmpIPredicate::eq, indexValue, caseValue); compareBuilder.create<CondBranchOp>( loc, eqValue, caseBlocks[i], i < caseCount - 1 ? compareBlocks[i + 1] : exitBlock); @@ -305,12 +307,12 @@ Value shapeValue, Value listValue, OpBuilder &builder) { builder.create<IREE::Util::ListResizeOp>( loc, listValue, - builder.createOrFold<ConstantIndexOp>(loc, shapeType.getRank())); + builder.createOrFold<arith::ConstantIndexOp>(loc, shapeType.getRank())); for (int i = 0; i < shapeType.getRank(); ++i) { auto dimValue = builder.createOrFold<Shape::RankedDimOp>(loc, shapeValue, i); builder.create<IREE::Util::ListSetOp>( - loc, listValue, builder.createOrFold<ConstantIndexOp>(loc, i), + loc, listValue, builder.createOrFold<arith::ConstantIndexOp>(loc, i), dimValue); } } @@ -323,7 +325,7 @@ if (!shapeType.isDimDynamic(i)) continue; dynamicDims.push_back(builder.createOrFold<IREE::Util::ListGetOp>( loc, builder.getIndexType(), listValue, - builder.createOrFold<ConstantIndexOp>(loc, i))); + builder.createOrFold<arith::ConstantIndexOp>(loc, i))); } return builder.createOrFold<Shape::MakeRankedShapeOp>(loc, shapeType, dynamicDims); @@ -400,7 +402,7 @@ // Set the dirty flag so that shapes get recalculated as needed. auto exitBuilder = OpBuilder::atBlockBegin(exitBlock); - auto trueValue = exitBuilder.createOrFold<ConstantIntOp>(loc, 1, 1); + auto trueValue = exitBuilder.createOrFold<arith::ConstantIntOp>(loc, 1, 1); exitBuilder.create<IREE::Util::GlobalStoreOp>(loc, trueValue, dirtyGlobalOp.getName()); exitBuilder.create<ReturnOp>(loc);
diff --git a/iree/compiler/Bindings/TFLite/Transforms/WrapEntryPoints.cpp b/iree/compiler/Bindings/TFLite/Transforms/WrapEntryPoints.cpp index 981cf82..d874d5a 100644 --- a/iree/compiler/Bindings/TFLite/Transforms/WrapEntryPoints.cpp +++ b/iree/compiler/Bindings/TFLite/Transforms/WrapEntryPoints.cpp
@@ -27,7 +27,7 @@ : public PassWrapper<WrapEntryPointsPass, OperationPass<ModuleOp>> { public: void getDependentDialects(DialectRegistry ®istry) const override { - registry.insert<StandardOpsDialect>(); + registry.insert<StandardOpsDialect, mlir::arith::ArithmeticDialect>(); } StringRef getArgument() const override {
diff --git a/iree/compiler/Bindings/TFLite/Transforms/test/materialize_shape_support.mlir b/iree/compiler/Bindings/TFLite/Transforms/test/materialize_shape_support.mlir index c2f0006..69be1d6 100644 --- a/iree/compiler/Bindings/TFLite/Transforms/test/materialize_shape_support.mlir +++ b/iree/compiler/Bindings/TFLite/Transforms/test/materialize_shape_support.mlir
@@ -9,7 +9,7 @@ // CHECK-DAG: util.global private mutable @_tflite_dynamicEntry_shapes_dirty = true // CHECK-LABEL: func private @_tflite_dynamicEntry_calculate_shapes() { -// CHECK-NEXT: %false = constant false +// CHECK-NEXT: %false = arith.constant false // CHECK-NEXT: %[[IS_DIRTY:.+]] = util.global.load @_tflite_dynamicEntry_shapes_dirty : i1 // CHECK-NEXT: cond_br %[[IS_DIRTY]], ^bb1, ^bb2 // CHECK-NEXT: ^bb1: @@ -32,7 +32,7 @@ // CHECK-LABEL: func @_tflite_dynamicEntry_query_input_shape // CHECK-SAME: (%[[INDEX:.+]]: index, %[[LIST:.+]]: !util.list<index>) -// CHECK: %[[IS_0:.+]] = cmpi eq, %[[INDEX]], %c0 : index +// CHECK: %[[IS_0:.+]] = arith.cmpi eq, %[[INDEX]], %c0 : index // CHECK-NEXT: cond_br %[[IS_0]], ^bb1, ^bb2 // CHECK-NEXT: ^bb1: // CHECK-NEXT: %[[IN0_SHAPE:.+]] = util.global.load @_tflite_dynamicEntry_input0_shape : !shapex.ranked_shape<[?,8,8,3]> @@ -44,7 +44,7 @@ // CHECK-NEXT: util.list.set %[[LIST]][%c3], %c3 : !util.list<index> // CHECK-NEXT: br ^bb4 // CHECK-NEXT: ^bb2: -// CHECK-NEXT: %[[IS_1:.+]] = cmpi eq, %[[INDEX]], %c1 : index +// CHECK-NEXT: %[[IS_1:.+]] = arith.cmpi eq, %[[INDEX]], %c1 : index // CHECK-NEXT: cond_br %[[IS_1]], ^bb3, ^bb4 // CHECK-NEXT: ^bb3: // CHECK-NEXT: %[[IN1_SHAPE:.+]] = util.global.load @_tflite_dynamicEntry_input1_shape : !shapex.ranked_shape<[?,8,8,3]> @@ -61,7 +61,7 @@ // CHECK-LABEL: func @_tflite_dynamicEntry_resize_input_shape // CHECK-SAME: (%[[INDEX:.+]]: index, %[[LIST:.+]]: !util.list<index>) -// CHECK: %[[IS_0:.+]] = cmpi eq, %[[INDEX]], %c0 : index +// CHECK: %[[IS_0:.+]] = arith.cmpi eq, %[[INDEX]], %c0 : index // CHECK-NEXT: cond_br %[[IS_0]], ^bb1, ^bb2 // CHECK-NEXT: ^bb1: // CHECK-NEXT: %[[IN0_D0:.+]] = util.list.get %[[LIST]][%c0] : !util.list<index> @@ -69,7 +69,7 @@ // CHECK-NEXT: util.global.store %[[IN0_SHAPE]], @_tflite_dynamicEntry_input0_shape : !shapex.ranked_shape<[?,8,8,3]> // CHECK-NEXT: br ^bb4 // CHECK-NEXT: ^bb2: -// CHECK-NEXT: %[[IS_1:.+]] = cmpi eq, %[[INDEX]], %c1 : index +// CHECK-NEXT: %[[IS_1:.+]] = arith.cmpi eq, %[[INDEX]], %c1 : index // CHECK-NEXT: cond_br %[[IS_1]], ^bb3, ^bb4 // CHECK-NEXT: ^bb3: // CHECK-NEXT: %[[IN1_D0:.+]] = util.list.get %[[LIST]][%c0] : !util.list<index> @@ -84,7 +84,7 @@ // CHECK-LABEL: func @_tflite_dynamicEntry_query_output_shape // CHECK-SAME: (%[[INDEX:.+]]: index, %[[LIST:.+]]: !util.list<index>) // CHECK: call @_tflite_dynamicEntry_calculate_shapes() : () -> () -// CHECK-NEXT: %[[IS_0:.+]] = cmpi eq, %[[INDEX]], %c0 : index +// CHECK-NEXT: %[[IS_0:.+]] = arith.cmpi eq, %[[INDEX]], %c0 : index // CHECK-NEXT: cond_br %[[IS_0]], ^bb1, ^bb2 // CHECK-NEXT: ^bb1: // CHECK-NEXT: %[[OUT0_SHAPE:.+]] = util.global.load @_tflite_dynamicEntry_output0_shape : !shapex.ranked_shape<[?,8,8,3]> @@ -96,7 +96,7 @@ // CHECK-NEXT: util.list.set %[[LIST]][%c3], %c3 : !util.list<index> // CHECK-NEXT: br ^bb4 // CHECK-NEXT: ^bb2: -// CHECK-NEXT: %[[IS_1:.+]] = cmpi eq, %[[INDEX]], %c1 : index +// CHECK-NEXT: %[[IS_1:.+]] = arith.cmpi eq, %[[INDEX]], %c1 : index // CHECK-NEXT: cond_br %[[IS_1]], ^bb3, ^bb4 // CHECK-NEXT: ^bb3: // CHECK-NEXT: %[[OUT1_SHAPE:.+]] = util.global.load @_tflite_dynamicEntry_output1_shape : !shapex.ranked_shape<[?,8,8,3]>
diff --git a/iree/compiler/Codegen/Common/FlattenMemRefSubspanPass.cpp b/iree/compiler/Codegen/Common/FlattenMemRefSubspanPass.cpp index 154deec..e32e1cb 100644 --- a/iree/compiler/Codegen/Common/FlattenMemRefSubspanPass.cpp +++ b/iree/compiler/Codegen/Common/FlattenMemRefSubspanPass.cpp
@@ -108,7 +108,7 @@ if (type.hasStaticShape()) { assert(dynamicDims.empty()); - return builder.create<ConstantIndexOp>(loc, type.getNumElements()); + return builder.create<arith::ConstantIndexOp>(loc, type.getNumElements()); } int dynamicDimIndex = 0; @@ -120,7 +120,7 @@ if (ShapedType::isDynamic(shape[i])) { dims.push_back(dynamicDims[dynamicDimIndex++]); } else { - dims.push_back(builder.create<ConstantIndexOp>(loc, shape[i])); + dims.push_back(builder.create<arith::ConstantIndexOp>(loc, shape[i])); } } return makeComposedAffineApply(builder, loc, sizeExpr, dims); @@ -269,7 +269,7 @@ if (ShapedType::isDynamic(shape[i])) { dims.push_back(dynamicDims[dynamicDimIndex++]); } else { - dims.push_back(builder.create<ConstantIndexOp>(loc, shape[i])); + dims.push_back(builder.create<arith::ConstantIndexOp>(loc, shape[i])); } } }; @@ -489,7 +489,8 @@ // Create a new subspan op with zero byte offset at the original location. auto ip = rewriter.saveInsertionPoint(); rewriter.setInsertionPointAfter(subspanOp); - Value zero = rewriter.create<ConstantIndexOp>(op.memref().getLoc(), 0); + Value zero = + rewriter.create<arith::ConstantIndexOp>(op.memref().getLoc(), 0); Value newSubspan = rewriter.create<IREE::HAL::InterfaceBindingSubspanOp>( op.memref().getLoc(), subspanOp.getType(), subspanOp.binding(), zero, subspanOp.byte_length(), subspanOp.dynamic_dims()); @@ -501,8 +502,8 @@ auto addMap = AffineMap::get(0, 2, {sym0 + sym1}, context); auto divMap = AffineMap::get(0, 2, {sym0.floorDiv(sym1)}, context); - Value byteValue = rewriter.create<ConstantIndexOp>(op.memref().getLoc(), - numBytes.getValue()); + Value byteValue = rewriter.create<arith::ConstantIndexOp>( + op.memref().getLoc(), numBytes.getValue()); // We assume that upper layers guarantee the byte offset is perfectly // divisible by the element byte count so the content is well aligned. Value offset = rewriter.create<AffineApplyOp>(
diff --git a/iree/compiler/Codegen/Common/LinalgBufferizePass.cpp b/iree/compiler/Codegen/Common/LinalgBufferizePass.cpp index e663c77..abc8454 100644 --- a/iree/compiler/Codegen/Common/LinalgBufferizePass.cpp +++ b/iree/compiler/Codegen/Common/LinalgBufferizePass.cpp
@@ -196,7 +196,8 @@ .Default([&](Operation *op) { return false; }); } return TypeSwitch<Operation *, bool>(definingOp) - .Case<ConstantOp>([&](ConstantOp constantOp) { return true; }) + .Case<arith::ConstantOp>( + [&](arith::ConstantOp constantOp) { return true; }) .Case<linalg::TensorCollapseShapeOp, linalg::TensorExpandShapeOp>( [&](auto op) { return isFromReadOnlyTensor(op.src(), plan); }) .Case<tensor::ExtractSliceOp>([&](tensor::ExtractSliceOp sliceOp) { @@ -214,7 +215,7 @@ /// Adds the result of `std.constant` to its set (there is nothing to tie to /// here). -static LogicalResult analyseConstantOp(ConstantOp constantOp, +static LogicalResult analyseConstantOp(arith::ConstantOp constantOp, BufferizationPlan &plan) { if (!constantOp.getResult().getType().isa<ShapedType>()) return success(); plan.insert(constantOp.getResult()); @@ -285,7 +286,8 @@ auto targetInterfaceOp = getEquivalentOpOfType<IREE::HAL::InterfaceBindingSubspanOp>(target, plan); assert(targetInterfaceOp); - if (auto valueConstantOp = getEquivalentOpOfType<ConstantOp>(value, plan)) { + if (auto valueConstantOp = + getEquivalentOpOfType<arith::ConstantOp>(value, plan)) { return false; } if (auto valueInterfaceOp = @@ -592,7 +594,7 @@ static LogicalResult analyseOperations(FuncOp funcOp, BufferizationPlan &plan) { auto bufferMappingFn = [&](Operation *op) -> WalkResult { return TypeSwitch<Operation *, LogicalResult>(op) - .Case<ConstantOp>([&](ConstantOp constantOp) { + .Case<arith::ConstantOp>([&](arith::ConstantOp constantOp) { return analyseConstantOp(constantOp, plan); }) .Case<IREE::Flow::DispatchTensorLoadOp>( @@ -1249,7 +1251,8 @@ /// Constants that return tensor types can be handled natively by the /// backends. Here just provide a cast to memref to bridge the gap from tensors /// to memrefs. -static LogicalResult convertConstantOp(OpBuilder &b, ConstantOp constantOp, +static LogicalResult convertConstantOp(OpBuilder &b, + arith::ConstantOp constantOp, BlockAndValueMapping &bvm) { Value result = constantOp.getResult(); assert(!bvm.lookupOrNull(result)); @@ -1472,12 +1475,12 @@ *padTensorOp.region().getOps<linalg::YieldOp>().begin(); Value paddingValue = yeildOp.values()[0]; - auto constOp = paddingValue.getDefiningOp<ConstantOp>(); + auto constOp = paddingValue.getDefiningOp<arith::ConstantOp>(); if (!constOp) { return padTensorOp.emitError( "Converting linalg.pad_tensor with non-constant padding value"); } - if (constOp.getValue().isa<DenseElementsAttr>()) { + if (constOp.value().isa<DenseElementsAttr>()) { return padTensorOp.emitError( "Converting linalg.pad_tensor with non-scalar constant padding " "value"); @@ -1506,9 +1509,9 @@ public: LinalgBufferizePass(WorkgroupMemoryAllocationFn fn) : allocationFn(fn) {} void getDependentDialects(DialectRegistry ®istry) const override { - registry - .insert<IREE::Util::UtilDialect, linalg::LinalgDialect, - memref::MemRefDialect, scf::SCFDialect, StandardOpsDialect>(); + registry.insert<IREE::Util::UtilDialect, linalg::LinalgDialect, + memref::MemRefDialect, scf::SCFDialect, StandardOpsDialect, + mlir::math::MathDialect, mlir::arith::ArithmeticDialect>(); } void runOnOperation() override; @@ -1559,7 +1562,7 @@ // `memref`s. auto convertTensorProducingOps = [&](Operation *op) -> WalkResult { return TypeSwitch<Operation *, LogicalResult>(op) - .Case<ConstantOp>([&](ConstantOp constantOp) { + .Case<arith::ConstantOp>([&](arith::ConstantOp constantOp) { return convertConstantOp(b, constantOp, bvm); }) .Case<IREE::Flow::DispatchTensorStoreOp>(
diff --git a/iree/compiler/Codegen/Common/SetNumWorkgroupsPass.cpp b/iree/compiler/Codegen/Common/SetNumWorkgroupsPass.cpp index ad24dac..0f70605 100644 --- a/iree/compiler/Codegen/Common/SetNumWorkgroupsPass.cpp +++ b/iree/compiler/Codegen/Common/SetNumWorkgroupsPass.cpp
@@ -45,8 +45,8 @@ if (dim >= workloadPerWorkgroup.size()) { return failure(); } - rewriter.replaceOpWithNewOp<ConstantIndexOp>(workgroupSizeOp, - workloadPerWorkgroup[dim]); + rewriter.replaceOpWithNewOp<arith::ConstantIndexOp>( + workgroupSizeOp, workloadPerWorkgroup[dim]); return success(); } @@ -120,7 +120,7 @@ // set the number of workgroups to be 1, 1, 1 to have a single invocation. regionBuilder = [](OpBuilder &b, Location loc, std::array<Value, 3> workload) { - Value one = b.create<ConstantIndexOp>(loc, 1); + Value one = b.create<arith::ConstantIndexOp>(loc, 1); return std::array<Value, 3>{one, one, one}; }; } else { @@ -129,7 +129,7 @@ regionBuilder = [&currWorkloadPerWorkgroup]( OpBuilder &b, Location loc, std::array<Value, 3> workload) { - Value one = b.create<ConstantIndexOp>(loc, 1); + Value one = b.create<arith::ConstantIndexOp>(loc, 1); std::array<Value, 3> returnValues = {one, one, one}; for (auto ts : llvm::enumerate(currWorkloadPerWorkgroup)) { returnValues[ts.index()] = linalg::applyMapToValues(
diff --git a/iree/compiler/Codegen/Common/VectorizeConv.cpp b/iree/compiler/Codegen/Common/VectorizeConv.cpp index 4817562..2b31062 100644 --- a/iree/compiler/Codegen/Common/VectorizeConv.cpp +++ b/iree/compiler/Codegen/Common/VectorizeConv.cpp
@@ -117,7 +117,7 @@ VectorType::get({numInputChannels, numOutputChannels}, elementType); auto vector1x4Type = VectorType::get({1, 4}, elementType); auto inputVectorType = VectorType::get({1, numInputChannels}, elementType); - Value zero = rewriter.createOrFold<ConstantIndexOp>(loc, 0); + Value zero = rewriter.createOrFold<arith::ConstantIndexOp>(loc, 0); // Load the entire filter subview. SmallVector<Value, 4> filterIndices(4, zero); @@ -165,20 +165,22 @@ // input vector are used for computing all output channels so data can // be reused. SmallVector<Value, 4> inputIndices(4, zero); - inputIndices[1] = - rewriter.createOrFold<ConstantIndexOp>(loc, oh * heightStride); - inputIndices[2] = - rewriter.createOrFold<ConstantIndexOp>(loc, ow * widthStride); + inputIndices[1] = rewriter.createOrFold<arith::ConstantIndexOp>( + loc, oh * heightStride); + inputIndices[2] = rewriter.createOrFold<arith::ConstantIndexOp>( + loc, ow * widthStride); Value inputVector = rewriter.create<vector::TransferReadOp>( loc, inputVectorType, inputViewOp, inputIndices); for (int oc = 0; oc < numOutputChannels / 4; ++oc) { // Read in the initial value for this output vector. SmallVector<Value, 4> outputIndices(4, zero); - outputIndices[1] = rewriter.createOrFold<ConstantIndexOp>(loc, oh); - outputIndices[2] = rewriter.createOrFold<ConstantIndexOp>(loc, ow); + outputIndices[1] = + rewriter.createOrFold<arith::ConstantIndexOp>(loc, oh); + outputIndices[2] = + rewriter.createOrFold<arith::ConstantIndexOp>(loc, ow); outputIndices[3] = - rewriter.createOrFold<ConstantIndexOp>(loc, oc * 4); + rewriter.createOrFold<arith::ConstantIndexOp>(loc, oc * 4); Value outputVector = rewriter.create<vector::TransferReadOp>( loc, vector1x4Type, outputViewOp, outputIndices); @@ -282,7 +284,7 @@ Type elementType = filterViewOp.getType().getElementType(); auto vector4Type = VectorType::get(4, elementType); auto filterVectorType = VectorType::get({numChannels}, elementType); - Value zero = rewriter.createOrFold<ConstantIndexOp>(loc, 0); + Value zero = rewriter.createOrFold<arith::ConstantIndexOp>(loc, 0); // Load the entire filter subview. SmallVector<Value, 4> filterIndices(3, zero); @@ -300,20 +302,23 @@ for (int ow = 0; ow < numOutputWidths; ++ow) { // Read in the initial value for this output vector. SmallVector<Value, 4> outputIndices(4, zero); - outputIndices[1] = rewriter.createOrFold<ConstantIndexOp>(loc, oh); - outputIndices[2] = rewriter.createOrFold<ConstantIndexOp>(loc, ow); + outputIndices[1] = + rewriter.createOrFold<arith::ConstantIndexOp>(loc, oh); + outputIndices[2] = + rewriter.createOrFold<arith::ConstantIndexOp>(loc, ow); outputIndices[3] = - rewriter.createOrFold<ConstantIndexOp>(loc, oc * 4); + rewriter.createOrFold<arith::ConstantIndexOp>(loc, oc * 4); Value outputVector = rewriter.create<vector::TransferReadOp>( loc, vector4Type, outputViewOp, outputIndices); // Read in the input vector for these 4 input channels a a batch. SmallVector<Value, 4> inputIndices(4, zero); - inputIndices[1] = - rewriter.createOrFold<ConstantIndexOp>(loc, oh * heightStride); - inputIndices[2] = - rewriter.createOrFold<ConstantIndexOp>(loc, ow * widthStride); - inputIndices[3] = rewriter.createOrFold<ConstantIndexOp>(loc, oc * 4); + inputIndices[1] = rewriter.createOrFold<arith::ConstantIndexOp>( + loc, oh * heightStride); + inputIndices[2] = rewriter.createOrFold<arith::ConstantIndexOp>( + loc, ow * widthStride); + inputIndices[3] = + rewriter.createOrFold<arith::ConstantIndexOp>(loc, oc * 4); Value inputVector = rewriter.create<vector::TransferReadOp>( loc, vector4Type, inputViewOp, inputIndices);
diff --git a/iree/compiler/Codegen/Common/VectorizeMMT4d.cpp b/iree/compiler/Codegen/Common/VectorizeMMT4d.cpp index 5c0b928..34340ef 100644 --- a/iree/compiler/Codegen/Common/VectorizeMMT4d.cpp +++ b/iree/compiler/Codegen/Common/VectorizeMMT4d.cpp
@@ -56,7 +56,7 @@ int K0 = lhsType.getShape()[3]; auto loc = mmt4DOp.getLoc(); - auto c0 = rewriter.create<ConstantIndexOp>(loc, 0); + auto c0 = rewriter.create<arith::ConstantIndexOp>(loc, 0); auto lhsVecType = VectorType::get({1, 1, M0, K0}, rewriter.getF32Type()); auto rhsVecType = VectorType::get({1, 1, N0, K0}, rewriter.getF32Type());
diff --git a/iree/compiler/Codegen/Common/test/affinemin_canonicalization.mlir b/iree/compiler/Codegen/Common/test/affinemin_canonicalization.mlir index a829ea3..fff1d11 100644 --- a/iree/compiler/Codegen/Common/test/affinemin_canonicalization.mlir +++ b/iree/compiler/Codegen/Common/test/affinemin_canonicalization.mlir
@@ -3,16 +3,16 @@ // CHECK-LABEL: scf_for_distributed func @scf_for_distributed(%A : memref<i64>, %id1 : index, %count1 : index, %id2 : index, %count2 : index) { - %c1020 = constant 1020 : index - %c1024 = constant 1024 : index + %c1020 = arith.constant 1020 : index + %c1024 = arith.constant 1024 : index %0 = affine.apply affine_map<()[s0] -> (s0 * 32)>()[%id1] %1 = affine.apply affine_map<()[s0] -> (s0 * 32)>()[%count1] // CHECK: scf.for - // CHECK: %[[C32:.*]] = constant 32 : index + // CHECK: %[[C32:.*]] = arith.constant 32 : index // CHECK: scf.for %{{.*}} = %{{.*}} to %[[C32]] - // CHECK-NEXT: %[[C4:.*]] = constant 4 : index - // CHECK-NEXT: %[[C4I64:.*]] = index_cast %[[C4:.*]] + // CHECK-NEXT: %[[C4:.*]] = arith.constant 4 : index + // CHECK-NEXT: %[[C4I64:.*]] = arith.index_cast %[[C4:.*]] // CHECK-NEXT: memref.store %[[C4I64]], %{{.*}}[] : memref<i64> scf.for %arg0 = %0 to %c1024 step %1 { %2 = affine.min affine_map<(d0) -> (32, -d0 + 1024)>(%arg0) @@ -20,7 +20,7 @@ %4 = affine.apply affine_map<()[s0] -> (s0 * 4)>()[%count2] scf.for %arg1 = %3 to %2 step %4 { %5 = affine.min affine_map<(d0, d1) -> (4, d0 - d1)>(%2, %arg1) - %6 = index_cast %5: index to i64 + %6 = arith.index_cast %5: index to i64 memref.store %6, %A[]: memref<i64> } } @@ -31,8 +31,8 @@ // CHECK: scf.for // CHECK: %[[MIN:.*]] = affine.min // CHECK: scf.for %{{.*}} = %{{.*}} to %[[MIN]] - // CHECK-NEXT: %[[C4:.*]] = constant 4 : index - // CHECK-NEXT: %[[C4I64:.*]] = index_cast %[[C4:.*]] + // CHECK-NEXT: %[[C4:.*]] = arith.constant 4 : index + // CHECK-NEXT: %[[C4I64:.*]] = arith.index_cast %[[C4:.*]] // CHECK-NEXT: memref.store %[[C4I64]], %{{.*}}[] : memref<i64> scf.for %arg0 = %0 to %c1020 step %1 { %2 = affine.min affine_map<(d0) -> (32, -d0 + 1020)>(%arg0) @@ -40,7 +40,7 @@ %4 = affine.apply affine_map<()[s0] -> (s0 * 4)>()[%count2] scf.for %arg1 = %3 to %2 step %4 { %5 = affine.min affine_map<(d0, d1) -> (4, d0 - d1)>(%2, %arg1) - %6 = index_cast %5: index to i64 + %6 = arith.index_cast %5: index to i64 memref.store %6, %A[]: memref<i64> } } @@ -49,8 +49,8 @@ // CHECK: scf.for // CHECK: %[[MIN:.*]] = affine.min // CHECK: scf.parallel {{.*}} to (%[[MIN]]) - // CHECK-NEXT: %[[C4:.*]] = constant 4 : index - // CHECK-NEXT: %[[C4I64:.*]] = index_cast %[[C4:.*]] + // CHECK-NEXT: %[[C4:.*]] = arith.constant 4 : index + // CHECK-NEXT: %[[C4I64:.*]] = arith.index_cast %[[C4:.*]] // CHECK-NEXT: memref.store %[[C4I64]], %{{.*}}[] : memref<i64> scf.for %arg0 = %0 to %c1020 step %1 { %2 = affine.min affine_map<(d0) -> (32, -d0 + 1020)>(%arg0) @@ -58,7 +58,7 @@ %4 = affine.apply affine_map<()[s0] -> (s0 * 4)>()[%count2] scf.parallel (%arg1) = (%3) to (%2) step (%4) { %5 = affine.min affine_map<(d0, d1) -> (4, d0 - d1)>(%2, %arg1) - %6 = index_cast %5: index to i64 + %6 = arith.index_cast %5: index to i64 memref.store %6, %A[]: memref<i64> } }
diff --git a/iree/compiler/Codegen/Common/test/canonicalize_interface_load_store.mlir b/iree/compiler/Codegen/Common/test/canonicalize_interface_load_store.mlir index 4cdb10a..1f231c9 100644 --- a/iree/compiler/Codegen/Common/test/canonicalize_interface_load_store.mlir +++ b/iree/compiler/Codegen/Common/test/canonicalize_interface_load_store.mlir
@@ -2,9 +2,9 @@ // CHECK-LABEL: func @fold_reshape() func @fold_reshape() { - // CHECK: %[[C0:.+]] = constant 0 : index - %c0 = constant 0 : index - %c1 = constant 1 : index + // CHECK: %[[C0:.+]] = arith.constant 0 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index // CHECK: %[[ARG:.+]] = hal.interface.binding.subspan @interface_io::@arg0[%[[C0]]] : !flow.dispatch.tensor<readonly:3x3x96xf32> %1 = hal.interface.binding.subspan @interface_io::@arg0[%c0] : !flow.dispatch.tensor<readonly:3x3x1x96xf32> %2 = hal.interface.binding.subspan @interface_io::@ret0[%c0] : !flow.dispatch.tensor<writeonly:3x3x96xf32> @@ -27,10 +27,10 @@ // CHECK-LABEL: func @dont_fold_reshape_with_not_full_load() func @dont_fold_reshape_with_not_full_load() { - %c0 = constant 0 : index - %c1 = constant 1 : index - %c3 = constant 3 : index - %c96 = constant 96 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + %c3 = arith.constant 3 : index + %c96 = arith.constant 96 : index %1 = hal.interface.binding.subspan @interface_io::@arg0[%c0] : !flow.dispatch.tensor<readonly:6x3x1x96xf32> %2 = hal.interface.binding.subspan @interface_io::@ret0[%c0] : !flow.dispatch.tensor<writeonly:3x3x96xf32> %3 = flow.dispatch.tensor.load %1, offsets = [%c3, %c0, %c0, %c0], sizes = [%c3, %c3, %c1, %c96], strides = [%c1, %c1, %c1, %c1] : !flow.dispatch.tensor<readonly:6x3x1x96xf32> -> tensor<3x3x1x96xf32> @@ -51,8 +51,8 @@ // CHECK-LABEL: func @dont_fold_dynamic_reshape() func @dont_fold_dynamic_reshape() { - %c0 = constant 0 : index - %c1 = constant 1 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index %dim0 = hal.interface.load.constant offset = 0 : index %dim1 = hal.interface.load.constant offset = 1 : index %dim2 = hal.interface.load.constant offset = 2 : index
diff --git a/iree/compiler/Codegen/Common/test/dead_alloc.mlir b/iree/compiler/Codegen/Common/test/dead_alloc.mlir index 7367498..b10dcf1 100644 --- a/iree/compiler/Codegen/Common/test/dead_alloc.mlir +++ b/iree/compiler/Codegen/Common/test/dead_alloc.mlir
@@ -5,8 +5,8 @@ %0 = memref.alloc() : memref<8x64xf32, 3> %1 = memref.subview %0[0, 0] [8, 4] [1, 1] : memref<8x64xf32, 3> to memref<8x4xf32, affine_map<(d0, d1) -> (d0 * 64 + d1)>, 3> - %c0 = constant 0 : index - %cst_0 = constant dense<0.000000e+00> : vector<1x4xf32> + %c0 = arith.constant 0 : index + %cst_0 = arith.constant dense<0.000000e+00> : vector<1x4xf32> vector.transfer_write %cst_0, %1[%c0, %c0] {in_bounds = [true, true]} : vector<1x4xf32>, memref<8x4xf32, affine_map<(d0, d1) -> (d0 * 64 + d1)>, 3> return
diff --git a/iree/compiler/Codegen/Common/test/flatten_memref_subspan.mlir b/iree/compiler/Codegen/Common/test/flatten_memref_subspan.mlir index 3b9f6a9..cd56e4c 100644 --- a/iree/compiler/Codegen/Common/test/flatten_memref_subspan.mlir +++ b/iree/compiler/Codegen/Common/test/flatten_memref_subspan.mlir
@@ -13,8 +13,8 @@ // CHECK: #[[MAP:.+]] = affine_map<()[s0, s1, s2, s3] -> (s0 * 56 + s1 * 8 + s2 + s3 floordiv 4)> // CHECK: func @load_subspan_with_offset // CHECK-SAME: (%[[OFFSET:.+]]: index, %[[I0:.+]]: index, %[[I1:.+]]: index, %[[I2:.+]]: index) -// CHECK-DAG: %[[ZERO:.+]] = constant 0 : index -// CHECK-DAG: %[[C336:.+]] = constant 336 : index +// CHECK-DAG: %[[ZERO:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C336:.+]] = arith.constant 336 : index // CHECK: %[[SUBSPAN:.+]] = hal.interface.binding.subspan @io::@s0b0_ro_constant[%[[ZERO]]] : memref<?xf32>{%[[C336]]} // CHECK: %[[INDEX:.+]] = affine.apply #[[MAP]]()[%[[I0]], %[[I1]], %[[I2]], %[[OFFSET]]] // CHECK: %[[LOAD:.+]] = memref.load %[[SUBSPAN]][%[[INDEX]]] @@ -35,8 +35,8 @@ // CHECK: #[[MAP:.+]] = affine_map<()[s0, s1, s2, s3] -> (s0 * 12 + s1 * 4 + s2 + s3 floordiv 4)> // CHECK: func @store_subspan_with_offset // CHECK-SAME: (%[[VALUE:.+]]: f32, %[[OFFSET:.+]]: index, %[[I0:.+]]: index, %[[I1:.+]]: index, %[[I2:.+]]: index) -// CHECK-DAG: %[[ZERO:.+]] = constant 0 : index -// CHECK-DAG: %[[C24:.+]] = constant 24 : index +// CHECK-DAG: %[[ZERO:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C24:.+]] = arith.constant 24 : index // CHECK: %[[SUBSPAN:.+]] = hal.interface.binding.subspan @io::@s0b0_xw_external[%[[ZERO]]] : memref<?xf32>{%[[C24]]} // CHECK: %[[INDEX:.+]] = affine.apply #[[MAP]]()[%[[I0]], %[[I1]], %[[I2]], %[[OFFSET]]] // CHECK: memref.store %[[VALUE]], %[[SUBSPAN]][%[[INDEX]]] : memref<?xf32> @@ -90,7 +90,7 @@ // CHECK: #[[OFFSET_MAP:.+]] = affine_map<()[s0, s1, s2, s3] -> (s0 * 12 + s1 * 4 + s2 + s3 floordiv 4)> // CHECK: func @store_subspan_with_leading_dynamic_dim // CHECK-SAME: (%[[VALUE:.+]]: f32, %[[OFFSET:.+]]: index, %[[I0:.+]]: index, %[[I1:.+]]: index, %[[I2:.+]]: index) -// CHECK: %[[C0:.+]] = constant 0 : index +// CHECK: %[[C0:.+]] = arith.constant 0 : index // CHECK: %[[DIM:.+]] = hal.interface.load.constant offset = 0 : index // CHECK: %[[SIZE:.+]] = affine.apply #[[SIZE_MAP]]()[%[[DIM]]] // CHECK: %[[DST:.+]] = hal.interface.binding.subspan @io::@s0b0_xw_external[%[[C0]]] : memref<?xf32>{%[[SIZE]]} @@ -117,7 +117,7 @@ // CHECK: #[[OFFSET_MAP:.+]] = affine_map<()[s0, s1, s2, s3, s4, s5, s6, s7] -> (((s4 * s5 + s6) * s2 + s3) * s0 + s1 + s7 floordiv 4)> // CHECK: func @store_subspan_with_all_dynamic_dim // CHECK-SAME: (%[[VALUE:.+]]: f32, %[[OFFSET:.+]]: index, %[[I0:.+]]: index, %[[I1:.+]]: index, %[[I2:.+]]: index, %[[I3:.+]]: index) -// CHECK: %[[C0:.+]] = constant 0 : index +// CHECK: %[[C0:.+]] = arith.constant 0 : index // CHECK: %[[DIM0:.+]] = hal.interface.load.constant offset = 0 : index // CHECK: %[[DIM1:.+]] = hal.interface.load.constant offset = 1 : index // CHECK: %[[DIM2:.+]] = hal.interface.load.constant offset = 2 : index @@ -145,7 +145,7 @@ // CHECK: #[[OFFSET_MAP:.+]] = affine_map<()[s0, s1, s2, s3, s4, s5] -> (((s3 * 4 + s4) * s1 + s2) * 8 + s0 + s5 floordiv 4)> // CHECK: func @store_subspan_with_mixed_dynamic_dim // CHECK-SAME: (%[[VALUE:.+]]: f32, %[[OFFSET:.+]]: index, %[[I0:.+]]: index, %[[I1:.+]]: index, %[[I2:.+]]: index, %[[I3:.+]]: index) -// CHECK: %[[C0:.+]] = constant 0 : index +// CHECK: %[[C0:.+]] = arith.constant 0 : index // CHECK: %[[DIM0:.+]] = hal.interface.load.constant offset = 0 : index // CHECK: %[[DIM2:.+]] = hal.interface.load.constant offset = 1 : index // CHECK: %[[SIZE:.+]] = affine.apply #[[SIZE_MAP]]()[%[[DIM0]], %[[DIM2]]] @@ -172,7 +172,7 @@ // CHECK: #[[OFFSET_MAP:.+]] = affine_map<()[s0, s1, s2, s3] -> (s0 * 12 + s1 * 4 + s2 + s3 floordiv 4)> // CHECK: func @store_subspan_with_flow_control // CHECK-SAME: (%[[VALUE:.+]]: f32, %[[OFFSET:.+]]: index, %[[I0:.+]]: index, %[[I1:.+]]: index, %[[I2:.+]]: index) -// CHECK: %[[C0:.+]] = constant 0 : index +// CHECK: %[[C0:.+]] = arith.constant 0 : index // CHECK: %[[DIM:.+]] = hal.interface.load.constant offset = 0 : index // CHECK: %[[SIZE:.+]] = affine.apply #[[SIZE_MAP]]()[%[[DIM]]] // CHECK: %[[DST:.+]] = hal.interface.binding.subspan @io::@s0b0_xw_external[%[[C0]]] : memref<?xf32>{%[[SIZE]]} @@ -237,7 +237,7 @@ // CHECK: #[[MAP:.+]] = affine_map<()[s0, s1] -> (s0 + s1 floordiv 4)> // CHECK: func @use_subspan_with_unrealized_conversion_cast // CHECK-SAME: (%[[OFFSET:.+]]: index, %[[I:.+]]: index) -// CHECK: %[[C0:.+]] = constant 0 : index +// CHECK: %[[C0:.+]] = arith.constant 0 : index // CHECK: %[[SUBSPAN:.+]] = hal.interface.binding.subspan @io::@s0b0_ro_constant[%[[C0]]] : memref<?xf32> // CHECK: %[[INDEX:.+]] = affine.apply #[[MAP]]()[%[[I]], %[[OFFSET]]] // CHECK: memref.load %[[SUBSPAN]][%[[INDEX]]] @@ -265,7 +265,7 @@ func @transfer_read_subspan_with_offset( %arg0 : index, %arg1: index, %arg2: index, %arg3: index) -> vector<4xf32> { %subspan = hal.interface.binding.subspan @io::@ro[%arg0] : memref<6x7x8xf32> - %cst = constant 0.0 : f32 + %cst = arith.constant 0.0 : f32 %val = vector.transfer_read %subspan[%arg1, %arg2, %arg3], %cst {in_bounds = [true]} : memref<6x7x8xf32>, vector<4xf32> return %val: vector<4xf32> }
diff --git a/iree/compiler/Codegen/Common/test/fold_tensor_extract_op.mlir b/iree/compiler/Codegen/Common/test/fold_tensor_extract_op.mlir index 6bb444e..92f2e96 100644 --- a/iree/compiler/Codegen/Common/test/fold_tensor_extract_op.mlir +++ b/iree/compiler/Codegen/Common/test/fold_tensor_extract_op.mlir
@@ -2,15 +2,15 @@ func @fold_tensor_extract(%arg0 : memref<2x3xi32>) -> i32 { - %c1 = constant 1 : index - %c2 = constant 2 : index + %c1 = arith.constant 1 : index + %c2 = arith.constant 2 : index %0 = memref.tensor_load %arg0 : memref<2x3xi32> %1 = tensor.extract %0[%c1, %c2] : tensor<2x3xi32> return %1 : i32 } // CHECK: func @fold_tensor_extract // CHECK-SAME: %[[ARG0:.+]]: memref<2x3xi32> -// CHECK-DAG: %[[C1:.+]] = constant 1 : index -// CHECK-DAG: %[[C2:.+]] = constant 2 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index +// CHECK-DAG: %[[C2:.+]] = arith.constant 2 : index // CHECK: %[[SCALAR:.+]] = memref.load %[[ARG0]][%[[C1]], %[[C2]]] // CHECK: return %[[SCALAR]]
diff --git a/iree/compiler/Codegen/Common/test/forop_canonicalization.mlir b/iree/compiler/Codegen/Common/test/forop_canonicalization.mlir index a619273..0ef0a35 100644 --- a/iree/compiler/Codegen/Common/test/forop_canonicalization.mlir +++ b/iree/compiler/Codegen/Common/test/forop_canonicalization.mlir
@@ -1,16 +1,16 @@ // RUN: iree-opt %s -split-input-file -iree-codegen-canonicalize-scf-for | FileCheck %s func @loop_carried_vector_shape_cast(%arg0: vector<4xf32>, %arg1: vector<4xf32>) -> (vector<4xf32>, vector<4xf32>) { - %c0 = constant 0 : index - %c1 = constant 1 : index - %c10 = constant 10 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + %c10 = arith.constant 10 : index %0 = vector.shape_cast %arg0 : vector<4xf32> to vector<1x4xf32> %1 = vector.shape_cast %arg1 : vector<4xf32> to vector<1x4xf32> %20:2 = scf.for %arg3 = %c0 to %c10 step %c1 iter_args(%arg4 = %0, %arg5 = %1) -> (vector<1x4xf32>, vector<1x4xf32>) { %a = vector.shape_cast %arg4 : vector<1x4xf32> to vector<4xf32> %b = vector.shape_cast %arg5 : vector<1x4xf32> to vector<4xf32> - %c = addf %a, %b : vector<4xf32> - %d = mulf %a, %b : vector<4xf32> + %c = arith.addf %a, %b : vector<4xf32> + %d = arith.mulf %a, %b : vector<4xf32> %cc = vector.shape_cast %c : vector<4xf32> to vector<1x4xf32> %dc = vector.shape_cast %d : vector<4xf32> to vector<1x4xf32> scf.yield %cc, %dc : vector<1x4xf32>, vector<1x4xf32> @@ -32,16 +32,16 @@ // ----- func @loop_carried_unrealized_conversion_cast(%arg0: vector<4xf32>, %arg1: vector<4xf32>) -> (vector<4xf32>, vector<4xf32>) { - %c0 = constant 0 : index - %c1 = constant 1 : index - %c10 = constant 10 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + %c10 = arith.constant 10 : index %0 = builtin.unrealized_conversion_cast %arg0 : vector<4xf32> to vector<1x4xf32> %1 = builtin.unrealized_conversion_cast %arg1 : vector<4xf32> to vector<1x4xf32> %20:2 = scf.for %arg3 = %c0 to %c10 step %c1 iter_args(%arg4 = %0, %arg5 = %1) -> (vector<1x4xf32>, vector<1x4xf32>) { %a = builtin.unrealized_conversion_cast %arg4 : vector<1x4xf32> to vector<4xf32> %b = builtin.unrealized_conversion_cast %arg5 : vector<1x4xf32> to vector<4xf32> - %c = addf %a, %b : vector<4xf32> - %d = mulf %a, %b : vector<4xf32> + %c = arith.addf %a, %b : vector<4xf32> + %d = arith.mulf %a, %b : vector<4xf32> %cc = builtin.unrealized_conversion_cast %c : vector<4xf32> to vector<1x4xf32> %dc = builtin.unrealized_conversion_cast %d : vector<4xf32> to vector<1x4xf32> scf.yield %cc, %dc : vector<1x4xf32>, vector<1x4xf32> @@ -63,13 +63,13 @@ // ----- func @loop_carried_extract(%arg0: f32) -> f32 { - %c0 = constant 0 : index - %c1 = constant 1 : index - %c10 = constant 10 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + %c10 = arith.constant 10 : index %0 = vector.broadcast %arg0 : f32 to vector<4xf32> %20 = scf.for %arg3 = %c0 to %c10 step %c1 iter_args(%arg4 = %0) -> (vector<4xf32>) { %a = vector.extract %arg4[0] : vector<4xf32> - %c = addf %a, %a : f32 + %c = arith.addf %a, %a : f32 %bc = vector.broadcast %c : f32 to vector<4xf32> scf.yield %bc : vector<4xf32> } @@ -91,14 +91,14 @@ func @loop_pack_v8f16(%arg0: vector<8xf16>, %arg1: vector<8xf16>, %arg2: vector<4xf16>) -> (vector<8xf16>, vector<8xf16>, vector<4xf16>) { - %c0 = constant 0 : index - %c1 = constant 1 : index - %c10 = constant 10 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + %c10 = arith.constant 10 : index %0:3 = scf.for %iv = %c0 to %c10 step %c1 iter_args(%forarg0 = %arg0, %forarg1 = %arg1, %forarg2 = %arg2) -> (vector<8xf16>, vector<8xf16>, vector<4xf16>) { - %add = addf %forarg0, %forarg1: vector<8xf16> + %add = arith.addf %forarg0, %forarg1: vector<8xf16> scf.yield %add, %forarg1, %forarg2: vector<8xf16>, vector<8xf16>, vector<4xf16> } @@ -112,7 +112,7 @@ // CHECK: %[[FOR:.+]]:3 = scf.for %{{.+}} = %{{.+}} to %{{.+}} step %{{.+}} iter_args(%[[FOR_ARG0:.+]] = %[[CAST_ARG0]], %[[FOR_ARG1:.+]] = %[[CAST_ARG1]], %[[FOR_ARG2:.+]] = %[[ARG2]]) -> (vector<4xf32>, vector<4xf32>, vector<4xf16>) { // CHECK: %[[CAST_FOR_ARG0:.+]] = vector.bitcast %[[FOR_ARG0]] : vector<4xf32> to vector<8xf16> // CHECK: %[[CAST_FOR_ARG1:.+]] = vector.bitcast %[[FOR_ARG1]] : vector<4xf32> to vector<8xf16> -// CHECK: %[[ADD:.+]] = addf %[[CAST_FOR_ARG0]], %[[CAST_FOR_ARG1]] : vector<8xf16> +// CHECK: %[[ADD:.+]] = arith.addf %[[CAST_FOR_ARG0]], %[[CAST_FOR_ARG1]] : vector<8xf16> // CHECK: %[[CAST_ADD:.+]] = vector.bitcast %[[ADD]] : vector<8xf16> to vector<4xf32> // CHECK: scf.yield %[[CAST_ADD]], %[[FOR_ARG1]], %[[FOR_ARG2]] : vector<4xf32>, vector<4xf32>, vector<4xf16> // CHECK: }
diff --git a/iree/compiler/Codegen/Common/test/linalg_bufferize.mlir b/iree/compiler/Codegen/Common/test/linalg_bufferize.mlir index fad78a0..c844f50 100644 --- a/iree/compiler/Codegen/Common/test/linalg_bufferize.mlir +++ b/iree/compiler/Codegen/Common/test/linalg_bufferize.mlir
@@ -1,11 +1,11 @@ // RUN: iree-opt %s --iree-codegen-linalg-bufferize -canonicalize -cse -split-input-file | IreeFileCheck %s func @tile_from_tensor_load() { - %c0 = constant 0 : index - %c2 = constant 2 : index - %c4 = constant 4 : index - %c1 = constant 1 : index - %c3 = constant 3 : index + %c0 = arith.constant 0 : index + %c2 = arith.constant 2 : index + %c4 = arith.constant 4 : index + %c1 = arith.constant 1 : index + %c3 = arith.constant 3 : index %M = hal.interface.load.constant offset = 0 : index %N = hal.interface.load.constant offset = 1 : index %K = hal.interface.load.constant offset = 2 : index @@ -52,11 +52,11 @@ // ----- func @tile_from_tensor_load_inplace() { - %c0 = constant 0 : index - %c2 = constant 2 : index - %c4 = constant 4 : index - %c1 = constant 1 : index - %c3 = constant 3 : index + %c0 = arith.constant 0 : index + %c2 = arith.constant 2 : index + %c4 = arith.constant 4 : index + %c1 = arith.constant 1 : index + %c3 = arith.constant 3 : index %M = hal.interface.load.constant offset = 0 : index %N = hal.interface.load.constant offset = 1 : index %K = hal.interface.load.constant offset = 2 : index @@ -98,11 +98,11 @@ // ----- func @tile_from_tensor_load_inplace_and_copy() { - %c0 = constant 0 : index - %c2 = constant 2 : index - %c4 = constant 4 : index - %c1 = constant 1 : index - %c3 = constant 3 : index + %c0 = arith.constant 0 : index + %c2 = arith.constant 2 : index + %c4 = arith.constant 4 : index + %c1 = arith.constant 1 : index + %c3 = arith.constant 3 : index %M = hal.interface.load.constant offset = 0 : index %N = hal.interface.load.constant offset = 1 : index %K = hal.interface.load.constant offset = 2 : index @@ -151,11 +151,11 @@ #map = affine_map<(d0, d1) -> (d0, d1)> func @tile_from_pointwise_lhs() { - %c0 = constant 0 : index - %c2 = constant 2 : index - %c4 = constant 4 : index - %c1 = constant 1 : index - %c3 = constant 3 : index + %c0 = arith.constant 0 : index + %c2 = arith.constant 2 : index + %c4 = arith.constant 4 : index + %c1 = arith.constant 1 : index + %c3 = arith.constant 3 : index %M = hal.interface.load.constant offset = 0 : index %N = hal.interface.load.constant offset = 1 : index %K = hal.interface.load.constant offset = 2 : index @@ -212,11 +212,11 @@ #map = affine_map<(d0, d1) -> (d0, d1)> func @tile_from_pointwise_lhs_inplace() { - %c0 = constant 0 : index - %c2 = constant 2 : index - %c4 = constant 4 : index - %c1 = constant 1 : index - %c3 = constant 3 : index + %c0 = arith.constant 0 : index + %c2 = arith.constant 2 : index + %c4 = arith.constant 4 : index + %c1 = arith.constant 1 : index + %c3 = arith.constant 3 : index %M = hal.interface.load.constant offset = 0 : index %N = hal.interface.load.constant offset = 1 : index %K = hal.interface.load.constant offset = 2 : index @@ -269,11 +269,11 @@ #map = affine_map<(d0, d1) -> (d0, d1)> func @tile_from_pointwise_outs() { - %c0 = constant 0 : index - %c2 = constant 2 : index - %c4 = constant 4 : index - %c1 = constant 1 : index - %c3 = constant 3 : index + %c0 = arith.constant 0 : index + %c2 = arith.constant 2 : index + %c4 = arith.constant 4 : index + %c1 = arith.constant 1 : index + %c3 = arith.constant 3 : index %M = hal.interface.load.constant offset = 0 : index %N = hal.interface.load.constant offset = 1 : index %K = hal.interface.load.constant offset = 2 : index @@ -328,11 +328,11 @@ #map = affine_map<(d0, d1) -> (d0, d1)> func @tile_from_pointwise_outs_inplace() { - %c0 = constant 0 : index - %c2 = constant 2 : index - %c4 = constant 4 : index - %c1 = constant 1 : index - %c3 = constant 3 : index + %c0 = arith.constant 0 : index + %c2 = arith.constant 2 : index + %c4 = arith.constant 4 : index + %c1 = arith.constant 1 : index + %c3 = arith.constant 3 : index %M = hal.interface.load.constant offset = 0 : index %N = hal.interface.load.constant offset = 1 : index %K = hal.interface.load.constant offset = 2 : index @@ -382,11 +382,11 @@ #map = affine_map<(d0, d1) -> (d0, d1)> func @tile_from_matmul_outs() { - %c0 = constant 0 : index - %c2 = constant 2 : index - %c4 = constant 4 : index - %c1 = constant 1 : index - %c3 = constant 3 : index + %c0 = arith.constant 0 : index + %c2 = arith.constant 2 : index + %c4 = arith.constant 4 : index + %c1 = arith.constant 1 : index + %c3 = arith.constant 3 : index %M = hal.interface.load.constant offset = 0 : index %N = hal.interface.load.constant offset = 1 : index %K = hal.interface.load.constant offset = 2 : index @@ -436,11 +436,11 @@ #map = affine_map<(d0, d1) -> (d0, d1)> func @tile_from_matmul_outs_inplace() { - %c0 = constant 0 : index - %c2 = constant 2 : index - %c4 = constant 4 : index - %c1 = constant 1 : index - %c3 = constant 3 : index + %c0 = arith.constant 0 : index + %c2 = arith.constant 2 : index + %c4 = arith.constant 4 : index + %c1 = arith.constant 1 : index + %c3 = arith.constant 3 : index %M = hal.interface.load.constant offset = 0 : index %N = hal.interface.load.constant offset = 1 : index %K = hal.interface.load.constant offset = 2 : index @@ -484,8 +484,8 @@ // ----- func @bufferize_dynamic() { - %c0 = constant 0 : index - %c1 = constant 1 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index %4 = hal.interface.load.constant offset = 0 : index %5 = hal.interface.load.constant offset = 1 : index %6 = hal.interface.load.constant offset = 2 : index @@ -504,11 +504,11 @@ %workgroup_count_x = hal.interface.workgroup.count[0] : index %workgroup_id_y = hal.interface.workgroup.id[1] : index %workgroup_count_y = hal.interface.workgroup.count[1] : index - %20 = muli %workgroup_size_y, %workgroup_id_y : index - %21 = muli %workgroup_size_y, %workgroup_count_y : index + %20 = arith.muli %workgroup_size_y, %workgroup_id_y : index + %21 = arith.muli %workgroup_size_y, %workgroup_count_y : index scf.for %arg0 = %20 to %4 step %21 { - %22 = muli %workgroup_size_x, %workgroup_id_x : index - %23 = muli %workgroup_size_x, %workgroup_count_x : index + %22 = arith.muli %workgroup_size_x, %workgroup_id_x : index + %23 = arith.muli %workgroup_size_x, %workgroup_count_x : index scf.for %arg1 = %22 to %7 step %23 { %24 = affine.min affine_map<(d0)[s0, s1] -> (s1, -d0 + s0)>(%arg0)[%4, %workgroup_size_y] %25 = flow.dispatch.tensor.load %LHS, offsets = [%arg0, %c0], sizes = [%24, %5], strides = [%c1, %c1] : !flow.dispatch.tensor<readonly:?x?xf32> -> tensor<?x?xf32> @@ -564,8 +564,8 @@ // ----- func @bufferize_dynamic_inplace() { - %c0 = constant 0 : index - %c1 = constant 1 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index %4 = hal.interface.load.constant offset = 0 : index %5 = hal.interface.load.constant offset = 1 : index %6 = hal.interface.load.constant offset = 2 : index @@ -581,11 +581,11 @@ %workgroup_count_x = hal.interface.workgroup.count[0] : index %workgroup_id_y = hal.interface.workgroup.id[1] : index %workgroup_count_y = hal.interface.workgroup.count[1] : index - %20 = muli %workgroup_size_y, %workgroup_id_y : index - %21 = muli %workgroup_size_y, %workgroup_count_y : index + %20 = arith.muli %workgroup_size_y, %workgroup_id_y : index + %21 = arith.muli %workgroup_size_y, %workgroup_count_y : index scf.for %arg0 = %20 to %4 step %21 { - %22 = muli %workgroup_size_x, %workgroup_id_x : index - %23 = muli %workgroup_size_x, %workgroup_count_x : index + %22 = arith.muli %workgroup_size_x, %workgroup_id_x : index + %23 = arith.muli %workgroup_size_x, %workgroup_count_x : index scf.for %arg1 = %22 to %7 step %23 { %24 = affine.min affine_map<(d0)[s0, s1] -> (s1, -d0 + s0)>(%arg0)[%4, %workgroup_size_y] %25 = flow.dispatch.tensor.load %LHS, offsets = [%arg0, %c0], sizes = [%24, %5], strides = [%c1, %c1] : !flow.dispatch.tensor<readonly:?x?xf32> -> tensor<?x?xf32> @@ -635,11 +635,11 @@ // ----- func @reshape_simple() { - %c0 = constant 0 : index - %c1 = constant 1 : index - %c3 = constant 3 : index - %c4 = constant 4 : index - %c12 = constant 12 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + %c3 = arith.constant 3 : index + %c4 = arith.constant 4 : index + %c12 = arith.constant 12 : index %0 = hal.interface.binding.subspan @io::@arg0[%c0] : !flow.dispatch.tensor<readonly:12xi32> %1 = hal.interface.binding.subspan @io::@ret0[%c0] : !flow.dispatch.tensor<writeonly:3x4xi32> %2 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:12xi32> -> tensor<12xi32> @@ -660,11 +660,11 @@ // ----- func @reshape_fused_source() { - %c0 = constant 0 : index - %c1 = constant 1 : index - %c3 = constant 3 : index - %c4 = constant 4 : index - %c12 = constant 12 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + %c3 = arith.constant 3 : index + %c4 = arith.constant 4 : index + %c12 = arith.constant 12 : index %0 = hal.interface.binding.subspan @io::@arg0[%c0] : !flow.dispatch.tensor<readonly:12xi32> %1 = hal.interface.binding.subspan @io::@ret0[%c0] : !flow.dispatch.tensor<writeonly:3x4xi32> %2 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:12xi32> -> tensor<12xi32> @@ -675,7 +675,7 @@ iterator_types = ["parallel", "parallel"]} ins(%3 : tensor<3x4xi32>) outs(%4 : tensor<3x4xi32>) { ^bb0(%arg0 : i32, %arg1 : i32): - %6 = addi %arg0, %arg0 : i32 + %6 = arith.addi %arg0, %arg0 : i32 linalg.yield %6 : i32 } -> tensor<3x4xi32> flow.dispatch.tensor.store %5, %1, offsets = [], sizes = [], strides = [] : tensor<3x4xi32> -> !flow.dispatch.tensor<writeonly:3x4xi32> @@ -686,7 +686,7 @@ hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard" } // CHECK: func @reshape_fused_source() -// CHECK: %[[C0:.+]] = constant 0 +// CHECK: %[[C0:.+]] = arith.constant 0 // CHECK-DAG: %[[ARG0:.+]] = hal.interface.binding.subspan @io::@arg0[%[[C0]]] : memref<12xi32> // CHECK-DAG: %[[RET0:.+]] = hal.interface.binding.subspan @io::@ret0[%[[C0]]] : memref<3x4xi32> // CHECK: %[[RESHAPE:.+]] = memref.expand_shape %[[ARG0]] {{\[}}[0, 1]] @@ -697,11 +697,11 @@ // ----- func @reshape_fused_source_and_copyout() { - %c0 = constant 0 : index - %c1 = constant 1 : index - %c3 = constant 3 : index - %c4 = constant 4 : index - %c12 = constant 12 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + %c3 = arith.constant 3 : index + %c4 = arith.constant 4 : index + %c12 = arith.constant 12 : index %0 = hal.interface.binding.subspan @io::@arg0[%c0] : !flow.dispatch.tensor<readonly:12xi32> %1 = hal.interface.binding.subspan @io::@ret0[%c0] : !flow.dispatch.tensor<writeonly:3x4xi32> %2 = hal.interface.binding.subspan @io::@ret1[%c0] : !flow.dispatch.tensor<writeonly:3x4xi32> @@ -713,7 +713,7 @@ iterator_types = ["parallel", "parallel"]} ins(%4 : tensor<3x4xi32>) outs(%5 : tensor<3x4xi32>) { ^bb0(%arg0 : i32, %arg1 : i32): - %7 = addi %arg0, %arg0 : i32 + %7 = arith.addi %arg0, %arg0 : i32 linalg.yield %7 : i32 } -> tensor<3x4xi32> flow.dispatch.tensor.store %6, %1, offsets = [], sizes = [], strides = [] : tensor<3x4xi32> -> !flow.dispatch.tensor<writeonly:3x4xi32> @@ -726,7 +726,7 @@ hal.interface.binding @ret1, set=0, binding=2, type="StorageBuffer", access="Write|Discard" } // CHECK: func @reshape_fused_source_and_copyout() -// CHECK: %[[C0:.+]] = constant 0 +// CHECK: %[[C0:.+]] = arith.constant 0 // CHECK-DAG: %[[ARG0:.+]] = hal.interface.binding.subspan @io::@arg0[%[[C0]]] : memref<12xi32> // CHECK-DAG: %[[RET0:.+]] = hal.interface.binding.subspan @io::@ret0[%[[C0]]] : memref<3x4xi32> // CHECK-DAG: %[[RET1:.+]] = hal.interface.binding.subspan @io::@ret1[%[[C0]]] : memref<3x4xi32> @@ -739,11 +739,11 @@ // ----- func @reshape_fused_target() { - %c0 = constant 0 : index - %c1 = constant 1 : index - %c3 = constant 3 : index - %c4 = constant 4 : index - %c12 = constant 12 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + %c3 = arith.constant 3 : index + %c4 = arith.constant 4 : index + %c12 = arith.constant 12 : index %0 = hal.interface.binding.subspan @io::@arg0[%c0] : !flow.dispatch.tensor<readonly:3x4xi32> %1 = hal.interface.binding.subspan @io::@ret0[%c0] : !flow.dispatch.tensor<writeonly:12xi32> %2 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:3x4xi32> -> tensor<3x4xi32> @@ -753,7 +753,7 @@ iterator_types = ["parallel", "parallel"]} ins(%2 : tensor<3x4xi32>) outs(%3 : tensor<3x4xi32>) { ^bb0(%arg0 : i32, %arg1 : i32): - %5 = addi %arg0, %arg0 : i32 + %5 = arith.addi %arg0, %arg0 : i32 linalg.yield %5 : i32 } -> tensor<3x4xi32> %5 = linalg.tensor_collapse_shape %4 [[0, 1]] : tensor<3x4xi32> into tensor<12xi32> @@ -765,7 +765,7 @@ hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard" } // CHECK: func @reshape_fused_target() -// CHECK: %[[C0:.+]] = constant 0 +// CHECK: %[[C0:.+]] = arith.constant 0 // CHECK-DAG: %[[ARG0:.+]] = hal.interface.binding.subspan @io::@arg0[%[[C0]]] : memref<3x4xi32> // CHECK-DAG: %[[RET0:.+]] = hal.interface.binding.subspan @io::@ret0[%[[C0]]] : memref<12xi32> // CHECK: %[[RESHAPE:.+]] = memref.expand_shape %[[RET0]] {{\[}}[0, 1]] @@ -776,11 +776,11 @@ // ----- func @dot_general_lowering() { - %cst = constant 0.000000e+00 : f32 - %c3 = constant 3 : index - %c0 = constant 0 : index - %c2 = constant 2 : index - %c1 = constant 1 : index + %cst = arith.constant 0.000000e+00 : f32 + %c3 = arith.constant 3 : index + %c0 = arith.constant 0 : index + %c2 = arith.constant 2 : index + %c1 = arith.constant 1 : index %0 = hal.interface.binding.subspan @io::@arg0[%c0] : !flow.dispatch.tensor<readonly:1x1x2xf32> %1 = hal.interface.binding.subspan @io::@arg1[%c0] : !flow.dispatch.tensor<readonly:2x3xf32> %2 = hal.interface.binding.subspan @io::@ret0[%c0] : !flow.dispatch.tensor<writeonly:1x3xf32> @@ -792,11 +792,11 @@ %workgroup_count_x = hal.interface.workgroup.count[0] : index %workgroup_id_y = hal.interface.workgroup.id[1] : index %workgroup_count_y = hal.interface.workgroup.count[1] : index - %5 = muli %workgroup_size_y, %workgroup_id_y : index - %6 = muli %workgroup_size_y, %workgroup_count_y : index + %5 = arith.muli %workgroup_size_y, %workgroup_id_y : index + %6 = arith.muli %workgroup_size_y, %workgroup_count_y : index scf.for %arg0 = %5 to %c1 step %6 { - %7 = muli %workgroup_size_x, %workgroup_id_x : index - %8 = muli %workgroup_size_x, %workgroup_count_x : index + %7 = arith.muli %workgroup_size_x, %workgroup_id_x : index + %8 = arith.muli %workgroup_size_x, %workgroup_count_x : index scf.for %arg1 = %7 to %c3 step %8 { %9 = affine.min affine_map<(d0)[s0] -> (s0, -d0 + 1)>(%arg0)[%workgroup_size_y] %10 = tensor.extract_slice %4[%arg0, 0] [%9, 2] [1, 1] : tensor<1x2xf32> to tensor<?x2xf32> @@ -833,7 +833,7 @@ // ----- func @slice() { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %2 = hal.interface.load.constant offset = 0 : index %3 = hal.interface.load.constant offset = 1 : index %4 = hal.interface.load.constant offset = 2 : index @@ -858,7 +858,7 @@ // ----- func @slice_rank_reducing() { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %2 = hal.interface.load.constant offset = 0 : index %3 = hal.interface.load.constant offset = 1 : index %4 = hal.interface.load.constant offset = 2 : index @@ -884,7 +884,7 @@ // ----- func @slice_multiple_copy() { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %3 = hal.interface.load.constant offset = 0 : index %4 = hal.interface.load.constant offset = 1 : index %5 = hal.interface.load.constant offset = 2 : index @@ -923,7 +923,7 @@ // ----- func @slice_in_place() { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %2 = hal.interface.load.constant offset = 0 : index %3 = hal.interface.load.constant offset = 1 : index %0 = hal.interface.binding.subspan @io::@arg0[%c0] : !flow.dispatch.tensor<readwrite:?x?xi32>{%2, %3} @@ -941,7 +941,7 @@ // ----- func @slice_whole_stride_dispatch_0() { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %dim0 = hal.interface.load.constant offset = 0 : index %dim1 = hal.interface.load.constant offset = 1 : index %dim2 = hal.interface.load.constant offset = 2 : index @@ -966,8 +966,8 @@ // ----- func @subtensor_insert() { - %c0 = constant 0 : index - %c1 = constant 1 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index %dim0 = hal.interface.load.constant offset = 0 : index %dim1 = hal.interface.load.constant offset = 1 : index %dim2 = hal.interface.load.constant offset = 2 : index @@ -991,8 +991,8 @@ hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard" } // CHECK-LABEL: func @subtensor_insert() -// CHECK-DAG: %[[C0:.+]] = constant 0 -// CHECK-DAG: %[[C1:.+]] = constant 1 +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 // CHECK-DAG: %[[ARG0:.+]] = hal.interface.binding.subspan @io::@arg0 // CHECK-DAG: %[[ARG1:.+]] = hal.interface.binding.subspan @io::@arg1 // CHECK-DAG: %[[RET0:.+]] = hal.interface.binding.subspan @io::@ret0 @@ -1005,7 +1005,7 @@ // ----- func @tensor_extract() { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %0 = hal.interface.binding.subspan @io::@arg0[%c0] : !flow.dispatch.tensor<readonly:i32> %1 = hal.interface.binding.subspan @io::@ret0[%c0] : !flow.dispatch.tensor<writeonly:3x9xi32> %2 = linalg.init_tensor [3, 9] : tensor<3x9xi32> @@ -1028,7 +1028,7 @@ // ----- func @load_to_store() { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %1 = hal.interface.binding.subspan @io::@ret0[%c0] : !flow.dispatch.tensor<writeonly:3x4xi32> %2 = hal.interface.binding.subspan @io::@arg0[%c0] : !flow.dispatch.tensor<readonly:3x4xi32> %3 = flow.dispatch.tensor.load %2, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:3x4xi32> -> tensor<3x4xi32> @@ -1049,14 +1049,14 @@ // ----- func @constant() { - %c0 = constant 0 : index - %cst = constant dense<[[[1, 2, 3], [4, 5, 6]], [[7, 8, 9], [10, 11, 12]]]> : tensor<2x2x3xi32> + %c0 = arith.constant 0 : index + %cst = arith.constant dense<[[[1, 2, 3], [4, 5, 6]], [[7, 8, 9], [10, 11, 12]]]> : tensor<2x2x3xi32> %0 = hal.interface.binding.subspan @io::@ret0[%c0] : !flow.dispatch.tensor<writeonly:2x2x3xi32> flow.dispatch.tensor.store %cst, %0, offsets = [], sizes = [], strides = [] : tensor<2x2x3xi32> -> !flow.dispatch.tensor<writeonly:2x2x3xi32> return } // CHECK-LABEL: func @constant() -// CHECK: %[[CST:.+]] = constant {{.+}} : tensor<2x2x3xi32> +// CHECK: %[[CST:.+]] = arith.constant {{.+}} : tensor<2x2x3xi32> // CHECK: %[[MEMREF:.+]] = memref.buffer_cast %[[CST]] : memref<2x2x3xi32> // CHECK: %[[RESULT:.+]] = hal.interface.binding.subspan @io::@ret0 // CHECK: linalg.copy(%[[MEMREF]], %[[RESULT]]) @@ -1064,11 +1064,11 @@ // ----- func @rhs_non_splat_constant() { - %c0 = constant 0 : index - %cst = constant dense<[[0.706495285, -0.567672312, 0.483717591, 0.522725761, 0.7563259], [-0.0899272263, -0.283501834, -0.350822538, -0.351515919, -0.337136656], [-0.451804549, 0.372324884, -0.620518147, 0.235451385, 0.851095855]]> : tensor<3x5xf32> - %cst_0 = constant 0.000000e+00 : f32 - %c5 = constant 5 : index - %c1 = constant 1 : index + %c0 = arith.constant 0 : index + %cst = arith.constant dense<[[0.706495285, -0.567672312, 0.483717591, 0.522725761, 0.7563259], [-0.0899272263, -0.283501834, -0.350822538, -0.351515919, -0.337136656], [-0.451804549, 0.372324884, -0.620518147, 0.235451385, 0.851095855]]> : tensor<3x5xf32> + %cst_0 = arith.constant 0.000000e+00 : f32 + %c5 = arith.constant 5 : index + %c1 = arith.constant 1 : index %0 = hal.interface.binding.subspan @io::@arg0[%c0] : !flow.dispatch.tensor<readonly:1x5x3x1xf32> %1 = hal.interface.binding.subspan @io::@ret0[%c0] : !flow.dispatch.tensor<writeonly:5x5xf32> %2 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:1x5x3x1xf32> -> tensor<1x5x3x1xf32> @@ -1102,7 +1102,7 @@ hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard" } // CHECK-LABEL: func @rhs_non_splat_constant -// CHECK-DAG: %[[CONSTANT:.+]] = constant {{.+}} : tensor<3x5xf32> +// CHECK-DAG: %[[CONSTANT:.+]] = arith.constant {{.+}} : tensor<3x5xf32> // CHECK-DAG: %[[RHS:.+]] = memref.buffer_cast %[[CONSTANT]] // CHECK-DAG: %[[LHS_INPUT:.+]] = hal.interface.binding.subspan @io::@arg0[%{{.+}}] : memref<1x5x3x1xf32> // CHECK-DAG: %[[RETURN:.+]] = hal.interface.binding.subspan @io::@ret0[%{{.+}}] : memref<5x5xf32> @@ -1120,8 +1120,8 @@ // ----- func @gather() { - %c0 = constant 0 : index - %c1 = constant 1 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index %dim0 = hal.interface.load.constant offset = 0 : index %dim1 = hal.interface.load.constant offset = 1 : index %dim2 = hal.interface.load.constant offset = 2 : index @@ -1138,7 +1138,7 @@ %7 = linalg.generic {indexing_maps = [affine_map<(d0, d1) -> (d0)>, affine_map<(d0, d1) -> (d0, d1)>], iterator_types = ["parallel", "parallel"]} ins(%5 : tensor<?xi32>) outs(%3 : tensor<?x?xf32>) { ^bb0( %arg2: i32, %arg3: f32): // no predecessors %iv1 = linalg.index 1 : index - %8 = index_cast %arg2 : i32 to index + %8 = arith.index_cast %arg2 : i32 to index %9 = tensor.extract %4[%8, %iv1] : tensor<?x?xf32> linalg.yield %9 : f32 } -> tensor<?x?xf32> @@ -1161,9 +1161,9 @@ // ----- func @pooling_nhwc_sum() { - %c2 = constant 2 : index - %c0 = constant 0 : index - %c1 = constant 1 : index + %c2 = arith.constant 2 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index %0 = hal.interface.binding.subspan @io::@ro0[%c0] : !flow.dispatch.tensor<readonly:f32> %1 = hal.interface.binding.subspan @io::@ro1[%c0] : !flow.dispatch.tensor<readonly:1x4x6x1xf32> %2 = hal.interface.binding.subspan @io::@wo2[%c0] : !flow.dispatch.tensor<writeonly:1x2x2x1xf32> @@ -1202,8 +1202,8 @@ // ----- func @read_only_subtensor() { - %c0 = constant 0 : index - %c1 = constant 1 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index %pc0 = hal.interface.load.constant offset = 0 : index %pc1 = hal.interface.load.constant offset = 1 : index %pc2 = hal.interface.load.constant offset = 2 : index @@ -1246,9 +1246,9 @@ %23 = linalg.init_tensor [%21, %22] : tensor<?x?xf32> %24 = linalg.generic {indexing_maps = [affine_map<(d0, d1) -> (d0, d1)>, affine_map<(d0, d1) -> (d0, d1)>, affine_map<(d0, d1) -> (d0, d1)>, affine_map<(d0, d1) -> (d0, d1)>, affine_map<(d0, d1) -> (d0, d1)>], iterator_types = ["parallel", "parallel"]} ins(%11, %14, %17, %20 : tensor<?x?xf32>, tensor<?x?xf32>, tensor<?x?xf32>, tensor<?x?xf32>) outs(%23 : tensor<?x?xf32>) attrs = {__internal_linalg_transform__ = "workgroup"} { ^bb0(%arg2: f32, %arg3: f32, %arg4: f32, %arg5: f32, %arg6: f32): // no predecessors - %25 = mulf %arg4, %arg5 : f32 - %26 = mulf %arg2, %arg3 : f32 - %27 = addf %26, %25 : f32 + %25 = arith.mulf %arg4, %arg5 : f32 + %26 = arith.mulf %arg2, %arg3 : f32 + %27 = arith.addf %26, %25 : f32 %28 = math.sqrt %27 : f32 linalg.yield %28 : f32 } -> tensor<?x?xf32> @@ -1277,7 +1277,7 @@ // ----- func @reshape_read_only() { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %dim0 = hal.interface.load.constant offset = 0 : index %dim1 = hal.interface.load.constant offset = 1 : index %dim2 = hal.interface.load.constant offset = 2 : index @@ -1293,7 +1293,7 @@ iterator_types = ["parallel"]} ins(%3 : tensor<?xf32>) outs(%5 : tensor<?xf32>) { ^bb0(%arg0 : f32, %arg1 : f32): - %7 = addf %arg0, %arg0 : f32 + %7 = arith.addf %arg0, %arg0 : f32 linalg.yield %7 : f32 } -> tensor<?xf32> flow.dispatch.tensor.store %6, %1, offsets = [], sizes = [], strides = []: tensor<?xf32> -> !flow.dispatch.tensor<writeonly:?xf32> @@ -1310,7 +1310,7 @@ // ----- func @use_buffer_for_operand_when_output_tensor_not_used() { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %input_subspan = hal.interface.binding.subspan @interface_io::@ro0[%c0] : !flow.dispatch.tensor<readonly:1x225x225x16xf32> %filter_subspan = hal.interface.binding.subspan @interface_io::@ro1[%c0] : !flow.dispatch.tensor<readonly:3x3x16x32xf32> @@ -1321,7 +1321,7 @@ %filter = flow.dispatch.tensor.load %filter_subspan, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:3x3x16x32xf32> -> tensor<3x3x16x32xf32> %offset = flow.dispatch.tensor.load %offset_subspan, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:32xf32> -> tensor<32xf32> - %cst = constant 0.0 : f32 + %cst = arith.constant 0.0 : f32 %0 = linalg.init_tensor [1, 112, 112, 32] : tensor<1x112x112x32xf32> %1 = linalg.fill(%cst, %0) : f32, tensor<1x112x112x32xf32> -> tensor<1x112x112x32xf32> %2 = linalg.conv_2d_nhwc_hwcf @@ -1338,7 +1338,7 @@ ins(%2, %offset: tensor<1x112x112x32xf32>, tensor<32xf32>) outs(%0 : tensor<1x112x112x32xf32>) { ^bb0(%a: f32, %b: f32, %c: f32): - %sub = subf %a, %b : f32 + %sub = arith.subf %a, %b : f32 linalg.yield %sub : f32 } -> tensor<1x112x112x32xf32> flow.dispatch.tensor.store %3, %output_subspan, offsets = [], sizes = [], strides = [] : tensor<1x112x112x32xf32> -> !flow.dispatch.tensor<writeonly:1x112x112x32xf32> @@ -1366,7 +1366,7 @@ // ----- func @dont_use_buffer_for_operand_when_output_tensor_used() { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %input_subspan = hal.interface.binding.subspan @interface_io::@ro0[%c0] : !flow.dispatch.tensor<readonly:1x225x225x16xf32> %filter_subspan = hal.interface.binding.subspan @interface_io::@ro1[%c0] : !flow.dispatch.tensor<readonly:3x3x16x32xf32> @@ -1377,8 +1377,8 @@ %filter = flow.dispatch.tensor.load %filter_subspan, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:3x3x16x32xf32> -> tensor<3x3x16x32xf32> %offset = flow.dispatch.tensor.load %offset_subspan, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:32xf32> -> tensor<32xf32> - %cst0 = constant 0.0 : f32 - %cst1 = constant 1.0 : f32 + %cst0 = arith.constant 0.0 : f32 + %cst1 = arith.constant 1.0 : f32 %0 = linalg.init_tensor [1, 112, 112, 32] : tensor<1x112x112x32xf32> %1 = linalg.fill(%cst0, %0) : f32, tensor<1x112x112x32xf32> -> tensor<1x112x112x32xf32> %2 = linalg.conv_2d_nhwc_hwcf @@ -1396,8 +1396,8 @@ ins(%2, %offset: tensor<1x112x112x32xf32>, tensor<32xf32>) outs(%3 : tensor<1x112x112x32xf32>) { ^bb0(%a: f32, %b: f32, %c: f32): - %sub = subf %a, %b : f32 - %add = addf %sub, %c : f32 + %sub = arith.subf %a, %b : f32 + %add = arith.addf %sub, %c : f32 linalg.yield %add : f32 } -> tensor<1x112x112x32xf32> flow.dispatch.tensor.store %4, %output_subspan, offsets = [], sizes = [], strides = []: tensor<1x112x112x32xf32> -> !flow.dispatch.tensor<writeonly:1x112x112x32xf32> @@ -1418,10 +1418,10 @@ // ----- func @bufferize_cst_output_tensor() { - %c0 = constant 0 : index - %cst1 = constant dense<-2147483648> : tensor<i32> - %zero = constant 0.000000e+00 : f32 - %cst5 = constant dense<[1, 2, 3, 4, 5]> : tensor<5xi32> + %c0 = arith.constant 0 : index + %cst1 = arith.constant dense<-2147483648> : tensor<i32> + %zero = arith.constant 0.000000e+00 : f32 + %cst5 = arith.constant dense<[1, 2, 3, 4, 5]> : tensor<5xi32> %input = hal.interface.binding.subspan @io::@ro0[%c0] : !flow.dispatch.tensor<readonly:5xf32> %output = hal.interface.binding.subspan @io::@wo1[%c0] : !flow.dispatch.tensor<writeonly:i32> %1 = flow.dispatch.tensor.load %input, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:5xf32> -> tensor<5xf32> @@ -1431,10 +1431,10 @@ ins(%1, %cst5 : tensor<5xf32>, tensor<5xi32>) outs(%cst1 : tensor<i32>) { ^bb0(%arg0: f32, %arg1: i32, %arg2: i32): - %8 = cmpf oeq, %arg0, %zero : f32 - %9 = zexti %8 : i1 to i32 - %10 = muli %9, %arg1 : i32 - %11 = cmpi sgt, %10, %arg2 : i32 + %8 = arith.cmpf oeq, %arg0, %zero : f32 + %9 = arith.extui %8 : i1 to i32 + %10 = arith.muli %9, %arg1 : i32 + %11 = arith.cmpi sgt, %10, %arg2 : i32 %12 = select %11, %10, %arg2 : i32 linalg.yield %12 : i32 } -> tensor<i32> @@ -1449,8 +1449,8 @@ // CHECK-LABEL: func @bufferize_cst_output_tensor() -// CHECK-DAG: %[[CST1:.+]] = constant dense<-2147483648> : tensor<i32> -// CHECK-DAG: %[[CST5:.+]] = constant dense<[1, 2, 3, 4, 5]> : tensor<5xi32> +// CHECK-DAG: %[[CST1:.+]] = arith.constant dense<-2147483648> : tensor<i32> +// CHECK-DAG: %[[CST5:.+]] = arith.constant dense<[1, 2, 3, 4, 5]> : tensor<5xi32> // CHECK: %[[CAST1:.+]] = memref.buffer_cast %[[CST1]] : memref<i32> // CHECK: %[[CAST5:.+]] = memref.buffer_cast %[[CST5]] : memref<5xi32> // CHECK: %[[INPUT:.+]] = hal.interface.binding.subspan @io::@ro0[%c0] : memref<5xf32> @@ -1463,12 +1463,12 @@ // ----- func @cast_follwed_by_store() { - %c0 = constant 0 : index - %cst = constant 0.000000e+00 : f32 - %c4 = constant 4 : index - %c64 = constant 64 : index - %c1 = constant 1 : index - %c32 = constant 32 : index + %c0 = arith.constant 0 : index + %cst = arith.constant 0.000000e+00 : f32 + %c4 = arith.constant 4 : index + %c64 = arith.constant 64 : index + %c1 = arith.constant 1 : index + %c32 = arith.constant 32 : index %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : !flow.dispatch.tensor<readonly:4x32x1024xf32> %1 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : !flow.dispatch.tensor<readonly:4x1024x64xf32> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : !flow.dispatch.tensor<writeonly:4x32x64xf32> @@ -1499,7 +1499,7 @@ } // CHECK-LABEL: func @cast_follwed_by_store() -// CHECK-DAG: %[[ZERO:.+]] = constant 0.000000e+00 : f32 +// CHECK-DAG: %[[ZERO:.+]] = arith.constant 0.000000e+00 : f32 // CHECK-DAG: %[[LHS:.+]] = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : memref<4x32x1024xf32> // CHECK-DAG: %[[RHS:.+]] = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : memref<4x1024x64xf32> // CHECK-DAG: %[[RESULT:.+]] = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : memref<4x32x64xf32> @@ -1512,9 +1512,9 @@ // ----- func @rank_reduced_subtensor_insert() { - %c0 = constant 0 : index - %c1 = constant 1 : index - %c2 = constant 2 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + %c2 = arith.constant 2 : index %dim0 = hal.interface.load.constant offset = 0 : index %dim1 = hal.interface.load.constant offset = 1 : index %dim2 = hal.interface.load.constant offset = 2 : index @@ -1546,11 +1546,11 @@ #map1 = affine_map<(d0, d1, d2) -> (d2, d1)> #map2 = affine_map<(d0, d1, d2) -> (d0, d1)> func @bufferize_transfer_op() { - %c3 = constant 3 : index - %cst = constant 0.000000e+00 : f32 - %c0 = constant 0 : index - %c2 = constant 2 : index - %c1 = constant 1 : index + %c3 = arith.constant 3 : index + %cst = arith.constant 0.000000e+00 : f32 + %c0 = arith.constant 0 : index + %c2 = arith.constant 2 : index + %c1 = arith.constant 1 : index %0 = hal.interface.binding.subspan @io::@arg0[%c0] : !flow.dispatch.tensor<readonly:2x3xf32> %1 = hal.interface.binding.subspan @io::@arg1[%c0] : !flow.dispatch.tensor<readonly:3x4xf32> %2 = hal.interface.binding.subspan @io::@arg2[%c0] : !flow.dispatch.tensor<readonly:2x4xf32> @@ -1608,11 +1608,11 @@ #map1 = affine_map<(d0, d1, d2) -> (d2, d1)> #map2 = affine_map<(d0, d1, d2) -> (d0, d1)> func @bufferize_transfer_op_inplace() { - %c3 = constant 3 : index - %cst = constant 0.000000e+00 : f32 - %c0 = constant 0 : index - %c2 = constant 2 : index - %c1 = constant 1 : index + %c3 = arith.constant 3 : index + %cst = arith.constant 0.000000e+00 : f32 + %c0 = arith.constant 0 : index + %c2 = arith.constant 2 : index + %c1 = arith.constant 1 : index %0 = hal.interface.binding.subspan @io::@arg0[%c0] : !flow.dispatch.tensor<readonly:2x3xf32> %1 = hal.interface.binding.subspan @io::@arg1[%c0] : !flow.dispatch.tensor<readonly:3x4xf32> %3 = hal.interface.binding.subspan @io::@ret0[%c0] : !flow.dispatch.tensor<readwrite:2x4xf32> @@ -1664,11 +1664,11 @@ #map = affine_map<(d0, d1) -> (d0, d1)> func @multi_result() { - %c0 = constant 0 : index - %c2 = constant 2 : index - %c4 = constant 4 : index - %c1 = constant 1 : index - %c3 = constant 3 : index + %c0 = arith.constant 0 : index + %c2 = arith.constant 2 : index + %c4 = arith.constant 4 : index + %c1 = arith.constant 1 : index + %c3 = arith.constant 3 : index %dim0 = hal.interface.load.constant offset = 0 : index %dim1 = hal.interface.load.constant offset = 1 : index %dim2 = hal.interface.load.constant offset = 2 : index @@ -1691,10 +1691,10 @@ %11 = hal.interface.workgroup.count[1] : index %12 = hal.interface.workgroup.size[0] : index %13 = hal.interface.workgroup.size[1] : index - %14 = muli %9, %13 : index - %15 = muli %11, %13 : index - %16 = muli %8, %12 : index - %17 = muli %10, %12 : index + %14 = arith.muli %9, %13 : index + %15 = arith.muli %11, %13 : index + %16 = arith.muli %8, %12 : index + %17 = arith.muli %10, %12 : index scf.for %arg0 = %14 to %4 step %15 { scf.for %arg1 = %16 to %5 step %17 { %18 = affine.min affine_map<(d0)[s0, s1] -> (s1, -d0 + s0)>(%arg0)[%4, %13] @@ -1706,8 +1706,8 @@ ins(%20, %21 : tensor<?x?xf32>, tensor<?x?xf32>) outs(%shape, %shape : tensor<?x?xf32>, tensor<?x?xf32>) { ^bb0(%arg2: f32, %arg3 : f32, %arg4 : f32, %arg5 : f32): // no predecessors - %23 = mulf %arg2, %arg3 : f32 - %24 = addf %arg2, %arg3 : f32 + %23 = arith.mulf %arg2, %arg3 : f32 + %24 = arith.addf %arg2, %arg3 : f32 linalg.yield %23, %24 : f32, f32 } -> (tensor<?x?xf32>, tensor<?x?xf32>) flow.dispatch.tensor.store %22#0, %2, offsets = [%arg0, %arg1], sizes = [%18, %19], strides = [%c1, %c1] : tensor<?x?xf32> -> !flow.dispatch.tensor<writeonly:?x?xf32> @@ -1741,11 +1741,11 @@ #map1 = affine_map<()[s0] -> (s0 * 16)> module { func @padded_matmul() { - %c0 = constant 0 : index - %c12544 = constant 12544 : index - %c64 = constant 64 : index - %c16 = constant 16 : index - %cst = constant 0.000000e+00 : f32 + %c0 = arith.constant 0 : index + %c12544 = arith.constant 12544 : index + %c64 = arith.constant 64 : index + %c16 = arith.constant 16 : index + %cst = arith.constant 0.000000e+00 : f32 %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : !flow.dispatch.tensor<readonly:12544x27xf32> %1 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : !flow.dispatch.tensor<readonly:27x16xf32> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : !flow.dispatch.tensor<writeonly:12544x16xf32> @@ -1783,7 +1783,7 @@ // CHECK-LABEL: func @padded_matmul() // CHECK-DAG: %[[LHS_PADDED:.+]] = memref.alloc() : memref<64x32xf32> // CHECK-DAG: %[[RHS_PADDED:.+]] = memref.alloc() : memref<32x16xf32> -// CHECK-DAG: %[[C0:.+]] = constant 0.000000e+00 : f32 +// CHECK-DAG: %[[C0:.+]] = arith.constant 0.000000e+00 : f32 // CHECK-DAG: %[[LHS:.+]] = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : memref<12544x27xf32> // CHECK-DAG: %[[RHS:.+]] = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : memref<27x16xf32> // CHECK-DAG: %[[DST:.+]] = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : memref<12544x16xf32> @@ -1802,10 +1802,10 @@ // ----- func @dot_general_padded() { - %c0 = constant 0 : index - %c3 = constant 3 : index - %c1 = constant 1 : index - %cst = constant 0.000000e+00 : f32 + %c0 = arith.constant 0 : index + %c3 = arith.constant 3 : index + %c1 = arith.constant 1 : index + %cst = arith.constant 0.000000e+00 : f32 %m = hal.interface.load.constant offset = 0 : index %n = hal.interface.load.constant offset = 1 : index %k = hal.interface.load.constant offset = 2 : index @@ -1882,12 +1882,12 @@ // ----- func @im2col() { - %c0 = constant 0 : index - %cst = constant 0.000000e+00 : f32 - %c112 = constant 112 : index - %c32 = constant 32 : index - %c16 = constant 16 : index - %c4 = constant 4 : index + %c0 = arith.constant 0 : index + %cst = arith.constant 0.000000e+00 : f32 + %c112 = arith.constant 112 : index + %c32 = arith.constant 32 : index + %c16 = arith.constant 16 : index + %c4 = arith.constant 4 : index %0 = hal.interface.binding.subspan @io::@arg0[%c0] : !flow.dispatch.tensor<readonly:1x225x225x8xf32> %1 = hal.interface.binding.subspan @io::@arg1[%c0] : !flow.dispatch.tensor<readonly:3x3x8x32xf32> %2 = hal.interface.binding.subspan @io::@ret0[%c0] : !flow.dispatch.tensor<writeonly:1x112x112x32xf32> @@ -1965,10 +1965,10 @@ // ----- func @multi_result_reduce() { - %c0 = constant 0 : index - %c0_i32 = constant 0 : i32 - %c-2147483648_i32 = constant -2147483648 : i32 - %c2 = constant 2 : index + %c0 = arith.constant 0 : index + %c0_i32 = arith.constant 0 : i32 + %c-2147483648_i32 = arith.constant -2147483648 : i32 + %c2 = arith.constant 2 : index %d0 = hal.interface.load.constant offset = 0 : index %d1 = hal.interface.load.constant offset = 1 : index %d2 = hal.interface.load.constant offset = 2 : index @@ -1989,10 +1989,10 @@ %17 = linalg.fill(%c0_i32, %13) {__internal_linalg_transform__ = "workgroup", lowering.config = {tileSizes = [[128]]}} : i32, tensor<?xi32> -> tensor<?xi32> %18:2 = linalg.generic {indexing_maps = [affine_map<(d0, d1) -> (d1, d0)>, affine_map<(d0, d1) -> (d1, d0)>, affine_map<(d0, d1) -> (d0)>, affine_map<(d0, d1) -> (d0)>], iterator_types = ["parallel", "reduction"]} ins(%7, %9 : tensor<?x?xi32>, tensor<?x?xi32>) outs(%14, %17 : tensor<?xi32>, tensor<?xi32>) attrs = {__internal_linalg_transform__ = "workgroup", lowering.config = {tileSizes = [[128]]}} { ^bb0(%arg1: i32, %arg2: i32, %arg3: i32, %arg4: i32): // no predecessors - %19 = cmpi sge, %arg1, %arg3 : i32 + %19 = arith.cmpi sge, %arg1, %arg3 : i32 %20 = select %19, %arg1, %arg3 : i32 - %21 = cmpi eq, %arg1, %arg3 : i32 - %22 = cmpi slt, %arg2, %arg4 : i32 + %21 = arith.cmpi eq, %arg1, %arg3 : i32 + %22 = arith.cmpi slt, %arg2, %arg4 : i32 %23 = select %22, %arg2, %arg4 : i32 %24 = select %19, %arg2, %arg4 : i32 %25 = select %21, %23, %24 : i32 @@ -2038,14 +2038,14 @@ #map6 = affine_map<(d0, d1) -> (32, d0 - d1)> module { func @l1_tiled_matmul_no_fill() { - %cst = constant 0.000000e+00 : f32 - %c32 = constant 32 : index - %c24 = constant 24 : index - %c144 = constant 144 : index - %c0 = constant 0 : index - %c1 = constant 1 : index - %c250 = constant 250 : index - %c370 = constant 370 : index + %cst = arith.constant 0.000000e+00 : f32 + %c32 = arith.constant 32 : index + %c24 = arith.constant 24 : index + %c144 = arith.constant 144 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + %c250 = arith.constant 250 : index + %c370 = arith.constant 370 : index %0 = hal.interface.binding.subspan @io::@ro1[%c0] : !flow.dispatch.tensor<readonly:250x144xf32> %1 = hal.interface.binding.subspan @io::@ro2[%c0] : !flow.dispatch.tensor<readonly:144x370xf32> %init = hal.interface.binding.subspan @io::@ro3[%c0] : !flow.dispatch.tensor<readonly:250x370xf32> @@ -2100,11 +2100,11 @@ } // CHECK-LABEL: l1_tiled_matmul_no_fill -// CHECK-DAG: %[[M:.+]] = constant 250 : index -// CHECK-DAG: %[[N:.+]] = constant 370 : index -// CHECK-DAG: %[[K:.+]] = constant 144 : index -// CHECK-DAG: %[[L1_MN_SIZE:.+]] = constant 32 : index -// CHECK-DAG: %[[L1_K_SIZE:.+]] = constant 24 : index +// CHECK-DAG: %[[M:.+]] = arith.constant 250 : index +// CHECK-DAG: %[[N:.+]] = arith.constant 370 : index +// CHECK-DAG: %[[K:.+]] = arith.constant 144 : index +// CHECK-DAG: %[[L1_MN_SIZE:.+]] = arith.constant 32 : index +// CHECK-DAG: %[[L1_K_SIZE:.+]] = arith.constant 24 : index // CHECK-DAG: %[[LHS:.+]] = hal.interface.binding.subspan @io::@ro1[%{{.*}}] : memref<250x144xf32> // CHECK-DAG: %[[RHS:.+]] = hal.interface.binding.subspan @io::@ro2[%{{.*}}] : memref<144x370xf32> // CHECK-DAG: %[[INIT:.+]] = hal.interface.binding.subspan @io::@ro3[%{{.*}}] : memref<250x370xf32> @@ -2144,14 +2144,14 @@ #map6 = affine_map<(d0, d1) -> (32, d0 - d1)> module { func @l1_tiled_matmul_no_fill_readwrite() { - %cst = constant 0.000000e+00 : f32 - %c32 = constant 32 : index - %c24 = constant 24 : index - %c144 = constant 144 : index - %c0 = constant 0 : index - %c1 = constant 1 : index - %c250 = constant 250 : index - %c370 = constant 370 : index + %cst = arith.constant 0.000000e+00 : f32 + %c32 = arith.constant 32 : index + %c24 = arith.constant 24 : index + %c144 = arith.constant 144 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + %c250 = arith.constant 250 : index + %c370 = arith.constant 370 : index %0 = hal.interface.binding.subspan @io::@ro1[%c0] : !flow.dispatch.tensor<readonly:250x144xf32> %1 = hal.interface.binding.subspan @io::@ro2[%c0] : !flow.dispatch.tensor<readonly:144x370xf32> %2 = hal.interface.binding.subspan @io::@wo[%c0] : !flow.dispatch.tensor<readwrite:250x370xf32> @@ -2204,11 +2204,11 @@ } // CHECK-LABEL: l1_tiled_matmul_no_fill_readwrite -// CHECK-DAG: %[[M:.+]] = constant 250 : index -// CHECK-DAG: %[[N:.+]] = constant 370 : index -// CHECK-DAG: %[[K:.+]] = constant 144 : index -// CHECK-DAG: %[[L1_MN_SIZE:.+]] = constant 32 : index -// CHECK-DAG: %[[L1_K_SIZE:.+]] = constant 24 : index +// CHECK-DAG: %[[M:.+]] = arith.constant 250 : index +// CHECK-DAG: %[[N:.+]] = arith.constant 370 : index +// CHECK-DAG: %[[K:.+]] = arith.constant 144 : index +// CHECK-DAG: %[[L1_MN_SIZE:.+]] = arith.constant 32 : index +// CHECK-DAG: %[[L1_K_SIZE:.+]] = arith.constant 24 : index // CHECK-DAG: %[[LHS:.+]] = hal.interface.binding.subspan @io::@ro1[%{{.*}}] : memref<250x144xf32> // CHECK-DAG: %[[RHS:.+]] = hal.interface.binding.subspan @io::@ro2[%{{.*}}] : memref<144x370xf32> // CHECK-DAG: %[[DST:.+]] = hal.interface.binding.subspan @io::@wo[%{{.*}}] : memref<250x370xf32> @@ -2244,14 +2244,14 @@ #map6 = affine_map<(d0, d1) -> (32, d0 - d1)> module { func @l1_tiled_matmul() { - %cst = constant 0.000000e+00 : f32 - %c32 = constant 32 : index - %c24 = constant 24 : index - %c144 = constant 144 : index - %c0 = constant 0 : index - %c1 = constant 1 : index - %c250 = constant 250 : index - %c370 = constant 370 : index + %cst = arith.constant 0.000000e+00 : f32 + %c32 = arith.constant 32 : index + %c24 = arith.constant 24 : index + %c144 = arith.constant 144 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + %c250 = arith.constant 250 : index + %c370 = arith.constant 370 : index %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : !flow.dispatch.tensor<readonly:250x144xf32> %1 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : !flow.dispatch.tensor<readonly:144x370xf32> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : !flow.dispatch.tensor<writeonly:250x370xf32> @@ -2305,11 +2305,11 @@ } // CHECK-LABEL: l1_tiled_matmul -// CHECK-DAG: %[[M:.+]] = constant 250 : index -// CHECK-DAG: %[[N:.+]] = constant 370 : index -// CHECK-DAG: %[[K:.+]] = constant 144 : index -// CHECK-DAG: %[[L1_MN_SIZE:.+]] = constant 32 : index -// CHECK-DAG: %[[L1_K_SIZE:.+]] = constant 24 : index +// CHECK-DAG: %[[M:.+]] = arith.constant 250 : index +// CHECK-DAG: %[[N:.+]] = arith.constant 370 : index +// CHECK-DAG: %[[K:.+]] = arith.constant 144 : index +// CHECK-DAG: %[[L1_MN_SIZE:.+]] = arith.constant 32 : index +// CHECK-DAG: %[[L1_K_SIZE:.+]] = arith.constant 24 : index // CHECK-DAG: %[[LHS:.+]] = hal.interface.binding.subspan @io::@s0b0_ro_external[%{{.*}}] : memref<250x144xf32> // CHECK-DAG: %[[RHS:.+]] = hal.interface.binding.subspan @io::@s0b1_ro_external[%{{.*}}] : memref<144x370xf32> // CHECK-DAG: %[[DST:.+]] = hal.interface.binding.subspan @io::@s0b2_xw_external[%{{.*}}] : memref<250x370xf32> @@ -2335,19 +2335,19 @@ // ----- func @sort1D() { - %c4 = constant 4 : index - %c3 = constant 3 : index - %c0 = constant 0 : index - %c1 = constant 1 : index + %c4 = arith.constant 4 : index + %c3 = arith.constant 3 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index %0 = hal.interface.binding.subspan @io::@ro[%c0] : !flow.dispatch.tensor<readonly:4xi32> %1 = hal.interface.binding.subspan @io::@wo[%c0] : !flow.dispatch.tensor<writeonly:4xi32> %2 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:4xi32> -> tensor<4xi32> %3 = scf.for %arg0 = %c0 to %c4 step %c1 iter_args(%arg1 = %2) -> (tensor<4xi32>) { %4 = scf.for %arg2 = %c0 to %c3 step %c1 iter_args(%arg3 = %arg1) -> (tensor<4xi32>) { - %5 = addi %arg2, %c1 : index + %5 = arith.addi %arg2, %c1 : index %6 = tensor.extract %arg3[%arg2] : tensor<4xi32> %7 = tensor.extract %arg3[%5] : tensor<4xi32> - %8 = cmpi sgt, %6, %7 : i32 + %8 = arith.cmpi sgt, %6, %7 : i32 %11 = scf.if %8 -> (tensor<4xi32>) { %12 = tensor.insert %6 into %arg3[%5] : tensor<4xi32> %13 = tensor.insert %7 into %12[%arg2] : tensor<4xi32> @@ -2372,7 +2372,7 @@ // CHECK: linalg.copy(%[[INPUT]], %[[OUTPUT]]) // CHECK: scf.for %[[ARG0:.+]] = // CHECK: scf.for %[[ARG1:.+]] = -// CHECK-DAG: %[[P1:.+]] = addi %[[ARG1]] +// CHECK-DAG: %[[P1:.+]] = arith.addi %[[ARG1]] // CHECK-DAG: %[[V1:.+]] = memref.load %[[OUTPUT]][%[[ARG1]]] // CHECK-DAG: %[[V2:.+]] = memref.load %[[OUTPUT]][%[[P1]]] // CHECK: scf.if @@ -2383,18 +2383,18 @@ // ----- func @sort1D_inplace() { - %c4 = constant 4 : index - %c3 = constant 3 : index - %c0 = constant 0 : index - %c1 = constant 1 : index + %c4 = arith.constant 4 : index + %c3 = arith.constant 3 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index %0 = hal.interface.binding.subspan @io::@rw[%c0] : !flow.dispatch.tensor<readwrite:4xi32> %2 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readwrite:4xi32> -> tensor<4xi32> %3 = scf.for %arg0 = %c0 to %c4 step %c1 iter_args(%arg1 = %2) -> (tensor<4xi32>) { %4 = scf.for %arg2 = %c0 to %c3 step %c1 iter_args(%arg3 = %arg1) -> (tensor<4xi32>) { - %5 = addi %arg2, %c1 : index + %5 = arith.addi %arg2, %c1 : index %6 = tensor.extract %arg3[%arg2] : tensor<4xi32> %7 = tensor.extract %arg3[%5] : tensor<4xi32> - %8 = cmpi sgt, %6, %7 : i32 + %8 = arith.cmpi sgt, %6, %7 : i32 %11 = scf.if %8 -> (tensor<4xi32>) { %12 = tensor.insert %6 into %arg3[%5] : tensor<4xi32> %13 = tensor.insert %7 into %12[%arg2] : tensor<4xi32> @@ -2416,7 +2416,7 @@ // CHECK-DAG: %[[INOUT:.+]] = hal.interface.binding.subspan @io::@rw // CHECK: scf.for %[[ARG0:.+]] = // CHECK: scf.for %[[ARG1:.+]] = -// CHECK-DAG: %[[P1:.+]] = addi %[[ARG1]] +// CHECK-DAG: %[[P1:.+]] = arith.addi %[[ARG1]] // CHECK-DAG: %[[V1:.+]] = memref.load %[[INOUT]][%[[ARG1]]] // CHECK-DAG: %[[V2:.+]] = memref.load %[[INOUT]][%[[P1]]] // CHECK: scf.if @@ -2426,12 +2426,12 @@ // ----- func @linalg_ext_sort_1d() { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %0 = hal.interface.binding.subspan @io::@rw[%c0] : !flow.dispatch.tensor<readwrite:128xi32> %1 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readwrite:128xi32> -> tensor<128xi32> %2 = linalg_ext.sort dimension(0) outs(%1 : tensor<128xi32>) { ^bb0(%arg0: i32, %arg1: i32): // no predecessors - %3 = cmpi sgt, %arg0, %arg1 : i32 + %3 = arith.cmpi sgt, %arg0, %arg1 : i32 linalg_ext.yield %3 : i1 } -> tensor<128xi32> flow.dispatch.tensor.store %2, %0, offsets = [], sizes = [], strides = [] : tensor<128xi32> -> !flow.dispatch.tensor<readwrite:128xi32> @@ -2446,7 +2446,7 @@ // ----- builtin.func @tensor_insert_slice() { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %1 = hal.interface.load.constant offset = 0 : index %2 = hal.interface.load.constant offset = 1 : index %d0 = hal.interface.load.constant offset = 2 : index @@ -2499,9 +2499,9 @@ // ----- builtin.func @dynamic_update_slice() { - %c0 = constant 0 : index - %c3 = constant 3 : index - %c0_i32 = constant 0 : i32 + %c0 = arith.constant 0 : index + %c3 = arith.constant 3 : index + %c0_i32 = arith.constant 0 : i32 %d0 = hal.interface.load.constant offset = 0 : index %d1 = hal.interface.load.constant offset = 1 : index %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : !flow.dispatch.tensor<readonly:?xi32>{%d0} @@ -2509,11 +2509,11 @@ %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : !flow.dispatch.tensor<writeonly:?x?xi32>{%d1, %d0} %3 = flow.dispatch.tensor.load %1, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:i32> -> tensor<i32> %4 = tensor.extract %3[] : tensor<i32> - %5 = cmpi slt, %4, %c0_i32 : i32 + %5 = arith.cmpi slt, %4, %c0_i32 : i32 %6 = select %5, %4, %c0_i32 : i32 - %7 = cmpi sgt, %6, %c0_i32 : i32 + %7 = arith.cmpi sgt, %6, %c0_i32 : i32 %8 = select %7, %6, %c0_i32 : i32 - %9 = index_cast %8 : i32 to index + %9 = arith.index_cast %8 : i32 to index %workgroup_id_x = hal.interface.workgroup.id[0] : index %workgroup_count_x = hal.interface.workgroup.count[0] : index %10 = affine.apply affine_map<()[s0] -> (s0 * 64)>()[%workgroup_id_x] @@ -2546,11 +2546,11 @@ // ----- func @multi_level_tile_fuse() { - %c4 = constant 4 : index - %c0 = constant 0 : index - %c1 = constant 1 : index - %c2 = constant 2 : index - %cst = constant 0.000000e+00 : f32 + %c4 = arith.constant 4 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + %c2 = arith.constant 2 : index + %cst = arith.constant 0.000000e+00 : f32 %m = hal.interface.load.constant offset = 0 : index %n = hal.interface.load.constant offset = 1 : index %k = hal.interface.load.constant offset = 2 : index @@ -2605,7 +2605,7 @@ } %19 = linalg.generic {indexing_maps = [affine_map<(d0, d1) -> ()>, affine_map<(d0, d1) -> (d0, d1)>, affine_map<(d0, d1) -> (d0, d1)>], iterator_types = ["parallel", "parallel"]} ins(%4, %18 : tensor<f32>, tensor<?x?xf32>) outs(%11 : tensor<?x?xf32>) attrs = {__internal_linalg_transform__ = "workgroup"} { ^bb0(%arg2: f32, %arg3: f32, %arg4: f32): // no predecessors - %20 = addf %arg2, %arg3 : f32 + %20 = arith.addf %arg2, %arg3 : f32 linalg.yield %20 : f32 } -> tensor<?x?xf32> flow.dispatch.tensor.store %19, %3, offsets = [%arg0, %arg1], sizes = [%9, %10], strides = [1, 1] : tensor<?x?xf32> -> !flow.dispatch.tensor<writeonly:?x?xf32> @@ -2620,7 +2620,7 @@ hal.interface.binding @arg3, set=0, binding=3, type="StorageBuffer", access="Write|Discard" } // CHECK-LABEL: func @multi_level_tile_fuse() -// CHECK-DAG: %[[C0:.+]] = constant 0 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index // CHECK-DAG: %[[M:.+]] = hal.interface.load.constant offset = 0 // CHECK-DAG: %[[N:.+]] = hal.interface.load.constant offset = 1 // CHECK-DAG: %[[K:.+]] = hal.interface.load.constant offset = 2 @@ -2652,11 +2652,11 @@ // ----- func @operand_fusion() { - %c4 = constant 4 : index - %c0 = constant 0 : index - %c1 = constant 1 : index - %c2 = constant 2 : index - %cst = constant 0.000000e+00 : f32 + %c4 = arith.constant 4 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + %c2 = arith.constant 2 : index + %cst = arith.constant 0.000000e+00 : f32 %m = hal.interface.load.constant offset = 0 : index %n = hal.interface.load.constant offset = 1 : index %k = hal.interface.load.constant offset = 2 : index @@ -2689,7 +2689,7 @@ outs(%17: tensor<?x?xf32>) -> tensor<?x?xf32> %19 = linalg.generic {indexing_maps = [affine_map<(d0, d1) -> ()>, affine_map<(d0, d1) -> (d0, d1)>, affine_map<(d0, d1) -> (d0, d1)>], iterator_types = ["parallel", "parallel"]} ins(%4, %18 : tensor<f32>, tensor<?x?xf32>) outs(%11 : tensor<?x?xf32>) attrs = {__internal_linalg_transform__ = "workgroup"} { ^bb0(%arg2: f32, %arg3: f32, %arg4: f32): // no predecessors - %20 = addf %arg2, %arg3 : f32 + %20 = arith.addf %arg2, %arg3 : f32 linalg.yield %20 : f32 } -> tensor<?x?xf32> flow.dispatch.tensor.store %19, %3, offsets = [%arg0, %arg1], sizes = [%9, %10], strides = [1, 1] : tensor<?x?xf32> -> !flow.dispatch.tensor<writeonly:?x?xf32> @@ -2704,7 +2704,7 @@ hal.interface.binding @arg3, set=0, binding=3, type="StorageBuffer", access="Write|Discard" } // CHECK-LABEL: func @operand_fusion() -// CHECK-DAG: %[[C0:.+]] = constant 0 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index // CHECK-DAG: %[[M:.+]] = hal.interface.load.constant offset = 0 // CHECK-DAG: %[[N:.+]] = hal.interface.load.constant offset = 1 // CHECK-DAG: %[[K:.+]] = hal.interface.load.constant offset = 2
diff --git a/iree/compiler/Codegen/Common/test/vectorize_linalg_conv.mlir b/iree/compiler/Codegen/Common/test/vectorize_linalg_conv.mlir index d86a553..8b661d9 100644 --- a/iree/compiler/Codegen/Common/test/vectorize_linalg_conv.mlir +++ b/iree/compiler/Codegen/Common/test/vectorize_linalg_conv.mlir
@@ -20,7 +20,7 @@ // CHECK-SAME: %[[INPUT_ARG:.+]]: memref<2x2x2x3xf32>, // CHECK-SAME: %[[OUTPUT_ARG:.+]]: memref<2x2x2x4xf32> -// CHECK: %[[FLOAT_ZERO:.+]] = constant 0.000000e+00 : f32 +// CHECK: %[[FLOAT_ZERO:.+]] = arith.constant 0.000000e+00 : f32 // CHECK-DAG: %[[FILTER_SUBVIEW:.+]] = memref.subview %[[FILTER_ARG]]{{.*}} to memref<1x1x3x4xf32> // CHECK-DAG: %[[INPUT_SUBVIEW:.+]] = memref.subview %[[INPUT_ARG]]{{.*}} to memref<1x2x2x3xf32> @@ -152,7 +152,7 @@ // CHECK-SAME: %[[FILTER_ARG:.+]]: memref<2x1x8xf32>, // CHECK-SAME: %[[OUTPUT_ARG:.+]]: memref<2x2x2x8xf32> -// CHECK: %[[FLOAT_ZERO:.+]] = constant 0.000000e+00 : f32 +// CHECK: %[[FLOAT_ZERO:.+]] = arith.constant 0.000000e+00 : f32 // CHECK-DAG: %[[INPUT_SUBVIEW:.+]] = memref.subview %[[INPUT_ARG]]{{.*}} to memref<1x3x3x8xf32> // CHECK-DAG: %[[FILTER_SUBVIEW:.+]] = memref.subview %[[FILTER_ARG]]{{.*}} to memref<1x1x8xf32>
diff --git a/iree/compiler/Codegen/LLVMCPU/BUILD b/iree/compiler/Codegen/LLVMCPU/BUILD index 8c439dd..0d41055 100644 --- a/iree/compiler/Codegen/LLVMCPU/BUILD +++ b/iree/compiler/Codegen/LLVMCPU/BUILD
@@ -42,6 +42,7 @@ "@llvm-project//llvm:Support", "@llvm-project//mlir:AffineToStandardTransforms", "@llvm-project//mlir:Analysis", + "@llvm-project//mlir:ArithmeticToLLVM", "@llvm-project//mlir:CFGTransforms", "@llvm-project//mlir:DialectUtils", "@llvm-project//mlir:IR",
diff --git a/iree/compiler/Codegen/LLVMCPU/CMakeLists.txt b/iree/compiler/Codegen/LLVMCPU/CMakeLists.txt index 726abba..50207fd 100644 --- a/iree/compiler/Codegen/LLVMCPU/CMakeLists.txt +++ b/iree/compiler/Codegen/LLVMCPU/CMakeLists.txt
@@ -29,6 +29,7 @@ LLVMSupport MLIRAffineToStandard MLIRAnalysis + MLIRArithmeticToLLVM MLIRIR MLIRLLVMCommonConversion MLIRLLVMIR
diff --git a/iree/compiler/Codegen/LLVMCPU/ConvertToLLVM.cpp b/iree/compiler/Codegen/LLVMCPU/ConvertToLLVM.cpp index a31cd40..e0dae14 100644 --- a/iree/compiler/Codegen/LLVMCPU/ConvertToLLVM.cpp +++ b/iree/compiler/Codegen/LLVMCPU/ConvertToLLVM.cpp
@@ -16,6 +16,7 @@ #include "llvm/Support/raw_ostream.h" #include "mlir/Analysis/DataLayoutAnalysis.h" #include "mlir/Conversion/AffineToStandard/AffineToStandard.h" +#include "mlir/Conversion/ArithmeticToLLVM/ArithmeticToLLVM.h" #include "mlir/Conversion/LLVMCommon/ConversionTarget.h" #include "mlir/Conversion/LLVMCommon/LoweringOptions.h" #include "mlir/Conversion/LLVMCommon/Pattern.h" @@ -405,8 +406,8 @@ for (auto returnOp : llvm::make_early_inc_range(llvmFuncOp.getOps<mlir::ReturnOp>())) { rewriter.setInsertionPoint(returnOp); - auto returnValue = - rewriter.createOrFold<mlir::ConstantIntOp>(returnOp.getLoc(), 0, 32); + auto returnValue = rewriter.createOrFold<mlir::arith::ConstantIntOp>( + returnOp.getLoc(), 0, 32); rewriter.replaceOpWithNewOp<mlir::ReturnOp>(returnOp, returnValue); } @@ -676,6 +677,7 @@ populateMathToLLVMConversionPatterns(converter, patterns); populateMemRefToLLVMConversionPatterns(converter, patterns); populateStdToLLVMConversionPatterns(converter, patterns); + arith::populateArithmeticToLLVMConversionPatterns(converter, patterns); populateVectorToSCFConversionPatterns(patterns); populateVectorToLLVMMatrixConversionPatterns(converter, patterns); populateVectorToLLVMConversionPatterns(converter, patterns); @@ -700,6 +702,7 @@ target.addLegalOp<ModuleOp, IREE::HAL::InterfaceOp, IREE::HAL::InterfaceBindingOp, IREE::HAL::InterfaceEndOp>(); target.addIllegalDialect<ShapeDialect, StandardOpsDialect, + mlir::arith::ArithmeticDialect, IREE::Util::UtilDialect, IREE::HAL::HALDialect, math::MathDialect, tosa::TosaDialect>(); target.addIllegalOp<UnrealizedConversionCastOp>(); @@ -709,15 +712,15 @@ if (isEntryPoint(funcOp)) return false; return true; }); - target.addDynamicallyLegalDialect<ShapeDialect, StandardOpsDialect, - IREE::Util::UtilDialect, - IREE::HAL::HALDialect, math::MathDialect>( - [&](Operation *op) { - auto funcParent = op->getParentOfType<FuncOp>(); - if (!funcParent) return false; - if (isEntryPoint(funcParent)) return false; - return true; - }); + target.addDynamicallyLegalDialect< + ShapeDialect, StandardOpsDialect, mlir::math::MathDialect, + mlir::arith::ArithmeticDialect, IREE::Util::UtilDialect, + IREE::HAL::HALDialect, math::MathDialect>([&](Operation *op) { + auto funcParent = op->getParentOfType<FuncOp>(); + if (!funcParent) return false; + if (isEntryPoint(funcParent)) return false; + return true; + }); if (failed(applyPartialConversion(module, target, std::move(patterns)))) { signalPassFailure();
diff --git a/iree/compiler/Codegen/LLVMCPU/VectorContractToAArch64InlineAsmOp.cpp b/iree/compiler/Codegen/LLVMCPU/VectorContractToAArch64InlineAsmOp.cpp index 9dd3d36..a6e675e 100644 --- a/iree/compiler/Codegen/LLVMCPU/VectorContractToAArch64InlineAsmOp.cpp +++ b/iree/compiler/Codegen/LLVMCPU/VectorContractToAArch64InlineAsmOp.cpp
@@ -52,12 +52,12 @@ } auto getI8Value = [&](Value v) -> Value { - if (auto parentOp = v.getDefiningOp<SignExtendIOp>()) { - if (parentOp.value().getType().cast<VectorType>().getElementType() != + if (auto parentOp = v.getDefiningOp<arith::ExtSIOp>()) { + if (parentOp.in().getType().cast<VectorType>().getElementType() != I8Type) { return nullptr; } else { - return parentOp.value(); + return parentOp.in(); } } return nullptr; @@ -123,7 +123,7 @@ auto int32x4x4xVType = VectorType::get({4, 4}, I32Type); Value result; - result = rewriter.create<ConstantOp>( + result = rewriter.create<arith::ConstantOp>( loc, int32x4x4xVType, DenseIntElementsAttr::get(int32x4x4xVType, 0)); for (int i = 0; i < 4; ++i) { result = rewriter.create<vector::InsertOp>(loc, resVec[i], result,
diff --git a/iree/compiler/Codegen/LLVMCPU/test/hal_interface_bindings.mlir b/iree/compiler/Codegen/LLVMCPU/test/hal_interface_bindings.mlir index 11adb9a..210fd6d 100644 --- a/iree/compiler/Codegen/LLVMCPU/test/hal_interface_bindings.mlir +++ b/iree/compiler/Codegen/LLVMCPU/test/hal_interface_bindings.mlir
@@ -5,7 +5,7 @@ // CHECK-LABEL: llvm.func internal @binding_ptrs func @binding_ptrs() { // CHECK-DAG: %[[C72:.+]] = llvm.mlir.constant(72 : index) : i64 - %c72 = constant 72 : index + %c72 = arith.constant 72 : index // CHECK: %[[STATE:.+]] = llvm.load %arg0 : !llvm.ptr<struct<"iree_hal_executable_dispatch_state_v0_t", (array<3 x i32>, array<3 x i32>, i64, ptr<i32>, i64, ptr<ptr<i8>>, ptr<i64>)>> // CHECK: %[[PC:.+]] = llvm.extractvalue %[[STATE]][3] : !llvm.struct<"iree_hal_executable_dispatch_state_v0_t", (array<3 x i32>, array<3 x i32>, i64, ptr<i32>, i64, ptr<ptr<i8>>, ptr<i64>)> @@ -39,7 +39,7 @@ %memref = hal.interface.binding.subspan @io::@ret0[%c72] : memref<?x2xf32>{%dim} // CHECK: %[[VAL:.+]] = llvm.load - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %val = memref.load %memref[%c0, %c0] : memref<?x2xf32> // CHECK: llvm.call @sink(%[[VAL]])
diff --git a/iree/compiler/Codegen/LLVMCPU/test/hal_interface_constants.mlir b/iree/compiler/Codegen/LLVMCPU/test/hal_interface_constants.mlir index 0902ffc..3bc2811 100644 --- a/iree/compiler/Codegen/LLVMCPU/test/hal_interface_constants.mlir +++ b/iree/compiler/Codegen/LLVMCPU/test/hal_interface_constants.mlir
@@ -12,7 +12,7 @@ // CHECK: %[[V64:.+]] = llvm.zext %[[V32]] : i32 to i64 %v1 = hal.interface.load.constant offset = 1 : index // CHECK-NOT: unrealized_conversion_cast - %v2 = index_cast %v1 : index to i64 + %v2 = arith.index_cast %v1 : index to i64 // CHECK: llvm.call @sink llvm.call @sink(%v2) : (i64) -> () return
diff --git a/iree/compiler/Codegen/LLVMCPU/test/hal_interface_workgroup_info.mlir b/iree/compiler/Codegen/LLVMCPU/test/hal_interface_workgroup_info.mlir index 1d49973..2a9e959 100644 --- a/iree/compiler/Codegen/LLVMCPU/test/hal_interface_workgroup_info.mlir +++ b/iree/compiler/Codegen/LLVMCPU/test/hal_interface_workgroup_info.mlir
@@ -9,7 +9,7 @@ // CHECK: %[[Z64:.+]] = llvm.zext %[[Z32]] : i32 to i64 %workgroup_id_z = hal.interface.workgroup.id[2] : index // CHECK-NEXT: llvm.call @sink(%[[Z64]]) - %val = index_cast %workgroup_id_z : index to i64 + %val = arith.index_cast %workgroup_id_z : index to i64 llvm.call @sink(%val) : (i64) -> () return } @@ -26,7 +26,7 @@ // CHECK: %[[Z64:.+]] = llvm.zext %[[Z32]] : i32 to i64 %workgroup_size_z = hal.interface.workgroup.size[2] : index // CHECK-NEXT: llvm.call @sink(%[[Z64]]) - %val = index_cast %workgroup_size_z : index to i64 + %val = arith.index_cast %workgroup_size_z : index to i64 llvm.call @sink(%val) : (i64) -> () return } @@ -43,7 +43,7 @@ // CHECK: %[[Z64:.+]] = llvm.zext %[[Z32]] : i32 to i64 %workgroup_count_z = hal.interface.workgroup.count[2] : index // CHECK-NEXT: llvm.call @sink(%[[Z64]]) - %val = index_cast %workgroup_count_z : index to i64 + %val = arith.index_cast %workgroup_count_z : index to i64 llvm.call @sink(%val) : (i64) -> () return }
diff --git a/iree/compiler/Codegen/LLVMCPU/test/materialize_launch_configuration.mlir b/iree/compiler/Codegen/LLVMCPU/test/materialize_launch_configuration.mlir index 0a07d23..91728ce 100644 --- a/iree/compiler/Codegen/LLVMCPU/test/materialize_launch_configuration.mlir +++ b/iree/compiler/Codegen/LLVMCPU/test/materialize_launch_configuration.mlir
@@ -17,8 +17,8 @@ } builtin.module { func @matmul_tensors() { - %c0 = constant 0 : index - %c1 = constant 1 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index %pcM = hal.interface.load.constant offset = 0 : index %pcN = hal.interface.load.constant offset = 1 : index %pcK = hal.interface.load.constant offset = 2 : index @@ -66,7 +66,7 @@ // CHECK-NEXT: (%[[ARG0:[a-zA-Z0-9_]+]]: index // CHECK-SAME: %[[ARG1:[a-zA-Z0-9_]+]]: index // CHECK-SAME: %[[ARG2:[a-zA-Z0-9_]+]]: index) -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK-DAG: %[[D0:.+]] = affine.apply #[[MAP0]]()[%[[ARG0]]] // CHECK-DAG: %[[D1:.+]] = affine.apply #[[MAP0]]()[%[[ARG1]]] // CHECK: hal.return %[[D0]], %[[D1]], %[[C1]] : index, index, index @@ -92,7 +92,7 @@ } builtin.module { func @add_no_config() { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %dim0 = hal.interface.load.constant offset = 0 : index %dim1 = hal.interface.load.constant offset = 1 : index %0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<?x?xf32>{%dim0, %dim1} @@ -105,7 +105,7 @@ iterator_types = ["parallel", "parallel"]} ins(%0, %1 : memref<?x?xf32>, memref<?xf32>) outs(%2 : memref<?x?xf32>) { ^bb0(%arg0: f32, %arg1: f32, %arg2: f32): // no predecessors - %3 = addf %arg0, %arg1 : f32 + %3 = arith.addf %arg0, %arg1 : f32 linalg.yield %3 : f32 } return @@ -143,8 +143,8 @@ !flow.dispatch.tensor<writeonly:?x?xf32>) -> ()} builtin.module { func @add() { - %c0 = constant 0 : index - %c1 = constant 1 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index %dim0 = hal.interface.load.constant offset = 0 : index %dim1 = hal.interface.load.constant offset = 1 : index %0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<?x?xf32>{%dim0, %dim1} @@ -177,7 +177,7 @@ iterator_types = ["parallel", "parallel"]} ins(%14, %15 : memref<?x?xf32, affine_map<(d0, d1)[s0, s1] -> (d0 * s1 + s0 + d1)>>, memref<?xf32, affine_map<(d0)[s0] -> (d0 + s0)>>) outs(%16 : memref<?x?xf32, affine_map<(d0, d1)[s0, s1] -> (d0 * s1 + s0 + d1)>>) { ^bb0(%arg2: f32, %arg3: f32, %arg4: f32): // no predecessors - %3 = addf %arg2, %arg3 : f32 + %3 = arith.addf %arg2, %arg3 : f32 linalg.yield %3 : f32 } } @@ -198,7 +198,7 @@ // CHECK-NEXT: (%[[ARG0:[a-zA-Z0-9_]+]]: index // CHECK-SAME: %[[ARG1:[a-zA-Z0-9_]+]]: index // CHECK-SAME: %[[ARG2:[a-zA-Z0-9_]+]]: index) -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK-DAG: %[[D0:.+]] = affine.apply #[[MAP0]]()[%[[ARG0]]] // CHECK-DAG: %[[D1:.+]] = affine.apply #[[MAP0]]()[%[[ARG1]]] // CHECK: hal.return %[[D0]], %[[D1]], %[[C1]] : index, index, index @@ -223,10 +223,10 @@ !flow.dispatch.tensor<writeonly:?x?x?x?xf32>) -> ()} builtin.module { func @add4D() { - %c0 = constant 0 : index - %c1 = constant 1 : index - %c2 = constant 2 : index - %c3 = constant 3 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + %c2 = arith.constant 2 : index + %c3 = arith.constant 3 : index %dim0 = hal.interface.load.constant offset = 0 : index %dim1 = hal.interface.load.constant offset = 1 : index %dim2 = hal.interface.load.constant offset = 2 : index @@ -279,7 +279,7 @@ memref<?x?x?x?xf32, affine_map<(d0, d1, d2, d3)[s0, s1, s2, s3] -> (d0 * s1 + s0 + d1 * s2 + d2 * s3 + d3)>>) outs(%16 : memref<?x?x?x?xf32, affine_map<(d0, d1, d2, d3)[s0, s1, s2, s3] -> (d0 * s1 + s0 + d1 * s2 + d2 * s3 + d3)>>) { ^bb0(%arg2: f32, %arg3: f32, %arg4: f32): // no predecessors - %3 = addf %arg2, %arg3 : f32 + %3 = arith.addf %arg2, %arg3 : f32 linalg.yield %3 : f32 } } @@ -326,9 +326,9 @@ } builtin.module { func @batch_matmul_tensors() { - %c0 = constant 0 : index - %c1 = constant 1 : index - %c2 = constant 2 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + %c2 = arith.constant 2 : index %pcB = hal.interface.load.constant offset = 0 : index %pcM = hal.interface.load.constant offset = 1 : index %pcN = hal.interface.load.constant offset = 2 : index @@ -384,7 +384,7 @@ // CHECK-NEXT: (%[[ARG0:[a-zA-Z0-9]+]]: index // CHECK-SAME: %[[ARG1:[a-zA-Z0-9]+]]: index // CHECK-SAME: %[[ARG2:[a-zA-Z0-9]+]]: index) -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK-DAG: %[[D0:.+]] = affine.apply #[[MAP0]]()[%[[ARG0]]] // CHECK-DAG: %[[D1:.+]] = affine.apply #[[MAP0]]()[%[[ARG1]]] // CHECK: hal.return %[[D0]], %[[D1]], %[[C1]] @@ -398,10 +398,10 @@ hal.executable.entry_point @preset_config attributes {interface = @io, ordinal = 0 : index} builtin.module { builtin.func @preset_config() { - %c0 = constant 0 : index - %c512 = constant 512 : index - %c128 = constant 128 : index - %cst = constant 0.000000e+00 : f32 + %c0 = arith.constant 0 : index + %c512 = arith.constant 512 : index + %c128 = arith.constant 128 : index + %cst = arith.constant 0.000000e+00 : f32 %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : !flow.dispatch.tensor<readonly:128x256xf32> %1 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : !flow.dispatch.tensor<readonly:256x512xf32> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : !flow.dispatch.tensor<writeonly:128x512xf32> @@ -447,7 +447,7 @@ // CHECK: hal.executable.entry_point // CHECK-SAME: translation.info = {passPipeline = "CPUVectorization", workloadPerWorkgroup = [32, 32]} // CHECK-NEXT: ^bb0(%[[ARG0:[a-zA-Z0-9]+]]: index, %[[ARG1:[a-zA-Z0-9]+]]: index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK-DAG: %[[NWG_X:.+]] = affine.apply #[[MAP0]]()[%[[ARG0]]] // CHECK-DAG: %[[NWG_Y:.+]] = affine.apply #[[MAP0]]()[%[[ARG1]]] // CHECK: return %[[NWG_X]], %[[NWG_Y]], %[[C1]] @@ -475,7 +475,7 @@ hal.executable.entry_point @tensor_insert_slice attributes {interface = @io, ordinal = 0 : index} builtin.module { builtin.func @tensor_insert_slice() { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : !flow.dispatch.tensor<readonly:?x?xi32> %1 = hal.interface.load.constant offset = 0 : index %2 = hal.interface.load.constant offset = 1 : index @@ -516,7 +516,7 @@ // CHECK-SAME: translation.info = {passPipeline = "CPUDefault", workloadPerWorkgroup = [64, 64]} // CHECK-NEXT: %[[ARG0:[a-zA-Z0-9_]+]]: index // CHECK-SAME: %[[ARG1:[a-zA-Z0-9_]+]]: index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK-DAG: %[[NWGSX:.+]] = affine.apply #[[MAP0]]()[%[[ARG0]]] // CHECK-DAG: %[[NWGSY:.+]] = affine.apply #[[MAP0]]()[%[[ARG1]]] // CHECK: hal.return %[[NWGSX]], %[[NWGSY]], %[[C1]] @@ -532,10 +532,10 @@ hal.executable.entry_point @static_1d_fft_stage2 attributes {interface = @io, ordinal = 0 : index} builtin.module { builtin.func @static_1d_fft_stage2() { - %c0 = constant 0 : index - %c2 = constant 2 : index - %cst = constant dense<[1.000000e+00, 6.12323426E-17]> : tensor<2xf32> - %cst_0 = constant dense<[-0.000000e+00, -1.000000e+00]> : tensor<2xf32> + %c0 = arith.constant 0 : index + %c2 = arith.constant 2 : index + %cst = arith.constant dense<[1.000000e+00, 6.12323426E-17]> : tensor<2xf32> + %cst_0 = arith.constant dense<[-0.000000e+00, -1.000000e+00]> : tensor<2xf32> %0 = hal.interface.binding.subspan @io::@s0b0_rw_external[%c0] : !flow.dispatch.tensor<readwrite:32xf32> %1 = hal.interface.binding.subspan @io::@s0b1_rw_external[%c0] : !flow.dispatch.tensor<readwrite:32xf32> %2 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readwrite:32xf32> -> tensor<32xf32> @@ -555,7 +555,7 @@ // CHECK-SAME: passPipeline = "CPUDefault" // CHECK-SAME: workloadPerWorkgroup = [64]} // CHECK-NEXT: ^{{.+}}(%[[ARG0:.+]]: index, %[[ARG1:.+]]: index, %[[ARG2:.+]]: index): -// CHECK-NEXT: %[[C1:.+]] = constant 1 : index +// CHECK-NEXT: %[[C1:.+]] = arith.constant 1 : index // CHECK-NEXT: %[[T0:.+]] = affine.apply #[[MAP0]]()[%[[ARG0]]] // CHECK-NEXT: hal.return %[[T0]], %[[C1]], %[[C1]] @@ -575,13 +575,13 @@ hal.executable.entry_point @static_3d_fft_stage3 attributes {interface = @io, ordinal = 0 : index} builtin.module { builtin.func @static_3d_fft_stage3() { - %c0 = constant 0 : index - %c3 = constant 3 : index - %c64 = constant 64 : index - %c128 = constant 128 : index - %c32 = constant 32 : index - %cst = constant dense<[1.000000e+00, 0.707106769, 6.12323426E-17, -0.707106769]> : tensor<4xf32> - %cst_0 = constant dense<[-0.000000e+00, -0.707106769, -1.000000e+00, -0.707106769]> : tensor<4xf32> + %c0 = arith.constant 0 : index + %c3 = arith.constant 3 : index + %c64 = arith.constant 64 : index + %c128 = arith.constant 128 : index + %c32 = arith.constant 32 : index + %cst = arith.constant dense<[1.000000e+00, 0.707106769, 6.12323426E-17, -0.707106769]> : tensor<4xf32> + %cst_0 = arith.constant dense<[-0.000000e+00, -0.707106769, -1.000000e+00, -0.707106769]> : tensor<4xf32> %0 = memref.buffer_cast %cst_0 : memref<4xf32> %1 = memref.buffer_cast %cst : memref<4xf32> %2 = hal.interface.binding.subspan @io::@s0b0_rw_external[%c0] : memref<64x128x32xf32> @@ -648,8 +648,8 @@ hal.executable.entry_point @outs_fusion_fn attributes {interface = @io, ordinal = 0 : index} builtin.module { builtin.func @outs_fusion_fn() { - %c0 = constant 0 : index - %cst = constant 0.0 : f32 + %c0 = arith.constant 0 : index + %cst = arith.constant 0.0 : f32 %0 = hal.interface.binding.subspan @io::@arg0[%c0] : !flow.dispatch.tensor<readonly:?x?xf32> %1 = hal.interface.binding.subspan @io::@arg1[%c0] : !flow.dispatch.tensor<readonly:?x?xf32> %2 = hal.interface.load.constant offset = 0 : index @@ -688,8 +688,8 @@ ins(%lhs, %rhs : tensor<?x?xf32>, tensor<?x?xf32>) outs(%fill : tensor<?x?xf32>) { ^bb0(%arg0: f32, %arg1: f32, %arg2: f32): - %6 = mulf %arg0, %arg1 : f32 - %7 = addf %6, %arg2 : f32 + %6 = arith.mulf %arg0, %arg1 : f32 + %7 = arith.addf %6, %arg2 : f32 linalg.yield %6 : f32 } -> tensor<?x?xf32> flow.dispatch.tensor.store %gemm, %5, offsets = [%iv0, %iv1], sizes = [%tile_m, %tile_n], strides = [1, 1] : tensor<?x?xf32> -> !flow.dispatch.tensor<writeonly:?x?xf32> @@ -710,7 +710,7 @@ hal.executable.entry_point public @conv attributes {interface = @io, ordinal = 0 : index} builtin.module { func @conv() { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %0 = hal.interface.load.constant offset = 0 : index %1 = hal.interface.load.constant offset = 1 : index %2 = hal.interface.load.constant offset = 2 : index @@ -786,10 +786,10 @@ hal.executable.entry_point public @conv_static attributes {interface = @io, ordinal = 0 : index} builtin.module { func @conv_static() { - %cst = constant 0.000000e+00 : f32 - %c80 = constant 80 : index - %c96 = constant 96 : index - %c0 = constant 0 : index + %cst = arith.constant 0.000000e+00 : f32 + %c80 = arith.constant 80 : index + %c96 = arith.constant 96 : index + %c0 = arith.constant 0 : index %0 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : !flow.dispatch.tensor<readonly:1x161x161x96xf32> %1 = hal.interface.binding.subspan @io::@s0b0_ro_constant[%c0] : !flow.dispatch.tensor<readonly:3x3x96xf32> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : !flow.dispatch.tensor<writeonly:1x80x80x96xf32> @@ -861,9 +861,9 @@ hal.executable.entry_point public @generic_static attributes {interface = @io, ordinal = 0 : index} builtin.module { func @generic_static() { - %c16 = constant 16 : index - %c96 = constant 96 : index - %c0 = constant 0 : index + %c16 = arith.constant 16 : index + %c96 = arith.constant 96 : index + %c0 = arith.constant 0 : index %0 = hal.interface.binding.subspan @io::@s0b0_ro_constant[%c0] : !flow.dispatch.tensor<readonly:96x16xf32> %1 = hal.interface.binding.subspan @io::@s0b1_xw_external[%c0] : !flow.dispatch.tensor<writeonly:16x96xf32> %workgroup_size_x = hal.interface.workgroup.size[0] : index @@ -905,7 +905,7 @@ // CHECK: hal.executable.entry_point public @generic_static attributes // CHECK-SAME: translation.info = {passPipeline = "CPUDefault", workloadPerWorkgroup = [32, 8]} // CHECK-NEXT: ^bb0(%[[ARG0:[a-zA-Z0-9]+]]: index, %[[ARG1:[a-zA-Z0-9]+]]: index, %[[ARG2:[a-zA-Z0-9]+]]: index) -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK-DAG: %[[D0:.+]] = affine.apply #[[MAP0]]()[%[[ARG0]] // CHECK-DAG: %[[D1:.+]] = affine.apply #[[MAP1]]()[%[[ARG1]] // CHECK: hal.return %[[D0]], %[[D1]], %[[C1]] @@ -919,12 +919,12 @@ hal.executable.entry_point public @matmul_static attributes {interface = @io, ordinal = 0 : index} builtin.module { func @matmul_static() { - %cst = constant 0.000000e+00 : f32 - %c196 = constant 196 : index - %c40 = constant 40 : index - %c0 = constant 0 : index - %c8 = constant 8 : index - %c28 = constant 28 : index + %cst = arith.constant 0.000000e+00 : f32 + %c196 = arith.constant 196 : index + %c40 = arith.constant 40 : index + %c0 = arith.constant 0 : index + %c8 = arith.constant 8 : index + %c28 = arith.constant 28 : index %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : !flow.dispatch.tensor<readonly:196x240xf32> %1 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : !flow.dispatch.tensor<readonly:240x40xf32> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : !flow.dispatch.tensor<writeonly:196x40xf32> @@ -966,7 +966,7 @@ // CHECK: hal.executable.entry_point public @matmul_static attributes // CHECK-SAME: translation.info = {passPipeline = "CPUTensorToVectors", workloadPerWorkgroup = [8, 28]} // CHECK-NEXT: ^bb0(%[[ARG0:[a-zA-Z0-9]+]]: index, %[[ARG1:[a-zA-Z0-9]+]]: index, %[[ARG2:[a-zA-Z0-9]+]]: index) -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK-DAG: %[[D0:.+]] = affine.apply #[[MAP0]]()[%[[ARG0]]] // CHECK-DAG: %[[D1:.+]] = affine.apply #[[MAP1]]()[%[[ARG1]]] // CHECK: hal.return %[[D0]], %[[D1]], %[[C1]] @@ -980,12 +980,12 @@ hal.executable.entry_point public @restrict_num_workgroups attributes {interface = @io, ordinal = 0 : index} builtin.module { func @restrict_num_workgroups() { - %cst = constant 0.000000e+00 : f32 - %c7 = constant 7 : index - %c576 = constant 576 : index - %c0 = constant 0 : index - %c64 = constant 64 : index - %c2 = constant 2 : index + %cst = arith.constant 0.000000e+00 : f32 + %c7 = arith.constant 7 : index + %c576 = arith.constant 576 : index + %c0 = arith.constant 0 : index + %c64 = arith.constant 64 : index + %c2 = arith.constant 2 : index %0 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : !flow.dispatch.tensor<readonly:1x11x11x576xf32> %1 = hal.interface.binding.subspan @io::@s0b0_ro_constant[%c0] : !flow.dispatch.tensor<readonly:5x5x576xf32> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : !flow.dispatch.tensor<writeonly:1x7x7x576xf32> @@ -1050,7 +1050,7 @@ hal.executable.entry_point public @test_exp_0 attributes {interface = @io, ordinal = 0 : index} builtin.module { func @test_exp_0() { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %size = hal.interface.workgroup.size[0] : index %count = hal.interface.workgroup.count[0] : index %id = hal.interface.workgroup.id[0] : index @@ -1078,7 +1078,7 @@ // CHECK: hal.executable.entry_point public @test_exp_0 attributes // CHECK-SAME: translation.info = {passPipeline = "CPUDefault", workloadPerWorkgroup = [64]} // CHECK-NEXT: ^bb0(%[[ARG0:[a-zA-Z0-9]+]]: index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK-DAG: %[[D0:.+]] = affine.apply #[[MAP]]()[%[[ARG0]]] // CHECK: hal.return %[[D0]], %[[C1]], %[[C1]] @@ -1089,7 +1089,7 @@ hal.executable.entry_point public @test_exp_1 attributes {interface = @io, ordinal = 0 : index} builtin.module { func @test_exp_1() { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %size = hal.interface.workgroup.size[0] : index %count = hal.interface.workgroup.count[0] : index %id = hal.interface.workgroup.id[0] : index @@ -1117,7 +1117,7 @@ // CHECK: hal.executable.entry_point public @test_exp_1 attributes // CHECK-SAME: translation.info = {passPipeline = "CPUDefault", workloadPerWorkgroup = [64]} // CHECK-NEXT: ^bb0(%[[ARG0:[a-zA-Z0-9]+]]: index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK-DAG: %[[D0:.+]] = affine.apply #[[MAP]]()[%[[ARG0]]] // CHECK: hal.return %[[D0]], %[[C1]], %[[C1]] @@ -1128,7 +1128,7 @@ hal.executable.entry_point public @test_exp_2 attributes {interface = @io, ordinal = 0 : index} builtin.module { func @test_exp_2() { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %size = hal.interface.workgroup.size[0] : index %count = hal.interface.workgroup.count[0] : index %id = hal.interface.workgroup.id[0] : index @@ -1156,7 +1156,7 @@ // CHECK: hal.executable.entry_point public @test_exp_2 attributes // CHECK-SAME: translation.info = {passPipeline = "CPUDefault", workloadPerWorkgroup = [64]} // CHECK-NEXT: ^bb0(%[[ARG0:[a-zA-Z0-9]+]]: index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK-DAG: %[[D0:.+]] = affine.apply #[[MAP]]()[%[[ARG0]]] // CHECK: hal.return %[[D0]], %[[C1]], %[[C1]] @@ -1167,7 +1167,7 @@ hal.executable.entry_point public @test_exp_3 attributes {interface = @io, ordinal = 0 : index} builtin.module { func @test_exp_3() { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %size = hal.interface.workgroup.size[0] : index %count = hal.interface.workgroup.count[0] : index %id = hal.interface.workgroup.id[0] : index @@ -1195,7 +1195,7 @@ // CHECK: hal.executable.entry_point public @test_exp_3 attributes // CHECK-SAME: translation.info = {passPipeline = "CPUDefault", workloadPerWorkgroup = [64]} // CHECK-NEXT: ^bb0(%[[ARG0:[a-zA-Z0-9]+]]: index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK-DAG: %[[D0:.+]] = affine.apply #[[MAP]]()[%[[ARG0]]] // CHECK: hal.return %[[D0]], %[[C1]], %[[C1]] @@ -1206,7 +1206,7 @@ hal.executable.entry_point public @test_exp_4 attributes {interface = @io, ordinal = 0 : index} builtin.module { func @test_exp_4() { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %size = hal.interface.workgroup.size[0] : index %count = hal.interface.workgroup.count[0] : index %id = hal.interface.workgroup.id[0] : index @@ -1234,7 +1234,7 @@ // CHECK: hal.executable.entry_point public @test_exp_4 attributes // CHECK-SAME: translation.info = {passPipeline = "CPUDefault", workloadPerWorkgroup = [64]} // CHECK-NEXT: ^bb0(%[[ARG0:[a-zA-Z0-9]+]]: index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK-DAG: %[[D0:.+]] = affine.apply #[[MAP]]()[%[[ARG0]]] // CHECK: hal.return %[[D0]], %[[C1]], %[[C1]] @@ -1245,7 +1245,7 @@ hal.executable.entry_point public @test_exp_5 attributes {interface = @io, ordinal = 0 : index} builtin.module { func @test_exp_5() { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %size = hal.interface.workgroup.size[0] : index %count = hal.interface.workgroup.count[0] : index %id = hal.interface.workgroup.id[0] : index @@ -1273,6 +1273,6 @@ // CHECK: hal.executable.entry_point public @test_exp_5 attributes // CHECK-SAME: translation.info = {passPipeline = "CPUDefault", workloadPerWorkgroup = [64]} // CHECK-NEXT: ^bb0(%[[ARG0:[a-zA-Z0-9]+]]: index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK-DAG: %[[D0:.+]] = affine.apply #[[MAP]]()[%[[ARG0]]] // CHECK: hal.return %[[D0]], %[[C1]], %[[C1]]
diff --git a/iree/compiler/Codegen/LLVMCPU/test/matmul_vectorization.mlir b/iree/compiler/Codegen/LLVMCPU/test/matmul_vectorization.mlir index 53d7bae..e491217 100644 --- a/iree/compiler/Codegen/LLVMCPU/test/matmul_vectorization.mlir +++ b/iree/compiler/Codegen/LLVMCPU/test/matmul_vectorization.mlir
@@ -15,8 +15,8 @@ } builtin.module { func @matmul_128x128x128() { - %c0 = constant 0 : index - %c128 = constant 128 : index + %c0 = arith.constant 0 : index + %c128 = arith.constant 128 : index %0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<128x128xf32> %1 = hal.interface.binding.subspan @io::@arg1[%c0] : memref<128x128xf32> %2 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<128x128xf32> @@ -45,11 +45,11 @@ // CHECK-DAG: %[[ARG0:.+]] = hal.interface.binding.subspan @io::@arg0 // CHECK-DAG: %[[ARG1:.+]] = hal.interface.binding.subspan @io::@arg1 // CHECK-DAG: %[[RET0:.+]] = hal.interface.binding.subspan @io::@ret0 -// CHECK-DAG: %[[START:.+]] = constant 0 : index -// CHECK-DAG: %[[WORGKROUP_SIZE:.+]] = constant 64 -// CHECK-DAG: %[[VECTOR_SIZE:.+]] = constant 4 -// CHECK-DAG: %[[L1_SIZE:.+]] = constant 32 -// CHECK-DAG: %[[KDIM_SIZE:.+]] = constant 128 +// CHECK-DAG: %[[START:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[WORGKROUP_SIZE:.+]] = arith.constant 64 +// CHECK-DAG: %[[VECTOR_SIZE:.+]] = arith.constant 4 +// CHECK-DAG: %[[L1_SIZE:.+]] = arith.constant 32 +// CHECK-DAG: %[[KDIM_SIZE:.+]] = arith.constant 128 // CHECK: scf.for // CHECK: scf.for // CHECK: scf.for {{.*}} = %[[START]] to %[[WORGKROUP_SIZE]] step %[[L1_SIZE]] { @@ -74,14 +74,14 @@ // CHECK-PROMOTED: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 * 64)> // CHECK-PROMOTED: #[[MAP1:.+]] = affine_map<(d0, d1)[s0] -> (d0 * 128 + s0 + d1)> // CHECK-PROMOTED: func @matmul_128x128x128 -// CHECK-PROMOTED-DAG: %[[KDIM_SIZE:.+]] = constant 128 : index -// CHECK-PROMOTED-DAG: %[[WORGKROUP_SIZE:.+]] = constant 64 : index -// CHECK-PROMOTED-DAG: %[[VECTOR_SIZE:.+]] = constant 4 : index -// CHECK-PROMOTED-DAG: %[[L1_SIZE:.+]] = constant 32 : index -// CHECK-PROMOTED-DAG: %[[START:.+]] = constant 0 : index -// CHECK-PROMOTED-DAG: %[[C1:.+]] = constant 1 : index -// CHECK-PROMOTED-DAG: %[[C1:.+]] = constant 2 : index -// CHECK-PROMOTED-DAG: %[[C1:.+]] = constant 3 : index +// CHECK-PROMOTED-DAG: %[[KDIM_SIZE:.+]] = arith.constant 128 : index +// CHECK-PROMOTED-DAG: %[[WORGKROUP_SIZE:.+]] = arith.constant 64 : index +// CHECK-PROMOTED-DAG: %[[VECTOR_SIZE:.+]] = arith.constant 4 : index +// CHECK-PROMOTED-DAG: %[[L1_SIZE:.+]] = arith.constant 32 : index +// CHECK-PROMOTED-DAG: %[[START:.+]] = arith.constant 0 : index +// CHECK-PROMOTED-DAG: %[[C1:.+]] = arith.constant 1 : index +// CHECK-PROMOTED-DAG: %[[C1:.+]] = arith.constant 2 : index +// CHECK-PROMOTED-DAG: %[[C1:.+]] = arith.constant 3 : index // CHECK-PROMOTED-DAG: %[[C_PROMOTED_TILE:.+]] = memref.alloca() : memref<64x64xf32> // CHECK-PROMOTED-DAG: %[[B_PROMOTED_TILE:.+]] = memref.alloca() : memref<128x64xf32> // CHECK-PROMOTED-DAG: %[[A_PROMOTED_TILE:.+]] = memref.alloca() : memref<64x128xf32> @@ -128,8 +128,8 @@ } builtin.module { func @matmul_i8_i8_i32_128x128x128() { - %c0 = constant 0 : index - %c128 = constant 128 : index + %c0 = arith.constant 0 : index + %c128 = arith.constant 128 : index %0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<128x128xi8> %1 = hal.interface.binding.subspan @io::@arg1[%c0] : memref<128x128xi8> %2 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<128x128xi32> @@ -158,11 +158,11 @@ // CHECK-DAG: %[[ARG0:.+]] = hal.interface.binding.subspan @io::@arg0 // CHECK-DAG: %[[ARG1:.+]] = hal.interface.binding.subspan @io::@arg1 // CHECK-DAG: %[[RET0:.+]] = hal.interface.binding.subspan @io::@ret0 -// CHECK-DAG: %[[START:.+]] = constant 0 : index -// CHECK-DAG: %[[WORGKROUP_SIZE:.+]] = constant 64 -// CHECK-DAG: %[[VECTOR_SIZE:.+]] = constant 4 -// CHECK-DAG: %[[L1_SIZE:.+]] = constant 32 -// CHECK-DAG: %[[KDIM_SIZE:.+]] = constant 128 +// CHECK-DAG: %[[START:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[WORGKROUP_SIZE:.+]] = arith.constant 64 +// CHECK-DAG: %[[VECTOR_SIZE:.+]] = arith.constant 4 +// CHECK-DAG: %[[L1_SIZE:.+]] = arith.constant 32 +// CHECK-DAG: %[[KDIM_SIZE:.+]] = arith.constant 128 // CHECK: scf.for // CHECK: scf.for // CHECK: scf.for {{.*}} = %[[START]] to %[[WORGKROUP_SIZE]] step %[[L1_SIZE]] {
diff --git a/iree/compiler/Codegen/LLVMCPU/test/tile_and_vectorize.mlir b/iree/compiler/Codegen/LLVMCPU/test/tile_and_vectorize.mlir index a0906cb..ddc4b8f 100644 --- a/iree/compiler/Codegen/LLVMCPU/test/tile_and_vectorize.mlir +++ b/iree/compiler/Codegen/LLVMCPU/test/tile_and_vectorize.mlir
@@ -9,10 +9,10 @@ #map4 = affine_map<(d0) -> (-d0 + 513, 64)> module { func @dot_383x383x513_dispatch_0() { - %c0 = constant 0 : index - %c513 = constant 513 : index - %c383 = constant 383 : index - %cst = constant 0.000000e+00 : f32 + %c0 = arith.constant 0 : index + %c513 = arith.constant 513 : index + %c383 = arith.constant 383 : index + %cst = arith.constant 0.000000e+00 : f32 %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : !flow.dispatch.tensor<readonly:383x383xf32> %1 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : !flow.dispatch.tensor<readonly:383x513xf32> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : !flow.dispatch.tensor<writeonly:383x513xf32> @@ -51,12 +51,12 @@ // CHECK: #[[MAP5:.+]] = affine_map<(d0, d1) -> (32, -d0 + d1)> // CHECK: #[[MAP6:.+]] = affine_map<(d0) -> (32, -d0 + 383)> // CHECK: @dot_383x383x513_dispatch_0 -// CHECK-DAG: %[[CST:.+]] = constant 0.000000e+00 : f32 -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C4:.+]] = constant 4 : index -// CHECK-DAG: %[[C383:.+]] = constant 383 : index -// CHECK-DAG: %[[C513:.+]] = constant 513 : index -// CHECK-DAG: %[[C32:.+]] = constant 32 : index +// CHECK-DAG: %[[CST:.+]] = arith.constant 0.000000e+00 : f32 +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C4:.+]] = arith.constant 4 : index +// CHECK-DAG: %[[C383:.+]] = arith.constant 383 : index +// CHECK-DAG: %[[C513:.+]] = arith.constant 513 : index +// CHECK-DAG: %[[C32:.+]] = arith.constant 32 : index // CHECK: %[[LHS:.+]] = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : !flow.dispatch.tensor<readonly:383x383xf32> // CHECK: %[[RHS:.+]] = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : !flow.dispatch.tensor<readonly:383x513xf32> // CHECK: %[[DST:.+]] = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : !flow.dispatch.tensor<writeonly:383x513xf32>
diff --git a/iree/compiler/Codegen/LLVMCPU/test/vector_contract_to_aarch64_asm.mlir b/iree/compiler/Codegen/LLVMCPU/test/vector_contract_to_aarch64_asm.mlir index 185bac5..d3f86e3 100644 --- a/iree/compiler/Codegen/LLVMCPU/test/vector_contract_to_aarch64_asm.mlir +++ b/iree/compiler/Codegen/LLVMCPU/test/vector_contract_to_aarch64_asm.mlir
@@ -1,14 +1,14 @@ // RUN: iree-opt -iree-llvmcpu-vector-to-aarch64-inline-asm %s | IreeFileCheck %s func @vector_matmul_to_aarch64_asm_vec_dot(%lhs: memref<4x4xi8>, %rhs: memref<4x4xi8>, %dst: memref<4x4xi32>) { - %c0 = constant 0 : index - %cst_i8_0 = constant 0 : i8 - %cst_i32_0 = constant 0 : i32 + %c0 = arith.constant 0 : index + %cst_i8_0 = arith.constant 0 : i8 + %cst_i32_0 = arith.constant 0 : i32 %0 = vector.transfer_read %lhs[%c0, %c0] , %cst_i8_0 {in_bounds = [false, false]} : memref<4x4xi8>, vector<4x4xi8> %1 = vector.transfer_read %rhs[%c0, %c0] , %cst_i8_0 {in_bounds = [false, false]} : memref<4x4xi8>, vector<4x4xi8> %2 = vector.transfer_read %dst[%c0, %c0], %cst_i32_0 {in_bounds = [false, false]} : memref<4x4xi32>, vector<4x4xi32> - %3 = sexti %0 : vector<4x4xi8> to vector<4x4xi32> - %4 = sexti %1 : vector<4x4xi8> to vector<4x4xi32> + %3 = arith.extsi %0 : vector<4x4xi8> to vector<4x4xi32> + %4 = arith.extsi %1 : vector<4x4xi8> to vector<4x4xi32> %5 = vector.contract { indexing_maps = [ affine_map<(d0, d1, d2) -> (d0, d2)>, @@ -20,7 +20,7 @@ return } // CHEC-LABEL: @vector_matmul_to_aarch64_asm_vec_dot -// CHECK-DAG: %[[RES_2D:.+]] = constant dense<0> : vector<4x4xi32> +// CHECK-DAG: %[[RES_2D:.+]] = arith.constant dense<0> : vector<4x4xi32> // CHECK-DAG: %[[LHS_2D:.+]] = vector.transfer_read // CHECK-DAG: %[[RHS_2d:.+]] = vector.transfer_read // CHECK-DAG: %[[DST:.+]] = vector.transfer_read
diff --git a/iree/compiler/Codegen/LLVMGPU/BUILD b/iree/compiler/Codegen/LLVMGPU/BUILD index 0012dc6..e9cc75d 100644 --- a/iree/compiler/Codegen/LLVMGPU/BUILD +++ b/iree/compiler/Codegen/LLVMGPU/BUILD
@@ -45,6 +45,7 @@ "@llvm-project//llvm:Support", "@llvm-project//mlir:Affine", "@llvm-project//mlir:AffineToStandard", + "@llvm-project//mlir:ArithmeticToLLVM", "@llvm-project//mlir:GPUDialect", "@llvm-project//mlir:GPUToNVVMTransforms", "@llvm-project//mlir:GPUToROCDLTransforms",
diff --git a/iree/compiler/Codegen/LLVMGPU/CMakeLists.txt b/iree/compiler/Codegen/LLVMGPU/CMakeLists.txt index 9540af2..cf63b73 100644 --- a/iree/compiler/Codegen/LLVMGPU/CMakeLists.txt +++ b/iree/compiler/Codegen/LLVMGPU/CMakeLists.txt
@@ -34,6 +34,7 @@ LLVMSupport MLIRAffine MLIRAffineToStandard + MLIRArithmeticToLLVM MLIRGPUOps MLIRGPUToNVVMTransforms MLIRGPUToROCDLTransforms
diff --git a/iree/compiler/Codegen/LLVMGPU/ConvertToLLVM.cpp b/iree/compiler/Codegen/LLVMGPU/ConvertToLLVM.cpp index 211643d..0d2f57d 100644 --- a/iree/compiler/Codegen/LLVMGPU/ConvertToLLVM.cpp +++ b/iree/compiler/Codegen/LLVMGPU/ConvertToLLVM.cpp
@@ -37,7 +37,7 @@ auto vecType = mathOp.getType().template dyn_cast<VectorType>(); if (!vecType) return failure(); Location loc = mathOp.getLoc(); - Value newVector = rewriter.create<ConstantOp>( + Value newVector = rewriter.create<arith::ConstantOp>( loc, vecType, rewriter.getZeroAttr(vecType)); for (int64_t element : llvm::seq(int64_t(0), vecType.getNumElements())) { @@ -402,11 +402,11 @@ } void populateScalarizeMathOps(RewritePatternSet &patterns) { - patterns.add<ScalarizeMathOp<math::SqrtOp>, ScalarizeMathOp<AbsFOp>, + patterns.add<ScalarizeMathOp<math::SqrtOp>, ScalarizeMathOp<math::AbsOp>, ScalarizeMathOp<math::AtanOp>, ScalarizeMathOp<math::Atan2Op>, - ScalarizeMathOp<CeilFOp>, ScalarizeMathOp<math::CosOp>, + ScalarizeMathOp<math::CeilOp>, ScalarizeMathOp<math::CosOp>, ScalarizeMathOp<math::ExpOp>, ScalarizeMathOp<math::ExpM1Op>, - ScalarizeMathOp<FloorFOp>, ScalarizeMathOp<math::LogOp>, + ScalarizeMathOp<math::FloorOp>, ScalarizeMathOp<math::LogOp>, ScalarizeMathOp<math::Log1pOp>, ScalarizeMathOp<math::Log10Op>, ScalarizeMathOp<math::Log2Op>, ScalarizeMathOp<math::PowFOp>, ScalarizeMathOp<math::RsqrtOp>, ScalarizeMathOp<math::SinOp>,
diff --git a/iree/compiler/Codegen/LLVMGPU/ConvertToNVVM.cpp b/iree/compiler/Codegen/LLVMGPU/ConvertToNVVM.cpp index e22f057..5cf1c24 100644 --- a/iree/compiler/Codegen/LLVMGPU/ConvertToNVVM.cpp +++ b/iree/compiler/Codegen/LLVMGPU/ConvertToNVVM.cpp
@@ -9,6 +9,7 @@ #include "iree/compiler/Codegen/Passes.h" #include "iree/compiler/Codegen/Utils/Utils.h" #include "iree/compiler/Dialect/Util/IR/UtilOps.h" +#include "mlir/Conversion/ArithmeticToLLVM/ArithmeticToLLVM.h" #include "mlir/Conversion/GPUToNVVM/GPUToNVVMPass.h" #include "mlir/Conversion/LLVMCommon/ConversionTarget.h" #include "mlir/Conversion/LLVMCommon/LoweringOptions.h" @@ -77,6 +78,8 @@ populateMathToLLVMConversionPatterns(converter, llvmPatterns); populateMemRefToLLVMConversionPatterns(converter, llvmPatterns); populateStdToLLVMConversionPatterns(converter, llvmPatterns); + arith::populateArithmeticToLLVMConversionPatterns(converter, + llvmPatterns); populateVectorToLLVMConversionPatterns(converter, llvmPatterns); populateGpuToNVVMConversionPatterns(converter, llvmPatterns); LLVMConversionTarget target(getContext());
diff --git a/iree/compiler/Codegen/LLVMGPU/ConvertToROCDL.cpp b/iree/compiler/Codegen/LLVMGPU/ConvertToROCDL.cpp index afaaa52..465db82 100644 --- a/iree/compiler/Codegen/LLVMGPU/ConvertToROCDL.cpp +++ b/iree/compiler/Codegen/LLVMGPU/ConvertToROCDL.cpp
@@ -9,6 +9,7 @@ #include "iree/compiler/Codegen/Passes.h" #include "iree/compiler/Codegen/Utils/Utils.h" #include "iree/compiler/Dialect/Util/IR/UtilOps.h" +#include "mlir/Conversion/ArithmeticToLLVM/ArithmeticToLLVM.h" #include "mlir/Conversion/GPUToROCDL/GPUToROCDLPass.h" #include "mlir/Conversion/LLVMCommon/ConversionTarget.h" #include "mlir/Conversion/LLVMCommon/LoweringOptions.h" @@ -77,6 +78,8 @@ populateMathToLLVMConversionPatterns(converter, llvmPatterns); populateMemRefToLLVMConversionPatterns(converter, llvmPatterns); populateStdToLLVMConversionPatterns(converter, llvmPatterns); + arith::populateArithmeticToLLVMConversionPatterns(converter, + llvmPatterns); populateVectorToLLVMConversionPatterns(converter, llvmPatterns); populateGpuToROCDLConversionPatterns(converter, llvmPatterns); LLVMConversionTarget target(getContext());
diff --git a/iree/compiler/Codegen/LLVMGPU/LLVMGPUDistributeSharedMemoryCopy.cpp b/iree/compiler/Codegen/LLVMGPU/LLVMGPUDistributeSharedMemoryCopy.cpp index 091c36b..a07d472 100644 --- a/iree/compiler/Codegen/LLVMGPU/LLVMGPUDistributeSharedMemoryCopy.cpp +++ b/iree/compiler/Codegen/LLVMGPU/LLVMGPUDistributeSharedMemoryCopy.cpp
@@ -45,10 +45,10 @@ for (unsigned i = 0; i < rank - 1; i++) { int64_t t = (rank - i) <= kNumGPUDims ? 1 : 0; tileSizesVal.push_back( - builder.create<ConstantIndexOp>(operation->getLoc(), t)); + builder.create<arith::ConstantIndexOp>(operation->getLoc(), t)); } - tileSizesVal.push_back( - builder.create<ConstantIndexOp>(operation->getLoc(), copyTileSize)); + tileSizesVal.push_back(builder.create<arith::ConstantIndexOp>( + operation->getLoc(), copyTileSize)); return tileSizesVal; }; auto getCopyThreadProcInfoFn = [workgroupSize](
diff --git a/iree/compiler/Codegen/LLVMGPU/LLVMGPUTileAndDistribute.cpp b/iree/compiler/Codegen/LLVMGPU/LLVMGPUTileAndDistribute.cpp index 9c0b895..26e0da5 100644 --- a/iree/compiler/Codegen/LLVMGPU/LLVMGPUTileAndDistribute.cpp +++ b/iree/compiler/Codegen/LLVMGPU/LLVMGPUTileAndDistribute.cpp
@@ -42,9 +42,9 @@ Location loc = op->getLoc(); auto tileSizesVal = llvm::to_vector<4>(llvm::map_range(tileSizes, [&](int64_t v) -> Value { - return builder.create<ConstantIndexOp>(loc, v); + return builder.create<arith::ConstantIndexOp>(loc, v); })); - auto zero = builder.create<ConstantIndexOp>(loc, 0); + auto zero = builder.create<arith::ConstantIndexOp>(loc, 0); for (unsigned depth : partitionedLoops) { if (depth < tileSizesVal.size()) { tileSizesVal[depth] = zero; @@ -88,7 +88,7 @@ for (auto val : llvm::enumerate(tileSizes)) { int64_t useTileSize = partitionedLoopsSet.count(val.index()) ? val.value() : 0; - tileSizesVal.push_back(builder.create<ConstantIndexOp>( + tileSizesVal.push_back(builder.create<arith::ConstantIndexOp>( operation->getLoc(), useTileSize)); } return tileSizesVal;
diff --git a/iree/compiler/Codegen/LLVMGPU/LLVMGPUUtils.h b/iree/compiler/Codegen/LLVMGPU/LLVMGPUUtils.h index 16ddcab..856b78c 100644 --- a/iree/compiler/Codegen/LLVMGPU/LLVMGPUUtils.h +++ b/iree/compiler/Codegen/LLVMGPU/LLVMGPUUtils.h
@@ -24,7 +24,7 @@ mlir::StringAttr attr = builder.getStringAttr(dimAttr[i]); procInfo[numDims - 1 - i] = { builder.create<mlir::gpu::ThreadIdOp>(loc, indexType, attr), - builder.create<mlir::ConstantOp>( + builder.create<mlir::arith::ConstantOp>( loc, builder.getIndexAttr(workgroupSize[i]))}; } return procInfo;
diff --git a/iree/compiler/Codegen/LLVMGPU/test/convert_to_nvvm.mlir b/iree/compiler/Codegen/LLVMGPU/test/convert_to_nvvm.mlir index 401a810..645a84e 100644 --- a/iree/compiler/Codegen/LLVMGPU/test/convert_to_nvvm.mlir +++ b/iree/compiler/Codegen/LLVMGPU/test/convert_to_nvvm.mlir
@@ -2,20 +2,20 @@ // Test that that standard and GPU ops are converted to LLVM and NVVM. func @abs_ex_dispatch_0() { - %c0 = constant 0 : index - %c128 = constant 128 : index + %c0 = arith.constant 0 : index + %c128 = arith.constant 128 : index %0 = hal.interface.binding.subspan @io::@arg0[%c128] : memref<16xf32> %1 = hal.interface.binding.subspan @io::@arg1[%c0] : memref<16xi32> %2 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<16xf32> %3 = "gpu.block_id"() {dimension = "x"} : () -> index %4 = "gpu.block_dim"() {dimension = "x"} : () -> index %5 = "gpu.thread_id"() {dimension = "x"} : () -> index - %6 = muli %3, %4 : index - %7 = addi %6, %5 : index + %6 = arith.muli %3, %4 : index + %7 = arith.addi %6, %5 : index %9 = memref.load %0[%7] : memref<16xf32> %10 = memref.load %1[%7] : memref<16xi32> - %11 = sitofp %10 : i32 to f32 - %12 = addf %9, %11 : f32 + %11 = arith.sitofp %10 : i32 to f32 + %12 = arith.addf %9, %11 : f32 memref.store %12, %2[%7] : memref<16xf32> return } @@ -41,8 +41,8 @@ // ----- func @abs_dynamic() { - %c0 = constant 0 : index - %c128 = constant 128 : index + %c0 = arith.constant 0 : index + %c128 = arith.constant 128 : index %s = hal.interface.load.constant offset = 1 : index %0 = hal.interface.binding.subspan @io::@arg0[%c128] : memref<?xf32>{%s} %1 = hal.interface.binding.subspan @io::@arg1[%c0] : memref<16xi32> @@ -50,12 +50,12 @@ %3 = "gpu.block_id"() {dimension = "x"} : () -> index %4 = "gpu.block_dim"() {dimension = "x"} : () -> index %5 = "gpu.thread_id"() {dimension = "x"} : () -> index - %6 = muli %3, %4 : index - %7 = addi %6, %5 : index + %6 = arith.muli %3, %4 : index + %7 = arith.addi %6, %5 : index %9 = memref.load %0[%7] : memref<?xf32> %10 = memref.load %1[%7] : memref<16xi32> - %11 = sitofp %10 : i32 to f32 - %12 = addf %9, %11 : f32 + %11 = arith.sitofp %10 : i32 to f32 + %12 = arith.addf %9, %11 : f32 memref.store %12, %2[%7] : memref<16xf32> return } @@ -83,18 +83,18 @@ // Test that we handle correctly the case where a symbol is dead. func @dead_symbol() { - %c0 = constant 0 : index - %c128 = constant 128 : index + %c0 = arith.constant 0 : index + %c128 = arith.constant 128 : index %1 = hal.interface.binding.subspan @io::@arg1[%c0] : memref<16xi32> %2 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<16xf32> %3 = "gpu.block_id"() {dimension = "x"} : () -> index %4 = "gpu.block_dim"() {dimension = "x"} : () -> index %5 = "gpu.thread_id"() {dimension = "x"} : () -> index - %6 = muli %3, %4 : index - %7 = addi %6, %5 : index + %6 = arith.muli %3, %4 : index + %7 = arith.addi %6, %5 : index %10 = memref.load %1[%7] : memref<16xi32> - %11 = sitofp %10 : i32 to f32 - %12 = addf %11, %11 : f32 + %11 = arith.sitofp %10 : i32 to f32 + %12 = arith.addf %11, %11 : f32 memref.store %12, %2[%7] : memref<16xf32> return } @@ -116,20 +116,20 @@ // A single binding may contain different data types. // Test that we cast pointers correctly. func @mixed_type() { - %c0 = constant 0 : index - %c128 = constant 128 : index + %c0 = arith.constant 0 : index + %c128 = arith.constant 128 : index %0 = hal.interface.binding.subspan @io::@arg0[%c128] : memref<16xf32> %1 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<16xi32> %2 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<16xf32> %3 = "gpu.block_id"() {dimension = "x"} : () -> index %4 = "gpu.block_dim"() {dimension = "x"} : () -> index %5 = "gpu.thread_id"() {dimension = "x"} : () -> index - %6 = muli %3, %4 : index - %7 = addi %6, %5 : index + %6 = arith.muli %3, %4 : index + %7 = arith.addi %6, %5 : index %9 = memref.load %0[%7] : memref<16xf32> %10 = memref.load %1[%7] : memref<16xi32> - %11 = sitofp %10 : i32 to f32 - %12 = addf %9, %11 : f32 + %11 = arith.sitofp %10 : i32 to f32 + %12 = arith.addf %9, %11 : f32 memref.store %12, %2[%7] : memref<16xf32> return }
diff --git a/iree/compiler/Codegen/LLVMGPU/test/convert_to_rocdl.mlir b/iree/compiler/Codegen/LLVMGPU/test/convert_to_rocdl.mlir index 3dd2b57..11a91ae 100644 --- a/iree/compiler/Codegen/LLVMGPU/test/convert_to_rocdl.mlir +++ b/iree/compiler/Codegen/LLVMGPU/test/convert_to_rocdl.mlir
@@ -2,18 +2,18 @@ // Test that that standard and GPU ops are converted to LLVM and NVVM. func @abs_ex_dispatch_0() { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<16xf32> %1 = hal.interface.binding.subspan @io::@arg1[%c0] : memref<16xf32> %2 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<16xf32> %3 = "gpu.block_id"() {dimension = "x"} : () -> index %4 = "gpu.block_dim"() {dimension = "x"} : () -> index %5 = "gpu.thread_id"() {dimension = "x"} : () -> index - %6 = muli %3, %4 : index - %7 = addi %6, %5 : index + %6 = arith.muli %3, %4 : index + %7 = arith.addi %6, %5 : index %9 = memref.load %1[%7] : memref<16xf32> %10 = memref.load %2[%7] : memref<16xf32> - %11 = addf %9, %10 : f32 + %11 = arith.addf %9, %10 : f32 memref.store %11, %0[%7] : memref<16xf32> return }
diff --git a/iree/compiler/Codegen/LLVMGPU/test/distribute_to_thread.mlir b/iree/compiler/Codegen/LLVMGPU/test/distribute_to_thread.mlir index 03cc0c5..2b7e6da 100644 --- a/iree/compiler/Codegen/LLVMGPU/test/distribute_to_thread.mlir +++ b/iree/compiler/Codegen/LLVMGPU/test/distribute_to_thread.mlir
@@ -16,9 +16,9 @@ workgroup_size = [64 : index, 1 : index, 1 : index]} builtin.module { builtin.func @dot_dispatch_0() { - %cst = constant 0.000000e+00 : f32 - %c0 = constant 0 : index - %c1024 = constant 1024 : index + %cst = arith.constant 0.000000e+00 : f32 + %c0 = arith.constant 0 : index + %c1024 = arith.constant 1024 : index %0 = hal.interface.binding.subspan @legacy_io::@ro0[%c0] : memref<1024x1024xf32> %1 = hal.interface.binding.subspan @legacy_io::@ro1[%c0] : memref<1024x1024xf32> %2 = hal.interface.binding.subspan @legacy_io::@wo2[%c0] : memref<1024x1024xf32> @@ -64,11 +64,11 @@ // CHECK: hal.executable.variant public @cuda // CHECK: memref.global "private" @{{.*}} : memref<4x256xf32, 3> // CHECK: memref.global "private" @{{.*}} : memref<2x4xf32, 3> -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C2:.+]] = constant 2 : index -// CHECK-DAG: %[[C4:.+]] = constant 4 : index -// CHECK-DAG: %[[C256:.+]] = constant 256 : index -// CHECK-DAG: %[[C1024:.+]] = constant 1024 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C2:.+]] = arith.constant 2 : index +// CHECK-DAG: %[[C4:.+]] = arith.constant 4 : index +// CHECK-DAG: %[[C256:.+]] = arith.constant 256 : index +// CHECK-DAG: %[[C1024:.+]] = arith.constant 1024 : index // CHECK: scf.for %[[K:.+]] = %[[C0]] to %[[C1024]] step %[[C4]] { // CHECK: gpu.barrier // CHECK: linalg.copy(%{{.*}}, %{{.*}}) {__internal_linalg_transform__ = "copy_to_workgroup_memory"} : memref<2x4xf32, #{{.*}}>, memref<2x4xf32, #{{.*}}, 3> @@ -97,17 +97,17 @@ workgroup_size = [1: index, 1: index, 1: index]} builtin.module { builtin.func @predict_dispatch_153() { - %c0 = constant 0 : index - %cst = constant 0x7FC00000 : f32 - %cst_0 = constant 0xFF800000 : f32 + %c0 = arith.constant 0 : index + %cst = arith.constant 0x7FC00000 : f32 + %cst_0 = arith.constant 0xFF800000 : f32 %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : memref<1000xf32> %1 = hal.interface.binding.subspan @io::@s0b1_xw_external[%c0] : memref<f32> linalg.fill(%cst_0, %1) {lowering.config = #config} : f32, memref<f32> linalg.generic {indexing_maps = [affine_map<(d0) -> (d0)>, affine_map<(d0) -> ()>], iterator_types = ["reduction"]} ins(%0 : memref<1000xf32>) outs(%1 : memref<f32>) attrs = {lowering.config = #config} { ^bb0(%arg0: f32, %arg1: f32): // no predecessors - %2 = cmpf ogt, %arg0, %arg1 : f32 + %2 = arith.cmpf ogt, %arg0, %arg1 : f32 %3 = select %2, %arg0, %arg1 : f32 - %4 = cmpf uno, %arg0, %arg1 : f32 + %4 = arith.cmpf uno, %arg0, %arg1 : f32 %5 = select %4, %cst, %3 : f32 linalg.yield %5 : f32 }
diff --git a/iree/compiler/Codegen/LLVMGPU/test/distribute_wg_copy.mlir b/iree/compiler/Codegen/LLVMGPU/test/distribute_wg_copy.mlir index 47e7b94..c75628a 100644 --- a/iree/compiler/Codegen/LLVMGPU/test/distribute_wg_copy.mlir +++ b/iree/compiler/Codegen/LLVMGPU/test/distribute_wg_copy.mlir
@@ -24,9 +24,9 @@ %sm1 = memref.get_global @__shared_memory___0 : memref<256x4xf32, 3> %sm2 = memref.get_global @__shared_memory___1 : memref<3x512xf32, 3> gpu.barrier - // CHECK-DAG: %[[C2:.*]] = constant 2 : index - // CHECK-DAG: %[[C1:.*]] = constant 1 : index - // CHECK-DAG: %[[C0:.*]] = constant 0 : index + // CHECK-DAG: %[[C2:.*]] = arith.constant 2 : index + // CHECK-DAG: %[[C1:.*]] = arith.constant 1 : index + // CHECK-DAG: %[[C0:.*]] = arith.constant 0 : index // CHECK-DAG: %[[TX:.*]] = "gpu.thread_id"() {dimension = "x"} : () -> index // CHECK-DAG: %[[TY:.*]] = "gpu.thread_id"() {dimension = "y"} : () -> index // CHECK-DAG: %[[TZ:.*]] = "gpu.thread_id"() {dimension = "z"} : () -> index
diff --git a/iree/compiler/Codegen/LLVMGPU/test/gpu_set_num_workgroups.mlir b/iree/compiler/Codegen/LLVMGPU/test/gpu_set_num_workgroups.mlir index d72539c..ec80d39 100644 --- a/iree/compiler/Codegen/LLVMGPU/test/gpu_set_num_workgroups.mlir +++ b/iree/compiler/Codegen/LLVMGPU/test/gpu_set_num_workgroups.mlir
@@ -9,7 +9,7 @@ hal.executable.entry_point @add_dispatch_0 attributes {interface = @io, ordinal = 0 : index} builtin.module { func @add_dispatch_0() { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %0 = hal.interface.binding.subspan @io::@arg0[%c0] : !flow.dispatch.tensor<readonly:16384xf32> %1 = hal.interface.binding.subspan @io::@arg1[%c0] : !flow.dispatch.tensor<readonly:16384xf32> %2 = hal.interface.binding.subspan @io::@ret0[%c0] : !flow.dispatch.tensor<writeonly:16384xf32> @@ -18,7 +18,7 @@ %5 = flow.dispatch.tensor.load %1, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:16384xf32> -> tensor<16384xf32> %6 = linalg.generic {indexing_maps = [affine_map<(d0) -> (d0)>, affine_map<(d0) -> (d0)>, affine_map<(d0) -> (d0)>], iterator_types = ["parallel"]} ins(%4, %5 : tensor<16384xf32>, tensor<16384xf32>) outs(%3 : tensor<16384xf32>) { ^bb0(%arg0: f32, %arg1: f32, %arg2: f32): // no predecessors - %7 = addf %arg0, %arg1 : f32 + %7 = arith.addf %arg0, %arg1 : f32 linalg.yield %7 : f32 } -> tensor<16384xf32> flow.dispatch.tensor.store %6, %2, offsets=[], sizes=[], strides=[] : tensor<16384xf32> -> !flow.dispatch.tensor<writeonly:16384xf32> @@ -40,7 +40,7 @@ // CHECK-SAME: workloadPerWorkgroup = [128] // CHECK-SAME: workgroup_size = [32 : index, 1 : index, 1 : index] // CHECK-NEXT: ^bb0(%[[ARG0:[a-zA-Z0-9]+]]: index, -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK-DAG: %[[NWGS_X:.+]] = affine.apply #[[MAP0]]()[%[[ARG0]]] // CHECK: hal.return %[[NWGS_X]], %[[C1]], %[[C1]] // CHECK: func @add_dispatch_0 @@ -54,10 +54,10 @@ hal.executable.entry_point @dot_dispatch_1 attributes {interface = @legacy_io, ordinal = 0 : index} builtin.module { func @dot_dispatch_1() { - %c0 = constant 0 : index - %c4 = constant 4 : index - %c2 = constant 2 : index - %cst = constant 0.000000e+00 : f32 + %c0 = arith.constant 0 : index + %c4 = arith.constant 4 : index + %c2 = arith.constant 2 : index + %cst = arith.constant 0.000000e+00 : f32 %0 = hal.interface.binding.subspan @io::@ro0[%c0] : memref<2x3xf32> %1 = hal.interface.binding.subspan @io::@ro1[%c0] : memref<3x4xf32> %2 = hal.interface.binding.subspan @io::@wo2[%c0] : memref<2x4xf32> @@ -100,7 +100,7 @@ // CHECK-SAME: workloadPerWorkgroup = [2, 4] // CHECK-SAME: workgroup_size = [2 : index, 4 : index, 1 : index] // CHECK-NEXT: ^bb0(%[[ARG0:[a-zA-Z0-9]+]]: index, %[[ARG1:[a-zA-Z0-9]+]]: index, -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK-DAG: %[[NWGS_X:.+]] = affine.apply #[[MAP0]]()[%[[ARG0]]] // CHECK-DAG: %[[NWGS_Y:.+]] = affine.apply #[[MAP1]]()[%[[ARG1]]] // CHECK: hal.return %[[NWGS_X]], %[[NWGS_Y]], %[[C1]] @@ -119,17 +119,17 @@ ordinal = 0 : index} builtin.module { func @predict_dispatch_153() { - %c0 = constant 0 : index - %cst = constant 0x7FC00000 : f32 - %cst_0 = constant 0xFF800000 : f32 + %c0 = arith.constant 0 : index + %cst = arith.constant 0x7FC00000 : f32 + %cst_0 = arith.constant 0xFF800000 : f32 %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : memref<1000xf32> %1 = hal.interface.binding.subspan @io::@s0b1_xw_external[%c0] : memref<f32> linalg.fill(%cst_0, %1) : f32, memref<f32> linalg.generic {indexing_maps = [affine_map<(d0) -> (d0)>, affine_map<(d0) -> ()>], iterator_types = ["reduction"]} ins(%0 : memref<1000xf32>) outs(%1 : memref<f32>) { ^bb0(%arg0: f32, %arg1: f32): // no predecessors - %2 = cmpf ogt, %arg0, %arg1 : f32 + %2 = arith.cmpf ogt, %arg0, %arg1 : f32 %3 = select %2, %arg0, %arg1 : f32 - %4 = cmpf uno, %arg0, %arg1 : f32 + %4 = arith.cmpf uno, %arg0, %arg1 : f32 %5 = select %4, %cst, %3 : f32 linalg.yield %5 : f32 } @@ -149,7 +149,7 @@ // CHECK-SAME: translation.info = #[[CONFIG0]] // CHECK-SAME: workgroup_size = [1 : index, 1 : index, 1 : index] // CHECK-NEXT: ^bb0(%[[ARG0:[a-zA-Z0-9]+]]: index, -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK: hal.return %[[C1]], %[[C1]], %[[C1]] // CHECK: linalg.fill // CHECK-SAME: lowering.config = #[[CONFIG1]] @@ -163,7 +163,7 @@ hal.executable.entry_point @tensor_insert_slice attributes {interface = @io, ordinal = 0 : index} builtin.module { builtin.func @tensor_insert_slice() { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : !flow.dispatch.tensor<readonly:?x?xi32> %1 = hal.interface.load.constant offset = 0 : index %2 = hal.interface.load.constant offset = 1 : index @@ -204,7 +204,7 @@ // CHECK-SAME: translation.info = {passPipeline = "LLVMGPUDistribute", workloadPerWorkgroup = [128, 1]} // CHECK-NEXT: %[[ARG0:[a-zA-Z0-9_]+]]: index // CHECK-SAME: %[[ARG1:[a-zA-Z0-9_]+]]: index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK-DAG: %[[NWGSX:.+]] = affine.apply #[[MAP0]]()[%[[ARG0]]] // CHECK: hal.return %[[NWGSX]], %[[ARG1]], %[[C1]] @@ -215,7 +215,7 @@ hal.executable.entry_point @tensor_insert_slice attributes {interface = @io, ordinal = 0 : index} builtin.module { builtin.func @tensor_insert_slice() { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %d0 = hal.interface.load.constant offset = 0 : index %d1 = hal.interface.load.constant offset = 1 : index %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : memref<?x?xi32>{%d0, %d1} @@ -252,7 +252,7 @@ // CHECK-SAME: translation.info = {passPipeline = "LLVMGPUVectorize", workloadPerWorkgroup = [128, 1]} // CHECK-NEXT: %[[ARG0:[a-zA-Z0-9_]+]]: index // CHECK-SAME: %[[ARG1:[a-zA-Z0-9_]+]]: index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK-DAG: %[[NWGSX:.+]] = affine.apply #[[MAP0]]()[%[[ARG0]]] // CHECK: hal.return %[[NWGSX]], %[[ARG1]], %[[C1]] // CHECK: linalg.copy @@ -269,10 +269,10 @@ hal.executable.entry_point @static_1d_fft_stage2 attributes {interface = @io, ordinal = 0 : index} builtin.module { builtin.func @static_1d_fft_stage2() { - %c0 = constant 0 : index - %c2 = constant 2 : index - %cst = constant dense<[1.000000e+00, 6.12323426E-17]> : tensor<2xf32> - %cst_0 = constant dense<[-0.000000e+00, -1.000000e+00]> : tensor<2xf32> + %c0 = arith.constant 0 : index + %c2 = arith.constant 2 : index + %cst = arith.constant dense<[1.000000e+00, 6.12323426E-17]> : tensor<2xf32> + %cst_0 = arith.constant dense<[-0.000000e+00, -1.000000e+00]> : tensor<2xf32> %0 = hal.interface.binding.subspan @io::@s0b0_rw_external[%c0] : !flow.dispatch.tensor<readwrite:32xf32> %1 = hal.interface.binding.subspan @io::@s0b1_rw_external[%c0] : !flow.dispatch.tensor<readwrite:32xf32> %2 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readwrite:32xf32> -> tensor<32xf32> @@ -293,7 +293,7 @@ // CHECK-SAME: workloadPerWorkgroup = [4]} // CHECK-SAME: workgroup_size = [32 : index, 1 : index, 1 : index] // CHECK-NEXT: ^{{.+}}(%[[ARG0:.+]]: index, %{{.+}}: index, %{{.+}}: index): -// CHECK-NEXT: %[[ONE:.+]] = constant 1 : index +// CHECK-NEXT: %[[ONE:.+]] = arith.constant 1 : index // CHECK-NEXT: %[[T:.+]] = affine.apply #[[MAP0]]()[%[[ARG0]]] // CHECK-NEXT: hal.return %[[T]], %[[ONE]], %[[ONE]] @@ -313,13 +313,13 @@ hal.executable.entry_point @static_3d_fft_stage3 attributes {interface = @io, ordinal = 0 : index} builtin.module { builtin.func @static_3d_fft_stage3() { - %c0 = constant 0 : index - %c3 = constant 3 : index - %c64 = constant 64 : index - %c128 = constant 128 : index - %c32 = constant 32 : index - %cst = constant dense<[1.000000e+00, 0.707106769, 6.12323426E-17, -0.707106769]> : tensor<4xf32> - %cst_0 = constant dense<[-0.000000e+00, -0.707106769, -1.000000e+00, -0.707106769]> : tensor<4xf32> + %c0 = arith.constant 0 : index + %c3 = arith.constant 3 : index + %c64 = arith.constant 64 : index + %c128 = arith.constant 128 : index + %c32 = arith.constant 32 : index + %cst = arith.constant dense<[1.000000e+00, 0.707106769, 6.12323426E-17, -0.707106769]> : tensor<4xf32> + %cst_0 = arith.constant dense<[-0.000000e+00, -0.707106769, -1.000000e+00, -0.707106769]> : tensor<4xf32> %0 = memref.buffer_cast %cst_0 : memref<4xf32> %1 = memref.buffer_cast %cst : memref<4xf32> %2 = hal.interface.binding.subspan @io::@s0b0_rw_external[%c0] : memref<64x128x32xf32> @@ -372,10 +372,10 @@ hal.executable.entry_point public @_lowering_config_test_dispatch_1 attributes {interface = @io, ordinal = 0 : index} builtin.module { func @_lowering_config_test_dispatch_1() { - %cst = constant 0.000000e+00 : f32 - %c128 = constant 128 : index - %c1024 = constant 1024 : index - %c0 = constant 0 : index + %cst = arith.constant 0.000000e+00 : f32 + %c128 = arith.constant 128 : index + %c1024 = arith.constant 1024 : index + %c0 = arith.constant 0 : index %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : !flow.dispatch.tensor<readonly:128x256xf32> %1 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : !flow.dispatch.tensor<readonly:256x1024xf32> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : !flow.dispatch.tensor<writeonly:128x1024xf32>
diff --git a/iree/compiler/Codegen/LLVMGPU/test/legalize.mlir b/iree/compiler/Codegen/LLVMGPU/test/legalize.mlir index d2df6b9..75f258f 100644 --- a/iree/compiler/Codegen/LLVMGPU/test/legalize.mlir +++ b/iree/compiler/Codegen/LLVMGPU/test/legalize.mlir
@@ -43,7 +43,7 @@ // CHECK: memref.store %{{.*}}, %[[A]][%{{.*}}, %{{.*}}] : memref<16x16xf32, 3> func @allocation(%arg0: f32) { %0 = memref.alloc() : memref<16x16xf32, 3> - %c0 = constant 0 : index + %c0 = arith.constant 0 : index memref.store %arg0, %0[%c0, %c0] : memref<16x16xf32, 3> return }
diff --git a/iree/compiler/Codegen/LLVMGPU/test/nvvm_pipeline_test.mlir b/iree/compiler/Codegen/LLVMGPU/test/nvvm_pipeline_test.mlir index c6a720a..81ed3e3 100644 --- a/iree/compiler/Codegen/LLVMGPU/test/nvvm_pipeline_test.mlir +++ b/iree/compiler/Codegen/LLVMGPU/test/nvvm_pipeline_test.mlir
@@ -12,7 +12,7 @@ hal.executable.entry_point @add_dispatch_0 attributes {interface = @io, ordinal = 0 : index} builtin.module { func @add_dispatch_0() { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %0 = hal.interface.binding.subspan @io::@arg0[%c0] : !flow.dispatch.tensor<readonly:16xf32> %1 = hal.interface.binding.subspan @io::@arg1[%c0] : !flow.dispatch.tensor<readonly:16xf32> %2 = hal.interface.binding.subspan @io::@ret0[%c0] : !flow.dispatch.tensor<writeonly:16xf32> @@ -21,7 +21,7 @@ %5 = flow.dispatch.tensor.load %1, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32> %6 = linalg.generic {indexing_maps = [affine_map<(d0) -> (d0)>, affine_map<(d0) -> (d0)>, affine_map<(d0) -> (d0)>], iterator_types = ["parallel"]} ins(%4, %5 : tensor<16xf32>, tensor<16xf32>) outs(%3 : tensor<16xf32>) { ^bb0(%arg0: f32, %arg1: f32, %arg2: f32): // no predecessors - %7 = addf %arg0, %arg1 : f32 + %7 = arith.addf %arg0, %arg1 : f32 linalg.yield %7 : f32 } -> tensor<16xf32> flow.dispatch.tensor.store %6, %2, offsets=[], sizes=[], strides=[] : tensor<16xf32> -> !flow.dispatch.tensor<writeonly:16xf32> @@ -55,10 +55,10 @@ hal.executable.entry_point @dot_dispatch_0 attributes {interface = @io, ordinal = 0 : index} builtin.module { func @dot_dispatch_0() { - %cst = constant 0.000000e+00 : f32 - %c0 = constant 0 : index - %c1024 = constant 1024 : index - %c1 = constant 1 : index + %cst = arith.constant 0.000000e+00 : f32 + %c0 = arith.constant 0 : index + %c1024 = arith.constant 1024 : index + %c1 = arith.constant 1 : index %0 = hal.interface.binding.subspan @io::@ro0[%c0] : !flow.dispatch.tensor<readonly:1024x1024xf32> %1 = hal.interface.binding.subspan @io::@ro1[%c0] : !flow.dispatch.tensor<readonly:1024x1024xf32> %2 = hal.interface.binding.subspan @io::@wo2[%c0] : !flow.dispatch.tensor<writeonly:1024x1024xf32> @@ -143,10 +143,10 @@ hal.executable.entry_point @dot_dispatch_0 attributes {interface = @io, ordinal = 0 : index} builtin.module { func @dot_dispatch_0() { - %cst = constant 0.000000e+00 : f32 - %c0 = constant 0 : index - %c1024 = constant 1024 : index - %c1 = constant 1 : index + %cst = arith.constant 0.000000e+00 : f32 + %c0 = arith.constant 0 : index + %c1024 = arith.constant 1024 : index + %c1 = arith.constant 1 : index %0 = hal.interface.binding.subspan @io::@ro0[%c0] : !flow.dispatch.tensor<readonly:1024x1024xf32> %1 = hal.interface.binding.subspan @io::@ro1[%c0] : !flow.dispatch.tensor<readonly:1024x1024xf32> %2 = hal.interface.binding.subspan @io::@wo2[%c0] : !flow.dispatch.tensor<writeonly:1024x1024xf32> @@ -174,8 +174,8 @@ %16 = linalg.fill(%cst, %15) : f32, tensor<?x?xf32> -> tensor<?x?xf32> %17 = linalg.generic #matmul_trait {__internal_linalg_transform__ = "workgroup"} ins(%8, %10 : tensor<?x1024xf32>, tensor<1024x?xf32>) outs(%16 : tensor<?x?xf32>) { ^bb(%a: f32, %b: f32, %c: f32) : - %d = mulf %a, %b: f32 - %e = addf %c, %d: f32 + %d = arith.mulf %a, %b: f32 + %e = arith.addf %c, %d: f32 linalg.yield %e : f32 } -> (tensor<?x?xf32>) flow.dispatch.tensor.store %17, %2, offsets = [%arg0, %arg1], sizes = [%11, %12], strides = [%c1, %c1] : tensor<?x?xf32> -> !flow.dispatch.tensor<writeonly:1024x1024xf32> @@ -211,11 +211,11 @@ hal.executable.entry_point @conv2d_dispatch_0 attributes {interface = @io, ordinal = 0 : index} builtin.module { func @conv2d_dispatch_0() { - %c0 = constant 0 : index - %cst = constant 0.000000e+00 : f32 - %c2 = constant 2 : index - %c3 = constant 3 : index - %c1 = constant 1 : index + %c0 = arith.constant 0 : index + %cst = arith.constant 0.000000e+00 : f32 + %c2 = arith.constant 2 : index + %c3 = arith.constant 3 : index + %c1 = arith.constant 1 : index %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : !flow.dispatch.tensor<readonly:1x4x4x2xf32> %1 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : !flow.dispatch.tensor<readonly:3x2x2x1xf32> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : !flow.dispatch.tensor<writeonly:1x2x3x1xf32> @@ -284,15 +284,15 @@ hal.executable.entry_point @add_dispatch_0 attributes {interface = @io, ordinal = 0 : index} builtin.module { func @add_dispatch_0() { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %0 = hal.interface.binding.subspan @io::@arg0[%c0] : !flow.dispatch.tensor<readonly:16xf32> %2 = hal.interface.binding.subspan @io::@ret0[%c0] : !flow.dispatch.tensor<writeonly:16xf32> %3 = linalg.init_tensor [16] : tensor<16xf32> %4 = flow.dispatch.tensor.load %0, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32> - %5 = constant dense<[1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0, 12.0, 13.0, 14.0, 15.0, 16.0]> : tensor<16xf32> + %5 = arith.constant dense<[1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0, 12.0, 13.0, 14.0, 15.0, 16.0]> : tensor<16xf32> %6 = linalg.generic {indexing_maps = [affine_map<(d0) -> (d0)>, affine_map<(d0) -> (d0)>, affine_map<(d0) -> (d0)>], iterator_types = ["parallel"]} ins(%4, %5 : tensor<16xf32>, tensor<16xf32>) outs(%3 : tensor<16xf32>) { ^bb0(%arg0: f32, %arg1: f32, %arg2: f32): // no predecessors - %7 = addf %arg0, %arg1 : f32 + %7 = arith.addf %arg0, %arg1 : f32 linalg.yield %7 : f32 } -> tensor<16xf32> flow.dispatch.tensor.store %6, %2, offsets=[], sizes=[], strides=[] : tensor<16xf32> -> !flow.dispatch.tensor<writeonly:16xf32> @@ -318,9 +318,9 @@ hal.executable.entry_point @reduction attributes {interface = @io, ordinal = 0 : index} builtin.module { func @reduction() { - %c0 = constant 0 : index - %cst = constant 0.000000e+00 : f32 - %c96 = constant 96 : index + %c0 = arith.constant 0 : index + %cst = arith.constant 0.000000e+00 : f32 + %c96 = arith.constant 96 : index %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : !flow.dispatch.tensor<readonly:14x14x96xf32> %1 = hal.interface.binding.subspan @io::@s0b1_xw_external[%c0] : !flow.dispatch.tensor<writeonly:96xf32> %workgroup_size_x = hal.interface.workgroup.size[0] : index @@ -337,7 +337,7 @@ %9 = linalg.fill(%cst, %8) : f32, tensor<?xf32> -> tensor<?xf32> %10 = linalg.generic {indexing_maps = [affine_map<(d0, d1, d2) -> (d1, d2, d0)>, affine_map<(d0, d1, d2) -> (d0)>], iterator_types = ["parallel", "reduction", "reduction"]} ins(%5 : tensor<14x14x?xf32>) outs(%9 : tensor<?xf32>) attrs = {__internal_linalg_transform__ = "workgroup"} { ^bb0(%arg1: f32, %arg2: f32): // no predecessors - %11 = addf %arg1, %arg2 : f32 + %11 = arith.addf %arg1, %arg2 : f32 linalg.yield %11 : f32 } -> tensor<?xf32> flow.dispatch.tensor.store %10, %1, offsets = [%arg0], sizes = [%6], strides = [1] : tensor<?xf32> -> !flow.dispatch.tensor<writeonly:96xf32> @@ -363,8 +363,8 @@ hal.executable.entry_point @vector_add_dispatch attributes {interface = @io, ordinal = 0 : index} builtin.module { builtin.func @vector_add_dispatch() { - %c0 = constant 0 : index - %c16384 = constant 16384 : index + %c0 = arith.constant 0 : index + %c16384 = arith.constant 16384 : index %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : !flow.dispatch.tensor<readonly:16384xf32> %1 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : !flow.dispatch.tensor<readonly:16384xf32> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : !flow.dispatch.tensor<writeonly:16384xf32> @@ -382,7 +382,7 @@ %10 = linalg.init_tensor [%9] : tensor<?xf32> %11 = linalg.generic {indexing_maps = [affine_map<(d0) -> (d0)>, affine_map<(d0) -> (d0)>, affine_map<(d0) -> (d0)>], iterator_types = ["parallel"]} ins(%6, %8 : tensor<?xf32>, tensor<?xf32>) outs(%10 : tensor<?xf32>) attrs = {__internal_linalg_transform__ = "workgroup"} { ^bb0(%arg1: f32, %arg2: f32, %arg3: f32): // no predecessors - %12 = addf %arg1, %arg2 : f32 + %12 = arith.addf %arg1, %arg2 : f32 linalg.yield %12 : f32 } -> tensor<?xf32> flow.dispatch.tensor.store %11, %2, offsets = [%arg0], sizes = [%9], strides = [1] : tensor<?xf32> -> !flow.dispatch.tensor<writeonly:16384xf32> @@ -416,9 +416,9 @@ hal.executable.entry_point @vector_reduction_dispatch attributes {interface = @io, ordinal = 0 : index} builtin.module { builtin.func @vector_reduction_dispatch() { - %c0 = constant 0 : index - %c16384 = constant 16384 : index - %cst = constant 1.000000e+00 : f32 + %c0 = arith.constant 0 : index + %c16384 = arith.constant 16384 : index + %cst = arith.constant 1.000000e+00 : f32 %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : !flow.dispatch.tensor<readonly:512x16384xf32> %1 = hal.interface.binding.subspan @io::@s0b1_xw_external[%c0] : !flow.dispatch.tensor<writeonly:16384xf32> %workgroup_size_x = hal.interface.workgroup.size[0] : index @@ -435,7 +435,7 @@ %9 = linalg.fill(%cst, %8) : f32, tensor<?xf32> -> tensor<?xf32> %10 = linalg.generic {indexing_maps = [#map3, #map4], iterator_types = ["parallel", "reduction"]} ins(%5 : tensor<512x?xf32>) outs(%9 : tensor<?xf32>) attrs = {__internal_linalg_transform__ = "workgroup"} { ^bb0(%arg1: f32, %arg2: f32): // no predecessors - %11 = addf %arg1, %arg2 : f32 + %11 = arith.addf %arg1, %arg2 : f32 linalg.yield %11 : f32 } -> tensor<?xf32> flow.dispatch.tensor.store %10, %1, offsets = [%arg0], sizes = [%6], strides = [1] : tensor<?xf32> -> !flow.dispatch.tensor<writeonly:16384xf32> @@ -452,5 +452,5 @@ // CHECK-LABEL: hal.executable public @vector_reduction_dispatch // CHECK: hal.executable.variant public @cuda -// CHECK-COUNT-4: llvm.fadd +// CHECK: llvm.fadd %{{.*}}, %{{.*}} : vector<4xf32> // CHECK: llvm.store %{{.*}} : !llvm.ptr<vector<4xf32>>
diff --git a/iree/compiler/Codegen/LLVMGPU/test/remove_loops.mlir b/iree/compiler/Codegen/LLVMGPU/test/remove_loops.mlir index 59bb295..011e2e9 100644 --- a/iree/compiler/Codegen/LLVMGPU/test/remove_loops.mlir +++ b/iree/compiler/Codegen/LLVMGPU/test/remove_loops.mlir
@@ -9,10 +9,10 @@ workgroup_size = [64: index, 1: index, 1:index]} builtin.module { builtin.func @dispatch_0() { - %c2 = constant 2 : index - %c256 = constant 256 : index - // CHECK: %[[C250:.+]] = constant 250 : index - %c250 = constant 250 : index + %c2 = arith.constant 2 : index + %c256 = arith.constant 256 : index + // CHECK: %[[C250:.+]] = arith.constant 250 : index + %c250 = arith.constant 250 : index %tidx = "gpu.thread_id"() {dimension = "x"} : () -> index %tidy = "gpu.thread_id"() {dimension = "y"} : () -> index // CHECK-NOT: scf.for
diff --git a/iree/compiler/Codegen/LLVMGPU/test/rocdl_pipeline_test.mlir b/iree/compiler/Codegen/LLVMGPU/test/rocdl_pipeline_test.mlir index 1f39a19..7578982 100644 --- a/iree/compiler/Codegen/LLVMGPU/test/rocdl_pipeline_test.mlir +++ b/iree/compiler/Codegen/LLVMGPU/test/rocdl_pipeline_test.mlir
@@ -12,7 +12,7 @@ hal.executable.entry_point @add_dispatch_0 attributes {interface = @io, ordinal = 0 : index, signature = (!flow.dispatch.tensor<readonly:16xf32>, !flow.dispatch.tensor<readonly:16xf32>, !flow.dispatch.tensor<writeonly:16xf32>) -> ()} builtin.module { func @add_dispatch_0() { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %0 = hal.interface.binding.subspan @io::@arg0[%c0] : !flow.dispatch.tensor<readonly:16xf32> %1 = hal.interface.binding.subspan @io::@arg1[%c0] : !flow.dispatch.tensor<readonly:16xf32> %2 = hal.interface.binding.subspan @io::@ret0[%c0] : !flow.dispatch.tensor<writeonly:16xf32> @@ -21,7 +21,7 @@ %5 = flow.dispatch.tensor.load %1, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32> %6 = linalg.generic {indexing_maps = [affine_map<(d0) -> (d0)>, affine_map<(d0) -> (d0)>, affine_map<(d0) -> (d0)>], iterator_types = ["parallel"]} ins(%4, %5 : tensor<16xf32>, tensor<16xf32>) outs(%3 : tensor<16xf32>) { ^bb0(%arg0: f32, %arg1: f32, %arg2: f32): // no predecessors - %7 = addf %arg0, %arg1 : f32 + %7 = arith.addf %arg0, %arg1 : f32 linalg.yield %7 : f32 } -> tensor<16xf32> flow.dispatch.tensor.store %6, %2, offsets=[], sizes=[], strides=[] : tensor<16xf32> -> !flow.dispatch.tensor<writeonly:16xf32> @@ -55,10 +55,10 @@ hal.executable.entry_point @dot_dispatch_0 attributes {interface = @io, ordinal = 0 : index, signature = (!flow.dispatch.tensor<readonly:1024x1024xf32>, !flow.dispatch.tensor<readonly:1024x1024xf32>, !flow.dispatch.tensor<writeonly:1024x1024xf32>) -> ()} builtin.module { func @dot_dispatch_0() { - %cst = constant 0.000000e+00 : f32 - %c0 = constant 0 : index - %c1024 = constant 1024 : index - %c1 = constant 1 : index + %cst = arith.constant 0.000000e+00 : f32 + %c0 = arith.constant 0 : index + %c1024 = arith.constant 1024 : index + %c1 = arith.constant 1 : index %0 = hal.interface.binding.subspan @io::@ro0[%c0] : !flow.dispatch.tensor<readonly:1024x1024xf32> %1 = hal.interface.binding.subspan @io::@ro1[%c0] : !flow.dispatch.tensor<readonly:1024x1024xf32> %2 = hal.interface.binding.subspan @io::@wo2[%c0] : !flow.dispatch.tensor<writeonly:1024x1024xf32>
diff --git a/iree/compiler/Codegen/LLVMGPU/test/vectorization.mlir b/iree/compiler/Codegen/LLVMGPU/test/vectorization.mlir index 6d1cd27..b822dc2 100644 --- a/iree/compiler/Codegen/LLVMGPU/test/vectorization.mlir +++ b/iree/compiler/Codegen/LLVMGPU/test/vectorization.mlir
@@ -9,7 +9,7 @@ ins(%arg0, %arg1 : memref<1x8x4xf32>, memref<1x4x8xf32>) outs(%arg2 : memref<1x4x8xf32>) attrs = {__internal_linalg_transform__ = "vectorize"} { ^bb0(%arg6: f32, %arg7: f32, %arg8: f32): // no predecessors - %19 = addf %arg6, %arg7 : f32 + %19 = arith.addf %arg6, %arg7 : f32 linalg.yield %19 : f32 } return
diff --git a/iree/compiler/Codegen/SPIRV/BUILD b/iree/compiler/Codegen/SPIRV/BUILD index acd92ca..d683287 100644 --- a/iree/compiler/Codegen/SPIRV/BUILD +++ b/iree/compiler/Codegen/SPIRV/BUILD
@@ -52,6 +52,7 @@ "@llvm-project//mlir:Affine", "@llvm-project//mlir:AffineToStandardTransforms", "@llvm-project//mlir:Analysis", + "@llvm-project//mlir:ArithmeticToSPIRV", "@llvm-project//mlir:DialectUtils", "@llvm-project//mlir:GPUDialect", "@llvm-project//mlir:GPUToSPIRV",
diff --git a/iree/compiler/Codegen/SPIRV/CMakeLists.txt b/iree/compiler/Codegen/SPIRV/CMakeLists.txt index d312ebb..532032a 100644 --- a/iree/compiler/Codegen/SPIRV/CMakeLists.txt +++ b/iree/compiler/Codegen/SPIRV/CMakeLists.txt
@@ -39,6 +39,7 @@ MLIRAffine MLIRAffineToStandard MLIRAnalysis + MLIRArithmeticToSPIRV MLIRGPUOps MLIRGPUToSPIRV MLIRGPUTransforms
diff --git a/iree/compiler/Codegen/SPIRV/ConvertToSPIRVPass.cpp b/iree/compiler/Codegen/SPIRV/ConvertToSPIRVPass.cpp index 5e6d2d1..a956c04 100644 --- a/iree/compiler/Codegen/SPIRV/ConvertToSPIRVPass.cpp +++ b/iree/compiler/Codegen/SPIRV/ConvertToSPIRVPass.cpp
@@ -26,6 +26,7 @@ #include "llvm/ADT/DenseSet.h" #include "llvm/ADT/STLExtras.h" #include "llvm/Support/FormatVariadic.h" +#include "mlir/Conversion/ArithmeticToSPIRV/ArithmeticToSPIRV.h" #include "mlir/Conversion/GPUToSPIRV/GPUToSPIRV.h" #include "mlir/Conversion/MathToSPIRV/MathToSPIRV.h" #include "mlir/Conversion/MemRefToSPIRV/MemRefToSPIRV.h" @@ -347,6 +348,7 @@ populateMemRefToSPIRVPatterns(typeConverter, patterns); // Pull in standard/math patterns to convert arithmetic ops and others. + arith::populateArithmeticToSPIRVPatterns(typeConverter, patterns); populateStandardToSPIRVPatterns(typeConverter, patterns); populateMathToSPIRVPatterns(typeConverter, patterns);
diff --git a/iree/compiler/Codegen/SPIRV/KernelConfig.cpp b/iree/compiler/Codegen/SPIRV/KernelConfig.cpp index 7371950..27bccf4 100644 --- a/iree/compiler/Codegen/SPIRV/KernelConfig.cpp +++ b/iree/compiler/Codegen/SPIRV/KernelConfig.cpp
@@ -50,7 +50,7 @@ expr = expr.ceilDiv(workgroupSizeX); Value numWorkgroupsX = linalg::applyMapToValues( b, loc, AffineMap::get(0, 3, expr), workload)[0]; - Value one = b.create<ConstantIndexOp>(loc, 1); + Value one = b.create<arith::ConstantIndexOp>(loc, 1); return std::array<Value, 3>{numWorkgroupsX, one, one}; }; return defineWorkgroupCountRegion(builder, funcOp, numWorkgroupsFn); @@ -73,7 +73,7 @@ int64_t count = outputShape[i] / workgroupTileSizes[i]; // This is meant for perfectly tilable cases. Double check that. assert(outputShape[i] % workgroupTileSizes[i] == 0 && count != 0); - xyz[2 - i] = b.create<ConstantIndexOp>(loc, count); + xyz[2 - i] = b.create<arith::ConstantIndexOp>(loc, count); } return xyz; };
diff --git a/iree/compiler/Codegen/SPIRV/SPIRVCopyToWorkgroupMemory.cpp b/iree/compiler/Codegen/SPIRV/SPIRVCopyToWorkgroupMemory.cpp index 0aa3596..bb11ffc 100644 --- a/iree/compiler/Codegen/SPIRV/SPIRVCopyToWorkgroupMemory.cpp +++ b/iree/compiler/Codegen/SPIRV/SPIRVCopyToWorkgroupMemory.cpp
@@ -57,12 +57,12 @@ linearized.procId = procInfo[0].procId; linearized.nprocs = procInfo[0].nprocs; for (unsigned i = 0; i < kNumGPUDims - 1; ++i) { - linearized.procId = - rewriter.create<MulIOp>(loc, linearized.procId, procInfo[i + 1].nprocs); - linearized.procId = - rewriter.create<AddIOp>(loc, linearized.procId, procInfo[i + 1].procId); - linearized.nprocs = - rewriter.create<MulIOp>(loc, linearized.nprocs, procInfo[i + 1].nprocs); + linearized.procId = rewriter.create<arith::MulIOp>(loc, linearized.procId, + procInfo[i + 1].nprocs); + linearized.procId = rewriter.create<arith::AddIOp>(loc, linearized.procId, + procInfo[i + 1].procId); + linearized.nprocs = rewriter.create<arith::MulIOp>(loc, linearized.nprocs, + procInfo[i + 1].nprocs); } return linearized; } @@ -91,11 +91,11 @@ auto lbs = pLoopOp.lowerBound(), ubs = pLoopOp.upperBound(), steps = pLoopOp.step(); for (unsigned i : llvm::seq<unsigned>(0, procInfo.size())) { - Value mappedLb = rewriter.create<AddIOp>( + Value mappedLb = rewriter.create<arith::AddIOp>( loc, lbs[i], - rewriter.create<MulIOp>(loc, steps[i], procInfo[i].procId)); + rewriter.create<arith::MulIOp>(loc, steps[i], procInfo[i].procId)); Value mappedStep = - rewriter.create<MulIOp>(loc, steps[i], procInfo[i].nprocs); + rewriter.create<arith::MulIOp>(loc, steps[i], procInfo[i].nprocs); forBounds.push_back({mappedLb, ubs[i], mappedStep}); permutation.push_back(i); }
diff --git a/iree/compiler/Codegen/SPIRV/SPIRVFoldGPUProcessorIDUses.cpp b/iree/compiler/Codegen/SPIRV/SPIRVFoldGPUProcessorIDUses.cpp index 52d000c..061ce5d 100644 --- a/iree/compiler/Codegen/SPIRV/SPIRVFoldGPUProcessorIDUses.cpp +++ b/iree/compiler/Codegen/SPIRV/SPIRVFoldGPUProcessorIDUses.cpp
@@ -240,8 +240,8 @@ LLVM_DEBUG(llvm::dbgs() << "map by subtracting expr '" << candidate << "': " << subMap << "\n"); if (llvm::all_of(subMap.getResults(), isNonNegativeCstExpr)) { - rewriter.replaceOpWithNewOp<ConstantIndexOp>(minOp, - candidate.getValue()); + rewriter.replaceOpWithNewOp<arith::ConstantIndexOp>( + minOp, candidate.getValue()); return success(); } }
diff --git a/iree/compiler/Codegen/SPIRV/SPIRVTileAndDistribute.cpp b/iree/compiler/Codegen/SPIRV/SPIRVTileAndDistribute.cpp index a006306..b151cef 100644 --- a/iree/compiler/Codegen/SPIRV/SPIRVTileAndDistribute.cpp +++ b/iree/compiler/Codegen/SPIRV/SPIRVTileAndDistribute.cpp
@@ -85,13 +85,13 @@ // subgroupID // = id.z * nsubgroups.y * nsubgroups.x + id.y * nsubgroups.x + id.x for (size_t i = 0, e = numSubgroups.size(); i != e; ++i) { - Value nprocs = builder.create<ConstantIndexOp>(loc, numSubgroups[i]); + Value nprocs = builder.create<arith::ConstantIndexOp>(loc, numSubgroups[i]); AffineExpr d0 = getAffineDimExpr(0, builder.getContext()); AffineExpr s0 = getAffineSymbolExpr(0, builder.getContext()); Value procId = makeComposedAffineApply(builder, loc, d0 % s0, {subgroupId, nprocs}); procInfo[e - i - 1] = linalg::ProcInfo{procId, nprocs}; - subgroupId = builder.create<SignedDivIOp>(loc, subgroupId, nprocs); + subgroupId = builder.create<arith::DivSIOp>(loc, subgroupId, nprocs); } return procInfo; } @@ -117,7 +117,7 @@ SmallVector<int64_t> tileSizes = getTileSizes(operation, 1); return llvm::to_vector<4>( llvm::map_range(tileSizes, [&](int64_t v) -> Value { - return builder.create<ConstantIndexOp>(operation->getLoc(), v); + return builder.create<arith::ConstantIndexOp>(operation->getLoc(), v); })); }; @@ -157,7 +157,8 @@ SmallVector<int64_t> tileSizes = getTileSizes(operation, 2); return llvm::to_vector<4>( llvm::map_range(tileSizes, [&](int64_t v) -> Value { - return builder.create<ConstantIndexOp>(operation->getLoc(), v); + return builder.create<arith::ConstantIndexOp>(operation->getLoc(), + v); })); }; @@ -239,7 +240,7 @@ SmallVector<int64_t> tileSizes = getTileSizes(op, 3); return llvm::to_vector<4>( llvm::map_range(tileSizes, [&](int64_t v) -> Value { - return builder.create<ConstantIndexOp>(op->getLoc(), v); + return builder.create<arith::ConstantIndexOp>(op->getLoc(), v); })); };
diff --git a/iree/compiler/Codegen/SPIRV/SPIRVVectorToCooperativeMatrix.cpp b/iree/compiler/Codegen/SPIRV/SPIRVVectorToCooperativeMatrix.cpp index 019a312..ff56f27 100644 --- a/iree/compiler/Codegen/SPIRV/SPIRVVectorToCooperativeMatrix.cpp +++ b/iree/compiler/Codegen/SPIRV/SPIRVVectorToCooperativeMatrix.cpp
@@ -85,7 +85,7 @@ auto contract = dyn_cast<vector::ContractionOp>(op); if (contract == nullptr) return; auto hasVectorDest = [](Operation *op) { - if (isa<ConstantOp, memref::AllocOp>(op)) return false; + if (isa<arith::ConstantOp, memref::AllocOp>(op)) return false; for (auto resultType : op->getResultTypes()) { if (resultType.isa<VectorType>()) return true; }
diff --git a/iree/compiler/Codegen/SPIRV/SPIRVVectorizeLoadStore.cpp b/iree/compiler/Codegen/SPIRV/SPIRVVectorizeLoadStore.cpp index 90e1164..6b42fde 100644 --- a/iree/compiler/Codegen/SPIRV/SPIRVVectorizeLoadStore.cpp +++ b/iree/compiler/Codegen/SPIRV/SPIRVVectorizeLoadStore.cpp
@@ -215,8 +215,9 @@ unsigned ratio = *vectorMemrefElemSize / *scalarMemrefElemSize; SmallVector<Value, 4> indices(adaptor.indices().begin(), adaptor.indices().end()); - indices.back() = rewriter.create<SignedDivIOp>( - loc, indices.back(), rewriter.create<ConstantIndexOp>(loc, ratio)); + indices.back() = rewriter.create<arith::DivSIOp>( + loc, indices.back(), + rewriter.create<arith::ConstantIndexOp>(loc, ratio)); // If the transfer_read can be replaced by a load after vectorization use // LoadOp and cast back to the original type. @@ -271,8 +272,9 @@ unsigned ratio = *vectorMemrefElemSize / *scalarMemrefElemSize; SmallVector<Value, 4> indices(adaptor.indices()); - indices.back() = rewriter.create<SignedDivIOp>( - loc, indices.back(), rewriter.create<ConstantIndexOp>(loc, ratio)); + indices.back() = rewriter.create<arith::DivSIOp>( + loc, indices.back(), + rewriter.create<arith::ConstantIndexOp>(loc, ratio)); // If the transfer_write can be replaced by a store after vectorization cast // the original value and use StoreOp. @@ -386,13 +388,14 @@ return failure(); Location loc = readOp.getLoc(); - Value newVector = rewriter.create<ConstantOp>( + Value newVector = rewriter.create<arith::ConstantOp>( loc, vectorType, rewriter.getZeroAttr(vectorType)); for (int i = 0; i < vectorType.getDimSize(0); ++i) { SmallVector<Value, 4> indices(readOp.indices().begin(), readOp.indices().end()); - indices.back() = rewriter.createOrFold<AddIOp>( - loc, indices.back(), rewriter.createOrFold<ConstantIndexOp>(loc, i)); + indices.back() = rewriter.createOrFold<arith::AddIOp>( + loc, indices.back(), + rewriter.createOrFold<arith::ConstantIndexOp>(loc, i)); Value scalar = rewriter.create<memref::LoadOp>(loc, scalarType, readOp.source(), indices); newVector = rewriter.create<vector::InsertOp>(loc, scalar, newVector, i); @@ -417,8 +420,9 @@ for (int i = 0; i < vectorType.getDimSize(0); ++i) { SmallVector<Value, 4> indices(writeOp.indices().begin(), writeOp.indices().end()); - indices.back() = rewriter.createOrFold<AddIOp>( - loc, indices.back(), rewriter.createOrFold<ConstantIndexOp>(loc, i)); + indices.back() = rewriter.createOrFold<arith::AddIOp>( + loc, indices.back(), + rewriter.createOrFold<arith::ConstantIndexOp>(loc, i)); Value scalar = rewriter.create<vector::ExtractOp>(loc, writeOp.vector(), i); rewriter.create<memref::StoreOp>(loc, scalar, writeOp.source(), indices);
diff --git a/iree/compiler/Codegen/SPIRV/Utils.cpp b/iree/compiler/Codegen/SPIRV/Utils.cpp index 917992c..8b36bca 100644 --- a/iree/compiler/Codegen/SPIRV/Utils.cpp +++ b/iree/compiler/Codegen/SPIRV/Utils.cpp
@@ -133,9 +133,10 @@ // ordering issue cause it assumes that the workgroup size has already been // set. If using affine_map can help, make sure that the workgroup size is set // before. - return {builder.create<AddIOp>( - loc, builder.create<MulIOp>(loc, blockId, blockDim), threadId), - builder.create<MulIOp>(loc, blockDim, gridDim)}; + return { + builder.create<arith::AddIOp>( + loc, builder.create<arith::MulIOp>(loc, blockId, blockDim), threadId), + builder.create<arith::MulIOp>(loc, blockDim, gridDim)}; } template <typename GPUIdOp, typename GPUCountOp> @@ -179,7 +180,7 @@ // Compute the number of iterations of each loops starting from the innermost. Location loc = pLoopOp.getLoc(); - Value totalNumIterations = rewriter.create<ConstantIndexOp>(loc, 1); + Value totalNumIterations = rewriter.create<arith::ConstantIndexOp>(loc, 1); // Track the "stride" of each loop, i.e. product of the total number of // iterations of the inner loops. @@ -190,17 +191,17 @@ auto steps = pLoopOp.step(); for (int i = numLoops - 1; i >= 0; --i) { Value lb = lbs[i], ub = ubs[i], step = steps[i]; - Value iterCount = rewriter.create<SignedDivIOp>( - loc, rewriter.create<SubIOp>(loc, ub, lb), step); + Value iterCount = rewriter.create<arith::DivSIOp>( + loc, rewriter.create<arith::SubIOp>(loc, ub, lb), step); iterationStride[i] = totalNumIterations; totalNumIterations = - rewriter.create<MulIOp>(loc, totalNumIterations, iterCount); + rewriter.create<arith::MulIOp>(loc, totalNumIterations, iterCount); } // Create the collapsed parallel loop op with lowerbound 0, step 1 and upper // bound being the totalNumIterations. - Value newLb = rewriter.create<ConstantIndexOp>(loc, 0); - Value newStep = rewriter.create<ConstantIndexOp>(loc, 1); + Value newLb = rewriter.create<arith::ConstantIndexOp>(loc, 0); + Value newStep = rewriter.create<arith::ConstantIndexOp>(loc, 1); scf::ParallelOp newPLoopOp = rewriter.create<scf::ParallelOp>(loc, newLb, totalNumIterations, newStep); @@ -216,14 +217,14 @@ BlockAndValueMapping map; for (int i : llvm::seq<int>(0, numLoops)) { Value iterNum = - rewriter.create<SignedDivIOp>(loc, loopIv, iterationStride[i]); + rewriter.create<arith::DivSIOp>(loc, loopIv, iterationStride[i]); AffineExpr d0, d1; bindDims(rewriter.getContext(), d0, d1); AffineExpr s0 = getAffineSymbolExpr(0, rewriter.getContext()); Value newIv = makeComposedAffineApply(rewriter, loc, d0 + d1 * s0, {lbs[i], iterNum, steps[i]}); map.map(pLoopBody.getArgument(i), newIv); - loopIv = rewriter.create<SignedRemIOp>(loc, loopIv, iterationStride[i]); + loopIv = rewriter.create<arith::RemSIOp>(loc, loopIv, iterationStride[i]); } for (Operation &op : pLoopBody.without_terminator()) { rewriter.clone(op, map); @@ -325,8 +326,9 @@ auto step = pLoopOp.step(); SmallVector<Value, 2> ivReplacements; for (unsigned i : llvm::seq<unsigned>(0, numLoops)) { - Value iterValue = rewriter.create<AddIOp>( - loc, lbs[i], rewriter.create<MulIOp>(loc, procInfo[i].procId, step[i])); + Value iterValue = rewriter.create<arith::AddIOp>( + loc, lbs[i], + rewriter.create<arith::MulIOp>(loc, procInfo[i].procId, step[i])); ivReplacements.push_back(iterValue); } Region &pLoopOpRegion = pLoopOp.getLoopBody(); @@ -336,9 +338,9 @@ Value cond = nullptr; auto ubs = pLoopOp.upperBound(); for (unsigned i : llvm::seq<unsigned>(0, numLoops)) { - Value cmp = rewriter.create<CmpIOp>(loc, CmpIPredicate::slt, - ivReplacements[i], ubs[i]); - cond = (cond ? rewriter.create<AndOp>(loc, cond, cmp) : cmp); + Value cmp = rewriter.create<arith::CmpIOp>(loc, arith::CmpIPredicate::slt, + ivReplacements[i], ubs[i]); + cond = (cond ? rewriter.create<arith::AndIOp>(loc, cond, cmp) : cmp); signatureConverter.remapInput(i, ivReplacements[i]); } rewriter.applySignatureConversion(&pLoopOpRegion, signatureConverter);
diff --git a/iree/compiler/Codegen/SPIRV/test/config_adreno_conv.mlir b/iree/compiler/Codegen/SPIRV/test/config_adreno_conv.mlir index 6d8dff3..1cfcf1a 100644 --- a/iree/compiler/Codegen/SPIRV/test/config_adreno_conv.mlir +++ b/iree/compiler/Codegen/SPIRV/test/config_adreno_conv.mlir
@@ -18,10 +18,10 @@ hal.executable.entry_point public @conv_112x112x512 attributes {interface = @io, ordinal = 0 : index} builtin.module { func @conv_112x112x512() { - %c0 = constant 0 : index - %c512 = constant 512 : index - %c112 = constant 112 : index - %cst = constant 0.000000e+00 : f32 + %c0 = arith.constant 0 : index + %c512 = arith.constant 512 : index + %c112 = arith.constant 112 : index + %cst = arith.constant 0.000000e+00 : f32 %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : !flow.dispatch.tensor<readonly:1x225x225x3xf32> %1 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : !flow.dispatch.tensor<readonly:3x3x3x512xf32> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : !flow.dispatch.tensor<writeonly:1x112x112x512xf32> @@ -78,9 +78,9 @@ // CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [256, 8, 1]} // CHECK-SAME: workgroup_size = [64 : index, 1 : index, 1 : index] // CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index): -// CHECK-NEXT: %[[C2:.+]] = constant 2 : index -// CHECK-NEXT: %[[C14:.+]] = constant 14 : index -// CHECK-NEXT: %[[C112:.+]] = constant 112 : index +// CHECK-NEXT: %[[C2:.+]] = arith.constant 2 : index +// CHECK-NEXT: %[[C14:.+]] = arith.constant 14 : index +// CHECK-NEXT: %[[C112:.+]] = arith.constant 112 : index // CHECK-NEXT: hal.return %[[C2]], %[[C14]], %[[C112]] // CHECK: func @conv_112x112x512() @@ -107,10 +107,10 @@ hal.executable.entry_point public @conv_112x112x32 attributes {interface = @io, ordinal = 0 : index} builtin.module { func @conv_112x112x32() { - %c0 = constant 0 : index - %c32 = constant 32 : index - %c112 = constant 112 : index - %cst = constant 0.000000e+00 : f32 + %c0 = arith.constant 0 : index + %c32 = arith.constant 32 : index + %c112 = arith.constant 112 : index + %cst = arith.constant 0.000000e+00 : f32 %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : !flow.dispatch.tensor<readonly:1x225x225x3xf32> %1 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : !flow.dispatch.tensor<readonly:3x3x3x32xf32> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : !flow.dispatch.tensor<writeonly:1x112x112x32xf32> @@ -167,9 +167,9 @@ // CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [32, 16, 4]} // CHECK-SAME: workgroup_size = [8 : index, 8 : index, 1 : index] // CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index): -// CHECK-NEXT: %[[C1:.+]] = constant 1 : index -// CHECK-NEXT: %[[C7:.+]] = constant 7 : index -// CHECK-NEXT: %[[C28:.+]] = constant 28 : index +// CHECK-NEXT: %[[C1:.+]] = arith.constant 1 : index +// CHECK-NEXT: %[[C7:.+]] = arith.constant 7 : index +// CHECK-NEXT: %[[C28:.+]] = arith.constant 28 : index // CHECK-NEXT: hal.return %[[C1]], %[[C7]], %[[C28]] // CHECK: func @conv_112x112x32() @@ -196,9 +196,9 @@ hal.executable.entry_point public @conv_16x16x16 attributes {interface = @io, ordinal = 0 : index} builtin.module { func @conv_16x16x16() { - %c0 = constant 0 : index - %c16 = constant 16 : index - %cst = constant 0.000000e+00 : f32 + %c0 = arith.constant 0 : index + %c16 = arith.constant 16 : index + %cst = arith.constant 0.000000e+00 : f32 %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : !flow.dispatch.tensor<readonly:1x33x33x3xf32> %1 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : !flow.dispatch.tensor<readonly:3x3x3x16xf32> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : !flow.dispatch.tensor<writeonly:1x16x16x16xf32> @@ -255,8 +255,8 @@ // CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [16, 8, 8]} // CHECK-SAME: workgroup_size = [4 : index, 4 : index, 4 : index] // CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index): -// CHECK-NEXT: %[[C1:.+]] = constant 1 : index -// CHECK-NEXT: %[[C2:.+]] = constant 2 : index +// CHECK-NEXT: %[[C1:.+]] = arith.constant 1 : index +// CHECK-NEXT: %[[C2:.+]] = arith.constant 2 : index // CHECK-NEXT: hal.return %[[C1]], %[[C2]], %[[C2]] // CHECK: func @conv_16x16x16() @@ -283,10 +283,10 @@ hal.executable.entry_point public @dwconv_28x28x144 attributes {interface = @io, ordinal = 0 : index} builtin.module { func @dwconv_28x28x144() { - %c0 = constant 0 : index - %c144 = constant 144 : index - %c28 = constant 28 : index - %cst = constant 0.000000e+00 : f32 + %c0 = arith.constant 0 : index + %c144 = arith.constant 144 : index + %c28 = arith.constant 28 : index + %cst = arith.constant 0.000000e+00 : f32 %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : !flow.dispatch.tensor<readonly:1x57x57x144xf32> %1 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : !flow.dispatch.tensor<readonly:3x3x144xf32> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : !flow.dispatch.tensor<writeonly:1x28x28x144xf32> @@ -344,8 +344,8 @@ // CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [16, 4, 4]} // CHECK-SAME: workgroup_size = [4 : index, 4 : index, 4 : index] // CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index): -// CHECK-NEXT: %[[C9:.+]] = constant 9 : index -// CHECK-NEXT: %[[C7:.+]] = constant 7 : index +// CHECK-NEXT: %[[C9:.+]] = arith.constant 9 : index +// CHECK-NEXT: %[[C7:.+]] = arith.constant 7 : index // CHECK-NEXT: hal.return %[[C9]], %[[C7]], %[[C7]] // CHECK: func @dwconv_28x28x144() @@ -372,10 +372,10 @@ hal.executable.entry_point public @dwconv_4x4x8 attributes {interface = @io, ordinal = 0 : index} builtin.module { func @dwconv_4x4x8() { - %c0 = constant 0 : index - %c8 = constant 8 : index - %c4 = constant 4 : index - %cst = constant 0.000000e+00 : f32 + %c0 = arith.constant 0 : index + %c8 = arith.constant 8 : index + %c4 = arith.constant 4 : index + %cst = arith.constant 0.000000e+00 : f32 %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : !flow.dispatch.tensor<readonly:1x9x9x8xf32> %1 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : !flow.dispatch.tensor<readonly:3x3x8xf32> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : !flow.dispatch.tensor<writeonly:1x4x4x8xf32> @@ -433,7 +433,7 @@ // CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [8, 4, 4]} // CHECK-SAME: workgroup_size = [2 : index, 4 : index, 4 : index] // CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index): -// CHECK-NEXT: %[[C1:.+]] = constant 1 : index +// CHECK-NEXT: %[[C1:.+]] = arith.constant 1 : index // CHECK-NEXT: hal.return %[[C1]], %[[C1]], %[[C1]] // CHECK: func @dwconv_4x4x8()
diff --git a/iree/compiler/Codegen/SPIRV/test/config_adreno_matmul.mlir b/iree/compiler/Codegen/SPIRV/test/config_adreno_matmul.mlir index aaa7218..8d5b022 100644 --- a/iree/compiler/Codegen/SPIRV/test/config_adreno_matmul.mlir +++ b/iree/compiler/Codegen/SPIRV/test/config_adreno_matmul.mlir
@@ -18,10 +18,10 @@ hal.executable.entry_point @matmul_1024x2048x512 attributes {interface = @io, ordinal = 0 : index} builtin.module { func @matmul_1024x2048x512() { - %c0 = constant 0 : index - %c2048 = constant 2048 : index - %c1024 = constant 1024 : index - %cst = constant 0.000000e+00 : f32 + %c0 = arith.constant 0 : index + %c2048 = arith.constant 2048 : index + %c1024 = arith.constant 1024 : index + %cst = arith.constant 0.000000e+00 : f32 %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : !flow.dispatch.tensor<readonly:1024x512xf32> %1 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : !flow.dispatch.tensor<readonly:512x2048xf32> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : !flow.dispatch.tensor<writeonly:1024x2048xf32> @@ -66,7 +66,7 @@ // CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [128, 32]} // CHECK-SAME: workgroup_size = [32 : index, 2 : index, 1 : index] // CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index): -// CHECK-NEXT: %[[ONE:.+]] = constant 1 : index +// CHECK-NEXT: %[[ONE:.+]] = arith.constant 1 : index // CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 128)>()[%[[X]]] // CHECK-NEXT: %[[Y_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 32)>()[%[[Y]]] // CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y_COUNT]], %[[ONE]] @@ -95,10 +95,10 @@ hal.executable.entry_point @matmul_3136x24x96 attributes {interface = @io, ordinal = 0 : index} builtin.module { func @matmul_3136x24x96() { - %c0 = constant 0 : index - %c24 = constant 24 : index - %c3136 = constant 3136 : index - %cst = constant 0.000000e+00 : f32 + %c0 = arith.constant 0 : index + %c24 = arith.constant 24 : index + %c3136 = arith.constant 3136 : index + %cst = arith.constant 0.000000e+00 : f32 %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : !flow.dispatch.tensor<readonly:3136x96xf32> %1 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : !flow.dispatch.tensor<readonly:96x24xf32> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : !flow.dispatch.tensor<writeonly:3136x24xf32> @@ -143,7 +143,7 @@ // CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [8, 448]} // CHECK-SAME: workgroup_size = [2 : index, 32 : index, 1 : index] // CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index): -// CHECK-NEXT: %[[ONE:.+]] = constant 1 : index +// CHECK-NEXT: %[[ONE:.+]] = arith.constant 1 : index // CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 8)>()[%[[X]]] // CHECK-NEXT: %[[Y_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 448)>()[%[[Y]]] // CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y_COUNT]], %[[ONE]] @@ -172,10 +172,10 @@ hal.executable.entry_point @matmul_196x64x192 attributes {interface = @io, ordinal = 0 : index} builtin.module { func @matmul_196x64x192() { - %c0 = constant 0 : index - %c64 = constant 64 : index - %c196 = constant 196 : index - %cst = constant 0.000000e+00 : f32 + %c0 = arith.constant 0 : index + %c64 = arith.constant 64 : index + %c196 = arith.constant 196 : index + %cst = arith.constant 0.000000e+00 : f32 %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : !flow.dispatch.tensor<readonly:196x192xf32> %1 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : !flow.dispatch.tensor<readonly:192x64xf32> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : !flow.dispatch.tensor<writeonly:196x64xf32> @@ -220,7 +220,7 @@ // CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [64, 28]} // CHECK-SAME: workgroup_size = [16 : index, 4 : index, 1 : index] // CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index): -// CHECK-NEXT: %[[ONE:.+]] = constant 1 : index +// CHECK-NEXT: %[[ONE:.+]] = arith.constant 1 : index // CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 64)>()[%[[X]]] // CHECK-NEXT: %[[Y_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 28)>()[%[[Y]]] // CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y_COUNT]], %[[ONE]] @@ -249,10 +249,10 @@ hal.executable.entry_point @matmul_12544x96x16 attributes {interface = @io, ordinal = 0 : index} builtin.module { func @matmul_12544x96x16() { - %c0 = constant 0 : index - %c96 = constant 96 : index - %c12544 = constant 12544 : index - %cst = constant 0.000000e+00 : f32 + %c0 = arith.constant 0 : index + %c96 = arith.constant 96 : index + %c12544 = arith.constant 12544 : index + %cst = arith.constant 0.000000e+00 : f32 %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : memref<12544x16xf32> %1 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : memref<16x96xf32> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : memref<12544x96xf32> @@ -292,7 +292,7 @@ // CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [32, 128]} // CHECK-SAME: workgroup_size = [8 : index, 8 : index, 1 : index] // CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index): -// CHECK-NEXT: %[[ONE:.+]] = constant 1 : index +// CHECK-NEXT: %[[ONE:.+]] = arith.constant 1 : index // CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 32)>()[%[[X]]] // CHECK-NEXT: %[[Y_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 128)>()[%[[Y]]] // CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y_COUNT]], %[[ONE]] @@ -321,10 +321,10 @@ hal.executable.entry_point @matmul_49x160x576 attributes {interface = @io, ordinal = 0 : index} builtin.module { func @matmul_49x160x576() { - %c0 = constant 0 : index - %c160 = constant 160 : index - %c49 = constant 49 : index - %cst = constant 0.000000e+00 : f32 + %c0 = arith.constant 0 : index + %c160 = arith.constant 160 : index + %c49 = arith.constant 49 : index + %cst = arith.constant 0.000000e+00 : f32 %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : !flow.dispatch.tensor<readonly:49x576xf32> %1 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : !flow.dispatch.tensor<readonly:576x160xf32> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : !flow.dispatch.tensor<writeonly:49x160xf32> @@ -369,7 +369,7 @@ // CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [32, 7]} // CHECK-SAME: workgroup_size = [8 : index, 1 : index, 1 : index] // CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index): -// CHECK-NEXT: %[[ONE:.+]] = constant 1 : index +// CHECK-NEXT: %[[ONE:.+]] = arith.constant 1 : index // CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 32)>()[%[[X]]] // CHECK-NEXT: %[[Y_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 7)>()[%[[Y]]] // CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y_COUNT]], %[[ONE]] @@ -398,10 +398,10 @@ hal.executable.entry_point @batch_matmul_4x384x384 attributes {interface = @io, ordinal = 0 : index} builtin.module { func @batch_matmul_4x384x384() { - %c0 = constant 0 : index - %c384 = constant 384 : index - %c4 = constant 4 : index - %cst = constant 0.000000e+00 : f32 + %c0 = arith.constant 0 : index + %c384 = arith.constant 384 : index + %c4 = arith.constant 4 : index + %cst = arith.constant 0.000000e+00 : f32 %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : !flow.dispatch.tensor<readonly:4x384x32xf32> %1 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : !flow.dispatch.tensor<readonly:4x32x384xf32> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : !flow.dispatch.tensor<writeonly:4x384x384xf32> @@ -485,10 +485,10 @@ hal.executable.entry_point @batch_matmul_4x8x8 attributes {interface = @io, ordinal = 0 : index} builtin.module { func @batch_matmul_4x8x8() { - %c0 = constant 0 : index - %c8 = constant 8 : index - %c4 = constant 4 : index - %cst = constant 0.000000e+00 : f32 + %c0 = arith.constant 0 : index + %c8 = arith.constant 8 : index + %c4 = arith.constant 4 : index + %cst = arith.constant 0.000000e+00 : f32 %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : !flow.dispatch.tensor<readonly:4x8x32xf32> %1 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : !flow.dispatch.tensor<readonly:4x32x8xf32> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : !flow.dispatch.tensor<writeonly:4x8x8xf32>
diff --git a/iree/compiler/Codegen/SPIRV/test/config_default_matmul.mlir b/iree/compiler/Codegen/SPIRV/test/config_default_matmul.mlir index 6cbdd91..fea465c 100644 --- a/iree/compiler/Codegen/SPIRV/test/config_default_matmul.mlir +++ b/iree/compiler/Codegen/SPIRV/test/config_default_matmul.mlir
@@ -18,11 +18,11 @@ hal.executable.entry_point public @batch_matmul_1x3x32 attributes {interface = @io, ordinal = 0 : index} builtin.module { func @batch_matmul_1x3x32() { - %c0 = constant 0 : index - %c32 = constant 32 : index - %c3 = constant 3 : index - %c1 = constant 1 : index - %cst = constant 0.000000e+00 : f32 + %c0 = arith.constant 0 : index + %c32 = arith.constant 32 : index + %c3 = arith.constant 3 : index + %c1 = arith.constant 1 : index + %cst = arith.constant 0.000000e+00 : f32 %0 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : !flow.dispatch.tensor<readonly:1x3x3xf32> %1 = hal.interface.binding.subspan @io::@s0b0_ro_constant[%c0] : !flow.dispatch.tensor<readonly:1x3x32xf32> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : !flow.dispatch.tensor<writeonly:1x3x32xf32> @@ -105,10 +105,10 @@ hal.executable.entry_point public @matmul_64x16 attributes {interface = @io, ordinal = 0 : index} builtin.module { func @matmul_64x16() { - %c0 = constant 0 : index - %c16 = constant 16 : index - %c64 = constant 64 : index - %c0_i32 = constant 0 : i32 + %c0 = arith.constant 0 : index + %c16 = arith.constant 16 : index + %c64 = arith.constant 64 : index + %c0_i32 = arith.constant 0 : i32 %0 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : !flow.dispatch.tensor<readonly:64x32xi8> %1 = hal.interface.binding.subspan @io::@s0b0_ro_constant[%c0] : !flow.dispatch.tensor<readonly:32x16xi8> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : !flow.dispatch.tensor<writeonly:64x16xi32> @@ -153,7 +153,7 @@ // CHECK-SAME: translation.info = {passPipeline = "SPIRVDistribute", workloadPerWorkgroup = [4, 1]} // CHECK-SAME: workgroup_size = [4 : index, 1 : index, 1 : index] // CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index): -// CHECK-NEXT: %[[ONE:.+]] = constant 1 : index +// CHECK-NEXT: %[[ONE:.+]] = arith.constant 1 : index // CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 4)>()[%[[X]]] // CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y]], %[[ONE]]
diff --git a/iree/compiler/Codegen/SPIRV/test/config_linalg_ext_ops.mlir b/iree/compiler/Codegen/SPIRV/test/config_linalg_ext_ops.mlir index 199975a..2cd4a8d 100644 --- a/iree/compiler/Codegen/SPIRV/test/config_linalg_ext_ops.mlir +++ b/iree/compiler/Codegen/SPIRV/test/config_linalg_ext_ops.mlir
@@ -14,12 +14,12 @@ hal.executable.entry_point @static_1d_sort attributes {interface = @io, ordinal = 0 : index} builtin.module { builtin.func @static_1d_sort() { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %0 = hal.interface.binding.subspan @io::@s0b0_rw_external[%c0] : !flow.dispatch.tensor<readwrite:1000xi32> %1 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readwrite:1000xi32> -> tensor<1000xi32> %2 = linalg_ext.sort dimension(0) {__internal_linalg_transform__ = "workgroup"} outs(%1 : tensor<1000xi32>) { ^bb0(%arg0: i32, %arg1: i32): // no predecessors - %3 = cmpi slt, %arg0, %arg1 : i32 + %3 = arith.cmpi slt, %arg0, %arg1 : i32 linalg_ext.yield %3 : i1 } -> tensor<1000xi32> flow.dispatch.tensor.store %2, %0, offsets = [], sizes = [], strides = [] : tensor<1000xi32> -> !flow.dispatch.tensor<readwrite:1000xi32> @@ -38,7 +38,7 @@ // CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize"} // CHECK-SAME: workgroup_size = [1 : index, 1 : index, 1 : index] // CHECK-NEXT: ^{{.+}}(%{{.+}}: index, %{{.+}}: index, %{{.+}}: index): -// CHECK-NEXT: %[[ONE:.+]] = constant 1 : index +// CHECK-NEXT: %[[ONE:.+]] = arith.constant 1 : index // CHECK-NEXT: hal.return %[[ONE]], %[[ONE]], %[[ONE]] // CHECK: func @static_1d_sort() @@ -62,9 +62,9 @@ hal.executable.entry_point @static_3d_sort attributes {interface = @io, ordinal = 0 : index} builtin.module { builtin.func @static_3d_sort() { - %c64 = constant 64 : index - %c128 = constant 128 : index - %c0 = constant 0 : index + %c64 = arith.constant 64 : index + %c128 = arith.constant 128 : index + %c0 = arith.constant 0 : index %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : memref<64x32x128xi32> %1 = hal.interface.binding.subspan @io::@s0b1_xw_external[%c0] : memref<64x32x128xi32> %workgroup_size_x = hal.interface.workgroup.size[0] : index @@ -87,7 +87,7 @@ linalg.copy(%9, %10) : memref<?x?x?xi32>, memref<?x32x?xi32, affine_map<(d0, d1, d2)[s0] -> (d0 * 4096 + s0 + d1 * 128 + d2)>> linalg_ext.sort dimension(1) {__internal_linalg_transform__ = "workgroup"} outs(%10 : memref<?x32x?xi32, affine_map<(d0, d1, d2)[s0] -> (d0 * 4096 + s0 + d1 * 128 + d2)>>) { ^bb0(%arg2: i32, %arg3: i32): // no predecessors - %11 = cmpi slt, %arg2, %arg3 : i32 + %11 = arith.cmpi slt, %arg2, %arg3 : i32 linalg_ext.yield %11 : i1 } } @@ -102,7 +102,7 @@ // CHECK-SAME: translation.info = {passPipeline = "SPIRVDistribute", workloadPerWorkgroup = [16, 1]} // CHECK-SAME: workgroup_size = [16 : index, 1 : index, 1 : index] // CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index): -// CHECK-NEXT: %[[ONE:.+]] = constant 1 : index +// CHECK-NEXT: %[[ONE:.+]] = arith.constant 1 : index // CHECK-NEXT: %[[DIV:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 16)>()[%[[X]]] // CHECK-NEXT: hal.return %[[DIV]], %[[Y]], %[[ONE]] @@ -127,10 +127,10 @@ hal.executable.entry_point @static_1d_fft_stage2 attributes {interface = @io, ordinal = 0 : index} builtin.module { builtin.func @static_1d_fft_stage2() { - %c0 = constant 0 : index - %c2 = constant 2 : index - %cst = constant dense<[1.000000e+00, 6.12323426E-17]> : tensor<2xf32> - %cst_0 = constant dense<[-0.000000e+00, -1.000000e+00]> : tensor<2xf32> + %c0 = arith.constant 0 : index + %c2 = arith.constant 2 : index + %cst = arith.constant dense<[1.000000e+00, 6.12323426E-17]> : tensor<2xf32> + %cst_0 = arith.constant dense<[-0.000000e+00, -1.000000e+00]> : tensor<2xf32> %0 = hal.interface.binding.subspan @io::@s0b0_rw_external[%c0] : !flow.dispatch.tensor<readwrite:32xf32> %1 = hal.interface.binding.subspan @io::@s0b1_rw_external[%c0] : !flow.dispatch.tensor<readwrite:32xf32> %2 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readwrite:32xf32> -> tensor<32xf32> @@ -149,7 +149,7 @@ // CHECK-SAME: workloadPerWorkgroup = [4]} // CHECK-SAME: workgroup_size = [16 : index, 1 : index, 1 : index] // CHECK-NEXT: ^{{.+}}(%[[ARG0:.+]]: index, %{{.+}}: index, %{{.+}}: index): -// CHECK-NEXT: %[[ONE:.+]] = constant 1 : index +// CHECK-NEXT: %[[ONE:.+]] = arith.constant 1 : index // CHECK-NEXT: %[[T:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 4)>()[%[[ARG0]]] // CHECK-NEXT: hal.return %[[T]], %[[ONE]], %[[ONE]] @@ -175,13 +175,13 @@ hal.executable.entry_point @static_3d_fft_stage3 attributes {interface = @io, ordinal = 0 : index} builtin.module { builtin.func @static_3d_fft_stage3() { - %c0 = constant 0 : index - %c3 = constant 3 : index - %c64 = constant 64 : index - %c128 = constant 128 : index - %c32 = constant 32 : index - %cst = constant dense<[1.000000e+00, 0.707106769, 6.12323426E-17, -0.707106769]> : tensor<4xf32> - %cst_0 = constant dense<[-0.000000e+00, -0.707106769, -1.000000e+00, -0.707106769]> : tensor<4xf32> + %c0 = arith.constant 0 : index + %c3 = arith.constant 3 : index + %c64 = arith.constant 64 : index + %c128 = arith.constant 128 : index + %c32 = arith.constant 32 : index + %cst = arith.constant dense<[1.000000e+00, 0.707106769, 6.12323426E-17, -0.707106769]> : tensor<4xf32> + %cst_0 = arith.constant dense<[-0.000000e+00, -0.707106769, -1.000000e+00, -0.707106769]> : tensor<4xf32> %0 = memref.buffer_cast %cst_0 : memref<4xf32> %1 = memref.buffer_cast %cst : memref<4xf32> %2 = hal.interface.binding.subspan @io::@s0b0_rw_external[%c0] : memref<64x128x32xf32>
diff --git a/iree/compiler/Codegen/SPIRV/test/config_linalg_ops.mlir b/iree/compiler/Codegen/SPIRV/test/config_linalg_ops.mlir index 9b97a1a..f6304bc 100644 --- a/iree/compiler/Codegen/SPIRV/test/config_linalg_ops.mlir +++ b/iree/compiler/Codegen/SPIRV/test/config_linalg_ops.mlir
@@ -11,7 +11,7 @@ hal.executable.entry_point @tensor_insert_slice attributes {interface = @io, ordinal = 0 : index} builtin.module { builtin.func @tensor_insert_slice() { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : !flow.dispatch.tensor<readonly:?x?xi32> %1 = hal.interface.load.constant offset = 0 : index %2 = hal.interface.load.constant offset = 1 : index @@ -52,7 +52,7 @@ // CHECK-SAME: translation.info = {passPipeline = "SPIRVDistribute", workloadPerWorkgroup = [64, 1]} // CHECK-NEXT: %[[ARG0:[a-zA-Z0-9_]+]]: index // CHECK-SAME: %[[ARG1:[a-zA-Z0-9_]+]]: index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK-DAG: %[[NWGSX:.+]] = affine.apply #[[MAP]]()[%[[ARG0]]] // CHECK: hal.return %[[NWGSX]], %[[ARG1]], %[[C1]] @@ -69,7 +69,7 @@ hal.executable.entry_point @tensor_insert_slice attributes {interface = @io, ordinal = 0 : index} builtin.module { builtin.func @tensor_insert_slice() { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %d0 = hal.interface.load.constant offset = 0 : index %d1 = hal.interface.load.constant offset = 1 : index %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : memref<?x?xi32>{%d0, %d1} @@ -106,7 +106,7 @@ // CHECK-SAME: translation.info = {passPipeline = "SPIRVDistribute", workloadPerWorkgroup = [16, 1]} // CHECK-NEXT: %[[ARG0:[a-zA-Z0-9_]+]]: index // CHECK-SAME: %[[ARG1:[a-zA-Z0-9_]+]]: index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK-DAG: %[[NWGSX:.+]] = affine.apply #[[MAP]]()[%[[ARG0]]] // CHECK: hal.return %[[NWGSX]], %[[ARG1]], %[[C1]] // CHECK: linalg.copy
diff --git a/iree/compiler/Codegen/SPIRV/test/config_mali_conv.mlir b/iree/compiler/Codegen/SPIRV/test/config_mali_conv.mlir index cc1ba0a..9dfd104 100644 --- a/iree/compiler/Codegen/SPIRV/test/config_mali_conv.mlir +++ b/iree/compiler/Codegen/SPIRV/test/config_mali_conv.mlir
@@ -18,10 +18,10 @@ hal.executable.entry_point public @conv_112x112x512 attributes {interface = @io, ordinal = 0 : index} builtin.module { func @conv_112x112x512() { - %c0 = constant 0 : index - %c512 = constant 512 : index - %c112 = constant 112 : index - %cst = constant 0.000000e+00 : f32 + %c0 = arith.constant 0 : index + %c512 = arith.constant 512 : index + %c112 = arith.constant 112 : index + %cst = arith.constant 0.000000e+00 : f32 %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : !flow.dispatch.tensor<readonly:1x225x225x3xf32> %1 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : !flow.dispatch.tensor<readonly:3x3x3x512xf32> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : !flow.dispatch.tensor<writeonly:1x112x112x512xf32> @@ -78,9 +78,9 @@ // CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [64, 4, 1]} // CHECK-SAME: workgroup_size = [16 : index, 1 : index, 1 : index] // CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index): -// CHECK-NEXT: %[[C8:.+]] = constant 8 : index -// CHECK-NEXT: %[[C28:.+]] = constant 28 : index -// CHECK-NEXT: %[[C112:.+]] = constant 112 : index +// CHECK-NEXT: %[[C8:.+]] = arith.constant 8 : index +// CHECK-NEXT: %[[C28:.+]] = arith.constant 28 : index +// CHECK-NEXT: %[[C112:.+]] = arith.constant 112 : index // CHECK-NEXT: hal.return %[[C8]], %[[C28]], %[[C112]] // CHECK: func @conv_112x112x512() @@ -107,10 +107,10 @@ hal.executable.entry_point public @conv_112x112x32 attributes {interface = @io, ordinal = 0 : index} builtin.module { func @conv_112x112x32() { - %c0 = constant 0 : index - %c32 = constant 32 : index - %c112 = constant 112 : index - %cst = constant 0.000000e+00 : f32 + %c0 = arith.constant 0 : index + %c32 = arith.constant 32 : index + %c112 = arith.constant 112 : index + %cst = arith.constant 0.000000e+00 : f32 %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : !flow.dispatch.tensor<readonly:1x225x225x3xf32> %1 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : !flow.dispatch.tensor<readonly:3x3x3x32xf32> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : !flow.dispatch.tensor<writeonly:1x112x112x32xf32> @@ -167,9 +167,9 @@ // CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [32, 8, 1]} // CHECK-SAME: workgroup_size = [8 : index, 2 : index, 1 : index] // CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index): -// CHECK-NEXT: %[[C1:.+]] = constant 1 : index -// CHECK-NEXT: %[[C14:.+]] = constant 14 : index -// CHECK-NEXT: %[[C112:.+]] = constant 112 : index +// CHECK-NEXT: %[[C1:.+]] = arith.constant 1 : index +// CHECK-NEXT: %[[C14:.+]] = arith.constant 14 : index +// CHECK-NEXT: %[[C112:.+]] = arith.constant 112 : index // CHECK-NEXT: hal.return %[[C1]], %[[C14]], %[[C112]] // CHECK: func @conv_112x112x32() @@ -196,9 +196,9 @@ hal.executable.entry_point public @conv_16x16x16 attributes {interface = @io, ordinal = 0 : index} builtin.module { func @conv_16x16x16() { - %c0 = constant 0 : index - %c16 = constant 16 : index - %cst = constant 0.000000e+00 : f32 + %c0 = arith.constant 0 : index + %c16 = arith.constant 16 : index + %cst = arith.constant 0.000000e+00 : f32 %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : !flow.dispatch.tensor<readonly:1x33x33x3xf32> %1 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : !flow.dispatch.tensor<readonly:3x3x3x16xf32> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : !flow.dispatch.tensor<writeonly:1x16x16x16xf32> @@ -255,8 +255,8 @@ // CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [16, 4, 4]} // CHECK-SAME: workgroup_size = [4 : index, 2 : index, 2 : index] // CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index): -// CHECK-NEXT: %[[C1:.+]] = constant 1 : index -// CHECK-NEXT: %[[C4:.+]] = constant 4 : index +// CHECK-NEXT: %[[C1:.+]] = arith.constant 1 : index +// CHECK-NEXT: %[[C4:.+]] = arith.constant 4 : index // CHECK-NEXT: hal.return %[[C1]], %[[C4]], %[[C4]] // CHECK: func @conv_16x16x16() @@ -283,10 +283,10 @@ hal.executable.entry_point public @dwconv_28x28x144 attributes {interface = @io, ordinal = 0 : index} builtin.module { func @dwconv_28x28x144() { - %c0 = constant 0 : index - %c144 = constant 144 : index - %c28 = constant 28 : index - %cst = constant 0.000000e+00 : f32 + %c0 = arith.constant 0 : index + %c144 = arith.constant 144 : index + %c28 = arith.constant 28 : index + %cst = arith.constant 0.000000e+00 : f32 %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : !flow.dispatch.tensor<readonly:1x57x57x144xf32> %1 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : !flow.dispatch.tensor<readonly:3x3x144xf32> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : !flow.dispatch.tensor<writeonly:1x28x28x144xf32> @@ -344,8 +344,8 @@ // CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [16, 4, 4]} // CHECK-SAME: workgroup_size = [4 : index, 2 : index, 2 : index] // CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index): -// CHECK-NEXT: %[[C9:.+]] = constant 9 : index -// CHECK-NEXT: %[[C7:.+]] = constant 7 : index +// CHECK-NEXT: %[[C9:.+]] = arith.constant 9 : index +// CHECK-NEXT: %[[C7:.+]] = arith.constant 7 : index // CHECK-NEXT: hal.return %[[C9]], %[[C7]], %[[C7]] // CHECK: func @dwconv_28x28x144() @@ -372,11 +372,11 @@ hal.executable.entry_point public @dwconv_1x2x8 attributes {interface = @io, ordinal = 0 : index} builtin.module { func @dwconv_1x2x8() { - %c0 = constant 0 : index - %c8 = constant 8 : index - %c2 = constant 2 : index - %c1 = constant 1 : index - %cst = constant 0.000000e+00 : f32 + %c0 = arith.constant 0 : index + %c8 = arith.constant 8 : index + %c2 = arith.constant 2 : index + %c1 = arith.constant 1 : index + %cst = arith.constant 0.000000e+00 : f32 %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : !flow.dispatch.tensor<readonly:1x3x5x8xf32> %1 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : !flow.dispatch.tensor<readonly:3x3x8xf32> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : !flow.dispatch.tensor<writeonly:1x1x2x8xf32> @@ -434,7 +434,7 @@ // CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [8, 2, 1]} // CHECK-SAME: workgroup_size = [2 : index, 2 : index, 1 : index] // CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index): -// CHECK-NEXT: %[[C1:.+]] = constant 1 : index +// CHECK-NEXT: %[[C1:.+]] = arith.constant 1 : index // CHECK-NEXT: hal.return %[[C1]], %[[C1]], %[[C1]] // CHECK: func @dwconv_1x2x8()
diff --git a/iree/compiler/Codegen/SPIRV/test/config_mali_matmul.mlir b/iree/compiler/Codegen/SPIRV/test/config_mali_matmul.mlir index a6e99a1..efa37b4 100644 --- a/iree/compiler/Codegen/SPIRV/test/config_mali_matmul.mlir +++ b/iree/compiler/Codegen/SPIRV/test/config_mali_matmul.mlir
@@ -18,10 +18,10 @@ hal.executable.entry_point @matmul_1024x2048x512 attributes {interface = @io, ordinal = 0 : index} builtin.module { func @matmul_1024x2048x512() { - %c0 = constant 0 : index - %c2048 = constant 2048 : index - %c1024 = constant 1024 : index - %cst = constant 0.000000e+00 : f32 + %c0 = arith.constant 0 : index + %c2048 = arith.constant 2048 : index + %c1024 = arith.constant 1024 : index + %cst = arith.constant 0.000000e+00 : f32 %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : !flow.dispatch.tensor<readonly:1024x512xf32> %1 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : !flow.dispatch.tensor<readonly:512x2048xf32> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : !flow.dispatch.tensor<writeonly:1024x2048xf32> @@ -66,7 +66,7 @@ // CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [32, 8]} // CHECK-SAME: workgroup_size = [8 : index, 2 : index, 1 : index] // CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index): -// CHECK-NEXT: %[[ONE:.+]] = constant 1 : index +// CHECK-NEXT: %[[ONE:.+]] = arith.constant 1 : index // CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 32)>()[%[[X]]] // CHECK-NEXT: %[[Y_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 8)>()[%[[Y]]] // CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y_COUNT]], %[[ONE]] @@ -95,10 +95,10 @@ hal.executable.entry_point @matmul_3136x24x96 attributes {interface = @io, ordinal = 0 : index} builtin.module { func @matmul_3136x24x96() { - %c0 = constant 0 : index - %c24 = constant 24 : index - %c3136 = constant 3136 : index - %cst = constant 0.000000e+00 : f32 + %c0 = arith.constant 0 : index + %c24 = arith.constant 24 : index + %c3136 = arith.constant 3136 : index + %cst = arith.constant 0.000000e+00 : f32 %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : !flow.dispatch.tensor<readonly:3136x96xf32> %1 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : !flow.dispatch.tensor<readonly:96x24xf32> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : !flow.dispatch.tensor<writeonly:3136x24xf32> @@ -143,7 +143,7 @@ // CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [8, 32]} // CHECK-SAME: workgroup_size = [2 : index, 8 : index, 1 : index] // CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index): -// CHECK-NEXT: %[[ONE:.+]] = constant 1 : index +// CHECK-NEXT: %[[ONE:.+]] = arith.constant 1 : index // CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 8)>()[%[[X]]] // CHECK-NEXT: %[[Y_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 32)>()[%[[Y]]] // CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y_COUNT]], %[[ONE]] @@ -172,10 +172,10 @@ hal.executable.entry_point @matmul_196x64x192 attributes {interface = @io, ordinal = 0 : index} builtin.module { func @matmul_196x64x192() { - %c0 = constant 0 : index - %c64 = constant 64 : index - %c196 = constant 196 : index - %cst = constant 0.000000e+00 : f32 + %c0 = arith.constant 0 : index + %c64 = arith.constant 64 : index + %c196 = arith.constant 196 : index + %cst = arith.constant 0.000000e+00 : f32 %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : !flow.dispatch.tensor<readonly:196x192xf32> %1 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : !flow.dispatch.tensor<readonly:192x64xf32> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : !flow.dispatch.tensor<writeonly:196x64xf32> @@ -220,7 +220,7 @@ // CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [32, 4]} // CHECK-SAME: workgroup_size = [8 : index, 2 : index, 1 : index] // CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index): -// CHECK-NEXT: %[[ONE:.+]] = constant 1 : index +// CHECK-NEXT: %[[ONE:.+]] = arith.constant 1 : index // CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 32)>()[%[[X]]] // CHECK-NEXT: %[[Y_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 4)>()[%[[Y]]] // CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y_COUNT]], %[[ONE]] @@ -249,10 +249,10 @@ hal.executable.entry_point @matmul_12544x96x16 attributes {interface = @io, ordinal = 0 : index} builtin.module { func @matmul_12544x96x16() { - %c0 = constant 0 : index - %c96 = constant 96 : index - %c12544 = constant 12544 : index - %cst = constant 0.000000e+00 : f32 + %c0 = arith.constant 0 : index + %c96 = arith.constant 96 : index + %c12544 = arith.constant 12544 : index + %cst = arith.constant 0.000000e+00 : f32 %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : memref<12544x16xf32> %1 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : memref<16x96xf32> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : memref<12544x96xf32> @@ -292,7 +292,7 @@ // CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [32, 8]} // CHECK-SAME: workgroup_size = [8 : index, 2 : index, 1 : index] // CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index): -// CHECK-NEXT: %[[ONE:.+]] = constant 1 : index +// CHECK-NEXT: %[[ONE:.+]] = arith.constant 1 : index // CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 32)>()[%[[X]]] // CHECK-NEXT: %[[Y_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 8)>()[%[[Y]]] // CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y_COUNT]], %[[ONE]] @@ -321,10 +321,10 @@ hal.executable.entry_point @matmul_49x160x576 attributes {interface = @io, ordinal = 0 : index} builtin.module { func @matmul_49x160x576() { - %c0 = constant 0 : index - %c160 = constant 160 : index - %c49 = constant 49 : index - %cst = constant 0.000000e+00 : f32 + %c0 = arith.constant 0 : index + %c160 = arith.constant 160 : index + %c49 = arith.constant 49 : index + %cst = arith.constant 0.000000e+00 : f32 %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : !flow.dispatch.tensor<readonly:49x576xf32> %1 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : !flow.dispatch.tensor<readonly:576x160xf32> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : !flow.dispatch.tensor<writeonly:49x160xf32> @@ -369,7 +369,7 @@ // CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [32, 1]} // CHECK-SAME: workgroup_size = [8 : index, 1 : index, 1 : index] // CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index): -// CHECK-NEXT: %[[ONE:.+]] = constant 1 : index +// CHECK-NEXT: %[[ONE:.+]] = arith.constant 1 : index // CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 32)>()[%[[X]]] // CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y]], %[[ONE]] @@ -397,10 +397,10 @@ hal.executable.entry_point @batch_matmul_4x384x384 attributes {interface = @io, ordinal = 0 : index} builtin.module { func @batch_matmul_4x384x384() { - %c0 = constant 0 : index - %c384 = constant 384 : index - %c4 = constant 4 : index - %cst = constant 0.000000e+00 : f32 + %c0 = arith.constant 0 : index + %c384 = arith.constant 384 : index + %c4 = arith.constant 4 : index + %cst = arith.constant 0.000000e+00 : f32 %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : !flow.dispatch.tensor<readonly:4x384x32xf32> %1 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : !flow.dispatch.tensor<readonly:4x32x384xf32> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : !flow.dispatch.tensor<writeonly:4x384x384xf32> @@ -484,11 +484,11 @@ hal.executable.entry_point @batch_matmul_4x2x8 attributes {interface = @io, ordinal = 0 : index} builtin.module { func @batch_matmul_4x2x8() { - %c0 = constant 0 : index - %c8 = constant 8 : index - %c2 = constant 2 : index - %c4 = constant 4 : index - %cst = constant 0.000000e+00 : f32 + %c0 = arith.constant 0 : index + %c8 = arith.constant 8 : index + %c2 = arith.constant 2 : index + %c4 = arith.constant 4 : index + %cst = arith.constant 0.000000e+00 : f32 %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : !flow.dispatch.tensor<readonly:4x2x32xf32> %1 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : !flow.dispatch.tensor<readonly:4x32x8xf32> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : !flow.dispatch.tensor<writeonly:4x2x8xf32>
diff --git a/iree/compiler/Codegen/SPIRV/test/convert_to_spirv.mlir b/iree/compiler/Codegen/SPIRV/test/convert_to_spirv.mlir index 64e7dc4..4d9ef5d 100644 --- a/iree/compiler/Codegen/SPIRV/test/convert_to_spirv.mlir +++ b/iree/compiler/Codegen/SPIRV/test/convert_to_spirv.mlir
@@ -55,7 +55,7 @@ // CHECK: spv.GlobalVariable @[[RET0:.+]] bind(3, 4) : !spv.ptr<!spv.struct<(!spv.array<16 x f32, stride=4> [0])>, StorageBuffer> // CHECK: spv.func @resource_bindings_in_same_entry_func() func @resource_bindings_in_same_entry_func() { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index // Same type // CHECK: spv.mlir.addressof @[[ARG0]] @@ -120,7 +120,7 @@ func @resource_bindings_in_entry_func1() { // CHECK: spv.mlir.addressof @[[FUNC1_ARG]] // CHECK: spv.mlir.addressof @[[FUNC1_RET]] - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<4x4xf32> %1 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<4xvector<4xf32>> @@ -134,7 +134,7 @@ func @resource_bindings_in_entry_func2() { // CHECK: spv.mlir.addressof @[[FUNC2_ARG]] // CHECK: spv.mlir.addressof @[[FUNC2_RET]] - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<4x4xf32> // Same type as previous function %1 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<4x4xf32> // Different type as previous function @@ -168,7 +168,7 @@ } builtin.module { func @interface_binding() { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<8x5xf32> %1 = hal.interface.binding.subspan @io::@arg1[%c0] : memref<5xf32> %2 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<8x5xf32>
diff --git a/iree/compiler/Codegen/SPIRV/test/distribute_to_global_id.mlir b/iree/compiler/Codegen/SPIRV/test/distribute_to_global_id.mlir index bafed53..c7d9c88 100644 --- a/iree/compiler/Codegen/SPIRV/test/distribute_to_global_id.mlir +++ b/iree/compiler/Codegen/SPIRV/test/distribute_to_global_id.mlir
@@ -11,7 +11,7 @@ hal.executable.entry_point @parallel_4D attributes {interface = @io, ordinal = 0 : index} builtin.module { func @parallel_4D() { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %dim0 = hal.interface.load.constant offset = 0 : index %dim1 = hal.interface.load.constant offset = 1 : index %dim2 = hal.interface.load.constant offset = 2 : index @@ -25,7 +25,7 @@ ins(%arg0, %arg1 : memref<?x?x?x?xf32>, memref<?x?x?x?xf32>) outs(%arg2 : memref<?x?x?x?xf32>) { ^bb0(%arg3 : f32, %arg4 : f32, %arg5 : f32): - %0 = addf %arg3, %arg4 : f32 + %0 = arith.addf %arg3, %arg4 : f32 linalg.yield %0 : f32 } return @@ -42,30 +42,30 @@ } } // CHECK-LABEL: func @parallel_4D -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index -// CHECK-DAG: %[[C2:.+]] = constant 2 : index -// CHECK-DAG: %[[C3:.+]] = constant 3 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index +// CHECK-DAG: %[[C2:.+]] = arith.constant 2 : index +// CHECK-DAG: %[[C3:.+]] = arith.constant 3 : index // CHECK-DAG: %[[UB0:.+]] = memref.dim %{{.+}}, %[[C0]] // CHECK-DAG: %[[UB1:.+]] = memref.dim %{{.+}}, %[[C1]] // CHECK-DAG: %[[UB2:.+]] = memref.dim %{{.+}}, %[[C2]] // CHECK-DAG: %[[UB3:.+]] = memref.dim %{{.+}}, %[[C3]] -// CHECK: %[[T4:.+]] = muli %[[UB3]], %[[UB2]] -// CHECK: %[[T5:.+]] = muli %[[T4]], %[[UB1]] -// CHECK: %[[UB:.+]] = muli %[[T5]], %[[UB0]] +// CHECK: %[[T4:.+]] = arith.muli %[[UB3]], %[[UB2]] +// CHECK: %[[T5:.+]] = arith.muli %[[T4]], %[[UB1]] +// CHECK: %[[UB:.+]] = arith.muli %[[T5]], %[[UB0]] // CHECK-DAG: %[[BID:.+]] = "gpu.block_id"() {dimension = "x"} // CHECK-DAG: %[[BDIM:.+]] = "gpu.block_dim"() {dimension = "x"} // CHECK-DAG: %[[TID:.+]] = "gpu.thread_id"() {dimension = "x"} -// CHECK: %[[BOFFSET:.+]] = muli %[[BID]], %[[BDIM]] -// CHECK: %[[IV:.+]] = addi %[[BOFFSET]], %[[TID]] -// CHECK: %[[COND:.+]] = cmpi slt, %[[IV]], %[[UB]] +// CHECK: %[[BOFFSET:.+]] = arith.muli %[[BID]], %[[BDIM]] +// CHECK: %[[IV:.+]] = arith.addi %[[BOFFSET]], %[[TID]] +// CHECK: %[[COND:.+]] = arith.cmpi slt, %[[IV]], %[[UB]] // CHECK: scf.if %[[COND]] -// CHECK: %[[IV0:.+]] = divi_signed %[[IV]], %[[T5]] -// CHECK: %[[T14:.+]] = remi_signed %[[IV]], %[[T5]] -// CHECK: %[[IV1:.+]] = divi_signed %[[T14]], %[[T4]] -// CHECK: %[[T16:.+]] = remi_signed %[[T14]], %[[T4]] -// CHECK: %[[IV2:.+]] = divi_signed %[[T16]], %[[UB3]] -// CHECK: %[[IV3:.+]] = remi_signed %[[T16]], %[[UB3]] +// CHECK: %[[IV0:.+]] = arith.divsi %[[IV]], %[[T5]] +// CHECK: %[[T14:.+]] = arith.remsi %[[IV]], %[[T5]] +// CHECK: %[[IV1:.+]] = arith.divsi %[[T14]], %[[T4]] +// CHECK: %[[T16:.+]] = arith.remsi %[[T14]], %[[T4]] +// CHECK: %[[IV2:.+]] = arith.divsi %[[T16]], %[[UB3]] +// CHECK: %[[IV3:.+]] = arith.remsi %[[T16]], %[[UB3]] // CHECK: load %{{.+}}[%[[IV0]], %[[IV1]], %[[IV2]], %[[IV3]]] // CHECK: load %{{.+}}[%[[IV0]], %[[IV1]], %[[IV2]], %[[IV3]]] // CHECK: store %{{.+}}[%[[IV0]], %[[IV1]], %[[IV2]], %[[IV3]]] @@ -83,7 +83,7 @@ hal.executable.entry_point @parallel_4D_static attributes {interface = @io, ordinal = 0 : index} builtin.module { func @parallel_4D_static() { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %arg0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<3x4x5x6xf32> %arg1 = hal.interface.binding.subspan @io::@arg1[%c0] : memref<3x4x5x6xf32> %arg2 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<3x4x5x6xf32> @@ -93,7 +93,7 @@ ins(%arg0, %arg1 : memref<3x4x5x6xf32>, memref<3x4x5x6xf32>) outs(%arg2 : memref<3x4x5x6xf32>) { ^bb0(%arg3 : f32, %arg4 : f32, %arg5 : f32): - %0 = addf %arg3, %arg4 : f32 + %0 = arith.addf %arg3, %arg4 : f32 linalg.yield %0 : f32 } return @@ -107,23 +107,23 @@ } } // CHECK-LABEL: func @parallel_4D_static() -// CHECK-DAG: %[[C360:.+]] = constant 360 : index -// CHECK-DAG: %[[C120:.+]] = constant 120 : index -// CHECK-DAG: %[[C30:.+]] = constant 30 : index -// CHECK-DAG: %[[C6:.+]] = constant 6 : index +// CHECK-DAG: %[[C360:.+]] = arith.constant 360 : index +// CHECK-DAG: %[[C120:.+]] = arith.constant 120 : index +// CHECK-DAG: %[[C30:.+]] = arith.constant 30 : index +// CHECK-DAG: %[[C6:.+]] = arith.constant 6 : index // CHECK-DAG: %[[BID:.+]] = "gpu.block_id"() {dimension = "x"} // CHECK-DAG: %[[BDIM:.+]] = "gpu.block_dim"() {dimension = "x"} // CHECK-DAG: %[[TID:.+]] = "gpu.thread_id"() {dimension = "x"} -// CHECK: %[[BOFFSET:.+]] = muli %[[BID]], %[[BDIM]] -// CHECK: %[[IV:.+]] = addi %[[BOFFSET]], %[[TID]] -// CHECK: %[[COND:.+]] = cmpi slt, %[[IV]], %[[C360]] +// CHECK: %[[BOFFSET:.+]] = arith.muli %[[BID]], %[[BDIM]] +// CHECK: %[[IV:.+]] = arith.addi %[[BOFFSET]], %[[TID]] +// CHECK: %[[COND:.+]] = arith.cmpi slt, %[[IV]], %[[C360]] // CHECK: scf.if %[[COND]] -// CHECK: %[[IV0:.+]] = divi_signed %[[IV]], %[[C120]] -// CHECK: %[[T14:.+]] = remi_signed %[[IV]], %[[C120]] -// CHECK: %[[IV1:.+]] = divi_signed %[[T14]], %[[C30]] -// CHECK: %[[T16:.+]] = remi_signed %[[T14]], %[[C30]] -// CHECK: %[[IV2:.+]] = divi_signed %[[T16]], %[[C6]] -// CHECK: %[[IV3:.+]] = remi_signed %[[T16]], %[[C6]] +// CHECK: %[[IV0:.+]] = arith.divsi %[[IV]], %[[C120]] +// CHECK: %[[T14:.+]] = arith.remsi %[[IV]], %[[C120]] +// CHECK: %[[IV1:.+]] = arith.divsi %[[T14]], %[[C30]] +// CHECK: %[[T16:.+]] = arith.remsi %[[T14]], %[[C30]] +// CHECK: %[[IV2:.+]] = arith.divsi %[[T16]], %[[C6]] +// CHECK: %[[IV3:.+]] = arith.remsi %[[T16]], %[[C6]] // CHECK: load %{{.+}}[%[[IV0]], %[[IV1]], %[[IV2]], %[[IV3]]] // CHECK: load %{{.+}}[%[[IV0]], %[[IV1]], %[[IV2]], %[[IV3]]] // CHECK: store %{{.+}}[%[[IV0]], %[[IV1]], %[[IV2]], %[[IV3]]] @@ -147,7 +147,7 @@ hal.executable.entry_point @scalar_add attributes {interface = @io, ordinal = 0 : index} builtin.module { func @scalar_add() attributes {hal.num_workgroups_fn = @scalar_add__num_workgroups__} { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %arg0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<f32> %arg1 = hal.interface.binding.subspan @io::@arg1[%c0] : memref<f32> %arg2 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<f32> @@ -155,7 +155,7 @@ ins(%arg0, %arg1 : memref<f32>, memref<f32>) outs(%arg2 : memref<f32>) { ^bb0(%arg3 : f32, %arg4 : f32, %arg5 : f32): - %0 = addf %arg3, %arg4 : f32 + %0 = arith.addf %arg3, %arg4 : f32 linalg.yield %0 : f32 } return @@ -191,7 +191,7 @@ } builtin.module { func @reduce_sum() { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %arg0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<40x50x75xf32> %arg1 = hal.interface.binding.subspan @io::@arg1[%c0] : memref<f32> %arg2 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<40xf32> @@ -205,12 +205,12 @@ ^bb0(%arg6: f32, %arg7: f32, %arg8: f32): // no predecessors %idx1 = linalg.index 1 : index %idx2 = linalg.index 2 : index - %zero = constant 0 : index - %0 = cmpi eq, %idx2, %zero : index - %1 = cmpi eq, %idx1, %zero : index - %2 = and %0, %1 : i1 + %zero = arith.constant 0 : index + %0 = arith.cmpi eq, %idx2, %zero : index + %1 = arith.cmpi eq, %idx1, %zero : index + %2 = arith.andi %0, %1 : i1 %3 = select %2, %arg7, %arg8 : f32 - %4 = addf %arg6, %3 : f32 + %4 = arith.addf %arg6, %3 : f32 linalg.yield %4 : f32 } return @@ -224,13 +224,13 @@ } } //CHECK-LABEL: func @reduce_sum -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C40:.+]] = constant 40 : index -// CHECK-DAG: %[[C50:.+]] = constant 50 : index -// CHECK-DAG: %[[C75:.+]] = constant 75 : index -// CHECK: %[[COND:.+]] = cmpi slt, %{{.+}}, %[[C40]] +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C40:.+]] = arith.constant 40 : index +// CHECK-DAG: %[[C50:.+]] = arith.constant 50 : index +// CHECK-DAG: %[[C75:.+]] = arith.constant 75 : index +// CHECK: %[[COND:.+]] = arith.cmpi slt, %{{.+}}, %[[C40]] // CHECK: scf.if %[[COND]] // CHECK: scf.for %[[IV0:.+]] = %{{.+}} to %[[C50]] // CHECK: scf.for %[[IV1:.+]] = %{{.+}} to %[[C75]] -// CHECK-DAG: %[[ISZERO0:.+]] = cmpi eq, %[[IV0]], %[[C0]] -// CHECK-DAG: %[[ISZERO1:.+]] = cmpi eq, %[[IV1]], %[[C0]] +// CHECK-DAG: %[[ISZERO0:.+]] = arith.cmpi eq, %[[IV0]], %[[C0]] +// CHECK-DAG: %[[ISZERO1:.+]] = arith.cmpi eq, %[[IV1]], %[[C0]]
diff --git a/iree/compiler/Codegen/SPIRV/test/fold_gpu_procid_uses.mlir b/iree/compiler/Codegen/SPIRV/test/fold_gpu_procid_uses.mlir index 3c76311..a80e168 100644 --- a/iree/compiler/Codegen/SPIRV/test/fold_gpu_procid_uses.mlir +++ b/iree/compiler/Codegen/SPIRV/test/fold_gpu_procid_uses.mlir
@@ -9,9 +9,9 @@ ordinal = 0 : index } { ^bb0(%arg0 : index, %arg1 : index, %arg2 : index): - %x = constant 112: index - %y = constant 42: index - %z = constant 1: index + %x = arith.constant 112: index + %y = arith.constant 42: index + %z = arith.constant 1: index hal.return %x, %y, %z: index, index, index } builtin.module { @@ -28,9 +28,9 @@ } } // CHECK-LABEL: func @fold_block_id() -// CHECK-DAG: %[[C3:.+]] = constant 3 -// CHECK-DAG: %[[C8:.+]] = constant 8 -// CHECK-DAG: %[[C11:.+]] = constant 11 +// CHECK-DAG: %[[C3:.+]] = arith.constant 3 +// CHECK-DAG: %[[C8:.+]] = arith.constant 8 +// CHECK-DAG: %[[C11:.+]] = arith.constant 11 // CHECK-DAG: return %[[C3]], %[[C8]], %[[C11]] // ----- @@ -44,9 +44,9 @@ ordinal = 0 : index } { ^bb0(%arg0 : index, %arg1 : index, %arg2 : index): - %x = constant 112: index - %y = constant 42: index - %z = constant 1: index + %x = arith.constant 112: index + %y = arith.constant 42: index + %z = arith.constant 1: index hal.return %x, %y, %z: index, index, index } builtin.module { @@ -63,9 +63,9 @@ } } // CHECK-LABEL: func @fold_interface_workgroup_id() -// CHECK-DAG: %[[C3:.+]] = constant 3 -// CHECK-DAG: %[[C8:.+]] = constant 8 -// CHECK-DAG: %[[C11:.+]] = constant 11 +// CHECK-DAG: %[[C3:.+]] = arith.constant 3 +// CHECK-DAG: %[[C8:.+]] = arith.constant 8 +// CHECK-DAG: %[[C11:.+]] = arith.constant 11 // CHECK-DAG: return %[[C3]], %[[C8]], %[[C11]] // ----- @@ -93,9 +93,9 @@ } } // CHECK-LABEL: func @fold_thread_id() -// CHECK-DAG: %[[C7:.+]] = constant 7 -// CHECK-DAG: %[[C11:.+]] = constant 11 -// CHECK-DAG: %[[C21:.+]] = constant 21 +// CHECK-DAG: %[[C7:.+]] = arith.constant 7 +// CHECK-DAG: %[[C11:.+]] = arith.constant 11 +// CHECK-DAG: %[[C21:.+]] = arith.constant 21 // CHECK-DAG: return %[[C7]], %[[C11]], %[[C21]] // -----
diff --git a/iree/compiler/Codegen/SPIRV/test/pipeline_matmul_cooperative_matrix.mlir b/iree/compiler/Codegen/SPIRV/test/pipeline_matmul_cooperative_matrix.mlir index 4306690..2698b9a 100644 --- a/iree/compiler/Codegen/SPIRV/test/pipeline_matmul_cooperative_matrix.mlir +++ b/iree/compiler/Codegen/SPIRV/test/pipeline_matmul_cooperative_matrix.mlir
@@ -31,9 +31,9 @@ } builtin.module { func @matmul_cooperative_matrix() { - %c32 = constant 32 : index - %c4096 = constant 4096 : index - %c0 = constant 0 : index + %c32 = arith.constant 32 : index + %c4096 = arith.constant 4096 : index + %c0 = arith.constant 0 : index %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : memref<4096x4096xf16> %1 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : memref<4096x4096xf16> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : memref<4096x4096xf16>
diff --git a/iree/compiler/Codegen/SPIRV/test/pipeline_matmul_vectorization.mlir b/iree/compiler/Codegen/SPIRV/test/pipeline_matmul_vectorization.mlir index eb7e645..53da2f5 100644 --- a/iree/compiler/Codegen/SPIRV/test/pipeline_matmul_vectorization.mlir +++ b/iree/compiler/Codegen/SPIRV/test/pipeline_matmul_vectorization.mlir
@@ -17,9 +17,9 @@ } builtin.module { func @fuse_and_vectorize_fill_matmul() { - %c0 = constant 0 : index - %cst = constant 0.000000e+00 : f32 - %c4096 = constant 4096 : index + %c0 = arith.constant 0 : index + %cst = arith.constant 0.000000e+00 : f32 + %c4096 = arith.constant 4096 : index %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : !flow.dispatch.tensor<readonly:4096x4096xf32> %1 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : !flow.dispatch.tensor<readonly:4096x4096xf32> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : !flow.dispatch.tensor<writeonly:4096x4096xf32> @@ -87,10 +87,10 @@ } builtin.module { func @fuse_and_vectorize_matmul_add() { - %c0 = constant 0 : index - %cst = constant 0.000000e+00 : f32 - %c1024 = constant 1024 : index - %c256 = constant 256 : index + %c0 = arith.constant 0 : index + %cst = arith.constant 0.000000e+00 : f32 + %c1024 = arith.constant 1024 : index + %c256 = arith.constant 256 : index %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : !flow.dispatch.tensor<readonly:1024x256xf32> %1 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : !flow.dispatch.tensor<readonly:1024x512xf32> %2 = hal.interface.binding.subspan @io::@s0b2_ro_external[%c0] : !flow.dispatch.tensor<readonly:512x256xf32> @@ -124,7 +124,7 @@ %22 = linalg.matmul {__internal_linalg_transform__ = "workgroup", lowering.config = #config} ins(%15, %17 : tensor<?x512xf32>, tensor<512x?xf32>) outs(%21 : tensor<?x?xf32>) -> tensor<?x?xf32> %23 = linalg.generic {indexing_maps = [affine_map<(d0, d1) -> (d0, d1)>, affine_map<(d0, d1) -> (d0, d1)>, affine_map<(d0, d1) -> (d0, d1)>], iterator_types = ["parallel", "parallel"]} ins(%22, %10 : tensor<?x?xf32>, tensor<?x?xf32>) outs(%13 : tensor<?x?xf32>) attrs = {__internal_linalg_transform__ = "workgroup", lowering.config = #config} { ^bb0(%arg2: f32, %arg3: f32, %arg4: f32): // no predecessors - %24 = addf %arg2, %arg3 : f32 + %24 = arith.addf %arg2, %arg3 : f32 linalg.yield %24 : f32 } -> tensor<?x?xf32> flow.dispatch.tensor.store %23, %3, offsets = [%arg0, %arg1], sizes = [%11, %12], strides = [1, 1] : tensor<?x?xf32> -> !flow.dispatch.tensor<writeonly:1024x256xf32>
diff --git a/iree/compiler/Codegen/SPIRV/test/promote_workgroup_memory.mlir b/iree/compiler/Codegen/SPIRV/test/promote_workgroup_memory.mlir index c1d0dc8..27aa8c0 100644 --- a/iree/compiler/Codegen/SPIRV/test/promote_workgroup_memory.mlir +++ b/iree/compiler/Codegen/SPIRV/test/promote_workgroup_memory.mlir
@@ -14,9 +14,9 @@ } builtin.module { func @matmul_promote_workgroup_memory() { - %c32 = constant 32 : index - %c50 = constant 50 : index - %c0 = constant 0 : index + %c32 = arith.constant 32 : index + %c50 = arith.constant 50 : index + %c0 = arith.constant 0 : index %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : memref<25x50xf32> %1 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : memref<50x75xf32> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : memref<25x75xf32> @@ -88,7 +88,7 @@ } builtin.module { func @conv_promote_workgroup_memory() { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : memref<3x4x6x14xf32> %1 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : memref<2x15x14x6xf32> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : memref<2x13x11x14xf32>
diff --git a/iree/compiler/Codegen/SPIRV/test/remove_one_trip_tiled_loop.mlir b/iree/compiler/Codegen/SPIRV/test/remove_one_trip_tiled_loop.mlir index 1318797..2dd847e 100644 --- a/iree/compiler/Codegen/SPIRV/test/remove_one_trip_tiled_loop.mlir +++ b/iree/compiler/Codegen/SPIRV/test/remove_one_trip_tiled_loop.mlir
@@ -14,10 +14,10 @@ } builtin.module { builtin.func @static_shaped_conv() { - %cst = constant 0.000000e+00 : f32 - %c112 = constant 112 : index - %c32 = constant 32 : index - %c0 = constant 0 : index + %cst = arith.constant 0.000000e+00 : f32 + %c112 = arith.constant 112 : index + %c32 = arith.constant 32 : index + %c0 = arith.constant 0 : index %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : memref<1x225x225x3xf32> %1 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : memref<3x3x3x32xf32> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : memref<1x112x112x32xf32>
diff --git a/iree/compiler/Codegen/SPIRV/test/tile_and_distribute_scatter.mlir b/iree/compiler/Codegen/SPIRV/test/tile_and_distribute_scatter.mlir index a248b46..7c14a45 100644 --- a/iree/compiler/Codegen/SPIRV/test/tile_and_distribute_scatter.mlir +++ b/iree/compiler/Codegen/SPIRV/test/tile_and_distribute_scatter.mlir
@@ -16,9 +16,9 @@ builtin.module { builtin.func @static_scatter_update_slice() { - %c40 = constant 40 : index - %c500 = constant 500 : index - %c0 = constant 0 : index + %c40 = arith.constant 40 : index + %c500 = arith.constant 500 : index + %c0 = arith.constant 0 : index %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : memref<40x500xi32> %1 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : memref<40x1xi32> %2 = hal.interface.binding.subspan @io::@s0b2_rw_external[%c0] : memref<100x500xi32>
diff --git a/iree/compiler/Codegen/SPIRV/test/tile_and_distribute_sort.mlir b/iree/compiler/Codegen/SPIRV/test/tile_and_distribute_sort.mlir index a9c25a0..9c93afe 100644 --- a/iree/compiler/Codegen/SPIRV/test/tile_and_distribute_sort.mlir +++ b/iree/compiler/Codegen/SPIRV/test/tile_and_distribute_sort.mlir
@@ -13,9 +13,9 @@ } builtin.module { builtin.func @static_3d_sort() { - %c64 = constant 64 : index - %c128 = constant 128 : index - %c0 = constant 0 : index + %c64 = arith.constant 64 : index + %c128 = arith.constant 128 : index + %c0 = arith.constant 0 : index %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : memref<64x32x128xi32> %1 = hal.interface.binding.subspan @io::@s0b1_xw_external[%c0] : memref<64x32x128xi32> %workgroup_id_x = hal.interface.workgroup.id[0] : index @@ -33,7 +33,7 @@ linalg.copy(%5, %6) {__internal_linalg_transform__ = "workgroup", lowering.config = {tileSizes = [[1, 0, 16], [], [1, 0, 1]]}} : memref<?x?x?xi32>, memref<1x32x16xi32, affine_map<(d0, d1, d2)[s0] -> (d0 * 4096 + s0 + d1 * 128 + d2)>> linalg_ext.sort dimension(1) {__internal_linalg_transform__ = "workgroup", lowering.config = {tileSizes = [[1, 0, 16], [], [1, 0, 1]]}} outs(%7 : memref<?x32x?xi32, affine_map<(d0, d1, d2)[s0] -> (d0 * 4096 + s0 + d1 * 128 + d2)>>) { ^bb0(%arg2: i32, %arg3: i32): // no predecessors - %8 = cmpi slt, %arg2, %arg3 : i32 + %8 = arith.cmpi slt, %arg2, %arg3 : i32 linalg_ext.yield %8 : i1 } }
diff --git a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize.mlir b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize.mlir index 03a188a..aa9fb4f 100644 --- a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize.mlir +++ b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize.mlir
@@ -24,15 +24,15 @@ } builtin.module { func @matmul() { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %M = hal.interface.load.constant offset = 0 : index %N = hal.interface.load.constant offset = 1 : index %K = hal.interface.load.constant offset = 2 : index %arg0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<?x?xf32>{%M, %K} %arg1 = hal.interface.binding.subspan @io::@arg1[%c0] : memref<?x?xf32>{%K, %N} %arg2 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<?x?xf32>{%M, %N} - %c4 = constant 4 : index - %c1 = constant 1 : index + %c4 = arith.constant 4 : index + %c1 = arith.constant 1 : index %0 = memref.dim %arg0, %c1 : memref<?x?xf32> %1 = "gpu.block_id"() {dimension = "x"} : () -> index %2 = "gpu.block_id"() {dimension = "y"} : () -> index @@ -68,8 +68,8 @@ } } // CHECK-LABEL: func @matmul -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK-DAG: %[[TIDX:.+]] = "gpu.thread_id"() {dimension = "x"} // CHECK-DAG: %[[TIDY:.+]] = "gpu.thread_id"() {dimension = "y"} // CHECK-DAG: %[[BDIMX:.+]] = "gpu.block_dim"() {dimension = "x"} @@ -98,8 +98,8 @@ } builtin.module { func @conv_1d() { - %cst = constant 0.000000e+00 : f32 - %c0 = constant 0 : index + %cst = arith.constant 0.000000e+00 : f32 + %c0 = arith.constant 0 : index %0 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<3x6x1xf32> %1 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<3x8x1xf32> %2 = hal.interface.binding.subspan @io::@arg1[%c0] : memref<3x1x1xf32> @@ -131,7 +131,7 @@ } // CHECK-LABEL: func @conv_1d -// CHECK: %[[C0:.+]] = constant 0 : index +// CHECK: %[[C0:.+]] = arith.constant 0 : index // CHECK: %[[RET:.+]] = hal.interface.binding.subspan @io::@ret0 // CHECK: %[[ARG0:.+]] = hal.interface.binding.subspan @io::@arg0 // CHECK: %[[ARG1:.+]] = hal.interface.binding.subspan @io::@arg1 @@ -181,7 +181,7 @@ } builtin.module { func @conv_no_padding() { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %n = hal.interface.load.constant offset = 0 : index %oh = hal.interface.load.constant offset = 1 : index %ow = hal.interface.load.constant offset = 2 : index @@ -194,9 +194,9 @@ %arg0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<?x?x?x?xf32>{%n, %ih, %iw, %ic} %arg1 = hal.interface.binding.subspan @io::@arg1[%c0] : memref<?x?x?x?xf32>{%fh, %fw, %ic, %oc} %arg2 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<?x?x?x?xf32>{%n, %oh, %ow, %oc} - %c2 = constant 2 : index - %c3 = constant 3 : index - %c1 = constant 1 : index + %c2 = arith.constant 2 : index + %c3 = arith.constant 3 : index + %c1 = arith.constant 1 : index %0 = memref.dim %arg0, %c0 : memref<?x?x?x?xf32> %1 = memref.dim %arg0, %c1 : memref<?x?x?x?xf32> %2 = memref.dim %arg1, %c0 : memref<?x?x?x?xf32> @@ -256,8 +256,8 @@ // CHECK-DAG: %[[ARG0:.+]] = hal.interface.binding.subspan @io::@arg0 // CHECK-DAG: %[[ARG1:.+]] = hal.interface.binding.subspan @io::@arg1 // CHECK-DAG: %[[RET0:.+]] = hal.interface.binding.subspan @io::@ret0 -// CHECK-DAG: %[[C1:.+]] = constant 1 -// CHECK-DAG: %[[C2:.+]] = constant 2 +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 +// CHECK-DAG: %[[C2:.+]] = arith.constant 2 // CHECK-DAG: %[[N:.+]] = memref.dim %[[ARG1]], %[[C0]] // CHECK-DAG: %[[P:.+]] = memref.dim %[[RET0]], %[[C1]] // CHECK-DAG: %[[Q:.+]] = memref.dim %[[RET0]], %[[C2]] @@ -308,8 +308,8 @@ } builtin.module { func @conv_3d() { - %cst = constant 0.000000e+00 : f32 - %c0 = constant 0 : index + %cst = arith.constant 0.000000e+00 : f32 + %c0 = arith.constant 0 : index %0 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<2x7x7x7x2xf32> %1 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<2x8x8x8x3xf32> %2 = hal.interface.binding.subspan @io::@arg1[%c0] : memref<2x2x2x3x2xf32> @@ -382,7 +382,7 @@ } builtin.module { func @pooling_nhwc_max() { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<2x16x16x6xf32> %1 = hal.interface.binding.subspan @io::@arg1[%c0] : memref<3x4xf32> %2 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<2x14x13x6xf32>
diff --git a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_batch_matmul.mlir b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_batch_matmul.mlir index 704a0a1..aec3f66 100644 --- a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_batch_matmul.mlir +++ b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_batch_matmul.mlir
@@ -16,9 +16,9 @@ } builtin.module { func @batch_matmul_static_shape() { - %c0 = constant 0 : index - %c4 = constant 4 : index - %c1024 = constant 1024 : index + %c0 = arith.constant 0 : index + %c4 = arith.constant 4 : index + %c1024 = arith.constant 1024 : index %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : memref<4x1024x1024xf32> %1 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : memref<4x1024x1024xf32> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : memref<4x1024x1024xf32> @@ -70,14 +70,14 @@ // CHECK-DAG: %[[ARG0:.+]] = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] // CHECK-DAG: %[[ARG1:.+]] = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] // CHECK-DAG: %[[RET0:.+]] = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index -// CHECK-DAG: %[[C2:.+]] = constant 2 : index -// CHECK-DAG: %[[C3:.+]] = constant 3 : index -// CHECK-DAG: %[[C4:.+]] = constant 4 : index -// CHECK-DAG: %[[C5:.+]] = constant 5 : index -// CHECK-DAG: %[[C6:.+]] = constant 6 : index -// CHECK-DAG: %[[C7:.+]] = constant 7 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index +// CHECK-DAG: %[[C2:.+]] = arith.constant 2 : index +// CHECK-DAG: %[[C3:.+]] = arith.constant 3 : index +// CHECK-DAG: %[[C4:.+]] = arith.constant 4 : index +// CHECK-DAG: %[[C5:.+]] = arith.constant 5 : index +// CHECK-DAG: %[[C6:.+]] = arith.constant 6 : index +// CHECK-DAG: %[[C7:.+]] = arith.constant 7 : index // CHECK: %[[BIDX:.+]] = hal.interface.workgroup.id[0] // CHECK: %[[BCNTX:.+]] = hal.interface.workgroup.count[0] // CHECK: %[[BIDY:.+]] = hal.interface.workgroup.id[1] @@ -386,10 +386,10 @@ } builtin.module { func @fused_fill_batch_matmul() { - %zero = constant 0.0 : f32 - %c0 = constant 0 : index - %c4 = constant 4 : index - %c1024 = constant 1024 : index + %zero = arith.constant 0.0 : f32 + %c0 = arith.constant 0 : index + %c4 = arith.constant 4 : index + %c1024 = arith.constant 1024 : index %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : memref<4x1024x1024xf32> %1 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : memref<4x1024x1024xf32> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : memref<4x1024x1024xf32>
diff --git a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_conv.mlir b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_conv.mlir index c74b7bc..8f8bd96 100644 --- a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_conv.mlir +++ b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_conv.mlir
@@ -16,17 +16,17 @@ translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [16, 4, 4]} } { ^bb0(%arg0 : index, %arg1 : index, %arg2 : index): - %x = constant 2: index - %y = constant 28: index - %z = constant 28: index + %x = arith.constant 2: index + %y = arith.constant 28: index + %z = arith.constant 28: index hal.return %x, %y, %z: index, index, index } builtin.module { func @conv_static_shape_f32() { - %cst = constant 0.000000e+00 : f32 - %c32 = constant 32 : index - %c112 = constant 112 : index - %c0 = constant 0 : index + %cst = arith.constant 0.000000e+00 : f32 + %c32 = arith.constant 32 : index + %c112 = arith.constant 112 : index + %c0 = arith.constant 0 : index %0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<1x225x225x16xf32> %1 = hal.interface.binding.subspan @io::@arg1[%c0] : memref<3x3x16x32xf32> %2 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<1x112x112x32xf32> @@ -115,17 +115,17 @@ translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [16, 4, 4]} } { ^bb0(%arg0 : index, %arg1 : index, %arg2 : index): - %x = constant 6: index - %y = constant 14: index - %z = constant 14: index + %x = arith.constant 6: index + %y = arith.constant 14: index + %z = arith.constant 14: index hal.return %x, %y, %z: index, index, index } builtin.module { func @depthwise_conv_static_shape_f32() { - %cst = constant 0.000000e+00 : f32 - %c96 = constant 96 : index - %c56 = constant 56 : index - %c0 = constant 0 : index + %cst = arith.constant 0.000000e+00 : f32 + %c96 = arith.constant 96 : index + %c56 = arith.constant 56 : index + %c0 = arith.constant 0 : index %0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<1x113x113x96xf32> %1 = hal.interface.binding.subspan @io::@arg1[%c0] : memref<3x3x1x96xf32> %2 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<1x56x56x96xf32>
diff --git a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_matmul.mlir b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_matmul.mlir index cae1c5a..dc3cf72 100644 --- a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_matmul.mlir +++ b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_matmul.mlir
@@ -16,9 +16,9 @@ } builtin.module { func @matmul_static_shape_f16() { - %cst = constant 0.000000e+00 : f16 - %c0 = constant 0 : index - %c4096 = constant 4096 : index + %cst = arith.constant 0.000000e+00 : f16 + %c0 = arith.constant 0 : index + %c4096 = arith.constant 4096 : index %0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<4096x4096xf16> %1 = hal.interface.binding.subspan @io::@arg1[%c0] : memref<4096x4096xf16> %2 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<4096x4096xf16> @@ -82,9 +82,9 @@ } builtin.module { func @matmul_static_shape_f32() { - %c0 = constant 0 : index - %cst = constant 0.000000e+00 : f32 - %c4096 = constant 4096 : index + %c0 = arith.constant 0 : index + %cst = arith.constant 0.000000e+00 : f32 + %c4096 = arith.constant 4096 : index %0 = hal.interface.binding.subspan @io::@s0b0_ro_external[%c0] : memref<4096x4096xf32> %1 = hal.interface.binding.subspan @io::@s0b1_ro_external[%c0] : memref<4096x4096xf32> %2 = hal.interface.binding.subspan @io::@s0b2_xw_external[%c0] : memref<4096x4096xf32>
diff --git a/iree/compiler/Codegen/SPIRV/test/vector_to_cooperative_matrix.mlir b/iree/compiler/Codegen/SPIRV/test/vector_to_cooperative_matrix.mlir index 4172812..bcc7185 100644 --- a/iree/compiler/Codegen/SPIRV/test/vector_to_cooperative_matrix.mlir +++ b/iree/compiler/Codegen/SPIRV/test/vector_to_cooperative_matrix.mlir
@@ -14,9 +14,9 @@ builtin.module { // CHECK-LABEL: func @kernel_matmul func @kernel_matmul(%arg0: memref<8x32xi8>, %arg1: memref<32x8xi8>, %arg2: memref<8x8xi32>) { - %c0 = constant 0 : index - %cst = constant 0 : i32 - %cst_i8 = constant 0 : i8 + %c0 = arith.constant 0 : index + %cst = arith.constant 0 : i32 + %cst_i8 = arith.constant 0 : i8 %0 = vector.transfer_read %arg0[%c0, %c0], %cst_i8 : memref<8x32xi8>, vector<8x32xi8> %1 = vector.transfer_read %arg1[%c0, %c0], %cst_i8 : memref<32x8xi8>, vector<32x8xi8> %2 = vector.transfer_read %arg2[%c0, %c0], %cst : memref<8x8xi32>, vector<8x8xi32> @@ -49,11 +49,11 @@ builtin.module { // CHECK-LABEL: func @kernel_matmul_licm func @kernel_matmul_licm(%arg0: memref<4096x4096xi8>, %arg1: memref<4096x4096xi8>, %arg2: memref<4096x4096xi32>) { - %c32 = constant 32 : index - %c4096 = constant 4096 : index - %c0 = constant 0 : index - %c0_i32 = constant 0 : i32 - %c0_i8 = constant 0 : i8 + %c32 = arith.constant 32 : index + %c4096 = arith.constant 4096 : index + %c0 = arith.constant 0 : index + %c0_i32 = arith.constant 0 : i32 + %c0_i8 = arith.constant 0 : i8 // CHECK: %[[C:.+]] = spv.CooperativeMatrixLoadNV %4 = vector.transfer_read %arg2[%c0, %c0], %c0_i32 {in_bounds = [true, true]} : memref<4096x4096xi32>, vector<16x16xi32> // CHECK: %[[INIT:.+]] = builtin.unrealized_conversion_cast %[[C]] : !spv.coopmatrix<16x16xi32, Subgroup> to vector<16x16xi32> @@ -95,10 +95,10 @@ builtin.module { // CHECK-LABEL: func @kernel_matmul_vector_memref func @kernel_matmul_vector_memref(%arg0: memref<4096x256xvector<4xi32>>, %arg1: memref<4096x256xvector<4xi32>>, %arg2: memref<4096x1024xvector<4xi32>>) { - %c32 = constant 32 : index - %c4096 = constant 4096 : index - %c0 = constant 0 : index - %cst = constant dense<0> : vector<4xi32> + %c32 = arith.constant 32 : index + %c4096 = arith.constant 4096 : index + %c0 = arith.constant 0 : index + %cst = arith.constant dense<0> : vector<4xi32> // CHECK: %[[C:.+]] = spv.CooperativeMatrixLoadNV %4 = vector.transfer_read %arg2[%c0, %c0], %cst : memref<4096x1024xvector<4xi32>>, vector<16x16xi32> // CHECK: scf.for
diff --git a/iree/compiler/Codegen/SPIRV/test/vectorize_copy_to_workgroup_memory.mlir b/iree/compiler/Codegen/SPIRV/test/vectorize_copy_to_workgroup_memory.mlir index de35a94..f169936 100644 --- a/iree/compiler/Codegen/SPIRV/test/vectorize_copy_to_workgroup_memory.mlir +++ b/iree/compiler/Codegen/SPIRV/test/vectorize_copy_to_workgroup_memory.mlir
@@ -8,16 +8,16 @@ module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>} { func @copy(%arg0: memref<4096x4096xf32>) attributes {spv.entry_point_abi = {local_size = dense<[128, 1, 1]> : vector<3xi32>}} { %a = memref.alloc() : memref<128x32xf32, 3> - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %sv = memref.subview %arg0[%c0, %c0] [128, 32] [1, 1] : memref<4096x4096xf32> to memref<128x32xf32, #map0> linalg.copy(%sv, %a) {__internal_linalg_transform__ = "copy_to_workgroup_memory"} : memref<128x32xf32, #map0>, memref<128x32xf32, 3> return } // CHECK: #[[MAP1:.+]] = affine_map<(d0) -> (d0 * 4)> - // CHECK-DAG: %[[C1024:.+]] = constant 1024 : index - // CHECK-DAG: %[[C8:.+]] = constant 8 : index - // CHECK-DAG: %[[C0:.+]] = constant 0 : index + // CHECK-DAG: %[[C1024:.+]] = arith.constant 1024 : index + // CHECK-DAG: %[[C8:.+]] = arith.constant 8 : index + // CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index // CHECK: %[[ALLOC:.+]] = memref.alloc() : memref<128x32xf32, 3> // CHECK: %[[DST:.+]] = memref.subview %{{.+}}[0, 0] [128, 32] [1, 1] : memref<4096x4096xf32> to memref<128x32xf32, #map0> // CHECK: %[[TIDx:.+]] = "gpu.thread_id"() {dimension = "x"} : () -> index @@ -26,15 +26,15 @@ // CHECK: %[[DIMy:.+]] = "gpu.block_dim"() {dimension = "y"} : () -> index // CHECK: %[[TIDz:.+]] = "gpu.thread_id"() {dimension = "z"} : () -> index // CHECK: %[[DIMz:.+]] = "gpu.block_dim"() {dimension = "z"} : () -> index - // CHECK: %[[LIDz:.+]] = muli %[[TIDz]], %[[DIMy]] : index - // CHECK: %[[LIDzy:.+]] = addi %[[LIDz]], %[[TIDy]] : index - // CHECK: %[[DIMzy:.+]] = muli %[[DIMz]], %[[DIMy]] : index - // CHECK: %[[LIDzyx:.+]] = muli %[[LIDzy]], %[[DIMx]] : index - // CHECK: %[[LID:.+]] = addi %[[LIDzyx]], %[[TIDx]] : index - // CHECK: %[[DIMzyx:.+]] = muli %[[DIMzy]], %[[DIMx]] : index + // CHECK: %[[LIDz:.+]] = arith.muli %[[TIDz]], %[[DIMy]] : index + // CHECK: %[[LIDzy:.+]] = arith.addi %[[LIDz]], %[[TIDy]] : index + // CHECK: %[[DIMzy:.+]] = arith.muli %[[DIMz]], %[[DIMy]] : index + // CHECK: %[[LIDzyx:.+]] = arith.muli %[[LIDzy]], %[[DIMx]] : index + // CHECK: %[[LID:.+]] = arith.addi %[[LIDzyx]], %[[TIDx]] : index + // CHECK: %[[DIMzyx:.+]] = arith.muli %[[DIMzy]], %[[DIMx]] : index // CHECK: scf.for %[[IV:.+]] = %[[LID]] to %[[C1024]] step %[[DIMzyx]] { - // CHECK: %[[SIZEx:.+]] = divi_signed %[[IV]], %[[C8]] : index - // CHECK: %[[MOD:.+]] = remi_signed %[[IV]], %[[C8]] : index + // CHECK: %[[SIZEx:.+]] = arith.divsi %[[IV]], %[[C8]] : index + // CHECK: %[[MOD:.+]] = arith.remsi %[[IV]], %[[C8]] : index // CHECK: %[[SIZEy:.+]] = affine.apply #[[MAP1]](%[[MOD]]) // CHECK: %[[SVs:.+]] = memref.subview %[[DST]][%[[SIZEx]], %[[SIZEy]]] [1, 4] [1, 1] : memref<128x32xf32, #map0> to memref<1x4xf32 // CHECK: %[[SVd:.+]] = memref.subview %[[ALLOC]][%[[SIZEx]], %[[SIZEy]]] [1, 4] [1, 1] : memref<128x32xf32, 3> to memref<1x4xf32
diff --git a/iree/compiler/Codegen/SPIRV/test/vectorize_elementwise_ops.mlir b/iree/compiler/Codegen/SPIRV/test/vectorize_elementwise_ops.mlir index de367ba..948b944 100644 --- a/iree/compiler/Codegen/SPIRV/test/vectorize_elementwise_ops.mlir +++ b/iree/compiler/Codegen/SPIRV/test/vectorize_elementwise_ops.mlir
@@ -18,7 +18,7 @@ } builtin.module { func @elementwise_static_shape() { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %arg0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<128xf32> %arg1 = hal.interface.binding.subspan @io::@arg1[%c0] : memref<128xf32> %ret0 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<128xf32> @@ -32,7 +32,7 @@ } ins(%arg0, %arg1 : memref<128xf32>, memref<128xf32>) outs(%ret0 : memref<128xf32>) { ^bb0(%a : f32, %b : f32, %c : f32): - %add = addf %a, %b : f32 + %add = arith.addf %a, %b : f32 linalg.yield %add : f32 } return @@ -67,7 +67,7 @@ } builtin.module { func @elementwise_transpose() { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %arg0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<128x8xf32> %arg1 = hal.interface.binding.subspan @io::@arg1[%c0] : memref<128xf32> %ret0 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<128x8xf32> @@ -81,7 +81,7 @@ } ins(%arg0, %arg1 : memref<128x8xf32>, memref<128xf32>) outs(%ret0 : memref<128x8xf32>) { ^bb0(%a : f32, %b : f32, %c : f32): - %add = addf %a, %b : f32 + %add = arith.addf %a, %b : f32 linalg.yield %add : f32 } return
diff --git a/iree/compiler/Codegen/SPIRV/test/vectorize_load_store.mlir b/iree/compiler/Codegen/SPIRV/test/vectorize_load_store.mlir index 4bfc6e7..8d01659 100644 --- a/iree/compiler/Codegen/SPIRV/test/vectorize_load_store.mlir +++ b/iree/compiler/Codegen/SPIRV/test/vectorize_load_store.mlir
@@ -9,7 +9,7 @@ // CHECK: vector.transfer_write %[[MAT]], %[[ALLOC]][%{{.*}}, %{{.*}}] : vector<32x8xf32>, memref<128x8xvector<4xf32>, 3> // CHECK: memref.dealloc %[[ALLOC]] : memref<128x8xvector<4xf32>, 3> func @copy(%arg0: memref<4096x4096xf32>, %x: index, %y: index) { - %cst = constant 0.000000e+00 : f32 + %cst = arith.constant 0.000000e+00 : f32 %0 = memref.alloc() : memref<128x32xf32, 3> %v = vector.transfer_read %arg0[%x, %y], %cst : memref<4096x4096xf32>, vector<1x4xf32> vector.transfer_write %v, %0[%x, %y] : vector<1x4xf32>, memref<128x32xf32, 3> @@ -25,7 +25,7 @@ // CHECK-LABEL: func @copy // CHECK-SAME: %[[ARG0:.+]]: memref<4096x4096xf32> func @copy(%arg0: memref<4096x4096xf32>, %x: index, %y: index) { - %cst = constant 0.000000e+00 : f32 + %cst = arith.constant 0.000000e+00 : f32 %0 = memref.alloc() : memref<128x32xf32, 3> %s = memref.load %arg0[%x, %y] : memref<4096x4096xf32> memref.store %s, %0[%x, %y] : memref<128x32xf32, 3> @@ -43,8 +43,8 @@ // CHECK: %[[MAT:.+]] = vector.transfer_read %[[A]][%{{.*}}, %{{.*}}], %{{.*}} : memref<4096x1024xvector<4xf32>>, vector<32x8xf32> // CHECK: vector.transfer_write %[[MAT]], %[[B]][%{{.*}}, %{{.*}}] {{.*}} : vector<32x8xf32>, memref<4096x1024xvector<4xf32>> func @resource_copy() { - %cst = constant 0.000000e+00 : f32 - %c0 = constant 0 : index + %cst = arith.constant 0.000000e+00 : f32 + %c0 = arith.constant 0 : index %0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<4096x4096xf32> %1 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<4096x4096xf32> %v = vector.transfer_read %0[%c0, %c0], %cst : memref<4096x4096xf32>, vector<1x4xf32> @@ -69,8 +69,8 @@ // CHECK: %[[MAT:.+]] = vector.transfer_read %[[A]][%{{.*}}, %{{.*}}], %{{.*}} : memref<4096x1024xvector<4xf16>>, vector<32x8xf16> // CHECK: vector.transfer_write %[[MAT]], %[[B]][%{{.*}}, %{{.*}}] {{.*}} : vector<32x8xf16>, memref<4096x1024xvector<4xf16>> func @resource_copy_f16() { - %cst = constant 0.000000e+00 : f16 - %c0 = constant 0 : index + %cst = arith.constant 0.000000e+00 : f16 + %c0 = arith.constant 0 : index %0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<4096x4096xf16> %1 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<4096x4096xf16> %v = vector.transfer_read %0[%c0, %c0], %cst : memref<4096x4096xf16>, vector<1x4xf16> @@ -95,8 +95,8 @@ // CHECK: %[[MAT:.+]] = vector.transfer_read %[[A]][%{{.*}}, %{{.*}}], %{{.*}} : memref<4096x512xvector<4xf32>>, vector<32x8xf16> // CHECK: vector.transfer_write %[[MAT]], %[[B]][%{{.*}}, %{{.*}}] {{.*}} : vector<32x8xf16>, memref<4096x512xvector<4xf32>> func @resource_copy_8xf16() { - %cst = constant 0.000000e+00 : f16 - %c0 = constant 0 : index + %cst = arith.constant 0.000000e+00 : f16 + %c0 = arith.constant 0 : index %0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<4096x4096xf16> %1 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<4096x4096xf16> %v = vector.transfer_read %0[%c0, %c0], %cst : memref<4096x4096xf16>, vector<1x8xf16> @@ -115,8 +115,8 @@ // CHECK-LABEL: func @resource_copy_dynamic_shape() func @resource_copy_dynamic_shape() { - %cst = constant 0.000000e+00 : f32 - %c0 = constant 0 : index + %cst = arith.constant 0.000000e+00 : f32 + %c0 = arith.constant 0 : index // CHECK: %[[DIM0:.+]] = hal.interface.load.constant offset = 0 : index // CHECK: %[[DIM1:.+]] = hal.interface.load.constant offset = 1 : index %dim0 = hal.interface.load.constant offset = 0 : index @@ -148,8 +148,8 @@ // CHECK-LABEL: func @resource_copy_dynamic_last_dim() func @resource_copy_dynamic_last_dim() { - %cst = constant 0.000000e+00 : f32 - %c0 = constant 0 : index + %cst = arith.constant 0.000000e+00 : f32 + %c0 = arith.constant 0 : index %dim = hal.interface.load.constant offset = 0 : index // CHECK: hal.interface.binding.subspan @io::@arg0[{{.+}}] : memref<4096x?xf32> // CHECK: hal.interface.binding.subspan @io::@ret0[{{.+}}] : memref<4096x?xf32> @@ -169,8 +169,8 @@ // CHECK-LABEL: func @do_not_vectorize_odd_vector_size func @do_not_vectorize_odd_vector_size() { - %cst = constant 0.0 : f32 - %c0 = constant 0 : index + %cst = arith.constant 0.0 : f32 + %c0 = arith.constant 0 : index // CHECK: hal.interface.binding.subspan // CHECK-SAME: memref<4x3xf32> %0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<4x3xf32> @@ -190,8 +190,8 @@ // ----- func @vectorize_binding_subspan() { - %cst = constant 0.000000e+00 : f32 - %c0 = constant 0 : index + %cst = arith.constant 0.000000e+00 : f32 + %c0 = arith.constant 0 : index // CHECK: hal.interface.binding.subspan @io::@arg0[%c0] // CHECK-SAME: memref<4096x1024xvector<4xf32>> // CHECK: hal.interface.binding.subspan @io::@ret0[%c0] @@ -212,15 +212,15 @@ // CHECK-LABEL: func @scalarize_vector_transfer_op func @scalarize_vector_transfer_op(%arg: vector<3xf32>) -> (vector<3xf32>) { - %c0 = constant 0: index - %c3 = constant 3: index - %f0 = constant 0.0 : f32 + %c0 = arith.constant 0: index + %c3 = arith.constant 3: index + %f0 = arith.constant 0.0 : f32 %0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<20xf32> %2 = hal.interface.binding.subspan @io::@ret1[%c0] : memref<20xf32> - // CHECK-DAG: %[[INDEX0:.+]] = constant 3 : index - // CHECK-DAG: %[[INDEX1:.+]] = constant 4 : index - // CHECK-DAG: %[[INDEX2:.+]] = constant 5 : index - // CHECK-DAG: %[[CST:.+]] = constant dense<0.000000e+00> : vector<3xf32> + // CHECK-DAG: %[[INDEX0:.+]] = arith.constant 3 : index + // CHECK-DAG: %[[INDEX1:.+]] = arith.constant 4 : index + // CHECK-DAG: %[[INDEX2:.+]] = arith.constant 5 : index + // CHECK-DAG: %[[CST:.+]] = arith.constant dense<0.000000e+00> : vector<3xf32> // CHECK: %[[ELEM0:.+]] = memref.load %{{.+}}[%[[INDEX0]]] // CHECK: %[[V0:.+]] = vector.insert %[[ELEM0]], %[[CST]] [0] : f32 into vector<3xf32>
diff --git a/iree/compiler/Codegen/SPIRV/test/vectorize_matmul.mlir b/iree/compiler/Codegen/SPIRV/test/vectorize_matmul.mlir index 3f48a5f..0b505fb 100644 --- a/iree/compiler/Codegen/SPIRV/test/vectorize_matmul.mlir +++ b/iree/compiler/Codegen/SPIRV/test/vectorize_matmul.mlir
@@ -30,9 +30,9 @@ } builtin.module { func @matmul_static_shape() { - %c32 = constant 32 : index - %c4096 = constant 4096 : index - %c0 = constant 0 : index + %c32 = arith.constant 32 : index + %c4096 = arith.constant 4096 : index + %c0 = arith.constant 0 : index %0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<4096x4096xf16> %1 = hal.interface.binding.subspan @io::@arg1[%c0] : memref<4096x4096xf16> %2 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<4096x4096xf16> @@ -63,11 +63,11 @@ // CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 * 64)> // CHECK: func @matmul_static_shape -// CHECK-DAG: %[[CST:.+]] = constant 0.0 -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C16:.+]] = constant 16 : index -// CHECK-DAG: %[[C32:.+]] = constant 32 : index -// CHECK-DAG: %[[C48:.+]] = constant 48 : index +// CHECK-DAG: %[[CST:.+]] = arith.constant 0.0 +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C16:.+]] = arith.constant 16 : index +// CHECK-DAG: %[[C32:.+]] = arith.constant 32 : index +// CHECK-DAG: %[[C48:.+]] = arith.constant 48 : index // CHECK-DAG: %[[ARG0:.+]] = hal.interface.binding.subspan @io::@arg0[%[[C0]]] // CHECK-DAG: %[[ARG1:.+]] = hal.interface.binding.subspan @io::@arg1[%[[C0]]] // CHECK-DAG: %[[RET0:.+]] = hal.interface.binding.subspan @io::@ret0[%[[C0]]] @@ -283,9 +283,9 @@ } builtin.module { func @matmul_static_shape() { - %c32 = constant 32 : index - %c4096 = constant 4096 : index - %c0 = constant 0 : index + %c32 = arith.constant 32 : index + %c4096 = arith.constant 4096 : index + %c0 = arith.constant 0 : index %0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<4096x4096xf16> %1 = hal.interface.binding.subspan @io::@arg1[%c0] : memref<4096x4096xf16> %2 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<4096x4096xf16> @@ -316,11 +316,11 @@ // PROMOTE-DAG: #[[MAP4:.+]] = affine_map<()[s0] -> (s0 * 64 - (s0 floordiv 2) * 128)> // PROMOTE: func @matmul_static_shape -// PROMOTE-DAG: %[[C0:.+]] = constant 0 : index -// PROMOTE-DAG: %[[C2:.+]] = constant 2 -// PROMOTE-DAG: %[[C16:.+]] = constant 16 -// PROMOTE-DAG: %[[C32:.+]] = constant 32 -// PROMOTE-DAG: %[[C48:.+]] = constant 48 +// PROMOTE-DAG: %[[C0:.+]] = arith.constant 0 : index +// PROMOTE-DAG: %[[C2:.+]] = arith.constant 2 +// PROMOTE-DAG: %[[C16:.+]] = arith.constant 16 +// PROMOTE-DAG: %[[C32:.+]] = arith.constant 32 +// PROMOTE-DAG: %[[C48:.+]] = arith.constant 48 // PROMOTE-DAG: %[[ALLOC1:.+]] = memref.alloc() : memref<128x32xf16, 3> // PROMOTE-DAG: %[[ALLOC2:.+]] = memref.alloc() : memref<32x128xf16, 3> // PROMOTE-DAG: %[[ARG0:.+]] = hal.interface.binding.subspan @io::@arg0[%[[C0]]] @@ -331,7 +331,7 @@ // PROMOTE: %[[WGMEM_LHS_SUBVIEW:.+]] = memref.subview %[[ALLOC1]][0, 0] [128, 32] [1, 1] // PROMOTE: %[[WGMEM_RHS_SUBVIEW:.+]] = memref.subview %[[ALLOC2]][0, 0] [32, 128] [1, 1] // PROMOTE: %[[SG_X:.+]] = gpu.subgroup_id -// PROMOTE: %[[SG_Y:.+]] = divi_signed %[[SG_X]], %[[C2]] +// PROMOTE: %[[SG_Y:.+]] = arith.divsi %[[SG_X]], %[[C2]] // PROMOTE: %[[SGOFFSET_Y:.+]] = affine.apply #[[MAP4]]()[%[[SG_Y]]] // PROMOTE: %[[SG_LHS_SUBVIEW:.+]] = memref.subview %[[WGMEM_LHS_SUBVIEW]][%[[SGOFFSET_Y]], 0] // PROMOTE: %[[SGOFFSET_X:.+]] = affine.apply #[[MAP4]]()[%[[SG_X]]]
diff --git a/iree/compiler/Codegen/Transforms/AffineMinCanonicalization.cpp b/iree/compiler/Codegen/Transforms/AffineMinCanonicalization.cpp index e2ac924..6fcea32 100644 --- a/iree/compiler/Codegen/Transforms/AffineMinCanonicalization.cpp +++ b/iree/compiler/Codegen/Transforms/AffineMinCanonicalization.cpp
@@ -37,21 +37,21 @@ SmallVectorImpl<Value> &dims, SmallVectorImpl<Value> &symbols) { MLIRContext *ctx = forOp.getContext(); - auto lbConstant = forOp.lowerBound().getDefiningOp<ConstantIndexOp>(); - AffineExpr lb = lbConstant ? getAffineConstantExpr(lbConstant.getValue(), ctx) + auto lbConstant = forOp.lowerBound().getDefiningOp<arith::ConstantIndexOp>(); + AffineExpr lb = lbConstant ? getAffineConstantExpr(lbConstant.value(), ctx) : getAffineDimExpr(dims.size(), ctx); - auto stepConstant = forOp.step().getDefiningOp<ConstantIndexOp>(); + auto stepConstant = forOp.step().getDefiningOp<arith::ConstantIndexOp>(); AffineExpr step = stepConstant - ? getAffineConstantExpr(stepConstant.getValue(), ctx) + ? getAffineConstantExpr(stepConstant.value(), ctx) : getAffineSymbolExpr(symbols.size(), ctx); if (!lbConstant) dims.push_back(forOp.lowerBound()); if (!stepConstant) symbols.push_back(forOp.step()); exprs.push_back(lb + step * getAffineDimExpr(dims.size(), ctx)); - auto ubConstant = forOp.upperBound().getDefiningOp<ConstantIndexOp>(); - AffineExpr ub = ubConstant ? getAffineConstantExpr(ubConstant.getValue(), ctx) + auto ubConstant = forOp.upperBound().getDefiningOp<arith::ConstantIndexOp>(); + AffineExpr ub = ubConstant ? getAffineConstantExpr(ubConstant.value(), ctx) : getAffineDimExpr(dims.size(), ctx); if (!ubConstant) dims.push_back(forOp.upperBound()); exprs.push_back(ub); @@ -136,7 +136,7 @@ })) return failure(); - rewriter.replaceOpWithNewOp<ConstantIndexOp>(minOp, min); + rewriter.replaceOpWithNewOp<arith::ConstantIndexOp>(minOp, min); return success(); }
diff --git a/iree/compiler/Codegen/Transforms/AffineMinDistributedSCFCanonicalization.cpp b/iree/compiler/Codegen/Transforms/AffineMinDistributedSCFCanonicalization.cpp index 5ca0812..d86c609 100644 --- a/iree/compiler/Codegen/Transforms/AffineMinDistributedSCFCanonicalization.cpp +++ b/iree/compiler/Codegen/Transforms/AffineMinDistributedSCFCanonicalization.cpp
@@ -77,8 +77,8 @@ } if (!ubDim) { - if (auto cstUb = ub.getDefiningOp<ConstantIndexOp>()) - ubDim = getAffineConstantExpr(cstUb.getValue(), minOp.getContext()); + if (auto cstUb = ub.getDefiningOp<arith::ConstantIndexOp>()) + ubDim = getAffineConstantExpr(cstUb.value(), minOp.getContext()); else return false; } @@ -128,7 +128,7 @@ /// %affine.min affine_map<(d0, d1) -> (N, d0 - d1)>(%ub, %iv) /// ``` /// With N a compile time constant. This operations can be replace by -/// `%cN = constant N : index` if we can prove that %lb, %step and %ub are +/// `%cN = arith.constant N : index` if we can prove that %lb, %step and %ub are /// divisible by N. static Optional<int64_t> foldAffineMin(AffineMinOp minOp) { AffineMap map = minOp.getAffineMap(); @@ -153,7 +153,8 @@ mlir::AffineMinOp minOp, mlir::PatternRewriter &rewriter) const override { Optional<int64_t> cst = foldAffineMin(minOp); if (!cst) return failure(); - rewriter.replaceOpWithNewOp<ConstantOp>(minOp, rewriter.getIndexAttr(*cst)); + rewriter.replaceOpWithNewOp<arith::ConstantOp>(minOp, + rewriter.getIndexAttr(*cst)); return failure(); } };
diff --git a/iree/compiler/Codegen/Transforms/Transforms.h b/iree/compiler/Codegen/Transforms/Transforms.h index e8270de..6fbc563 100644 --- a/iree/compiler/Codegen/Transforms/Transforms.h +++ b/iree/compiler/Codegen/Transforms/Transforms.h
@@ -40,7 +40,7 @@ /// %affine.min affine_map<(d0, d1) -> (N, d0 - d1)>(%ub, %iv) /// ``` /// With N a compile time constant. This operations can be replace by -/// `%cN = constant N : index` if we can prove that %lb, %step and %ub are +/// `%cN = arith.constant N : index` if we can prove that %lb, %step and %ub are /// divisible by N. void populateAffineMinSCFCanonicalizationPattern(RewritePatternSet &patterns);
diff --git a/iree/compiler/Codegen/Utils/Utils.cpp b/iree/compiler/Codegen/Utils/Utils.cpp index 1294900..2897b73 100644 --- a/iree/compiler/Codegen/Utils/Utils.cpp +++ b/iree/compiler/Codegen/Utils/Utils.cpp
@@ -436,7 +436,7 @@ /// Checks if the `forOp` is a tiled + distributed op. Looks for the op of this /// form /// ``` -/// %dim = constant ... : index +/// %dim = arith.constant ... : index /// %id = flow.dispatch.workgroup.id[%dim] /// %count = flow.dispatch.workgroup.count[%dim] /// %size = flow.dispatch.workgroup.size[%dim]
diff --git a/iree/compiler/Dialect/Flow/Conversion/TensorToFlow/ConvertTensorToFlow.cpp b/iree/compiler/Dialect/Flow/Conversion/TensorToFlow/ConvertTensorToFlow.cpp index b02d07c..3127412 100644 --- a/iree/compiler/Dialect/Flow/Conversion/TensorToFlow/ConvertTensorToFlow.cpp +++ b/iree/compiler/Dialect/Flow/Conversion/TensorToFlow/ConvertTensorToFlow.cpp
@@ -75,8 +75,8 @@ SmallVector<Value, 4> values; for (auto valueOrAttr : valueOrAttrList) { if (auto attr = valueOrAttr.dyn_cast<Attribute>()) { - values.push_back( - b.create<ConstantIndexOp>(loc, attr.cast<IntegerAttr>().getInt())); + values.push_back(b.create<arith::ConstantIndexOp>( + loc, attr.cast<IntegerAttr>().getInt())); } else { values.push_back(valueOrAttr.get<Value>()); } @@ -267,7 +267,8 @@ int64_t dimSize = !inputType.isDynamicDim(position) ? inputType.getDimSize(position) : resultType.getDimSize(position); - dimSizes[position] = rewriter.create<ConstantIndexOp>(loc, dimSize); + dimSizes[position] = + rewriter.create<arith::ConstantIndexOp>(loc, dimSize); } else { // Dynamic dim. dimSizes[position] = @@ -314,7 +315,7 @@ auto loc = op.getLoc(); SmallVector<Value> dimSizes(1); - dimSizes[0] = rewriter.create<ConstantIndexOp>(loc, 1); + dimSizes[0] = rewriter.create<arith::ConstantIndexOp>(loc, 1); rewriter.replaceOpWithNewOp<IREE::Flow::TensorSplatOp>( op, op.getType(), op.getOperand(0), dimSizes); return success();
diff --git a/iree/compiler/Dialect/Flow/Conversion/TensorToFlow/test/cast.mlir b/iree/compiler/Dialect/Flow/Conversion/TensorToFlow/test/cast.mlir index ef90610..9243c67 100644 --- a/iree/compiler/Dialect/Flow/Conversion/TensorToFlow/test/cast.mlir +++ b/iree/compiler/Dialect/Flow/Conversion/TensorToFlow/test/cast.mlir
@@ -1,7 +1,7 @@ // RUN: iree-opt -allow-unregistered-dialect -split-input-file -iree-flow-convert-to-flow-before-dispatch-formation %s | IreeFileCheck %s func @static_tensor_cast_to_dynamic(%arg0: tensor<4x4xf32>) -> tensor<?x?xf32> { - // CHECK-DAG: %[[C4:.*]] = constant 4 : index + // CHECK-DAG: %[[C4:.*]] = arith.constant 4 : index // CHECK-DAG: %[[RESULT:.*]] = flow.tensor.reshape %arg0 : tensor<4x4xf32> -> tensor<?x?xf32>{%[[C4]], %[[C4]]} // CHECK: return %[[RESULT]] %0 = tensor.cast %arg0 : tensor<4x4xf32> to tensor<?x?xf32> @@ -10,7 +10,7 @@ // ----- func @dynamic_tensor_cast_to_static(%arg0: tensor<?xf32>) -> tensor<4xf32> { - // CHECK: %[[C4:.*]] = constant 4 : index + // CHECK: %[[C4:.*]] = arith.constant 4 : index // CHECK: %[[RESULT:.*]] = flow.tensor.reshape %arg0 : tensor<?xf32>{%[[C4]]} -> tensor<4xf32> // CHECK: return %[[RESULT]] %0 = tensor.cast %arg0 : tensor<?xf32> to tensor<4xf32> @@ -19,9 +19,9 @@ // ----- func @dynamic_tensor_cast_to_dynamic(%arg0: tensor<?x?xf32>) -> tensor<?x3xf32> { - // CHECK-DAG: %[[C0:.*]] = constant 0 : index + // CHECK-DAG: %[[C0:.*]] = arith.constant 0 : index // CHECK-DAG: %[[D0:.*]] = tensor.dim %arg0, %[[C0]] : tensor<?x?xf32> - // CHECK-DAG: %[[C3:.*]] = constant 3 : index + // CHECK-DAG: %[[C3:.*]] = arith.constant 3 : index // CHECK: %[[RESULT:.*]] = flow.tensor.reshape %arg0 : tensor<?x?xf32>{%[[D0]], %[[C3]]} -> tensor<?x3xf32>{%[[D0]]} // CHECK: return %[[RESULT]] %0 = tensor.cast %arg0 : tensor<?x?xf32> to tensor<?x3xf32> @@ -30,7 +30,7 @@ // ----- func @tensor_cast_within_dispatch_workgroups_not_converted() -> tensor<f32> { - %x = constant 100 : index + %x = arith.constant 100 : index %0 = flow.dispatch.workgroups[%x]() : () -> (tensor<f32>) = () { // CHECK: = tensor.cast %[[source:.+]] : tensor<4x4xf32> to tensor<?x?xf32> %1 = "test.source"() : () -> (tensor<4x4xf32>)
diff --git a/iree/compiler/Dialect/Flow/Conversion/TensorToFlow/test/extract_slice.mlir b/iree/compiler/Dialect/Flow/Conversion/TensorToFlow/test/extract_slice.mlir index dc5b0d2..cf4e44b 100644 --- a/iree/compiler/Dialect/Flow/Conversion/TensorToFlow/test/extract_slice.mlir +++ b/iree/compiler/Dialect/Flow/Conversion/TensorToFlow/test/extract_slice.mlir
@@ -7,10 +7,10 @@ } // CHECK-LABEL: func @extract_slice1( // CHECK-SAME: %[[ARG0:.+]]: tensor<5x24x48xf32>) -// CHECK-DAG: %[[C2:.+]] = constant 2 : index -// CHECK-DAG: %[[C3:.+]] = constant 3 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index -// CHECK-DAG: %[[C4:.+]] = constant 4 : index +// CHECK-DAG: %[[C2:.+]] = arith.constant 2 : index +// CHECK-DAG: %[[C3:.+]] = arith.constant 3 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index +// CHECK-DAG: %[[C4:.+]] = arith.constant 4 : index // CHECK: %[[SLICE:.+]] = flow.tensor.slice %[[ARG0]][%[[C2]], %[[C3]], %[[C4]] for %[[C1]], %[[C1]], %[[C4]]] // CHECK: %[[RESULT:.+]] = flow.tensor.reshape %[[SLICE]] // CHECK: return %[[RESULT]] @@ -24,11 +24,11 @@ } // CHECK-LABEL: func @extract_slice2 // CHECK-SAME: %[[ARG0:.+]]: tensor<5x24x48xf32>) -// CHECK-DAG: %[[C3:.+]] = constant 3 : index -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index -// CHECK-DAG: %[[C2:.+]] = constant 2 : index -// CHECK-DAG: %[[C48:.+]] = constant 48 : index +// CHECK-DAG: %[[C3:.+]] = arith.constant 3 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index +// CHECK-DAG: %[[C2:.+]] = arith.constant 2 : index +// CHECK-DAG: %[[C48:.+]] = arith.constant 48 : index // CHECK: %[[SLICE:.+]] = flow.tensor.slice %[[ARG0]][%[[C2]], %[[C3]], %[[C0]] for %[[C1]], %[[C2]], %[[C48]]] // CHECK: %[[RESULT:.+]] = flow.tensor.reshape %[[SLICE]] // CHECK: return %[[RESULT]] @@ -83,11 +83,11 @@ // CHECK-LABEL: func @extract_slice7( // CHECK-SAME: %[[ARG0:.+]]: tensor<5x?x48xf32> // CHECK-SAME: %[[ARG1:.+]]: index) -// CHECK-DAG: %[[C2:.+]] = constant 2 : index -// CHECK-DAG: %[[C3:.+]] = constant 3 : index -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index -// CHECK-DAG: %[[C48:.+]] = constant 48 : index +// CHECK-DAG: %[[C2:.+]] = arith.constant 2 : index +// CHECK-DAG: %[[C3:.+]] = arith.constant 3 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index +// CHECK-DAG: %[[C48:.+]] = arith.constant 48 : index // CHECK-DAG: %[[DIM:.+]] = tensor.dim %[[ARG0]], %[[C1]] : tensor<5x?x48xf32> // CHECK: %[[SLICE:.+]] = flow.tensor.slice %[[ARG0]][%[[C2]], %[[C3]], %[[C0]] for %[[C1]], %[[C2]], %[[C48]]] // CHECK: %[[RESULT:.+]] = flow.tensor.reshape %[[SLICE]] @@ -101,10 +101,10 @@ } // CHECK-LABEL: func @rank_reducing_extract_slice // CHECK-SAME: %[[ARG0:[a-zA-Z0-9_]+]] -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index -// CHECK-DAG: %[[C4:.+]] = constant 4 : index -// CHECK-DAG: %[[C513:.+]] = constant 513 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index +// CHECK-DAG: %[[C4:.+]] = arith.constant 4 : index +// CHECK-DAG: %[[C513:.+]] = arith.constant 513 : index // CHECK: %[[DIM:.+]] = tensor.dim %[[ARG0]], %[[C0]] // CHECK: %[[SLICE:.+]] = flow.tensor.slice %[[ARG0]] // CHECK-SAME: [%[[C4]], %[[C0]] for %[[C1]], %[[C513]]] @@ -120,17 +120,17 @@ return %0 : tensor<49x20xf32> } // CHECK-LABEL: func @rank_reducing_extract_slice_trailing_unit_dims -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index -// CHECK-DAG: %[[C49:.+]] = constant 49 : index -// CHECK-DAG: %[[C20:.+]] = constant 20 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index +// CHECK-DAG: %[[C49:.+]] = arith.constant 49 : index +// CHECK-DAG: %[[C20:.+]] = arith.constant 20 : index // CHECK: %[[extract_slice:.+]] = flow.tensor.slice %{{.+}}[%[[C0]], %[[C1]], %[[C0]], %[[C0]] for %[[C1]], %[[C49]], %[[C20]], %[[C1]]] : tensor<1x50x20x1xf32> -> tensor<1x49x20x1xf32> // CHECK: flow.tensor.reshape %[[extract_slice]] : tensor<1x49x20x1xf32> -> tensor<49x20xf32> // ----- func @extract_slice_within_dispatch_workgroups_not_converted() -> tensor<f32> { - %x = constant 100 : index + %x = arith.constant 100 : index %0 = flow.dispatch.workgroups[%x]() : () -> (tensor<f32>) = () { // CHECK: = tensor.extract_slice %[[source:.+]][2, 3, 4] [1, 1, 4] [1, 1, 1] : tensor<5x24x48xf32> to tensor<4xf32> %1 = "test.source"() : () -> (tensor<5x24x48xf32>)
diff --git a/iree/compiler/Dialect/Flow/Conversion/TensorToFlow/test/from_elements.mlir b/iree/compiler/Dialect/Flow/Conversion/TensorToFlow/test/from_elements.mlir index a2953b8..2537227 100644 --- a/iree/compiler/Dialect/Flow/Conversion/TensorToFlow/test/from_elements.mlir +++ b/iree/compiler/Dialect/Flow/Conversion/TensorToFlow/test/from_elements.mlir
@@ -12,8 +12,8 @@ // ----- // CHECK: func @tensor.from_elements__not_convertible(%[[arg0:.*]]: i8) func @tensor.from_elements__not_convertible(%arg0: i8) -> (i8) { - // CHECK: %[[c0:.*]] = constant 0 - %c0 = constant 0 : index + // CHECK: %[[c0:.*]] = arith.constant 0 + %c0 = arith.constant 0 : index // CHECK: %[[res:.*]] = tensor.from_elements %[[arg0]], %[[arg0]] : tensor<2xi8> %0 = tensor.from_elements %arg0, %arg0 : tensor<2xi8> // CHECK: flow.tensor.load %[[res]][%[[c0]]] @@ -23,7 +23,7 @@ // ----- func @tensor.from_elements__within_dispatch_workgroups_not_converted() -> tensor<f32> { - %x = constant 100 : index + %x = arith.constant 100 : index %0 = flow.dispatch.workgroups[%x]() : () -> (tensor<f32>) = () { // CHECK: = tensor.from_elements %[[source:.+]] : tensor<1xi8> %1 = "test.source"() : () -> (i8)
diff --git a/iree/compiler/Dialect/Flow/Conversion/TensorToFlow/test/insert_slice.mlir b/iree/compiler/Dialect/Flow/Conversion/TensorToFlow/test/insert_slice.mlir index 70f0fff..5d08aff 100644 --- a/iree/compiler/Dialect/Flow/Conversion/TensorToFlow/test/insert_slice.mlir +++ b/iree/compiler/Dialect/Flow/Conversion/TensorToFlow/test/insert_slice.mlir
@@ -3,7 +3,7 @@ func @insert_slice_convert (%arg0 : tensor<?x24x48xf32>, %arg1 : tensor<1x4x48xf32>) -> tensor<?x24x48xf32> { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %0 = tensor.insert_slice %arg1 into %arg0[4, 2, 0] [1, 4, 48] [1, 1, 1] : tensor<1x4x48xf32> into tensor<?x24x48xf32> return %0 : tensor<?x24x48xf32> @@ -11,9 +11,9 @@ // CHECK-LABEL: func @insert_slice_convert // CHECK-SAME: %[[ARG0:[a-zA-Z0-9_]+]] // CHECK-SAME: %[[ARG1:[a-zA-Z0-9_]+]] -// CHECK-DAG: %[[C0:.+]] = constant 0 -// CHECK-DAG: %[[C2:.+]] = constant 2 -// CHECK-DAG: %[[C4:.+]] = constant 4 +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 +// CHECK-DAG: %[[C2:.+]] = arith.constant 2 +// CHECK-DAG: %[[C4:.+]] = arith.constant 4 // CHECK-DAG: %[[DIM0:.+]] = tensor.dim %[[ARG0]], %[[C0]] // CHECK: %[[UPDATE:.+]] = flow.tensor.update %[[ARG1]], %[[ARG0]][%[[C4]], %[[C2]], %[[C0]]] // CHECK-SAME: : tensor<1x4x48xf32> -> %[[ARG0]] as tensor<?x24x48xf32>{%[[DIM0]]} @@ -23,7 +23,7 @@ func @insert_slice_convert_rank_reducing (%arg0 : tensor<?x24x48xf32>, %arg1 : tensor<4x48xf32>) -> tensor<?x24x48xf32> { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %0 = tensor.insert_slice %arg1 into %arg0[4, 2, 0] [1, 4, 48] [1, 1, 1] : tensor<4x48xf32> into tensor<?x24x48xf32> return %0 : tensor<?x24x48xf32> @@ -31,9 +31,9 @@ // CHECK-LABEL: func @insert_slice_convert_rank_reducing // CHECK-SAME: %[[ARG0:[a-zA-Z0-9_]+]] // CHECK-SAME: %[[ARG1:[a-zA-Z0-9_]+]] -// CHECK-DAG: %[[C0:.+]] = constant 0 -// CHECK-DAG: %[[C2:.+]] = constant 2 -// CHECK-DAG: %[[C4:.+]] = constant 4 +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 +// CHECK-DAG: %[[C2:.+]] = arith.constant 2 +// CHECK-DAG: %[[C4:.+]] = arith.constant 4 // CHECK-DAG: %[[RESHAPE:.+]] = flow.tensor.reshape %[[ARG1]] : tensor<4x48xf32> -> tensor<1x4x48xf32> // CHECK-DAG: %[[DIM:.+]] = tensor.dim %[[ARG0]], %[[C0]] // CHECK: %[[UPDATE:.+]] = flow.tensor.update %[[RESHAPE]], %[[ARG0]][%[[C4]], %[[C2]], %[[C0]]] @@ -47,8 +47,8 @@ return %0 : tensor<1x50x20x1xf32> } // CHECK-LABEL: func @rank_reducing_insert_slice_trailing_unit_dims -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK: %[[RESHAPE:.+]] = flow.tensor.reshape %{{.+}} : tensor<49x20xf32> -> tensor<1x49x20x1xf32> // CHECK: flow.tensor.update %[[RESHAPE]], %{{.+}}[%[[C0]], %[[C1]], %[[C0]], %[[C0]]] : tensor<1x49x20x1xf32> -> %{{.+}} as tensor<1x50x20x1xf32> @@ -56,7 +56,7 @@ // ----- func @insert_slice_within_dispatch_workgroups_not_converted() -> tensor<f32> { - %x = constant 100 : index + %x = arith.constant 100 : index %0 = flow.dispatch.workgroups[%x]() : () -> (tensor<f32>) = () { // CHECK: = tensor.insert_slice %[[source2:.+]] into %[[source1:.+]][4, 2, 0] [1, 4, 48] [1, 1, 1] : tensor<1x4x48xf32> into tensor<?x24x48xf32> %1 = "test.source1"() : () -> (tensor<?x24x48xf32>)
diff --git a/iree/compiler/Dialect/Flow/IR/FlowDialect.cpp b/iree/compiler/Dialect/Flow/IR/FlowDialect.cpp index 60fa352..8e33e36 100644 --- a/iree/compiler/Dialect/Flow/IR/FlowDialect.cpp +++ b/iree/compiler/Dialect/Flow/IR/FlowDialect.cpp
@@ -75,8 +75,8 @@ Operation *FlowDialect::materializeConstant(OpBuilder &builder, Attribute value, Type type, Location loc) { - if (ConstantOp::isBuildableWith(value, type)) - return builder.create<ConstantOp>(loc, type, value); + if (arith::ConstantOp::isBuildableWith(value, type)) + return builder.create<arith::ConstantOp>(loc, type, value); return nullptr; }
diff --git a/iree/compiler/Dialect/Flow/IR/FlowOpFolders.cpp b/iree/compiler/Dialect/Flow/IR/FlowOpFolders.cpp index ce99460..863b089 100644 --- a/iree/compiler/Dialect/Flow/IR/FlowOpFolders.cpp +++ b/iree/compiler/Dialect/Flow/IR/FlowOpFolders.cpp
@@ -73,7 +73,7 @@ // Dimension has changed to be dynamic; insert a constant to use. // This sometimes happens during folding of casts and usually is cleaned // up pretty quickly. - newDims.push_back(rewriter.createOrFold<ConstantIndexOp>( + newDims.push_back(rewriter.createOrFold<arith::ConstantIndexOp>( op->getLoc(), oldShapedType.getDimSize(i))); } } @@ -287,7 +287,7 @@ *constantIndex); } else { // Tensor tile : if (loadOp.getMixedSizes()[*constantIndex].is<Attribute>()) { - rewriter.replaceOpWithNewOp<ConstantOp>( + rewriter.replaceOpWithNewOp<arith::ConstantOp>( op, loadOp.getMixedSizes()[*constantIndex] .get<Attribute>() .dyn_cast<IntegerAttr>());
diff --git a/iree/compiler/Dialect/Flow/IR/FlowOps.cpp b/iree/compiler/Dialect/Flow/IR/FlowOps.cpp index 4b51012..cee79f4 100644 --- a/iree/compiler/Dialect/Flow/IR/FlowOps.cpp +++ b/iree/compiler/Dialect/Flow/IR/FlowOps.cpp
@@ -298,8 +298,8 @@ // Inline operations that the dispatch region can handle natively. static bool canDispatchRegionContainOp(Operation *op) { // Inline constant operations that are splat or small constants. - if (auto constantOp = dyn_cast<ConstantOp>(op)) { - auto constantValueAttr = constantOp.getValue(); + if (auto constantOp = dyn_cast<arith::ConstantOp>(op)) { + auto constantValueAttr = constantOp.value(); auto constantType = constantOp.getType(); if (constantValueAttr.isa<SplatElementsAttr>()) { return true; @@ -916,7 +916,7 @@ bool ExStreamFragmentOp::canClosureContainOp(Operation *op) { // NOTE: we widen support on new stream ops only - the legacy path isn't worth // upgrading to support more. - if (auto constantOp = dyn_cast<ConstantOp>(op)) { + if (auto constantOp = dyn_cast<arith::ConstantOp>(op)) { return constantOp.getType().isIntOrIndexOrFloat(); } if (auto loadOp = dyn_cast<IREE::Util::GlobalLoadOp>(op)) {
diff --git a/iree/compiler/Dialect/Flow/IR/test/dispatch_ops.mlir b/iree/compiler/Dialect/Flow/IR/test/dispatch_ops.mlir index 8250528..f7416b2 100644 --- a/iree/compiler/Dialect/Flow/IR/test/dispatch_ops.mlir +++ b/iree/compiler/Dialect/Flow/IR/test/dispatch_ops.mlir
@@ -11,8 +11,8 @@ // CHECK-LABEL: @dispatch func @dispatch(%arg0 : tensor<4xf32>) -> tensor<4xf32> { - // CHECK: %[[CST:.+]] = constant - %cst = constant 4 : index + // CHECK: %[[CST:.+]] = arith.constant + %cst = arith.constant 4 : index // CHECK: %0 = flow.dispatch @ex0::@dispatch_fn[%[[CST]]](%[[CST]], %arg0) : (index, tensor<4xf32>) -> tensor<4xf32> %0 = flow.dispatch @ex0::@dispatch_fn[%cst](%cst, %arg0) : (index, tensor<4xf32>) -> tensor<4xf32> return %0 : tensor<4xf32> @@ -22,8 +22,8 @@ // CHECK-LABEL: @inplaceDispatch func @inplaceDispatch(%arg0 : tensor<4xf32>, %arg1 : tensor<8xf32>) -> (tensor<4xf32>, tensor<8xf32>) { - // CHECK: %[[CST:.+]] = constant - %cst = constant 4 : index + // CHECK: %[[CST:.+]] = arith.constant + %cst = arith.constant 4 : index // CHECK: %0:2 = flow.dispatch @ex0::@dispatch_fn[%[[CST]]](%[[CST]], %arg0, %arg1) : (index, tensor<4xf32>, tensor<8xf32>) -> (%arg0, %arg1) %0, %1 = flow.dispatch @ex0::@dispatch_fn[%cst](%cst, %arg0, %arg1) : (index, tensor<4xf32>, tensor<8xf32>) -> (%arg0, %arg1) return %0, %1 : tensor<4xf32>, tensor<8xf32> @@ -33,12 +33,12 @@ // CHECK-LABEL: @inplaceDynamicDispatch func @inplaceDynamicDispatch(%arg0 : tensor<4x?xf32>, %arg1 : tensor<8x?xf32>) -> (tensor<4x?xf32>, tensor<8x?xf32>) { - // CHECK-DAG: %[[CST:.+]] = constant 4 - %cst = constant 4 : index - // CHECK-DAG: %[[DIM0:.+]] = constant 100 - %dim0 = constant 100 : index - // CHECK-DAG: %[[DIM1:.+]] = constant 200 - %dim1 = constant 200 : index + // CHECK-DAG: %[[CST:.+]] = arith.constant 4 + %cst = arith.constant 4 : index + // CHECK-DAG: %[[DIM0:.+]] = arith.constant 100 + %dim0 = arith.constant 100 : index + // CHECK-DAG: %[[DIM1:.+]] = arith.constant 200 + %dim1 = arith.constant 200 : index // CHECK: %0:2 = flow.dispatch @ex0::@dispatch_fn[%[[CST]]](%[[CST]], %arg0, %arg1) : (index, tensor<4x?xf32>{%[[DIM0]]}, tensor<8x?xf32>{%[[DIM1]]}) -> (%arg0{%[[DIM1]]}, %arg1{%[[DIM0]]}) %0, %1 = flow.dispatch @ex0::@dispatch_fn[%cst](%cst, %arg0, %arg1) : (index, tensor<4x?xf32>{%dim0}, tensor<8x?xf32>{%dim1}) -> (%arg0{%dim1}, %arg1{%dim0}) return %0, %1 : tensor<4x?xf32>, tensor<8x?xf32> @@ -49,10 +49,10 @@ // CHECK-LABEL: @inplaceTypeChange // CHECK-SAME: (%[[ARG0:.+]]: tensor<4x?xf32>) func @inplaceTypeChange(%arg0: tensor<4x?xf32>) -> tensor<?x4xf32> { - // CHECK-DAG: %[[CST:.+]] = constant 4 - %cst = constant 4 : index - // CHECK-DAG: %[[DIM0:.+]] = constant 100 - %dim0 = constant 100 : index + // CHECK-DAG: %[[CST:.+]] = arith.constant 4 + %cst = arith.constant 4 : index + // CHECK-DAG: %[[DIM0:.+]] = arith.constant 100 + %dim0 = arith.constant 100 : index // CHECK: %0 = flow.dispatch @ex0::@dispatch_fn[%[[CST]]](%[[ARG0]]) : (tensor<4x?xf32>{%[[DIM0]]}) -> %arg0 as tensor<?x4xf32>{%[[DIM0]]} %0 = flow.dispatch @ex0::@dispatch_fn[%cst](%arg0) : (tensor<4x?xf32>{%dim0}) -> %arg0 as tensor<?x4xf32>{%dim0} return %0 : tensor<?x4xf32>
diff --git a/iree/compiler/Dialect/Flow/IR/test/dispatch_tensor_folding.mlir b/iree/compiler/Dialect/Flow/IR/test/dispatch_tensor_folding.mlir index a29ce18..77364e7 100644 --- a/iree/compiler/Dialect/Flow/IR/test/dispatch_tensor_folding.mlir +++ b/iree/compiler/Dialect/Flow/IR/test/dispatch_tensor_folding.mlir
@@ -1,9 +1,9 @@ // RUN: iree-opt -allow-unregistered-dialect -split-input-file -canonicalize %s | IreeFileCheck %s func @canonicalizeStaticOperands(%arg0: !flow.dispatch.tensor<readonly:4x4xf32>) { - %c0 = constant 0 : index - %c1 = constant 1 : index - %c2 = constant 2 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + %c2 = arith.constant 2 : index %0 = flow.dispatch.tensor.load %arg0, offsets=[%c0, %c0], sizes=[%c2, %c2], strides=[%c1, %c1] : !flow.dispatch.tensor<readonly:4x4xf32> -> tensor<?x?xf32> "test.sink"(%0) : (tensor<?x?xf32>) -> () return @@ -20,9 +20,9 @@ // ----- func @canonicalizePartiallyStaticOperands(%arg0: !flow.dispatch.tensor<readonly:4x4xf32>, %offset: index, %size: index, %stride: index) { - %c0 = constant 0 : index - %c1 = constant 1 : index - %c2 = constant 2 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + %c2 = arith.constant 2 : index %0 = flow.dispatch.tensor.load %arg0, offsets=[%offset, %c0], sizes=[%size, %c2], strides=[%stride, %c1] : !flow.dispatch.tensor<readonly:4x4xf32> -> tensor<?x?xf32> "test.sink"(%0) : (tensor<?x?xf32>) -> () return @@ -40,7 +40,7 @@ // ----- func @canonicalizeDimOfTensorTile(%arg0: !flow.dispatch.tensor<readonly:250x1024xf32>, %arg1 : index, %arg2: index) { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %0 = affine.min affine_map<(d0) -> (64, -d0 + 250)>(%arg1) %1 = flow.dispatch.tensor.load %arg0, offsets = [%arg2, 0], sizes = [%0, 1024], strides = [1, 1] : !flow.dispatch.tensor<readonly:250x1024xf32> -> tensor<?x1024xf32> %2 = tensor.dim %1, %c0 : tensor<?x1024xf32>
diff --git a/iree/compiler/Dialect/Flow/IR/test/dispatch_workgroups.mlir b/iree/compiler/Dialect/Flow/IR/test/dispatch_workgroups.mlir index 9231857..01dd874 100644 --- a/iree/compiler/Dialect/Flow/IR/test/dispatch_workgroups.mlir +++ b/iree/compiler/Dialect/Flow/IR/test/dispatch_workgroups.mlir
@@ -6,11 +6,11 @@ %arg0 : tensor<?x4xf32>, // CHECK-SAME: %[[ARG1:.+]]: index %arg1 : index) -> tensor<4x?xf32> { - %c128 = constant 128 : index - // CHECK-DAG: %[[WORKGROUP_COUNT_X:.+]] = constant 100 - %x = constant 100 : index - // CHECK-DAG: %[[WORKGROUP_COUNT_Y:.+]] = constant 50 - %y = constant 50 : index + %c128 = arith.constant 128 : index + // CHECK-DAG: %[[WORKGROUP_COUNT_X:.+]] = arith.constant 100 + %x = arith.constant 100 : index + // CHECK-DAG: %[[WORKGROUP_COUNT_Y:.+]] = arith.constant 50 + %y = arith.constant 50 : index // CHECK: %[[OUTER_RET0:.+]] = flow.dispatch.workgroups[ // CHECK-SAME: %[[WORKGROUP_COUNT_X]], %[[WORKGROUP_COUNT_Y]] // CHECK-SAME: ](%[[ARG0]], %[[ARG1]]) @@ -91,11 +91,11 @@ %arg0: tensor<?x4xf32>, // CHECK-SAME: %[[ARG1:.+]]: index %arg1: index) -> tensor<?x4xf32> { - %c128 = constant 128 : index - // CHECK-DAG: %[[WORKGROUP_COUNT_X:.+]] = constant 100 - %x = constant 100 : index - // CHECK-DAG: %[[WORKGROUP_COUNT_Y:.+]] = constant 50 - %y = constant 50 : index + %c128 = arith.constant 128 : index + // CHECK-DAG: %[[WORKGROUP_COUNT_X:.+]] = arith.constant 100 + %x = arith.constant 100 : index + // CHECK-DAG: %[[WORKGROUP_COUNT_Y:.+]] = arith.constant 50 + %y = arith.constant 50 : index // CHECK: %[[OUTER_RET0:.+]] = flow.dispatch.workgroups[ // CHECK-SAME: %[[WORKGROUP_COUNT_X]], %[[WORKGROUP_COUNT_Y]] // CHECK-SAME: ](%[[ARG0]], %[[ARG1]])
diff --git a/iree/compiler/Dialect/Flow/IR/test/dispatch_workgroups_folding.mlir b/iree/compiler/Dialect/Flow/IR/test/dispatch_workgroups_folding.mlir index d2999fe..8cd80cf 100644 --- a/iree/compiler/Dialect/Flow/IR/test/dispatch_workgroups_folding.mlir +++ b/iree/compiler/Dialect/Flow/IR/test/dispatch_workgroups_folding.mlir
@@ -2,9 +2,9 @@ // CHECK-LABEL: @workgroupStaticShapeDims func @workgroupStaticShapeDims(%arg0 : tensor<?x4xf32>) -> tensor<4x?xf32> { - %c128 = constant 128 : index - %x = constant 100 : index - %y = constant 50 : index + %c128 = arith.constant 128 : index + %x = arith.constant 100 : index + %y = arith.constant 50 : index // CHECK: flow.dispatch.workgroups %0 = flow.dispatch.workgroups[%x, %y](%arg0) : (tensor<?x4xf32>{%c128}) -> (tensor<4x?xf32>{%c128}) = ( // CHECK-NEXT: (%[[ARG0:.+]]: !flow.dispatch.tensor<readonly:?x4xf32>, @@ -12,7 +12,7 @@ // CHECK-SAME: %[[RET0:.+]]: !flow.dispatch.tensor<writeonly:4x?xf32>) %ret0: !flow.dispatch.tensor<writeonly:4x?xf32> ) { - // CHECK: %[[DIM_4:.+]] = constant 4 : index + // CHECK: %[[DIM_4:.+]] = arith.constant 4 : index // CHECK: %[[ARG0_SHAPE:.+]] = flow.dispatch.shape %[[ARG0]] %arg0_shape = flow.dispatch.shape %arg0_capture : !flow.dispatch.tensor<readonly:?x4xf32> -> !shapex.ranked_shape<[?,4]> @@ -39,15 +39,15 @@ // CHECK-LABEL: @workgroupRankFolding func @workgroupRankFolding(%arg0 : tensor<?x4xf32>) -> tensor<4x?xf32> { - %c128 = constant 128 : index - %x = constant 100 : index - %y = constant 50 : index + %c128 = arith.constant 128 : index + %x = arith.constant 100 : index + %y = arith.constant 50 : index // CHECK: flow.dispatch.workgroups %0 = flow.dispatch.workgroups[%x, %y](%arg0) : (tensor<?x4xf32>{%c128}) -> (tensor<4x?xf32>{%c128}) = ( %arg0_capture: !flow.dispatch.tensor<readonly:?x4xf32>, %ret0: !flow.dispatch.tensor<writeonly:4x?xf32> ) { - // CHECK: %[[RANK:.+]] = constant 2 : index + // CHECK: %[[RANK:.+]] = arith.constant 2 : index %workgroup_rank = flow.dispatch.workgroup.rank : index // CHECK-NEXT: "test.sink"(%[[RANK]]) "test.sink"(%workgroup_rank) : (index) -> () @@ -65,7 +65,7 @@ // CHECK-NEXT: %[[DIM:.*]] = shapex.ranked_dim %[[RANKED_SHAPE]][0] // CHECK-NEXT: "test.sink"(%[[DIM]]) : (index) -> () %tensor = flow.dispatch.tensor.load %arg0, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:?xf32> -> tensor<?xf32> - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %dim = tensor.dim %tensor, %c0 : tensor<?xf32> "test.sink"(%dim) : (index) -> () return @@ -77,18 +77,18 @@ // CHECK-SAME: (%[[ARG0:.+]]: tensor<1x4xf32>) func @inlineWithTiedResults1(%arg0: tensor<1x4xf32>) -> tensor<1x4xf32> { // CHECK-NOT: constant 128 - %cst = constant 128 : index - // CHECK-DAG: %[[X:.+]] = constant 100 - %x = constant 100 : index - // CHECK-DAG: %[[Y:.+]] = constant 50 - %y = constant 50 : index + %cst = arith.constant 128 : index + // CHECK-DAG: %[[X:.+]] = arith.constant 100 + %x = arith.constant 100 : index + // CHECK-DAG: %[[Y:.+]] = arith.constant 50 + %y = arith.constant 50 : index // CHECK: flow.dispatch.workgroups[%[[X]], %[[Y]]](%[[ARG0]]) : (tensor<1x4xf32>) -> %[[ARG0]] = // CHECK-NEXT: (%[[ARG0_INNER:.+]]: !flow.dispatch.tensor<readwrite:1x4xf32>) %0 = flow.dispatch.workgroups[%x, %y](%cst, %arg0) : (index, tensor<1x4xf32>) -> %arg0 = ( %cst_capture: index, %arg0_capture: !flow.dispatch.tensor<readwrite:1x4xf32> ) { - // CHECK: %[[INLINED_CST:.+]] = constant 128 : index + // CHECK: %[[INLINED_CST:.+]] = arith.constant 128 : index // CHECK-NEXT: "test.sink"(%[[INLINED_CST]]) "test.sink"(%cst_capture) : (index) -> () // CHECK-NEXT: "test.sink"(%[[ARG0_INNER]]) @@ -104,18 +104,18 @@ // CHECK-SAME: (%[[ARG0:.+]]: tensor<1x4xf32>) func @inlineWithTiedResults2(%arg0: tensor<1x4xf32>) -> tensor<1x4xf32> { // CHECK-NOT: constant 128 - %cst = constant 128 : index - // CHECK-DAG: %[[X:.+]] = constant 100 - %x = constant 100 : index - // CHECK-DAG: %[[Y:.+]] = constant 50 - %y = constant 50 : index + %cst = arith.constant 128 : index + // CHECK-DAG: %[[X:.+]] = arith.constant 100 + %x = arith.constant 100 : index + // CHECK-DAG: %[[Y:.+]] = arith.constant 50 + %y = arith.constant 50 : index // CHECK: flow.dispatch.workgroups[%[[X]], %[[Y]]](%[[ARG0]]) : (tensor<1x4xf32>) -> %[[ARG0]] = // CHECK-NEXT: (%[[ARG0_INNER:.+]]: !flow.dispatch.tensor<readwrite:1x4xf32>) %0 = flow.dispatch.workgroups[%x, %y](%arg0, %cst) : (tensor<1x4xf32>, index) -> %arg0 = ( %arg0_capture: !flow.dispatch.tensor<readwrite:1x4xf32>, %cst_capture: index ) { - // CHECK: %[[INLINED_CST:.+]] = constant 128 : index + // CHECK: %[[INLINED_CST:.+]] = arith.constant 128 : index // CHECK-NEXT: "test.sink"(%[[INLINED_CST]]) "test.sink"(%cst_capture) : (index) -> () // CHECK-NEXT: "test.sink"(%[[ARG0_INNER]]) @@ -130,10 +130,10 @@ // CHECK-LABEL: func @dontInlineReadWrite // CHECK-SAME: (%[[ARG0:.+]]: tensor<1x4xf32>) func @dontInlineReadWrite(%arg0: tensor<1x4xf32>) -> tensor<4x8xf32> { - // CHECK: %[[CST:.+]] = constant dense<0.000000e+00> : tensor<4x8xf32> - %cst = constant dense<0.0> : tensor<4x8xf32> - %x = constant 100 : index - %y = constant 50 : index + // CHECK: %[[CST:.+]] = arith.constant dense<0.000000e+00> : tensor<4x8xf32> + %cst = arith.constant dense<0.0> : tensor<4x8xf32> + %x = arith.constant 100 : index + %y = arith.constant 50 : index // CHECK: flow.dispatch.workgroups[{{.+}}](%[[ARG0]], %[[CST]]) : (tensor<1x4xf32>, tensor<4x8xf32>) -> %cst // CHECK-NEXT: (%{{.+}}: !flow.dispatch.tensor<readonly:1x4xf32>, %{{.+}}: !flow.dispatch.tensor<readwrite:4x8xf32>) %0 = flow.dispatch.workgroups[%x, %y](%arg0, %cst) : (tensor<1x4xf32>, tensor<4x8xf32>) -> %cst = ( @@ -153,15 +153,15 @@ // CHECK-LABEL: func @remove_unused_result func @remove_unused_result(%arg0 : tensor<9xi32>, %arg1 : tensor<9xi32>) -> (tensor<i32>) { - %c1 = constant 1 : index + %c1 = arith.constant 1 : index // CHECK: flow.dispatch.workgroups[%c1, %c1, %c1]() : () -> tensor<i32> = // CHECK-NEXT: (%{{.+}}: !flow.dispatch.tensor<writeonly:i32>) // CHECK: flow.dispatch.tensor.store // CHECK-NOT: flow.dispatch.tensor.store %0:2 = flow.dispatch.workgroups[%c1, %c1, %c1](%arg0, %arg1) : (tensor<9xi32>, tensor<9xi32>) -> (tensor<i32>, tensor<i32>) = (%arg0: !flow.dispatch.tensor<readonly:9xi32>, %arg1: !flow.dispatch.tensor<readonly:9xi32>, %arg2: !flow.dispatch.tensor<writeonly:i32>, %arg3: !flow.dispatch.tensor<writeonly:i32>) { - %c0_i32 = constant 0 : i32 - %c-2147483648_i32 = constant -2147483648 : i32 + %c0_i32 = arith.constant 0 : i32 + %c-2147483648_i32 = arith.constant -2147483648 : i32 %0 = flow.dispatch.tensor.load %arg0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:9xi32> -> tensor<9xi32> %1 = flow.dispatch.tensor.load %arg1, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:9xi32> -> tensor<9xi32> %2 = linalg.init_tensor [] : tensor<i32> @@ -178,15 +178,15 @@ // CHECK-LABEL: func @remove_unused_read_write_result func @remove_unused_read_write_result(%arg0 : tensor<9xi32>, %arg1 : tensor<9xi32>) -> (tensor<i32>) { - %c1 = constant 1 : index + %c1 = arith.constant 1 : index // CHECK: flow.dispatch.workgroups[%c1, %c1, %c1]() : () -> tensor<i32> = // CHECK-NEXT: (%{{.+}}: !flow.dispatch.tensor<writeonly:i32>) // CHECK: flow.dispatch.tensor.store %{{.+}}, // CHECK-NOT: flow.dispatch.tensor.store %0:2 = flow.dispatch.workgroups[%c1, %c1, %c1](%arg0, %arg1) : (tensor<9xi32>, tensor<9xi32>) -> (tensor<i32>, tensor<i32>) = (%arg0: !flow.dispatch.tensor<readonly:9xi32>, %arg1: !flow.dispatch.tensor<readonly:9xi32>, %arg2: !flow.dispatch.tensor<writeonly:i32>, %arg3: !flow.dispatch.tensor<readwrite:i32>) { - %c0_i32 = constant 0 : i32 - %c-2147483648_i32 = constant -2147483648 : i32 + %c0_i32 = arith.constant 0 : i32 + %c-2147483648_i32 = arith.constant -2147483648 : i32 %0 = flow.dispatch.tensor.load %arg0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:9xi32> -> tensor<9xi32> %1 = flow.dispatch.tensor.load %arg1, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:9xi32> -> tensor<9xi32> %2 = linalg.init_tensor [] : tensor<i32> @@ -203,12 +203,12 @@ // CHECK-LABEL: func @keep_used_read_write_result func @keep_used_read_write_result(%arg0 : tensor<9xi32>, %arg1 : tensor<9xi32>) -> (tensor<i32>) { - %c1 = constant 1 : index + %c1 = arith.constant 1 : index // CHECK: flow.dispatch.workgroups[%c1, %c1, %c1]() : () -> (tensor<i32>, tensor<i32>) = // CHECK-NEXT: (%{{.+}}: !flow.dispatch.tensor<writeonly:i32>, %{{.+}}: !flow.dispatch.tensor<readwrite:i32>) %0:2 = flow.dispatch.workgroups[%c1, %c1, %c1](%arg0, %arg1) : (tensor<9xi32>, tensor<9xi32>) -> (tensor<i32>, tensor<i32>) = (%arg0: !flow.dispatch.tensor<readonly:9xi32>, %arg1: !flow.dispatch.tensor<readonly:9xi32>, %arg2: !flow.dispatch.tensor<writeonly:i32>, %arg3: !flow.dispatch.tensor<readwrite:i32>) { - %c-2147483648_i32 = constant -2147483648 : i32 + %c-2147483648_i32 = arith.constant -2147483648 : i32 %0 = flow.dispatch.tensor.load %arg3, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readwrite:i32> -> tensor<i32> %val = tensor.extract %0[] : tensor<i32> %1 = flow.dispatch.tensor.load %arg1, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:9xi32> -> tensor<9xi32> @@ -226,16 +226,16 @@ // CHECK-LABEL: func @inline_cst_and_remove_unused_read_write_result func @inline_cst_and_remove_unused_read_write_result() -> tensor<i32> { - %cst = constant dense<0> : tensor<i32> - %c1 = constant 1 : index + %cst = arith.constant dense<0> : tensor<i32> + %c1 = arith.constant 1 : index // CHECK: flow.dispatch.workgroups[%c1, %c1, %c1]() : () -> tensor<i32> = // CHECK-NEXT: (%{{.+}}: !flow.dispatch.tensor<readwrite:i32>) %0:2 = flow.dispatch.workgroups[%c1, %c1, %c1](%cst) : (tensor<i32>) -> (tensor<i32>, tensor<i32>) = (%arg0: !flow.dispatch.tensor<readonly:i32>, %arg1: !flow.dispatch.tensor<writeonly:i32>, %arg2: !flow.dispatch.tensor<readwrite:i32>) { %1 = flow.dispatch.tensor.load %arg0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:i32> -> tensor<i32> "test.sink"(%1) : (tensor<i32>) -> () - %c0_i32 = constant 0 : i32 - %c-2147483648_i32 = constant -2147483648 : i32 + %c0_i32 = arith.constant 0 : i32 + %c-2147483648_i32 = arith.constant -2147483648 : i32 %2 = linalg.init_tensor [] : tensor<i32> %3 = linalg.fill(%c-2147483648_i32, %2) : i32, tensor<i32> -> tensor<i32> %4 = linalg.fill(%c0_i32, %2) : i32, tensor<i32> -> tensor<i32>
diff --git a/iree/compiler/Dialect/Flow/IR/test/stream_folding.mlir b/iree/compiler/Dialect/Flow/IR/test/stream_folding.mlir index 93e5259..0044db6 100644 --- a/iree/compiler/Dialect/Flow/IR/test/stream_folding.mlir +++ b/iree/compiler/Dialect/Flow/IR/test/stream_folding.mlir
@@ -2,11 +2,11 @@ // CHECK-LABEL: func @inlineConstant func @inlineConstant() -> index { - %cst = constant 4 : index + %cst = arith.constant 4 : index // CHECK: flow.ex.stream.fragment() %0 = flow.ex.stream.fragment(%cst) : (index) -> index = (%arg0: index) -> index { - // CHECK: %[[C:.+]] = constant 4 : index + // CHECK: %[[C:.+]] = arith.constant 4 : index // CHECK-NEXT: return %[[C]] flow.return %arg0 : index } @@ -18,13 +18,13 @@ // CHECK-LABEL: func @removeUnusedCapture // CHECK-SAME: (%[[ARG:.+]]: index) func @removeUnusedCapture(%arg: index) -> index { - %unused = constant 5 : index + %unused = arith.constant 5 : index // CHECK: flow.ex.stream.fragment(%[[ARG]]) %0 = flow.ex.stream.fragment(%arg, %unused) : (index, index) -> index = // CHECK-NEXT: (%[[INNER_ARG:.+]]: index) -> index { (%arg0: index, %arg1: index) -> index { - // CHECK-NEXT: %[[T:.+]] = addi %[[INNER_ARG]], %[[INNER_ARG]] - %t = addi %arg0, %arg0 : index + // CHECK-NEXT: %[[T:.+]] = arith.addi %[[INNER_ARG]], %[[INNER_ARG]] + %t = arith.addi %arg0, %arg0 : index // CHECK-NEXT: flow.return %[[T]] flow.return %t : index } @@ -39,7 +39,7 @@ // CHECK: flow.ex.stream.fragment(%[[ARG]]) %0 = flow.ex.stream.fragment(%arg, %arg) : (index, index) -> index = (%arg0: index, %arg1: index) -> index { - %t = addi %arg0, %arg0 : index + %t = arith.addi %arg0, %arg0 : index flow.return %t : index } return %0 : index @@ -53,9 +53,9 @@ // CHECK: flow.ex.stream.fragment(%[[ARG0]]) : (index) -> index = %0:2 = flow.ex.stream.fragment(%arg0) : (index) -> (index, index) = (%arg0_in: index) -> (index, index) { - // CHECK: %[[T:.+]] = addi - %t = addi %arg0_in, %arg0_in : index - %unused = muli %arg0_in, %arg0_in : index + // CHECK: %[[T:.+]] = arith.addi + %t = arith.addi %arg0_in, %arg0_in : index + %unused = arith.muli %arg0_in, %arg0_in : index // CHECK: flow.return %[[T]] : index flow.return %t, %unused : index, index } @@ -70,8 +70,8 @@ // CHECK: flow.ex.stream.fragment(%[[ARG1]]) %0:2 = flow.ex.stream.fragment(%arg0, %arg1) : (index, index) -> (index, index) = (%unused: index, %arg1_in: index) -> (index, index) { - // CHECK: %[[T:.+]] = addi - %t = addi %arg1_in, %arg1_in : index + // CHECK: %[[T:.+]] = arith.addi + %t = arith.addi %arg1_in, %arg1_in : index // CHECK: flow.return %[[T]] : index flow.return %t, %unused : index, index } @@ -110,9 +110,9 @@ (tensor<2x4xi32>, tensor<1x1xi32>) -> tensor<2x4xi32> = // CHECK-NEXT: (%[[TARGET:.+]]: tensor<2x4xi32>, %[[UPDATE:.+]]: tensor<1x1xi32>) (%stream_target: tensor<2x4xi32>, %stream_update: tensor<1x1xi32>) -> tensor<2x4xi32> { - %start0 = constant 0 : index - %start1 = constant 1 : index - %workload = constant 8 : index + %start0 = arith.constant 0 : index + %start1 = arith.constant 1 : index + %workload = arith.constant 8 : index // CHECK: %[[TARGET_CLONE:.+]] = flow.tensor.clone %[[TARGET]] : tensor<2x4xi32> // CHECK: %[[UPDATED:.+]] = flow.tensor.update %[[UPDATE]], %[[TARGET]] %t0 = flow.tensor.update %stream_update, %stream_target[%start0, %start1] : tensor<1x1xi32> -> %stream_target as tensor<2x4xi32> @@ -137,11 +137,11 @@ func @dagImmutability(%arg0: tensor<1xi32>) -> (tensor<i32>, tensor<1xi32>, tensor<3xi32>) { %0:3 = flow.ex.stream.fragment(%arg0) : (tensor<1xi32>) -> (tensor<i32>, tensor<1xi32>, tensor<3xi32>) = (%arg1: tensor<1xi32>) -> (tensor<i32>, tensor<1xi32>, tensor<3xi32>) { - %c9 = constant 9 : index - %c1 = constant 1 : index - %c18 = constant 18 : index - %c0 = constant 0 : index - %c3 = constant 3 : index + %c9 = arith.constant 9 : index + %c1 = arith.constant 1 : index + %c18 = arith.constant 18 : index + %c0 = arith.constant 0 : index + %c3 = arith.constant 3 : index %1 = flow.dispatch @_run_dispatch_1::@_run_dispatch_1[%c1, %c1, %c1]() : () -> tensor<i32> %2 = flow.dispatch @_run_dispatch_2::@_run_dispatch_2[%c9, %c1, %c1](%1) : (tensor<i32>) -> tensor<9xi32> %3 = flow.tensor.reshape %1 : tensor<i32> -> tensor<1xi32> @@ -159,8 +159,8 @@ func @insertCloneForUpdatedConstant(%input: tensor<2x2xi32>) -> tensor<4x4xi32> { %4 = flow.ex.stream.fragment(%input) : (tensor<2x2xi32>) -> tensor<4x4xi32> = (%arg0: tensor<2x2xi32>) -> tensor<4x4xi32> { - %c4 = constant 4 : index - %c1 = constant 1 : index + %c4 = arith.constant 4 : index + %c1 = arith.constant 1 : index // CHECK: %[[LOAD:.+]] = util.global.load @_large_const %5 = util.global.load @_large_const : tensor<4x4xi32> // CHECK: %[[CLONE:.+]] = flow.tensor.clone %[[LOAD]] @@ -179,7 +179,7 @@ func @insertCloneForUpdatedConstant(%input: tensor<2xi32>) -> tensor<7xi32> { %4 = flow.ex.stream.fragment(%input) : (tensor<2xi32>) -> tensor<7xi32> = (%arg0: tensor<2xi32>) -> tensor<7xi32> { - %c3 = constant 3 : index + %c3 = arith.constant 3 : index // CHECK: %[[LOAD:.+]] = util.global.load @_large_const %5 = util.global.load @_large_const : tensor<7xi32> // CHECK: %[[CLONE:.+]] = flow.tensor.clone %[[LOAD]]
diff --git a/iree/compiler/Dialect/Flow/IR/test/stream_ops.mlir b/iree/compiler/Dialect/Flow/IR/test/stream_ops.mlir index 89f69bc..a0e4bca 100644 --- a/iree/compiler/Dialect/Flow/IR/test/stream_ops.mlir +++ b/iree/compiler/Dialect/Flow/IR/test/stream_ops.mlir
@@ -12,8 +12,8 @@ // CHECK-LABEL: func @fragment func @fragment(%arg0 : tensor<4xf32>) -> (tensor<4xf32>, tensor<4xf32>) { - // CHECK: %[[WORKLOAD:.+]] = constant - %cst = constant 4 : index + // CHECK: %[[WORKLOAD:.+]] = arith.constant + %cst = arith.constant 4 : index // CHECK: %0:2 = flow.ex.stream.fragment(%[[WORKLOAD]], %arg0) : (index, tensor<4xf32>) -> (tensor<4xf32>, tensor<4xf32>) = // CHECK-NEXT: (%arg1: index, %arg2: tensor<4xf32>) -> (tensor<4xf32>, tensor<4xf32>) { %0:2 = flow.ex.stream.fragment(%cst, %arg0) : (index, tensor<4xf32>) -> (tensor<4xf32>, tensor<4xf32>) =
diff --git a/iree/compiler/Dialect/Flow/IR/test/tensor_folding.mlir b/iree/compiler/Dialect/Flow/IR/test/tensor_folding.mlir index 724b8e5..42828c1 100644 --- a/iree/compiler/Dialect/Flow/IR/test/tensor_folding.mlir +++ b/iree/compiler/Dialect/Flow/IR/test/tensor_folding.mlir
@@ -71,10 +71,10 @@ // CHECK-LABEL: @loadConst func @loadConst() -> i32 { - %0 = constant dense<[[0, 1], [2, 3]]> : tensor<2x2xi32> - %c0 = constant 0 : index - %c1 = constant 1 : index - // CHECK-NEXT: %[[C2:.+]] = constant 2 : i32 + %0 = arith.constant dense<[[0, 1], [2, 3]]> : tensor<2x2xi32> + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + // CHECK-NEXT: %[[C2:.+]] = arith.constant 2 : i32 %2 = flow.tensor.load %0[%c1, %c0] : tensor<2x2xi32> // CHECK-NEXT: return %[[C2]] return %2 : i32 @@ -84,8 +84,8 @@ // CHECK-LABEL: @loadConstScalar func @loadConstScalar() -> i32 { - %0 = constant dense<4> : tensor<i32> - // CHECK-NEXT: %[[C4:.+]] = constant 4 : i32 + %0 = arith.constant dense<4> : tensor<i32> + // CHECK-NEXT: %[[C4:.+]] = arith.constant 4 : i32 %1 = flow.tensor.load %0 : tensor<i32> // CHECK-NEXT: return %[[C4]] return %1 : i32 @@ -95,11 +95,11 @@ // CHECK-LABEL: @storeConst func @storeConst() -> tensor<2x2xi32> { - %0 = constant dense<[[0, 1], [2, 3]]> : tensor<2x2xi32> - %c0 = constant 0 : index - %c1 = constant 1 : index - %c4 = constant 4 : i32 - // CHECK-NEXT: %[[C:.+]] = constant dense<[ + %0 = arith.constant dense<[[0, 1], [2, 3]]> : tensor<2x2xi32> + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + %c4 = arith.constant 4 : i32 + // CHECK-NEXT: %[[C:.+]] = arith.constant dense<[ // CHECK-SAME: [0, 1], [4, 3] // CHECK-SAME: ]> : tensor<2x2xi32> %1 = flow.tensor.store %c4, %0[%c1, %c0] : tensor<2x2xi32> @@ -111,9 +111,9 @@ // CHECK-LABEL: @storeConstScalar func @storeConstScalar() -> tensor<i32> { - %0 = constant dense<0> : tensor<i32> - %1 = constant 4 : i32 - // CHECK-NEXT: %[[C:.+]] = constant dense<4> : tensor<i32> + %0 = arith.constant dense<0> : tensor<i32> + %1 = arith.constant 4 : i32 + // CHECK-NEXT: %[[C:.+]] = arith.constant dense<4> : tensor<i32> %2 = flow.tensor.store %1, %0 : tensor<i32> // CHECK-NEXT: return %[[C]] return %2 : tensor<i32> @@ -123,8 +123,8 @@ // CHECK-LABEL: @splatConst func @splatConst() -> tensor<4xi32> { - %0 = constant 4 : i32 - // CHECK-NEXT: %[[C:.+]] = constant dense<4> : tensor<4xi32> + %0 = arith.constant 4 : i32 + // CHECK-NEXT: %[[C:.+]] = arith.constant dense<4> : tensor<4xi32> %1 = flow.tensor.splat %0 : tensor<4xi32> // CHECK-NEXT: return %[[C]] return %1 : tensor<4xi32> @@ -134,8 +134,8 @@ // CHECK-LABEL: @splatConstScalar func @splatConstScalar() -> tensor<i32> { - %0 = constant 4 : i32 - // CHECK-NEXT: %[[C:.+]] = constant dense<4> : tensor<i32> + %0 = arith.constant 4 : i32 + // CHECK-NEXT: %[[C:.+]] = arith.constant dense<4> : tensor<i32> %1 = flow.tensor.splat %0 : tensor<i32> // CHECK-NEXT: return %[[C]] return %1 : tensor<i32> @@ -146,8 +146,8 @@ // CHECK-LABEL: @splatDynamicShape // CHECK-SAME: (%[[DIM0:.+]]: index, %[[DIM1:.+]]: index) func @splatDynamicShape(%dim0: index, %dim1: index) -> tensor<?x?xi32> { - // CHECK: %[[FOUR:.+]] = constant 4 : i32 - %four = constant 4 : i32 + // CHECK: %[[FOUR:.+]] = arith.constant 4 : i32 + %four = arith.constant 4 : i32 // CHECK: %[[SPLAT:.+]] = flow.tensor.splat %[[FOUR]] : tensor<?x?xi32>{%[[DIM0]], %[[DIM1]]} %1 = flow.tensor.splat %four : tensor<?x?xi32>{%dim0, %dim1} // CHECK: return %[[SPLAT]] @@ -158,8 +158,8 @@ // CHECK-LABEL: @cloneConst func @cloneConst() -> tensor<4xi32> { - %0 = constant dense<[0, 1, 2, 3]> : tensor<4xi32> - // CHECK-NEXT: %[[C:.+]] = constant dense<[0, 1, 2, 3]> : tensor<4xi32> + %0 = arith.constant dense<[0, 1, 2, 3]> : tensor<4xi32> + // CHECK-NEXT: %[[C:.+]] = arith.constant dense<[0, 1, 2, 3]> : tensor<4xi32> %1 = flow.tensor.clone %0 : tensor<4xi32> // CHECK-NEXT: return %[[C]] return %1 : tensor<4xi32> @@ -169,8 +169,8 @@ // CHECK-LABEL: @sliceConst0D func @sliceConst0D() -> tensor<i32> { - %0 = constant dense<0> : tensor<i32> - // CHECK-NEXT: %[[C:.+]] = constant dense<0> : tensor<i32> + %0 = arith.constant dense<0> : tensor<i32> + // CHECK-NEXT: %[[C:.+]] = arith.constant dense<0> : tensor<i32> %1 = flow.tensor.slice %0[for] : tensor<i32> -> tensor<i32> // CHECK-NEXT: return %[[C]] return %1 : tensor<i32> @@ -180,10 +180,10 @@ // CHECK-LABEL: @sliceConst1D func @sliceConst1D() -> tensor<1xi32> { - %0 = constant dense<0> : tensor<1xi32> - %c0 = constant 0 : index - %c1 = constant 1 : index - // CHECK-NEXT: %[[C:.+]] = constant dense<0> : tensor<1xi32> + %0 = arith.constant dense<0> : tensor<1xi32> + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + // CHECK-NEXT: %[[C:.+]] = arith.constant dense<0> : tensor<1xi32> %1 = flow.tensor.slice %0[%c0 for %c1] : tensor<1xi32> -> tensor<1xi32> // CHECK-NEXT: return %[[C]] return %1 : tensor<1xi32> @@ -193,9 +193,9 @@ // CHECK-LABEL: @sliceConst1DZeroLength func @sliceConst1DZeroLength() -> tensor<0xi32> { - %0 = constant dense<0> : tensor<1xi32> - %c0 = constant 0 : index - // CHECK-NEXT: %[[C:.+]] = constant dense<> : tensor<0xi32> + %0 = arith.constant dense<0> : tensor<1xi32> + %c0 = arith.constant 0 : index + // CHECK-NEXT: %[[C:.+]] = arith.constant dense<> : tensor<0xi32> %1 = flow.tensor.slice %0[%c0 for %c0] : tensor<1xi32> -> tensor<0xi32> // CHECK-NEXT: return %[[C]] return %1 : tensor<0xi32> @@ -205,11 +205,11 @@ // CHECK-LABEL: @sliceConst2D func @sliceConst2D() -> tensor<1x2xi32> { - %0 = constant dense<[[0, 1, 2], [3, 4, 5]]> : tensor<2x3xi32> - %c0 = constant 0 : index - %c1 = constant 1 : index - %c2 = constant 2 : index - // CHECK-NEXT: %[[C:.+]] = constant dense<[ + %0 = arith.constant dense<[[0, 1, 2], [3, 4, 5]]> : tensor<2x3xi32> + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + %c2 = arith.constant 2 : index + // CHECK-NEXT: %[[C:.+]] = arith.constant dense<[ // CHECK-SAME: [1, 2] // CHECK-SAME: ]> : tensor<1x2xi32> %1 = flow.tensor.slice %0[%c0, %c1 for %c1, %c2] : tensor<2x3xi32> -> tensor<1x2xi32> @@ -221,10 +221,10 @@ // CHECK-LABEL: @sliceConst2DZeroLength1 func @sliceConst2DZeroLength1() -> tensor<1x0xi32> { - %0 = constant dense<[[0, 1, 2], [3, 4, 5]]> : tensor<2x3xi32> - %c0 = constant 0 : index - %c1 = constant 1 : index - // CHECK-NEXT: %[[C:.+]] = constant dense<> : tensor<1x0xi32> + %0 = arith.constant dense<[[0, 1, 2], [3, 4, 5]]> : tensor<2x3xi32> + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + // CHECK-NEXT: %[[C:.+]] = arith.constant dense<> : tensor<1x0xi32> %1 = flow.tensor.slice %0[%c0, %c0 for %c1, %c0] : tensor<2x3xi32> -> tensor<1x0xi32> // CHECK-NEXT: return %[[C]] return %1 : tensor<1x0xi32> @@ -234,9 +234,9 @@ // CHECK-LABEL: @sliceConst2DZeroLength01 func @sliceConst2DZeroLength01() -> tensor<0x0xi32> { - %0 = constant dense<[[0, 1, 2], [3, 4, 5]]> : tensor<2x3xi32> - %c0 = constant 0 : index - // CHECK-NEXT: %[[C:.+]] = constant dense<> : tensor<0x0xi32> + %0 = arith.constant dense<[[0, 1, 2], [3, 4, 5]]> : tensor<2x3xi32> + %c0 = arith.constant 0 : index + // CHECK-NEXT: %[[C:.+]] = arith.constant dense<> : tensor<0x0xi32> %1 = flow.tensor.slice %0[%c0, %c0 for %c0, %c0] : tensor<2x3xi32> -> tensor<0x0xi32> // CHECK-NEXT: return %[[C]] return %1 : tensor<0x0xi32> @@ -246,12 +246,12 @@ // CHECK-LABEL: @sliceConst3D func @sliceConst3D() -> tensor<1x2x3xi32> { - %0 = constant dense<[[[0, 1, 2], [3, 4, 5], [6, 7, 8]], [[9, 10, 11], [12, 13, 14], [15, 16, 17]]]> : tensor<2x3x3xi32> - %c0 = constant 0 : index - %c1 = constant 1 : index - %c2 = constant 2 : index - %c3 = constant 3 : index - // CHECK-NEXT: %[[C:.+]] = constant dense<[ + %0 = arith.constant dense<[[[0, 1, 2], [3, 4, 5], [6, 7, 8]], [[9, 10, 11], [12, 13, 14], [15, 16, 17]]]> : tensor<2x3x3xi32> + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + %c2 = arith.constant 2 : index + %c3 = arith.constant 3 : index + // CHECK-NEXT: %[[C:.+]] = arith.constant dense<[ // CHECK-SAME: [ // CHECK-SAME: [3, 4, 5], [6, 7, 8]]]> : tensor<1x2x3xi32> %1 = flow.tensor.slice %0[%c0, %c1, %c0 for %c1, %c2, %c3] : tensor<2x3x3xi32> -> tensor<1x2x3xi32> @@ -263,9 +263,9 @@ // CHECK-LABEL: @updateConst0D func @updateConst0D() -> tensor<i32> { - %0 = constant dense<0> : tensor<i32> - %1 = constant dense<1> : tensor<i32> - // CHECK-NEXT: %[[C:.+]] = constant dense<0> : tensor<i32> + %0 = arith.constant dense<0> : tensor<i32> + %1 = arith.constant dense<1> : tensor<i32> + // CHECK-NEXT: %[[C:.+]] = arith.constant dense<0> : tensor<i32> %2 = flow.tensor.update %0, %1[] : tensor<i32> -> tensor<i32> // CHECK-NEXT: return %[[C]] return %2 : tensor<i32> @@ -275,10 +275,10 @@ // CHECK-LABEL: @updateConst1D func @updateConst1D() -> tensor<1xi32> { - %0 = constant dense<0> : tensor<1xi32> - %1 = constant dense<1> : tensor<1xi32> - %c0 = constant 0 : index - // CHECK-NEXT: %[[C:.+]] = constant dense<0> : tensor<1xi32> + %0 = arith.constant dense<0> : tensor<1xi32> + %1 = arith.constant dense<1> : tensor<1xi32> + %c0 = arith.constant 0 : index + // CHECK-NEXT: %[[C:.+]] = arith.constant dense<0> : tensor<1xi32> %2 = flow.tensor.update %0, %1[%c0] : tensor<1xi32> -> tensor<1xi32> // CHECK-NEXT: return %[[C]] return %2 : tensor<1xi32> @@ -288,10 +288,10 @@ // CHECK-LABEL: @updateConst1DUpdateZeroSize func @updateConst1DUpdateZeroSize() -> tensor<1xi32> { - %0 = constant dense<> : tensor<0xi32> - %1 = constant dense<1> : tensor<1xi32> - %c0 = constant 0 : index - // CHECK-NEXT: %[[C:.+]] = constant dense<1> : tensor<1xi32> + %0 = arith.constant dense<> : tensor<0xi32> + %1 = arith.constant dense<1> : tensor<1xi32> + %c0 = arith.constant 0 : index + // CHECK-NEXT: %[[C:.+]] = arith.constant dense<1> : tensor<1xi32> %2 = flow.tensor.update %0, %1[%c0] : tensor<0xi32> -> tensor<1xi32> // CHECK-NEXT: return %[[C]] return %2 : tensor<1xi32> @@ -301,11 +301,11 @@ // CHECK-LABEL: @updateConst2DUpdate1x1 func @updateConst2DUpdate1x1() -> tensor<3x4xi32> { - %0 = constant dense<[[12]]> : tensor<1x1xi32> - %1 = constant dense<[[0, 1, 2, 3], [4, 5, 6, 7], [8, 9, 10, 11]]> : tensor<3x4xi32> - %c0 = constant 0 : index - %c1 = constant 1 : index - // CHECK-NEXT: %[[C:.+]] = constant dense<[ + %0 = arith.constant dense<[[12]]> : tensor<1x1xi32> + %1 = arith.constant dense<[[0, 1, 2, 3], [4, 5, 6, 7], [8, 9, 10, 11]]> : tensor<3x4xi32> + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + // CHECK-NEXT: %[[C:.+]] = arith.constant dense<[ // CHECK-SAME: [0, 12, 2, 3], [4, 5, 6, 7], [8, 9, 10, 11]]> : tensor<3x4xi32> %2 = flow.tensor.update %0, %1[%c0, %c1] : tensor<1x1xi32> -> tensor<3x4xi32> // CHECK-NEXT: return %[[C]] @@ -316,11 +316,11 @@ // CHECK-LABEL: @updateConst2DUpdate2x2 func @updateConst2DUpdate2x2() -> tensor<3x4xi32> { - %0 = constant dense<[[12, 13], [14, 15]]> : tensor<2x2xi32> - %1 = constant dense<[[0, 1, 2, 3], [4, 5, 6, 7], [8, 9, 10, 11]]> : tensor<3x4xi32> - %c0 = constant 0 : index - %c1 = constant 1 : index - // CHECK-NEXT: %[[C:.+]] = constant dense<[ + %0 = arith.constant dense<[[12, 13], [14, 15]]> : tensor<2x2xi32> + %1 = arith.constant dense<[[0, 1, 2, 3], [4, 5, 6, 7], [8, 9, 10, 11]]> : tensor<3x4xi32> + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + // CHECK-NEXT: %[[C:.+]] = arith.constant dense<[ // CHECK-SAME: [0, 12, 13, 3], [4, 14, 15, 7], [8, 9, 10, 11]]> : tensor<3x4xi32> %2 = flow.tensor.update %0, %1[%c0, %c1] : tensor<2x2xi32> -> tensor<3x4xi32> // CHECK-NEXT: return %[[C]] @@ -331,11 +331,11 @@ // CHECK-LABEL: @updateConst3DUpdate1x2x3 func @updateConst3DUpdate1x2x3() -> tensor<2x3x3xi32> { - %0 = constant dense<[[[18, 19, 20], [21, 22, 23]]]> : tensor<1x2x3xi32> - %1 = constant dense<[[[0, 1, 2], [3, 4, 5], [6, 7, 8]], [[9, 10, 11], [12, 13, 14], [15, 16, 17]]]> : tensor<2x3x3xi32> - %c0 = constant 0 : index - %c1 = constant 1 : index - // CHECK-NEXT: %[[C:.+]] = constant dense<[ + %0 = arith.constant dense<[[[18, 19, 20], [21, 22, 23]]]> : tensor<1x2x3xi32> + %1 = arith.constant dense<[[[0, 1, 2], [3, 4, 5], [6, 7, 8]], [[9, 10, 11], [12, 13, 14], [15, 16, 17]]]> : tensor<2x3x3xi32> + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + // CHECK-NEXT: %[[C:.+]] = arith.constant dense<[ // CHECK-SAME: [ // CHECK-SAME: [0, 1, 2], [18, 19, 20], [21, 22, 23]], [ // CHECK-SAME: [9, 10, 11], [12, 13, 14], [15, 16, 17]]]> : tensor<2x3x3xi32> @@ -348,11 +348,11 @@ // CHECK-LABEL: @updateConst3DUpdate2x3x2 func @updateConst3DUpdate2x3x2() -> tensor<2x3x3xi32> { - %0 = constant dense<[[[18, 19], [20, 21], [22, 23]], [[24, 25], [26, 27], [28, 29]]]> : tensor<2x3x2xi32> - %1 = constant dense<[[[0, 1, 2], [3, 4, 5], [6, 7, 8]], [[9, 10, 11], [12, 13, 14], [15, 16, 17]]]> : tensor<2x3x3xi32> - %c0 = constant 0 : index - %c1 = constant 0 : index - // CHECK-NEXT: %[[C:.+]] = constant dense<[ + %0 = arith.constant dense<[[[18, 19], [20, 21], [22, 23]], [[24, 25], [26, 27], [28, 29]]]> : tensor<2x3x2xi32> + %1 = arith.constant dense<[[[0, 1, 2], [3, 4, 5], [6, 7, 8]], [[9, 10, 11], [12, 13, 14], [15, 16, 17]]]> : tensor<2x3x3xi32> + %c0 = arith.constant 0 : index + %c1 = arith.constant 0 : index + // CHECK-NEXT: %[[C:.+]] = arith.constant dense<[ // CHECK-SAME: [ // CHECK-SAME: [18, 19, 2], [20, 21, 5], [22, 23, 8]], [ // CHECK-SAME: [24, 25, 11], [26, 27, 14], [28, 29, 17]]]> : tensor<2x3x3xi32> @@ -365,7 +365,7 @@ // CHECK-LABEL: @updateReplace func @updateReplace(%arg0 : tensor<4xi32>, %arg1 : tensor<4xi32>) -> tensor<4xi32> { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %0 = flow.tensor.update %arg0, %arg1[%c0] : tensor<4xi32> -> tensor<4xi32> // CHECK-NEXT: return %arg0 return %0 : tensor<4xi32> @@ -375,10 +375,10 @@ // CHECK-LABEL: @propogateStaticShapeOfTarget func @propogateStaticShapeOfTarget(%arg0 : tensor<?x?xf32>, %arg1 : f32) -> tensor<?x?xf32> { - %c21 = constant 21 : index - %c42 = constant 42 : index - %c2 = constant 2 : index - %c4 = constant 4 : index + %c21 = arith.constant 21 : index + %c42 = arith.constant 42 : index + %c2 = arith.constant 2 : index + %c4 = arith.constant 4 : index // CHECK: %[[TARGET:.+]] = tensor.generate { // CHECK: } : tensor<21x42xf32> %0 = tensor.generate %c21, %c42 { @@ -396,10 +396,10 @@ // CHECK-LABEL: @propogateStaticShapeOfUpdate func @propogateStaticShapeOfUpdate(%arg0 : tensor<?x?xf32>, %arg1 : f32) -> tensor<?x?xf32> { - %c21 = constant 21 : index - %c42 = constant 42 : index - %c2 = constant 2 : index - %c4 = constant 4 : index + %c21 = arith.constant 21 : index + %c42 = arith.constant 42 : index + %c2 = arith.constant 2 : index + %c4 = arith.constant 4 : index // CHECK: %[[UPDATE:.+]] = tensor.generate { // CHECK: } : tensor<21x42xf32> %0 = tensor.generate %c21, %c42 {
diff --git a/iree/compiler/Dialect/Flow/IR/test/tensor_ops.mlir b/iree/compiler/Dialect/Flow/IR/test/tensor_ops.mlir index 12f2481..d8f0171 100644 --- a/iree/compiler/Dialect/Flow/IR/test/tensor_ops.mlir +++ b/iree/compiler/Dialect/Flow/IR/test/tensor_ops.mlir
@@ -18,8 +18,8 @@ // CHECK-LABEL: @tensorReshapeDynamic func @tensorReshapeDynamic(%arg0 : tensor<?x4xf32>) -> tensor<?x2xf32> { - %c4 = constant 4 : index - %c8 = constant 8 : index + %c4 = arith.constant 4 : index + %c8 = arith.constant 8 : index // CHECK: %0 = flow.tensor.reshape %arg0 : tensor<?x4xf32>{%c4} -> tensor<?x2xf32>{%c8} %0 = flow.tensor.reshape %arg0 : tensor<?x4xf32>{%c4} -> tensor<?x2xf32>{%c8} return %0 : tensor<?x2xf32> @@ -43,7 +43,7 @@ // CHECK-LABEL: @tensorLoadDynamic func @tensorLoadDynamic(%arg0 : tensor<?x4xf32>, %arg1 : index, %arg2 : index) -> f32 { - %c4 = constant 4 : index + %c4 = arith.constant 4 : index // CHECK: %0 = flow.tensor.load %arg0[%arg1, %arg2] : tensor<?x4xf32>{%c4} %0 = flow.tensor.load %arg0[%arg1, %arg2] : tensor<?x4xf32>{%c4} return %0 : f32 @@ -67,7 +67,7 @@ // CHECK-LABEL: @tensorStoreDynamic func @tensorStoreDynamic(%arg0 : tensor<?x4xf32>, %arg1 : index, %arg2 : index, %arg3 : f32) -> tensor<?x4xf32> { - %c4 = constant 4 : index + %c4 = arith.constant 4 : index // CHECK: %0 = flow.tensor.store %arg3, %arg0[%arg1, %arg2] : tensor<?x4xf32>{%c4} %0 = flow.tensor.store %arg3, %arg0[%arg1, %arg2] : tensor<?x4xf32>{%c4} return %0 : tensor<?x4xf32> @@ -91,7 +91,7 @@ // CHECK-LABEL: @tensorSplatDynamic func @tensorSplatDynamic(%arg0 : f32) -> tensor<?x4xf32> { - %c4 = constant 4 : index + %c4 = arith.constant 4 : index // CHECK: %0 = flow.tensor.splat %arg0 : tensor<?x4xf32>{%c4} %0 = flow.tensor.splat %arg0 : tensor<?x4xf32>{%c4} return %0 : tensor<?x4xf32> @@ -115,7 +115,7 @@ // CHECK-LABEL: @tensorCloneDynamic func @tensorCloneDynamic(%arg0 : tensor<?x4xf32>) -> tensor<?x4xf32> { - %c4 = constant 4 : index + %c4 = arith.constant 4 : index // CHECK: %0 = flow.tensor.clone %arg0 : tensor<?x4xf32>{%c4} %0 = flow.tensor.clone %arg0 : tensor<?x4xf32>{%c4} return %0 : tensor<?x4xf32> @@ -132,8 +132,8 @@ // CHECK-LABEL: @tensorSliceDynamic func @tensorSliceDynamic(%arg0 : tensor<?x4xf32>, %arg1 : index, %arg2 : index) -> tensor<?x2xf32> { - %c2 = constant 2 : index - %c4 = constant 4 : index + %c2 = arith.constant 2 : index + %c4 = arith.constant 4 : index // CHECK: %0 = flow.tensor.slice %arg0[%arg1, %arg2 for %arg2, %arg1] : tensor<?x4xf32>{%c4} -> tensor<?x2xf32>{%c2} %0 = flow.tensor.slice %arg0[%arg1, %arg2 for %arg2, %arg1] : tensor<?x4xf32>{%c4} -> tensor<?x2xf32>{%c2} return %0 : tensor<?x2xf32> @@ -150,9 +150,9 @@ // CHECK-LABEL: @tensorUpdateDynamic func @tensorUpdateDynamic(%arg0 : tensor<?x?xf32>, %arg1 : tensor<?x4xf32>, %arg2 : index, %arg3 : index) -> tensor<?x4xf32> { - %c1 = constant 1 : index - %c2 = constant 2 : index - %c3 = constant 3 : index + %c1 = arith.constant 1 : index + %c2 = arith.constant 2 : index + %c3 = arith.constant 3 : index // CHECK: %0 = flow.tensor.update %arg0, %arg1[%arg2, %arg3] : tensor<?x?xf32>{%c1, %c2} -> %arg1 as tensor<?x4xf32>{%c3} %0 = flow.tensor.update %arg0, %arg1[%arg2, %arg3] : tensor<?x?xf32>{%c1, %c2} -> %arg1 as tensor<?x4xf32>{%c3} return %0 : tensor<?x4xf32>
diff --git a/iree/compiler/Dialect/Flow/Transforms/ConvertLinalgTensorOps.cpp b/iree/compiler/Dialect/Flow/Transforms/ConvertLinalgTensorOps.cpp index 455a7ef..869143a 100644 --- a/iree/compiler/Dialect/Flow/Transforms/ConvertLinalgTensorOps.cpp +++ b/iree/compiler/Dialect/Flow/Transforms/ConvertLinalgTensorOps.cpp
@@ -102,7 +102,8 @@ void getDependentDialects(DialectRegistry ®istry) const override { registry.insert<IREE::Flow::FlowDialect, tensor::TensorDialect, - linalg::LinalgDialect, mlir::StandardOpsDialect>(); + linalg::LinalgDialect, mlir::StandardOpsDialect, + mlir::arith::ArithmeticDialect, mlir::math::MathDialect>(); } void runOnOperation() override { auto funcOp = getOperation();
diff --git a/iree/compiler/Dialect/Flow/Transforms/ConvertTensorOps.cpp b/iree/compiler/Dialect/Flow/Transforms/ConvertTensorOps.cpp index a4ba75b..d2d5a28 100644 --- a/iree/compiler/Dialect/Flow/Transforms/ConvertTensorOps.cpp +++ b/iree/compiler/Dialect/Flow/Transforms/ConvertTensorOps.cpp
@@ -30,6 +30,7 @@ : public ConvertTensorOpsBase<ConvertTensorOpsPass> { void getDependentDialects(DialectRegistry ®istry) const override { registry.insert<IREE::Flow::FlowDialect, mlir::StandardOpsDialect, + mlir::arith::ArithmeticDialect, mlir::math::MathDialect, tensor::TensorDialect>(); } void runOnOperation() override {
diff --git a/iree/compiler/Dialect/Flow/Transforms/ConvertToFlow.cpp b/iree/compiler/Dialect/Flow/Transforms/ConvertToFlow.cpp index 95671d6..993871f 100644 --- a/iree/compiler/Dialect/Flow/Transforms/ConvertToFlow.cpp +++ b/iree/compiler/Dialect/Flow/Transforms/ConvertToFlow.cpp
@@ -96,7 +96,8 @@ ConvertToFlowBeforeDispatchFormation> { void getDependentDialects(DialectRegistry ®istry) const override { registry.insert<IREE::Flow::FlowDialect, tensor::TensorDialect, - linalg::LinalgDialect, mlir::StandardOpsDialect>(); + linalg::LinalgDialect, mlir::StandardOpsDialect, + mlir::arith::ArithmeticDialect, mlir::math::MathDialect>(); } void runOnOperation() override { auto funcOp = getOperation(); @@ -122,7 +123,8 @@ ConvertToFlowAfterDispatchFormation> { void getDependentDialects(DialectRegistry ®istry) const override { registry.insert<IREE::Flow::FlowDialect, tensor::TensorDialect, - linalg::LinalgDialect, mlir::StandardOpsDialect>(); + linalg::LinalgDialect, mlir::StandardOpsDialect, + mlir::arith::ArithmeticDialect, mlir::math::MathDialect>(); } void runOnOperation() override { auto funcOp = getOperation();
diff --git a/iree/compiler/Dialect/Flow/Transforms/DeduplicateExecutables.cpp b/iree/compiler/Dialect/Flow/Transforms/DeduplicateExecutables.cpp index 39fbeb4..8989fbc 100644 --- a/iree/compiler/Dialect/Flow/Transforms/DeduplicateExecutables.cpp +++ b/iree/compiler/Dialect/Flow/Transforms/DeduplicateExecutables.cpp
@@ -48,12 +48,12 @@ // // Example: // func @lhs(%arg0 : index) -> index { -// %c1 = constant 1 : index +// %c1 = arith.constant 1 : index // %0 = add %arg0, %c1 : index // return %0 : index // } // func @rhs(%arg0 : index) -> index { -// %c1 = constant 1 : index +// %c1 = arith.constant 1 : index // %0 = add %arg0, %c1 : index // return %0 : index // }
diff --git a/iree/compiler/Dialect/Flow/Transforms/DispatchLinalgOnTensors.cpp b/iree/compiler/Dialect/Flow/Transforms/DispatchLinalgOnTensors.cpp index 691a6ba..8905b71 100644 --- a/iree/compiler/Dialect/Flow/Transforms/DispatchLinalgOnTensors.cpp +++ b/iree/compiler/Dialect/Flow/Transforms/DispatchLinalgOnTensors.cpp
@@ -134,7 +134,7 @@ assert(shape.size() <= kNumMaxParallelDims && "workload cannot be more than 3D for now"); SmallVector<Value, 4> workload = llvm::to_vector<4>(llvm::reverse(shape)); - Value one = b.create<ConstantIndexOp>(loc, 1); + Value one = b.create<arith::ConstantIndexOp>(loc, 1); workload.resize(kNumMaxParallelDims, one); return workload; } @@ -179,10 +179,10 @@ } static bool isAlwaysClonedIntoDispatchOp(Operation *op) { - if (isa<IndexCastOp, linalg::InitTensorOp, tensor::ExtractOp>(op)) { + if (isa<arith::IndexCastOp, linalg::InitTensorOp, tensor::ExtractOp>(op)) { return true; } - if (auto constantOp = dyn_cast<ConstantOp>(op)) { + if (auto constantOp = dyn_cast<arith::ConstantOp>(op)) { return constantOp.getResult().getType().isIntOrIndexOrFloat(); } if (llvm::all_of(op->getOperands(), @@ -749,7 +749,7 @@ if (!shapedType.hasStaticShape()) return llvm::None; outputShapes.push_back(llvm::to_vector<4>( llvm::map_range(shapedType.getShape(), [&](int64_t dim) -> Value { - return builder.create<ConstantIndexOp>(op->getLoc(), dim); + return builder.create<arith::ConstantIndexOp>(op->getLoc(), dim); }))); continue; } @@ -1281,13 +1281,13 @@ ArrayRef<int64_t>(tileSizes).take_front( std::min<size_t>(tileSizes.size(), maxDepth)), [&](int64_t t) -> Value { - return builder.create<ConstantIndexOp>(op->getLoc(), t); + return builder.create<arith::ConstantIndexOp>(op->getLoc(), t); })); } // Set all loops not partitioned to tile size 0. and those partitioned to // `flow.workgroup.size`. - auto zero = builder.create<ConstantIndexOp>(op->getLoc(), 0); + auto zero = builder.create<arith::ConstantIndexOp>(op->getLoc(), 0); SmallVector<Value, 4> useTileSizes(maxDepth, zero); llvm::DenseSet<unsigned> partitionedLoopsSet; partitionedLoopsSet.insert(partitionedLoops.begin(),
diff --git a/iree/compiler/Dialect/Flow/Transforms/FusionOfTensorOps.cpp b/iree/compiler/Dialect/Flow/Transforms/FusionOfTensorOps.cpp index c5eb2c2..ba0869b 100644 --- a/iree/compiler/Dialect/Flow/Transforms/FusionOfTensorOps.cpp +++ b/iree/compiler/Dialect/Flow/Transforms/FusionOfTensorOps.cpp
@@ -99,7 +99,7 @@ // simplistic heuristic to avoid duplicating ops that may be // expensive. // TODO: Add a cost model to allow ops to be duplicated. - if (!isBroadcast && !isa<ConstantOp>(producer) && + if (!isBroadcast && !isa<arith::ConstantOp>(producer) && !llvm::hasSingleElement(producerResult.getUsers())) return false; return llvm::all_of(producerResult.getUsers(), [](Operation *user) {
diff --git a/iree/compiler/Dialect/Flow/Transforms/OutlineLargeConstants.cpp b/iree/compiler/Dialect/Flow/Transforms/OutlineLargeConstants.cpp index 674397c..f80b73f 100644 --- a/iree/compiler/Dialect/Flow/Transforms/OutlineLargeConstants.cpp +++ b/iree/compiler/Dialect/Flow/Transforms/OutlineLargeConstants.cpp
@@ -27,7 +27,7 @@ // Returns true if |constantOp| is large enough to be considered for pooling. // Some constants are small enough that inlining them into the ringbuffer is // more efficient and fewer bindings. -static bool isConstantLarge(ConstantOp constantOp, +static bool isConstantLarge(arith::ConstantOp constantOp, size_t minLargeConstantSize) { auto type = constantOp.getType(); if (auto shapedType = type.dyn_cast<RankedTensorType>()) { @@ -43,12 +43,12 @@ // Returns a list of all large constants in the module. // Only walks top-level functions and ops to avoid pulling constants out of // executables. -static std::vector<ConstantOp> findLargeConstantsInModule( +static std::vector<arith::ConstantOp> findLargeConstantsInModule( ModuleOp moduleOp, size_t minLargeConstantSize) { - std::vector<ConstantOp> largeConstantOps; + std::vector<arith::ConstantOp> largeConstantOps; for (auto funcOp : moduleOp.getOps<mlir::FuncOp>()) { for (auto &block : funcOp.getBlocks()) { - for (auto constantOp : block.getOps<ConstantOp>()) { + for (auto constantOp : block.getOps<arith::ConstantOp>()) { if (isConstantLarge(constantOp, minLargeConstantSize)) { largeConstantOps.push_back(constantOp); } @@ -80,7 +80,8 @@ // Create all top-level util.globals from large constants in the module. OpBuilder moduleBuilder(&moduleOp.getBody()->front()); - std::vector<std::pair<ConstantOp, IREE::Util::GlobalOp>> replacements; + std::vector<std::pair<arith::ConstantOp, IREE::Util::GlobalOp>> + replacements; for (auto &largeConstantOp : findLargeConstantsInModule(moduleOp, minLargeConstantSize)) { std::string name; @@ -89,7 +90,7 @@ } while (moduleSymbols.lookup(name) != nullptr); auto globalOp = moduleBuilder.create<IREE::Util::GlobalOp>( largeConstantOp.getLoc(), name, /*isMutable=*/false, - largeConstantOp.getType(), largeConstantOp.getValue()); + largeConstantOp.getType(), largeConstantOp.value()); globalOp.setPrivate(); replacements.emplace_back(largeConstantOp, globalOp);
diff --git a/iree/compiler/Dialect/Flow/Transforms/PadLinalgOps.cpp b/iree/compiler/Dialect/Flow/Transforms/PadLinalgOps.cpp index 48cba52..356e268 100644 --- a/iree/compiler/Dialect/Flow/Transforms/PadLinalgOps.cpp +++ b/iree/compiler/Dialect/Flow/Transforms/PadLinalgOps.cpp
@@ -63,10 +63,10 @@ auto rhsPaddedType = RankedTensorType::get({newKSize, newNSize}, rhsType.getElementType()); - Value lhsPaddingValue = rewriter.create<ConstantOp>( + Value lhsPaddingValue = rewriter.create<arith::ConstantOp>( loc, rewriter.getZeroAttr(lhsType.getElementType())); - Value rhsPaddingValue = rewriter.create<ConstantOp>( + Value rhsPaddingValue = rewriter.create<arith::ConstantOp>( loc, rewriter.getZeroAttr(rhsType.getElementType())); auto createPadding = [&](ArrayRef<int64_t> padding) { @@ -103,7 +103,7 @@ } else { auto newResultType = RankedTensorType::get({newMSize, newNSize}, resultType.getElementType()); - auto resultPaddingValue = rewriter.create<ConstantOp>( + auto resultPaddingValue = rewriter.create<arith::ConstantOp>( loc, rewriter.getZeroAttr(resultType.getElementType())); Value paddedResult = linalg::PadTensorOp::createPadScalarOp( newResultType, result, resultPaddingValue, createPadding({0, 0}),
diff --git a/iree/compiler/Dialect/Flow/Transforms/PadTensorToSubTensorInsert.cpp b/iree/compiler/Dialect/Flow/Transforms/PadTensorToSubTensorInsert.cpp index d6c5ec7..7f3ad37 100644 --- a/iree/compiler/Dialect/Flow/Transforms/PadTensorToSubTensorInsert.cpp +++ b/iree/compiler/Dialect/Flow/Transforms/PadTensorToSubTensorInsert.cpp
@@ -79,7 +79,7 @@ expr = addValueOrAttr(expr, highPad[dim]); Value v = linalg::applyMapToValues( rewriter, loc, AffineMap::get(1, numSymbols, expr), mapValues)[0]; - if (auto cst = v.getDefiningOp<ConstantOp>()) { + if (auto cst = v.getDefiningOp<arith::ConstantOp>()) { outputShape.push_back(cst.value()); } else { outputShape.push_back(v); @@ -100,7 +100,8 @@ : public PadTensorToSubTensorInsertBase<PadTensorToSubTensorInsertPass> { void getDependentDialects(DialectRegistry ®istry) const override { registry.insert<linalg::LinalgDialect, memref::MemRefDialect, - StandardOpsDialect>(); + StandardOpsDialect, mlir::math::MathDialect, + mlir::arith::ArithmeticDialect>(); } void runOnOperation() override {
diff --git a/iree/compiler/Dialect/Flow/Transforms/PromoteI1ToI8Pass.cpp b/iree/compiler/Dialect/Flow/Transforms/PromoteI1ToI8Pass.cpp index 1deced1..37f8d60 100644 --- a/iree/compiler/Dialect/Flow/Transforms/PromoteI1ToI8Pass.cpp +++ b/iree/compiler/Dialect/Flow/Transforms/PromoteI1ToI8Pass.cpp
@@ -22,11 +22,12 @@ // Legalizes boolean (i1) constants to i8 with a linalg.generic operation // downcasting to i1. This occurs as IREE does not currently support tightly // packing and unpacking i1 buffers. -class ConvertBoolConstantPattern : public OpRewritePattern<mlir::ConstantOp> { +class ConvertBoolConstantPattern + : public OpRewritePattern<mlir::arith::ConstantOp> { public: - using OpRewritePattern<mlir::ConstantOp>::OpRewritePattern; + using OpRewritePattern<mlir::arith::ConstantOp>::OpRewritePattern; - LogicalResult matchAndRewrite(mlir::ConstantOp op, + LogicalResult matchAndRewrite(mlir::arith::ConstantOp op, PatternRewriter &rewriter) const override { Location loc = op.getLoc(); auto resultTy = op.getType().dyn_cast<ShapedType>(); @@ -43,7 +44,7 @@ if (!attr) return failure(); // Create a new ConstantOp that contains the same values as an int8. - auto newConst = rewriter.createOrFold<ConstantOp>( + auto newConst = rewriter.createOrFold<arith::ConstantOp>( loc, attr.mapValues(rewriter.getIntegerType(8), [&](APInt src) { return src.zext(8); })); @@ -78,7 +79,7 @@ getParallelIteratorTypeName()), [&](OpBuilder &nestedBuilder, Location nestedLoc, ValueRange blockArgs) { - auto cast = rewriter.create<TruncateIOp>( + auto cast = rewriter.create<arith::TruncIOp>( nestedLoc, rewriter.getIntegerType(1), blockArgs[0]); rewriter.create<linalg::YieldOp>(nestedLoc, cast->getResult(0)); @@ -95,7 +96,8 @@ class PromoteI1ToI8Pass : public PromoteI1ToI8Base<PromoteI1ToI8Pass> { public: void getDependentDialects(DialectRegistry ®istry) const override { - registry.insert<linalg::LinalgDialect, mlir::StandardOpsDialect>(); + registry.insert<linalg::LinalgDialect, mlir::StandardOpsDialect, + mlir::math::MathDialect, mlir::arith::ArithmeticDialect>(); } void runOnOperation() override {
diff --git a/iree/compiler/Dialect/Flow/Transforms/PromoteTensorLoads.cpp b/iree/compiler/Dialect/Flow/Transforms/PromoteTensorLoads.cpp index 2e6d158..1114875 100644 --- a/iree/compiler/Dialect/Flow/Transforms/PromoteTensorLoads.cpp +++ b/iree/compiler/Dialect/Flow/Transforms/PromoteTensorLoads.cpp
@@ -37,12 +37,12 @@ if (tensorType.getElementType().isInteger(1)) { auto i1Type = rewriter.getI1Type(); auto i8Type = rewriter.getIntegerType(8); - auto convertedOperand = rewriter.createOrFold<ZeroExtendIOp>( + auto convertedOperand = rewriter.createOrFold<arith::ExtUIOp>( op.getLoc(), args[0], RankedTensorType::get(tensorType.getShape(), i8Type)); auto i8Value = rewriter.createOrFold<IREE::Flow::TensorLoadOp>( op.getLoc(), i8Type, convertedOperand, op.indices()); - rewriter.replaceOpWithNewOp<TruncateIOp>(op, i1Type, i8Value); + rewriter.replaceOpWithNewOp<arith::TruncIOp>(op, i1Type, i8Value); } else { rewriter.replaceOpWithNewOp<IREE::Flow::TensorLoadOp>( op, tensorType.getElementType(), op.tensor(), op.indices()); @@ -67,7 +67,9 @@ : public PromoteTensorLoadsBase<PromoteTensorLoadsPass> { public: void getDependentDialects(DialectRegistry ®istry) const override { - registry.insert<FlowDialect, StandardOpsDialect, tensor::TensorDialect>(); + registry + .insert<FlowDialect, StandardOpsDialect, mlir::arith::ArithmeticDialect, + mlir::math::MathDialect, tensor::TensorDialect>(); } void runOnOperation() override { @@ -76,7 +78,9 @@ OwningRewritePatternList conversionPatterns(&getContext()); conversionTarget.addLegalDialect<IREE::Flow::FlowDialect>(); - conversionTarget.addLegalDialect<StandardOpsDialect>(); + conversionTarget + .addLegalDialect<StandardOpsDialect, mlir::arith::ArithmeticDialect, + mlir::math::MathDialect>(); setupStandardToFlowTensorLoadLegality(context, conversionTarget); populateStandardToFlowTensorLoadPatterns(context, conversionPatterns);
diff --git a/iree/compiler/Dialect/Flow/Transforms/test/convert_linalg_tensor_ops_after.mlir b/iree/compiler/Dialect/Flow/Transforms/test/convert_linalg_tensor_ops_after.mlir index a2a906d..475ecb8 100644 --- a/iree/compiler/Dialect/Flow/Transforms/test/convert_linalg_tensor_ops_after.mlir +++ b/iree/compiler/Dialect/Flow/Transforms/test/convert_linalg_tensor_ops_after.mlir
@@ -1,8 +1,8 @@ // RUN: iree-opt -iree-flow-convert-to-flow-after-dispatch-formation -canonicalize -cse -split-input-file %s | IreeFileCheck %s func @turn_fill_into_splat(%arg0: tensor<?x?xf32>, %arg1: tensor<f32>, %arg2: index, %arg3: index, %arg4: index, %arg5: index) -> tensor<?x?xf32> { - %c0 = constant 0 : index - %c1 = constant 1 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index %0 = tensor.extract %arg1[] : tensor<f32> %1 = tensor.dim %arg0, %c0 : tensor<?x?xf32> %2 = tensor.dim %arg0, %c1 : tensor<?x?xf32> @@ -22,8 +22,8 @@ // CHECK-SAME: %[[ARG3:[a-zA-Z0-9]+]]: index // CHECK-SAME: %[[ARG4:[a-zA-Z0-9]+]]: index // CHECK-SAME: %[[ARG5:[a-zA-Z0-9]+]]: index -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK: %[[VAL:.+]] = flow.tensor.load %[[ARG1]] : tensor<f32> // CHECK-DAG: %[[D0:.+]] = tensor.dim %[[ARG0]], %[[C0]] // CHECK-DAG: %[[D1:.+]] = tensor.dim %[[ARG0]], %[[C1]]
diff --git a/iree/compiler/Dialect/Flow/Transforms/test/deduplicate_executables.mlir b/iree/compiler/Dialect/Flow/Transforms/test/deduplicate_executables.mlir index e58e05b..1e2e929 100644 --- a/iree/compiler/Dialect/Flow/Transforms/test/deduplicate_executables.mlir +++ b/iree/compiler/Dialect/Flow/Transforms/test/deduplicate_executables.mlir
@@ -12,7 +12,7 @@ } // CHECK-LABEL: func @single_executable func @single_executable(%arg0: tensor<4xf32>) -> tensor<4xf32> { - %c4 = constant 4 : index + %c4 = arith.constant 4 : index // CHECK: %0 = flow.dispatch @single_executable_ex_0::@single_executable_entry_0[%c4](%arg0) : (tensor<4xf32>) -> tensor<4xf32> %0 = flow.dispatch @single_executable_ex_0::@single_executable_entry_0[%c4] (%arg0) : (tensor<4xf32>) -> tensor<4xf32> return %0 : tensor<4xf32> @@ -52,7 +52,7 @@ } // CHECK-LABEL: func @duplicate_executables func @duplicate_executables(%arg0: tensor<4xf32>) -> tensor<4xf32> { - %c4 = constant 4 : index + %c4 = arith.constant 4 : index // CHECK: %0 = flow.dispatch @duplicate_executables_ex_0::@duplicate_executables_entry_0[%c4](%arg0) : (tensor<4xf32>) -> tensor<4xf32> %0 = flow.dispatch @duplicate_executables_ex_0::@duplicate_executables_entry_0[%c4] (%arg0) : (tensor<4xf32>) -> tensor<4xf32> // CHECK: %1 = flow.dispatch @duplicate_executables_ex_0::@duplicate_executables_entry_0[%c4](%arg0) : (tensor<4xf32>) -> tensor<4xf32> @@ -86,7 +86,7 @@ } // CHECK-LABEL: func @same_ops_diff_operands func @same_ops_diff_operands(%arg0: tensor<2xi32>, %arg1: tensor<2xi32>) -> tensor<2xi32> { - %c4 = constant 4 : index + %c4 = arith.constant 4 : index // CHECK: %0 = flow.dispatch @same_ops_diff_operands_ex_0::@entry_0[%c4](%arg0, %arg1) : (tensor<2xi32>, tensor<2xi32>) -> tensor<2xi32> %0 = flow.dispatch @same_ops_diff_operands_ex_0::@entry_0[%c4] (%arg0, %arg1) : (tensor<2xi32>, tensor<2xi32>) -> tensor<2xi32> // CHECK: %1 = flow.dispatch @same_ops_diff_operands_ex_1::@entry_1[%c4](%arg0, %arg1) : (tensor<2xi32>, tensor<2xi32>) -> tensor<2xi32> @@ -128,7 +128,7 @@ } // CHECK-LABEL: func @multiple_entry_points func @multiple_entry_points(%arg0: tensor<4xf32>) -> tensor<4xf32> { - %c4 = constant 4 : index + %c4 = arith.constant 4 : index // CHECK: %0 = flow.dispatch @multiple_entry_points_ex_0::@multiple_entry_points_0_entry_0[%c4](%arg0) : (tensor<4xf32>) -> tensor<4xf32> %0 = flow.dispatch @multiple_entry_points_ex_0::@multiple_entry_points_0_entry_0[%c4] (%arg0) : (tensor<4xf32>) -> tensor<4xf32> // CHECK: %1 = flow.dispatch @multiple_entry_points_ex_0::@multiple_entry_points_0_entry_1[%c4](%arg0) : (tensor<4xf32>) -> tensor<4xf32> @@ -164,7 +164,7 @@ } // CHECK-LABEL: func @different_types func @different_types(%arg0: tensor<4xf32>) -> tensor<4xi1> { - %c4 = constant 4 : index + %c4 = arith.constant 4 : index // CHECK: %0 = flow.dispatch @different_types_float_ex::@different_types_float_entry[%c4](%arg0) : (tensor<4xf32>) -> tensor<4xi1> %0 = flow.dispatch @different_types_float_ex::@different_types_float_entry[%c4] (%arg0) : (tensor<4xf32>) -> tensor<4xi1> // CHECK: %1 = flow.dispatch @different_types_int_ex::@different_types_int_entry[%c4](%arg0) : (tensor<4xf32>) -> tensor<4xi1> @@ -179,7 +179,7 @@ flow.dispatch.entry @nested_ops_entry_0 builtin.module { func @nested_ops_entry_0(%input: tensor<1x4xi32>) -> tensor<1xi32> { - %0 = constant dense<0> : tensor<i32> + %0 = arith.constant dense<0> : tensor<i32> %1 = "mhlo.reduce"(%input, %0) ( { ^bb0(%arg0: tensor<i32>, %arg1: tensor<i32>): // no predecessors %3 = "mhlo.add"(%arg0, %arg1) : (tensor<i32>, tensor<i32>) -> tensor<i32> @@ -194,7 +194,7 @@ flow.dispatch.entry @nested_ops_entry_1 builtin.module { func @nested_ops_entry_1(%input: tensor<1x4xi32>) -> tensor<1xi32> { - %0 = constant dense<0> : tensor<i32> + %0 = arith.constant dense<0> : tensor<i32> %1 = "mhlo.reduce"(%input, %0) ( { ^bb0(%arg0: tensor<i32>, %arg1: tensor<i32>): // no predecessors %3 = "mhlo.add"(%arg0, %arg1) : (tensor<i32>, tensor<i32>) -> tensor<i32> @@ -209,7 +209,7 @@ flow.dispatch.entry @nested_ops_entry_2 builtin.module { func @nested_ops_entry_2(%input: tensor<1x4xi32>) -> tensor<1xi32> { - %0 = constant dense<0> : tensor<i32> + %0 = arith.constant dense<0> : tensor<i32> %1 = "mhlo.reduce"(%input, %0) ( { ^bb0(%arg0: tensor<i32>, %arg1: tensor<i32>): // no predecessors %3 = "mhlo.maximum"(%arg0, %arg1) : (tensor<i32>, tensor<i32>) -> tensor<i32> @@ -221,7 +221,7 @@ } // CHECK-LABEL: func @nested_ops func @nested_ops(%arg0: tensor<1x4xi32>) -> tensor<1xi32> { - %c4 = constant 4 : index + %c4 = arith.constant 4 : index // CHECK: %0 = flow.dispatch @nested_ops_ex_0::@nested_ops_entry_0[%c4](%arg0) : (tensor<1x4xi32>) -> tensor<1xi32> %0 = flow.dispatch @nested_ops_ex_0::@nested_ops_entry_0[%c4] (%arg0) : (tensor<1x4xi32>) -> tensor<1xi32> // CHECK: %1 = flow.dispatch @nested_ops_ex_0::@nested_ops_entry_0[%c4](%arg0) : (tensor<1x4xi32>) -> tensor<1xi32> @@ -238,7 +238,7 @@ flow.dispatch.entry @attributes_entry_0 builtin.module { func @attributes_entry_0(%input: tensor<1x4xi32>) -> tensor<1xi32> { - %0 = constant dense<0> : tensor<i32> + %0 = arith.constant dense<0> : tensor<i32> %1 = "mhlo.reduce"(%input, %0) ( { ^bb0(%arg0: tensor<i32>, %arg1: tensor<i32>): // no predecessors %3 = "mhlo.maximum"(%arg0, %arg1) : (tensor<i32>, tensor<i32>) -> tensor<i32> @@ -254,7 +254,7 @@ flow.dispatch.entry @attributes_entry_1 builtin.module { func @attributes_entry_1(%input: tensor<1x4xi32>) -> tensor<1xi32> { - %0 = constant dense<0> : tensor<i32> + %0 = arith.constant dense<0> : tensor<i32> %1 = "mhlo.reduce"(%input, %0) ( { ^bb0(%arg0: tensor<i32>, %arg1: tensor<i32>): // no predecessors %3 = "mhlo.maximum"(%arg0, %arg1) : (tensor<i32>, tensor<i32>) -> tensor<i32> @@ -271,7 +271,7 @@ flow.dispatch.entry @attributes_entry_2 builtin.module { func @attributes_entry_2(%input: tensor<1x4xi32>) -> tensor<1xi32> { - %0 = constant dense<0> : tensor<i32> + %0 = arith.constant dense<0> : tensor<i32> %1 = "mhlo.reduce"(%input, %0) ( { ^bb0(%arg0: tensor<i32>, %arg1: tensor<i32>): // no predecessors %3 = "mhlo.maximum"(%arg0, %arg1) : (tensor<i32>, tensor<i32>) -> tensor<i32> @@ -289,9 +289,9 @@ flow.dispatch.entry @entry_0 builtin.module { func @entry_0(%arg0: i32, %arg1: i32) -> i32 { - %c0 = constant 0 : i32 - %c1 = constant 1 : i32 - %eqz = cmpi eq, %arg0, %arg1 : i32 + %c0 = arith.constant 0 : i32 + %c1 = arith.constant 1 : i32 + %eqz = arith.cmpi eq, %arg0, %arg1 : i32 cond_br %eqz, ^bb_a(%c0 : i32), ^bb_b(%c1 : i32) ^bb_a(%bb_a_arg0 : i32): return %bb_a_arg0 : i32 @@ -305,9 +305,9 @@ flow.dispatch.entry @entry_1 builtin.module { func @entry_0(%arg0: i32, %arg1: i32) -> i32 { - %c0 = constant 0 : i32 - %c1 = constant 1 : i32 - %eqz = cmpi eq, %arg0, %arg1 : i32 + %c0 = arith.constant 0 : i32 + %c1 = arith.constant 1 : i32 + %eqz = arith.cmpi eq, %arg0, %arg1 : i32 cond_br %eqz, ^bb_b(%c0 : i32), ^bb_b(%c1 : i32) ^bb_a(%bb_a_arg0 : i32): return %bb_a_arg0 : i32
diff --git a/iree/compiler/Dialect/Flow/Transforms/test/dispatch_linalg_on_tensors.mlir b/iree/compiler/Dialect/Flow/Transforms/test/dispatch_linalg_on_tensors.mlir index 732502e..a4b0747 100644 --- a/iree/compiler/Dialect/Flow/Transforms/test/dispatch_linalg_on_tensors.mlir +++ b/iree/compiler/Dialect/Flow/Transforms/test/dispatch_linalg_on_tensors.mlir
@@ -45,8 +45,8 @@ // ----- func @tile_generic_op_alone(%A: tensor<?x?xf32>, %B: tensor<?xf32>) -> tensor<?x?xf32> { - %c0 = constant 0 : index - %c1 = constant 1 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index %d0 = tensor.dim %A, %c0 : tensor<?x?xf32> %d1 = tensor.dim %A, %c1 : tensor<?x?xf32> %0 = linalg.init_tensor [%d0, %d1] : tensor<?x?xf32> @@ -58,7 +58,7 @@ ins (%A, %B: tensor<?x?xf32>, tensor<?xf32>) outs (%0 : tensor<?x?xf32>) { ^bb0(%arg0 : f32, %arg1 : f32, %arg2 : f32): - %2 = addf %arg0, %arg1 : f32 + %2 = arith.addf %arg0, %arg1 : f32 linalg.yield %2 : f32 } -> tensor<?x?xf32> return %1 : tensor<?x?xf32> @@ -66,8 +66,8 @@ // CHECK: func @tile_generic_op_alone // CHECK-SAME: %[[ARG0:[a-zA-Z0-9_]+]]: tensor<?x?xf32> // CHECK-SAME: %[[ARG1:[a-zA-Z0-9_]+]]: tensor<?xf32> -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK-DAG: %[[D0:.+]] = tensor.dim %[[ARG0]], %[[C0]] // CHECK-DAG: %[[D1:.+]] = tensor.dim %[[ARG0]], %[[C1]] // CHECK: flow.dispatch.workgroups @@ -86,9 +86,9 @@ // ----- func @fuse_matmul_with_fill(%A : tensor<?x?xf32>, %B : tensor<?x?xf32>) -> tensor<?x?xf32> { - %zero = constant 0.0 : f32 - %c0 = constant 0 : index - %c1 = constant 1 : index + %zero = arith.constant 0.0 : f32 + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index %M = tensor.dim %A, %c0 : tensor<?x?xf32> %N = tensor.dim %B, %c1 : tensor<?x?xf32> %0 = linalg.init_tensor [%M, %N] : tensor<?x?xf32> @@ -100,8 +100,8 @@ // CHECK: func @fuse_matmul_with_fill // CHECK-SAME: %[[ARG0:[a-zA-Z0-9_]+]]: tensor<?x?xf32> // CHECK-SAME: %[[ARG1:[a-zA-Z0-9_]+]]: tensor<?x?xf32> -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK: %[[M:.+]] = tensor.dim %[[ARG0]], %[[C0]] // CHECK: %[[N:.+]] = tensor.dim %[[ARG1]], %[[C1]] // CHECK: flow.dispatch.workgroups[%[[N]], %[[M]], %[[C1]]] @@ -111,7 +111,7 @@ // CHECK-SAME: %[[ARG2:[a-zA-Z0-9_]+]]: index // CHECK-SAME: %[[ARG3:[a-zA-Z0-9_]+]]: index // CHECK-SAME: %[[ARG6:[a-zA-Z0-9_]+]]: !flow.dispatch.tensor<writeonly:?x?xf32>) { -// CHECK: %[[ZERO:.+]] = constant 0.000000e+00 : f32 +// CHECK: %[[ZERO:.+]] = arith.constant 0.000000e+00 : f32 // CHECK: scf.for // CHECK: scf.for // CHECK-DAG: %[[LHS_TILE:.+]] = flow.dispatch.tensor.load %[[ARG4]], {{.*}} @@ -128,10 +128,10 @@ // ----- func @keep_separate_dispatches_for_producer(%A : tensor<?x?xf32>, %B : tensor<?x?xf32>) -> tensor<?x?xf32> { - %zero = constant 0.0 : f32 - %one = constant 1.0 : f32 - %c0 = constant 0 : index - %c1 = constant 1 : index + %zero = arith.constant 0.0 : f32 + %one = arith.constant 1.0 : f32 + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index %M = tensor.dim %A, %c0 : tensor<?x?xf32> %N = tensor.dim %B, %c1 : tensor<?x?xf32> %K = tensor.dim %A, %c1 : tensor<?x?xf32> @@ -144,7 +144,7 @@ iterator_types = ["parallel", "parallel"]} ins(%A : tensor<?x?xf32>) outs(%2 : tensor<?x?xf32>) { ^bb0(%arg0 : f32, %arg1 : f32): - %4 = addf %arg0, %one : f32 + %4 = arith.addf %arg0, %one : f32 linalg.yield %4 : f32 } -> tensor<?x?xf32> %4 = linalg.matmul ins(%3, %B : tensor<?x?xf32>, tensor<?x?xf32>) @@ -154,8 +154,8 @@ // CHECK: func @keep_separate_dispatches_for_producer // CHECK-SAME: %[[ARG0:[a-zA-Z0-9_]+]]: tensor<?x?xf32> // CHECK-SAME: %[[ARG1:[a-zA-Z0-9_]+]]: tensor<?x?xf32> -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK-DAG: %[[M:.+]] = tensor.dim %[[ARG0]], %[[C0]] // CHECK-DAG: %[[N:.+]] = tensor.dim %[[ARG1]], %[[C1]] // CHECK-DAG: %[[K:.+]] = tensor.dim %[[ARG0]], %[[C1]] @@ -163,7 +163,7 @@ // CHECK-SAME: (%[[ARG0]]) // CHECK-NEXT: (%[[ARG2:[a-zA-Z0-9_]+]]: !flow.dispatch.tensor<readonly:?x?xf32> // CHECK-SAME: %[[ARG5:[a-zA-Z0-9_]+]]: !flow.dispatch.tensor<writeonly:?x?xf32>) { -// CHECK: %[[ONE:.+]] = constant 1.0 +// CHECK: %[[ONE:.+]] = arith.constant 1.0 // CHECK-DAG: %[[INPUT:.+]] = flow.dispatch.tensor.load %[[ARG2]] // CHECK-DAG: %[[INIT:.+]] = linalg.init_tensor // CHECK: %[[RESULT:.+]] = linalg.generic @@ -173,7 +173,7 @@ // CHECK: flow.return // CHECK: } // CHECK: flow.dispatch.workgroups[%[[N]], %[[M]], %[[C1]]] -// CHECK: %[[ZERO:.+]] = constant 0.0 +// CHECK: %[[ZERO:.+]] = arith.constant 0.0 // CHECK: scf.for // CHECK: scf.for // CHECK: %[[INIT_TILE:.+]] = linalg.init_tensor @@ -189,7 +189,7 @@ // NOCHECK-SAME: %[[ARG5:[a-zA-Z0-9_]+]]: !flow.dispatch.tensor<readonly:?x?xf32> // NOCHECK-SAME: %[[ARG6:[a-zA-Z0-9_]+]]: !flow.dispatch.tensor<readonly:?x?xf32> // NOCHECK-SAME: %[[ARG7:[a-zA-Z0-9_]+]]: !flow.dispatch.tensor<writeonly:?x?xf32>) { -// NOCHECK: %[[ZERO:.+]] = constant 0.0 +// NOCHECK: %[[ZERO:.+]] = arith.constant 0.0 // NOCHECK: scf.for // NOCHECK: scf.for // NOCHECK-DAG: %[[LHS_TILE_2:.+]] = flow.dispatch.tensor.load %[[ARG6]], {{.*}} @@ -213,8 +213,8 @@ // CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0, s1] -> (s0 * s1)> // CHECK: func @fuse_reshape_op // CHECK-SAME: (%[[ARG0:.+]]: tensor<?x?xf32>) -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK-DAG: %[[D0:.+]] = tensor.dim %[[ARG0]], %[[C0]] // CHECK-DAG: %[[D1:.+]] = tensor.dim %[[ARG0]], %[[C1]] // CHECK: %[[WORKLOAD:.+]] = affine.apply #[[MAP0]]()[%[[D0]], %[[D1]]] @@ -230,10 +230,10 @@ func @tile_4d_generic_op_alone (%A: tensor<?x?x?x?xf32>, %B: tensor<?x?x?x?xf32>) -> tensor<?x?x?x?xf32> { - %c0 = constant 0 : index - %c1 = constant 1 : index - %c2 = constant 2 : index - %c3 = constant 3 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + %c2 = arith.constant 2 : index + %c3 = arith.constant 3 : index %d0 = tensor.dim %A, %c0 : tensor<?x?x?x?xf32> %d1 = tensor.dim %A, %c1 : tensor<?x?x?x?xf32> %d2 = tensor.dim %A, %c2 : tensor<?x?x?x?xf32> @@ -247,7 +247,7 @@ ins (%A, %B: tensor<?x?x?x?xf32>, tensor<?x?x?x?xf32>) outs (%0 : tensor<?x?x?x?xf32>) { ^bb0(%arg0 : f32, %arg1 : f32, %arg2 : f32): - %2 = addf %arg0, %arg1 : f32 + %2 = arith.addf %arg0, %arg1 : f32 linalg.yield %2 : f32 } -> tensor<?x?x?x?xf32> return %1 : tensor<?x?x?x?xf32> @@ -257,10 +257,10 @@ // CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0, s1] -> (s0 * s1)> // CHECK: func @tile_4d_generic_op_alone // CHECK-SAME: %[[ARG0:[a-zA-Z0-9_]+]]: tensor<?x?x?x?xf32> -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index -// CHECK-DAG: %[[C2:.+]] = constant 2 : index -// CHECK-DAG: %[[C3:.+]] = constant 3 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index +// CHECK-DAG: %[[C2:.+]] = arith.constant 2 : index +// CHECK-DAG: %[[C3:.+]] = arith.constant 3 : index // CHECK-DAG: %[[D0:.+]] = tensor.dim %[[ARG0]], %[[C0]] // CHECK-DAG: %[[D1:.+]] = tensor.dim %[[ARG0]], %[[C1]] // CHECK-DAG: %[[D2:.+]] = tensor.dim %[[ARG0]], %[[C2]] @@ -276,9 +276,9 @@ (%lhs : tensor<?xf32>, %rhs1 : tensor<4x?xf32>, %rhs2 : tensor<4x?xf32>) -> (tensor<?x?xf32>, tensor<?x?xf32>) { - %cst = constant 0.0 : f32 - %c0 = constant 0 : index - %c1 = constant 1 : index + %cst = arith.constant 0.0 : f32 + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index %0 = linalg.tensor_expand_shape %lhs [[0, 1]] : tensor<?xf32> into tensor<?x4xf32> %m = tensor.dim %0, %c0 : tensor<?x4xf32> @@ -302,8 +302,8 @@ // CHECK-SAME: %[[ARG0:[a-zA-Z0-9_]+]]: tensor<?xf32> // CHECK-SAME: %[[RHS1:[a-zA-Z0-9_]+]]: tensor<4x?xf32> // CHECK-SAME: %[[RHS2:[a-zA-Z0-9_]+]]: tensor<4x?xf32> -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK: %[[D0:.+]] = tensor.dim %[[ARG0]], %[[C0]] // CHECK-DAG: %[[M:.+]] = affine.apply #[[MAP]]()[%[[D0]]] // CHECK-DAG: %[[N1:.+]] = tensor.dim %[[ARG1]], %[[C1]] @@ -321,8 +321,8 @@ %arg0: tensor<?x?xf32>, %arg1: tensor<f32>, %arg2: index, %arg3: index, %arg4: index, %arg5: index) -> tensor<?x?xf32> { - %c0 = constant 0 : index - %c1 = constant 1 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index %0 = tensor.extract %arg1[] : tensor<f32> %1 = tensor.dim %arg0, %c0 : tensor<?x?xf32> %2 = tensor.dim %arg0, %c1 : tensor<?x?xf32> @@ -342,8 +342,8 @@ // CHECK-LABEL: func @pass_constant_through() func @pass_constant_through() -> tensor<2x2x3xi32> { - // CHECK: %[[CST:.+]] = constant dense<{{.+}}> : tensor<2x2x3xi32> - %cst = constant dense<[[[1, 2, 3], [4, 5, 6]], [[7, 8, 9], [10, 11, 12]]]> : tensor<2x2x3xi32> + // CHECK: %[[CST:.+]] = arith.constant dense<{{.+}}> : tensor<2x2x3xi32> + %cst = arith.constant dense<[[[1, 2, 3], [4, 5, 6]], [[7, 8, 9], [10, 11, 12]]]> : tensor<2x2x3xi32> // CHECK: return %[[CST]] return %cst : tensor<2x2x3xi32> } @@ -353,7 +353,7 @@ // CHECK-LABEL: func @fuse_matmul_with_generic_op func @fuse_matmul_with_generic_op(%A: tensor<?x?xf32>, %B: tensor<?x?xf32>, %C: tensor<?x?xf32>) -> tensor<?x?xf32> { - %f12 = constant 12.0 : f32 + %f12 = arith.constant 12.0 : f32 // linalg.generic is fused inside the dispatch region and becomes dead. // CHECK-NOT: generic @@ -384,9 +384,9 @@ // CHECK-SAME: %[[ARG2:[a-zA-Z0-9_]+]]: tensor<?x?xf32> func @keep_original_producer_uses(%A: tensor<?x?xf32>, %B: tensor<?x?xf32>, %C: tensor<?x?xf32>) -> (tensor<?x?xf32>, tensor<?x?xf32>) { - %f12 = constant 12.0 : f32 - // CHECK-DAG: %[[C0:.+]] = constant 0 : index - // CHECK-DAG: %[[C1:.+]] = constant 1 : index + %f12 = arith.constant 12.0 : f32 + // CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index + // CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK-DAG: %[[D0:.+]] = tensor.dim %[[ARG2]], %[[C0]] // CHECK-DAG: %[[D1:.+]] = tensor.dim %[[ARG2]], %[[C1]] // CHECK: %[[origCC:.+]] = flow.dispatch.workgroups[%[[D1]], %[[D0]], %[[C1]]](%[[ARG2]]) @@ -425,7 +425,7 @@ func @conv2d(%input: tensor<1x225x225x16xf32>, %filter: tensor<3x3x16x32xf32>) -> tensor<1x112x112x32xf32> { %0 = linalg.init_tensor [1, 112, 112, 32] : tensor<1x112x112x32xf32> - %cst = constant 0.000000e+00 : f32 + %cst = arith.constant 0.000000e+00 : f32 %1 = linalg.fill(%cst, %0) : f32, tensor<1x112x112x32xf32> -> tensor<1x112x112x32xf32> %2 = linalg.conv_2d_nhwc_hwcf {dilations = dense<1> : tensor<2xi64>, strides = dense<2> : tensor<2xi64>} @@ -443,7 +443,7 @@ // ----- func @depthwise_conv2d(%input: tensor<1x113x113x96xf32>, %filter: tensor<3x3x96xf32>) -> tensor<1x56x56x96xf32> { - %cst = constant 0.000000e+00 : f32 + %cst = arith.constant 0.000000e+00 : f32 %1 = linalg.init_tensor [1, 56, 56, 96] : tensor<1x56x56x96xf32> %2 = linalg.fill(%cst, %1) : f32, tensor<1x56x56x96xf32> -> tensor<1x56x56x96xf32> %4 = linalg.depthwise_conv2D_nhw {dilations = dense<1> : tensor<2xi64>, strides = dense<2> : tensor<2xi64>} ins(%input, %filter : tensor<1x113x113x96xf32>, tensor<3x3x96xf32>) outs(%2 : tensor<1x56x56x96xf32>) -> tensor<1x56x56x96xf32> @@ -472,8 +472,8 @@ // CHECK-SAME: %[[ARG1:[a-zA-Z0-9_]+]]: tensor<?x?xf32> // CHECK-SAME: %[[ARG2:[a-zA-Z0-9_]+]]: index // CHECK-SAME: %[[ARG3:[a-zA-Z0-9_]+]]: index -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK-DAG: %[[D0:.+]] = tensor.dim %[[ARG0]], %[[C0]] // CHECK-DAG: %[[D1:.+]] = tensor.dim %[[ARG0]], %[[C1]] // CHECK-DAG: %[[D2:.+]] = tensor.dim %[[ARG1]], %[[C0]] @@ -514,7 +514,7 @@ // ----- func @fuse_non_tiled_reduction_fill(%input1: tensor<1000xf32>, %input2: tensor<1000xf32>, %offset: tensor<f32>) -> tensor<f32> { - %zero = constant 0.0 : f32 + %zero = arith.constant 0.0 : f32 %init = linalg.init_tensor [] : tensor<f32> %fill = linalg.fill(%zero, %init) : f32, tensor<f32> -> tensor<f32> %reduce = linalg.generic { @@ -523,10 +523,10 @@ ins(%input1, %input2, %offset : tensor<1000xf32>, tensor<1000xf32>, tensor<f32>) outs(%fill : tensor<f32>) { ^bb0(%arg1: f32, %arg2: f32, %arg3: f32, %arg4: f32): - %555 = addf %arg1, %arg2 : f32 - %556 = subf %555, %arg3 : f32 + %555 = arith.addf %arg1, %arg2 : f32 + %556 = arith.subf %555, %arg3 : f32 %557 = math.exp %556 : f32 - %558 = addf %557, %arg4 : f32 + %558 = arith.addf %557, %arg4 : f32 linalg.yield %558 : f32 } -> tensor<f32> return %reduce : tensor<f32> @@ -534,7 +534,7 @@ // CHECK-LABEL: func @fuse_non_tiled_reduction_fill -// CHECK: %[[C1:.+]] = constant 1 : index +// CHECK: %[[C1:.+]] = arith.constant 1 : index // CHECK: flow.dispatch.workgroups[%[[C1]], %[[C1]], %[[C1]]]({{.+}}) : (tensor<1000xf32>, tensor<1000xf32>, tensor<f32>) -> tensor<f32> = // CHECK-NEXT: (%[[INPUT1:[a-z0-9]+]]: !flow.dispatch.tensor<readonly:1000xf32>, // CHECK-SAME: %[[INPUT2:[a-z0-9]+]]: !flow.dispatch.tensor<readonly:1000xf32>, @@ -572,11 +572,11 @@ tensor<?xf32>, tensor<?xf32>, tensor<?xf32>) outs(%8 : tensor<?xf32>) { ^bb0(%arg4: i32, %arg5: f32, %arg6: f32, %arg7: f32, %arg8: f32, %arg9: f32): - %10 = addf %arg5, %arg6 : f32 - %11 = addf %arg7, %arg8 : f32 - %12 = addf %10, %11 : f32 - %13 = sitofp %arg4 : i32 to f32 - %14 = addf %12, %13 : f32 + %10 = arith.addf %arg5, %arg6 : f32 + %11 = arith.addf %arg7, %arg8 : f32 + %12 = arith.addf %10, %11 : f32 + %13 = arith.sitofp %arg4 : i32 to f32 + %14 = arith.addf %12, %13 : f32 linalg.yield %14 : f32 } -> tensor<?xf32> return %9 : tensor<?xf32> @@ -627,11 +627,11 @@ tensor<?xf32>, tensor<?xf32>, tensor<?xf32>) outs(%8 : tensor<?xf32>) { ^bb0(%arg4: i32, %arg5: f32, %arg6: f32, %arg7: f32, %arg8: f32, %arg9: f32): - %10 = addf %arg5, %arg6 : f32 - %11 = addf %arg7, %arg8 : f32 - %12 = addf %10, %11 : f32 - %13 = sitofp %arg4 : i32 to f32 - %14 = addf %12, %13 : f32 + %10 = arith.addf %arg5, %arg6 : f32 + %11 = arith.addf %arg7, %arg8 : f32 + %12 = arith.addf %10, %11 : f32 + %13 = arith.sitofp %arg4 : i32 to f32 + %14 = arith.addf %12, %13 : f32 linalg.yield %14 : f32 } -> tensor<?xf32> return %9 : tensor<?xf32> @@ -660,17 +660,17 @@ // ----- func @inline_dag_3(%240 : tensor<9xi32>, %244 : tensor<18xi32>, %247 : tensor<i32>) -> tensor<9xi1> { - %c9 = constant 9 : index - %c5_i32 = constant 5 : i32 - %c0_i32 = constant 0 : i32 - %c9_i32 = constant 9 : i32 + %c9 = arith.constant 9 : index + %c5_i32 = arith.constant 5 : i32 + %c0_i32 = arith.constant 0 : i32 + %c9_i32 = arith.constant 9 : i32 %245 = flow.tensor.update %240, %244[%c9] : tensor<9xi32> -> %244 as tensor<18xi32> %248 = tensor.extract %247[] : tensor<i32> - %249 = cmpi slt, %248, %c9_i32 : i32 + %249 = arith.cmpi slt, %248, %c9_i32 : i32 %250 = select %249, %248, %c9_i32 : i32 - %251 = cmpi sgt, %250, %c0_i32 : i32 + %251 = arith.cmpi sgt, %250, %c0_i32 : i32 %252 = select %251, %250, %c0_i32 : i32 - %253 = index_cast %252 : i32 to index + %253 = arith.index_cast %252 : i32 to index %254 = tensor.extract_slice %245[%253] [9] [1] : tensor<18xi32> to tensor<9xi32> %255 = linalg.init_tensor [9] : tensor<9xi1> %256 = linalg.generic { @@ -678,7 +678,7 @@ iterator_types = ["parallel"]} ins(%254 : tensor<9xi32>) outs(%255 : tensor<9xi1>) { ^bb0(%arg20: i32, %arg21: i1): // no predecessors - %849 = cmpi eq, %arg20, %c5_i32 : i32 + %849 = arith.cmpi eq, %arg20, %c5_i32 : i32 linalg.yield %849 : i1 } -> tensor<9xi1> return %256 : tensor<9xi1> @@ -695,16 +695,16 @@ // CHECK-SAME: (%[[ARG2]], %[[UPDATE]]) // CHECK-NEXT: (%[[ARG3:.+]]: !flow.dispatch.tensor<readonly:i32>, // CHECK-SAME: %[[ARG4:.+]]: !flow.dispatch.tensor<readonly:18xi32>, -// CHECK-DAG: %[[C5:.+]] = constant 5 : i32 -// CHECK-DAG: %[[C0:.+]] = constant 0 : i32 -// CHECK-DAG: %[[C9:.+]] = constant 9 : i32 +// CHECK-DAG: %[[C5:.+]] = arith.constant 5 : i32 +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : i32 +// CHECK-DAG: %[[C9:.+]] = arith.constant 9 : i32 // CHECK-DAG: %[[ARG3V:.+]] = flow.dispatch.tensor.load %[[ARG3]] // CHECK-DAG: %[[EXTRACT:.+]] = tensor.extract %[[ARG3V]] -// CHECK-DAG: %[[CMP1:.+]] = cmpi slt, %[[EXTRACT]] +// CHECK-DAG: %[[CMP1:.+]] = arith.cmpi slt, %[[EXTRACT]] // CHECK-DAG: %[[SELECT1:.+]] = select %[[CMP1]], %[[EXTRACT]], %[[C9]] -// CHECK-DAG: %[[CMP2:.+]] = cmpi sgt, %[[SELECT1]], %[[C0]] +// CHECK-DAG: %[[CMP2:.+]] = arith.cmpi sgt, %[[SELECT1]], %[[C0]] // CHECK-DAG: %[[SELECT2:.+]] = select %[[CMP2]], %[[SELECT1]], %[[C0]] -// CHECK-DAG: %[[INDEX_CAST:.+]] = index_cast %[[SELECT2]] +// CHECK-DAG: %[[INDEX_CAST:.+]] = arith.index_cast %[[SELECT2]] // CHECK: scf.for %[[IV0:.+]] = // CHECK: %[[OFFSET:.+]] = affine.apply #[[MAP0]](%[[IV0]])[%[[INDEX_CAST]] // CHECK: %[[ARG4V:.+]] = flow.dispatch.tensor.load %[[ARG4]], offsets = [%[[OFFSET]] @@ -714,21 +714,21 @@ #map = affine_map<() -> ()> func @inline_dag_4(%arg0: tensor<4xi32>, %arg1: tensor<i32>) -> tensor<i16> { - %c3_i32 = constant 3 : i32 - %c0_i32 = constant 0 : i32 + %c3_i32 = arith.constant 3 : i32 + %c0_i32 = arith.constant 0 : i32 %0 = tensor.extract %arg1[] : tensor<i32> - %1 = cmpi slt, %0, %c3_i32 : i32 + %1 = arith.cmpi slt, %0, %c3_i32 : i32 %2 = select %1, %0, %c3_i32 : i32 - %3 = cmpi sgt, %2, %c0_i32 : i32 + %3 = arith.cmpi sgt, %2, %c0_i32 : i32 %4 = select %3, %2, %c0_i32 : i32 - %5 = index_cast %4 : i32 to index + %5 = arith.index_cast %4 : i32 to index %6 = tensor.extract_slice %arg0[%5] [1] [1] : tensor<4xi32> to tensor<i32> br ^bb1 ^bb1: // pred: ^bb0 %7 = linalg.init_tensor [] : tensor<i16> %8 = linalg.generic {indexing_maps = [#map, #map], iterator_types = []} ins(%6 : tensor<i32>) outs(%7 : tensor<i16>) { ^bb0(%arg2: i32, %arg3: i16): // no predecessors - %9 = trunci %arg2 : i32 to i16 + %9 = arith.trunci %arg2 : i32 to i16 linalg.yield %9 : i16 } -> tensor<i16> return %8 : tensor<i16> @@ -741,23 +741,23 @@ // CHECK-NEXT: (%[[ARG2:.+]]: !flow.dispatch.tensor<readonly:4xi32> // CHECK-SAME: %[[ARG3:.+]]: !flow.dispatch.tensor<readonly:i32> // CHECK-SAME: %[[ARG4:.+]]: !flow.dispatch.tensor<writeonly:i16> -// CHECK-DAG: %[[C0:.+]] = constant 0 : i32 -// CHECK-DAG: %[[C3:.+]] = constant 3 : i32 +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : i32 +// CHECK-DAG: %[[C3:.+]] = arith.constant 3 : i32 // CHECK: %[[LEAF1:.+]] = flow.dispatch.tensor.load %[[ARG2]], {{.*}} // CHECK: %[[LEAF2:.+]] = flow.dispatch.tensor.load %[[ARG3]], {{.*}} // CHECK: %[[INIT:.+]] = linalg.init_tensor [] : tensor<i16> // CHECK: %[[OP1:.+]] = tensor.extract %[[LEAF2]][] : tensor<i32> -// CHECK: %[[OP2:.+]] = cmpi slt, %[[OP1]], %[[C3]] : i32 +// CHECK: %[[OP2:.+]] = arith.cmpi slt, %[[OP1]], %[[C3]] : i32 // CHECK: %[[OP3:.+]] = select %[[OP2]], %[[OP1]], %[[C3]] : i32 -// CHECK: %[[OP4:.+]] = cmpi sgt, %[[OP3]], %[[C0]] : i32 +// CHECK: %[[OP4:.+]] = arith.cmpi sgt, %[[OP3]], %[[C0]] : i32 // CHECK: %[[OP5:.+]] = select %[[OP4]], %[[OP3]], %[[C0]] : i32 -// CHECK: %[[OP6:.+]] = index_cast %[[OP5]] : i32 to index +// CHECK: %[[OP6:.+]] = arith.index_cast %[[OP5]] : i32 to index // CHECK: %[[OP7:.+]] = tensor.extract_slice %[[LEAF1]][%[[OP6]]] [1] [1] : tensor<4xi32> to tensor<i32> // CHECK: %[[RES:.+]] = linalg.generi // CHECK-SAME: ins(%[[OP7]] : tensor<i32>) // CHECK-SAME: outs(%[[INIT]] : tensor<i16>) { // CHECK: ^bb0(%[[ARG5:.+]]: i32, %{{.+}}: i16): // no predecessors -// CHECK: %[[TRUNC:.+]] = trunci %[[ARG5]] : i32 to i16 +// CHECK: %[[TRUNC:.+]] = arith.trunci %[[ARG5]] : i32 to i16 // CHECK: linalg.yield %[[TRUNC]] : i16 // CHECK: } -> tensor<i16> // CHECK: flow.dispatch.tensor.store %[[RES]], %[[ARG4]] @@ -765,9 +765,9 @@ // ----- func @multi_result(%arg0: tensor<?x?xi32>, %arg1: tensor<?x?xi32>) -> (tensor<?xi32>, tensor<?xi32>) { - %cmin = constant -2147483648 : i32 - %c0_i32 = constant 0 : i32 - %c0 = constant 0 : index + %cmin = arith.constant -2147483648 : i32 + %c0_i32 = arith.constant 0 : i32 + %c0 = arith.constant 0 : index %0 = tensor.dim %arg0, %c0 : tensor<?x?xi32> %1 = linalg.init_tensor [%0] : tensor<?xi32> %2 = linalg.fill(%cmin, %1) : i32, tensor<?xi32> -> tensor<?xi32> @@ -781,10 +781,10 @@ ins(%arg0, %arg1 : tensor<?x?xi32>, tensor<?x?xi32>) outs(%1, %2 : tensor<?xi32>, tensor<?xi32>) { ^bb0(%arg2: i32, %arg3: i32, %arg4: i32, %arg5: i32): // no predecessors - %5 = cmpi sge, %arg2, %arg4 : i32 + %5 = arith.cmpi sge, %arg2, %arg4 : i32 %6 = select %5, %arg2, %arg4 : i32 - %7 = cmpi eq, %arg2, %arg4 : i32 - %8 = cmpi slt, %arg3, %arg5 : i32 + %7 = arith.cmpi eq, %arg2, %arg4 : i32 + %8 = arith.cmpi slt, %arg3, %arg5 : i32 %9 = select %8, %arg3, %arg5 : i32 %10 = select %5, %arg3, %arg5 : i32 %11 = select %7, %9, %10 : i32 @@ -806,8 +806,8 @@ func @multi_result_fallback(%arg0: tensor<?x10xi32>, %arg1: tensor<?x10xi32>) -> (tensor<?x10xi32>, tensor<?x10xi32>) { - %c0 = constant 0 : index - %c1 = constant 1 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index %0 = tensor.dim %arg0, %c0 : tensor<?x10xi32> %1 = linalg.init_tensor [%0, 10] : tensor<?x10xi32> %2:2 = linalg.generic { @@ -836,20 +836,20 @@ // ----- func @dynamic_slice(%arg0: tensor<?x?xi32>, %arg1: tensor<i32>, %arg2: tensor<i32>, %arg3 : index) -> tensor<1x?xi32> { - %c1_i32 = constant 1 : i32 - %c0_i32 = constant 0 : i32 + %c1_i32 = arith.constant 1 : i32 + %c0_i32 = arith.constant 0 : i32 %0 = tensor.extract %arg1[] : tensor<i32> - %1 = cmpi slt, %0, %c1_i32 : i32 + %1 = arith.cmpi slt, %0, %c1_i32 : i32 %2 = select %1, %0, %c1_i32 : i32 - %3 = cmpi sgt, %2, %c0_i32 : i32 + %3 = arith.cmpi sgt, %2, %c0_i32 : i32 %4 = select %3, %2, %c0_i32 : i32 - %5 = index_cast %4 : i32 to index + %5 = arith.index_cast %4 : i32 to index %6 = tensor.extract %arg2[] : tensor<i32> - %7 = cmpi slt, %6, %c0_i32 : i32 + %7 = arith.cmpi slt, %6, %c0_i32 : i32 %8 = select %7, %6, %c0_i32 : i32 - %9 = cmpi sgt, %8, %c0_i32 : i32 + %9 = arith.cmpi sgt, %8, %c0_i32 : i32 %10 = select %9, %8, %c0_i32 : i32 - %11 = index_cast %10 : i32 to index + %11 = arith.index_cast %10 : i32 to index %12 = tensor.extract_slice %arg0[%5, %11] [1, %arg3] [1, 1] : tensor<?x?xi32> to tensor<1x?xi32> return %12 : tensor<1x?xi32> } @@ -858,7 +858,7 @@ // CHECK-SAME: %[[ARG1:[a-zA-Z0-9_]+]]: tensor<i32> // CHECK-SAME: %[[ARG2:[a-zA-Z0-9_]+]]: tensor<i32> // CHECK-SAME: %[[ARG3:.+]]: index -// CHECK: %[[C1:.+]] = constant 1 : index +// CHECK: %[[C1:.+]] = arith.constant 1 : index // CHECK: %[[RESULT:.+]] = flow.dispatch.workgroups // CHECK-SAME: [%[[ARG3]], %[[C1]], %[[C1]]] // CHECK-SAME: (%[[ARG0]], %[[ARG1]], %[[ARG2]], %[[ARG3]]) @@ -879,9 +879,9 @@ // ----- func @dynamic_dot() -> !hal.buffer_view attributes {iree.abi.stub} { - %c0 = constant 0 : index - %c1 = constant 1 : index - %cst = constant 0.000000e+00 : f32 + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + %cst = arith.constant 0.000000e+00 : f32 %0 = util.dynamic_shape_constant dense<[[1.500000e+01, 1.400000e+01, 1.300000e+01], [1.200000e+01, 1.100000e+01, 1.000000e+01], [9.000000e+00, 8.000000e+00, 7.000000e+00], [6.000000e+00, 5.000000e+00, 4.000000e+00], [3.000000e+00, 2.000000e+00, 1.000000e+00]]> : tensor<5x3xf32> -> tensor<?x?xf32> %1 = util.dynamic_shape_constant dense<[[1.500000e+01, 1.400000e+01, 1.300000e+01, 1.200000e+01, 1.100000e+01], [1.000000e+01, 9.000000e+00, 8.000000e+00, 7.000000e+00, 6.000000e+00], [5.000000e+00, 4.000000e+00, 3.000000e+00, 2.000000e+00, 1.000000e+00]]> : tensor<3x5xf32> -> tensor<?x?xf32> %2 = tensor.dim %0, %c0 : tensor<?x?xf32> @@ -914,7 +914,7 @@ ins(%update, %indices : tensor<?x?xf32>, tensor<?x1xi32>) outs(%original : tensor<?x?xf32>) { ^bb0(%arg0: f32, %arg1: f32): - %1 = addf %arg0, %arg1 : f32 + %1 = arith.addf %arg0, %arg1 : f32 linalg_ext.yield %1 : f32 } -> tensor<?x?xf32> return %0 : tensor<?x?xf32> @@ -925,8 +925,8 @@ // CHECK-SAME: %[[ARG0:[a-zA-Z0-9_]+]]: tensor<?x?xf32> // CHECK-SAME: %[[ARG1:[a-zA-Z0-9_]+]]: tensor<?x1xi32> // CHECK-SAME: %[[ARG2:[a-zA-Z0-9_]+]]: tensor<?x?xf32> -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK-DAG: %[[WORKLOADY:.+]] = tensor.dim %[[ARG2]], %[[C0]] // CHECK-DAG: %[[WORKLOADX:.+]] = tensor.dim %[[ARG2]], %[[C1]] // CHECK: %[[RESULT:.+]] = flow.dispatch.workgroups[%[[WORKLOADX]], %[[WORKLOADY]], %[[C1]]] @@ -963,7 +963,7 @@ %0, %1 = linalg_ext.sort dimension(0) outs(%arg0, %arg1 : tensor<?x?x?xi32>, tensor<?x?x?xf32>) { ^bb0(%arg2: i32, %arg3: i32, %arg4 : f32, %arg5 : f32): // no predecessors - %2 = cmpf ogt, %arg4, %arg5 : f32 + %2 = arith.cmpf ogt, %arg4, %arg5 : f32 linalg_ext.yield %2 : i1 } -> tensor<?x?x?xi32>, tensor<?x?x?xf32> return %0, %1 : tensor<?x?x?xi32>, tensor<?x?x?xf32> @@ -973,8 +973,8 @@ // CHECK: func @sort_3d( // CHECK-SAME: %[[ARG0:[a-zA-Z0-9_]+]]: tensor<?x?x?xi32> // CHECK-SAME: %[[ARG1:[a-zA-Z0-9_]+]]: tensor<?x?x?xf32> -// CHECK-DAG: %[[C1:.+]] = constant 1 : index -// CHECK-DAG: %[[C2:.+]] = constant 2 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index +// CHECK-DAG: %[[C2:.+]] = arith.constant 2 : index // CHECK-DAG: %[[WLY:.+]] = tensor.dim %[[ARG0]], %[[C1]] // CHECK-DAG: %[[WLX:.+]] = tensor.dim %[[ARG0]], %[[C2]] // CHECK: %[[RESULT:.+]]:2 = flow.dispatch.workgroups[%[[WLX]], %[[WLY]], %[[C1]]] @@ -1027,7 +1027,7 @@ %0, %1 = linalg_ext.sort dimension(0) outs(%arg0, %arg1 : tensor<?xi32>, tensor<?xf32>) { ^bb0(%arg2: i32, %arg3: i32, %arg4 : f32, %arg5 : f32): // no predecessors - %2 = cmpf ogt, %arg4, %arg5 : f32 + %2 = arith.cmpf ogt, %arg4, %arg5 : f32 linalg_ext.yield %2 : i1 } -> tensor<?xi32>, tensor<?xf32> return %0, %1 : tensor<?xi32>, tensor<?xf32> @@ -1035,7 +1035,7 @@ // CHECK: func @sort_1d( // CHECK-SAME: %[[ARG0:[a-zA-Z0-9_]+]]: tensor<?xi32> // CHECK-SAME: %[[ARG1:[a-zA-Z0-9_]+]]: tensor<?xf32> -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK: %[[RESULT:.+]]:2 = flow.dispatch.workgroups[%[[C1]], %[[C1]], %[[C1]]] // CHECK-SAME: (%[[ARG0]], %[[ARG1]]) // CHECK-NEXT: (%[[ARG2:[a-zA-Z0-9_]+]]: !flow.dispatch.tensor<readwrite:?xi32> @@ -1055,10 +1055,10 @@ func @scatter_static(%arg0 : tensor<4xi32>, %arg1 : tensor<4x1xi32>, %arg2 : tensor<8xi32>) -> tensor<8xi32>{ - %cst = constant dense<[0, 9, 0, 10, 11, 0, 0, 12]> : tensor<8xi32> - %cst_0 = constant dense<[9, 10, 11, 12]> : tensor<4xi32> - %cst_1 = constant dense<[[1], [3], [4], [7]]> : tensor<4x1xi32> - %cst_2 = constant dense<0> : tensor<8xi32> + %cst = arith.constant dense<[0, 9, 0, 10, 11, 0, 0, 12]> : tensor<8xi32> + %cst_0 = arith.constant dense<[9, 10, 11, 12]> : tensor<4xi32> + %cst_1 = arith.constant dense<[[1], [3], [4], [7]]> : tensor<4x1xi32> + %cst_2 = arith.constant dense<0> : tensor<8xi32> %0 = linalg_ext.scatter ins(%arg0, %arg1 : tensor<4xi32>, tensor<4x1xi32>) outs(%arg2 : tensor<8xi32>) { @@ -1086,7 +1086,7 @@ // Check that we are distributing along the last three dimensions for NHWC-output pooling op. func @pooling_nwhc_sum_static(%input: tensor<1x33x33x160xf32>) -> tensor<1x1x1x160xf32> { - %cst = constant 0.0 : f32 + %cst = arith.constant 0.0 : f32 %1 = linalg.init_tensor [1, 1, 1, 160] : tensor<1x1x1x160xf32> %2 = linalg.fill(%cst, %1) : f32, tensor<1x1x1x160xf32> -> tensor<1x1x1x160xf32> %3 = linalg.init_tensor [33, 33] : tensor<33x33xf32> @@ -1106,11 +1106,11 @@ // ----- func @named_op_outs_fusion(%arg0 : tensor<?x?xf32>, %arg1 : tensor<?x?xf32>) -> tensor<?x?xf32> { - %c0 = constant 0 : index - %c1 = constant 1 : index - %cst1 = constant -1.0 : f64 - %cstm1 = constant 1.0 : f64 - %c12345 = constant 12345 : i32 + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + %cst1 = arith.constant -1.0 : f64 + %cstm1 = arith.constant 1.0 : f64 + %c12345 = arith.constant 12345 : i32 %d0 = tensor.dim %arg0, %c0 : tensor<?x?xf32> %d1 = tensor.dim %arg1, %c1 : tensor<?x?xf32> %init = linalg.init_tensor [%d0, %d1] : tensor<?x?xf32> @@ -1130,19 +1130,19 @@ func @dynamic_slice(%arg0 : i32, %arg1 : i32, %arg2 : tensor<?xi32>, %arg3 : tensor<?x?xi32>) -> tensor<?x?xi32>{ - %c0 = constant 0 : index - %c0_i32 = constant 0 : i32 - %c2_i32 = constant 2 : i32 - %5 = cmpi slt, %arg0, %c2_i32 : i32 + %c0 = arith.constant 0 : index + %c0_i32 = arith.constant 0 : i32 + %c2_i32 = arith.constant 2 : i32 + %5 = arith.cmpi slt, %arg0, %c2_i32 : i32 %6 = select %5, %arg0, %c2_i32 : i32 - %7 = cmpi sgt, %6, %c0_i32 : i32 + %7 = arith.cmpi sgt, %6, %c0_i32 : i32 %8 = select %7, %6, %c0_i32 : i32 - %9 = index_cast %8 : i32 to index - %11 = cmpi slt, %arg1, %c0_i32 : i32 + %9 = arith.index_cast %8 : i32 to index + %11 = arith.cmpi slt, %arg1, %c0_i32 : i32 %12 = select %11, %arg1, %c0_i32 : i32 - %13 = cmpi sgt, %12, %c0_i32 : i32 + %13 = arith.cmpi sgt, %12, %c0_i32 : i32 %14 = select %13, %12, %c0_i32 : i32 - %15 = index_cast %14 : i32 to index + %15 = arith.index_cast %14 : i32 to index %d0 = tensor.dim %arg2, %c0 : tensor<?xi32> %17 = tensor.insert_slice %arg2 into %arg3[%9, %15] [1, %d0] [1, 1] : tensor<?xi32> into tensor<?x?xi32> @@ -1153,8 +1153,8 @@ // CHECK-SAME: %[[ARG1:[a-zA-Z0-9]+]]: i32 // CHECK-SAME: %[[ARG2:.+]]: tensor<?xi32> // CHECK-SAME: %[[ARG3:.+]]: tensor<?x?xi32> -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK-DAG: %[[D0:.+]] = tensor.dim %[[ARG2]], %[[C0]] // CHECK-DAG: %[[D1:.+]] = tensor.dim %[[ARG3]], %[[C0]] // CHECK-DAG: %[[D2:.+]] = tensor.dim %[[ARG3]], %[[C1]]
diff --git a/iree/compiler/Dialect/Flow/Transforms/test/dispatch_linalg_on_tensors_elementwise.mlir b/iree/compiler/Dialect/Flow/Transforms/test/dispatch_linalg_on_tensors_elementwise.mlir index 466c7e5..b34ab02 100644 --- a/iree/compiler/Dialect/Flow/Transforms/test/dispatch_linalg_on_tensors_elementwise.mlir +++ b/iree/compiler/Dialect/Flow/Transforms/test/dispatch_linalg_on_tensors_elementwise.mlir
@@ -1,8 +1,8 @@ // RUN: iree-opt -split-input-file -verify-diagnostics -iree-flow-dispatch-linalg-on-tensors-pass -canonicalize -cse %s | IreeFileCheck %s func @tile_generic_op_alone(%A: tensor<?x?xf32>, %B: tensor<?xf32>) -> tensor<?x?xf32> { - %c0 = constant 0 : index - %c1 = constant 1 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index %d0 = tensor.dim %A, %c0 : tensor<?x?xf32> %d1 = tensor.dim %A, %c1 : tensor<?x?xf32> %0 = linalg.init_tensor [%d0, %d1] : tensor<?x?xf32> @@ -14,7 +14,7 @@ ins (%A, %B: tensor<?x?xf32>, tensor<?xf32>) outs (%0 : tensor<?x?xf32>) { ^bb0(%arg0 : f32, %arg1 : f32, %arg2 : f32): - %2 = addf %arg0, %arg1 : f32 + %2 = arith.addf %arg0, %arg1 : f32 linalg.yield %2 : f32 } -> tensor<?x?xf32> return %1 : tensor<?x?xf32> @@ -23,8 +23,8 @@ // CHECK: func @tile_generic_op_alone // CHECK-SAME: %[[ARG0:[a-zA-Z0-9_]+]]: tensor<?x?xf32> // CHECK-SAME: %[[ARG1:[a-zA-Z0-9_]+]]: tensor<?xf32> -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK-DAG: %[[D0:.+]] = tensor.dim %[[ARG0]], %[[C0]] // CHECK-DAG: %[[D1:.+]] = tensor.dim %[[ARG0]], %[[C1]] // CHECK: flow.dispatch.workgroups @@ -61,10 +61,10 @@ func @tile_4d_generic_op_alone (%A: tensor<?x?x?x?xf32>, %B: tensor<?x?x?x?xf32>) -> tensor<?x?x?x?xf32> { - %c0 = constant 0 : index - %c1 = constant 1 : index - %c2 = constant 2 : index - %c3 = constant 3 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + %c2 = arith.constant 2 : index + %c3 = arith.constant 3 : index %d0 = tensor.dim %A, %c0 : tensor<?x?x?x?xf32> %d1 = tensor.dim %A, %c1 : tensor<?x?x?x?xf32> %d2 = tensor.dim %A, %c2 : tensor<?x?x?x?xf32> @@ -78,7 +78,7 @@ ins (%A, %B: tensor<?x?x?x?xf32>, tensor<?x?x?x?xf32>) outs (%0 : tensor<?x?x?x?xf32>) { ^bb0(%arg0 : f32, %arg1 : f32, %arg2 : f32): - %2 = addf %arg0, %arg1 : f32 + %2 = arith.addf %arg0, %arg1 : f32 linalg.yield %2 : f32 } -> tensor<?x?x?x?xf32> return %1 : tensor<?x?x?x?xf32> @@ -88,10 +88,10 @@ // CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0, s1] -> (s0 * s1)> // CHECK: func @tile_4d_generic_op_alone // CHECK-SAME: %[[ARG0:[a-zA-Z0-9_]+]]: tensor<?x?x?x?xf32> -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index -// CHECK-DAG: %[[C2:.+]] = constant 2 : index -// CHECK-DAG: %[[C3:.+]] = constant 3 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index +// CHECK-DAG: %[[C2:.+]] = arith.constant 2 : index +// CHECK-DAG: %[[C3:.+]] = arith.constant 3 : index // CHECK-DAG: %[[D0:.+]] = tensor.dim %[[ARG0]], %[[C0]] // CHECK-DAG: %[[D1:.+]] = tensor.dim %[[ARG0]], %[[C1]] // CHECK-DAG: %[[D2:.+]] = tensor.dim %[[ARG0]], %[[C2]] @@ -104,12 +104,12 @@ // ----- func @tile_parallel_reduction(%arg0: tensor<7x7x1280xf32>) -> tensor<1280xf32> { - %cst = constant 0.000000e+00 : f32 + %cst = arith.constant 0.000000e+00 : f32 %0 = linalg.init_tensor [1280] : tensor<1280xf32> %1 = linalg.fill(%cst, %0) : f32, tensor<1280xf32> -> tensor<1280xf32> %2 = linalg.generic {indexing_maps = [affine_map<(d0, d1, d2) -> (d1, d2, d0)>, affine_map<(d0, d1, d2) -> (d0)>], iterator_types = ["parallel", "reduction", "reduction"]} ins(%arg0 : tensor<7x7x1280xf32>) outs(%1 : tensor<1280xf32>) { ^bb0(%arg1: f32, %arg2: f32): - %3 = addf %arg1, %arg2 : f32 + %3 = arith.addf %arg1, %arg2 : f32 linalg.yield %3 : f32 } -> tensor<1280xf32> return %2 : tensor<1280xf32> @@ -121,8 +121,8 @@ // CHECK: func @tile_parallel_reduction // CHECK-SAME: (%[[INPUT:.+]]: tensor<7x7x1280xf32>) -// CHECK-DAG: %[[C1:.+]] = constant 1 : index -// CHECK-DAG: %[[C1280:.+]] = constant 1280 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index +// CHECK-DAG: %[[C1280:.+]] = arith.constant 1280 : index // CHECK: %[[REDUCE:.+]] = flow.dispatch.workgroups[%[[C1280]], %[[C1]], %[[C1]]](%[[INPUT]]) : (tensor<7x7x1280xf32>) -> tensor<1280xf32> = // CHECK-NEXT: (%[[ARG1:.+]]: !flow.dispatch.tensor<readonly:7x7x1280xf32>, %[[ARG2:.+]]: !flow.dispatch.tensor<writeonly:1280xf32>) { // CHECK: %[[WG_SIZE0:.+]] = flow.dispatch.workgroup.size[0] : index
diff --git a/iree/compiler/Dialect/Flow/Transforms/test/dispatch_linalg_on_tensors_fusion.mlir b/iree/compiler/Dialect/Flow/Transforms/test/dispatch_linalg_on_tensors_fusion.mlir index 63aa5ff..080b76a 100644 --- a/iree/compiler/Dialect/Flow/Transforms/test/dispatch_linalg_on_tensors_fusion.mlir +++ b/iree/compiler/Dialect/Flow/Transforms/test/dispatch_linalg_on_tensors_fusion.mlir
@@ -1,7 +1,7 @@ // RUN: iree-opt -split-input-file -verify-diagnostics -iree-flow-dispatch-linalg-on-tensors-pass -canonicalize -cse %s | IreeFileCheck %s func @fuse_conv2d_elementwise(%input: tensor<1x225x225x16xf32>, %filter: tensor<3x3x16x32xf32>, %offset: tensor<32xf32>) -> tensor<1x112x112x32xf32> { - %cst = constant 0.000000e+00 : f32 + %cst = arith.constant 0.000000e+00 : f32 %0 = linalg.init_tensor [1, 112, 112, 32] : tensor<1x112x112x32xf32> %1 = linalg.fill(%cst, %0) : f32, tensor<1x112x112x32xf32> -> tensor<1x112x112x32xf32> %2 = linalg.conv_2d_nhwc_hwcf @@ -18,7 +18,7 @@ ins(%2, %offset: tensor<1x112x112x32xf32>, tensor<32xf32>) outs(%1 : tensor<1x112x112x32xf32>) { ^bb0(%a: f32, %b: f32, %c: f32): - %sub = subf %a, %b : f32 + %sub = arith.subf %a, %b : f32 linalg.yield %sub : f32 } -> tensor<1x112x112x32xf32> return %3 : tensor<1x112x112x32xf32> @@ -47,7 +47,7 @@ func @dont_fuse_conv2d_with_multiple_uses(%input: tensor<1x225x225x16xf32>, %filter: tensor<3x3x16x32xf32>, %offset: tensor<32xf32>) -> (tensor<1x112x112x32xf32>, tensor<1x112x112x32xf32>) { - %cst = constant 0.000000e+00 : f32 + %cst = arith.constant 0.000000e+00 : f32 %0 = linalg.init_tensor [1, 112, 112, 32] : tensor<1x112x112x32xf32> %1 = linalg.fill(%cst, %0) : f32, tensor<1x112x112x32xf32> -> tensor<1x112x112x32xf32> %2 = linalg.conv_2d_nhwc_hwcf @@ -64,7 +64,7 @@ ins(%2, %offset: tensor<1x112x112x32xf32>, tensor<32xf32>) outs(%1 : tensor<1x112x112x32xf32>) { ^bb0(%a: f32, %b: f32, %c: f32): - %sub = subf %a, %b : f32 + %sub = arith.subf %a, %b : f32 linalg.yield %sub : f32 } -> tensor<1x112x112x32xf32> return %3, %2 : tensor<1x112x112x32xf32>, tensor<1x112x112x32xf32> @@ -83,7 +83,7 @@ // ----- func @dont_fuse_conv2d_with_non_identity_map(%input: tensor<1x225x225x16xf32>, %filter: tensor<3x3x16x32xf32>, %offset: tensor<32xf32>) -> tensor<1x112x112x32xf32> { - %cst = constant 0.000000e+00 : f32 + %cst = arith.constant 0.000000e+00 : f32 %0 = linalg.init_tensor [1, 112, 112, 32] : tensor<1x112x112x32xf32> %1 = linalg.fill(%cst, %0) : f32, tensor<1x112x112x32xf32> -> tensor<1x112x112x32xf32> %2 = linalg.conv_2d_nhwc_hwcf @@ -100,7 +100,7 @@ ins(%2, %offset: tensor<1x112x112x32xf32>, tensor<32xf32>) outs(%1 : tensor<1x112x112x32xf32>) { ^bb0(%a: f32, %b: f32, %c: f32): - %sub = subf %a, %b : f32 + %sub = arith.subf %a, %b : f32 linalg.yield %sub : f32 } -> tensor<1x112x112x32xf32> return %3 : tensor<1x112x112x32xf32>
diff --git a/iree/compiler/Dialect/Flow/Transforms/test/form_streams.mlir b/iree/compiler/Dialect/Flow/Transforms/test/form_streams.mlir index 7e73103..4f5cff0 100644 --- a/iree/compiler/Dialect/Flow/Transforms/test/form_streams.mlir +++ b/iree/compiler/Dialect/Flow/Transforms/test/form_streams.mlir
@@ -2,12 +2,12 @@ // CHECK-LABEL: func @outsideTieShape func @outsideTieShape(%arg0: tensor<?xi32> {iree.reflection = {}}, %arg1: !shapex.ranked_shape<[?]> {iree.reflection = {}}) -> (tensor<?xi32> {iree.reflection = {}}) { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index // CHECK-DAG: %[[DIM:.+]] = shapex.ranked_dim %arg1[0] %dim = shapex.ranked_dim %arg1[0] : !shapex.ranked_shape<[?]> -> index // CHECK-NEXT: %[[RET:.+]] = flow.ex.stream.fragment(%[[DIM]], %arg0) : (index, tensor<?xi32>{%[[DIM]]}) -> tensor<?xi32>{%[[DIM]]} = // CHECK-NEXT: (%[[INNER_DIM:.+]]: index, %[[CAPTURE:.+]]: tensor<?xi32>) -> tensor<?xi32> { - // CHECK-NEXT: %[[WORKLOAD0:.+]] = constant 0 : index + // CHECK-NEXT: %[[WORKLOAD0:.+]] = arith.constant 0 : index // CHECK-NEXT: %[[INNER_RET:.+]] = flow.dispatch @main_ex_dispatch_1::@main_ex_dispatch_1[%[[WORKLOAD0]]](%[[INNER_DIM]], %[[CAPTURE]]) : (index, tensor<?xi32>{%[[INNER_DIM]]}) -> tensor<?xi32>{%[[INNER_DIM]]} // CHECK-NEXT: flow.return %[[INNER_RET]] : tensor<?xi32> // CHECK-NEXT: } @@ -31,18 +31,18 @@ } // CHECK-LABEL: func @outerOps func @outerOps(%arg0: tensor<4xf32>) -> tensor<4xf32> { - // CHECK: %0 = addf %arg0, %arg0 : tensor<4xf32> - %0 = addf %arg0, %arg0 : tensor<4xf32> - %cst = constant 4 : index + // CHECK: %0 = arith.addf %arg0, %arg0 : tensor<4xf32> + %0 = arith.addf %arg0, %arg0 : tensor<4xf32> + %cst = arith.constant 4 : index // CHECK-NEXT: %1 = flow.ex.stream.fragment(%0) : (tensor<4xf32>) -> tensor<4xf32> = // CHECK-NEXT: (%[[INNER_ARG:.+]]: tensor<4xf32>) -> tensor<4xf32> { - // CHECK-NEXT: %[[WORKLOAD:.+]] = constant 4 : index + // CHECK-NEXT: %[[WORKLOAD:.+]] = arith.constant 4 : index // CHECK-NEXT: %3 = flow.dispatch @outerOps_ex_dispatch_0::@outerOps_rgn_dispatch_0[%[[WORKLOAD]]](%[[INNER_ARG]]) : (tensor<4xf32>) -> tensor<4xf32> // CHECK-NEXT: flow.return %3 : tensor<4xf32> // CHECK-NEXT: } %1 = flow.dispatch @outerOps_ex_dispatch_0::@outerOps_rgn_dispatch_0[%cst](%0) : (tensor<4xf32>) -> tensor<4xf32> - // CHECK: %2 = addf %1, %1 : tensor<4xf32> - %2 = addf %1, %1 : tensor<4xf32> + // CHECK: %2 = arith.addf %1, %1 : tensor<4xf32> + %2 = arith.addf %1, %1 : tensor<4xf32> // CHECK-NEXT: return %2 : tensor<4xf32> return %2 : tensor<4xf32> } @@ -51,20 +51,20 @@ // CHECK-LABEL: func @nondependentOuterOps( func @nondependentOuterOps(%arg0: tensor<4xf32>) -> tensor<4xf32> { - %cst = constant 4 : index - // CHECK-NEXT: %[[ADD1:.+]] = addf %arg0, %arg0 : tensor<4xf32> - %add1 = addf %arg0, %arg0 : tensor<4xf32> + %cst = arith.constant 4 : index + // CHECK-NEXT: %[[ADD1:.+]] = arith.addf %arg0, %arg0 : tensor<4xf32> + %add1 = arith.addf %arg0, %arg0 : tensor<4xf32> // CHECK-NEXT: %[[S:.+]] = flow.ex.stream.fragment(%arg0, %[[ADD1]]) : (tensor<4xf32>, tensor<4xf32>) -> tensor<4xf32> = // CHECK-NEXT: (%arg1: tensor<4xf32>, %arg2: tensor<4xf32>) -> tensor<4xf32> { - // CHECK-NEXT: %[[WORKLOAD:.+]] = constant 4 : index + // CHECK-NEXT: %[[WORKLOAD:.+]] = arith.constant 4 : index // CHECK-NEXT: %[[D1:.+]] = flow.dispatch @dispatch_1::@dispatch_1[%[[WORKLOAD]]](%arg1, %arg1) : (tensor<4xf32>, tensor<4xf32>) -> tensor<4xf32> %d1 = flow.dispatch @dispatch_1::@dispatch_1[%cst](%arg0, %arg0) : (tensor<4xf32>, tensor<4xf32>) -> tensor<4xf32> // CHECK-NEXT: %[[D2:.+]] = flow.dispatch @dispatch_2::@dispatch_2[%[[WORKLOAD]]](%[[D1]], %arg2) : (tensor<4xf32>, tensor<4xf32>) -> tensor<4xf32> %d2 = flow.dispatch @dispatch_2::@dispatch_2[%cst](%d1, %add1) : (tensor<4xf32>, tensor<4xf32>) -> tensor<4xf32> // CHECK-NEXT: flow.return %[[D2]] : tensor<4xf32> // CHECK-NEXT: } - // CHECK-NEXT: %[[ADD2:.+]] = addf %[[S]], %arg0 : tensor<4xf32> - %add2 = addf %d2, %arg0 : tensor<4xf32> + // CHECK-NEXT: %[[ADD2:.+]] = arith.addf %[[S]], %arg0 : tensor<4xf32> + %add2 = arith.addf %d2, %arg0 : tensor<4xf32> // CHECK-NEXT: return %[[ADD2]] : tensor<4xf32> return %add2 : tensor<4xf32> } @@ -84,19 +84,19 @@ } // CHECK-LABEL: func @interleavedOuterOps( func @interleavedOuterOps(%arg0: tensor<4xf32>) -> tensor<4xf32> { - %cst = constant 4 : index + %cst = arith.constant 4 : index // CHECK-NEXT: %0 = flow.ex.stream.fragment(%arg0) : (tensor<4xf32>) -> tensor<4xf32> = // CHECK-NEXT: (%arg1: tensor<4xf32>) -> tensor<4xf32> - // CHECK-NEXT: %[[WORKLOAD1:.+]] = constant 4 : index + // CHECK-NEXT: %[[WORKLOAD1:.+]] = arith.constant 4 : index // CHECK-NEXT: %3 = flow.dispatch @interleavedOuterOps_ex_dispatch_0::@interleavedOuterOps_rgn_dispatch_0[%[[WORKLOAD1]]](%arg1) : (tensor<4xf32>) -> tensor<4xf32> // CHECK-NEXT: flow.return %3 : tensor<4xf32> // CHECK-NEXT: } %0 = flow.dispatch @interleavedOuterOps_ex_dispatch_0::@interleavedOuterOps_rgn_dispatch_0[%cst](%arg0) : (tensor<4xf32>) -> tensor<4xf32> - // CHECK-NEXT: %1 = addf %0, %0 : tensor<4xf32> - %1 = addf %0, %0 : tensor<4xf32> + // CHECK-NEXT: %1 = arith.addf %0, %0 : tensor<4xf32> + %1 = arith.addf %0, %0 : tensor<4xf32> // CHECK-NEXT: %2 = flow.ex.stream.fragment(%1) : (tensor<4xf32>) -> tensor<4xf32> = // CHECK-NEXT: (%arg1: tensor<4xf32>) -> tensor<4xf32> { - // CHECK-NEXT: %[[WORKLOAD2:.+]] = constant 4 : index + // CHECK-NEXT: %[[WORKLOAD2:.+]] = arith.constant 4 : index // CHECK-NEXT: %3 = flow.dispatch @interleavedOuterOps_ex_dispatch_0::@interleavedOuterOps_rgn_dispatch_0[%[[WORKLOAD2]]](%arg1) : (tensor<4xf32>) -> tensor<4xf32> // CHECK-NEXT: flow.return %3 : tensor<4xf32> // CHECK-NEXT: } @@ -117,10 +117,10 @@ } // CHECK-LABEL: func @independentOps( func @independentOps(%arg0: tensor<4xf32>) -> (tensor<4xf32>, tensor<4xf32>) { - %cst = constant 4 : index + %cst = arith.constant 4 : index // CHECK-NEXT: %0:2 = flow.ex.stream.fragment(%arg0) : (tensor<4xf32>) -> (tensor<4xf32>, tensor<4xf32>) = // CHECK-NEXT: (%arg1: tensor<4xf32>) -> (tensor<4xf32>, tensor<4xf32>) { - // CHECK-NEXT: %[[WORKLOAD:.+]] = constant 4 : index + // CHECK-NEXT: %[[WORKLOAD:.+]] = arith.constant 4 : index // CHECK-DAG: = flow.dispatch @independentOps_ex_dispatch_0::@independentOps_rgn_dispatch_0[%[[WORKLOAD]]](%arg1) %0 = flow.dispatch @independentOps_ex_dispatch_0::@independentOps_rgn_dispatch_0[%cst](%arg0) : (tensor<4xf32>) -> tensor<4xf32> // CHECK-DAG: = flow.dispatch @independentOps_ex_dispatch_0::@independentOps_rgn_dispatch_1[%[[WORKLOAD]]](%arg1) @@ -168,10 +168,10 @@ } // CHECK-LABEL: func @interleavedDot( func @interleavedDot(%arg0: tensor<4x4xf32>) -> tensor<4x4xf32> { - %cst = constant 16 : index + %cst = arith.constant 16 : index // CHECK-NEXT: %0 = flow.ex.stream.fragment(%arg0) : (tensor<4x4xf32>) -> tensor<4x4xf32> = // CHECK-NEXT: (%arg1: tensor<4x4xf32>) -> tensor<4x4xf32> { - // CHECK-NEXT: %[[WORKLOAD:.+]] = constant 16 : index + // CHECK-NEXT: %[[WORKLOAD:.+]] = arith.constant 16 : index // CHECK-NEXT: %1 = flow.dispatch @interleavedDot_ex_dispatch_0::@interleavedDot_rgn_dispatch_0[%[[WORKLOAD]]](%arg1) : (tensor<4x4xf32>) -> tensor<4x4xf32> // CHECK-NEXT: %2 = flow.dispatch @interleavedDot_ex_dispatch_1::@interleavedDot_rgn_dispatch_1[%[[WORKLOAD]]](%1, %arg1) : (tensor<4x4xf32>, tensor<4x4xf32>) -> tensor<4x4xf32> // CHECK-NEXT: %3 = flow.dispatch @interleavedDot_ex_dispatch_2::@interleavedDot_rgn_dispatch_2[%[[WORKLOAD]]](%2, %arg1) : (tensor<4x4xf32>, tensor<4x4xf32>) -> tensor<4x4xf32> @@ -210,10 +210,10 @@ } // CHECK-LABEL: func @caller( func @caller(%arg0: tensor<4xf32>) -> tensor<4xf32> { - %cst = constant 4 : index + %cst = arith.constant 4 : index // CHECK-NEXT: %0 = flow.ex.stream.fragment(%arg0) : (tensor<4xf32>) -> tensor<4xf32> = // CHECK-NEXT: (%arg1: tensor<4xf32>) -> tensor<4xf32> { - // CHECK-NEXT: %[[WORKLOAD1:.+]] = constant 4 : index + // CHECK-NEXT: %[[WORKLOAD1:.+]] = arith.constant 4 : index // CHECK-NEXT: %3 = flow.dispatch @caller_ex_dispatch_0::@caller_rgn_dispatch_0[%[[WORKLOAD1]]](%arg1) : (tensor<4xf32>) -> tensor<4xf32> // CHECK-NEXT: flow.return %3 : tensor<4xf32> // CHECK-NEXT: } @@ -222,7 +222,7 @@ %1 = call @callee(%0) : (tensor<4xf32>) -> tensor<4xf32> // CHECK-NEXT: %2 = flow.ex.stream.fragment(%arg0, %1) : (tensor<4xf32>, tensor<4xf32>) -> tensor<4xf32> = // CHECK-NEXT: (%arg1: tensor<4xf32>, %arg2: tensor<4xf32>) -> tensor<4xf32> { - // CHECK-NEXT: %[[WORKLOAD2:.+]] = constant 4 : index + // CHECK-NEXT: %[[WORKLOAD2:.+]] = arith.constant 4 : index // CHECK-NEXT: %3 = flow.dispatch @caller_ex_dispatch_1::@caller_rgn_dispatch_1[%[[WORKLOAD2]]](%arg1, %arg2) : (tensor<4xf32>, tensor<4xf32>) -> tensor<4xf32> // CHECK-NEXT: flow.return %3 : tensor<4xf32> // CHECK-NEXT: } @@ -241,10 +241,10 @@ } // CHECK-LABEL: func @callee( func @callee(%arg0: tensor<4xf32>) -> tensor<4xf32> { - %cst = constant 4 : index + %cst = arith.constant 4 : index // CHECK-NEXT: %0 = flow.ex.stream.fragment(%arg0) : (tensor<4xf32>) -> tensor<4xf32> = // CHECK-NEXT: (%arg1: tensor<4xf32>) -> tensor<4xf32> { - // CHECK-NEXT: %[[WORKLOAD:.+]] = constant 4 : index + // CHECK-NEXT: %[[WORKLOAD:.+]] = arith.constant 4 : index // CHECK-NEXT: %1 = flow.dispatch @callee_ex_dispatch_0::@callee_rgn_dispatch_0[%[[WORKLOAD]]](%arg1) : (tensor<4xf32>) -> tensor<4xf32> // CHECK-NEXT: flow.return %1 : tensor<4xf32> // CHECK-NEXT: } @@ -263,8 +263,8 @@ %dim0 = shapex.ranked_dim %arg1[0] : !shapex.ranked_shape<[?,?]> -> index // CHECK-DAG: %[[DIM1:.+]] = shapex.ranked_dim %arg1[1] %dim1 = shapex.ranked_dim %arg1[1] : !shapex.ranked_shape<[?,?]> -> index - // CHECK: %[[SZ:.+]] = muli - %2 = muli %dim0, %dim1 : index + // CHECK: %[[SZ:.+]] = arith.muli + %2 = arith.muli %dim0, %dim1 : index // Verify that the fragment captures the tie_shapes and marshals the indices // in as loose index values (not as ranked_shape types). // CHECK: %[[S:.+]] = flow.ex.stream.fragment(%[[SZ]], %[[A0]], %[[DIM0]], %[[DIM1]]) : (index, tensor<?x?xf32>{%[[DIM0]], %[[DIM1]]}, index, index) -> tensor<?x?xf32>{%[[DIM0]], %[[DIM1]]} = @@ -281,13 +281,13 @@ // CHECK-LABEL: @bad_input_ordering func @bad_input_ordering() -> (tensor<i32>, tensor<f32>) { // CHECK: %[[S:.+]] = flow.ex.stream.fragment - // CHECK: = constant 1 : index + // CHECK: = arith.constant 1 : index // CHECK: %[[D1:.+]] = flow.dispatch @dispatch_1::@dispatch_1 - %workload = constant 1 : index + %workload = arith.constant 1 : index %0 = flow.dispatch @dispatch_1::@dispatch_1[%workload]() : () -> tensor<i32> - // CHECK: %[[C2:.+]] = constant 2 : i32 + // CHECK: %[[C2:.+]] = arith.constant 2 : i32 // CHECK-DAG: %[[D2:.+]] = flow.dispatch @dispatch_2::@dispatch_2 - %c2 = constant 2 : i32 + %c2 = arith.constant 2 : i32 %1 = flow.dispatch @dispatch_2::@dispatch_2[%workload](%c2) : (i32) -> tensor<f32> // CHECK: flow.return return %0, %1 : tensor<i32>, tensor<f32> @@ -297,9 +297,9 @@ // CHECK-LABEL: @interstream_readback func @interstream_readback() -> (tensor<i32>, tensor<f32>, tensor<2xf32>) { - %w = constant 1 : index + %w = arith.constant 1 : index // CHECK: %[[S1:.+]]:2 = flow.ex.stream.fragment - // CHECK: %[[W:.+]] = constant 1 : index + // CHECK: %[[W:.+]] = arith.constant 1 : index // CHECK: %[[D1:.+]] = flow.dispatch @dispatch_1::@dispatch_1 // CHECK-DAG: %[[D2:.+]] = flow.dispatch @dispatch_2::@dispatch_2 // Could be returned in either order @@ -323,17 +323,17 @@ // CHECK-LABEL: @ordering func @ordering(%w : index) -> (tensor<i32>, tensor<f32>, tensor<i32>) { - %c1 = constant 1 : i32 + %c1 = arith.constant 1 : i32 // CHECK: %[[S1:.+]] = flow.ex.stream.fragment - // CHECK: %[[C1:.+]] = constant 1 + // CHECK: %[[C1:.+]] = arith.constant 1 // CHECK-DAG: %[[D1:.+]] = flow.dispatch @dispatch_1::@dispatch_1 // CHECK-NEXT: flow.return %[[D1]] %d1 = flow.dispatch @dispatch_1::@dispatch_1[%w](%c1) : (i32) -> tensor<i32> // CHECK: %[[SE_USER:.+]] = util.do_not_optimize(%[[S1]]) %side_effecting_user = util.do_not_optimize(%d1) : tensor<i32> - %c2 = constant 2 : i32 + %c2 = arith.constant 2 : i32 // CHECK: %[[S2:.+]] = flow.ex.stream.fragment - // CHECK: %[[C2:.+]] = constant 2 + // CHECK: %[[C2:.+]] = arith.constant 2 // CHECK-DAG: %[[D2:.+]] = flow.dispatch @dispatch_2::@dispatch_2 // CHECK-NEXT: flow.return %[[D2]] %d2 = flow.dispatch @dispatch_2::@dispatch_2[%w](%c2) : (i32) -> tensor<f32> @@ -346,7 +346,7 @@ // CHECK-LABEL: @metadata_only func @metadata_only(%t: tensor<?xf32>) -> (tensor<?xf32>, !shapex.ranked_shape<[?]>) { // CHECK-NOT: flow.ex.stream.fragment - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %4 = tensor.dim %t, %c0 : tensor<?xf32> %5 = shapex.make_ranked_shape %4 : (index) -> !shapex.ranked_shape<[?]> %6 = shapex.tie_shape %t, %5 : tensor<?xf32>, !shapex.ranked_shape<[?]>
diff --git a/iree/compiler/Dialect/Flow/Transforms/test/hoist_unstreamable_ops.mlir b/iree/compiler/Dialect/Flow/Transforms/test/hoist_unstreamable_ops.mlir index 1566d74..04767f0 100644 --- a/iree/compiler/Dialect/Flow/Transforms/test/hoist_unstreamable_ops.mlir +++ b/iree/compiler/Dialect/Flow/Transforms/test/hoist_unstreamable_ops.mlir
@@ -2,7 +2,7 @@ // CHECK-LABEL: @constants( func @constants() { - // CHECK-DAG: %[[W:.+]] = constant 1 : index + // CHECK-DAG: %[[W:.+]] = arith.constant 1 : index // CHECK-DAG: constant 2 : index // CHECK-DAG: constant 3 : index // CHECK-DAG: constant 4 : index @@ -14,17 +14,17 @@ // CHECK: flow.dispatch @dispatch3::@dispatch3[%[[W]]]() : () -> tensor<f32> // CHECK: flow.dispatch @dispatch4::@dispatch4[%[[W]]]() : () -> tensor<f32> // CHECK: flow.dispatch @dispatch5::@dispatch5[%[[W]]]() : () -> tensor<f32> - %w = constant 1 : index + %w = arith.constant 1 : index %d0 = flow.dispatch @dispatch0::@dispatch0[%w]() : () -> tensor<f32> - %c2 = constant 2 : index + %c2 = arith.constant 2 : index %d1 = flow.dispatch @dispatch1::@dispatch1[%w]() : () -> tensor<f32> - %c3 = constant 3 : index + %c3 = arith.constant 3 : index %d2 = flow.dispatch @dispatch2::@dispatch2[%w]() : () -> tensor<f32> - %c4 = constant 4 : index + %c4 = arith.constant 4 : index %d3 = flow.dispatch @dispatch3::@dispatch3[%w]() : () -> tensor<f32> - %c5 = constant 5 : index + %c5 = arith.constant 5 : index %d4 = flow.dispatch @dispatch4::@dispatch4[%w]() : () -> tensor<f32> - %c6 = constant 6 : index + %c6 = arith.constant 6 : index %d5 = flow.dispatch @dispatch5::@dispatch5[%w]() : () -> tensor<f32> return } @@ -35,11 +35,11 @@ // CHECK-SAME: %[[INPUT:[a-zA-Z0-9$._-]+]] // CHECK-SAME: %[[SHAPE:[a-zA-Z0-9$._-]+]] func @dynamic_tensor(%input: tensor<?x?xf32>, %shape: !shapex.ranked_shape<[?,?]>) -> (tensor<?x?xf32>, !shapex.ranked_shape<[?,?]>) { - // CHECK-DAG: %[[W:.+]] = constant 1 + // CHECK-DAG: %[[W:.+]] = arith.constant 1 // CHECK-DAG: %[[DIM0:.+]] shapex.ranked_dim %[[SHAPE]][0] // CHECK-DAG: %[[DIM1:.+]] shapex.ranked_dim %[[SHAPE]][1] // CHECK: %[[D:.+]] = flow.dispatch - %w = constant 1 : index + %w = arith.constant 1 : index %dim0 = shapex.ranked_dim %shape[0] : !shapex.ranked_shape<[?,?]> -> index %dim1 = shapex.ranked_dim %shape[1] : !shapex.ranked_shape<[?,?]> -> index %d = flow.dispatch @dispatch::@dispatch[%w](%input, %dim0, %dim1) : (tensor<?x?xf32>{%dim0, %dim1}, index, index) -> tensor<?x?xf32>{%dim0, %dim1} @@ -52,10 +52,10 @@ func @side_effects() { // CHECK-DAG: constant 1 // CHECK-DAG: constant 2 - %c1 = constant 1 : index + %c1 = arith.constant 1 : index // CHECK: util.do_not_optimize() util.do_not_optimize() - %c2 = constant 2 : index + %c2 = arith.constant 2 : index return } @@ -65,10 +65,10 @@ func @dependencies() { // CHECK-DAG: constant 1 // CHECK-DAG: constant 2 - %c1 = constant 1 : index - %c2 = constant 2 : index + %c1 = arith.constant 1 : index + %c2 = arith.constant 2 : index // CHECK: addi - %add0 = addi %c1, %c2 : index + %add0 = arith.addi %c1, %c2 : index return } @@ -76,15 +76,15 @@ // CHECK-LABEL: @dependencies_with_dispatch( func @dependencies_with_dispatch() { - // CHECK-DAG: %[[W:.+]] = constant 1 + // CHECK-DAG: %[[W:.+]] = arith.constant 1 // CHECK-DAG: constant 2 // CHECK-DAG: constant dense<3> - %w = constant 1 : index - %c2 = constant 2 : index - %ct3 = constant dense<3> : tensor<i32> + %w = arith.constant 1 : index + %c2 = arith.constant 2 : index + %ct3 = arith.constant dense<3> : tensor<i32> // CHECK: flow.dispatch %d0 = flow.dispatch @dispatch0::@dispatch0[%w]() : () -> tensor<i32> // CHECK: addi - %add0 = addi %d0, %ct3 : tensor<i32> + %add0 = arith.addi %d0, %ct3 : tensor<i32> return }
diff --git a/iree/compiler/Dialect/Flow/Transforms/test/inject_dispatch_tracing.mlir b/iree/compiler/Dialect/Flow/Transforms/test/inject_dispatch_tracing.mlir index 5993201..3c23a71 100644 --- a/iree/compiler/Dialect/Flow/Transforms/test/inject_dispatch_tracing.mlir +++ b/iree/compiler/Dialect/Flow/Transforms/test/inject_dispatch_tracing.mlir
@@ -3,7 +3,7 @@ // CHECK-LABEL: func @singleDispatch // CHECK-SAME: (%[[ARG0:.+]]: tensor<4xf32>) func @singleDispatch(%arg0: tensor<4xf32>) -> tensor<4xf32> { - %c4 = constant 4 : index + %c4 = arith.constant 4 : index // CHECK: flow.tensor.trace {key = "ex::entry0 inputs"} %[[ARG0]] : tensor<4xf32> // CHECK-NEXT: %[[RET0:.+]] = flow.dispatch @ex::@entry0[%c4](%[[ARG0]]) : (tensor<4xf32>) -> tensor<4xf32> %0 = flow.dispatch @ex::@entry0[%c4](%arg0) : (tensor<4xf32>) -> tensor<4xf32> @@ -17,7 +17,7 @@ // CHECK-LABEL: func @multiDispatch // CHECK-SAME: (%[[ARG0:.+]]: tensor<4xf32>) func @multiDispatch(%arg0: tensor<4xf32>) -> tensor<4xf32> { - %c4 = constant 4 : index + %c4 = arith.constant 4 : index // CHECK: flow.tensor.trace {key = "ex::entry0 inputs"} %[[ARG0]] : tensor<4xf32> // CHECK-NEXT: %[[RET0:.+]] = flow.dispatch @ex::@entry0[%c4](%[[ARG0]]) : (tensor<4xf32>) -> tensor<4xf32>
diff --git a/iree/compiler/Dialect/Flow/Transforms/test/insert_constant_clones.mlir b/iree/compiler/Dialect/Flow/Transforms/test/insert_constant_clones.mlir index 06f8b78..8bc3f24 100644 --- a/iree/compiler/Dialect/Flow/Transforms/test/insert_constant_clones.mlir +++ b/iree/compiler/Dialect/Flow/Transforms/test/insert_constant_clones.mlir
@@ -2,10 +2,10 @@ // CHECK-LABEL: @function_return func @function_return() -> (tensor<8xf32>, i32) { - // CHECK-DAG: %[[SCALAR:.+]] = constant 5 - %1 = constant 5 : i32 - // CHECK-DAG: %[[CST:.+]] = constant dense - %cst = constant dense<1.200000e+00> : tensor<8xf32> + // CHECK-DAG: %[[SCALAR:.+]] = arith.constant 5 + %1 = arith.constant 5 : i32 + // CHECK-DAG: %[[CST:.+]] = arith.constant dense + %cst = arith.constant dense<1.200000e+00> : tensor<8xf32> // CHECK-NEXT: %[[RESHAPE:.+]] = flow.tensor.reshape %[[CST]] %0 = flow.tensor.reshape %cst : tensor<8xf32> -> tensor<8xf32> // CHECK-NEXT: %[[CLONE:.+]] = flow.tensor.clone %[[RESHAPE]] : tensor<8xf32> @@ -17,8 +17,8 @@ // CHECK-LABEL: @branch_argument func @branch_argument() -> tensor<8xf32> { - // CHECK: %[[CST:.+]] = constant dense - %cst = constant dense<1.200000e+00> : tensor<8xf32> + // CHECK: %[[CST:.+]] = arith.constant dense + %cst = arith.constant dense<1.200000e+00> : tensor<8xf32> // CHECK-NEXT: %[[CLONE:.+]] = flow.tensor.clone %[[CST]] : tensor<8xf32> // CHECK-NEXT: br ^[[EXIT:.+]](%[[CLONE]] : tensor<8xf32> br ^exit(%cst : tensor<8xf32>) @@ -34,8 +34,8 @@ // CHECK-LABEL: @branch_argument_reuse func @branch_argument_reuse(%cond : i1) -> tensor<8xf32> { - // CHECK: %[[CST:.+]] = constant dense - %cst = constant dense<1.200000e+00> : tensor<8xf32> + // CHECK: %[[CST:.+]] = arith.constant dense + %cst = arith.constant dense<1.200000e+00> : tensor<8xf32> // CHECK-NEXT: %[[CLONE:.+]] = flow.tensor.clone %[[CST]] : tensor<8xf32> // CHECK-NEXT: cond_br %{{.+}}, ^[[BBT:.+]](%[[CLONE]] : tensor<8xf32>), ^[[BBF:.+]](%[[CLONE]] : tensor<8xf32>) cond_br %cond, ^exit_t(%cst : tensor<8xf32>), ^exit_f(%cst : tensor<8xf32>)
diff --git a/iree/compiler/Dialect/Flow/Transforms/test/interchange_generic_ops.mlir b/iree/compiler/Dialect/Flow/Transforms/test/interchange_generic_ops.mlir index d15b20e..3118fc0 100644 --- a/iree/compiler/Dialect/Flow/Transforms/test/interchange_generic_ops.mlir +++ b/iree/compiler/Dialect/Flow/Transforms/test/interchange_generic_ops.mlir
@@ -15,8 +15,8 @@ ins(%arg0, %arg1 : tensor<?x?x?xf32>, tensor<?x?x?xf32>) outs(%arg2 : tensor<?x?x?xf32>) { ^bb0(%arg3: f32, %arg4: f32, %arg5: f32): // no predecessors - %m = mulf %arg3, %arg4 : f32 - %a = addf %arg5, %m : f32 + %m = arith.mulf %arg3, %arg4 : f32 + %a = arith.addf %arg5, %m : f32 linalg.yield %a : f32 } -> tensor<?x?x?xf32> return %0 : tensor<?x?x?xf32>
diff --git a/iree/compiler/Dialect/Flow/Transforms/test/matmul_to_mmt4d.mlir b/iree/compiler/Dialect/Flow/Transforms/test/matmul_to_mmt4d.mlir index 7aa4f40..df0af24 100644 --- a/iree/compiler/Dialect/Flow/Transforms/test/matmul_to_mmt4d.mlir +++ b/iree/compiler/Dialect/Flow/Transforms/test/matmul_to_mmt4d.mlir
@@ -53,7 +53,7 @@ // ----- func @check_mmt4d_with_init_tensor_and_fill(%arg0: tensor<24x8xf32>, %arg1: tensor<8x32xf32>) -> tensor<24x32xf32> { - %c0 = constant 0.0 : f32 + %c0 = arith.constant 0.0 : f32 %0 = linalg.init_tensor [24, 32] : tensor<24x32xf32> %1 = linalg.fill(%c0, %0) : f32, tensor<24x32xf32> -> tensor<24x32xf32> %2 = linalg.matmul ins(%arg0, %arg1 : tensor<24x8xf32>, tensor<8x32xf32>) outs(%1 : tensor<24x32xf32>) -> tensor<24x32xf32> @@ -63,7 +63,7 @@ // CHECK-DAG:#[[MAP1:.+]] = affine_map<(d0, d1, d2, d3) -> (d0, d1, d2, d3)> // CHECK-DAG:#[[MAP2:.+]] = affine_map<(d0, d1, d2, d3) -> (d1, d3, d0, d2)> // CHECK: @check_mmt4d_with_init_tensor_and_fill(%[[LHS:.+]]: tensor<24x8xf32>, %[[RHS:.+]]: tensor<8x32xf32>) -// CHECK: %[[ZERO:.+]] = constant 0.000000e+00 : f32 +// CHECK: %[[ZERO:.+]] = arith.constant 0.000000e+00 : f32 // CHECK: %[[LHS4D:.+]] = linalg.tensor_expand_shape %[[LHS]] // CHECK-SAME: tensor<24x8xf32> into tensor<3x8x4x2xf32> // CHECK: %[[RHS4D:.+]] = linalg.tensor_expand_shape %[[RHS]]
diff --git a/iree/compiler/Dialect/Flow/Transforms/test/outline_dispatch_regions.mlir b/iree/compiler/Dialect/Flow/Transforms/test/outline_dispatch_regions.mlir index 6e744be..e98c492 100644 --- a/iree/compiler/Dialect/Flow/Transforms/test/outline_dispatch_regions.mlir +++ b/iree/compiler/Dialect/Flow/Transforms/test/outline_dispatch_regions.mlir
@@ -17,10 +17,10 @@ // CHECK-LABEL: func @staticShapeDispatch( // CHECK-SAME: %[[ARG0:.+]]: tensor<8x4xf32>) func @staticShapeDispatch(%arg0 : tensor<8x4xf32>) -> tensor<4x8xf32> { - // CHECK-DAG: %[[X:.+]] = constant 100 - %x = constant 100 : index - // CHECK-DAG: %[[Y:.+]] = constant 50 - %y = constant 50 : index + // CHECK-DAG: %[[X:.+]] = arith.constant 100 + %x = arith.constant 100 : index + // CHECK-DAG: %[[Y:.+]] = arith.constant 50 + %y = arith.constant 50 : index // CHECK: %[[RET:.+]] = flow.dispatch @staticShapeDispatch_dispatch_0::@staticShapeDispatch_dispatch_0[ // CHECK-SAME: %[[X]], %[[Y]] // CHECK-SAME: ](%[[ARG0]]) : (tensor<8x4xf32>) -> tensor<4x8xf32> @@ -53,10 +53,10 @@ // CHECK-LABEL: func @dispatchFnMuli( // CHECK-SAME: %[[ARG0:.+]]: tensor<8x4xf32>) func @dispatchFnMuli(%arg0 : tensor<8x4xf32>) -> tensor<8x4xf32> { - // CHECK-DAG: %[[X:.+]] = constant 100 - %x = constant 100 : index - // CHECK-DAG: %[[Y:.+]] = constant 50 - %y = constant 50 : index + // CHECK-DAG: %[[X:.+]] = arith.constant 100 + %x = arith.constant 100 : index + // CHECK-DAG: %[[Y:.+]] = arith.constant 50 + %y = arith.constant 50 : index // CHECK: %[[RET0:.+]] = flow.dispatch @dispatchFnMuli_dispatch_0::@dispatchFnMuli_dispatch_0[ // CHECK-SAME: %[[X]], %[[Y]] // CHECK-SAME: ](%[[ARG0]]) : (tensor<8x4xf32>) -> tensor<4x8xf32> @@ -93,8 +93,8 @@ // CHECK-LABEL: func @dispatchFn1 func @dispatchFn1(%arg0 : tensor<8x4xf32>) -> tensor<4x8xf32> { - %x = constant 100 : index - %y = constant 50 : index + %x = arith.constant 100 : index + %y = arith.constant 50 : index // CHECK: flow.dispatch @dispatchFn1_dispatch_0::@dispatchFn1_dispatch_0 %0 = flow.dispatch.workgroups[%x, %y](%arg0) : (tensor<8x4xf32>) -> (tensor<4x8xf32>) = ( %arg: !flow.dispatch.tensor<readonly:8x4xf32>, %ret: !flow.dispatch.tensor<writeonly:4x8xf32> @@ -108,8 +108,8 @@ // CHECK-LABEL: func @dispatchFn2 func @dispatchFn2(%arg0 : tensor<8x4xf32>) -> tensor<4x8xf32> { - %x = constant 100 : index - %y = constant 50 : index + %x = arith.constant 100 : index + %y = arith.constant 50 : index // CHECK: flow.dispatch @dispatchFn2_dispatch_0::@dispatchFn2_dispatch_0 %0 = flow.dispatch.workgroups[%x, %y](%arg0) : (tensor<8x4xf32>) -> (tensor<4x8xf32>) = ( %arg: !flow.dispatch.tensor<readonly:8x4xf32>, %ret: !flow.dispatch.tensor<writeonly:4x8xf32> @@ -154,16 +154,16 @@ // CHECK-LABEL: func @dynamicShapeDispatch( // CHECK-SAME: %[[ARG0:.+]]: tensor<7x?x24x?xf32>) func @dynamicShapeDispatch(%arg0 : tensor<7x?x24x?xf32>) -> tensor<?x?x1024xf32> { - %c1 = constant 1 : index - %c3 = constant 3 : index + %c1 = arith.constant 1 : index + %c3 = arith.constant 3 : index // CHECK-DAG: %[[ARG0_DIM1:.+]] = tensor.dim %[[ARG0]], %c1 %dim1 = tensor.dim %arg0, %c1 : tensor<7x?x24x?xf32> // CHECK-DAG: %[[ARG0_DIM3:.+]] = tensor.dim %[[ARG0]], %c3 %dim3 = tensor.dim %arg0, %c3 : tensor<7x?x24x?xf32> - // CHECK-DAG: %[[X:.+]] = constant 1024 - %x = constant 1024 : index - // CHECK-DAG: %[[Y:.+]] = constant 512 - %y = constant 512 : index + // CHECK-DAG: %[[X:.+]] = arith.constant 1024 + %x = arith.constant 1024 : index + // CHECK-DAG: %[[Y:.+]] = arith.constant 512 + %y = arith.constant 512 : index // CHECK-NEXT: %[[ARG0_SHAPE:.+]] = shapex.make_ranked_shape %[[ARG0_DIM1]], %[[ARG0_DIM3]] // CHECK-DAG: %[[IN_ARG0_DIM1:.+]] = shapex.ranked_dim %[[ARG0_SHAPE]][1] // CHECK-DAG: %[[IN_ARG0_DIM3:.+]] = shapex.ranked_dim %[[ARG0_SHAPE]][3]
diff --git a/iree/compiler/Dialect/Flow/Transforms/test/outline_large_constants.mlir b/iree/compiler/Dialect/Flow/Transforms/test/outline_large_constants.mlir index 3d749dc..cdc299f 100644 --- a/iree/compiler/Dialect/Flow/Transforms/test/outline_large_constants.mlir +++ b/iree/compiler/Dialect/Flow/Transforms/test/outline_large_constants.mlir
@@ -2,9 +2,9 @@ // CHECK: util.global private @[[LARGE_VARIABLE:.+]] {noinline} = dense<1.200000e+00> : tensor<512x128xf32> func @fn1() -> (tensor<2xf32>, tensor<512x128xf32>) { - // CHECK-DAG: %[[SMALL_VALUE:.+]] = constant dense<{{.+}}> : tensor<2xf32> - %cst_0 = constant dense<[0.0287729427, 0.0297581609]> : tensor<2xf32> + // CHECK-DAG: %[[SMALL_VALUE:.+]] = arith.constant dense<{{.+}}> : tensor<2xf32> + %cst_0 = arith.constant dense<[0.0287729427, 0.0297581609]> : tensor<2xf32> // CHECK-DAG: %[[LARGE_VALUE:.+]] = util.global.load @[[LARGE_VARIABLE]] : tensor<512x128xf32> - %cst_1 = constant dense<1.2> : tensor<512x128xf32> + %cst_1 = arith.constant dense<1.2> : tensor<512x128xf32> return %cst_0, %cst_1 : tensor<2xf32>, tensor<512x128xf32> }
diff --git a/iree/compiler/Dialect/Flow/Transforms/test/pad_tensor_to_tensor.mlir b/iree/compiler/Dialect/Flow/Transforms/test/pad_tensor_to_tensor.mlir index e788b79..3c6d0a3 100644 --- a/iree/compiler/Dialect/Flow/Transforms/test/pad_tensor_to_tensor.mlir +++ b/iree/compiler/Dialect/Flow/Transforms/test/pad_tensor_to_tensor.mlir
@@ -2,9 +2,9 @@ module { func @pad_tensor(%arg0 : tensor<?x?xf32>, %arg1 : tensor<f32>, %arg2 : index, %arg3 : index) -> tensor<?x?xf32> { - %c0 = constant 0 : index - %c4 = constant 4 : index - %c3 = constant 3 : index + %c0 = arith.constant 0 : index + %c4 = arith.constant 4 : index + %c3 = arith.constant 3 : index %0 = tensor.extract %arg1[] : tensor<f32> %1 = linalg.pad_tensor %arg0 low[%c4, %arg2] high[%arg3, %c3] { ^bb0(%arg4: index, %arg5: index): // no predecessors @@ -20,8 +20,8 @@ // CHECK-SAME: %[[ARG1:[a-zA-Z0-9_]+]]: tensor<f32> // CHECK-SAME: %[[ARG2:[a-zA-Z0-9_]+]]: index // CHECK-SAME: %[[ARG3:[a-zA-Z0-9_]+]]: index -// CHECK-DAG: %[[C0:.+]] = constant 0 -// CHECK-DAG: %[[C1:.+]] = constant 1 +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 // CHECK-DAG: %[[VAL:.+]] = tensor.extract %[[ARG1]] // CHECK-DAG: %[[D0:.+]] = tensor.dim %[[ARG0]], %[[C0]] // CHECK-DAG: %[[D1:.+]] = tensor.dim %[[ARG0]], %[[C1]] @@ -36,10 +36,10 @@ module { func @pad_tensor_static(%arg0: tensor<12x4xf32>, %arg1: tensor<f32>) -> tensor<18x12xf32> { - %c4 = constant 4 : index - %c2 = constant 2 : index - %c5 = constant 5 : index - %c3 = constant 3 : index + %c4 = arith.constant 4 : index + %c2 = arith.constant 2 : index + %c5 = arith.constant 5 : index + %c3 = arith.constant 3 : index %0 = tensor.extract %arg1[] : tensor<f32> %1 = linalg.pad_tensor %arg0 low[%c4, %c5] high[%c2, %c3] { ^bb0(%arg2: index, %arg3: index): // no predecessors
diff --git a/iree/compiler/Dialect/Flow/Transforms/test/promote_i1_to_i8.mlir b/iree/compiler/Dialect/Flow/Transforms/test/promote_i1_to_i8.mlir index b351f8a..384ea13 100644 --- a/iree/compiler/Dialect/Flow/Transforms/test/promote_i1_to_i8.mlir +++ b/iree/compiler/Dialect/Flow/Transforms/test/promote_i1_to_i8.mlir
@@ -5,14 +5,14 @@ // CHECK-LABEL: boolean_const func @boolean_const() -> (tensor<4xi1>) { - // CHECK: [[CONST:%.+]] = constant dense<[1, 1, 0, 1]> : tensor<4xi8> + // CHECK: [[CONST:%.+]] = arith.constant dense<[1, 1, 0, 1]> : tensor<4xi8> // CHECK: [[INIT:%.+]] = linalg.init_tensor [4] : tensor<4xi1> // CHECK: [[GENERIC:%.+]] = linalg.generic {indexing_maps = [#[[$MAP]], #[[$MAP]]], iterator_types = ["parallel"]} ins([[CONST]] : tensor<4xi8>) outs([[INIT]] : tensor<4xi1>) // CHECK: ^bb0(%arg0: i8, %arg1: i1): - // CHECK: [[TRUNC:%.+]] = trunci %arg0 : i8 to i1 + // CHECK: [[TRUNC:%.+]] = arith.trunci %arg0 : i8 to i1 // CHECK: linalg.yield [[TRUNC]] // CHECK: return [[GENERIC]] - %0 = constant dense<[true, true, false, true]> : tensor<4xi1> + %0 = arith.constant dense<[true, true, false, true]> : tensor<4xi1> return %0 : tensor<4xi1> } @@ -22,14 +22,14 @@ // CHECK-LABEL: boolean_const func @boolean_const() -> (tensor<4xi1>, tensor<4xi8>) { - // CHECK: [[CONST:%.+]] = constant dense<[1, 1, 0, 1]> : tensor<4xi8> + // CHECK: [[CONST:%.+]] = arith.constant dense<[1, 1, 0, 1]> : tensor<4xi8> // CHECK: [[INIT:%.+]] = linalg.init_tensor [4] : tensor<4xi1> // CHECK: [[GENERIC:%.+]] = linalg.generic {indexing_maps = [#[[$MAP]], #[[$MAP]]], iterator_types = ["parallel"]} ins([[CONST]] : tensor<4xi8>) outs([[INIT]] : tensor<4xi1>) // CHECK: ^bb0(%arg0: i8, %arg1: i1): - // CHECK: [[TRUNC:%.+]] = trunci %arg0 : i8 to i1 + // CHECK: [[TRUNC:%.+]] = arith.trunci %arg0 : i8 to i1 // CHECK: linalg.yield [[TRUNC]] // CHECK: return [[GENERIC]], [[CONST]] - %0 = constant dense<[true, true, false, true]> : tensor<4xi1> - %1 = constant dense<[1, 1, 0, 1]> : tensor<4xi8> + %0 = arith.constant dense<[true, true, false, true]> : tensor<4xi1> + %1 = arith.constant dense<[1, 1, 0, 1]> : tensor<4xi8> return %0, %1 : tensor<4xi1>, tensor<4xi8> }
diff --git a/iree/compiler/Dialect/Flow/Transforms/test/strip_signedness.mlir b/iree/compiler/Dialect/Flow/Transforms/test/strip_signedness.mlir index cd87396..c48abdd 100644 --- a/iree/compiler/Dialect/Flow/Transforms/test/strip_signedness.mlir +++ b/iree/compiler/Dialect/Flow/Transforms/test/strip_signedness.mlir
@@ -15,7 +15,7 @@ func @strip_signedness_const() -> (tensor<4xi8>) { // CHECK: constant // CHECK-SAME: tensor<4xi8> - %0 = constant dense<[0, 2, 3, 7]> : tensor<4xi8> + %0 = arith.constant dense<[0, 2, 3, 7]> : tensor<4xi8> // CHECK: return // CHECK-SAME: tensor<4xi8> return %0 : tensor<4xi8>
diff --git a/iree/compiler/Dialect/Flow/Transforms/test/transformation.mlir b/iree/compiler/Dialect/Flow/Transforms/test/transformation.mlir index 89846e2..b4496af 100644 --- a/iree/compiler/Dialect/Flow/Transforms/test/transformation.mlir +++ b/iree/compiler/Dialect/Flow/Transforms/test/transformation.mlir
@@ -20,14 +20,14 @@ // CHECK-NEXT: module { // CHECK-NEXT: func @hloElementwiseOps_dispatch_0(%arg0: !flow.dispatch.tensor<readonly:4xf32>, %arg1: !flow.dispatch.tensor<writeonly:4xf32>) { // CHECK: %{{.+}} = linalg.generic -// CHECK: %{{.+}} = addf %{{.+}}, %{{.+}} : f32 -// CHECK-NEXT: %{{.+}} = subf %{{.+}}, %{{.+}} : f32 -// CHECK-NEXT: %{{.+}} = mulf %{{.+}}, %{{.+}} : f32 +// CHECK: %{{.+}} = arith.addf %{{.+}}, %{{.+}} : f32 +// CHECK-NEXT: %{{.+}} = arith.subf %{{.+}}, %{{.+}} : f32 +// CHECK-NEXT: %{{.+}} = arith.mulf %{{.+}}, %{{.+}} : f32 // CHECK: func @hloElementwiseOps(%arg0: tensor<4xf32>) -> tensor<4xf32> { // CHECK-NEXT: %0 = flow.ex.stream.fragment(%arg0) : (tensor<4xf32>) -> tensor<4xf32> = // CHECK-NEXT: (%arg1: tensor<4xf32>) -> tensor<4xf32> { -// CHECK-DAG: %[[C1:.+]] = constant 1 : index -// CHECK-DAG: %[[C4:.+]] = constant 4 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index +// CHECK-DAG: %[[C4:.+]] = arith.constant 4 : index // CHECK-NEXT: %1 = flow.dispatch @hloElementwiseOps_dispatch_0::@hloElementwiseOps_dispatch_0[%[[C4]], %[[C1]], %[[C1]]](%arg1) : (tensor<4xf32>) -> tensor<4xf32> // CHECK-NEXT: flow.return %1 : tensor<4xf32> // CHECK-NEXT: } @@ -48,19 +48,19 @@ // CHECK-NEXT: module { // CHECK-NEXT: func @interleavedDot_dispatch_0 // CHECK: %{{.+}} = linalg.generic -// CHECK: %{{.+}} = addf %{{.+}}, %{{.+}} : f32 +// CHECK: %{{.+}} = arith.addf %{{.+}}, %{{.+}} : f32 // CHECK: flow.executable private @interleavedDot_dispatch_1 { // CHECK-NEXT: flow.dispatch.entry public @interleavedDot_dispatch_1 // CHECK-NEXT: module { // CHECK-NEXT: func @interleavedDot_dispatch_1 // CHECK: %{{.+}} = linalg.matmul // CHECK: %{{.+}} = linalg.generic -// CHECK: %{{.+}} = mulf %{{.+}}, %{{.+}} : f32 +// CHECK: %{{.+}} = arith.mulf %{{.+}}, %{{.+}} : f32 // CHECK: func @interleavedDot(%arg0: tensor<4x4xf32>) -> tensor<4x4xf32> { // CHECK-NEXT: %0 = flow.ex.stream.fragment(%arg0) : (tensor<4x4xf32>) -> tensor<4x4xf32> = // CHECK-NEXT: (%arg1: tensor<4x4xf32>) -> tensor<4x4xf32> { -// CHECK-DAG: %[[C1:.+]] = constant 1 : index -// CHECK-DAG: %[[C4:.+]] = constant 4 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index +// CHECK-DAG: %[[C4:.+]] = arith.constant 4 : index // CHECK-NEXT: %1 = flow.dispatch @interleavedDot_dispatch_0::@interleavedDot_dispatch_0[%[[C4]], %[[C4]], %[[C1]]](%arg1) : (tensor<4x4xf32>) -> tensor<4x4xf32> // CHECK-NEXT: %2 = flow.dispatch @interleavedDot_dispatch_1::@interleavedDot_dispatch_1[%[[C4]], %[[C4]], %[[C1]]](%arg1, %1) : (tensor<4x4xf32>, tensor<4x4xf32>) -> tensor<4x4xf32> // CHECK-NEXT: flow.return %2 : tensor<4x4xf32> @@ -71,7 +71,7 @@ // ----- func @reduction(%arg0 : tensor<4x8xf32>) -> tensor<4xf32> { - %0 = constant dense<0.0> : tensor<f32> + %0 = arith.constant dense<0.0> : tensor<f32> %1 = "mhlo.reduce"(%arg0, %0) ( { ^bb0(%arg1 : tensor<f32>, %arg2 : tensor<f32>): %2 = mhlo.add %arg1, %arg2 : tensor<f32> @@ -85,12 +85,12 @@ // CHECK-NEXT: module { // CHECK-NEXT: func @reduction_dispatch_0 // CHECK: %{{.+}} = linalg.generic -// CHECK: %{{.+}} = addf %{{.+}}, %{{.+}} : f32 +// CHECK: %{{.+}} = arith.addf %{{.+}}, %{{.+}} : f32 // CHECK: func @reduction(%arg0: tensor<4x8xf32>) -> tensor<4xf32> { // CHECK-NEXT: %0 = flow.ex.stream.fragment(%arg0) : (tensor<4x8xf32>) -> tensor<4xf32> = // CHECK-NEXT: (%arg1: tensor<4x8xf32>) -> tensor<4xf32> { -// CHECK-DAG: %[[C1:.+]] = constant 1 : index -// CHECK-DAG: %[[C4:.+]] = constant 4 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index +// CHECK-DAG: %[[C4:.+]] = arith.constant 4 : index // CHECK-NEXT: %1 = flow.dispatch @reduction_dispatch_0::@reduction_dispatch_0[%[[C4]], %[[C1]], %[[C1]]](%arg1) : (tensor<4x8xf32>) -> tensor<4xf32> // CHECK-NEXT: flow.return %1 : tensor<4xf32> // CHECK-NEXT: }
diff --git a/iree/compiler/Dialect/Flow/Transforms/test/verify_input_ir.mlir b/iree/compiler/Dialect/Flow/Transforms/test/verify_input_ir.mlir index 9aa95d3..6e13f41 100644 --- a/iree/compiler/Dialect/Flow/Transforms/test/verify_input_ir.mlir +++ b/iree/compiler/Dialect/Flow/Transforms/test/verify_input_ir.mlir
@@ -40,7 +40,7 @@ ins(%conv, %bias : tensor<1x112x112x16xf32>, tensor<16xf32>) outs(%init : tensor<1x112x112x16xf32>) { ^bb0(%arg0 : f32, %arg1 : f32, %arg2 : f32): - %0 = addf %arg0, %arg1 : f32 + %0 = arith.addf %arg0, %arg1 : f32 linalg.yield %0 : f32 } -> tensor<1x112x112x16xf32> return %result : tensor<1x112x112x16xf32>
diff --git a/iree/compiler/Dialect/HAL/Conversion/FlowToHAL/ConvertGlobalOps.cpp b/iree/compiler/Dialect/HAL/Conversion/FlowToHAL/ConvertGlobalOps.cpp index 419d7c5..748821a 100644 --- a/iree/compiler/Dialect/HAL/Conversion/FlowToHAL/ConvertGlobalOps.cpp +++ b/iree/compiler/Dialect/HAL/Conversion/FlowToHAL/ConvertGlobalOps.cpp
@@ -32,7 +32,7 @@ rewriter.setInsertionPointToStart(initializerOp.addEntryBlock()); // Create const and store ops. - auto constValue = rewriter.create<ConstantOp>(loc, immediateElements); + auto constValue = rewriter.create<arith::ConstantOp>(loc, immediateElements); rewriter.create<IREE::Util::GlobalStoreOp>(loc, constValue.getResult(), globalOp.getName());
diff --git a/iree/compiler/Dialect/HAL/Conversion/FlowToHAL/ConvertStreamOps.cpp b/iree/compiler/Dialect/HAL/Conversion/FlowToHAL/ConvertStreamOps.cpp index bc34353..16676dd 100644 --- a/iree/compiler/Dialect/HAL/Conversion/FlowToHAL/ConvertStreamOps.cpp +++ b/iree/compiler/Dialect/HAL/Conversion/FlowToHAL/ConvertStreamOps.cpp
@@ -264,11 +264,12 @@ Value device() { return device_; } Value allocator() { return allocator_; } - // Returns a ConstantIndexOp of |value|. + // Returns a arith::ConstantIndexOp of |value|. Value lookupOrCreateIndex(int64_t value, OpBuilder &builder) { auto it = indexConstantMap.find(value); if (it != indexConstantMap.end()) return it->second; - auto constantValue = builder.createOrFold<ConstantIndexOp>(loc, value); + auto constantValue = + builder.createOrFold<arith::ConstantIndexOp>(loc, value); indexConstantMap.insert(std::make_pair(value, constantValue)); return constantValue; } @@ -401,8 +402,8 @@ if (it != memoizedElementTypesConstants.end()) return it->second; auto i32Value = IREE::HAL::getElementTypeValue(elementType); assert(i32Value.hasValue() && "unhandled element type for allocation"); - auto constantValue = - builder.createOrFold<ConstantIntOp>(loc, i32Value.getValue(), 32); + auto constantValue = builder.createOrFold<arith::ConstantIntOp>( + loc, i32Value.getValue(), 32); memoizedElementTypesConstants[elementType] = constantValue; return constantValue; } @@ -412,8 +413,8 @@ if (it != memoizedEncodingTypesConstants.end()) return it->second; auto i32Value = IREE::HAL::getEncodingTypeValue(encodingType); assert(i32Value.hasValue() && "unhandled encoding type for allocation"); - auto constantValue = - builder.createOrFold<ConstantIntOp>(loc, i32Value.getValue(), 32); + auto constantValue = builder.createOrFold<arith::ConstantIntOp>( + loc, i32Value.getValue(), 32); memoizedEncodingTypesConstants[encodingType] = constantValue; return constantValue; } @@ -758,7 +759,7 @@ // Need an explicit index cast to i32 since the // CommandBufferPushConstantsOp is intrinsically i32 based. if (inputValue.getType().isa<IndexType>()) { - pushConstantValue = rewriter.create<mlir::IndexCastOp>( + pushConstantValue = rewriter.create<mlir::arith::IndexCastOp>( dispatchOp.getLoc(), rewriter.getIntegerType(32), pushConstantValue); } @@ -797,9 +798,9 @@ // When no workgroup size is specified we just assume [1,1,1]. // This yields a workgroup count that models the extents of the workload. return { - builder.createOrFold<mlir::ConstantIndexOp>(loc, 1), - builder.createOrFold<mlir::ConstantIndexOp>(loc, 1), - builder.createOrFold<mlir::ConstantIndexOp>(loc, 1), + builder.createOrFold<mlir::arith::ConstantIndexOp>(loc, 1), + builder.createOrFold<mlir::arith::ConstantIndexOp>(loc, 1), + builder.createOrFold<mlir::arith::ConstantIndexOp>(loc, 1), }; } @@ -828,40 +829,42 @@ const std::array<Value, 3> &workgroupSize, OpBuilder &builder) { std::array<Value, 3> result; - auto constantOne = builder.createOrFold<mlir::ConstantIndexOp>(loc, 1); + auto constantOne = builder.createOrFold<mlir::arith::ConstantIndexOp>(loc, 1); if (workload.size() <= 3) { // 1-D to 3-D are easy (pad 2 to 0 dimensions) and divide by workgroup size. for (int i = 0; i < 3; ++i) { // Round up: (workload[i] + workgroup_size - 1) / workgroup_size; Value workloadI = i < workload.size() ? workload[i] : constantOne; - workloadI = builder.createOrFold<mlir::SubIOp>( + workloadI = builder.createOrFold<mlir::arith::SubIOp>( loc, - builder.createOrFold<mlir::AddIOp>(loc, workloadI, workgroupSize[i]), + builder.createOrFold<mlir::arith::AddIOp>(loc, workloadI, + workgroupSize[i]), constantOne); - result[i] = builder.createOrFold<UnsignedDivIOp>(loc, workloadI, + result[i] = builder.createOrFold<arith::DivUIOp>(loc, workloadI, workgroupSize[i]); } } else { // TODO(#4140): remapping of N-D to 3-D: this is not how you do this! Value flatWorkload = constantOne; for (auto workloadI : workload) { - flatWorkload = builder.createOrFold<MulIOp>(loc, flatWorkload, workloadI); + flatWorkload = + builder.createOrFold<arith::MulIOp>(loc, flatWorkload, workloadI); } for (int i = 0; i < 3; ++i) { // Round up: (workload[i] + workgroup_size - 1) / workgroup_size; - auto rounded = builder.createOrFold<mlir::SubIOp>( + auto rounded = builder.createOrFold<mlir::arith::SubIOp>( loc, - builder.createOrFold<mlir::AddIOp>(loc, flatWorkload, - workgroupSize[i]), + builder.createOrFold<mlir::arith::AddIOp>(loc, flatWorkload, + workgroupSize[i]), constantOne); - auto workgroupCountI = builder.createOrFold<mlir::UnsignedDivIOp>( + auto workgroupCountI = builder.createOrFold<mlir::arith::DivUIOp>( loc, rounded, workgroupSize[i]); result[i] = workgroupCountI; // Multiply back out and subtract from invocations. - flatWorkload = builder.createOrFold<SubIOp>( + flatWorkload = builder.createOrFold<arith::SubIOp>( loc, flatWorkload, - builder.createOrFold<MulIOp>(loc, workgroupCountI, rounded)); + builder.createOrFold<arith::MulIOp>(loc, workgroupCountI, rounded)); } } @@ -969,32 +972,32 @@ OpBuilder &builder) { // Bitcast to an integer, then use integer math for the rest of the pattern. auto baseBitWidth = baseValue.getType().getIntOrFloatBitWidth(); - baseValue = builder.createOrFold<BitcastOp>( + baseValue = builder.createOrFold<arith::BitcastOp>( loc, builder.getIntegerType(baseBitWidth), baseValue); switch (baseBitWidth) { case 8: { // (v << 24) | (v << 16) | (v << 8) | v - auto b0 = builder.createOrFold<ZeroExtendIOp>(loc, baseValue, - builder.getIntegerType(32)); - auto c8 = builder.create<ConstantIntOp>(loc, 8, 32); - auto b1 = builder.createOrFold<ShiftLeftOp>(loc, b0, c8); - auto c16 = builder.create<ConstantIntOp>(loc, 16, 32); - auto b2 = builder.createOrFold<ShiftLeftOp>(loc, b0, c16); - auto c24 = builder.create<ConstantIntOp>(loc, 24, 32); - auto b3 = builder.createOrFold<ShiftLeftOp>(loc, b0, c24); - return builder.createOrFold<OrOp>( + auto b0 = builder.createOrFold<arith::ExtUIOp>( + loc, baseValue, builder.getIntegerType(32)); + auto c8 = builder.create<arith::ConstantIntOp>(loc, 8, 32); + auto b1 = builder.createOrFold<arith::ShLIOp>(loc, b0, c8); + auto c16 = builder.create<arith::ConstantIntOp>(loc, 16, 32); + auto b2 = builder.createOrFold<arith::ShLIOp>(loc, b0, c16); + auto c24 = builder.create<arith::ConstantIntOp>(loc, 24, 32); + auto b3 = builder.createOrFold<arith::ShLIOp>(loc, b0, c24); + return builder.createOrFold<arith::OrIOp>( loc, b0, - builder.createOrFold<OrOp>(loc, b1, - builder.createOrFold<OrOp>(loc, b2, b3))); + builder.createOrFold<arith::OrIOp>( + loc, b1, builder.createOrFold<arith::OrIOp>(loc, b2, b3))); } case 16: { // (v << 16) | v - auto c16 = builder.create<ConstantIntOp>(loc, 16, 32); - auto b0 = builder.createOrFold<ZeroExtendIOp>(loc, baseValue, - builder.getIntegerType(32)); - auto b1 = builder.createOrFold<ShiftLeftOp>(loc, b0, c16); - return builder.createOrFold<OrOp>(loc, b0, b1); + auto c16 = builder.create<arith::ConstantIntOp>(loc, 16, 32); + auto b0 = builder.createOrFold<arith::ExtUIOp>( + loc, baseValue, builder.getIntegerType(32)); + auto b1 = builder.createOrFold<arith::ShLIOp>(loc, b0, c16); + return builder.createOrFold<arith::OrIOp>(loc, b0, b1); } case 32: return baseValue; @@ -1133,7 +1136,7 @@ static void hoistConstants(Block &streamBlock, ConversionPatternRewriter &rewriter) { for (auto &op : streamBlock) { - if (isa<ConstantOp>(op)) { + if (isa<arith::ConstantOp>(op)) { auto newOp = rewriter.clone(op); op.replaceAllUsesWith(newOp); } @@ -1172,7 +1175,7 @@ } } else if (auto returnOp = dyn_cast<IREE::Flow::ReturnOp>(op)) { // No-op; handled by the buffer allocation. - } else if (isa<ConstantOp>(op)) { + } else if (isa<arith::ConstantOp>(op)) { // Note that even though constants were hoisted early, they can be // materialized as part of various conversions so do it again to get // any new ones.
diff --git a/iree/compiler/Dialect/HAL/Conversion/FlowToHAL/test/stream_ops.mlir b/iree/compiler/Dialect/HAL/Conversion/FlowToHAL/test/stream_ops.mlir index 806e0d7..b8acb8e 100644 --- a/iree/compiler/Dialect/HAL/Conversion/FlowToHAL/test/stream_ops.mlir +++ b/iree/compiler/Dialect/HAL/Conversion/FlowToHAL/test/stream_ops.mlir
@@ -19,9 +19,9 @@ // CHECK-LABEL: func @multipleDispatches // CHECK-SAME: %[[INPUT_BUF:.+]]: !hal.buffer func @multipleDispatches(%input: tensor<128xf32>) -> tensor<128xf32> { - // CHECK-DAG: %[[C0:.+]] = constant 0 - // CHECK-DAG: %[[C128:.+]] = constant 128 - %cst = constant 128 : index + // CHECK-DAG: %[[C0:.+]] = arith.constant 0 + // CHECK-DAG: %[[C128:.+]] = arith.constant 128 + %cst = arith.constant 128 : index // CHECK: %[[RET_BUF:.+]] = hal.allocator.allocate // CHECK-SAME: type("HostVisible|DeviceVisible|DeviceLocal") // CHECK-SAME: usage("Transfer|Mapping|Dispatch") @@ -184,7 +184,7 @@ // CHECK-NEXT: hal.command_buffer.begin<%[[CMD]] %0 = flow.ex.stream.fragment(%arg0) : (tensor<4x4x2xf32>) -> (tensor<4x4x1x2xf32>) = (%source: tensor<4x4x2xf32>) -> (tensor<4x4x1x2xf32>) { - %c1 = constant 1 : index + %c1 = arith.constant 1 : index %r = flow.tensor.reshape %source : tensor<4x4x2xf32> -> tensor<4x4x1x2xf32> // CHECK: hal.command_buffer.push_descriptor_set<%[[CMD]] // CHECK: %c0 = (%[[SRC_BUF]] : !hal.buffer)[%c0, %c128], @@ -218,7 +218,7 @@ %1 = hal.tensor.cast %arg0 : !hal.buffer_view -> tensor<?xf32>{%0} %2 = flow.ex.stream.fragment(%1) : (tensor<?xf32>{%0}) -> tensor<4xf32> = (%arg1: tensor<?xf32>) -> tensor<4xf32> { - %c4 = constant 4 : index + %c4 = arith.constant 4 : index %4 = flow.tensor.reshape %arg1 : tensor<?xf32>{%c4} -> tensor<4xf32> flow.return %4 : tensor<4xf32> } @@ -235,11 +235,11 @@ // CHECK-LABEL: @tensorSlice // CHECK-SAME: (%[[SBUF:.+]]:{{.+}}) func @tensorSlice(%arg0 : tensor<5x24x48xf32>) -> tensor<3x24x48xf32> { - %c0 = constant 0 : index - %c2 = constant 2 : index - %c3 = constant 3 : index - %c24 = constant 24 : index - %c48 = constant 48 : index + %c0 = arith.constant 0 : index + %c2 = arith.constant 2 : index + %c3 = arith.constant 3 : index + %c24 = arith.constant 24 : index + %c48 = arith.constant 48 : index // CHECK: %[[RET_BUF:.+]] = hal.allocator.allocate // CHECK: %[[CMD:.+]] = hal.command_buffer.create // CHECK-NEXT: hal.command_buffer.begin<%[[CMD]] @@ -267,8 +267,8 @@ // CHECK-LABEL: @tensorUpdate // CHECK-SAME: (%[[UBUF:.+]]:{{.+}}, %[[TBUF:.+]]:{{.+}}) func @tensorUpdate(%arg0 : tensor<1x1x10xf32>, %arg1 : tensor<5x1x10xf32>) -> tensor<5x1x10xf32> { - %c4 = constant 4 : index - %c1 = constant 1 : index + %c4 = arith.constant 4 : index + %c1 = arith.constant 1 : index // CHECK: %[[RET_BUF:.+]] = hal.allocator.allocate // CHECK: %[[CMD:.+]] = hal.command_buffer.create // CHECK-NEXT: hal.command_buffer.begin<%[[CMD]] @@ -315,12 +315,12 @@ // CHECK-LABEL: func @dispatchWithShapeTies // CHECK-SAME: (%[[T:.+]]:{{.+}}, %[[BS:.+]]:{{.+}}) func @dispatchWithShapeTies(%arg0: tensor<?x128xf32>, %bs : index) -> tensor<?x128xf32> { - %cst = constant 128 : index + %cst = arith.constant 128 : index // Verify that size computation derives from the passed dynamic index. - // CHECK-DAG: %[[BS4:.+]] = muli %[[BS]], %c4 : index - // CHECK-DAG: = muli %[[BS4]], %c128 : index + // CHECK-DAG: %[[BS4:.+]] = arith.muli %[[BS]], %c4 : index + // CHECK-DAG: = arith.muli %[[BS4]], %c128 : index // Verify that an i32 is pushed. - // CHECK: %[[CAST_BS:.+]] = index_cast %[[BS]] : index to i32 + // CHECK: %[[CAST_BS:.+]] = arith.index_cast %[[BS]] : index to i32 // CHECK: hal.command_buffer.push_constants // CHECK-SAME: layout({{.+}} : !hal.executable_layout) // CHECK-SAME: offset(0) @@ -378,8 +378,8 @@ // CHECK-LABEL: func @staticTiledDispatch // CHECK-SAME: %[[INPUT:.+]]: !hal.buffer func @staticTiledDispatch(%input: tensor<7x4x24xf32>) -> tensor<4x7x1024xf32> { - %c1024 = constant 1024 : index - %c512 = constant 512 : index + %c1024 = arith.constant 1024 : index + %c512 = arith.constant 512 : index // CHECK: %[[CMD:.+]] = hal.command_buffer.create // CHECK-NEXT: hal.command_buffer.begin<%[[CMD]] %1 = flow.ex.stream.fragment(%input, %c1024, %c512) : (tensor<7x4x24xf32>, index, index) -> tensor<4x7x1024xf32> = @@ -426,8 +426,8 @@ // CHECK-LABEL: func @dynamicTiledDispatch // CHECK-SAME: %[[INPUT:.+]]: !hal.buffer func @dynamicTiledDispatch(%arg0: tensor<7x?x24x?xf32>, %arg1: index, %arg2: index) -> tensor<?x?x1024xf32> { - %c1024 = constant 1024 : index - %c512 = constant 512 : index + %c1024 = arith.constant 1024 : index + %c512 = arith.constant 512 : index // CHECK: %[[CMD:.+]] = hal.command_buffer.create // CHECK-NEXT: hal.command_buffer.begin<%[[CMD]] %2 = flow.ex.stream.fragment(%arg0, %arg1, %arg2, %c1024, %c512) : (tensor<7x?x24x?xf32>{%arg1, %arg2}, index, index, index, index) -> tensor<?x?x1024xf32>{%arg2, %arg1} = @@ -505,9 +505,9 @@ // CHECK-SAME: usage("Transfer|Mapping|Dispatch") %0 = flow.ex.stream.fragment(%fill, %input) : (tensor<i32>, tensor<2x3xi32>) -> tensor<3x9xi32> = (%arg0: tensor<i32>, %arg1: tensor<2x3xi32>) -> tensor<3x9xi32> { - %c9 = constant 9 : index - %c3 = constant 3 : index - %c1 = constant 1 : index + %c9 = arith.constant 9 : index + %c3 = arith.constant 3 : index + %c1 = arith.constant 1 : index // CHECK: %[[LAYOUT0:.+]] = hal.executable_layout.lookup // CHECK-SAME: layouts([ @@ -563,10 +563,10 @@ func @cloneFromLargeBufferToSmallBuffer(%input: tensor<2xi32>) -> tensor<7xi32> { %1 = flow.ex.stream.fragment(%input) : (tensor<2xi32>) -> tensor<7xi32> = (%arg0: tensor<2xi32>) -> tensor<7xi32> { - %c3 = constant 3 : index + %c3 = arith.constant 3 : index %const_span = hal.constant.subspan @_const_pool_splats[#util.byte_range<0, 32>] : tensor<7xi32> - // CHECK-DAG: %[[C0:.+]] = constant 0 : index - // CHECK-DAG: %[[C28:.+]] = constant 28 : index + // CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index + // CHECK-DAG: %[[C28:.+]] = arith.constant 28 : index // CHECK: %[[DSTBUF:.+]] = hal.allocator.allocate<%{{.+}} : !hal.allocator> type("HostVisible|DeviceVisible|DeviceLocal") usage("Transfer|Mapping|Dispatch") : !hal.buffer{%[[C28]]} // CHECK: %[[CSTBUF:.+]] = util.global.load @_const_pool_splats : !hal.buffer // CHECK: hal.command_buffer.copy_buffer<%cmd : !hal.command_buffer> source(%[[CSTBUF]] : !hal.buffer)[%[[C0]]] target(%[[DSTBUF]] : !hal.buffer)[%[[C0]]] length(%[[C28]]) @@ -605,17 +605,17 @@ // CHECK-LABEL: func @tensorSplatDynamic // CHECK-SAME: (%[[VALUE:.+]]: i8, %[[DIM:.+]]: index) func @tensorSplatDynamic(%value: i8, %dim: index) -> tensor<?x128xi8> { - // CHECK: %[[SIZE:.+]] = muli %[[DIM]], %c128 : index + // CHECK: %[[SIZE:.+]] = arith.muli %[[DIM]], %c128 : index // CHECK: %[[BUFFER:.+]] = hal.allocator.allocate<%allocator : !hal.allocator> type("HostVisible|DeviceVisible|DeviceLocal") usage("Transfer|Mapping|Dispatch") : !hal.buffer{%[[SIZE]]} %0 = flow.ex.stream.fragment(%value, %dim) : (i8, index) -> tensor<?x128xi8>{%dim} = (%arg0: i8, %arg1: index) -> tensor<?x128xi8> { - // CHECK-DAG: %[[B0:.+]] = zexti %[[VALUE]] : i8 to i32 - // CHECK-DAG: %[[B1:.+]] = shift_left %[[B0]], %c8 - // CHECK-DAG: %[[B2:.+]] = shift_left %[[B0]], %c16 - // CHECK-DAG: %[[B3:.+]] = shift_left %[[B0]], %c24 - // CHECK-DAG: %[[ORA:.+]] = or %[[B2]], %[[B3]] - // CHECK-DAG: %[[ORB:.+]] = or %[[B1]], %[[ORA]] - // CHECK-DAG: %[[PATTERN:.+]] = or %[[B0]], %[[ORB]] + // CHECK-DAG: %[[B0:.+]] = arith.extui %[[VALUE]] : i8 to i32 + // CHECK-DAG: %[[B1:.+]] = arith.shli %[[B0]], %c8 + // CHECK-DAG: %[[B2:.+]] = arith.shli %[[B0]], %c16 + // CHECK-DAG: %[[B3:.+]] = arith.shli %[[B0]], %c24 + // CHECK-DAG: %[[ORA:.+]] = arith.ori %[[B2]], %[[B3]] + // CHECK-DAG: %[[ORB:.+]] = arith.ori %[[B1]], %[[ORA]] + // CHECK-DAG: %[[PATTERN:.+]] = arith.ori %[[B0]], %[[ORB]] // CHECK-NEXT: hal.command_buffer.fill_buffer<%cmd : !hal.command_buffer> target(%[[BUFFER]] : !hal.buffer)[%c0, %[[SIZE]]] pattern(%[[PATTERN]] : i32) %1 = flow.tensor.splat %arg0 : tensor<?x128xi8>{%arg1} flow.return %1 : tensor<?x128xi8> @@ -635,7 +635,7 @@ // CHECK: %[[BUFFER:.+]] = hal.allocator.allocate<%allocator : !hal.allocator> type("HostVisible|DeviceVisible|DeviceLocal") usage("Transfer|Mapping|Dispatch") : !hal.buffer{%c1024} %0 = flow.ex.stream.fragment(%value) : (f32) -> tensor<2x128xf32> = (%arg0: f32) -> tensor<2x128xf32> { - // CHECK-DAG: %[[PATTERN:.+]] = bitcast %[[VALUE]] : f32 to i32 + // CHECK-DAG: %[[PATTERN:.+]] = arith.bitcast %[[VALUE]] : f32 to i32 // CHECK: hal.command_buffer.fill_buffer<%cmd : !hal.command_buffer> target(%[[BUFFER]] : !hal.buffer)[%c0, %c1024] pattern(%[[PATTERN]] : i32) %1 = flow.tensor.splat %arg0 : tensor<2x128xf32> flow.return %1 : tensor<2x128xf32> @@ -655,10 +655,10 @@ // CHECK: %[[BUFFER:.+]] = hal.allocator.allocate<%allocator : !hal.allocator> type("HostVisible|DeviceVisible|DeviceLocal") usage("Transfer|Mapping|Dispatch") : !hal.buffer{%c512} %0 = flow.ex.stream.fragment(%value) : (f16) -> tensor<2x128xf16> = (%arg0: f16) -> tensor<2x128xf16> { - // CHECK-DAG: %[[BITCAST:.+]] = bitcast %[[VALUE]] : f16 to i16 - // CHECK-DAG: %[[B0:.+]] = zexti %[[BITCAST]] : i16 to i32 - // CHECK-DAG: %[[B1:.+]] = shift_left %[[B0]], %c16 - // CHECK-DAG: %[[PATTERN:.+]] = or %[[B0]], %[[B1]] + // CHECK-DAG: %[[BITCAST:.+]] = arith.bitcast %[[VALUE]] : f16 to i16 + // CHECK-DAG: %[[B0:.+]] = arith.extui %[[BITCAST]] : i16 to i32 + // CHECK-DAG: %[[B1:.+]] = arith.shli %[[B0]], %c16 + // CHECK-DAG: %[[PATTERN:.+]] = arith.ori %[[B0]], %[[B1]] // CHECK: hal.command_buffer.fill_buffer<%cmd : !hal.command_buffer> target(%[[BUFFER]] : !hal.buffer)[%c0, %c512] pattern(%[[PATTERN]] : i32) %1 = flow.tensor.splat %arg0 : tensor<2x128xf16> flow.return %1 : tensor<2x128xf16>
diff --git a/iree/compiler/Dialect/HAL/Conversion/FlowToHAL/test/tensor_ops.mlir b/iree/compiler/Dialect/HAL/Conversion/FlowToHAL/test/tensor_ops.mlir index ced0c86..9dc0e47 100644 --- a/iree/compiler/Dialect/HAL/Conversion/FlowToHAL/test/tensor_ops.mlir +++ b/iree/compiler/Dialect/HAL/Conversion/FlowToHAL/test/tensor_ops.mlir
@@ -3,12 +3,12 @@ // CHECK-LABEL: @tensorLoad // CHECK-SAME: %[[BUFFER:.+]]: !hal.buffer func @tensorLoad(%tensor : tensor<2x3xi32>) { - // CHECK-DAG: %[[C0:.+]] = constant 0 : index - // CHECK-DAG: %[[C1:.+]] = constant 1 : index - // CHECK-DAG: %[[C2:.+]] = constant 2 : index - // CHECK-DAG: %[[C3:.+]] = constant 3 : index - %i0 = constant 0 : index - %i1 = constant 1 : index + // CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index + // CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index + // CHECK-DAG: %[[C2:.+]] = arith.constant 2 : index + // CHECK-DAG: %[[C3:.+]] = arith.constant 3 : index + %i0 = arith.constant 0 : index + %i1 = arith.constant 1 : index // CHECK: %[[OFF:.+]] = hal.allocator.compute_offset<%allocator : !hal.allocator> // CHECK-SAME: indices([%[[C0]], %[[C1]]]) // CHECK-SAME: shape([%[[C2]], %[[C3]]]) @@ -37,14 +37,14 @@ // CHECK-LABEL: @tensorStore // CHECK-SAME: %[[BUFFER:.+]]: !hal.buffer func @tensorStore(%tensor : tensor<2x3xi32>) { - // CHECK-DAG: %[[C0:.+]] = constant 0 : index - // CHECK-DAG: %[[C1:.+]] = constant 1 : index - // CHECK-DAG: %[[C9:.+]] = constant 9 : i32 - // CHECK-DAG: %[[C2:.+]] = constant 2 : index - // CHECK-DAG: %[[C3:.+]] = constant 3 : index - %i0 = constant 0 : index - %i1 = constant 1 : index - %c9 = constant 9 : i32 + // CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index + // CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index + // CHECK-DAG: %[[C9:.+]] = arith.constant 9 : i32 + // CHECK-DAG: %[[C2:.+]] = arith.constant 2 : index + // CHECK-DAG: %[[C3:.+]] = arith.constant 3 : index + %i0 = arith.constant 0 : index + %i1 = arith.constant 1 : index + %c9 = arith.constant 9 : i32 // CHECK: %[[OFF:.+]] = hal.allocator.compute_offset<%allocator : !hal.allocator> // CHECK-SAME: indices([%[[C0]], %[[C1]]]) // CHECK-SAME: shape([%[[C2]], %[[C3]]]) @@ -59,8 +59,8 @@ // CHECK-LABEL: @tensorStore1 // CHECK-SAME: %[[BUFFER:.+]]: !hal.buffer func @tensorStore1(%tensor : tensor<i1>) { - // CHECK-DAG: %[[C1:.+]] = constant true - %c1 = constant true + // CHECK-DAG: %[[C1:.+]] = arith.constant true + %c1 = arith.constant true // CHECK: %[[OFF:.+]] = hal.allocator.compute_offset<%allocator : !hal.allocator> // CHECK-SAME: indices([]) // CHECK-SAME: shape([])
diff --git a/iree/compiler/Dialect/HAL/Conversion/HALToHAL/ConvertConstantOps.cpp b/iree/compiler/Dialect/HAL/Conversion/HALToHAL/ConvertConstantOps.cpp index a4e6688..aa90e2c 100644 --- a/iree/compiler/Dialect/HAL/Conversion/HALToHAL/ConvertConstantOps.cpp +++ b/iree/compiler/Dialect/HAL/Conversion/HALToHAL/ConvertConstantOps.cpp
@@ -26,9 +26,9 @@ auto bufferValue = rewriter.createOrFold<IREE::Util::GlobalLoadOp>( op.getLoc(), IREE::HAL::BufferType::get(rewriter.getContext()), op.runtime_buffer().getLeafReference().getValue()); - auto offsetValue = rewriter.createOrFold<mlir::ConstantIndexOp>( + auto offsetValue = rewriter.createOrFold<mlir::arith::ConstantIndexOp>( op.getLoc(), op.runtime_range().getOffset()); - auto lengthValue = rewriter.createOrFold<mlir::ConstantIndexOp>( + auto lengthValue = rewriter.createOrFold<mlir::arith::ConstantIndexOp>( op.getLoc(), op.runtime_range().getLength()); rewriter.replaceOpWithNewOp<IREE::HAL::BufferSubspanOp>( op, bufferValue.getType(), bufferValue, offsetValue, lengthValue);
diff --git a/iree/compiler/Dialect/HAL/Conversion/HALToHAL/test/constant_ops.mlir b/iree/compiler/Dialect/HAL/Conversion/HALToHAL/test/constant_ops.mlir index b66f3f4..2a8761d 100644 --- a/iree/compiler/Dialect/HAL/Conversion/HALToHAL/test/constant_ops.mlir +++ b/iree/compiler/Dialect/HAL/Conversion/HALToHAL/test/constant_ops.mlir
@@ -3,8 +3,8 @@ // CHECK-LABEL: func @constant_subspan func @constant_subspan() { // CHECK-DAG: %[[BUFFER:.+]] = util.global.load @pool_buffer : !hal.buffer - // CHECK-DAG: %[[OFFSET:.+]] = constant 123 : index - // CHECK-DAG: %[[LENGTH:.+]] = constant 16 : index + // CHECK-DAG: %[[OFFSET:.+]] = arith.constant 123 : index + // CHECK-DAG: %[[LENGTH:.+]] = arith.constant 16 : index // CHECK-NEXT: = hal.buffer.subspan<%[[BUFFER]] : !hal.buffer>[%[[OFFSET]], %[[LENGTH]]] : !hal.buffer %cst0 = hal.constant.subspan @pool_buffer[#util.byte_range<123, 16>] : tensor<4xf32> return
diff --git a/iree/compiler/Dialect/HAL/Conversion/HALToVM/ConvertBufferOps.cpp b/iree/compiler/Dialect/HAL/Conversion/HALToVM/ConvertBufferOps.cpp index e351cd4..e1cfaad 100644 --- a/iree/compiler/Dialect/HAL/Conversion/HALToVM/ConvertBufferOps.cpp +++ b/iree/compiler/Dialect/HAL/Conversion/HALToVM/ConvertBufferOps.cpp
@@ -27,7 +27,7 @@ ConversionPatternRewriter &rewriter) const override { IREE::HAL::BufferLoadOp::Adaptor adaptor(operands); auto importType = importOp.getType(); - auto sizeConst = rewriter.createOrFold<mlir::ConstantOp>( + auto sizeConst = rewriter.createOrFold<mlir::arith::ConstantOp>( op.getLoc(), rewriter.getI32IntegerAttr( IREE::HAL::getRoundedElementByteWidth(op.getResult().getType()))); @@ -44,7 +44,8 @@ if (newResultType == callResult.getType()) { rewriter.replaceOp(op, {callResult}); } else { - rewriter.replaceOpWithNewOp<BitcastOp>(op, newResultType, callResult); + rewriter.replaceOpWithNewOp<arith::BitcastOp>(op, newResultType, + callResult); } return success(); @@ -69,7 +70,7 @@ ConversionPatternRewriter &rewriter) const override { IREE::HAL::BufferStoreOp::Adaptor adaptor(operands); auto importType = importOp.getType(); - auto sizeConst = rewriter.createOrFold<mlir::ConstantOp>( + auto sizeConst = rewriter.createOrFold<mlir::arith::ConstantOp>( op.getLoc(), rewriter.getI32IntegerAttr( IREE::HAL::getRoundedElementByteWidth(op.value().getType())));
diff --git a/iree/compiler/Dialect/HAL/Conversion/HALToVM/ConvertDeviceOps.cpp b/iree/compiler/Dialect/HAL/Conversion/HALToVM/ConvertDeviceOps.cpp index 55d1ceb..37a3002 100644 --- a/iree/compiler/Dialect/HAL/Conversion/HALToVM/ConvertDeviceOps.cpp +++ b/iree/compiler/Dialect/HAL/Conversion/HALToVM/ConvertDeviceOps.cpp
@@ -43,19 +43,19 @@ // Truncate or extend based on the target type. if (targetType.isIndex()) { // i32 -> index cast. - value = - rewriter.createOrFold<IndexCastOp>(op.getLoc(), targetType, value); + value = rewriter.createOrFold<arith::IndexCastOp>(op.getLoc(), targetType, + value); } else if (targetType.isa<IntegerType>()) { // i32 -> {integer} cast. if (targetType.getIntOrFloatBitWidth() < value.getType().getIntOrFloatBitWidth()) { // i32 -> narrowing cast. - value = - rewriter.createOrFold<TruncateIOp>(op.getLoc(), targetType, value); + value = rewriter.createOrFold<arith::TruncIOp>(op.getLoc(), targetType, + value); } else { // i32 -> widening cast. - value = rewriter.createOrFold<ZeroExtendIOp>(op.getLoc(), targetType, - value); + value = rewriter.createOrFold<arith::ExtUIOp>(op.getLoc(), targetType, + value); } } @@ -65,8 +65,8 @@ // already handled the error case. value = rewriter.createOrFold<SelectOp>( op.getLoc(), ok, value, - rewriter.createOrFold<ConstantOp>(op.getLoc(), - op.default_valueAttr())); + rewriter.createOrFold<arith::ConstantOp>(op.getLoc(), + op.default_valueAttr())); ok = rewriter.createOrFold<IREE::VM::ConstI32Op>(op.getLoc(), 1); }
diff --git a/iree/compiler/Dialect/HAL/Conversion/HALToVM/test/allocator_ops.mlir b/iree/compiler/Dialect/HAL/Conversion/HALToVM/test/allocator_ops.mlir index 1dd1bba..bf7eb81 100644 --- a/iree/compiler/Dialect/HAL/Conversion/HALToVM/test/allocator_ops.mlir +++ b/iree/compiler/Dialect/HAL/Conversion/HALToVM/test/allocator_ops.mlir
@@ -4,9 +4,9 @@ func @allocatorComputeSizeFoldsAway(%arg0 : !hal.allocator) -> index { // CHECK: %c4194304 = vm.const.i32 4194304 : i32 // CHECK-NOT: hal.allocator.compute_size - %c1024 = constant 1024 : index - %c1_i32 = constant 1 : i32 - %c32_i32 = constant 32 : i32 + %c1024 = arith.constant 1024 : index + %c1_i32 = arith.constant 1 : i32 + %c32_i32 = arith.constant 32 : i32 %0 = hal.allocator.compute_size<%arg0 : !hal.allocator> shape([%c1024, %c1024]) type(%c32_i32) encoding(%c1_i32) : index return %0 : index @@ -16,7 +16,7 @@ // CHECK-LABEL: vm.func private @allocatorAllocate func @allocatorAllocate(%arg0 : !hal.allocator) -> !hal.buffer { - %c1024 = constant 1024 : index + %c1024 = arith.constant 1024 : index // CHECK: %ref = vm.call @hal.allocator.allocate(%arg0, %c6, %c14, %c1024) : (!vm.ref<!hal.allocator>, i32, i32, i32) -> !vm.ref<!hal.buffer> %0 = hal.allocator.allocate<%arg0 : !hal.allocator> type("HostLocal") usage("All") : !hal.buffer{%c1024} return %0 : !hal.buffer @@ -26,8 +26,8 @@ // CHECK-LABEL: vm.func private @allocatorMapByteBuffer func @allocatorMapByteBuffer(%arg0 : !hal.allocator, %arg1 : !util.byte_buffer) -> !hal.buffer { - %offset = constant 128 : index - %length = constant 256 : index + %offset = arith.constant 128 : index + %length = arith.constant 256 : index // CHECK: = vm.call @hal.allocator.wrap.byte_buffer(%arg0, %c6, %c2, %arg1, %c128, %c256) : (!vm.ref<!hal.allocator>, i32, i32, !vm.buffer, i32, i32) -> !vm.ref<!hal.buffer> %buffer = hal.allocator.map<%arg0 : !hal.allocator> source(%arg1 : !util.byte_buffer)[%offset, %length] type("HostVisible|HostCoherent") usage(Transfer) : !hal.buffer return %buffer : !hal.buffer
diff --git a/iree/compiler/Dialect/HAL/Conversion/HALToVM/test/buffer_ops.mlir b/iree/compiler/Dialect/HAL/Conversion/HALToVM/test/buffer_ops.mlir index ae78f27..e943bb0 100644 --- a/iree/compiler/Dialect/HAL/Conversion/HALToVM/test/buffer_ops.mlir +++ b/iree/compiler/Dialect/HAL/Conversion/HALToVM/test/buffer_ops.mlir
@@ -2,8 +2,8 @@ // CHECK-LABEL: @buffer_subspan func @buffer_subspan(%arg0 : !hal.buffer) -> !hal.buffer { - %c42 = constant 42 : index - %c43 = constant 43 : index + %c42 = arith.constant 42 : index + %c43 = arith.constant 43 : index // CHECK: %ref = vm.call @hal.buffer.subspan(%arg0, %c42, %c43) : (!vm.ref<!hal.buffer>, i32, i32) -> !vm.ref<!hal.buffer> %buffer = hal.buffer.subspan<%arg0 : !hal.buffer>[%c42, %c43] : !hal.buffer return %buffer : !hal.buffer @@ -13,11 +13,11 @@ // CHECK-LABEL: @buffer_load func @buffer_load(%arg0 : !hal.buffer) -> (i8, i16, i32, f16, f32) { - %c42 = constant 42 : index - %c43 = constant 43 : index - %c44 = constant 44 : index - %c45 = constant 45 : index - %c46 = constant 46 : index + %c42 = arith.constant 42 : index + %c43 = arith.constant 43 : index + %c44 = arith.constant 44 : index + %c45 = arith.constant 45 : index + %c46 = arith.constant 46 : index // CHECK: %0 = vm.call @hal.buffer.load(%arg0, %c42, %c1) : (!vm.ref<!hal.buffer>, i32, i32) -> i32 %0 = hal.buffer.load<%arg0 : !hal.buffer>[%c42] : i8 // CHECK: %1 = vm.call @hal.buffer.load(%arg0, %c43, %c2) : (!vm.ref<!hal.buffer>, i32, i32) -> i32 @@ -35,11 +35,11 @@ // CHECK-LABEL: @buffer_store func @buffer_store(%arg0 : !hal.buffer, %arg1 : i8, %arg2 : i16, %arg3 : i32, %arg4 : f16, %arg5 : f32) { - %c42 = constant 42 : index - %c43 = constant 43 : index - %c44 = constant 44 : index - %c45 = constant 45 : index - %c46 = constant 46 : index + %c42 = arith.constant 42 : index + %c43 = arith.constant 43 : index + %c44 = arith.constant 44 : index + %c45 = arith.constant 45 : index + %c46 = arith.constant 46 : index // CHECK: vm.call @hal.buffer.store(%arg1, %arg0, %c42, %c1) : (i32, !vm.ref<!hal.buffer>, i32, i32) -> () hal.buffer.store<%arg0 : !hal.buffer>[%c42] value(%arg1 : i8) // CHECK: vm.call @hal.buffer.store(%arg2, %arg0, %c43, %c2) : (i32, !vm.ref<!hal.buffer>, i32, i32) -> ()
diff --git a/iree/compiler/Dialect/HAL/Conversion/HALToVM/test/command_buffer_ops.mlir b/iree/compiler/Dialect/HAL/Conversion/HALToVM/test/command_buffer_ops.mlir index c7c8174..4699c05 100644 --- a/iree/compiler/Dialect/HAL/Conversion/HALToVM/test/command_buffer_ops.mlir +++ b/iree/compiler/Dialect/HAL/Conversion/HALToVM/test/command_buffer_ops.mlir
@@ -40,9 +40,9 @@ %arg0: !hal.command_buffer, %arg1: !hal.buffer ) { - %c100 = constant 100 : index - %c200 = constant 200 : index - %c300 = constant 300 : i32 + %c100 = arith.constant 100 : index + %c200 = arith.constant 200 : index + %c300 = arith.constant 300 : i32 // CHECK: vm.call @hal.command_buffer.fill_buffer(%arg0, %arg1, %c100, %c200, %c300) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.buffer>, i32, i32, i32) -> () hal.command_buffer.fill_buffer<%arg0 : !hal.command_buffer> target(%arg1 : !hal.buffer)[%c100, %c200] @@ -57,9 +57,9 @@ %arg0: !hal.command_buffer, %arg1: !hal.buffer ) { - %c100 = constant 100 : index - %c200 = constant 200 : index - %c300 = constant 300 : index + %c100 = arith.constant 100 : index + %c200 = arith.constant 200 : index + %c300 = arith.constant 300 : index // CHECK: vm.call @hal.command_buffer.copy_buffer(%arg0, %arg1, %c100, %arg1, %c200, %c300) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.buffer>, i32, !vm.ref<!hal.buffer>, i32, i32) -> () hal.command_buffer.copy_buffer<%arg0 : !hal.command_buffer> source(%arg1 : !hal.buffer)[%c100] @@ -76,8 +76,8 @@ %arg1: !hal.executable_layout, %arg2: !hal.descriptor_set ) { - %c0 = constant 0 : index - %c100 = constant 100 : index + %c0 = arith.constant 0 : index + %c100 = arith.constant 100 : index // CHECK: vm.call.variadic @hal.command_buffer.bind_descriptor_set(%arg0, %arg1, %zero, %arg2, []) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable_layout>, i32, !vm.ref<!hal.descriptor_set>, i32 ...) hal.command_buffer.bind_descriptor_set<%arg0 : !hal.command_buffer> layout(%arg1 : !hal.executable_layout)[%c0] @@ -97,9 +97,9 @@ %arg0: !hal.command_buffer, %arg1: !hal.executable ) { - %c100 = constant 100 : index - %c200 = constant 200 : index - %c300 = constant 300 : index + %c100 = arith.constant 100 : index + %c200 = arith.constant 200 : index + %c300 = arith.constant 300 : index // CHECK: vm.call @hal.command_buffer.dispatch(%arg0, %arg1, %zero, %c100, %c200, %c300) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable>, i32, i32, i32, i32) -> () hal.command_buffer.dispatch<%arg0 : !hal.command_buffer> target(%arg1 : !hal.executable)[0] @@ -115,7 +115,7 @@ %arg1: !hal.executable, %arg2: !hal.buffer ) { - %c100 = constant 100 : index + %c100 = arith.constant 100 : index // CHECK: vm.call @hal.command_buffer.dispatch.indirect(%arg0, %arg1, %zero, %arg2, %c100) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable>, i32, !vm.ref<!hal.buffer>, i32) -> () hal.command_buffer.dispatch.indirect<%arg0 : !hal.command_buffer> target(%arg1 : !hal.executable)[0]
diff --git a/iree/compiler/Dialect/HAL/Conversion/HALToVM/test/constant_ops.mlir b/iree/compiler/Dialect/HAL/Conversion/HALToVM/test/constant_ops.mlir index 424790a..c814c9d 100644 --- a/iree/compiler/Dialect/HAL/Conversion/HALToVM/test/constant_ops.mlir +++ b/iree/compiler/Dialect/HAL/Conversion/HALToVM/test/constant_ops.mlir
@@ -15,8 +15,8 @@ util.global private @pool_storage0_buffer : !hal.buffer // CHECK-NEXT: vm.initializer { util.initializer { - %c0 = constant 0 : index - %c16 = constant 16 : index + %c0 = arith.constant 0 : index + %c16 = arith.constant 16 : index %dev = hal.ex.shared_device : !hal.device %allocator = hal.device.allocator<%dev : !hal.device> : !hal.allocator // CHECK: %[[STORAGE_REF:.+]] = vm.const.ref.rodata @pool_storage0 : !vm.buffer @@ -38,12 +38,12 @@ util.global private @pool_splats : !hal.buffer // CHECK: vm.initializer { util.initializer { - %c64 = constant 64 : index - %c0 = constant 0 : index - %c4 = constant 4 : index - %c1065353216_i32 = constant 1065353216 : i32 - %c32 = constant 32 : index - %c1234567890_i32 = constant 1234567890 : i32 + %c64 = arith.constant 64 : index + %c0 = arith.constant 0 : index + %c4 = arith.constant 4 : index + %c1065353216_i32 = arith.constant 1065353216 : i32 + %c32 = arith.constant 32 : index + %c1234567890_i32 = arith.constant 1234567890 : i32 %dev = hal.ex.shared_device : !hal.device %allocator = hal.device.allocator<%dev : !hal.device> : !hal.allocator // CHECK: [[BUFFER:%.+]] = vm.call @hal.allocator.allocate({{.+}}, %c50, %c15, %c64)
diff --git a/iree/compiler/Dialect/HAL/Conversion/StandardToHAL/ConvertConstantOps.cpp b/iree/compiler/Dialect/HAL/Conversion/StandardToHAL/ConvertConstantOps.cpp index 199cfbe..bb91412 100644 --- a/iree/compiler/Dialect/HAL/Conversion/StandardToHAL/ConvertConstantOps.cpp +++ b/iree/compiler/Dialect/HAL/Conversion/StandardToHAL/ConvertConstantOps.cpp
@@ -23,13 +23,13 @@ namespace { class ConstantTensorOpConversion - : public OpConversionPattern<mlir::ConstantOp> { + : public OpConversionPattern<mlir::arith::ConstantOp> { public: ConstantTensorOpConversion(MLIRContext *ctx, TypeConverter &converter) : OpConversionPattern(ctx) {} LogicalResult matchAndRewrite( - mlir::ConstantOp constantOp, llvm::ArrayRef<Value> newOperands, + mlir::arith::ConstantOp constantOp, llvm::ArrayRef<Value> newOperands, ConversionPatternRewriter &rewriter) const override { if (!constantOp.getType().isa<TensorType>()) return failure(); @@ -46,7 +46,7 @@ IREE::HAL::BufferUsageBitfield::All | IREE::HAL::BufferUsageBitfield::Constant; - auto elementsAttr = constantOp.getValue().cast<ElementsAttr>(); + auto elementsAttr = constantOp.value().cast<ElementsAttr>(); auto elementsTy = elementsAttr.getType().cast<ShapedType>(); // Expand boolean elements to the minimum bit widht supported by the HAL
diff --git a/iree/compiler/Dialect/HAL/Conversion/StandardToHAL/ConvertStandardToHAL.cpp b/iree/compiler/Dialect/HAL/Conversion/StandardToHAL/ConvertStandardToHAL.cpp index 2eb1ac7..634d2d4 100644 --- a/iree/compiler/Dialect/HAL/Conversion/StandardToHAL/ConvertStandardToHAL.cpp +++ b/iree/compiler/Dialect/HAL/Conversion/StandardToHAL/ConvertStandardToHAL.cpp
@@ -78,8 +78,8 @@ if (sourceType.isDynamicDim(i)) { shapeDims[i] = sourceDims[sourceDimsIndex++]; } else { - shapeDims[i] = - builder.create<ConstantIndexOp>(loc, sourceType.getDimSize(i)); + shapeDims[i] = builder.create<arith::ConstantIndexOp>( + loc, sourceType.getDimSize(i)); } } return shapeDims;
diff --git a/iree/compiler/Dialect/HAL/Conversion/StandardToHAL/test/constant_ops.mlir b/iree/compiler/Dialect/HAL/Conversion/StandardToHAL/test/constant_ops.mlir index cb56984..291c652 100644 --- a/iree/compiler/Dialect/HAL/Conversion/StandardToHAL/test/constant_ops.mlir +++ b/iree/compiler/Dialect/HAL/Conversion/StandardToHAL/test/constant_ops.mlir
@@ -8,7 +8,7 @@ // CHECK-SAME: usage("Constant|Transfer|Mapping|Dispatch") // CHECK-SAME: : !hal.buffer // CHECK-SAME: = dense<[1, 2]> : tensor<2xi32> - %0 = constant dense<[1, 2]> : tensor<2xi32> + %0 = arith.constant dense<[1, 2]> : tensor<2xi32> return } @@ -22,6 +22,6 @@ // CHECK-SAME: usage("Constant|Transfer|Mapping|Dispatch") // CHECK-SAME: : !hal.buffer // CHECK-SAME: = dense<[1, 0]> : tensor<2xi8> - %0 = constant dense<[1, 0]> : tensor<2xi1> + %0 = arith.constant dense<[1, 0]> : tensor<2xi1> return }
diff --git a/iree/compiler/Dialect/HAL/Conversion/StandardToHAL/test/structural_ops.mlir b/iree/compiler/Dialect/HAL/Conversion/StandardToHAL/test/structural_ops.mlir index 86abd72..c8f919a 100644 --- a/iree/compiler/Dialect/HAL/Conversion/StandardToHAL/test/structural_ops.mlir +++ b/iree/compiler/Dialect/HAL/Conversion/StandardToHAL/test/structural_ops.mlir
@@ -32,12 +32,12 @@ // Just anchor on the end of the function that creates a new buffer view. // CHECK: hal.ex.submit_and_wait - // CHECK-DAG: %[[C3:.*]] = constant 3 : index - // CHECK-DAG: %[[C2:.*]] = constant 2 : index - // CHECK-DAG: %[[C1_1:.*]] = constant 1 : index - // CHECK-DAG: %[[C4:.*]] = constant 4 : index - // CHECK-DAG: %[[C1_2:.*]] = constant 1 : index - // CHECK-DAG: %[[ET:.*]] = constant 50331680 : i32 + // CHECK-DAG: %[[C3:.*]] = arith.constant 3 : index + // CHECK-DAG: %[[C2:.*]] = arith.constant 2 : index + // CHECK-DAG: %[[C1_1:.*]] = arith.constant 1 : index + // CHECK-DAG: %[[C4:.*]] = arith.constant 4 : index + // CHECK-DAG: %[[C1_2:.*]] = arith.constant 1 : index + // CHECK-DAG: %[[ET:.*]] = arith.constant 50331680 : i32 // CHECK: %[[VIEW:.*]] = hal.buffer_view.create // CHECK-SAME: shape([%[[C3]], %[[C2]], %[[C1_1]], %[[C4]], %[[C1_2]]]) // CHECK-SAME: type(%[[ET]])
diff --git a/iree/compiler/Dialect/HAL/Conversion/UtilToHAL/ConvertUtilToHAL.cpp b/iree/compiler/Dialect/HAL/Conversion/UtilToHAL/ConvertUtilToHAL.cpp index 908520c..49d2a6a 100644 --- a/iree/compiler/Dialect/HAL/Conversion/UtilToHAL/ConvertUtilToHAL.cpp +++ b/iree/compiler/Dialect/HAL/Conversion/UtilToHAL/ConvertUtilToHAL.cpp
@@ -61,7 +61,7 @@ SmallVector<Value, 4> shape; if (shapedType.getRank() >= 1) { for (auto dim : shapedType.getShape()) { - shape.push_back(rewriter.createOrFold<mlir::ConstantIndexOp>( + shape.push_back(rewriter.createOrFold<mlir::arith::ConstantIndexOp>( constantOp.getLoc(), dim)); } }
diff --git a/iree/compiler/Dialect/HAL/IR/HALDialect.cpp b/iree/compiler/Dialect/HAL/IR/HALDialect.cpp index b4ba817..7b38e4a 100644 --- a/iree/compiler/Dialect/HAL/IR/HALDialect.cpp +++ b/iree/compiler/Dialect/HAL/IR/HALDialect.cpp
@@ -129,7 +129,7 @@ if (type.isa<IndexType>()) { // Some folders materialize raw index types, which just become std // constants. - return builder.create<mlir::ConstantIndexOp>( + return builder.create<mlir::arith::ConstantIndexOp>( loc, value.cast<IntegerAttr>().getValue().getSExtValue()); } return nullptr;
diff --git a/iree/compiler/Dialect/HAL/IR/HALOpFolders.cpp b/iree/compiler/Dialect/HAL/IR/HALOpFolders.cpp index bc72746..8a8fab4 100644 --- a/iree/compiler/Dialect/HAL/IR/HALOpFolders.cpp +++ b/iree/compiler/Dialect/HAL/IR/HALOpFolders.cpp
@@ -44,7 +44,7 @@ ValueRange shapeDims, OpBuilder &builder) { Value value = baseValue; for (auto dim : shapeDims) { - value = builder.createOrFold<mlir::MulIOp>(loc, value, dim); + value = builder.createOrFold<mlir::arith::MulIOp>(loc, value, dim); } return value; } @@ -94,7 +94,8 @@ // TODO(#6762): switch based on op.encoding(). - auto offset = rewriter.createOrFold<mlir::ConstantIndexOp>(op.getLoc(), 0); + auto offset = + rewriter.createOrFold<mlir::arith::ConstantIndexOp>(op.getLoc(), 0); for (size_t i = 0; i < op.indices().size(); ++i) { // TODO(benvanik): check error case in debug builds. // if (indices[i] >= shape[i]) { @@ -104,16 +105,16 @@ // } auto axisOffset = op.indices()[i]; for (size_t j = i + 1; j < op.shape().size(); ++j) { - axisOffset = rewriter.createOrFold<mlir::MulIOp>( + axisOffset = rewriter.createOrFold<mlir::arith::MulIOp>( op.getLoc(), axisOffset, op.shape()[j]); } - offset = - rewriter.createOrFold<mlir::AddIOp>(op.getLoc(), offset, axisOffset); + offset = rewriter.createOrFold<mlir::arith::AddIOp>(op.getLoc(), offset, + axisOffset); } auto elementSize = getElementByteCount(op.getLoc(), op.element_type(), rewriter); - auto byteOffset = - rewriter.createOrFold<mlir::MulIOp>(op.getLoc(), offset, elementSize); + auto byteOffset = rewriter.createOrFold<mlir::arith::MulIOp>( + op.getLoc(), offset, elementSize); rewriter.replaceOp(op, {byteOffset}); return success(); @@ -140,12 +141,13 @@ BufferConstraintsAdaptor bufferConstraints(op.getLoc(), op.allocator()); SmallVector<Value, 6> endIndices(op.shape().size()); - auto one = rewriter.createOrFold<mlir::ConstantIndexOp>(op.getLoc(), 1); + auto one = + rewriter.createOrFold<mlir::arith::ConstantIndexOp>(op.getLoc(), 1); for (size_t i = 0; i < endIndices.size(); ++i) { - endIndices[i] = rewriter.createOrFold<mlir::SubIOp>( + endIndices[i] = rewriter.createOrFold<mlir::arith::SubIOp>( op.getLoc(), - rewriter.createOrFold<mlir::AddIOp>(op.getLoc(), op.indices()[i], - op.lengths()[i]), + rewriter.createOrFold<mlir::arith::AddIOp>( + op.getLoc(), op.indices()[i], op.lengths()[i]), one); } @@ -158,10 +160,10 @@ auto elementSize = getElementByteCount(op.getLoc(), op.element_type(), rewriter); - auto offsetLength = rewriter.createOrFold<mlir::AddIOp>( + auto offsetLength = rewriter.createOrFold<mlir::arith::AddIOp>( op.getLoc(), - rewriter.createOrFold<mlir::SubIOp>(op.getLoc(), endByteOffset, - startByteOffset), + rewriter.createOrFold<mlir::arith::SubIOp>(op.getLoc(), endByteOffset, + startByteOffset), elementSize); rewriter.replaceOp(op, {startByteOffset, offsetLength}); @@ -209,8 +211,10 @@ auto hostBuffer = rewriter.createOrFold<IREE::Util::ByteBufferConstantOp>( op.getLoc(), IREE::Util::ByteBufferType::get(rewriter.getContext()), op.value()); - auto zero = rewriter.createOrFold<mlir::ConstantIndexOp>(op.getLoc(), 0); - auto neg1 = rewriter.createOrFold<mlir::ConstantIndexOp>(op.getLoc(), -1); + auto zero = + rewriter.createOrFold<mlir::arith::ConstantIndexOp>(op.getLoc(), 0); + auto neg1 = + rewriter.createOrFold<mlir::arith::ConstantIndexOp>(op.getLoc(), -1); auto deviceBuffer = rewriter.createOrFold<AllocatorMapOp>( op.getLoc(), bufferType, op.allocator(), memoryTypes, bufferUsage, hostBuffer, zero, neg1); @@ -220,8 +224,8 @@ SmallVector<Value, 4> shape; if (shapedType.getRank() >= 1) { for (auto dim : shapedType.getShape()) { - shape.push_back( - rewriter.createOrFold<mlir::ConstantIndexOp>(op.getLoc(), dim)); + shape.push_back(rewriter.createOrFold<mlir::arith::ConstantIndexOp>( + op.getLoc(), dim)); } } auto bufferView = rewriter.createOrFold<BufferViewCreateOp>( @@ -289,9 +293,9 @@ // Zero offsets don't do anything and can just be removed so we can avoid // inserting a bunch of additional IR. - if (auto constantOp = - dyn_cast_or_null<ConstantIndexOp>(baseOffset.getDefiningOp())) { - if (constantOp.getValue() == 0) { + if (auto constantOp = dyn_cast_or_null<arith::ConstantIndexOp>( + baseOffset.getDefiningOp())) { + if (constantOp.value() == 0) { return success(); } } @@ -299,8 +303,8 @@ // Propagate the offset to all returned slice offsets. rewriter.setInsertionPointAfter(op); for (auto sliceOffset : op.packed_offsets()) { - auto addOp = - rewriter.create<mlir::AddIOp>(op.getLoc(), baseOffset, sliceOffset); + auto addOp = rewriter.create<mlir::arith::AddIOp>(op.getLoc(), baseOffset, + sliceOffset); SmallPtrSet<Operation *, 1> exclusions; exclusions.insert(addOp); sliceOffset.replaceAllUsesExcept(addOp.result(), exclusions); @@ -509,7 +513,7 @@ if (auto subspanOp = dyn_cast_or_null<BufferSubspanOp>( op.target_buffer().getDefiningOp())) { newTargetBuffer = subspanOp.source_buffer(); - newTargetOffset = rewriter.createOrFold<mlir::AddIOp>( + newTargetOffset = rewriter.createOrFold<mlir::arith::AddIOp>( subspanOp.getLoc(), subspanOp.source_offset(), op.target_offset()); needsUpdate = true; } @@ -547,7 +551,7 @@ if (auto subspanOp = dyn_cast_or_null<BufferSubspanOp>( op.source_buffer().getDefiningOp())) { newSourceBuffer = subspanOp.source_buffer(); - newSourceOffset = rewriter.createOrFold<mlir::AddIOp>( + newSourceOffset = rewriter.createOrFold<mlir::arith::AddIOp>( subspanOp.getLoc(), subspanOp.source_offset(), op.source_offset()); needsUpdate = true; } @@ -556,7 +560,7 @@ if (auto subspanOp = dyn_cast_or_null<BufferSubspanOp>( op.target_buffer().getDefiningOp())) { newTargetBuffer = subspanOp.source_buffer(); - newTargetOffset = rewriter.createOrFold<mlir::AddIOp>( + newTargetOffset = rewriter.createOrFold<mlir::arith::AddIOp>( subspanOp.getLoc(), subspanOp.source_offset(), op.target_offset()); needsUpdate = true; } @@ -600,7 +604,7 @@ if (auto subspanOp = dyn_cast<BufferSubspanOp>(definingOp)) { needsUpdate = true; bindingBuffers[i] = subspanOp.source_buffer(); - bindingOffsets[i] = rewriter.createOrFold<mlir::AddIOp>( + bindingOffsets[i] = rewriter.createOrFold<mlir::arith::AddIOp>( subspanOp.getLoc(), subspanOp.source_offset(), bindingOffsets[i]); } }
diff --git a/iree/compiler/Dialect/HAL/IR/HALOps.cpp b/iree/compiler/Dialect/HAL/IR/HALOps.cpp index 3652cf6..35d883a 100644 --- a/iree/compiler/Dialect/HAL/IR/HALOps.cpp +++ b/iree/compiler/Dialect/HAL/IR/HALOps.cpp
@@ -257,8 +257,10 @@ Value allocator, ValueRange shape, int32_t elementType, int32_t encodingType) { build(builder, state, allocator, shape, - builder.createOrFold<ConstantIntOp>(state.location, elementType, 32), - builder.createOrFold<ConstantIntOp>(state.location, encodingType, 32)); + builder.createOrFold<arith::ConstantIntOp>(state.location, elementType, + 32), + builder.createOrFold<arith::ConstantIntOp>(state.location, encodingType, + 32)); } void AllocatorComputeSizeOp::build(OpBuilder &builder, OperationState &state, @@ -285,8 +287,10 @@ int32_t elementType, int32_t encodingType, ValueRange indices) { build(builder, state, allocator, shape, - builder.createOrFold<ConstantIntOp>(state.location, elementType, 32), - builder.createOrFold<ConstantIntOp>(state.location, encodingType, 32), + builder.createOrFold<arith::ConstantIntOp>(state.location, elementType, + 32), + builder.createOrFold<arith::ConstantIntOp>(state.location, encodingType, + 32), indices); } @@ -316,8 +320,10 @@ int32_t elementType, int32_t encodingType, ValueRange indices, ValueRange lengths) { build(builder, state, allocator, shape, - builder.createOrFold<ConstantIntOp>(state.location, elementType, 32), - builder.createOrFold<ConstantIntOp>(state.location, encodingType, 32), + builder.createOrFold<arith::ConstantIntOp>(state.location, elementType, + 32), + builder.createOrFold<arith::ConstantIntOp>(state.location, encodingType, + 32), indices, lengths); } @@ -455,8 +461,10 @@ Value buffer, int32_t elementType, int32_t encodingType, ValueRange shape) { build(builder, state, buffer, - builder.createOrFold<ConstantIntOp>(state.location, elementType, 32), - builder.createOrFold<ConstantIntOp>(state.location, encodingType, 32), + builder.createOrFold<arith::ConstantIntOp>(state.location, elementType, + 32), + builder.createOrFold<arith::ConstantIntOp>(state.location, encodingType, + 32), shape); } @@ -562,7 +570,8 @@ Value executableLayout, int64_t set, ArrayRef<DescriptorSetBindingValue> bindings) { build(builder, state, commandBuffer, executableLayout, - builder.createOrFold<ConstantIndexOp>(state.location, set), bindings); + builder.createOrFold<arith::ConstantIndexOp>(state.location, set), + bindings); } void CommandBufferPushDescriptorSetOp::build(
diff --git a/iree/compiler/Dialect/HAL/IR/HALOps.td b/iree/compiler/Dialect/HAL/IR/HALOps.td index 1fc4cdf..263e303 100644 --- a/iree/compiler/Dialect/HAL/IR/HALOps.td +++ b/iree/compiler/Dialect/HAL/IR/HALOps.td
@@ -1681,9 +1681,9 @@ return the same types. ```mlir - %c0 = constant 0 : i32 - %c1 = constant 1 : i32 - %c2 = constant 2 : i32 + %c0 = arith.constant 0 : i32 + %c1 = arith.constant 1 : i32 + %c2 = arith.constant 2 : i32 %device = ... : !hal.device %0 = hal.device.switch<%device : !hal.device> -> i32 #hal.device.match.id<"vulkan-v1.?-*"> {
diff --git a/iree/compiler/Dialect/HAL/IR/HALTypes.cpp b/iree/compiler/Dialect/HAL/IR/HALTypes.cpp index 17afd26..6e73569 100644 --- a/iree/compiler/Dialect/HAL/IR/HALTypes.cpp +++ b/iree/compiler/Dialect/HAL/IR/HALTypes.cpp
@@ -338,11 +338,11 @@ } Value getElementBitCount(Location loc, Value elementType, OpBuilder &builder) { - return builder.createOrFold<AndOp>( + return builder.createOrFold<arith::AndIOp>( loc, - builder.createOrFold<IndexCastOp>(loc, builder.getIndexType(), - elementType), - builder.createOrFold<ConstantIndexOp>(loc, 0xFF)); + builder.createOrFold<arith::IndexCastOp>(loc, builder.getIndexType(), + elementType), + builder.createOrFold<arith::ConstantIndexOp>(loc, 0xFF)); } size_t getElementByteCount(IntegerAttr elementType) { @@ -350,13 +350,13 @@ } Value getElementByteCount(Location loc, Value elementType, OpBuilder &builder) { - auto c1 = builder.createOrFold<ConstantIndexOp>(loc, 1); - auto c8 = builder.createOrFold<ConstantIndexOp>(loc, 8); + auto c1 = builder.createOrFold<arith::ConstantIndexOp>(loc, 1); + auto c8 = builder.createOrFold<arith::ConstantIndexOp>(loc, 8); auto bitCount = getElementBitCount(loc, elementType, builder); - return builder.createOrFold<UnsignedDivIOp>( + return builder.createOrFold<arith::DivUIOp>( loc, - builder.createOrFold<SubIOp>( - loc, builder.createOrFold<AddIOp>(loc, bitCount, c8), c1), + builder.createOrFold<arith::SubIOp>( + loc, builder.createOrFold<arith::AddIOp>(loc, bitCount, c8), c1), c8); } @@ -462,23 +462,23 @@ } Value BufferConstraintsAdaptor::getMaxAllocationSize(OpBuilder &builder) { - return builder.createOrFold<mlir::ConstantOp>( + return builder.createOrFold<mlir::arith::ConstantOp>( loc_, bufferConstraints_.max_allocation_sizeAttr()); } Value BufferConstraintsAdaptor::getMinBufferOffsetAlignment( OpBuilder &builder) { - return builder.createOrFold<mlir::ConstantOp>( + return builder.createOrFold<mlir::arith::ConstantOp>( loc_, bufferConstraints_.min_buffer_offset_alignmentAttr()); } Value BufferConstraintsAdaptor::getMaxBufferRange(OpBuilder &builder) { - return builder.createOrFold<mlir::ConstantOp>( + return builder.createOrFold<mlir::arith::ConstantOp>( loc_, bufferConstraints_.max_buffer_rangeAttr()); } Value BufferConstraintsAdaptor::getMinBufferRangeAlignment(OpBuilder &builder) { - return builder.createOrFold<mlir::ConstantOp>( + return builder.createOrFold<mlir::arith::ConstantOp>( loc_, bufferConstraints_.min_buffer_range_alignmentAttr()); } @@ -761,7 +761,8 @@ Value MatchAlwaysAttr::buildConditionExpression(Location loc, Value value, OpBuilder builder) const { // #hal.match.always -> true - return builder.createOrFold<ConstantIntOp>(loc, /*value=*/1, /*width=*/1); + return builder.createOrFold<arith::ConstantIntOp>(loc, /*value=*/1, + /*width=*/1); } static ArrayAttr parseMultiMatchAttrArray(DialectAsmParser &p) { @@ -807,7 +808,7 @@ // #hal.match.any<[a, b, c]> -> or(or(a, b), c) if (getConditions().empty()) { // Empty returns false (no conditions match). - return builder.create<ConstantIntOp>(loc, /*value=*/0, /*width=*/1); + return builder.create<arith::ConstantIntOp>(loc, /*value=*/0, /*width=*/1); } auto conditionValues = llvm::map_range(getConditions(), [&](MatchAttrInterface attr) { @@ -815,8 +816,8 @@ }); Value resultValue; for (auto conditionValue : conditionValues) { - resultValue = resultValue ? builder.createOrFold<OrOp>(loc, resultValue, - conditionValue) + resultValue = resultValue ? builder.createOrFold<arith::OrIOp>( + loc, resultValue, conditionValue) : conditionValue; } return resultValue; @@ -837,7 +838,7 @@ // #hal.match.all<[a, b, c]> -> and(and(a, b), c) if (getConditions().empty()) { // Empty returns true (all 0 conditions match). - return builder.create<ConstantIntOp>(loc, /*value=*/1, /*width=*/1); + return builder.create<arith::ConstantIntOp>(loc, /*value=*/1, /*width=*/1); } auto conditionValues = llvm::map_range(getConditions(), [&](MatchAttrInterface attr) { @@ -845,8 +846,8 @@ }); Value resultValue; for (auto conditionValue : conditionValues) { - resultValue = resultValue ? builder.createOrFold<AndOp>(loc, resultValue, - conditionValue) + resultValue = resultValue ? builder.createOrFold<arith::AndIOp>( + loc, resultValue, conditionValue) : conditionValue; } return resultValue;
diff --git a/iree/compiler/Dialect/HAL/IR/LoweringConfig.cpp b/iree/compiler/Dialect/HAL/IR/LoweringConfig.cpp index 9ec9570..4bc7e4e 100644 --- a/iree/compiler/Dialect/HAL/IR/LoweringConfig.cpp +++ b/iree/compiler/Dialect/HAL/IR/LoweringConfig.cpp
@@ -137,7 +137,7 @@ unsigned level) { return llvm::to_vector<4>( llvm::map_range(getTileSizes(op, level), [&](int64_t t) -> Value { - return b.create<ConstantIndexOp>(op->getLoc(), t); + return b.create<arith::ConstantIndexOp>(op->getLoc(), t); })); }
diff --git a/iree/compiler/Dialect/HAL/IR/test/allocator_op_folding.mlir b/iree/compiler/Dialect/HAL/IR/test/allocator_op_folding.mlir index 7f5cbb5..e31e854 100644 --- a/iree/compiler/Dialect/HAL/IR/test/allocator_op_folding.mlir +++ b/iree/compiler/Dialect/HAL/IR/test/allocator_op_folding.mlir
@@ -44,7 +44,7 @@ // CHECK-LABEL: @allocator_pack_zero_offset func @allocator_pack_zero_offset(%allocator: !hal.allocator, %size : index) -> index { // CHECK-NOT: constant 0 - %base_offset = constant 0 : index + %base_offset = arith.constant 0 : index // CHECK: hal.allocator.pack<{{.+}}> slices({ %total_length, %offset_0, %offset_1 = hal.allocator.pack<%allocator : !hal.allocator> @@ -62,7 +62,7 @@ // CHECK-LABEL: @allocator_pack_no_slices func @allocator_pack_no_slices(%allocator: !hal.allocator) -> index { - // CHECK-NEXT: %[[ZERO_LENGTH:.+]] = constant 0 + // CHECK-NEXT: %[[ZERO_LENGTH:.+]] = arith.constant 0 %total_length = hal.allocator.pack<%allocator : !hal.allocator> slices({}) : index // CHECK-NEXT: return %[[ZERO_LENGTH]] @@ -96,7 +96,7 @@ // CHECK-LABEL: @allocator_pack_drop_zero_offset func @allocator_pack_drop_zero_offset(%allocator: !hal.allocator, %size : index) -> (index, index, index) { // CHECK-NEXT: = hal.allocator.pack<{{.+}}> slices({ - %base_offset = constant 0 : index + %base_offset = arith.constant 0 : index %total_length, %offset_0, %offset_1 = hal.allocator.pack<%allocator : !hal.allocator> offset(%base_offset) @@ -125,8 +125,8 @@ [0, 4] = %size, [1, 2] = %size, }) : index - // CHECK: %[[ADJUSTED_0:.+]] = addi %[[BASE_OFFSET]], %[[PACKED]]#1 - // CHECK-NEXT: %[[ADJUSTED_1:.+]] = addi %[[BASE_OFFSET]], %[[PACKED]]#2 + // CHECK: %[[ADJUSTED_0:.+]] = arith.addi %[[BASE_OFFSET]], %[[PACKED]]#1 + // CHECK-NEXT: %[[ADJUSTED_1:.+]] = arith.addi %[[BASE_OFFSET]], %[[PACKED]]#2 // CHECK-NEXT: return %[[PACKED]]#0, %[[ADJUSTED_0]], %[[ADJUSTED_1]] return %total_length, %offset_0, %offset_1 : index, index, index }
diff --git a/iree/compiler/Dialect/HAL/IR/test/allocator_ops.mlir b/iree/compiler/Dialect/HAL/IR/test/allocator_ops.mlir index 4d6be0f..0da2669 100644 --- a/iree/compiler/Dialect/HAL/IR/test/allocator_ops.mlir +++ b/iree/compiler/Dialect/HAL/IR/test/allocator_ops.mlir
@@ -2,14 +2,14 @@ // CHECK-LABEL: @allocator_compute_size func @allocator_compute_size(%arg0: !hal.allocator) -> index { - // CHECK-DAG: %[[DIM0:.+]] = constant 100 - %dim0 = constant 100 : index - // CHECK-DAG: %[[DIM1:.+]] = constant 200 - %dim1 = constant 200 : index - // CHECK-DAG: %[[TYPE:.+]] = constant 32 - %type = constant 32 : i32 - // CHECK-DAG: %[[ENCODING:.+]] = constant 1 - %encoding = constant 1 : i32 + // CHECK-DAG: %[[DIM0:.+]] = arith.constant 100 + %dim0 = arith.constant 100 : index + // CHECK-DAG: %[[DIM1:.+]] = arith.constant 200 + %dim1 = arith.constant 200 : index + // CHECK-DAG: %[[TYPE:.+]] = arith.constant 32 + %type = arith.constant 32 : i32 + // CHECK-DAG: %[[ENCODING:.+]] = arith.constant 1 + %encoding = arith.constant 1 : i32 // CHECK: %[[SIZE:.+]] = hal.allocator.compute_size<%arg0 : !hal.allocator> // CHECK-SAME: shape([%[[DIM0]], %[[DIM1]]]) // CHECK-SAME: type(%[[TYPE]]) @@ -26,18 +26,18 @@ // CHECK-LABEL: @allocator_compute_offset func @allocator_compute_offset(%arg0: !hal.allocator) -> index { - // CHECK-DAG: %[[IDX0:.+]] = constant 10 - %idx0 = constant 10 : index - // CHECK-DAG: %[[IDX1:.+]] = constant 20 - %idx1 = constant 20 : index - // CHECK-DAG: %[[DIM0:.+]] = constant 100 - %dim0 = constant 100 : index - // CHECK-DAG: %[[DIM1:.+]] = constant 200 - %dim1 = constant 200 : index - // CHECK-DAG: %[[TYPE:.+]] = constant 32 - %type = constant 32 : i32 - // CHECK-DAG: %[[ENCODING:.+]] = constant 1 - %encoding = constant 1 : i32 + // CHECK-DAG: %[[IDX0:.+]] = arith.constant 10 + %idx0 = arith.constant 10 : index + // CHECK-DAG: %[[IDX1:.+]] = arith.constant 20 + %idx1 = arith.constant 20 : index + // CHECK-DAG: %[[DIM0:.+]] = arith.constant 100 + %dim0 = arith.constant 100 : index + // CHECK-DAG: %[[DIM1:.+]] = arith.constant 200 + %dim1 = arith.constant 200 : index + // CHECK-DAG: %[[TYPE:.+]] = arith.constant 32 + %type = arith.constant 32 : i32 + // CHECK-DAG: %[[ENCODING:.+]] = arith.constant 1 + %encoding = arith.constant 1 : i32 // CHECK: %[[OFFSET:.+]] = hal.allocator.compute_offset<%arg0 : !hal.allocator> // CHECK-SAME: indices([%[[IDX0]], %[[IDX1]]]) // CHECK-SAME: shape([%[[DIM0]], %[[DIM1]]]) @@ -56,22 +56,22 @@ // CHECK-LABEL: @allocator_compute_range func @allocator_compute_range(%arg0: !hal.allocator) -> (index, index) { - // CHECK-DAG: %[[IDX0:.+]] = constant 10 - %idx0 = constant 10 : index - // CHECK-DAG: %[[IDX1:.+]] = constant 20 - %idx1 = constant 20 : index - // CHECK-DAG: %[[LEN0:.+]] = constant 11 - %len0 = constant 11 : index - // CHECK-DAG: %[[LEN1:.+]] = constant 21 - %len1 = constant 21 : index - // CHECK-DAG: %[[DIM0:.+]] = constant 100 - %dim0 = constant 100 : index - // CHECK-DAG: %[[DIM1:.+]] = constant 200 - %dim1 = constant 200 : index - // CHECK-DAG: %[[TYPE:.+]] = constant 32 - %type = constant 32 : i32 - // CHECK-DAG: %[[ENCODING:.+]] = constant 1 - %encoding = constant 1 : i32 + // CHECK-DAG: %[[IDX0:.+]] = arith.constant 10 + %idx0 = arith.constant 10 : index + // CHECK-DAG: %[[IDX1:.+]] = arith.constant 20 + %idx1 = arith.constant 20 : index + // CHECK-DAG: %[[LEN0:.+]] = arith.constant 11 + %len0 = arith.constant 11 : index + // CHECK-DAG: %[[LEN1:.+]] = arith.constant 21 + %len1 = arith.constant 21 : index + // CHECK-DAG: %[[DIM0:.+]] = arith.constant 100 + %dim0 = arith.constant 100 : index + // CHECK-DAG: %[[DIM1:.+]] = arith.constant 200 + %dim1 = arith.constant 200 : index + // CHECK-DAG: %[[TYPE:.+]] = arith.constant 32 + %type = arith.constant 32 : i32 + // CHECK-DAG: %[[ENCODING:.+]] = arith.constant 1 + %encoding = arith.constant 1 : i32 // CHECK: = hal.allocator.compute_range<%arg0 : !hal.allocator> // CHECK-SAME: indices([%[[IDX0]], %[[IDX1]]]) // CHECK-SAME: lengths([%[[LEN0]], %[[LEN1]]]) @@ -92,8 +92,8 @@ // CHECK-LABEL: @allocator_allocate // CHECK-SAME: (%[[ALLOCATOR:.+]]: !hal.allocator) func @allocator_allocate(%allocator: !hal.allocator) { - // CHECK-DAG: %[[SIZE:.+]] = constant 123 - %size = constant 123 : index + // CHECK-DAG: %[[SIZE:.+]] = arith.constant 123 + %size = arith.constant 123 : index // CHECK: %[[REF:.+]] = hal.allocator.allocate<%[[ALLOCATOR]] : !hal.allocator> // CHECK-SAME: type("HostVisible|HostCoherent") // CHECK-SAME: usage(Transfer) @@ -123,10 +123,10 @@ // CHECK-LABEL: @allocator_map_byte_buffer // CHECK-SAME: %[[ALLOCATOR:.+]]: !hal.allocator func @allocator_map_byte_buffer(%allocator: !hal.allocator, %arg1: !util.byte_buffer) { - // CHECK-DAG: %[[OFFSET:.+]] = constant 100 - %offset = constant 100 : index - // CHECK-DAG: %[[LENGTH:.+]] = constant 200 - %length = constant 200 : index + // CHECK-DAG: %[[OFFSET:.+]] = arith.constant 100 + %offset = arith.constant 100 : index + // CHECK-DAG: %[[LENGTH:.+]] = arith.constant 200 + %length = arith.constant 200 : index // CHECK: = hal.allocator.map<%[[ALLOCATOR]] : !hal.allocator> // CHECK-SAME: source(%arg1 : !util.byte_buffer)[%[[OFFSET]], %[[LENGTH]]] // CHECK-SAME: type("DeviceVisible|DeviceLocal") @@ -143,14 +143,14 @@ // CHECK-LABEL: @allocator_pack // CHECK-SAME: %[[ALLOCATOR:.+]]: !hal.allocator func @allocator_pack(%allocator: !hal.allocator) { - // CHECK-DAG: %[[BASE_OFFSET:.+]] = constant 999 - %base_offset = constant 999 : index - // CHECK-DAG: %[[SIZE_0:.+]] = constant 100 - %size_0 = constant 100 : index - // CHECK-DAG: %[[SIZE_1:.+]] = constant 101 - %size_1 = constant 101 : index - // CHECK-DAG: %[[SIZE_2:.+]] = constant 102 - %size_2 = constant 102 : index + // CHECK-DAG: %[[BASE_OFFSET:.+]] = arith.constant 999 + %base_offset = arith.constant 999 : index + // CHECK-DAG: %[[SIZE_0:.+]] = arith.constant 100 + %size_0 = arith.constant 100 : index + // CHECK-DAG: %[[SIZE_1:.+]] = arith.constant 101 + %size_1 = arith.constant 101 : index + // CHECK-DAG: %[[SIZE_2:.+]] = arith.constant 102 + %size_2 = arith.constant 102 : index // CHECK-NEXT: %{{.+}}:4 = %total_length, %offset_0, %offset_1, %offset_2 = // CHECK-SAME: hal.allocator.pack<%[[ALLOCATOR]] : !hal.allocator>
diff --git a/iree/compiler/Dialect/HAL/IR/test/buffer_folding.mlir b/iree/compiler/Dialect/HAL/IR/test/buffer_folding.mlir index f6a749d..ecd59ee 100644 --- a/iree/compiler/Dialect/HAL/IR/test/buffer_folding.mlir +++ b/iree/compiler/Dialect/HAL/IR/test/buffer_folding.mlir
@@ -3,7 +3,7 @@ // CHECK-LABEL: @skip_buffer_allocator // CHECK-SAME: (%[[ALLOCATOR:.+]]: !hal.allocator) func @skip_buffer_allocator(%allocator: !hal.allocator) -> !hal.allocator { - %sz = constant 4 : index + %sz = arith.constant 4 : index %buffer = hal.allocator.allocate<%allocator : !hal.allocator> type("HostVisible|HostCoherent") usage(Transfer) : !hal.buffer{%sz} @@ -17,9 +17,9 @@ // CHECK-LABEL: @skip_subspan_buffer_allocator // CHECK-SAME: (%[[ALLOCATOR:.+]]: !hal.allocator) func @skip_subspan_buffer_allocator(%allocator: !hal.allocator) -> !hal.allocator { - %c0 = constant 0 : index - %c184 = constant 184 : index - %c384 = constant 384 : index + %c0 = arith.constant 0 : index + %c184 = arith.constant 184 : index + %c384 = arith.constant 384 : index %source_buffer = hal.allocator.allocate<%allocator : !hal.allocator> type("HostVisible|HostCoherent") usage(Transfer) : !hal.buffer{%c384}
diff --git a/iree/compiler/Dialect/HAL/IR/test/buffer_ops.mlir b/iree/compiler/Dialect/HAL/IR/test/buffer_ops.mlir index 02b0f5b..2888d16 100644 --- a/iree/compiler/Dialect/HAL/IR/test/buffer_ops.mlir +++ b/iree/compiler/Dialect/HAL/IR/test/buffer_ops.mlir
@@ -13,10 +13,10 @@ // CHECK-LABEL: @buffer_subspan func @buffer_subspan(%arg0: !hal.buffer) -> !hal.buffer { - // CHECK-DAG: %[[OFFSET:.+]] = constant 100 - %offset = constant 100 : index - // CHECK-DAG: %[[LENGTH:.+]] = constant 200 - %length = constant 200 : index + // CHECK-DAG: %[[OFFSET:.+]] = arith.constant 100 + %offset = arith.constant 100 : index + // CHECK-DAG: %[[LENGTH:.+]] = arith.constant 200 + %length = arith.constant 200 : index // CHECK: %buffer = hal.buffer.subspan<%arg0 : !hal.buffer>[%[[OFFSET]], %[[LENGTH]]] : !hal.buffer %buffer = hal.buffer.subspan<%arg0 : !hal.buffer>[%offset, %length] : !hal.buffer return %buffer : !hal.buffer @@ -35,8 +35,8 @@ // CHECK-LABEL: @buffer_load func @buffer_load(%arg0: !hal.buffer) -> i32 { - // CHECK-DAG: %[[SRC_OFFSET:.+]] = constant 100 - %src_offset = constant 100 : index + // CHECK-DAG: %[[SRC_OFFSET:.+]] = arith.constant 100 + %src_offset = arith.constant 100 : index // CHECK: %[[VAL:.+]] = hal.buffer.load<%arg0 : !hal.buffer>[%[[SRC_OFFSET]]] : i32 %1 = hal.buffer.load<%arg0 : !hal.buffer>[%src_offset] : i32 // CHECK-NEXT: return %[[VAL]] @@ -47,8 +47,8 @@ // CHECK-LABEL: @buffer_store func @buffer_store(%arg0: !hal.buffer, %arg1: i32) { - // CHECK-DAG: %[[DST_OFFSET:.+]] = constant 100 - %dst_offset = constant 100 : index + // CHECK-DAG: %[[DST_OFFSET:.+]] = arith.constant 100 + %dst_offset = arith.constant 100 : index // CHECK: hal.buffer.store<%arg0 : !hal.buffer>[%[[DST_OFFSET]]] value(%arg1 : i32) hal.buffer.store<%arg0 : !hal.buffer>[%dst_offset] value(%arg1 : i32) return
diff --git a/iree/compiler/Dialect/HAL/IR/test/buffer_view_folding.mlir b/iree/compiler/Dialect/HAL/IR/test/buffer_view_folding.mlir index 6348da4..a2a01a7 100644 --- a/iree/compiler/Dialect/HAL/IR/test/buffer_view_folding.mlir +++ b/iree/compiler/Dialect/HAL/IR/test/buffer_view_folding.mlir
@@ -3,10 +3,10 @@ // CHECK-LABEL: func @skip_buffer_view_buffer // CHECK-SAME: %[[BUFFER:.+]]: !hal.buffer func @skip_buffer_view_buffer(%buffer : !hal.buffer) -> !hal.buffer { - %c1 = constant 1 : i32 - %c10 = constant 10 : index - %c11 = constant 11 : index - %c32 = constant 32 : i32 + %c1 = arith.constant 1 : i32 + %c10 = arith.constant 10 : index + %c11 = arith.constant 11 : index + %c32 = arith.constant 32 : i32 %view = hal.buffer_view.create buffer(%buffer : !hal.buffer) shape([%c10, %c11]) type(%c32)
diff --git a/iree/compiler/Dialect/HAL/IR/test/buffer_view_ops.mlir b/iree/compiler/Dialect/HAL/IR/test/buffer_view_ops.mlir index a689425..dd36399 100644 --- a/iree/compiler/Dialect/HAL/IR/test/buffer_view_ops.mlir +++ b/iree/compiler/Dialect/HAL/IR/test/buffer_view_ops.mlir
@@ -2,8 +2,8 @@ // CHECK-LABEL: @buffer_view_create func @buffer_view_create(%arg0: !hal.buffer, %arg1: index, %arg2: index) -> !hal.buffer_view { - %c1 = constant 1 : i32 - %c32 = constant 32 : i32 + %c1 = arith.constant 1 : i32 + %c32 = arith.constant 32 : i32 // CHECK: %view = hal.buffer_view.create // CHECK-SAME: buffer(%arg0 : !hal.buffer) // CHECK-SAME: shape([%arg1, %arg2])
diff --git a/iree/compiler/Dialect/HAL/IR/test/command_buffer_folding.mlir b/iree/compiler/Dialect/HAL/IR/test/command_buffer_folding.mlir index 6c64506..58837ab 100644 --- a/iree/compiler/Dialect/HAL/IR/test/command_buffer_folding.mlir +++ b/iree/compiler/Dialect/HAL/IR/test/command_buffer_folding.mlir
@@ -27,11 +27,11 @@ %cmd: !hal.command_buffer, %buffer: !hal.buffer ) { - %c0 = constant 0 : index - %c8192 = constant 8192 : index - %c100000 = constant 100000 : index - %c262144 = constant 262144 : index - %c1234_i32 = constant 1234 : i32 + %c0 = arith.constant 0 : index + %c8192 = arith.constant 8192 : index + %c100000 = arith.constant 100000 : index + %c262144 = arith.constant 262144 : index + %c1234_i32 = arith.constant 1234 : i32 %target_subspan = hal.buffer.subspan<%buffer : !hal.buffer>[%c8192, %c262144] : !hal.buffer // CHECK: hal.command_buffer.fill_buffer hal.command_buffer.fill_buffer<%cmd : !hal.command_buffer> @@ -50,11 +50,11 @@ %cmd: !hal.command_buffer, %buffer: !hal.buffer ) { - %c0 = constant 0 : index - %c4096 = constant 4096 : index - %c8192 = constant 8192 : index - %c100000 = constant 100000 : index - %c262144 = constant 262144 : index + %c0 = arith.constant 0 : index + %c4096 = arith.constant 4096 : index + %c8192 = arith.constant 8192 : index + %c100000 = arith.constant 100000 : index + %c262144 = arith.constant 262144 : index %source_subspan = hal.buffer.subspan<%buffer : !hal.buffer>[%c4096, %c262144] : !hal.buffer %target_subspan = hal.buffer.subspan<%buffer : !hal.buffer>[%c8192, %c262144] : !hal.buffer // CHECK: hal.command_buffer.copy_buffer @@ -78,14 +78,14 @@ %layout: !hal.executable_layout, %buffer: !hal.buffer ) { - %c0 = constant 0 : index - %c1 = constant 1 : index - %c2 = constant 2 : index - %c4 = constant 4 : index - %c4096 = constant 4096 : index - %c8000 = constant 8000 : index - %c262140 = constant 262140 : index - %c262144 = constant 262144 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + %c2 = arith.constant 2 : index + %c4 = arith.constant 4 : index + %c4096 = arith.constant 4096 : index + %c8000 = arith.constant 8000 : index + %c262140 = arith.constant 262140 : index + %c262144 = arith.constant 262144 : index %subspan = hal.buffer.subspan<%buffer : !hal.buffer>[%c4096, %c262144] : !hal.buffer // CHECK: hal.command_buffer.push_descriptor_set // CHECK-SAME: bindings([
diff --git a/iree/compiler/Dialect/HAL/IR/test/command_buffer_ops.mlir b/iree/compiler/Dialect/HAL/IR/test/command_buffer_ops.mlir index 00e57d3..5ec309a 100644 --- a/iree/compiler/Dialect/HAL/IR/test/command_buffer_ops.mlir +++ b/iree/compiler/Dialect/HAL/IR/test/command_buffer_ops.mlir
@@ -112,8 +112,8 @@ %set: !hal.descriptor_set, %offset: index ) { - // CHECK: %[[SET_IDX:.+]] = constant 0 - %c0 = constant 0 : index + // CHECK: %[[SET_IDX:.+]] = arith.constant 0 + %c0 = arith.constant 0 : index // CHECK: hal.command_buffer.bind_descriptor_set<%[[CMD]] : !hal.command_buffer> // CHECK-SAME: layout(%[[LAYOUT]] : !hal.executable_layout)[%[[SET_IDX]]] // CHECK-SAME: set(%[[SET]] : !hal.descriptor_set)
diff --git a/iree/compiler/Dialect/HAL/IR/test/device_ops.mlir b/iree/compiler/Dialect/HAL/IR/test/device_ops.mlir index a1bf7e3..2cb8738 100644 --- a/iree/compiler/Dialect/HAL/IR/test/device_ops.mlir +++ b/iree/compiler/Dialect/HAL/IR/test/device_ops.mlir
@@ -13,12 +13,12 @@ // CHECK-LABEL: @device_switch // CHECK-SAME: (%[[DEVICE:.+]]: !hal.device) func @device_switch(%device: !hal.device) -> i32 { - // CHECK-DAG: %[[C0:.+]] = constant 0 - %c0 = constant 0 : i32 - // CHECK-DAG: %[[C1:.+]] = constant 1 - %c1 = constant 1 : i32 - // CHECK-DAG: %[[C2:.+]] = constant 2 - %c2 = constant 2 : i32 + // CHECK-DAG: %[[C0:.+]] = arith.constant 0 + %c0 = arith.constant 0 : i32 + // CHECK-DAG: %[[C1:.+]] = arith.constant 1 + %c1 = arith.constant 1 : i32 + // CHECK-DAG: %[[C2:.+]] = arith.constant 2 + %c2 = arith.constant 2 : i32 // CHECK: = hal.device.switch<%[[DEVICE]] : !hal.device> -> i32 %0 = hal.device.switch<%device : !hal.device> -> i32 // CHECK-NEXT: #hal.device.match.id<"vulkan-v1.?-*"> {
diff --git a/iree/compiler/Dialect/HAL/IR/test/interface_ops.mlir b/iree/compiler/Dialect/HAL/IR/test/interface_ops.mlir index 194cb24..76f497f 100644 --- a/iree/compiler/Dialect/HAL/IR/test/interface_ops.mlir +++ b/iree/compiler/Dialect/HAL/IR/test/interface_ops.mlir
@@ -16,8 +16,8 @@ // CHECK-LABEL: @interface_io_subspan // CHECK-SAME: (%[[DIM0:.+]]: index, %[[DIM2:.+]]: index) func @interface_io_subspan(%dim0: index, %dim2: index) { - %c8 = constant 8 : index - %c16 = constant 16 : index + %c8 = arith.constant 8 : index + %c16 = arith.constant 16 : index // CHECK: = hal.interface.binding.subspan @interface::@s0b0[%c8] : memref<?x4x?x16xi8>{%[[DIM0]], %[[DIM2]]} %0 = hal.interface.binding.subspan @interface::@s0b0[%c8] : memref<?x4x?x16xi8>{%dim0, %dim2} @@ -31,7 +31,7 @@ // ----- func @interface_io_subspan_wrong_dynamic_dim(%dim: index) { - %c8 = constant 8 : index + %c8 = arith.constant 8 : index // expected-error @+1{{result type 'memref<?x4x?x16xi8>' has 2 dynamic dimensions but 1 associated dimension SSA values}} %0 = hal.interface.binding.subspan @interface::@s0b0[%c8] : memref<?x4x?x16xi8>{%dim}
diff --git a/iree/compiler/Dialect/HAL/IR/test/semaphore_ops.mlir b/iree/compiler/Dialect/HAL/IR/test/semaphore_ops.mlir index a43fc3e..8d40811 100644 --- a/iree/compiler/Dialect/HAL/IR/test/semaphore_ops.mlir +++ b/iree/compiler/Dialect/HAL/IR/test/semaphore_ops.mlir
@@ -4,7 +4,7 @@ // CHECK-LABEL: @semaphore_create func @semaphore_create(%arg0 : !hal.device) -> !hal.semaphore { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index // CHECK: %semaphore = hal.semaphore.create device(%arg0 : !hal.device) initial(%c0) : !hal.semaphore %semaphore = hal.semaphore.create device(%arg0 : !hal.device) initial(%c0) : !hal.semaphore return %semaphore : !hal.semaphore @@ -23,7 +23,7 @@ // CHECK-LABEL: @semaphore_signal func @semaphore_signal(%arg0 : !hal.semaphore) { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index // CHECK: hal.semaphore.signal<%arg0 : !hal.semaphore> value(%c0) hal.semaphore.signal<%arg0 : !hal.semaphore> value(%c0) return @@ -33,8 +33,8 @@ // CHECK-LABEL: @semaphore_fail func @semaphore_fail(%arg0 : !hal.semaphore) { - // CHECK: %[[C0:.+]] = constant 0 - %c0 = constant 0 : i32 + // CHECK: %[[C0:.+]] = arith.constant 0 + %c0 = arith.constant 0 : i32 // CHECK: hal.semaphore.fail<%arg0 : !hal.semaphore> status(%[[C0]]) hal.semaphore.fail<%arg0 : !hal.semaphore> status(%c0) return @@ -44,8 +44,8 @@ // CHECK-LABEL: @semaphore_await func @semaphore_await(%arg0 : !hal.semaphore) { - // CHECK: %[[C0:.+]] = constant 0 - %c0 = constant 0 : index + // CHECK: %[[C0:.+]] = arith.constant 0 + %c0 = arith.constant 0 : index // CHECK: = hal.semaphore.await<%arg0 : !hal.semaphore> until(%[[C0]]) : i32 %0 = hal.semaphore.await<%arg0 : !hal.semaphore> until(%c0) : i32 return
diff --git a/iree/compiler/Dialect/HAL/Target/CUDA/test/smoketest.mlir b/iree/compiler/Dialect/HAL/Target/CUDA/test/smoketest.mlir index 3d2a07b..483e3d8 100644 --- a/iree/compiler/Dialect/HAL/Target/CUDA/test/smoketest.mlir +++ b/iree/compiler/Dialect/HAL/Target/CUDA/test/smoketest.mlir
@@ -24,7 +24,7 @@ %2 = flow.dispatch.tensor.load %arg1, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32> %3 = linalg.generic {indexing_maps = [#map, #map, #map], iterator_types = ["parallel"]} ins(%1, %2 : tensor<16xf32>, tensor<16xf32>) outs(%0 : tensor<16xf32>) { ^bb0(%arg3: f32, %arg4: f32, %arg5: f32): // no predecessors - %4 = addf %arg3, %arg4 : f32 + %4 = arith.addf %arg3, %arg4 : f32 %5 = math.sqrt %4 : f32 linalg.yield %5 : f32 } -> tensor<16xf32>
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/test/smoketest.mlir b/iree/compiler/Dialect/HAL/Target/LLVM/test/smoketest.mlir index 38d9e3c..10b74e4 100644 --- a/iree/compiler/Dialect/HAL/Target/LLVM/test/smoketest.mlir +++ b/iree/compiler/Dialect/HAL/Target/LLVM/test/smoketest.mlir
@@ -23,7 +23,7 @@ %2 = flow.dispatch.tensor.load %arg1, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32> %3 = linalg.generic {indexing_maps = [#map, #map, #map], iterator_types = ["parallel"]} ins(%1, %2 : tensor<16xf32>, tensor<16xf32>) outs(%0 : tensor<16xf32>) { ^bb0(%arg3: f32, %arg4: f32, %arg5: f32): // no predecessors - %4 = addf %arg3, %arg4 : f32 + %4 = arith.addf %arg3, %arg4 : f32 linalg.yield %4 : f32 } -> tensor<16xf32> flow.dispatch.tensor.store %3, %arg2, offsets=[], sizes=[], strides=[] : tensor<16xf32> -> !flow.dispatch.tensor<writeonly:16xf32>
diff --git a/iree/compiler/Dialect/HAL/Target/MetalSPIRV/test/smoketest.mlir b/iree/compiler/Dialect/HAL/Target/MetalSPIRV/test/smoketest.mlir index fd10ea9..269bb35 100644 --- a/iree/compiler/Dialect/HAL/Target/MetalSPIRV/test/smoketest.mlir +++ b/iree/compiler/Dialect/HAL/Target/MetalSPIRV/test/smoketest.mlir
@@ -25,7 +25,7 @@ %2 = flow.dispatch.tensor.load %arg1, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32> %3 = linalg.generic {indexing_maps = [#map, #map, #map], iterator_types = ["parallel"]} ins(%1, %2 : tensor<16xf32>, tensor<16xf32>) outs(%0 : tensor<16xf32>) { ^bb0(%arg3: f32, %arg4: f32, %arg5: f32): // no predecessors - %4 = addf %arg3, %arg4 : f32 + %4 = arith.addf %arg3, %arg4 : f32 linalg.yield %4 : f32 } -> tensor<16xf32> flow.dispatch.tensor.store %3, %arg2, offsets=[], sizes=[], strides=[] : tensor<16xf32> -> !flow.dispatch.tensor<writeonly:16xf32>
diff --git a/iree/compiler/Dialect/HAL/Target/ROCM/test/smoketest.mlir b/iree/compiler/Dialect/HAL/Target/ROCM/test/smoketest.mlir index a8fc4be..0c9d3d7 100644 --- a/iree/compiler/Dialect/HAL/Target/ROCM/test/smoketest.mlir +++ b/iree/compiler/Dialect/HAL/Target/ROCM/test/smoketest.mlir
@@ -24,7 +24,7 @@ %2 = flow.dispatch.tensor.load %arg1, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32> %3 = linalg.generic {indexing_maps = [#map, #map, #map], iterator_types = ["parallel"]} ins(%1, %2 : tensor<16xf32>, tensor<16xf32>) outs(%0 : tensor<16xf32>) { ^bb0(%arg3: f32, %arg4: f32, %arg5: f32): // no predecessors - %4 = addf %arg3, %arg4 : f32 + %4 = arith.addf %arg3, %arg4 : f32 linalg.yield %4 : f32 } -> tensor<16xf32> flow.dispatch.tensor.store %3, %arg2, offsets=[], sizes=[], strides=[] : tensor<16xf32> -> !flow.dispatch.tensor<writeonly:16xf32>
diff --git a/iree/compiler/Dialect/HAL/Target/VMVX/test/linking.mlir b/iree/compiler/Dialect/HAL/Target/VMVX/test/linking.mlir index 846ef60..c6abd8d 100644 --- a/iree/compiler/Dialect/HAL/Target/VMVX/test/linking.mlir +++ b/iree/compiler/Dialect/HAL/Target/VMVX/test/linking.mlir
@@ -60,7 +60,7 @@ func @basic_linking() -> () { %device = hal.ex.shared_device : !hal.device %cmd = hal.command_buffer.create device(%device : !hal.device) mode("OneShot") categories("Transfer|Dispatch") : !hal.command_buffer - %c1 = constant 1 : index + %c1 = arith.constant 1 : index hal.command_buffer.dispatch.symbol<%cmd : !hal.command_buffer> target(@dispatch_0::@vmvx::@dispatch_0) workgroups([%c1, %c1, %c1]) hal.command_buffer.dispatch.symbol<%cmd : !hal.command_buffer> target(@dispatch_1::@vmvx::@dispatch_1) workgroups([%c1, %c1, %c1]) hal.command_buffer.dispatch.symbol<%cmd : !hal.command_buffer> target(@dispatch_2::@vmvx::@dispatch_2) workgroups([%c1, %c1, %c1]) @@ -167,13 +167,13 @@ %cmd = hal.command_buffer.create device(%device : !hal.device) mode("OneShot") categories("Transfer|Dispatch") : !hal.command_buffer hal.device.switch<%device : !hal.device> #hal.device.match.executable.format<"vmvx-bytecode-fb"> { - %c1 = constant 1 : index + %c1 = arith.constant 1 : index hal.command_buffer.dispatch.symbol<%cmd : !hal.command_buffer> target(@dispatch_0::@vmvx::@dispatch_0) workgroups([%c1, %c1, %c1]) hal.command_buffer.dispatch.symbol<%cmd : !hal.command_buffer> target(@dispatch_1::@vmvx::@dispatch_1) workgroups([%c1, %c1, %c1]) hal.return }, #hal.device.match.executable.format<"cuda-nvptx-fb"> { - %c1 = constant 1 : index + %c1 = arith.constant 1 : index hal.command_buffer.dispatch.symbol<%cmd : !hal.command_buffer> target(@dispatch_0::@otherdispatch::@dispatch_0) workgroups([%c1, %c1, %c1]) hal.command_buffer.dispatch.symbol<%cmd : !hal.command_buffer> target(@dispatch_1::@otherdispatch::@dispatch_1) workgroups([%c1, %c1, %c1]) hal.return @@ -222,13 +222,13 @@ // CHECK: func @other_targets() { // CHECK: hal.device.switch<%device : !hal.device> // CHECK-NEXT: #hal.device.match.executable.format<"vmvx-bytecode-fb"> { -// CHECK-NEXT: %c1 = constant 1 : index +// CHECK-NEXT: %c1 = arith.constant 1 : index // CHECK-NEXT: hal.command_buffer.dispatch.symbol<%cmd : !hal.command_buffer> target(@vmvx_linked::@vmvx_bytecode_fb::@dispatch_0) workgroups([%c1, %c1, %c1]) // CHECK-NEXT: hal.command_buffer.dispatch.symbol<%cmd : !hal.command_buffer> target(@vmvx_linked::@vmvx_bytecode_fb::@dispatch_1) workgroups([%c1, %c1, %c1]) // CHECK-NEXT: hal.return // CHECK-NEXT: }, // CHECK-NEXT: #hal.device.match.executable.format<"cuda-nvptx-fb"> { -// CHECK-NEXT: %c1 = constant 1 : index +// CHECK-NEXT: %c1 = arith.constant 1 : index // CHECK-NEXT: hal.command_buffer.dispatch.symbol<%cmd : !hal.command_buffer> target(@dispatch_0::@otherdispatch::@dispatch_0) workgroups([%c1, %c1, %c1]) // CHECK-NEXT: hal.command_buffer.dispatch.symbol<%cmd : !hal.command_buffer> target(@dispatch_1::@otherdispatch::@dispatch_1) workgroups([%c1, %c1, %c1]) // CHECK-NEXT: hal.return
diff --git a/iree/compiler/Dialect/HAL/Target/VMVX/test/smoketest.mlir b/iree/compiler/Dialect/HAL/Target/VMVX/test/smoketest.mlir index 65da95e..cb3afbd 100644 --- a/iree/compiler/Dialect/HAL/Target/VMVX/test/smoketest.mlir +++ b/iree/compiler/Dialect/HAL/Target/VMVX/test/smoketest.mlir
@@ -23,7 +23,7 @@ %2 = flow.dispatch.tensor.load %arg1, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32> %3 = linalg.generic {indexing_maps = [#map, #map, #map], iterator_types = ["parallel"]} ins(%1, %2 : tensor<16xf32>, tensor<16xf32>) outs(%0 : tensor<16xf32>) { ^bb0(%arg3: f32, %arg4: f32, %arg5: f32): // no predecessors - %4 = addf %arg3, %arg4 : f32 + %4 = arith.addf %arg3, %arg4 : f32 linalg.yield %4 : f32 } -> tensor<16xf32> flow.dispatch.tensor.store %3, %arg2, offsets=[], sizes=[], strides=[] : tensor<16xf32> -> !flow.dispatch.tensor<writeonly:16xf32>
diff --git a/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/test/linking.mlir b/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/test/linking.mlir index b94284f..63fcfd2 100644 --- a/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/test/linking.mlir +++ b/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/test/linking.mlir
@@ -82,9 +82,9 @@ hal.executable.variant @vulkan_spirv_fb, target = #executable_target_vulkan_spirv_fb { hal.executable.entry_point @call_dispatch_3 attributes {interface = @io, ordinal = 0 : index} { ^bb0(%arg0: index, %arg1: index, %arg2: index): // no predecessors - %c1 = constant 1 : index - %c56 = constant 56 : index - %c56_0 = constant 56 : index + %c1 = arith.constant 1 : index + %c56 = arith.constant 56 : index + %c56_0 = arith.constant 56 : index hal.return %c1, %c56, %c56_0 : index, index, index } builtin.module {
diff --git a/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/test/smoketest.mlir b/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/test/smoketest.mlir index a02179b..d3d6598 100644 --- a/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/test/smoketest.mlir +++ b/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/test/smoketest.mlir
@@ -25,7 +25,7 @@ %2 = flow.dispatch.tensor.load %arg1, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32> %3 = linalg.generic {indexing_maps = [#map, #map, #map], iterator_types = ["parallel"]} ins(%1, %2 : tensor<16xf32>, tensor<16xf32>) outs(%0 : tensor<16xf32>) { ^bb0(%arg3: f32, %arg4: f32, %arg5: f32): // no predecessors - %4 = addf %arg3, %arg4 : f32 + %4 = arith.addf %arg3, %arg4 : f32 linalg.yield %4 : f32 } -> tensor<16xf32> flow.dispatch.tensor.store %3, %arg2, offsets=[], sizes=[], strides=[] : tensor<16xf32> -> !flow.dispatch.tensor<writeonly:16xf32>
diff --git a/iree/compiler/Dialect/HAL/Transforms/BenchmarkBatchDispatches.cpp b/iree/compiler/Dialect/HAL/Transforms/BenchmarkBatchDispatches.cpp index 3bbf27d..d1458ed 100644 --- a/iree/compiler/Dialect/HAL/Transforms/BenchmarkBatchDispatches.cpp +++ b/iree/compiler/Dialect/HAL/Transforms/BenchmarkBatchDispatches.cpp
@@ -24,7 +24,8 @@ : repeatCount_(repeatCount) {} void getDependentDialects(DialectRegistry& registry) const override { - registry.insert<HALDialect, StandardOpsDialect>(); + registry.insert<HALDialect, StandardOpsDialect, + mlir::arith::ArithmeticDialect>(); } StringRef getArgument() const override {
diff --git a/iree/compiler/Dialect/HAL/Transforms/ConvertToHAL.cpp b/iree/compiler/Dialect/HAL/Transforms/ConvertToHAL.cpp index 75ea007..54810d7 100644 --- a/iree/compiler/Dialect/HAL/Transforms/ConvertToHAL.cpp +++ b/iree/compiler/Dialect/HAL/Transforms/ConvertToHAL.cpp
@@ -43,7 +43,7 @@ void getDependentDialects(DialectRegistry ®istry) const override { registry.insert<IREE::Util::UtilDialect>(); registry.insert<HALDialect>(); - registry.insert<StandardOpsDialect>(); + registry.insert<StandardOpsDialect, mlir::arith::ArithmeticDialect>(); } StringRef getArgument() const override { return "iree-convert-to-hal"; }
diff --git a/iree/compiler/Dialect/HAL/Transforms/IdentifyConstantPools.cpp b/iree/compiler/Dialect/HAL/Transforms/IdentifyConstantPools.cpp index fd0b92f..e439db4 100644 --- a/iree/compiler/Dialect/HAL/Transforms/IdentifyConstantPools.cpp +++ b/iree/compiler/Dialect/HAL/Transforms/IdentifyConstantPools.cpp
@@ -32,7 +32,7 @@ IdentifyConstantPoolsPass() = default; void getDependentDialects(DialectRegistry ®istry) const override { - registry.insert<mlir::StandardOpsDialect>(); + registry.insert<mlir::StandardOpsDialect, mlir::arith::ArithmeticDialect>(); registry.insert<IREE::Flow::FlowDialect>(); registry.insert<IREE::HAL::HALDialect>(); registry.insert<IREE::Util::UtilDialect>();
diff --git a/iree/compiler/Dialect/HAL/Transforms/MaterializeConstantPoolBuffers.cpp b/iree/compiler/Dialect/HAL/Transforms/MaterializeConstantPoolBuffers.cpp index 044bb24..9493480 100644 --- a/iree/compiler/Dialect/HAL/Transforms/MaterializeConstantPoolBuffers.cpp +++ b/iree/compiler/Dialect/HAL/Transforms/MaterializeConstantPoolBuffers.cpp
@@ -28,7 +28,7 @@ OperationPass<ModuleOp>> { public: void getDependentDialects(DialectRegistry ®istry) const override { - registry.insert<mlir::StandardOpsDialect>(); + registry.insert<mlir::StandardOpsDialect, mlir::arith::ArithmeticDialect>(); registry.insert<IREE::Util::UtilDialect>(); registry.insert<IREE::HAL::HALDialect>(); } @@ -136,14 +136,15 @@ builder.getContext(), storageOp->getParentOfType<ConstantPoolOp>().getName(), {SymbolRefAttr::get(storageOp)})); - auto offsetValue = builder.createOrFold<mlir::ConstantIndexOp>(loc, 0); + auto offsetValue = + builder.createOrFold<mlir::arith::ConstantIndexOp>(loc, 0); auto storageValueAttr = storageOp.value().cast<IREE::Util::SerializableAttrInterface>(); uint64_t runtimeLength = align(storageValueAttr.getStorageSize(), bufferConstraints.min_buffer_range_alignment()); auto lengthValue = - builder.createOrFold<mlir::ConstantIndexOp>(loc, runtimeLength); + builder.createOrFold<mlir::arith::ConstantIndexOp>(loc, runtimeLength); auto memoryType = IREE::HAL::MemoryTypeBitfield::DeviceLocal | IREE::HAL::MemoryTypeBitfield::HostVisible; auto bufferUsage = IREE::HAL::BufferUsageBitfield::Constant | @@ -241,7 +242,7 @@ auto bufferUsage = IREE::HAL::BufferUsageBitfield::Constant | IREE::HAL::BufferUsageBitfield::All; auto allocationSizeValue = - builder.createOrFold<mlir::ConstantIndexOp>(loc, bufferLength); + builder.createOrFold<mlir::arith::ConstantIndexOp>(loc, bufferLength); auto bufferValue = builder.createOrFold<IREE::HAL::AllocatorAllocateOp>( loc, IREE::HAL::BufferType::get(context), allocatorValue, memoryType, bufferUsage, allocationSizeValue); @@ -258,13 +259,13 @@ builder.create<IREE::HAL::CommandBufferBeginOp>(loc, commandBufferValue); for (auto splatOp : splatOps) { auto runtimeRange = splatOp.runtime_range().getValue(); - auto offsetValue = builder.createOrFold<mlir::ConstantIndexOp>( + auto offsetValue = builder.createOrFold<mlir::arith::ConstantIndexOp>( splatOp.getLoc(), runtimeRange.getOffset()); - auto lengthValue = builder.createOrFold<mlir::ConstantIndexOp>( + auto lengthValue = builder.createOrFold<mlir::arith::ConstantIndexOp>( splatOp.getLoc(), runtimeRange.getLength()); uint32_t pattern = makePatternFromSplatValue( splatOp.value().cast<SplatElementsAttr>().getSplatValue()); - auto patternValue = builder.createOrFold<mlir::ConstantIntOp>( + auto patternValue = builder.createOrFold<mlir::arith::ConstantIntOp>( loc, static_cast<int64_t>(pattern), 32); builder.create<IREE::HAL::CommandBufferFillBufferOp>( splatOp.getLoc(), commandBufferValue, bufferValue, offsetValue,
diff --git a/iree/compiler/Dialect/HAL/Transforms/MaterializeInterfaces.cpp b/iree/compiler/Dialect/HAL/Transforms/MaterializeInterfaces.cpp index ac5a5d6..e307167 100644 --- a/iree/compiler/Dialect/HAL/Transforms/MaterializeInterfaces.cpp +++ b/iree/compiler/Dialect/HAL/Transforms/MaterializeInterfaces.cpp
@@ -567,7 +567,7 @@ OpBuilder entryBuilder = OpBuilder::atBlockBegin(entryBlock); // +0 offset is used a lot. - auto zeroOffset = entryBuilder.createOrFold<mlir::ConstantIndexOp>( + auto zeroOffset = entryBuilder.createOrFold<mlir::arith::ConstantIndexOp>( clonedFuncOp.getLoc(), 0); for (auto regionOperand : llvm::enumerate(regionOperands)) { @@ -587,7 +587,7 @@ Value offset = zeroOffset; if (value.bindingOffset.staticOffset != BindingOffset::INVALID_OFFSET && value.bindingOffset.staticOffset != 0) { - offset = entryBuilder.createOrFold<mlir::ConstantIndexOp>( + offset = entryBuilder.createOrFold<mlir::arith::ConstantIndexOp>( clonedFuncOp.getLoc(), value.bindingOffset.staticOffset); }
diff --git a/iree/compiler/Dialect/HAL/Transforms/PackAllocations.cpp b/iree/compiler/Dialect/HAL/Transforms/PackAllocations.cpp index b8b896e..c1d2f43 100644 --- a/iree/compiler/Dialect/HAL/Transforms/PackAllocations.cpp +++ b/iree/compiler/Dialect/HAL/Transforms/PackAllocations.cpp
@@ -33,7 +33,7 @@ PackAllocationsPass() = default; void getDependentDialects(DialectRegistry ®istry) const override { - registry.insert<mlir::StandardOpsDialect>(); + registry.insert<mlir::StandardOpsDialect, mlir::arith::ArithmeticDialect>(); registry.insert<IREE::HAL::HALDialect>(); } @@ -64,7 +64,8 @@ staticSlices.reserve(allSlices.size()); dynamicSlices.reserve(allSlices.size()); for (auto &slice : allSlices) { - if (isa_and_nonnull<ConstantOp>(slice.dynamicSize.getDefiningOp())) { + if (isa_and_nonnull<arith::ConstantOp>( + slice.dynamicSize.getDefiningOp())) { staticSlices.push_back(slice); } else { dynamicSlices.push_back(slice); @@ -75,9 +76,10 @@ // First pack all static slices as these are entirely knowable here at // compile time. - auto offset = packOp.offset() ? packOp.offset() - : builder.createOrFold<ConstantIndexOp>( - packOp.getLoc(), 0); + auto offset = packOp.offset() + ? packOp.offset() + : builder.createOrFold<arith::ConstantIndexOp>( + packOp.getLoc(), 0); if (!staticSlices.empty()) { offset = packStaticSlicesGreedily(packOp, offset, staticSlices, bufferConstraints, builder); @@ -125,8 +127,9 @@ for (auto &slice : slices) { auto sliceSize = align(loc, slice.dynamicSize, rangeAlignment, builder); slice.packedOffset.replaceAllUsesWith(offset); - offset = align(loc, builder.createOrFold<AddIOp>(loc, offset, sliceSize), - offsetAlignment, builder); + offset = align( + loc, builder.createOrFold<arith::AddIOp>(loc, offset, sliceSize), + offsetAlignment, builder); } return align(loc, offset, rangeAlignment, builder); @@ -172,8 +175,8 @@ int64_t bestOffset = UNASSIGNED; int64_t bestOffsetFit = UNASSIGNED; int64_t staticSize = - dyn_cast<ConstantIndexOp>(slice.dynamicSize.getDefiningOp()) - .getValue(); + dyn_cast<arith::ConstantIndexOp>(slice.dynamicSize.getDefiningOp()) + .value(); int64_t alignedSize = align(staticSize, rangeAlignment); // Iterate through reservations (sorted by ascending offset) and identify @@ -213,9 +216,10 @@ ++insertionIt; } reservations.insert(insertionIt, reservation); - slice.packedOffset.replaceAllUsesWith(builder.createOrFold<AddIOp>( + slice.packedOffset.replaceAllUsesWith(builder.createOrFold<arith::AddIOp>( packOp.getLoc(), baseOffset, - builder.createOrFold<ConstantIndexOp>(packOp.getLoc(), bestOffset))); + builder.createOrFold<arith::ConstantIndexOp>(packOp.getLoc(), + bestOffset))); // Update highwater mark indicating how much memory needs to be allocated // for the entire slab. @@ -223,9 +227,10 @@ } highwaterMark = align(highwaterMark, rangeAlignment); - return builder.createOrFold<AddIOp>( + return builder.createOrFold<arith::AddIOp>( packOp.getLoc(), baseOffset, - builder.createOrFold<ConstantIndexOp>(packOp.getLoc(), highwaterMark)); + builder.createOrFold<arith::ConstantIndexOp>(packOp.getLoc(), + highwaterMark)); } // Packs a set of dynamically-sized slices based on the structural information @@ -303,9 +308,9 @@ // Allocate a new bin for this slice. bins.push_back({offset, {}}); targetBin = &bins.back(); - offset = - align(loc, builder.createOrFold<AddIOp>(loc, offset, sliceSize), - offsetAlignment, builder); + offset = align( + loc, builder.createOrFold<arith::AddIOp>(loc, offset, sliceSize), + offsetAlignment, builder); } targetBin->slices.push_back(slice); slice->packedOffset.replaceAllUsesWith(targetBin->offset);
diff --git a/iree/compiler/Dialect/HAL/Transforms/PropagateConstantWorkgroupInfo.cpp b/iree/compiler/Dialect/HAL/Transforms/PropagateConstantWorkgroupInfo.cpp index f753f68..08ffabf 100644 --- a/iree/compiler/Dialect/HAL/Transforms/PropagateConstantWorkgroupInfo.cpp +++ b/iree/compiler/Dialect/HAL/Transforms/PropagateConstantWorkgroupInfo.cpp
@@ -41,7 +41,7 @@ funcOp.getOps<IREE::HAL::InterfaceWorkgroupSizeOp>()); for (auto workgroupSizeOp : workgroupSizeOps) { OpBuilder builder(workgroupSizeOp); - auto dimValue = builder.createOrFold<ConstantIndexOp>( + auto dimValue = builder.createOrFold<arith::ConstantIndexOp>( workgroupSizeOp.getLoc(), workgroupSizeAttr[workgroupSizeOp.dimension().getZExtValue()] .cast<IntegerAttr>()
diff --git a/iree/compiler/Dialect/HAL/Transforms/test/benchmark_batch_dispatches.mlir b/iree/compiler/Dialect/HAL/Transforms/test/benchmark_batch_dispatches.mlir index 902d989..ccbadeb 100644 --- a/iree/compiler/Dialect/HAL/Transforms/test/benchmark_batch_dispatches.mlir +++ b/iree/compiler/Dialect/HAL/Transforms/test/benchmark_batch_dispatches.mlir
@@ -9,8 +9,8 @@ // CHECK: %[[EXE:.+]] = util.global.load @_executable %exe = util.global.load @_executable : !hal.executable - %c1 = constant 1 : index - %c2 = constant 2 : index + %c1 = arith.constant 1 : index + %c2 = arith.constant 2 : index hal.command_buffer.dispatch<%cmd1 : !hal.command_buffer> target(%exe : !hal.executable)[0] workgroups([%c1, %c1, %c1]) hal.command_buffer.execution_barrier<%cmd1 : !hal.command_buffer> source("Dispatch|CommandRetire") target("CommandIssue|Dispatch") flags("None") hal.command_buffer.dispatch<%cmd1 : !hal.command_buffer> target(%exe : !hal.executable)[1] workgroups([%c2, %c2, %c2])
diff --git a/iree/compiler/Dialect/HAL/Transforms/test/inline_device_switches.mlir b/iree/compiler/Dialect/HAL/Transforms/test/inline_device_switches.mlir index 0073e26..c5a1319 100644 --- a/iree/compiler/Dialect/HAL/Transforms/test/inline_device_switches.mlir +++ b/iree/compiler/Dialect/HAL/Transforms/test/inline_device_switches.mlir
@@ -4,14 +4,14 @@ // CHECK-SAME: %[[DEVICE:.+]]: !hal.device // CHECK-SAME: %[[ARG:.+]]: i32 func @simple_constants(%device : !hal.device, %arg : i32) -> i32 { - // CHECK-DAG: %[[C0:.+]] = constant 0 - %c0 = constant 0 : i32 - // CHECK-DAG: %[[C1:.+]] = constant 1 - %c1 = constant 1 : i32 - // CHECK-DAG: %[[C2:.+]] = constant 2 - %c2 = constant 2 : i32 - // CHECK-DAG: %[[C3:.+]] = constant 3 - // CHECK-DAG: %[[C4:.+]] = constant 4 + // CHECK-DAG: %[[C0:.+]] = arith.constant 0 + %c0 = arith.constant 0 : i32 + // CHECK-DAG: %[[C1:.+]] = arith.constant 1 + %c1 = arith.constant 1 : i32 + // CHECK-DAG: %[[C2:.+]] = arith.constant 2 + %c2 = arith.constant 2 : i32 + // CHECK-DAG: %[[C3:.+]] = arith.constant 3 + // CHECK-DAG: %[[C4:.+]] = arith.constant 4 %0 = hal.device.switch<%device : !hal.device> -> i32 // CHECK-NEXT: %{{.+}}, %[[IS0:.+]] = hal.device.query<%[[DEVICE]] : !hal.device> key("hal.device.id" :: "vulkan-v1.?-*") : i1, i1 = false // CHECK-NEXT: cond_br %[[IS0]], ^bb3(%[[C1]] : i32), ^bb1 @@ -21,19 +21,19 @@ // CHECK-NEXT: ^bb1: // CHECK-NEXT: %{{.+}}, %[[IS1L:.+]] = hal.device.query<%[[DEVICE]] : !hal.device> key("hal.device.id" :: "vmvx") : i1, i1 = false // CHECK-NEXT: %{{.+}}, %[[IS1R:.+]] = hal.device.query<%[[DEVICE]] : !hal.device> key("hal.device.id" :: "vulkan-*") : i1, i1 = false - // CHECK-NEXT: %[[IS1:.+]] = or %[[IS1L]], %[[IS1R]] : i1 + // CHECK-NEXT: %[[IS1:.+]] = arith.ori %[[IS1L]], %[[IS1R]] : i1 // CHECK-NEXT: cond_br %[[IS1]], ^bb2, ^bb3(%[[C0]] : i32) // CHECK-NEXT: ^bb2: - // CHECK-NEXT: %[[EQZ:.+]] = cmpi eq, %[[ARG]], %[[C2]] : i32 + // CHECK-NEXT: %[[EQZ:.+]] = arith.cmpi eq, %[[ARG]], %[[C2]] : i32 // CHECK-NEXT: cond_br %[[EQZ]], ^bb3(%[[C3]] : i32), ^bb3(%[[C4]] : i32) #hal.match.any<[#hal.device.match.id<"vmvx">, #hal.device.match.id<"vulkan-*">]> { - %eqz = cmpi eq, %arg, %c2 : i32 + %eqz = arith.cmpi eq, %arg, %c2 : i32 cond_br %eqz, ^bb_true, ^bb_false ^bb_true: - %c3 = constant 3 : i32 + %c3 = arith.constant 3 : i32 hal.return %c3 : i32 ^bb_false: - %c4 = constant 4 : i32 + %c4 = arith.constant 4 : i32 hal.return %c4 : i32 }, #hal.match.always { @@ -62,7 +62,7 @@ // CHECK-NEXT: ^bb2: // CHECK-NEXT: %{{.+}}, %[[IS1L:.+]] = hal.device.query<%[[DEVICE]] : !hal.device> key("hal.device.id" :: "vmvx") : i1, i1 = false // CHECK-NEXT: %{{.+}}, %[[IS1R:.+]] = hal.device.query<%[[DEVICE]] : !hal.device> key("hal.device.id" :: "vulkan-*") : i1, i1 = false - // CHECK-NEXT: %[[IS1:.+]] = or %[[IS1L]], %[[IS1R]] : i1 + // CHECK-NEXT: %[[IS1:.+]] = arith.ori %[[IS1L]], %[[IS1R]] : i1 // CHECK-NEXT: cond_br %[[IS1]], ^bb3, ^bb4 // CHECK-NEXT: ^bb3: // CHECK-NEXT: "some.op_b"()
diff --git a/iree/compiler/Dialect/HAL/Transforms/test/materialize_interfaces.mlir b/iree/compiler/Dialect/HAL/Transforms/test/materialize_interfaces.mlir index 6f04508..e84f144 100644 --- a/iree/compiler/Dialect/HAL/Transforms/test/materialize_interfaces.mlir +++ b/iree/compiler/Dialect/HAL/Transforms/test/materialize_interfaces.mlir
@@ -22,7 +22,7 @@ builtin.module { // CHECK-NEXT: func @entry() { func @entry(%arg: !flow.dispatch.tensor<readonly:8x4xf32>, %ret: !flow.dispatch.tensor<writeonly:4x8xf32>) { - // CHECK-NEXT: %c0 = constant 0 : index + // CHECK-NEXT: %c0 = arith.constant 0 : index // CHECK-NEXT: %[[ARG:.+]] = hal.interface.binding.subspan @[[IO]]::@[[S0B0]][%c0] : !flow.dispatch.tensor<readonly:8x4xf32> // CHECK-NEXT: %[[RET:.+]] = hal.interface.binding.subspan @[[IO]]::@[[S0B1]][%c0] : !flow.dispatch.tensor<writeonly:4x8xf32> @@ -39,7 +39,7 @@ func @usage(%func_arg: tensor<8x4xf32>) -> tensor<4x8xf32> { %0 = flow.ex.stream.fragment(%func_arg) : (tensor<8x4xf32>) -> tensor<4x8xf32> = (%stream_arg: tensor<8x4xf32>) -> tensor<4x8xf32> { - %c1 = constant 1 : index + %c1 = arith.constant 1 : index // CHECK: = flow.dispatch @static_tiled_dispatch::@entry // CHECK-SAME: hal.bindings = [ // CHECK-SAME: #hal.ex.operand_buffer<"[[S0B0]]", 0 : index>, @@ -82,7 +82,7 @@ builtin.module { // CHECK-NEXT: func @entry() { func @entry(%arg: !flow.dispatch.tensor<readonly:8x4xf32>, %ret: !flow.dispatch.tensor<writeonly:4x8xf32>) { - // CHECK-NEXT: %c0 = constant 0 : index + // CHECK-NEXT: %c0 = arith.constant 0 : index // CHECK-NEXT: %[[ARG:.+]] = hal.interface.binding.subspan @[[IO]]::@[[S0B0]][%c0] : !flow.dispatch.tensor<readonly:8x4xf32> // CHECK-NEXT: %[[RET:.+]] = hal.interface.binding.subspan @[[IO]]::@[[S0B1]][%c0] : !flow.dispatch.tensor<writeonly:4x8xf32> @@ -104,7 +104,7 @@ func @usage(%func_arg: tensor<8x4xf32>) -> tensor<4x8xf32> { %0 = flow.ex.stream.fragment(%func_arg) : (tensor<8x4xf32>) -> tensor<4x8xf32> = (%stream_arg: tensor<8x4xf32>) -> tensor<4x8xf32> { - %c1 = constant 1 : index + %c1 = arith.constant 1 : index // CHECK: = flow.dispatch @multi_target_ex::@entry // CHECK-SAME: hal.bindings = [ // CHECK-SAME: #hal.ex.operand_buffer<"[[S0B0]]", 0 : index>, @@ -141,7 +141,7 @@ builtin.module { // CHECK-NEXT: func @entry() { func @entry( - // CHECK-NEXT: %c0 = constant 0 : index + // CHECK-NEXT: %c0 = arith.constant 0 : index %arg: !flow.dispatch.tensor<readonly:7x?x24x?xf32>, %ret: !flow.dispatch.tensor<writeonly:?x?x1024xf32>, // CHECK-DAG: %[[ARG_DIM1:.+]] = hal.interface.load.constant offset = 0 : index @@ -171,11 +171,11 @@ } } func @usage(%func_arg: tensor<7x?x24x?xf32>) -> tensor<?x?x1024xf32> { - %d0 = constant 100 : index - %d1 = constant 200 : index + %d0 = arith.constant 100 : index + %d1 = arith.constant 200 : index %0 = flow.ex.stream.fragment(%func_arg, %d0, %d1) : (tensor<7x?x24x?xf32>{%d0, %d1}, index, index) -> tensor<?x?x1024xf32>{%d1, %d0} = (%stream_arg: tensor<7x?x24x?xf32>, %stream_d0: index, %stream_d1: index) -> tensor<?x?x1024xf32> { - %c1 = constant 1 : index + %c1 = arith.constant 1 : index // CHECK: = flow.dispatch @dynamic_tiled_dispatch::@entry // CHECK-SAME: hal.bindings = [ // CHECK-SAME: #hal.ex.push_constant<0 : index, 1 : index>, @@ -258,7 +258,7 @@ builtin.module { // CHECK: func @entry() { func @entry(%arg: !flow.dispatch.tensor<readonly:8x4xf32>, %ret: !flow.dispatch.tensor<readwrite:4x8xf32>) { - // CHECK-NEXT: %c0 = constant 0 : index + // CHECK-NEXT: %c0 = arith.constant 0 : index // CHECK-NEXT: %[[ARG:.+]] = hal.interface.binding.subspan @[[IO]]::@[[S0B0]][%c0] : !flow.dispatch.tensor<readonly:8x4xf32> // CHECK-NEXT: %[[RET:.+]] = hal.interface.binding.subspan @[[IO]]::@[[S0B1]][%c0] : !flow.dispatch.tensor<readwrite:4x8xf32> // CHECK-NEXT: %[[ARG0_TILE:.+]] = flow.dispatch.tensor.load %[[ARG]] @@ -276,7 +276,7 @@ func @usage(%func_arg: tensor<8x4xf32>, %func_ret: tensor<4x8xf32>) -> tensor<4x8xf32> { %0 = flow.ex.stream.fragment(%func_arg, %func_ret) : (tensor<8x4xf32>, tensor<4x8xf32>) -> %func_ret = (%stream_arg: tensor<8x4xf32>, %stream_ret: tensor<4x8xf32>) -> tensor<4x8xf32> { - %c1 = constant 1 : index + %c1 = arith.constant 1 : index // CHECK: = flow.dispatch @static_tied_result::@entry // CHECK-SAME: hal.bindings = [ // CHECK-SAME: #hal.ex.operand_buffer<"[[S0B0]]", 0 : index>, @@ -355,7 +355,7 @@ %const_span_0b = hal.constant.subspan @storage0[#util.byte_range<128, 128>] : tensor<8x4xf32> %const_span_0c = hal.constant.subspan @storage0[#util.byte_range<256, 128>] : tensor<8x4xf32> %const_span_1a = hal.constant.subspan @storage1[#util.byte_range<0, 128>] : tensor<8x4xf32> - %c1 = constant 1 : index + %c1 = arith.constant 1 : index // CHECK: = flow.dispatch @constant_dispatch::@entry // CHECK-SAME: hal.bindings = [ // CHECK-SAME: #hal.ex.constant_storage<"[[S0B0]]", "storage0", 0 : index, 256 : index>, @@ -399,7 +399,7 @@ flow.dispatch.entry @entry attributes {workgroup_rank = 2 : index} builtin.module { func @entry(%unused_arg: !flow.dispatch.tensor<readonly:8x4xf32>, %ret: !flow.dispatch.tensor<writeonly:4x8xf32>) { - %val = constant dense<4.2> : tensor<4x8xf32> + %val = arith.constant dense<4.2> : tensor<4x8xf32> // CHECK-NOT: hal.interface.binding.subspan @[[IO]]::@[[S0B0]] // CHECK: hal.interface.binding.subspan @[[IO]]::@[[S0B1]] flow.dispatch.tensor.store %val, %ret, offsets=[], sizes=[], strides=[] : tensor<4x8xf32> -> !flow.dispatch.tensor<writeonly:4x8xf32> @@ -410,7 +410,7 @@ func @usage(%func_arg: tensor<8x4xf32>) -> tensor<4x8xf32> { %0 = flow.ex.stream.fragment(%func_arg) : (tensor<8x4xf32>) -> tensor<4x8xf32> = (%stream_arg: tensor<8x4xf32>) -> tensor<4x8xf32> { - %c1 = constant 1 : index + %c1 = arith.constant 1 : index %1 = flow.dispatch @unsued_arg::@entry[%c1, %c1, %c1](%stream_arg) : (tensor<8x4xf32>) -> tensor<4x8xf32> flow.return %1 : tensor<4x8xf32> }
diff --git a/iree/compiler/Dialect/HAL/Transforms/test/pack_allocations.mlir b/iree/compiler/Dialect/HAL/Transforms/test/pack_allocations.mlir index 4b761b8..4eeac27 100644 --- a/iree/compiler/Dialect/HAL/Transforms/test/pack_allocations.mlir +++ b/iree/compiler/Dialect/HAL/Transforms/test/pack_allocations.mlir
@@ -12,8 +12,8 @@ // CHECK-SAME: %[[ALLOCATOR:.+]]: !hal.allocator func @packStatic(%allocator: !hal.allocator) -> (index, index, index, index, index, index, index) { - %c100 = constant 100 : index - %c200 = constant 200 : index + %c100 = arith.constant 100 : index + %c200 = arith.constant 200 : index %t:7 = hal.allocator.pack<%allocator : !hal.allocator> slices({ [0, 1] = %c100, // +0 [1, 2] = %c100, // +112 (100 align 16) @@ -56,20 +56,20 @@ // Right now this is too verbose to really test against with anything but // a change detector like this. - // CHECK-DAG: %c0 = constant 0 : index - // CHECK-DAG: %c-16 = constant -16 : index - // CHECK-DAG: %c15 = constant 15 : index - // CHECK-DAG: %0 = addi %arg1, %c15 : index - // CHECK-DAG: %1 = and %0, %c-16 : index - // CHECK-DAG: %2 = addi %1, %c15 : index - // CHECK-DAG: %3 = and %2, %c-16 : index - // CHECK-DAG: %4 = addi %arg2, %c15 : index - // CHECK-DAG: %5 = and %4, %c-16 : index - // CHECK-DAG: %6 = addi %3, %5 : index - // CHECK-DAG: %7 = addi %6, %c15 : index - // CHECK-DAG: %8 = and %7, %c-16 : index - // CHECK-DAG: %9 = addi %8, %c15 : index - // CHECK-DAG: %10 = and %9, %c-16 : index + // CHECK-DAG: %c0 = arith.constant 0 : index + // CHECK-DAG: %c-16 = arith.constant -16 : index + // CHECK-DAG: %c15 = arith.constant 15 : index + // CHECK-DAG: %0 = arith.addi %arg1, %c15 : index + // CHECK-DAG: %1 = arith.andi %0, %c-16 : index + // CHECK-DAG: %2 = arith.addi %1, %c15 : index + // CHECK-DAG: %3 = arith.andi %2, %c-16 : index + // CHECK-DAG: %4 = arith.addi %arg2, %c15 : index + // CHECK-DAG: %5 = arith.andi %4, %c-16 : index + // CHECK-DAG: %6 = arith.addi %3, %5 : index + // CHECK-DAG: %7 = arith.addi %6, %c15 : index + // CHECK-DAG: %8 = arith.andi %7, %c-16 : index + // CHECK-DAG: %9 = arith.addi %8, %c15 : index + // CHECK-DAG: %10 = arith.andi %9, %c-16 : index // CHECK-DAG: return %10, %c0, %3, %c0 return %t#0, %t#1, %t#2, %t#3 : index, index, index, index @@ -92,8 +92,8 @@ // CHECK-SAME: %[[SIZE_A:.+]]: index, %[[SIZE_B:.+]]: index func @packMixedStaticDynamic(%allocator: !hal.allocator, %size_a: index, %size_b: index) -> (index, index, index, index, index) { - %c100 = constant 100 : index - %c200 = constant 200 : index + %c100 = arith.constant 100 : index + %c200 = arith.constant 200 : index %t:5 = hal.allocator.pack<%allocator : !hal.allocator> slices({ [0, 1] = %c100, [1, 2] = %size_a, @@ -106,21 +106,21 @@ // Right now this is too verbose to really test against with anything but // a change detector like this. - // CHECK-DAG: %c0 = constant 0 : index - // CHECK-DAG: %c208 = constant 208 : index - // CHECK-DAG: %c-16 = constant -16 : index - // CHECK-DAG: %c15 = constant 15 : index - // CHECK-DAG: %0 = addi %arg1, %c15 : index - // CHECK-DAG: %1 = and %0, %c-16 : index - // CHECK-DAG: %2 = addi %1, %c223 : index - // CHECK-DAG: %3 = and %2, %c-16 : index - // CHECK-DAG: %4 = addi %arg2, %c15 : index - // CHECK-DAG: %5 = and %4, %c-16 : index - // CHECK-DAG: %6 = addi %3, %5 : index - // CHECK-DAG: %7 = addi %6, %c15 : index - // CHECK-DAG: %8 = and %7, %c-16 : index - // CHECK-DAG: %9 = addi %8, %c15 : index - // CHECK-DAG: %10 = and %9, %c-16 : index + // CHECK-DAG: %c0 = arith.constant 0 : index + // CHECK-DAG: %c208 = arith.constant 208 : index + // CHECK-DAG: %c-16 = arith.constant -16 : index + // CHECK-DAG: %c15 = arith.constant 15 : index + // CHECK-DAG: %0 = arith.addi %arg1, %c15 : index + // CHECK-DAG: %1 = arith.andi %0, %c-16 : index + // CHECK-DAG: %2 = arith.addi %1, %c223 : index + // CHECK-DAG: %3 = arith.andi %2, %c-16 : index + // CHECK-DAG: %4 = arith.addi %arg2, %c15 : index + // CHECK-DAG: %5 = arith.andi %4, %c-16 : index + // CHECK-DAG: %6 = arith.addi %3, %5 : index + // CHECK-DAG: %7 = arith.addi %6, %c15 : index + // CHECK-DAG: %8 = arith.andi %7, %c-16 : index + // CHECK-DAG: %9 = arith.addi %8, %c15 : index + // CHECK-DAG: %10 = arith.andi %9, %c-16 : index // CHECK-DAG: return %10, %c0, %c208, %3, %c0 return %t#0, %t#1, %t#2, %t#3, %t#4 : index, index, index, index, index
diff --git a/iree/compiler/Dialect/HAL/Transforms/test/resolve_entry_point_ordinals.mlir b/iree/compiler/Dialect/HAL/Transforms/test/resolve_entry_point_ordinals.mlir index 6d2822f..773d566 100644 --- a/iree/compiler/Dialect/HAL/Transforms/test/resolve_entry_point_ordinals.mlir +++ b/iree/compiler/Dialect/HAL/Transforms/test/resolve_entry_point_ordinals.mlir
@@ -17,9 +17,9 @@ // CHECK-LABEL: @dispatch_with_nested_references // CHECK-SAME: %[[CMD:.+]]: !hal.command_buffer func @dispatch_with_nested_references(%cmd : !hal.command_buffer) { - %c10 = constant 10 : index - %c11 = constant 11 : index - %c12 = constant 12 : index + %c10 = arith.constant 10 : index + %c11 = arith.constant 11 : index + %c12 = arith.constant 12 : index // CHECK: %[[DEVICE:.+]] = hal.command_buffer.device<%[[CMD]] // CHECK: %[[EXE:.+]] = hal.executable.lookup // CHECK-SAME: device(%[[DEVICE]] : !hal.device) @@ -42,9 +42,9 @@ // CHECK-SAME: %[[EXE:.+]]: !hal.executable %exe: !hal.executable ) { - %c10 = constant 10 : index - %c11 = constant 11 : index - %c12 = constant 12 : index + %c10 = arith.constant 10 : index + %c11 = arith.constant 11 : index + %c12 = arith.constant 12 : index // CHECK: hal.command_buffer.dispatch<%[[CMD]] : !hal.command_buffer> // CHECK-SAME: target(%[[EXE]] : !hal.executable)[2] // CHECK-SAME: workgroups([%c10, %c11, %c12]) @@ -77,7 +77,7 @@ // CHECK-SAME: %[[BUF:.+]]: !hal.buffer %buf: !hal.buffer ) { - %c10 = constant 10 : index + %c10 = arith.constant 10 : index // CHECK: %[[DEVICE:.+]] = hal.command_buffer.device<%[[CMD]] // CHECK: %[[EXE:.+]] = hal.executable.lookup device(%[[DEVICE]] : !hal.device) executable(@exe) // CHECK: hal.command_buffer.dispatch.indirect<%[[CMD]] : !hal.command_buffer> @@ -100,7 +100,7 @@ // CHECK-SAME: %[[BUF:.+]]: !hal.buffer %buf: !hal.buffer ) { - %c10 = constant 10 : index + %c10 = arith.constant 10 : index // CHECK: hal.command_buffer.dispatch.indirect<%[[CMD]] : !hal.command_buffer> // CHECK-SAME: target(%[[EXE]] : !hal.executable)[0] // CHECK-SAME: workgroups(%[[BUF]] : !hal.buffer)[%c10]
diff --git a/iree/compiler/Dialect/HAL/Utils/TypeUtils.cpp b/iree/compiler/Dialect/HAL/Utils/TypeUtils.cpp index 20296da..d63cb01 100644 --- a/iree/compiler/Dialect/HAL/Utils/TypeUtils.cpp +++ b/iree/compiler/Dialect/HAL/Utils/TypeUtils.cpp
@@ -27,12 +27,12 @@ Value align(Location loc, Value value, int64_t alignment, OpBuilder &builder) { // (value + (alignment - 1)) & ~(alignment - 1) - return builder.createOrFold<AndOp>( + return builder.createOrFold<arith::AndIOp>( loc, - builder.createOrFold<AddIOp>( + builder.createOrFold<arith::AddIOp>( loc, value, - builder.createOrFold<ConstantIndexOp>(loc, alignment - 1)), - builder.createOrFold<ConstantIndexOp>(loc, ~(alignment - 1))); + builder.createOrFold<arith::ConstantIndexOp>(loc, alignment - 1)), + builder.createOrFold<arith::ConstantIndexOp>(loc, ~(alignment - 1))); } int32_t getRoundedElementByteWidth(Type type) { @@ -46,7 +46,8 @@ SmallVector<Value, 4> shape; if (shapedType.getRank() >= 1) { for (auto dim : shapedType.getShape()) { - shape.push_back(builder.createOrFold<mlir::ConstantIndexOp>(loc, dim)); + shape.push_back( + builder.createOrFold<mlir::arith::ConstantIndexOp>(loc, dim)); } } return shape;
diff --git a/iree/compiler/Dialect/LinalgExt/IR/LinalgExtOps.cpp b/iree/compiler/Dialect/LinalgExt/IR/LinalgExtOps.cpp index 917af2c..c70db7f 100644 --- a/iree/compiler/Dialect/LinalgExt/IR/LinalgExtOps.cpp +++ b/iree/compiler/Dialect/LinalgExt/IR/LinalgExtOps.cpp
@@ -193,8 +193,8 @@ SmallVector<Range> ScatterOp::getLoopBounds(OpBuilder &builder) { Location loc = getLoc(); - Value zero = builder.create<ConstantIndexOp>(loc, 0); - Value one = builder.create<ConstantIndexOp>(loc, 1); + Value zero = builder.create<arith::ConstantIndexOp>(loc, 0); + Value one = builder.create<arith::ConstantIndexOp>(loc, 1); SmallVector<Range> ranges; for (auto dim : llvm::seq<int64_t>(0, getUpdateType().getRank())) { Value ub = getDimValue(builder, loc, updates(), dim); @@ -278,9 +278,9 @@ loadIndices.push_back(ivs.front()); loadIndices.push_back(Value()); for (auto i : llvm::seq<unsigned>(0, indexDepth)) { - loadIndices.back() = b.create<ConstantIndexOp>(loc, i); + loadIndices.back() = b.create<arith::ConstantIndexOp>(loc, i); Value idx = b.create<memref::LoadOp>(loc, indices(), loadIndices); - starts.push_back(b.create<IndexCastOp>(loc, b.getIndexType(), idx)); + starts.push_back(b.create<arith::IndexCastOp>(loc, b.getIndexType(), idx)); } starts.append(std::next(ivs.begin()), ivs.end()); Value init = b.create<memref::LoadOp>(loc, original(), starts); @@ -378,8 +378,8 @@ int64_t operandRank = getOperandRank(); SmallVector<Range> loopBounds(operandRank); Location loc = getLoc(); - Value zero = builder.create<ConstantIndexOp>(loc, 0); - Value one = builder.create<ConstantIndexOp>(loc, 1); + Value zero = builder.create<arith::ConstantIndexOp>(loc, 0); + Value one = builder.create<arith::ConstantIndexOp>(loc, 1); Value source = operand(0); for (auto dim : llvm::seq<int64_t>(0, operandRank)) { loopBounds[dim].offset = zero; @@ -443,20 +443,21 @@ SmallVector<Value> indices, sortBlkArgs; indices.append(ivs.begin(), ivs.end()); // Bubble sort innermost loop. - Value zero = b.create<ConstantIndexOp>(loc, 0); - Value one = b.create<ConstantIndexOp>(loc, 1); + Value zero = b.create<arith::ConstantIndexOp>(loc, 0); + Value one = b.create<arith::ConstantIndexOp>(loc, 1); Value ub; if (getOperandType(0).isDynamicDim(sortDim)) { ub = b.create<memref::DimOp>(loc, operand(0), sortDim); } else { - ub = b.create<ConstantIndexOp>(loc, getOperandType(0).getDimSize(sortDim)); + ub = b.create<arith::ConstantIndexOp>( + loc, getOperandType(0).getDimSize(sortDim)); } - ub = b.create<SubIOp>(loc, ub, one); + ub = b.create<arith::SubIOp>(loc, ub, one); auto scfFor = b.create<scf::ForOp>( loc, zero, ub, one, ValueRange{}, [&](OpBuilder &b, Location loc, Value iv, ValueRange iters) { SmallVector<Value> indices(ivs); - Value ivPlusOne = b.create<AddIOp>(loc, iv, one); + Value ivPlusOne = b.create<arith::AddIOp>(loc, iv, one); for (auto output : getOutputOperands()) { indices[sortDim] = iv; sortBlkArgs.push_back( @@ -494,7 +495,8 @@ [&](OpBuilder &b, Location loc) { // Swap the pairs if false. SmallVector<Value> indices(ivs.begin(), ivs.end()); - Value ivPlusOne = b.create<AddIOp>(loc, scfFor.getInductionVar(), one); + Value ivPlusOne = + b.create<arith::AddIOp>(loc, scfFor.getInductionVar(), one); for (int i = 0, e = getNumOutputs(); i < e; ++i) { Value v1 = sortBlkArgs[i * 2]; Value v2 = sortBlkArgs[i * 2 + 1]; @@ -551,20 +553,20 @@ SmallVector<Range> FftOp::getLoopBounds(OpBuilder &builder) { SmallVector<Range> res; Location loc = getLoc(); - Value zero = builder.create<ConstantIndexOp>(loc, 0); - Value one = builder.create<ConstantIndexOp>(loc, 1); + Value zero = builder.create<arith::ConstantIndexOp>(loc, 0); + Value one = builder.create<arith::ConstantIndexOp>(loc, 1); for (auto en : llvm::enumerate(getOperandShape().drop_back())) { Value size; if (en.value() == ShapedType::kDynamicSize) { size = getDimValue(builder, loc, getReal(), en.index()); } else { - size = builder.create<ConstantIndexOp>(loc, en.value()); + size = builder.create<arith::ConstantIndexOp>(loc, en.value()); } res.emplace_back(Range{/*offset=*/zero, size, /*stride=*/one}); } Value size = getDimValue(builder, loc, getReal(), getOperandRank() - 1); - Value stride = builder.create<ShiftLeftOp>(loc, one, getStage()); + Value stride = builder.create<arith::ShLIOp>(loc, one, getStage()); res.emplace_back(Range{/*offset=*/zero, size, /*stride=*/stride}); return res; } @@ -577,15 +579,15 @@ auto f32Type = b.getF32Type(); auto indexToF32 = [](OpBuilder &builder, Location loc, Value v) -> Value { - v = builder.create<IndexCastOp>(loc, builder.getI32Type(), v); - return builder.create<SIToFPOp>(loc, builder.getF32Type(), v); + v = builder.create<arith::IndexCastOp>(loc, builder.getI32Type(), v); + return builder.create<arith::SIToFPOp>(loc, builder.getF32Type(), v); }; // We will need exp(-2 * PI * j / m * I), compute "-2 * PI / m" for imag part // first. - Value coeff = b.create<ConstantFloatOp>( + Value coeff = b.create<arith::ConstantFloatOp>( loc, llvm::APFloat(static_cast<float>(-2 * acos(-1))), f32Type); - coeff = b.create<DivFOp>(loc, coeff, indexToF32(b, loc, wholeSize)); + coeff = b.create<arith::DivFOp>(loc, coeff, indexToF32(b, loc, wholeSize)); b.create<linalg::GenericOp>( loc, TypeRange{}, ValueRange{}, operands, maps, getLoopIteratorTypes(), @@ -596,7 +598,7 @@ Value rhsImag = args[3]; // Compute "-2 * PI / m * j" - Value w = b.create<MulFOp>( + Value w = b.create<arith::MulFOp>( loc, coeff, indexToF32(b, loc, b.create<linalg::IndexOp>(loc, rank - 1))); Value wReal = b.create<math::CosOp>(loc, w); @@ -604,20 +606,20 @@ // t = w * a[k + j + mh]; // -> (x + yi)(u + vi) = (xu - yv) + (xv + yu)i - Value xu = b.create<MulFOp>(loc, wReal, rhsReal); - Value yv = b.create<MulFOp>(loc, wImag, rhsImag); - Value xv = b.create<MulFOp>(loc, wReal, rhsImag); - Value yu = b.create<MulFOp>(loc, wImag, rhsReal); - Value tReal = b.create<SubFOp>(loc, xu, yv); - Value tImag = b.create<AddFOp>(loc, xv, yu); + Value xu = b.create<arith::MulFOp>(loc, wReal, rhsReal); + Value yv = b.create<arith::MulFOp>(loc, wImag, rhsImag); + Value xv = b.create<arith::MulFOp>(loc, wReal, rhsImag); + Value yu = b.create<arith::MulFOp>(loc, wImag, rhsReal); + Value tReal = b.create<arith::SubFOp>(loc, xu, yv); + Value tImag = b.create<arith::AddFOp>(loc, xv, yu); // cplx u = a[k + j]; // a[k + j] = u + t; // a[k + j + mh] = u - t; - Value r1 = b.create<AddFOp>(loc, lhsReal, tReal); - Value r2 = b.create<AddFOp>(loc, lhsImag, tImag); - Value r3 = b.create<SubFOp>(loc, lhsReal, tReal); - Value r4 = b.create<SubFOp>(loc, lhsImag, tImag); + Value r1 = b.create<arith::AddFOp>(loc, lhsReal, tReal); + Value r2 = b.create<arith::AddFOp>(loc, lhsImag, tImag); + Value r3 = b.create<arith::SubFOp>(loc, lhsReal, tReal); + Value r4 = b.create<arith::SubFOp>(loc, lhsImag, tImag); b.create<linalg::YieldOp>(loc, ValueRange{r1, r2, r3, r4}); }); } @@ -645,20 +647,20 @@ // t = w * a[k + j + mh]; // -> (x + yi)(u + vi) = (xu - yv) + (xv + yu)i - Value xu = b.create<MulFOp>(loc, wReal, rhsReal); - Value yv = b.create<MulFOp>(loc, wImag, rhsImag); - Value xv = b.create<MulFOp>(loc, wReal, rhsImag); - Value yu = b.create<MulFOp>(loc, wImag, rhsReal); - Value tReal = b.create<SubFOp>(loc, xu, yv); - Value tImag = b.create<AddFOp>(loc, xv, yu); + Value xu = b.create<arith::MulFOp>(loc, wReal, rhsReal); + Value yv = b.create<arith::MulFOp>(loc, wImag, rhsImag); + Value xv = b.create<arith::MulFOp>(loc, wReal, rhsImag); + Value yu = b.create<arith::MulFOp>(loc, wImag, rhsReal); + Value tReal = b.create<arith::SubFOp>(loc, xu, yv); + Value tImag = b.create<arith::AddFOp>(loc, xv, yu); // cplx u = a[k + j]; // a[k + j] = u + t; // a[k + j + mh] = u - t; - Value r1 = b.create<AddFOp>(loc, lhsReal, tReal); - Value r2 = b.create<AddFOp>(loc, lhsImag, tImag); - Value r3 = b.create<SubFOp>(loc, lhsReal, tReal); - Value r4 = b.create<SubFOp>(loc, lhsImag, tImag); + Value r1 = b.create<arith::AddFOp>(loc, lhsReal, tReal); + Value r2 = b.create<arith::AddFOp>(loc, lhsImag, tImag); + Value r3 = b.create<arith::SubFOp>(loc, lhsReal, tReal); + Value r4 = b.create<arith::SubFOp>(loc, lhsImag, tImag); b.create<linalg::YieldOp>(loc, ValueRange{r1, r2, r3, r4}); }); } @@ -682,9 +684,9 @@ Value real = getReal(); Value imag = getImag(); Value stage = getStage(); - Value one = b.create<ConstantIndexOp>(loc, 1); - Value wholeSize = b.create<ShiftLeftOp>(loc, one, stage); - Value halfSize = b.create<SignedShiftRightOp>(loc, wholeSize, one); + Value one = b.create<arith::ConstantIndexOp>(loc, 1); + Value wholeSize = b.create<arith::ShLIOp>(loc, one, stage); + Value halfSize = b.create<arith::ShRSIOp>(loc, wholeSize, one); auto rank = getOperandRank(); SmallVector<Value> operands; @@ -698,7 +700,8 @@ b.create<memref::SubViewOp>(loc, imag, lhsIvs, sizes, ones)); SmallVector<OpFoldResult> rhsIvs(ivs.begin(), ivs.end()); - rhsIvs.back() = b.create<AddIOp>(loc, ivs.back(), halfSize).getResult(); + rhsIvs.back() = + b.create<arith::AddIOp>(loc, ivs.back(), halfSize).getResult(); operands.push_back( b.create<memref::SubViewOp>(loc, real, rhsIvs, sizes, ones)); operands.push_back( @@ -819,8 +822,8 @@ SmallVector<Range> ReverseOp::getLoopBounds(OpBuilder &builder) { Location loc = getLoc(); - Value zero = builder.create<ConstantIndexOp>(loc, 0); - Value one = builder.create<ConstantIndexOp>(loc, 1); + Value zero = builder.create<arith::ConstantIndexOp>(loc, 0); + Value one = builder.create<arith::ConstantIndexOp>(loc, 1); SmallVector<Range> ranges; for (auto dim : llvm::seq<int64_t>(0, getOperandRank())) { Value ub = getDimValue(builder, loc, input(), dim); @@ -835,8 +838,9 @@ SmallVector<Value> mirrorIndices(ivs.begin(), ivs.end()); for (auto dim : dims()) { auto size = getDimValue(b, loc, input(), dim); - size = b.create<SubIOp>(loc, size, b.create<ConstantIndexOp>(loc, 1)); - mirrorIndices[dim] = b.create<SubIOp>(loc, size, mirrorIndices[dim]); + size = b.create<arith::SubIOp>(loc, size, + b.create<arith::ConstantIndexOp>(loc, 1)); + mirrorIndices[dim] = b.create<arith::SubIOp>(loc, size, mirrorIndices[dim]); } Value val = b.create<memref::LoadOp>(loc, input(), ivs); b.create<memref::StoreOp>(loc, val, output(), mirrorIndices);
diff --git a/iree/compiler/Dialect/LinalgExt/IR/TiledOpInterface.cpp b/iree/compiler/Dialect/LinalgExt/IR/TiledOpInterface.cpp index 2cd8a14..f6bc021 100644 --- a/iree/compiler/Dialect/LinalgExt/IR/TiledOpInterface.cpp +++ b/iree/compiler/Dialect/LinalgExt/IR/TiledOpInterface.cpp
@@ -25,8 +25,8 @@ static Value getValue(OpBuilder &builder, Location loc, OpFoldResult valueOrAttr) { if (auto attr = valueOrAttr.dyn_cast<Attribute>()) { - return builder.create<ConstantIndexOp>(loc, - attr.cast<IntegerAttr>().getInt()); + return builder.create<arith::ConstantIndexOp>( + loc, attr.cast<IntegerAttr>().getInt()); } return valueOrAttr.get<Value>(); } @@ -56,8 +56,8 @@ Value source = insertSliceOp.source(); RankedTensorType sourceType = insertSliceOp.getSourceType(); Location loc = op->getLoc(); - Value zero = b.create<ConstantIndexOp>(loc, 0); - Value one = b.create<ConstantIndexOp>(loc, 1); + Value zero = b.create<arith::ConstantIndexOp>(loc, 0); + Value one = b.create<arith::ConstantIndexOp>(loc, 1); SmallVector<Range> loopBounds(sourceType.getRank(), Range{zero, nullptr, one}); for (auto dim :
diff --git a/iree/compiler/Dialect/LinalgExt/IR/test/invalid.mlir b/iree/compiler/Dialect/LinalgExt/IR/test/invalid.mlir index 2c76c95..5266e64 100644 --- a/iree/compiler/Dialect/LinalgExt/IR/test/invalid.mlir +++ b/iree/compiler/Dialect/LinalgExt/IR/test/invalid.mlir
@@ -5,7 +5,7 @@ %0 = linalg_ext.sort dimension(1) outs(%arg0 : tensor<128xi32>) { ^bb0(%arg1: i32, %arg2: i32): // no predecessors - %1 = cmpi sgt, %arg1, %arg2 : i32 + %1 = arith.cmpi sgt, %arg1, %arg2 : i32 linalg_ext.yield %1 : i1 } -> tensor<128xi32> return %0 : tensor<128xi32> @@ -18,7 +18,7 @@ %0 = linalg_ext.sort outs(%arg0 : tensor<3x4xi32>) { ^bb0(%arg1: i32, %arg2: i32): // no predecessors - %1 = cmpi sgt, %arg1, %arg2 : i32 + %1 = arith.cmpi sgt, %arg1, %arg2 : i32 linalg_ext.yield %1 : i1 } -> tensor<3x4xi32> return %0 : tensor<3x4xi32> @@ -32,7 +32,7 @@ %0:2 = linalg_ext.sort dimension(0) outs(%arg0, %arg1 : tensor<?x?xi32>, tensor<?xf32>) { ^bb0(%arg2: i32, %arg3: i32, %arg4 : f32, %arg5 : f32): // no predecessors - %1 = cmpf ogt, %arg4, %arg5 : f32 + %1 = arith.cmpf ogt, %arg4, %arg5 : f32 linalg_ext.yield %1 : i1 } -> tensor<?x?xi32>, tensor<?xf32> return %0#0, %0#1 : tensor<?x?xi32>, tensor<?xf32> @@ -46,7 +46,7 @@ %0:2 = linalg_ext.sort dimension(0) outs(%arg0, %arg1 : tensor<?xi32>, tensor<42xf32>) { ^bb0(%arg2: i32, %arg3: i32, %arg4 : f32, %arg5 : f32): // no predecessors - %1 = cmpf ogt, %arg4, %arg5 : f32 + %1 = arith.cmpf ogt, %arg4, %arg5 : f32 linalg_ext.yield %1 : i1 } -> tensor<?xi32>, tensor<42xf32> return %0#0, %0#1 : tensor<?xi32>, tensor<42xf32> @@ -62,7 +62,7 @@ ins(%update, %indices : memref<?x?xf32>, tensor<?x1xi32>) outs(%original : tensor<?x?xf32>) { ^bb0(%arg1: f32, %arg2: f32): - %1 = addf %arg1, %arg2 : f32 + %1 = arith.addf %arg1, %arg2 : f32 linalg_ext.yield %1 : f32 } -> tensor<?x?xf32> return %0 : tensor<?x?xf32> @@ -78,7 +78,7 @@ ins(%update, %indices : tensor<?x?xf32>, memref<?x1xi32>) outs(%original : tensor<?x?xf32>) { ^bb0(%arg1: f32, %arg2: f32): - %1 = addf %arg1, %arg2 : f32 + %1 = arith.addf %arg1, %arg2 : f32 linalg_ext.yield %1 : f32 } -> tensor<?x?xf32> return %0 : tensor<?x?xf32> @@ -94,7 +94,7 @@ ins(%update, %indices : tensor<?x?xf32>, tensor<?x1xi32>) outs(%original : tensor<?x?xf32>) { ^bb0(%arg1: f32, %arg2: f32): - %1 = addf %arg1, %arg2 : f32 + %1 = arith.addf %arg1, %arg2 : f32 linalg_ext.yield %1 : f32 } -> tensor<?x?xf32>, tensor<?x?xf32> return %0, %1 : tensor<?x?xf32>, tensor<?x?xf32> @@ -110,7 +110,7 @@ ins(%update, %indices : tensor<?x?xf32>, tensor<?x1xi32>) outs(%original : memref<?x?xf32>) { ^bb0(%arg1: f32, %arg2: f32): - %1 = addf %arg1, %arg2 : f32 + %1 = arith.addf %arg1, %arg2 : f32 linalg_ext.yield %1 : f32 } -> tensor<?x?xf32> return %0 : tensor<?x?xf32> @@ -126,7 +126,7 @@ ins(%update, %indices : tensor<?x?xf32>, tensor<?x1xi32>) outs(%original : tensor<?x?xf32>) { ^bb0(%arg1: f32, %arg2: f32): - %1 = addf %arg1, %arg2 : f32 + %1 = arith.addf %arg1, %arg2 : f32 linalg_ext.yield %1 : f32 } -> memref<?x?xf32> return %0 : memref<?x?xf32> @@ -142,7 +142,7 @@ ins(%update, %indices : memref<?x?xf32>, tensor<?x1xi32>) outs(%original : memref<?x?xf32>) { ^bb0(%arg1: f32, %arg2: f32): - %1 = addf %arg1, %arg2 : f32 + %1 = arith.addf %arg1, %arg2 : f32 linalg_ext.yield %1 : f32 } return @@ -158,7 +158,7 @@ ins(%update, %indices : memref<?x?xf32>, memref<?x1xi32>) outs(%original : tensor<?x?xf32>) { ^bb0(%arg1: f32, %arg2: f32): - %1 = addf %arg1, %arg2 : f32 + %1 = arith.addf %arg1, %arg2 : f32 linalg_ext.yield %1 : f32 } return @@ -174,7 +174,7 @@ ins(%update, %indices : tensor<?x?xf32>, tensor<48x1xi32>) outs(%original : tensor<?x?xf32>) { ^bb0(%arg1: f32, %arg2: f32): - %1 = addf %arg1, %arg2 : f32 + %1 = arith.addf %arg1, %arg2 : f32 linalg_ext.yield %1 : f32 } -> tensor<?x?xf32> return %0 : tensor<?x?xf32> @@ -190,7 +190,7 @@ ins(%update, %indices : tensor<64x?xf32>, tensor<48x1xi32>) outs(%original : tensor<?x?xf32>) { ^bb0(%arg1: f32, %arg2: f32): - %1 = addf %arg1, %arg2 : f32 + %1 = arith.addf %arg1, %arg2 : f32 linalg_ext.yield %1 : f32 } -> tensor<?x?xf32> return %0 : tensor<?x?xf32> @@ -206,7 +206,7 @@ ins(%update, %indices : tensor<?x?x?xf32>, tensor<?x1xi32>) outs(%original : tensor<?x?xf32>) { ^bb0(%arg1: f32, %arg2: f32): - %1 = addf %arg1, %arg2 : f32 + %1 = arith.addf %arg1, %arg2 : f32 linalg_ext.yield %1 : f32 } -> tensor<?x?xf32> return %0 : tensor<?x?xf32> @@ -222,7 +222,7 @@ ins(%update, %indices : tensor<?x4xf32>, tensor<?x1xi32>) outs(%original : tensor<?x?xf32>) { ^bb0(%arg1: f32, %arg2: f32): - %1 = addf %arg1, %arg2 : f32 + %1 = arith.addf %arg1, %arg2 : f32 linalg_ext.yield %1 : f32 } -> tensor<?x?xf32> return %0 : tensor<?x?xf32> @@ -238,8 +238,8 @@ ins(%update, %indices : tensor<?x?xi32>, tensor<?x1xi32>) outs(%original : tensor<?x?xi32>) { ^bb0(%arg1: index, %arg2: index): - %1 = addi %arg1, %arg2 : index - %2 = index_cast %1 : index to i32 + %1 = arith.addi %arg1, %arg2 : index + %2 = arith.index_cast %1 : index to i32 linalg_ext.yield %2 : i32 } -> tensor<?x?xi32> return %0 : tensor<?x?xi32> @@ -255,8 +255,8 @@ ins(%update, %indices : tensor<?x?xi32>, tensor<?x1xi32>) outs(%original : tensor<?x?xi32>) { ^bb0(%arg1: i64, %arg2: i32): - %1 = trunci %arg1 : i64 to i32 - %2 = addi %1, %arg2 : i32 + %1 = arith.trunci %arg1 : i64 to i32 + %2 = arith.addi %1, %arg2 : i32 linalg_ext.yield %2 : i32 } -> tensor<?x?xi32> return %0 : tensor<?x?xi32> @@ -272,8 +272,8 @@ ins(%update, %indices : tensor<?x?xi32>, tensor<?x1xi32>) outs(%original : tensor<?x?xi32>) { ^bb0(%arg1: i32, %arg2: i64): - %1 = trunci %arg2 : i64 to i32 - %2 = addi %1, %arg1 : i32 + %1 = arith.trunci %arg2 : i64 to i32 + %2 = arith.addi %1, %arg1 : i32 linalg_ext.yield %2 : i32 } -> tensor<?x?xi32> return %0 : tensor<?x?xi32> @@ -289,8 +289,8 @@ ins(%update, %indices : tensor<?x?xi32>, tensor<?x1xi32>) outs(%original : tensor<?x?xi64>) { ^bb0(%arg1: i32, %arg2: i64): - %1 = sexti %arg1 : i32 to i64 - %2 = addi %1, %arg2 : i64 + %1 = arith.extsi %arg1 : i32 to i64 + %2 = arith.addi %1, %arg2 : i64 linalg_ext.yield %2 : i64 } -> tensor<?x?xi64> return %0 : tensor<?x?xi64> @@ -306,7 +306,7 @@ ins(%update, %indices : tensor<?x?xi64>, tensor<?x1xi32>) outs(%original : tensor<?x?xi64>) { ^bb0(%arg1: i64, %arg2: i64, %arg3 : i64): - %1 = addi %arg1, %arg2 : i64 + %1 = arith.addi %arg1, %arg2 : i64 linalg_ext.yield %1 : i64 } -> tensor<?x?xi64> return %0 : tensor<?x?xi64> @@ -322,8 +322,8 @@ ins(%update, %indices : tensor<?x?xi64>, tensor<?x1xi32>) outs(%original : tensor<?x?xi64>) { ^bb0(%arg1: i64, %arg2: i64): - %1 = addi %arg1, %arg2 : i64 - %2 = trunci %1 : i64 to i32 + %1 = arith.addi %arg1, %arg2 : i64 + %2 = arith.trunci %1 : i64 to i32 // expected-error @+1 {{mismatch in type of yielded value 'i32' and argument of the region 'i64'}} linalg_ext.yield %2 : i32 } -> tensor<?x?xi64> @@ -339,8 +339,8 @@ ins(%update, %indices : tensor<?x?xi64>, tensor<?x1xi32>) outs(%original : tensor<?x?xi64>) { ^bb0(%arg1: i64, %arg2: i64): - %1 = addi %arg1, %arg2 : i64 - %2 = trunci %1 : i64 to i32 + %1 = arith.addi %arg1, %arg2 : i64 + %2 = arith.trunci %1 : i64 to i32 // expected-error @+1 {{expected region to yield a single value}} linalg_ext.yield %1, %2 : i64, i32 } -> tensor<?x?xi64> @@ -357,8 +357,8 @@ ins(%update, %indices : tensor<?x?xi64>, tensor<?x?xi32>) outs(%original : tensor<?x?xi64>) { ^bb0(%arg1: i64, %arg2: i64): - %1 = addi %arg1, %arg2 : i64 - %2 = trunci %1 : i64 to i32 + %1 = arith.addi %arg1, %arg2 : i64 + %2 = arith.trunci %1 : i64 to i32 linalg_ext.yield %1, %2 : i64, i32 } -> tensor<?x?xi64> return %0 : tensor<?x?xi64> @@ -374,8 +374,8 @@ ins(%update, %indices : tensor<?x?xi64>, tensor<?x2xi32>) outs(%original : tensor<?x?xi64>) { ^bb0(%arg1: i64, %arg2: i64): - %1 = addi %arg1, %arg2 : i64 - %2 = trunci %1 : i64 to i32 + %1 = arith.addi %arg1, %arg2 : i64 + %2 = arith.trunci %1 : i64 to i32 linalg_ext.yield %1, %2 : i64, i32 } -> tensor<?x?xi64> return %0 : tensor<?x?xi64>
diff --git a/iree/compiler/Dialect/LinalgExt/IR/test/roundtrip.mlir b/iree/compiler/Dialect/LinalgExt/IR/test/roundtrip.mlir index cf19294..d818a90 100644 --- a/iree/compiler/Dialect/LinalgExt/IR/test/roundtrip.mlir +++ b/iree/compiler/Dialect/LinalgExt/IR/test/roundtrip.mlir
@@ -8,7 +8,7 @@ %0 = linalg_ext.sort outs(%arg0 : tensor<128xi32>) { ^bb0(%arg1: i32, %arg2: i32): // no predecessors - %1 = cmpi sgt, %arg1, %arg2 : i32 + %1 = arith.cmpi sgt, %arg1, %arg2 : i32 linalg_ext.yield %1 : i1 } -> tensor<128xi32> return %0 : tensor<128xi32> @@ -24,7 +24,7 @@ linalg_ext.sort dimension(0) outs(%arg0 : memref<128xi32>) { ^bb0(%arg1: i32, %arg2: i32): // no predecessors - %0 = cmpi sgt, %arg1, %arg2 : i32 + %0 = arith.cmpi sgt, %arg1, %arg2 : i32 linalg_ext.yield %0 : i1 } return @@ -38,7 +38,7 @@ %0:2 = linalg_ext.sort dimension(0) outs(%arg0, %arg1 : tensor<?x?xi32>, tensor<?x?xf32>) { ^bb0(%arg2: i32, %arg3: i32, %arg4 : f32, %arg5 : f32): // no predecessors - %1 = cmpf ogt, %arg4, %arg5 : f32 + %1 = arith.cmpf ogt, %arg4, %arg5 : f32 linalg_ext.yield %1 : i1 } -> tensor<?x?xi32>, tensor<?x?xf32> return %0#0, %0#1 : tensor<?x?xi32>, tensor<?x?xf32> @@ -57,7 +57,7 @@ linalg_ext.sort dimension(0) outs(%arg0, %arg1 : memref<?x?xi32>, memref<?x?xf32>) { ^bb0(%arg2: i32, %arg3: i32, %arg4 : f32, %arg5 : f32): // no predecessors - %1 = cmpf ogt, %arg4, %arg5 : f32 + %1 = arith.cmpf ogt, %arg4, %arg5 : f32 linalg_ext.yield %1 : i1 } return @@ -77,7 +77,7 @@ ins(%update, %indices : tensor<?x?xf32>, tensor<?x1xi32>) outs(%original: tensor<?x?xf32>) { ^bb0(%arg1: f32, %arg2: f32): - %1 = addf %arg1, %arg2 : f32 + %1 = arith.addf %arg1, %arg2 : f32 linalg_ext.yield %1 : f32 } -> tensor<?x?xf32> return %0 : tensor<?x?xf32> @@ -101,7 +101,7 @@ ins(%update, %indices : tensor<48x3xf32>, tensor<48x1xi32>) outs(%original: tensor<128x3xf32>) { ^bb0(%arg1: f32, %arg2: f32): - %1 = addf %arg1, %arg2 : f32 + %1 = arith.addf %arg1, %arg2 : f32 linalg_ext.yield %1 : f32 } -> tensor<128x3xf32> return %0 : tensor<128x3xf32> @@ -125,7 +125,7 @@ ins(%update, %indices : tensor<48x3xf32>, tensor<48x2xi32>) outs(%original: tensor<1x128x3xf32>) { ^bb0(%arg1: f32, %arg2: f32): - %1 = addf %arg1, %arg2 : f32 + %1 = arith.addf %arg1, %arg2 : f32 linalg_ext.yield %1 : f32 } -> tensor<1x128x3xf32> return %0 : tensor<1x128x3xf32> @@ -149,7 +149,7 @@ ins(%update, %indices : memref<?x?xf32>, memref<?x1xi32>) outs(%original: memref<?x?xf32>) { ^bb0(%arg1: f32, %arg2: f32): - %1 = addf %arg1, %arg2 : f32 + %1 = arith.addf %arg1, %arg2 : f32 linalg_ext.yield %1 : f32 } return @@ -173,7 +173,7 @@ ins(%update, %indices : memref<48x3xf32>, memref<48x1xi32>) outs(%original: memref<128x3xf32>) { ^bb0(%arg1: f32, %arg2: f32): - %1 = addf %arg1, %arg2 : f32 + %1 = arith.addf %arg1, %arg2 : f32 linalg_ext.yield %1 : f32 } return @@ -197,7 +197,7 @@ ins(%update, %indices : memref<48x3xf32>, memref<48x2xi32>) outs(%original: memref<1x128x3xf32>) { ^bb0(%arg1: f32, %arg2: f32): - %1 = addf %arg1, %arg2 : f32 + %1 = arith.addf %arg1, %arg2 : f32 linalg_ext.yield %1 : f32 } return @@ -285,7 +285,7 @@ func @fft_tensor(%arg0: tensor<1024xf32>, %arg1: tensor<1024xf32>) -> (tensor<1024xf32>, tensor<1024xf32>) { - %cst1 = constant 1 : index + %cst1 = arith.constant 1 : index %0:2 = linalg_ext.fft ins(%cst1: index) outs(%arg0, %arg1: tensor<1024xf32>, tensor<1024xf32>) @@ -295,7 +295,7 @@ // CHECK-LABEL: func @fft_tensor( // CHECK-SAME: %[[REAL:[a-zA-Z0-9_]+]] // CHECK-SAME: %[[IMAG:[a-zA-Z0-9_]+]] -// CHECK: %[[CST:.+]] = constant 1 : index +// CHECK: %[[CST:.+]] = arith.constant 1 : index // CHECK: %[[RES:.+]]:2 = linalg_ext.fft // CHECK-SAME: ins(%[[CST]] : index) // CHECK-SAME: outs(%[[REAL]], %[[IMAG]] : tensor<1024xf32>, tensor<1024xf32>) @@ -305,7 +305,7 @@ // ----- func @fft_memref(%arg0: memref<1024xf32>, %arg1: memref<1024xf32>) { - %cst1 = constant 1 : index + %cst1 = arith.constant 1 : index linalg_ext.fft ins(%cst1: index) outs(%arg0, %arg1: memref<1024xf32>, memref<1024xf32>) @@ -314,7 +314,7 @@ // CHECK-LABEL: func @fft_memref( // CHECK-SAME: %[[REAL:[a-zA-Z0-9_]+]] // CHECK-SAME: %[[IMAG:[a-zA-Z0-9_]+]] -// CHECK: %[[CST:.+]] = constant 1 : index +// CHECK: %[[CST:.+]] = arith.constant 1 : index // CHECK: linalg_ext.fft // CHECK-SAME: ins(%[[CST]] : index) // CHECK-SAME: outs(%[[REAL]], %[[IMAG]] : memref<1024xf32>, memref<1024xf32>) @@ -324,7 +324,7 @@ func @fft_tensor_coef(%arg0: tensor<1024xf32>, %arg1: tensor<1024xf32>, %arg2: tensor<1xf32>, %arg3: tensor<1xf32>) -> (tensor<1024xf32>, tensor<1024xf32>) { - %cst1 = constant 1 : index + %cst1 = arith.constant 1 : index %0:2 = linalg_ext.fft ins(%cst1, %arg2, %arg3: index, tensor<1xf32>, tensor<1xf32>) outs(%arg0, %arg1: tensor<1024xf32>, tensor<1024xf32>) @@ -336,7 +336,7 @@ // CHECK-SAME: %[[IMAG:[a-zA-Z0-9_]+]] // CHECK-SAME: %[[COEF_REAL:[a-zA-Z0-9_]+]] // CHECK-SAME: %[[COEF_IMAG:[a-zA-Z0-9_]+]] -// CHECK: %[[CST:.+]] = constant 1 : index +// CHECK: %[[CST:.+]] = arith.constant 1 : index // CHECK: %[[RES:.+]]:2 = linalg_ext.fft // CHECK-SAME: ins(%[[CST]], %[[COEF_REAL]], %[[COEF_IMAG]] : index, tensor<1xf32>, tensor<1xf32>) // CHECK-SAME: outs(%[[REAL]], %[[IMAG]] : tensor<1024xf32>, tensor<1024xf32>) @@ -347,7 +347,7 @@ func @fft_memref_coef(%arg0: memref<1024xf32>, %arg1: memref<1024xf32>, %arg2: memref<1xf32>, %arg3: memref<1xf32>) { - %cst1 = constant 1 : index + %cst1 = arith.constant 1 : index linalg_ext.fft ins(%cst1, %arg2, %arg3: index, memref<1xf32>, memref<1xf32>) outs(%arg0, %arg1: memref<1024xf32>, memref<1024xf32>) @@ -358,7 +358,7 @@ // CHECK-SAME: %[[IMAG:[a-zA-Z0-9_]+]] // CHECK-SAME: %[[COEF_REAL:[a-zA-Z0-9_]+]] // CHECK-SAME: %[[COEF_IMAG:[a-zA-Z0-9_]+]] -// CHECK: %[[CST:.+]] = constant 1 : index +// CHECK: %[[CST:.+]] = arith.constant 1 : index // CHECK: linalg_ext.fft // CHECK-SAME: ins(%[[CST]], %[[COEF_REAL]], %[[COEF_IMAG]] : index, memref<1xf32>, memref<1xf32>) // CHECK-SAME: outs(%[[REAL]], %[[IMAG]] : memref<1024xf32>, memref<1024xf32>) @@ -369,7 +369,7 @@ // The size of coefficient tensor is 2^(stage-1). func @fft_tensor_coef_stage_5(%arg0: tensor<1024xf32>, %arg1: tensor<1024xf32>, %arg2: tensor<16xf32>, %arg3: tensor<16xf32>) -> (tensor<1024xf32>, tensor<1024xf32>) { - %cst1 = constant 5 : index + %cst1 = arith.constant 5 : index %0:2 = linalg_ext.fft ins(%cst1, %arg2, %arg3: index, tensor<16xf32>, tensor<16xf32>) outs(%arg0, %arg1: tensor<1024xf32>, tensor<1024xf32>) @@ -381,7 +381,7 @@ // CHECK-SAME: %[[IMAG:[a-zA-Z0-9_]+]] // CHECK-SAME: %[[COEF_REAL:[a-zA-Z0-9_]+]] // CHECK-SAME: %[[COEF_IMAG:[a-zA-Z0-9_]+]] -// CHECK: %[[CST:.+]] = constant 5 : index +// CHECK: %[[CST:.+]] = arith.constant 5 : index // CHECK: %[[RES:.+]]:2 = linalg_ext.fft // CHECK-SAME: ins(%[[CST]], %[[COEF_REAL]], %[[COEF_IMAG]] : index, tensor<16xf32>, tensor<16xf32>) // CHECK-SAME: outs(%[[REAL]], %[[IMAG]] : tensor<1024xf32>, tensor<1024xf32>) @@ -426,8 +426,8 @@ // ----- func @reverse_dynamic_tensor(%arg0: tensor<?x?xi32>) -> tensor<?x?xi32> { - %c0 = constant 0 : index - %c1 = constant 1 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index %d0 = tensor.dim %arg0, %c0 : tensor<?x?xi32> %d1 = tensor.dim %arg0, %c1 : tensor<?x?xi32> %init = linalg.init_tensor [%d0, %d1] : tensor<?x?xi32> @@ -439,8 +439,8 @@ } // CHECK-LABEL: func @reverse_dynamic_tensor // CHECK-SAME: %[[ARG0:[a-zA-Z0-9]+]]: tensor<?x?xi32> -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK-DAG: %[[D0:.+]] = tensor.dim %[[ARG0]], %[[C0]] // CHECK-DAG: %[[D1:.+]] = tensor.dim %[[ARG0]], %[[C1]] // CHECK: %[[INIT:.+]] = linalg.init_tensor [%[[D0]], %[[D1]]] @@ -452,8 +452,8 @@ // ----- func @reverse_static_dynamic_tensor(%arg0: tensor<3x5xi32>) -> tensor<?x?xi32> { - %c0 = constant 0 : index - %c1 = constant 1 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index %d0 = tensor.dim %arg0, %c0 : tensor<3x5xi32> %d1 = tensor.dim %arg0, %c1 : tensor<3x5xi32> %init = linalg.init_tensor [%d0, %d1] : tensor<?x?xi32> @@ -465,8 +465,8 @@ } // CHECK-LABEL: func @reverse_static_dynamic_tensor // CHECK-SAME: %[[ARG0:[a-zA-Z0-9]+]]: tensor<3x5xi32> -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK-DAG: %[[D0:.+]] = tensor.dim %[[ARG0]], %[[C0]] // CHECK-DAG: %[[D1:.+]] = tensor.dim %[[ARG0]], %[[C1]] // CHECK: %[[INIT:.+]] = linalg.init_tensor [%[[D0]], %[[D1]]]
diff --git a/iree/compiler/Dialect/LinalgExt/Transforms/ConvertToLoops.cpp b/iree/compiler/Dialect/LinalgExt/Transforms/ConvertToLoops.cpp index a3da0b7..6b886dd 100644 --- a/iree/compiler/Dialect/LinalgExt/Transforms/ConvertToLoops.cpp +++ b/iree/compiler/Dialect/LinalgExt/Transforms/ConvertToLoops.cpp
@@ -91,9 +91,9 @@ struct LinalgExtToLoopsPass : public LinalgExtToLoopsBase<LinalgExtToLoopsPass> { void getDependentDialects(DialectRegistry ®istry) const override { - registry - .insert<linalg::LinalgDialect, StandardOpsDialect, math::MathDialect, - memref::MemRefDialect, scf::SCFDialect>(); + registry.insert<linalg::LinalgDialect, StandardOpsDialect, + mlir::arith::ArithmeticDialect, math::MathDialect, + memref::MemRefDialect, scf::SCFDialect>(); } void runOnOperation() override {
diff --git a/iree/compiler/Dialect/LinalgExt/Transforms/Tiling.cpp b/iree/compiler/Dialect/LinalgExt/Transforms/Tiling.cpp index ae62abe..3224687 100644 --- a/iree/compiler/Dialect/LinalgExt/Transforms/Tiling.cpp +++ b/iree/compiler/Dialect/LinalgExt/Transforms/Tiling.cpp
@@ -64,8 +64,8 @@ static Value getValue(OpBuilder &builder, Location loc, OpFoldResult valueOrAttr) { if (auto attr = valueOrAttr.dyn_cast<Attribute>()) { - return builder.create<ConstantIndexOp>(loc, - attr.cast<IntegerAttr>().getInt()); + return builder.create<arith::ConstantIndexOp>( + loc, attr.cast<IntegerAttr>().getInt()); } return valueOrAttr.get<Value>(); } @@ -265,7 +265,8 @@ registry .insert<AffineDialect, IREE::Flow::FlowDialect, linalg::LinalgDialect, linalg_ext::LinalgExtDialect, memref::MemRefDialect, - StandardOpsDialect, tensor::TensorDialect, scf::SCFDialect>(); + StandardOpsDialect, mlir::arith::ArithmeticDialect, + math::MathDialect, tensor::TensorDialect, scf::SCFDialect>(); } void runOnOperation() override; };
diff --git a/iree/compiler/Dialect/LinalgExt/Transforms/test/convert_to_loops.mlir b/iree/compiler/Dialect/LinalgExt/Transforms/test/convert_to_loops.mlir index d3fc5a4..2298ced 100644 --- a/iree/compiler/Dialect/LinalgExt/Transforms/test/convert_to_loops.mlir +++ b/iree/compiler/Dialect/LinalgExt/Transforms/test/convert_to_loops.mlir
@@ -4,26 +4,26 @@ linalg_ext.sort dimension(0) outs(%arg0 : memref<128xi32>) { ^bb0(%arg2: i32, %arg3: i32): // no predecessors - %0 = cmpi sgt, %arg2, %arg3 : i32 + %0 = arith.cmpi sgt, %arg2, %arg3 : i32 linalg_ext.yield %0 : i1 } return } // CHECK-LABEL: func @sort_1d // CHECK-SAME: %[[BUF:[a-zA-Z0-9]+]] -// CHECK-DAG: %[[C128:.+]] = constant 128 : index -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index -// CHECK-DAG: %[[C127:.+]] = constant 127 : index +// CHECK-DAG: %[[C128:.+]] = arith.constant 128 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index +// CHECK-DAG: %[[C127:.+]] = arith.constant 127 : index // CHECK: scf.for %[[ARG1:.+]] = %[[C0]] to %[[C128]] step %[[C1]] // CHECK: scf.for %[[ARG2:.+]] = %[[C0]] to %[[C127]] step %[[C1]] -// CHECK: %[[T1:.+]] = addi %[[ARG2]], %[[C1]] : index +// CHECK: %[[T1:.+]] = arith.addi %[[ARG2]], %[[C1]] : index // CHECK: %[[V1:.+]] = memref.load %[[BUF]][%[[ARG2]]] // CHECK: %[[V2:.+]] = memref.load %[[BUF]][%[[T1]]] -// CHECK: %[[COND:.+]] = cmpi sgt, %[[V1]], %[[V2]] : i32 +// CHECK: %[[COND:.+]] = arith.cmpi sgt, %[[V1]], %[[V2]] : i32 // CHECK: scf.if %[[COND]] { // CHECK: } else { -// CHECK: %[[T2:.+]] = addi %[[ARG2]], %[[C1]] : index +// CHECK: %[[T2:.+]] = arith.addi %[[ARG2]], %[[C1]] : index // CHECK: memref.store %[[V2]], %[[BUF]][%[[ARG2]]] // CHECK: memref.store %[[V1]], %[[BUF]][%[[T2]]] // CHECK: } @@ -34,28 +34,28 @@ linalg_ext.sort dimension(0) outs(%arg0 : memref<16x32xi32>) { ^bb0(%arg2: i32, %arg3: i32): // no predecessors - %0 = cmpi sgt, %arg2, %arg3 : i32 + %0 = arith.cmpi sgt, %arg2, %arg3 : i32 linalg_ext.yield %0 : i1 } return } // CHECK-LABEL: func @sort_2d // CHECK-SAME: %[[BUF:[a-zA-Z0-9]+]] -// CHECK-DAG: %[[C16:.+]] = constant 16 : index -// CHECK-DAG: %[[C32:.+]] = constant 32 : index -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index -// CHECK-DAG: %[[C15:.+]] = constant 15 : index +// CHECK-DAG: %[[C16:.+]] = arith.constant 16 : index +// CHECK-DAG: %[[C32:.+]] = arith.constant 32 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index +// CHECK-DAG: %[[C15:.+]] = arith.constant 15 : index // CHECK: scf.for %[[ARG1:.+]] = %[[C0]] to %[[C16]] step %[[C1]] // CHECK: scf.for %[[ARG2:.+]] = %[[C0]] to %[[C32]] step %[[C1]] // CHECK: scf.for %[[ARG3:.+]] = %[[C0]] to %[[C15]] step %[[C1]] -// CHECK: %[[T1:.+]] = addi %[[ARG3]], %[[C1]] : index +// CHECK: %[[T1:.+]] = arith.addi %[[ARG3]], %[[C1]] : index // CHECK: %[[V1:.+]] = memref.load %[[BUF]][%[[ARG3]], %[[ARG2]]] // CHECK: %[[V2:.+]] = memref.load %[[BUF]][%[[T1]], %[[ARG2]]] -// CHECK: %[[COND:.+]] = cmpi sgt, %[[V1]], %[[V2]] : i32 +// CHECK: %[[COND:.+]] = arith.cmpi sgt, %[[V1]], %[[V2]] : i32 // CHECK: scf.if %[[COND]] { // CHECK: } else { -// CHECK: %[[T2:.+]] = addi %[[ARG3]], %[[C1]] : index +// CHECK: %[[T2:.+]] = arith.addi %[[ARG3]], %[[C1]] : index // CHECK: memref.store %[[V2]], %[[BUF]][%[[ARG3]], %[[ARG2]]] // CHECK: memref.store %[[V1]], %[[BUF]][%[[T2]], %[[ARG2]]] // CHECK: } @@ -66,7 +66,7 @@ linalg_ext.sort outs(%arg0, %arg1 : memref<128xf32>, memref<128xi32>) { ^bb0(%arg2: f32, %arg3: f32, %arg4: i32, %arg5: i32): // no predecessors - %0 = cmpf ogt, %arg2, %arg3 : f32 + %0 = arith.cmpf ogt, %arg2, %arg3 : f32 linalg_ext.yield %0 : i1 } return @@ -74,21 +74,21 @@ // CHECK-LABEL: func @sort_multi // CHECK-SAME: %[[BUF1:[a-zA-Z0-9]+]] // CHECK-SAME: %[[BUF2:[a-zA-Z0-9]+]] -// CHECK-DAG: %[[C128:.+]] = constant 128 : index -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index -// CHECK-DAG: %[[C127:.+]] = constant 127 : index +// CHECK-DAG: %[[C128:.+]] = arith.constant 128 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index +// CHECK-DAG: %[[C127:.+]] = arith.constant 127 : index // CHECK: scf.for %[[ARG1:.+]] = %[[C0]] to %[[C128]] step %[[C1]] // CHECK: scf.for %[[ARG2:.+]] = %[[C0]] to %[[C127]] step %[[C1]] -// CHECK: %[[T1:.+]] = addi %[[ARG2]], %[[C1]] : index +// CHECK: %[[T1:.+]] = arith.addi %[[ARG2]], %[[C1]] : index // CHECK: %[[V1:.+]] = memref.load %[[BUF1]][%[[ARG2]]] // CHECK: %[[V2:.+]] = memref.load %[[BUF1]][%[[T1]]] // CHECK: %[[V3:.+]] = memref.load %[[BUF2]][%[[ARG2]]] // CHECK: %[[V4:.+]] = memref.load %[[BUF2]][%[[T1]]] -// CHECK: %[[COND:.+]] = cmpf ogt, %[[V1]], %[[V2]] : f32 +// CHECK: %[[COND:.+]] = arith.cmpf ogt, %[[V1]], %[[V2]] : f32 // CHECK: scf.if %[[COND]] { // CHECK: } else { -// CHECK: %[[T2:.+]] = addi %[[ARG2]], %[[C1]] : index +// CHECK: %[[T2:.+]] = arith.addi %[[ARG2]], %[[C1]] : index // CHECK: memref.store %[[V2]], %[[BUF1]][%[[ARG2]]] // CHECK: memref.store %[[V1]], %[[BUF1]][%[[T2]]] // CHECK: memref.store %[[V4]], %[[BUF2]][%[[ARG2]]] @@ -112,13 +112,13 @@ // CHECK-SAME: %[[ORIGINAL:[a-zA-Z0-9]+]] // CHECK-SAME: %[[INDICES:[a-zA-Z0-9]+]] // CHECK-SAME: %[[UPDATES:[a-zA-Z0-9]+]] -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index -// CHECK-DAG: %[[C3:.+]] = constant 3 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index +// CHECK-DAG: %[[C3:.+]] = arith.constant 3 : index // CHECK: scf.for %[[I:.+]] = %[[C0]] to %[[C3]] step %[[C1]] { // CHECK: %[[T1:.+]] = memref.load %[[UPDATES]][%[[I]]] : memref<3xi32> // CHECK: %[[T2:.+]] = memref.load %[[INDICES]][%[[I]], %[[C0]]] : memref<3x1xi32> -// CHECK: %[[IDX:.+]] = index_cast %[[T2]] : i32 to index +// CHECK: %[[IDX:.+]] = arith.index_cast %[[T2]] : i32 to index // CHECK: memref.store %[[T1]], %[[ORIGINAL]][%[[IDX]]] // ----- @@ -130,7 +130,7 @@ ins(%updates, %indices : memref<3xi32>, memref<3x2xi32>) outs(%original : memref<4x3xi32>) { ^bb0(%arg0: i32, %arg1: i32): // no predecessors - %0 = addi %arg1, %arg0 : i32 + %0 = arith.addi %arg1, %arg0 : i32 linalg_ext.yield %0 : i32 } return @@ -139,17 +139,17 @@ // CHECK-SAME: %[[ORIGINAL:[a-zA-Z0-9]+]] // CHECK-SAME: %[[INDICES:[a-zA-Z0-9]+]] // CHECK-SAME: %[[UPDATES:[a-zA-Z0-9]+]] -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index -// CHECK-DAG: %[[C3:.+]] = constant 3 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index +// CHECK-DAG: %[[C3:.+]] = arith.constant 3 : index // CHECK: scf.for %[[I:.+]] = %[[C0]] to %[[C3]] step %[[C1]] { // CHECK: %[[T1:.+]] = memref.load %[[UPDATES]][%[[I]]] : memref<3xi32> // CHECK: %[[T2:.+]] = memref.load %[[INDICES]][%[[I]], %[[C0]]] : memref<3x2xi32> -// CHECK: %[[IDX1:.+]] = index_cast %[[T2]] : i32 to index +// CHECK: %[[IDX1:.+]] = arith.index_cast %[[T2]] : i32 to index // CHECK: %[[T3:.+]] = memref.load %[[INDICES]][%[[I]], %[[C1]]] : memref<3x2xi32> -// CHECK: %[[IDX2:.+]] = index_cast %[[T3]] : i32 to index +// CHECK: %[[IDX2:.+]] = arith.index_cast %[[T3]] : i32 to index // CHECK: %[[ORI:.+]] = memref.load %[[ORIGINAL]][%[[IDX1]], %[[IDX2]]] : memref<4x3xi32> -// CHECK: %[[ADD:.+]] = addi %[[ORI]], %[[T1]] : i32 +// CHECK: %[[ADD:.+]] = arith.addi %[[ORI]], %[[T1]] : i32 // CHECK: memref.store %[[ADD]], %[[ORIGINAL]][%[[IDX1]], %[[IDX2]]] // ----- @@ -169,15 +169,15 @@ // CHECK-SAME: %[[ORIGINAL:[a-zA-Z0-9]+]] // CHECK-SAME: %[[INDICES:[a-zA-Z0-9]+]] // CHECK-SAME: %[[UPDATES:[a-zA-Z0-9]+]] -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index -// CHECK-DAG: %[[C2:.+]] = constant 2 : index -// CHECK-DAG: %[[C3:.+]] = constant 3 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index +// CHECK-DAG: %[[C2:.+]] = arith.constant 2 : index +// CHECK-DAG: %[[C3:.+]] = arith.constant 3 : index // CHECK: scf.for %[[I:.+]] = %[[C0]] to %[[C2]] step %[[C1]] { // CHECK: scf.for %[[J:.+]] = %[[C0]] to %[[C3]] step %[[C1]] { // CHECK: %[[UPDATE:.+]] = memref.load %[[UPDATES]][%[[I]], %[[J]]] // CHECK: %[[INDEX:.+]] = memref.load %[[INDICES]][%[[I]], %[[C0]]] -// CHECK: %[[LOC:.+]] = index_cast %[[INDEX]] : i32 to index +// CHECK: %[[LOC:.+]] = arith.index_cast %[[INDEX]] : i32 to index // CHECK: memref.store %[[UPDATE]], %[[ORIGINAL]][%[[LOC]], %[[J]]] // CHECK: } // CHECK: } @@ -191,7 +191,7 @@ ins(%updates, %indices : memref<3xi32>, memref<3x1xi32>) outs(%original : memref<8xi32>) { ^bb0(%arg0: i32, %arg1: i32): // no predecessors - %0 = addi %arg1, %arg0 : i32 + %0 = arith.addi %arg1, %arg0 : i32 linalg_ext.yield %0 : i32 } return @@ -200,15 +200,15 @@ // CHECK-SAME: %[[ORIGINAL:[a-zA-Z0-9]+]] // CHECK-SAME: %[[INDICES:[a-zA-Z0-9]+]] // CHECK-SAME: %[[UPDATES:[a-zA-Z0-9]+]] -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index -// CHECK-DAG: %[[C3:.+]] = constant 3 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index +// CHECK-DAG: %[[C3:.+]] = arith.constant 3 : index // CHECK: scf.for %[[I:.+]] = %[[C0]] to %[[C3]] step %[[C1]] { // CHECK: %[[T1:.+]] = memref.load %[[UPDATES]][%[[I]]] : memref<3xi32> // CHECK: %[[T2:.+]] = memref.load %[[INDICES]][%[[I]], %[[C0]]] : memref<3x1xi32> -// CHECK: %[[IDX:.+]] = index_cast %[[T2]] : i32 to index +// CHECK: %[[IDX:.+]] = arith.index_cast %[[T2]] : i32 to index // CHECK: %[[ORI:.+]] = memref.load %[[ORIGINAL]][%[[IDX]]] : memref<8xi32> -// CHECK: %[[ADD:.+]] = addi %[[ORI]], %[[T1]] : i32 +// CHECK: %[[ADD:.+]] = arith.addi %[[ORI]], %[[T1]] : i32 // CHECK: memref.store %[[ADD]], %[[ORIGINAL]][%[[IDX]]] // ----- @@ -220,7 +220,7 @@ ins(%updates, %indices : memref<2x3xi32>, memref<2x1xi32>) outs(%original : memref<4x3xi32>) { ^bb0(%arg0: i32, %arg1: i32): // no predecessors - %0 = addi %arg1, %arg0 : i32 + %0 = arith.addi %arg1, %arg0 : i32 linalg_ext.yield %0 : i32 } return @@ -229,16 +229,16 @@ // CHECK-SAME: %[[ORIGINAL:[a-zA-Z0-9]+]] // CHECK-SAME: %[[INDICES:[a-zA-Z0-9]+]] // CHECK-SAME: %[[UPDATES:[a-zA-Z0-9]+]] -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index -// CHECK-DAG: %[[C2:.+]] = constant 2 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index +// CHECK-DAG: %[[C2:.+]] = arith.constant 2 : index // CHECK: scf.for %[[I:.+]] = %[[C0]] to %[[C2]] step %[[C1]] { // CHECK: scf.for %[[J:.+]] = %[[C0]] to %[[C3]] step %[[C1]] { // CHECK: %[[UPDATEVAL:.+]] = memref.load %[[UPDATES]][%[[I]], %[[J]]] // CHECK: %[[INDEXVAL:.+]] = memref.load %[[INDICES]][%[[I]], %[[C0]]] -// CHECK: %[[INDEX:.+]] = index_cast %[[INDEXVAL]] : i32 to index +// CHECK: %[[INDEX:.+]] = arith.index_cast %[[INDEXVAL]] : i32 to index // CHECK: %[[ORIGINALVAL:.+]] = memref.load %[[ORIGINAL]][%[[INDEX]], %[[J]]] -// CHECK: %[[STOREVAL:.+]] = addi %[[ORIGINALVAL]], %[[UPDATEVAL]] +// CHECK: %[[STOREVAL:.+]] = arith.addi %[[ORIGINALVAL]], %[[UPDATEVAL]] // CHECK: memref.store %[[STOREVAL]], %[[ORIGINAL]][%[[INDEX]], %[[J]]] // ----- @@ -258,13 +258,13 @@ // CHECK-SAME: %[[ORIGINAL:[a-zA-Z0-9]+]] // CHECK-SAME: %[[INDICES:[a-zA-Z0-9]+]] // CHECK-SAME: %[[UPDATES:[a-zA-Z0-9]+]] -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK-DAG: %[[UB:.+]] = memref.dim %[[UPDATES]], %[[C0]] : memref<?xi32> // CHECK: scf.for %[[I:.+]] = %[[C0]] to %[[UB]] step %[[C1]] { // CHECK: %[[T1:.+]] = memref.load %[[UPDATES]][%[[I]]] : memref<?xi32> // CHECK: %[[T2:.+]] = memref.load %[[INDICES]][%[[I]], %[[C0]]] : memref<?x1xi32> -// CHECK: %[[IDX:.+]] = index_cast %[[T2]] : i32 to index +// CHECK: %[[IDX:.+]] = arith.index_cast %[[T2]] : i32 to index // CHECK: memref.store %[[T1]], %[[ORIGINAL]][%[[IDX]]] // ----- @@ -276,7 +276,7 @@ ins(%updates, %indices : memref<?xi32>, memref<?x2xi32>) outs(%original : memref<?x?xi32>) { ^bb0(%arg0: i32, %arg1: i32): // no predecessors - %0 = addi %arg1, %arg0 : i32 + %0 = arith.addi %arg1, %arg0 : i32 linalg_ext.yield %0 : i32 } return @@ -285,17 +285,17 @@ // CHECK-SAME: %[[ORIGINAL:[a-zA-Z0-9]+]] // CHECK-SAME: %[[INDICES:[a-zA-Z0-9]+]] // CHECK-SAME: %[[UPDATES:[a-zA-Z0-9]+]] -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK-DAG: %[[UB:.+]] = memref.dim %[[UPDATES]], %[[C0]] : memref<?xi32> // CHECK: scf.for %[[I:.+]] = %[[C0]] to %[[UB]] step %[[C1]] { // CHECK: %[[T1:.+]] = memref.load %[[UPDATES]][%[[I]]] : memref<?xi32> // CHECK: %[[T2:.+]] = memref.load %[[INDICES]][%[[I]], %[[C0]]] : memref<?x2xi32> -// CHECK: %[[IDX1:.+]] = index_cast %[[T2]] : i32 to index +// CHECK: %[[IDX1:.+]] = arith.index_cast %[[T2]] : i32 to index // CHECK: %[[T3:.+]] = memref.load %[[INDICES]][%[[I]], %[[C1]]] : memref<?x2xi32> -// CHECK: %[[IDX2:.+]] = index_cast %[[T3]] : i32 to index +// CHECK: %[[IDX2:.+]] = arith.index_cast %[[T3]] : i32 to index // CHECK: %[[ORI:.+]] = memref.load %[[ORIGINAL]][%[[IDX1]], %[[IDX2]]] : memref<?x?xi32> -// CHECK: %[[ADD:.+]] = addi %[[ORI]], %[[T1]] : i32 +// CHECK: %[[ADD:.+]] = arith.addi %[[ORI]], %[[T1]] : i32 // CHECK: memref.store %[[ADD]], %[[ORIGINAL]][%[[IDX1]], %[[IDX2]]] // ----- @@ -315,21 +315,21 @@ // CHECK-SAME: %[[ORIGINAL:[a-zA-Z0-9]+]] // CHECK-SAME: %[[INDICES:[a-zA-Z0-9]+]] // CHECK-SAME: %[[UPDATES:[a-zA-Z0-9]+]] -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK-DAG: %[[UB1:.+]] = memref.dim %[[UPDATES]], %[[C0]] : memref<?x?xi32> // CHECK-DAG: %[[UB2:.+]] = memref.dim %[[UPDATES]], %[[C1]] : memref<?x?xi32> // CHECK: scf.for %[[I:.+]] = %[[C0]] to %[[UB1]] step %[[C1]] { // CHECK: scf.for %[[J:.+]] = %[[C0]] to %[[UB2]] step %[[C1]] { // CHECK: %[[UPDATEVAL:.+]] = memref.load %[[UPDATES]][%[[I]], %[[J]]] // CHECK: %[[INDEXVAL:.+]] = memref.load %[[INDICES]][%[[I]], %[[C0]]] -// CHECK: %[[INDEX:.+]] = index_cast %[[INDEXVAL]] : i32 to index +// CHECK: %[[INDEX:.+]] = arith.index_cast %[[INDEXVAL]] : i32 to index // CHECK: memref.store %[[UPDATEVAL]], %[[ORIGINAL]][%[[INDEX]], %[[J]]] // ----- func @fft_1D(%real: memref<16xf32>, %imag: memref<16xf32>) { - %stage = constant 1 : index + %stage = arith.constant 1 : index linalg_ext.fft ins(%stage: index) outs(%real, %imag: memref<16xf32>, memref<16xf32>) @@ -340,22 +340,22 @@ // CHECK: func @fft_1D // CHECK-SAME: %[[REAL:[a-zA-Z0-9]+]] // CHECK-SAME: %[[IMAG:[a-zA-Z0-9]+]] -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index -// CHECK-DAG: %[[C16:.+]] = constant 16 : index -// CHECK-DAG: %[[SCALE:.+]] = constant -6.28318548 : f32 -// CHECK-DAG: %[[NODE_RNG:.+]] = shift_left %[[C1]], %[[C1]] : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index +// CHECK-DAG: %[[C16:.+]] = arith.constant 16 : index +// CHECK-DAG: %[[SCALE:.+]] = arith.constant -6.28318548 : f32 +// CHECK-DAG: %[[NODE_RNG:.+]] = arith.shli %[[C1]], %[[C1]] : index // CHECK: scf.for %[[K:.+]] = %[[C0]] to %[[C16]] step %[[NODE_RNG]] -// CHECK-DAG: %[[M:.+]] = shift_left %[[C1]], %[[C1]] : index -// CHECK-DAG: %[[HM:.+]] = shift_right_signed %[[M]], %[[C1]] : index +// CHECK-DAG: %[[M:.+]] = arith.shli %[[C1]], %[[C1]] : index +// CHECK-DAG: %[[HM:.+]] = arith.shrsi %[[M]], %[[C1]] : index // CHECK: %[[L_REAL_SLICE:.+]] = memref.subview %[[REAL]][%[[K]]] [%[[HM]]] [1] // CHECK: %[[L_IMAG_SLICE:.+]] = memref.subview %[[IMAG]][%[[K]]] [%[[HM]]] [1] -// CHECK: %[[R_OFFSET:.+]] = addi %[[K]], %[[HM]] : index +// CHECK: %[[R_OFFSET:.+]] = arith.addi %[[K]], %[[HM]] : index // CHECK: %[[R_REAL_SLICE:.+]] = memref.subview %[[REAL]][%[[R_OFFSET]]] [%[[HM]]] [1] // CHECK: %[[R_IMAG_SLICE:.+]] = memref.subview %[[IMAG]][%[[R_OFFSET]]] [%[[HM]]] [1] -// CHECK: %[[M_I32:.+]] = index_cast %[[M]] : index to i32 -// CHECK: %[[M_F32:.+]] = sitofp %[[M_I32]] : i32 to f32 -// CHECK: %[[COEFF:.+]] = divf %[[SCALE]], %[[M_F32]] +// CHECK: %[[M_I32:.+]] = arith.index_cast %[[M]] : index to i32 +// CHECK: %[[M_F32:.+]] = arith.sitofp %[[M_I32]] : i32 to f32 +// CHECK: %[[COEFF:.+]] = arith.divf %[[SCALE]], %[[M_F32]] // CHECK: linalg.generic // CHECK-SAME: indexing_maps = [#[[MAP1]], #[[MAP1]], #[[MAP1]], #[[MAP1]]] // CHECK-SAME: iterator_types = ["parallel"] @@ -364,35 +364,35 @@ // // Compute exp coeff. // CHECK: %[[J_IDX:.+]] = linalg.index 0 : index -// CHECK: %[[J_I32:.+]] = index_cast %[[J_IDX]] : index to i32 -// CHECK: %[[J_F32:.+]] = sitofp %[[J_I32]] : i32 to f32 -// CHECK: %[[EXP_COEF:.+]] = mulf %[[COEFF]], %[[J_F32]] : f32 +// CHECK: %[[J_I32:.+]] = arith.index_cast %[[J_IDX]] : index to i32 +// CHECK: %[[J_F32:.+]] = arith.sitofp %[[J_I32]] : i32 to f32 +// CHECK: %[[EXP_COEF:.+]] = arith.mulf %[[COEFF]], %[[J_F32]] : f32 // CHECK: %[[W_REAL:.+]] = math.cos %[[EXP_COEF]] // CHECK: %[[W_IMAG:.+]] = math.sin %[[EXP_COEF]] // // Compute "t = w * a[k + j + mh]" by expanding // (x + yi)(u + vi) = (xu - yv) + (xv + yu)i -// CHECK-DAG: %[[XU:.+]] = mulf %[[W_REAL]], %[[R_REAL]] -// CHECK-DAG: %[[YV:.+]] = mulf %[[W_IMAG]], %[[R_IMAG]] -// CHECK-DAG: %[[XV:.+]] = mulf %[[W_REAL]], %[[R_IMAG]] -// CHECK-DAG: %[[YU:.+]] = mulf %[[W_IMAG]], %[[R_REAL]] -// CHECK: %[[T_REAL:.+]] = subf %[[XU]], %[[YV]] -// CHECK: %[[T_IMAG:.+]] = addf %[[XV]], %[[YU]] +// CHECK-DAG: %[[XU:.+]] = arith.mulf %[[W_REAL]], %[[R_REAL]] +// CHECK-DAG: %[[YV:.+]] = arith.mulf %[[W_IMAG]], %[[R_IMAG]] +// CHECK-DAG: %[[XV:.+]] = arith.mulf %[[W_REAL]], %[[R_IMAG]] +// CHECK-DAG: %[[YU:.+]] = arith.mulf %[[W_IMAG]], %[[R_REAL]] +// CHECK: %[[T_REAL:.+]] = arith.subf %[[XU]], %[[YV]] +// CHECK: %[[T_IMAG:.+]] = arith.addf %[[XV]], %[[YU]] // // Compute the results. // u = a[k + j]; // a[k + j] = u + t; // a[k + j + mh] = u - t; -// CHECK: %[[RES1:.+]] = addf %[[L_REAL]], %[[T_REAL]] -// CHECK: %[[RES2:.+]] = addf %[[L_IMAG]], %[[T_IMAG]] -// CHECK: %[[RES3:.+]] = subf %[[L_REAL]], %[[T_REAL]] -// CHECK: %[[RES4:.+]] = subf %[[L_IMAG]], %[[T_IMAG]] +// CHECK: %[[RES1:.+]] = arith.addf %[[L_REAL]], %[[T_REAL]] +// CHECK: %[[RES2:.+]] = arith.addf %[[L_IMAG]], %[[T_IMAG]] +// CHECK: %[[RES3:.+]] = arith.subf %[[L_REAL]], %[[T_REAL]] +// CHECK: %[[RES4:.+]] = arith.subf %[[L_IMAG]], %[[T_IMAG]] // CHECK: linalg.yield %[[RES1]], %[[RES2]], %[[RES3]], %[[RES4]] // ----- func @fft_2D(%real: memref<?x16xf32>, %imag: memref<?x16xf32>) { - %stage = constant 2 : index + %stage = arith.constant 2 : index linalg_ext.fft ins(%stage: index) outs(%real, %imag: memref<?x16xf32>, memref<?x16xf32>) @@ -403,18 +403,18 @@ // CHECK: func @fft_2D // CHECK-SAME: %[[REAL:[a-zA-Z0-9]+]] // CHECK-SAME: %[[IMAG:[a-zA-Z0-9]+]] -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index -// CHECK-DAG: %[[C2:.+]] = constant 2 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index +// CHECK-DAG: %[[C2:.+]] = arith.constant 2 : index // CHECK-DAG: %[[D0:.+]] = memref.dim %[[REAL]], %[[C0]] : memref<?x16xf32> -// CHECK-DAG: %[[NODE_RNG:.+]] = shift_left %[[C1]], %[[C2]] : index +// CHECK-DAG: %[[NODE_RNG:.+]] = arith.shli %[[C1]], %[[C2]] : index // CHECK: scf.for %[[I:.+]] = %[[C0]] to %[[D0]] step %[[C1]] // CHECK: scf.for %[[K:.+]] = %[[C0]] to %[[C16]] step %[[NODE_RNG]] -// CHECK-DAG: %[[M:.+]] = shift_left %[[C1]], %[[C2]] : index -// CHECK-DAG: %[[HM:.+]] = shift_right_signed %[[M]], %[[C1]] : index +// CHECK-DAG: %[[M:.+]] = arith.shli %[[C1]], %[[C2]] : index +// CHECK-DAG: %[[HM:.+]] = arith.shrsi %[[M]], %[[C1]] : index // CHECK: %[[L_REAL_SLICE:.+]] = memref.subview %[[REAL]][%[[I]], %[[K]]] [1, %[[HM]]] [1, 1] // CHECK: %[[L_IMAG_SLICE:.+]] = memref.subview %[[IMAG]][%[[I]], %[[K]]] [1, %[[HM]]] [1, 1] -// CHECK: %[[R_OFFSET:.+]] = addi %[[K]], %[[HM]] : index +// CHECK: %[[R_OFFSET:.+]] = arith.addi %[[K]], %[[HM]] : index // CHECK: %[[R_REAL_SLICE:.+]] = memref.subview %[[REAL]][%[[I]], %[[R_OFFSET]]] [1, %[[HM]]] [1, 1] // CHECK: %[[R_IMAG_SLICE:.+]] = memref.subview %[[IMAG]][%[[I]], %[[R_OFFSET]]] [1, %[[HM]]] [1, 1] // CHECK: linalg.generic @@ -430,7 +430,7 @@ func @fft_2D_coef_buf(%real: memref<?x16xf32>, %imag: memref<?x16xf32>, %coef_real: memref<1xf32>, %coef_imag: memref<1xf32>) { - %stage = constant 1 : index + %stage = arith.constant 1 : index linalg_ext.fft ins(%stage, %coef_real, %coef_imag: index, memref<1xf32>, memref<1xf32>) outs(%real, %imag: memref<?x16xf32>, memref<?x16xf32>) @@ -444,17 +444,17 @@ // CHECK-SAME: %[[IMAG:[a-zA-Z0-9]+]] // CHECK-SAME: %[[COEF_REAL:[a-zA-Z0-9]+]] // CHECK-SAME: %[[COEF_IMAG:[a-zA-Z0-9]+]] -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK-DAG: %[[D0:.+]] = memref.dim %[[REAL]], %[[C0]] : memref<?x16xf32> -// CHECK-DAG: %[[NODE_RNG:.+]] = shift_left %[[C1]], %[[C1]] : index +// CHECK-DAG: %[[NODE_RNG:.+]] = arith.shli %[[C1]], %[[C1]] : index // CHECK: scf.for %[[I:.+]] = %[[C0]] to %[[D0]] step %[[C1]] // CHECK: scf.for %[[K:.+]] = %[[C0]] to %[[C16]] step %[[NODE_RNG]] -// CHECK-DAG: %[[M:.+]] = shift_left %[[C1]], %[[C1]] : index -// CHECK-DAG: %[[HM:.+]] = shift_right_signed %[[M]], %[[C1]] : index +// CHECK-DAG: %[[M:.+]] = arith.shli %[[C1]], %[[C1]] : index +// CHECK-DAG: %[[HM:.+]] = arith.shrsi %[[M]], %[[C1]] : index // CHECK: %[[L_REAL_SLICE:.+]] = memref.subview %[[REAL]][%[[I]], %[[K]]] [1, %[[HM]]] [1, 1] // CHECK: %[[L_IMAG_SLICE:.+]] = memref.subview %[[IMAG]][%[[I]], %[[K]]] [1, %[[HM]]] [1, 1] -// CHECK: %[[R_OFFSET:.+]] = addi %[[K]], %[[HM]] : index +// CHECK: %[[R_OFFSET:.+]] = arith.addi %[[K]], %[[HM]] : index // CHECK: %[[R_REAL_SLICE:.+]] = memref.subview %[[REAL]][%[[I]], %[[R_OFFSET]]] [1, %[[HM]]] [1, 1] // CHECK: %[[R_IMAG_SLICE:.+]] = memref.subview %[[IMAG]][%[[I]], %[[R_OFFSET]]] [1, %[[HM]]] [1, 1] // CHECK: linalg.generic @@ -465,21 +465,21 @@ // CHECK: ^bb0(%[[W_REAL:.+]]: f32, %[[W_IMAG:.+]]: f32, %[[L_REAL:.+]]: f32, %[[L_IMAG:.+]]: f32, %[[R_REAL:.+]]: f32, %[[R_IMAG:.+]]: f32) // Compute "t = w * a[k + j + mh]" by expanding // (x + yi)(u + vi) = (xu - yv) + (xv + yu)i -// CHECK-DAG: %[[XU:.+]] = mulf %[[W_REAL]], %[[R_REAL]] -// CHECK-DAG: %[[YV:.+]] = mulf %[[W_IMAG]], %[[R_IMAG]] -// CHECK-DAG: %[[XV:.+]] = mulf %[[W_REAL]], %[[R_IMAG]] -// CHECK-DAG: %[[YU:.+]] = mulf %[[W_IMAG]], %[[R_REAL]] -// CHECK: %[[T_REAL:.+]] = subf %[[XU]], %[[YV]] -// CHECK: %[[T_IMAG:.+]] = addf %[[XV]], %[[YU]] +// CHECK-DAG: %[[XU:.+]] = arith.mulf %[[W_REAL]], %[[R_REAL]] +// CHECK-DAG: %[[YV:.+]] = arith.mulf %[[W_IMAG]], %[[R_IMAG]] +// CHECK-DAG: %[[XV:.+]] = arith.mulf %[[W_REAL]], %[[R_IMAG]] +// CHECK-DAG: %[[YU:.+]] = arith.mulf %[[W_IMAG]], %[[R_REAL]] +// CHECK: %[[T_REAL:.+]] = arith.subf %[[XU]], %[[YV]] +// CHECK: %[[T_IMAG:.+]] = arith.addf %[[XV]], %[[YU]] // // Compute the results. // u = a[k + j]; // a[k + j] = u + t; // a[k + j + mh] = u - t; -// CHECK: %[[RES1:.+]] = addf %[[L_REAL]], %[[T_REAL]] -// CHECK: %[[RES2:.+]] = addf %[[L_IMAG]], %[[T_IMAG]] -// CHECK: %[[RES3:.+]] = subf %[[L_REAL]], %[[T_REAL]] -// CHECK: %[[RES4:.+]] = subf %[[L_IMAG]], %[[T_IMAG]] +// CHECK: %[[RES1:.+]] = arith.addf %[[L_REAL]], %[[T_REAL]] +// CHECK: %[[RES2:.+]] = arith.addf %[[L_IMAG]], %[[T_IMAG]] +// CHECK: %[[RES3:.+]] = arith.subf %[[L_REAL]], %[[T_REAL]] +// CHECK: %[[RES4:.+]] = arith.subf %[[L_IMAG]], %[[T_IMAG]] // CHECK: linalg.yield %[[RES1]], %[[RES2]], %[[RES3]], %[[RES4]] // ----- @@ -494,14 +494,14 @@ // CHECK-LABEL: func @reverse_dim_0 // CHECK-SAME: %[[IN:[a-zA-Z0-9]+]] // CHECK-SAME: %[[OUT:[a-zA-Z0-9]+]] -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK-DAG: %[[D0:.+]] = memref.dim %arg0, %c0 : memref<?x?xi32> // CHECK-DAG: %[[D1:.+]] = memref.dim %arg0, %c1 : memref<?x?xi32> // CHECK: scf.for %[[I:.+]] = %[[C0]] to %[[D0]] step %[[C1]] // CHECK: scf.for %[[J:.+]] = %[[C0]] to %[[D1]] step %[[C1]] // CHECK: %[[T0:.+]] = memref.dim %[[IN]], %[[C0]] -// CHECK: %[[T1:.+]] = subi %[[T0]], %[[C1]] : index -// CHECK: %[[T2:.+]] = subi %[[T1]], %[[I]] : index +// CHECK: %[[T1:.+]] = arith.subi %[[T0]], %[[C1]] : index +// CHECK: %[[T2:.+]] = arith.subi %[[T1]], %[[I]] : index // CHECK: %[[V0:.+]] = memref.load %[[IN]][%[[I]], %[[J]]] // CHECK: memref.store %[[V0]], %[[OUT]][%[[T2]], %[[J]]] : memref<?x?xi32>
diff --git a/iree/compiler/Dialect/LinalgExt/Transforms/test/tiling.mlir b/iree/compiler/Dialect/LinalgExt/Transforms/test/tiling.mlir index f632cf6..f269e75 100644 --- a/iree/compiler/Dialect/LinalgExt/Transforms/test/tiling.mlir +++ b/iree/compiler/Dialect/LinalgExt/Transforms/test/tiling.mlir
@@ -8,7 +8,7 @@ ins(%update, %indices : tensor<?x?xf32>, tensor<?x1xi32>) outs(%original : tensor<?x?xf32>) { ^bb0(%arg1: f32, %arg2: f32): - %1 = addf %arg1, %arg2 : f32 + %1 = arith.addf %arg1, %arg2 : f32 linalg_ext.yield %1 : f32 } -> tensor<?x?xf32> return %0 : tensor<?x?xf32> @@ -19,10 +19,10 @@ // CHECK-SAME: %[[ORIGINAL:[a-zA-Z0-9_]+]]: tensor<?x?xf32> // CHECK-SAME: %[[INDICES:[a-zA-Z0-9_]+]]: tensor<?x1xi32> // CHECK-SAME: %[[UPDATES:[a-zA-Z0-9_]+]]: tensor<?x?xf32> -// CHECK-DAG: %[[TILESIZEY:.+]] = constant 10 : index -// CHECK-DAG: %[[TILESIZEX:.+]] = constant 20 : index -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[TILESIZEY:.+]] = arith.constant 10 : index +// CHECK-DAG: %[[TILESIZEX:.+]] = arith.constant 20 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK-DAG: %[[D0:.+]] = tensor.dim %[[UPDATES]], %[[C0]] // CHECK-DAG: %[[D1:.+]] = tensor.dim %[[UPDATES]], %[[C1]] // CHECK: %[[RESULT:.+]] = scf.for %[[IV0:.+]] = %[[C0]] to %[[D0]] step %[[TILESIZEY]] @@ -58,7 +58,7 @@ ins(%update, %indices : memref<?x?xf32>, memref<?x1xi32>) outs(%original : memref<?x?xf32>) { ^bb0(%arg1: f32, %arg2: f32): - %1 = addf %arg1, %arg2 : f32 + %1 = arith.addf %arg1, %arg2 : f32 linalg_ext.yield %1 : f32 } return @@ -69,10 +69,10 @@ // CHECK-SAME: %[[ORIGINAL:[a-zA-Z0-9_]+]]: memref<?x?xf32> // CHECK-SAME: %[[INDICES:[a-zA-Z0-9_]+]]: memref<?x1xi32> // CHECK-SAME: %[[UPDATES:[a-zA-Z0-9_]+]]: memref<?x?xf32> -// CHECK-DAG: %[[TILESIZEY:.+]] = constant 10 : index -// CHECK-DAG: %[[TILESIZEX:.+]] = constant 20 : index -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[TILESIZEY:.+]] = arith.constant 10 : index +// CHECK-DAG: %[[TILESIZEX:.+]] = arith.constant 20 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK-DAG: %[[D0:.+]] = memref.dim %[[UPDATES]], %[[C0]] // CHECK-DAG: %[[D1:.+]] = memref.dim %[[UPDATES]], %[[C1]] // CHECK: scf.for %[[IV0:.+]] = %[[C0]] to %[[D0]] step %[[TILESIZEY]] @@ -101,7 +101,7 @@ ins(%update, %indices : tensor<?x?xf32>, tensor<?x1xi32>) outs(%original : tensor<?x?xf32>) { ^bb0(%arg1: f32, %arg2: f32): - %1 = addf %arg1, %arg2 : f32 + %1 = arith.addf %arg1, %arg2 : f32 linalg_ext.yield %1 : f32 } -> tensor<?x?xf32> return %0 : tensor<?x?xf32> @@ -112,9 +112,9 @@ // CHECK-SAME: %[[ORIGINAL:[a-zA-Z0-9_]+]]: tensor<?x?xf32> // CHECK-SAME: %[[INDICES:[a-zA-Z0-9_]+]]: tensor<?x1xi32> // CHECK-SAME: %[[UPDATES:[a-zA-Z0-9_]+]]: tensor<?x?xf32> -// CHECK-DAG: %[[C1:.+]] = constant 1 : index -// CHECK-DAG: %[[TILESIZE:.+]] = constant 10 : index -// CHECK-DAG: %[[C0:.+]] = constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index +// CHECK-DAG: %[[TILESIZE:.+]] = arith.constant 10 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index // CHECK-DAG: %[[D0:.+]] = tensor.dim %[[UPDATES]], %[[C0]] // CHECK-DAG: %[[D1:.+]] = tensor.dim %[[UPDATES]], %[[C1]] // CHECK-DAG: %[[ID:.+]] = flow.dispatch.workgroup.id[0] @@ -149,7 +149,7 @@ ins(%update, %indices : tensor<?x?xf32>, tensor<?x1xi32>) outs(%original : tensor<?x?xf32>) { ^bb0(%arg1: f32, %arg2: f32): - %1 = addf %arg1, %arg2 : f32 + %1 = arith.addf %arg1, %arg2 : f32 linalg_ext.yield %1 : f32 } -> tensor<?x?xf32> return %0 : tensor<?x?xf32> @@ -171,7 +171,7 @@ {__internal_linalg_transform__ = "outer_reduce_input"} outs(%arg0 : tensor<?xi32>) { ^bb0(%arg2: i32, %arg3: i32): // no predecessors - %0 = cmpi sgt, %arg2, %arg3 : i32 + %0 = arith.cmpi sgt, %arg2, %arg3 : i32 linalg_ext.yield %0 : i1 } -> tensor<?xi32> return %0 : tensor<?xi32> @@ -190,7 +190,7 @@ {__internal_linalg_transform__ = "inner_reduce_input"} outs(%arg0 : tensor<?x?xi32>) { ^bb0(%arg2: i32, %arg3: i32): // no predecessors - %0 = cmpi sgt, %arg2, %arg3 : i32 + %0 = arith.cmpi sgt, %arg2, %arg3 : i32 linalg_ext.yield %0 : i1 } -> tensor<?x?xi32> return %0 : tensor<?x?xi32> @@ -198,9 +198,9 @@ // CHECK: #[[MAP:.+]] = affine_map<(d0)[s0, s1] -> (10, -d0 + s1)> // CHECK: func @sort_2d( // CHECK-SAME: %[[OPERAND:.+]]: tensor<?x?xi32> -// CHECK-DAG: %[[TILESIZE:.+]] = constant 10 : index -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[TILESIZE:.+]] = arith.constant 10 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK-DAG: %[[D0:.+]] = tensor.dim %[[OPERAND]], %[[C0]] // CHECK-DAG: %[[D1:.+]] = tensor.dim %[[OPERAND]], %[[C1]] // CHECK: %[[RESULT:.+]] = scf.for %[[IV:.+]] = %[[C0]] to %[[D0]] step %[[TILESIZE]] @@ -223,7 +223,7 @@ {__internal_linalg_transform__ = "outer_reduce_input"} outs(%arg0 : tensor<?x?xi32>) { ^bb0(%arg2: i32, %arg3: i32): // no predecessors - %0 = cmpi sgt, %arg2, %arg3 : i32 + %0 = arith.cmpi sgt, %arg2, %arg3 : i32 linalg_ext.yield %0 : i1 } -> tensor<?x?xi32> return %0 : tensor<?x?xi32> @@ -231,9 +231,9 @@ // CHECK: #[[MAP:.+]] = affine_map<(d0)[s0, s1] -> (20, -d0 + s1)> // CHECK: func @sort_2d_inner_parallel( // CHECK-SAME: %[[OPERAND:.+]]: tensor<?x?xi32> -// CHECK-DAG: %[[TILESIZE:.+]] = constant 20 : index -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[TILESIZE:.+]] = arith.constant 20 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK-DAG: %[[D0:.+]] = tensor.dim %[[OPERAND]], %[[C0]] // CHECK-DAG: %[[D1:.+]] = tensor.dim %[[OPERAND]], %[[C1]] // CHECK: %[[RESULT:.+]] = scf.for %[[IV:.+]] = %[[C0]] to %[[D1]] step %[[TILESIZE]] @@ -258,7 +258,7 @@ {__internal_linalg_transform__ = "inner_reduce_input"} outs(%arg0, %arg1 : tensor<?x?xi32>, tensor<?x?xf32>) { ^bb0(%arg2: i32, %arg3: i32, %arg4 : f32, %arg5 : f32): // no predecessors - %1 = cmpf ogt, %arg4, %arg5 : f32 + %1 = arith.cmpf ogt, %arg4, %arg5 : f32 linalg_ext.yield %1 : i1 } -> tensor<?x?xi32>, tensor<?x?xf32> return %0#0, %0#1 : tensor<?x?xi32>, tensor<?x?xf32> @@ -267,9 +267,9 @@ // CHECK: func @sort_2d_multi_result( // CHECK-SAME: %[[OPERAND1:.+]]: tensor<?x?xi32> // CHECK-SAME: %[[OPERAND2:.+]]: tensor<?x?xf32> -// CHECK-DAG: %[[TILESIZE:.+]] = constant 10 : index -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[TILESIZE:.+]] = arith.constant 10 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK-DAG: %[[D0:.+]] = tensor.dim %[[OPERAND1]], %[[C0]] // CHECK-DAG: %[[D1:.+]] = tensor.dim %[[OPERAND1]], %[[C1]] // CHECK: %[[RESULT:.+]]:2 = scf.for %[[IV:.+]] = %[[C0]] to %[[D0]] step %[[TILESIZE]] @@ -297,7 +297,7 @@ {__internal_linalg_transform__ = "outer_reduce_input"} outs(%arg0, %arg1 : memref<?x?xi32>, memref<?x?xf32>) { ^bb0(%arg2: i32, %arg3: i32, %arg4 : f32, %arg5 : f32): // no predecessors - %0 = cmpf ogt, %arg4, %arg5 : f32 + %0 = arith.cmpf ogt, %arg4, %arg5 : f32 linalg_ext.yield %0 : i1 } return @@ -306,9 +306,9 @@ // CHECK: func @sort_2d_multi_result_memref( // CHECK-SAME: %[[OPERAND1:.+]]: memref<?x?xi32> // CHECK-SAME: %[[OPERAND2:.+]]: memref<?x?xf32> -// CHECK-DAG: %[[TILESIZE:.+]] = constant 20 : index -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[TILESIZE:.+]] = arith.constant 20 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK-DAG: %[[D0:.+]] = memref.dim %[[OPERAND1]], %[[C0]] // CHECK-DAG: %[[D1:.+]] = memref.dim %[[OPERAND1]], %[[C1]] // CHECK: scf.for %[[IV:.+]] = %[[C0]] to %[[D1]] step %[[TILESIZE]] @@ -330,7 +330,7 @@ {__internal_linalg_transform__ = "distribute_input"} outs(%arg0, %arg1 : tensor<?x?x?xi32>, tensor<?x?x?xf32>) { ^bb0(%arg2: i32, %arg3: i32, %arg4 : f32, %arg5 : f32): // no predecessors - %2 = cmpf ogt, %arg4, %arg5 : f32 + %2 = arith.cmpf ogt, %arg4, %arg5 : f32 linalg_ext.yield %2 : i1 } -> tensor<?x?x?xi32>, tensor<?x?x?xf32> return %0, %1 : tensor<?x?x?xi32>, tensor<?x?x?xf32> @@ -342,11 +342,11 @@ // CHECK: func @sort_3d_multi_result_distribute( // CHECK-SAME: %[[OPERAND1:[a-zA-Z0-9_]+]]: tensor<?x?x?xi32> // CHECK-SAME: %[[OPERAND2:[a-zA-Z0-9_]+]]: tensor<?x?x?xf32> -// CHECK-DAG: %[[TILESIZE1:.+]] = constant 10 : index -// CHECK-DAG: %[[TILESIZE2:.+]] = constant 30 : index -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index -// CHECK-DAG: %[[C2:.+]] = constant 2 : index +// CHECK-DAG: %[[TILESIZE1:.+]] = arith.constant 10 : index +// CHECK-DAG: %[[TILESIZE2:.+]] = arith.constant 30 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index +// CHECK-DAG: %[[C2:.+]] = arith.constant 2 : index // CHECK-DAG: %[[D0:.+]] = tensor.dim %[[OPERAND1]], %[[C0]] // CHECK-DAG: %[[D1:.+]] = tensor.dim %[[OPERAND1]], %[[C1]] // CHECK-DAG: %[[D2:.+]] = tensor.dim %[[OPERAND1]], %[[C2]] @@ -387,7 +387,7 @@ {__internal_linalg_transform__ = "distribute_input"} outs(%arg0, %arg1 : memref<?x?x?xi32>, memref<?x?x?xf32>) { ^bb0(%arg2: i32, %arg3: i32, %arg4 : f32, %arg5 : f32): // no predecessors - %0 = cmpf ogt, %arg4, %arg5 : f32 + %0 = arith.cmpf ogt, %arg4, %arg5 : f32 linalg_ext.yield %0 : i1 } return @@ -399,11 +399,11 @@ // CHECK: func @sort_3d_multi_result_distribute_memref( // CHECK-SAME: %[[OPERAND1:[a-zA-Z0-9_]+]]: memref<?x?x?xi32> // CHECK-SAME: %[[OPERAND2:[a-zA-Z0-9_]+]]: memref<?x?x?xf32> -// CHECK-DAG: %[[TILESIZE1:.+]] = constant 10 : index -// CHECK-DAG: %[[TILESIZE2:.+]] = constant 30 : index -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index -// CHECK-DAG: %[[C2:.+]] = constant 2 : index +// CHECK-DAG: %[[TILESIZE1:.+]] = arith.constant 10 : index +// CHECK-DAG: %[[TILESIZE2:.+]] = arith.constant 30 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index +// CHECK-DAG: %[[C2:.+]] = arith.constant 2 : index // CHECK-DAG: %[[D0:.+]] = memref.dim %[[OPERAND1]], %[[C0]] // CHECK-DAG: %[[D1:.+]] = memref.dim %[[OPERAND1]], %[[C1]] // CHECK-DAG: %[[D2:.+]] = memref.dim %[[OPERAND1]], %[[C2]] @@ -431,8 +431,8 @@ func @slice_insert(%source :tensor<?x?xf32>, %dest: tensor<?x?xf32>, %idx0 : index, %idx1 : index) -> tensor<?x?xf32> { - %c0 = constant 0 : index - %c1 = constant 1 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index %0 = tensor.dim %source, %c0 : tensor<?x?xf32> %1 = tensor.dim %source, %c1 : tensor<?x?xf32> %2 = tensor.insert_slice %source into %dest[%idx0, %idx1] [%0, %1] [1, 1] @@ -462,8 +462,8 @@ func @slice_insert_rank_reduce(%source :tensor<?x?xf32>, %dest: tensor<?x?x?xf32>, %idx0 : index, %idx1 : index) -> tensor<?x?x?xf32> { - %c0 = constant 0 : index - %c1 = constant 1 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index %0 = tensor.dim %source, %c0 : tensor<?x?xf32> %1 = tensor.dim %source, %c1 : tensor<?x?xf32> %2 = tensor.insert_slice %source into %dest[%idx0, 0, %idx1] [%0, 1, %1] [1, 1, 1] @@ -493,7 +493,7 @@ func @fft_1d_stage_5(%arg0: tensor<1024xf32>, %arg1: tensor<1024xf32>, %arg2: tensor<16xf32>, %arg3: tensor<16xf32>) -> (tensor<1024xf32>, tensor<1024xf32>) { - %cst1 = constant 5 : index + %cst1 = arith.constant 5 : index %0:2 = linalg_ext.fft {__internal_linalg_transform__ = "tiling_1d_stage5_fft_input"} ins(%cst1, %arg2, %arg3: index, tensor<16xf32>, tensor<16xf32>) @@ -507,10 +507,10 @@ // CHECK-SAME: %[[ARG1:[a-zA-Z0-9_]+]] // CHECK-SAME: %[[COEF_REAL:[a-zA-Z0-9_]+]] // CHECK-SAME: %[[COEF_IMAG:[a-zA-Z0-9_]+]] -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C5:.+]] = constant 5 : index -// CHECK-DAG: %[[C32:.+]] = constant 32 : index -// CHECK-DAG: %[[C1024:.+]] = constant 1024 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C5:.+]] = arith.constant 5 : index +// CHECK-DAG: %[[C32:.+]] = arith.constant 32 : index +// CHECK-DAG: %[[C1024:.+]] = arith.constant 1024 : index // CHECK: %[[RES:.+]]:2 = scf.for %[[I:.+]] = %[[C0]] to %[[C1024]] step %[[C32]] // CHECK-SAME: iter_args(%[[ARG5:.+]] = %[[ARG0]], %[[ARG6:.+]] = %[[ARG1]]) // CHECK-SAME: -> (tensor<1024xf32>, tensor<1024xf32>) { @@ -530,7 +530,7 @@ func @fft_2d_stage_5(%arg0: tensor<3x1024xf32>, %arg1: tensor<3x1024xf32>, %arg2: tensor<16xf32>, %arg3: tensor<16xf32>) -> (tensor<3x1024xf32>, tensor<3x1024xf32>) { - %cst1 = constant 5 : index + %cst1 = arith.constant 5 : index %0:2 = linalg_ext.fft {__internal_linalg_transform__ = "tiling_2d_stage5_fft_input"} ins(%cst1, %arg2, %arg3: index, tensor<16xf32>, tensor<16xf32>) @@ -545,12 +545,12 @@ // CHECK-SAME: %[[ARG1:[a-zA-Z0-9_]+]] // CHECK-SAME: %[[COEF_REAL:[a-zA-Z0-9_]+]] // CHECK-SAME: %[[COEF_IMAG:[a-zA-Z0-9_]+]] -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C3:.+]] = constant 3 : index -// CHECK-DAG: %[[C5:.+]] = constant 5 : index -// CHECK-DAG: %[[C10:.+]] = constant 10 : index -// CHECK-DAG: %[[C32:.+]] = constant 32 : index -// CHECK-DAG: %[[C1024:.+]] = constant 1024 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C3:.+]] = arith.constant 3 : index +// CHECK-DAG: %[[C5:.+]] = arith.constant 5 : index +// CHECK-DAG: %[[C10:.+]] = arith.constant 10 : index +// CHECK-DAG: %[[C32:.+]] = arith.constant 32 : index +// CHECK-DAG: %[[C1024:.+]] = arith.constant 1024 : index // CHECK: %[[RES:.+]]:2 = scf.for %[[I:.+]] = %[[C0]] to %[[C3]] step %[[C10]] // CHECK-SAME: iter_args(%[[ARG5:.+]] = %[[ARG0]], %[[ARG6:.+]] = %[[ARG1]]) // CHECK-SAME: -> (tensor<3x1024xf32>, tensor<3x1024xf32>) { @@ -572,7 +572,7 @@ func @fft_1d_stage_5_memref(%arg0: memref<1024xf32>, %arg1: memref<1024xf32>, %arg2: memref<16xf32>, %arg3: memref<16xf32>) { - %cst1 = constant 5 : index + %cst1 = arith.constant 5 : index linalg_ext.fft {__internal_linalg_transform__ = "tiling_1d_stage5_fft_input"} ins(%cst1, %arg2, %arg3: index, memref<16xf32>, memref<16xf32>) @@ -586,10 +586,10 @@ // CHECK-SAME: %[[ARG1:[a-zA-Z0-9_]+]] // CHECK-SAME: %[[COEF_REAL:[a-zA-Z0-9_]+]] // CHECK-SAME: %[[COEF_IMAG:[a-zA-Z0-9_]+]] -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C5:.+]] = constant 5 : index -// CHECK-DAG: %[[C32:.+]] = constant 32 : index -// CHECK-DAG: %[[C1024:.+]] = constant 1024 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C5:.+]] = arith.constant 5 : index +// CHECK-DAG: %[[C32:.+]] = arith.constant 32 : index +// CHECK-DAG: %[[C1024:.+]] = arith.constant 1024 : index // CHECK: scf.for %[[I:.+]] = %[[C0]] to %[[C1024]] step %[[C32]] { // CHECK: %[[SZ:.+]] = affine.min #[[MAP0]](%[[I]])[%[[C32]], %[[C1024]]] // CHECK: %[[SUB1:.+]] = memref.subview %[[ARG0]][%[[I]]] [%[[SZ]]] [1] : memref<1024xf32> to memref<?xf32, #[[MAP1]]> @@ -615,8 +615,8 @@ // CHECK: func @reverse_memref( // CHECK-SAME: %[[ARG0:[a-zA-Z0-9_]+]] // CHECK-SAME: %[[ARG1:[a-zA-Z0-9_]+]] -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C10:.+]] = constant 10 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C10:.+]] = arith.constant 10 : index // CHECK-DAG: %[[D0:.+]] = memref.dim %[[ARG0]], %[[C0]] : memref<?xi32> // CHECK: scf.for %[[I:.+]] = %[[C0]] to %[[D0]] step %[[C10]] { // CHECK: %[[SIZE:.+]] = affine.min #[[MAP0]](%[[I]])[%[[C10]], %[[D0]]] @@ -633,8 +633,8 @@ // ----- func @reverse_tensor_multi_dim(%arg0: tensor<?x?xi32>) -> tensor<?x?xi32> { - %c0 = constant 0 : index - %c1 = constant 1 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index %d0 = tensor.dim %arg0, %c0 : tensor<?x?xi32> %d1 = tensor.dim %arg0, %c1 : tensor<?x?xi32> %init = linalg.init_tensor [%d0, %d1] : tensor<?x?xi32> @@ -650,10 +650,10 @@ // CHECK-DAG: #[[MAP2:.+]] = affine_map<()[s0, s1, s2] -> (s0 - s1 - s2)> // CHECK: func @reverse_tensor_multi_dim( // CHECK-SAME: %[[ARG0:[a-zA-Z0-9_]+]] -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index -// CHECK-DAG: %[[C10:.+]] = constant 10 : index -// CHECK-DAG: %[[C20:.+]] = constant 20 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index +// CHECK-DAG: %[[C10:.+]] = arith.constant 10 : index +// CHECK-DAG: %[[C20:.+]] = arith.constant 20 : index // CHECK-DAG: %[[D0:.+]] = tensor.dim %[[ARG0]], %[[C0]] : tensor<?x?xi32> // CHECK-DAG: %[[D1:.+]] = tensor.dim %[[ARG0]], %[[C1]] : tensor<?x?xi32> // CHECK: %[[INIT:.+]] = linalg.init_tensor [%[[D0]], %[[D1]]] : tensor<?x?xi32> @@ -688,7 +688,7 @@ func @dynamic_insert_slice(%arg0 : tensor<?xf32>, %arg1 : tensor<?x?xf32>, %arg2 : index, %arg3 : index) -> tensor<?x?xf32> { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %d0 = tensor.dim %arg0, %c0 : tensor<?xf32> %0 = tensor.insert_slice %arg0 into %arg1[%arg2, %arg3] [1, %d0] [1, 1] {__internal_linalg_transform__ = "tiling_input"} : tensor<?xf32> into tensor<?x?xf32> @@ -701,8 +701,8 @@ // CHECK-SAME: %[[ARG1:.+]]: tensor<?x?xf32> // CHECK-SAME: %[[ARG2:[a-zA-Z0-9_]+]]: index // CHECK-SAME: %[[ARG3:[a-zA-Z0-9_]+]]: index -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C10:.+]] = constant 10 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C10:.+]] = arith.constant 10 : index // CHECK: %[[D0:.+]] = tensor.dim %[[ARG0]], %[[C0]] : tensor<?xf32> // CHECK: %[[RESULT:.+]] = scf.for %[[ARG4:.+]] = %[[C0]] to %[[D0]] // CHECK-SAME: step %[[C10]] iter_args(%[[ARG5:.+]] = %[[ARG1]]) @@ -719,7 +719,7 @@ func @insert_slice_rank_reduced_inner(%arg0 : tensor<?xf32>, %arg1 : tensor<?x?x?xf32>, %arg2: index, %arg3 : index, %arg4 : index) -> tensor<?x?x?xf32> { - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %d0 = tensor.dim %arg0, %c0 : tensor<?xf32> %0 = tensor.insert_slice %arg0 into %arg1[%arg2, %arg3, %arg4] [1, %d0, 1] [1, 1, 1] {__internal_linalg_transform__ = "tiling_input"} : tensor<?xf32> into tensor<?x?x?xf32> @@ -733,8 +733,8 @@ // CHECK-SAME: %[[ARG2:[a-zA-Z0-9_]+]]: index // CHECK-SAME: %[[ARG3:[a-zA-Z0-9_]+]]: index // CHECK-SAME: %[[ARG4:[a-zA-Z0-9_]+]]: index -// CHECK-DAG: %[[LB:.+]] = constant 0 : index -// CHECK-DAG: %[[STEP:.+]] = constant 10 : index +// CHECK-DAG: %[[LB:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[STEP:.+]] = arith.constant 10 : index // CHECK: %[[UB:.+]] = tensor.dim %[[ARG0]] // CHECK: %[[RESULT:.+]] = scf.for %[[IV0:[a-zA-Z0-9_]+]] = %[[LB]] // CHECK-SAME: to %[[D0]] step %[[STEP]] iter_args(%[[ARG6:.+]] = %[[ARG1]])
diff --git a/iree/compiler/Dialect/Modules/Check/IR/CheckOps.cpp b/iree/compiler/Dialect/Modules/Check/IR/CheckOps.cpp index 67c4ad7..ed43857 100644 --- a/iree/compiler/Dialect/Modules/Check/IR/CheckOps.cpp +++ b/iree/compiler/Dialect/Modules/Check/IR/CheckOps.cpp
@@ -22,7 +22,7 @@ using OpRewritePattern<SrcOp>::OpRewritePattern; LogicalResult matchAndRewrite(SrcOp op, PatternRewriter &rewriter) const override { - auto rhs = rewriter.create<ConstantOp>(op.getLoc(), op.value()); + auto rhs = rewriter.create<arith::ConstantOp>(op.getLoc(), op.value()); rewriter.replaceOpWithNewOp<DstOp>(op, op.lhs(), rhs); return success(); }
diff --git a/iree/compiler/Dialect/Modules/Check/test/canonicalize.mlir b/iree/compiler/Dialect/Modules/Check/test/canonicalize.mlir index 3250712..e952d0b 100644 --- a/iree/compiler/Dialect/Modules/Check/test/canonicalize.mlir +++ b/iree/compiler/Dialect/Modules/Check/test/canonicalize.mlir
@@ -5,7 +5,7 @@ // CHECK-LABEL: @expect_eq_const // CHECK-SAME: %[[LHS:[a-zA-Z0-9$._-]+]] func @expect_eq_const(%lhs : tensor<2x2xi32>) { - // CHECK: %[[C:.+]] = constant dense<1> : tensor<2x2xi32> + // CHECK: %[[C:.+]] = arith.constant dense<1> : tensor<2x2xi32> // CHECK: check.expect_eq(%[[LHS]], %[[C]]) : tensor<2x2xi32> check.expect_eq_const(%lhs, dense<1> : tensor<2x2xi32>) : tensor<2x2xi32> return @@ -16,7 +16,7 @@ // CHECK-LABEL: @expect_almost_eq_const // CHECK-SAME: %[[LHS:[a-zA-Z0-9$._-]+]] func @expect_almost_eq_const(%lhs : tensor<2x2xf32>) { - // CHECK: %[[C:.+]] = constant dense<1.000000e+00> : tensor<2x2xf32> + // CHECK: %[[C:.+]] = arith.constant dense<1.000000e+00> : tensor<2x2xf32> // CHECK: check.expect_almost_eq(%[[LHS]], %[[C]]) : tensor<2x2xf32> check.expect_almost_eq_const(%lhs, dense<1.0> : tensor<2x2xf32>) : tensor<2x2xf32> return
diff --git a/iree/compiler/Dialect/Modules/VMVX/Conversion/HALToVMVX/ConvertHALToVMVX.cpp b/iree/compiler/Dialect/Modules/VMVX/Conversion/HALToVMVX/ConvertHALToVMVX.cpp index f62175c..38646b3 100644 --- a/iree/compiler/Dialect/Modules/VMVX/Conversion/HALToVMVX/ConvertHALToVMVX.cpp +++ b/iree/compiler/Dialect/Modules/VMVX/Conversion/HALToVMVX/ConvertHALToVMVX.cpp
@@ -188,11 +188,12 @@ auto resultType = getTypeConverter()->convertType(op.result().getType()); - auto constantOrdinal = rewriter.createOrFold<ConstantIndexOp>( + auto constantOrdinal = rewriter.createOrFold<arith::ConstantIndexOp>( op.getLoc(), op.offset().getZExtValue()); auto loadedValue = rewriter.createOrFold<memref::LoadOp>( op.getLoc(), constantType, constantsArg, ValueRange{constantOrdinal}); - rewriter.replaceOpWithNewOp<IndexCastOp>(op, loadedValue, resultType); + rewriter.replaceOpWithNewOp<arith::IndexCastOp>(op, loadedValue, + resultType); return success(); } }; @@ -232,7 +233,7 @@ bindingsArg.getType().cast<IREE::Util::ListType>().getElementType(); auto getOp = rewriter.create<IREE::Util::ListGetOp>( op.getLoc(), bindingType, bindingsArg, - rewriter.createOrFold<ConstantIndexOp>( + rewriter.createOrFold<arith::ConstantIndexOp>( op.getLoc(), interfaceBindingOp.binding().getZExtValue())); rewriter.replaceOpWithNewOp<UnrealizedConversionCastOp>( op,
diff --git a/iree/compiler/Dialect/Modules/VMVX/Conversion/HALToVMVX/test/interface_ops.mlir b/iree/compiler/Dialect/Modules/VMVX/Conversion/HALToVMVX/test/interface_ops.mlir index 51ecc83..9b09ba7 100644 --- a/iree/compiler/Dialect/Modules/VMVX/Conversion/HALToVMVX/test/interface_ops.mlir +++ b/iree/compiler/Dialect/Modules/VMVX/Conversion/HALToVMVX/test/interface_ops.mlir
@@ -22,10 +22,10 @@ // CHECK-SAME: %[[WORKGROUP_COUNT_Y:[a-z0-9]+]]: index, // CHECK-SAME: %[[WORKGROUP_COUNT_Z:[a-z0-9]+]]: index) { func @entry() { - %cst = constant 0.000000e+00 : f32 - %c5 = constant 5 : index - %c0 = constant 0 : index - %c1 = constant 1 : index + %cst = arith.constant 0.000000e+00 : f32 + %c5 = arith.constant 5 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index %0 = memref.get_global @__constant_5xi32 : memref<5xi32> // CHECK: %[[BINDING0_RAW:.+]] = util.list.get %[[BINDINGS]][%c0] : !util.list<memref<?xi8>> // CHECK-NEXT: %[[BINDING0:.+]] = builtin.unrealized_conversion_cast %[[BINDING0_RAW]] : memref<?xi8> to memref<5xf32> @@ -46,9 +46,9 @@ %6 = affine.apply affine_map<(d0)[s0] -> (d0 + s0)>(%arg1)[%arg0] %7 = memref.load %1[%6] : memref<5xf32> %8 = memref.load %0[%6] : memref<5xi32> - %9 = cmpf oeq, %7, %cst : f32 - %10 = zexti %9 : i1 to i32 - %11 = muli %10, %8 : i32 + %9 = arith.cmpf oeq, %7, %cst : f32 + %10 = arith.extui %9 : i1 to i32 + %11 = arith.muli %10, %8 : i32 memref.store %11, %2[%6] : memref<5xi32> } }
diff --git a/iree/compiler/Dialect/Modules/VMVX/Transforms/Conversion.cpp b/iree/compiler/Dialect/Modules/VMVX/Transforms/Conversion.cpp index ee4d50f..5b6775a 100644 --- a/iree/compiler/Dialect/Modules/VMVX/Transforms/Conversion.cpp +++ b/iree/compiler/Dialect/Modules/VMVX/Transforms/Conversion.cpp
@@ -67,7 +67,8 @@ conversionTarget.addIllegalDialect<tensor::TensorDialect>(); conversionTarget.addLegalDialect<IREE::Util::UtilDialect>(); conversionTarget.addLegalDialect<IREE::VMVX::VMVXDialect>(); - conversionTarget.addLegalDialect<mlir::StandardOpsDialect>(); + conversionTarget.addLegalDialect<mlir::StandardOpsDialect, + mlir::arith::ArithmeticDialect>(); conversionTarget.addLegalDialect<mlir::AffineDialect>(); conversionTarget.addLegalDialect<memref::MemRefDialect>(); conversionTarget.addLegalOp<mlir::UnrealizedConversionCastOp>();
diff --git a/iree/compiler/Dialect/Shape/IR/Builders.cpp b/iree/compiler/Dialect/Shape/IR/Builders.cpp index 1e53e2d..77e6bdb 100644 --- a/iree/compiler/Dialect/Shape/IR/Builders.cpp +++ b/iree/compiler/Dialect/Shape/IR/Builders.cpp
@@ -104,8 +104,8 @@ if (rsType.isDimDynamic(i)) { result.push_back(builder.createOrFold<Shape::RankedDimOp>(loc, rs, i)); } else { - result.push_back( - builder.create<ConstantIndexOp>(loc, rsType.getStaticDim(i))); + result.push_back(builder.create<arith::ConstantIndexOp>( + loc, rsType.getStaticDim(i))); } } } else { @@ -116,7 +116,7 @@ result.push_back(builder.createOrFold<tensor::DimOp>(loc, value, i)); } else { result.push_back( - builder.create<ConstantIndexOp>(loc, valueSt.getDimSize(i))); + builder.create<arith::ConstantIndexOp>(loc, valueSt.getDimSize(i))); } } }
diff --git a/iree/compiler/Dialect/Shape/IR/ShapeDialect.cpp b/iree/compiler/Dialect/Shape/IR/ShapeDialect.cpp index d6adb8d..0a0f34d 100644 --- a/iree/compiler/Dialect/Shape/IR/ShapeDialect.cpp +++ b/iree/compiler/Dialect/Shape/IR/ShapeDialect.cpp
@@ -79,8 +79,10 @@ auto rankedShape = typeAttr.getValue().cast<Shape::RankedShapeType>(); return builder.create<Shape::ConstRankedShapeOp>(loc, rankedShape); } - if (ConstantOp::isBuildableWith(value, type)) - return builder.create<ConstantOp>(loc, type, value); + if (arith::ConstantOp::isBuildableWith(value, type)) + return builder.create<arith::ConstantOp>(loc, type, value); + if (mlir::ConstantOp::isBuildableWith(value, type)) + return builder.create<mlir::ConstantOp>(loc, type, value); return nullptr; }
diff --git a/iree/compiler/Dialect/Shape/IR/test/canonicalize.mlir b/iree/compiler/Dialect/Shape/IR/test/canonicalize.mlir index fb1ddc3..b9c3913 100644 --- a/iree/compiler/Dialect/Shape/IR/test/canonicalize.mlir +++ b/iree/compiler/Dialect/Shape/IR/test/canonicalize.mlir
@@ -31,7 +31,7 @@ // CHECK-LABEL: @foldStaticRankedDim // CHECK-SAME: %[[SHAPE:[^:[:space:]]+]]: !shapex.ranked_shape<[1,?,2,?]> func @foldStaticRankedDim(%arg0: !shapex.ranked_shape<[1,?,2,?]>) -> (i32, i32) { - // CHECK-DAG: %[[D2:.+]] = constant 2 : i32 + // CHECK-DAG: %[[D2:.+]] = arith.constant 2 : i32 %0 = shapex.ranked_dim %arg0[2] : !shapex.ranked_shape<[1,?,2,?]> -> i32 // CHECK-DAG: %[[D1:.+]] = shapex.ranked_dim %[[SHAPE]][1] %1 = shapex.ranked_dim %arg0[1] : !shapex.ranked_shape<[1,?,2,?]> -> i32 @@ -65,8 +65,8 @@ %d1 = shapex.ranked_dim %rs[1] : !shapex.ranked_shape<[?,8,?,16]> -> index %d2 = shapex.ranked_dim %rs[2] : !shapex.ranked_shape<[?,8,?,16]> -> index %d3 = shapex.ranked_dim %rs[3] : !shapex.ranked_shape<[?,8,?,16]> -> index - // CHECK-DAG: %[[C8:.+]] = constant 8 : index - // CHECK-DAG: %[[C16:.+]] = constant 16 : index + // CHECK-DAG: %[[C8:.+]] = arith.constant 8 : index + // CHECK-DAG: %[[C16:.+]] = arith.constant 16 : index // CHECK-DAG: return %[[DD0]], %[[C8]], %[[DD1]], %[[C16]] return %d0, %d1, %d2, %d3 : index, index, index, index }
diff --git a/iree/compiler/Dialect/Shape/Transforms/test/expand_function_dynamic_dims.mlir b/iree/compiler/Dialect/Shape/Transforms/test/expand_function_dynamic_dims.mlir index be98e44..289eef8 100644 --- a/iree/compiler/Dialect/Shape/Transforms/test/expand_function_dynamic_dims.mlir +++ b/iree/compiler/Dialect/Shape/Transforms/test/expand_function_dynamic_dims.mlir
@@ -26,7 +26,7 @@ // CHECK-SAME: -> (tensor<?x3xf32>, !shapex.ranked_shape<[?,3]>) // Should insert function shape argument and result. func @dynamicReturnInBlock(%arg0 : tensor<2x3xf32>) -> (tensor<?x3xf32>) { - %0 = addf %arg0, %arg0 : tensor<2x3xf32> + %0 = arith.addf %arg0, %arg0 : tensor<2x3xf32> br ^bb1 // CHECK: ^bb1 ^bb1:
diff --git a/iree/compiler/Dialect/Shape/Transforms/test/expand_function_ranked_shape_dims.mlir b/iree/compiler/Dialect/Shape/Transforms/test/expand_function_ranked_shape_dims.mlir index 5e29f4b..a33c2e1 100644 --- a/iree/compiler/Dialect/Shape/Transforms/test/expand_function_ranked_shape_dims.mlir +++ b/iree/compiler/Dialect/Shape/Transforms/test/expand_function_ranked_shape_dims.mlir
@@ -25,8 +25,8 @@ // CHECK-LABEL: @expandResults // CHECK-SAME: () -> (index, index) func @expandResults() -> (!shapex.ranked_shape<[?,?]>) { - %idx0 = constant 5 : index - %idx1 = constant 6 : index + %idx0 = arith.constant 5 : index + %idx1 = arith.constant 6 : index // CHECK: %[[RS:.+]] = shapex.make_ranked_shape %rs = shapex.make_ranked_shape %idx0, %idx1 : (index, index) -> !shapex.ranked_shape<[?,?]> // CHECK-DAG: %[[CAST_IDX0:.+]] = shapex.ranked_dim %[[RS]][0]
diff --git a/iree/compiler/Dialect/Shape/Transforms/test/fold_dim_over_shape_carrying_op.mlir b/iree/compiler/Dialect/Shape/Transforms/test/fold_dim_over_shape_carrying_op.mlir index 84931a3..b44a067 100644 --- a/iree/compiler/Dialect/Shape/Transforms/test/fold_dim_over_shape_carrying_op.mlir +++ b/iree/compiler/Dialect/Shape/Transforms/test/fold_dim_over_shape_carrying_op.mlir
@@ -3,14 +3,14 @@ // CHECK: func @memrefDim // CHECK-SAME: (%[[DIM0:.+]]: index, %[[DIM1:.+]]: index) func @memrefDim(%d0: index, %d1: index) -> (index, index, index) { - %c0 = constant 0 : index - %c1 = constant 1 : index - %c2 = constant 2 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + %c2 = arith.constant 2 : index %subspan = hal.interface.binding.subspan @io::@s0b0_ro_constant[%c0] : memref<?x7x?xf32>{%d0, %d1} %dim0 = memref.dim %subspan, %c0 : memref<?x7x?xf32> %dim1 = memref.dim %subspan, %c1 : memref<?x7x?xf32> %dim2 = memref.dim %subspan, %c2 : memref<?x7x?xf32> - // CHECK: %[[C7:.+]] = constant 7 : index + // CHECK: %[[C7:.+]] = arith.constant 7 : index // CHECK: return %[[DIM0]], %[[C7]], %[[DIM1]] return %dim0, %dim1, %dim2 : index, index, index } @@ -24,14 +24,14 @@ // CHECK: func @tensorDim // CHECK-SAME: (%{{.+}}: f32, %[[DIM0:.+]]: index, %[[DIM1:.+]]: index) func @tensorDim(%value: f32, %d0: index, %d1: index) -> (index, index, index) { - %c0 = constant 0 : index - %c1 = constant 1 : index - %c2 = constant 2 : index + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + %c2 = arith.constant 2 : index %splat = flow.tensor.splat %value : tensor<?x8x?xf32>{%d0, %d1} %dim0 = tensor.dim %splat, %c0 : tensor<?x8x?xf32> %dim1 = tensor.dim %splat, %c1 : tensor<?x8x?xf32> %dim2 = tensor.dim %splat, %c2 : tensor<?x8x?xf32> - // CHECK: %[[C8:.+]] = constant 8 : index + // CHECK: %[[C8:.+]] = arith.constant 8 : index // CHECK: return %[[DIM0]], %[[C8]], %[[DIM1]] return %dim0, %dim1, %dim2 : index, index, index }
diff --git a/iree/compiler/Dialect/Util/Conversion/test/hint_ops.mlir b/iree/compiler/Dialect/Util/Conversion/test/hint_ops.mlir index e611cd9..ad4bfa3 100644 --- a/iree/compiler/Dialect/Util/Conversion/test/hint_ops.mlir +++ b/iree/compiler/Dialect/Util/Conversion/test/hint_ops.mlir
@@ -5,8 +5,8 @@ // CHECK: util.do_not_optimize() util.do_not_optimize() - // CHECK: %[[C:.+]] = constant 2 - %c = constant 2 : i32 + // CHECK: %[[C:.+]] = arith.constant 2 + %c = arith.constant 2 : i32 // CHECK: util.do_not_optimize(%[[C]]) util.do_not_optimize(%c) : i32 return
diff --git a/iree/compiler/Dialect/Util/IR/BUILD b/iree/compiler/Dialect/Util/IR/BUILD index 757be03..c602293 100644 --- a/iree/compiler/Dialect/Util/IR/BUILD +++ b/iree/compiler/Dialect/Util/IR/BUILD
@@ -70,6 +70,7 @@ ":UtilInterfacesGen", ":UtilOpsGen", "@llvm-project//llvm:Support", + "@llvm-project//mlir:ArithmeticDialect", "@llvm-project//mlir:ControlFlowInterfaces", "@llvm-project//mlir:IR", "@llvm-project//mlir:Parser",
diff --git a/iree/compiler/Dialect/Util/IR/CMakeLists.txt b/iree/compiler/Dialect/Util/IR/CMakeLists.txt index 9696c1b..cb1cc88 100644 --- a/iree/compiler/Dialect/Util/IR/CMakeLists.txt +++ b/iree/compiler/Dialect/Util/IR/CMakeLists.txt
@@ -42,6 +42,7 @@ ::UtilInterfacesGen ::UtilOpsGen LLVMSupport + MLIRArithmetic MLIRControlFlowInterfaces MLIRIR MLIRParser
diff --git a/iree/compiler/Dialect/Util/IR/ClosureOpUtils.cpp b/iree/compiler/Dialect/Util/IR/ClosureOpUtils.cpp index 7febab2..4246a73 100644 --- a/iree/compiler/Dialect/Util/IR/ClosureOpUtils.cpp +++ b/iree/compiler/Dialect/Util/IR/ClosureOpUtils.cpp
@@ -117,7 +117,7 @@ // This is also still at a fairly high level (flow dialect): once the closures // are expanded out in lower dialects things like CSE have a chance to once // again get at the constants and dedupe them if they survive. -static bool isConstantSmall(ConstantOp constantOp) { +static bool isConstantSmall(arith::ConstantOp constantOp) { // We could tune this/take it as a configuration setting. // The current value is chosen based on what is known to be reasonable to // inline into command buffers way down in the HAL, which is not great but at @@ -125,7 +125,7 @@ // constants or inlining megabytes. static constexpr int kMaxInlinedConstantBytes = 256; - auto constantValueAttr = constantOp.getValue(); + auto constantValueAttr = constantOp.value(); auto constantType = constantOp.getType(); if (constantValueAttr.isa<SplatElementsAttr>()) { // Splats are always small and can often have special handling when we @@ -154,7 +154,7 @@ // returns true for. static bool shouldInlineIntoClosure(Value value) { auto definingOp = value.getDefiningOp(); - if (auto constantOp = dyn_cast<ConstantOp>(definingOp)) { + if (auto constantOp = dyn_cast<arith::ConstantOp>(definingOp)) { // Constants are perfect! return isConstantSmall(constantOp); } else if (auto loadOp = dyn_cast<IREE::Util::GlobalLoadOp>(definingOp)) {
diff --git a/iree/compiler/Dialect/Util/IR/UtilDialect.cpp b/iree/compiler/Dialect/Util/IR/UtilDialect.cpp index e97c5b4..3476362 100644 --- a/iree/compiler/Dialect/Util/IR/UtilDialect.cpp +++ b/iree/compiler/Dialect/Util/IR/UtilDialect.cpp
@@ -10,6 +10,7 @@ #include "iree/compiler/Dialect/Util/IR/UtilTypes.h" #include "llvm/ADT/SmallVector.h" #include "llvm/Support/SourceMgr.h" +#include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/IR/Attributes.h" #include "mlir/IR/DialectImplementation.h" #include "mlir/IR/OpDefinition.h" @@ -73,6 +74,13 @@ >(); } +Operation *UtilDialect::materializeConstant(OpBuilder &builder, Attribute value, + Type type, Location loc) { + if (arith::ConstantOp::isBuildableWith(value, type)) + return builder.create<arith::ConstantOp>(loc, value, type); + return nullptr; +} + } // namespace Util } // namespace IREE } // namespace iree_compiler
diff --git a/iree/compiler/Dialect/Util/IR/UtilDialect.h b/iree/compiler/Dialect/Util/IR/UtilDialect.h index 4e240d5..468d3ed 100644 --- a/iree/compiler/Dialect/Util/IR/UtilDialect.h +++ b/iree/compiler/Dialect/Util/IR/UtilDialect.h
@@ -26,6 +26,9 @@ Type parseType(DialectAsmParser& parser) const override; void printType(Type type, DialectAsmPrinter& os) const override; + Operation* materializeConstant(OpBuilder& builder, Attribute value, Type type, + Location loc) override; + private: void registerAttributes(); void registerTypes();
diff --git a/iree/compiler/Dialect/Util/IR/UtilOps.cpp b/iree/compiler/Dialect/Util/IR/UtilOps.cpp index 619cdde..bbbca81 100644 --- a/iree/compiler/Dialect/Util/IR/UtilOps.cpp +++ b/iree/compiler/Dialect/Util/IR/UtilOps.cpp
@@ -554,7 +554,7 @@ using OpRewritePattern<IREE::Util::UnfoldableConstantOp>::OpRewritePattern; LogicalResult matchAndRewrite(UnfoldableConstantOp op, PatternRewriter &rewriter) const override { - auto stdConst = rewriter.create<ConstantOp>(op.getLoc(), op.value()); + auto stdConst = rewriter.create<arith::ConstantOp>(op.getLoc(), op.value()); rewriter.replaceOpWithNewOp<DoNotOptimizeOp>(op, stdConst.getResult()); return success(); }
diff --git a/iree/compiler/Dialect/Util/IR/UtilOps.td b/iree/compiler/Dialect/Util/IR/UtilOps.td index b188fe4..6730ef6 100644 --- a/iree/compiler/Dialect/Util/IR/UtilOps.td +++ b/iree/compiler/Dialect/Util/IR/UtilOps.td
@@ -136,7 +136,7 @@ ```mlir ^bb0: - %true = constant true + %true = arith.constant true cond_br %true, ^bb2, ^bb1 ^bb1: // Indicates that this branch should never be taken.
diff --git a/iree/compiler/Dialect/Util/IR/test/global_folding.mlir b/iree/compiler/Dialect/Util/IR/test/global_folding.mlir index 2f22f9f..0ba2f0e 100644 --- a/iree/compiler/Dialect/Util/IR/test/global_folding.mlir +++ b/iree/compiler/Dialect/Util/IR/test/global_folding.mlir
@@ -4,7 +4,7 @@ util.global private @v_initialized : tensor<4xi32> // CHECK-NOT: util.initializer util.initializer { - %0 = constant dense<4> : tensor<4xi32> + %0 = arith.constant dense<4> : tensor<4xi32> util.global.store %0, @v_initialized : tensor<4xi32> util.initializer.return } @@ -24,7 +24,7 @@ util.global private @v_const = dense<1.0> : tensor<8xf32> // CHECK-LABEL: @fold_immutable_const func @fold_immutable_const() -> tensor<8xf32> { - // CHECK-NEXT: %[[CONST:.+]] = constant dense<1.{{.+}}> : tensor<8xf32> + // CHECK-NEXT: %[[CONST:.+]] = arith.constant dense<1.{{.+}}> : tensor<8xf32> %0 = util.global.load @v_const : tensor<8xf32> // CHECK-NEXT: return %[[CONST]] : tensor<8xf32> return %0 : tensor<8xf32>
diff --git a/iree/compiler/Dialect/Util/IR/test/global_ops.mlir b/iree/compiler/Dialect/Util/IR/test/global_ops.mlir index 58612e3..6af4b03 100644 --- a/iree/compiler/Dialect/Util/IR/test/global_ops.mlir +++ b/iree/compiler/Dialect/Util/IR/test/global_ops.mlir
@@ -54,8 +54,8 @@ util.global private mutable @v_stored : tensor<4xi32> // CHECK-LABEL: @stored func @stored() { - // CHECK-NEXT: %[[VAL:.+]] = constant - %cst = constant dense<5> : tensor<4xi32> + // CHECK-NEXT: %[[VAL:.+]] = arith.constant + %cst = arith.constant dense<5> : tensor<4xi32> // CHECK-NEXT: util.global.store %[[VAL]], @v_stored : tensor<4xi32> util.global.store %cst, @v_stored : tensor<4xi32> return
diff --git a/iree/compiler/Dialect/Util/IR/test/hint_folding.mlir b/iree/compiler/Dialect/Util/IR/test/hint_folding.mlir index 9a38ac7..4a84dec 100644 --- a/iree/compiler/Dialect/Util/IR/test/hint_folding.mlir +++ b/iree/compiler/Dialect/Util/IR/test/hint_folding.mlir
@@ -3,7 +3,7 @@ // CHECK-LABEL: @no_fold_constant func @no_fold_constant() -> (i32) { // CHECK: constant 1 : i32 - %0 = constant 1 : i32 + %0 = arith.constant 1 : i32 // CHECK: util.do_not_optimize %1 = "util.do_not_optimize"(%0) : (i32) -> i32 return %1 : i32 @@ -55,7 +55,7 @@ // CHECK-LABEL: @canonicalize_unfoldable_constant func @canonicalize_unfoldable_constant() -> i32 { - // CHECK-NEXT: %[[C:.+]] = constant 42 : i32 + // CHECK-NEXT: %[[C:.+]] = arith.constant 42 : i32 // CHECK-NEXT: %[[R:.+]] = util.do_not_optimize(%[[C]]) : i32 %c42 = util.unfoldable_constant 42 : i32 // CHECK-NEXT: return %[[R]]
diff --git a/iree/compiler/Dialect/Util/IR/test/list_ops.mlir b/iree/compiler/Dialect/Util/IR/test/list_ops.mlir index 7a7d13c..320e679 100644 --- a/iree/compiler/Dialect/Util/IR/test/list_ops.mlir +++ b/iree/compiler/Dialect/Util/IR/test/list_ops.mlir
@@ -2,16 +2,16 @@ // CHECK-LABEL: @list_init_ops func @list_init_ops() { - // CHECK: %[[CAPACITY:.+]] = constant 5 - %capacity = constant 5 : index + // CHECK: %[[CAPACITY:.+]] = arith.constant 5 + %capacity = arith.constant 5 : index // CHECK: = util.list.create %[[CAPACITY]] : !util.list<?> %list_initial_capacity = util.list.create %capacity : !util.list<?> // CHECK: %[[LIST:.+]] = util.list.create : !util.list<?> %list = util.list.create : !util.list<?> - // CHECK: %[[NEW_SIZE:.+]] = constant 100 - %new_size = constant 100 : index + // CHECK: %[[NEW_SIZE:.+]] = arith.constant 100 + %new_size = arith.constant 100 : index // CHECK: util.list.resize %[[LIST]], %[[NEW_SIZE]] : !util.list<?> util.list.resize %list, %new_size : !util.list<?> @@ -23,15 +23,15 @@ // CHECK-LABEL: @list_access // CHECK-SAME: (%[[LIST:.+]]: !util.list<i32>) func @list_access(%list: !util.list<i32>) { - %c10 = constant 10 : index + %c10 = arith.constant 10 : index // CHECK: = util.list.get %[[LIST]][%c10] : !util.list<i32> %0 = util.list.get %list[%c10] : !util.list<i32> // CHECK: = util.list.get %[[LIST]][%c10] : !util.list<i32> %1 = util.list.get %list[%c10] : !util.list<i32> -> i32 - // CHECK: %[[NEW_VALUE:.+]] = constant 100 : i32 - %new_value = constant 100 : i32 + // CHECK: %[[NEW_VALUE:.+]] = arith.constant 100 : i32 + %new_value = arith.constant 100 : i32 // CHECK: util.list.set %[[LIST]][%c10], %[[NEW_VALUE]] : !util.list<i32> util.list.set %list[%c10], %new_value : !util.list<i32> @@ -43,13 +43,13 @@ // CHECK-LABEL: @list_access_tensor // CHECK-SAME: (%[[LIST:.+]]: !util.list<tensor<*xf32>>) func @list_access_tensor(%list: !util.list<tensor<*xf32>>) { - %c10 = constant 10 : index + %c10 = arith.constant 10 : index // CHECK: = util.list.get %[[LIST]][%c10] : !util.list<tensor<*xf32>> -> tensor<?xf32> %0 = util.list.get %list[%c10] : !util.list<tensor<*xf32>> -> tensor<?xf32> - // CHECK: %[[NEW_VALUE:.+]] = constant dense<1> : tensor<5xi32> - %new_value = constant dense<1> : tensor<5xi32> + // CHECK: %[[NEW_VALUE:.+]] = arith.constant dense<1> : tensor<5xi32> + %new_value = arith.constant dense<1> : tensor<5xi32> // CHECK: util.list.set %[[LIST]][%c10], %[[NEW_VALUE]] : tensor<5xi32> -> !util.list<tensor<*xf32>> util.list.set %list[%c10], %new_value : tensor<5xi32> -> !util.list<tensor<*xf32>> @@ -61,22 +61,22 @@ // CHECK-LABEL: @list_access_variant // CHECK-SAME: (%[[LIST:.+]]: !util.list<?>) func @list_access_variant(%list: !util.list<?>) { - %c10 = constant 10 : index - %c11 = constant 11 : index + %c10 = arith.constant 10 : index + %c11 = arith.constant 11 : index // CHECK: = util.list.get %[[LIST]][%c10] : !util.list<?> -> i32 %0 = util.list.get %list[%c10] : !util.list<?> -> i32 - // CHECK: %[[NEW_I32_VALUE:.+]] = constant 100 : i32 - %new_i32_value = constant 100 : i32 + // CHECK: %[[NEW_I32_VALUE:.+]] = arith.constant 100 : i32 + %new_i32_value = arith.constant 100 : i32 // CHECK: util.list.set %[[LIST]][%c10], %[[NEW_I32_VALUE]] : i32 -> !util.list<?> util.list.set %list[%c10], %new_i32_value : i32 -> !util.list<?> // CHECK: = util.list.get %[[LIST]][%c11] : !util.list<?> -> tensor<5xf32> %1 = util.list.get %list[%c11] : !util.list<?> -> tensor<5xf32> - // CHECK: %[[NEW_TENSOR_VALUE:.+]] = constant dense<1> : tensor<5xi32> - %new_tensor_value = constant dense<1> : tensor<5xi32> + // CHECK: %[[NEW_TENSOR_VALUE:.+]] = arith.constant dense<1> : tensor<5xi32> + %new_tensor_value = arith.constant dense<1> : tensor<5xi32> // CHECK: util.list.set %[[LIST]][%c11], %[[NEW_TENSOR_VALUE]] : tensor<5xi32> -> !util.list<?> util.list.set %list[%c11], %new_tensor_value : tensor<5xi32> -> !util.list<?>
diff --git a/iree/compiler/Dialect/Util/IR/test/structural_ops.mlir b/iree/compiler/Dialect/Util/IR/test/structural_ops.mlir index 13c5950..2efae61 100644 --- a/iree/compiler/Dialect/Util/IR/test/structural_ops.mlir +++ b/iree/compiler/Dialect/Util/IR/test/structural_ops.mlir
@@ -20,8 +20,8 @@ // CHECK: util.initializer { util.initializer { - // CHECK-NEXT: %[[ZERO:.+]] = constant 0 : i32 - %zero = constant 0 : i32 + // CHECK-NEXT: %[[ZERO:.+]] = arith.constant 0 : i32 + %zero = arith.constant 0 : i32 // CHECK-NEXT: br ^bb1(%[[ZERO]] : i32) br ^bb1(%zero: i32) // CHECK-NEXT: ^bb1(%0: i32):
diff --git a/iree/compiler/Dialect/Util/Transforms/test/drop_compiler_hints.mlir b/iree/compiler/Dialect/Util/Transforms/test/drop_compiler_hints.mlir index e158da0..d1ab191 100644 --- a/iree/compiler/Dialect/Util/Transforms/test/drop_compiler_hints.mlir +++ b/iree/compiler/Dialect/Util/Transforms/test/drop_compiler_hints.mlir
@@ -5,8 +5,8 @@ // CHECK-LABEL: @constant func @constant() -> i32 { - // CHECK-NEXT: %[[C1:.+]] = constant 1 - %c1 = constant 1 : i32 + // CHECK-NEXT: %[[C1:.+]] = arith.constant 1 + %c1 = arith.constant 1 : i32 %0 = util.do_not_optimize(%c1) : i32 // CHECK-NEXT: return %[[C1]] return %0 : i32 @@ -16,12 +16,12 @@ // CHECK-LABEL: @multiple func @multiple() -> (i32, i32) { - // CHECK-NEXT: %[[C1:.+]] = constant 1 - %c1 = constant 1 : i32 + // CHECK-NEXT: %[[C1:.+]] = arith.constant 1 + %c1 = arith.constant 1 : i32 %0 = util.do_not_optimize(%c1) : i32 %1 = util.do_not_optimize(%0) : i32 - // CHECK-NEXT: %[[C2:.+]] = constant 2 - %c2 = constant 2 : i32 + // CHECK-NEXT: %[[C2:.+]] = arith.constant 2 + %c2 = arith.constant 2 : i32 %2 = util.do_not_optimize(%1) : i32 %3 = util.do_not_optimize(%c2) : i32 // CHECK-NEXT: return %[[C1]], %[[C2]] @@ -32,10 +32,10 @@ // CHECK-LABEL: @multiple_operands func @multiple_operands() -> (i32, i32) { - // CHECK-NEXT: %[[C1:.+]] = constant 1 - %c1 = constant 1 : i32 - // CHECK-NEXT: %[[C2:.+]] = constant 2 - %c2 = constant 2 : i32 + // CHECK-NEXT: %[[C1:.+]] = arith.constant 1 + %c1 = arith.constant 1 : i32 + // CHECK-NEXT: %[[C2:.+]] = arith.constant 2 + %c2 = arith.constant 2 : i32 %0, %1 = util.do_not_optimize(%c1, %c2) : i32, i32 // CHECK-NEXT: return %[[C1]], %[[C2]] return %0, %1 : i32, i32
diff --git a/iree/compiler/Dialect/Util/Transforms/test/simplify_global_accesses.mlir b/iree/compiler/Dialect/Util/Transforms/test/simplify_global_accesses.mlir index 2b9c3a9..300532c 100644 --- a/iree/compiler/Dialect/Util/Transforms/test/simplify_global_accesses.mlir +++ b/iree/compiler/Dialect/Util/Transforms/test/simplify_global_accesses.mlir
@@ -8,7 +8,7 @@ // CHECK-NEXT: %[[VAR_A:.+]] = util.global.load @varA : tensor<2xi32> // CHECK-NEXT: %[[VAR_B:.+]] = util.global.load @varB : tensor<2x4xi32> // CHECK-NEXT: constant 10 - %w = constant 10 : index + %w = arith.constant 10 : index %varA = util.global.load @varA : tensor<2xi32> // CHECK-NEXT: %[[T:.+]] = flow.dispatch @ex::@dispatch0{{.+}}(%[[VAR_A]]) %d0 = flow.dispatch @ex::@dispatch0[%w](%varA) : (tensor<2xi32>) -> tensor<2xi32> @@ -31,22 +31,22 @@ br ^bb1(%start : i32) // CHECK: ^bb1(%[[BB1_ARG:.+]]: i32): ^bb1(%2: i32): - %cmp = cmpi slt, %2, %bound : i32 + %cmp = arith.cmpi slt, %2, %bound : i32 cond_br %cmp, ^bb2(%2 : i32), ^bb3(%2 : i32) // CHECK: ^bb2(%[[BB2_ARG:.+]]: i32): ^bb2(%5: i32): %6 = util.global.load @varA : i32 - // CHECK-NEXT: = addi %[[BB2_ARG]], %[[VAR_A]] : i32 - %7 = addi %5, %6 : i32 + // CHECK-NEXT: = arith.addi %[[BB2_ARG]], %[[VAR_A]] : i32 + %7 = arith.addi %5, %6 : i32 br ^bb1(%7 : i32) // CHECK: ^bb3(%[[BB3_ARG:.+]]: i32): ^bb3(%8: i32): %9 = util.global.load @varA : i32 - // CHECK-NEXT: %[[T0:.+]] = muli %[[BB3_ARG]], %[[VAR_A]] : i32 - %10 = muli %8, %9 : i32 + // CHECK-NEXT: %[[T0:.+]] = arith.muli %[[BB3_ARG]], %[[VAR_A]] : i32 + %10 = arith.muli %8, %9 : i32 %11 = util.global.load @varB : i32 - // CHECK-NEXT: %[[T1:.+]] = subi %[[T0]], %[[VAR_B]] - %12 = subi %10, %11 : i32 + // CHECK-NEXT: %[[T1:.+]] = arith.subi %[[T0]], %[[VAR_B]] + %12 = arith.subi %10, %11 : i32 // CHECK-NEXT: return %[[T1]] return %12 : i32 } @@ -61,7 +61,7 @@ // CHECK-DAG: %[[VAR_A:.+]] = util.global.load @varA : tensor<2xi32> // CHECK-DAG: %[[VAR_B:.+]] = util.global.load @varB : tensor<2x4xi32> // CHECK-NEXT: constant 10 - %w = constant 10 : index + %w = arith.constant 10 : index %varA = util.global.load @varA : tensor<2xi32> // CHECK-NEXT: %[[T0:.+]] = flow.dispatch @ex::@dispatch0{{.+}}(%[[VAR_A]]) %d0 = flow.dispatch @ex::@dispatch0[%w](%varA) : (tensor<2xi32>) -> tensor<2xi32>
diff --git a/iree/compiler/Dialect/VM/Conversion/ImportUtils.cpp b/iree/compiler/Dialect/VM/Conversion/ImportUtils.cpp index 95f8985..72c4a35 100644 --- a/iree/compiler/Dialect/VM/Conversion/ImportUtils.cpp +++ b/iree/compiler/Dialect/VM/Conversion/ImportUtils.cpp
@@ -137,7 +137,7 @@ // conversions can do their job. If we want to remove the dependency // from standard ops in the future we could instead go directly to // one of the vm constant ops. - auto constValue = rewriter.createOrFold<mlir::ConstantOp>( + auto constValue = rewriter.createOrFold<mlir::arith::ConstantOp>( loc, inputType, IntegerAttr::get(inputType, APInt(32, static_cast<int32_t>(intAttr.getInt())))); @@ -146,7 +146,7 @@ SmallVector<Value, 4> elementValues; elementValues.reserve(elementsAttr.getNumElements()); for (auto intAttr : elementsAttr.getValues<Attribute>()) { - elementValues.push_back(rewriter.createOrFold<mlir::ConstantOp>( + elementValues.push_back(rewriter.createOrFold<mlir::arith::ConstantOp>( loc, elementsAttr.getType().getElementType(), intAttr)); } return elementValues;
diff --git a/iree/compiler/Dialect/VM/Conversion/MemRefToVM/ConvertMemRefToVM.cpp b/iree/compiler/Dialect/VM/Conversion/MemRefToVM/ConvertMemRefToVM.cpp index 4e2440c..fc3422e 100644 --- a/iree/compiler/Dialect/VM/Conversion/MemRefToVM/ConvertMemRefToVM.cpp +++ b/iree/compiler/Dialect/VM/Conversion/MemRefToVM/ConvertMemRefToVM.cpp
@@ -59,7 +59,7 @@ auto memrefType = memrefValue.getType().cast<ShapedType>(); if (memrefType.getRank() == 0) { // Rank 0 buffers (like memref<i32>) have only a single valid offset at 0. - return rewriter.createOrFold<ConstantIndexOp>(loc, 0); + return rewriter.createOrFold<arith::ConstantIndexOp>(loc, 0); } assert(memrefType.getRank() == 1 && "memrefs should have been flattened");
diff --git a/iree/compiler/Dialect/VM/Conversion/StandardToVM/BUILD b/iree/compiler/Dialect/VM/Conversion/StandardToVM/BUILD index 4793441..f7fb054 100644 --- a/iree/compiler/Dialect/VM/Conversion/StandardToVM/BUILD +++ b/iree/compiler/Dialect/VM/Conversion/StandardToVM/BUILD
@@ -25,6 +25,7 @@ "//iree/compiler/Dialect/VM/Conversion", "//iree/compiler/Dialect/VM/IR", "@llvm-project//mlir:IR", + "@llvm-project//mlir:MathDialect", "@llvm-project//mlir:Pass", "@llvm-project//mlir:StandardOps", "@llvm-project//mlir:TransformUtils",
diff --git a/iree/compiler/Dialect/VM/Conversion/StandardToVM/CMakeLists.txt b/iree/compiler/Dialect/VM/Conversion/StandardToVM/CMakeLists.txt index 603e51e..16addb1 100644 --- a/iree/compiler/Dialect/VM/Conversion/StandardToVM/CMakeLists.txt +++ b/iree/compiler/Dialect/VM/Conversion/StandardToVM/CMakeLists.txt
@@ -20,6 +20,7 @@ "ConvertStandardToVMTest.cpp" DEPS MLIRIR + MLIRMath MLIRPass MLIRStandard MLIRTransformUtils
diff --git a/iree/compiler/Dialect/VM/Conversion/StandardToVM/ConvertStandardToVM.cpp b/iree/compiler/Dialect/VM/Conversion/StandardToVM/ConvertStandardToVM.cpp index f46fb38..dca93f3 100644 --- a/iree/compiler/Dialect/VM/Conversion/StandardToVM/ConvertStandardToVM.cpp +++ b/iree/compiler/Dialect/VM/Conversion/StandardToVM/ConvertStandardToVM.cpp
@@ -11,6 +11,7 @@ #include "iree/compiler/Dialect/VM/Conversion/TargetOptions.h" #include "iree/compiler/Dialect/VM/Conversion/TypeConverter.h" #include "iree/compiler/Dialect/VM/IR/VMOps.h" +#include "mlir/Dialect/Math/IR/Math.h" #include "mlir/Dialect/StandardOps/IR/Ops.h" #include "mlir/IR/Attributes.h" #include "mlir/IR/Builders.h" @@ -144,14 +145,14 @@ } }; -struct ConstantOpConversion : public OpConversionPattern<ConstantOp> { +struct ConstantOpConversion : public OpConversionPattern<arith::ConstantOp> { ConstantOpConversion(MLIRContext *context, TypeConverter &typeConverter) : OpConversionPattern(context), typeConverter(typeConverter) {} TypeConverter &typeConverter; LogicalResult matchAndRewrite( - ConstantOp srcOp, ArrayRef<Value> operands, + arith::ConstantOp srcOp, ArrayRef<Value> operands, ConversionPatternRewriter &rewriter) const override { auto targetType = typeConverter.convertType(srcOp.getType()); if (!targetType) { @@ -159,7 +160,7 @@ << " (check -iree-vm-target-* options)"; } if (targetType.isa<IntegerType>()) { - auto integerAttr = srcOp.getValue().dyn_cast<IntegerAttr>(); + auto integerAttr = srcOp.value().dyn_cast<IntegerAttr>(); if (!integerAttr) { return srcOp.emitRemark() << "unsupported const type for dialect"; } @@ -187,7 +188,7 @@ << "unsupported const integer bit width for dialect"; } } else if (targetType.isa<FloatType>()) { - auto floatAttr = srcOp.getValue().dyn_cast<FloatAttr>(); + auto floatAttr = srcOp.value().dyn_cast<FloatAttr>(); if (!floatAttr) { return srcOp.emitRemark() << "unsupported const type for dialect"; } @@ -217,52 +218,52 @@ } }; -class CmpIOpConversion : public OpConversionPattern<CmpIOp> { +class CmpIOpConversion : public OpConversionPattern<arith::CmpIOp> { using OpConversionPattern::OpConversionPattern; LogicalResult matchAndRewrite( - CmpIOp srcOp, ArrayRef<Value> operands, + arith::CmpIOp srcOp, ArrayRef<Value> operands, ConversionPatternRewriter &rewriter) const override { - CmpIOp::Adaptor srcAdaptor(operands); + arith::CmpIOp::Adaptor srcAdaptor(operands); auto returnType = rewriter.getIntegerType(32); switch (srcOp.getPredicate()) { - case CmpIPredicate::eq: + case arith::CmpIPredicate::eq: rewriter.replaceOpWithNewOp<IREE::VM::CmpEQI32Op>( srcOp, returnType, srcAdaptor.lhs(), srcAdaptor.rhs()); return success(); - case CmpIPredicate::ne: + case arith::CmpIPredicate::ne: rewriter.replaceOpWithNewOp<IREE::VM::CmpNEI32Op>( srcOp, returnType, srcAdaptor.lhs(), srcAdaptor.rhs()); return success(); - case CmpIPredicate::slt: + case arith::CmpIPredicate::slt: rewriter.replaceOpWithNewOp<IREE::VM::CmpLTI32SOp>( srcOp, returnType, srcAdaptor.lhs(), srcAdaptor.rhs()); return success(); - case CmpIPredicate::sle: + case arith::CmpIPredicate::sle: rewriter.replaceOpWithNewOp<IREE::VM::CmpLTEI32SOp>( srcOp, returnType, srcAdaptor.lhs(), srcAdaptor.rhs()); return success(); - case CmpIPredicate::sgt: + case arith::CmpIPredicate::sgt: rewriter.replaceOpWithNewOp<IREE::VM::CmpGTI32SOp>( srcOp, returnType, srcAdaptor.lhs(), srcAdaptor.rhs()); return success(); - case CmpIPredicate::sge: + case arith::CmpIPredicate::sge: rewriter.replaceOpWithNewOp<IREE::VM::CmpGTEI32SOp>( srcOp, returnType, srcAdaptor.lhs(), srcAdaptor.rhs()); return success(); - case CmpIPredicate::ult: + case arith::CmpIPredicate::ult: rewriter.replaceOpWithNewOp<IREE::VM::CmpLTI32UOp>( srcOp, returnType, srcAdaptor.lhs(), srcAdaptor.rhs()); return success(); - case CmpIPredicate::ule: + case arith::CmpIPredicate::ule: rewriter.replaceOpWithNewOp<IREE::VM::CmpLTEI32UOp>( srcOp, returnType, srcAdaptor.lhs(), srcAdaptor.rhs()); return success(); - case CmpIPredicate::ugt: + case arith::CmpIPredicate::ugt: rewriter.replaceOpWithNewOp<IREE::VM::CmpGTI32UOp>( srcOp, returnType, srcAdaptor.lhs(), srcAdaptor.rhs()); return success(); - case CmpIPredicate::uge: + case arith::CmpIPredicate::uge: rewriter.replaceOpWithNewOp<IREE::VM::CmpGTEI32UOp>( srcOp, returnType, srcAdaptor.lhs(), srcAdaptor.rhs()); return success(); @@ -272,22 +273,22 @@ } }; -class CmpFOpConversion : public OpConversionPattern<CmpFOp> { +class CmpFOpConversion : public OpConversionPattern<arith::CmpFOp> { using OpConversionPattern::OpConversionPattern; LogicalResult matchAndRewrite( - CmpFOp srcOp, ArrayRef<Value> operands, + arith::CmpFOp srcOp, ArrayRef<Value> operands, ConversionPatternRewriter &rewriter) const override { - CmpFOp::Adaptor srcAdaptor(operands); + arith::CmpFOp::Adaptor srcAdaptor(operands); auto returnType = rewriter.getIntegerType(32); switch (srcOp.getPredicate()) { - case CmpFPredicate::AlwaysFalse: // 0 + case arith::CmpFPredicate::AlwaysFalse: // 0 rewriter.replaceOpWithNewOp<IREE::VM::ConstI32ZeroOp>(srcOp); break; - case CmpFPredicate::AlwaysTrue: // 1 + case arith::CmpFPredicate::AlwaysTrue: // 1 rewriter.replaceOpWithNewOp<IREE::VM::ConstI32Op>(srcOp, 1); break; - case CmpFPredicate::UNO: // isnan(lhs) || isnan(rhs) + case arith::CmpFPredicate::UNO: // isnan(lhs) || isnan(rhs) rewriter.replaceOpWithNewOp<IREE::VM::OrI32Op>( srcOp, returnType, rewriter.createOrFold<IREE::VM::CmpNaNF32Op>( @@ -295,7 +296,7 @@ rewriter.createOrFold<IREE::VM::CmpNaNF32Op>( srcOp.getLoc(), returnType, srcAdaptor.rhs())); break; - case CmpFPredicate::ORD: // !(isnan(lhs) || isnan(rhs)) + case arith::CmpFPredicate::ORD: // !(isnan(lhs) || isnan(rhs)) rewriter.replaceOpWithNewOp<IREE::VM::XorI32Op>( srcOp, returnType, rewriter.createOrFold<IREE::VM::ConstI32Op>(srcOp.getLoc(), 1), @@ -306,56 +307,57 @@ rewriter.createOrFold<IREE::VM::CmpNaNF32Op>( srcOp.getLoc(), returnType, srcAdaptor.rhs()))); break; - case CmpFPredicate::OEQ: // ordered and equal + case arith::CmpFPredicate::OEQ: // ordered and equal rewriter.replaceOpWithNewOp<IREE::VM::CmpEQF32OOp>( srcOp, returnType, srcAdaptor.lhs(), srcAdaptor.rhs()); break; - case CmpFPredicate::OGT: // ordered and greater than + case arith::CmpFPredicate::OGT: // ordered and greater than rewriter.replaceOpWithNewOp<IREE::VM::CmpGTF32OOp>( srcOp, returnType, srcAdaptor.lhs(), srcAdaptor.rhs()); break; - case CmpFPredicate::OGE: // ordered and greater than or equal + case arith::CmpFPredicate::OGE: // ordered and greater than or equal rewriter.replaceOpWithNewOp<IREE::VM::CmpGTEF32OOp>( srcOp, returnType, srcAdaptor.lhs(), srcAdaptor.rhs()); break; - case CmpFPredicate::OLT: // ordered and less than + case arith::CmpFPredicate::OLT: // ordered and less than rewriter.replaceOpWithNewOp<IREE::VM::CmpLTF32OOp>( srcOp, returnType, srcAdaptor.lhs(), srcAdaptor.rhs()); break; - case CmpFPredicate::OLE: // ordered and less than or equal + case arith::CmpFPredicate::OLE: // ordered and less than or equal rewriter.replaceOpWithNewOp<IREE::VM::CmpLTEF32OOp>( srcOp, returnType, srcAdaptor.lhs(), srcAdaptor.rhs()); break; - case CmpFPredicate::ONE: // ordered and not equal + case arith::CmpFPredicate::ONE: // ordered and not equal rewriter.replaceOpWithNewOp<IREE::VM::CmpNEF32OOp>( srcOp, returnType, srcAdaptor.lhs(), srcAdaptor.rhs()); break; - case CmpFPredicate::UEQ: // unordered or equal + case arith::CmpFPredicate::UEQ: // unordered or equal rewriter.replaceOpWithNewOp<IREE::VM::CmpEQF32UOp>( srcOp, returnType, srcAdaptor.lhs(), srcAdaptor.rhs()); break; - case CmpFPredicate::UGT: // unordered or greater than + case arith::CmpFPredicate::UGT: // unordered or greater than rewriter.replaceOpWithNewOp<IREE::VM::CmpGTF32UOp>( srcOp, returnType, srcAdaptor.lhs(), srcAdaptor.rhs()); break; - case CmpFPredicate::UGE: // unordered or greater than or equal + case arith::CmpFPredicate::UGE: // unordered or greater than or equal rewriter.replaceOpWithNewOp<IREE::VM::CmpGTEF32UOp>( srcOp, returnType, srcAdaptor.lhs(), srcAdaptor.rhs()); break; - case CmpFPredicate::ULT: // unordered or less than + case arith::CmpFPredicate::ULT: // unordered or less than rewriter.replaceOpWithNewOp<IREE::VM::CmpLTF32UOp>( srcOp, returnType, srcAdaptor.lhs(), srcAdaptor.rhs()); break; - case CmpFPredicate::ULE: // unordered or less than or equal + case arith::CmpFPredicate::ULE: // unordered or less than or equal rewriter.replaceOpWithNewOp<IREE::VM::CmpLTEF32UOp>( srcOp, returnType, srcAdaptor.lhs(), srcAdaptor.rhs()); break; - case CmpFPredicate::UNE: // unordered or not equal + case arith::CmpFPredicate::UNE: // unordered or not equal rewriter.replaceOpWithNewOp<IREE::VM::CmpNEF32UOp>( srcOp, returnType, srcAdaptor.lhs(), srcAdaptor.rhs()); break; default: - return rewriter.notifyMatchFailure(srcOp, "unhandled CmpFPredicate"); + return rewriter.notifyMatchFailure(srcOp, + "unhandled arith::CmpFPredicate"); } return success(); } @@ -422,7 +424,7 @@ Value amount = srcAdaptor.rhs(); if (amount.getType().getIntOrFloatBitWidth() > 32) { // Shift amounts are always 32-bit in the VM. - amount = rewriter.createOrFold<TruncateIOp>( + amount = rewriter.createOrFold<arith::TruncIOp>( srcOp.getLoc(), rewriter.getI32Type(), amount); } switch (srcAdaptor.lhs().getType().getIntOrFloatBitWidth()) { @@ -453,35 +455,37 @@ } }; -class IndexCastOpConversion : public OpConversionPattern<IndexCastOp> { +class IndexCastOpConversion : public OpConversionPattern<arith::IndexCastOp> { using OpConversionPattern::OpConversionPattern; LogicalResult matchAndRewrite( - IndexCastOp srcOp, ArrayRef<Value> rawOperands, + arith::IndexCastOp srcOp, ArrayRef<Value> rawOperands, ConversionPatternRewriter &rewriter) const override { - IndexCastOpAdaptor operands(rawOperands); + arith::IndexCastOpAdaptor operands(rawOperands); auto srcType = operands.in().getType(); auto dstType = getTypeConverter()->convertType(srcOp.getResult().getType()); if (srcType == dstType) { rewriter.replaceOp(srcOp, rawOperands); } else if (srcType.getIntOrFloatBitWidth() < dstType.getIntOrFloatBitWidth()) { - rewriter.replaceOpWithNewOp<ZeroExtendIOp>(srcOp, dstType, operands.in()); + rewriter.replaceOpWithNewOp<arith::ExtUIOp>(srcOp, dstType, + operands.in()); } else { - rewriter.replaceOpWithNewOp<TruncateIOp>(srcOp, dstType, operands.in()); + rewriter.replaceOpWithNewOp<arith::TruncIOp>(srcOp, dstType, + operands.in()); } return success(); } }; -class ZeroExtendIOpConversion : public OpConversionPattern<ZeroExtendIOp> { +class ZeroExtendIOpConversion : public OpConversionPattern<arith::ExtUIOp> { using OpConversionPattern::OpConversionPattern; LogicalResult matchAndRewrite( - ZeroExtendIOp srcOp, ArrayRef<Value> rawOperands, + arith::ExtUIOp srcOp, ArrayRef<Value> rawOperands, ConversionPatternRewriter &rewriter) const override { - ZeroExtendIOpAdaptor operands(rawOperands); - auto srcType = srcOp.value().getType(); + arith::ExtUIOpAdaptor operands(rawOperands); + auto srcType = srcOp.in().getType(); auto dstType = getTypeConverter()->convertType(srcOp.getResult().getType()); if (srcType.isInteger(1) && dstType.isInteger(32)) { // This may not be needed but ensures that the input was treated as a @@ -489,17 +493,17 @@ // NOTE: this may not be required - if we know that the i1 is never able // to have more than bit 0 manipulated then this is wasted work. rewriter.replaceOpWithNewOp<IREE::VM::AndI32Op>( - srcOp, dstType, operands.value(), + srcOp, dstType, operands.in(), rewriter.createOrFold<IREE::VM::ConstI32Op>(srcOp.getLoc(), 1)); } else if (srcType.isInteger(8) && dstType.isInteger(32)) { rewriter.replaceOpWithNewOp<IREE::VM::ExtI8I32UOp>(srcOp, dstType, - operands.value()); + operands.in()); } else if (srcType.isInteger(16) && dstType.isInteger(32)) { rewriter.replaceOpWithNewOp<IREE::VM::ExtI16I32UOp>(srcOp, dstType, - operands.value()); + operands.in()); } else if (srcType.isInteger(32) && dstType.isInteger(64)) { rewriter.replaceOpWithNewOp<IREE::VM::ExtI32I64UOp>(srcOp, dstType, - operands.value()); + operands.in()); } else { // TODO(benvanik): we should be building a sequence of extensions for // things like i8 -> i64. @@ -509,24 +513,24 @@ } }; -class SignExtendIOpConversion : public OpConversionPattern<SignExtendIOp> { +class SignExtendIOpConversion : public OpConversionPattern<arith::ExtSIOp> { using OpConversionPattern::OpConversionPattern; LogicalResult matchAndRewrite( - SignExtendIOp srcOp, ArrayRef<Value> rawOperands, + arith::ExtSIOp srcOp, ArrayRef<Value> rawOperands, ConversionPatternRewriter &rewriter) const override { - SignExtendIOpAdaptor operands(rawOperands); - auto srcType = srcOp.value().getType(); + arith::ExtSIOpAdaptor operands(rawOperands); + auto srcType = srcOp.in().getType(); auto dstType = getTypeConverter()->convertType(srcOp.getResult().getType()); if (srcType.isInteger(8) && dstType.isInteger(32)) { rewriter.replaceOpWithNewOp<IREE::VM::ExtI8I32SOp>(srcOp, dstType, - operands.value()); + operands.in()); } else if (srcType.isInteger(16) && dstType.isInteger(32)) { rewriter.replaceOpWithNewOp<IREE::VM::ExtI16I32SOp>(srcOp, dstType, - operands.value()); + operands.in()); } else if (srcType.isInteger(32) && dstType.isInteger(64)) { rewriter.replaceOpWithNewOp<IREE::VM::ExtI32I64SOp>(srcOp, dstType, - operands.value()); + operands.in()); } else { // TODO(benvanik): we should be building a sequence of extensions for // things like i8 -> i64. @@ -536,21 +540,21 @@ } }; -class TruncateIOpConversion : public OpConversionPattern<TruncateIOp> { +class TruncateIOpConversion : public OpConversionPattern<arith::TruncIOp> { using OpConversionPattern::OpConversionPattern; LogicalResult matchAndRewrite( - TruncateIOp srcOp, ArrayRef<Value> rawOperands, + arith::TruncIOp srcOp, ArrayRef<Value> rawOperands, ConversionPatternRewriter &rewriter) const override { - TruncateIOpAdaptor operands(rawOperands); - auto srcType = srcOp.value().getType(); + arith::TruncIOpAdaptor operands(rawOperands); + auto srcType = srcOp.in().getType(); auto resultType = srcOp.getResult().getType(); auto dstType = getTypeConverter()->convertType(resultType); if (resultType.isInteger(1)) { // i1 is represented as i32, so just mask off the bit and truncate as // normal. Note that if we started as i64 we need to first get that into // an i32 that we can work with. - auto value = operands.value(); + auto value = operands.in(); if (srcType.isInteger(64)) { value = rewriter.createOrFold<IREE::VM::TruncI64I32Op>(srcOp.getLoc(), dstType, value); @@ -560,19 +564,19 @@ rewriter.createOrFold<IREE::VM::ConstI32Op>(srcOp.getLoc(), 1)); } else if (srcType.isInteger(32) && resultType.isInteger(8)) { rewriter.replaceOpWithNewOp<IREE::VM::TruncI32I8Op>(srcOp, dstType, - operands.value()); + operands.in()); } else if (srcType.isInteger(32) && resultType.isInteger(16)) { rewriter.replaceOpWithNewOp<IREE::VM::TruncI32I16Op>(srcOp, dstType, - operands.value()); + operands.in()); } else if (srcType.isInteger(64) && resultType.isInteger(8)) { rewriter.replaceOpWithNewOp<IREE::VM::TruncI64I8Op>(srcOp, dstType, - operands.value()); + operands.in()); } else if (srcType.isInteger(64) && resultType.isInteger(16)) { rewriter.replaceOpWithNewOp<IREE::VM::TruncI64I16Op>(srcOp, dstType, - operands.value()); + operands.in()); } else if (srcType.isInteger(64) && resultType.isInteger(32)) { rewriter.replaceOpWithNewOp<IREE::VM::TruncI64I32Op>(srcOp, dstType, - operands.value()); + operands.in()); } else { return rewriter.notifyMatchFailure(srcOp, "unsupported truncation"); } @@ -580,13 +584,13 @@ } }; -class SIToFPOpConversion : public OpConversionPattern<SIToFPOp> { +class SIToFPOpConversion : public OpConversionPattern<arith::SIToFPOp> { using OpConversionPattern::OpConversionPattern; LogicalResult matchAndRewrite( - SIToFPOp srcOp, ArrayRef<Value> operands, + arith::SIToFPOp srcOp, ArrayRef<Value> operands, ConversionPatternRewriter &rewriter) const override { - SIToFPOpAdaptor srcAdaptor(operands); + arith::SIToFPOpAdaptor srcAdaptor(operands); auto srcType = operands[0].getType(); auto dstType = getTypeConverter()->convertType(srcOp.getResult().getType()); if (srcType.isSignlessInteger(32) || srcType.isSignedInteger(32)) { @@ -600,13 +604,13 @@ } }; -class UIToFPOpConversion : public OpConversionPattern<UIToFPOp> { +class UIToFPOpConversion : public OpConversionPattern<arith::UIToFPOp> { using OpConversionPattern::OpConversionPattern; LogicalResult matchAndRewrite( - UIToFPOp srcOp, ArrayRef<Value> operands, + arith::UIToFPOp srcOp, ArrayRef<Value> operands, ConversionPatternRewriter &rewriter) const override { - UIToFPOpAdaptor srcAdaptor(operands); + arith::UIToFPOpAdaptor srcAdaptor(operands); auto srcType = operands[0].getType(); auto dstType = getTypeConverter()->convertType(srcOp.getResult().getType()); if (srcType.isUnsignedInteger(32)) { @@ -620,13 +624,13 @@ } }; -class FPToSIOpConversion : public OpConversionPattern<FPToSIOp> { +class FPToSIOpConversion : public OpConversionPattern<arith::FPToSIOp> { using OpConversionPattern::OpConversionPattern; LogicalResult matchAndRewrite( - FPToSIOp srcOp, ArrayRef<Value> operands, + arith::FPToSIOp srcOp, ArrayRef<Value> operands, ConversionPatternRewriter &rewriter) const override { - FPToSIOpAdaptor srcAdaptor(operands); + arith::FPToSIOpAdaptor srcAdaptor(operands); auto srcType = operands[0].getType(); auto dstType = getTypeConverter()->convertType(srcOp.getResult().getType()); if (srcType.isF32()) { @@ -640,13 +644,13 @@ } }; -class FPToUIOpConversion : public OpConversionPattern<FPToUIOp> { +class FPToUIOpConversion : public OpConversionPattern<arith::FPToUIOp> { using OpConversionPattern::OpConversionPattern; LogicalResult matchAndRewrite( - FPToUIOp srcOp, ArrayRef<Value> operands, + arith::FPToUIOp srcOp, ArrayRef<Value> operands, ConversionPatternRewriter &rewriter) const override { - FPToUIOpAdaptor srcAdaptor(operands); + arith::FPToUIOpAdaptor srcAdaptor(operands); auto srcType = operands[0].getType(); auto dstType = getTypeConverter()->convertType(srcOp.getResult().getType()); if (srcType.isF32()) { @@ -660,11 +664,11 @@ } }; -class BitcastOpConversion : public OpConversionPattern<BitcastOp> { +class BitcastOpConversion : public OpConversionPattern<arith::BitcastOp> { using OpConversionPattern::OpConversionPattern; LogicalResult matchAndRewrite( - BitcastOp srcOp, ArrayRef<Value> operands, + arith::BitcastOp srcOp, ArrayRef<Value> operands, ConversionPatternRewriter &rewriter) const override { auto srcType = operands[0].getType(); auto dstType = getTypeConverter()->convertType(srcOp.getResult().getType()); @@ -819,47 +823,50 @@ context); // Integer arithmetic ops. - patterns.insert< - BinaryArithmeticOpConversion<AddIOp, IREE::VM::AddI32Op, - IREE::VM::AddI64Op>, - BinaryArithmeticOpConversion<SignedDivIOp, IREE::VM::DivI32SOp, - IREE::VM::DivI64SOp>, - BinaryArithmeticOpConversion<UnsignedDivIOp, IREE::VM::DivI32UOp, - IREE::VM::DivI64UOp>, - BinaryArithmeticOpConversion<MulIOp, IREE::VM::MulI32Op, - IREE::VM::MulI64Op>, - BinaryArithmeticOpConversion<SignedRemIOp, IREE::VM::RemI32SOp, - IREE::VM::RemI64SOp>, - BinaryArithmeticOpConversion<UnsignedRemIOp, IREE::VM::RemI32UOp, - IREE::VM::RemI64UOp>, - BinaryArithmeticOpConversion<SubIOp, IREE::VM::SubI32Op, - IREE::VM::SubI64Op>, - BinaryArithmeticOpConversion<AndOp, IREE::VM::AndI32Op, - IREE::VM::AndI64Op>, - BinaryArithmeticOpConversion<OrOp, IREE::VM::OrI32Op, IREE::VM::OrI64Op>, - BinaryArithmeticOpConversion<XOrOp, IREE::VM::XorI32Op, - IREE::VM::XorI64Op>>(typeConverter, context); + patterns + .insert<BinaryArithmeticOpConversion<arith::AddIOp, IREE::VM::AddI32Op, + IREE::VM::AddI64Op>, + BinaryArithmeticOpConversion<arith::DivSIOp, IREE::VM::DivI32SOp, + IREE::VM::DivI64SOp>, + BinaryArithmeticOpConversion<arith::DivUIOp, IREE::VM::DivI32UOp, + IREE::VM::DivI64UOp>, + BinaryArithmeticOpConversion<arith::MulIOp, IREE::VM::MulI32Op, + IREE::VM::MulI64Op>, + BinaryArithmeticOpConversion<arith::RemSIOp, IREE::VM::RemI32SOp, + IREE::VM::RemI64SOp>, + BinaryArithmeticOpConversion<arith::RemUIOp, IREE::VM::RemI32UOp, + IREE::VM::RemI64UOp>, + BinaryArithmeticOpConversion<arith::SubIOp, IREE::VM::SubI32Op, + IREE::VM::SubI64Op>, + BinaryArithmeticOpConversion<arith::AndIOp, IREE::VM::AndI32Op, + IREE::VM::AndI64Op>, + BinaryArithmeticOpConversion<arith::OrIOp, IREE::VM::OrI32Op, + IREE::VM::OrI64Op>, + BinaryArithmeticOpConversion<arith::XOrIOp, IREE::VM::XorI32Op, + IREE::VM::XorI64Op>>(typeConverter, + context); // Floating-point arithmetic ops. - patterns.insert<UnaryArithmeticOpConversion<AbsFOp, IREE::VM::AbsF32Op, - IREE::VM::AbsF64Op>, - BinaryArithmeticOpConversion<AddFOp, IREE::VM::AddF32Op, - IREE::VM::AddF64Op>, - UnaryArithmeticOpConversion<CeilFOp, IREE::VM::CeilF32Op, - IREE::VM::CeilF64Op>, - UnaryArithmeticOpConversion<FloorFOp, IREE::VM::FloorF32Op, - IREE::VM::FloorF64Op>, - BinaryArithmeticOpConversion<DivFOp, IREE::VM::DivF32Op, - IREE::VM::DivF64Op>, - BinaryArithmeticOpConversion<MulFOp, IREE::VM::MulF32Op, - IREE::VM::MulF64Op>, - UnaryArithmeticOpConversion<NegFOp, IREE::VM::NegF32Op, - IREE::VM::NegF64Op>, - BinaryArithmeticOpConversion<RemFOp, IREE::VM::RemF32Op, - IREE::VM::RemF64Op>, - BinaryArithmeticOpConversion<SubFOp, IREE::VM::SubF32Op, - IREE::VM::SubF64Op>>( - typeConverter, context); + patterns + .insert<UnaryArithmeticOpConversion<math::AbsOp, IREE::VM::AbsF32Op, + IREE::VM::AbsF64Op>, + BinaryArithmeticOpConversion<arith::AddFOp, IREE::VM::AddF32Op, + IREE::VM::AddF64Op>, + UnaryArithmeticOpConversion<math::CeilOp, IREE::VM::CeilF32Op, + IREE::VM::CeilF64Op>, + UnaryArithmeticOpConversion<math::FloorOp, IREE::VM::FloorF32Op, + IREE::VM::FloorF64Op>, + BinaryArithmeticOpConversion<arith::DivFOp, IREE::VM::DivF32Op, + IREE::VM::DivF64Op>, + BinaryArithmeticOpConversion<arith::MulFOp, IREE::VM::MulF32Op, + IREE::VM::MulF64Op>, + UnaryArithmeticOpConversion<arith::NegFOp, IREE::VM::NegF32Op, + IREE::VM::NegF64Op>, + BinaryArithmeticOpConversion<arith::RemFOp, IREE::VM::RemF32Op, + IREE::VM::RemF64Op>, + BinaryArithmeticOpConversion<arith::SubFOp, IREE::VM::SubF32Op, + IREE::VM::SubF64Op>>(typeConverter, + context); // Floating-point conversion ops. patterns.insert<SIToFPOpConversion, UIToFPOpConversion, FPToSIOpConversion, @@ -867,13 +874,14 @@ context); // Shift ops. - patterns.insert< - ShiftArithmeticOpConversion<ShiftLeftOp, IREE::VM::ShlI32Op, - IREE::VM::ShlI64Op>, - ShiftArithmeticOpConversion<SignedShiftRightOp, IREE::VM::ShrI32SOp, - IREE::VM::ShrI64SOp>, - ShiftArithmeticOpConversion<UnsignedShiftRightOp, IREE::VM::ShrI32UOp, - IREE::VM::ShrI64UOp>>(typeConverter, context); + patterns + .insert<ShiftArithmeticOpConversion<arith::ShLIOp, IREE::VM::ShlI32Op, + IREE::VM::ShlI64Op>, + ShiftArithmeticOpConversion<arith::ShRSIOp, IREE::VM::ShrI32SOp, + IREE::VM::ShrI64SOp>, + ShiftArithmeticOpConversion<arith::ShRUIOp, IREE::VM::ShrI32UOp, + IREE::VM::ShrI64UOp>>(typeConverter, + context); } } // namespace iree_compiler
diff --git a/iree/compiler/Dialect/VM/Conversion/StandardToVM/ConvertStandardToVMTest.cpp b/iree/compiler/Dialect/VM/Conversion/StandardToVM/ConvertStandardToVMTest.cpp index 1938dea..c3ee86c 100644 --- a/iree/compiler/Dialect/VM/Conversion/StandardToVM/ConvertStandardToVMTest.cpp +++ b/iree/compiler/Dialect/VM/Conversion/StandardToVM/ConvertStandardToVMTest.cpp
@@ -36,7 +36,8 @@ void runOnOperation() override { ConversionTarget target(getContext()); target.addLegalDialect<IREE::VM::VMDialect>(); - target.addIllegalDialect<StandardOpsDialect>(); + target.addIllegalDialect<StandardOpsDialect, + mlir::arith::ArithmeticDialect>(); IREE::VM::TypeConverter typeConverter( IREE::VM::getTargetOptionsFromFlags());
diff --git a/iree/compiler/Dialect/VM/Conversion/StandardToVM/test/arithmetic_ops.mlir b/iree/compiler/Dialect/VM/Conversion/StandardToVM/test/arithmetic_ops.mlir index 46e00c2..fb455bb 100644 --- a/iree/compiler/Dialect/VM/Conversion/StandardToVM/test/arithmetic_ops.mlir +++ b/iree/compiler/Dialect/VM/Conversion/StandardToVM/test/arithmetic_ops.mlir
@@ -10,7 +10,7 @@ // CHECK-SAME: %[[ARG1:[a-zA-Z0-9$._-]+]] func @my_fn(%arg0: i32, %arg1: i32) -> (i32) { // CHECK: vm.add.i32 %[[ARG0]], %[[ARG1]] - %0 = addi %arg0, %arg1 : i32 + %0 = arith.addi %arg0, %arg1 : i32 return %0 : i32 } } @@ -27,7 +27,7 @@ // CHECK-SAME: %[[ARG1:[a-zA-Z0-9$._-]+]] func @my_fn(%arg0: i32, %arg1: i32) -> (i32) { // CHECK: vm.div.i32.s %[[ARG0]], %[[ARG1]] - %0 = divi_signed %arg0, %arg1 : i32 + %0 = arith.divsi %arg0, %arg1 : i32 return %0 : i32 } } @@ -44,7 +44,7 @@ // CHECK-SAME: %[[ARG1:[a-zA-Z0-9$._-]+]] func @my_fn(%arg0: i32, %arg1: i32) -> (i32) { // CHECK: vm.div.i32.u %[[ARG0]], %[[ARG1]] - %0 = divi_unsigned %arg0, %arg1 : i32 + %0 = arith.divui %arg0, %arg1 : i32 return %0 : i32 } } @@ -61,7 +61,7 @@ // CHECK-SAME: %[[ARG1:[a-zA-Z0-9$._-]+]] func @my_fn(%arg0: i32, %arg1: i32) -> (i32) { // CHECK: vm.mul.i32 %[[ARG0]], %[[ARG1]] - %0 = muli %arg0, %arg1 : i32 + %0 = arith.muli %arg0, %arg1 : i32 return %0 : i32 } } @@ -78,7 +78,7 @@ // CHECK-SAME: %[[ARG1:[a-zA-Z0-9$._-]+]] func @my_fn(%arg0: i32, %arg1: i32) -> (i32) { // CHECK: vm.rem.i32.s %[[ARG0]], %[[ARG1]] - %0 = remi_signed %arg0, %arg1 : i32 + %0 = arith.remsi %arg0, %arg1 : i32 return %0 : i32 } } @@ -95,7 +95,7 @@ // CHECK-SAME: %[[ARG1:[a-zA-Z0-9$._-]+]] func @my_fn(%arg0: i32, %arg1: i32) -> (i32) { // CHECK: vm.rem.i32.u %[[ARG0]], %[[ARG1]] - %0 = remi_unsigned %arg0, %arg1 : i32 + %0 = arith.remui %arg0, %arg1 : i32 return %0 : i32 } } @@ -112,7 +112,7 @@ // CHECK-SAME: %[[ARG1:[a-zA-Z0-9$._-]+]] func @my_fn(%arg0: i32, %arg1: i32) -> (i32) { // CHECK: vm.sub.i32 %[[ARG0]], %[[ARG1]] - %0 = subi %arg0, %arg1 : i32 + %0 = arith.subi %arg0, %arg1 : i32 return %0 : i32 } } @@ -129,7 +129,7 @@ // CHECK-SAME: %[[ARG1:[a-zA-Z0-9$._-]+]] func @my_fn(%arg0: i32, %arg1: i32) -> (i32) { // CHECK: vm.and.i32 %[[ARG0]], %[[ARG1]] - %0 = and %arg0, %arg1 : i32 + %0 = arith.andi %arg0, %arg1 : i32 return %0 : i32 } } @@ -146,7 +146,7 @@ // CHECK-SAME: %[[ARG1:[a-zA-Z0-9$._-]+]] func @my_fn(%arg0: i32, %arg1: i32) -> (i32) { // CHECK: vm.or.i32 %[[ARG0]], %[[ARG1]] - %0 = or %arg0, %arg1 : i32 + %0 = arith.ori %arg0, %arg1 : i32 return %0 : i32 } } @@ -163,7 +163,7 @@ // CHECK-SAME: %[[ARG1:[a-zA-Z0-9$._-]+]] func @my_fn(%arg0: i32, %arg1: i32) -> (i32) { // CHECK: vm.xor.i32 %[[ARG0]], %[[ARG1]] - %0 = xor %arg0, %arg1 : i32 + %0 = arith.xori %arg0, %arg1 : i32 return %0 : i32 } } @@ -178,9 +178,9 @@ // CHECK: vm.func private @my_fn // CHECK-SAME: %[[ARG0:[a-zA-Z0-9$._-]+]] func @my_fn(%arg0: i32) -> (i32) { - %c3 = constant 3 : i32 + %c3 = arith.constant 3 : i32 // CHECK: vm.shl.i32 %[[ARG0]], %c3 : i32 - %1 = shift_left %arg0, %c3 : i32 + %1 = arith.shli %arg0, %c3 : i32 return %1 : i32 } } @@ -195,11 +195,11 @@ // CHECK: vm.func private @my_fn // CHECK-SAME: %[[ARG0:[a-zA-Z0-9$._-]+]] func @my_fn(%arg0: i32) -> (i32) { - %c3 = constant 3 : i32 + %c3 = arith.constant 3 : i32 // CHECK: %[[T:.+]] = vm.shr.i32.s %[[ARG0]], %c3 : i32 - %1 = shift_right_signed %arg0, %c3 : i32 + %1 = arith.shrsi %arg0, %c3 : i32 // CHECK: vm.shr.i32.u %[[T]], %c3 : i32 - %2 = shift_right_unsigned %1, %c3 : i32 + %2 = arith.shrui %1, %c3 : i32 return %2 : i32 } }
diff --git a/iree/compiler/Dialect/VM/Conversion/StandardToVM/test/assignment_ops.mlir b/iree/compiler/Dialect/VM/Conversion/StandardToVM/test/assignment_ops.mlir index fd67a99..c0ceaa2 100644 --- a/iree/compiler/Dialect/VM/Conversion/StandardToVM/test/assignment_ops.mlir +++ b/iree/compiler/Dialect/VM/Conversion/StandardToVM/test/assignment_ops.mlir
@@ -14,7 +14,7 @@ // Note that in std, cmp returns an i1 and this relies on the dialect // conversion framework promoting that to i32. // CHECK: %[[CMP:[a-zA-Z0-9$._-]+]] = vm.cmp.eq.i32 - %1 = cmpi eq, %arg0, %arg1 : i32 + %1 = arith.cmpi eq, %arg0, %arg1 : i32 // CHECK: vm.select.i32 %[[CMP]], %[[ARG2]], %[[ARG3]] : i32 %2 = select %1, %arg2, %arg3 : i32 return %2 : i32 @@ -37,7 +37,7 @@ // Note that in std, cmp returns an i1 and this relies on the dialect // conversion framework promoting that to i32. // CHECK: %[[CMP:[a-zA-Z0-9$._-]+]] = vm.cmp.eq.i32 - %1 = cmpi eq, %arg0, %arg1 : index + %1 = arith.cmpi eq, %arg0, %arg1 : index // CHECK: vm.select.i32 %[[CMP]], %[[ARG2]], %[[ARG3]] : i32 %2 = select %1, %arg2, %arg3 : index return %2 : index
diff --git a/iree/compiler/Dialect/VM/Conversion/StandardToVM/test/comparison_ops.mlir b/iree/compiler/Dialect/VM/Conversion/StandardToVM/test/comparison_ops.mlir index 3d3a6f5..2ea6516 100644 --- a/iree/compiler/Dialect/VM/Conversion/StandardToVM/test/comparison_ops.mlir +++ b/iree/compiler/Dialect/VM/Conversion/StandardToVM/test/comparison_ops.mlir
@@ -10,7 +10,7 @@ // CHECK-SAME: %[[ARG1:[a-zA-Z0-9$._-]+]] func @my_fn(%arg0: i32, %arg1 : i32) -> (i1) { // CHECK: vm.cmp.eq.i32 %[[ARG0]], %[[ARG1]] : i32 - %1 = cmpi eq, %arg0, %arg1 : i32 + %1 = arith.cmpi eq, %arg0, %arg1 : i32 return %1 : i1 } } @@ -27,7 +27,7 @@ // CHECK-SAME: %[[ARG1:[a-zA-Z0-9$._-]+]] func @my_fn(%arg0: i32, %arg1 : i32) -> (i1) { // CHECK: vm.cmp.ne.i32 %[[ARG0]], %[[ARG1]] : i32 - %1 = cmpi ne, %arg0, %arg1 : i32 + %1 = arith.cmpi ne, %arg0, %arg1 : i32 return %1 : i1 } } @@ -44,7 +44,7 @@ // CHECK-SAME: %[[ARG1:[a-zA-Z0-9$._-]+]] func @my_fn(%arg0: i32, %arg1 : i32) -> (i1) { // CHECK: vm.cmp.lt.i32.s %[[ARG0]], %[[ARG1]] : i32 - %1 = cmpi slt, %arg0, %arg1 : i32 + %1 = arith.cmpi slt, %arg0, %arg1 : i32 return %1 : i1 } } @@ -61,7 +61,7 @@ // CHECK-SAME: %[[ARG1:[a-zA-Z0-9$._-]+]] func @my_fn(%arg0: i32, %arg1 : i32) -> (i1) { // CHECK: vm.cmp.lte.i32.s %[[ARG0]], %[[ARG1]] : i32 - %1 = cmpi sle, %arg0, %arg1 : i32 + %1 = arith.cmpi sle, %arg0, %arg1 : i32 return %1 : i1 } } @@ -78,7 +78,7 @@ // CHECK-SAME: %[[ARG1:[a-zA-Z0-9$._-]+]] func @my_fn(%arg0: i32, %arg1 : i32) -> (i1) { // CHECK: vm.cmp.gt.i32.s %[[ARG0]], %[[ARG1]] : i32 - %1 = cmpi sgt, %arg0, %arg1 : i32 + %1 = arith.cmpi sgt, %arg0, %arg1 : i32 return %1 : i1 } } @@ -95,7 +95,7 @@ // CHECK-SAME: %[[ARG1:[a-zA-Z0-9$._-]+]] func @my_fn(%arg0: i32, %arg1 : i32) -> (i1) { // CHECK: vm.cmp.gte.i32.s %[[ARG0]], %[[ARG1]] : i32 - %1 = cmpi sge, %arg0, %arg1 : i32 + %1 = arith.cmpi sge, %arg0, %arg1 : i32 return %1 : i1 } } @@ -112,7 +112,7 @@ // CHECK-SAME: %[[ARG1:[a-zA-Z0-9$._-]+]] func @my_fn(%arg0: i32, %arg1 : i32) -> (i1) { // CHECK: vm.cmp.lt.i32.u %[[ARG0]], %[[ARG1]] : i32 - %1 = cmpi ult, %arg0, %arg1 : i32 + %1 = arith.cmpi ult, %arg0, %arg1 : i32 return %1 : i1 } } @@ -129,7 +129,7 @@ // CHECK-SAME: %[[ARG1:[a-zA-Z0-9$._-]+]] func @my_fn(%arg0: i32, %arg1 : i32) -> (i1) { // CHECK: vm.cmp.lte.i32.u %[[ARG0]], %[[ARG1]] : i32 - %1 = cmpi ule, %arg0, %arg1 : i32 + %1 = arith.cmpi ule, %arg0, %arg1 : i32 return %1 : i1 } } @@ -146,7 +146,7 @@ // CHECK-SAME: %[[ARG1:[a-zA-Z0-9$._-]+]] func @my_fn(%arg0: i32, %arg1 : i32) -> (i1) { // CHECK: vm.cmp.gt.i32.u %[[ARG0]], %[[ARG1]] : i32 - %1 = cmpi ugt, %arg0, %arg1 : i32 + %1 = arith.cmpi ugt, %arg0, %arg1 : i32 return %1 : i1 } } @@ -163,7 +163,7 @@ // CHECK-SAME: %[[ARG1:[a-zA-Z0-9$._-]+]] func @my_fn(%arg0: i32, %arg1 : i32) -> (i1) { // CHECK: vm.cmp.gte.i32.u %[[ARG0]], %[[ARG1]] : i32 - %1 = cmpi uge, %arg0, %arg1 : i32 + %1 = arith.cmpi uge, %arg0, %arg1 : i32 return %1 : i1 } }
diff --git a/iree/compiler/Dialect/VM/Conversion/StandardToVM/test/const_ops.mlir b/iree/compiler/Dialect/VM/Conversion/StandardToVM/test/const_ops.mlir index 97f2d18..38b4def 100644 --- a/iree/compiler/Dialect/VM/Conversion/StandardToVM/test/const_ops.mlir +++ b/iree/compiler/Dialect/VM/Conversion/StandardToVM/test/const_ops.mlir
@@ -7,7 +7,7 @@ module { func @non_zero() -> (i32) { // CHECK: vm.const.i32 1 : i32 - %1 = constant 1 : i32 + %1 = arith.constant 1 : i32 return %1 : i32 } } @@ -21,7 +21,7 @@ module { func @zero() -> (i32) { // CHECK: vm.const.i32.zero : i32 - %1 = constant 0 : i32 + %1 = arith.constant 0 : i32 return %1 : i32 } } @@ -35,7 +35,7 @@ module { func @non_zero() -> (f32) { // CHECK: vm.const.f32 1.000000e+00 : f32 - %1 = constant 1. : f32 + %1 = arith.constant 1. : f32 return %1 : f32 } } @@ -49,7 +49,7 @@ module { func @zero() -> (f32) { // CHECK: vm.const.f32.zero : f32 - %1 = constant 0. : f32 + %1 = arith.constant 0. : f32 return %1 : f32 } }
diff --git a/iree/compiler/Dialect/VM/Conversion/StandardToVM/test/control_flow_ops.mlir b/iree/compiler/Dialect/VM/Conversion/StandardToVM/test/control_flow_ops.mlir index 95d653a..8e01870 100644 --- a/iree/compiler/Dialect/VM/Conversion/StandardToVM/test/control_flow_ops.mlir +++ b/iree/compiler/Dialect/VM/Conversion/StandardToVM/test/control_flow_ops.mlir
@@ -114,14 +114,14 @@ module { // CHECK: vm.func private @my_fn func @my_fn(%arg0: i32) -> i32 { - %zero = constant 0 : i32 + %zero = arith.constant 0 : i32 // CHECK: %[[COND:.+]] = vm.cmp.ne.i32 - %cond = cmpi ne, %arg0, %zero : i32 + %cond = arith.cmpi ne, %arg0, %zero : i32 // CHECK: %[[STATUS:.+]] = vm.const.i32 9 // CHECK: %[[INVCOND:.+]] = vm.xor.i32 %[[COND]], %c1 // CHECK: vm.cond_fail %[[INVCOND]], %[[STATUS]], "Assertion failed" assert %cond, "Assertion failed" - %sum = addi %arg0, %arg0 : i32 + %sum = arith.addi %arg0, %arg0 : i32 return %sum : i32 } }
diff --git a/iree/compiler/Dialect/VM/Conversion/StandardToVM/test/conversion_ops.mlir b/iree/compiler/Dialect/VM/Conversion/StandardToVM/test/conversion_ops.mlir index 4bb5eb0..a016f09 100644 --- a/iree/compiler/Dialect/VM/Conversion/StandardToVM/test/conversion_ops.mlir +++ b/iree/compiler/Dialect/VM/Conversion/StandardToVM/test/conversion_ops.mlir
@@ -7,7 +7,7 @@ // CHECK-SAME: %[[ARG0:.+]]: i32 func @my_fn(%arg0 : i32) -> (f32) { // CHECK: vm.bitcast.i32.f32 %[[ARG0]] : i32 -> f32 - %1 = bitcast %arg0 : i32 to f32 + %1 = arith.bitcast %arg0 : i32 to f32 return %1 : f32 } } @@ -22,7 +22,7 @@ // CHECK-SAME: %[[ARG0:.+]]: f32 func @my_fn(%arg0 : f32) -> (i32) { // CHECK: vm.bitcast.f32.i32 %[[ARG0]] : f32 -> i32 - %1 = bitcast %arg0 : f32 to i32 + %1 = arith.bitcast %arg0 : f32 to i32 return %1 : i32 } }
diff --git a/iree/compiler/Dialect/VM/Conversion/TypeConverter.cpp b/iree/compiler/Dialect/VM/Conversion/TypeConverter.cpp index 66c0cb8..b119353 100644 --- a/iree/compiler/Dialect/VM/Conversion/TypeConverter.cpp +++ b/iree/compiler/Dialect/VM/Conversion/TypeConverter.cpp
@@ -140,7 +140,7 @@ if (inputs.size() != 1 || !inputs.front().getType().isa<IntegerType>()) { return nullptr; } - return builder.create<IndexCastOp>(loc, type, inputs.front()); + return builder.create<arith::IndexCastOp>(loc, type, inputs.front()); }); addTargetMaterialization(
diff --git a/iree/compiler/Dialect/VM/Conversion/UtilToVM/test/list_ops.mlir b/iree/compiler/Dialect/VM/Conversion/UtilToVM/test/list_ops.mlir index eb2727f..06e2b76 100644 --- a/iree/compiler/Dialect/VM/Conversion/UtilToVM/test/list_ops.mlir +++ b/iree/compiler/Dialect/VM/Conversion/UtilToVM/test/list_ops.mlir
@@ -6,23 +6,23 @@ // CHECK-SAME: (%[[BUFFER_VIEW:.+]]: !vm.ref<!hal.buffer_view>) func @my_fn(%buffer_view: !hal.buffer_view) { // CHECK: %[[CAPACITY:.+]] = vm.const.i32 5 - %capacity = constant 5 : index + %capacity = arith.constant 5 : index // CHECK: %[[LIST:.+]] = vm.list.alloc %[[CAPACITY]] : (i32) -> !vm.list<?> %list = util.list.create %capacity : !util.list<?> // CHECK: %[[NEW_SIZE:.+]] = vm.const.i32 100 - %new_size = constant 100 : index + %new_size = arith.constant 100 : index // CHECK: vm.list.resize %[[LIST]], %[[NEW_SIZE]] : (!vm.list<?>, i32) util.list.resize %list, %new_size : !util.list<?> - %c10 = constant 10 : index - %c11 = constant 11 : index + %c10 = arith.constant 10 : index + %c11 = arith.constant 11 : index // CHECK: = vm.list.get.i32 %[[LIST]], %c10 : (!vm.list<?>, i32) -> i32 %0 = util.list.get %list[%c10] : !util.list<?> -> i32 // CHECK: %[[NEW_I32_VALUE:.+]] = vm.const.i32 101 - %new_i32_value = constant 101 : i32 + %new_i32_value = arith.constant 101 : i32 // CHECK: vm.list.set.i32 %[[LIST]], %c10, %[[NEW_I32_VALUE]] : (!vm.list<?>, i32, i32) util.list.set %list[%c10], %new_i32_value : i32 -> !util.list<?>
diff --git a/iree/compiler/Dialect/VM/Conversion/UtilToVM/test/status_ops.mlir b/iree/compiler/Dialect/VM/Conversion/UtilToVM/test/status_ops.mlir index 925e9fd..9303dce 100644 --- a/iree/compiler/Dialect/VM/Conversion/UtilToVM/test/status_ops.mlir +++ b/iree/compiler/Dialect/VM/Conversion/UtilToVM/test/status_ops.mlir
@@ -3,7 +3,7 @@ // CHECK-LABEL: vm.func private @status_check_ok func @status_check_ok() { // CHECK: %[[CODE:.+]] = - %statusCode = constant 1 : i32 + %statusCode = arith.constant 1 : i32 // CHECK: vm.cond_fail %[[CODE]] util.status.check_ok %statusCode return @@ -14,7 +14,7 @@ // CHECK-LABEL: vm.func private @status_check_ok_with_message func @status_check_ok_with_message() { // CHECK: %[[CODE:.+]] = - %statusCode = constant 1 : i32 + %statusCode = arith.constant 1 : i32 // CHECK: vm.cond_fail %[[CODE]], "failure message" util.status.check_ok %statusCode, "failure message" return
diff --git a/iree/compiler/Dialect/VM/Conversion/VMToEmitC/ConvertVMToEmitC.cpp b/iree/compiler/Dialect/VM/Conversion/VMToEmitC/ConvertVMToEmitC.cpp index 8c4c1fd..a4f49e3 100644 --- a/iree/compiler/Dialect/VM/Conversion/VMToEmitC/ConvertVMToEmitC.cpp +++ b/iree/compiler/Dialect/VM/Conversion/VMToEmitC/ConvertVMToEmitC.cpp
@@ -13,6 +13,7 @@ #include "iree/compiler/Dialect/VM/Utils/CallingConvention.h" #include "llvm/ADT/TypeSwitch.h" #include "mlir/Dialect/EmitC/IR/EmitC.h" +#include "mlir/Dialect/Math/IR/Math.h" #include "mlir/Dialect/StandardOps/IR/Ops.h" #include "mlir/IR/Builders.h" #include "mlir/IR/BuiltinDialect.h" @@ -3244,7 +3245,8 @@ OperationPass<IREE::VM::ModuleOp>> { void getDependentDialects(DialectRegistry ®istry) const override { registry.insert<mlir::emitc::EmitCDialect, mlir::BuiltinDialect, - mlir::StandardOpsDialect, IREE::Util::UtilDialect>(); + mlir::StandardOpsDialect, mlir::arith::ArithmeticDialect, + mlir::math::MathDialect, IREE::Util::UtilDialect>(); } StringRef getArgument() const override { return "iree-convert-vm-to-emitc"; } @@ -3291,8 +3293,9 @@ populateVMToEmitCPatterns(&getContext(), target, typeConverter, patterns, vmAnalysisCache); - target.addLegalDialect<emitc::EmitCDialect, mlir::BuiltinDialect, - mlir::StandardOpsDialect>(); + target.addLegalDialect< + emitc::EmitCDialect, mlir::BuiltinDialect, mlir::StandardOpsDialect, + mlir::arith::ArithmeticDialect, mlir::math::MathDialect>(); target.addDynamicallyLegalOp<mlir::FuncOp>([&](mlir::FuncOp op) { return typeConverter.isSignatureLegal(op.getType());
diff --git a/iree/compiler/Dialect/VM/IR/test/control_flow_folding.mlir b/iree/compiler/Dialect/VM/IR/test/control_flow_folding.mlir index 69a8115..cdc37ee 100644 --- a/iree/compiler/Dialect/VM/IR/test/control_flow_folding.mlir +++ b/iree/compiler/Dialect/VM/IR/test/control_flow_folding.mlir
@@ -86,8 +86,8 @@ // CHECK-LABEL: @cond_fail_to_cond_br_fail // CHECK-SAME: %[[COND:.+]]: vm.func @cond_fail_to_cond_br_fail(%cond : i32) { - // CHECK-DAG: %[[CODE2:.+]] = constant 2 - %code2 = constant 2 : i32 + // CHECK-DAG: %[[CODE2:.+]] = arith.constant 2 + %code2 = arith.constant 2 : i32 // CHECK: vm.cond_br %[[COND]], ^bb2(%[[CODE2]] : i32), ^bb1 vm.cond_fail %cond, %code2, "message" // CHECK-NEXT: ^bb1:
diff --git a/iree/compiler/Dialect/VM/IR/test/control_flow_ops.mlir b/iree/compiler/Dialect/VM/IR/test/control_flow_ops.mlir index dc37a31..2445c45 100644 --- a/iree/compiler/Dialect/VM/IR/test/control_flow_ops.mlir +++ b/iree/compiler/Dialect/VM/IR/test/control_flow_ops.mlir
@@ -161,15 +161,15 @@ vm.module @my_module { // CHECK-LABEL: @fail vm.func @fail() { - // CHECK-DAG: %[[CODE1:.+]] = constant 1 - %code1 = constant 1 : i32 + // CHECK-DAG: %[[CODE1:.+]] = arith.constant 1 + %code1 = arith.constant 1 : i32 // CHECK: vm.fail %[[CODE1]] vm.fail %code1 } // CHECK-LABEL: @fail_message vm.func @fail_message() { - // CHECK-DAG: %[[CODE2:.+]] = constant 2 - %code2 = constant 2 : i32 + // CHECK-DAG: %[[CODE2:.+]] = arith.constant 2 + %code2 = arith.constant 2 : i32 // CHECK: vm.fail %[[CODE2]], "message" vm.fail %code2, "message" } @@ -181,8 +181,8 @@ // CHECK-LABEL: @cond_fail // CHECK-SAME: %[[COND:.+]]: vm.func @cond_fail(%cond : i32) { - // CHECK-DAG: %[[CODE1:.+]] = constant 1 - %code1 = constant 1 : i32 + // CHECK-DAG: %[[CODE1:.+]] = arith.constant 1 + %code1 = arith.constant 1 : i32 // CHECK: vm.cond_fail %[[COND]], %[[CODE1]] vm.cond_fail %cond, %code1 vm.return @@ -190,24 +190,24 @@ // CHECK-LABEL: @cond_fail_message // CHECK-SAME: %[[COND:.+]]: vm.func @cond_fail_message(%cond : i32) { - // CHECK-DAG: %[[CODE2:.+]] = constant 2 - %code2 = constant 2 : i32 + // CHECK-DAG: %[[CODE2:.+]] = arith.constant 2 + %code2 = arith.constant 2 : i32 // CHECK: vm.cond_fail %[[COND]], %[[CODE2]], "message" vm.cond_fail %cond, %code2, "message" vm.return } // CHECK-LABEL: @cond_fail_no_condition vm.func @cond_fail_no_condition() { - // CHECK-DAG: %[[CODE3:.+]] = constant 3 - %code3 = constant 3 : i32 + // CHECK-DAG: %[[CODE3:.+]] = arith.constant 3 + %code3 = arith.constant 3 : i32 // CHECK: vm.cond_fail %[[CODE3]] vm.cond_fail %code3 vm.return } // CHECK-LABEL: @cond_fail_no_condition_with_message vm.func @cond_fail_no_condition_with_message() { - // CHECK-DAG: %[[CODE4:.+]] = constant 4 - %code4 = constant 4 : i32 + // CHECK-DAG: %[[CODE4:.+]] = arith.constant 4 + %code4 = arith.constant 4 : i32 // CHECK: vm.cond_fail %[[CODE4]] vm.cond_fail %code4, "message" vm.return
diff --git a/iree/compiler/Dialect/VM/Transforms/Conversion.cpp b/iree/compiler/Dialect/VM/Transforms/Conversion.cpp index fe293b3..3cfbb39 100644 --- a/iree/compiler/Dialect/VM/Transforms/Conversion.cpp +++ b/iree/compiler/Dialect/VM/Transforms/Conversion.cpp
@@ -84,8 +84,8 @@ void getDependentDialects(DialectRegistry ®istry) const override { registry.insert<IREE::Util::UtilDialect, IREE::VM::VMDialect, - StandardOpsDialect, math::MathDialect, AffineDialect, - memref::MemRefDialect>(); + StandardOpsDialect, mlir::arith::ArithmeticDialect, + math::MathDialect, AffineDialect, memref::MemRefDialect>(); } StringRef getArgument() const override { return "iree-vm-conversion"; } @@ -141,7 +141,8 @@ populateAffineToStdConversionPatterns(conversionPatterns); conversionPatterns.insert<ElideTieShapeOp>(context); - conversionTarget.addIllegalDialect<StandardOpsDialect>(); + conversionTarget.addIllegalDialect<StandardOpsDialect, + mlir::arith::ArithmeticDialect>(); conversionTarget.addIllegalDialect<AffineDialect>(); conversionTarget.addIllegalDialect<math::MathDialect>();
diff --git a/iree/compiler/Dialect/Vulkan/Utils/test/target_env_conversion.mlir b/iree/compiler/Dialect/Vulkan/Utils/test/target_env_conversion.mlir index 4987e18..0b9ac53 100644 --- a/iree/compiler/Dialect/Vulkan/Utils/test/target_env_conversion.mlir +++ b/iree/compiler/Dialect/Vulkan/Utils/test/target_env_conversion.mlir
@@ -18,9 +18,9 @@ flow.dispatch.entry @simpleMath_dispatch_0 attributes {workgroup_rank = 3 : index} builtin.module { func @simpleMath_dispatch_0(%arg0: !flow.dispatch.tensor<readonly:4xf32>, %arg1: !flow.dispatch.tensor<writeonly:4xf32>) { - %c4 = constant 4 : index - %c1 = constant 1 : index - %c0 = constant 0 : index + %c4 = arith.constant 4 : index + %c1 = arith.constant 1 : index + %c0 = arith.constant 0 : index %0 = flow.dispatch.tensor.load %arg0, offsets = [%c0], sizes = [%c4], strides = [%c1] : !flow.dispatch.tensor<readonly:4xf32> -> tensor<4xf32> %1 = linalg.init_tensor [4] : tensor<4xf32> %2 = linalg.generic { @@ -29,7 +29,7 @@ ins(%0 : tensor<4xf32>) outs(%1 : tensor<4xf32>) { ^bb0(%arg3: f32, %arg4: f32): // no predecessors - %3 = addf %arg3, %arg3 : f32 + %3 = arith.addf %arg3, %arg3 : f32 linalg.yield %3 : f32 } -> tensor<4xf32> flow.dispatch.tensor.store %2, %arg1, offsets = [%c0], sizes = [%c4], strides = [%c1] : tensor<4xf32> -> !flow.dispatch.tensor<writeonly:4xf32>
diff --git a/iree/compiler/InputConversion/Common/IREEImportPublic.cpp b/iree/compiler/InputConversion/Common/IREEImportPublic.cpp index 15efbed..c815be2 100644 --- a/iree/compiler/InputConversion/Common/IREEImportPublic.cpp +++ b/iree/compiler/InputConversion/Common/IREEImportPublic.cpp
@@ -45,7 +45,7 @@ void getDependentDialects(DialectRegistry ®istry) const override { registry.insert<mlir::iree::IREEDialect, IREE::Flow::FlowDialect, IREE::HAL::HALDialect, IREE::Util::UtilDialect, - mlir::StandardOpsDialect>(); + mlir::StandardOpsDialect, mlir::arith::ArithmeticDialect>(); } void runOnOperation() override; };
diff --git a/iree/compiler/InputConversion/Common/test/iree_import_public.mlir b/iree/compiler/InputConversion/Common/test/iree_import_public.mlir index 4cf1713..8a978a1 100644 --- a/iree/compiler/InputConversion/Common/test/iree_import_public.mlir +++ b/iree/compiler/InputConversion/Common/test/iree_import_public.mlir
@@ -212,8 +212,8 @@ builtin.module @global_store { iree.global private mutable @v_stored : tensor<4xi32> func @stored() { - // CHECK: %[[CST:.*]] = constant - %cst = constant dense<5> : tensor<4xi32> + // CHECK: %[[CST:.*]] = arith.constant + %cst = arith.constant dense<5> : tensor<4xi32> // CHECK: util.global.store %[[CST]], @v_stored : tensor<4xi32> iree.global.store %cst, @v_stored : tensor<4xi32> return
diff --git a/iree/compiler/InputConversion/Common/test/top_level_scf_to_cfg.mlir b/iree/compiler/InputConversion/Common/test/top_level_scf_to_cfg.mlir index c06f57c..4b04669 100644 --- a/iree/compiler/InputConversion/Common/test/top_level_scf_to_cfg.mlir +++ b/iree/compiler/InputConversion/Common/test/top_level_scf_to_cfg.mlir
@@ -7,13 +7,13 @@ // (multiple basic blocks). The -iree-top-level-scf-to-cfg should not touch it. #map = affine_map<(d0, d1, d2, d3) -> (d0, d1, d2, d3)> func @generic_nested_for(%arg0: tensor<?x?x?x?xi32>, %arg1: tensor<?x?x?x?xi32>, %out0: tensor<?x?x?x?xi32>) -> tensor<?x?x?x?xi32> { - %c0 = constant 0 : index - %c1 = constant 1 : index - %c6 = constant 6 : index - %c-1_i32 = constant -1 : i32 - %c0_i32 = constant 0 : i32 - %c1_i32 = constant 1 : i32 - %c2_i32 = constant 2 : i32 + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + %c6 = arith.constant 6 : index + %c-1_i32 = arith.constant -1 : i32 + %c0_i32 = arith.constant 0 : i32 + %c1_i32 = arith.constant 1 : i32 + %c2_i32 = arith.constant 2 : i32 // CHECK: linalg.generic // CHECK: scf.for // CHECK: linalg.yield @@ -21,19 +21,19 @@ ins(%arg0, %arg1 : tensor<?x?x?x?xi32>, tensor<?x?x?x?xi32>) outs(%out0 : tensor<?x?x?x?xi32>) { ^bb0(%arg2: i32, %arg3: i32, %arg4: i32): // no predecessors %18:3 = scf.for %arg5 = %c0 to %c6 step %c1 iter_args(%arg6 = %c1_i32, %arg7 = %arg2, %arg8 = %arg3) -> (i32, i32, i32) { - %28 = and %arg8, %c1_i32 : i32 - %29 = cmpi eq, %28, %c1_i32 : i32 - %30 = muli %arg6, %arg7 : i32 + %28 = arith.andi %arg8, %c1_i32 : i32 + %29 = arith.cmpi eq, %28, %c1_i32 : i32 + %30 = arith.muli %arg6, %arg7 : i32 %31 = select %29, %30, %arg6 : i32 - %32 = muli %arg7, %arg7 : i32 - %33 = shift_right_unsigned %arg8, %c1_i32 : i32 + %32 = arith.muli %arg7, %arg7 : i32 + %33 = arith.shrui %arg8, %c1_i32 : i32 scf.yield %31, %32, %33 : i32, i32, i32 } - %19 = remi_signed %arg3, %c2_i32 : i32 - %20 = cmpi eq, %19, %c0_i32 : i32 - %21 = cmpi slt, %arg3, %c0_i32 : i32 - %22 = cmpi eq, %arg2, %c1_i32 : i32 - %23 = cmpi eq, %arg2, %c-1_i32 : i32 + %19 = arith.remsi %arg3, %c2_i32 : i32 + %20 = arith.cmpi eq, %19, %c0_i32 : i32 + %21 = arith.cmpi slt, %arg3, %c0_i32 : i32 + %22 = arith.cmpi eq, %arg2, %c1_i32 : i32 + %23 = arith.cmpi eq, %arg2, %c-1_i32 : i32 %24 = select %22, %c1_i32, %c0_i32 : i32 %25 = select %20, %c1_i32, %c-1_i32 : i32 %26 = select %23, %25, %24 : i32
diff --git a/iree/compiler/InputConversion/MHLO/BroadcastingToLinalgPatterns.cpp b/iree/compiler/InputConversion/MHLO/BroadcastingToLinalgPatterns.cpp index 2f88c85..1e50720 100644 --- a/iree/compiler/InputConversion/MHLO/BroadcastingToLinalgPatterns.cpp +++ b/iree/compiler/InputConversion/MHLO/BroadcastingToLinalgPatterns.cpp
@@ -66,7 +66,7 @@ Value convertToValue(OpBuilder &builder, Location loc) { if (!isStatic()) return getValue(); - return builder.create<ConstantIndexOp>(loc, getStatic()); + return builder.create<arith::ConstantIndexOp>(loc, getStatic()); } private: @@ -192,8 +192,8 @@ Value lhsExtentValue = lhsDim.convertToValue(builder, loc); Value rhsExtentValue = rhsDim.convertToValue(builder, loc); - Value isEqual = builder.create<CmpIOp>(loc, CmpIPredicate::eq, lhsExtentValue, - rhsExtentValue); + Value isEqual = builder.create<arith::CmpIOp>(loc, arith::CmpIPredicate::eq, + lhsExtentValue, rhsExtentValue); builder.create<AssertOp>( loc, isEqual, builder.getStringAttr("mismatched dynamic broadcast extents")); @@ -269,8 +269,8 @@ // Dynamic check. Value cmpLhsValue = cmpLhs.convertToValue(builder, loc); Value cmpRhsValue = cmpRhs.convertToValue(builder, loc); - Value isEqual = builder.create<CmpIOp>(loc, CmpIPredicate::eq, cmpLhsValue, - cmpRhsValue); + Value isEqual = builder.create<arith::CmpIOp>(loc, arith::CmpIPredicate::eq, + cmpLhsValue, cmpRhsValue); builder.create<AssertOp>( loc, isEqual, builder.getStringAttr("mismatched dynamic broadcast extents")); @@ -730,7 +730,7 @@ "mismatched rank"); for (int i = 0, e = resultType.getRank(); i < e; ++i) { if (resultType.isDynamicDim(i)) { - Value index = rewriter.create<ConstantIndexOp>(loc, i); + Value index = rewriter.create<arith::ConstantIndexOp>(loc, i); targetDims.push_back( rewriter.create<tensor::ExtractOp>(loc, outputShape, index)); }
diff --git a/iree/compiler/InputConversion/MHLO/ConvertComplexToReal.cpp b/iree/compiler/InputConversion/MHLO/ConvertComplexToReal.cpp index fcb3baf..43f2ad1 100644 --- a/iree/compiler/InputConversion/MHLO/ConvertComplexToReal.cpp +++ b/iree/compiler/InputConversion/MHLO/ConvertComplexToReal.cpp
@@ -400,7 +400,8 @@ target.addLegalDialect<mhlo::MhloDialect>(); target.addLegalDialect<chlo::HloClientDialect>(); - target.addLegalDialect<StandardOpsDialect>(); + target + .addLegalDialect<StandardOpsDialect, mlir::arith::ArithmeticDialect>(); // For the test, require that casts fully convert. target.addIllegalOp<mhlo::ComplexOp>();
diff --git a/iree/compiler/InputConversion/MHLO/ConvertMHLOToFlow.cpp b/iree/compiler/InputConversion/MHLO/ConvertMHLOToFlow.cpp index d6a9504..b562208 100644 --- a/iree/compiler/InputConversion/MHLO/ConvertMHLOToFlow.cpp +++ b/iree/compiler/InputConversion/MHLO/ConvertMHLOToFlow.cpp
@@ -26,7 +26,7 @@ using OpRewritePattern::OpRewritePattern; LogicalResult matchAndRewrite(mhlo::ConstOp op, PatternRewriter &rewriter) const override { - rewriter.replaceOpWithNewOp<ConstantOp>(op, op.value()); + rewriter.replaceOpWithNewOp<arith::ConstantOp>(op, op.value()); return success(); } };
diff --git a/iree/compiler/InputConversion/MHLO/ConvertMHLOToLinalgExt.cpp b/iree/compiler/InputConversion/MHLO/ConvertMHLOToLinalgExt.cpp index cd616bb..15ed9af 100644 --- a/iree/compiler/InputConversion/MHLO/ConvertMHLOToLinalgExt.cpp +++ b/iree/compiler/InputConversion/MHLO/ConvertMHLOToLinalgExt.cpp
@@ -326,7 +326,8 @@ values.push_back(b.getI32IntegerAttr(r)); } auto type = RankedTensorType::get({fftLength}, b.getI32Type()); - return b.create<ConstantOp>(type, DenseIntElementsAttr::get(type, values)); + return b.create<arith::ConstantOp>(type, + DenseIntElementsAttr::get(type, values)); } static SmallVector<Value> getBitReversalOrder(ImplicitLocOpBuilder &b, @@ -357,13 +358,14 @@ for (auto i : llvm::seq<unsigned>(0, rank - 1)) { ivs.push_back(b.create<linalg::IndexOp>(loc, i)); } - ivs.push_back(b.create<IndexCastOp>(loc, args[0], b.getIndexType())); + ivs.push_back( + b.create<arith::IndexCastOp>(loc, args[0], b.getIndexType())); b.create<linalg::YieldOp>( loc, b.create<tensor::ExtractOp>(loc, real, ivs).getResult()); }); return { genericOp.getResult(0), - b.create<ConstantOp>( + b.create<arith::ConstantOp>( realType, DenseFPElementsAttr::get( realType, b.getF32FloatAttr(0.0).cast<Attribute>()))}; } @@ -380,8 +382,10 @@ imag.push_back(b.getF32FloatAttr(v.imag())); } auto type = RankedTensorType::get({mh}, b.getF32Type()); - return {b.create<ConstantOp>(type, DenseFPElementsAttr::get(type, real)), - b.create<ConstantOp>(type, DenseFPElementsAttr::get(type, imag))}; + return { + b.create<arith::ConstantOp>(type, DenseFPElementsAttr::get(type, real)), + b.create<arith::ConstantOp>(type, + DenseFPElementsAttr::get(type, imag))}; } LogicalResult matchAndRewrite( @@ -406,7 +410,7 @@ int lognPlus1 = std::log(fftLength) / std::log(2) + 1; for (auto s : llvm::seq<unsigned>(1, lognPlus1)) { SmallVector<Value> inputs; - inputs.push_back(b.create<ConstantIndexOp>(s)); + inputs.push_back(b.create<arith::ConstantIndexOp>(s)); inputs.append(getCoeffConstants(b, s)); auto fft = b.create<linalg_ext::FftOp>( TypeRange{results[0].getType(), results[1].getType()}, inputs, @@ -477,6 +481,7 @@ void getDependentDialects(DialectRegistry ®istry) const override { registry.insert<linalg_ext::LinalgExtDialect, linalg::LinalgDialect, IREE::Flow::FlowDialect, StandardOpsDialect, + mlir::math::MathDialect, mlir::arith::ArithmeticDialect, complex::ComplexDialect, tensor::TensorDialect>(); } @@ -538,6 +543,8 @@ ConversionTarget target(getContext()); target.addLegalDialect<linalg_ext::LinalgExtDialect, linalg::LinalgDialect, IREE::Flow::FlowDialect, StandardOpsDialect, + mlir::math::MathDialect, + mlir::arith::ArithmeticDialect, tensor::TensorDialect, complex::ComplexDialect>(); target.addIllegalOp<mhlo::SortOp, mhlo::ScatterOp, mhlo::FftOp, mhlo::ReverseOp>();
diff --git a/iree/compiler/InputConversion/MHLO/LegalizeInputTypes.cpp b/iree/compiler/InputConversion/MHLO/LegalizeInputTypes.cpp index f531114..d765b77 100644 --- a/iree/compiler/InputConversion/MHLO/LegalizeInputTypes.cpp +++ b/iree/compiler/InputConversion/MHLO/LegalizeInputTypes.cpp
@@ -94,7 +94,8 @@ } } - if (llvm::isa<mlir::ConstantOp>(oldOp) || llvm::isa<mhlo::ConstOp>(oldOp) || + if (llvm::isa<mlir::arith::ConstantOp>(oldOp) || + llvm::isa<mhlo::ConstOp>(oldOp) || llvm::isa<IREE::Util::GlobalOp>(oldOp)) { for (auto attr : oldOp->getAttrs()) { auto newAttr =
diff --git a/iree/compiler/InputConversion/MHLO/MHLOToLinalgOnTensors.cpp b/iree/compiler/InputConversion/MHLO/MHLOToLinalgOnTensors.cpp index 0c5e4fc..969f0a5 100644 --- a/iree/compiler/InputConversion/MHLO/MHLOToLinalgOnTensors.cpp +++ b/iree/compiler/InputConversion/MHLO/MHLOToLinalgOnTensors.cpp
@@ -72,30 +72,30 @@ int rank = resultType.getRank(); SmallVector<Value, 3> offsets, sizes, strides; for (int i = 0; i < rank; ++i) { - offsets.push_back(rewriter.create<ConstantIndexOp>(loc, 0)); + offsets.push_back(rewriter.create<arith::ConstantIndexOp>(loc, 0)); sizes.push_back(rewriter.create<tensor::DimOp>(loc, args[0], i)); - strides.push_back(rewriter.create<ConstantIndexOp>(loc, 1)); + strides.push_back(rewriter.create<arith::ConstantIndexOp>(loc, 1)); } - Value resultDimSize = rewriter.create<ConstantIndexOp>(loc, 0); + Value resultDimSize = rewriter.create<arith::ConstantIndexOp>(loc, 0); for (auto arg : args) { auto size = rewriter.create<tensor::DimOp>(loc, arg, dim); - resultDimSize = rewriter.create<AddIOp>(loc, resultDimSize, size); + resultDimSize = rewriter.create<arith::AddIOp>(loc, resultDimSize, size); } sizes[dim] = resultDimSize; auto initTensor = rewriter.create<linalg::InitTensorOp>( loc, resultType.getShape(), resultType.getElementType()); auto zeroAttr = rewriter.getZeroAttr(resultType.getElementType()); - Value zero = rewriter.create<ConstantOp>(loc, zeroAttr); + Value zero = rewriter.create<arith::ConstantOp>(loc, zeroAttr); Value result = rewriter.create<linalg::FillOp>(loc, zero, initTensor).getResult(0); - Value accBound = rewriter.create<ConstantIndexOp>(loc, 0); + Value accBound = rewriter.create<arith::ConstantIndexOp>(loc, 0); for (auto arg : args) { offsets[dim] = accBound; sizes[dim] = rewriter.create<tensor::DimOp>(loc, arg, dim); result = rewriter.create<tensor::InsertSliceOp>(loc, arg, result, offsets, sizes, strides); - accBound = rewriter.create<AddIOp>(loc, accBound, sizes[dim]); + accBound = rewriter.create<arith::AddIOp>(loc, accBound, sizes[dim]); } rewriter.replaceOp(op, result); return success(); @@ -129,15 +129,15 @@ values.push_back(b.getF32FloatAttr(v)); } } - return b.create<ConstantOp>(loc, matrixType, - DenseFPElementsAttr::get(matrixType, values)); + return b.create<arith::ConstantOp>( + loc, matrixType, DenseFPElementsAttr::get(matrixType, values)); } Value createLinalgMatmulOnTensors(OpBuilder b, Location loc, RankedTensorType resultType, Value lhs, Value rhs) { - Value zero = - b.create<ConstantOp>(loc, b.getZeroAttr(resultType.getElementType())); + Value zero = b.create<arith::ConstantOp>( + loc, b.getZeroAttr(resultType.getElementType())); auto initTensor = b.create<linalg::InitTensorOp>( loc, /*dyn_size=*/ValueRange{}, resultType.getShape(), resultType.getElementType()); @@ -214,8 +214,9 @@ ConvertMHLOToLinalgOnTensorsPass> { void getDependentDialects(DialectRegistry ®istry) const override { registry.insert<IREE::Flow::FlowDialect, linalg::LinalgDialect, - mhlo::MhloDialect, ShapeDialect, math::MathDialect, - memref::MemRefDialect, complex::ComplexDialect>(); + mhlo::MhloDialect, shape::ShapeDialect, ShapeDialect, + math::MathDialect, memref::MemRefDialect, + complex::ComplexDialect>(); } void runOnOperation() override { @@ -263,7 +264,7 @@ newValueAttr = valueAttr.cast<DenseIntOrFPElementsAttr>().mapValues( newElType, [](const APInt &oldEl) { return oldEl; }); } - rewriter.replaceOpWithNewOp<ConstantOp>(op, newValueAttr); + rewriter.replaceOpWithNewOp<arith::ConstantOp>(op, newValueAttr); return success(); } };
diff --git a/iree/compiler/InputConversion/MHLO/MHLOToMHLOPreprocessing.cpp b/iree/compiler/InputConversion/MHLO/MHLOToMHLOPreprocessing.cpp index 1f5084e..588a9de 100644 --- a/iree/compiler/InputConversion/MHLO/MHLOToMHLOPreprocessing.cpp +++ b/iree/compiler/InputConversion/MHLO/MHLOToMHLOPreprocessing.cpp
@@ -14,6 +14,7 @@ #include "mlir-hlo/Dialect/mhlo/IR/chlo_ops.h" #include "mlir-hlo/Dialect/mhlo/IR/hlo_ops.h" #include "mlir-hlo/Dialect/mhlo/transforms/rewriters.h" +#include "mlir/Dialect/Math/IR/Math.h" #include "mlir/Dialect/Shape/IR/Shape.h" #include "mlir/Dialect/StandardOps/IR/Ops.h" #include "mlir/Dialect/Tensor/IR/Tensor.h" @@ -92,7 +93,7 @@ auto type = op.operand().getType().cast<TensorType>(); DenseElementsAttr attr = DenseElementsAttr::get(type, rewriter.getF32FloatAttr(1.0)); - auto one = rewriter.create<ConstantOp>(loc, attr); + auto one = rewriter.create<arith::ConstantOp>(loc, attr); auto x = rewriter.create<mhlo::AddOp>(loc, op.operand(), one); rewriter.replaceOpWithNewOp<mhlo::LogOp>(op, x); return success(); @@ -109,7 +110,7 @@ auto type = op.operand().getType().cast<TensorType>(); DenseElementsAttr attr = DenseElementsAttr::get(type, rewriter.getF32FloatAttr(1.0)); - auto one = rewriter.create<ConstantOp>(loc, attr); + auto one = rewriter.create<arith::ConstantOp>(loc, attr); auto x = rewriter.create<mhlo::ExpOp>(loc, op.operand()); rewriter.replaceOpWithNewOp<mhlo::SubOp>(op, x, one); return success(); @@ -165,7 +166,7 @@ RankedTensorType::get(shape, inputType.getElementType()); Attribute zeroAttr = rewriter.getZeroAttr( RankedTensorType::get({}, inputType.getElementType())); - auto zero = rewriter.create<ConstantOp>(loc, zeroAttr); + auto zero = rewriter.create<arith::ConstantOp>(loc, zeroAttr); auto padOp = rewriter.create<mhlo::PadOp>( loc, padResultType, op.lhs(), zero, toDenseAttr(paddingLow), toDenseAttr(paddingHigh), toDenseAttr(interiorPadding)); @@ -894,7 +895,8 @@ // chlo::PopulateLegalizeChloToHloPatterns(context, &conversionPatterns); conversionTarget.addLegalDialect< shape::ShapeDialect, chlo::HloClientDialect, mhlo::MhloDialect, - mlir::StandardOpsDialect, mlir::tensor::TensorDialect>(); + math::MathDialect, mlir::StandardOpsDialect, + mlir::arith::ArithmeticDialect, mlir::tensor::TensorDialect>(); // conversionTarget.addIllegalDialect<chlo::HloClientDialect>(); if (failed(applyPartialConversion(getOperation(), conversionTarget, std::move(conversionPatterns)))) {
diff --git a/iree/compiler/InputConversion/MHLO/test/broadcasting.mlir b/iree/compiler/InputConversion/MHLO/test/broadcasting.mlir index ed0272e..d657341 100644 --- a/iree/compiler/InputConversion/MHLO/test/broadcasting.mlir +++ b/iree/compiler/InputConversion/MHLO/test/broadcasting.mlir
@@ -23,13 +23,13 @@ func @dynamicBroadcast(%arg0: tensor<?xf32>, %arg1: tensor<?x?xf32>) -> tensor<?x?xf32> { // Should broadcast %arg0 -> %arg1 and assert on dynamic expansion. - // CHECK: %[[C0_0:.*]] = constant 0 : index + // CHECK: %[[C0_0:.*]] = arith.constant 0 : index // CHECK: %[[ARG0_D0:.*]] = tensor.dim %arg0, %[[C0_0]] - // CHECK: %[[C0_1:.*]] = constant 0 : index + // CHECK: %[[C0_1:.*]] = arith.constant 0 : index // CHECK: %[[ARG1_D0:.*]] = tensor.dim %arg1, %[[C0_1]] : tensor<?x?xf32> - // CHECK: %[[C1_0:.*]] = constant 1 : index + // CHECK: %[[C1_0:.*]] = arith.constant 1 : index // CHECK: %[[ARG1_D1:.*]] = tensor.dim %arg1, %[[C1_0]] : tensor<?x?xf32> - // CHECK: %[[EQ:.*]] = cmpi eq, %[[ARG0_D0]], %[[ARG1_D1]] : index + // CHECK: %[[EQ:.*]] = arith.cmpi eq, %[[ARG0_D0]], %[[ARG1_D1]] : index // CHECK: assert %[[EQ]], "mismatched dynamic broadcast extents" // CHECK: %[[INIT_0:.*]] = linalg.init_tensor [%[[ARG1_D0]], %[[ARG0_D0]]] : tensor<?x?xf32> @@ -174,7 +174,7 @@ // CHECK: #map3 = affine_map<(d0, d1, d2) -> (0, 0, d2)> // CHECK-LABEL: func @selectv2_broadcast_dyn_pred func @selectv2_broadcast_dyn_pred(%arg0: tensor<?x1x1xi1>, %arg1: tensor<1x8x1xi32>, %arg2: tensor<1x1x8xi32>) -> tensor<?x8x8xi32> { - // CHECK: %[[C0_0:.*]] = constant 0 : index + // CHECK: %[[C0_0:.*]] = arith.constant 0 : index // CHECK: %[[DIM_PRED_0:.*]] = tensor.dim %arg0, %[[C0_0]] // CHECK: %[[INIT_PRED:.*]] = linalg.init_tensor [%[[DIM_PRED_0]], 8, 8] // CHECK: %[[BCAST_PRED:.*]] = linalg.generic @@ -188,9 +188,10 @@ // CHECK: %[[BCAST_ELSE:.*]] = linalg.generic // CHECK-SAME: indexing_maps = [#map3, #map1] // CHECK-SAME: ins(%arg2 : tensor<1x1x8xi32>) outs(%[[INIT_ELSE]] : tensor<?x8x8xi32>) - // CHECK: %[[C0_1:.*]] = constant 0 : index - // CHECK: %[[DIM_BCAST_PRED_0:.*]] = tensor.dim %[[BCAST_PRED]], %[[C0_1]] - // CHECK: %[[INIT_RESULT:.*]] = linalg.init_tensor [%[[DIM_BCAST_PRED_0]], 8, 8] + // CHECK: %[[SHAPE_BCAST_THEN:.*]] = shape.shape_of %[[BCAST_THEN]] + // CHECK: %[[C0_1:.*]] = arith.constant 0 : index + // CHECK: %[[DIM_BCAST_THEN_0:.*]] = tensor.extract %[[SHAPE_BCAST_THEN]][%[[C0_1]]] + // CHECK: %[[INIT_RESULT:.*]] = linalg.init_tensor [%[[DIM_BCAST_THEN_0]], 8, 8] // CHECK: linalg.generic // CHECK-SAME: ins(%[[BCAST_PRED]], %[[BCAST_THEN]], %[[BCAST_ELSE]] : tensor<?x8x8xi1>, tensor<?x8x8xi32>, tensor<?x8x8xi32>) outs(%[[INIT_RESULT]] : tensor<?x8x8xi32>) %0 = "chlo.broadcast_select"(%arg0, %arg1, %arg2) : (tensor<?x1x1xi1>, tensor<1x8x1xi32>, tensor<1x1x8xi32>) -> tensor<?x8x8xi32> @@ -200,7 +201,7 @@ // ----- // CHECK-LABEL: func @selectv2_broadcast_dyn_then func @selectv2_broadcast_dyn_then(%arg0: tensor<8x1x1xi1>, %arg1: tensor<1x?x1xi32>, %arg2: tensor<1x1x8xi32>) -> tensor<8x?x8xi32> { - // CHECK: %[[C1_0:.*]] = constant 1 : index + // CHECK: %[[C1_0:.*]] = arith.constant 1 : index // CHECK: %[[DIM_THEN_1:.*]] = tensor.dim %arg1, %[[C1_0]] // CHECK: %[[INIT_PRED:.*]] = linalg.init_tensor [8, %[[DIM_THEN_1]], 8] // CHECK: %[[BCAST_PRED:.*]] = linalg.generic @@ -214,11 +215,12 @@ // CHECK: %[[BCAST_ELSE:.*]] = linalg.generic // CHECK-SAME: indexing_maps = [#map3, #map1] // CHECK-SAME: ins(%arg2 : tensor<1x1x8xi32>) outs(%[[INIT_ELSE]] : tensor<8x?x8xi32>) - // CHECK: %[[C1_1:.*]] = constant 1 : index - // CHECK: %[[DIM_BCAST_PRED_1:.*]] = tensor.dim %[[BCAST_PRED]], %[[C1_1]] - // CHECK: %[[INIT_RESULT:.*]] = linalg.init_tensor [8, %[[DIM_BCAST_PRED_1]], 8] + // CHECK: %[[SHAPE_BCAST_THEN:.*]] = shape.shape_of %[[BCAST_THEN]] + // CHECK: %[[C1_1:.*]] = arith.constant 1 : index + // CHECK: %[[DIM_BCAST_THEN_1:.*]] = tensor.extract %[[SHAPE_BCAST_THEN]][%[[C1_1]]] + // CHECK: %[[INIT_RESULT:.*]] = linalg.init_tensor [8, %[[DIM_BCAST_THEN_1]], 8] // CHECK: linalg.generic - // CHECK-SAME: ins(%2, %4, %6 : tensor<8x?x8xi1>, tensor<8x?x8xi32>, tensor<8x?x8xi32>) outs(%8 : tensor<8x?x8xi32>) + // CHECK-SAME: ins(%[[BCAST_PRED]], %[[BCAST_THEN]], %[[BCAST_ELSE]] : tensor<8x?x8xi1>, tensor<8x?x8xi32>, tensor<8x?x8xi32>) outs(%[[INIT_RESULT]] : tensor<8x?x8xi32>) %0 = "chlo.broadcast_select"(%arg0, %arg1, %arg2) : (tensor<8x1x1xi1>, tensor<1x?x1xi32>, tensor<1x1x8xi32>) -> tensor<8x?x8xi32> return %0: tensor<8x?x8xi32> } @@ -226,7 +228,7 @@ // ----- // CHECK-LABEL: func @selectv2_broadcast_dyn_else func @selectv2_broadcast_dyn_else(%arg0: tensor<8x1x1xi1>, %arg1: tensor<1x8x1xi32>, %arg2: tensor<1x1x?xi32>) -> tensor<8x8x?xi32> { - // CHECK: %[[C2_0:.*]] = constant 2 : index + // CHECK: %[[C2_0:.*]] = arith.constant 2 : index // CHECK: %[[DIM_ELSE_2:.*]] = tensor.dim %arg2, %[[C2_0]] // CHECK: %[[INIT_PRED:.*]] = linalg.init_tensor [8, 8, %[[DIM_ELSE_2]]] // CHECK: %[[BCAST_PRED:.*]] = linalg.generic @@ -241,11 +243,12 @@ // CHECK: %[[BCAST_ELSE:.*]] = linalg.generic // CHECK-SAME: indexing_maps = [#map3, #map1] // CHECK-SAME: ins(%arg2 : tensor<1x1x?xi32>) outs(%[[INIT_ELSE]] : tensor<8x8x?xi32>) - // CHECK: %[[C2_1:.*]] = constant 2 : index - // CHECK: %[[DIM_BCAST_PRED_1:.*]] = tensor.dim %[[BCAST_PRED]], %[[C2_1]] - // CHECK: %[[INIT_RESULT:.*]] = linalg.init_tensor [8, 8, %[[DIM_BCAST_PRED_1]]] + // CHECK: %[[SHAPE_BCAST_THEN:.*]] = shape.shape_of %[[BCAST_THEN]] + // CHECK: %[[C2_1:.*]] = arith.constant 2 : index + // CHECK: %[[DIM_BCAST_THEN_1:.*]] = tensor.extract %[[SHAPE_BCAST_THEN]][%[[C2_1]]] + // CHECK: %[[INIT_RESULT:.*]] = linalg.init_tensor [8, 8, %[[DIM_BCAST_THEN_1]]] // CHECK: linalg.generic - // CHECK-SAME: ins(%2, %4, %6 : tensor<8x8x?xi1>, tensor<8x8x?xi32>, tensor<8x8x?xi32>) outs(%8 : tensor<8x8x?xi32>) + // CHECK-SAME: ins(%[[BCAST_PRED]], %[[BCAST_THEN]], %[[BCAST_ELSE]] : tensor<8x8x?xi1>, tensor<8x8x?xi32>, tensor<8x8x?xi32>) outs(%[[INIT_RESULT]] : tensor<8x8x?xi32>) %0 = "chlo.broadcast_select"(%arg0, %arg1, %arg2) : (tensor<8x1x1xi1>, tensor<1x8x1xi32>, tensor<1x1x?xi32>) -> tensor<8x8x?xi32> return %0: tensor<8x8x?xi32> } @@ -253,17 +256,17 @@ // ----- // CHECK-LABEL: func @selectv2_broadcast_dyn_all func @selectv2_broadcast_dyn_all(%arg0: tensor<?x1x1xi1>, %arg1: tensor<?x8x1xi32>, %arg2: tensor<?x1x?xi32>) -> tensor<?x8x?xi32> { - // CHECK: %[[C0:.*]] = constant 0 : index + // CHECK: %[[C0:.*]] = arith.constant 0 : index // CHECK: %[[PRED_D0:.*]] = tensor.dim %arg0, %[[C0]] : tensor<?x1x1xi1> - // CHECK: %[[C0_0:.*]] = constant 0 : index + // CHECK: %[[C0_0:.*]] = arith.constant 0 : index // CHECK: %[[THEN_D0:.*]] = tensor.dim %arg1, %[[C0_0]] : tensor<?x8x1xi32> - // CHECK: %[[C0_1:.*]] = constant 0 : index + // CHECK: %[[C0_1:.*]] = arith.constant 0 : index // CHECK: %[[ELSE_D0:.*]] = tensor.dim %arg2, %[[C0_1]] : tensor<?x1x?xi32> - // CHECK: %[[C2:.*]] = constant 2 : index + // CHECK: %[[C2:.*]] = arith.constant 2 : index // CHECK: %[[ELSE_D2:.*]] = tensor.dim %arg2, %[[C2]] : tensor<?x1x?xi32> - // CHECK: %[[CMP_0:.*]] = cmpi eq, %[[PRED_D0]], %[[THEN_D0]] : index + // CHECK: %[[CMP_0:.*]] = arith.cmpi eq, %[[PRED_D0]], %[[THEN_D0]] : index // CHECK: assert %[[CMP_0]], "mismatched dynamic broadcast extents" - // CHECK: %[[CMP_1:.*]] = cmpi eq, %[[PRED_D0]], %[[ELSE_D0]] : index + // CHECK: %[[CMP_1:.*]] = arith.cmpi eq, %[[PRED_D0]], %[[ELSE_D0]] : index // CHECK: assert %[[CMP_1]], "mismatched dynamic broadcast extents" // Only two asserts are needed. The rest are statically verified. // CHECK-NOT: assert @@ -409,15 +412,15 @@ // CHECK-LABEL: @fallbackDynamicReshape func @fallbackDynamicReshape(%arg0 : tensor<4x?x3x?xui32>, %arg1 : tensor<5xindex>) -> tensor<12x?x?x1x?xui32> { // CHECK: %[[INPUT:.*]] = builtin.unrealized_conversion_cast %arg0 : tensor<4x?x3x?xui32> to tensor<4x?x3x?xi32> - // CHECK-DAG: %[[C1:.*]] = constant 1 : index + // CHECK-DAG: %[[C1:.*]] = arith.constant 1 : index // CHECK-DAG: %[[RESULT_D1:.*]] = tensor.extract %arg1[%[[C1]]] : tensor<5xindex> - // CHECK-DAG: %[[C2:.*]] = constant 2 : index + // CHECK-DAG: %[[C2:.*]] = arith.constant 2 : index // CHECK-DAG: %[[RESULT_D2:.*]] = tensor.extract %arg1[%[[C2]]] : tensor<5xindex> - // CHECK-DAG: %[[C4:.*]] = constant 4 : index + // CHECK-DAG: %[[C4:.*]] = arith.constant 4 : index // CHECK-DAG: %[[RESULT_D4:.*]] = tensor.extract %arg1[%[[C4]]] : tensor<5xindex> - // CHECK-DAG: %[[INDEX1:.*]] = constant 1 : index + // CHECK-DAG: %[[INDEX1:.*]] = arith.constant 1 : index // CHECK-DAG: %[[ARG_D1:.*]] = tensor.dim %[[INPUT]], %[[INDEX1]] : tensor<4x?x3x?xi32> - // CHECK-DAG: %[[INDEX3:.*]] = constant 3 : index + // CHECK-DAG: %[[INDEX3:.*]] = arith.constant 3 : index // CHECK-DAG: %[[ARG_D3:.*]] = tensor.dim %[[INPUT]], %[[INDEX3]] : tensor<4x?x3x?xi32> // CHECK-DAG: %[[RESULT:.*]] = flow.tensor.reshape %[[INPUT]] : tensor<4x?x3x?xi32>{%[[ARG_D1]], %[[ARG_D3]]} -> tensor<12x?x?x1x?xi32>{%[[RESULT_D1]], %[[RESULT_D2]], %[[RESULT_D4]]} %0 = "mhlo.dynamic_reshape"(%arg0, %arg1) : (tensor<4x?x3x?xui32>, tensor<5xindex>) -> tensor<12x?x?x1x?xui32>
diff --git a/iree/compiler/InputConversion/MHLO/test/convert_mhlo_to_linalg_ext.mlir b/iree/compiler/InputConversion/MHLO/test/convert_mhlo_to_linalg_ext.mlir index 1cee567..d0e9225 100644 --- a/iree/compiler/InputConversion/MHLO/test/convert_mhlo_to_linalg_ext.mlir +++ b/iree/compiler/InputConversion/MHLO/test/convert_mhlo_to_linalg_ext.mlir
@@ -17,7 +17,7 @@ // CHECK-SAME: dimension(0) // CHECK-SAME: outs(%[[ARG0]] : tensor<128xi32>) // CHECK: ^bb0(%[[ARG1:.+]]: i32, %[[ARG2:.+]]: i32) -// CHECK: %[[CMP:.+]] = cmpi sgt, %[[ARG1]], %[[ARG2]] +// CHECK: %[[CMP:.+]] = arith.cmpi sgt, %[[ARG1]], %[[ARG2]] // CHECK: linalg_ext.yield %[[CMP]] // CHECK: return %[[SORT]] @@ -38,8 +38,8 @@ // CHECK-SAME: ) // CHECK: %[[SORT:.+]] = linalg_ext.sort dimension(1) outs(%[[ARG0]] : tensor<1x10xi32>) { // CHECK: ^bb0(%[[ARG1:.+]]: i32, %{{.*}}: i32) -// CHECK: %[[SCALAR:.+]] = constant 0 : i32 -// CHECK: %[[RES:.+]] = cmpi slt, %[[ARG1]], %[[SCALAR]] : i32 +// CHECK: %[[SCALAR:.+]] = arith.constant 0 : i32 +// CHECK: %[[RES:.+]] = arith.cmpi slt, %[[ARG1]], %[[SCALAR]] : i32 // CHECK: linalg_ext.yield %[[RES]] : i1 // CHECK: } -> tensor<1x10xi32> // CHECK: return %[[SORT]] @@ -62,7 +62,7 @@ // CHECK: %[[SORT:.+]] = linalg_ext.sort dimension(1) outs(%[[ARG0]] : tensor<1x10xi32>) { // CHECK: ^bb0(%[[ARG2:.+]]: i32, %{{.*}}: i32) // CHECK: %[[SCALAR:.+]] = tensor.extract %[[ARG1]][] : tensor<i32> -// CHECK: %[[RES:.+]] = cmpi slt, %[[ARG2]], %[[SCALAR]] : i32 +// CHECK: %[[RES:.+]] = arith.cmpi slt, %[[ARG2]], %[[SCALAR]] : i32 // CHECK: linalg_ext.yield %[[RES]] : i1 // CHECK: } -> tensor<1x10xi32> // CHECK: return %[[SORT]] @@ -84,7 +84,7 @@ // CHECK-SAME: dimension(0) // CHECK-SAME: outs(%[[ARG0]] : tensor<16x32xi32>) // CHECK: ^bb0(%[[ARG1:.+]]: i32, %[[ARG2:.+]]: i32) -// CHECK: %[[CMP:.+]] = cmpi sgt, %[[ARG1]], %[[ARG2]] +// CHECK: %[[CMP:.+]] = arith.cmpi sgt, %[[ARG1]], %[[ARG2]] // CHECK: linalg_ext.yield %[[CMP]] // CHECK: return %[[SORT]] @@ -108,9 +108,9 @@ // CHECK-SAME: dimension(1) // CHECK-SAME: outs(%[[ARG0]] : tensor<1x5xf32>) // CHECK: ^bb0(%[[ARG1:.+]]: f32, %[[ARG2:.+]]: f32) -// CHECK: %[[CAST1:.+]] = bitcast %[[ARG1]] : f32 to i32 -// CHECK: %[[CAST2:.+]] = bitcast %[[ARG2]] : f32 to i32 -// CHECK: %[[CMP:.+]] = cmpi ult, %[[CAST1]], %[[CAST2]] : i32 +// CHECK: %[[CAST1:.+]] = arith.bitcast %[[ARG1]] : f32 to i32 +// CHECK: %[[CAST2:.+]] = arith.bitcast %[[ARG2]] : f32 to i32 +// CHECK: %[[CMP:.+]] = arith.cmpi ult, %[[CAST1]], %[[CAST2]] : i32 // CHECK: linalg_ext.yield %[[CMP]] // CHECK: return %[[SORT]] @@ -135,10 +135,10 @@ // CHECK-SAME: dimension(1) // CHECK-SAME: outs(%[[ARG0]] : tensor<1x5xf32>) // CHECK: ^bb0(%[[ARG1:.+]]: f32, %[[ARG2:.+]]: f32) -// CHECK: %[[CAST1:.+]] = bitcast %[[ARG1]] : f32 to i32 +// CHECK: %[[CAST1:.+]] = arith.bitcast %[[ARG1]] : f32 to i32 // CHECK: %[[CONVERSION_CAST_CST:.+]] = builtin.unrealized_conversion_cast %[[UI32]] : tensor<ui32> to tensor<i32> // CHECK: %[[EXTRACT_CST:.+]] = tensor.extract %[[CONVERSION_CAST_CST]][] : tensor<i32> -// CHECK: %[[CMP:.+]] = cmpi ult, %[[CAST1]], %[[EXTRACT_CST]] : i32 +// CHECK: %[[CMP:.+]] = arith.cmpi ult, %[[CAST1]], %[[EXTRACT_CST]] : i32 // CHECK: linalg_ext.yield %[[CMP]] // CHECK: return %[[SORT]] @@ -168,7 +168,7 @@ // CHECK-SAME: outs(%[[ARG0]] : tensor<1x5xf32>) // CHECK: ^bb0(%[[ARG1:.+]]: f32, %[[ARG2:.+]]: f32) // CHECK-NOT: mhlo.complex -// CHECK: %[[CMP:.+]] = cmpf olt, %{{.+}}, %{{.+}} : f32 +// CHECK: %[[CMP:.+]] = arith.cmpf olt, %{{.+}}, %{{.+}} : f32 // CHECK: linalg_ext.yield %[[CMP]] // CHECK: return %[[SORT]] @@ -189,7 +189,7 @@ // CHECK-SAME: dimension(0) // CHECK-SAME: outs(%[[ARG0]], %[[ARG1]] : tensor<128xi32>, tensor<128xi32>) // CHECK: ^bb0(%[[ARG2:.+]]: i32, %[[ARG3:.+]]: i32, %{{.*}}: i32, %{{.*}}: i32) -// CHECK: %[[CMP:.+]] = cmpi sgt, %[[ARG2]], %[[ARG3]] +// CHECK: %[[CMP:.+]] = arith.cmpi sgt, %[[ARG2]], %[[ARG3]] // CHECK: linalg_ext.yield %[[CMP]] // CHECK: return %[[SORT]]#0 @@ -310,7 +310,7 @@ // CHECK: ^bb0(%[[V1:.+]]: i32, %[[V2:.+]]: i32): // no predecessors // // The order is reverse. -// CHECK: %[[V3:.+]] = addi %[[V2]], %[[V1]] +// CHECK: %[[V3:.+]] = arith.addi %[[V2]], %[[V1]] // CEECK: linalg.yield %[[V3]] // CHECK: return %[[SCATTER]] @@ -393,7 +393,7 @@ // CHECK-DAG: #[[MAP:.+]] = affine_map<(d0) -> (d0)> // CHECK: func @rfft_1d // CHECK-SAME: %[[REAL:[a-zA-Z0-9]+]] -// CHECK-DAG: %[[INDICES:.+]] = constant dense<[0, 4, 2, 6, 1, 5, 3, 7]> : tensor<8xi32> +// CHECK-DAG: %[[INDICES:.+]] = arith.constant dense<[0, 4, 2, 6, 1, 5, 3, 7]> : tensor<8xi32> // CHECK-DAG: %[[INIT_TENSOR:.+]] = linalg.init_tensor [8] : tensor<8xf32> // CHECK: %[[REORDERED:.+]] = linalg.generic // CHECK-SAME: {indexing_maps = [#[[MAP]], #[[MAP]]] @@ -401,25 +401,25 @@ // CHECK-SAME: ins(%[[INDICES]] // CHECK-SAME: outs(%[[INIT_TENSOR]] // CHECK: ^bb0(%[[IDX:.+]]: i32, %{{.+}}: f32): -// CHECK: %[[IDXVAL:.+]] = index_cast %[[IDX]] : i32 to index +// CHECK: %[[IDXVAL:.+]] = arith.index_cast %[[IDX]] : i32 to index // CHECK: %[[LOAD:.+]] = tensor.extract %[[REAL]][%[[IDXVAL]]] : tensor<8xf32> // CHECK: linalg.yield %[[LOAD]] : f32 -// CHECK-DAG: %[[IMAG:.+]] = constant dense<0.000000e+00> : tensor<8xf32> -// CHECK-DAG: %[[C1:.+]] = constant 1 : index -// CHECK-DAG: %[[COEF_REAL:.+]] = constant dense<{{.+}}> : tensor<1xf32> -// CHECK-DAG: %[[COEF_IMAG:.+]] = constant dense<{{.+}}> : tensor<1xf32> +// CHECK-DAG: %[[IMAG:.+]] = arith.constant dense<0.000000e+00> : tensor<8xf32> +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index +// CHECK-DAG: %[[COEF_REAL:.+]] = arith.constant dense<{{.+}}> : tensor<1xf32> +// CHECK-DAG: %[[COEF_IMAG:.+]] = arith.constant dense<{{.+}}> : tensor<1xf32> // CHECK: %[[R1:.+]]:2 = linalg_ext.fft // CHECK-SAME: ins(%[[C1]], %[[COEF_REAL]], %[[COEF_IMAG]] // CHECK-SAME: outs(%[[REORDERED]], %[[IMAG]] -// CHECK-DAG: %[[C2:.+]] = constant 2 : index -// CHECK-DAG: %[[COEF_REAL:.+]] = constant dense<{{.+}}> : tensor<2xf32> -// CHECK-DAG: %[[COEF_IMAG:.+]] = constant dense<{{.+}}> : tensor<2xf32> +// CHECK-DAG: %[[C2:.+]] = arith.constant 2 : index +// CHECK-DAG: %[[COEF_REAL:.+]] = arith.constant dense<{{.+}}> : tensor<2xf32> +// CHECK-DAG: %[[COEF_IMAG:.+]] = arith.constant dense<{{.+}}> : tensor<2xf32> // CHECK: %[[R2:.+]]:2 = linalg_ext.fft // CHECK-SAME: ins(%[[C2]], %[[COEF_REAL]], %[[COEF_IMAG]] // CHECK-SAME: outs(%[[R1]]#0, %[[R1]]#1 -// CHECK-DAG: %[[C3:.+]] = constant 3 : index -// CHECK-DAG: %[[COEF_REAL:.+]] = constant dense<{{.+}}> : tensor<4xf32> -// CHECK-DAG: %[[COEF_IMAG:.+]] = constant dense<{{.+}}> : tensor<4xf32> +// CHECK-DAG: %[[C3:.+]] = arith.constant 3 : index +// CHECK-DAG: %[[COEF_REAL:.+]] = arith.constant dense<{{.+}}> : tensor<4xf32> +// CHECK-DAG: %[[COEF_IMAG:.+]] = arith.constant dense<{{.+}}> : tensor<4xf32> // CHECK: %[[R3:.+]]:2 = linalg_ext.fft // CHECK-SAME: ins(%[[C3]], %[[COEF_REAL]], %[[COEF_IMAG]] // CHECK-SAME: outs(%[[R2]]#0, %[[R2]]#1 @@ -441,7 +441,7 @@ // CHECK-DAG: #[[MAP1:.+]] = affine_map<(d0, d1) -> (d0, d1)> // CHECK: func @rfft_2d // CHECK-SAME: %[[REAL:[a-zA-Z0-9]+]] -// CHECK-DAG: %[[INDICES:.+]] = constant dense<[0, 4, 2, 6, 1, 5, 3, 7]> : tensor<8xi32> +// CHECK-DAG: %[[INDICES:.+]] = arith.constant dense<[0, 4, 2, 6, 1, 5, 3, 7]> : tensor<8xi32> // CHECK-DAG: %[[INIT_TENSOR:.+]] = linalg.init_tensor [4, 8] : tensor<4x8xf32> // CHECK: %[[REORDERED:.+]] = linalg.generic // CHECK-SAME: {indexing_maps = [#[[MAP0]], #[[MAP1]]] @@ -450,25 +450,25 @@ // CHECK-SAME: outs(%[[INIT_TENSOR]] // CHECK: ^bb0(%[[IDX:.+]]: i32, %{{.+}}: f32): // CHECK: %[[I:.+]] = linalg.index 0 -// CHECK: %[[IDXVAL:.+]] = index_cast %[[IDX]] : i32 to index +// CHECK: %[[IDXVAL:.+]] = arith.index_cast %[[IDX]] : i32 to index // CHECK: %[[LOAD:.+]] = tensor.extract %[[REAL]][%[[I]], %[[IDXVAL]]] : tensor<4x8xf32> // CHECK: linalg.yield %[[LOAD]] : f32 -// CHECK-DAG: %[[IMAG:.+]] = constant dense<0.000000e+00> : tensor<4x8xf32> -// CHECK-DAG: %[[C1:.+]] = constant 1 : index -// CHECK-DAG: %[[COEF_REAL:.+]] = constant dense<{{.+}}> : tensor<1xf32> -// CHECK-DAG: %[[COEF_IMAG:.+]] = constant dense<{{.+}}> : tensor<1xf32> +// CHECK-DAG: %[[IMAG:.+]] = arith.constant dense<0.000000e+00> : tensor<4x8xf32> +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index +// CHECK-DAG: %[[COEF_REAL:.+]] = arith.constant dense<{{.+}}> : tensor<1xf32> +// CHECK-DAG: %[[COEF_IMAG:.+]] = arith.constant dense<{{.+}}> : tensor<1xf32> // CHECK: %[[R1:.+]]:2 = linalg_ext.fft // CHECK-SAME: ins(%[[C1]], %[[COEF_REAL]], %[[COEF_IMAG]] // CHECK-SAME: outs(%[[REORDERED]], %[[IMAG]] -// CHECK-DAG: %[[C2:.+]] = constant 2 : index -// CHECK-DAG: %[[COEF_REAL:.+]] = constant dense<{{.+}}> : tensor<2xf32> -// CHECK-DAG: %[[COEF_IMAG:.+]] = constant dense<{{.+}}> : tensor<2xf32> +// CHECK-DAG: %[[C2:.+]] = arith.constant 2 : index +// CHECK-DAG: %[[COEF_REAL:.+]] = arith.constant dense<{{.+}}> : tensor<2xf32> +// CHECK-DAG: %[[COEF_IMAG:.+]] = arith.constant dense<{{.+}}> : tensor<2xf32> // CHECK: %[[R2:.+]]:2 = linalg_ext.fft // CHECK-SAME: ins(%[[C2]], %[[COEF_REAL]], %[[COEF_IMAG]] // CHECK-SAME: outs(%[[R1]]#0, %[[R1]]#1 -// CHECK-DAG: %[[C3:.+]] = constant 3 : index -// CHECK-DAG: %[[COEF_REAL:.+]] = constant dense<{{.+}}> : tensor<4xf32> -// CHECK-DAG: %[[COEF_IMAG:.+]] = constant dense<{{.+}}> : tensor<4xf32> +// CHECK-DAG: %[[C3:.+]] = arith.constant 3 : index +// CHECK-DAG: %[[COEF_REAL:.+]] = arith.constant dense<{{.+}}> : tensor<4xf32> +// CHECK-DAG: %[[COEF_IMAG:.+]] = arith.constant dense<{{.+}}> : tensor<4xf32> // CHECK: %[[R3:.+]]:2 = linalg_ext.fft // CHECK-SAME: ins(%[[C3]], %[[COEF_REAL]], %[[COEF_IMAG]] // CHECK-SAME: outs(%[[R2]]#0, %[[R2]]#1 @@ -503,8 +503,8 @@ } // CHECK-LABEL: func @reverse_multi_dim // CHECK-SAME: %[[IN:[a-zA-Z0-9]+]] -// CHECK-DAG: %[[C0:.+]] = constant 0 : index -// CHECK-DAG: %[[C1:.+]] = constant 1 : index +// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index +// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index // CHECK-DAG: %[[D0:.+]] = tensor.dim %[[IN]], %[[C0]] // CHECK-DAG: %[[D1:.+]] = tensor.dim %[[IN]], %[[C1]] // CHECK: %[[INIT:.+]] = linalg.init_tensor [%[[D0]], %[[D1]]] : tensor<?x?xi32>
diff --git a/iree/compiler/InputConversion/MHLO/test/dynamic_shape.mlir b/iree/compiler/InputConversion/MHLO/test/dynamic_shape.mlir index ec32c98..2d22d3f 100644 --- a/iree/compiler/InputConversion/MHLO/test/dynamic_shape.mlir +++ b/iree/compiler/InputConversion/MHLO/test/dynamic_shape.mlir
@@ -9,10 +9,11 @@ // CHECK: #[[MAP0:.+]] = affine_map<(d0, d1) -> (d0, d1)> // CHECK: func @dynamic_shape // CHECK-SAME: %[[ARG0:.+]]: tensor<?x?xf32> -// CHECK: %[[C0:.+]] = constant 0 : index -// CHECK: %[[T0:.+]] = tensor.dim %[[ARG0]], %[[C0]] -// CHECK: %[[C1:.+]] = constant 1 : index -// CHECK: %[[T1:.+]] = tensor.dim %[[ARG0]], %[[C1]] +// CHECK: %[[SHAPE:.+]] = shape.shape_of %[[ARG0]] +// CHECK: %[[C0:.+]] = arith.constant 0 : index +// CHECK: %[[T0:.+]] = tensor.extract %[[SHAPE]][%[[C0]]] +// CHECK: %[[C1:.+]] = arith.constant 1 : index +// CHECK: %[[T1:.+]] = tensor.extract %[[SHAPE]][%[[C1]]] // CHECK: %[[T2:.+]] = linalg.init_tensor [%[[T0]], %[[T1]]] // CHECK: %[[T3:.+]] = linalg.generic // CHECK-SAME: indexing_maps = [#[[MAP0]], #[[MAP0]]]
diff --git a/iree/compiler/InputConversion/MHLO/test/fft.mlir b/iree/compiler/InputConversion/MHLO/test/fft.mlir index 8d07c2a..0567433 100644 --- a/iree/compiler/InputConversion/MHLO/test/fft.mlir +++ b/iree/compiler/InputConversion/MHLO/test/fft.mlir
@@ -10,9 +10,9 @@ } // CHECK: func @rfft_1d // CHECK-SAME: %[[Arg0:[a-zA-Z0-9_]*]] -// CHECK-DAG: %[[RealMatrix:.+]] = constant dense<"0x0000803F{{.*}}"> : tensor<32x17xf32> -// CHECK-DAG: %[[ImagMatrix:.+]] = constant dense<"0x00000080{{.*}}"> : tensor<32x17xf32> -// CHECK-DAG: %[[Zero:.+]] = constant 0.000000e+00 : f32 +// CHECK-DAG: %[[RealMatrix:.+]] = arith.constant dense<"0x0000803F{{.*}}"> : tensor<32x17xf32> +// CHECK-DAG: %[[ImagMatrix:.+]] = arith.constant dense<"0x00000080{{.*}}"> : tensor<32x17xf32> +// CHECK-DAG: %[[Zero:.+]] = arith.constant 0.000000e+00 : f32 // CHECK: %[[RealInit:.+]] = linalg.init_tensor [17] : tensor<17xf32> // CHECK: %[[RealFill:.+]] = linalg.fill(%[[Zero]], %[[RealInit]]) // CHECK: %[[RealRes:.+]] = linalg.vecmat @@ -37,9 +37,9 @@ } // CHECK: func @rfft_2d // CHECK-SAME: %[[Arg0:[a-zA-Z0-9_]*]] -// CHECK-DAG: %[[RealMatrix:.+]] = constant dense<"0x0000803F{{.*}}"> : tensor<32x17xf32> -// CHECK-DAG: %[[ImagMatrix:.+]] = constant dense<"0x00000080{{.*}}"> : tensor<32x17xf32> -// CHECK-DAG: %[[Zero:.+]] = constant 0.000000e+00 : f32 +// CHECK-DAG: %[[RealMatrix:.+]] = arith.constant dense<"0x0000803F{{.*}}"> : tensor<32x17xf32> +// CHECK-DAG: %[[ImagMatrix:.+]] = arith.constant dense<"0x00000080{{.*}}"> : tensor<32x17xf32> +// CHECK-DAG: %[[Zero:.+]] = arith.constant 0.000000e+00 : f32 // CHECK: %[[RealInit:.+]] = linalg.init_tensor [1, 17] : tensor<1x17xf32> // CHECK: %[[RealFill:.+]] = linalg.fill(%[[Zero]], %[[RealInit]]) // CHECK: %[[RealRes:.+]] = linalg.matmul
diff --git a/iree/compiler/InputConversion/MHLO/test/legalize_input_types.mlir b/iree/compiler/InputConversion/MHLO/test/legalize_input_types.mlir index 29a2a69..313df38 100644 --- a/iree/compiler/InputConversion/MHLO/test/legalize_input_types.mlir +++ b/iree/compiler/InputConversion/MHLO/test/legalize_input_types.mlir
@@ -4,7 +4,7 @@ // CHECK-SAME: () -> i32 func @constantI64() -> i64 { // CHECK-NEXT: constant 123 : i32 - %c123 = constant 123 : i64 + %c123 = arith.constant 123 : i64 return %c123 : i64 } @@ -36,7 +36,7 @@ // CHECK-SAME: () -> f32 func @constantF64() -> f64 { // CHECK-NEXT: constant 1.234000e+02 : f32 - %c1234 = constant 123.4 : f64 + %c1234 = arith.constant 123.4 : f64 return %c1234 : f64 } @@ -46,7 +46,7 @@ // CHECK-SAME: () -> tensor<4xi32> func @constantSplatTensorI64() -> tensor<4xi64> { // CHECK-NEXT: constant dense<123> : tensor<4xi32> - %c123 = constant dense<123> : tensor<4xi64> + %c123 = arith.constant dense<123> : tensor<4xi64> return %c123 : tensor<4xi64> } @@ -56,7 +56,7 @@ // CHECK-SAME: () -> tensor<4xi32> func @constantDenseTensorI64() -> tensor<4xi64> { // CHECK-NEXT: constant dense<[0, 1, 2, 3]> : tensor<4xi32> - %c123 = constant dense<[0, 1, 2, 3]> : tensor<4xi64> + %c123 = arith.constant dense<[0, 1, 2, 3]> : tensor<4xi64> return %c123 : tensor<4xi64> }
diff --git a/iree/compiler/Translation/test/do_not_optimize.mlir b/iree/compiler/Translation/test/do_not_optimize.mlir index b7037fe..57aa2ea 100644 --- a/iree/compiler/Translation/test/do_not_optimize.mlir +++ b/iree/compiler/Translation/test/do_not_optimize.mlir
@@ -2,10 +2,10 @@ // CHECK-LABEL: @add func @add() -> i32 { - %input = constant 1 : i32 + %input = arith.constant 1 : i32 %unf = util.do_not_optimize(%input) : i32 // CHECK: vm.add.i32 - %result = addi %unf, %unf : i32 + %result = arith.addi %unf, %unf : i32 return %result : i32 } @@ -14,9 +14,9 @@ // Ensure that add would normally be folded. // CHECK-LABEL: @add_folded func @add_folded() -> i32 { - %input = constant 1 : i32 + %input = arith.constant 1 : i32 // CHECK-NOT: vm.add.i32 - %result = addi %input, %input : i32 + %result = arith.addi %input, %input : i32 return %result : i32 } @@ -24,12 +24,12 @@ // CHECK-LABEL: @chained_add func @chained_add() -> i32 { - %input = constant 1 : i32 + %input = arith.constant 1 : i32 %unf = util.do_not_optimize(%input) : i32 // CHECK: vm.add.i32 - %int = addi %unf, %unf : i32 + %int = arith.addi %unf, %unf : i32 // CHECK: vm.add.i32 - %result = addi %int, %int : i32 + %result = arith.addi %int, %int : i32 return %result : i32 } @@ -39,7 +39,7 @@ func @unfoldable_constant() -> i32 { %input = util.unfoldable_constant 1 : i32 // CHECK: vm.add.i32 - %result = addi %input, %input : i32 + %result = arith.addi %input, %input : i32 return %result : i32 }
diff --git a/iree/compiler/Translation/test/smoketest.mlir b/iree/compiler/Translation/test/smoketest.mlir index ce3298f..e72c931 100644 --- a/iree/compiler/Translation/test/smoketest.mlir +++ b/iree/compiler/Translation/test/smoketest.mlir
@@ -31,10 +31,10 @@ // CHECK: "exported_functions": // CHECK: "local_name": "add" func @add() -> i32 { - %c1 = constant 1 : i32 + %c1 = arith.constant 1 : i32 %unf_c1 = util.do_not_optimize(%c1) : i32 %unf_c2 = util.unfoldable_constant 2 : i32 - %result = addi %unf_c1, %unf_c2 : i32 + %result = arith.addi %unf_c1, %unf_c2 : i32 return %result : i32 } }
diff --git a/iree/samples/custom_modules/custom_modules_test.mlir b/iree/samples/custom_modules/custom_modules_test.mlir index 72191be..fe16474 100644 --- a/iree/samples/custom_modules/custom_modules_test.mlir +++ b/iree/samples/custom_modules/custom_modules_test.mlir
@@ -11,7 +11,7 @@ // custom_modules/dialect/custom.imports.mlir for the import definitions. func @reverseAndPrint(%message : !custom.message, %count : i32) -> !custom.message attributes { iree.module.export, iree.abi.none } { - %c1 = constant 1 : i32 + %c1 = arith.constant 1 : i32 %0 = "custom.get_unique_message"() : () -> !custom.message "custom.print"(%0, %c1) : (!custom.message, i32) -> () %1 = call @reverse(%message) : (!custom.message) -> !custom.message @@ -29,7 +29,7 @@ func @printTensor(%tensor : tensor<2x4xf32>) -> !custom.message attributes { iree.module.export, iree.abi.none } { %0 = "custom.tensor_to_message"(%tensor) : (tensor<2x4xf32>) -> !custom.message - %c1 = constant 1 : i32 + %c1 = arith.constant 1 : i32 "custom.print"(%0, %c1) : (!custom.message, i32) -> () return %0 : !custom.message } @@ -40,7 +40,7 @@ %0 = "custom.tensor_to_message"(%tensor) : (tensor<2x4xf32>) -> !custom.message %1 = "custom.message_to_tensor"(%0) : (!custom.message) -> tensor<2x4xf32> %2 = "custom.tensor_to_message"(%1) : (tensor<2x4xf32>) -> !custom.message - %c1 = constant 1 : i32 + %c1 = arith.constant 1 : i32 "custom.print"(%2, %c1) : (!custom.message, i32) -> () return %0 : !custom.message }
diff --git a/iree/samples/custom_modules/dialect/test/conversion.mlir b/iree/samples/custom_modules/dialect/test/conversion.mlir index 50b5a20..a6c65ca 100644 --- a/iree/samples/custom_modules/dialect/test/conversion.mlir +++ b/iree/samples/custom_modules/dialect/test/conversion.mlir
@@ -22,7 +22,7 @@ // CHECK-SAME: ]) // CHECK-NEXT: [[MSG:%.+]] = vm.call @custom.buffer_to_message([[VIEW]]) {nosideeffects} : (!vm.ref<!hal.buffer_view>) -> !vm.ref<!custom.message> %0 = "custom.tensor_to_message"(%tensor) : (tensor<2x4xf32>) -> !custom.message - %c1 = constant 1 : i32 + %c1 = arith.constant 1 : i32 // CHECK: vm.call @custom.print([[MSG]] "custom.print"(%0, %c1) : (!custom.message, i32) -> () return @@ -40,7 +40,7 @@ %shape = shapex.make_ranked_shape %arg1, %arg2 : (index, index) -> !shapex.ranked_shape<[?, ?]> %shaped_tensor = shapex.tie_shape %arg0, %shape : tensor<?x?xf32>, !shapex.ranked_shape<[?, ?]> %0 = "custom.tensor_to_message"(%shaped_tensor) : (tensor<?x?xf32>) -> !custom.message - %c1 = constant 1 : i32 + %c1 = arith.constant 1 : i32 // CHECK: vm.call @custom.print([[MSG]] "custom.print"(%0, %c1) : (!custom.message, i32) -> () return @@ -57,7 +57,7 @@ // CHECK-NEXT: [[MSG:%.+]] = vm.call @custom.buffer_to_message([[VIEW]]) {nosideeffects} : (!vm.ref<!hal.buffer_view>) -> !vm.ref<!custom.message> %shaped_tensor = shapex.tie_shape %arg0, %arg1 : tensor<?x?xf32>, !shapex.ranked_shape<[?, ?]> %0 = "custom.tensor_to_message"(%shaped_tensor) : (tensor<?x?xf32>) -> !custom.message - %c1 = constant 1 : i32 + %c1 = arith.constant 1 : i32 // CHECK: vm.call @custom.print([[MSG]] "custom.print"(%0, %c1) : (!custom.message, i32) -> () return @@ -79,7 +79,7 @@ // CHECK-LABEL: @messageToTensorReturnDim func @messageToTensorReturnDim(%arg0 : !custom.message) -> index { %0 = "custom.message_to_tensor"(%arg0) : (!custom.message) -> tensor<?x4xf32> - %c0 = constant 0 : index + %c0 = arith.constant 0 : index %1 = tensor.dim %0, %c0 : tensor<?x4xf32> // CHECK: [[VIEW:%.+]] = vm.call @custom.message_to_buffer(%arg0) {nosideeffects} : (!vm.ref<!custom.message>) -> !vm.ref<!hal.buffer_view> // CHECK: [[BUFFER:%.+]] = vm.call @hal.buffer_view.buffer([[VIEW]]) @@ -107,7 +107,7 @@ // CHECK-LABEL: @printOp func @printOp(%arg0 : !custom.message) { - %c1_i32 = constant 1 : i32 + %c1_i32 = arith.constant 1 : i32 // CHECK: vm.call @custom.print(%arg0, %c1) : (!vm.ref<!custom.message>, i32) -> () "custom.print"(%arg0, %c1_i32) : (!custom.message, i32) -> () return
diff --git a/iree/samples/custom_modules/dialect/test/custom_ops.mlir b/iree/samples/custom_modules/dialect/test/custom_ops.mlir index 4589481..5a7c281 100644 --- a/iree/samples/custom_modules/dialect/test/custom_ops.mlir +++ b/iree/samples/custom_modules/dialect/test/custom_ops.mlir
@@ -12,7 +12,7 @@ // CHECK-LABEL: @printOp func @printOp(%arg0 : !custom.message) { - %c1_i32 = constant 1 : i32 + %c1_i32 = arith.constant 1 : i32 // CHECK: "custom.print"(%arg0, %c1_i32) : (!custom.message, i32) -> () "custom.print"(%arg0, %c1_i32) : (!custom.message, i32) -> () return
diff --git a/iree/samples/dynamic_shapes/dynamic_shapes.ipynb b/iree/samples/dynamic_shapes/dynamic_shapes.ipynb index 556641f..29a85af 100644 --- a/iree/samples/dynamic_shapes/dynamic_shapes.ipynb +++ b/iree/samples/dynamic_shapes/dynamic_shapes.ipynb
@@ -189,7 +189,7 @@ "#map3 = affine_map<(d0, d1) -> (d0)>\n", "module {\n", " func @add_one(%arg0: !hal.buffer_view) -> !hal.buffer_view attributes {iree.abi = \"{\\22a\\22:[[\\22ndarray\\22,\\22i32\\22,1,null]],\\22r\\22:[[\\22ndarray\\22,\\22i32\\22,1,null]],\\22v\\22:1}\"} {\n", - " %c0 = constant 0 : index\n", + " %c0 = arith.constant 0 : index\n", " %0 = hal.buffer_view.dim %arg0, 0 : index\n", " %1 = hal.tensor.cast %arg0 : !hal.buffer_view -> tensor<?xi32>{%0}\n", " %2 = call @__inference_add_one_70(%1) : (tensor<?xi32>) -> tensor<?xi32>\n", @@ -198,8 +198,8 @@ " return %4 : !hal.buffer_view\n", " }\n", " func private @__inference_add_one_70(%arg0: tensor<?xi32> {tf._user_specified_name = \"values\"}) -> tensor<?xi32> attributes {tf._construction_context = \"kEagerRuntime\", tf._input_shapes = [#tf.shape<?>]} {\n", - " %cst = constant dense<1> : tensor<i32>\n", - " %c0 = constant 0 : index\n", + " %cst = arith.constant dense<1> : tensor<i32>\n", + " %c0 = arith.constant 0 : index\n", " %0 = memref.dim %arg0, %c0 : tensor<?xi32>\n", " %1 = linalg.init_tensor [%0] : tensor<?xi32>\n", " %2 = linalg.generic {indexing_maps = [#map0, #map1], iterator_types = [\"parallel\"]} ins(%cst : tensor<i32>) outs(%1 : tensor<?xi32>) {\n", @@ -210,34 +210,34 @@ " %4 = linalg.init_tensor [%3] : tensor<?xi32>\n", " %5 = linalg.generic {indexing_maps = [#map1, #map1, #map1], iterator_types = [\"parallel\"]} ins(%arg0, %2 : tensor<?xi32>, tensor<?xi32>) outs(%4 : tensor<?xi32>) {\n", " ^bb0(%arg1: i32, %arg2: i32, %arg3: i32): // no predecessors\n", - " %6 = addi %arg1, %arg2 : i32\n", + " %6 = arith.addi %arg1, %arg2 : i32\n", " linalg.yield %6 : i32\n", " } -> tensor<?xi32>\n", " return %5 : tensor<?xi32>\n", " }\n", " func @reduce_sum_1d(%arg0: !hal.buffer_view) -> !hal.buffer_view attributes {iree.abi = \"{\\22a\\22:[[\\22ndarray\\22,\\22i32\\22,1,null]],\\22r\\22:[[\\22ndarray\\22,\\22i32\\22,0]],\\22v\\22:1}\"} {\n", - " %c0_i32 = constant 0 : i32\n", + " %c0_i32 = arith.constant 0 : i32\n", " %0 = hal.buffer_view.dim %arg0, 0 : index\n", " %1 = hal.tensor.cast %arg0 : !hal.buffer_view -> tensor<?xi32>{%0}\n", " %2 = linalg.init_tensor [] : tensor<i32>\n", " %3 = linalg.fill(%2, %c0_i32) : tensor<i32>, i32 -> tensor<i32> \n", " %4 = linalg.generic {indexing_maps = [#map1, #map0], iterator_types = [\"reduction\"]} ins(%1 : tensor<?xi32>) outs(%3 : tensor<i32>) {\n", " ^bb0(%arg1: i32, %arg2: i32): // no predecessors\n", - " %6 = addi %arg1, %arg2 : i32\n", + " %6 = arith.addi %arg1, %arg2 : i32\n", " linalg.yield %6 : i32\n", " } -> tensor<i32>\n", " %5 = hal.tensor.cast %4 : tensor<i32> -> !hal.buffer_view\n", " return %5 : !hal.buffer_view\n", " }\n", " func @reduce_sum_2d(%arg0: !hal.buffer_view) -> !hal.buffer_view attributes {iree.abi = \"{\\22a\\22:[[\\22ndarray\\22,\\22i32\\22,2,null,3]],\\22r\\22:[[\\22ndarray\\22,\\22i32\\22,1,3]],\\22v\\22:1}\"} {\n", - " %c0_i32 = constant 0 : i32\n", + " %c0_i32 = arith.constant 0 : i32\n", " %0 = hal.buffer_view.dim %arg0, 0 : index\n", " %1 = hal.tensor.cast %arg0 : !hal.buffer_view -> tensor<?x3xi32>{%0}\n", " %2 = linalg.init_tensor [3] : tensor<3xi32>\n", " %3 = linalg.fill(%2, %c0_i32) : tensor<3xi32>, i32 -> tensor<3xi32> \n", " %4 = linalg.generic {indexing_maps = [#map2, #map3], iterator_types = [\"parallel\", \"reduction\"]} ins(%1 : tensor<?x3xi32>) outs(%3 : tensor<3xi32>) {\n", " ^bb0(%arg1: i32, %arg2: i32): // no predecessors\n", - " %6 = addi %arg1, %arg2 : i32\n", + " %6 = arith.addi %arg1, %arg2 : i32\n", " linalg.yield %6 : i32\n", " } -> tensor<3xi32>\n", " %5 = hal.tensor.cast %4 : tensor<3xi32> -> !hal.buffer_view\n",
diff --git a/iree/samples/models/simple_abs.mlir b/iree/samples/models/simple_abs.mlir index d376406..5011a99 100644 --- a/iree/samples/models/simple_abs.mlir +++ b/iree/samples/models/simple_abs.mlir
@@ -1,4 +1,4 @@ func @abs(%input : tensor<f32>) -> (tensor<f32>) { - %result = absf %input : tensor<f32> + %result = math.abs %input : tensor<f32> return %result : tensor<f32> }
diff --git a/iree/samples/static_library/simple_mul.mlir b/iree/samples/static_library/simple_mul.mlir index 2b480bd..aefbb25 100644 --- a/iree/samples/static_library/simple_mul.mlir +++ b/iree/samples/static_library/simple_mul.mlir
@@ -1,5 +1,5 @@ func @simple_mul(%arg0: tensor<4xf32>, %arg1: tensor<4xf32>) -> tensor<4xf32> { - %0 = "std.mulf"(%arg0, %arg1) {name = "mul.1"} : (tensor<4xf32>, tensor<4xf32>) -> tensor<4xf32> + %0 = "arith.mulf"(%arg0, %arg1) {name = "mul.1"} : (tensor<4xf32>, tensor<4xf32>) -> tensor<4xf32> return %0 : tensor<4xf32> }
diff --git a/iree/samples/variables_and_state/variables_and_state.ipynb b/iree/samples/variables_and_state/variables_and_state.ipynb index 3c16e93..bacfc57 100644 --- a/iree/samples/variables_and_state/variables_and_state.ipynb +++ b/iree/samples/variables_and_state/variables_and_state.ipynb
@@ -201,7 +201,7 @@ " %1 = linalg.init_tensor [] : tensor<i32>\n", " %2 = linalg.generic {indexing_maps = [#map, #map, #map], iterator_types = []} ins(%0, %arg0 : tensor<i32>, tensor<i32>) outs(%1 : tensor<i32>) {\n", " ^bb0(%arg1: i32, %arg2: i32, %arg3: i32): // no predecessors\n", - " %3 = addi %arg1, %arg2 : i32\n", + " %3 = arith.addi %arg1, %arg2 : i32\n", " linalg.yield %3 : i32\n", " } -> tensor<i32>\n", " util.global.store %2, @counter : tensor<i32>\n", @@ -213,7 +213,7 @@ " return %1 : !hal.buffer_view\n", " }\n", " func @reset_value() attributes {iree.abi = \"{\\22a\\22:[],\\22r\\22:[],\\22v\\22:1}\"} {\n", - " %cst = constant dense<0> : tensor<i32>\n", + " %cst = arith.constant dense<0> : tensor<i32>\n", " util.global.store %cst, @counter : tensor<i32>\n", " return\n", " }\n",
diff --git a/iree/test/e2e/models/fullyconnected.mlir b/iree/test/e2e/models/fullyconnected.mlir index aa5c7da..0c64174 100644 --- a/iree/test/e2e/models/fullyconnected.mlir +++ b/iree/test/e2e/models/fullyconnected.mlir
@@ -8,48 +8,48 @@ %2 = "mhlo.reshape"(%1) {name = "reshape.42"} : (tensor<5x1xf32>) -> tensor<5x1x1xf32> %3 = "mhlo.reshape"(%2) {name = "reshape.55"} : (tensor<5x1x1xf32>) -> tensor<5x1xf32> %4 = "mhlo.broadcast_in_dim"(%3) {broadcast_dimensions = dense<[0, 1]> : tensor<2xi64>, name = "broadcast.56"} : (tensor<5x1xf32>) -> tensor<5x1x5xf32> - %cst = constant {name = "constant.22"} dense<1.000000e+00> : tensor<f32> + %cst = arith.constant {name = "constant.22"} dense<1.000000e+00> : tensor<f32> %5 = "mhlo.broadcast_in_dim"(%cst) {broadcast_dimensions = dense<[]> : tensor<0xi64>, name = "broadcast.23"} : (tensor<f32>) -> tensor<5x1x5xf32> %6 = mhlo.multiply %4, %5 {name = "multiply.57"} : tensor<5x1x5xf32> - %cst_0 = constant {name = "constant.58"} dense<0.000000e+00> : tensor<f32> + %cst_0 = arith.constant {name = "constant.58"} dense<0.000000e+00> : tensor<f32> %7 = "mhlo.broadcast_in_dim"(%cst_0) {broadcast_dimensions = dense<[]> : tensor<0xi64>, name = "broadcast.59"} : (tensor<f32>) -> tensor<5x1x5xf32> %8 = "mhlo.compare"(%6, %7) {comparison_direction = "GT", name = "compare.60"} : (tensor<5x1x5xf32>, tensor<5x1x5xf32>) -> tensor<5x1x5xi1> - %cst_1 = constant {name = "constant.24"} dense<0.000000e+00> : tensor<f32> + %cst_1 = arith.constant {name = "constant.24"} dense<0.000000e+00> : tensor<f32> %9 = "mhlo.broadcast_in_dim"(%cst_1) {broadcast_dimensions = dense<[]> : tensor<0xi64>, name = "broadcast.25"} : (tensor<f32>) -> tensor<5x1x5xf32> - %cst_2 = constant {name = "constant.90"} dense<0.000000e+00> : tensor<f32> + %cst_2 = arith.constant {name = "constant.90"} dense<0.000000e+00> : tensor<f32> %10 = "mhlo.broadcast_in_dim"(%cst_2) {broadcast_dimensions = dense<[]> : tensor<0xi64>, name = "broadcast.91"} : (tensor<f32>) -> tensor<5x5xf32> %11 = "mhlo.reshape"(%2) {name = "reshape.49"} : (tensor<5x1x1xf32>) -> tensor<5x1xf32> %12 = "mhlo.broadcast_in_dim"(%11) {broadcast_dimensions = dense<[0, 1]> : tensor<2xi64>, name = "broadcast.50"} : (tensor<5x1xf32>) -> tensor<5x1x5xf32> - %cst_3 = constant {name = "constant.15"} dense<1.000000e+00> : tensor<f32> + %cst_3 = arith.constant {name = "constant.15"} dense<1.000000e+00> : tensor<f32> %13 = "mhlo.broadcast_in_dim"(%cst_3) {broadcast_dimensions = dense<[]> : tensor<0xi64>, name = "broadcast.16"} : (tensor<f32>) -> tensor<5x1x5xf32> %14 = mhlo.multiply %12, %13 {name = "multiply.51"} : tensor<5x1x5xf32> - %cst_4 = constant {name = "constant.52"} dense<0.000000e+00> : tensor<f32> + %cst_4 = arith.constant {name = "constant.52"} dense<0.000000e+00> : tensor<f32> %15 = "mhlo.broadcast_in_dim"(%cst_4) {broadcast_dimensions = dense<[]> : tensor<0xi64>, name = "broadcast.53"} : (tensor<f32>) -> tensor<5x1x5xf32> %16 = "mhlo.compare"(%14, %15) {comparison_direction = "GT", name = "compare.54"} : (tensor<5x1x5xf32>, tensor<5x1x5xf32>) -> tensor<5x1x5xi1> - %cst_5 = constant {name = "constant.17"} dense<0.000000e+00> : tensor<f32> + %cst_5 = arith.constant {name = "constant.17"} dense<0.000000e+00> : tensor<f32> %17 = "mhlo.broadcast_in_dim"(%cst_5) {broadcast_dimensions = dense<[]> : tensor<0xi64>, name = "broadcast.18"} : (tensor<f32>) -> tensor<5x1x5xf32> - %cst_6 = constant {name = "constant.78"} dense<0.000000e+00> : tensor<f32> + %cst_6 = arith.constant {name = "constant.78"} dense<0.000000e+00> : tensor<f32> %18 = "mhlo.broadcast_in_dim"(%cst_6) {broadcast_dimensions = dense<[]> : tensor<0xi64>, name = "broadcast.79"} : (tensor<f32>) -> tensor<5x5xf32> %19 = "mhlo.reshape"(%2) {name = "reshape.43"} : (tensor<5x1x1xf32>) -> tensor<5x1xf32> %20 = "mhlo.broadcast_in_dim"(%19) {broadcast_dimensions = dense<[0, 1]> : tensor<2xi64>, name = "broadcast.44"} : (tensor<5x1xf32>) -> tensor<5x1x5xf32> - %cst_7 = constant {name = "constant.8"} dense<1.000000e+00> : tensor<f32> + %cst_7 = arith.constant {name = "constant.8"} dense<1.000000e+00> : tensor<f32> %21 = "mhlo.broadcast_in_dim"(%cst_7) {broadcast_dimensions = dense<[]> : tensor<0xi64>, name = "broadcast.9"} : (tensor<f32>) -> tensor<5x1x5xf32> %22 = mhlo.multiply %20, %21 {name = "multiply.45"} : tensor<5x1x5xf32> - %cst_8 = constant {name = "constant.46"} dense<0.000000e+00> : tensor<f32> + %cst_8 = arith.constant {name = "constant.46"} dense<0.000000e+00> : tensor<f32> %23 = "mhlo.broadcast_in_dim"(%cst_8) {broadcast_dimensions = dense<[]> : tensor<0xi64>, name = "broadcast.47"} : (tensor<f32>) -> tensor<5x1x5xf32> %24 = "mhlo.compare"(%22, %23) {comparison_direction = "GT", name = "compare.48"} : (tensor<5x1x5xf32>, tensor<5x1x5xf32>) -> tensor<5x1x5xi1> - %cst_9 = constant {name = "constant.10"} dense<0.000000e+00> : tensor<f32> + %cst_9 = arith.constant {name = "constant.10"} dense<0.000000e+00> : tensor<f32> %25 = "mhlo.broadcast_in_dim"(%cst_9) {broadcast_dimensions = dense<[]> : tensor<0xi64>, name = "broadcast.11"} : (tensor<f32>) -> tensor<5x1x5xf32> - %cst_10 = constant {name = "constant.66"} dense<0.000000e+00> : tensor<f32> + %cst_10 = arith.constant {name = "constant.66"} dense<0.000000e+00> : tensor<f32> %26 = "mhlo.broadcast_in_dim"(%cst_10) {broadcast_dimensions = dense<[]> : tensor<0xi64>, name = "broadcast.67"} : (tensor<f32>) -> tensor<5x5xf32> %27 = "mhlo.copy"(%arg1) {name = "copy.3"} : (tensor<1x5x3x1xf32>) -> tensor<1x5x3x1xf32> %28 = "mhlo.reshape"(%27) {name = "reshape.4"} : (tensor<1x5x3x1xf32>) -> tensor<1x5x3x1xf32> %29 = "mhlo.reshape"(%28) {name = "reshape.38"} : (tensor<1x5x3x1xf32>) -> tensor<1x5x3xf32> %30 = "mhlo.transpose"(%29) {name = "transpose.39", permutation = dense<[1, 0, 2]> : tensor<3xi64>} : (tensor<1x5x3xf32>) -> tensor<5x1x3xf32> %31 = "mhlo.reshape"(%30) {name = "reshape.40"} : (tensor<5x1x3xf32>) -> tensor<5x3xf32> - %cst_11 = constant {name = "constant.61"} dense<[[0.706495285, -0.567672312, 0.483717591, 0.522725761, 0.7563259], [-0.0899272263, -0.283501834, -0.350822538, -0.351515919, -0.337136656], [-0.451804549, 0.372324884, -0.620518147, 0.235451385, 0.851095855]]> : tensor<3x5xf32> + %cst_11 = arith.constant {name = "constant.61"} dense<[[0.706495285, -0.567672312, 0.483717591, 0.522725761, 0.7563259], [-0.0899272263, -0.283501834, -0.350822538, -0.351515919, -0.337136656], [-0.451804549, 0.372324884, -0.620518147, 0.235451385, 0.851095855]]> : tensor<3x5xf32> %32 = "mhlo.dot"(%31, %cst_11) {name = "dot.62", precision_config = ["DEFAULT", "DEFAULT"]} : (tensor<5x3xf32>, tensor<3x5xf32>) -> tensor<5x5xf32> - %cst_12 = constant {name = "constant.63"} dense<[0.000000e+00, 0.000000e+00, 0.000000e+00, 0.000000e+00, 0.000000e+00]> : tensor<5xf32> + %cst_12 = arith.constant {name = "constant.63"} dense<[0.000000e+00, 0.000000e+00, 0.000000e+00, 0.000000e+00, 0.000000e+00]> : tensor<5xf32> %33 = "mhlo.broadcast_in_dim"(%cst_12) {broadcast_dimensions = dense<[1]> : tensor<1xi64>, name = "broadcast.64"} : (tensor<5xf32>) -> tensor<5x5xf32> %34 = mhlo.add %32, %33 {name = "add.65"} : tensor<5x5xf32> %35 = mhlo.maximum %26, %34 {name = "maximum.68"} : tensor<5x5xf32> @@ -57,9 +57,9 @@ %37 = "mhlo.select"(%24, %25, %36) {name = "select.71"} : (tensor<5x1x5xi1>, tensor<5x1x5xf32>, tensor<5x1x5xf32>) -> tensor<5x1x5xf32> %38 = "mhlo.copy"(%37) {name = "copy.4"} : (tensor<5x1x5xf32>) -> tensor<5x1x5xf32> %39 = "mhlo.reshape"(%38) {name = "reshape.72"} : (tensor<5x1x5xf32>) -> tensor<5x5xf32> - %cst_13 = constant {name = "constant.73"} dense<[[-0.0118641369, -3.785000e-02, 0.489048243, 0.321015775, -0.702280283], [-0.280262798, -0.724645615, -0.00332254497, 0.392334729, 0.619746447], [-0.113318317, -0.180415511, -0.146743968, 0.250408649, -0.442881733], [0.115600757, 0.703136146, -0.00812680274, -0.225454301, -0.0835619792], [-0.136745885, -6.298570e-01, 0.43629986, -0.689790308, 0.230725273]]> : tensor<5x5xf32> + %cst_13 = arith.constant {name = "constant.73"} dense<[[-0.0118641369, -3.785000e-02, 0.489048243, 0.321015775, -0.702280283], [-0.280262798, -0.724645615, -0.00332254497, 0.392334729, 0.619746447], [-0.113318317, -0.180415511, -0.146743968, 0.250408649, -0.442881733], [0.115600757, 0.703136146, -0.00812680274, -0.225454301, -0.0835619792], [-0.136745885, -6.298570e-01, 0.43629986, -0.689790308, 0.230725273]]> : tensor<5x5xf32> %40 = "mhlo.dot"(%39, %cst_13) {name = "dot.74", precision_config = ["DEFAULT", "DEFAULT"]} : (tensor<5x5xf32>, tensor<5x5xf32>) -> tensor<5x5xf32> - %cst_14 = constant {name = "constant.75"} dense<[0.000000e+00, 0.000000e+00, 0.000000e+00, 0.000000e+00, 0.000000e+00]> : tensor<5xf32> + %cst_14 = arith.constant {name = "constant.75"} dense<[0.000000e+00, 0.000000e+00, 0.000000e+00, 0.000000e+00, 0.000000e+00]> : tensor<5xf32> %41 = "mhlo.broadcast_in_dim"(%cst_14) {broadcast_dimensions = dense<[1]> : tensor<1xi64>, name = "broadcast.76"} : (tensor<5xf32>) -> tensor<5x5xf32> %42 = mhlo.add %40, %41 {name = "add.77"} : tensor<5x5xf32> %43 = mhlo.maximum %18, %42 {name = "maximum.80"} : tensor<5x5xf32> @@ -67,9 +67,9 @@ %45 = "mhlo.select"(%16, %17, %44) {name = "select.83"} : (tensor<5x1x5xi1>, tensor<5x1x5xf32>, tensor<5x1x5xf32>) -> tensor<5x1x5xf32> %46 = "mhlo.copy"(%45) {name = "copy.5"} : (tensor<5x1x5xf32>) -> tensor<5x1x5xf32> %47 = "mhlo.reshape"(%46) {name = "reshape.84"} : (tensor<5x1x5xf32>) -> tensor<5x5xf32> - %cst_15 = constant {name = "constant.85"} dense<[[-0.136191264, -0.0401721969, 0.38497138, -5.850760e-01, 0.370910525], [-0.391011149, 0.0266356133, 0.309115469, -0.205079094, -0.559861302], [0.497760415, 0.689488232, 0.0759292394, -0.33134672, -0.237128958], [-0.53243047, 0.476418108, -0.371978909, 0.283265263, 0.63842845], [0.101761498, -0.218626946, 0.475128263, 0.042601984, 0.0988005772]]> : tensor<5x5xf32> + %cst_15 = arith.constant {name = "constant.85"} dense<[[-0.136191264, -0.0401721969, 0.38497138, -5.850760e-01, 0.370910525], [-0.391011149, 0.0266356133, 0.309115469, -0.205079094, -0.559861302], [0.497760415, 0.689488232, 0.0759292394, -0.33134672, -0.237128958], [-0.53243047, 0.476418108, -0.371978909, 0.283265263, 0.63842845], [0.101761498, -0.218626946, 0.475128263, 0.042601984, 0.0988005772]]> : tensor<5x5xf32> %48 = "mhlo.dot"(%47, %cst_15) {name = "dot.86", precision_config = ["DEFAULT", "DEFAULT"]} : (tensor<5x5xf32>, tensor<5x5xf32>) -> tensor<5x5xf32> - %cst_16 = constant {name = "constant.87"} dense<[0.000000e+00, 0.000000e+00, 0.000000e+00, 0.000000e+00, 0.000000e+00]> : tensor<5xf32> + %cst_16 = arith.constant {name = "constant.87"} dense<[0.000000e+00, 0.000000e+00, 0.000000e+00, 0.000000e+00, 0.000000e+00]> : tensor<5xf32> %49 = "mhlo.broadcast_in_dim"(%cst_16) {broadcast_dimensions = dense<[1]> : tensor<1xi64>, name = "broadcast.88"} : (tensor<5xf32>) -> tensor<5x5xf32> %50 = mhlo.add %48, %49 {name = "add.89"} : tensor<5x5xf32> %51 = mhlo.maximum %10, %50 {name = "maximum.92"} : tensor<5x5xf32>
diff --git a/iree/test/e2e/models/unidirectional_lstm.mlir b/iree/test/e2e/models/unidirectional_lstm.mlir index 6452cd7..b80d9a4 100644 --- a/iree/test/e2e/models/unidirectional_lstm.mlir +++ b/iree/test/e2e/models/unidirectional_lstm.mlir
@@ -17,23 +17,23 @@ return %0 : tensor<i1> } func private @Forward_o16DF3vQKaI__disable_call_shape_inference_true_.189(%arg0: tensor<1x10xf32>, %arg1: tensor<1x10xf32>, %arg2: tensor<5x1x64xf32>, %arg3: tensor<5x1x1xf32>, %arg4: tensor<5x1x1xf32>) -> (tensor<i64>, tensor<5xi64>, tensor<5x1x10xf32>, tensor<5x1x10xf32>, tensor<i64>, tensor<1x10xf32>, tensor<1x10xf32>) { - %cst = constant dense<0x7F800000> : tensor<f32> + %cst = arith.constant dense<0x7F800000> : tensor<f32> %0 = mhlo.constant dense<0.000000e+00> : tensor<5xf32> - %cst_0 = constant dense<[1, 2, 3, 4, 5]> : tensor<5xi32> - %cst_1 = constant dense<-2147483648> : tensor<i32> - %cst_2 = constant dense<5> : tensor<i32> + %cst_0 = arith.constant dense<[1, 2, 3, 4, 5]> : tensor<5xi32> + %cst_1 = arith.constant dense<-2147483648> : tensor<i32> + %cst_2 = arith.constant dense<5> : tensor<i32> %1 = mhlo.constant dense<0.000000e+00> : tensor<40xf32> - %cst_3 = constant dense<4.200000e-01> : tensor<74x40xf32> - %cst_4 = constant dense<0> : tensor<i64> + %cst_3 = arith.constant dense<4.200000e-01> : tensor<74x40xf32> + %cst_4 = arith.constant dense<0> : tensor<i64> %2 = mhlo.constant dense<0> : tensor<5xi64> %3 = mhlo.constant dense<0.000000e+00> : tensor<5x1x10xf32> - %cst_5 = constant dense<1> : tensor<i64> + %cst_5 = arith.constant dense<1> : tensor<i64> %4 = mhlo.constant dense<1.000000e+01> : tensor<1x10xf32> %5 = mhlo.constant dense<-1.000000e+01> : tensor<1x10xf32> %6 = mhlo.constant dense<1.000000e+00> : tensor<1x10xf32> %7 = mhlo.constant dense<0.000000e+00> : tensor<1x10xf32> %8 = mhlo.constant dense<5.000000e-01> : tensor<1x10xf32> - %cst_6 = constant dense<0> : tensor<i32> + %cst_6 = arith.constant dense<0> : tensor<i32> %9 = "mhlo.reduce"(%arg3, %cst) ( { ^bb0(%arg5: tensor<f32>, %arg6: tensor<f32>): // no predecessors %115 = mhlo.minimum %arg5, %arg6 : tensor<f32> @@ -143,7 +143,7 @@ } func @main(%arg0: tensor<1x5xf32>, %arg1: tensor<1x5x2x2xf32>) -> tensor<5x1x10xf32> { %0 = mhlo.constant dense<0.000000e+00> : tensor<1x10xf32> - %cst = constant dense<0.000000e+00> : tensor<f32> + %cst = arith.constant dense<0.000000e+00> : tensor<f32> %1 = mhlo.constant dense<0.000000e+00> : tensor<5x1x1xf32> %2 = "mhlo.reshape"(%arg1) : (tensor<1x5x2x2xf32>) -> tensor<1x5x4xf32> %3 = "mhlo.pad"(%2, %cst) {edge_padding_high = dense<[0, 0, 60]> : tensor<3xi64>, edge_padding_low = dense<0> : tensor<3xi64>, interior_padding = dense<0> : tensor<3xi64>} : (tensor<1x5x4xf32>, tensor<f32>) -> tensor<1x5x64xf32>
diff --git a/iree/test/e2e/regression/dynamic_linalg_matmul_on_tensors_fuse_0.mlir b/iree/test/e2e/regression/dynamic_linalg_matmul_on_tensors_fuse_0.mlir index a3ca848..8d533a3 100644 --- a/iree/test/e2e/regression/dynamic_linalg_matmul_on_tensors_fuse_0.mlir +++ b/iree/test/e2e/regression/dynamic_linalg_matmul_on_tensors_fuse_0.mlir
@@ -15,8 +15,8 @@ // TODO(nicolasvasilache): The constant ends up being hoisted and turns // into a pushconstant. But pushconstants must be integers, so we use // sitofp to temporarily circumvent the problem. - %im1 = constant -1 : i32 - %fm1 = sitofp %im1: i32 to f32 + %im1 = arith.constant -1 : i32 + %fm1 = arith.sitofp %im1: i32 to f32 linalg.yield %fm1 : f32 } -> tensor<?x?xf32> %D = linalg.matmul ins(%AA, %B: tensor<?x?xf32>, tensor<?x?xf32>)
diff --git a/iree/test/e2e/regression/dynamic_linalg_matmul_on_tensors_fuse_1.mlir b/iree/test/e2e/regression/dynamic_linalg_matmul_on_tensors_fuse_1.mlir index 6ee006e..9e6725c 100644 --- a/iree/test/e2e/regression/dynamic_linalg_matmul_on_tensors_fuse_1.mlir +++ b/iree/test/e2e/regression/dynamic_linalg_matmul_on_tensors_fuse_1.mlir
@@ -15,8 +15,8 @@ // TODO(nicolasvasilache): The constant ends up being hoisted and turns // into a pushconstant. But pushconstants must be integers, so we use // sitofp to temporarily circumvent the problem. - %im1 = constant -1 : i32 - %fm1 = sitofp %im1: i32 to f32 + %im1 = arith.constant -1 : i32 + %fm1 = arith.sitofp %im1: i32 to f32 linalg.yield %fm1 : f32 } -> tensor<?x?xf32> %D = linalg.matmul ins(%A, %BB: tensor<?x?xf32>, tensor<?x?xf32>)
diff --git a/iree/test/e2e/regression/dynamic_linalg_matmul_on_tensors_fuse_2.mlir b/iree/test/e2e/regression/dynamic_linalg_matmul_on_tensors_fuse_2.mlir index 13e2e79..139cb08 100644 --- a/iree/test/e2e/regression/dynamic_linalg_matmul_on_tensors_fuse_2.mlir +++ b/iree/test/e2e/regression/dynamic_linalg_matmul_on_tensors_fuse_2.mlir
@@ -15,8 +15,8 @@ // TODO(nicolasvasilache): The constant ends up being hoisted and turns // into a pushconstant. But pushconstants must be integers, so we use // sitofp to temporarily circumvent the problem. - %im1 = constant -1 : i32 - %fm1 = sitofp %im1: i32 to f32 + %im1 = arith.constant -1 : i32 + %fm1 = arith.sitofp %im1: i32 to f32 linalg.yield %fm1 : f32 } -> tensor<?x?xf32> %D = linalg.matmul ins(%A, %B: tensor<?x?xf32>, tensor<?x?xf32>)
diff --git a/iree/test/e2e/regression/globals.mlir b/iree/test/e2e/regression/globals.mlir index 1a20207..ceb5a8c 100644 --- a/iree/test/e2e/regression/globals.mlir +++ b/iree/test/e2e/regression/globals.mlir
@@ -14,7 +14,7 @@ // CHECK: EXEC @inc func @inc() -> tensor<f32> { %0 = util.global.load @counter : tensor<f32> - %c1 = constant dense<1.0> : tensor<f32> + %c1 = arith.constant dense<1.0> : tensor<f32> %1 = mhlo.add %0, %c1 : tensor<f32> util.global.store %1, @counter : tensor<f32> %2 = util.global.load @counter : tensor<f32>
diff --git a/iree/test/e2e/regression/linalg_ext_ops.mlir b/iree/test/e2e/regression/linalg_ext_ops.mlir index 39fb1aa..8dc426c 100644 --- a/iree/test/e2e/regression/linalg_ext_ops.mlir +++ b/iree/test/e2e/regression/linalg_ext_ops.mlir
@@ -2,7 +2,7 @@ %input = util.unfoldable_constant dense<[[5, 6], [3, 7]]> : tensor<2x2xi32> %0 = linalg_ext.sort dimension(0) outs(%input : tensor<2x2xi32>) { ^bb0(%arg2: i32, %arg3: i32): // no predecessors - %1 = cmpi slt, %arg2, %arg3 : i32 + %1 = arith.cmpi slt, %arg2, %arg3 : i32 linalg_ext.yield %1 : i1 } -> tensor<2x2xi32> check.expect_eq_const(%0, dense<[[3, 6], [5, 7]]> : tensor<2x2xi32>) : tensor<2x2xi32>
diff --git a/iree/test/e2e/regression/linalg_ops.mlir b/iree/test/e2e/regression/linalg_ops.mlir index 6b58310..98173af 100644 --- a/iree/test/e2e/regression/linalg_ops.mlir +++ b/iree/test/e2e/regression/linalg_ops.mlir
@@ -17,8 +17,8 @@ ins(%input1, %input2 : tensor<3x4xi32>, tensor<3x4xi32>) outs(%init, %init : tensor<3x4xi32>, tensor<3x4xi32>) { ^bb0(%arg0: i32, %arg1: i32, %arg2: i32, %arg3: i32) : - %1 = addi %arg0, %arg1 : i32 - %2 = muli %arg0, %arg1 : i32 + %1 = arith.addi %arg0, %arg1 : i32 + %2 = arith.muli %arg0, %arg1 : i32 linalg.yield %1, %2 : i32, i32 } -> (tensor<3x4xi32>, tensor<3x4xi32>) check.expect_eq_const(%0#0, dense<[ @@ -37,7 +37,7 @@ %filter = util.unfoldable_constant dense<1.0> : tensor<3x3x3x16xf32> %bias = util.unfoldable_constant dense<1.0> : tensor<16xf32> %init = linalg.init_tensor [1, 112, 112, 16] : tensor<1x112x112x16xf32> - %cst = constant 0.0 : f32 + %cst = arith.constant 0.0 : f32 %fill = linalg.fill(%cst, %init) : f32, tensor<1x112x112x16xf32> -> tensor<1x112x112x16xf32> %conv = linalg.conv_2d_nhwc_hwcf {dilations = dense<1> : tensor<2xi64>, strides = dense<2> : tensor<2xi64>} @@ -51,7 +51,7 @@ ins(%conv, %bias : tensor<1x112x112x16xf32>, tensor<16xf32>) outs(%init : tensor<1x112x112x16xf32>) { ^bb0(%arg0 : f32, %arg1 : f32, %arg2 : f32): - %0 = addf %arg0, %arg1 : f32 + %0 = arith.addf %arg0, %arg1 : f32 linalg.yield %0 : f32 } -> tensor<1x112x112x16xf32> check.expect_eq_const(%result, dense<28.0> : tensor<1x112x112x16xf32>) : tensor<1x112x112x16xf32>
diff --git a/iree/test/e2e/tosa_ops/gather.mlir b/iree/test/e2e/tosa_ops/gather.mlir index b64140c..9581332 100644 --- a/iree/test/e2e/tosa_ops/gather.mlir +++ b/iree/test/e2e/tosa_ops/gather.mlir
@@ -1,5 +1,5 @@ func @gather_float() { - %0 = constant dense<[[[1.0, 2.0], [3.0, 4.0]]]> : tensor<1x2x2xf32> + %0 = arith.constant dense<[[[1.0, 2.0], [3.0, 4.0]]]> : tensor<1x2x2xf32> %1 = "tosa.const"() { value = dense<[[1, 0]]> : tensor<1x2xi32> } : () -> (tensor<1x2xi32>) %2 = "tosa.gather"(%0, %1) : (tensor<1x2x2xf32>, tensor<1x2xi32>) -> (tensor<1x2x2xf32>) check.expect_eq_const(%2, dense<[[[3.0, 4.0], [1.0, 2.0]]]> : tensor<1x2x2xf32>) : tensor<1x2x2xf32> @@ -7,7 +7,7 @@ } func @gather_int() { - %0 = constant dense<[[[1, 2], [3, 4]]]> : tensor<1x2x2xi32> + %0 = arith.constant dense<[[[1, 2], [3, 4]]]> : tensor<1x2x2xi32> %1 = "tosa.const"() { value = dense<[[1, 0]]> : tensor<1x2xi32> } : () -> (tensor<1x2xi32>) %2 = "tosa.gather"(%0, %1) : (tensor<1x2x2xi32>, tensor<1x2xi32>) -> (tensor<1x2x2xi32>) check.expect_eq_const(%2, dense<[[[3, 4], [1, 2]]]> : tensor<1x2x2xi32>) : tensor<1x2x2xi32>
diff --git a/iree/test/e2e/tosa_ops/table.mlir b/iree/test/e2e/tosa_ops/table.mlir index 7411c47..a25d929 100644 --- a/iree/test/e2e/tosa_ops/table.mlir +++ b/iree/test/e2e/tosa_ops/table.mlir
@@ -1,5 +1,5 @@ func @table() { - %input = constant dense<[-5405, 15214, -14896, 22008, 12529, -13501]> : tensor<6xi16> + %input = arith.constant dense<[-5405, 15214, -14896, 22008, 12529, -13501]> : tensor<6xi16> // This generates [0, ... 512] for a constant value to avoid an excessively large constant. %init = linalg.init_tensor [513] : tensor<513xi16> @@ -9,7 +9,7 @@ outs(%init: tensor<513xi16>) { ^bb0(%arg1: i16): %i = linalg.index 0 : index - %0 = index_cast %i : index to i16 + %0 = arith.index_cast %i : index to i16 linalg.yield %0 : i16 } -> tensor<513xi16>
diff --git a/iree/test/e2e/tosa_ops/transpose.mlir b/iree/test/e2e/tosa_ops/transpose.mlir index 77019d3..544789e 100644 --- a/iree/test/e2e/tosa_ops/transpose.mlir +++ b/iree/test/e2e/tosa_ops/transpose.mlir
@@ -1,6 +1,6 @@ func @test_transpose() { - %0 = constant dense<[[[0, 1, 2], [3, 4, 5]]]> : tensor<1x2x3xi32> - %1 = constant dense<[1, 2, 0]> : tensor<3xi32> + %0 = arith.constant dense<[[[0, 1, 2], [3, 4, 5]]]> : tensor<1x2x3xi32> + %1 = arith.constant dense<[1, 2, 0]> : tensor<3xi32> %2 = "tosa.transpose"(%0, %1) : (tensor<1x2x3xi32>, tensor<3xi32>) -> (tensor<2x3x1xi32>) check.expect_eq_const(%2, dense<[[[0], [1], [2]], [[3], [4], [5]]]> : tensor<2x3x1xi32>) : tensor<2x3x1xi32> return
diff --git a/iree/test/e2e/vulkan_specific/reduce.mlir b/iree/test/e2e/vulkan_specific/reduce.mlir index 5b151d2..7d6c0c9 100644 --- a/iree/test/e2e/vulkan_specific/reduce.mlir +++ b/iree/test/e2e/vulkan_specific/reduce.mlir
@@ -14,7 +14,7 @@ // make sure this works as expected. func @reduce_dim_1_const() { %0 = util.unfoldable_constant dense<[[1, 2, 3, 4, 5], [6, 7, 8, 9, 10]]> : tensor<2x5xi32> - %1 = constant dense<10> : tensor<i32> + %1 = arith.constant dense<10> : tensor<i32> %2 = "mhlo.reduce"(%0, %1) ( { ^bb0(%arg0 : tensor<i32>, %arg1 : tensor<i32>): %3 = "mhlo.add"(%arg0, %arg1) : (tensor<i32>, tensor<i32>) -> tensor<i32>
diff --git a/iree/test/e2e/xla_ops/pad.mlir b/iree/test/e2e/xla_ops/pad.mlir index a5210a5..5a2ad13 100644 --- a/iree/test/e2e/xla_ops/pad.mlir +++ b/iree/test/e2e/xla_ops/pad.mlir
@@ -1,6 +1,6 @@ func @pad_test() { %input = util.unfoldable_constant dense<[[1, 2, 3], [4, 5, 6]]> : tensor<2x3xi32> - %c0 = constant dense<0> : tensor<i32> + %c0 = arith.constant dense<0> : tensor<i32> %res = "mhlo.pad"(%input, %c0) { edge_padding_low = dense<[0, 1]> : tensor<2xi64>, edge_padding_high = dense<[1, 5]> : tensor<2xi64>, @@ -15,7 +15,7 @@ func @pad_no_op() { %input = util.unfoldable_constant dense<[[1, 2, 3], [4, 5, 6]]> : tensor<2x3xi32> - %c0 = constant dense<0> : tensor<i32> + %c0 = arith.constant dense<0> : tensor<i32> %res = "mhlo.pad"(%input, %c0) {edge_padding_high = dense<[0, 0]> : tensor<2xi64>, edge_padding_low = dense<[0, 0]> : tensor<2xi64>, interior_padding = dense<0> : tensor<2xi64>} : (tensor<2x3xi32>, tensor<i32>) -> tensor<2x3xi32> check.expect_eq(%res, %input) : tensor<2x3xi32> return
diff --git a/iree/test/e2e/xla_ops/scatter.mlir b/iree/test/e2e/xla_ops/scatter.mlir index fb4901f..259a78e 100644 --- a/iree/test/e2e/xla_ops/scatter.mlir +++ b/iree/test/e2e/xla_ops/scatter.mlir
@@ -105,7 +105,7 @@ outs(%init : tensor<1400xi32>) { ^bb0(%arg0: i32): %0 = linalg.index 0 : index - %1 = index_cast %0 : index to i32 + %1 = arith.index_cast %0 : index to i32 linalg.yield %1 : i32 } -> tensor<1400xi32> %indices_reshaped = linalg.tensor_expand_shape %indices [[0, 1]] : @@ -136,7 +136,7 @@ outs(%init : tensor<200xi32>) { ^bb0(%arg0: i32): %0 = linalg.index 0 : index - %1 = index_cast %0 : index to i32 + %1 = arith.index_cast %0 : index to i32 linalg.yield %1 : i32 } -> tensor<200xi32> %indices_reshaped = linalg.tensor_expand_shape %indices [[0, 1]] :
diff --git a/iree/test/e2e/xla_ops/while.mlir b/iree/test/e2e/xla_ops/while.mlir index d4884c2..f3fbdb2 100644 --- a/iree/test/e2e/xla_ops/while.mlir +++ b/iree/test/e2e/xla_ops/while.mlir
@@ -2,7 +2,7 @@ func @while() { %start = util.unfoldable_constant dense<1> : tensor<i32> %bound = util.unfoldable_constant dense<3> : tensor<i32> - %cst_1 = constant dense<4> : tensor<i32> + %cst_1 = arith.constant dense<4> : tensor<i32> br ^bb1(%start : tensor<i32>) ^bb1(%2: tensor<i32>): %3 = "mhlo.compare"(%2, %bound) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1>
diff --git a/iree/tools/init_mlir_dialects.h b/iree/tools/init_mlir_dialects.h index 67853c3..2e8148a 100644 --- a/iree/tools/init_mlir_dialects.h +++ b/iree/tools/init_mlir_dialects.h
@@ -47,6 +47,7 @@ quant::QuantizationDialect, spirv::SPIRVDialect, StandardOpsDialect, + mlir::arith::ArithmeticDialect, vector::VectorDialect, tensor::TensorDialect, tosa::TosaDialect,
diff --git a/iree/tools/iree-run-mlir-main.cc b/iree/tools/iree-run-mlir-main.cc index e3a62a6..11e4841 100644 --- a/iree/tools/iree-run-mlir-main.cc +++ b/iree/tools/iree-run-mlir-main.cc
@@ -20,7 +20,7 @@ // // CHECK-LABEL: @foo // // CHECK: 1xf32: 2 // func @foo() -> tensor<f32> { -// %0 = constant dense<2.0> : tensor<f32> +// %0 = arith.constant dense<2.0> : tensor<f32> // return %0 : tensor<f32> // } //
diff --git a/iree/tools/test/iree-benchmark-module.mlir b/iree/tools/test/iree-benchmark-module.mlir index e1d8613..ded2caf 100644 --- a/iree/tools/test/iree-benchmark-module.mlir +++ b/iree/tools/test/iree-benchmark-module.mlir
@@ -4,6 +4,6 @@ // CHECK-LABEL: BM_abs func @abs(%input : tensor<f32>) -> (tensor<f32>) { - %result = absf %input : tensor<f32> + %result = math.abs %input : tensor<f32> return %result : tensor<f32> }
diff --git a/iree/tools/test/iree-run-mlir.mlir b/iree/tools/test/iree-run-mlir.mlir index 14c64c6..3de7f89 100644 --- a/iree/tools/test/iree-run-mlir.mlir +++ b/iree/tools/test/iree-run-mlir.mlir
@@ -4,7 +4,7 @@ // CHECK-LABEL: EXEC @abs func @abs(%input : tensor<f32>) -> (tensor<f32>) { - %result = absf %input : tensor<f32> + %result = math.abs %input : tensor<f32> return %result : tensor<f32> } // CHECK: f32=2
diff --git a/iree/tools/test/iree-run-module.mlir b/iree/tools/test/iree-run-module.mlir index 2810fef..907d754 100644 --- a/iree/tools/test/iree-run-module.mlir +++ b/iree/tools/test/iree-run-module.mlir
@@ -4,7 +4,7 @@ // CHECK-LABEL: EXEC @abs func @abs(%input : tensor<f32>) -> (tensor<f32>) { - %result = absf %input : tensor<f32> + %result = math.abs %input : tensor<f32> return %result : tensor<f32> } // CHECK: f32=2
diff --git a/llvm-external-projects/iree-dialects/lib/Dialect/IREEPyDM/Transforms/ToIREE/LoweringPatterns.cpp b/llvm-external-projects/iree-dialects/lib/Dialect/IREEPyDM/Transforms/ToIREE/LoweringPatterns.cpp index e537125..308702d 100644 --- a/llvm-external-projects/iree-dialects/lib/Dialect/IREEPyDM/Transforms/ToIREE/LoweringPatterns.cpp +++ b/llvm-external-projects/iree-dialects/lib/Dialect/IREEPyDM/Transforms/ToIREE/LoweringPatterns.cpp
@@ -17,10 +17,10 @@ using namespace mlir; using namespace mlir::iree_pydm; -namespace arith_d = mlir; namespace iree_d = mlir::iree; namespace builtin_d = mlir; namespace std_d = mlir; +namespace arith_d = mlir::arith; namespace pydm_d = mlir::iree_pydm; namespace { @@ -49,7 +49,7 @@ .Default([&](Type t) -> Value { auto attr = builder.getZeroAttr(t); assert(attr && "could not get zero attr for builtin type"); - return builder.create<std_d::ConstantOp>(loc, t, attr); + return builder.create<arith_d::ConstantOp>(loc, t, attr); }); } @@ -66,12 +66,12 @@ } static Value getSuccessStatusValue(Location loc, OpBuilder &builder) { - return builder.create<std_d::ConstantOp>(loc, builder.getI32IntegerAttr(0)); + return builder.create<arith_d::ConstantOp>(loc, builder.getI32IntegerAttr(0)); } static Value getFailureStatusValue(Location loc, OpBuilder &builder, ExceptionCode code) { - return builder.create<std_d::ConstantOp>( + return builder.create<arith_d::ConstantOp>( loc, builder.getI32IntegerAttr(static_cast<int>(code))); } @@ -85,13 +85,15 @@ // Note: The list can record optional runtime state at positions > 1, so // to truly reset, we have to resize. Low level optimizations should be able // to elide this if it turns out to be unnecessary. - auto size = builder.create<std_d::ConstantOp>(loc, builder.getIndexAttr(2)); + auto size = builder.create<arith_d::ConstantOp>(loc, builder.getIndexAttr(2)); builder.create<iree_d::ListResizeOp>(loc, list, size); - auto index0 = builder.create<std_d::ConstantOp>(loc, builder.getIndexAttr(0)); - Value typeCodeValue = builder.create<std_d::ConstantOp>( + auto index0 = + builder.create<arith_d::ConstantOp>(loc, builder.getIndexAttr(0)); + Value typeCodeValue = builder.create<arith_d::ConstantOp>( loc, builder.getI32IntegerAttr(typeCode)); builder.create<iree_d::ListSetOp>(loc, list, index0, typeCodeValue); - auto index1 = builder.create<std_d::ConstantOp>(loc, builder.getIndexAttr(1)); + auto index1 = + builder.create<arith_d::ConstantOp>(loc, builder.getIndexAttr(1)); builder.create<iree_d::ListSetOp>(loc, list, index1, data); } @@ -110,9 +112,9 @@ if (inputType.getWidth() == resultType.getWidth()) { return input; } else if (inputType.getWidth() < resultType.getWidth()) { - return builder.create<arith_d::SignExtendIOp>(loc, resultType, input); + return builder.create<arith_d::ExtSIOp>(loc, resultType, input); } else { - return builder.create<arith_d::TruncateIOp>(loc, resultType, input); + return builder.create<arith_d::TruncIOp>(loc, resultType, input); } } @@ -213,22 +215,22 @@ if (dunderName == "add") { return rewriter.create<arith_d::AddIOp>(loc, left, right); } else if (dunderName == "and") { - return rewriter.create<arith_d::AndOp>(loc, left, right); + return rewriter.create<arith_d::AndIOp>(loc, left, right); } else if (dunderName == "mul") { return rewriter.create<arith_d::MulIOp>(loc, left, right); } else if (dunderName == "lshift") { - return rewriter.create<arith_d::ShiftLeftOp>(loc, left, right); + return rewriter.create<arith_d::ShLIOp>(loc, left, right); } else if (dunderName == "or") { - return rewriter.create<arith_d::OrOp>(loc, left, right); + return rewriter.create<arith_d::OrIOp>(loc, left, right); } else if (dunderName == "rshift") { if (isSigned) - return rewriter.create<arith_d::SignedShiftRightOp>(loc, left, right); + return rewriter.create<arith_d::ShRSIOp>(loc, left, right); else - return rewriter.create<arith_d::UnsignedShiftRightOp>(loc, left, right); + return rewriter.create<arith_d::ShRUIOp>(loc, left, right); } else if (dunderName == "sub") { return rewriter.create<arith_d::SubIOp>(loc, left, right); } else if (dunderName == "xor") { - return rewriter.create<arith_d::XOrOp>(loc, left, right); + return rewriter.create<arith_d::XOrIOp>(loc, left, right); } return nullptr; } @@ -361,7 +363,8 @@ if (!newValue) return rewriter.notifyMatchFailure( srcOp, "constant cannot be represented as a standard constant"); - rewriter.replaceOpWithNewOp<std_d::ConstantOp>(srcOp, resultType, newValue); + rewriter.replaceOpWithNewOp<arith_d::ConstantOp>(srcOp, resultType, + newValue); return success(); } }; @@ -376,7 +379,7 @@ ConversionPatternRewriter &rewriter) const override { Type i32 = rewriter.getI32Type(); // '-3' == RuntimeError - rewriter.replaceOpWithNewOp<std_d::ConstantOp>( + rewriter.replaceOpWithNewOp<arith_d::ConstantOp>( srcOp, i32, rewriter.getIntegerAttr(i32, -3)); return success(); } @@ -446,7 +449,7 @@ "result type could not be converted"); Type i32Type = rewriter.getIntegerType(32); Value index0 = - rewriter.create<std_d::ConstantOp>(loc, rewriter.getIndexAttr(0)); + rewriter.create<arith_d::ConstantOp>(loc, rewriter.getIndexAttr(0)); Value typeCode = rewriter.create<iree_d::ListGetOp>( loc, i32Type, adaptor.value(), index0); rewriter.replaceOp( @@ -471,7 +474,7 @@ srcOp, "could not convert load_var result type"); auto list = operands[0]; auto index1 = - rewriter.create<std_d::ConstantOp>(loc, rewriter.getIndexAttr(1)); + rewriter.create<arith_d::ConstantOp>(loc, rewriter.getIndexAttr(1)); rewriter.replaceOpWithNewOp<iree_d::ListGetOp>(srcOp, resultType, list, index1); return success(); @@ -488,7 +491,7 @@ pydm_d::NoneOp srcOp, ArrayRef<Value> operands, ConversionPatternRewriter &rewriter) const override { Type i32 = rewriter.getI32Type(); - rewriter.replaceOpWithNewOp<std_d::ConstantOp>( + rewriter.replaceOpWithNewOp<arith_d::ConstantOp>( srcOp, i32, rewriter.getIntegerAttr(i32, 0)); return success(); } @@ -523,8 +526,8 @@ // Branch on success conditional. rewriter.setInsertionPointToEnd(entryBlock); Value successValue = getSuccessStatusValue(loc, rewriter); - Value isSuccess = rewriter.create<std_d::CmpIOp>( - loc, std_d::CmpIPredicate::eq, successValue, status); + Value isSuccess = rewriter.create<arith_d::CmpIOp>( + loc, arith_d::CmpIPredicate::eq, successValue, status); rewriter.create<std_d::CondBranchOp>(loc, isSuccess, continuationBlock, raiseAndReturnBlock); rewriter.eraseOp(srcOp); @@ -545,8 +548,8 @@ pydm_d::ReturnOp srcOp, ArrayRef<Value> operands, ConversionPatternRewriter &rewriter) const override { auto loc = srcOp.getLoc(); - auto zeroResult = - rewriter.create<std_d::ConstantOp>(loc, rewriter.getI32IntegerAttr(0)); + auto zeroResult = rewriter.create<arith_d::ConstantOp>( + loc, rewriter.getI32IntegerAttr(0)); rewriter.replaceOpWithNewOp<std_d::ReturnOp>( srcOp, ValueRange{zeroResult, operands[0]}); return success(); @@ -615,13 +618,13 @@ { rewriter.setInsertionPointToEnd(entryBlock); auto index0 = - rewriter.create<std_d::ConstantOp>(loc, rewriter.getIndexAttr(0)); - Value requiredTypeCodeValue = rewriter.create<std_d::ConstantOp>( + rewriter.create<arith_d::ConstantOp>(loc, rewriter.getIndexAttr(0)); + Value requiredTypeCodeValue = rewriter.create<arith_d::ConstantOp>( loc, rewriter.getI32IntegerAttr(typeCode)); Value actualTypeCodeValue = rewriter.create<iree_d::ListGetOp>( loc, rewriter.getI32Type(), list, index0); - Value typeCodeEqual = rewriter.create<std_d::CmpIOp>( - loc, std_d::CmpIPredicate::eq, requiredTypeCodeValue, + Value typeCodeEqual = rewriter.create<arith_d::CmpIOp>( + loc, arith_d::CmpIPredicate::eq, requiredTypeCodeValue, actualTypeCodeValue); rewriter.create<std_d::CondBranchOp>(loc, typeCodeEqual, typesMatchBlock, slowPathMismatchBlock); @@ -631,7 +634,7 @@ { rewriter.setInsertionPointToEnd(typesMatchBlock); auto index1 = - rewriter.create<std_d::ConstantOp>(loc, rewriter.getIndexAttr(1)); + rewriter.create<arith_d::ConstantOp>(loc, rewriter.getIndexAttr(1)); Value successResult = getSuccessStatusValue(loc, rewriter); Value unboxedValue = rewriter.create<iree_d::ListGetOp>( loc, targetUnboxedType, list, index1);
diff --git a/llvm-external-projects/iree-dialects/test/iree_pydm/canonicalize/booleans.mlir b/llvm-external-projects/iree-dialects/test/iree_pydm/canonicalize/booleans.mlir index a0c134d..6917bae 100644 --- a/llvm-external-projects/iree-dialects/test/iree_pydm/canonicalize/booleans.mlir +++ b/llvm-external-projects/iree-dialects/test/iree_pydm/canonicalize/booleans.mlir
@@ -2,7 +2,7 @@ // CHECK-LABEL: @fold_none iree_pydm.func @fold_none(%arg0 : !iree_pydm.none) -> (!iree_pydm.exception_result, !iree_pydm.bool) { - // CHECK: %[[F:.*]] = constant false -> !iree_pydm.bool + // CHECK: %[[F:.*]] = arith.constant false -> !iree_pydm.bool // CHECK: return %[[F]] %0 = as_bool %arg0 : !iree_pydm.none -> !iree_pydm.bool return %0 : !iree_pydm.bool @@ -19,9 +19,9 @@ // ----- // CHECK-LABEL: @as_bool_from_integer_to_compare iree_pydm.func @as_bool_from_integer_to_compare(%arg0 : !iree_pydm.integer) -> (!iree_pydm.exception_result, !iree_pydm.bool) { - // CHECK-DAG: %[[Z:.*]] = constant 0 : i64 -> !iree_pydm.integer - // CHECK-DAG: %[[T:.*]] = constant true -> !iree_pydm.bool - // CHECK-DAG: %[[F:.*]] = constant false -> !iree_pydm.bool + // CHECK-DAG: %[[Z:.*]] = arith.constant 0 : i64 -> !iree_pydm.integer + // CHECK-DAG: %[[T:.*]] = arith.constant true -> !iree_pydm.bool + // CHECK-DAG: %[[F:.*]] = arith.constant false -> !iree_pydm.bool // CHECK: %[[CMP:.*]] = apply_compare "eq", %arg0, %[[Z]] // CHECK: %[[SEL:.*]] = select %[[CMP]], %[[F]], %[[T]] : !iree_pydm.bool // CHECK: return %[[SEL]] @@ -32,9 +32,9 @@ // ----- // CHECK-LABEL: @as_bool_from_real_to_compare iree_pydm.func @as_bool_from_real_to_compare(%arg0 : !iree_pydm.real) -> (!iree_pydm.exception_result, !iree_pydm.bool) { - // CHECK-DAG: %[[Z:.*]] = constant 0.000000e+00 : f64 -> !iree_pydm.real - // CHECK-DAG: %[[T:.*]] = constant true -> !iree_pydm.bool - // CHECK-DAG: %[[F:.*]] = constant false -> !iree_pydm.bool + // CHECK-DAG: %[[Z:.*]] = arith.constant 0.000000e+00 : f64 -> !iree_pydm.real + // CHECK-DAG: %[[T:.*]] = arith.constant true -> !iree_pydm.bool + // CHECK-DAG: %[[F:.*]] = arith.constant false -> !iree_pydm.bool // CHECK: %[[CMP:.*]] = apply_compare "eq", %arg0, %[[Z]] // CHECK: %[[SEL:.*]] = select %[[CMP]], %[[F]], %[[T]] : !iree_pydm.bool // CHECK: return %[[SEL]] @@ -46,7 +46,7 @@ // CHECK-LABEL: @fold_bool_to_pred_from_constant iree_pydm.func @fold_bool_to_pred_from_constant() -> (!iree_pydm.exception_result, !iree_pydm.none) { %0 = iree_pydm.constant true -> !iree_pydm.bool - // CHECK: %[[P:.*]] = constant true -> i1 + // CHECK: %[[P:.*]] = arith.constant true -> i1 // CHECK: "custom.donotoptimize"(%[[P]]) %1 = bool_to_pred %0 "custom.donotoptimize"(%1) : (i1) -> () @@ -58,7 +58,7 @@ // CHECK-LABEL: @as_bool_from_real_to_compare iree_pydm.func @as_bool_from_real_to_compare(%arg0: !iree_pydm.integer, %arg1: !iree_pydm.integer) -> (!iree_pydm.exception_result, !iree_pydm.integer) { // CHECK: return %arg0 - %0 = constant true -> !iree_pydm.bool + %0 = arith.constant true -> !iree_pydm.bool %1 = select %0, %arg0, %arg1 : !iree_pydm.integer return %1 : !iree_pydm.integer }
diff --git a/llvm-external-projects/iree-dialects/test/iree_pydm/canonicalize/numerics.mlir b/llvm-external-projects/iree-dialects/test/iree_pydm/canonicalize/numerics.mlir index e302baa..43d5ca4 100644 --- a/llvm-external-projects/iree-dialects/test/iree_pydm/canonicalize/numerics.mlir +++ b/llvm-external-projects/iree-dialects/test/iree_pydm/canonicalize/numerics.mlir
@@ -53,9 +53,9 @@ // ----- // CHECK-LABEL: @fold_promote_numeric_true_to_integer iree_pydm.func @fold_promote_numeric_true_to_integer() -> (!iree_pydm.exception_result, !iree_pydm.none) { - // CHECK: %[[CST:.*]] = constant 1 : i64 -> !iree_pydm.integer + // CHECK: %[[CST:.*]] = arith.constant 1 : i64 -> !iree_pydm.integer // CHECK: "custom.donotoptimize"(%[[CST]]) - %0 = constant true -> !iree_pydm.bool + %0 = arith.constant true -> !iree_pydm.bool %1 = promote_numeric %0 : !iree_pydm.bool -> !iree_pydm.integer "custom.donotoptimize"(%1) : (!iree_pydm.integer) -> () %none = none @@ -65,9 +65,9 @@ // ----- // CHECK-LABEL: @fold_promote_numeric_false_to_integer iree_pydm.func @fold_promote_numeric_false_to_integer() -> (!iree_pydm.exception_result, !iree_pydm.none) { - // CHECK: %[[CST:.*]] = constant 0 : i64 -> !iree_pydm.integer + // CHECK: %[[CST:.*]] = arith.constant 0 : i64 -> !iree_pydm.integer // CHECK: "custom.donotoptimize"(%[[CST]]) - %0 = constant false -> !iree_pydm.bool + %0 = arith.constant false -> !iree_pydm.bool %1 = promote_numeric %0 : !iree_pydm.bool -> !iree_pydm.integer "custom.donotoptimize"(%1) : (!iree_pydm.integer) -> () %none = none @@ -77,9 +77,9 @@ // ----- // CHECK-LABEL: @fold_promote_numeric_true_to_real iree_pydm.func @fold_promote_numeric_true_to_real() -> (!iree_pydm.exception_result, !iree_pydm.none) { - // CHECK: %[[CST:.*]] = constant 1.000000e+00 : f64 -> !iree_pydm.real + // CHECK: %[[CST:.*]] = arith.constant 1.000000e+00 : f64 -> !iree_pydm.real // CHECK: "custom.donotoptimize"(%[[CST]]) - %0 = constant true -> !iree_pydm.bool + %0 = arith.constant true -> !iree_pydm.bool %1 = promote_numeric %0 : !iree_pydm.bool -> !iree_pydm.real "custom.donotoptimize"(%1) : (!iree_pydm.real) -> () %none = none @@ -89,9 +89,9 @@ // ----- // CHECK-LABEL: @fold_promote_numeric_false_to_real iree_pydm.func @fold_promote_numeric_false_to_real() -> (!iree_pydm.exception_result, !iree_pydm.none) { - // CHECK: %[[CST:.*]] = constant 0.000000e+00 : f64 -> !iree_pydm.real + // CHECK: %[[CST:.*]] = arith.constant 0.000000e+00 : f64 -> !iree_pydm.real // CHECK: "custom.donotoptimize"(%[[CST]]) - %0 = constant false -> !iree_pydm.bool + %0 = arith.constant false -> !iree_pydm.bool %1 = promote_numeric %0 : !iree_pydm.bool -> !iree_pydm.real "custom.donotoptimize"(%1) : (!iree_pydm.real) -> () %none = none @@ -101,9 +101,9 @@ // ----- // CHECK-LABEL: @fold_promote_numeric_integet_to_real iree_pydm.func @fold_promote_numeric_integet_to_real() -> (!iree_pydm.exception_result, !iree_pydm.none) { - // CHECK: %[[CST:.*]] = constant 2.000000e+00 : f64 -> !iree_pydm.real + // CHECK: %[[CST:.*]] = arith.constant 2.000000e+00 : f64 -> !iree_pydm.real // CHECK: "custom.donotoptimize"(%[[CST]]) - %0 = constant 2 : i64 -> !iree_pydm.integer + %0 = arith.constant 2 : i64 -> !iree_pydm.integer %1 = promote_numeric %0 : !iree_pydm.integer -> !iree_pydm.real "custom.donotoptimize"(%1) : (!iree_pydm.real) -> () %none = none
diff --git a/llvm-external-projects/iree-dialects/test/iree_pydm/to_iree/constants.mlir b/llvm-external-projects/iree-dialects/test/iree_pydm/to_iree/constants.mlir index d8fcb91..d0ea71b 100644 --- a/llvm-external-projects/iree-dialects/test/iree_pydm/to_iree/constants.mlir +++ b/llvm-external-projects/iree-dialects/test/iree_pydm/to_iree/constants.mlir
@@ -2,8 +2,8 @@ // CHECK-LABEL: @none_constant iree_pydm.func @none_constant() -> (!iree_pydm.exception_result, !iree_pydm.none) { - // CHECK: %[[CST0:.*]] = constant 0 : i32 - // CHECK: %[[CST1:.*]] = constant 0 : i32 + // CHECK: %[[CST0:.*]] = arith.constant 0 : i32 + // CHECK: %[[CST1:.*]] = arith.constant 0 : i32 // CHECK: return %[[CST1]], %[[CST0]] %0 = none return %0 : !iree_pydm.none @@ -12,13 +12,13 @@ // CHECK-LABEL: @constant_integer_trunc iree_pydm.func @constant_integer_trunc() -> (!iree_pydm.exception_result, !iree_pydm.integer) { // CHECK: constant -10 : i32 - %0 = constant -10 : i64 -> !iree_pydm.integer + %0 = arith.constant -10 : i64 -> !iree_pydm.integer return %0 : !iree_pydm.integer } // CHECK-LABEL: @constant_real_trunc iree_pydm.func @constant_real_trunc() -> (!iree_pydm.exception_result, !iree_pydm.real) { // CHECK: constant -2.000000e+00 : f32 - %0 = constant -2.0 : f64 -> !iree_pydm.real + %0 = arith.constant -2.0 : f64 -> !iree_pydm.real return %0 : !iree_pydm.real }
diff --git a/llvm-external-projects/iree-dialects/test/iree_pydm/to_iree/integer_compare.mlir b/llvm-external-projects/iree-dialects/test/iree_pydm/to_iree/integer_compare.mlir index dfd387c..c0dfc32 100644 --- a/llvm-external-projects/iree-dialects/test/iree_pydm/to_iree/integer_compare.mlir +++ b/llvm-external-projects/iree-dialects/test/iree_pydm/to_iree/integer_compare.mlir
@@ -2,7 +2,7 @@ // CHECK-LABEL: @lt iree_pydm.func @lt(%arg0 : !iree_pydm.integer, %arg1 : !iree_pydm.integer) -> (!iree_pydm.exception_result, !iree_pydm.bool) { - // CHECK: %[[R:.*]] = cmpi slt, %arg0, %arg1 : i32 + // CHECK: %[[R:.*]] = arith.cmpi slt, %arg0, %arg1 : i32 %0 = apply_compare "lt", %arg0, %arg1 : !iree_pydm.integer, !iree_pydm.integer // CHECK: return {{.*}}, %[[R]] return %0 : !iree_pydm.bool @@ -10,7 +10,7 @@ // CHECK-LABEL: @le iree_pydm.func @le(%arg0 : !iree_pydm.integer, %arg1 : !iree_pydm.integer) -> (!iree_pydm.exception_result, !iree_pydm.bool) { - // CHECK: %[[R:.*]] = cmpi sle, %arg0, %arg1 : i32 + // CHECK: %[[R:.*]] = arith.cmpi sle, %arg0, %arg1 : i32 %0 = apply_compare "le", %arg0, %arg1 : !iree_pydm.integer, !iree_pydm.integer // CHECK: return {{.*}}, %[[R]] return %0 : !iree_pydm.bool @@ -18,7 +18,7 @@ // CHECK-LABEL: @eq iree_pydm.func @eq(%arg0 : !iree_pydm.integer, %arg1 : !iree_pydm.integer) -> (!iree_pydm.exception_result, !iree_pydm.bool) { - // CHECK: %[[R:.*]] = cmpi eq, %arg0, %arg1 : i32 + // CHECK: %[[R:.*]] = arith.cmpi eq, %arg0, %arg1 : i32 %0 = apply_compare "eq", %arg0, %arg1 : !iree_pydm.integer, !iree_pydm.integer // CHECK: return {{.*}}, %[[R]] return %0 : !iree_pydm.bool @@ -26,7 +26,7 @@ // CHECK-LABEL: @is iree_pydm.func @is(%arg0 : !iree_pydm.integer, %arg1 : !iree_pydm.integer) -> (!iree_pydm.exception_result, !iree_pydm.bool) { - // CHECK: %[[R:.*]] = cmpi eq, %arg0, %arg1 : i32 + // CHECK: %[[R:.*]] = arith.cmpi eq, %arg0, %arg1 : i32 %0 = apply_compare "is", %arg0, %arg1 : !iree_pydm.integer, !iree_pydm.integer // CHECK: return {{.*}}, %[[R]] return %0 : !iree_pydm.bool @@ -34,7 +34,7 @@ // CHECK-LABEL: @ne iree_pydm.func @ne(%arg0 : !iree_pydm.integer, %arg1 : !iree_pydm.integer) -> (!iree_pydm.exception_result, !iree_pydm.bool) { - // CHECK: %[[R:.*]] = cmpi ne, %arg0, %arg1 : i32 + // CHECK: %[[R:.*]] = arith.cmpi ne, %arg0, %arg1 : i32 %0 = apply_compare "ne", %arg0, %arg1 : !iree_pydm.integer, !iree_pydm.integer // CHECK: return {{.*}}, %[[R]] return %0 : !iree_pydm.bool @@ -42,7 +42,7 @@ // CHECK-LABEL: @isnot iree_pydm.func @isnot(%arg0 : !iree_pydm.integer, %arg1 : !iree_pydm.integer) -> (!iree_pydm.exception_result, !iree_pydm.bool) { - // CHECK: %[[R:.*]] = cmpi ne, %arg0, %arg1 : i32 + // CHECK: %[[R:.*]] = arith.cmpi ne, %arg0, %arg1 : i32 %0 = apply_compare "isnot", %arg0, %arg1 : !iree_pydm.integer, !iree_pydm.integer // CHECK: return {{.*}}, %[[R]] return %0 : !iree_pydm.bool @@ -50,7 +50,7 @@ // CHECK-LABEL: @gt iree_pydm.func @gt(%arg0 : !iree_pydm.integer, %arg1 : !iree_pydm.integer) -> (!iree_pydm.exception_result, !iree_pydm.bool) { - // CHECK: %[[R:.*]] = cmpi sgt, %arg0, %arg1 : i32 + // CHECK: %[[R:.*]] = arith.cmpi sgt, %arg0, %arg1 : i32 %0 = apply_compare "gt", %arg0, %arg1 : !iree_pydm.integer, !iree_pydm.integer // CHECK: return {{.*}}, %[[R]] return %0 : !iree_pydm.bool @@ -58,7 +58,7 @@ // CHECK-LABEL: @ge iree_pydm.func @ge(%arg0 : !iree_pydm.integer, %arg1 : !iree_pydm.integer) -> (!iree_pydm.exception_result, !iree_pydm.bool) { - // CHECK: %[[R:.*]] = cmpi sge, %arg0, %arg1 : i32 + // CHECK: %[[R:.*]] = arith.cmpi sge, %arg0, %arg1 : i32 %0 = apply_compare "ge", %arg0, %arg1 : !iree_pydm.integer, !iree_pydm.integer // CHECK: return {{.*}}, %[[R]] return %0 : !iree_pydm.bool
diff --git a/llvm-external-projects/iree-dialects/test/iree_pydm/to_iree/real_compare.mlir b/llvm-external-projects/iree-dialects/test/iree_pydm/to_iree/real_compare.mlir index 83601c1..fb38be0 100644 --- a/llvm-external-projects/iree-dialects/test/iree_pydm/to_iree/real_compare.mlir +++ b/llvm-external-projects/iree-dialects/test/iree_pydm/to_iree/real_compare.mlir
@@ -2,7 +2,7 @@ // CHECK-LABEL: @lt iree_pydm.func @lt(%arg0 : !iree_pydm.real, %arg1 : !iree_pydm.real) -> (!iree_pydm.exception_result, !iree_pydm.bool) { - // CHECK: %[[R:.*]] = cmpf olt, %arg0, %arg1 : f32 + // CHECK: %[[R:.*]] = arith.cmpf olt, %arg0, %arg1 : f32 %0 = apply_compare "lt", %arg0, %arg1 : !iree_pydm.real, !iree_pydm.real // CHECK: return {{.*}}, %[[R]] return %0 : !iree_pydm.bool @@ -10,7 +10,7 @@ // CHECK-LABEL: @le iree_pydm.func @le(%arg0 : !iree_pydm.real, %arg1 : !iree_pydm.real) -> (!iree_pydm.exception_result, !iree_pydm.bool) { - // CHECK: %[[R:.*]] = cmpf ole, %arg0, %arg1 : f32 + // CHECK: %[[R:.*]] = arith.cmpf ole, %arg0, %arg1 : f32 %0 = apply_compare "le", %arg0, %arg1 : !iree_pydm.real, !iree_pydm.real // CHECK: return {{.*}}, %[[R]] return %0 : !iree_pydm.bool @@ -18,7 +18,7 @@ // CHECK-LABEL: @eq iree_pydm.func @eq(%arg0 : !iree_pydm.real, %arg1 : !iree_pydm.real) -> (!iree_pydm.exception_result, !iree_pydm.bool) { - // CHECK: %[[R:.*]] = cmpf oeq, %arg0, %arg1 : f32 + // CHECK: %[[R:.*]] = arith.cmpf oeq, %arg0, %arg1 : f32 %0 = apply_compare "eq", %arg0, %arg1 : !iree_pydm.real, !iree_pydm.real // CHECK: return {{.*}}, %[[R]] return %0 : !iree_pydm.bool @@ -26,7 +26,7 @@ // CHECK-LABEL: @is iree_pydm.func @is(%arg0 : !iree_pydm.real, %arg1 : !iree_pydm.real) -> (!iree_pydm.exception_result, !iree_pydm.bool) { - // CHECK: %[[R:.*]] = cmpf oeq, %arg0, %arg1 : f32 + // CHECK: %[[R:.*]] = arith.cmpf oeq, %arg0, %arg1 : f32 %0 = apply_compare "is", %arg0, %arg1 : !iree_pydm.real, !iree_pydm.real // CHECK: return {{.*}}, %[[R]] return %0 : !iree_pydm.bool @@ -34,7 +34,7 @@ // CHECK-LABEL: @ne iree_pydm.func @ne(%arg0 : !iree_pydm.real, %arg1 : !iree_pydm.real) -> (!iree_pydm.exception_result, !iree_pydm.bool) { - // CHECK: %[[R:.*]] = cmpf one, %arg0, %arg1 : f32 + // CHECK: %[[R:.*]] = arith.cmpf one, %arg0, %arg1 : f32 %0 = apply_compare "ne", %arg0, %arg1 : !iree_pydm.real, !iree_pydm.real // CHECK: return {{.*}}, %[[R]] return %0 : !iree_pydm.bool @@ -42,7 +42,7 @@ // CHECK-LABEL: @isnot iree_pydm.func @isnot(%arg0 : !iree_pydm.real, %arg1 : !iree_pydm.real) -> (!iree_pydm.exception_result, !iree_pydm.bool) { - // CHECK: %[[R:.*]] = cmpf one, %arg0, %arg1 : f32 + // CHECK: %[[R:.*]] = arith.cmpf one, %arg0, %arg1 : f32 %0 = apply_compare "isnot", %arg0, %arg1 : !iree_pydm.real, !iree_pydm.real // CHECK: return {{.*}}, %[[R]] return %0 : !iree_pydm.bool @@ -50,7 +50,7 @@ // CHECK-LABEL: @gt iree_pydm.func @gt(%arg0 : !iree_pydm.real, %arg1 : !iree_pydm.real) -> (!iree_pydm.exception_result, !iree_pydm.bool) { - // CHECK: %[[R:.*]] = cmpf ogt, %arg0, %arg1 : f32 + // CHECK: %[[R:.*]] = arith.cmpf ogt, %arg0, %arg1 : f32 %0 = apply_compare "gt", %arg0, %arg1 : !iree_pydm.real, !iree_pydm.real // CHECK: return {{.*}}, %[[R]] return %0 : !iree_pydm.bool @@ -58,7 +58,7 @@ // CHECK-LABEL: @ge iree_pydm.func @ge(%arg0 : !iree_pydm.real, %arg1 : !iree_pydm.real) -> (!iree_pydm.exception_result, !iree_pydm.bool) { - // CHECK: %[[R:.*]] = cmpf oge, %arg0, %arg1 : f32 + // CHECK: %[[R:.*]] = arith.cmpf oge, %arg0, %arg1 : f32 %0 = apply_compare "ge", %arg0, %arg1 : !iree_pydm.real, !iree_pydm.real // CHECK: return {{.*}}, %[[R]] return %0 : !iree_pydm.bool
diff --git a/llvm-external-projects/iree-dialects/test/iree_pydm/to_iree/structural.mlir b/llvm-external-projects/iree-dialects/test/iree_pydm/to_iree/structural.mlir index b54dae6..fbd944f 100644 --- a/llvm-external-projects/iree-dialects/test/iree_pydm/to_iree/structural.mlir +++ b/llvm-external-projects/iree-dialects/test/iree_pydm/to_iree/structural.mlir
@@ -29,14 +29,14 @@ // NOTE: "9" is the type code for integer iree_pydm.func @box(%arg0 : !iree_pydm.integer) -> (!iree_pydm.exception_result, !iree_pydm.object<!iree_pydm.integer>) { // CHECK: %[[LIST:.*]] = iree.list.create : !iree.list<!iree.variant> - // CHECK: %[[c2:.*]] = constant 2 : index + // CHECK: %[[c2:.*]] = arith.constant 2 : index // CHECK: iree.list.resize %[[LIST]], %c2 : !iree.list<!iree.variant> - // CHECK: %[[c0:.*]] = constant 0 : index - // CHECK: %[[c9:.*]] = constant 9 : i32 + // CHECK: %[[c0:.*]] = arith.constant 0 : index + // CHECK: %[[c9:.*]] = arith.constant 9 : i32 // CHECK: iree.list.set %[[LIST]][%[[c0]]], %[[c9]] : !iree.list<!iree.variant>, i32 - // CHECK: %[[c1:.*]] = constant 1 : index + // CHECK: %[[c1:.*]] = arith.constant 1 : index // CHECK: iree.list.set %[[LIST]][%[[c1]]], %arg0 : !iree.list<!iree.variant>, i32 - // CHECK: %[[c0_i32:.*]] = constant 0 : i32 + // CHECK: %[[c0_i32:.*]] = arith.constant 0 : i32 // return %[[c0_i32]], %[[LIST]] : i32, !iree.list<!iree.variant> %0 = box %arg0 : !iree_pydm.integer -> !iree_pydm.object<!iree_pydm.integer> return %0 : !iree_pydm.object<!iree_pydm.integer> @@ -48,15 +48,15 @@ iree_pydm.func @alloc_store_load_var(%arg0 : !iree_pydm.object) -> (!iree_pydm.exception_result, !iree_pydm.object) { // CHECK: %[[A:.*]] = iree.list.create : !iree.list<!iree.variant> %a = alloc_free_var "a" -> !iree_pydm.free_var_ref - // CHECK: %[[c2:.*]] = constant 2 : index + // CHECK: %[[c2:.*]] = arith.constant 2 : index // CHECK: iree.list.resize %[[A]], %[[c2]] : !iree.list<!iree.variant> - // CHECK: %[[c0:.*]] = constant 0 : index - // CHECK: %[[object_code:.*]] = constant 256 : i32 + // CHECK: %[[c0:.*]] = arith.constant 0 : index + // CHECK: %[[object_code:.*]] = arith.constant 256 : i32 // CHECK: iree.list.set %[[A]][%[[c0]]], %[[object_code]] - // CHECK: %[[c1:.*]] = constant 1 : index + // CHECK: %[[c1:.*]] = arith.constant 1 : index // CHECK: iree.list.set %[[A]][%[[c1]]], %arg0 : !iree.list<!iree.variant>, !iree.list<!iree.variant> store_var %a = %arg0 : !iree_pydm.free_var_ref, !iree_pydm.object - // CHECK: %[[c1_0:.*]] = constant 1 : index + // CHECK: %[[c1_0:.*]] = arith.constant 1 : index // CHECK: %[[LOADED:.*]] = iree.list.get %[[A]][%[[c1_0]]] : !iree.list<!iree.variant> -> !iree.list<!iree.variant> %0 = load_var %a : !iree_pydm.free_var_ref -> !iree_pydm.object // CHECK: return {{.*}}, %[[LOADED]] @@ -67,16 +67,16 @@ // CHECK-LABEL: @unbox // NOTE: "9" is the type code for integer iree_pydm.func @unbox(%arg0 : !iree_pydm.object) -> (!iree_pydm.exception_result, !iree_pydm.integer) { - // CHECK: %[[c0:.*]] = constant 0 : index - // CHECK: %[[NEEDED_TYPE_CODE:.*]] = constant 9 : i32 + // CHECK: %[[c0:.*]] = arith.constant 0 : index + // CHECK: %[[NEEDED_TYPE_CODE:.*]] = arith.constant 9 : i32 // CHECK: %[[TYPE_CODE:.*]] = iree.list.get %arg0[%[[c0]]] : !iree.list<!iree.variant> -> i32 // CHECK: %[[TYPE_EQ:.*]] = cmpi eq, %[[NEEDED_TYPE_CODE]], %[[TYPE_CODE]] : i32 // CHECK: cond_br %[[TYPE_EQ]], ^bb1, ^bb4 // bb1: On equal // CHECK: ^bb1: - // CHECK: %[[c1:.*]] = constant 1 : index - // CHECK: %[[c0_i32:.*]] = constant 0 : i32 + // CHECK: %[[c1:.*]] = arith.constant 1 : index + // CHECK: %[[c0_i32:.*]] = arith.constant 0 : i32 // CHECK: %[[CONTENTS:.*]] = iree.list.get %arg0[%[[c1]]] : !iree.list<!iree.variant> -> i32 // CHECK: br ^bb2(%[[c0_i32]], %[[CONTENTS]] : i32, i32) @@ -88,8 +88,8 @@ // bb4: Signal ValueError (-4 == ValueError) // CHECK: ^bb4: - // CHECK: %[[VALUE_ERROR_CODE:.*]] = constant -4 : i32 - // CHECK: %[[c0_i32_2:.*]] = constant 0 : i32 + // CHECK: %[[VALUE_ERROR_CODE:.*]] = arith.constant -4 : i32 + // CHECK: %[[c0_i32_2:.*]] = arith.constant 0 : i32 // CHECK: br ^bb2(%[[VALUE_ERROR_CODE]], %[[c0_i32_2]] : i32, i32) %status, %primitive = unbox %arg0 : !iree_pydm.object -> !iree_pydm.integer raise_on_failure %status : !iree_pydm.exception_result @@ -99,12 +99,12 @@ // ----- // CHECK-LABEL: @raise_on_failure_object_return iree_pydm.func @raise_on_failure_object_return(%arg0 : !iree_pydm.exception_result, %arg1: !iree_pydm.object) -> (!iree_pydm.exception_result, !iree_pydm.object) { - // CHECK: %[[c0_i32:.*]] = constant 0 : i32 + // CHECK: %[[c0_i32:.*]] = arith.constant 0 : i32 // CHECK: %[[CMP:.*]] = cmpi eq, %[[c0_i32]], %arg0 : i32 // CHECK: cond_br %[[CMP]], ^bb1, ^bb2 // bb1: success // CHECK: ^bb1: - // CHECK: %[[c0_i32_0:.*]] = constant 0 : i32 + // CHECK: %[[c0_i32_0:.*]] = arith.constant 0 : i32 // CHECK: return %[[c0_i32_0]], %arg1 : i32, !iree.list<!iree.variant> // bb2: failure // CHECK: ^bb2: @@ -119,7 +119,7 @@ iree_pydm.func @raise_on_failure_builtin(%arg0 : !iree_pydm.exception_result, %arg1: !iree_pydm.integer) -> (!iree_pydm.exception_result, !iree_pydm.integer) { // bb2: failure // CHECK: ^bb2: - // CHECK: %[[ZERO:.*]] = constant 0 : i32 + // CHECK: %[[ZERO:.*]] = arith.constant 0 : i32 // CHECK: return %arg0, %[[ZERO]] : i32, i32 raise_on_failure %arg0 : !iree_pydm.exception_result return %arg1 : !iree_pydm.integer @@ -141,7 +141,7 @@ // ----- // CHECK-LABEL: @get_type_code iree_pydm.func @get_type_code(%arg0 : !iree_pydm.object) -> (!iree_pydm.exception_result, !iree_pydm.integer) { - // CHECK: %[[c0:.*]] = constant 0 : index + // CHECK: %[[c0:.*]] = arith.constant 0 : index // CHECK: %[[R:.*]] = iree.list.get %arg0[%[[c0]]] : !iree.list<!iree.variant> -> i32 %0 = get_type_code %arg0 : !iree_pydm.object -> !iree_pydm.integer return %0 : !iree_pydm.integer
diff --git a/llvm-external-projects/iree-dialects/test/python/iree_pydm/importer/assignment.py b/llvm-external-projects/iree-dialects/test/python/iree_pydm/importer/assignment.py index 830195c..424a349 100644 --- a/llvm-external-projects/iree-dialects/test/python/iree_pydm/importer/assignment.py +++ b/llvm-external-projects/iree-dialects/test/python/iree_pydm/importer/assignment.py
@@ -5,7 +5,7 @@ # CHECK-LABEL: @assign_free_var_not_arg -# CHECK: %[[CST:.*]] = constant 1 +# CHECK: %[[CST:.*]] = arith.constant 1 # CHECK: %[[BOXED:.*]] = box %[[CST]] : !iree_pydm.integer -> !iree_pydm.object<!iree_pydm.integer> # CHECK: store_var %x = %[[BOXED]] @test_import_global
diff --git a/llvm-external-projects/iree-dialects/test/python/iree_pydm/importer/booleans.py b/llvm-external-projects/iree-dialects/test/python/iree_pydm/importer/booleans.py index 36da239..fd5f77f 100644 --- a/llvm-external-projects/iree-dialects/test/python/iree_pydm/importer/booleans.py +++ b/llvm-external-projects/iree-dialects/test/python/iree_pydm/importer/booleans.py
@@ -55,8 +55,8 @@ # CHECK-LABEL: func @logical_not # CHECK: %[[XVAL:.*]] = load_var %x # CHECK: %[[XBOOL:.*]] = as_bool %[[XVAL]] -# CHECK: %[[T:.*]] = constant true -# CHECK: %[[F:.*]] = constant false +# CHECK: %[[T:.*]] = arith.constant true +# CHECK: %[[F:.*]] = arith.constant false # CHECK: %[[R:.*]] = select %[[XBOOL]], %[[F]], %[[T]] @test_import_global def logical_not(): @@ -68,11 +68,11 @@ # CHECK: %[[XVAL:.*]] = load_var %x # CHECK: %[[XBOOL:.*]] = as_bool %[[XVAL]] # CHECK: %[[R1:.*]] = functional_if %[[XBOOL]] {{.*}} { -# CHECK: %[[TWOVAL:.*]] = constant 2 +# CHECK: %[[TWOVAL:.*]] = arith.constant 2 # CHECK: %[[TWOBOXED:.*]] = box %[[TWOVAL]] : !iree_pydm.integer -> !iree_pydm.object # CHECK: yield %[[TWOBOXED]] # CHECK: } else { -# CHECK: %[[THREEVAL:.*]] = constant 3 +# CHECK: %[[THREEVAL:.*]] = arith.constant 3 # CHECK: %[[THREEBOXED:.*]] = box %[[THREEVAL]] : !iree_pydm.integer -> !iree_pydm.object # CHECK: yield %[[THREEBOXED]] # CHECK: }
diff --git a/llvm-external-projects/iree-dialects/test/python/iree_pydm/importer/comparison.py b/llvm-external-projects/iree-dialects/test/python/iree_pydm/importer/comparison.py index 1a0efc7..fcd3e5f 100644 --- a/llvm-external-projects/iree-dialects/test/python/iree_pydm/importer/comparison.py +++ b/llvm-external-projects/iree-dialects/test/python/iree_pydm/importer/comparison.py
@@ -110,7 +110,7 @@ # CHECK-LABEL: @short_circuit -# CHECK-DAG: %[[FALSE:.*]] = constant false +# CHECK-DAG: %[[FALSE:.*]] = arith.constant false # CHECK-DAG: %[[X:.*]] = load_var %x # CHECK-DAG: %[[Y:.*]] = load_var %y # CHECK: %[[XP:.*]], %[[YP:.*]] = dynamic_binary_promote %[[X]], %[[Y]]
diff --git a/llvm-external-projects/iree-dialects/test/python/iree_pydm/importer/constants.py b/llvm-external-projects/iree-dialects/test/python/iree_pydm/importer/constants.py index 2c6823b..47f92ac 100644 --- a/llvm-external-projects/iree-dialects/test/python/iree_pydm/importer/constants.py +++ b/llvm-external-projects/iree-dialects/test/python/iree_pydm/importer/constants.py
@@ -5,28 +5,28 @@ # CHECK-LABEL: @const_integer -# CHECK: = constant 1 : i64 -> !iree_pydm.integer +# CHECK: = arith.constant 1 : i64 -> !iree_pydm.integer @test_import_global def const_integer(): return 1 # CHECK-LABEL: @const_float -# CHECK: = constant 2.200000e+00 : f64 -> !iree_pydm.real +# CHECK: = arith.constant 2.200000e+00 : f64 -> !iree_pydm.real @test_import_global def const_float(): return 2.2 # CHECK-LABEL: @const_str -# CHECK: = constant "Hello" -> !iree_pydm.str +# CHECK: = arith.constant "Hello" -> !iree_pydm.str @test_import_global def const_str(): return "Hello" # CHECK-LABEL: @const_bytes -# CHECK: = constant "Bonjour" -> !iree_pydm.bytes +# CHECK: = arith.constant "Bonjour" -> !iree_pydm.bytes @test_import_global def const_bytes(): return b"Bonjour" @@ -40,14 +40,14 @@ # CHECK-LABEL: @const_true -# CHECK: = constant true -> !iree_pydm.bool +# CHECK: = arith.constant true -> !iree_pydm.bool @test_import_global def const_true(): return True # CHECK-LABEL: @const_false -# CHECK: = constant false -> !iree_pydm.bool +# CHECK: = arith.constant false -> !iree_pydm.bool @test_import_global def const_false(): return False
diff --git a/llvm-external-projects/iree-dialects/test/python/iree_pydm/importer/intrinsics.py b/llvm-external-projects/iree-dialects/test/python/iree_pydm/importer/intrinsics.py index 8d0030c..cc63a67 100644 --- a/llvm-external-projects/iree-dialects/test/python/iree_pydm/importer/intrinsics.py +++ b/llvm-external-projects/iree-dialects/test/python/iree_pydm/importer/intrinsics.py
@@ -41,8 +41,8 @@ # CHECK-LABEL: @test_intrinsic_function_args -# CHECK: %[[ZERO:.*]] = constant 0 : i64 -> !iree_pydm.integer -# CHECK: %[[ONE:.*]] = constant 1 : i64 -> !iree_pydm.integer +# CHECK: %[[ZERO:.*]] = arith.constant 0 : i64 -> !iree_pydm.integer +# CHECK: %[[ONE:.*]] = arith.constant 1 : i64 -> !iree_pydm.integer # CHECK: dynamic_call @__return_first_true(%[[ZERO]], %[[ONE]]) : (!iree_pydm.integer, !iree_pydm.integer) -> (!iree_pydm.exception_result, !iree_pydm.object) # CHECK: func @__return_first_true @test_import_global @@ -68,7 +68,7 @@ # CHECK-LABEL: @test_intrinsic_macro_no_args -# CHECK: %[[ONE:.*]] = constant 1 +# CHECK: %[[ONE:.*]] = arith.constant 1 # CHECK: box %[[ONE]] : !iree_pydm.integer -> !iree_pydm.object<!iree_pydm.integer> @test_import_global def test_intrinsic_macro_no_args() -> int: @@ -95,7 +95,7 @@ # CHECK-LABEL: @test_pattern_call -# CHECK: %[[TRUE:.*]] = constant true +# CHECK: %[[TRUE:.*]] = arith.constant true # CHECK: pattern_match_call(%[[TRUE]]) : (!iree_pydm.bool) -> (!iree_pydm.exception_result, !iree_pydm.object) # CHECK-SAME: matching generic [@__logical_not_generic] specific [@__logical_not_bool] # CHECK-DAG: func @__logical_not_generic
diff --git a/third_party/llvm-project b/third_party/llvm-project index 57b40b5..6fbc812 160000 --- a/third_party/llvm-project +++ b/third_party/llvm-project
@@ -1 +1 @@ -Subproject commit 57b40b5f34383634949d1639e64a5c2acd0dc5f6 +Subproject commit 6fbc812883ab949d2ce8ee8768f986f2cb39ac35
diff --git a/third_party/mlir-hlo b/third_party/mlir-hlo index a04c82f..d343113 160000 --- a/third_party/mlir-hlo +++ b/third_party/mlir-hlo
@@ -1 +1 @@ -Subproject commit a04c82f16281fc6288edb8e34d9b2f2fad767ab2 +Subproject commit d343113a490691716eac87db61d77b07a7304436
diff --git a/third_party/tensorflow b/third_party/tensorflow index 0a1636b..55e85b4 160000 --- a/third_party/tensorflow +++ b/third_party/tensorflow
@@ -1 +1 @@ -Subproject commit 0a1636bb1667586c9109934625e15814aadadba1 +Subproject commit 55e85b4c921995b6130026a0988e2dec69ad244b