From 0d01d47017556f318b0b9cfd5c4c65f505d541d8 Mon Sep 17 00:00:00 2001 From: Joren Dumoulin Date: Wed, 8 Jan 2025 08:29:00 +0100 Subject: [PATCH] core: print attribute names without double quotes for bare identifiers (#3710) This PR makes sure that attribute names that are bare identifiers (matching `bare_identifier_regex`) are printed without double quotes, as it is done in MLIR. The actual change is quite minimal (in `xdsl/printer.py`), all other diffs come from all of the filechecks and other tests that change. most changes were made automatically with [this script](https://gist.github.com/jorendumoulin/855686adb8197ee207fb3224a480cfd6), using the following conservative regex: `r'(\/\/.*\{( *|[^[}\n]*, *))(\"[a-zA-Z_][a-zA-Z0-9_$.]*\")'` the remaining ones were done using a vim search and replace with manual confirmation `:%s/"\([a-zA-Z_][a-zA-Z0-9_$.]*\)"/\1/gc` --- docs/Toy/Toy_Ch2.ipynb | 14 +- docs/Toy/Toy_Ch3.ipynb | 34 +-- docs/Toy/examples/codegen.toy | 12 +- docs/Toy/examples/scalar.toy | 4 +- docs/Toy/examples/tests/inline.mlir | 18 +- docs/Toy/examples/tests/optimise_toy.mlir | 24 +- docs/database_example.ipynb | 6 +- docs/irdl.ipynb | 8 +- docs/xdsl-introduction.ipynb | 2 +- tests/dialects/test_func.py | 18 +- tests/dialects/test_transform.py | 28 +-- .../backend/riscv/convert_ptr_to_riscv.mlir | 4 +- .../backend/riscv/memref_to_riscv.mlir | 50 ++-- .../backend/riscv/memref_to_riscv_opt.mlir | 12 +- .../register-allocation/preallocated.mlir | 2 +- tests/filecheck/dialects/accfg/accfg_ops.mlir | 26 +- .../filecheck/dialects/affine/affine_ops.mlir | 20 +- tests/filecheck/dialects/affine/examples.mlir | 8 +- tests/filecheck/dialects/air/air_ops.mlir | 38 +-- tests/filecheck/dialects/arith/arith_ops.mlir | 2 +- tests/filecheck/dialects/arm/test_ops.mlir | 8 +- .../dialects/arm/test_registers.mlir | 4 +- .../dialects/arm_func/arm_func_ops.mlir | 6 +- .../bufferization/bufferization_ops.mlir | 10 +- tests/filecheck/dialects/builtin/attrs.mlir | 36 +-- tests/filecheck/dialects/builtin/module.mlir | 4 +- .../builtin/unrealized_conv_cast.mlir | 4 +- tests/filecheck/dialects/cf/cf_ops.mlir | 22 +- .../dialects/csl/csl-canonicalize.mlir | 4 +- .../csl/csl-stencil-canonicalize.mlir | 12 +- .../dialects/csl/csl-stencil-ops.mlir | 136 +++++----- .../dialects/csl/csl-wrapper-ops.mlir | 42 ++-- tests/filecheck/dialects/csl/ops.mlir | 138 +++++----- .../filecheck/dialects/dmp/canonicalize.mlir | 2 +- tests/filecheck/dialects/dmp/ops.mlir | 4 +- tests/filecheck/dialects/eqsat/eqsat_ops.mlir | 4 +- tests/filecheck/dialects/func/func_ops.mlir | 8 +- .../dialects/func/func_ops_generic.mlir | 4 +- tests/filecheck/dialects/gpu/ops.mlir | 60 ++--- tests/filecheck/dialects/hw/hw_instance.mlir | 16 +- tests/filecheck/dialects/hw/hw_module.mlir | 20 +- tests/filecheck/dialects/hw/hw_ops.mlir | 6 +- .../filecheck/dialects/linalg/linalg_ops.mlir | 44 ++-- tests/filecheck/dialects/llvm/arithmetic.mlir | 2 +- tests/filecheck/dialects/llvm/example.mlir | 4 +- tests/filecheck/dialects/llvm/global.mlir | 6 +- tests/filecheck/dialects/llvm/icmp.mlir | 40 +-- tests/filecheck/dialects/llvm/inline_asm.mlir | 4 +- tests/filecheck/dialects/llvm/pointers.mlir | 4 +- .../filecheck/dialects/memref/memref_ops.mlir | 30 +-- .../filecheck/dialects/memref_stream/ops.mlir | 70 +++--- .../dialects/mod_arith/mod_arith.mlir | 4 +- .../filecheck/dialects/mpi/memref_compat.mlir | 8 +- tests/filecheck/dialects/omp/ops.mlir | 44 ++-- tests/filecheck/dialects/onnx/onnx_ops.mlir | 46 ++-- .../dialects/printf/printf_basics.mlir | 2 +- .../dialects/printf/printf_to_llvm.mlir | 8 +- tests/filecheck/dialects/qref/ops.mlir | 2 +- tests/filecheck/dialects/qssa/ops.mlir | 2 +- tests/filecheck/dialects/quantum/attrs.mlir | 14 +- tests/filecheck/dialects/riscv/riscv_ops.mlir | 180 ++++++------- .../dialects/riscv_cf/canonicalize.mlir | 4 +- tests/filecheck/dialects/riscv_cf/ops.mlir | 16 +- .../dialects/riscv_debug/riscv_debug_ops.mlir | 2 +- .../dialects/riscv_func/riscv_func_ops.mlir | 2 +- .../filecheck/dialects/riscv_snitch/ops.mlir | 32 +-- tests/filecheck/dialects/scf/reduce.mlir | 2 +- tests/filecheck/dialects/scf/scf_ops.mlir | 2 +- tests/filecheck/dialects/seq/seq_ops.mlir | 4 +- .../filecheck/dialects/snitch/snitch_ops.mlir | 10 +- .../snitch/snitch_to_riscv_lowering.mlir | 18 +- .../snitch_runtime/snitch_runtime_ops.mlir | 14 +- .../convert_snitch_stream_to_snitch.mlir | 50 ++-- .../filecheck/dialects/snitch_stream/ops.mlir | 8 +- tests/filecheck/dialects/stablehlo/attrs.mlir | 4 +- tests/filecheck/dialects/stablehlo/ops.mlir | 2 +- .../stencil/oec-kernels/fvtp2d_qi.mlir | 18 +- .../dialects/stencil/stencil_ops.mlir | 2 +- tests/filecheck/dialects/stim/attrs.mlir | 2 +- tests/filecheck/dialects/stim/stim_ops.mlir | 22 +- tests/filecheck/dialects/tensor/ops.mlir | 8 +- tests/filecheck/dialects/tosa/ops.mlir | 4 +- .../dialects/transform/transform_types.mlir | 36 +-- .../filecheck/dialects/varith/varith_ops.mlir | 2 +- tests/filecheck/dialects/wasm/ops.mlir | 4 +- tests/filecheck/frontend/dialects/affine.py | 12 +- tests/filecheck/frontend/dialects/func.py | 10 +- tests/filecheck/frontend/dialects/scf.py | 30 +-- tests/filecheck/frontend/passes/desymref.mlir | 68 ++--- .../with-mlir/affinet_set.mlir | 10 +- .../with-mlir/dialects/affine/affine_ops.mlir | 20 +- .../with-mlir/dialects/arith/arith_attrs.mlir | 20 +- .../with-mlir/dialects/arith/arith_cmp.mlir | 6 +- .../dialects/arith/arith_fp_ops.mlir | 4 +- .../builtin/unrealized_conversion_cast.mlir | 4 +- .../with-mlir/dialects/func/func_ops.mlir | 6 +- .../dialects/func/func_ops_generic.mlir | 2 +- .../with-mlir/dialects/gpu/ops.mlir | 68 ++--- .../with-mlir/dialects/llvm/arithmetic.mlir | 2 +- .../with-mlir/dialects/llvm/global.mlir | 6 +- .../with-mlir/dialects/llvm/icmp.mlir | 40 +-- .../with-mlir/dialects/llvm/inline_asm.mlir | 4 +- .../with-mlir/dialects/llvm/llvm_func.mlir | 12 +- .../dialects/llvm/llvm_pointer_ops.mlir | 16 +- .../with-mlir/dialects/memref/matmul.mlir | 14 +- .../memref/memref_ops_mlir_conversion.mlir | 36 +-- .../with-mlir/dialects/omp/ops.mlir | 44 ++-- .../with-mlir/dialects/scf/for_generic.mlir | 10 +- .../scf/for_generic_non_index_iv.mlir | 10 +- .../with-mlir/dialects/scf/parallel.mlir | 8 +- .../dialects/scf/parallel_with_reduce.mlir | 14 +- .../with-mlir/dialects/stencil/ops.mlir | 2 +- .../mlir-conversion/with-mlir/mlir_opt.mlir | 4 +- .../parser-printer/attribute_names.mlir | 52 ++++ .../filecheck/parser-printer/affine_map.mlir | 32 +-- .../filecheck/parser-printer/affine_set.mlir | 12 +- tests/filecheck/parser-printer/aliases.mlir | 10 +- .../parser-printer/attribute_names.mlir | 2 +- .../parser-printer/builtin_attrs.mlir | 58 ++--- .../parser-printer/escaped_characters.mlir | 10 +- .../parser-printer/float_parsing.mlir | 14 +- .../operation_with_properties.mlir | 8 +- tests/filecheck/parser-printer/scope.mlir | 2 +- .../parser-printer/unregistered_dialect.mlir | 10 +- .../apply-pdl/apply_pdl_extra_file.mlir | 2 +- .../apply-pdl/apply_pdl_property_rewrite.mlir | 2 +- .../apply-pdl/apply_pdl_simple.mlir | 2 +- .../transforms/convert-scf-to-openmp.mlir | 42 ++-- .../convert-stencil-to-csl-stencil.mlir | 34 +-- .../convert-stencil-to-ll-mlir.mlir | 62 ++--- .../convert_ml_program_to_memref.mlir | 2 +- .../transforms/convert_onnx_to_linalg.mlir | 6 +- .../transforms/convert_riscv_to_llvm.mlir | 64 ++--- tests/filecheck/transforms/cse.mlir | 10 +- .../csl-stencil-handle-async-flow.mlir | 66 ++--- .../csl-stencil-materialize-stores.mlir | 30 +-- .../csl-stencil-to-csl-wrapper.mlir | 78 +++--- .../transforms/csl-wrapper-hoist-buffers.mlir | 12 +- .../transforms/csl_stencil_bufferize.mlir | 12 +- .../transforms/distribute-stencil.mlir | 8 +- .../filecheck/transforms/eqsat-add-costs.mlir | 34 +-- tests/filecheck/transforms/eqsat-extract.mlir | 2 +- .../transforms/function-constant-pinning.mlir | 20 +- .../function-persist-arg-names.mlir | 6 +- tests/filecheck/transforms/gpu-allocs.mlir | 12 +- .../transforms/gpu-map-parallel-loops.mlir | 10 +- ...-use-donated-arguments-remove-outputs.mlir | 2 +- .../transforms/jax-use-donated-arguments.mlir | 4 +- .../loop-hoist-memref/loop-nest.mlir | 2 +- .../transforms/lower-csl-stencil.mlir | 186 +++++++------- .../transforms/lower-csl-wrapper.mlir | 142 +++++------ tests/filecheck/transforms/memref-to-dsd.mlir | 22 +- .../memref_stream_tile_outer_loops.mlir | 4 +- tests/filecheck/transforms/riscv_cse.mlir | 6 +- .../scf-parallel-loop-tiling-partial.mlir | 38 +-- .../transforms/scf-parallel-loop-tiling.mlir | 26 +- .../transforms/stencil-bufferize.mlir | 16 +- .../transforms/stencil-inlining.mlir | 14 +- .../transforms/stencil-shape-inference.mlir | 6 +- .../stencil-tensorize-z-dimension.mlir | 40 +-- .../filecheck/transforms/stencil-unroll.mlir | 2 +- .../test-add-timers-to-top-level-funcs.mlir | 6 +- .../test_get_all_available_passes.py | 2 +- tests/interactive/test_rewrites.py | 4 +- .../irdl/test_declarative_assembly_format.py | 104 ++++---- .../pattern_rewriter/test_pattern_rewriter.py | 238 +++++++++--------- tests/test_printer.py | 64 ++--- tests/test_rewriter.py | 82 +++--- xdsl/printer.py | 10 +- 169 files changed, 1973 insertions(+), 1917 deletions(-) create mode 100644 tests/filecheck/mlir-conversion/with-mlir/parser-printer/attribute_names.mlir diff --git a/docs/Toy/Toy_Ch2.ipynb b/docs/Toy/Toy_Ch2.ipynb index 2c491bed34..b224082ef0 100644 --- a/docs/Toy/Toy_Ch2.ipynb +++ b/docs/Toy/Toy_Ch2.ipynb @@ -33,15 +33,15 @@ "text": [ "builtin.module {\n", " \"toy.func\"() ({\n", - " %0 = \"toy.constant\"() {\"value\" = dense<[[1.000000e+00, 2.000000e+00, 3.000000e+00], [4.000000e+00, 5.000000e+00, 6.000000e+00]]> : tensor<2x3xf64>} : () -> tensor<2x3xf64>\n", + " %0 = \"toy.constant\"() {value = dense<[[1.000000e+00, 2.000000e+00, 3.000000e+00], [4.000000e+00, 5.000000e+00, 6.000000e+00]]> : tensor<2x3xf64>} : () -> tensor<2x3xf64>\n", " %1 = \"toy.reshape\"(%0) : (tensor<2x3xf64>) -> tensor<2x3xf64>\n", - " %2 = \"toy.constant\"() {\"value\" = dense<[1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00, 5.000000e+00, 6.000000e+00]> : tensor<6xf64>} : () -> tensor<6xf64>\n", + " %2 = \"toy.constant\"() {value = dense<[1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00, 5.000000e+00, 6.000000e+00]> : tensor<6xf64>} : () -> tensor<6xf64>\n", " %3 = \"toy.reshape\"(%2) : (tensor<6xf64>) -> tensor<3x2xf64>\n", " %4 = \"toy.reshape\"(%3) : (tensor<3x2xf64>) -> tensor<2x3xf64>\n", " %5 = \"toy.add\"(%1, %4) : (tensor<2x3xf64>, tensor<2x3xf64>) -> tensor<2x3xf64>\n", " \"toy.print\"(%5) : (tensor<2x3xf64>) -> ()\n", " \"toy.return\"() : () -> ()\n", - " }) {\"sym_name\" = \"main\", \"function_type\" = () -> ()} : () -> ()\n", + " }) {sym_name = \"main\", function_type = () -> ()} : () -> ()\n", "}" ] } @@ -171,7 +171,7 @@ "name": "stdout", "output_type": "stream", "text": [ - "%0 = \"toy.constant\"() {\"value\" = dense<[[1.000000e+00, 2.000000e+00, 3.000000e+00], [4.000000e+00, 5.000000e+00, 6.000000e+00]]> : tensor<2x3xf64>} : () -> tensor<2x3xf64>" + "%0 = \"toy.constant\"() {value = dense<[[1.000000e+00, 2.000000e+00, 3.000000e+00], [4.000000e+00, 5.000000e+00, 6.000000e+00]]> : tensor<2x3xf64>} : () -> tensor<2x3xf64>" ] } ], @@ -313,15 +313,15 @@ "text": [ "builtin.module {\n", " \"toy.func\"() ({\n", - " %0 = \"toy.constant\"() {\"value\" = dense<[[1.000000e+00, 2.000000e+00, 3.000000e+00], [4.000000e+00, 5.000000e+00, 6.000000e+00]]> : tensor<2x3xf64>} : () -> tensor<2x3xf64>\n", + " %0 = \"toy.constant\"() {value = dense<[[1.000000e+00, 2.000000e+00, 3.000000e+00], [4.000000e+00, 5.000000e+00, 6.000000e+00]]> : tensor<2x3xf64>} : () -> tensor<2x3xf64>\n", " %1 = \"toy.reshape\"(%0) : (tensor<2x3xf64>) -> tensor<2x3xf64>\n", - " %2 = \"toy.constant\"() {\"value\" = dense<[1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00, 5.000000e+00, 6.000000e+00]> : tensor<6xf64>} : () -> tensor<6xf64>\n", + " %2 = \"toy.constant\"() {value = dense<[1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00, 5.000000e+00, 6.000000e+00]> : tensor<6xf64>} : () -> tensor<6xf64>\n", " %3 = \"toy.reshape\"(%2) : (tensor<6xf64>) -> tensor<3x2xf64>\n", " %4 = \"toy.reshape\"(%3) : (tensor<3x2xf64>) -> tensor<2x3xf64>\n", " %5 = \"toy.add\"(%1, %4) : (tensor<2x3xf64>, tensor<2x3xf64>) -> tensor<2x3xf64>\n", " \"toy.print\"(%5) : (tensor<2x3xf64>) -> ()\n", " \"toy.return\"() : () -> ()\n", - " }) {\"sym_name\" = \"main\", \"function_type\" = () -> ()} : () -> ()\n", + " }) {sym_name = \"main\", function_type = () -> ()} : () -> ()\n", "}\n" ] } diff --git a/docs/Toy/Toy_Ch3.ipynb b/docs/Toy/Toy_Ch3.ipynb index 3837b1efaf..e59b3ace60 100644 --- a/docs/Toy/Toy_Ch3.ipynb +++ b/docs/Toy/Toy_Ch3.ipynb @@ -30,15 +30,15 @@ "text": [ "builtin.module {\n", " \"toy.func\"() ({\n", - " %0 = \"toy.constant\"() {\"value\" = dense<[[1.000000e+00, 2.000000e+00, 3.000000e+00], [4.000000e+00, 5.000000e+00, 6.000000e+00]]> : tensor<2x3xf64>} : () -> tensor<2x3xf64>\n", + " %0 = \"toy.constant\"() {value = dense<[[1.000000e+00, 2.000000e+00, 3.000000e+00], [4.000000e+00, 5.000000e+00, 6.000000e+00]]> : tensor<2x3xf64>} : () -> tensor<2x3xf64>\n", " %1 = \"toy.reshape\"(%0) : (tensor<2x3xf64>) -> tensor<2x3xf64>\n", - " %2 = \"toy.constant\"() {\"value\" = dense<[1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00, 5.000000e+00, 6.000000e+00]> : tensor<6xf64>} : () -> tensor<6xf64>\n", + " %2 = \"toy.constant\"() {value = dense<[1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00, 5.000000e+00, 6.000000e+00]> : tensor<6xf64>} : () -> tensor<6xf64>\n", " %3 = \"toy.reshape\"(%2) : (tensor<6xf64>) -> tensor<6xf64>\n", " %4 = \"toy.reshape\"(%3) : (tensor<6xf64>) -> tensor<2x3xf64>\n", " %5 = \"toy.add\"(%1, %4) : (tensor<2x3xf64>, tensor<2x3xf64>) -> tensor<2x3xf64>\n", " \"toy.print\"(%5) : (tensor<2x3xf64>) -> ()\n", " \"toy.return\"() : () -> ()\n", - " }) {\"sym_name\" = \"main\", \"function_type\" = () -> ()} : () -> ()\n", + " }) {sym_name = \"main\", function_type = () -> ()} : () -> ()\n", "}\n" ] } @@ -84,15 +84,15 @@ "text": [ "builtin.module {\n", " \"toy.func\"() ({\n", - " %0 = \"toy.constant\"() {\"value\" = dense<[[1.000000e+00, 2.000000e+00, 3.000000e+00], [4.000000e+00, 5.000000e+00, 6.000000e+00]]> : tensor<2x3xf64>} : () -> tensor<2x3xf64>\n", + " %0 = \"toy.constant\"() {value = dense<[[1.000000e+00, 2.000000e+00, 3.000000e+00], [4.000000e+00, 5.000000e+00, 6.000000e+00]]> : tensor<2x3xf64>} : () -> tensor<2x3xf64>\n", " %1 = \"toy.reshape\"(%0) : (tensor<2x3xf64>) -> tensor<2x3xf64>\n", - " %2 = \"toy.constant\"() {\"value\" = dense<[1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00, 5.000000e+00, 6.000000e+00]> : tensor<6xf64>} : () -> tensor<6xf64>\n", + " %2 = \"toy.constant\"() {value = dense<[1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00, 5.000000e+00, 6.000000e+00]> : tensor<6xf64>} : () -> tensor<6xf64>\n", " %3 = \"toy.reshape\"(%2) : (tensor<6xf64>) -> tensor<6xf64>\n", " %4 = \"toy.reshape\"(%2) : (tensor<6xf64>) -> tensor<2x3xf64>\n", " %5 = \"toy.add\"(%1, %4) : (tensor<2x3xf64>, tensor<2x3xf64>) -> tensor<2x3xf64>\n", " \"toy.print\"(%5) : (tensor<2x3xf64>) -> ()\n", " \"toy.return\"() : () -> ()\n", - " }) {\"sym_name\" = \"main\", \"function_type\" = () -> ()} : () -> ()\n", + " }) {sym_name = \"main\", function_type = () -> ()} : () -> ()\n", "}" ] } @@ -162,14 +162,14 @@ "text": [ "builtin.module {\n", " \"toy.func\"() ({\n", - " %0 = \"toy.constant\"() {\"value\" = dense<[[1.000000e+00, 2.000000e+00, 3.000000e+00], [4.000000e+00, 5.000000e+00, 6.000000e+00]]> : tensor<2x3xf64>} : () -> tensor<2x3xf64>\n", + " %0 = \"toy.constant\"() {value = dense<[[1.000000e+00, 2.000000e+00, 3.000000e+00], [4.000000e+00, 5.000000e+00, 6.000000e+00]]> : tensor<2x3xf64>} : () -> tensor<2x3xf64>\n", " %1 = \"toy.reshape\"(%0) : (tensor<2x3xf64>) -> tensor<2x3xf64>\n", - " %2 = \"toy.constant\"() {\"value\" = dense<[1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00, 5.000000e+00, 6.000000e+00]> : tensor<6xf64>} : () -> tensor<6xf64>\n", + " %2 = \"toy.constant\"() {value = dense<[1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00, 5.000000e+00, 6.000000e+00]> : tensor<6xf64>} : () -> tensor<6xf64>\n", " %3 = \"toy.reshape\"(%2) : (tensor<6xf64>) -> tensor<2x3xf64>\n", " %4 = \"toy.add\"(%1, %3) : (tensor<2x3xf64>, tensor<2x3xf64>) -> tensor<2x3xf64>\n", " \"toy.print\"(%4) : (tensor<2x3xf64>) -> ()\n", " \"toy.return\"() : () -> ()\n", - " }) {\"sym_name\" = \"main\", \"function_type\" = () -> ()} : () -> ()\n", + " }) {sym_name = \"main\", function_type = () -> ()} : () -> ()\n", "}" ] } @@ -206,14 +206,14 @@ "text": [ "builtin.module {\n", " \"toy.func\"() ({\n", - " %0 = \"toy.constant\"() {\"value\" = dense<[[1.000000e+00, 2.000000e+00, 3.000000e+00], [4.000000e+00, 5.000000e+00, 6.000000e+00]]> : tensor<2x3xf64>} : () -> tensor<2x3xf64>\n", - " %1 = \"toy.constant\"() {\"value\" = dense<[[1.000000e+00, 2.000000e+00, 3.000000e+00], [4.000000e+00, 5.000000e+00, 6.000000e+00]]> : tensor<2x3xf64>} : () -> tensor<2x3xf64>\n", - " %2 = \"toy.constant\"() {\"value\" = dense<[1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00, 5.000000e+00, 6.000000e+00]> : tensor<6xf64>} : () -> tensor<6xf64>\n", - " %3 = \"toy.constant\"() {\"value\" = dense<[[1.000000e+00, 2.000000e+00, 3.000000e+00], [4.000000e+00, 5.000000e+00, 6.000000e+00]]> : tensor<2x3xf64>} : () -> tensor<2x3xf64>\n", + " %0 = \"toy.constant\"() {value = dense<[[1.000000e+00, 2.000000e+00, 3.000000e+00], [4.000000e+00, 5.000000e+00, 6.000000e+00]]> : tensor<2x3xf64>} : () -> tensor<2x3xf64>\n", + " %1 = \"toy.constant\"() {value = dense<[[1.000000e+00, 2.000000e+00, 3.000000e+00], [4.000000e+00, 5.000000e+00, 6.000000e+00]]> : tensor<2x3xf64>} : () -> tensor<2x3xf64>\n", + " %2 = \"toy.constant\"() {value = dense<[1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00, 5.000000e+00, 6.000000e+00]> : tensor<6xf64>} : () -> tensor<6xf64>\n", + " %3 = \"toy.constant\"() {value = dense<[[1.000000e+00, 2.000000e+00, 3.000000e+00], [4.000000e+00, 5.000000e+00, 6.000000e+00]]> : tensor<2x3xf64>} : () -> tensor<2x3xf64>\n", " %4 = \"toy.add\"(%1, %3) : (tensor<2x3xf64>, tensor<2x3xf64>) -> tensor<2x3xf64>\n", " \"toy.print\"(%4) : (tensor<2x3xf64>) -> ()\n", " \"toy.return\"() : () -> ()\n", - " }) {\"sym_name\" = \"main\", \"function_type\" = () -> ()} : () -> ()\n", + " }) {sym_name = \"main\", function_type = () -> ()} : () -> ()\n", "}" ] } @@ -265,12 +265,12 @@ "text": [ "builtin.module {\n", " \"toy.func\"() ({\n", - " %0 = \"toy.constant\"() {\"value\" = dense<[[1.000000e+00, 2.000000e+00, 3.000000e+00], [4.000000e+00, 5.000000e+00, 6.000000e+00]]> : tensor<2x3xf64>} : () -> tensor<2x3xf64>\n", - " %1 = \"toy.constant\"() {\"value\" = dense<[[1.000000e+00, 2.000000e+00, 3.000000e+00], [4.000000e+00, 5.000000e+00, 6.000000e+00]]> : tensor<2x3xf64>} : () -> tensor<2x3xf64>\n", + " %0 = \"toy.constant\"() {value = dense<[[1.000000e+00, 2.000000e+00, 3.000000e+00], [4.000000e+00, 5.000000e+00, 6.000000e+00]]> : tensor<2x3xf64>} : () -> tensor<2x3xf64>\n", + " %1 = \"toy.constant\"() {value = dense<[[1.000000e+00, 2.000000e+00, 3.000000e+00], [4.000000e+00, 5.000000e+00, 6.000000e+00]]> : tensor<2x3xf64>} : () -> tensor<2x3xf64>\n", " %2 = \"toy.add\"(%0, %1) : (tensor<2x3xf64>, tensor<2x3xf64>) -> tensor<2x3xf64>\n", " \"toy.print\"(%2) : (tensor<2x3xf64>) -> ()\n", " \"toy.return\"() : () -> ()\n", - " }) {\"sym_name\" = \"main\", \"function_type\" = () -> ()} : () -> ()\n", + " }) {sym_name = \"main\", function_type = () -> ()} : () -> ()\n", "}" ] } diff --git a/docs/Toy/examples/codegen.toy b/docs/Toy/examples/codegen.toy index 75e66e80af..2443044fb9 100644 --- a/docs/Toy/examples/codegen.toy +++ b/docs/Toy/examples/codegen.toy @@ -20,16 +20,16 @@ def main() { # CHECK-NEXT: %{{.*}} = "toy.transpose"(%{{.*}}) : (tensor<*xf64>) -> tensor<*xf64> # CHECK-NEXT: %{{.*}} = "toy.mul"(%{{.*}}, %{{.*}}) : (tensor<*xf64>, tensor<*xf64>) -> tensor<*xf64> # CHECK-NEXT: "toy.return"(%{{.*}}) : (tensor<*xf64>) -> () -# CHECK-NEXT: }) {"sym_name" = "multiply_transpose", "function_type" = (tensor<*xf64>, tensor<*xf64>) -> tensor<*xf64>, "sym_visibility" = "private"} : () -> () +# CHECK-NEXT: }) {sym_name = "multiply_transpose", function_type = (tensor<*xf64>, tensor<*xf64>) -> tensor<*xf64>, sym_visibility = "private"} : () -> () # CHECK-NEXT: "toy.func"() ({ -# CHECK-NEXT: %{{.*}} = "toy.constant"() {"value" = +# CHECK-NEXT: %{{.*}} = "toy.constant"() {value = # CHECK-SAME{LITERAL}: dense<[[1.000000e+00, 2.000000e+00, 3.000000e+00], [4.000000e+00, 5.000000e+00, 6.000000e+00]]> : tensor<2x3xf64>} : () -> tensor<2x3xf64> # CHECK-NEXT: %{{.*}} = "toy.reshape"(%{{.*}}) : (tensor<2x3xf64>) -> tensor<2x3xf64> -# CHECK-NEXT: %{{.*}} = "toy.constant"() {"value" = dense<[1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00, 5.000000e+00, 6.000000e+00]> : tensor<6xf64>} : () -> tensor<6xf64> +# CHECK-NEXT: %{{.*}} = "toy.constant"() {value = dense<[1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00, 5.000000e+00, 6.000000e+00]> : tensor<6xf64>} : () -> tensor<6xf64> # CHECK-NEXT: %{{.*}} = "toy.reshape"(%{{.*}}) : (tensor<6xf64>) -> tensor<2x3xf64> -# CHECK-NEXT: %{{.*}} = "toy.generic_call"(%{{.*}}, %{{.*}}) {"callee" = @multiply_transpose} : (tensor<2x3xf64>, tensor<2x3xf64>) -> tensor<*xf64> -# CHECK-NEXT: %{{.*}} = "toy.generic_call"(%{{.*}}, %{{.*}}) {"callee" = @multiply_transpose} : (tensor<2x3xf64>, tensor<2x3xf64>) -> tensor<*xf64> +# CHECK-NEXT: %{{.*}} = "toy.generic_call"(%{{.*}}, %{{.*}}) {callee = @multiply_transpose} : (tensor<2x3xf64>, tensor<2x3xf64>) -> tensor<*xf64> +# CHECK-NEXT: %{{.*}} = "toy.generic_call"(%{{.*}}, %{{.*}}) {callee = @multiply_transpose} : (tensor<2x3xf64>, tensor<2x3xf64>) -> tensor<*xf64> # CHECK-NEXT: "toy.print"(%{{.*}}) : (tensor<*xf64>) -> () # CHECK-NEXT: "toy.return"() : () -> () -# CHECK-NEXT: }) {"sym_name" = "main", "function_type" = () -> ()} : () -> () +# CHECK-NEXT: }) {sym_name = "main", function_type = () -> ()} : () -> () # CHECK-NEXT: } diff --git a/docs/Toy/examples/scalar.toy b/docs/Toy/examples/scalar.toy index bad9886c0d..9741e6db7a 100644 --- a/docs/Toy/examples/scalar.toy +++ b/docs/Toy/examples/scalar.toy @@ -6,8 +6,8 @@ def main() { } # CHECK: "toy.func"() ({ -# CHECK-NEXT: %0 = "toy.constant"() {"value" = dense<5.500000e+00> : tensor} : () -> tensor +# CHECK-NEXT: %0 = "toy.constant"() {value = dense<5.500000e+00> : tensor} : () -> tensor # CHECK-NEXT: %1 = "toy.reshape"(%0) : (tensor) -> tensor<2x2xf64> # CHECK-NEXT: "toy.print"(%1) : (tensor<2x2xf64>) -> () # CHECK-NEXT: "toy.return"() : () -> () -# CHECK-NEXT: }) {"sym_name" = "main", "function_type" = () -> ()} : () -> () +# CHECK-NEXT: }) {sym_name = "main", function_type = () -> ()} : () -> () diff --git a/docs/Toy/examples/tests/inline.mlir b/docs/Toy/examples/tests/inline.mlir index 00714d8df2..7607f581c0 100644 --- a/docs/Toy/examples/tests/inline.mlir +++ b/docs/Toy/examples/tests/inline.mlir @@ -9,21 +9,21 @@ builtin.module { %3 = "toy.transpose"(%1) : (tensor<*xf64>) -> tensor<*xf64> %4 = "toy.mul"(%2, %3) : (tensor<*xf64>, tensor<*xf64>) -> tensor<*xf64> "toy.return"(%4) : (tensor<*xf64>) -> () - }) {"sym_name" = "multiply_transpose", "function_type" = (tensor<*xf64>, tensor<*xf64>) -> tensor<*xf64>, "sym_visibility" = "private"} : () -> () + }) {sym_name = "multiply_transpose", function_type = (tensor<*xf64>, tensor<*xf64>) -> tensor<*xf64>, sym_visibility = "private"} : () -> () "toy.func"() ({ - %5 = "toy.constant"() {"value" = dense<[[1.000000e+00, 2.000000e+00, 3.000000e+00], [4.000000e+00, 5.000000e+00, 6.000000e+00]]> : tensor<2x3xf64>} : () -> tensor<2x3xf64> + %5 = "toy.constant"() {value = dense<[[1.000000e+00, 2.000000e+00, 3.000000e+00], [4.000000e+00, 5.000000e+00, 6.000000e+00]]> : tensor<2x3xf64>} : () -> tensor<2x3xf64> %6 = "toy.reshape"(%5) : (tensor<2x3xf64>) -> tensor<2x3xf64> - %7 = "toy.constant"() {"value" = dense<[1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00, 5.000000e+00, 6.000000e+00]> : tensor<6xf64>} : () -> tensor<6xf64> + %7 = "toy.constant"() {value = dense<[1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00, 5.000000e+00, 6.000000e+00]> : tensor<6xf64>} : () -> tensor<6xf64> %8 = "toy.reshape"(%7) : (tensor<6xf64>) -> tensor<2x3xf64> - %9 = "toy.generic_call"(%6, %8) {"callee" = @multiply_transpose} : (tensor<2x3xf64>, tensor<2x3xf64>) -> tensor<*xf64> - %10 = "toy.generic_call"(%8, %6) {"callee" = @multiply_transpose} : (tensor<2x3xf64>, tensor<2x3xf64>) -> tensor<*xf64> + %9 = "toy.generic_call"(%6, %8) {callee = @multiply_transpose} : (tensor<2x3xf64>, tensor<2x3xf64>) -> tensor<*xf64> + %10 = "toy.generic_call"(%8, %6) {callee = @multiply_transpose} : (tensor<2x3xf64>, tensor<2x3xf64>) -> tensor<*xf64> "toy.print"(%10) : (tensor<*xf64>) -> () "toy.return"() : () -> () - }) {"sym_name" = "main", "function_type" = () -> ()} : () -> () + }) {sym_name = "main", function_type = () -> ()} : () -> () // CHECK-NEXT: "toy.func"() ({ -// CHECK-NEXT{LITERAL}: %0 = "toy.constant"() {"value" = dense<[[1.000000e+00, 2.000000e+00, 3.000000e+00], [4.000000e+00, 5.000000e+00, 6.000000e+00]]> : tensor<2x3xf64>} : () -> tensor<2x3xf64> -// CHECK-NEXT{LITERAL}: %1 = "toy.constant"() {"value" = dense<[[1.000000e+00, 2.000000e+00, 3.000000e+00], [4.000000e+00, 5.000000e+00, 6.000000e+00]]> : tensor<2x3xf64>} : () -> tensor<2x3xf64> +// CHECK-NEXT{LITERAL}: %0 = "toy.constant"() {value = dense<[[1.000000e+00, 2.000000e+00, 3.000000e+00], [4.000000e+00, 5.000000e+00, 6.000000e+00]]> : tensor<2x3xf64>} : () -> tensor<2x3xf64> +// CHECK-NEXT{LITERAL}: %1 = "toy.constant"() {value = dense<[[1.000000e+00, 2.000000e+00, 3.000000e+00], [4.000000e+00, 5.000000e+00, 6.000000e+00]]> : tensor<2x3xf64>} : () -> tensor<2x3xf64> // CHECK-NEXT: %2 = "toy.cast"(%1) : (tensor<2x3xf64>) -> tensor<*xf64> // CHECK-NEXT: %3 = "toy.cast"(%0) : (tensor<2x3xf64>) -> tensor<*xf64> // CHECK-NEXT: %4 = "toy.transpose"(%2) : (tensor<*xf64>) -> tensor<*xf64> @@ -31,7 +31,7 @@ builtin.module { // CHECK-NEXT: %6 = "toy.mul"(%4, %5) : (tensor<*xf64>, tensor<*xf64>) -> tensor<*xf64> // CHECK-NEXT: "toy.print"(%6) : (tensor<*xf64>) -> () // CHECK-NEXT: "toy.return"() : () -> () -// CHECK-NEXT: }) {"sym_name" = "main", "function_type" = () -> ()} : () -> () +// CHECK-NEXT: }) {sym_name = "main", function_type = () -> ()} : () -> () } // CHECK-NEXT: } diff --git a/docs/Toy/examples/tests/optimise_toy.mlir b/docs/Toy/examples/tests/optimise_toy.mlir index fe69804b2d..d770fcc062 100644 --- a/docs/Toy/examples/tests/optimise_toy.mlir +++ b/docs/Toy/examples/tests/optimise_toy.mlir @@ -4,46 +4,46 @@ // CHECK: builtin.module { "toy.func"() ({ - %10 = "toy.constant"() {"value" = dense<[[1.0, 2.0, 3.0], [4.0, 5.0, 6.0]]> : tensor<2x3xf64>} : () -> tensor<2x3xf64> + %10 = "toy.constant"() {value = dense<[[1.0, 2.0, 3.0], [4.0, 5.0, 6.0]]> : tensor<2x3xf64>} : () -> tensor<2x3xf64> %11 = "toy.transpose"(%10) : (tensor<2x3xf64>) -> tensor<3x2xf64> %12 = "toy.transpose"(%11) : (tensor<3x2xf64>) -> tensor<2x3xf64> "toy.return"(%12) : (tensor<2x3xf64>) -> () - }) {"sym_name" = "redundant_transpose", "function_type" = () -> tensor<2x3xf64>} : () -> () + }) {sym_name = "redundant_transpose", function_type = () -> tensor<2x3xf64>} : () -> () // CHECK-NEXT: "toy.func"() ({ -// CHECK-NEXT: %{{.*}} = "toy.constant"() {"value" = +// CHECK-NEXT: %{{.*}} = "toy.constant"() {value = // CHECK-SAME{LITERAL}: dense<[[1.000000e+00, 2.000000e+00, 3.000000e+00], [4.000000e+00, 5.000000e+00, 6.000000e+00]]> : tensor<2x3xf64>} : () // CHECK-NEXT: "toy.return"(%{{.*}}) : (tensor<2x3xf64>) -> () -// CHECK-NEXT: }) {"sym_name" = "redundant_transpose", "function_type" = () -> tensor<2x3xf64>} : () -> () +// CHECK-NEXT: }) {sym_name = "redundant_transpose", function_type = () -> tensor<2x3xf64>} : () -> () "toy.func"() ({ - %30 = "toy.constant"() {"value" = dense<[[1.0, 2.0, 3.0], [4.0, 5.0, 6.0]]> : tensor<2x3xf64>} : () -> tensor<2x3xf64> + %30 = "toy.constant"() {value = dense<[[1.0, 2.0, 3.0], [4.0, 5.0, 6.0]]> : tensor<2x3xf64>} : () -> tensor<2x3xf64> %31 = "toy.reshape"(%30) : (tensor<2x3xf64>) -> tensor<6x1xf64> %32 = "toy.reshape"(%31) : (tensor<6x1xf64>) -> tensor<1x6xf64> %33 = "toy.reshape"(%32) : (tensor<1x6xf64>) -> tensor<2x3xf64> "toy.return"(%33) : (tensor<2x3xf64>) -> () - }) {"sym_name" = "redundant_reshape", "function_type" = () -> tensor<2x3xf64>} : () -> () + }) {sym_name = "redundant_reshape", function_type = () -> tensor<2x3xf64>} : () -> () // CHECK-NEXT: "toy.func"() ({ -// CHECK-NEXT: %{{.*}} = "toy.constant"() {"value" = +// CHECK-NEXT: %{{.*}} = "toy.constant"() {value = // CHECK-SAME{LITERAL}: dense<[[1.000000e+00, 2.000000e+00, 3.000000e+00], [4.000000e+00, 5.000000e+00, 6.000000e+00]]> : tensor<2x3xf64>} : () -> tensor<2x3xf64> // CHECK-NEXT: "toy.return"(%{{.*}}) : (tensor<2x3xf64>) -> () -// CHECK-NEXT: }) {"sym_name" = "redundant_reshape", "function_type" = () -> tensor<2x3xf64>} : () -> () +// CHECK-NEXT: }) {sym_name = "redundant_reshape", function_type = () -> tensor<2x3xf64>} : () -> () "toy.func"() ({ - %30 = "toy.constant"() {"value" = dense<[1.0, 2.0, 3.0, 4.0, 5.0, 6.0]> : tensor<6xf64>} : () -> tensor<6xf64> + %30 = "toy.constant"() {value = dense<[1.0, 2.0, 3.0, 4.0, 5.0, 6.0]> : tensor<6xf64>} : () -> tensor<6xf64> %31 = "toy.reshape"(%30) : (tensor<6xf64>) -> tensor<2x3xf64> "toy.return"(%31) : (tensor<2x3xf64>) -> () - }) {"sym_name" = "constant_reshape", "function_type" = () -> tensor<2x3xf64>} : () -> () + }) {sym_name = "constant_reshape", function_type = () -> tensor<2x3xf64>} : () -> () // CHECK-NEXT: "toy.func"() ({ -// CHECK-NEXT: %{{.*}} = "toy.constant"() {"value" = +// CHECK-NEXT: %{{.*}} = "toy.constant"() {value = // CHECK-SAME{LITERAL}: dense<[[1.000000e+00, 2.000000e+00, 3.000000e+00], [4.000000e+00, 5.000000e+00, 6.000000e+00]]> : tensor<2x3xf64>} : () -> tensor<2x3xf64> // CHECK-NEXT: "toy.return"(%{{.*}}) : (tensor<2x3xf64>) -> () -// CHECK-NEXT: }) {"sym_name" = "constant_reshape", "function_type" = () -> tensor<2x3xf64>} : () -> () +// CHECK-NEXT: }) {sym_name = "constant_reshape", function_type = () -> tensor<2x3xf64>} : () -> () }) : () -> () diff --git a/docs/database_example.ipynb b/docs/database_example.ipynb index 9bc68604cc..6ac268a4d5 100644 --- a/docs/database_example.ipynb +++ b/docs/database_example.ipynb @@ -132,7 +132,7 @@ "name": "stdout", "output_type": "stream", "text": [ - "%0 = \"sql.table\"() {\"table_name\" = \"T\"} : () -> #sql.bag" + "%0 = \"sql.table\"() {table_name = \"T\"} : () -> #sql.bag" ] } ], @@ -164,7 +164,7 @@ "output_type": "stream", "text": [ "builtin.module {\n", - " %0 = \"sql.table\"() {\"table_name\" = \"T\"} : () -> #sql.bag\n", + " %0 = \"sql.table\"() {table_name = \"T\"} : () -> #sql.bag\n", "}\n" ] } @@ -459,7 +459,7 @@ "output_type": "stream", "text": [ "builtin.module {\n", - " %0 = \"sql.table\"() {\"table_name\" = \"T\"} : () -> #sql.bag\n", + " %0 = \"sql.table\"() {table_name = \"T\"} : () -> #sql.bag\n", " \"sql.sink\"(%1) : (#sql.bag) -> ()\n", "}\n" ] diff --git a/docs/irdl.ipynb b/docs/irdl.ipynb index 9902bf977c..5fe233970b 100644 --- a/docs/irdl.ipynb +++ b/docs/irdl.ipynb @@ -1286,7 +1286,7 @@ "name": "stdout", "output_type": "stream", "text": [ - "\"string_attr_op\"() {\"value\" = \"ga\"} : () -> ()" + "\"string_attr_op\"() {value = \"ga\"} : () -> ()" ] } ], @@ -1381,7 +1381,7 @@ "name": "stdout", "output_type": "stream", "text": [ - "\"string_attr_op\"() {\"value\" = \"ga\", \"other_attr\" = #builtin.int<42>} : () -> ()" + "\"string_attr_op\"() {value = \"ga\", other_attr = #builtin.int<42>} : () -> ()" ] } ], @@ -1476,11 +1476,11 @@ "Operation does not verify: result at position 0 does not verify:\n", "attribute i32 expected from variable 'T', but got i64\n", "\n", - "%0 = \"arith.addi\"(%1, %1) <{\"overflowFlags\" = #arith.overflow}> : (i32, i32) -> i64\n", + "%0 = \"arith.addi\"(%1, %1) <{overflowFlags = #arith.overflow}> : (i32, i32) -> i64\n", "\n", "\n", "\n", - "%0 = \"arith.addi\"(%1, %1) <{\"overflowFlags\" = #arith.overflow}> : (i32, i32) -> i64\n", + "%0 = \"arith.addi\"(%1, %1) <{overflowFlags = #arith.overflow}> : (i32, i32) -> i64\n", "\n", "\n" ] diff --git a/docs/xdsl-introduction.ipynb b/docs/xdsl-introduction.ipynb index 279cde5f34..a07ccbf923 100644 --- a/docs/xdsl-introduction.ipynb +++ b/docs/xdsl-introduction.ipynb @@ -584,7 +584,7 @@ "name": "stdout", "output_type": "stream", "text": [ - "%0 = \"sql.select\"() {\"table_name\" = \"T\"} : () -> #sql.bag%1 = \"sql.filter\"(%0) ({\n", + "%0 = \"sql.select\"() {table_name = \"T\"} : () -> #sql.bag%1 = \"sql.filter\"(%0) ({\n", "^0(%2 : i32):\n", " %3 = arith.constant 5 : i32\n", " %4 = arith.constant 5 : i32\n", diff --git a/tests/dialects/test_func.py b/tests/dialects/test_func.py index 755e41b7c3..c0554af374 100644 --- a/tests/dialects/test_func.py +++ b/tests/dialects/test_func.py @@ -213,14 +213,14 @@ def test_call(): expected = """ "builtin.module"() ({ - "func.func"() <{"sym_name" = "func0", "function_type" = (i32, i32) -> i32}> ({ + "func.func"() <{sym_name = "func0", function_type = (i32, i32) -> i32}> ({ ^0(%0 : i32, %1 : i32): - %2 = "arith.addi"(%0, %1) <{"overflowFlags" = #arith.overflow}> : (i32, i32) -> i32 + %2 = "arith.addi"(%0, %1) <{overflowFlags = #arith.overflow}> : (i32, i32) -> i32 "func.return"(%2) : (i32) -> () }) : () -> () - %0 = "arith.constant"() <{"value" = 1 : i32}> : () -> i32 - %1 = "arith.constant"() <{"value" = 2 : i32}> : () -> i32 - %2 = "func.call"(%0, %1) <{"callee" = @func0}> : (i32, i32) -> i32 + %0 = "arith.constant"() <{value = 1 : i32}> : () -> i32 + %1 = "arith.constant"() <{value = 2 : i32}> : () -> i32 + %2 = "func.call"(%0, %1) <{callee = @func0}> : (i32, i32) -> i32 }) : () -> () """ # noqa assert len(call0.operands) == 2 @@ -259,13 +259,13 @@ def test_call_II(): expected = """ "builtin.module"() ({ - "func.func"() <{"sym_name" = "func1", "function_type" = (i32) -> i32}> ({ + "func.func"() <{sym_name = "func1", function_type = (i32) -> i32}> ({ ^0(%0 : i32): - %1 = "arith.addi"(%0, %0) <{"overflowFlags" = #arith.overflow}> : (i32, i32) -> i32 + %1 = "arith.addi"(%0, %0) <{overflowFlags = #arith.overflow}> : (i32, i32) -> i32 "func.return"(%1) : (i32) -> () }) : () -> () - %0 = "arith.constant"() <{"value" = 1 : i32}> : () -> i32 - %1 = "func.call"(%0) <{"callee" = @func1}> : (i32) -> i32 + %0 = "arith.constant"() <{value = 1 : i32}> : () -> i32 + %1 = "func.call"(%0) <{callee = @func1}> : (i32) -> i32 }) : () -> () """ # noqa assert len(call0.operands) == 1 diff --git a/tests/dialects/test_transform.py b/tests/dialects/test_transform.py index ee3fe6302a..b2112ab3be 100644 --- a/tests/dialects/test_transform.py +++ b/tests/dialects/test_transform.py @@ -53,7 +53,7 @@ def test_sequence_init(): extra_bindings=extra_bindings, ), """ - "transform.sequence"() <{"failure_propagation_mode" = 1 : i32, "operandSegmentSizes" = array}> ({ + "transform.sequence"() <{failure_propagation_mode = 1 : i32, operandSegmentSizes = array}> ({ ^0(%0 : !transform.any_value, %1 : !transform.op<"linalg.matmul">): }) : () -> () """, @@ -78,7 +78,7 @@ def test_tileop_init(): dynamic_sizes=[], static_sizes=static_sizes, ), - """%0, %1, %2 = "transform.structured.tile_using_for"(%3) <{"static_sizes" = array}> : (!transform.any_value) -> (!transform.any_op, !transform.any_op, !transform.any_op)""", + """%0, %1, %2 = "transform.structured.tile_using_for"(%3) <{static_sizes = array}> : (!transform.any_value) -> (!transform.any_op, !transform.any_op, !transform.any_op)""", None, ) @@ -88,7 +88,7 @@ def test_get_consumer_of_result(): target = test.TestOp(result_types=[transform.AnyOpType()]).results[0] assert_print_op( transform.GetConsumersOfResultOp(target=target, result_number=result_number), - """%0 = "transform.get_consumers_of_result"(%1) <{"result_number" = 0 : i64}> : (!transform.any_op) -> !transform.any_op""", + """%0 = "transform.get_consumers_of_result"(%1) <{result_number = 0 : i64}> : (!transform.any_op) -> !transform.any_op""", None, ) @@ -106,7 +106,7 @@ def test_get_parent_op(): target = test.TestOp(result_types=[transform.AnyOpType()]).results[0] assert_print_op( transform.GetParentOp(target=target), - """%0 = "transform.get_parent_op"(%1) <{"nth_parent" = 1 : i64}> : (!transform.any_op) -> !transform.any_op""", + """%0 = "transform.get_parent_op"(%1) <{nth_parent = 1 : i64}> : (!transform.any_op) -> !transform.any_op""", None, ) @@ -115,7 +115,7 @@ def test_get_producer_of_operand(): target = test.TestOp(result_types=[transform.AnyValueType()]).results[0] assert_print_op( transform.GetProducerOfOperandOp(operand_number=0, target=target), - """%0 = "transform.get_producer_of_operand"(%1) <{"operand_number" = 0 : i64}> : (!transform.any_value) -> !transform.any_op""", + """%0 = "transform.get_producer_of_operand"(%1) <{operand_number = 0 : i64}> : (!transform.any_value) -> !transform.any_op""", None, ) @@ -125,7 +125,7 @@ def test_get_result(): result_number = 0 assert_print_op( transform.GetResultOp(target=target, raw_position_list=[result_number]), - """%0 = "transform.get_result"(%1) <{"raw_position_list" = array}> : (!transform.any_op) -> !transform.any_value""", + """%0 = "transform.get_result"(%1) <{raw_position_list = array}> : (!transform.any_op) -> !transform.any_value""", None, ) @@ -146,7 +146,7 @@ def test_include(): transform.IncludeOp( target=target, failure_propagation_mode=0, operands_input=operands_input ), - """%0 = "transform.include"(%1) <{"target" = @foo, "failure_propagation_mode" = false}> : (!transform.any_value) -> !transform.any_value""", + """%0 = "transform.include"(%1) <{target = @foo, failure_propagation_mode = false}> : (!transform.any_value) -> !transform.any_value""", None, ) @@ -164,7 +164,7 @@ def test_match_name(): handle = test.TestOp(result_types=[transform.AnyOpType()]).results[0] assert_print_op( transform.MatchOperationNameOp(operand_handle=handle, op_names=["foo"]), - """ "transform.match.operation_name"(%0) <{"op_names" = ["foo"]}> : (!transform.any_op) -> () """, + """ "transform.match.operation_name"(%0) <{op_names = ["foo"]}> : (!transform.any_op) -> () """, None, ) @@ -177,7 +177,7 @@ def test_match_param(): transform.MatchParamCmpIOp( predicate=predicate, param=param, reference=reference ), - """ "transform.match.param.cmpi"(%0, %1) <{"predicate" = 0 : i64}> : (!transform.any_param, !transform.any_param) -> () """, + """ "transform.match.param.cmpi"(%0, %1) <{predicate = 0 : i64}> : (!transform.any_param, !transform.any_param) -> () """, None, ) @@ -186,7 +186,7 @@ def test_merge_handles(): handles = [test.TestOp(result_types=[transform.AnyOpType()]).results[0]] assert_print_op( transform.MergeHandlesOp(handles=handles, deduplicate=True), - """%0 = "transform.merge_handles"(%1) <{"deduplicate"}> : (!transform.any_op) -> !transform.any_op """, + """%0 = "transform.merge_handles"(%1) <{deduplicate}> : (!transform.any_op) -> !transform.any_op """, None, ) @@ -195,7 +195,7 @@ def test_param_const(): value = IntegerAttr(1, IntegerType(32)) assert_print_op( transform.ParamConstantOp(value=value, param_type=IntegerType(32)), - """ %0 = "transform.param.constant"() <{"value" = 1 : i32}> : () -> !transform.param """, + """ %0 = "transform.param.constant"() <{value = 1 : i32}> : () -> !transform.param """, None, ) @@ -210,7 +210,7 @@ def test_split_handle(): fail_on_payload_too_small=True, overflow_result=1, ), - """ %0, %1 = "transform.split_handle"(%2) <{"pass_through_empty_handle" = true, "fail_on_payload_too_small" = true, "overflow_result" = 1 : i64}> : (!transform.any_op) -> (!transform.any_op, !transform.any_op) """, + """ %0, %1 = "transform.split_handle"(%2) <{pass_through_empty_handle = true, fail_on_payload_too_small = true, overflow_result = 1 : i64}> : (!transform.any_op) -> (!transform.any_op, !transform.any_op) """, None, ) @@ -232,7 +232,7 @@ def test_amount_of_loops(): dynamic_sizes=[], static_sizes=static_sizes, ), - """%0, %1 = "transform.structured.tile_using_for"(%2) <{"static_sizes" = array}> : (!transform.any_value) -> (!transform.any_op, !transform.any_op)""", + """%0, %1 = "transform.structured.tile_using_for"(%2) <{static_sizes = array}> : (!transform.any_value) -> (!transform.any_op, !transform.any_op)""", None, ) @@ -245,6 +245,6 @@ def test_structured_match(): ops=[], op_attrs={}, ), - """ %0 = "transform.structured.match"(%1) <{"ops" = [], "op_attrs" = {}}> : (!transform.any_op) -> !transform.any_op """, + """ %0 = "transform.structured.match"(%1) <{ops = [], op_attrs = {}}> : (!transform.any_op) -> !transform.any_op """, None, ) diff --git a/tests/filecheck/backend/riscv/convert_ptr_to_riscv.mlir b/tests/filecheck/backend/riscv/convert_ptr_to_riscv.mlir index 9f1a5553e2..f5f46deebc 100644 --- a/tests/filecheck/backend/riscv/convert_ptr_to_riscv.mlir +++ b/tests/filecheck/backend/riscv/convert_ptr_to_riscv.mlir @@ -11,10 +11,10 @@ ptr_xdsl.store %v, %p : i32, !ptr_xdsl.ptr // CHECK-NEXT: %v_1 = builtin.unrealized_conversion_cast %v : i32 to !riscv.reg -// CHECK-NEXT: riscv.sw %p, %v_1, 0 {"comment" = "store int value to pointer"} : (!riscv.reg, !riscv.reg) -> () +// CHECK-NEXT: riscv.sw %p, %v_1, 0 {comment = "store int value to pointer"} : (!riscv.reg, !riscv.reg) -> () %r3 = ptr_xdsl.load %p : !ptr_xdsl.ptr -> i32 -// CHECK-NEXT: %r3 = riscv.lw %p, 0 {"comment" = "load word from pointer"} : (!riscv.reg) -> !riscv.reg +// CHECK-NEXT: %r3 = riscv.lw %p, 0 {comment = "load word from pointer"} : (!riscv.reg) -> !riscv.reg // CHECK-NEXT: %r3_1 = builtin.unrealized_conversion_cast %r3 : !riscv.reg to i32 // ----- diff --git a/tests/filecheck/backend/riscv/memref_to_riscv.mlir b/tests/filecheck/backend/riscv/memref_to_riscv.mlir index 3631e9947c..ae26dc6285 100644 --- a/tests/filecheck/backend/riscv/memref_to_riscv.mlir +++ b/tests/filecheck/backend/riscv/memref_to_riscv.mlir @@ -17,9 +17,9 @@ // CHECK-NEXT: %pointer_dim_offset = riscv.mul %r_1, %pointer_dim_stride : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-NEXT: %pointer_offset = riscv.add %pointer_dim_offset, %c_1 : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-NEXT: %bytes_per_element = riscv.li 4 : !riscv.reg -// CHECK-NEXT: %scaled_pointer_offset = riscv.mul %pointer_offset, %bytes_per_element {"comment" = "multiply by element size"} : (!riscv.reg, !riscv.reg) -> !riscv.reg +// CHECK-NEXT: %scaled_pointer_offset = riscv.mul %pointer_offset, %bytes_per_element {comment = "multiply by element size"} : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-NEXT: %offset_pointer = riscv.add %m_f32_1, %scaled_pointer_offset : (!riscv.reg, !riscv.reg) -> !riscv.reg -// CHECK-NEXT: riscv.fsw %offset_pointer, %v_f32_1, 0 {"comment" = "store float value to memref of shape (3, 2)"} : (!riscv.reg, !riscv.freg) -> () +// CHECK-NEXT: riscv.fsw %offset_pointer, %v_f32_1, 0 {comment = "store float value to memref of shape (3, 2)"} : (!riscv.reg, !riscv.freg) -> () memref.store %v_f32, %m_f32[%r, %c] {"nontemporal" = false} : memref<3x2xf32> // CHECK-NEXT: %m_f32_2 = builtin.unrealized_conversion_cast %m_f32 : memref<3x2xf32> to !riscv.reg @@ -29,9 +29,9 @@ memref.store %v_f32, %m_f32[%r, %c] {"nontemporal" = false} : memref<3x2xf32> // CHECK-NEXT: %pointer_dim_offset_1 = riscv.mul %r_2, %pointer_dim_stride_1 : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-NEXT: %pointer_offset_1 = riscv.add %pointer_dim_offset_1, %c_2 : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-NEXT: %bytes_per_element_1 = riscv.li 4 : !riscv.reg -// CHECK-NEXT: %scaled_pointer_offset_1 = riscv.mul %pointer_offset_1, %bytes_per_element_1 {"comment" = "multiply by element size"} : (!riscv.reg, !riscv.reg) -> !riscv.reg +// CHECK-NEXT: %scaled_pointer_offset_1 = riscv.mul %pointer_offset_1, %bytes_per_element_1 {comment = "multiply by element size"} : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-NEXT: %offset_pointer_1 = riscv.add %m_f32_2, %scaled_pointer_offset_1 : (!riscv.reg, !riscv.reg) -> !riscv.reg -// CHECK-NEXT: %x_f32 = riscv.flw %offset_pointer_1, 0 {"comment" = "load float from memref of shape (3, 2)"} : (!riscv.reg) -> !riscv.freg +// CHECK-NEXT: %x_f32 = riscv.flw %offset_pointer_1, 0 {comment = "load float from memref of shape (3, 2)"} : (!riscv.reg) -> !riscv.freg // CHECK-NEXT: %x_f32_1 = builtin.unrealized_conversion_cast %x_f32 : !riscv.freg to f32 %x_f32 = memref.load %m_f32[%r, %c] {"nontemporal" = false} : memref<3x2xf32> @@ -39,17 +39,17 @@ memref.store %v_f32, %m_f32[%r, %c] {"nontemporal" = false} : memref<3x2xf32> // CHECK-NEXT: %m_i32_1 = builtin.unrealized_conversion_cast %m_i32 : memref<3xi32> to !riscv.reg // CHECK-NEXT: %c_3 = builtin.unrealized_conversion_cast %c : index to !riscv.reg // CHECK-NEXT: %bytes_per_element_2 = riscv.li 4 : !riscv.reg -// CHECK-NEXT: %scaled_pointer_offset_2 = riscv.mul %c_3, %bytes_per_element_2 {"comment" = "multiply by element size"} : (!riscv.reg, !riscv.reg) -> !riscv.reg +// CHECK-NEXT: %scaled_pointer_offset_2 = riscv.mul %c_3, %bytes_per_element_2 {comment = "multiply by element size"} : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-NEXT: %offset_pointer_2 = riscv.add %m_i32_1, %scaled_pointer_offset_2 : (!riscv.reg, !riscv.reg) -> !riscv.reg -// CHECK-NEXT: riscv.sw %offset_pointer_2, %v_i32_1, 0 {"comment" = "store int value to memref of shape (3,)"} : (!riscv.reg, !riscv.reg) -> () +// CHECK-NEXT: riscv.sw %offset_pointer_2, %v_i32_1, 0 {comment = "store int value to memref of shape (3,)"} : (!riscv.reg, !riscv.reg) -> () memref.store %v_i32, %m_i32[%c] {"nontemporal" = false} : memref<3xi32> // CHECK-NEXT: %m_i32_2 = builtin.unrealized_conversion_cast %m_i32 : memref<3xi32> to !riscv.reg // CHECK-NEXT: %c_4 = builtin.unrealized_conversion_cast %c : index to !riscv.reg // CHECK-NEXT: %bytes_per_element_3 = riscv.li 4 : !riscv.reg -// CHECK-NEXT: %scaled_pointer_offset_3 = riscv.mul %c_4, %bytes_per_element_3 {"comment" = "multiply by element size"} : (!riscv.reg, !riscv.reg) -> !riscv.reg +// CHECK-NEXT: %scaled_pointer_offset_3 = riscv.mul %c_4, %bytes_per_element_3 {comment = "multiply by element size"} : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-NEXT: %offset_pointer_3 = riscv.add %m_i32_2, %scaled_pointer_offset_3 : (!riscv.reg, !riscv.reg) -> !riscv.reg -// CHECK-NEXT: %x_i32 = riscv.lw %offset_pointer_3, 0 {"comment" = "load word from memref of shape (3,)"} : (!riscv.reg) -> !riscv.reg +// CHECK-NEXT: %x_i32 = riscv.lw %offset_pointer_3, 0 {comment = "load word from memref of shape (3,)"} : (!riscv.reg) -> !riscv.reg // CHECK-NEXT: %x_i32_1 = builtin.unrealized_conversion_cast %x_i32 : !riscv.reg to i32 %x_i32 = memref.load %m_i32[%c] {"nontemporal" = false} : memref<3xi32> @@ -61,19 +61,19 @@ memref.store %v_i32, %m_i32[%c] {"nontemporal" = false} : memref<3xi32> // CHECK-NEXT: %pointer_dim_offset_2 = riscv.mul %r_3, %pointer_dim_stride_2 : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-NEXT: %pointer_offset_2 = riscv.add %pointer_dim_offset_2, %c_5 : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-NEXT: %bytes_per_element_4 = riscv.li 8 : !riscv.reg -// CHECK-NEXT: %scaled_pointer_offset_4 = riscv.mul %pointer_offset_2, %bytes_per_element_4 {"comment" = "multiply by element size"} : (!riscv.reg, !riscv.reg) -> !riscv.reg +// CHECK-NEXT: %scaled_pointer_offset_4 = riscv.mul %pointer_offset_2, %bytes_per_element_4 {comment = "multiply by element size"} : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-NEXT: %offset_pointer_4 = riscv.add %m_f64_1, %scaled_pointer_offset_4 : (!riscv.reg, !riscv.reg) -> !riscv.reg -// CHECK-NEXT: riscv.fsd %offset_pointer_4, %v_f64_1, 0 {"comment" = "store double value to memref of shape (3, 2)"} : (!riscv.reg, !riscv.freg) -> () +// CHECK-NEXT: riscv.fsd %offset_pointer_4, %v_f64_1, 0 {comment = "store double value to memref of shape (3, 2)"} : (!riscv.reg, !riscv.freg) -> () memref.store %v_f64, %m_f64[%r, %c] {"nontemporal" = false} : memref<3x2xf64> // CHECK-NEXT: %m_scalar_i32_1 = builtin.unrealized_conversion_cast %m_scalar_i32 : memref to !riscv.reg -// CHECK-NEXT: %scalar_x_i32 = riscv.lw %m_scalar_i32_1, 0 {"comment" = "load word from memref of shape ()"} : (!riscv.reg) -> !riscv.reg +// CHECK-NEXT: %scalar_x_i32 = riscv.lw %m_scalar_i32_1, 0 {comment = "load word from memref of shape ()"} : (!riscv.reg) -> !riscv.reg // CHECK-NEXT: %scalar_x_i32_1 = builtin.unrealized_conversion_cast %scalar_x_i32 : !riscv.reg to i32 %scalar_x_i32 = memref.load %m_scalar_i32[] {"nontemporal" = false} : memref // CHECK-NEXT: %scalar_x_i32_2 = builtin.unrealized_conversion_cast %scalar_x_i32_1 : i32 to !riscv.reg // CHECK-NEXT: %m_scalar_i32_2 = builtin.unrealized_conversion_cast %m_scalar_i32 : memref to !riscv.reg -// CHECK-NEXT: riscv.sw %m_scalar_i32_2, %scalar_x_i32_2, 0 {"comment" = "store int value to memref of shape ()"} : (!riscv.reg, !riscv.reg) -> () +// CHECK-NEXT: riscv.sw %m_scalar_i32_2, %scalar_x_i32_2, 0 {comment = "store int value to memref of shape ()"} : (!riscv.reg, !riscv.reg) -> () memref.store %scalar_x_i32, %m_scalar_i32[] {"nontemporal" = false} : memref // CHECK-NEXT: %m_f64_2 = builtin.unrealized_conversion_cast %m_f64 : memref<3x2xf64> to !riscv.reg @@ -83,9 +83,9 @@ memref.store %scalar_x_i32, %m_scalar_i32[] {"nontemporal" = false} : memref !riscv.reg // CHECK-NEXT: %pointer_offset_3 = riscv.add %pointer_dim_offset_3, %c_6 : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-NEXT: %bytes_per_element_5 = riscv.li 8 : !riscv.reg -// CHECK-NEXT: %scaled_pointer_offset_5 = riscv.mul %pointer_offset_3, %bytes_per_element_5 {"comment" = "multiply by element size"} : (!riscv.reg, !riscv.reg) -> !riscv.reg +// CHECK-NEXT: %scaled_pointer_offset_5 = riscv.mul %pointer_offset_3, %bytes_per_element_5 {comment = "multiply by element size"} : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-NEXT: %offset_pointer_5 = riscv.add %m_f64_2, %scaled_pointer_offset_5 : (!riscv.reg, !riscv.reg) -> !riscv.reg -// CHECK-NEXT: %x_f64 = riscv.fld %offset_pointer_5, 0 {"comment" = "load double from memref of shape (3, 2)"} : (!riscv.reg) -> !riscv.freg +// CHECK-NEXT: %x_f64 = riscv.fld %offset_pointer_5, 0 {comment = "load double from memref of shape (3, 2)"} : (!riscv.reg) -> !riscv.freg // CHECK-NEXT: %x_f64_1 = builtin.unrealized_conversion_cast %x_f64 : !riscv.freg to f64 %x_f64 = memref.load %m_f64[%r, %c] {"nontemporal" = false} : memref<3x2xf64> @@ -105,14 +105,14 @@ memref.store %scalar_x_i32, %m_scalar_i32[] {"nontemporal" = false} : memref !riscv.reg // CHECK-NEXT: %m0_2 = riscv_func.call @malloc(%m0_1) : (!riscv.reg) -> !riscv.reg // CHECK-NEXT: %m0_3 = riscv.mv %m0_2 : (!riscv.reg) -> !riscv.reg // CHECK-NEXT: %m0_4 = builtin.unrealized_conversion_cast %m0_3 : !riscv.reg to memref<1x1xf32> %m0 = memref.alloc() : memref<1x1xf32> -// CHECK-NEXT: %m1 = riscv.li 8 {"comment" = "memref alloc size"} : !riscv.reg +// CHECK-NEXT: %m1 = riscv.li 8 {comment = "memref alloc size"} : !riscv.reg // CHECK-NEXT: %m1_1 = riscv.mv %m1 : (!riscv.reg) -> !riscv.reg // CHECK-NEXT: %m1_2 = riscv_func.call @malloc(%m1_1) : (!riscv.reg) -> !riscv.reg // CHECK-NEXT: %m1_3 = riscv.mv %m1_2 : (!riscv.reg) -> !riscv.reg @@ -151,9 +151,9 @@ memref.store %scalar_x_i32, %m_scalar_i32[] {"nontemporal" = false} : memref to !riscv.reg // CHECK-NEXT: %d0_1 = builtin.unrealized_conversion_cast %d0 : index to !riscv.reg // CHECK-NEXT: %bytes_per_element = riscv.li 1 : !riscv.reg -// CHECK-NEXT: %scaled_pointer_offset = riscv.mul %d0_1, %bytes_per_element {"comment" = "multiply by element size"} : (!riscv.reg, !riscv.reg) -> !riscv.reg +// CHECK-NEXT: %scaled_pointer_offset = riscv.mul %d0_1, %bytes_per_element {comment = "multiply by element size"} : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-NEXT: %offset_pointer = riscv.add %m_1, %scaled_pointer_offset : (!riscv.reg, !riscv.reg) -> !riscv.reg -// CHECK-NEXT: riscv.sw %offset_pointer, %v_1, 0 {"comment" = "store int value to memref of shape (1,)"} : (!riscv.reg, !riscv.reg) -> () +// CHECK-NEXT: riscv.sw %offset_pointer, %v_1, 0 {comment = "store int value to memref of shape (1,)"} : (!riscv.reg, !riscv.reg) -> () memref.store %v, %m[%d0] {"nontemporal" = false} : memref<1xi8> // CHECK-NEXT: } @@ -169,9 +169,9 @@ memref.store %v, %m[%d0] {"nontemporal" = false} : memref<1xi8> // CHECK-NEXT: %m_1 = builtin.unrealized_conversion_cast %m : memref<1xi16> to !riscv.reg // CHECK-NEXT: %d0_1 = builtin.unrealized_conversion_cast %d0 : index to !riscv.reg // CHECK-NEXT: %bytes_per_element = riscv.li 2 : !riscv.reg -// CHECK-NEXT: %scaled_pointer_offset = riscv.mul %d0_1, %bytes_per_element {"comment" = "multiply by element size"} : (!riscv.reg, !riscv.reg) -> !riscv.reg +// CHECK-NEXT: %scaled_pointer_offset = riscv.mul %d0_1, %bytes_per_element {comment = "multiply by element size"} : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-NEXT: %offset_pointer = riscv.add %m_1, %scaled_pointer_offset : (!riscv.reg, !riscv.reg) -> !riscv.reg -// CHECK-NEXT: riscv.sw %offset_pointer, %v_1, 0 {"comment" = "store int value to memref of shape (1,)"} : (!riscv.reg, !riscv.reg) -> () +// CHECK-NEXT: riscv.sw %offset_pointer, %v_1, 0 {comment = "store int value to memref of shape (1,)"} : (!riscv.reg, !riscv.reg) -> () memref.store %v, %m[%d0] {"nontemporal" = false} : memref<1xi16> // CHECK-NEXT: } @@ -187,9 +187,9 @@ memref.store %v, %m[%d0] {"nontemporal" = false} : memref<1xi16> // CHECK-NEXT: %m_1 = builtin.unrealized_conversion_cast %m : memref<1xi64> to !riscv.reg // CHECK-NEXT: %d0_1 = builtin.unrealized_conversion_cast %d0 : index to !riscv.reg // CHECK-NEXT: %bytes_per_element = riscv.li 8 : !riscv.reg -// CHECK-NEXT: %scaled_pointer_offset = riscv.mul %d0_1, %bytes_per_element {"comment" = "multiply by element size"} : (!riscv.reg, !riscv.reg) -> !riscv.reg +// CHECK-NEXT: %scaled_pointer_offset = riscv.mul %d0_1, %bytes_per_element {comment = "multiply by element size"} : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-NEXT: %offset_pointer = riscv.add %m_1, %scaled_pointer_offset : (!riscv.reg, !riscv.reg) -> !riscv.reg -// CHECK-NEXT: riscv.sw %offset_pointer, %v_1, 0 {"comment" = "store int value to memref of shape (1,)"} : (!riscv.reg, !riscv.reg) -> () +// CHECK-NEXT: riscv.sw %offset_pointer, %v_1, 0 {comment = "store int value to memref of shape (1,)"} : (!riscv.reg, !riscv.reg) -> () memref.store %v, %m[%d0] {"nontemporal" = false} : memref<1xi64> // CHECK-NEXT: } @@ -209,7 +209,7 @@ memref.store %v, %m[%d0] {"nontemporal" = false} : memref<1xi64> %zero_subview = memref.subview %original[0, 0, 0][1, 3, 2][1, 1, 1] : memref<4x3x2xf64> to memref<3x2xf64> // CHECK-NEXT: %static_subview = builtin.unrealized_conversion_cast %original : memref<4x3x2xf64> to !riscv.reg -// CHECK-NEXT: %static_subview_1 = riscv.addi %static_subview, 48 {"comment" = "subview offset"} : (!riscv.reg) -> !riscv.reg +// CHECK-NEXT: %static_subview_1 = riscv.addi %static_subview, 48 {comment = "subview offset"} : (!riscv.reg) -> !riscv.reg // CHECK-NEXT: %static_subview_2 = builtin.unrealized_conversion_cast %static_subview_1 : !riscv.reg to memref<3x2xf64, strided<[2, 1], offset: 6>> %static_subview = memref.subview %original[1, 0, 0][1, 3, 2][1, 1, 1] : memref<4x3x2xf64> to memref<3x2xf64, strided<[2, 1], offset: 6>> @@ -225,7 +225,7 @@ memref.store %v, %m[%d0] {"nontemporal" = false} : memref<1xi64> // CHECK-NEXT: %pointer_offset = riscv.add %pointer_dim_offset, %pointer_dim_offset_1 : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-NEXT: %pointer_offset_1 = riscv.add %pointer_offset, %subview_dim_index_2 : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-NEXT: %bytes_per_element = riscv.li 8 : !riscv.reg -// CHECK-NEXT: %scaled_pointer_offset = riscv.mul %pointer_offset_1, %bytes_per_element {"comment" = "multiply by element size"} : (!riscv.reg, !riscv.reg) -> !riscv.reg +// CHECK-NEXT: %scaled_pointer_offset = riscv.mul %pointer_offset_1, %bytes_per_element {comment = "multiply by element size"} : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-NEXT: %offset_pointer = riscv.add %dynamic_subview, %scaled_pointer_offset : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-NEXT: %dynamic_subview_1 = builtin.unrealized_conversion_cast %offset_pointer : !riscv.reg to memref<3x2xf64, strided<[2, 1], offset: ?>> %dynamic_subview = memref.subview %original[%offset, 0, 0][1, 3, 2][1, 1, 1] : @@ -248,7 +248,7 @@ memref.store %v, %m[%d0] {"nontemporal" = false} : memref<1xi64> // CHECK-NEXT: %pointer_offset_3 = riscv.add %pointer_offset_2, %pointer_dim_offset_4 : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-NEXT: %pointer_offset_4 = riscv.add %pointer_offset_3, %subview_dim_index_6 : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-NEXT: %bytes_per_element_1 = riscv.li 8 : !riscv.reg -// CHECK-NEXT: %scaled_pointer_offset_1 = riscv.mul %pointer_offset_4, %bytes_per_element_1 {"comment" = "multiply by element size"} : (!riscv.reg, !riscv.reg) -> !riscv.reg +// CHECK-NEXT: %scaled_pointer_offset_1 = riscv.mul %pointer_offset_4, %bytes_per_element_1 {comment = "multiply by element size"} : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-NEXT: %offset_pointer_1 = riscv.add %larger_dynamic_subview, %scaled_pointer_offset_1 : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-NEXT: %larger_dynamic_subview_1 = builtin.unrealized_conversion_cast %offset_pointer_1 : !riscv.reg to memref<3x2xf64, strided<[2, 1], offset: ?>> %larger_dynamic_subview = memref.subview %larger_original[%offset, %offset, 0, 0][1, 1, 3, 2][1, 1, 1, 1] : diff --git a/tests/filecheck/backend/riscv/memref_to_riscv_opt.mlir b/tests/filecheck/backend/riscv/memref_to_riscv_opt.mlir index c520189128..41519937a9 100644 --- a/tests/filecheck/backend/riscv/memref_to_riscv_opt.mlir +++ b/tests/filecheck/backend/riscv/memref_to_riscv_opt.mlir @@ -30,12 +30,12 @@ builtin.module { // CHECK: builtin.module { // CHECK-NEXT: %vf_reg, %vd_reg, %vi_reg, %mf_reg, %md_reg, %mi_reg = "test.op"() : () -> (!riscv.freg, !riscv.freg, !riscv.reg, !riscv.reg, !riscv.reg, !riscv.reg) -// CHECK-NEXT: riscv.fsw %mf_reg, %vf_reg, 12 {"comment" = "store float value to memref of shape (3, 2)"} : (!riscv.reg, !riscv.freg) -> () -// CHECK-NEXT: %xf = riscv.flw %mf_reg, 12 {"comment" = "load float from memref of shape (3, 2)"} : (!riscv.reg) -> !riscv.freg -// CHECK-NEXT: riscv.fsd %md_reg, %vd_reg, 24 {"comment" = "store double value to memref of shape (3, 2)"} : (!riscv.reg, !riscv.freg) -> () -// CHECK-NEXT: %xd = riscv.fld %md_reg, 24 {"comment" = "load double from memref of shape (3, 2)"} : (!riscv.reg) -> !riscv.freg -// CHECK-NEXT: riscv.sw %mi_reg, %vi_reg, 12 {"comment" = "store int value to memref of shape (3, 2)"} : (!riscv.reg, !riscv.reg) -> () -// CHECK-NEXT: %xi = riscv.lw %mi_reg, 12 {"comment" = "load word from memref of shape (3, 2)"} : (!riscv.reg) -> !riscv.reg +// CHECK-NEXT: riscv.fsw %mf_reg, %vf_reg, 12 {comment = "store float value to memref of shape (3, 2)"} : (!riscv.reg, !riscv.freg) -> () +// CHECK-NEXT: %xf = riscv.flw %mf_reg, 12 {comment = "load float from memref of shape (3, 2)"} : (!riscv.reg) -> !riscv.freg +// CHECK-NEXT: riscv.fsd %md_reg, %vd_reg, 24 {comment = "store double value to memref of shape (3, 2)"} : (!riscv.reg, !riscv.freg) -> () +// CHECK-NEXT: %xd = riscv.fld %md_reg, 24 {comment = "load double from memref of shape (3, 2)"} : (!riscv.reg) -> !riscv.freg +// CHECK-NEXT: riscv.sw %mi_reg, %vi_reg, 12 {comment = "store int value to memref of shape (3, 2)"} : (!riscv.reg, !riscv.reg) -> () +// CHECK-NEXT: %xi = riscv.lw %mi_reg, 12 {comment = "load word from memref of shape (3, 2)"} : (!riscv.reg) -> !riscv.reg // CHECK-NEXT: "test.op"(%xf) : (!riscv.freg) -> () // CHECK-NEXT: "test.op"(%xd) : (!riscv.freg) -> () // CHECK-NEXT: "test.op"(%xi) : (!riscv.reg) -> () diff --git a/tests/filecheck/backend/riscv/register-allocation/preallocated.mlir b/tests/filecheck/backend/riscv/register-allocation/preallocated.mlir index 117b5f1b29..18395df95d 100644 --- a/tests/filecheck/backend/riscv/register-allocation/preallocated.mlir +++ b/tests/filecheck/backend/riscv/register-allocation/preallocated.mlir @@ -11,7 +11,7 @@ riscv_func.func @main() { } // LIVE-BNAIVE: builtin.module { -// LIVE-BNAIVE-NEXT: riscv.comment {"comment" = "Regalloc stats: {\"preallocated_float\": [\"ft0\"], \"preallocated_int\": [\"t0\"], \"allocated_float\": [\"ft0\", \"ft1\"], \"allocated_int\": [\"t0\", \"t1\"]}"} : () -> () +// LIVE-BNAIVE-NEXT: riscv.comment {comment = "Regalloc stats: {\"preallocated_float\": [\"ft0\"], \"preallocated_int\": [\"t0\"], \"allocated_float\": [\"ft0\", \"ft1\"], \"allocated_int\": [\"t0\", \"t1\"]}"} : () -> () // LIVE-BNAIVE-NEXT: riscv_func.func @main() { // LIVE-BNAIVE-NEXT: %0 = riscv.li 6 : !riscv.reg // LIVE-BNAIVE-NEXT: %1 = riscv.li 5 : !riscv.reg diff --git a/tests/filecheck/dialects/accfg/accfg_ops.mlir b/tests/filecheck/dialects/accfg/accfg_ops.mlir index fd4d4f5274..26a8e87940 100644 --- a/tests/filecheck/dialects/accfg/accfg_ops.mlir +++ b/tests/filecheck/dialects/accfg/accfg_ops.mlir @@ -39,32 +39,32 @@ func.func @test() { // CHECK-NEXT: builtin.module { -// CHECK-NEXT: "accfg.accelerator"() <{"name" = @acc1, "fields" = {"A" = 960 : i64, "B" = 961 : i64}, "launch_fields" = {"launch" = 975 : i64}, "barrier" = 1987 : i64}> : () -> () +// CHECK-NEXT: "accfg.accelerator"() <{name = @acc1, fields = {A = 960 : i64, B = 961 : i64}, launch_fields = {launch = 975 : i64}, barrier = 1987 : i64}> : () -> () // CHECK-NEXT: func.func @test() { // CHECK-NEXT: %one, %two = "test.op"() : () -> (i32, i32) // CHECK-NEXT: %zero = arith.constant 0 : i32 -// CHECK-NEXT: %state = accfg.setup "acc1" to ("A" = %one : i32, "B" = %two : i32) attrs {"test_attr" = 100 : i64} : !accfg.state<"acc1"> -// CHECK-NEXT: %token = "accfg.launch"(%zero, %state) <{"param_names" = ["launch"], "accelerator" = "acc1"}> : (i32, !accfg.state<"acc1">) -> !accfg.token<"acc1"> +// CHECK-NEXT: %state = accfg.setup "acc1" to ("A" = %one : i32, "B" = %two : i32) attrs {test_attr = 100 : i64} : !accfg.state<"acc1"> +// CHECK-NEXT: %token = "accfg.launch"(%zero, %state) <{param_names = ["launch"], accelerator = "acc1"}> : (i32, !accfg.state<"acc1">) -> !accfg.token<"acc1"> // CHECK-NEXT: %state2 = accfg.setup "acc1" from %state to ("A" = %one : i32, "B" = %two : i32) : !accfg.state<"acc1"> // CHECK-NEXT: "accfg.await"(%token) : (!accfg.token<"acc1">) -> () -// CHECK-NEXT: "test.op"() {"accfg.effects" = #accfg.effects} : () -> () -// CHECK-NEXT: "test.op"() {"accfg.effects" = #accfg.effects} : () -> () +// CHECK-NEXT: "test.op"() {accfg.effects = #accfg.effects} : () -> () +// CHECK-NEXT: "test.op"() {accfg.effects = #accfg.effects} : () -> () // CHECK-NEXT: func.return // CHECK-NEXT: } // CHECK-NEXT: } // CHECK-GENERIC-NEXT: "builtin.module"() ({ -// CHECK-GENERIC-NEXT: "accfg.accelerator"() <{"name" = @acc1, "fields" = {"A" = 960 : i64, "B" = 961 : i64}, "launch_fields" = {"launch" = 975 : i64}, "barrier" = 1987 : i64}> : () -> () -// CHECK-GENERIC-NEXT: "func.func"() <{"sym_name" = "test", "function_type" = () -> ()}> ({ +// CHECK-GENERIC-NEXT: "accfg.accelerator"() <{name = @acc1, fields = {A = 960 : i64, B = 961 : i64}, launch_fields = {launch = 975 : i64}, barrier = 1987 : i64}> : () -> () +// CHECK-GENERIC-NEXT: "func.func"() <{sym_name = "test", function_type = () -> ()}> ({ // CHECK-GENERIC-NEXT: %one, %two = "test.op"() : () -> (i32, i32) -// CHECK-GENERIC-NEXT: %zero = "arith.constant"() <{"value" = 0 : i32}> : () -> i32 -// CHECK-GENERIC-NEXT: %state = "accfg.setup"(%one, %two) <{"param_names" = ["A", "B"], "accelerator" = "acc1", "operandSegmentSizes" = array}> {"test_attr" = 100 : i64} : (i32, i32) -> !accfg.state<"acc1"> -// CHECK-GENERIC-NEXT: %token = "accfg.launch"(%zero, %state) <{"param_names" = ["launch"], "accelerator" = "acc1"}> : (i32, !accfg.state<"acc1">) -> !accfg.token<"acc1"> -// CHECK-GENERIC-NEXT: %state2 = "accfg.setup"(%one, %two, %state) <{"param_names" = ["A", "B"], "accelerator" = "acc1", "operandSegmentSizes" = array}> : (i32, i32, !accfg.state<"acc1">) -> !accfg.state<"acc1"> +// CHECK-GENERIC-NEXT: %zero = "arith.constant"() <{value = 0 : i32}> : () -> i32 +// CHECK-GENERIC-NEXT: %state = "accfg.setup"(%one, %two) <{param_names = ["A", "B"], accelerator = "acc1", operandSegmentSizes = array}> {test_attr = 100 : i64} : (i32, i32) -> !accfg.state<"acc1"> +// CHECK-GENERIC-NEXT: %token = "accfg.launch"(%zero, %state) <{param_names = ["launch"], accelerator = "acc1"}> : (i32, !accfg.state<"acc1">) -> !accfg.token<"acc1"> +// CHECK-GENERIC-NEXT: %state2 = "accfg.setup"(%one, %two, %state) <{param_names = ["A", "B"], accelerator = "acc1", operandSegmentSizes = array}> : (i32, i32, !accfg.state<"acc1">) -> !accfg.state<"acc1"> // CHECK-GENERIC-NEXT: "accfg.await"(%token) : (!accfg.token<"acc1">) -> () -// CHECK-GENERIC-NEXT: "test.op"() {"accfg.effects" = #accfg.effects} : () -> () -// CHECK-GENERIC-NEXT: "test.op"() {"accfg.effects" = #accfg.effects} : () -> () +// CHECK-GENERIC-NEXT: "test.op"() {accfg.effects = #accfg.effects} : () -> () +// CHECK-GENERIC-NEXT: "test.op"() {accfg.effects = #accfg.effects} : () -> () // CHECK-GENERIC-NEXT: "func.return"() : () -> () // CHECK-GENERIC-NEXT: }) : () -> () // CHECK-GENERIC-NEXT: }) : () -> () diff --git a/tests/filecheck/dialects/affine/affine_ops.mlir b/tests/filecheck/dialects/affine/affine_ops.mlir index fe818643ae..85ce8e29b8 100644 --- a/tests/filecheck/dialects/affine/affine_ops.mlir +++ b/tests/filecheck/dialects/affine/affine_ops.mlir @@ -9,7 +9,7 @@ "affine.yield"() : () -> () }) : () -> () - // CHECK: "affine.for"() <{"lowerBoundMap" = affine_map<() -> (0)>, "upperBoundMap" = affine_map<() -> (256)>, "step" = 1 : index, "operandSegmentSizes" = array}> ({ + // CHECK: "affine.for"() <{lowerBoundMap = affine_map<() -> (0)>, upperBoundMap = affine_map<() -> (256)>, step = 1 : index, operandSegmentSizes = array}> ({ // CHECK-NEXT: ^0(%{{.*}} : index): // CHECK-NEXT: "affine.yield"() : () -> () // CHECK-NEXT: }) : () -> () @@ -35,12 +35,12 @@ "affine.yield"() : () -> () }) : (index) -> () - // CHECK: %res = "affine.for"(%{{.*}}) <{"lowerBoundMap" = affine_map<() -> (-10)>, "upperBoundMap" = affine_map<() -> (10)>, "step" = 1 : index, "operandSegmentSizes" = array}> ({ + // CHECK: %res = "affine.for"(%{{.*}}) <{lowerBoundMap = affine_map<() -> (-10)>, upperBoundMap = affine_map<() -> (10)>, step = 1 : index, operandSegmentSizes = array}> ({ // CHECK-NEXT: ^1(%{{.*}} : index, %{{.*}} : !test.type<"int">): // CHECK-NEXT: %{{.*}} = "test.op"() : () -> !test.type<"int"> // CHECK-NEXT: "affine.yield"(%{{.*}}) : (!test.type<"int">) -> () // CHECK-NEXT: }) : (!test.type<"int">) -> !test.type<"int"> - // CHECK: "affine.parallel"(%N) <{"lowerBoundsMap" = affine_map<() -> (0)>, "lowerBoundsGroups" = dense<1> : vector<1xi32>, "upperBoundsMap" = affine_map<()[s0] -> (s0)>, "upperBoundsGroups" = dense<1> : vector<1xi32>, "steps" = [1 : i64], "reductions" = []}> ({ + // CHECK: "affine.parallel"(%N) <{lowerBoundsMap = affine_map<() -> (0)>, lowerBoundsGroups = dense<1> : vector<1xi32>, upperBoundsMap = affine_map<()[s0] -> (s0)>, upperBoundsGroups = dense<1> : vector<1xi32>, steps = [1 : i64], reductions = []}> ({ // CHECK-NEXT: ^{{.*}}(%{{.*}} : index): // CHECK-NEXT: "affine.yield"() : () -> () // CHECK-NEXT: }) : (index) -> () @@ -52,7 +52,7 @@ // CHECK: %memref = "test.op"() : () -> memref<2x3xf64> // CHECK-NEXT: %value = "test.op"() : () -> f64 - // CHECK-NEXT: "affine.store"(%value, %memref) <{"map" = affine_map<() -> (0, 0)>}> : (f64, memref<2x3xf64>) -> () + // CHECK-NEXT: "affine.store"(%value, %memref) <{map = affine_map<() -> (0, 0)>}> : (f64, memref<2x3xf64>) -> () %zero = "test.op"() : () -> index %2 = affine.apply affine_map<(d0)[s0] -> (((d0 + (s0 * 42)) + -1))> (%zero)[%zero] @@ -61,8 +61,8 @@ // CHECK: %zero = "test.op"() : () -> index // CHECK-NEXT: %{{.*}} = affine.apply affine_map<(d0)[s0] -> (((d0 + (s0 * 42)) + -1))> (%{{.*}})[%{{.*}}] - // CHECK-NEXT: %{{.*}} = "affine.min"(%{{.*}}) <{"map" = affine_map<(d0) -> ((d0 + 41), d0)>}> : (index) -> index - // CHECK-NEXT: %same_value = "affine.load"(%memref, %zero, %zero) <{"map" = affine_map<(d0, d1) -> (d0, d1)>}> : (memref<2x3xf64>, index, index) -> f64 + // CHECK-NEXT: %{{.*}} = "affine.min"(%{{.*}}) <{map = affine_map<(d0) -> ((d0 + 41), d0)>}> : (index) -> index + // CHECK-NEXT: %same_value = "affine.load"(%memref, %zero, %zero) <{map = affine_map<(d0, d1) -> (d0, d1)>}> : (memref<2x3xf64>, index, index) -> f64 func.func @empty() { "affine.for"() <{"lowerBoundMap" = affine_map<() -> (0)>, "step" = 1 : index, "upperBoundMap" = affine_map<() -> (10)>, "operandSegmentSizes" = array}> ({ @@ -82,19 +82,19 @@ func.return } // CHECK: func.func @empty() { -// CHECK-NEXT: "affine.for"() <{"lowerBoundMap" = affine_map<() -> (0)>, "step" = 1 : index, "upperBoundMap" = affine_map<() -> (10)>, "operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "affine.for"() <{lowerBoundMap = affine_map<() -> (0)>, step = 1 : index, upperBoundMap = affine_map<() -> (10)>, operandSegmentSizes = array}> ({ // CHECK-NEXT: ^{{.*}}(%{{.*}} : index): // CHECK-NEXT: "affine.yield"() : () -> () // CHECK-NEXT: }) : () -> () // CHECK-NEXT: "affine.if"() ({ // CHECK-NEXT: "affine.yield"() : () -> () // CHECK-NEXT: }, { -// CHECK-NEXT: }) {"condition" = affine_set<() : (0 == 0)>} : () -> () +// CHECK-NEXT: }) {condition = affine_set<() : (0 == 0)>} : () -> () // CHECK-NEXT: "affine.if"() ({ // CHECK-NEXT: "affine.yield"() : () -> () // CHECK-NEXT: }, { // CHECK-NEXT: "affine.yield"() : () -> () -// CHECK-NEXT: }) {"condition" = affine_set<() : (0 == 0)>} : () -> () +// CHECK-NEXT: }) {condition = affine_set<() : (0 == 0)>} : () -> () // CHECK-NEXT: func.return // CHECK-NEXT: } @@ -113,7 +113,7 @@ // CHECK-NEXT: "affine.yield"(%{{.*}}) : (f32) -> () // CHECK-NEXT: }, { // CHECK-NEXT: "affine.yield"(%{{.*}}) : (f32) -> () -// CHECK-NEXT: }) {"condition" = affine_set<() : (0 == 0)>} : () -> f32 +// CHECK-NEXT: }) {condition = affine_set<() : (0 == 0)>} : () -> f32 // CHECK-NEXT: func.return %{{.*}} : f32 // CHECK-NEXT: } diff --git a/tests/filecheck/dialects/affine/examples.mlir b/tests/filecheck/dialects/affine/examples.mlir index 6f907997e7..f099a53f04 100644 --- a/tests/filecheck/dialects/affine/examples.mlir +++ b/tests/filecheck/dialects/affine/examples.mlir @@ -18,7 +18,7 @@ // CHECK: func.func private @sum_vec(%{{.*}} : memref<128xi32>) -> i32 { // CHECK-NEXT: %{{.*}} = arith.constant 0 : i32 -// CHECK-NEXT: %{{.*}} = "affine.for"(%const0) <{"lowerBoundMap" = affine_map<() -> (0)>, "upperBoundMap" = affine_map<() -> (256)>, "step" = 1 : index, "operandSegmentSizes" = array}> ({ +// CHECK-NEXT: %{{.*}} = "affine.for"(%const0) <{lowerBoundMap = affine_map<() -> (0)>, upperBoundMap = affine_map<() -> (256)>, step = 1 : index, operandSegmentSizes = array}> ({ // CHECK-NEXT: ^{{.*}}(%{{.*}} : index, %{{.*}} : i32): // CHECK-NEXT: %{{.*}} = memref.load %{{.*}}[%{{.*}}] : memref<128xi32> // CHECK-NEXT: %{{.*}} = arith.addi %{{.*}}, %{{.*}} : i32 @@ -54,11 +54,11 @@ }) {"sym_name" = "affine_mm", "function_type" = (memref<256x256xf32>, memref<256x256xf32>, memref<256x256xf32>) -> memref<256x256xf32>, "sym_visibility" = "private"} : () -> () // CHECK: func.func private @affine_mm(%{{.*}} : memref<256x256xf32>, %{{.*}} : memref<256x256xf32>, %{{.*}} : memref<256x256xf32>) -> memref<256x256xf32> { -// CHECK-NEXT: "affine.for"() <{"lowerBoundMap" = affine_map<() -> (0)>, "upperBoundMap" = affine_map<() -> (256)>, "step" = 1 : index, "operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "affine.for"() <{lowerBoundMap = affine_map<() -> (0)>, upperBoundMap = affine_map<() -> (256)>, step = 1 : index, operandSegmentSizes = array}> ({ // CHECK-NEXT: ^{{.*}}(%{{.*}} : index): -// CHECK-NEXT: "affine.for"() <{"lowerBoundMap" = affine_map<() -> (0)>, "upperBoundMap" = affine_map<() -> (256)>, "step" = 1 : index, "operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "affine.for"() <{lowerBoundMap = affine_map<() -> (0)>, upperBoundMap = affine_map<() -> (256)>, step = 1 : index, operandSegmentSizes = array}> ({ // CHECK-NEXT: ^{{.*}}(%{{.*}} : index): -// CHECK-NEXT: "affine.for"() <{"lowerBoundMap" = affine_map<() -> (0)>, "upperBoundMap" = affine_map<() -> (256)>, "step" = 1 : index, "operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "affine.for"() <{lowerBoundMap = affine_map<() -> (0)>, upperBoundMap = affine_map<() -> (256)>, step = 1 : index, operandSegmentSizes = array}> ({ // CHECK-NEXT: ^{{.*}}(%{{.*}} : index): // CHECK-NEXT: %{{.*}} = memref.load %{{.*}}[%{{.*}}, %{{.*}}] : memref<256x256xf32> // CHECK-NEXT: %{{.*}} = memref.load %{{.*}}[%{{.*}}, %{{.*}}] : memref<256x256xf32> diff --git a/tests/filecheck/dialects/air/air_ops.mlir b/tests/filecheck/dialects/air/air_ops.mlir index edbdd36b2c..0e7698aaac 100644 --- a/tests/filecheck/dialects/air/air_ops.mlir +++ b/tests/filecheck/dialects/air/air_ops.mlir @@ -33,10 +33,10 @@ func.func @graph(%arg0 : memref<32x16xi32>, %arg1 : memref<32x16xi32>) -> () { // CHECK-NEXT: %c32 = arith.constant 32 : index // CHECK-NEXT: %c16 = arith.constant 16 : index // CHECK-NEXT: %c8 = arith.constant 8 : index -// CHECK-NEXT: %buf0 = memref.alloc() {"sym_name" = "scratch"} : memref<16x8xi32, 2 : i64> -// CHECK-NEXT: %buf1 = memref.alloc() {"sym_name" = "scratch_copy"} : memref<16x8xi32, 2 : i64> -// CHECK-NEXT: %1 = air.dma_memcpy_nd(%buf0[%c0, %c0] [%c8, %c16] [%c32, %c0], %ext0[%c8, %c0] [%c8, %c16] [%c32, %c0]) {"id" = 1 : i32} : (memref<16x8xi32, 2 : i64>, memref<32x16xi32>) -// CHECK-NEXT: %2 = air.dma_memcpy_nd(%ext1[%c8, %c0] [%c8, %c16] [%c32, %c0], %buf1[%c0, %c0] [%c8, %c16] [%c32, %c0]) {"id" = 2 : i32} : (memref<32x16xi32>, memref<16x8xi32, 2 : i64>) +// CHECK-NEXT: %buf0 = memref.alloc() {sym_name = "scratch"} : memref<16x8xi32, 2 : i64> +// CHECK-NEXT: %buf1 = memref.alloc() {sym_name = "scratch_copy"} : memref<16x8xi32, 2 : i64> +// CHECK-NEXT: %1 = air.dma_memcpy_nd(%buf0[%c0, %c0] [%c8, %c16] [%c32, %c0], %ext0[%c8, %c0] [%c8, %c16] [%c32, %c0]) {id = 1 : i32} : (memref<16x8xi32, 2 : i64>, memref<32x16xi32>) +// CHECK-NEXT: %2 = air.dma_memcpy_nd(%ext1[%c8, %c0] [%c8, %c16] [%c32, %c0], %buf1[%c0, %c0] [%c8, %c16] [%c32, %c0]) {id = 2 : i32} : (memref<32x16xi32>, memref<16x8xi32, 2 : i64>) // CHECK-NEXT: air.herd_terminator // CHECK-NEXT: } // CHECK-NEXT: func.return @@ -105,10 +105,10 @@ module { // CHECK: builtin.module { // CHECK-NEXT: air.channel @channel_4 [2 : i64, 2 : i64] -// CHECK-NEXT: air.channel @channel_3 [1 : i64, 1 : i64] {"broadcast_shape" = [2 : i64, 1 : i64]} -// CHECK-NEXT: air.channel @channel_2 [1 : i64, 1 : i64] {"broadcast_shape" = [2 : i64, 1 : i64]} -// CHECK-NEXT: air.channel @channel_1 [1 : i64, 1 : i64] {"broadcast_shape" = [1 : i64, 2 : i64]} -// CHECK-NEXT: air.channel @channel_0 [1 : i64, 1 : i64] {"broadcast_shape" = [1 : i64, 2 : i64]} +// CHECK-NEXT: air.channel @channel_3 [1 : i64, 1 : i64] {broadcast_shape = [2 : i64, 1 : i64]} +// CHECK-NEXT: air.channel @channel_2 [1 : i64, 1 : i64] {broadcast_shape = [2 : i64, 1 : i64]} +// CHECK-NEXT: air.channel @channel_1 [1 : i64, 1 : i64] {broadcast_shape = [1 : i64, 2 : i64]} +// CHECK-NEXT: air.channel @channel_0 [1 : i64, 1 : i64] {broadcast_shape = [1 : i64, 2 : i64]} // CHECK-NEXT: func.func @forward(%arg0 : memref<64x128xi32>, %arg1 : memref<128x64xi32>, %arg2 : memref<64x64xi32>) { // CHECK-NEXT: %c64 = arith.constant 64 : index // CHECK-NEXT: %c32 = arith.constant 32 : index @@ -118,42 +118,42 @@ module { // CHECK-NEXT: %c2 = arith.constant 2 : index // CHECK-NEXT: %c0_i32 = arith.constant 0 : i32 // CHECK-NEXT: %async_token, %results = "air.execute"() ({ -// CHECK-NEXT: %alloc = memref.alloc() {"alignment" = 64 : i64} : memref<64x64xi32> +// CHECK-NEXT: %alloc = memref.alloc() {alignment = 64 : i64} : memref<64x64xi32> // CHECK-NEXT: "air.execute_terminator"(%alloc) : (memref<64x64xi32>) -> () // CHECK-NEXT: }) : () -> (!air.async.token, memref<64x64xi32>) // CHECK-NEXT: %async_token_1 = "air.execute"(%async_token) ({ -// CHECK-NEXT: %alloc_1 = memref.alloc() {"alignment" = 64 : i64} : memref<64x64xi32> +// CHECK-NEXT: %alloc_1 = memref.alloc() {alignment = 64 : i64} : memref<64x64xi32> // CHECK-NEXT: "air.execute_terminator"(%alloc_1) : (memref<64x64xi32>) -> () // CHECK-NEXT: }) : (!air.async.token) -> !air.async.token // CHECK-NEXT: %async_token_2, %results_1 = "air.execute"() ({ -// CHECK-NEXT: %alloc_2 = memref.alloc() {"alignment" = 64 : i64} : memref<64x64xi32> +// CHECK-NEXT: %alloc_2 = memref.alloc() {alignment = 64 : i64} : memref<64x64xi32> // CHECK-NEXT: "air.execute_terminator"(%alloc_2) : (memref<64x64xi32>) -> () // CHECK-NEXT: }) : () -> (!air.async.token, memref<64x64xi32>) // CHECK-NEXT: %async_token_3 = "air.execute"(%async_token_2, %async_token_1) ({ -// CHECK-NEXT: %alloc_3 = memref.alloc() {"alignment" = 64 : i64} : memref<64x64xi32> +// CHECK-NEXT: %alloc_3 = memref.alloc() {alignment = 64 : i64} : memref<64x64xi32> // CHECK-NEXT: "air.execute_terminator"(%alloc_3) : (memref<64x64xi32>) -> () // CHECK-NEXT: }) : (!air.async.token, !air.async.token) -> !air.async.token // CHECK-NEXT: %0 = air.wait_all async // CHECK-NEXT: %1 = scf.for %arg3 = %c0 to %c128 step %c32 iter_args(%arg4 = %0) -> (!air.async.token) { -// CHECK-NEXT: %2 = air.channel.put async[%arg4] @channel_0[] (%arg0[%c0, %arg3] [%c32, %c32] [%c128, %c1]) {"id" = 1 : i32} : (memref<64x128xi32>) +// CHECK-NEXT: %2 = air.channel.put async[%arg4] @channel_0[] (%arg0[%c0, %arg3] [%c32, %c32] [%c128, %c1]) {id = 1 : i32} : (memref<64x128xi32>) // CHECK-NEXT: scf.yield %2 : !air.async.token // CHECK-NEXT: } // CHECK-NEXT: %3 = air.wait_all async // CHECK-NEXT: %4 = scf.for %arg3_1 = %c0 to %c128 step %c32 iter_args(%arg4_1 = %3) -> (!air.async.token) { -// CHECK-NEXT: %5 = air.channel.put async[%arg4_1] @channel_1[] (%arg0[%c32, %arg3_1] [%c32, %c32] [%c128, %c1]) {"id" = 2 : i32} : (memref<64x128xi32>) +// CHECK-NEXT: %5 = air.channel.put async[%arg4_1] @channel_1[] (%arg0[%c32, %arg3_1] [%c32, %c32] [%c128, %c1]) {id = 2 : i32} : (memref<64x128xi32>) // CHECK-NEXT: scf.yield %5 : !air.async.token // CHECK-NEXT: } // CHECK-NEXT: %6 = air.wait_all async // CHECK-NEXT: %7 = scf.for %arg3_2 = %c0 to %c128 step %c32 iter_args(%arg4_2 = %6) -> (!air.async.token) { -// CHECK-NEXT: %8 = air.channel.put async[%arg4_2] @channel_2[] (%arg1[%arg3_2, %c0] [%c32, %c32] [%c64, %c1]) {"id" = 3 : i32} : (memref<128x64xi32>) +// CHECK-NEXT: %8 = air.channel.put async[%arg4_2] @channel_2[] (%arg1[%arg3_2, %c0] [%c32, %c32] [%c64, %c1]) {id = 3 : i32} : (memref<128x64xi32>) // CHECK-NEXT: scf.yield %8 : !air.async.token // CHECK-NEXT: } // CHECK-NEXT: %9 = air.wait_all async // CHECK-NEXT: %10 = scf.for %arg3_3 = %c0 to %c128 step %c32 iter_args(%arg4_3 = %9) -> (!air.async.token) { -// CHECK-NEXT: %11 = air.channel.put async[%arg4_3] @channel_3[] (%arg1[%arg3_3, %c32] [%c32, %c32] [%c64, %c1]) {"id" = 4 : i32} : (memref<128x64xi32>) +// CHECK-NEXT: %11 = air.channel.put async[%arg4_3] @channel_3[] (%arg1[%arg3_3, %c32] [%c32, %c32] [%c64, %c1]) {id = 4 : i32} : (memref<128x64xi32>) // CHECK-NEXT: scf.yield %11 : !air.async.token // CHECK-NEXT: } -// CHECK-NEXT: %12 = air.channel.get async[%async_token_3] @channel_4[] (%results_1[] [%c32, %c32] [%c64, %c1]) {"id" = 5 : i32} : (memref<64x64xi32>) +// CHECK-NEXT: %12 = air.channel.get async[%async_token_3] @channel_4[] (%results_1[] [%c32, %c32] [%c64, %c1]) {id = 5 : i32} : (memref<64x64xi32>) // CHECK-NEXT: func.return // CHECK-NEXT: } // CHECK-NEXT: } @@ -238,7 +238,7 @@ module { // CHECK-NEXT: %output1 = "air.pipeline.stage"() ({ // CHECK-NEXT: %a = memref.alloc() : memref<1024xi32, 2 : i64> // CHECK-NEXT: %b = memref.alloc() : memref<1024xi32, 2 : i64> -// CHECK-NEXT: %1 = air.dma_memcpy_nd(%b[] [] [], %op1[%c0] [%c1024] [%c1_1]) {"id" = 2 : i32} : (memref<1024xi32, 2 : i64>, memref<1024xi32>) +// CHECK-NEXT: %1 = air.dma_memcpy_nd(%b[] [] [], %op1[%c0] [%c1024] [%c1_1]) {id = 2 : i32} : (memref<1024xi32, 2 : i64>, memref<1024xi32>) // CHECK-NEXT: %init = tensor.empty() : tensor<1024xi32> // CHECK-NEXT: air.pipeline.yield %init : tensor<1024xi32> // CHECK-NEXT: }) : () -> tensor<1024xi32> @@ -276,7 +276,7 @@ module { // CHECK-NEXT: air.pipeline.yield // CHECK-NEXT: }) : () -> () // CHECK-NEXT: air.pipeline.terminator -// CHECK-NEXT: }) {"direction" = "horiz"} : () -> () +// CHECK-NEXT: }) {direction = "horiz"} : () -> () // CHECK-NEXT: air.herd_terminator // CHECK-NEXT: } // CHECK-NEXT: func.return diff --git a/tests/filecheck/dialects/arith/arith_ops.mlir b/tests/filecheck/dialects/arith/arith_ops.mlir index 6c0d0c51f9..b02dbb7144 100644 --- a/tests/filecheck/dialects/arith/arith_ops.mlir +++ b/tests/filecheck/dialects/arith/arith_ops.mlir @@ -115,7 +115,7 @@ %addi = arith.addi %lhsi32, %rhsi32 {"hello" = "world"} : i32 - // CHECK-NEXT: %addi = arith.addi %lhsi32, %rhsi32 {"hello" = "world"} : i32 + // CHECK-NEXT: %addi = arith.addi %lhsi32, %rhsi32 {hello = "world"} : i32 %addf = arith.addf %lhsf32, %rhsf32 : f32 %addf_vector = arith.addf %lhsvec, %rhsvec : vector<4xf32> diff --git a/tests/filecheck/dialects/arm/test_ops.mlir b/tests/filecheck/dialects/arm/test_ops.mlir index 3a2020293b..0f3f0a5e86 100644 --- a/tests/filecheck/dialects/arm/test_ops.mlir +++ b/tests/filecheck/dialects/arm/test_ops.mlir @@ -9,14 +9,14 @@ // CHECK: %x2 = arm.get_register : !arm.reg %x2 = arm.get_register : !arm.reg -// CHECK: %ds_mov = arm.ds.mov %x1 {"comment" = "move contents of s to d"} : (!arm.reg) -> !arm.reg +// CHECK: %ds_mov = arm.ds.mov %x1 {comment = "move contents of s to d"} : (!arm.reg) -> !arm.reg // CHECK-ASM: mov x2, x1 # move contents of s to d %ds_mov = arm.ds.mov %x1 {"comment" = "move contents of s to d"} : (!arm.reg) -> !arm.reg -// CHECK: %dss_mul = arm.dss.mul %x1, %x2 {"comment" = "multiply s1 by s2"} : (!arm.reg, !arm.reg) -> !arm.reg +// CHECK: %dss_mul = arm.dss.mul %x1, %x2 {comment = "multiply s1 by s2"} : (!arm.reg, !arm.reg) -> !arm.reg // CHECK-ASM: mul x3, x1, x2 # multiply s1 by s2 %dss_mul = arm.dss.mul %x1, %x2 {"comment" = "multiply s1 by s2"} : (!arm.reg, !arm.reg) -> !arm.reg // CHECK-GENERIC: %x1 = "arm.get_register"() : () -> !arm.reg -// CHECK-GENERIC: %ds_mov = "arm.ds.mov"(%x1) {"comment" = "move contents of s to d"} : (!arm.reg) -> !arm.reg -// CHECK-GENERIC: %dss_mul = "arm.dss.mul"(%x1, %x2) {"comment" = "multiply s1 by s2"} : (!arm.reg, !arm.reg) -> !arm.reg +// CHECK-GENERIC: %ds_mov = "arm.ds.mov"(%x1) {comment = "move contents of s to d"} : (!arm.reg) -> !arm.reg +// CHECK-GENERIC: %dss_mul = "arm.dss.mul"(%x1, %x2) {comment = "multiply s1 by s2"} : (!arm.reg, !arm.reg) -> !arm.reg diff --git a/tests/filecheck/dialects/arm/test_registers.mlir b/tests/filecheck/dialects/arm/test_registers.mlir index 9a84b29d6f..3137477762 100644 --- a/tests/filecheck/dialects/arm/test_registers.mlir +++ b/tests/filecheck/dialects/arm/test_registers.mlir @@ -1,8 +1,8 @@ // RUN: XDSL_ROUNDTRIP -// CHECK: "test.op"() {"unallocated" = !arm.reg} : () -> () +// CHECK: "test.op"() {unallocated = !arm.reg} : () -> () "test.op"() {unallocated = !arm.reg} : () -> () -// CHECK: "test.op"() {"allocated" = !arm.reg} : () -> () +// CHECK: "test.op"() {allocated = !arm.reg} : () -> () "test.op"() {allocated = !arm.reg} : () -> () diff --git a/tests/filecheck/dialects/arm_func/arm_func_ops.mlir b/tests/filecheck/dialects/arm_func/arm_func_ops.mlir index 97da638e41..0f260005a4 100644 --- a/tests/filecheck/dialects/arm_func/arm_func_ops.mlir +++ b/tests/filecheck/dialects/arm_func/arm_func_ops.mlir @@ -3,7 +3,7 @@ // RUN: xdsl-opt -t arm-asm %s | filecheck %s --check-prefix=CHECK-ASM // CHECK: arm_func.func @noarg_void() { -// CHECK-NEXT: arm_func.return {"comment" = "this is a return instruction"} +// CHECK-NEXT: arm_func.return {comment = "this is a return instruction"} // CHECK-NEXT: } // CHECK-ASM: noarg_void: // CHECK-ASM: ret # this is a return instruction @@ -12,5 +12,5 @@ arm_func.func @noarg_void() { } // CHECK-GENERIC: "arm_func.func"() ({ -// CHECK-GENERIC-NEXT: "arm_func.return"() {"comment" = "this is a return instruction"} : () -> () -// CHECK-GENERIC-NEXT: }) {"sym_name" = "noarg_void", "function_type" = () -> ()} : () -> () +// CHECK-GENERIC-NEXT: "arm_func.return"() {comment = "this is a return instruction"} : () -> () +// CHECK-GENERIC-NEXT: }) {sym_name = "noarg_void", function_type = () -> ()} : () -> () diff --git a/tests/filecheck/dialects/bufferization/bufferization_ops.mlir b/tests/filecheck/dialects/bufferization/bufferization_ops.mlir index 9a2e5d1cf4..c07d2f30b0 100644 --- a/tests/filecheck/dialects/bufferization/bufferization_ops.mlir +++ b/tests/filecheck/dialects/bufferization/bufferization_ops.mlir @@ -6,7 +6,7 @@ %i0, %i1 = "test.op"() : () -> (index, index) %t0, %m0 = "test.op"() : () -> (tensor<10x20x30xf64>, memref<10x20x30xf64>) -// CHECK-NEXT: %t1 = bufferization.alloc_tensor(%i0, %i1) {"hello" = "world"} : tensor<10x20x?x?xf64> +// CHECK-NEXT: %t1 = bufferization.alloc_tensor(%i0, %i1) {hello = "world"} : tensor<10x20x?x?xf64> // CHECK-NEXT: %t2 = bufferization.alloc_tensor() copy(%t0) : tensor<10x20x30xf64> // CHECK-NEXT: %t3 = bufferization.alloc_tensor(%i0, %i1) size_hint = %i1 : tensor<10x20x?x?xf64> %t1 = bufferization.alloc_tensor(%i0, %i1) {"hello"="world"}: tensor<10x20x?x?xf64> @@ -26,11 +26,11 @@ bufferization.materialize_in_destination %t0 in %t0 : (tensor<10x20x30xf64>, ten // CHECK-GENERIC: "builtin.module"() ({ // CHECK-GENERIC-NEXT: %i0, %i1 = "test.op"() : () -> (index, index) // CHECK-GENERIC-NEXT: %t0, %m0 = "test.op"() : () -> (tensor<10x20x30xf64>, memref<10x20x30xf64>) -// CHECK-GENERIC-NEXT: %t1 = "bufferization.alloc_tensor"(%i0, %i1) <{"operandSegmentSizes" = array}> {"hello" = "world"} : (index, index) -> tensor<10x20x?x?xf64> -// CHECK-GENERIC-NEXT: %t2 = "bufferization.alloc_tensor"(%t0) <{"operandSegmentSizes" = array}> : (tensor<10x20x30xf64>) -> tensor<10x20x30xf64> -// CHECK-GENERIC-NEXT: %t3 = "bufferization.alloc_tensor"(%i0, %i1, %i1) <{"operandSegmentSizes" = array}> : (index, index, index) -> tensor<10x20x?x?xf64> +// CHECK-GENERIC-NEXT: %t1 = "bufferization.alloc_tensor"(%i0, %i1) <{operandSegmentSizes = array}> {hello = "world"} : (index, index) -> tensor<10x20x?x?xf64> +// CHECK-GENERIC-NEXT: %t2 = "bufferization.alloc_tensor"(%t0) <{operandSegmentSizes = array}> : (tensor<10x20x30xf64>) -> tensor<10x20x30xf64> +// CHECK-GENERIC-NEXT: %t3 = "bufferization.alloc_tensor"(%i0, %i1, %i1) <{operandSegmentSizes = array}> : (index, index, index) -> tensor<10x20x?x?xf64> // CHECK-GENERIC-NEXT: %m = "test.op"() : () -> memref<30x20x10xf32> // CHECK-GENERIC-NEXT: %c = "bufferization.clone"(%m) : (memref<30x20x10xf32>) -> memref<30x20x10xf32> -// CHECK-GENERIC-NEXT: "bufferization.materialize_in_destination"(%t0, %m0) <{"writable"}> : (tensor<10x20x30xf64>, memref<10x20x30xf64>) -> () +// CHECK-GENERIC-NEXT: "bufferization.materialize_in_destination"(%t0, %m0) <{writable}> : (tensor<10x20x30xf64>, memref<10x20x30xf64>) -> () // CHECK-GENERIC-NEXT: %0 = "bufferization.materialize_in_destination"(%t0, %t0) : (tensor<10x20x30xf64>, tensor<10x20x30xf64>) -> tensor<10x20x30xf64> // CHECK-GENERIC-NEXT: }) : () -> () diff --git a/tests/filecheck/dialects/builtin/attrs.mlir b/tests/filecheck/dialects/builtin/attrs.mlir index f0257933f4..66ab2167bd 100644 --- a/tests/filecheck/dialects/builtin/attrs.mlir +++ b/tests/filecheck/dialects/builtin/attrs.mlir @@ -3,41 +3,41 @@ // CHECK: module "builtin.module"() ({ "func.func"() ({ - // CHECK: "test" = dense<0> : tensor<1xi32> + // CHECK: test = dense<0> : tensor<1xi32> %x1 = "arith.constant"() {"value" = 0 : i64, "test" = dense<0> : tensor<1xi32>} : () -> i64 - // CHECK: "test" = dense<0.000000e+00> : tensor<1xf32> + // CHECK: test = dense<0.000000e+00> : tensor<1xf32> %x2 = "arith.constant"() {"value" = 0 : i64, "test" = dense<0.0> : tensor<1xf32>} : () -> i64 - // CHECK: "test" = true + // CHECK: test = true %x3 = "arith.constant"() {"value" = 0 : i64, "test" = true} : () -> i64 - // CHECK: "test" = false + // CHECK: test = false %x4 = "arith.constant"() {"value" = 0 : i64, "test" = false} : () -> i64 - // CHECK: "test" = true + // CHECK: test = true %x5 = "arith.constant"() {"value" = 0 : i64, "test" = true} : () -> i64 - // CHECK: "test" = false + // CHECK: test = false %x6 = "arith.constant"() {"value" = 0 : i64, "test" = false} : () -> i64 - // CHECK: "test" = array + // CHECK: test = array %x7 = "arith.constant"() {"value" = 0 : i64, "test" = array} : () -> i64 - // CHECK: "test" = array + // CHECK: test = array %x8 = "arith.constant"() {"value" = 0 : i64, "test" = array} : () -> i64 - // CHECK: "test" = #builtin.signedness + // CHECK: test = #builtin.signedness %x9 = "arith.constant"() {"value" = 0 : i64, "test" = #builtin.signedness} : () -> i64 - // CHECK: "test" = #builtin.signedness + // CHECK: test = #builtin.signedness %x10 = "arith.constant"() {"value" = 0 : i64, "test" = #builtin.signedness} : () -> i64 - // CHECK: "test" = #builtin.signedness + // CHECK: test = #builtin.signedness %x11 = "arith.constant"() {"value" = 0 : i64, "test" = #builtin.signedness} : () -> i64 - // CHECK: "test" = @foo + // CHECK: test = @foo %x12 = "arith.constant"() {"value" = 0 : i64, "test" = @foo} : () -> i64 - // CHECK: "test" = @foo::@bar + // CHECK: test = @foo::@bar %x13 = "arith.constant"() {"value" = 0 : i64, "test" = @foo::@bar} : () -> i64 - // CHECK: "test" = @foo::@bar::@baz + // CHECK: test = @foo::@bar::@baz %x14 = "arith.constant"() {"value" = 0 : i64, "test" = @foo::@bar::@baz} : () -> i64 - // CHECK: "test" = loc(unknown) + // CHECK: test = loc(unknown) %x15 = "arith.constant"() {"value" = 0 : i64, "test" = loc(unknown)} : () -> i64 - // CHECK: "test" = none + // CHECK: test = none %x16 = "arith.constant"() {"value" = 0 : i64, "test" = none} : () -> i64 - // CHECK: "test" = 0 : i0 + // CHECK: test = 0 : i0 %x17 = "arith.constant"() {"value" = 0 : i64, "test" = 0 : i0} : () -> i64 - // CHECK: %x18 = arith.constant {"array" = array, "dense" = dense<-1> : vector, "test" = -1 : i8} -1 : i8 + // CHECK: %x18 = arith.constant {array = array, dense = dense<-1> : vector, test = -1 : i8} -1 : i8 %x18 = arith.constant {"array" = array, "dense" = dense<255> : vector, "test" = 255 : i8} -1 : i8 "func.return"() : () -> () }) {"function_type" = () -> (), "sym_name" = "builtin"} : () -> () diff --git a/tests/filecheck/dialects/builtin/module.mlir b/tests/filecheck/dialects/builtin/module.mlir index 3ffb7254bc..195b50ec6c 100644 --- a/tests/filecheck/dialects/builtin/module.mlir +++ b/tests/filecheck/dialects/builtin/module.mlir @@ -7,13 +7,13 @@ builtin.module { // CHECK: builtin.module { // CHECK-NEXT: } builtin.module attributes {a = "foo", b = "bar", unit} {} - // CHECK-NEXT: builtin.module attributes {"a" = "foo", "b" = "bar", "unit"} { + // CHECK-NEXT: builtin.module attributes {a = "foo", b = "bar", unit} { // CHECK-NEXT: } builtin.module @moduleName {} // CHECK-NEXT: builtin.module @moduleName { // CHECK-NEXT: } builtin.module @otherModule attributes {dialect.attr} {} - // CHECK-NEXT: builtin.module @otherModule attributes {"dialect.attr"} { + // CHECK-NEXT: builtin.module @otherModule attributes {dialect.attr} { // CHECK-NEXT: } module {} // CHECK-NEXT: builtin.module { diff --git a/tests/filecheck/dialects/builtin/unrealized_conv_cast.mlir b/tests/filecheck/dialects/builtin/unrealized_conv_cast.mlir index 4bc9130db1..9ed020b309 100644 --- a/tests/filecheck/dialects/builtin/unrealized_conv_cast.mlir +++ b/tests/filecheck/dialects/builtin/unrealized_conv_cast.mlir @@ -11,13 +11,13 @@ %2 = "builtin.unrealized_conversion_cast"(%0) : (i64) -> i32 // CHECK: %3 = builtin.unrealized_conversion_cast %0 : i64 to i64 %3 = "builtin.unrealized_conversion_cast"(%0) : (i64) -> i64 - // CHECK: %4 = builtin.unrealized_conversion_cast to i64 {"comment" = "test"} + // CHECK: %4 = builtin.unrealized_conversion_cast to i64 {comment = "test"} %4 = "builtin.unrealized_conversion_cast"() {"comment" = "test"} : () -> i64 // CHECK: %5 = builtin.unrealized_conversion_cast %0, %0 : i64, i64 to f32 %5 = "builtin.unrealized_conversion_cast"(%0, %0) : (i64, i64) -> f32 // CHECK: %6, %7 = builtin.unrealized_conversion_cast %5 : f32 to i64, i64 %6, %7 = "builtin.unrealized_conversion_cast"(%5) : (f32) -> (i64, i64) - // CHECK: %8 = builtin.unrealized_conversion_cast to none {"comment" = "test"} + // CHECK: %8 = builtin.unrealized_conversion_cast to none {comment = "test"} %8 = "builtin.unrealized_conversion_cast"() {"comment" = "test"} : () -> none "func.return"() : () -> () }) {"function_type" = () -> (), "sym_name" = "builtin"} : () -> () diff --git a/tests/filecheck/dialects/cf/cf_ops.mlir b/tests/filecheck/dialects/cf/cf_ops.mlir index 8077cf3775..1e3361ade8 100644 --- a/tests/filecheck/dialects/cf/cf_ops.mlir +++ b/tests/filecheck/dialects/cf/cf_ops.mlir @@ -13,9 +13,9 @@ builtin.module { // CHECK-NEXT: func.return // CHECK-NEXT: } - // CHECK-GENERIC: "func.func"() <{"sym_name" = "assert", "function_type" = () -> (), "sym_visibility" = "private"}> ({ - // CHECK-GENERIC-NEXT: %{{.*}} = "arith.constant"() <{"value" = true}> : () -> i1 - // CHECK-GENERIC-NEXT: "cf.assert"(%{{.*}}) {"msg" = "some message"} : (i1) -> () + // CHECK-GENERIC: "func.func"() <{sym_name = "assert", function_type = () -> (), sym_visibility = "private"}> ({ + // CHECK-GENERIC-NEXT: %{{.*}} = "arith.constant"() <{value = true}> : () -> i1 + // CHECK-GENERIC-NEXT: "cf.assert"(%{{.*}}) {msg = "some message"} : (i1) -> () // CHECK-GENERIC-NEXT: "func.return"() : () -> () // CHECK-GENERIC-NEXT: } @@ -30,7 +30,7 @@ builtin.module { // CHECK-NEXT: cf.br ^0 // CHECK-NEXT: } - // CHECK-GENERIC: "func.func"() <{"sym_name" = "unconditional_br", "function_type" = () -> (), "sym_visibility" = "private"}> ({ + // CHECK-GENERIC: "func.func"() <{sym_name = "unconditional_br", function_type = () -> (), sym_visibility = "private"}> ({ // CHECK-GENERIC-NEXT: "cf.br"() [^{{.*}}] : () -> () // CHECK-GENERIC-NEXT: ^{{.*}}: // CHECK-GENERIC-NEXT: "cf.br"() [^{{.*}}] : () -> () @@ -47,7 +47,7 @@ builtin.module { // CHECK-NEXT: cf.br ^0(%1 : i32) // CHECK-NEXT: } - // CHECK-GENERIC: "func.func"() <{"sym_name" = "br", "function_type" = (i32) -> (), "sym_visibility" = "private"}> ({ + // CHECK-GENERIC: "func.func"() <{sym_name = "br", function_type = (i32) -> (), sym_visibility = "private"}> ({ // CHECK-GENERIC-NEXT: ^{{.*}}(%{{.*}} : i32): // CHECK-GENERIC-NEXT: "cf.br"(%{{.*}}) [^{{.*}}] : (i32) -> () // CHECK-GENERIC-NEXT: ^{{.*}}(%{{.*}} : i32): @@ -70,11 +70,11 @@ builtin.module { // CHECK-NEXT: func.return %4 : i32 // CHECK-NEXT: } - // CHECK-GENERIC: "func.func"() <{"sym_name" = "cond_br", "function_type" = (i1, i32) -> i32, "sym_visibility" = "private"}> ({ + // CHECK-GENERIC: "func.func"() <{sym_name = "cond_br", function_type = (i1, i32) -> i32, sym_visibility = "private"}> ({ // CHECK-GENERIC-NEXT: ^{{.*}}(%{{.*}} : i1, %{{.*}} : i32): // CHECK-GENERIC-NEXT: "cf.br"(%{{.*}}, %{{.*}}) [^{{.*}}] : (i1, i32) -> () // CHECK-GENERIC-NEXT: ^{{.*}}(%{{.*}} : i1, %{{.*}} : i32): - // CHECK-GENERIC-NEXT: "cf.cond_br"(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) [^{{.*}}, ^{{.*}}] <{"operandSegmentSizes" = array}> : (i1, i1, i32, i32, i32, i32) -> () + // CHECK-GENERIC-NEXT: "cf.cond_br"(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) [^{{.*}}, ^{{.*}}] <{operandSegmentSizes = array}> : (i1, i1, i32, i32, i32, i32) -> () // CHECK-GENERIC-NEXT: ^{{.*}}(%{{.*}} : i32, %{{.*}} : i32, %{{.*}} : i32): // CHECK-GENERIC-NEXT: "func.return"(%{{.*}}) : (i32) -> () // CHECK-GENERIC-NEXT: } @@ -111,11 +111,11 @@ builtin.module { // CHECK-NEXT: func.return // CHECK-NEXT: } - // CHECK-GENERIC: "func.func"() <{"sym_name" = "switch", "function_type" = (i32) -> ()}> ({ + // CHECK-GENERIC: "func.func"() <{sym_name = "switch", function_type = (i32) -> ()}> ({ // CHECK-GENERIC-NEXT: ^{{.*}}(%flag : i32): - // CHECK-GENERIC-NEXT: %a = "arith.constant"() <{"value" = 0 : i32}> : () -> i32 - // CHECK-GENERIC-NEXT: %b = "arith.constant"() <{"value" = 1 : i32}> : () -> i32 - // CHECK-GENERIC-NEXT: "cf.switch"(%flag, %a, %b, %b) [^[[#b0:]], ^[[#b1:]], ^[[#b2:]]] <{"case_operand_segments" = array, "case_values" = dense<[42, 43]> : vector<2xi32>, "operandSegmentSizes" = array}> : (i32, i32, i32, i32) -> () + // CHECK-GENERIC-NEXT: %a = "arith.constant"() <{value = 0 : i32}> : () -> i32 + // CHECK-GENERIC-NEXT: %b = "arith.constant"() <{value = 1 : i32}> : () -> i32 + // CHECK-GENERIC-NEXT: "cf.switch"(%flag, %a, %b, %b) [^[[#b0:]], ^[[#b1:]], ^[[#b2:]]] <{case_operand_segments = array, case_values = dense<[42, 43]> : vector<2xi32>, operandSegmentSizes = array}> : (i32, i32, i32, i32) -> () // CHECK-GENERIC-NEXT: ^[[#b0]](%{{.*}} : i32): // CHECK-GENERIC-NEXT: "func.return"() : () -> () // CHECK-GENERIC-NEXT: ^[[#b1]](%{{.*}} : i32, %{{.*}} : i32): diff --git a/tests/filecheck/dialects/csl/csl-canonicalize.mlir b/tests/filecheck/dialects/csl/csl-canonicalize.mlir index 7188986cb7..9c31534273 100644 --- a/tests/filecheck/dialects/csl/csl-canonicalize.mlir +++ b/tests/filecheck/dialects/csl/csl-canonicalize.mlir @@ -21,7 +21,7 @@ builtin.module { // CHECK-NEXT: %0 = "csl.zeros"() : () -> memref<512xf32> // CHECK-NEXT: %1 = arith.constant 510 : ui16 -// CHECK-NEXT: %2 = "csl.get_mem_dsd"(%0, %1) <{"tensor_access" = affine_map<(d0) -> (((d0 * 3) + 1))>}> : (memref<512xf32>, ui16) -> !csl +// CHECK-NEXT: %2 = "csl.get_mem_dsd"(%0, %1) <{tensor_access = affine_map<(d0) -> (((d0 * 3) + 1))>}> : (memref<512xf32>, ui16) -> !csl // CHECK-NEXT: "test.op"(%2) : (!csl) -> () @@ -42,7 +42,7 @@ builtin.module { // CHECK-NEXT: %3 = "test.op"() : () -> !csl // CHECK-NEXT: %4 = arith.constant 4 : si16 -// CHECK-NEXT: %5 = "csl.increment_dsd_offset"(%3, %4) <{"elem_type" = f32}> : (!csl, si16) -> !csl +// CHECK-NEXT: %5 = "csl.increment_dsd_offset"(%3, %4) <{elem_type = f32}> : (!csl, si16) -> !csl // CHECK-NEXT: %6 = arith.constant 511 : ui16 // CHECK-NEXT: %7 = "csl.set_dsd_length"(%5, %6) : (!csl, ui16) -> !csl // CHECK-NEXT: %8 = arith.constant 3 : si8 diff --git a/tests/filecheck/dialects/csl/csl-stencil-canonicalize.mlir b/tests/filecheck/dialects/csl/csl-stencil-canonicalize.mlir index cbefbb2493..d7bfbe7fdc 100644 --- a/tests/filecheck/dialects/csl/csl-stencil-canonicalize.mlir +++ b/tests/filecheck/dialects/csl/csl-stencil-canonicalize.mlir @@ -49,30 +49,30 @@ builtin.module { // CHECK-NEXT: func.func @gauss_seidel(%a : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>, %b : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>) { // CHECK-NEXT: %0 = stencil.load %a : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>> -> !stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>> // CHECK-NEXT: %1 = tensor.empty() : tensor<510xf32> -// CHECK-NEXT: %2 = csl_stencil.apply(%0 : !stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>>, %1 : tensor<510xf32>) -> (!stencil.temp<[0,1]x[0,1]xtensor<510xf32>>) <{"num_chunks" = 2 : i64, "topo" = #dmp.topo<1022x510>, "swaps" = [#csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange], "operandSegmentSizes" = array}> ({ +// CHECK-NEXT: %2 = csl_stencil.apply(%0 : !stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>>, %1 : tensor<510xf32>) -> (!stencil.temp<[0,1]x[0,1]xtensor<510xf32>>) <{num_chunks = 2 : i64, topo = #dmp.topo<1022x510>, swaps = [#csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange], operandSegmentSizes = array}> ({ // CHECK-NEXT: ^0(%3 : tensor<4x255xf32>, %4 : index, %5 : tensor<510xf32>): // CHECK-NEXT: %6 = csl_stencil.access %3[1, 0] : tensor<4x255xf32> -// CHECK-NEXT: %7 = "tensor.insert_slice"(%6, %5, %4) <{"static_offsets" = array, "static_sizes" = array, "static_strides" = array, "operandSegmentSizes" = array}> : (tensor<255xf32>, tensor<510xf32>, index) -> tensor<510xf32> +// CHECK-NEXT: %7 = "tensor.insert_slice"(%6, %5, %4) <{static_offsets = array, static_sizes = array, static_strides = array, operandSegmentSizes = array}> : (tensor<255xf32>, tensor<510xf32>, index) -> tensor<510xf32> // CHECK-NEXT: csl_stencil.yield %7 : tensor<510xf32> // CHECK-NEXT: }, { // CHECK-NEXT: ^1(%8 : !stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>>, %9 : tensor<510xf32>): // CHECK-NEXT: csl_stencil.yield %9 : tensor<510xf32> // CHECK-NEXT: }) // CHECK-NEXT: stencil.store %2 to %b(<[0, 0], [1, 1]>) : !stencil.temp<[0,1]x[0,1]xtensor<510xf32>> to !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>> -// CHECK-NEXT: %3 = csl_stencil.apply(%0 : !stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>>, %1 : tensor<510xf32>) -> (!stencil.temp<[0,1]x[0,1]xtensor<510xf32>>) <{"num_chunks" = 2 : i64, "topo" = #dmp.topo<1022x510>, "swaps" = [#csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange], "operandSegmentSizes" = array}> ({ +// CHECK-NEXT: %3 = csl_stencil.apply(%0 : !stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>>, %1 : tensor<510xf32>) -> (!stencil.temp<[0,1]x[0,1]xtensor<510xf32>>) <{num_chunks = 2 : i64, topo = #dmp.topo<1022x510>, swaps = [#csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange], operandSegmentSizes = array}> ({ // CHECK-NEXT: ^0(%4 : tensor<4x255xf32>, %5 : index, %6 : tensor<510xf32>): // CHECK-NEXT: %7 = csl_stencil.access %4[1, 0] : tensor<4x255xf32> -// CHECK-NEXT: %8 = "tensor.insert_slice"(%7, %6, %5) <{"static_offsets" = array, "static_sizes" = array, "static_strides" = array, "operandSegmentSizes" = array}> : (tensor<255xf32>, tensor<510xf32>, index) -> tensor<510xf32> +// CHECK-NEXT: %8 = "tensor.insert_slice"(%7, %6, %5) <{static_offsets = array, static_sizes = array, static_strides = array, operandSegmentSizes = array}> : (tensor<255xf32>, tensor<510xf32>, index) -> tensor<510xf32> // CHECK-NEXT: csl_stencil.yield %8 : tensor<510xf32> // CHECK-NEXT: }, { // CHECK-NEXT: ^1(%9 : !stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>>, %10 : tensor<510xf32>): // CHECK-NEXT: csl_stencil.yield %10 : tensor<510xf32> // CHECK-NEXT: }) // CHECK-NEXT: stencil.store %3 to %b(<[0, 0], [1, 1]>) : !stencil.temp<[0,1]x[0,1]xtensor<510xf32>> to !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>> -// CHECK-NEXT: %4 = csl_stencil.apply(%0 : !stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>>, %1 : tensor<510xf32>) -> (!stencil.temp<[0,1]x[0,1]xtensor<510xf32>>) <{"num_chunks" = 2 : i64, "topo" = #dmp.topo<1022x510>, "swaps" = [#csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange], "operandSegmentSizes" = array}> ({ +// CHECK-NEXT: %4 = csl_stencil.apply(%0 : !stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>>, %1 : tensor<510xf32>) -> (!stencil.temp<[0,1]x[0,1]xtensor<510xf32>>) <{num_chunks = 2 : i64, topo = #dmp.topo<1022x510>, swaps = [#csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange], operandSegmentSizes = array}> ({ // CHECK-NEXT: ^0(%5 : tensor<4x255xf32>, %6 : index, %7 : tensor<510xf32>): // CHECK-NEXT: %8 = csl_stencil.access %5[1, 0] : tensor<4x255xf32> -// CHECK-NEXT: %9 = "tensor.insert_slice"(%8, %7, %6) <{"static_offsets" = array, "static_sizes" = array, "static_strides" = array, "operandSegmentSizes" = array}> : (tensor<255xf32>, tensor<510xf32>, index) -> tensor<510xf32> +// CHECK-NEXT: %9 = "tensor.insert_slice"(%8, %7, %6) <{static_offsets = array, static_sizes = array, static_strides = array, operandSegmentSizes = array}> : (tensor<255xf32>, tensor<510xf32>, index) -> tensor<510xf32> // CHECK-NEXT: csl_stencil.yield %9 : tensor<510xf32> // CHECK-NEXT: }, { // CHECK-NEXT: ^1(%10 : !stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>>, %11 : tensor<510xf32>): diff --git a/tests/filecheck/dialects/csl/csl-stencil-ops.mlir b/tests/filecheck/dialects/csl/csl-stencil-ops.mlir index 24cc1c68c0..1d59128710 100644 --- a/tests/filecheck/dialects/csl/csl-stencil-ops.mlir +++ b/tests/filecheck/dialects/csl/csl-stencil-ops.mlir @@ -33,15 +33,15 @@ builtin.module { // CHECK: builtin.module { // CHECK-NEXT: func.func @gauss_seidel_func(%a : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>, %b : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>) { // CHECK-NEXT: %0 = stencil.load %a : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>> -> !stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>> -// CHECK-NEXT: %pref = "csl_stencil.prefetch"(%0) <{"topo" = #dmp.topo<1022x510>, "swaps" = [#csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange], "num_chunks" = 2 : i64}> : (!stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>>) -> tensor<4x510xf32> +// CHECK-NEXT: %pref = "csl_stencil.prefetch"(%0) <{topo = #dmp.topo<1022x510>, swaps = [#csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange], num_chunks = 2 : i64}> : (!stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>>) -> tensor<4x510xf32> // CHECK-NEXT: %1 = stencil.apply(%2 = %0 : !stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>>, %3 = %pref : tensor<4x510xf32>) -> (!stencil.temp<[0,1]x[0,1]xtensor<510xf32>>) { // CHECK-NEXT: %4 = arith.constant 1.666600e-01 : f32 // CHECK-NEXT: %5 = csl_stencil.access %3[1, 0] : tensor<4x510xf32> // CHECK-NEXT: %6 = csl_stencil.access %3[-1, 0] : tensor<4x510xf32> // CHECK-NEXT: %7 = csl_stencil.access %2[0, 0] : !stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>> // CHECK-NEXT: %8 = csl_stencil.access %2[0, 0] : !stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>> -// CHECK-NEXT: %9 = "tensor.extract_slice"(%7) <{"static_offsets" = array, "static_sizes" = array, "static_strides" = array, "operandSegmentSizes" = array}> : (tensor<512xf32>) -> tensor<510xf32> -// CHECK-NEXT: %10 = "tensor.extract_slice"(%8) <{"static_offsets" = array, "static_sizes" = array, "static_strides" = array, "operandSegmentSizes" = array}> : (tensor<512xf32>) -> tensor<510xf32> +// CHECK-NEXT: %9 = "tensor.extract_slice"(%7) <{static_offsets = array, static_sizes = array, static_strides = array, operandSegmentSizes = array}> : (tensor<512xf32>) -> tensor<510xf32> +// CHECK-NEXT: %10 = "tensor.extract_slice"(%8) <{static_offsets = array, static_sizes = array, static_strides = array, operandSegmentSizes = array}> : (tensor<512xf32>) -> tensor<510xf32> // CHECK-NEXT: %11 = csl_stencil.access %3[0, 1] : tensor<4x510xf32> // CHECK-NEXT: %12 = csl_stencil.access %3[0, -1] : tensor<4x510xf32> // CHECK-NEXT: %13 = arith.addf %12, %11 : tensor<510xf32> @@ -60,35 +60,35 @@ builtin.module { // CHECK-NEXT: } // CHECK-GENERIC: "builtin.module"() ({ -// CHECK-GENERIC-NEXT: "func.func"() <{"sym_name" = "gauss_seidel_func", "function_type" = (!stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>, !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>) -> ()}> ({ +// CHECK-GENERIC-NEXT: "func.func"() <{sym_name = "gauss_seidel_func", function_type = (!stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>, !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>) -> ()}> ({ // CHECK-GENERIC-NEXT: ^0(%a : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>, %b : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>): // CHECK-GENERIC-NEXT: %0 = "stencil.load"(%a) : (!stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>) -> !stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>> -// CHECK-GENERIC-NEXT: %pref = "csl_stencil.prefetch"(%0) <{"topo" = #dmp.topo<1022x510>, "swaps" = [#csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange], "num_chunks" = 2 : i64}> : (!stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>>) -> tensor<4x510xf32> -// CHECK-GENERIC-NEXT: %1 = "stencil.apply"(%0, %pref) <{"operandSegmentSizes" = array}> ({ +// CHECK-GENERIC-NEXT: %pref = "csl_stencil.prefetch"(%0) <{topo = #dmp.topo<1022x510>, swaps = [#csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange], num_chunks = 2 : i64}> : (!stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>>) -> tensor<4x510xf32> +// CHECK-GENERIC-NEXT: %1 = "stencil.apply"(%0, %pref) <{operandSegmentSizes = array}> ({ // CHECK-GENERIC-NEXT: ^1(%2 : !stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>>, %3 : tensor<4x510xf32>): -// CHECK-GENERIC-NEXT: %4 = "arith.constant"() <{"value" = 1.666600e-01 : f32}> : () -> f32 -// CHECK-GENERIC-NEXT: %5 = "csl_stencil.access"(%3) <{"offset" = #stencil.index<[1, 0]>, "offset_mapping" = #stencil.index<[0, 1]>}> : (tensor<4x510xf32>) -> tensor<510xf32> -// CHECK-GENERIC-NEXT: %6 = "csl_stencil.access"(%3) <{"offset" = #stencil.index<[-1, 0]>, "offset_mapping" = #stencil.index<[0, 1]>}> : (tensor<4x510xf32>) -> tensor<510xf32> -// CHECK-GENERIC-NEXT: %7 = "csl_stencil.access"(%2) <{"offset" = #stencil.index<[0, 0]>, "offset_mapping" = #stencil.index<[0, 1]>}> : (!stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>>) -> tensor<512xf32> -// CHECK-GENERIC-NEXT: %8 = "csl_stencil.access"(%2) <{"offset" = #stencil.index<[0, 0]>, "offset_mapping" = #stencil.index<[0, 1]>}> : (!stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>>) -> tensor<512xf32> -// CHECK-GENERIC-NEXT: %9 = "tensor.extract_slice"(%7) <{"static_offsets" = array, "static_sizes" = array, "static_strides" = array, "operandSegmentSizes" = array}> : (tensor<512xf32>) -> tensor<510xf32> -// CHECK-GENERIC-NEXT: %10 = "tensor.extract_slice"(%8) <{"static_offsets" = array, "static_sizes" = array, "static_strides" = array, "operandSegmentSizes" = array}> : (tensor<512xf32>) -> tensor<510xf32> -// CHECK-GENERIC-NEXT: %11 = "csl_stencil.access"(%3) <{"offset" = #stencil.index<[0, 1]>, "offset_mapping" = #stencil.index<[0, 1]>}> : (tensor<4x510xf32>) -> tensor<510xf32> -// CHECK-GENERIC-NEXT: %12 = "csl_stencil.access"(%3) <{"offset" = #stencil.index<[0, -1]>, "offset_mapping" = #stencil.index<[0, 1]>}> : (tensor<4x510xf32>) -> tensor<510xf32> -// CHECK-GENERIC-NEXT: %13 = "arith.addf"(%12, %11) <{"fastmath" = #arith.fastmath}> : (tensor<510xf32>, tensor<510xf32>) -> tensor<510xf32> -// CHECK-GENERIC-NEXT: %14 = "arith.addf"(%13, %10) <{"fastmath" = #arith.fastmath}> : (tensor<510xf32>, tensor<510xf32>) -> tensor<510xf32> -// CHECK-GENERIC-NEXT: %15 = "arith.addf"(%14, %9) <{"fastmath" = #arith.fastmath}> : (tensor<510xf32>, tensor<510xf32>) -> tensor<510xf32> -// CHECK-GENERIC-NEXT: %16 = "arith.addf"(%15, %6) <{"fastmath" = #arith.fastmath}> : (tensor<510xf32>, tensor<510xf32>) -> tensor<510xf32> -// CHECK-GENERIC-NEXT: %17 = "arith.addf"(%16, %5) <{"fastmath" = #arith.fastmath}> : (tensor<510xf32>, tensor<510xf32>) -> tensor<510xf32> +// CHECK-GENERIC-NEXT: %4 = "arith.constant"() <{value = 1.666600e-01 : f32}> : () -> f32 +// CHECK-GENERIC-NEXT: %5 = "csl_stencil.access"(%3) <{offset = #stencil.index<[1, 0]>, offset_mapping = #stencil.index<[0, 1]>}> : (tensor<4x510xf32>) -> tensor<510xf32> +// CHECK-GENERIC-NEXT: %6 = "csl_stencil.access"(%3) <{offset = #stencil.index<[-1, 0]>, offset_mapping = #stencil.index<[0, 1]>}> : (tensor<4x510xf32>) -> tensor<510xf32> +// CHECK-GENERIC-NEXT: %7 = "csl_stencil.access"(%2) <{offset = #stencil.index<[0, 0]>, offset_mapping = #stencil.index<[0, 1]>}> : (!stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>>) -> tensor<512xf32> +// CHECK-GENERIC-NEXT: %8 = "csl_stencil.access"(%2) <{offset = #stencil.index<[0, 0]>, offset_mapping = #stencil.index<[0, 1]>}> : (!stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>>) -> tensor<512xf32> +// CHECK-GENERIC-NEXT: %9 = "tensor.extract_slice"(%7) <{static_offsets = array, static_sizes = array, static_strides = array, operandSegmentSizes = array}> : (tensor<512xf32>) -> tensor<510xf32> +// CHECK-GENERIC-NEXT: %10 = "tensor.extract_slice"(%8) <{static_offsets = array, static_sizes = array, static_strides = array, operandSegmentSizes = array}> : (tensor<512xf32>) -> tensor<510xf32> +// CHECK-GENERIC-NEXT: %11 = "csl_stencil.access"(%3) <{offset = #stencil.index<[0, 1]>, offset_mapping = #stencil.index<[0, 1]>}> : (tensor<4x510xf32>) -> tensor<510xf32> +// CHECK-GENERIC-NEXT: %12 = "csl_stencil.access"(%3) <{offset = #stencil.index<[0, -1]>, offset_mapping = #stencil.index<[0, 1]>}> : (tensor<4x510xf32>) -> tensor<510xf32> +// CHECK-GENERIC-NEXT: %13 = "arith.addf"(%12, %11) <{fastmath = #arith.fastmath}> : (tensor<510xf32>, tensor<510xf32>) -> tensor<510xf32> +// CHECK-GENERIC-NEXT: %14 = "arith.addf"(%13, %10) <{fastmath = #arith.fastmath}> : (tensor<510xf32>, tensor<510xf32>) -> tensor<510xf32> +// CHECK-GENERIC-NEXT: %15 = "arith.addf"(%14, %9) <{fastmath = #arith.fastmath}> : (tensor<510xf32>, tensor<510xf32>) -> tensor<510xf32> +// CHECK-GENERIC-NEXT: %16 = "arith.addf"(%15, %6) <{fastmath = #arith.fastmath}> : (tensor<510xf32>, tensor<510xf32>) -> tensor<510xf32> +// CHECK-GENERIC-NEXT: %17 = "arith.addf"(%16, %5) <{fastmath = #arith.fastmath}> : (tensor<510xf32>, tensor<510xf32>) -> tensor<510xf32> // CHECK-GENERIC-NEXT: %18 = "tensor.empty"() : () -> tensor<510xf32> -// CHECK-GENERIC-NEXT: %19 = "linalg.fill"(%4, %18) <{"operandSegmentSizes" = array}> ({ +// CHECK-GENERIC-NEXT: %19 = "linalg.fill"(%4, %18) <{operandSegmentSizes = array}> ({ // CHECK-GENERIC-NEXT: ^2(%20 : f32, %21 : f32): // CHECK-GENERIC-NEXT: "linalg.yield"(%20) : (f32) -> () // CHECK-GENERIC-NEXT: }) : (f32, tensor<510xf32>) -> tensor<510xf32> -// CHECK-GENERIC-NEXT: %22 = "arith.mulf"(%17, %19) <{"fastmath" = #arith.fastmath}> : (tensor<510xf32>, tensor<510xf32>) -> tensor<510xf32> +// CHECK-GENERIC-NEXT: %22 = "arith.mulf"(%17, %19) <{fastmath = #arith.fastmath}> : (tensor<510xf32>, tensor<510xf32>) -> tensor<510xf32> // CHECK-GENERIC-NEXT: "stencil.return"(%22) : (tensor<510xf32>) -> () // CHECK-GENERIC-NEXT: }) : (!stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>>, tensor<4x510xf32>) -> !stencil.temp<[0,1]x[0,1]xtensor<510xf32>> -// CHECK-GENERIC-NEXT: "stencil.store"(%1, %b) {"bounds" = #stencil.bounds<[0, 0], [1, 1]>} : (!stencil.temp<[0,1]x[0,1]xtensor<510xf32>>, !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>) -> () +// CHECK-GENERIC-NEXT: "stencil.store"(%1, %b) {bounds = #stencil.bounds<[0, 0], [1, 1]>} : (!stencil.temp<[0,1]x[0,1]xtensor<510xf32>>, !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>) -> () // CHECK-GENERIC-NEXT: "func.return"() : () -> () // CHECK-GENERIC-NEXT: }) : () -> () // CHECK-GENERIC-NEXT: }) : () -> () @@ -142,7 +142,7 @@ builtin.module { // CHECK-NEXT: func.func @gauss_seidel(%a : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>, %b : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>) { // CHECK-NEXT: %0 = stencil.load %a : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>> -> !stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>> // CHECK-NEXT: %1 = tensor.empty() : tensor<510xf32> -// CHECK-NEXT: %2 = csl_stencil.apply(%0 : !stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>>, %1 : tensor<510xf32>) -> (!stencil.temp<[0,1]x[0,1]xtensor<510xf32>>) <{"num_chunks" = 2 : i64, "topo" = #dmp.topo<1022x510>, "swaps" = [#csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange], "operandSegmentSizes" = array}> ({ +// CHECK-NEXT: %2 = csl_stencil.apply(%0 : !stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>>, %1 : tensor<510xf32>) -> (!stencil.temp<[0,1]x[0,1]xtensor<510xf32>>) <{num_chunks = 2 : i64, topo = #dmp.topo<1022x510>, swaps = [#csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange], operandSegmentSizes = array}> ({ // CHECK-NEXT: ^0(%recv : tensor<4x255xf32>, %offset : index, %iter_arg : tensor<510xf32>): // CHECK-NEXT: %3 = csl_stencil.access %recv[1, 0] : tensor<4x255xf32> // CHECK-NEXT: %4 = csl_stencil.access %recv[-1, 0] : tensor<4x255xf32> @@ -151,14 +151,14 @@ builtin.module { // CHECK-NEXT: %7 = arith.addf %3, %4 : tensor<255xf32> // CHECK-NEXT: %8 = arith.addf %7, %5 : tensor<255xf32> // CHECK-NEXT: %9 = arith.addf %8, %6 : tensor<255xf32> -// CHECK-NEXT: %10 = "tensor.insert_slice"(%9, %iter_arg, %offset) <{"static_offsets" = array, "static_sizes" = array, "static_strides" = array, "operandSegmentSizes" = array}> : (tensor<255xf32>, tensor<510xf32>, index) -> tensor<510xf32> +// CHECK-NEXT: %10 = "tensor.insert_slice"(%9, %iter_arg, %offset) <{static_offsets = array, static_sizes = array, static_strides = array, operandSegmentSizes = array}> : (tensor<255xf32>, tensor<510xf32>, index) -> tensor<510xf32> // CHECK-NEXT: csl_stencil.yield %10 : tensor<510xf32> // CHECK-NEXT: }, { // CHECK-NEXT: ^1(%11 : !stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>>, %rcv : tensor<510xf32>): // CHECK-NEXT: %12 = csl_stencil.access %11[0, 0] : !stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>> // CHECK-NEXT: %13 = csl_stencil.access %11[0, 0] : !stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>> -// CHECK-NEXT: %14 = "tensor.extract_slice"(%12) <{"static_offsets" = array, "static_sizes" = array, "static_strides" = array, "operandSegmentSizes" = array}> : (tensor<512xf32>) -> tensor<510xf32> -// CHECK-NEXT: %15 = "tensor.extract_slice"(%13) <{"static_offsets" = array, "static_sizes" = array, "static_strides" = array, "operandSegmentSizes" = array}> : (tensor<512xf32>) -> tensor<510xf32> +// CHECK-NEXT: %14 = "tensor.extract_slice"(%12) <{static_offsets = array, static_sizes = array, static_strides = array, operandSegmentSizes = array}> : (tensor<512xf32>) -> tensor<510xf32> +// CHECK-NEXT: %15 = "tensor.extract_slice"(%13) <{static_offsets = array, static_sizes = array, static_strides = array, operandSegmentSizes = array}> : (tensor<512xf32>) -> tensor<510xf32> // CHECK-NEXT: %16 = arith.addf %rcv, %14 : tensor<510xf32> // CHECK-NEXT: %17 = arith.addf %16, %15 : tensor<510xf32> // CHECK-NEXT: %18 = arith.constant 1.666600e-01 : f32 @@ -173,39 +173,39 @@ builtin.module { // CHECK-NEXT: } // CHECK-GENERIC: "builtin.module"() ({ -// CHECK-GENERIC-NEXT: "func.func"() <{"sym_name" = "gauss_seidel", "function_type" = (!stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>, !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>) -> ()}> ({ +// CHECK-GENERIC-NEXT: "func.func"() <{sym_name = "gauss_seidel", function_type = (!stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>, !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>) -> ()}> ({ // CHECK-GENERIC-NEXT: ^0(%a : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>, %b : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>): // CHECK-GENERIC-NEXT: %0 = "stencil.load"(%a) : (!stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>) -> !stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>> // CHECK-GENERIC-NEXT: %1 = "tensor.empty"() : () -> tensor<510xf32> -// CHECK-GENERIC-NEXT: %2 = "csl_stencil.apply"(%0, %1) <{"num_chunks" = 2 : i64, "topo" = #dmp.topo<1022x510>, "swaps" = [#csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange], "operandSegmentSizes" = array}> ({ +// CHECK-GENERIC-NEXT: %2 = "csl_stencil.apply"(%0, %1) <{num_chunks = 2 : i64, topo = #dmp.topo<1022x510>, swaps = [#csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange], operandSegmentSizes = array}> ({ // CHECK-GENERIC-NEXT: ^1(%recv : tensor<4x255xf32>, %offset : index, %iter_arg : tensor<510xf32>): -// CHECK-GENERIC-NEXT: %3 = "csl_stencil.access"(%recv) <{"offset" = #stencil.index<[1, 0]>, "offset_mapping" = #stencil.index<[0, 1]>}> : (tensor<4x255xf32>) -> tensor<255xf32> -// CHECK-GENERIC-NEXT: %4 = "csl_stencil.access"(%recv) <{"offset" = #stencil.index<[-1, 0]>, "offset_mapping" = #stencil.index<[0, 1]>}> : (tensor<4x255xf32>) -> tensor<255xf32> -// CHECK-GENERIC-NEXT: %5 = "csl_stencil.access"(%recv) <{"offset" = #stencil.index<[0, 1]>, "offset_mapping" = #stencil.index<[0, 1]>}> : (tensor<4x255xf32>) -> tensor<255xf32> -// CHECK-GENERIC-NEXT: %6 = "csl_stencil.access"(%recv) <{"offset" = #stencil.index<[0, -1]>, "offset_mapping" = #stencil.index<[0, 1]>}> : (tensor<4x255xf32>) -> tensor<255xf32> -// CHECK-GENERIC-NEXT: %7 = "arith.addf"(%3, %4) <{"fastmath" = #arith.fastmath}> : (tensor<255xf32>, tensor<255xf32>) -> tensor<255xf32> -// CHECK-GENERIC-NEXT: %8 = "arith.addf"(%7, %5) <{"fastmath" = #arith.fastmath}> : (tensor<255xf32>, tensor<255xf32>) -> tensor<255xf32> -// CHECK-GENERIC-NEXT: %9 = "arith.addf"(%8, %6) <{"fastmath" = #arith.fastmath}> : (tensor<255xf32>, tensor<255xf32>) -> tensor<255xf32> -// CHECK-GENERIC-NEXT: %10 = "tensor.insert_slice"(%9, %iter_arg, %offset) <{"static_offsets" = array, "static_sizes" = array, "static_strides" = array, "operandSegmentSizes" = array}> : (tensor<255xf32>, tensor<510xf32>, index) -> tensor<510xf32> +// CHECK-GENERIC-NEXT: %3 = "csl_stencil.access"(%recv) <{offset = #stencil.index<[1, 0]>, offset_mapping = #stencil.index<[0, 1]>}> : (tensor<4x255xf32>) -> tensor<255xf32> +// CHECK-GENERIC-NEXT: %4 = "csl_stencil.access"(%recv) <{offset = #stencil.index<[-1, 0]>, offset_mapping = #stencil.index<[0, 1]>}> : (tensor<4x255xf32>) -> tensor<255xf32> +// CHECK-GENERIC-NEXT: %5 = "csl_stencil.access"(%recv) <{offset = #stencil.index<[0, 1]>, offset_mapping = #stencil.index<[0, 1]>}> : (tensor<4x255xf32>) -> tensor<255xf32> +// CHECK-GENERIC-NEXT: %6 = "csl_stencil.access"(%recv) <{offset = #stencil.index<[0, -1]>, offset_mapping = #stencil.index<[0, 1]>}> : (tensor<4x255xf32>) -> tensor<255xf32> +// CHECK-GENERIC-NEXT: %7 = "arith.addf"(%3, %4) <{fastmath = #arith.fastmath}> : (tensor<255xf32>, tensor<255xf32>) -> tensor<255xf32> +// CHECK-GENERIC-NEXT: %8 = "arith.addf"(%7, %5) <{fastmath = #arith.fastmath}> : (tensor<255xf32>, tensor<255xf32>) -> tensor<255xf32> +// CHECK-GENERIC-NEXT: %9 = "arith.addf"(%8, %6) <{fastmath = #arith.fastmath}> : (tensor<255xf32>, tensor<255xf32>) -> tensor<255xf32> +// CHECK-GENERIC-NEXT: %10 = "tensor.insert_slice"(%9, %iter_arg, %offset) <{static_offsets = array, static_sizes = array, static_strides = array, operandSegmentSizes = array}> : (tensor<255xf32>, tensor<510xf32>, index) -> tensor<510xf32> // CHECK-GENERIC-NEXT: "csl_stencil.yield"(%10) : (tensor<510xf32>) -> () // CHECK-GENERIC-NEXT: }, { // CHECK-GENERIC-NEXT: ^2(%11 : !stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>>, %rcv : tensor<510xf32>): -// CHECK-GENERIC-NEXT: %12 = "csl_stencil.access"(%11) <{"offset" = #stencil.index<[0, 0]>, "offset_mapping" = #stencil.index<[0, 1]>}> : (!stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>>) -> tensor<512xf32> -// CHECK-GENERIC-NEXT: %13 = "csl_stencil.access"(%11) <{"offset" = #stencil.index<[0, 0]>, "offset_mapping" = #stencil.index<[0, 1]>}> : (!stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>>) -> tensor<512xf32> -// CHECK-GENERIC-NEXT: %14 = "tensor.extract_slice"(%12) <{"static_offsets" = array, "static_sizes" = array, "static_strides" = array, "operandSegmentSizes" = array}> : (tensor<512xf32>) -> tensor<510xf32> -// CHECK-GENERIC-NEXT: %15 = "tensor.extract_slice"(%13) <{"static_offsets" = array, "static_sizes" = array, "static_strides" = array, "operandSegmentSizes" = array}> : (tensor<512xf32>) -> tensor<510xf32> -// CHECK-GENERIC-NEXT: %16 = "arith.addf"(%rcv, %14) <{"fastmath" = #arith.fastmath}> : (tensor<510xf32>, tensor<510xf32>) -> tensor<510xf32> -// CHECK-GENERIC-NEXT: %17 = "arith.addf"(%16, %15) <{"fastmath" = #arith.fastmath}> : (tensor<510xf32>, tensor<510xf32>) -> tensor<510xf32> -// CHECK-GENERIC-NEXT: %18 = "arith.constant"() <{"value" = 1.666600e-01 : f32}> : () -> f32 +// CHECK-GENERIC-NEXT: %12 = "csl_stencil.access"(%11) <{offset = #stencil.index<[0, 0]>, offset_mapping = #stencil.index<[0, 1]>}> : (!stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>>) -> tensor<512xf32> +// CHECK-GENERIC-NEXT: %13 = "csl_stencil.access"(%11) <{offset = #stencil.index<[0, 0]>, offset_mapping = #stencil.index<[0, 1]>}> : (!stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>>) -> tensor<512xf32> +// CHECK-GENERIC-NEXT: %14 = "tensor.extract_slice"(%12) <{static_offsets = array, static_sizes = array, static_strides = array, operandSegmentSizes = array}> : (tensor<512xf32>) -> tensor<510xf32> +// CHECK-GENERIC-NEXT: %15 = "tensor.extract_slice"(%13) <{static_offsets = array, static_sizes = array, static_strides = array, operandSegmentSizes = array}> : (tensor<512xf32>) -> tensor<510xf32> +// CHECK-GENERIC-NEXT: %16 = "arith.addf"(%rcv, %14) <{fastmath = #arith.fastmath}> : (tensor<510xf32>, tensor<510xf32>) -> tensor<510xf32> +// CHECK-GENERIC-NEXT: %17 = "arith.addf"(%16, %15) <{fastmath = #arith.fastmath}> : (tensor<510xf32>, tensor<510xf32>) -> tensor<510xf32> +// CHECK-GENERIC-NEXT: %18 = "arith.constant"() <{value = 1.666600e-01 : f32}> : () -> f32 // CHECK-GENERIC-NEXT: %19 = "tensor.empty"() : () -> tensor<510xf32> -// CHECK-GENERIC-NEXT: %20 = "linalg.fill"(%18, %19) <{"operandSegmentSizes" = array}> ({ +// CHECK-GENERIC-NEXT: %20 = "linalg.fill"(%18, %19) <{operandSegmentSizes = array}> ({ // CHECK-GENERIC-NEXT: ^3(%21 : f32, %22 : f32): // CHECK-GENERIC-NEXT: "linalg.yield"(%21) : (f32) -> () // CHECK-GENERIC-NEXT: }) : (f32, tensor<510xf32>) -> tensor<510xf32> -// CHECK-GENERIC-NEXT: %23 = "arith.mulf"(%17, %20) <{"fastmath" = #arith.fastmath}> : (tensor<510xf32>, tensor<510xf32>) -> tensor<510xf32> +// CHECK-GENERIC-NEXT: %23 = "arith.mulf"(%17, %20) <{fastmath = #arith.fastmath}> : (tensor<510xf32>, tensor<510xf32>) -> tensor<510xf32> // CHECK-GENERIC-NEXT: "csl_stencil.yield"(%23) : (tensor<510xf32>) -> () // CHECK-GENERIC-NEXT: }) : (!stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>>, tensor<510xf32>) -> !stencil.temp<[0,1]x[0,1]xtensor<510xf32>> -// CHECK-GENERIC-NEXT: "stencil.store"(%2, %b) {"bounds" = #stencil.bounds<[0, 0], [1, 1]>} : (!stencil.temp<[0,1]x[0,1]xtensor<510xf32>>, !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>) -> () +// CHECK-GENERIC-NEXT: "stencil.store"(%2, %b) {bounds = #stencil.bounds<[0, 0], [1, 1]>} : (!stencil.temp<[0,1]x[0,1]xtensor<510xf32>>, !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>) -> () // CHECK-GENERIC-NEXT: "func.return"() : () -> () // CHECK-GENERIC-NEXT: }) : () -> () // CHECK-GENERIC-NEXT: }) : () -> () @@ -245,7 +245,7 @@ builtin.module { //CHECK: builtin.module { //CHECK-NEXT: func.func @bufferized_stencil(%a : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>, %b : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>) { //CHECK-NEXT: %0 = tensor.empty() : tensor<510xf32> -//CHECK-NEXT: csl_stencil.apply(%a : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>, %0 : tensor<510xf32>) outs (%b : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>) <{"swaps" = [#csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange], "topo" = #dmp.topo<1022x510>, "num_chunks" = 2 : i64, "bounds" = #stencil.bounds<[0, 0], [1, 1]>, "operandSegmentSizes" = array}> ({ +//CHECK-NEXT: csl_stencil.apply(%a : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>, %0 : tensor<510xf32>) outs (%b : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>) <{swaps = [#csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange], topo = #dmp.topo<1022x510>, num_chunks = 2 : i64, bounds = #stencil.bounds<[0, 0], [1, 1]>, operandSegmentSizes = array}> ({ //CHECK-NEXT: ^0(%1 : tensor<4x255xf32>, %2 : index, %3 : tensor<510xf32>): //CHECK-NEXT: %4 = csl_stencil.access %1[1, 0] : tensor<4x255xf32> //CHECK-NEXT: %5 = csl_stencil.access %1[-1, 0] : tensor<4x255xf32> @@ -254,14 +254,14 @@ builtin.module { //CHECK-NEXT: %8 = arith.addf %7, %6 : tensor<255xf32> //CHECK-NEXT: %9 = arith.addf %8, %5 : tensor<255xf32> //CHECK-NEXT: %10 = arith.addf %9, %4 : tensor<255xf32> -//CHECK-NEXT: %11 = "tensor.insert_slice"(%10, %3, %2) <{"static_offsets" = array, "static_sizes" = array, "static_strides" = array, "operandSegmentSizes" = array}> : (tensor<255xf32>, tensor<510xf32>, index) -> tensor<510xf32> +//CHECK-NEXT: %11 = "tensor.insert_slice"(%10, %3, %2) <{static_offsets = array, static_sizes = array, static_strides = array, operandSegmentSizes = array}> : (tensor<255xf32>, tensor<510xf32>, index) -> tensor<510xf32> //CHECK-NEXT: csl_stencil.yield %11 : tensor<510xf32> //CHECK-NEXT: }, { //CHECK-NEXT: ^1(%12 : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>, %13 : tensor<510xf32>): //CHECK-NEXT: %14 = csl_stencil.access %12[0, 0] : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>> //CHECK-NEXT: %15 = arith.constant dense<1.666600e-01> : tensor<510xf32> -//CHECK-NEXT: %16 = "tensor.extract_slice"(%14) <{"static_offsets" = array, "static_sizes" = array, "static_strides" = array, "operandSegmentSizes" = array}> : (tensor<512xf32>) -> tensor<510xf32> -//CHECK-NEXT: %17 = "tensor.extract_slice"(%14) <{"static_offsets" = array, "static_sizes" = array, "static_strides" = array, "operandSegmentSizes" = array}> : (tensor<512xf32>) -> tensor<510xf32> +//CHECK-NEXT: %16 = "tensor.extract_slice"(%14) <{static_offsets = array, static_sizes = array, static_strides = array, operandSegmentSizes = array}> : (tensor<512xf32>) -> tensor<510xf32> +//CHECK-NEXT: %17 = "tensor.extract_slice"(%14) <{static_offsets = array, static_sizes = array, static_strides = array, operandSegmentSizes = array}> : (tensor<512xf32>) -> tensor<510xf32> //CHECK-NEXT: %18 = arith.addf %13, %17 : tensor<510xf32> //CHECK-NEXT: %19 = arith.addf %18, %16 : tensor<510xf32> //CHECK-NEXT: %20 = arith.mulf %19, %15 : tensor<510xf32> @@ -272,29 +272,29 @@ builtin.module { //CHECK-NEXT: } // CHECK-GENERIC: "builtin.module"() ({ -// CHECK-GENERIC-NEXT: "func.func"() <{"sym_name" = "bufferized_stencil", "function_type" = (!stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>, !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>) -> ()}> ({ +// CHECK-GENERIC-NEXT: "func.func"() <{sym_name = "bufferized_stencil", function_type = (!stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>, !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>) -> ()}> ({ // CHECK-GENERIC-NEXT: ^0(%a : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>, %b : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>): // CHECK-GENERIC-NEXT: %0 = "tensor.empty"() : () -> tensor<510xf32> -// CHECK-GENERIC-NEXT: "csl_stencil.apply"(%a, %0, %b) <{"swaps" = [#csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange], "topo" = #dmp.topo<1022x510>, "num_chunks" = 2 : i64, "bounds" = #stencil.bounds<[0, 0], [1, 1]>, "operandSegmentSizes" = array}> ({ +// CHECK-GENERIC-NEXT: "csl_stencil.apply"(%a, %0, %b) <{swaps = [#csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange], topo = #dmp.topo<1022x510>, num_chunks = 2 : i64, bounds = #stencil.bounds<[0, 0], [1, 1]>, operandSegmentSizes = array}> ({ // CHECK-GENERIC-NEXT: ^1(%1 : tensor<4x255xf32>, %2 : index, %3 : tensor<510xf32>): -// CHECK-GENERIC-NEXT: %4 = "csl_stencil.access"(%1) <{"offset" = #stencil.index<[1, 0]>, "offset_mapping" = #stencil.index<[0, 1]>}> : (tensor<4x255xf32>) -> tensor<255xf32> -// CHECK-GENERIC-NEXT: %5 = "csl_stencil.access"(%1) <{"offset" = #stencil.index<[-1, 0]>, "offset_mapping" = #stencil.index<[0, 1]>}> : (tensor<4x255xf32>) -> tensor<255xf32> -// CHECK-GENERIC-NEXT: %6 = "csl_stencil.access"(%1) <{"offset" = #stencil.index<[0, 1]>, "offset_mapping" = #stencil.index<[0, 1]>}> : (tensor<4x255xf32>) -> tensor<255xf32> -// CHECK-GENERIC-NEXT: %7 = "csl_stencil.access"(%1) <{"offset" = #stencil.index<[0, -1]>, "offset_mapping" = #stencil.index<[0, 1]>}> : (tensor<4x255xf32>) -> tensor<255xf32> -// CHECK-GENERIC-NEXT: %8 = "arith.addf"(%7, %6) <{"fastmath" = #arith.fastmath}> : (tensor<255xf32>, tensor<255xf32>) -> tensor<255xf32> -// CHECK-GENERIC-NEXT: %9 = "arith.addf"(%8, %5) <{"fastmath" = #arith.fastmath}> : (tensor<255xf32>, tensor<255xf32>) -> tensor<255xf32> -// CHECK-GENERIC-NEXT: %10 = "arith.addf"(%9, %4) <{"fastmath" = #arith.fastmath}> : (tensor<255xf32>, tensor<255xf32>) -> tensor<255xf32> -// CHECK-GENERIC-NEXT: %11 = "tensor.insert_slice"(%10, %3, %2) <{"static_offsets" = array, "static_sizes" = array, "static_strides" = array, "operandSegmentSizes" = array}> : (tensor<255xf32>, tensor<510xf32>, index) -> tensor<510xf32> +// CHECK-GENERIC-NEXT: %4 = "csl_stencil.access"(%1) <{offset = #stencil.index<[1, 0]>, offset_mapping = #stencil.index<[0, 1]>}> : (tensor<4x255xf32>) -> tensor<255xf32> +// CHECK-GENERIC-NEXT: %5 = "csl_stencil.access"(%1) <{offset = #stencil.index<[-1, 0]>, offset_mapping = #stencil.index<[0, 1]>}> : (tensor<4x255xf32>) -> tensor<255xf32> +// CHECK-GENERIC-NEXT: %6 = "csl_stencil.access"(%1) <{offset = #stencil.index<[0, 1]>, offset_mapping = #stencil.index<[0, 1]>}> : (tensor<4x255xf32>) -> tensor<255xf32> +// CHECK-GENERIC-NEXT: %7 = "csl_stencil.access"(%1) <{offset = #stencil.index<[0, -1]>, offset_mapping = #stencil.index<[0, 1]>}> : (tensor<4x255xf32>) -> tensor<255xf32> +// CHECK-GENERIC-NEXT: %8 = "arith.addf"(%7, %6) <{fastmath = #arith.fastmath}> : (tensor<255xf32>, tensor<255xf32>) -> tensor<255xf32> +// CHECK-GENERIC-NEXT: %9 = "arith.addf"(%8, %5) <{fastmath = #arith.fastmath}> : (tensor<255xf32>, tensor<255xf32>) -> tensor<255xf32> +// CHECK-GENERIC-NEXT: %10 = "arith.addf"(%9, %4) <{fastmath = #arith.fastmath}> : (tensor<255xf32>, tensor<255xf32>) -> tensor<255xf32> +// CHECK-GENERIC-NEXT: %11 = "tensor.insert_slice"(%10, %3, %2) <{static_offsets = array, static_sizes = array, static_strides = array, operandSegmentSizes = array}> : (tensor<255xf32>, tensor<510xf32>, index) -> tensor<510xf32> // CHECK-GENERIC-NEXT: "csl_stencil.yield"(%11) : (tensor<510xf32>) -> () // CHECK-GENERIC-NEXT: }, { // CHECK-GENERIC-NEXT: ^2(%12 : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>, %13 : tensor<510xf32>): -// CHECK-GENERIC-NEXT: %14 = "csl_stencil.access"(%12) <{"offset" = #stencil.index<[0, 0]>, "offset_mapping" = #stencil.index<[0, 1]>}> : (!stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>) -> tensor<512xf32> -// CHECK-GENERIC-NEXT: %15 = "arith.constant"() <{"value" = dense<1.666600e-01> : tensor<510xf32>}> : () -> tensor<510xf32> -// CHECK-GENERIC-NEXT: %16 = "tensor.extract_slice"(%14) <{"static_offsets" = array, "static_sizes" = array, "static_strides" = array, "operandSegmentSizes" = array}> : (tensor<512xf32>) -> tensor<510xf32> -// CHECK-GENERIC-NEXT: %17 = "tensor.extract_slice"(%14) <{"static_offsets" = array, "static_sizes" = array, "static_strides" = array, "operandSegmentSizes" = array}> : (tensor<512xf32>) -> tensor<510xf32> -// CHECK-GENERIC-NEXT: %18 = "arith.addf"(%13, %17) <{"fastmath" = #arith.fastmath}> : (tensor<510xf32>, tensor<510xf32>) -> tensor<510xf32> -// CHECK-GENERIC-NEXT: %19 = "arith.addf"(%18, %16) <{"fastmath" = #arith.fastmath}> : (tensor<510xf32>, tensor<510xf32>) -> tensor<510xf32> -// CHECK-GENERIC-NEXT: %20 = "arith.mulf"(%19, %15) <{"fastmath" = #arith.fastmath}> : (tensor<510xf32>, tensor<510xf32>) -> tensor<510xf32> +// CHECK-GENERIC-NEXT: %14 = "csl_stencil.access"(%12) <{offset = #stencil.index<[0, 0]>, offset_mapping = #stencil.index<[0, 1]>}> : (!stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>) -> tensor<512xf32> +// CHECK-GENERIC-NEXT: %15 = "arith.constant"() <{value = dense<1.666600e-01> : tensor<510xf32>}> : () -> tensor<510xf32> +// CHECK-GENERIC-NEXT: %16 = "tensor.extract_slice"(%14) <{static_offsets = array, static_sizes = array, static_strides = array, operandSegmentSizes = array}> : (tensor<512xf32>) -> tensor<510xf32> +// CHECK-GENERIC-NEXT: %17 = "tensor.extract_slice"(%14) <{static_offsets = array, static_sizes = array, static_strides = array, operandSegmentSizes = array}> : (tensor<512xf32>) -> tensor<510xf32> +// CHECK-GENERIC-NEXT: %18 = "arith.addf"(%13, %17) <{fastmath = #arith.fastmath}> : (tensor<510xf32>, tensor<510xf32>) -> tensor<510xf32> +// CHECK-GENERIC-NEXT: %19 = "arith.addf"(%18, %16) <{fastmath = #arith.fastmath}> : (tensor<510xf32>, tensor<510xf32>) -> tensor<510xf32> +// CHECK-GENERIC-NEXT: %20 = "arith.mulf"(%19, %15) <{fastmath = #arith.fastmath}> : (tensor<510xf32>, tensor<510xf32>) -> tensor<510xf32> // CHECK-GENERIC-NEXT: "csl_stencil.yield"(%20) : (tensor<510xf32>) -> () // CHECK-GENERIC-NEXT: }) : (!stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>, tensor<510xf32>, !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>) -> () // CHECK-GENERIC-NEXT: "func.return"() : () -> () diff --git a/tests/filecheck/dialects/csl/csl-wrapper-ops.mlir b/tests/filecheck/dialects/csl/csl-wrapper-ops.mlir index 6cf85c5bdb..fc604e9635 100644 --- a/tests/filecheck/dialects/csl/csl-wrapper-ops.mlir +++ b/tests/filecheck/dialects/csl/csl-wrapper-ops.mlir @@ -40,14 +40,14 @@ builtin.module { // CHECK: builtin.module { -// CHECK-NEXT: "csl_wrapper.module"() <{"width" = 10 : i16, "height" = 10 : i16, "params" = [#csl_wrapper.param<"z_dim" default=4 : i16>, #csl_wrapper.param<"pattern" : i16>]}> ({ +// CHECK-NEXT: "csl_wrapper.module"() <{width = 10 : i16, height = 10 : i16, params = [#csl_wrapper.param<"z_dim" default=4 : i16>, #csl_wrapper.param<"pattern" : i16>]}> ({ // CHECK-NEXT: ^0(%x : i16, %y : i16, %width : i16, %height : i16, %z_dim : i16, %pattern : i16): // CHECK-NEXT: %0 = arith.constant 0 : i16 // CHECK-NEXT: %1 = "csl.get_color"(%0) : (i16) -> !csl.color -// CHECK-NEXT: %routes = "csl_wrapper.import"(%pattern, %width, %height) <{"module" = "routes.csl", "fields" = ["pattern", "peWidth", "peHeight"]}> : (i16, i16, i16) -> !csl.imported_module -// CHECK-NEXT: %memcpy = "csl_wrapper.import"(%width, %height, %1) <{"module" = "", "fields" = ["width", "height", "LAUNCH"]}> : (i16, i16, !csl.color) -> !csl.imported_module -// CHECK-NEXT: %compute_all_routes = "csl.member_call"(%routes, %x, %y, %height, %width, %pattern) <{"field" = "computeAllRoutes"}> : (!csl.imported_module, i16, i16, i16, i16, i16) -> !csl.comptime_struct -// CHECK-NEXT: %memcpy_params = "csl.member_call"(%memcpy, %x) <{"field" = "get_params"}> : (!csl.imported_module, i16) -> !csl.comptime_struct +// CHECK-NEXT: %routes = "csl_wrapper.import"(%pattern, %width, %height) <{module = "routes.csl", fields = ["pattern", "peWidth", "peHeight"]}> : (i16, i16, i16) -> !csl.imported_module +// CHECK-NEXT: %memcpy = "csl_wrapper.import"(%width, %height, %1) <{module = "", fields = ["width", "height", "LAUNCH"]}> : (i16, i16, !csl.color) -> !csl.imported_module +// CHECK-NEXT: %compute_all_routes = "csl.member_call"(%routes, %x, %y, %height, %width, %pattern) <{field = "computeAllRoutes"}> : (!csl.imported_module, i16, i16, i16, i16, i16) -> !csl.comptime_struct +// CHECK-NEXT: %memcpy_params = "csl.member_call"(%memcpy, %x) <{field = "get_params"}> : (!csl.imported_module, i16) -> !csl.comptime_struct // CHECK-NEXT: %2 = arith.constant 1 : i16 // CHECK-NEXT: %3 = arith.minsi %pattern, %2 : i16 // CHECK-NEXT: %4 = arith.minsi %width, %x : i16 @@ -59,43 +59,43 @@ builtin.module { // CHECK-NEXT: %10 = arith.ori %6, %7 : i1 // CHECK-NEXT: %11 = arith.ori %10, %8 : i1 // CHECK-NEXT: %is_border_region_pe = arith.ori %11, %9 : i1 -// CHECK-NEXT: "csl_wrapper.yield"(%memcpy_params, %compute_all_routes, %is_border_region_pe) <{"fields" = ["memcpy_params", "stencil_comms_params", "isBorderRegionPE"]}> : (!csl.comptime_struct, !csl.comptime_struct, i1) -> () +// CHECK-NEXT: "csl_wrapper.yield"(%memcpy_params, %compute_all_routes, %is_border_region_pe) <{fields = ["memcpy_params", "stencil_comms_params", "isBorderRegionPE"]}> : (!csl.comptime_struct, !csl.comptime_struct, i1) -> () // CHECK-NEXT: }, { // CHECK-NEXT: ^1(%width_1 : i16, %height_1 : i16, %z_dim_1 : i16, %pattern_1 : i16, %memcpy_params_1 : !csl.comptime_struct, %stencil_comms_params : !csl.comptime_struct, %is_border_region_pe_1 : i1): // CHECK-NEXT: func.func @gauss_seidel() { // CHECK-NEXT: func.return // CHECK-NEXT: } -// CHECK-NEXT: "csl_wrapper.yield"() <{"fields" = []}> : () -> () +// CHECK-NEXT: "csl_wrapper.yield"() <{fields = []}> : () -> () // CHECK-NEXT: }) : () -> () // CHECK-NEXT: } // CHECK-GENERIC: "builtin.module"() ({ -// CHECK-GENERIC-NEXT: "csl_wrapper.module"() <{"width" = 10 : i16, "height" = 10 : i16, "params" = [#csl_wrapper.param<"z_dim" default=4 : i16>, #csl_wrapper.param<"pattern" : i16>]}> ({ +// CHECK-GENERIC-NEXT: "csl_wrapper.module"() <{width = 10 : i16, height = 10 : i16, params = [#csl_wrapper.param<"z_dim" default=4 : i16>, #csl_wrapper.param<"pattern" : i16>]}> ({ // CHECK-GENERIC-NEXT: ^0(%x : i16, %y : i16, %width : i16, %height : i16, %z_dim : i16, %pattern : i16): -// CHECK-GENERIC-NEXT: %0 = "arith.constant"() <{"value" = 0 : i16}> : () -> i16 +// CHECK-GENERIC-NEXT: %0 = "arith.constant"() <{value = 0 : i16}> : () -> i16 // CHECK-GENERIC-NEXT: %1 = "csl.get_color"(%0) : (i16) -> !csl.color -// CHECK-GENERIC-NEXT: %routes = "csl_wrapper.import"(%pattern, %width, %height) <{"module" = "routes.csl", "fields" = ["pattern", "peWidth", "peHeight"]}> : (i16, i16, i16) -> !csl.imported_module -// CHECK-GENERIC-NEXT: %memcpy = "csl_wrapper.import"(%width, %height, %1) <{"module" = "", "fields" = ["width", "height", "LAUNCH"]}> : (i16, i16, !csl.color) -> !csl.imported_module -// CHECK-GENERIC-NEXT: %compute_all_routes = "csl.member_call"(%routes, %x, %y, %height, %width, %pattern) <{"field" = "computeAllRoutes"}> : (!csl.imported_module, i16, i16, i16, i16, i16) -> !csl.comptime_struct -// CHECK-GENERIC-NEXT: %memcpy_params = "csl.member_call"(%memcpy, %x) <{"field" = "get_params"}> : (!csl.imported_module, i16) -> !csl.comptime_struct -// CHECK-GENERIC-NEXT: %2 = "arith.constant"() <{"value" = 1 : i16}> : () -> i16 +// CHECK-GENERIC-NEXT: %routes = "csl_wrapper.import"(%pattern, %width, %height) <{module = "routes.csl", fields = ["pattern", "peWidth", "peHeight"]}> : (i16, i16, i16) -> !csl.imported_module +// CHECK-GENERIC-NEXT: %memcpy = "csl_wrapper.import"(%width, %height, %1) <{module = "", fields = ["width", "height", "LAUNCH"]}> : (i16, i16, !csl.color) -> !csl.imported_module +// CHECK-GENERIC-NEXT: %compute_all_routes = "csl.member_call"(%routes, %x, %y, %height, %width, %pattern) <{field = "computeAllRoutes"}> : (!csl.imported_module, i16, i16, i16, i16, i16) -> !csl.comptime_struct +// CHECK-GENERIC-NEXT: %memcpy_params = "csl.member_call"(%memcpy, %x) <{field = "get_params"}> : (!csl.imported_module, i16) -> !csl.comptime_struct +// CHECK-GENERIC-NEXT: %2 = "arith.constant"() <{value = 1 : i16}> : () -> i16 // CHECK-GENERIC-NEXT: %3 = "arith.minsi"(%pattern, %2) : (i16, i16) -> i16 // CHECK-GENERIC-NEXT: %4 = "arith.minsi"(%width, %x) : (i16, i16) -> i16 // CHECK-GENERIC-NEXT: %5 = "arith.minsi"(%height, %y) : (i16, i16) -> i16 -// CHECK-GENERIC-NEXT: %6 = "arith.cmpi"(%x, %3) <{"predicate" = 2 : i64}> : (i16, i16) -> i1 -// CHECK-GENERIC-NEXT: %7 = "arith.cmpi"(%y, %3) <{"predicate" = 2 : i64}> : (i16, i16) -> i1 -// CHECK-GENERIC-NEXT: %8 = "arith.cmpi"(%4, %pattern) <{"predicate" = 2 : i64}> : (i16, i16) -> i1 -// CHECK-GENERIC-NEXT: %9 = "arith.cmpi"(%5, %pattern) <{"predicate" = 2 : i64}> : (i16, i16) -> i1 +// CHECK-GENERIC-NEXT: %6 = "arith.cmpi"(%x, %3) <{predicate = 2 : i64}> : (i16, i16) -> i1 +// CHECK-GENERIC-NEXT: %7 = "arith.cmpi"(%y, %3) <{predicate = 2 : i64}> : (i16, i16) -> i1 +// CHECK-GENERIC-NEXT: %8 = "arith.cmpi"(%4, %pattern) <{predicate = 2 : i64}> : (i16, i16) -> i1 +// CHECK-GENERIC-NEXT: %9 = "arith.cmpi"(%5, %pattern) <{predicate = 2 : i64}> : (i16, i16) -> i1 // CHECK-GENERIC-NEXT: %10 = "arith.ori"(%6, %7) : (i1, i1) -> i1 // CHECK-GENERIC-NEXT: %11 = "arith.ori"(%10, %8) : (i1, i1) -> i1 // CHECK-GENERIC-NEXT: %is_border_region_pe = "arith.ori"(%11, %9) : (i1, i1) -> i1 -// CHECK-GENERIC-NEXT: "csl_wrapper.yield"(%memcpy_params, %compute_all_routes, %is_border_region_pe) <{"fields" = ["memcpy_params", "stencil_comms_params", "isBorderRegionPE"]}> : (!csl.comptime_struct, !csl.comptime_struct, i1) -> () +// CHECK-GENERIC-NEXT: "csl_wrapper.yield"(%memcpy_params, %compute_all_routes, %is_border_region_pe) <{fields = ["memcpy_params", "stencil_comms_params", "isBorderRegionPE"]}> : (!csl.comptime_struct, !csl.comptime_struct, i1) -> () // CHECK-GENERIC-NEXT: }, { // CHECK-GENERIC-NEXT: ^1(%width_1 : i16, %height_1 : i16, %z_dim_1 : i16, %pattern_1 : i16, %memcpy_params_1 : !csl.comptime_struct, %stencil_comms_params : !csl.comptime_struct, %is_border_region_pe_1 : i1): -// CHECK-GENERIC-NEXT: "func.func"() <{"sym_name" = "gauss_seidel", "function_type" = () -> ()}> ({ +// CHECK-GENERIC-NEXT: "func.func"() <{sym_name = "gauss_seidel", function_type = () -> ()}> ({ // CHECK-GENERIC-NEXT: "func.return"() : () -> () // CHECK-GENERIC-NEXT: }) : () -> () -// CHECK-GENERIC-NEXT: "csl_wrapper.yield"() <{"fields" = []}> : () -> () +// CHECK-GENERIC-NEXT: "csl_wrapper.yield"() <{fields = []}> : () -> () // CHECK-GENERIC-NEXT: }) : () -> () // CHECK-GENERIC-NEXT: }) : () -> () diff --git a/tests/filecheck/dialects/csl/ops.mlir b/tests/filecheck/dialects/csl/ops.mlir index 94ca709154..2901df0ad4 100644 --- a/tests/filecheck/dialects/csl/ops.mlir +++ b/tests/filecheck/dialects/csl/ops.mlir @@ -339,45 +339,45 @@ csl.func @builtins() { // CHECK-NEXT: builtin.module { -// CHECK-NEXT: "csl.module"() <{"kind" = #csl}> ({ -// CHECK-NEXT: %thing = "csl.import_module"() <{"module" = ""}> : () -> !csl.imported_module +// CHECK-NEXT: "csl.module"() <{kind = #csl}> ({ +// CHECK-NEXT: %thing = "csl.import_module"() <{module = ""}> : () -> !csl.imported_module // CHECK-NEXT: csl.func @func_with_args(%arg1 : i32, %arg2 : i16) -> i32 { // CHECK-NEXT: csl.return %arg1 : i32 // CHECK-NEXT: } // CHECK-NEXT: %zero = arith.constant 0 : i32 // CHECK-NEXT: %c100 = arith.constant 100 : i32 // CHECK-NEXT: %zeros = "csl.constants"(%zero, %c100) : (i32, i32) -> memref -// CHECK-NEXT: %zeros2 = "csl.constants"(%zero, %c100) <{"is_const"}> : (i32, i32) -> memref -// CHECK-NEXT: csl.task @local_task() attributes {"kind" = #csl, "id" = 0 : ui5}{ +// CHECK-NEXT: %zeros2 = "csl.constants"(%zero, %c100) <{is_const}> : (i32, i32) -> memref +// CHECK-NEXT: csl.task @local_task() attributes {kind = #csl, id = 0 : ui5}{ // CHECK-NEXT: csl.return // CHECK-NEXT: } -// CHECK-NEXT: csl.task @data_task(%a : i32) attributes {"kind" = #csl, "id" = 1 : ui5}{ +// CHECK-NEXT: csl.task @data_task(%a : i32) attributes {kind = #csl, id = 1 : ui5}{ // CHECK-NEXT: csl.return // CHECK-NEXT: } -// CHECK-NEXT: csl.task @control_task() attributes {"kind" = #csl, "id" = 2 : ui6}{ +// CHECK-NEXT: csl.task @control_task() attributes {kind = #csl, id = 2 : ui6}{ // CHECK-NEXT: csl.return // CHECK-NEXT: } -// CHECK-NEXT: csl.task @control_task_args(%a_1 : i32) attributes {"kind" = #csl, "id" = 2 : ui6}{ +// CHECK-NEXT: csl.task @control_task_args(%a_1 : i32) attributes {kind = #csl, id = 2 : ui6}{ // CHECK-NEXT: csl.return // CHECK-NEXT: } -// CHECK-NEXT: csl.task @runtime_bound_local_task() attributes {"kind" = #csl}{ +// CHECK-NEXT: csl.task @runtime_bound_local_task() attributes {kind = #csl}{ // CHECK-NEXT: csl.return // CHECK-NEXT: } // CHECK-NEXT: csl.func @initialize() { // CHECK-NEXT: %lb, %ub = "test.op"() : () -> (i16, i16) -// CHECK-NEXT: "csl.member_call"(%thing, %lb, %ub) <{"field" = "some_func"}> : (!csl.imported_module, i16, i16) -> () -// CHECK-NEXT: %res = "csl.member_call"(%thing, %lb, %ub) <{"field" = "some_func"}> : (!csl.imported_module, i16, i16) -> i32 -// CHECK-NEXT: %0 = "csl.member_access"(%thing) <{"field" = "some_field"}> : (!csl.imported_module) -> !csl.comptime_struct +// CHECK-NEXT: "csl.member_call"(%thing, %lb, %ub) <{field = "some_func"}> : (!csl.imported_module, i16, i16) -> () +// CHECK-NEXT: %res = "csl.member_call"(%thing, %lb, %ub) <{field = "some_func"}> : (!csl.imported_module, i16, i16) -> i32 +// CHECK-NEXT: %0 = "csl.member_access"(%thing) <{field = "some_field"}> : (!csl.imported_module) -> !csl.comptime_struct // CHECK-NEXT: %single_const = "test.op"() : () -> !csl.ptr, #csl> // CHECK-NEXT: %single_var = "test.op"() : () -> !csl.ptr, #csl> // CHECK-NEXT: %many_const = "test.op"() : () -> !csl.ptr, #csl> // CHECK-NEXT: %many_var = "test.op"() : () -> !csl.ptr, #csl> // CHECK-NEXT: %col = "test.op"() : () -> !csl.color // CHECK-NEXT: %arg1_1, %arg2_1 = "test.op"() : () -> (i32, i16) -// CHECK-NEXT: %call_res = "csl.call"(%arg1_1, %arg2_1) <{"callee" = @func_with_args}> : (i32, i16) -> i32 -// CHECK-NEXT: %attr_struct = "csl.const_struct"() <{"items" = {"i" = 42 : i32, "f" = 3.700000e+00 : f32}}> : () -> !csl.comptime_struct -// CHECK-NEXT: %ssa_struct = "csl.const_struct"(%arg1_1, %arg2_1, %col) <{"ssa_fields" = ["i32_", "i16_", "col"]}> : (i32, i16, !csl.color) -> !csl.comptime_struct -// CHECK-NEXT: %mixed_struct = "csl.const_struct"(%arg1_1, %arg2_1, %col) <{"ssa_fields" = ["i32_", "i16_", "col"], "items" = {"i" = 42 : i32, "f" = 3.700000e+00 : f32}}> : (i32, i16, !csl.color) -> !csl.comptime_struct +// CHECK-NEXT: %call_res = "csl.call"(%arg1_1, %arg2_1) <{callee = @func_with_args}> : (i32, i16) -> i32 +// CHECK-NEXT: %attr_struct = "csl.const_struct"() <{items = {i = 42 : i32, f = 3.700000e+00 : f32}}> : () -> !csl.comptime_struct +// CHECK-NEXT: %ssa_struct = "csl.const_struct"(%arg1_1, %arg2_1, %col) <{ssa_fields = ["i32_", "i16_", "col"]}> : (i32, i16, !csl.color) -> !csl.comptime_struct +// CHECK-NEXT: %mixed_struct = "csl.const_struct"(%arg1_1, %arg2_1, %col) <{ssa_fields = ["i32_", "i16_", "col"], items = {i = 42 : i32, f = 3.700000e+00 : f32}}> : (i32, i16, !csl.color) -> !csl.comptime_struct // CHECK-NEXT: %concat = "csl.concat_structs"(%attr_struct, %ssa_struct) : (!csl.comptime_struct, !csl.comptime_struct) -> !csl.comptime_struct // CHECK-NEXT: %three = arith.constant 3 : i16 // CHECK-NEXT: %col_1 = "csl.get_color"(%three) : (i16) -> !csl.color @@ -389,26 +389,26 @@ csl.func @builtins() { // CHECK-NEXT: %single_arr_ptr = "csl.addressof"(%arr) : (memref<10xf32>) -> !csl.ptr, #csl, #csl> // CHECK-NEXT: %ptrcast = "csl.ptrcast"(%scalar_ptr) : (!csl.ptr, #csl>) -> !csl.ptr, #csl, #csl> // CHECK-NEXT: %ptrcast_many = "csl.ptrcast"(%many_arr_ptr) : (!csl.ptr, #csl>) -> !csl.ptr, #csl, #csl> -// CHECK-NEXT: %function_ptr = "csl.addressof_fn"() <{"fn_name" = @initialize}> : () -> !csl.ptr<() -> (), #csl, #csl> -// CHECK-NEXT: %dir = "csl.get_dir"() <{"dir" = #csl}> : () -> !csl.direction +// CHECK-NEXT: %function_ptr = "csl.addressof_fn"() <{fn_name = @initialize}> : () -> !csl.ptr<() -> (), #csl, #csl> +// CHECK-NEXT: %dir = "csl.get_dir"() <{dir = #csl}> : () -> !csl.direction // CHECK-NEXT: %dsd_1d = "csl.get_mem_dsd"(%arr, %scalar) : (memref<10xf32>, i32) -> !csl -// CHECK-NEXT: %dsd_2d = "csl.get_mem_dsd"(%arr, %scalar, %scalar) <{"tensor_access" = affine_map<(d0, d1) -> (((d0 * 3) + 1), ((d1 * 4) + 2))>}> : (memref<10xf32>, i32, i32) -> !csl +// CHECK-NEXT: %dsd_2d = "csl.get_mem_dsd"(%arr, %scalar, %scalar) <{tensor_access = affine_map<(d0, d1) -> (((d0 * 3) + 1), ((d1 * 4) + 2))>}> : (memref<10xf32>, i32, i32) -> !csl // CHECK-NEXT: %dsd_3d = "csl.get_mem_dsd"(%arr, %scalar, %scalar, %scalar) : (memref<10xf32>, i32, i32, i32) -> !csl // CHECK-NEXT: %dsd_4d = "csl.get_mem_dsd"(%arr, %scalar, %scalar, %scalar, %scalar) : (memref<10xf32>, i32, i32, i32, i32) -> !csl // CHECK-NEXT: %dsd_1d1 = "csl.set_dsd_base_addr"(%dsd_1d, %many_arr_ptr) : (!csl, !csl.ptr, #csl>) -> !csl // CHECK-NEXT: %dsd_1d2 = "csl.set_dsd_base_addr"(%dsd_1d, %arr) : (!csl, memref<10xf32>) -> !csl -// CHECK-NEXT: %dsd_1d3 = "csl.increment_dsd_offset"(%dsd_1d2, %int16) <{"elem_type" = f32}> : (!csl, si16) -> !csl +// CHECK-NEXT: %dsd_1d3 = "csl.increment_dsd_offset"(%dsd_1d2, %int16) <{elem_type = f32}> : (!csl, si16) -> !csl // CHECK-NEXT: %dsd_1d4 = "csl.set_dsd_length"(%dsd_1d3, %u16) : (!csl, ui16) -> !csl // CHECK-NEXT: %dsd_1d5 = "csl.set_dsd_stride"(%dsd_1d4, %int8) : (!csl, si8) -> !csl // CHECK-NEXT: %tensor_dsd1 = "csl.get_mem_dsd"(%tens, %scalar) : (tensor<510xf32>, i32) -> !csl // CHECK-NEXT: %tensor_dsd2 = "csl.set_dsd_base_addr"(%dsd_1d, %tens) : (!csl, tensor<510xf32>) -> !csl -// CHECK-NEXT: %fabin_dsd = "csl.get_fab_dsd"(%scalar) <{"fabric_color" = 2 : ui5, "queue_id" = 0 : i3}> : (i32) -> !csl -// CHECK-NEXT: %fabout_dsd = "csl.get_fab_dsd"(%scalar) <{"fabric_color" = 3 : ui5, "queue_id" = 1 : i3, "control" = true, "wavelet_index_offset" = false}> : (i32) -> !csl +// CHECK-NEXT: %fabin_dsd = "csl.get_fab_dsd"(%scalar) <{fabric_color = 2 : ui5, queue_id = 0 : i3}> : (i32) -> !csl +// CHECK-NEXT: %fabout_dsd = "csl.get_fab_dsd"(%scalar) <{fabric_color = 3 : ui5, queue_id = 1 : i3, control = true, wavelet_index_offset = false}> : (i32) -> !csl // CHECK-NEXT: %f16_ptr, %f16_val, %f32_ptr = "test.op"() : () -> (!csl.ptr, #csl>, f16, !csl.ptr, #csl>) // CHECK-NEXT: "csl.faddh"(%dsd_1d1, %dsd_1d2, %dsd_1d3) : (!csl, !csl, !csl) -> () // CHECK-NEXT: "csl.faddh"(%f16_ptr, %f16_val, %dsd_1d3) : (!csl.ptr, #csl>, f16, !csl) -> () // CHECK-NEXT: %one = "test.op"() : () -> i32 -// CHECK-NEXT: %variable_with_default = "csl.variable"() <{"default" = 42 : i32}> : () -> !csl.var +// CHECK-NEXT: %variable_with_default = "csl.variable"() <{default = 42 : i32}> : () -> !csl.var // CHECK-NEXT: %variable = "csl.variable"() : () -> !csl.var // CHECK-NEXT: %value = "csl.load_var"(%variable_with_default) : (!csl.var) -> i32 // CHECK-NEXT: %new_value = arith.addi %value, %one : i32 @@ -564,69 +564,69 @@ csl.func @builtins() { // CHECK-NEXT: csl.return // CHECK-NEXT: } // CHECK-NEXT: %global_ptr = "test.op"() : () -> !csl.ptr, #csl> -// CHECK-NEXT: "csl.export"() <{"var_name" = @initialize, "type" = () -> ()}> : () -> () -// CHECK-NEXT: "csl.export"(%global_ptr) <{"var_name" = "some_name", "type" = !csl.ptr, #csl>}> : (!csl.ptr, #csl>) -> () +// CHECK-NEXT: "csl.export"() <{var_name = @initialize, type = () -> ()}> : () -> () +// CHECK-NEXT: "csl.export"(%global_ptr) <{var_name = "some_name", type = !csl.ptr, #csl>}> : (!csl.ptr, #csl>) -> () // CHECK-NEXT: %rpc_col = "test.op"() : () -> !csl.color // CHECK-NEXT: "csl.rpc"(%rpc_col) : (!csl.color) -> () -// CHECK-NEXT: }) {"sym_name" = "program"} : () -> () -// CHECK-NEXT: "csl.module"() <{"kind" = #csl}> ({ -// CHECK-NEXT: %comp_const = "csl.param"() <{"param_name" = "comp_constant"}> : () -> i32 +// CHECK-NEXT: }) {sym_name = "program"} : () -> () +// CHECK-NEXT: "csl.module"() <{kind = #csl}> ({ +// CHECK-NEXT: %comp_const = "csl.param"() <{param_name = "comp_constant"}> : () -> i32 // CHECK-NEXT: %init = arith.constant 3.140620e+00 : f16 -// CHECK-NEXT: %p2 = "csl.param"(%init) <{"param_name" = "param_2"}> : (f16) -> f16 +// CHECK-NEXT: %p2 = "csl.param"(%init) <{param_name = "param_2"}> : (f16) -> f16 // CHECK-NEXT: csl.layout { // CHECK-NEXT: %x_dim, %y_dim = "test.op"() : () -> (i32, i32) // CHECK-NEXT: "csl.set_rectangle"(%x_dim, %y_dim) : (i32, i32) -> () // CHECK-NEXT: %x_coord, %y_coord, %params = "test.op"() : () -> (i32, i32, !csl.comptime_struct) -// CHECK-NEXT: "csl.set_tile_code"(%x_coord, %y_coord, %params) <{"file" = "pe_program.csl"}> : (i32, i32, !csl.comptime_struct) -> () +// CHECK-NEXT: "csl.set_tile_code"(%x_coord, %y_coord, %params) <{file = "pe_program.csl"}> : (i32, i32, !csl.comptime_struct) -> () // CHECK-NEXT: } -// CHECK-NEXT: }) {"sym_name" = "layout"} : () -> () +// CHECK-NEXT: }) {sym_name = "layout"} : () -> () // CHECK-NEXT: } // CHECK-GENERIC-NEXT: "builtin.module"() ({ -// CHECK-GENERIC-NEXT: "csl.module"() <{"kind" = #csl}> ({ -// CHECK-GENERIC-NEXT: %thing = "csl.import_module"() <{"module" = ""}> : () -> !csl.imported_module -// CHECK-GENERIC-NEXT: "csl.func"() <{"sym_name" = "func_with_args", "function_type" = (i32, i16) -> i32}> ({ +// CHECK-GENERIC-NEXT: "csl.module"() <{kind = #csl}> ({ +// CHECK-GENERIC-NEXT: %thing = "csl.import_module"() <{module = ""}> : () -> !csl.imported_module +// CHECK-GENERIC-NEXT: "csl.func"() <{sym_name = "func_with_args", function_type = (i32, i16) -> i32}> ({ // CHECK-GENERIC-NEXT: ^0(%arg1 : i32, %arg2 : i16): // CHECK-GENERIC-NEXT: "csl.return"(%arg1) : (i32) -> () // CHECK-GENERIC-NEXT: }) : () -> () -// CHECK-GENERIC-NEXT: %zero = "arith.constant"() <{"value" = 0 : i32}> : () -> i32 -// CHECK-GENERIC-NEXT: %c100 = "arith.constant"() <{"value" = 100 : i32}> : () -> i32 +// CHECK-GENERIC-NEXT: %zero = "arith.constant"() <{value = 0 : i32}> : () -> i32 +// CHECK-GENERIC-NEXT: %c100 = "arith.constant"() <{value = 100 : i32}> : () -> i32 // CHECK-GENERIC-NEXT: %zeros = "csl.constants"(%zero, %c100) : (i32, i32) -> memref -// CHECK-GENERIC-NEXT: %zeros2 = "csl.constants"(%zero, %c100) <{"is_const"}> : (i32, i32) -> memref -// CHECK-GENERIC-NEXT: "csl.task"() <{"sym_name" = "local_task", "function_type" = () -> (), "kind" = #csl, "id" = 0 : ui5}> ({ +// CHECK-GENERIC-NEXT: %zeros2 = "csl.constants"(%zero, %c100) <{is_const}> : (i32, i32) -> memref +// CHECK-GENERIC-NEXT: "csl.task"() <{sym_name = "local_task", function_type = () -> (), kind = #csl, id = 0 : ui5}> ({ // CHECK-GENERIC-NEXT: "csl.return"() : () -> () // CHECK-GENERIC-NEXT: }) : () -> () -// CHECK-GENERIC-NEXT: "csl.task"() <{"sym_name" = "data_task", "function_type" = (i32) -> (), "kind" = #csl, "id" = 1 : ui5}> ({ +// CHECK-GENERIC-NEXT: "csl.task"() <{sym_name = "data_task", function_type = (i32) -> (), kind = #csl, id = 1 : ui5}> ({ // CHECK-GENERIC-NEXT: ^1(%a : i32): // CHECK-GENERIC-NEXT: "csl.return"() : () -> () // CHECK-GENERIC-NEXT: }) : () -> () -// CHECK-GENERIC-NEXT: "csl.task"() <{"sym_name" = "control_task", "function_type" = () -> (), "kind" = #csl, "id" = 2 : ui6}> ({ +// CHECK-GENERIC-NEXT: "csl.task"() <{sym_name = "control_task", function_type = () -> (), kind = #csl, id = 2 : ui6}> ({ // CHECK-GENERIC-NEXT: "csl.return"() : () -> () // CHECK-GENERIC-NEXT: }) : () -> () -// CHECK-GENERIC-NEXT: "csl.task"() <{"sym_name" = "control_task_args", "function_type" = (i32) -> (), "kind" = #csl, "id" = 2 : ui6}> ({ +// CHECK-GENERIC-NEXT: "csl.task"() <{sym_name = "control_task_args", function_type = (i32) -> (), kind = #csl, id = 2 : ui6}> ({ // CHECK-GENERIC-NEXT: ^2(%a_1 : i32): // CHECK-GENERIC-NEXT: "csl.return"() : () -> () // CHECK-GENERIC-NEXT: }) : () -> () -// CHECK-GENERIC-NEXT: "csl.task"() <{"sym_name" = "runtime_bound_local_task", "function_type" = () -> (), "kind" = #csl}> ({ +// CHECK-GENERIC-NEXT: "csl.task"() <{sym_name = "runtime_bound_local_task", function_type = () -> (), kind = #csl}> ({ // CHECK-GENERIC-NEXT: "csl.return"() : () -> () // CHECK-GENERIC-NEXT: }) : () -> () -// CHECK-GENERIC-NEXT: "csl.func"() <{"sym_name" = "initialize", "function_type" = () -> ()}> ({ +// CHECK-GENERIC-NEXT: "csl.func"() <{sym_name = "initialize", function_type = () -> ()}> ({ // CHECK-GENERIC-NEXT: %lb, %ub = "test.op"() : () -> (i16, i16) -// CHECK-GENERIC-NEXT: "csl.member_call"(%thing, %lb, %ub) <{"field" = "some_func"}> : (!csl.imported_module, i16, i16) -> () -// CHECK-GENERIC-NEXT: %res = "csl.member_call"(%thing, %lb, %ub) <{"field" = "some_func"}> : (!csl.imported_module, i16, i16) -> i32 -// CHECK-GENERIC-NEXT: %0 = "csl.member_access"(%thing) <{"field" = "some_field"}> : (!csl.imported_module) -> !csl.comptime_struct +// CHECK-GENERIC-NEXT: "csl.member_call"(%thing, %lb, %ub) <{field = "some_func"}> : (!csl.imported_module, i16, i16) -> () +// CHECK-GENERIC-NEXT: %res = "csl.member_call"(%thing, %lb, %ub) <{field = "some_func"}> : (!csl.imported_module, i16, i16) -> i32 +// CHECK-GENERIC-NEXT: %0 = "csl.member_access"(%thing) <{field = "some_field"}> : (!csl.imported_module) -> !csl.comptime_struct // CHECK-GENERIC-NEXT: %single_const = "test.op"() : () -> !csl.ptr, #csl> // CHECK-GENERIC-NEXT: %single_var = "test.op"() : () -> !csl.ptr, #csl> // CHECK-GENERIC-NEXT: %many_const = "test.op"() : () -> !csl.ptr, #csl> // CHECK-GENERIC-NEXT: %many_var = "test.op"() : () -> !csl.ptr, #csl> // CHECK-GENERIC-NEXT: %col = "test.op"() : () -> !csl.color // CHECK-GENERIC-NEXT: %arg1_1, %arg2_1 = "test.op"() : () -> (i32, i16) -// CHECK-GENERIC-NEXT: %call_res = "csl.call"(%arg1_1, %arg2_1) <{"callee" = @func_with_args}> : (i32, i16) -> i32 -// CHECK-GENERIC-NEXT: %attr_struct = "csl.const_struct"() <{"items" = {"i" = 42 : i32, "f" = 3.700000e+00 : f32}}> : () -> !csl.comptime_struct -// CHECK-GENERIC-NEXT: %ssa_struct = "csl.const_struct"(%arg1_1, %arg2_1, %col) <{"ssa_fields" = ["i32_", "i16_", "col"]}> : (i32, i16, !csl.color) -> !csl.comptime_struct -// CHECK-GENERIC-NEXT: %mixed_struct = "csl.const_struct"(%arg1_1, %arg2_1, %col) <{"ssa_fields" = ["i32_", "i16_", "col"], "items" = {"i" = 42 : i32, "f" = 3.700000e+00 : f32}}> : (i32, i16, !csl.color) -> !csl.comptime_struct +// CHECK-GENERIC-NEXT: %call_res = "csl.call"(%arg1_1, %arg2_1) <{callee = @func_with_args}> : (i32, i16) -> i32 +// CHECK-GENERIC-NEXT: %attr_struct = "csl.const_struct"() <{items = {i = 42 : i32, f = 3.700000e+00 : f32}}> : () -> !csl.comptime_struct +// CHECK-GENERIC-NEXT: %ssa_struct = "csl.const_struct"(%arg1_1, %arg2_1, %col) <{ssa_fields = ["i32_", "i16_", "col"]}> : (i32, i16, !csl.color) -> !csl.comptime_struct +// CHECK-GENERIC-NEXT: %mixed_struct = "csl.const_struct"(%arg1_1, %arg2_1, %col) <{ssa_fields = ["i32_", "i16_", "col"], items = {i = 42 : i32, f = 3.700000e+00 : f32}}> : (i32, i16, !csl.color) -> !csl.comptime_struct // CHECK-GENERIC-NEXT: %concat = "csl.concat_structs"(%attr_struct, %ssa_struct) : (!csl.comptime_struct, !csl.comptime_struct) -> !csl.comptime_struct -// CHECK-GENERIC-NEXT: %three = "arith.constant"() <{"value" = 3 : i16}> : () -> i16 +// CHECK-GENERIC-NEXT: %three = "arith.constant"() <{value = 3 : i16}> : () -> i16 // CHECK-GENERIC-NEXT: %col_1 = "csl.get_color"(%three) : (i16) -> !csl.color // CHECK-GENERIC-NEXT: %three_ui = "csl.mlir.signedness_cast"(%three) : (i16) -> ui16 // CHECK-GENERIC-NEXT: %arr, %scalar, %tens = "test.op"() : () -> (memref<10xf32>, i32, tensor<510xf32>) @@ -636,34 +636,34 @@ csl.func @builtins() { // CHECK-GENERIC-NEXT: %single_arr_ptr = "csl.addressof"(%arr) : (memref<10xf32>) -> !csl.ptr, #csl, #csl> // CHECK-GENERIC-NEXT: %ptrcast = "csl.ptrcast"(%scalar_ptr) : (!csl.ptr, #csl>) -> !csl.ptr, #csl, #csl> // CHECK-GENERIC-NEXT: %ptrcast_many = "csl.ptrcast"(%many_arr_ptr) : (!csl.ptr, #csl>) -> !csl.ptr, #csl, #csl> -// CHECK-GENERIC-NEXT: %function_ptr = "csl.addressof_fn"() <{"fn_name" = @initialize}> : () -> !csl.ptr<() -> (), #csl, #csl> -// CHECK-GENERIC-NEXT: %dir = "csl.get_dir"() <{"dir" = #csl}> : () -> !csl.direction +// CHECK-GENERIC-NEXT: %function_ptr = "csl.addressof_fn"() <{fn_name = @initialize}> : () -> !csl.ptr<() -> (), #csl, #csl> +// CHECK-GENERIC-NEXT: %dir = "csl.get_dir"() <{dir = #csl}> : () -> !csl.direction // CHECK-GENERIC-NEXT: %dsd_1d = "csl.get_mem_dsd"(%arr, %scalar) : (memref<10xf32>, i32) -> !csl -// CHECK-GENERIC-NEXT: %dsd_2d = "csl.get_mem_dsd"(%arr, %scalar, %scalar) <{"tensor_access" = affine_map<(d0, d1) -> (((d0 * 3) + 1), ((d1 * 4) + 2))>}> : (memref<10xf32>, i32, i32) -> !csl +// CHECK-GENERIC-NEXT: %dsd_2d = "csl.get_mem_dsd"(%arr, %scalar, %scalar) <{tensor_access = affine_map<(d0, d1) -> (((d0 * 3) + 1), ((d1 * 4) + 2))>}> : (memref<10xf32>, i32, i32) -> !csl // CHECK-GENERIC-NEXT: %dsd_3d = "csl.get_mem_dsd"(%arr, %scalar, %scalar, %scalar) : (memref<10xf32>, i32, i32, i32) -> !csl // CHECK-GENERIC-NEXT: %dsd_4d = "csl.get_mem_dsd"(%arr, %scalar, %scalar, %scalar, %scalar) : (memref<10xf32>, i32, i32, i32, i32) -> !csl // CHECK-GENERIC-NEXT: %dsd_1d1 = "csl.set_dsd_base_addr"(%dsd_1d, %many_arr_ptr) : (!csl, !csl.ptr, #csl>) -> !csl // CHECK-GENERIC-NEXT: %dsd_1d2 = "csl.set_dsd_base_addr"(%dsd_1d, %arr) : (!csl, memref<10xf32>) -> !csl -// CHECK-GENERIC-NEXT: %dsd_1d3 = "csl.increment_dsd_offset"(%dsd_1d2, %int16) <{"elem_type" = f32}> : (!csl, si16) -> !csl +// CHECK-GENERIC-NEXT: %dsd_1d3 = "csl.increment_dsd_offset"(%dsd_1d2, %int16) <{elem_type = f32}> : (!csl, si16) -> !csl // CHECK-GENERIC-NEXT: %dsd_1d4 = "csl.set_dsd_length"(%dsd_1d3, %u16) : (!csl, ui16) -> !csl // CHECK-GENERIC-NEXT: %dsd_1d5 = "csl.set_dsd_stride"(%dsd_1d4, %int8) : (!csl, si8) -> !csl // CHECK-GENERIC-NEXT: %tensor_dsd1 = "csl.get_mem_dsd"(%tens, %scalar) : (tensor<510xf32>, i32) -> !csl // CHECK-GENERIC-NEXT: %tensor_dsd2 = "csl.set_dsd_base_addr"(%dsd_1d, %tens) : (!csl, tensor<510xf32>) -> !csl -// CHECK-GENERIC-NEXT: %fabin_dsd = "csl.get_fab_dsd"(%scalar) <{"fabric_color" = 2 : ui5, "queue_id" = 0 : i3}> : (i32) -> !csl -// CHECK-GENERIC-NEXT: %fabout_dsd = "csl.get_fab_dsd"(%scalar) <{"fabric_color" = 3 : ui5, "queue_id" = 1 : i3, "control" = true, "wavelet_index_offset" = false}> : (i32) -> !csl +// CHECK-GENERIC-NEXT: %fabin_dsd = "csl.get_fab_dsd"(%scalar) <{fabric_color = 2 : ui5, queue_id = 0 : i3}> : (i32) -> !csl +// CHECK-GENERIC-NEXT: %fabout_dsd = "csl.get_fab_dsd"(%scalar) <{fabric_color = 3 : ui5, queue_id = 1 : i3, control = true, wavelet_index_offset = false}> : (i32) -> !csl // CHECK-GENERIC-NEXT: %f16_ptr, %f16_val, %f32_ptr = "test.op"() : () -> (!csl.ptr, #csl>, f16, !csl.ptr, #csl>) // CHECK-GENERIC-NEXT: "csl.faddh"(%dsd_1d1, %dsd_1d2, %dsd_1d3) : (!csl, !csl, !csl) -> () // CHECK-GENERIC-NEXT: "csl.faddh"(%f16_ptr, %f16_val, %dsd_1d3) : (!csl.ptr, #csl>, f16, !csl) -> () // CHECK-GENERIC-NEXT: %one = "test.op"() : () -> i32 -// CHECK-GENERIC-NEXT: %variable_with_default = "csl.variable"() <{"default" = 42 : i32}> : () -> !csl.var +// CHECK-GENERIC-NEXT: %variable_with_default = "csl.variable"() <{default = 42 : i32}> : () -> !csl.var // CHECK-GENERIC-NEXT: %variable = "csl.variable"() : () -> !csl.var // CHECK-GENERIC-NEXT: %value = "csl.load_var"(%variable_with_default) : (!csl.var) -> i32 -// CHECK-GENERIC-NEXT: %new_value = "arith.addi"(%value, %one) <{"overflowFlags" = #arith.overflow}> : (i32, i32) -> i32 +// CHECK-GENERIC-NEXT: %new_value = "arith.addi"(%value, %one) <{overflowFlags = #arith.overflow}> : (i32, i32) -> i32 // CHECK-GENERIC-NEXT: "csl.store_var"(%variable_with_default, %new_value) : (!csl.var, i32) -> () // CHECK-GENERIC-NEXT: "csl.store_var"(%variable, %new_value) : (!csl.var, i32) -> () // CHECK-GENERIC-NEXT: "csl.return"() : () -> () // CHECK-GENERIC-NEXT: }) : () -> () -// CHECK-GENERIC-NEXT: "csl.func"() <{"sym_name" = "builtins", "function_type" = () -> ()}> ({ +// CHECK-GENERIC-NEXT: "csl.func"() <{sym_name = "builtins", function_type = () -> ()}> ({ // CHECK-GENERIC-NEXT: %i16_value, %i32_value, %u16_value, %u32_value, %f16_value, %f32_value = "test.op"() : () -> (si16, si32, ui16, ui32, f16, f32) // CHECK-GENERIC-NEXT: %i16_pointer, %i32_pointer = "test.op"() : () -> (!csl.ptr, #csl>, !csl.ptr, #csl>) // CHECK-GENERIC-NEXT: %u16_pointer, %u32_pointer = "test.op"() : () -> (!csl.ptr, #csl>, !csl.ptr, #csl>) @@ -807,24 +807,24 @@ csl.func @builtins() { // CHECK-GENERIC-NEXT: "csl.xp162fs"(%dest_dsd, %src_dsd1) : (!csl, !csl) -> () // CHECK-GENERIC-NEXT: "csl.xp162fs"(%dest_dsd, %i16_value) : (!csl, si16) -> () // CHECK-GENERIC-NEXT: "csl.xp162fs"(%dest_dsd, %u16_value) : (!csl, ui16) -> () -// CHECK-GENERIC-NEXT: "csl.activate"() <{"kind" = #csl, "id" = 0 : ui6}> : () -> () +// CHECK-GENERIC-NEXT: "csl.activate"() <{kind = #csl, id = 0 : ui6}> : () -> () // CHECK-GENERIC-NEXT: "csl.return"() : () -> () // CHECK-GENERIC-NEXT: }) : () -> () // CHECK-GENERIC-NEXT: %global_ptr = "test.op"() : () -> !csl.ptr, #csl> -// CHECK-GENERIC-NEXT: "csl.export"() <{"var_name" = @initialize, "type" = () -> ()}> : () -> () -// CHECK-GENERIC-NEXT: "csl.export"(%global_ptr) <{"var_name" = "some_name", "type" = !csl.ptr, #csl>}> : (!csl.ptr, #csl>) -> () +// CHECK-GENERIC-NEXT: "csl.export"() <{var_name = @initialize, type = () -> ()}> : () -> () +// CHECK-GENERIC-NEXT: "csl.export"(%global_ptr) <{var_name = "some_name", type = !csl.ptr, #csl>}> : (!csl.ptr, #csl>) -> () // CHECK-GENERIC-NEXT: %rpc_col = "test.op"() : () -> !csl.color // CHECK-GENERIC-NEXT: "csl.rpc"(%rpc_col) : (!csl.color) -> () -// CHECK-GENERIC-NEXT: }) {"sym_name" = "program"} : () -> () -// CHECK-GENERIC-NEXT: "csl.module"() <{"kind" = #csl}> ({ -// CHECK-GENERIC-NEXT: %comp_const = "csl.param"() <{"param_name" = "comp_constant"}> : () -> i32 -// CHECK-GENERIC-NEXT: %init = "arith.constant"() <{"value" = 3.140620e+00 : f16}> : () -> f16 -// CHECK-GENERIC-NEXT: %p2 = "csl.param"(%init) <{"param_name" = "param_2"}> : (f16) -> f16 +// CHECK-GENERIC-NEXT: }) {sym_name = "program"} : () -> () +// CHECK-GENERIC-NEXT: "csl.module"() <{kind = #csl}> ({ +// CHECK-GENERIC-NEXT: %comp_const = "csl.param"() <{param_name = "comp_constant"}> : () -> i32 +// CHECK-GENERIC-NEXT: %init = "arith.constant"() <{value = 3.140620e+00 : f16}> : () -> f16 +// CHECK-GENERIC-NEXT: %p2 = "csl.param"(%init) <{param_name = "param_2"}> : (f16) -> f16 // CHECK-GENERIC-NEXT: "csl.layout"() ({ // CHECK-GENERIC-NEXT: %x_dim, %y_dim = "test.op"() : () -> (i32, i32) // CHECK-GENERIC-NEXT: "csl.set_rectangle"(%x_dim, %y_dim) : (i32, i32) -> () // CHECK-GENERIC-NEXT: %x_coord, %y_coord, %params = "test.op"() : () -> (i32, i32, !csl.comptime_struct) -// CHECK-GENERIC-NEXT: "csl.set_tile_code"(%x_coord, %y_coord, %params) <{"file" = "pe_program.csl"}> : (i32, i32, !csl.comptime_struct) -> () +// CHECK-GENERIC-NEXT: "csl.set_tile_code"(%x_coord, %y_coord, %params) <{file = "pe_program.csl"}> : (i32, i32, !csl.comptime_struct) -> () // CHECK-GENERIC-NEXT: }) : () -> () -// CHECK-GENERIC-NEXT: }) {"sym_name" = "layout"} : () -> () +// CHECK-GENERIC-NEXT: }) {sym_name = "layout"} : () -> () // CHECK-GENERIC-NEXT: }) : () -> () diff --git a/tests/filecheck/dialects/dmp/canonicalize.mlir b/tests/filecheck/dialects/dmp/canonicalize.mlir index e2fdb737b0..3b834a3eaf 100644 --- a/tests/filecheck/dialects/dmp/canonicalize.mlir +++ b/tests/filecheck/dialects/dmp/canonicalize.mlir @@ -12,7 +12,7 @@ builtin.module { ] } : (!stencil.field<[0,1024]x[0,1024]xf32>) -> () - // CHECK: "dmp.swap"(%ref) {"strategy" = #dmp.grid_slice_2d<#dmp.topo<2x2>, false>, "swaps" = [#dmp.exchange]} : (!stencil.field<[0,1024]x[0,1024]xf32>) -> () + // CHECK: "dmp.swap"(%ref) {strategy = #dmp.grid_slice_2d<#dmp.topo<2x2>, false>, swaps = [#dmp.exchange]} : (!stencil.field<[0,1024]x[0,1024]xf32>) -> () %swap_val = "dmp.swap"(%val) { strategy = #dmp.grid_slice_2d<#dmp.topo<2x2>, false>, diff --git a/tests/filecheck/dialects/dmp/ops.mlir b/tests/filecheck/dialects/dmp/ops.mlir index 41d8515207..11dfd001c3 100644 --- a/tests/filecheck/dialects/dmp/ops.mlir +++ b/tests/filecheck/dialects/dmp/ops.mlir @@ -21,6 +21,6 @@ builtin.module { ] } : (!stencil.temp<[0,1024]x[0,1024]xf32>) -> (!stencil.temp<[0,1024]x[0,1024]xf32>) - // CHECK: "dmp.swap"(%ref) {"strategy" = #dmp.grid_slice_2d<#dmp.topo<2x2>, false>, "swaps" = [#dmp.exchange, #dmp.exchange]} : (!stencil.field<[0,1024]x[0,1024]xf32>) -> () - // CHECK-NEXT: %swap_val = "dmp.swap"(%val) {"strategy" = #dmp.grid_slice_2d<#dmp.topo<2x2>, false>, "swaps" = [#dmp.exchange, #dmp.exchange]} : (!stencil.temp<[0,1024]x[0,1024]xf32>) -> !stencil.temp<[0,1024]x[0,1024]xf32> + // CHECK: "dmp.swap"(%ref) {strategy = #dmp.grid_slice_2d<#dmp.topo<2x2>, false>, swaps = [#dmp.exchange, #dmp.exchange]} : (!stencil.field<[0,1024]x[0,1024]xf32>) -> () + // CHECK-NEXT: %swap_val = "dmp.swap"(%val) {strategy = #dmp.grid_slice_2d<#dmp.topo<2x2>, false>, swaps = [#dmp.exchange, #dmp.exchange]} : (!stencil.temp<[0,1024]x[0,1024]xf32>) -> !stencil.temp<[0,1024]x[0,1024]xf32> } diff --git a/tests/filecheck/dialects/eqsat/eqsat_ops.mlir b/tests/filecheck/dialects/eqsat/eqsat_ops.mlir index 0c20e98568..16abdf1390 100644 --- a/tests/filecheck/dialects/eqsat/eqsat_ops.mlir +++ b/tests/filecheck/dialects/eqsat/eqsat_ops.mlir @@ -7,7 +7,7 @@ %v0, %v1 = "test.op"() : () -> (index, index) // CHECK-NEXT: %r0 = eqsat.eclass %v0 : index -// CHECK-NEXT: %r1 = eqsat.eclass %v0, %v1 {"hello" = "world"} : index +// CHECK-NEXT: %r1 = eqsat.eclass %v0, %v1 {hello = "world"} : index %r0 = eqsat.eclass %v0 : index %r1 = eqsat.eclass %v0, %v1 {"hello"="world"} : index @@ -16,5 +16,5 @@ // CHECK-GENERIC: "builtin.module"() ({ // CHECK-GENERIC-NEXT: %v0, %v1 = "test.op"() : () -> (index, index) // CHECK-GENERIC-NEXT: %r0 = "eqsat.eclass"(%v0) : (index) -> index -// CHECK-GENERIC-NEXT: %r1 = "eqsat.eclass"(%v0, %v1) {"hello" = "world"} : (index, index) -> index +// CHECK-GENERIC-NEXT: %r1 = "eqsat.eclass"(%v0, %v1) {hello = "world"} : (index, index) -> index // CHECK-GENERIC-NEXT: }) : () -> () diff --git a/tests/filecheck/dialects/func/func_ops.mlir b/tests/filecheck/dialects/func/func_ops.mlir index d5ef668d25..a33480921d 100644 --- a/tests/filecheck/dialects/func/func_ops.mlir +++ b/tests/filecheck/dialects/func/func_ops.mlir @@ -26,7 +26,7 @@ builtin.module { } // CHECK: func.func @call_void_attributes() { - // CHECK-NEXT: func.call @call_void_attributes() {"hello" = "world"} : () -> () + // CHECK-NEXT: func.call @call_void_attributes() {hello = "world"} : () -> () // CHECK-NEXT: func.return // CHECK-NEXT: } @@ -68,7 +68,7 @@ builtin.module { return %X : tensor<8x8xf64> } - // CHECK: func.func public @arg_attrs(%{{.*}} : tensor<8x8xf64> {"llvm.noalias"}, %{{.*}} : tensor<8x8xf64> {"llvm.noalias"}, %{{.*}} : tensor<8x8xf64> {"llvm.noalias"}) -> tensor<8x8xf64> { + // CHECK: func.func public @arg_attrs(%{{.*}} : tensor<8x8xf64> {llvm.noalias}, %{{.*}} : tensor<8x8xf64> {llvm.noalias}, %{{.*}} : tensor<8x8xf64> {llvm.noalias}) -> tensor<8x8xf64> { // CHECK-NEXT: return %{{.*}} : tensor<8x8xf64> // CHECK-NEXT: } @@ -77,7 +77,7 @@ builtin.module { return %r1, %r2 : f32, f32 } - // CHECK: func.func @output_attributes() -> (f32 {"dialect.a" = 0 : i32}, f32 {"dialect.b" = 0 : i32, "dialect.c" = 1 : i64}) { + // CHECK: func.func @output_attributes() -> (f32 {dialect.a = 0 : i32}, f32 {dialect.b = 0 : i32, dialect.c = 1 : i64}) { // CHECK-NEXT: %r1, %r2 = "test.op"() : () -> (f32, f32) // CHECK-NEXT: func.return %r1, %r2 : f32, f32 // CHECK-NEXT: } @@ -87,7 +87,7 @@ builtin.module { return %r1: f32 } - // CHECK: func.func @output_attribute_single() -> (f32 {"dialect.a" = 0 : i32}) { + // CHECK: func.func @output_attribute_single() -> (f32 {dialect.a = 0 : i32}) { // CHECK-NEXT: %r1 = "test.op"() : () -> f32 // CHECK-NEXT: func.return %r1 : f32 // CHECK-NEXT: } diff --git a/tests/filecheck/dialects/func/func_ops_generic.mlir b/tests/filecheck/dialects/func/func_ops_generic.mlir index 471a0b117a..5c368b1ed4 100644 --- a/tests/filecheck/dialects/func/func_ops_generic.mlir +++ b/tests/filecheck/dialects/func/func_ops_generic.mlir @@ -6,7 +6,7 @@ }) : () -> () -// CHECK: "func.func"() <{"function_type" = (tensor<8x8xf64>, tensor<8x8xf64>) -> (tensor<8x8xf64>, tensor<8x8xf64>), "res_attrs" = [{"llvm.noalias"}, {"llvm.noalias"}], "sym_name" = "arg_attrs", "sym_visibility" = "public"}> ({ +// CHECK: "func.func"() <{function_type = (tensor<8x8xf64>, tensor<8x8xf64>) -> (tensor<8x8xf64>, tensor<8x8xf64>), res_attrs = [{llvm.noalias}, {llvm.noalias}], sym_name = "arg_attrs", sym_visibility = "public"}> ({ // CHECK-NEXT: ^0(%arg0 : tensor<8x8xf64>, %arg1 : tensor<8x8xf64>): // CHECK-NEXT: "func.return"(%arg0, %arg1) : (tensor<8x8xf64>, tensor<8x8xf64>) -> () // CHECK-NEXT: }) : () -> () @@ -16,7 +16,7 @@ func.func @output_attributes() -> (f32 {dialect.a = 0 : i32}, f32 {dialect.b = 0 return %r1, %r2 : f32, f32 } -// CHECK: "func.func"() <{"sym_name" = "output_attributes", "function_type" = () -> (f32, f32), "res_attrs" = [{"dialect.a" = 0 : i32}, {"dialect.b" = 0 : i32, "dialect.c" = 1 : i64}]}> ({ +// CHECK: "func.func"() <{sym_name = "output_attributes", function_type = () -> (f32, f32), res_attrs = [{dialect.a = 0 : i32}, {dialect.b = 0 : i32, dialect.c = 1 : i64}]}> ({ // CHECK-NEXT: %r1, %r2 = "test.op"() : () -> (f32, f32) // CHECK-NEXT: "func.return"(%r1, %r2) : (f32, f32) -> () // CHECK-NEXT: }) : () -> () diff --git a/tests/filecheck/dialects/gpu/ops.mlir b/tests/filecheck/dialects/gpu/ops.mlir index 5395fc89f4..15b8cb6790 100644 --- a/tests/filecheck/dialects/gpu/ops.mlir +++ b/tests/filecheck/dialects/gpu/ops.mlir @@ -80,45 +80,45 @@ builtin.module attributes {"gpu.container_module"} { }) : () -> () } -// CHECK: builtin.module attributes {"gpu.container_module"} { -// CHECK-NEXT: "gpu.module"() <{"sym_name" = "gpu"}> ({ +// CHECK: builtin.module attributes {gpu.container_module} { +// CHECK-NEXT: "gpu.module"() <{sym_name = "gpu"}> ({ // CHECK-NEXT: func.func @kernel() { -// CHECK-NEXT: %{{.*}} = arith.constant {"proc" = #gpu} 13 : index -// CHECK-NEXT: %{{.*}} = arith.constant {"loopdim" = #gpu.loop_dim_map (d0), bound = (d0) -> (d0)>} 1 : index +// CHECK-NEXT: %{{.*}} = arith.constant {proc = #gpu} 13 : index +// CHECK-NEXT: %{{.*}} = arith.constant {loopdim = #gpu.loop_dim_map (d0), bound = (d0) -> (d0)>} 1 : index -// CHECK-NEXT: %{{.*}} = memref.alloc() {"alignment" = 0 : i64} : memref<10x10xi32> +// CHECK-NEXT: %{{.*}} = memref.alloc() {alignment = 0 : i64} : memref<10x10xi32> // CHECK-NEXT: %{{.*}} = "memref.cast"(%{{.*}}) : (memref<10x10xi32>) -> memref<*xi32> // CHECK-NEXT: "gpu.host_register"(%{{.*}}) : (memref<*xi32>) -> () // CHECK-NEXT: "gpu.host_unregister"(%{{.*}}) : (memref<*xi32>) -> () // CHECK-NEXT: %{{.*}} = "gpu.wait"() : () -> !gpu.async.token -// CHECK-NEXT: %{{.*}} = "gpu.thread_id"() <{"dimension" = #gpu}> : () -> index -// CHECK-NEXT: %{{.*}} = "gpu.thread_id"() <{"dimension" = #gpu}> : () -> index -// CHECK-NEXT: %{{.*}} = "gpu.thread_id"() <{"dimension" = #gpu}> : () -> index +// CHECK-NEXT: %{{.*}} = "gpu.thread_id"() <{dimension = #gpu}> : () -> index +// CHECK-NEXT: %{{.*}} = "gpu.thread_id"() <{dimension = #gpu}> : () -> index +// CHECK-NEXT: %{{.*}} = "gpu.thread_id"() <{dimension = #gpu}> : () -> index -// CHECK-NEXT: %{{.*}} = "gpu.block_dim"() <{"dimension" = #gpu}> : () -> index -// CHECK-NEXT: %{{.*}} = "gpu.block_dim"() <{"dimension" = #gpu}> : () -> index -// CHECK-NEXT: %{{.*}} = "gpu.block_dim"() <{"dimension" = #gpu}> : () -> index +// CHECK-NEXT: %{{.*}} = "gpu.block_dim"() <{dimension = #gpu}> : () -> index +// CHECK-NEXT: %{{.*}} = "gpu.block_dim"() <{dimension = #gpu}> : () -> index +// CHECK-NEXT: %{{.*}} = "gpu.block_dim"() <{dimension = #gpu}> : () -> index -// CHECK-NEXT: %{{.*}} = "gpu.block_id"() <{"dimension" = #gpu}> : () -> index -// CHECK-NEXT: %{{.*}} = "gpu.block_id"() <{"dimension" = #gpu}> : () -> index -// CHECK-NEXT: %{{.*}} = "gpu.block_id"() <{"dimension" = #gpu}> : () -> index +// CHECK-NEXT: %{{.*}} = "gpu.block_id"() <{dimension = #gpu}> : () -> index +// CHECK-NEXT: %{{.*}} = "gpu.block_id"() <{dimension = #gpu}> : () -> index +// CHECK-NEXT: %{{.*}} = "gpu.block_id"() <{dimension = #gpu}> : () -> index -// CHECK-NEXT: %{{.*}} = "gpu.global_id"() <{"dimension" = #gpu}> : () -> index -// CHECK-NEXT: %{{.*}} = "gpu.global_id"() <{"dimension" = #gpu}> : () -> index -// CHECK-NEXT: %{{.*}} = "gpu.global_id"() <{"dimension" = #gpu}> : () -> index +// CHECK-NEXT: %{{.*}} = "gpu.global_id"() <{dimension = #gpu}> : () -> index +// CHECK-NEXT: %{{.*}} = "gpu.global_id"() <{dimension = #gpu}> : () -> index +// CHECK-NEXT: %{{.*}} = "gpu.global_id"() <{dimension = #gpu}> : () -> index -// CHECK-NEXT: %{{.*}} = "gpu.grid_dim"() <{"dimension" = #gpu}> : () -> index -// CHECK-NEXT: %{{.*}} = "gpu.grid_dim"() <{"dimension" = #gpu}> : () -> index -// CHECK-NEXT: %{{.*}} = "gpu.grid_dim"() <{"dimension" = #gpu}> : () -> index +// CHECK-NEXT: %{{.*}} = "gpu.grid_dim"() <{dimension = #gpu}> : () -> index +// CHECK-NEXT: %{{.*}} = "gpu.grid_dim"() <{dimension = #gpu}> : () -> index +// CHECK-NEXT: %{{.*}} = "gpu.grid_dim"() <{dimension = #gpu}> : () -> index -// CHECK-NEXT: %gmemref = "gpu.alloc"() <{"operandSegmentSizes" = array}> : () -> memref<10x10xi32> -// CHECK-NEXT: %gdmemref = "gpu.alloc"(%griddimx, %griddimy, %griddimz) <{"operandSegmentSizes" = array}> : (index, index, index) -> memref +// CHECK-NEXT: %gmemref = "gpu.alloc"() <{operandSegmentSizes = array}> : () -> memref<10x10xi32> +// CHECK-NEXT: %gdmemref = "gpu.alloc"(%griddimx, %griddimy, %griddimz) <{operandSegmentSizes = array}> : (index, index, index) -> memref -// CHECK-NEXT: "gpu.memcpy"(%memref, %gmemref) {"operandSegmentSizes" = array} : (memref<10x10xi32>, memref<10x10xi32>) -> () +// CHECK-NEXT: "gpu.memcpy"(%memref, %gmemref) {operandSegmentSizes = array} : (memref<10x10xi32>, memref<10x10xi32>) -> () -// CHECK-NEXT: "gpu.dealloc"(%gdmemref) {"operandSegmentSizes" = array} : (memref) -> () +// CHECK-NEXT: "gpu.dealloc"(%gdmemref) {operandSegmentSizes = array} : (memref) -> () // CHECK-NEXT: %{{.*}} = "gpu.lane_id"() : () -> index // CHECK-NEXT: %{{.*}} = "gpu.num_subgroups"() : () -> index @@ -129,7 +129,7 @@ builtin.module attributes {"gpu.container_module"} { // CHECK-NEXT: %{{.*}} = "gpu.subgroup_id"() : () -> index // CHECK-NEXT: %{{.*}} = "gpu.subgroup_size"() : () -> index -// CHECK-NEXT: %{{.*}} = "gpu.all_reduce"(%{{.*}}) <{"op" = #gpu}> ({ +// CHECK-NEXT: %{{.*}} = "gpu.all_reduce"(%{{.*}}) <{op = #gpu}> ({ // CHECK-NEXT: }) : (index) -> index // CHECK-NEXT: %{{.*}} = "gpu.all_reduce"(%{{.*}}) ({ @@ -138,24 +138,24 @@ builtin.module attributes {"gpu.container_module"} { // CHECK-NEXT: "gpu.yield"(%{{.*}}) : (index) -> () // CHECK-NEXT: }) : (index) -> index -// CHECK-NEXT: "gpu.launch"(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "gpu.launch"(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) <{operandSegmentSizes = array}> ({ // CHECK-NEXT: ^{{\S+}}(%{{\S+}} : index, %{{\S+}} : index, %{{\S+}} : index, // CHECK-SAME: %{{\S+}} : index, %{{\S+}} : index, %{{\S+}} : index, // CHECK-SAME: %{{\S+}} : index, %{{\S+}} : index, %{{\S+}} : index, // CHECK-SAME: %{{\S+}} : index, %{{\S+}} : index, %{{\S+}} : index): -// CHECK-NEXT: %{{.*}} = "gpu.all_reduce"(%{{.*}}) <{"op" = #gpu}> ({ +// CHECK-NEXT: %{{.*}} = "gpu.all_reduce"(%{{.*}}) <{op = #gpu}> ({ // CHECK-NEXT: }) : (index) -> index // CHECK-NEXT: %{{.*}} = arith.muli %{{.*}}, %{{.*}} : index // CHECK-NEXT: "gpu.terminator"() : () -> () // CHECK-NEXT: }) : (index, index, index, index, index, index) -> () -// CHECK-NEXT: "gpu.launch_func"(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) <{"kernel" = @gpu::@foo, "operandSegmentSizes" = array}> : (index, index, index, index, index, index, i32, index) -> () +// CHECK-NEXT: "gpu.launch_func"(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) <{kernel = @gpu::@foo, operandSegmentSizes = array}> : (index, index, index, index, index, index, i32, index) -> () // CHECK-NEXT: func.return // CHECK-NEXT: } -// CHECK-NEXT: "gpu.func"() <{"function_type" = (index) -> (), "kernel"}> ({ +// CHECK-NEXT: "gpu.func"() <{function_type = (index) -> (), kernel}> ({ // CHECK-NEXT: ^{{.*}}(%{{.*}}: index): // CHECK-NEXT: "gpu.return"() : () -> () -// CHECK-NEXT: }) {"sym_name" = "foo", "gpu.known_block_size" = array, "gpu.known_grid_size" = array} : () -> () +// CHECK-NEXT: }) {sym_name = "foo", gpu.known_block_size = array, gpu.known_grid_size = array} : () -> () // CHECK-NEXT: "gpu.module_end"() : () -> () // CHECK-NEXT: }) : () -> () diff --git a/tests/filecheck/dialects/hw/hw_instance.mlir b/tests/filecheck/dialects/hw/hw_instance.mlir index af1a60d894..62664d11ea 100644 --- a/tests/filecheck/dialects/hw/hw_instance.mlir +++ b/tests/filecheck/dialects/hw/hw_instance.mlir @@ -18,27 +18,27 @@ builtin.module { %ci1 = "test.op"() : () -> (i1) // CHECK: hw.instance "empty_instance" @test_empty() -> () - // CHECK-GENERIC: "hw.instance"() {"instanceName" = "empty_instance", "moduleName" = @test_empty, "argNames" = [], "resultNames" = []} : () -> () + // CHECK-GENERIC: "hw.instance"() {instanceName = "empty_instance", moduleName = @test_empty, argNames = [], resultNames = []} : () -> () hw.instance "empty_instance" @test_empty() -> () - // CHECK: hw.instance "empty_instance_with_attrs" @test_empty() -> () {"foo" = 4 : i32, "bar"} - // CHECK-GENERIC: "hw.instance"() {"foo" = 4 : i32, "bar", "instanceName" = "empty_instance_with_attrs", "moduleName" = @test_empty, "argNames" = [], "resultNames" = []} : () -> () - hw.instance "empty_instance_with_attrs" @test_empty() -> () {"foo" = 4 : i32, "bar"} + // CHECK: hw.instance "empty_instance_with_attrs" @test_empty() -> () {foo = 4 : i32, bar} + // CHECK-GENERIC: "hw.instance"() {foo = 4 : i32, bar, instanceName = "empty_instance_with_attrs", moduleName = @test_empty, argNames = [], resultNames = []} : () -> () + hw.instance "empty_instance_with_attrs" @test_empty() -> () {foo = 4 : i32, bar} // CHECK: hw.instance "empty_instance_with_inner_sym" sym #hw @test_empty() -> () - // CHECK-GENERIC: "hw.instance"() {"instanceName" = "empty_instance_with_inner_sym", "moduleName" = @test_empty, "argNames" = [], "resultNames" = [], "inner_sym" = #hw} : () -> () + // CHECK-GENERIC: "hw.instance"() {instanceName = "empty_instance_with_inner_sym", moduleName = @test_empty, argNames = [], resultNames = [], inner_sym = #hw} : () -> () hw.instance "empty_instance_with_inner_sym" sym #hw @test_empty() -> () // CHECK: %{{.*}} = hw.instance "basic_instance" @test_basic(testIn: %{{.*}}: i32) -> (testOut: i32) - // CHECK-GENERIC: %{{.*}} = "hw.instance"(%{{.*}}) {"instanceName" = "basic_instance", "moduleName" = @test_basic, "argNames" = ["testIn"], "resultNames" = ["testOut"]} : (i32) -> i32 + // CHECK-GENERIC: %{{.*}} = "hw.instance"(%{{.*}}) {instanceName = "basic_instance", moduleName = @test_basic, argNames = ["testIn"], resultNames = ["testOut"]} : (i32) -> i32 %res = hw.instance "basic_instance" @test_basic(testIn: %ci32: i32) -> (testOut: i32) // CHECK: %{{.*}} = hw.instance "multiple_instance" @test_multiple(testIn1: %{{.*}}: i32, testIn2: %{{.*}}: i1) -> (testOut1: i1, testOut2: i32) - // CHECK-GENERIC: %{{.*}} = "hw.instance"(%{{.*}}, %{{.*}}) {"instanceName" = "multiple_instance", "moduleName" = @test_multiple, "argNames" = ["testIn1", "testIn2"], "resultNames" = ["testOut1", "testOut2"]} : (i32, i1) -> (i1, i32) + // CHECK-GENERIC: %{{.*}} = "hw.instance"(%{{.*}}, %{{.*}}) {instanceName = "multiple_instance", moduleName = @test_multiple, argNames = ["testIn1", "testIn2"], resultNames = ["testOut1", "testOut2"]} : (i32, i1) -> (i1, i32) %res1, %res2 = hw.instance "multiple_instance" @test_multiple(testIn1: %ci32: i32, testIn2: %ci1: i1) -> (testOut1: i1, testOut2: i32) // CHECK: %{{.*}} = hw.instance "multiple_instance_swap" @test_multiple(testIn2: %{{.*}}: i1, testIn1: %{{.*}}: i32) -> (testOut2: i32, testOut1: i1) - // CHECK-GENERIC: %{{.*}} = "hw.instance"(%{{.*}}, %{{.*}}) {"instanceName" = "multiple_instance_swap", "moduleName" = @test_multiple, "argNames" = ["testIn2", "testIn1"], "resultNames" = ["testOut2", "testOut1"]} : (i1, i32) -> (i32, i1) + // CHECK-GENERIC: %{{.*}} = "hw.instance"(%{{.*}}, %{{.*}}) {instanceName = "multiple_instance_swap", moduleName = @test_multiple, argNames = ["testIn2", "testIn1"], resultNames = ["testOut2", "testOut1"]} : (i1, i32) -> (i32, i1) %res2s, %res1s = hw.instance "multiple_instance_swap" @test_multiple(testIn2: %ci1: i1, testIn1: %ci32: i32) -> (testOut2: i32, testOut1: i1) // Assert types. diff --git a/tests/filecheck/dialects/hw/hw_module.mlir b/tests/filecheck/dialects/hw/hw_module.mlir index 357a8a2c5a..ec308f76cb 100644 --- a/tests/filecheck/dialects/hw/hw_module.mlir +++ b/tests/filecheck/dialects/hw/hw_module.mlir @@ -9,7 +9,7 @@ // CHECK-GENERIC: "hw.module"() ({ // CHECK-GENERIC-NEXT: ^{{[^(]*}}(%{{[^ ]*}} : i1, %{{[^ ]*}} : i1, %{{[^ ]*}} : i32): // CHECK-GENERIC-NEXT: "hw.output"(%{{[^ ]*}}, %{{[^ ]*}}) : (i1, i1) -> () -// CHECK-GENERIC-NEXT: }) {"sym_name" = "custom_basic", "module_type" = !hw.modty, "parameters" = []} : () -> () +// CHECK-GENERIC-NEXT: }) {sym_name = "custom_basic", module_type = !hw.modty, parameters = []} : () -> () hw.module @custom_basic(in %a_foo a: i1, out nameOfPortInSV: i1, out "": i1, in %b customName: i1, in %c "very custom name": i32) { hw.output %a_foo, %b: i1, i1 } @@ -19,7 +19,7 @@ hw.module @custom_basic(in %a_foo a: i1, out nameOfPortInSV: i1, out "": i1, in // CHECK-NEXT: } // CHECK-GENERIC: "hw.module"() ({ // CHECK-GENERIC-NEXT: "hw.output"() : () -> () -// CHECK-GENERIC-NEXT: }) {"sym_name" = "custom_param", "module_type" = !hw.modty<>, "parameters" = [#hw.param.decl<"p1": i42>, #hw.param.decl<"p2 with wack name": i1>]} : () -> () +// CHECK-GENERIC-NEXT: }) {sym_name = "custom_param", module_type = !hw.modty<>, parameters = [#hw.param.decl<"p1": i42>, #hw.param.decl<"p2 with wack name": i1>]} : () -> () hw.module @custom_param() { hw.output } @@ -29,7 +29,7 @@ hw.module @custom_param() { // CHECK-NEXT: } // CHECK-GENERIC: "hw.module"() ({ // CHECK-GENERIC-NEXT: "hw.output"() : () -> () -// CHECK-GENERIC-NEXT: }) {"sym_name" = "private_module", "module_type" = !hw.modty<>, "parameters" = [], "sym_visibility" = "private"} : () -> () +// CHECK-GENERIC-NEXT: }) {sym_name = "private_module", module_type = !hw.modty<>, parameters = [], sym_visibility = "private"} : () -> () hw.module private @private_module() { hw.output } @@ -39,25 +39,25 @@ hw.module private @private_module() { // CHECK-NEXT: } // CHECK-GENERIC: "hw.module"() ({ // CHECK-GENERIC-NEXT: "hw.output"() : () -> () -// CHECK-GENERIC-NEXT: }) {"sym_name" = "wack name!!", "module_type" = !hw.modty<>, "parameters" = []} : () -> () +// CHECK-GENERIC-NEXT: }) {sym_name = "wack name!!", module_type = !hw.modty<>, parameters = []} : () -> () hw.module @"wack name!!"() { hw.output } // CHECK: hw.module @name_preserve(in %preserve: i8) -// CHECK-GENERIC: {"sym_name" = "name_preserve", "module_type" = !hw.modty, "parameters" = []} : () -> () +// CHECK-GENERIC: {sym_name = "name_preserve", module_type = !hw.modty, parameters = []} : () -> () hw.module @name_preserve(in %preserve: i8) { hw.output } -// CHECK: hw.module @more_attrs() attributes {"foo"} { -// CHECK-GENERIC: }) {"sym_name" = "more_attrs", "module_type" = !hw.modty<>, "parameters" = [], "foo"} : () -> () +// CHECK: hw.module @more_attrs() attributes {foo} { +// CHECK-GENERIC: }) {sym_name = "more_attrs", module_type = !hw.modty<>, parameters = [], foo} : () -> () hw.module @more_attrs() attributes {"foo"} { hw.output } // CHECK: hw.module.extern @extern(in %{{.*}} foo: i32, in %{{.*}} hello: i16, out baz: i8) -// CHECK-GENERIC: "hw.module.extern"() {"sym_name" = "extern", "module_type" = !hw.modty, "parameters" = []} : () -> () +// CHECK-GENERIC: "hw.module.extern"() {sym_name = "extern", module_type = !hw.modty, parameters = []} : () -> () hw.module.extern @extern(in %foo: i32, in %bar "hello": i16, out baz: i8) @@ -69,7 +69,7 @@ hw.module.extern @extern(in %foo: i32, in %bar "hello": i16, out baz: i8) // CHECK-GENERIC: "hw.module"() ({ // CHECK-GENERIC-NEXT: ^{{[^(]*}}(%{{[^ ]*}} : i1, %{{[^ ]*}} : i1, %{{[^ ]*}} : i32): // CHECK-GENERIC-NEXT: "hw.output"(%{{[^ ]*}}, %{{[^ ]*}}) : (i1, i1) -> () -// CHECK-GENERIC-NEXT: }) {"sym_name" = "generic_basic", "module_type" = !hw.modty, "parameters" = []} : () -> () +// CHECK-GENERIC-NEXT: }) {sym_name = "generic_basic", module_type = !hw.modty, parameters = []} : () -> () "hw.module"() ({ ^0(%a_foo : i1, %b : i1, %c: i32): "hw.output"(%a_foo, %b) : (i1, i1) -> () @@ -84,7 +84,7 @@ hw.module.extern @extern(in %foo: i32, in %bar "hello": i16, out baz: i8) // CHECK-NEXT: } // CHECK-GENERIC: "hw.module"() ({ // CHECK-GENERIC-NEXT: "hw.output"() : () -> () -// CHECK-GENERIC-NEXT: }) {"sym_name" = "generic_param", "module_type" = !hw.modty<>, "parameters" = [#hw.param.decl<"p1": i42>, #hw.param.decl<"p2 with wack name": i1>]} : () -> () +// CHECK-GENERIC-NEXT: }) {sym_name = "generic_param", module_type = !hw.modty<>, parameters = [#hw.param.decl<"p1": i42>, #hw.param.decl<"p2 with wack name": i1>]} : () -> () "hw.module"() ({ "hw.output"() : () -> () }) { diff --git a/tests/filecheck/dialects/hw/hw_ops.mlir b/tests/filecheck/dialects/hw/hw_ops.mlir index 16b1aad56c..0f85829d53 100644 --- a/tests/filecheck/dialects/hw/hw_ops.mlir +++ b/tests/filecheck/dialects/hw/hw_ops.mlir @@ -2,13 +2,13 @@ "builtin.module"() ({ "test.op"() { "test" = #hw.innerNameRef<@Foo::@Bar> } : () -> () - // CHECK: "test.op"() {"test" = #hw.innerNameRef<@Foo::@Bar>} : () -> () + // CHECK: "test.op"() {test = #hw.innerNameRef<@Foo::@Bar>} : () -> () "test.op"() { "test" = #hw, <@y,5,public>]> } : () -> () - // CHECK-NEXT: "test.op"() {"test" = #hw, <@y,5,public>]>} : () -> () + // CHECK-NEXT: "test.op"() {test = #hw, <@y,5,public>]>} : () -> () "test.op"() { "test" = #hw } : () -> () - // CHECK-NEXT: "test.op"() {"test" = #hw} : () -> () + // CHECK-NEXT: "test.op"() {test = #hw} : () -> () "test.op"() { "test" = #hw.direction } : () -> () }) : () -> () diff --git a/tests/filecheck/dialects/linalg/linalg_ops.mlir b/tests/filecheck/dialects/linalg/linalg_ops.mlir index 3888c7c1fd..dd9a9d3485 100644 --- a/tests/filecheck/dialects/linalg/linalg_ops.mlir +++ b/tests/filecheck/dialects/linalg/linalg_ops.mlir @@ -53,7 +53,7 @@ linalg.quantized_matmul ins(%5, %6, %7, %8 : tensor<64x9216xi8>, tensor<9216x409 // CHECK-NEXT: ^1(%{{.*}} : f32, %{{.*}} : f32): // CHECK-NEXT: linalg.yield %{{.*}} : f32 // CHECK-NEXT: } -// CHECK-NEXT: linalg.generic {indexing_maps = [affine_map<(d0, d1) -> ()>, affine_map<(d0, d1) -> (d0, d1)>], iterator_types = ["parallel", "parallel"]} ins(%{{.*}} : f32) outs(%{{.*}} : memref<1x256xf32>) attrs = {"hello" = "world"} { +// CHECK-NEXT: linalg.generic {indexing_maps = [affine_map<(d0, d1) -> ()>, affine_map<(d0, d1) -> (d0, d1)>], iterator_types = ["parallel", "parallel"]} ins(%{{.*}} : f32) outs(%{{.*}} : memref<1x256xf32>) attrs = {hello = "world"} { // CHECK-NEXT: ^{{.*}}(%{{.*}} f32, %{{.*}} f32): // CHECK-NEXT: linalg.yield %{{.*}} : f32 // CHECK-NEXT: } @@ -65,7 +65,7 @@ linalg.quantized_matmul ins(%5, %6, %7, %8 : tensor<64x9216xi8>, tensor<9216x409 // CHECK-NEXT: linalg.mul ins(%{{.*}} %{{.*}} : memref<4x16xf32>, memref<4x16xf32>) outs(%{{.*}} : memref<4x16xf32>) // CHECK-NEXT: %{{.*}} %{{.*}} = "test.op"() : () -> (memref<64x9216xf32>, memref<9216x4096xf32>) // CHECK-NEXT: %{{.*}} = "test.op"() : () -> memref<64x4096xf32> -// CHECK-NEXT: linalg.matmul {"id"} ins(%{{.*}} %{{.*}} : memref<64x9216xf32>, memref<9216x4096xf32>) outs(%{{.*}} : memref<64x4096xf32>) +// CHECK-NEXT: linalg.matmul {id} ins(%{{.*}} %{{.*}} : memref<64x9216xf32>, memref<9216x4096xf32>) outs(%{{.*}} : memref<64x4096xf32>) // CHECK-NEXT: %{{.*}} = linalg.fill ins(%{{.*}} : f32) outs(%{{.*}} : tensor<4x16xf32>) -> tensor<4x16xf32> // CHECK-NEXT: linalg.fill ins(%{{.*}} : f32) outs(%{{.*}} : memref<4x16xf32>) // CHECK-NEXT: %5, %6 = "test.op"() : () -> (tensor<64x9216xi8>, tensor<9216x4096xi8>) @@ -75,42 +75,42 @@ linalg.quantized_matmul ins(%5, %6, %7, %8 : tensor<64x9216xi8>, tensor<9216x409 // CHECK-NEXT: linalg.quantized_matmul ins(%5, %6, %7, %8 : tensor<64x9216xi8>, tensor<9216x4096xi8>, i32, i32) outs(%9 : tensor<64x4096xi32>) -> tensor<64x4096xi32> // CHECK-NEXT: } -// CHECK-GENERIC: "linalg.generic"(%{{.*}} %{{.*}} <{"indexing_maps" = [affine_map<(d0, d1) -> ()>, affine_map<(d0, d1) -> (d0, d1)>], "iterator_types" = [#linalg.iterator_type, #linalg.iterator_type], "operandSegmentSizes" = array}> ({ +// CHECK-GENERIC: "linalg.generic"(%{{.*}} %{{.*}} <{indexing_maps = [affine_map<(d0, d1) -> ()>, affine_map<(d0, d1) -> (d0, d1)>], iterator_types = [#linalg.iterator_type, #linalg.iterator_type], operandSegmentSizes = array}> ({ // CHECK-GENERIC-NEXT: ^0(%{{.*}} f32, %{{.*}} f32): // CHECK-GENERIC-NEXT: "linalg.yield"(%{{.*}} : (f32) -> () // CHECK-GENERIC-NEXT: }) : (f32, memref<1x256xf32>) -> () -// CHECK-GENERIC-NEXT: "linalg.generic"(%{{.*}} %{{.*}} <{"indexing_maps" = [affine_map<(d0, d1) -> ()>, affine_map<(d0, d1) -> (d0, d1)>], "iterator_types" = [#linalg.iterator_type, #linalg.iterator_type], "doc" = "a_docstring", "library_call" = "a_library_call", "operandSegmentSizes" = array}> ({ +// CHECK-GENERIC-NEXT: "linalg.generic"(%{{.*}} %{{.*}} <{indexing_maps = [affine_map<(d0, d1) -> ()>, affine_map<(d0, d1) -> (d0, d1)>], iterator_types = [#linalg.iterator_type, #linalg.iterator_type], doc = "a_docstring", library_call = "a_library_call", operandSegmentSizes = array}> ({ // CHECK-GENERIC-NEXT: ^1(%{{.*}} : f32, %{{.*}} : f32): // CHECK-GENERIC-NEXT: "linalg.yield"(%{{.*}} : (f32) -> () // CHECK-GENERIC-NEXT: }) : (f32, memref<1x256xf32>) -> () -// CHECK-GENERIC: "linalg.generic"(%{{.*}} %{{.*}} <{"indexing_maps" = [affine_map<(d0, d1) -> ()>, affine_map<(d0, d1) -> (d0, d1)>], "iterator_types" = [#linalg.iterator_type, #linalg.iterator_type], "operandSegmentSizes" = array}> ({ +// CHECK-GENERIC: "linalg.generic"(%{{.*}} %{{.*}} <{indexing_maps = [affine_map<(d0, d1) -> ()>, affine_map<(d0, d1) -> (d0, d1)>], iterator_types = [#linalg.iterator_type, #linalg.iterator_type], operandSegmentSizes = array}> ({ // CHECK-GENERIC-NEXT: ^{{.*}}(%{{.*}} f32, %{{.*}} f32): // CHECK-GENERIC-NEXT: "linalg.yield"(%{{.*}} : (f32) -> () -// CHECK-GENERIC-NEXT: }) {"hello" = "world"} : (f32, memref<1x256xf32>) -> () +// CHECK-GENERIC-NEXT: }) {hello = "world"} : (f32, memref<1x256xf32>) -> () // CHECK-GENERIC-NEXT: %{{.*}} %{{.*}} %{{.*}} = "test.op"() : () -> (tensor<4x16xf32>, tensor<4x16xf32>, tensor<4x16xf32>) // CHECK-GENERIC-NEXT: %{{.*}} %{{.*}} %{{.*}} = "test.op"() : () -> (memref<4x16xf32>, memref<4x16xf32>, memref<4x16xf32>) -// CHECK-GENERIC-NEXT: %{{.*}} = "linalg.add"(%{{.*}} %{{.*}} %{{.*}} <{"operandSegmentSizes" = array}> ({ +// CHECK-GENERIC-NEXT: %{{.*}} = "linalg.add"(%{{.*}} %{{.*}} %{{.*}} <{operandSegmentSizes = array}> ({ // CHECK-GENERIC-NEXT: ^3(%{{.*}} : f32, %{{.*}} : f32, %{{.*}} : f32): // CHECK-GENERIC-NEXT: %{{.*}} = "arith.addf"(%{{.*}} %{{.*}} : (f32, f32) -> f32 // CHECK-GENERIC-NEXT: "linalg.yield"(%{{.*}} : (f32) -> () // CHECK-GENERIC-NEXT: }) : (tensor<4x16xf32>, tensor<4x16xf32>, tensor<4x16xf32>) -> tensor<4x16xf32> -// CHECK-GENERIC-NEXT: "linalg.add"(%{{.*}} %{{.*}} %{{.*}} <{"operandSegmentSizes" = array}> ({ +// CHECK-GENERIC-NEXT: "linalg.add"(%{{.*}} %{{.*}} %{{.*}} <{operandSegmentSizes = array}> ({ // CHECK-GENERIC-NEXT: ^4(%{{.*}} : f32, %{{.*}} : f32, %{{.*}} : f32): // CHECK-GENERIC-NEXT: %{{.*}} = "arith.addf"(%{{.*}} %{{.*}} : (f32, f32) -> f32 // CHECK-GENERIC-NEXT: "linalg.yield"(%{{.*}} : (f32) -> () // CHECK-GENERIC-NEXT: }) : (memref<4x16xf32>, memref<4x16xf32>, memref<4x16xf32>) -> () -// CHECK-GENERIC-NEXT: %{{.*}} = "linalg.mul"(%{{.*}} %{{.*}} %{{.*}} <{"operandSegmentSizes" = array}> ({ +// CHECK-GENERIC-NEXT: %{{.*}} = "linalg.mul"(%{{.*}} %{{.*}} %{{.*}} <{operandSegmentSizes = array}> ({ // CHECK-GENERIC-NEXT: ^5(%{{.*}} : f32, %{{.*}} : f32, %{{.*}} : f32): // CHECK-GENERIC-NEXT: %{{.*}} = "arith.mulf"(%{{.*}} %{{.*}} : (f32, f32) -> f32 // CHECK-GENERIC-NEXT: "linalg.yield"(%{{.*}} : (f32) -> () // CHECK-GENERIC-NEXT: }) : (tensor<4x16xf32>, tensor<4x16xf32>, tensor<4x16xf32>) -> tensor<4x16xf32> -// CHECK-GENERIC-NEXT: "linalg.mul"(%{{.*}} %{{.*}} %{{.*}} <{"operandSegmentSizes" = array}> ({ +// CHECK-GENERIC-NEXT: "linalg.mul"(%{{.*}} %{{.*}} %{{.*}} <{operandSegmentSizes = array}> ({ // CHECK-GENERIC-NEXT: ^6(%{{.*}} : f32, %{{.*}} : f32, %{{.*}} : f32): // CHECK-GENERIC-NEXT: %{{.*}} = "arith.mulf"(%{{.*}} %{{.*}} : (f32, f32) -> f32 // CHECK-GENERIC-NEXT: "linalg.yield"(%{{.*}} : (f32) -> () @@ -119,35 +119,35 @@ linalg.quantized_matmul ins(%5, %6, %7, %8 : tensor<64x9216xi8>, tensor<9216x409 // CHECK-GENERIC-NEXT: %{{.*}} %{{.*}} = "test.op"() : () -> (memref<64x9216xf32>, memref<9216x4096xf32>) // CHECK-GENERIC-NEXT: %{{.*}} = "test.op"() : () -> memref<64x4096xf32> -// CHECK-GENERIC-NEXT: "linalg.matmul"(%{{.*}} %{{.*}} %{{.*}} <{"operandSegmentSizes" = array}> ({ +// CHECK-GENERIC-NEXT: "linalg.matmul"(%{{.*}} %{{.*}} %{{.*}} <{operandSegmentSizes = array}> ({ // CHECK-GENERIC-NEXT: ^7(%{{.*}} : f32, %{{.*}} : f32, %{{.*}} : f32): // CHECK-GENERIC-NEXT: %{{.*}} = "arith.mulf"(%{{.*}}, %{{.*}} : (f32, f32) -> f32 // CHECK-GENERIC-NEXT: %{{.*}} = "arith.addf"(%{{.*}}, %{{.*}} : (f32, f32) -> f32 // CHECK-GENERIC-NEXT: "linalg.yield"(%{{.*}} : (f32) -> () -// CHECK-GENERIC-NEXT: }) {"id", "linalg.memoized_indexing_maps" = [affine_map<(d0, d1, d2) -> (d0, d2)>, affine_map<(d0, d1, d2) -> (d2, d1)>, affine_map<(d0, d1, d2) -> (d0, d1)>]} : (memref<64x9216xf32>, memref<9216x4096xf32>, memref<64x4096xf32>) -> () +// CHECK-GENERIC-NEXT: }) {id, linalg.memoized_indexing_maps = [affine_map<(d0, d1, d2) -> (d0, d2)>, affine_map<(d0, d1, d2) -> (d2, d1)>, affine_map<(d0, d1, d2) -> (d0, d1)>]} : (memref<64x9216xf32>, memref<9216x4096xf32>, memref<64x4096xf32>) -> () -// CHECK-GENERIC-NEXT: %{{.*}} = "linalg.fill"(%{{.*}}, %{{.*}} <{"operandSegmentSizes" = array}> ({ +// CHECK-GENERIC-NEXT: %{{.*}} = "linalg.fill"(%{{.*}}, %{{.*}} <{operandSegmentSizes = array}> ({ // CHECK-GENERIC-NEXT: ^8(%{{.*}} : f32, %{{.*}} : f32): // CHECK-GENERIC-NEXT: "linalg.yield"(%{{.*}} : (f32) -> () // CHECK-GENERIC-NEXT: }) : (f32, tensor<4x16xf32>) -> tensor<4x16xf32> -// CHECK-GENERIC-NEXT: "linalg.fill"(%{{.*}}, %{{.*}} <{"operandSegmentSizes" = array}> ({ +// CHECK-GENERIC-NEXT: "linalg.fill"(%{{.*}}, %{{.*}} <{operandSegmentSizes = array}> ({ // CHECK-GENERIC-NEXT: ^9(%{{.*}} : f32, %{{.*}} : f32): // CHECK-GENERIC-NEXT: "linalg.yield"(%{{.*}} : (f32) -> () // CHECK-GENERIC-NEXT: }) : (f32, memref<4x16xf32>) -> () // CHECK-GENERIC-NEXT: %{{.*}}, %{{.*}} = "test.op"() : () -> (tensor<64x9216xi8>, tensor<9216x4096xi8>) -// CHECK-GENERIC-NEXT: %{{.*}} = "arith.constant"() <{"value" = 0 : i32}> : () -> i32 -// CHECK-GENERIC-NEXT: %{{.*}} = "arith.constant"() <{"value" = 0 : i32}> : () -> i32 +// CHECK-GENERIC-NEXT: %{{.*}} = "arith.constant"() <{value = 0 : i32}> : () -> i32 +// CHECK-GENERIC-NEXT: %{{.*}} = "arith.constant"() <{value = 0 : i32}> : () -> i32 // CHECK-GENERIC-NEXT: %{{.*}} = "test.op"() : () -> tensor<64x4096xi32> -// CHECK-GENERIC-NEXT: %{{.*}} = "linalg.quantized_matmul"(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) <{"operandSegmentSizes" = array}> ({ +// CHECK-GENERIC-NEXT: %{{.*}} = "linalg.quantized_matmul"(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) <{operandSegmentSizes = array}> ({ // CHECK-GENERIC-NEXT: ^10(%36 : i8, %37 : i8, %38 : i32, %39 : i32, %40 : i32): // CHECK-GENERIC-NEXT: %41 = "arith.extsi"(%36) : (i8) -> i32 -// CHECK-GENERIC-NEXT: %42 = "arith.subi"(%41, %38) <{"overflowFlags" = #arith.overflow}> : (i32, i32) -> i32 +// CHECK-GENERIC-NEXT: %42 = "arith.subi"(%41, %38) <{overflowFlags = #arith.overflow}> : (i32, i32) -> i32 // CHECK-GENERIC-NEXT: %43 = "arith.extsi"(%37) : (i8) -> i32 -// CHECK-GENERIC-NEXT: %44 = "arith.subi"(%43, %39) <{"overflowFlags" = #arith.overflow}> : (i32, i32) -> i32 -// CHECK-GENERIC-NEXT: %45 = "arith.muli"(%42, %44) <{"overflowFlags" = #arith.overflow}> : (i32, i32) -> i32 -// CHECK-GENERIC-NEXT: %46 = "arith.addi"(%40, %45) <{"overflowFlags" = #arith.overflow}> : (i32, i32) -> i32 +// CHECK-GENERIC-NEXT: %44 = "arith.subi"(%43, %39) <{overflowFlags = #arith.overflow}> : (i32, i32) -> i32 +// CHECK-GENERIC-NEXT: %45 = "arith.muli"(%42, %44) <{overflowFlags = #arith.overflow}> : (i32, i32) -> i32 +// CHECK-GENERIC-NEXT: %46 = "arith.addi"(%40, %45) <{overflowFlags = #arith.overflow}> : (i32, i32) -> i32 // CHECK-GENERIC-NEXT: "linalg.yield"(%46) : (i32) -> () -// CHECK-GENERIC-NEXT: }) {"linalg.memoized_indexing_maps" = [affine_map<(d0, d1, d2) -> (d0, d2)>, affine_map<(d0, d1, d2) -> (d2, d1)>, affine_map<(d0, d1, d2) -> (d0, d1)>]} : (tensor<64x9216xi8>, tensor<9216x4096xi8>, i32, i32, tensor<64x4096xi32>) -> tensor<64x4096xi32> +// CHECK-GENERIC-NEXT: }) {linalg.memoized_indexing_maps = [affine_map<(d0, d1, d2) -> (d0, d2)>, affine_map<(d0, d1, d2) -> (d2, d1)>, affine_map<(d0, d1, d2) -> (d0, d1)>]} : (tensor<64x9216xi8>, tensor<9216x4096xi8>, i32, i32, tensor<64x4096xi32>) -> tensor<64x4096xi32> diff --git a/tests/filecheck/dialects/llvm/arithmetic.mlir b/tests/filecheck/dialects/llvm/arithmetic.mlir index ada534f4da..d305c61ab0 100644 --- a/tests/filecheck/dialects/llvm/arithmetic.mlir +++ b/tests/filecheck/dialects/llvm/arithmetic.mlir @@ -3,7 +3,7 @@ %arg0, %arg1 = "test.op"() : () -> (i32, i32) %add_both = llvm.add %arg0, %arg1 {"overflowFlags" = #llvm.overflow} : i32 -// CHECK: %add_both = llvm.add %arg0, %arg1 {"overflowFlags" = #llvm.overflow} : i32 +// CHECK: %add_both = llvm.add %arg0, %arg1 {overflowFlags = #llvm.overflow} : i32 %add_both_pretty = llvm.add %arg0, %arg1 overflow : i32 // CHECK: %add_both_pretty = llvm.add %arg0, %arg1 overflow : i32 diff --git a/tests/filecheck/dialects/llvm/example.mlir b/tests/filecheck/dialects/llvm/example.mlir index 2b30e57c64..1a8ee0e7a2 100644 --- a/tests/filecheck/dialects/llvm/example.mlir +++ b/tests/filecheck/dialects/llvm/example.mlir @@ -78,8 +78,8 @@ builtin.module { // CHECK: func.func public @main() { // CHECK-NEXT: %0 = arith.constant 1 : i32 // CHECK-NEXT: %1 = "llvm.mlir.undef"() : () -> !llvm.struct<(i32)> -// CHECK-NEXT: %2 = "llvm.insertvalue"(%1, %0) <{"position" = array}> : (!llvm.struct<(i32)>, i32) -> !llvm.struct<(i32)> -// CHECK-NEXT: %3 = "llvm.extractvalue"(%2) <{"position" = array}> : (!llvm.struct<(i32)>) -> i32 +// CHECK-NEXT: %2 = "llvm.insertvalue"(%1, %0) <{position = array}> : (!llvm.struct<(i32)>, i32) -> !llvm.struct<(i32)> +// CHECK-NEXT: %3 = "llvm.extractvalue"(%2) <{position = array}> : (!llvm.struct<(i32)>) -> i32 // CHECK-NEXT: %4 = llvm.mlir.zero : !llvm.struct<(i32, f32)> // CHECK-NEXT: func.return // CHECK-NEXT: } diff --git a/tests/filecheck/dialects/llvm/global.mlir b/tests/filecheck/dialects/llvm/global.mlir index 15f45af3ae..b9775e42e7 100644 --- a/tests/filecheck/dialects/llvm/global.mlir +++ b/tests/filecheck/dialects/llvm/global.mlir @@ -7,8 +7,8 @@ builtin.module { } // CHECK: builtin.module { -// CHECK-NEXT: "llvm.mlir.global"() <{"global_type" = !llvm.array<13 x i8>, "constant", "sym_name" = "str0", "linkage" = #llvm.linkage<"internal">, "value" = "Hello world!\n", "addr_space" = 0 : i32, "unnamed_addr" = 0 : i64}> ({ +// CHECK-NEXT: "llvm.mlir.global"() <{global_type = !llvm.array<13 x i8>, constant, sym_name = "str0", linkage = #llvm.linkage<"internal">, value = "Hello world!\n", addr_space = 0 : i32, unnamed_addr = 0 : i64}> ({ // CHECK-NEXT: }) : () -> () -// CHECK-NEXT: %0 = "llvm.mlir.addressof"() <{"global_name" = @str0}> : () -> !llvm.ptr -// CHECK-NEXT: %1 = "llvm.getelementptr"(%0) <{"elem_type" = !llvm.array<13 x i8>, "rawConstantIndices" = array}> : (!llvm.ptr) -> !llvm.ptr +// CHECK-NEXT: %0 = "llvm.mlir.addressof"() <{global_name = @str0}> : () -> !llvm.ptr +// CHECK-NEXT: %1 = "llvm.getelementptr"(%0) <{elem_type = !llvm.array<13 x i8>, rawConstantIndices = array}> : (!llvm.ptr) -> !llvm.ptr // CHECK-NEXT: } diff --git a/tests/filecheck/dialects/llvm/icmp.mlir b/tests/filecheck/dialects/llvm/icmp.mlir index dc558e3010..a135858bbf 100644 --- a/tests/filecheck/dialects/llvm/icmp.mlir +++ b/tests/filecheck/dialects/llvm/icmp.mlir @@ -4,50 +4,50 @@ %icmp_eq_p = "llvm.icmp"(%arg0, %arg1) <{predicate = 0 : i64}> : (i32, i32) -> i1 %icmp_eq = llvm.icmp "eq" %arg0, %arg1 : i32 -// CHECK: %icmp_eq_p = "llvm.icmp"(%arg0, %arg1) <{"predicate" = 0 : i64}> : (i32, i32) -> i1 -// CHECK-NEXT: %icmp_eq = "llvm.icmp"(%arg0, %arg1) <{"predicate" = 0 : i64}> : (i32, i32) -> i1 +// CHECK: %icmp_eq_p = "llvm.icmp"(%arg0, %arg1) <{predicate = 0 : i64}> : (i32, i32) -> i1 +// CHECK-NEXT: %icmp_eq = "llvm.icmp"(%arg0, %arg1) <{predicate = 0 : i64}> : (i32, i32) -> i1 %icmp_ne_p = "llvm.icmp"(%arg0, %arg1) <{predicate = 1 : i64}> : (i32, i32) -> i1 %icmp_ne = llvm.icmp "ne" %arg0, %arg1 : i32 -// CHECK-NEXT: %icmp_ne_p = "llvm.icmp"(%arg0, %arg1) <{"predicate" = 1 : i64}> : (i32, i32) -> i1 -// CHECK-NEXT: %icmp_ne = "llvm.icmp"(%arg0, %arg1) <{"predicate" = 1 : i64}> : (i32, i32) -> i1 +// CHECK-NEXT: %icmp_ne_p = "llvm.icmp"(%arg0, %arg1) <{predicate = 1 : i64}> : (i32, i32) -> i1 +// CHECK-NEXT: %icmp_ne = "llvm.icmp"(%arg0, %arg1) <{predicate = 1 : i64}> : (i32, i32) -> i1 %icmp_slt_p = "llvm.icmp"(%arg0, %arg1) <{predicate = 2 : i64}> : (i32, i32) -> i1 %icmp_slt = llvm.icmp "slt" %arg0, %arg1 : i32 -// CHECK-NEXT: %icmp_slt_p = "llvm.icmp"(%arg0, %arg1) <{"predicate" = 2 : i64}> : (i32, i32) -> i1 -// CHECK-NEXT: %icmp_slt = "llvm.icmp"(%arg0, %arg1) <{"predicate" = 2 : i64}> : (i32, i32) -> i1 +// CHECK-NEXT: %icmp_slt_p = "llvm.icmp"(%arg0, %arg1) <{predicate = 2 : i64}> : (i32, i32) -> i1 +// CHECK-NEXT: %icmp_slt = "llvm.icmp"(%arg0, %arg1) <{predicate = 2 : i64}> : (i32, i32) -> i1 %icmp_sle_p = "llvm.icmp"(%arg0, %arg1) <{predicate = 3 : i64}> : (i32, i32) -> i1 %icmp_sle = llvm.icmp "sle" %arg0, %arg1 : i32 -// CHECK-NEXT: %icmp_sle_p = "llvm.icmp"(%arg0, %arg1) <{"predicate" = 3 : i64}> : (i32, i32) -> i1 -// CHECK-NEXT: %icmp_sle = "llvm.icmp"(%arg0, %arg1) <{"predicate" = 3 : i64}> : (i32, i32) -> i1 +// CHECK-NEXT: %icmp_sle_p = "llvm.icmp"(%arg0, %arg1) <{predicate = 3 : i64}> : (i32, i32) -> i1 +// CHECK-NEXT: %icmp_sle = "llvm.icmp"(%arg0, %arg1) <{predicate = 3 : i64}> : (i32, i32) -> i1 %icmp_sgt_p = "llvm.icmp"(%arg0, %arg1) <{predicate = 4 : i64}> : (i32, i32) -> i1 %icmp_sgt = llvm.icmp "sgt" %arg0, %arg1 : i32 -// CHECK-NEXT: %icmp_sgt_p = "llvm.icmp"(%arg0, %arg1) <{"predicate" = 4 : i64}> : (i32, i32) -> i1 -// CHECK-NEXT: %icmp_sgt = "llvm.icmp"(%arg0, %arg1) <{"predicate" = 4 : i64}> : (i32, i32) -> i1 +// CHECK-NEXT: %icmp_sgt_p = "llvm.icmp"(%arg0, %arg1) <{predicate = 4 : i64}> : (i32, i32) -> i1 +// CHECK-NEXT: %icmp_sgt = "llvm.icmp"(%arg0, %arg1) <{predicate = 4 : i64}> : (i32, i32) -> i1 %icmp_sge_p = "llvm.icmp"(%arg0, %arg1) <{predicate = 5 : i64}> : (i32, i32) -> i1 %icmp_sge = llvm.icmp "sge" %arg0, %arg1 : i32 -// CHECK-NEXT: %icmp_sge_p = "llvm.icmp"(%arg0, %arg1) <{"predicate" = 5 : i64}> : (i32, i32) -> i1 -// CHECK-NEXT: %icmp_sge = "llvm.icmp"(%arg0, %arg1) <{"predicate" = 5 : i64}> : (i32, i32) -> i1 +// CHECK-NEXT: %icmp_sge_p = "llvm.icmp"(%arg0, %arg1) <{predicate = 5 : i64}> : (i32, i32) -> i1 +// CHECK-NEXT: %icmp_sge = "llvm.icmp"(%arg0, %arg1) <{predicate = 5 : i64}> : (i32, i32) -> i1 %icmp_ult_p = "llvm.icmp"(%arg0, %arg1) <{predicate = 6 : i64}> : (i32, i32) -> i1 %icmp_ult = llvm.icmp "ult" %arg0, %arg1 : i32 -// CHECK-NEXT: %icmp_ult_p = "llvm.icmp"(%arg0, %arg1) <{"predicate" = 6 : i64}> : (i32, i32) -> i1 -// CHECK-NEXT: %icmp_ult = "llvm.icmp"(%arg0, %arg1) <{"predicate" = 6 : i64}> : (i32, i32) -> i1 +// CHECK-NEXT: %icmp_ult_p = "llvm.icmp"(%arg0, %arg1) <{predicate = 6 : i64}> : (i32, i32) -> i1 +// CHECK-NEXT: %icmp_ult = "llvm.icmp"(%arg0, %arg1) <{predicate = 6 : i64}> : (i32, i32) -> i1 %icmp_ule_p = "llvm.icmp"(%arg0, %arg1) <{predicate = 7 : i64}> : (i32, i32) -> i1 %icmp_ule = llvm.icmp "ule" %arg0, %arg1 : i32 -// CHECK-NEXT: %icmp_ule_p = "llvm.icmp"(%arg0, %arg1) <{"predicate" = 7 : i64}> : (i32, i32) -> i1 -// CHECK-NEXT: %icmp_ule = "llvm.icmp"(%arg0, %arg1) <{"predicate" = 7 : i64}> : (i32, i32) -> i1 +// CHECK-NEXT: %icmp_ule_p = "llvm.icmp"(%arg0, %arg1) <{predicate = 7 : i64}> : (i32, i32) -> i1 +// CHECK-NEXT: %icmp_ule = "llvm.icmp"(%arg0, %arg1) <{predicate = 7 : i64}> : (i32, i32) -> i1 %icmp_ugt_p = "llvm.icmp"(%arg0, %arg1) <{predicate = 8 : i64}> : (i32, i32) -> i1 %icmp_ugt = llvm.icmp "ugt" %arg0, %arg1 : i32 -// CHECK-NEXT: %icmp_ugt_p = "llvm.icmp"(%arg0, %arg1) <{"predicate" = 8 : i64}> : (i32, i32) -> i1 -// CHECK-NEXT: %icmp_ugt = "llvm.icmp"(%arg0, %arg1) <{"predicate" = 8 : i64}> : (i32, i32) -> i1 +// CHECK-NEXT: %icmp_ugt_p = "llvm.icmp"(%arg0, %arg1) <{predicate = 8 : i64}> : (i32, i32) -> i1 +// CHECK-NEXT: %icmp_ugt = "llvm.icmp"(%arg0, %arg1) <{predicate = 8 : i64}> : (i32, i32) -> i1 %icmp_uge_p = "llvm.icmp"(%arg0, %arg1) <{predicate = 9 : i64}> : (i32, i32) -> i1 %icmp_uge = llvm.icmp "uge" %arg0, %arg1 : i32 -// CHECK-NEXT: %icmp_uge_p = "llvm.icmp"(%arg0, %arg1) <{"predicate" = 9 : i64}> : (i32, i32) -> i1 -// CHECK-NEXT: %icmp_uge = "llvm.icmp"(%arg0, %arg1) <{"predicate" = 9 : i64}> : (i32, i32) -> i1 +// CHECK-NEXT: %icmp_uge_p = "llvm.icmp"(%arg0, %arg1) <{predicate = 9 : i64}> : (i32, i32) -> i1 +// CHECK-NEXT: %icmp_uge = "llvm.icmp"(%arg0, %arg1) <{predicate = 9 : i64}> : (i32, i32) -> i1 diff --git a/tests/filecheck/dialects/llvm/inline_asm.mlir b/tests/filecheck/dialects/llvm/inline_asm.mlir index cace16fbef..bce625382c 100644 --- a/tests/filecheck/dialects/llvm/inline_asm.mlir +++ b/tests/filecheck/dialects/llvm/inline_asm.mlir @@ -10,7 +10,7 @@ // CHECK: %0 = "test.op"() : () -> i32 // CHECK-NEXT: 1 = "test.op"() : () -> i32 -// CHECK-NEXT: "llvm.inline_asm"(%0, %1) <{"asm_string" = "csrw $0, $1", "constraints" = "i, r", "has_side_effects"}> : (i32, i32) -> () +// CHECK-NEXT: "llvm.inline_asm"(%0, %1) <{asm_string = "csrw $0, $1", constraints = "i, r", has_side_effects}> : (i32, i32) -> () // CHECK-NEXT: %2 = "test.op"() : () -> vector<8xf32> // CHECK-NEXT: %3 = "test.op"() : () -> vector<8xf32> -// CHECK-NEXT: %4 = "llvm.inline_asm"(%2, %3) <{"asm_dialect" = 1 : i64, "asm_string" = "vaddps $0, $1, $2", "constraints" = "=x,x,x"}> : (vector<8xf32>, vector<8xf32>) -> vector<8xf32> +// CHECK-NEXT: %4 = "llvm.inline_asm"(%2, %3) <{asm_dialect = 1 : i64, asm_string = "vaddps $0, $1, $2", constraints = "=x,x,x"}> : (vector<8xf32>, vector<8xf32>) -> vector<8xf32> diff --git a/tests/filecheck/dialects/llvm/pointers.mlir b/tests/filecheck/dialects/llvm/pointers.mlir index b4590109af..f1b15c1d53 100644 --- a/tests/filecheck/dialects/llvm/pointers.mlir +++ b/tests/filecheck/dialects/llvm/pointers.mlir @@ -15,8 +15,8 @@ builtin.module { // CHECK-NEXT: %1 = "llvm.inttoptr"(%0) : (i64) -> !llvm.ptr // CHECK-NEXT: %2 = "llvm.load"(%1) : (!llvm.ptr) -> i32 // CHECK-NEXT: %3 = "llvm.mlir.null"() : () -> !llvm.ptr -// CHECK-NEXT: %4 = "llvm.alloca"(%0) <{"alignment" = 32 : i64}> : (i64) -> !llvm.ptr +// CHECK-NEXT: %4 = "llvm.alloca"(%0) <{alignment = 32 : i64}> : (i64) -> !llvm.ptr // CHECK-NEXT: %5 = "llvm.load"(%4) : (!llvm.ptr) -> index -// CHECK-NEXT: %6 = "llvm.getelementptr"(%4, %0) <{"elem_type" = i32, "rawConstantIndices" = array}> : (!llvm.ptr, i64) -> !llvm.ptr +// CHECK-NEXT: %6 = "llvm.getelementptr"(%4, %0) <{elem_type = i32, rawConstantIndices = array}> : (!llvm.ptr, i64) -> !llvm.ptr // CHECK-NEXT: %7 = "llvm.alloca"(%0) : (i64) -> !llvm.ptr // CHECK-NEXT: } diff --git a/tests/filecheck/dialects/memref/memref_ops.mlir b/tests/filecheck/dialects/memref/memref_ops.mlir index f5695f1c94..cffbd0a1fe 100644 --- a/tests/filecheck/dialects/memref/memref_ops.mlir +++ b/tests/filecheck/dialects/memref/memref_ops.mlir @@ -56,32 +56,32 @@ builtin.module { // CHECK-NEXT: }) : () -> () // CHECK-NEXT: func.return // CHECK-NEXT: } -// CHECK-NEXT: "memref.global"() <{"sym_name" = "g", "sym_visibility" = "public", "type" = memref<1xindex>, "initial_value" = dense<0> : tensor<1xindex>}> : () -> () -// CHECK-NEXT: "memref.global"() <{"sym_name" = "g_constant", "sym_visibility" = "public", "type" = memref<1xindex>, "initial_value" = dense<0> : tensor<1xindex>, "constant"}> : () -> () -// CHECK-NEXT: "memref.global"() <{"sym_name" = "g_with_alignment", "sym_visibility" = "public", "type" = memref<1xindex>, "initial_value" = dense<0> : tensor<1xindex>, "alignment" = 64 : i64}> : () -> () +// CHECK-NEXT: "memref.global"() <{sym_name = "g", sym_visibility = "public", type = memref<1xindex>, initial_value = dense<0> : tensor<1xindex>}> : () -> () +// CHECK-NEXT: "memref.global"() <{sym_name = "g_constant", sym_visibility = "public", type = memref<1xindex>, initial_value = dense<0> : tensor<1xindex>, constant}> : () -> () +// CHECK-NEXT: "memref.global"() <{sym_name = "g_with_alignment", sym_visibility = "public", type = memref<1xindex>, initial_value = dense<0> : tensor<1xindex>, alignment = 64 : i64}> : () -> () // CHECK-NEXT: func.func private @memref_test() { // CHECK-NEXT: %{{.*}} = memref.get_global @g : memref<1xindex> // CHECK-NEXT: %{{.*}} = arith.constant 0 : index -// CHECK-NEXT: %{{.*}} = "memref.alloca"() <{"alignment" = 0 : i64, "operandSegmentSizes" = array}> : () -> memref<1xindex> +// CHECK-NEXT: %{{.*}} = "memref.alloca"() <{alignment = 0 : i64, operandSegmentSizes = array}> : () -> memref<1xindex> // CHECK-NEXT: %{{.*}} = arith.constant 42 : index // CHECK-NEXT: memref.store %{{.*}}, %{{.*}}[%{{.*}}] : memref<1xindex> -// CHECK-NEXT: memref.store %{{.*}}, %{{.*}}[%{{.*}}] {"nontemporal" = true} : memref<1xindex> +// CHECK-NEXT: memref.store %{{.*}}, %{{.*}}[%{{.*}}] {nontemporal = true} : memref<1xindex> // CHECK-NEXT: %{{.*}} = memref.load %{{.*}}[%{{.*}}] : memref<1xindex> -// CHECK-NEXT: %{{.*}} = memref.load %{{.*}}[%{{.*}}] {"nontemporal" = false} : memref<1xindex> -// CHECK-NEXT: %{{.*}} = memref.alloc() {"alignment" = 0 : i64} : memref<10x2xindex> +// CHECK-NEXT: %{{.*}} = memref.load %{{.*}}[%{{.*}}] {nontemporal = false} : memref<1xindex> +// CHECK-NEXT: %{{.*}} = memref.alloc() {alignment = 0 : i64} : memref<10x2xindex> // CHECK-NEXT: memref.store %{{.*}}, %{{.*}}[%{{.*}}, %{{.*}}] : memref<10x2xindex> -// CHECK-NEXT: %{{.*}} = memref.subview %{{.*}}[0, 0] [1, 1] [1, 1] attributes {"hello" = "world"} : memref<10x2xindex> to memref<1x1xindex> +// CHECK-NEXT: %{{.*}} = memref.subview %{{.*}}[0, 0] [1, 1] [1, 1] attributes {hello = "world"} : memref<10x2xindex> to memref<1x1xindex> // CHECK-NEXT: %{{.*}} = "memref.cast"(%{{.*}}) : (memref<10x2xindex>) -> memref -// CHECK-NEXT: %{{.*}} = "memref.alloca"() <{"operandSegmentSizes" = array}> : () -> memref<1xindex> +// CHECK-NEXT: %{{.*}} = "memref.alloca"() <{operandSegmentSizes = array}> : () -> memref<1xindex> // CHECK-NEXT: %{{.*}} = "memref.memory_space_cast"(%{{.*}}) : (memref<10x2xindex>) -> memref<10x2xindex, 1 : i32> // CHECK-NEXT: %{{.*}} = memref.alloc() : memref<64x64xindex, strided<[2, 4], offset: 6>, 2 : i32> -// CHECK-NEXT: %{{.*}} = "memref.alloca"() <{"operandSegmentSizes" = array}> : () -> memref<64x64xindex, strided<[2, 4], offset: 6>, 2 : i32> +// CHECK-NEXT: %{{.*}} = "memref.alloca"() <{operandSegmentSizes = array}> : () -> memref<64x64xindex, strided<[2, 4], offset: 6>, 2 : i32> // CHECK-NEXT: %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}} = "memref.extract_strided_metadata"(%{{.*}}) : (memref<64x64xindex, strided<[2, 4], offset: 6>, 2 : i32>) -> (memref, index, index, index, index, index) // CHECK-NEXT: %{{.*}}, %{{.*}}, %{{.*}} = "test.op"() : () -> (index, index, index) -// CHECK-NEXT: %{{.*}} = memref.alloc(%{{.*}}) {"alignment" = 0 : i64} : memref -// CHECK-NEXT: %{{.*}} = memref.alloc(%{{.*}}, %{{.*}}, %{{.*}}) {"alignment" = 0 : i64} : memref -// CHECK-NEXT: %{{.*}} = "memref.alloca"(%{{.*}}) <{"alignment" = 0 : i64, "operandSegmentSizes" = array}> : (index) -> memref -// CHECK-NEXT: %{{.*}} = "memref.alloca"(%{{.*}}, %{{.*}}, %{{.*}}) <{"alignment" = 0 : i64, "operandSegmentSizes" = array}> : (index, index, index) -> memref +// CHECK-NEXT: %{{.*}} = memref.alloc(%{{.*}}) {alignment = 0 : i64} : memref +// CHECK-NEXT: %{{.*}} = memref.alloc(%{{.*}}, %{{.*}}, %{{.*}}) {alignment = 0 : i64} : memref +// CHECK-NEXT: %{{.*}} = "memref.alloca"(%{{.*}}) <{alignment = 0 : i64, operandSegmentSizes = array}> : (index) -> memref +// CHECK-NEXT: %{{.*}} = "memref.alloca"(%{{.*}}, %{{.*}}, %{{.*}}) <{alignment = 0 : i64, operandSegmentSizes = array}> : (index, index, index) -> memref // CHECK-NEXT: %{{.*}} = memref.collapse_shape %{{\S*}} // CHECK-SAME{LITERAL}: [[0 : i64, 1 : i64]] : memref<10x2xindex> into memref<20xindex> // CHECK-NEXT: %{{.*}} = arith.constant 2 : index @@ -94,7 +94,7 @@ builtin.module { // CHECK-NEXT: memref.dealloc %{{.*}} : memref<64x64xindex, strided<[2, 4], offset: 6>, 2 : i32> // CHECK-NEXT: %{{.*}} = "test.op"() : () -> memref<32x32xf32> // CHECK-NEXT: %{{.*}} = "test.op"() : () -> f32 -// CHECK-NEXT: %{{.*}} = "memref.atomic_rmw"(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) <{"kind" = 0 : i64}> : (f32, memref<32x32xf32>, index, index) -> f32 +// CHECK-NEXT: %{{.*}} = "memref.atomic_rmw"(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) <{kind = 0 : i64}> : (f32, memref<32x32xf32>, index, index) -> f32 // CHECK-NEXT: func.return // CHECK-NEXT: } // CHECK-NEXT: } diff --git a/tests/filecheck/dialects/memref_stream/ops.mlir b/tests/filecheck/dialects/memref_stream/ops.mlir index 3815da1f75..26e0eb1595 100644 --- a/tests/filecheck/dialects/memref_stream/ops.mlir +++ b/tests/filecheck/dialects/memref_stream/ops.mlir @@ -37,16 +37,16 @@ memref_stream.streaming_region { // CHECK-NEXT: #memref_stream.stride_pattern (d1)>, // CHECK-NEXT: #memref_stream.stride_pattern (d0, d1)> // CHECK-NEXT: ] -// CHECK-NEXT: } ins(%A, %B : memref<2xf32>, memref<3xf32>) outs(%C : memref<3x2xf64>) attrs = {"hello" = "world"} { +// CHECK-NEXT: } ins(%A, %B : memref<2xf32>, memref<3xf32>) outs(%C : memref<3x2xf64>) attrs = {hello = "world"} { // CHECK-NEXT: ^0(%a : !memref_stream.readable, %b : !memref_stream.readable, %c : !memref_stream.writable): // CHECK-NEXT: "test.op"(%a, %b, %c) : (!memref_stream.readable, !memref_stream.readable, !memref_stream.writable) -> () // CHECK-NEXT: } // CHECK-GENERIC-NEXT: %A, %B, %C, %D = "test.op"() : () -> (memref<2xf32>, memref<3xf32>, memref<3x2xf64>, f64) -// CHECK-GENERIC-NEXT: "memref_stream.streaming_region"(%A, %B, %C) <{"patterns" = [#memref_stream.stride_pattern (d0)>, #memref_stream.stride_pattern (d1)>, #memref_stream.stride_pattern (d0, d1)>], "operandSegmentSizes" = array}> ({ +// CHECK-GENERIC-NEXT: "memref_stream.streaming_region"(%A, %B, %C) <{patterns = [#memref_stream.stride_pattern (d0)>, #memref_stream.stride_pattern (d1)>, #memref_stream.stride_pattern (d0, d1)>], operandSegmentSizes = array}> ({ // CHECK-GENERIC-NEXT: ^0(%a : !memref_stream.readable, %b : !memref_stream.readable, %c : !memref_stream.writable): // CHECK-GENERIC-NEXT: "test.op"(%a, %b, %c) : (!memref_stream.readable, !memref_stream.readable, !memref_stream.writable) -> () -// CHECK-GENERIC-NEXT: }) {"hello" = "world"} : (memref<2xf32>, memref<3xf32>, memref<3x2xf64>) -> () +// CHECK-GENERIC-NEXT: }) {hello = "world"} : (memref<2xf32>, memref<3xf32>, memref<3x2xf64>) -> () memref_stream.generic { @@ -74,15 +74,15 @@ memref_stream.generic { // CHECK-NEXT: iterator_types = ["parallel", "parallel"] // CHECK-NEXT: doc = "documentation string", // CHECK-NEXT: library_call = "library call" -// CHECK-NEXT: } ins(%A, %B : memref<2xf32>, memref<3xf32>) outs(%C : memref<3x2xf64>) attrs = {"hello" = "world"} { +// CHECK-NEXT: } ins(%A, %B : memref<2xf32>, memref<3xf32>) outs(%C : memref<3x2xf64>) attrs = {hello = "world"} { // CHECK-NEXT: ^1(%arg3 : f32, %arg4 : f32, %arg5 : f32): // CHECK-NEXT: memref_stream.yield %arg3 : f32 // CHECK-NEXT: } -// CHECK-GENERIC-NEXT: "memref_stream.generic"(%A, %B, %C) <{"bounds" = [3 : index, 2 : index], "init_indices" = [], "indexing_maps" = [affine_map<(d0, d1) -> (d0)>, affine_map<(d0, d1) -> (d1)>, affine_map<(d0, d1) -> (d0, d1)>], "iterator_types" = [#memref_stream.iterator_type, #memref_stream.iterator_type], "doc" = "documentation string", "library_call" = "library call", "operandSegmentSizes" = array}> ({ +// CHECK-GENERIC-NEXT: "memref_stream.generic"(%A, %B, %C) <{bounds = [3 : index, 2 : index], init_indices = [], indexing_maps = [affine_map<(d0, d1) -> (d0)>, affine_map<(d0, d1) -> (d1)>, affine_map<(d0, d1) -> (d0, d1)>], iterator_types = [#memref_stream.iterator_type, #memref_stream.iterator_type], doc = "documentation string", library_call = "library call", operandSegmentSizes = array}> ({ // CHECK-GENERIC-NEXT: ^1(%arg3 : f32, %arg4 : f32, %arg5 : f32): // CHECK-GENERIC-NEXT: "memref_stream.yield"(%arg3) : (f32) -> () -// CHECK-GENERIC-NEXT: }) {"hello" = "world"} : (memref<2xf32>, memref<3xf32>, memref<3x2xf64>) -> () +// CHECK-GENERIC-NEXT: }) {hello = "world"} : (memref<2xf32>, memref<3xf32>, memref<3x2xf64>) -> () memref_stream.generic { bounds = [3, 2], @@ -108,7 +108,7 @@ memref_stream.generic { // CHECK-NEXT: memref_stream.yield %d : f64 // CHECK-NEXT: } -// CHECK-GENERIC-NEXT: "memref_stream.generic"(%D, %C) <{"bounds" = [3 : index, 2 : index], "init_indices" = [], "indexing_maps" = [affine_map<(d0, d1) -> ()>, affine_map<(d0, d1) -> (d0, d1)>], "iterator_types" = [#memref_stream.iterator_type, #memref_stream.iterator_type], "operandSegmentSizes" = array}> ({ +// CHECK-GENERIC-NEXT: "memref_stream.generic"(%D, %C) <{bounds = [3 : index, 2 : index], init_indices = [], indexing_maps = [affine_map<(d0, d1) -> ()>, affine_map<(d0, d1) -> (d0, d1)>], iterator_types = [#memref_stream.iterator_type, #memref_stream.iterator_type], operandSegmentSizes = array}> ({ // CHECK-GENERIC-NEXT: ^2(%d : f64, %c_1 : f64): // CHECK-GENERIC-NEXT: "memref_stream.yield"(%d) : (f64) -> () // CHECK-GENERIC-NEXT: }) : (f64, memref<3x2xf64>) -> () @@ -147,10 +147,10 @@ memref_stream.generic { // CHECK-NEXT: linalg.yield %{{.*}} : f64 // CHECK-NEXT: } -// CHECK-GENERIC-NEXT: "memref_stream.generic"(%{{.*}}, %{{.*}}, %{{.*}}, %H) <{"bounds" = [4 : index, 2 : index, 3 : index], "init_indices" = [#builtin.int<0>], "indexing_maps" = [affine_map<(d0, d1, d2) -> (d0, d2)>, affine_map<(d0, d1, d2) -> (d2, d1)>, affine_map<(d0, d1) -> (d0, d1)>], "iterator_types" = [#memref_stream.iterator_type, #memref_stream.iterator_type, #memref_stream.iterator_type], "operandSegmentSizes" = array}> ({ +// CHECK-GENERIC-NEXT: "memref_stream.generic"(%{{.*}}, %{{.*}}, %{{.*}}, %H) <{bounds = [4 : index, 2 : index, 3 : index], init_indices = [#builtin.int<0>], indexing_maps = [affine_map<(d0, d1, d2) -> (d0, d2)>, affine_map<(d0, d1, d2) -> (d2, d1)>, affine_map<(d0, d1) -> (d0, d1)>], iterator_types = [#memref_stream.iterator_type, #memref_stream.iterator_type, #memref_stream.iterator_type], operandSegmentSizes = array}> ({ // CHECK-GENERIC-NEXT: ^{{.*}}(%{{.*}} : f64, %{{.*}} : f64, %{{.*}} : f64): -// CHECK-GENERIC-NEXT: %{{.*}} = "arith.mulf"(%{{.*}}, %{{.*}}) <{"fastmath" = #arith.fastmath}> : (f64, f64) -> f64 -// CHECK-GENERIC-NEXT: %{{.*}} = "arith.addf"(%{{.*}}, %{{.*}}) <{"fastmath" = #arith.fastmath}> : (f64, f64) -> f64 +// CHECK-GENERIC-NEXT: %{{.*}} = "arith.mulf"(%{{.*}}, %{{.*}}) <{fastmath = #arith.fastmath}> : (f64, f64) -> f64 +// CHECK-GENERIC-NEXT: %{{.*}} = "arith.addf"(%{{.*}}, %{{.*}}) <{fastmath = #arith.fastmath}> : (f64, f64) -> f64 // CHECK-GENERIC-NEXT: "linalg.yield"(%{{.*}}) : (f64) -> () // CHECK-GENERIC-NEXT: }) : (memref<4x2xf64>, memref<2x3xf64>, memref<4x3xf64>, f64) -> () @@ -192,11 +192,11 @@ memref_stream.generic { // CHECK-NEXT: linalg.yield %{{.*}}, %{{.*}} : f64, f64 // CHECK-NEXT: } -// CHECK-GENERIC-NEXT: "memref_stream.generic"(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) <{"bounds" = [4 : index, 2 : index, 3 : index], "init_indices" = [#builtin.int<0>], "indexing_maps" = [affine_map<(d0, d1, d2) -> (d0, d2)>, affine_map<(d0, d1, d2) -> (d2, d1)>, affine_map<(d0, d1) -> (d0, d1)>, affine_map<(d0, d1) -> (d0, d1)>], "iterator_types" = [#memref_stream.iterator_type, #memref_stream.iterator_type, #memref_stream.iterator_type], "operandSegmentSizes" = array}> ({ +// CHECK-GENERIC-NEXT: "memref_stream.generic"(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) <{bounds = [4 : index, 2 : index, 3 : index], init_indices = [#builtin.int<0>], indexing_maps = [affine_map<(d0, d1, d2) -> (d0, d2)>, affine_map<(d0, d1, d2) -> (d2, d1)>, affine_map<(d0, d1) -> (d0, d1)>, affine_map<(d0, d1) -> (d0, d1)>], iterator_types = [#memref_stream.iterator_type, #memref_stream.iterator_type, #memref_stream.iterator_type], operandSegmentSizes = array}> ({ // CHECK-GENERIC-NEXT: ^{{.*}}(%{{.*}} : f64, %{{.*}} : f64, %{{.*}} : f64, %{{.*}} : f64): -// CHECK-GENERIC-NEXT: %{{.*}} = "arith.mulf"(%{{.*}}, %{{.*}}) <{"fastmath" = #arith.fastmath}> : (f64, f64) -> f64 -// CHECK-GENERIC-NEXT: %{{.*}} = "arith.addf"(%{{.*}}, %{{.*}}) <{"fastmath" = #arith.fastmath}> : (f64, f64) -> f64 -// CHECK-GENERIC-NEXT: %{{.*}} = "arith.addf"(%{{.*}}, %{{.*}}) <{"fastmath" = #arith.fastmath}> : (f64, f64) -> f64 +// CHECK-GENERIC-NEXT: %{{.*}} = "arith.mulf"(%{{.*}}, %{{.*}}) <{fastmath = #arith.fastmath}> : (f64, f64) -> f64 +// CHECK-GENERIC-NEXT: %{{.*}} = "arith.addf"(%{{.*}}, %{{.*}}) <{fastmath = #arith.fastmath}> : (f64, f64) -> f64 +// CHECK-GENERIC-NEXT: %{{.*}} = "arith.addf"(%{{.*}}, %{{.*}}) <{fastmath = #arith.fastmath}> : (f64, f64) -> f64 // CHECK-GENERIC-NEXT: "linalg.yield"(%{{.*}}, %{{.*}}) : (f64, f64) -> () // CHECK-GENERIC-NEXT: }) : (memref<4x2xf64>, memref<2x3xf64>, memref<4x3xf64>, memref<4x3xf64>, f64) -> () @@ -256,18 +256,18 @@ func.func @interleaved_no_init(%A0 : memref<3x5xf64>, %B0 : memref<5x8xf64>, %C0 // CHECK-NEXT: func.return %{{.*}} : memref<3x8xf64> // CHECK-NEXT: } -// CHECK-GENERIC-NEXT: "func.func"() <{"sym_name" = "interleaved_no_init", "function_type" = (memref<3x5xf64>, memref<5x8xf64>, memref<3x8xf64>) -> memref<3x8xf64>}> ({ +// CHECK-GENERIC-NEXT: "func.func"() <{sym_name = "interleaved_no_init", function_type = (memref<3x5xf64>, memref<5x8xf64>, memref<3x8xf64>) -> memref<3x8xf64>}> ({ // CHECK-GENERIC-NEXT: ^{{.*}}(%{{.*}} : memref<3x5xf64>, %{{.*}} : memref<5x8xf64>, %{{.*}} : memref<3x8xf64>): -// CHECK-GENERIC-NEXT: "memref_stream.generic"(%{{.*}}, %{{.*}}, %{{.*}}) <{"bounds" = [3 : index, 2 : index, 5 : index, 4 : index], "init_indices" = [], "indexing_maps" = [affine_map<(d0, d1, d2, d3) -> (d0, d2)>, affine_map<(d0, d1, d2, d3) -> (d2, ((d1 * 4) + d3))>, affine_map<(d0, d1, d2) -> (d0, ((d1 * 4) + d2))>], "iterator_types" = [#memref_stream.iterator_type, #memref_stream.iterator_type, #memref_stream.iterator_type, #memref_stream.iterator_type], "operandSegmentSizes" = array}> ({ +// CHECK-GENERIC-NEXT: "memref_stream.generic"(%{{.*}}, %{{.*}}, %{{.*}}) <{bounds = [3 : index, 2 : index, 5 : index, 4 : index], init_indices = [], indexing_maps = [affine_map<(d0, d1, d2, d3) -> (d0, d2)>, affine_map<(d0, d1, d2, d3) -> (d2, ((d1 * 4) + d3))>, affine_map<(d0, d1, d2) -> (d0, ((d1 * 4) + d2))>], iterator_types = [#memref_stream.iterator_type, #memref_stream.iterator_type, #memref_stream.iterator_type, #memref_stream.iterator_type], operandSegmentSizes = array}> ({ // CHECK-GENERIC-NEXT: ^{{.*}}(%{{.*}} : f64, %{{.*}} : f64, %{{.*}} : f64, %{{.*}} : f64, %{{.*}} : f64, %{{.*}} : f64, %{{.*}} : f64, %{{.*}} : f64, %{{.*}} : f64, %{{.*}} : f64, %{{.*}} : f64, %{{.*}} : f64): -// CHECK-GENERIC-NEXT: %{{.*}} = "arith.mulf"(%{{.*}}, %{{.*}}) <{"fastmath" = #arith.fastmath}> : (f64, f64) -> f64 -// CHECK-GENERIC-NEXT: %{{.*}} = "arith.mulf"(%{{.*}}, %{{.*}}) <{"fastmath" = #arith.fastmath}> : (f64, f64) -> f64 -// CHECK-GENERIC-NEXT: %{{.*}} = "arith.mulf"(%{{.*}}, %{{.*}}) <{"fastmath" = #arith.fastmath}> : (f64, f64) -> f64 -// CHECK-GENERIC-NEXT: %{{.*}} = "arith.mulf"(%{{.*}}, %{{.*}}) <{"fastmath" = #arith.fastmath}> : (f64, f64) -> f64 -// CHECK-GENERIC-NEXT: %{{.*}} = "arith.addf"(%{{.*}}, %{{.*}}) <{"fastmath" = #arith.fastmath}> : (f64, f64) -> f64 -// CHECK-GENERIC-NEXT: %{{.*}} = "arith.addf"(%{{.*}}, %{{.*}}) <{"fastmath" = #arith.fastmath}> : (f64, f64) -> f64 -// CHECK-GENERIC-NEXT: %{{.*}} = "arith.addf"(%{{.*}}, %{{.*}}) <{"fastmath" = #arith.fastmath}> : (f64, f64) -> f64 -// CHECK-GENERIC-NEXT: %{{.*}} = "arith.addf"(%{{.*}}, %{{.*}}) <{"fastmath" = #arith.fastmath}> : (f64, f64) -> f64 +// CHECK-GENERIC-NEXT: %{{.*}} = "arith.mulf"(%{{.*}}, %{{.*}}) <{fastmath = #arith.fastmath}> : (f64, f64) -> f64 +// CHECK-GENERIC-NEXT: %{{.*}} = "arith.mulf"(%{{.*}}, %{{.*}}) <{fastmath = #arith.fastmath}> : (f64, f64) -> f64 +// CHECK-GENERIC-NEXT: %{{.*}} = "arith.mulf"(%{{.*}}, %{{.*}}) <{fastmath = #arith.fastmath}> : (f64, f64) -> f64 +// CHECK-GENERIC-NEXT: %{{.*}} = "arith.mulf"(%{{.*}}, %{{.*}}) <{fastmath = #arith.fastmath}> : (f64, f64) -> f64 +// CHECK-GENERIC-NEXT: %{{.*}} = "arith.addf"(%{{.*}}, %{{.*}}) <{fastmath = #arith.fastmath}> : (f64, f64) -> f64 +// CHECK-GENERIC-NEXT: %{{.*}} = "arith.addf"(%{{.*}}, %{{.*}}) <{fastmath = #arith.fastmath}> : (f64, f64) -> f64 +// CHECK-GENERIC-NEXT: %{{.*}} = "arith.addf"(%{{.*}}, %{{.*}}) <{fastmath = #arith.fastmath}> : (f64, f64) -> f64 +// CHECK-GENERIC-NEXT: %{{.*}} = "arith.addf"(%{{.*}}, %{{.*}}) <{fastmath = #arith.fastmath}> : (f64, f64) -> f64 // CHECK-GENERIC-NEXT: "memref_stream.yield"(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) : (f64, f64, f64, f64) -> () // CHECK-GENERIC-NEXT: }) : (memref<3x5xf64>, memref<5x8xf64>, memref<3x8xf64>) -> () // CHECK-GENERIC-NEXT: "func.return"(%{{.*}}) : (memref<3x8xf64>) -> () @@ -330,19 +330,19 @@ func.func @interleaved_init(%A1 : memref<3x5xf64>, %B1 : memref<5x8xf64>, %C1 : // CHECK-NEXT: func.return %{{.*}} : memref<3x8xf64> // CHECK-NEXT: } -// CHECK-GENERIC-NEXT: "func.func"() <{"sym_name" = "interleaved_init", "function_type" = (memref<3x5xf64>, memref<5x8xf64>, memref<3x8xf64>) -> memref<3x8xf64>}> ({ +// CHECK-GENERIC-NEXT: "func.func"() <{sym_name = "interleaved_init", function_type = (memref<3x5xf64>, memref<5x8xf64>, memref<3x8xf64>) -> memref<3x8xf64>}> ({ // CHECK-GENERIC-NEXT: ^{{.*}}(%{{.*}} : memref<3x5xf64>, %{{.*}} : memref<5x8xf64>, %{{.*}} : memref<3x8xf64>): -// CHECK-GENERIC-NEXT: %{{.*}} = "arith.constant"() <{"value" = 0.000000e+00 : f64}> : () -> f64 -// CHECK-GENERIC-NEXT: "memref_stream.generic"(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) <{"bounds" = [3 : index, 2 : index, 5 : index, 4 : index], "init_indices" = [#builtin.int<0>], "indexing_maps" = [affine_map<(d0, d1, d2, d3) -> (d0, d2)>, affine_map<(d0, d1, d2, d3) -> (d2, ((d1 * 4) + d3))>, affine_map<(d0, d1, d2) -> (d0, ((d1 * 4) + d2))>], "iterator_types" = [#memref_stream.iterator_type, #memref_stream.iterator_type, #memref_stream.iterator_type, #memref_stream.iterator_type], "operandSegmentSizes" = array}> ({ +// CHECK-GENERIC-NEXT: %{{.*}} = "arith.constant"() <{value = 0.000000e+00 : f64}> : () -> f64 +// CHECK-GENERIC-NEXT: "memref_stream.generic"(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) <{bounds = [3 : index, 2 : index, 5 : index, 4 : index], init_indices = [#builtin.int<0>], indexing_maps = [affine_map<(d0, d1, d2, d3) -> (d0, d2)>, affine_map<(d0, d1, d2, d3) -> (d2, ((d1 * 4) + d3))>, affine_map<(d0, d1, d2) -> (d0, ((d1 * 4) + d2))>], iterator_types = [#memref_stream.iterator_type, #memref_stream.iterator_type, #memref_stream.iterator_type, #memref_stream.iterator_type], operandSegmentSizes = array}> ({ // CHECK-GENERIC-NEXT: ^{{.*}}(%{{.*}} : f64, %{{.*}} : f64, %{{.*}} : f64, %{{.*}} : f64, %{{.*}} : f64, %{{.*}} : f64, %{{.*}} : f64, %{{.*}} : f64, %{{.*}} : f64, %{{.*}} : f64, %{{.*}} : f64, %{{.*}} : f64): -// CHECK-GENERIC-NEXT: %{{.*}} = "arith.mulf"(%{{.*}}, %{{.*}}) <{"fastmath" = #arith.fastmath}> : (f64, f64) -> f64 -// CHECK-GENERIC-NEXT: %{{.*}} = "arith.mulf"(%{{.*}}, %{{.*}}) <{"fastmath" = #arith.fastmath}> : (f64, f64) -> f64 -// CHECK-GENERIC-NEXT: %{{.*}} = "arith.mulf"(%{{.*}}, %{{.*}}) <{"fastmath" = #arith.fastmath}> : (f64, f64) -> f64 -// CHECK-GENERIC-NEXT: %{{.*}} = "arith.mulf"(%{{.*}}, %{{.*}}) <{"fastmath" = #arith.fastmath}> : (f64, f64) -> f64 -// CHECK-GENERIC-NEXT: %{{.*}} = "arith.addf"(%{{.*}}, %{{.*}}) <{"fastmath" = #arith.fastmath}> : (f64, f64) -> f64 -// CHECK-GENERIC-NEXT: %{{.*}} = "arith.addf"(%{{.*}}, %{{.*}}) <{"fastmath" = #arith.fastmath}> : (f64, f64) -> f64 -// CHECK-GENERIC-NEXT: %{{.*}} = "arith.addf"(%{{.*}}, %{{.*}}) <{"fastmath" = #arith.fastmath}> : (f64, f64) -> f64 -// CHECK-GENERIC-NEXT: %{{.*}} = "arith.addf"(%{{.*}}, %{{.*}}) <{"fastmath" = #arith.fastmath}> : (f64, f64) -> f64 +// CHECK-GENERIC-NEXT: %{{.*}} = "arith.mulf"(%{{.*}}, %{{.*}}) <{fastmath = #arith.fastmath}> : (f64, f64) -> f64 +// CHECK-GENERIC-NEXT: %{{.*}} = "arith.mulf"(%{{.*}}, %{{.*}}) <{fastmath = #arith.fastmath}> : (f64, f64) -> f64 +// CHECK-GENERIC-NEXT: %{{.*}} = "arith.mulf"(%{{.*}}, %{{.*}}) <{fastmath = #arith.fastmath}> : (f64, f64) -> f64 +// CHECK-GENERIC-NEXT: %{{.*}} = "arith.mulf"(%{{.*}}, %{{.*}}) <{fastmath = #arith.fastmath}> : (f64, f64) -> f64 +// CHECK-GENERIC-NEXT: %{{.*}} = "arith.addf"(%{{.*}}, %{{.*}}) <{fastmath = #arith.fastmath}> : (f64, f64) -> f64 +// CHECK-GENERIC-NEXT: %{{.*}} = "arith.addf"(%{{.*}}, %{{.*}}) <{fastmath = #arith.fastmath}> : (f64, f64) -> f64 +// CHECK-GENERIC-NEXT: %{{.*}} = "arith.addf"(%{{.*}}, %{{.*}}) <{fastmath = #arith.fastmath}> : (f64, f64) -> f64 +// CHECK-GENERIC-NEXT: %{{.*}} = "arith.addf"(%{{.*}}, %{{.*}}) <{fastmath = #arith.fastmath}> : (f64, f64) -> f64 // CHECK-GENERIC-NEXT: "memref_stream.yield"(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) : (f64, f64, f64, f64) -> () // CHECK-GENERIC-NEXT: }) : (memref<3x5xf64>, memref<5x8xf64>, memref<3x8xf64>, f64) -> () // CHECK-GENERIC-NEXT: "func.return"(%{{.*}}) : (memref<3x8xf64>) -> () diff --git a/tests/filecheck/dialects/mod_arith/mod_arith.mlir b/tests/filecheck/dialects/mod_arith/mod_arith.mlir index 4a929c817d..bbc0c1eedb 100644 --- a/tests/filecheck/dialects/mod_arith/mod_arith.mlir +++ b/tests/filecheck/dialects/mod_arith/mod_arith.mlir @@ -3,6 +3,6 @@ %lhsi1, %rhsi1 = "test.op"() : () -> (i1, i1) -// CHECK: %add_res = mod_arith.add %lhsi1, %rhsi1 {"modulus" = 17 : i64} : i1 -// CHECK-GENERIC: %add_res = "mod_arith.add"(%lhsi1, %rhsi1) <{"modulus" = 17 : i64}> : (i1, i1) -> i1 +// CHECK: %add_res = mod_arith.add %lhsi1, %rhsi1 {modulus = 17 : i64} : i1 +// CHECK-GENERIC: %add_res = "mod_arith.add"(%lhsi1, %rhsi1) <{modulus = 17 : i64}> : (i1, i1) -> i1 %add_res = "mod_arith.add"(%lhsi1, %rhsi1) {modulus=17} : (i1, i1) -> i1 diff --git a/tests/filecheck/dialects/mpi/memref_compat.mlir b/tests/filecheck/dialects/mpi/memref_compat.mlir index fa6a1116b0..11f865a86b 100644 --- a/tests/filecheck/dialects/mpi/memref_compat.mlir +++ b/tests/filecheck/dialects/mpi/memref_compat.mlir @@ -7,11 +7,11 @@ // CHECK: "builtin.module"() ({ -// CHECK-NEXT: %ref = "memref.alloc"() {"alignment" = 32 : i64, "operandSegmentSizes" = array} : () -> memref<100x14x14xf64> +// CHECK-NEXT: %ref = "memref.alloc"() {alignment = 32 : i64, operandSegmentSizes = array} : () -> memref<100x14x14xf64> // CHECK-NEXT: %buff = "memref.extract_aligned_pointer_as_index"(%ref) : (memref<100x14x14xf64>) -> index // CHECK-NEXT: %buff1 = "arith.index_cast"(%buff) : (index) -> i64 // CHECK-NEXT: %buff2 = "llvm.inttoptr"(%buff1) : (i64) -> !llvm.ptr -// CHECK-NEXT: %buff3 = "arith.constant"() {"value" = 128 : i32} : () -> i32 -// CHECK-NEXT: %buff4 = "arith.constant"() {"value" = 1275070475 : i32} : () -> i32 -// CHECK-NEXT: %0 = "arith.constant"() {"value" = 1275069445 : i32} : () -> i32 +// CHECK-NEXT: %buff3 = "arith.constant"() {value = 128 : i32} : () -> i32 +// CHECK-NEXT: %buff4 = "arith.constant"() {value = 1275070475 : i32} : () -> i32 +// CHECK-NEXT: %0 = "arith.constant"() {value = 1275069445 : i32} : () -> i32 // CHECK-NEXT: }) : () -> () diff --git a/tests/filecheck/dialects/omp/ops.mlir b/tests/filecheck/dialects/omp/ops.mlir index 0da9a788b0..b83f72d6fb 100644 --- a/tests/filecheck/dialects/omp/ops.mlir +++ b/tests/filecheck/dialects/omp/ops.mlir @@ -151,39 +151,39 @@ builtin.module { // CHECK: builtin.module { // CHECK-NEXT: func.func @omp_parallel(%arg0 : memref<1xi32>, %arg1 : i1, %arg2 : i32) { -// CHECK-NEXT: "omp.parallel"(%arg1, %arg2, %arg0, %arg0) <{"operandSegmentSizes" = array, "proc_bind_val" = #omp}> ({ -// CHECK-NEXT: "omp.parallel"(%arg2, %arg0, %arg0) <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "omp.parallel"(%arg1, %arg2, %arg0, %arg0) <{operandSegmentSizes = array, proc_bind_val = #omp}> ({ +// CHECK-NEXT: "omp.parallel"(%arg2, %arg0, %arg0) <{operandSegmentSizes = array}> ({ // CHECK-NEXT: "omp.terminator"() : () -> () // CHECK-NEXT: }) : (i32, memref<1xi32>, memref<1xi32>) -> () -// CHECK-NEXT: "omp.parallel"(%arg1, %arg0, %arg0) <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "omp.parallel"(%arg1, %arg0, %arg0) <{operandSegmentSizes = array}> ({ // CHECK-NEXT: "omp.terminator"() : () -> () // CHECK-NEXT: }) : (i1, memref<1xi32>, memref<1xi32>) -> () -// CHECK-NEXT: "omp.parallel"(%arg1, %arg2) <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "omp.parallel"(%arg1, %arg2) <{operandSegmentSizes = array}> ({ // CHECK-NEXT: "omp.terminator"() : () -> () // CHECK-NEXT: }) : (i1, i32) -> () // CHECK-NEXT: "omp.terminator"() : () -> () // CHECK-NEXT: }) : (i1, i32, memref<1xi32>, memref<1xi32>) -> () -// CHECK-NEXT: "omp.parallel"(%arg0, %arg0) <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "omp.parallel"(%arg0, %arg0) <{operandSegmentSizes = array}> ({ // CHECK-NEXT: "omp.terminator"() : () -> () // CHECK-NEXT: }) : (memref<1xi32>, memref<1xi32>) -> () // CHECK-NEXT: func.return // CHECK-NEXT: } // CHECK-NEXT: func.func @omp_ordered(%arg0 : i32, %arg1 : i32, %arg2 : i32, %arg3 : i64, %arg4 : i64, %arg5 : i64, %arg6 : i64) { -// CHECK-NEXT: "omp.wsloop"() <{"operandSegmentSizes" = array, "ordered_val" = 0 : i64}> ({ +// CHECK-NEXT: "omp.wsloop"() <{operandSegmentSizes = array, ordered_val = 0 : i64}> ({ // CHECK-NEXT: "omp.loop_nest"(%arg0, %arg1, %arg2) ({ // CHECK-NEXT: ^0(%arg7 : i32): // CHECK-NEXT: omp.yield // CHECK-NEXT: }) : (i32, i32, i32) -> () // CHECK-NEXT: "omp.terminator"() : () -> () // CHECK-NEXT: }) : () -> () -// CHECK-NEXT: "omp.wsloop"() <{"operandSegmentSizes" = array, "ordered_val" = 1 : i64}> ({ +// CHECK-NEXT: "omp.wsloop"() <{operandSegmentSizes = array, ordered_val = 1 : i64}> ({ // CHECK-NEXT: "omp.loop_nest"(%arg0, %arg1, %arg2) ({ // CHECK-NEXT: ^1(%arg7_1 : i32): // CHECK-NEXT: omp.yield // CHECK-NEXT: }) : (i32, i32, i32) -> () // CHECK-NEXT: "omp.terminator"() : () -> () // CHECK-NEXT: }) : () -> () -// CHECK-NEXT: "omp.wsloop"() <{"operandSegmentSizes" = array, "ordered_val" = 2 : i64}> ({ +// CHECK-NEXT: "omp.wsloop"() <{operandSegmentSizes = array, ordered_val = 2 : i64}> ({ // CHECK-NEXT: "omp.loop_nest"(%arg0, %arg1, %arg2) ({ // CHECK-NEXT: ^2(%arg7_2 : i32): // CHECK-NEXT: omp.yield @@ -193,35 +193,35 @@ builtin.module { // CHECK-NEXT: func.return // CHECK-NEXT: } // CHECK-NEXT: func.func @omp_wsloop(%arg0 : index, %arg1 : index, %arg2 : index, %arg3 : memref<1xi32>, %arg4 : i32, %arg5 : i32) { -// CHECK-NEXT: "omp.wsloop"() <{"operandSegmentSizes" = array, "ordered_val" = 1 : i64}> ({ +// CHECK-NEXT: "omp.wsloop"() <{operandSegmentSizes = array, ordered_val = 1 : i64}> ({ // CHECK-NEXT: "omp.loop_nest"(%arg0, %arg1, %arg2) ({ // CHECK-NEXT: ^0(%arg6 : index): // CHECK-NEXT: omp.yield // CHECK-NEXT: }) : (index, index, index) -> () // CHECK-NEXT: "omp.terminator"() : () -> () // CHECK-NEXT: }) : () -> () -// CHECK-NEXT: "omp.wsloop"(%arg3, %arg4) <{"operandSegmentSizes" = array, "schedule_val" = #omp}> ({ +// CHECK-NEXT: "omp.wsloop"(%arg3, %arg4) <{operandSegmentSizes = array, schedule_val = #omp}> ({ // CHECK-NEXT: "omp.loop_nest"(%arg0, %arg1, %arg2) ({ // CHECK-NEXT: ^1(%arg6_1 : index): // CHECK-NEXT: omp.yield // CHECK-NEXT: }) : (index, index, index) -> () // CHECK-NEXT: "omp.terminator"() : () -> () // CHECK-NEXT: }) : (memref<1xi32>, i32) -> () -// CHECK-NEXT: "omp.wsloop"(%arg3, %arg3, %arg4, %arg4) <{"operandSegmentSizes" = array, "schedule_val" = #omp}> ({ +// CHECK-NEXT: "omp.wsloop"(%arg3, %arg3, %arg4, %arg4) <{operandSegmentSizes = array, schedule_val = #omp}> ({ // CHECK-NEXT: "omp.loop_nest"(%arg0, %arg1, %arg2) ({ // CHECK-NEXT: ^2(%arg6_2 : index): // CHECK-NEXT: omp.yield // CHECK-NEXT: }) : (index, index, index) -> () // CHECK-NEXT: "omp.terminator"() : () -> () // CHECK-NEXT: }) : (memref<1xi32>, memref<1xi32>, i32, i32) -> () -// CHECK-NEXT: "omp.wsloop"(%arg3, %arg4, %arg5) <{"operandSegmentSizes" = array, "ordered_val" = 2 : i64, "schedule_val" = #omp}> ({ +// CHECK-NEXT: "omp.wsloop"(%arg3, %arg4, %arg5) <{operandSegmentSizes = array, ordered_val = 2 : i64, schedule_val = #omp}> ({ // CHECK-NEXT: "omp.loop_nest"(%arg0, %arg1, %arg2) ({ // CHECK-NEXT: ^3(%arg6_3 : index): // CHECK-NEXT: omp.yield // CHECK-NEXT: }) : (index, index, index) -> () // CHECK-NEXT: "omp.terminator"() : () -> () // CHECK-NEXT: }) : (memref<1xi32>, i32, i32) -> () -// CHECK-NEXT: "omp.wsloop"() <{"nowait", "operandSegmentSizes" = array, "schedule_val" = #omp}> ({ +// CHECK-NEXT: "omp.wsloop"() <{nowait, operandSegmentSizes = array, schedule_val = #omp}> ({ // CHECK-NEXT: "omp.loop_nest"(%arg0, %arg1, %arg2) ({ // CHECK-NEXT: ^4(%arg6_4 : index): // CHECK-NEXT: omp.yield @@ -231,63 +231,63 @@ builtin.module { // CHECK-NEXT: func.return // CHECK-NEXT: } // CHECK-NEXT: func.func @omp_wsloop_pretty(%arg0 : index, %arg1 : index, %arg2 : index, %arg3 : memref<1xi32>, %arg4 : i32, %arg5 : i32, %arg6 : i16) { -// CHECK-NEXT: "omp.wsloop"() <{"operandSegmentSizes" = array, "ordered_val" = 2 : i64}> ({ +// CHECK-NEXT: "omp.wsloop"() <{operandSegmentSizes = array, ordered_val = 2 : i64}> ({ // CHECK-NEXT: "omp.loop_nest"(%arg0, %arg1, %arg2) ({ // CHECK-NEXT: ^0(%arg7 : index): // CHECK-NEXT: omp.yield // CHECK-NEXT: }) : (index, index, index) -> () // CHECK-NEXT: "omp.terminator"() : () -> () // CHECK-NEXT: }) : () -> () -// CHECK-NEXT: "omp.wsloop"(%arg3, %arg4) <{"operandSegmentSizes" = array, "schedule_val" = #omp}> ({ +// CHECK-NEXT: "omp.wsloop"(%arg3, %arg4) <{operandSegmentSizes = array, schedule_val = #omp}> ({ // CHECK-NEXT: "omp.loop_nest"(%arg0, %arg1, %arg2) ({ // CHECK-NEXT: ^1(%arg7_1 : index): // CHECK-NEXT: omp.yield // CHECK-NEXT: }) : (index, index, index) -> () // CHECK-NEXT: "omp.terminator"() : () -> () // CHECK-NEXT: }) : (memref<1xi32>, i32) -> () -// CHECK-NEXT: "omp.wsloop"(%arg3, %arg4, %arg5) <{"operandSegmentSizes" = array, "ordered_val" = 2 : i64, "schedule_val" = #omp}> ({ +// CHECK-NEXT: "omp.wsloop"(%arg3, %arg4, %arg5) <{operandSegmentSizes = array, ordered_val = 2 : i64, schedule_val = #omp}> ({ // CHECK-NEXT: "omp.loop_nest"(%arg0, %arg1, %arg2) ({ // CHECK-NEXT: ^2(%arg7_2 : index): // CHECK-NEXT: omp.yield // CHECK-NEXT: }) : (index, index, index) -> () // CHECK-NEXT: "omp.terminator"() : () -> () // CHECK-NEXT: }) : (memref<1xi32>, i32, i32) -> () -// CHECK-NEXT: "omp.wsloop"(%arg3, %arg4, %arg5) <{"operandSegmentSizes" = array, "ordered_val" = 2 : i64, "schedule_modifier" = #omp, "schedule_val" = #omp}> ({ +// CHECK-NEXT: "omp.wsloop"(%arg3, %arg4, %arg5) <{operandSegmentSizes = array, ordered_val = 2 : i64, schedule_modifier = #omp, schedule_val = #omp}> ({ // CHECK-NEXT: "omp.loop_nest"(%arg0, %arg1, %arg2) ({ // CHECK-NEXT: ^3(%arg7_3 : index): // CHECK-NEXT: omp.yield // CHECK-NEXT: }) : (index, index, index) -> () // CHECK-NEXT: "omp.terminator"() : () -> () // CHECK-NEXT: }) : (memref<1xi32>, i32, i32) -> () -// CHECK-NEXT: "omp.wsloop"(%arg3, %arg4, %arg6) <{"operandSegmentSizes" = array, "ordered_val" = 2 : i64, "schedule_modifier" = #omp, "schedule_val" = #omp}> ({ +// CHECK-NEXT: "omp.wsloop"(%arg3, %arg4, %arg6) <{operandSegmentSizes = array, ordered_val = 2 : i64, schedule_modifier = #omp, schedule_val = #omp}> ({ // CHECK-NEXT: "omp.loop_nest"(%arg0, %arg1, %arg2) ({ // CHECK-NEXT: ^4(%arg7_4 : index): // CHECK-NEXT: omp.yield // CHECK-NEXT: }) : (index, index, index) -> () // CHECK-NEXT: "omp.terminator"() : () -> () // CHECK-NEXT: }) : (memref<1xi32>, i32, i16) -> () -// CHECK-NEXT: "omp.wsloop"() <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "omp.wsloop"() <{operandSegmentSizes = array}> ({ // CHECK-NEXT: "omp.loop_nest"(%arg0, %arg1, %arg2) ({ // CHECK-NEXT: ^5(%arg7_5 : index): // CHECK-NEXT: omp.yield // CHECK-NEXT: }) : (index, index, index) -> () // CHECK-NEXT: "omp.terminator"() : () -> () // CHECK-NEXT: }) : () -> () -// CHECK-NEXT: "omp.wsloop"() <{"inclusive", "operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "omp.wsloop"() <{inclusive, operandSegmentSizes = array}> ({ // CHECK-NEXT: "omp.loop_nest"(%arg0, %arg1, %arg2) ({ // CHECK-NEXT: ^6(%arg7_6 : index): // CHECK-NEXT: omp.yield // CHECK-NEXT: }) : (index, index, index) -> () // CHECK-NEXT: "omp.terminator"() : () -> () // CHECK-NEXT: }) : () -> () -// CHECK-NEXT: "omp.wsloop"() <{"nowait", "operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "omp.wsloop"() <{nowait, operandSegmentSizes = array}> ({ // CHECK-NEXT: "omp.loop_nest"(%arg0, %arg1, %arg2) ({ // CHECK-NEXT: ^7(%arg7_7 : index): // CHECK-NEXT: omp.yield // CHECK-NEXT: }) : (index, index, index) -> () // CHECK-NEXT: "omp.terminator"() : () -> () // CHECK-NEXT: }) : () -> () -// CHECK-NEXT: "omp.wsloop"() <{"nowait", "operandSegmentSizes" = array, "order_val" = #omp}> ({ +// CHECK-NEXT: "omp.wsloop"() <{nowait, operandSegmentSizes = array, order_val = #omp}> ({ // CHECK-NEXT: "omp.loop_nest"(%arg0, %arg1, %arg2) ({ // CHECK-NEXT: ^8(%arg7_8 : index): // CHECK-NEXT: omp.yield diff --git a/tests/filecheck/dialects/onnx/onnx_ops.mlir b/tests/filecheck/dialects/onnx/onnx_ops.mlir index 2d6a068bb6..ed12de520a 100644 --- a/tests/filecheck/dialects/onnx/onnx_ops.mlir +++ b/tests/filecheck/dialects/onnx/onnx_ops.mlir @@ -16,58 +16,58 @@ %t27, %t28, %t29 = "test.op"(): () -> (tensor<1x320xf32>, tensor<50x320xf32>, tensor<50xf32>) %res_add = "onnx.Add"(%t0, %t1) {onnx_node_name = "/Add"} : (tensor<1x2x6xf32>, tensor<1x2x6xf32>) -> tensor<1x2x6xf32> -// CHECK: %res_add = onnx.Add(%t0, %t1) {"onnx_node_name" = "/Add"} : (tensor<1x2x6xf32>, tensor<1x2x6xf32>) -> tensor<1x2x6xf32> +// CHECK: %res_add = onnx.Add(%t0, %t1) {onnx_node_name = "/Add"} : (tensor<1x2x6xf32>, tensor<1x2x6xf32>) -> tensor<1x2x6xf32> %res_sub = "onnx.Sub"(%t2, %t3) {onnx_node_name = "/Sub"}: (tensor<3x2xf32>, tensor<1x2xf32>) -> tensor<3x2xf32> -// CHECK: %res_sub = onnx.Sub(%t2, %t3) {"onnx_node_name" = "/Sub"} : (tensor<3x2xf32>, tensor<1x2xf32>) -> tensor<3x2xf32> +// CHECK: %res_sub = onnx.Sub(%t2, %t3) {onnx_node_name = "/Sub"} : (tensor<3x2xf32>, tensor<1x2xf32>) -> tensor<3x2xf32> %res_mul = "onnx.Mul"(%t4, %t5) {onnx_node_name = "/Mul"}: (tensor<3x1x2xf32>, tensor<3x4x1xf32>) -> tensor<3x4x2xf32> -// CHECK: %res_mul = onnx.Mul(%t4, %t5) {"onnx_node_name" = "/Mul"} : (tensor<3x1x2xf32>, tensor<3x4x1xf32>) -> tensor<3x4x2xf32> +// CHECK: %res_mul = onnx.Mul(%t4, %t5) {onnx_node_name = "/Mul"} : (tensor<3x1x2xf32>, tensor<3x4x1xf32>) -> tensor<3x4x2xf32> %res_div = "onnx.Div"(%t6, %t7) {onnx_node_name = "/Div"}: (tensor<1x5x1x3xf32>, tensor<2x1x6x3xf32>) -> tensor<2x5x6x3xf32> -// CHECK: %res_div = onnx.Div(%t6, %t7) {"onnx_node_name" = "/Div"} : (tensor<1x5x1x3xf32>, tensor<2x1x6x3xf32>) -> tensor<2x5x6x3xf32> +// CHECK: %res_div = onnx.Div(%t6, %t7) {onnx_node_name = "/Div"} : (tensor<1x5x1x3xf32>, tensor<2x1x6x3xf32>) -> tensor<2x5x6x3xf32> %res_relu = "onnx.Relu"(%t8) {onnx_node_name = "/Relu"}: (tensor<3x4xf32>) -> tensor<3x4xf32> -// CHECK: %res_relu = onnx.Relu(%t8) {"onnx_node_name" = "/Relu"} : (tensor<3x4xf32>) -> tensor<3x4xf32> +// CHECK: %res_relu = onnx.Relu(%t8) {onnx_node_name = "/Relu"} : (tensor<3x4xf32>) -> tensor<3x4xf32> %res_gemm = "onnx.Gemm"(%t9, %t10, %t11) {onnx_node_name = "/Gemm"}: (tensor<2x2xf32>, tensor<2x2xf32>, tensor<2x2xf32>) -> tensor<2x2xf32> -// CHECK: %res_gemm = onnx.Gemm(%t9, %t10, %t11) {"onnx_node_name" = "/Gemm"} : (tensor<2x2xf32>, tensor<2x2xf32>, tensor<2x2xf32>) -> tensor<2x2xf32> +// CHECK: %res_gemm = onnx.Gemm(%t9, %t10, %t11) {onnx_node_name = "/Gemm"} : (tensor<2x2xf32>, tensor<2x2xf32>, tensor<2x2xf32>) -> tensor<2x2xf32> %res_gemm_1 = "onnx.Gemm"(%t12, %t13, %t14) {onnx_node_name = "/Gemm"}: (tensor<5x3xf32>, tensor<3x2xf32>, tensor<5x2xf32>) -> tensor<5x2xf32> -// CHECK: %res_gemm_1 = onnx.Gemm(%t12, %t13, %t14) {"onnx_node_name" = "/Gemm"} : (tensor<5x3xf32>, tensor<3x2xf32>, tensor<5x2xf32>) -> tensor<5x2xf32> +// CHECK: %res_gemm_1 = onnx.Gemm(%t12, %t13, %t14) {onnx_node_name = "/Gemm"} : (tensor<5x3xf32>, tensor<3x2xf32>, tensor<5x2xf32>) -> tensor<5x2xf32> %res_reshape = "onnx.Reshape"(%t15, %t16) {onnx_node_name = "/Reshape", "allowzero" = 1 : i64}: (tensor<48x256x64xf32>, tensor<3xi64>) -> tensor<48x256x64xf32> -//CHECK: res_reshape = onnx.Reshape(%t15, %t16) {"onnx_node_name" = "/Reshape", "allowzero" = 1 : i64} : (tensor<48x256x64xf32>, tensor<3xi64>) -> tensor<48x256x64xf32> +//CHECK: res_reshape = onnx.Reshape(%t15, %t16) {onnx_node_name = "/Reshape", allowzero = 1 : i64} : (tensor<48x256x64xf32>, tensor<3xi64>) -> tensor<48x256x64xf32> %res_reshape_1 = "onnx.Reshape"(%t17, %t18) {onnx_node_name = "/Reshape"}: (tensor<1x2x3x4x5xf32>, tensor<5xi64>) -> tensor<1x120xf32> -//CHECK: %res_reshape_1 = onnx.Reshape(%t17, %t18) {"onnx_node_name" = "/Reshape"} : (tensor<1x2x3x4x5xf32>, tensor<5xi64>) -> tensor<1x120xf32> +//CHECK: %res_reshape_1 = onnx.Reshape(%t17, %t18) {onnx_node_name = "/Reshape"} : (tensor<1x2x3x4x5xf32>, tensor<5xi64>) -> tensor<1x120xf32> %res_abs = "onnx.Abs"(%t19) {onnx_node_name = "/Abs"}: (tensor<10x10xf32>) -> tensor<10x10xf32> -// CHECK: %res_abs = onnx.Abs(%t19) {"onnx_node_name" = "/Abs"} : (tensor<10x10xf32>) -> tensor<10x10xf32> +// CHECK: %res_abs = onnx.Abs(%t19) {onnx_node_name = "/Abs"} : (tensor<10x10xf32>) -> tensor<10x10xf32> %res_conv = "onnx.Conv"(%t20, %t21, %t22) {onnx_node_name = "/Conv", "auto_pad" = "NOTSET", "group" = 1 : i64, "kernel_shape" = [3 : i64, 3 : i64], "dilations" = [1 : i64, 1 : i64], "strides" = [1 : i64, 1 : i64], "pads" = [1 : i64, 1 : i64, 1: i64, 1 : i64]}: (tensor<1x1x5x5xf32>, tensor<1x1x3x3xf32>, none) -> tensor<1x1x5x5xf32> -//CHECK: %res_conv = onnx.Conv(%t20, %t21, %t22) {"onnx_node_name" = "/Conv", "auto_pad" = "NOTSET", "group" = 1 : i64, "kernel_shape" = [3 : i64, 3 : i64], "dilations" = [1 : i64, 1 : i64], "strides" = [1 : i64, 1 : i64], "pads" = [1 : i64, 1 : i64, 1 : i64, 1 : i64]} : (tensor<1x1x5x5xf32>, tensor<1x1x3x3xf32>, none) -> tensor<1x1x5x5xf32> +//CHECK: %res_conv = onnx.Conv(%t20, %t21, %t22) {onnx_node_name = "/Conv", auto_pad = "NOTSET", group = 1 : i64, kernel_shape = [3 : i64, 3 : i64], dilations = [1 : i64, 1 : i64], strides = [1 : i64, 1 : i64], pads = [1 : i64, 1 : i64, 1 : i64, 1 : i64]} : (tensor<1x1x5x5xf32>, tensor<1x1x3x3xf32>, none) -> tensor<1x1x5x5xf32> %res_conv_1 = "onnx.Conv"(%t20, %t21, %t22) {onnx_node_name = "/Conv", "auto_pad" = "NOTSET", "group" = 1 : i64, "kernel_shape" = [3 : i64, 3 : i64], "dilations" = [1 : i64, 1 : i64], "strides" = [1 : i64, 1 : i64], "pads" = [0 : i64, 0 : i64, 0: i64, 0 : i64]}: (tensor<1x1x5x5xf32>, tensor<1x1x3x3xf32>, none) -> tensor<1x1x3x3xf32> -//CHECK: %res_conv_1 = onnx.Conv(%t20, %t21, %t22) {"onnx_node_name" = "/Conv", "auto_pad" = "NOTSET", "group" = 1 : i64, "kernel_shape" = [3 : i64, 3 : i64], "dilations" = [1 : i64, 1 : i64], "strides" = [1 : i64, 1 : i64], "pads" = [0 : i64, 0 : i64, 0 : i64, 0 : i64]} : (tensor<1x1x5x5xf32>, tensor<1x1x3x3xf32>, none) -> tensor<1x1x3x3xf32> +//CHECK: %res_conv_1 = onnx.Conv(%t20, %t21, %t22) {onnx_node_name = "/Conv", auto_pad = "NOTSET", group = 1 : i64, kernel_shape = [3 : i64, 3 : i64], dilations = [1 : i64, 1 : i64], strides = [1 : i64, 1 : i64], pads = [0 : i64, 0 : i64, 0 : i64, 0 : i64]} : (tensor<1x1x5x5xf32>, tensor<1x1x3x3xf32>, none) -> tensor<1x1x3x3xf32> %res_conv_2 = "onnx.Conv"(%t20, %t21, %t22) {onnx_node_name = "/Conv", "auto_pad" = "SAME_LOWER", "group" = 1 : i64, "kernel_shape" = [3 : i64, 3 : i64], "dilations" = [1 : i64, 1 : i64], "strides" = [2 : i64, 2 : i64], "pads" = [0 : i64, 0 : i64, 0: i64, 0 : i64]}: (tensor<1x1x5x5xf32>, tensor<1x1x3x3xf32>, none) -> tensor<1x1x3x3xf32> -//CHECK: %res_conv_2 = onnx.Conv(%t20, %t21, %t22) {"onnx_node_name" = "/Conv", "auto_pad" = "SAME_LOWER", "group" = 1 : i64, "kernel_shape" = [3 : i64, 3 : i64], "dilations" = [1 : i64, 1 : i64], "strides" = [2 : i64, 2 : i64], "pads" = [0 : i64, 0 : i64, 0 : i64, 0 : i64]} : (tensor<1x1x5x5xf32>, tensor<1x1x3x3xf32>, none) -> tensor<1x1x3x3xf32> +//CHECK: %res_conv_2 = onnx.Conv(%t20, %t21, %t22) {onnx_node_name = "/Conv", auto_pad = "SAME_LOWER", group = 1 : i64, kernel_shape = [3 : i64, 3 : i64], dilations = [1 : i64, 1 : i64], strides = [2 : i64, 2 : i64], pads = [0 : i64, 0 : i64, 0 : i64, 0 : i64]} : (tensor<1x1x5x5xf32>, tensor<1x1x3x3xf32>, none) -> tensor<1x1x3x3xf32> %res_conv_3 = "onnx.Conv"(%t23, %t24, %t25) {onnx_node_name = "/Conv", "auto_pad" = "NOTSET", "group" = 1 : i64, "kernel_shape" = [3 : i64, 3 : i64], "dilations" = [1 : i64, 1 : i64], "strides" = [2 : i64, 2 : i64], "pads" = [1 : i64, 1 : i64, 1 : i64, 1 : i64]}: (tensor<1x1x7x5xf32>, tensor<1x1x3x3xf32>, none) -> tensor<1x1x4x3xf32> -//CHECK: %res_conv_3 = onnx.Conv(%t23, %t24, %t25) {"onnx_node_name" = "/Conv", "auto_pad" = "NOTSET", "group" = 1 : i64, "kernel_shape" = [3 : i64, 3 : i64], "dilations" = [1 : i64, 1 : i64], "strides" = [2 : i64, 2 : i64], "pads" = [1 : i64, 1 : i64, 1 : i64, 1 : i64]} : (tensor<1x1x7x5xf32>, tensor<1x1x3x3xf32>, none) -> tensor<1x1x4x3xf32> +//CHECK: %res_conv_3 = onnx.Conv(%t23, %t24, %t25) {onnx_node_name = "/Conv", auto_pad = "NOTSET", group = 1 : i64, kernel_shape = [3 : i64, 3 : i64], dilations = [1 : i64, 1 : i64], strides = [2 : i64, 2 : i64], pads = [1 : i64, 1 : i64, 1 : i64, 1 : i64]} : (tensor<1x1x7x5xf32>, tensor<1x1x3x3xf32>, none) -> tensor<1x1x4x3xf32> %res_conv_4 = "onnx.Conv"(%t23, %t24, %t25) {onnx_node_name = "/Conv", "auto_pad" = "NOTSET", "group" = 1 : i64, "kernel_shape" = [3 : i64, 3 : i64], "dilations" = [1 : i64, 1 : i64], "strides" = [2 : i64, 2 : i64], "pads" = [0 : i64, 0 : i64, 0 : i64, 0 : i64]}: (tensor<1x1x7x5xf32>, tensor<1x1x3x3xf32>, none) -> tensor<1x1x3x2xf32> -//CHECK: %res_conv_4 = onnx.Conv(%t23, %t24, %t25) {"onnx_node_name" = "/Conv", "auto_pad" = "NOTSET", "group" = 1 : i64, "kernel_shape" = [3 : i64, 3 : i64], "dilations" = [1 : i64, 1 : i64], "strides" = [2 : i64, 2 : i64], "pads" = [0 : i64, 0 : i64, 0 : i64, 0 : i64]} : (tensor<1x1x7x5xf32>, tensor<1x1x3x3xf32>, none) -> tensor<1x1x3x2xf32> +//CHECK: %res_conv_4 = onnx.Conv(%t23, %t24, %t25) {onnx_node_name = "/Conv", auto_pad = "NOTSET", group = 1 : i64, kernel_shape = [3 : i64, 3 : i64], dilations = [1 : i64, 1 : i64], strides = [2 : i64, 2 : i64], pads = [0 : i64, 0 : i64, 0 : i64, 0 : i64]} : (tensor<1x1x7x5xf32>, tensor<1x1x3x3xf32>, none) -> tensor<1x1x3x2xf32> %res_conv_5 = "onnx.Conv"(%t23, %t24, %t25) {onnx_node_name = "/Conv", "auto_pad" = "NOTSET", "group" = 1 : i64, "kernel_shape" = [3 : i64, 3 : i64], "dilations" = [1 : i64, 1 : i64], "strides" = [2 : i64, 2 : i64], "pads" = [1 : i64, 0 : i64, 1 : i64, 0 : i64]}: (tensor<1x1x7x5xf32>, tensor<1x1x3x3xf32>, none) -> tensor<1x1x4x2xf32> -//CHECK: %res_conv_5 = onnx.Conv(%t23, %t24, %t25) {"onnx_node_name" = "/Conv", "auto_pad" = "NOTSET", "group" = 1 : i64, "kernel_shape" = [3 : i64, 3 : i64], "dilations" = [1 : i64, 1 : i64], "strides" = [2 : i64, 2 : i64], "pads" = [1 : i64, 0 : i64, 1 : i64, 0 : i64]} : (tensor<1x1x7x5xf32>, tensor<1x1x3x3xf32>, none) -> tensor<1x1x4x2xf32> +//CHECK: %res_conv_5 = onnx.Conv(%t23, %t24, %t25) {onnx_node_name = "/Conv", auto_pad = "NOTSET", group = 1 : i64, kernel_shape = [3 : i64, 3 : i64], dilations = [1 : i64, 1 : i64], strides = [2 : i64, 2 : i64], pads = [1 : i64, 0 : i64, 1 : i64, 0 : i64]} : (tensor<1x1x7x5xf32>, tensor<1x1x3x3xf32>, none) -> tensor<1x1x4x2xf32> %res_max_pool_single_out = "onnx.MaxPoolSingleOut"(%t26) {onnx_node_name = "/MaxPoolSingleOut", "auto_pad" = "VALID", "ceil_mode" = 0 : i64, "kernel_shape" = [3 : i64, 3 : i64], "dilations" = [1 : i64, 1 : i64], "pads" = [0 : i64, 0 : i64, 0 : i64, 0 : i64], "storage_order" = 0 : i64, "strides" = [1 : i64, 1 : i64]}: (tensor<5x5x32x32xf32>) -> tensor<5x5x30x30xf32> -//CHECK: %res_max_pool_single_out = onnx.MaxPoolSingleOut(%t26) {"onnx_node_name" = "/MaxPoolSingleOut", "auto_pad" = "VALID", "ceil_mode" = 0 : i64, "kernel_shape" = [3 : i64, 3 : i64], "dilations" = [1 : i64, 1 : i64], "pads" = [0 : i64, 0 : i64, 0 : i64, 0 : i64], "storage_order" = 0 : i64, "strides" = [1 : i64, 1 : i64]} : (tensor<5x5x32x32xf32>) -> tensor<5x5x30x30xf32> +//CHECK: %res_max_pool_single_out = onnx.MaxPoolSingleOut(%t26) {onnx_node_name = "/MaxPoolSingleOut", auto_pad = "VALID", ceil_mode = 0 : i64, kernel_shape = [3 : i64, 3 : i64], dilations = [1 : i64, 1 : i64], pads = [0 : i64, 0 : i64, 0 : i64, 0 : i64], storage_order = 0 : i64, strides = [1 : i64, 1 : i64]} : (tensor<5x5x32x32xf32>) -> tensor<5x5x30x30xf32> "onnx.EntryPoint"() {onnx_node_name = "/EntryPoint", "func" = @main_graph} : () -> () -//CHECK: "onnx.EntryPoint"() {"onnx_node_name" = "/EntryPoint", "func" = @main_graph} : () -> () +//CHECK: "onnx.EntryPoint"() {onnx_node_name = "/EntryPoint", func = @main_graph} : () -> () %res_constant = onnx.Constant dense<1> : tensor<1xi64> //CHECK: %res_constant = onnx.Constant dense<1> : tensor<1xi64> @@ -76,16 +76,16 @@ //CHECK: %res_constant_1 = onnx.Constant dense<[5, 5, 16, 2]> : tensor<4xi64> %res_gemm_2 = "onnx.Gemm"(%t27, %t28, %t29) {onnx_node_name = "/Gemm", "alpha" = 1.000000e+00 : f32, "beta" = 1.000000e+00 : f32, "transA" = 0 : si64, "transB" = 1 : si64}: (tensor<1x320xf32>, tensor<50x320xf32>, tensor<50xf32>) -> tensor<1x50xf32> -// CHECK: %res_gemm_2 = onnx.Gemm(%t27, %t28, %t29) {"onnx_node_name" = "/Gemm", "alpha" = 1.000000e+00 : f32, "beta" = 1.000000e+00 : f32, "transA" = 0 : si64, "transB" = 1 : si64} : (tensor<1x320xf32>, tensor<50x320xf32>, tensor<50xf32>) -> tensor<1x50xf32> +// CHECK: %res_gemm_2 = onnx.Gemm(%t27, %t28, %t29) {onnx_node_name = "/Gemm", alpha = 1.000000e+00 : f32, beta = 1.000000e+00 : f32, transA = 0 : si64, transB = 1 : si64} : (tensor<1x320xf32>, tensor<50x320xf32>, tensor<50xf32>) -> tensor<1x50xf32> %res_matmul = "onnx.MatMul"(%t9, %t10) {onnx_node_name = "/MatMul"}: (tensor<2x2xf32>, tensor<2x2xf32>) -> tensor<2x2xf32> -// CHECK: %res_matmul = onnx.MatMul(%t9, %t10) {"onnx_node_name" = "/MatMul"} : (tensor<2x2xf32>, tensor<2x2xf32>) -> tensor<2x2xf32> +// CHECK: %res_matmul = onnx.MatMul(%t9, %t10) {onnx_node_name = "/MatMul"} : (tensor<2x2xf32>, tensor<2x2xf32>) -> tensor<2x2xf32> %res_transpose = "onnx.Transpose"(%t8) {onnx_node_name = "/Transpose", "perm" = [1 : i64, 0 : i64]}: (tensor<3x4xf32>) -> tensor<4x3xf32> -// CHECK: %res_transpose = onnx.Transpose(%t8) {"onnx_node_name" = "/Transpose", "perm" = [1 : i64, 0 : i64]} : (tensor<3x4xf32>) -> tensor<4x3xf32> +// CHECK: %res_transpose = onnx.Transpose(%t8) {onnx_node_name = "/Transpose", perm = [1 : i64, 0 : i64]} : (tensor<3x4xf32>) -> tensor<4x3xf32> %res_squeeze = "onnx.Squeeze"(%t0) {onnx_node_name = "/Squeeze", "axes" = 0}: (tensor<1x2x6xf32>) -> tensor<2x6xf32> -// CHECK: %res_squeeze = onnx.Squeeze(%t0) {"onnx_node_name" = "/Squeeze", "axes" = 0 : i64} : (tensor<1x2x6xf32>) -> tensor<2x6xf32> +// CHECK: %res_squeeze = onnx.Squeeze(%t0) {onnx_node_name = "/Squeeze", axes = 0 : i64} : (tensor<1x2x6xf32>) -> tensor<2x6xf32> %res_sigmoid = "onnx.Sigmoid"(%t8) {onnx_node_name = "/Sigmoid"}: (tensor<3x4xf32>) -> tensor<3x4xf32> -// CHECK: %res_sigmoid = onnx.Sigmoid(%t8) {"onnx_node_name" = "/Sigmoid"} : (tensor<3x4xf32>) -> tensor<3x4xf32> +// CHECK: %res_sigmoid = onnx.Sigmoid(%t8) {onnx_node_name = "/Sigmoid"} : (tensor<3x4xf32>) -> tensor<3x4xf32> diff --git a/tests/filecheck/dialects/printf/printf_basics.mlir b/tests/filecheck/dialects/printf/printf_basics.mlir index aca6cfd6e3..63b977f03a 100644 --- a/tests/filecheck/dialects/printf/printf_basics.mlir +++ b/tests/filecheck/dialects/printf/printf_basics.mlir @@ -19,6 +19,6 @@ builtin.module { // CHECK-NEXT: %byte = "test.op"() : () -> i8 // CHECK-NEXT: printf.print_format "Uses vals twice {} {} {} {}", %1 : i32, %0 : i32, %1 : i32, %0 : i32 // CHECK-NEXT: printf.print_format "{}", %0 : i32 -// CHECK-NEXT: printf.print_format "{}", %0 : i32 {"unit"} +// CHECK-NEXT: printf.print_format "{}", %0 : i32 {unit} // CHECK-NEXT: "printf.print_char"(%{{.*}}) : (i8) -> () // CHECK-NEXT: "printf.print_int"(%{{.*}}) : (i32) -> () diff --git a/tests/filecheck/dialects/printf/printf_to_llvm.mlir b/tests/filecheck/dialects/printf/printf_to_llvm.mlir index 37a371ae7a..4e210807ca 100644 --- a/tests/filecheck/dialects/printf/printf_to_llvm.mlir +++ b/tests/filecheck/dialects/printf/printf_to_llvm.mlir @@ -11,14 +11,14 @@ builtin.module { }) {sym_name = "main", function_type=() -> ()} : () -> () } -// CHECK: %{{\d+}} = "llvm.mlir.addressof"() <{"global_name" = @Hello_f_842f9d94ff2eba9703926bef3c2bc5f427db9871}> : () -> !llvm.ptr +// CHECK: %{{\d+}} = "llvm.mlir.addressof"() <{global_name = @Hello_f_842f9d94ff2eba9703926bef3c2bc5f427db9871}> : () -> !llvm.ptr -// CHECK: "llvm.call"(%{{\d+}}, %{{\d+}}, %{{\d+}}) <{"callee" = @printf{{.*}}}> : (!llvm.ptr, f64, i32) -> () +// CHECK: "llvm.call"(%{{\d+}}, %{{\d+}}, %{{\d+}}) <{callee = @printf{{.*}}}> : (!llvm.ptr, f64, i32) -> () -// CHECK: "llvm.func"() <{"sym_name" = "printf", "function_type" = !llvm.func, "CConv" = #llvm.cconv, "linkage" = #llvm.linkage<"external">, "visibility_" = 0 : i64}> ({ +// CHECK: "llvm.func"() <{sym_name = "printf", function_type = !llvm.func, CConv = #llvm.cconv, linkage = #llvm.linkage<"external">, visibility_ = 0 : i64}> ({ // CHECK-NEXT: }) : () -> () -// CHECK: "llvm.mlir.global"() <{"global_type" = !llvm.array<14 x i8>, "sym_name" = "Hello_f_842f9d94ff2eba9703926bef3c2bc5f427db9871", "linkage" = #llvm.linkage<"internal">, "addr_space" = 0 : i32, "constant", "value" = dense<[72, 101, 108, 108, 111, 58, 32, 37, 102, 32, 37, 105, 10, 0]> : tensor<14xi8>}> ({ +// CHECK: "llvm.mlir.global"() <{global_type = !llvm.array<14 x i8>, sym_name = "Hello_f_842f9d94ff2eba9703926bef3c2bc5f427db9871", linkage = #llvm.linkage<"internal">, addr_space = 0 : i32, constant, value = dense<[72, 101, 108, 108, 111, 58, 32, 37, 102, 32, 37, 105, 10, 0]> : tensor<14xi8>}> ({ // CHECK-NEXT: }) : () -> () // ----- diff --git a/tests/filecheck/dialects/qref/ops.mlir b/tests/filecheck/dialects/qref/ops.mlir index 16c933f0f0..4239d48de6 100644 --- a/tests/filecheck/dialects/qref/ops.mlir +++ b/tests/filecheck/dialects/qref/ops.mlir @@ -23,6 +23,6 @@ qref.cnot %q0, %q1 // CHECK-GENERIC: %q0, %q1 = "qref.alloc"() : () -> (!qref.qubit, !qref.qubit) // CHECK-GENERIC-NEXT: "qref.h"(%q0) : (!qref.qubit) -> () -// CHECK-GENERIC-NEXT: "qref.rz"(%q1) <{"angle" = !quantum.angle}> : (!qref.qubit) -> () +// CHECK-GENERIC-NEXT: "qref.rz"(%q1) <{angle = !quantum.angle}> : (!qref.qubit) -> () // CHECK-GENERIC-NEXT: "qref.cnot"(%q0, %q1) : (!qref.qubit, !qref.qubit) -> () // CHECK-GENERIC-NEXT: %0 = "qref.measure"(%q0) : (!qref.qubit) -> i1 diff --git a/tests/filecheck/dialects/qssa/ops.mlir b/tests/filecheck/dialects/qssa/ops.mlir index 9b9183a311..545cc93bf1 100644 --- a/tests/filecheck/dialects/qssa/ops.mlir +++ b/tests/filecheck/dialects/qssa/ops.mlir @@ -23,6 +23,6 @@ // CHECK-GENERIC: %q0, %q1 = "qssa.alloc"() : () -> (!qssa.qubit, !qssa.qubit) // CHECK-GENERIC-NEXT: %q2 = "qssa.h"(%q0) : (!qssa.qubit) -> !qssa.qubit -// CHECK-GENERIC-NEXT: %q3 = "qssa.rz"(%q1) <{"angle" = !quantum.angle}> : (!qssa.qubit) -> !qssa.qubit +// CHECK-GENERIC-NEXT: %q3 = "qssa.rz"(%q1) <{angle = !quantum.angle}> : (!qssa.qubit) -> !qssa.qubit // CHECK-GENERIC-NEXT: %q4, %q5 = "qssa.cnot"(%q2, %q3) : (!qssa.qubit, !qssa.qubit) -> (!qssa.qubit, !qssa.qubit) // CHECK-GENERIC-NEXT: %0 = "qssa.measure"(%q4) : (!qssa.qubit) -> i1 diff --git a/tests/filecheck/dialects/quantum/attrs.mlir b/tests/filecheck/dialects/quantum/attrs.mlir index bda781a364..d669db5ed5 100644 --- a/tests/filecheck/dialects/quantum/attrs.mlir +++ b/tests/filecheck/dialects/quantum/attrs.mlir @@ -2,28 +2,28 @@ "test.op"() {"angle" = !quantum.angle<0>} : () -> () -// CHECK: "test.op"() {"angle" = !quantum.angle<0>} : () -> () +// CHECK: "test.op"() {angle = !quantum.angle<0>} : () -> () "test.op"() {"angle" = !quantum.angle} : () -> () -// CHECK-NEXT: "test.op"() {"angle" = !quantum.angle} : () -> () +// CHECK-NEXT: "test.op"() {angle = !quantum.angle} : () -> () "test.op"() {"angle" = !quantum.angle<2pi>} : () -> () -// CHECK-NEXT: "test.op"() {"angle" = !quantum.angle<0>} : () -> () +// CHECK-NEXT: "test.op"() {angle = !quantum.angle<0>} : () -> () "test.op"() {"angle" = !quantum.angle} : () -> () -// CHECK-NEXT: "test.op"() {"angle" = !quantum.angle} : () -> () +// CHECK-NEXT: "test.op"() {angle = !quantum.angle} : () -> () "test.op"() {"angle" = !quantum.angle<3pi:2>} : () -> () -// CHECK-NEXT: "test.op"() {"angle" = !quantum.angle<3pi:2>} : () -> () +// CHECK-NEXT: "test.op"() {angle = !quantum.angle<3pi:2>} : () -> () "test.op"() {"angle" = !quantum.angle<5pi:2>} : () -> () -// CHECK-NEXT: "test.op"() {"angle" = !quantum.angle} : () -> () +// CHECK-NEXT: "test.op"() {angle = !quantum.angle} : () -> () "test.op"() {"angle" = !quantum.angle<-pi>} : () -> () -// CHECK-NEXT: "test.op"() {"angle" = !quantum.angle} : () -> () +// CHECK-NEXT: "test.op"() {angle = !quantum.angle} : () -> () diff --git a/tests/filecheck/dialects/riscv/riscv_ops.mlir b/tests/filecheck/dialects/riscv/riscv_ops.mlir index 8823bec412..eda5589682 100644 --- a/tests/filecheck/dialects/riscv/riscv_ops.mlir +++ b/tests/filecheck/dialects/riscv/riscv_ops.mlir @@ -187,7 +187,7 @@ riscv.assembly_section ".text" attributes {"foo" = i32} { %nested_li = riscv.li 1 : !riscv.reg } - // CHECK-NEXT: riscv.assembly_section ".text" attributes {"foo" = i32} { + // CHECK-NEXT: riscv.assembly_section ".text" attributes {foo = i32} { // CHECK-NEXT: %{{.*}} = riscv.li 1 : !riscv.reg // CHECK-NEXT: } @@ -200,7 +200,7 @@ // Custom instruction %custom0, %custom1 = riscv.custom_assembly_instruction %0, %1 {"instruction_name" = "hello"} : (!riscv.reg, !riscv.reg) -> (!riscv.reg, !riscv.reg) - // CHECK-NEXT: %custom0, %custom1 = riscv.custom_assembly_instruction %0, %1 {"instruction_name" = "hello"} : (!riscv.reg, !riscv.reg) -> (!riscv.reg, !riscv.reg) + // CHECK-NEXT: %custom0, %custom1 = riscv.custom_assembly_instruction %0, %1 {instruction_name = "hello"} : (!riscv.reg, !riscv.reg) -> (!riscv.reg, !riscv.reg) // RV32F: 8 “F” Standard Extension for Single-Precision Floating-Point, Version 2.0 %f0 = riscv.get_float_register : !riscv.freg @@ -356,17 +356,17 @@ // CHECK-GENERIC-NEXT: "riscv_func.func"() ({ // CHECK-GENERIC-NEXT: %0 = "riscv.get_register"() : () -> !riscv.reg // CHECK-GENERIC-NEXT: %1 = "riscv.get_register"() : () -> !riscv.reg -// CHECK-GENERIC-NEXT: %addi = "riscv.addi"(%0) {"immediate" = 1 : si12} : (!riscv.reg) -> !riscv.reg -// CHECK-GENERIC-NEXT: %slti = "riscv.slti"(%0) {"immediate" = 1 : si12} : (!riscv.reg) -> !riscv.reg -// CHECK-GENERIC-NEXT: %sltiu = "riscv.sltiu"(%0) {"immediate" = 1 : si12} : (!riscv.reg) -> !riscv.reg -// CHECK-GENERIC-NEXT: %andi = "riscv.andi"(%0) {"immediate" = 1 : si12} : (!riscv.reg) -> !riscv.reg -// CHECK-GENERIC-NEXT: %ori = "riscv.ori"(%0) {"immediate" = 1 : si12} : (!riscv.reg) -> !riscv.reg -// CHECK-GENERIC-NEXT: %xori = "riscv.xori"(%0) {"immediate" = 1 : si12} : (!riscv.reg) -> !riscv.reg -// CHECK-GENERIC-NEXT: %slli = "riscv.slli"(%0) {"immediate" = 1 : ui5} : (!riscv.reg) -> !riscv.reg -// CHECK-GENERIC-NEXT: %srli = "riscv.srli"(%0) {"immediate" = 1 : ui5} : (!riscv.reg) -> !riscv.reg -// CHECK-GENERIC-NEXT: %srai = "riscv.srai"(%0) {"immediate" = 1 : ui5} : (!riscv.reg) -> !riscv.reg -// CHECK-GENERIC-NEXT: %lui = "riscv.lui"() {"immediate" = 1 : i20} : () -> !riscv.reg -// CHECK-GENERIC-NEXT: %auipc = "riscv.auipc"() {"immediate" = 1 : i20} : () -> !riscv.reg +// CHECK-GENERIC-NEXT: %addi = "riscv.addi"(%0) {immediate = 1 : si12} : (!riscv.reg) -> !riscv.reg +// CHECK-GENERIC-NEXT: %slti = "riscv.slti"(%0) {immediate = 1 : si12} : (!riscv.reg) -> !riscv.reg +// CHECK-GENERIC-NEXT: %sltiu = "riscv.sltiu"(%0) {immediate = 1 : si12} : (!riscv.reg) -> !riscv.reg +// CHECK-GENERIC-NEXT: %andi = "riscv.andi"(%0) {immediate = 1 : si12} : (!riscv.reg) -> !riscv.reg +// CHECK-GENERIC-NEXT: %ori = "riscv.ori"(%0) {immediate = 1 : si12} : (!riscv.reg) -> !riscv.reg +// CHECK-GENERIC-NEXT: %xori = "riscv.xori"(%0) {immediate = 1 : si12} : (!riscv.reg) -> !riscv.reg +// CHECK-GENERIC-NEXT: %slli = "riscv.slli"(%0) {immediate = 1 : ui5} : (!riscv.reg) -> !riscv.reg +// CHECK-GENERIC-NEXT: %srli = "riscv.srli"(%0) {immediate = 1 : ui5} : (!riscv.reg) -> !riscv.reg +// CHECK-GENERIC-NEXT: %srai = "riscv.srai"(%0) {immediate = 1 : ui5} : (!riscv.reg) -> !riscv.reg +// CHECK-GENERIC-NEXT: %lui = "riscv.lui"() {immediate = 1 : i20} : () -> !riscv.reg +// CHECK-GENERIC-NEXT: %auipc = "riscv.auipc"() {immediate = 1 : i20} : () -> !riscv.reg // CHECK-GENERIC-NEXT: %mv = "riscv.mv"(%0) : (!riscv.reg) -> !riscv.reg // CHECK-GENERIC-NEXT: %add = "riscv.add"(%0, %1) : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-GENERIC-NEXT: %slt = "riscv.slt"(%0, %1) : (!riscv.reg, !riscv.reg) -> !riscv.reg @@ -379,42 +379,42 @@ // CHECK-GENERIC-NEXT: %sub = "riscv.sub"(%0, %1) : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-GENERIC-NEXT: %sra = "riscv.sra"(%0, %1) : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-GENERIC-NEXT: "riscv.nop"() : () -> () -// CHECK-GENERIC-NEXT: "riscv.jal"() {"immediate" = 1 : si20} : () -> () -// CHECK-GENERIC-NEXT: "riscv.jal"() {"immediate" = 1 : si20, "rd" = !riscv.reg} : () -> () -// CHECK-GENERIC-NEXT: "riscv.jal"() {"immediate" = #riscv.label<"label">} : () -> () -// CHECK-GENERIC-NEXT: "riscv.j"() {"immediate" = 1 : si20} : () -> () -// CHECK-GENERIC-NEXT: "riscv.j"() {"immediate" = #riscv.label<"label">} : () -> () -// CHECK-GENERIC-NEXT: "riscv.jalr"(%0) {"immediate" = 1 : si12} : (!riscv.reg) -> () -// CHECK-GENERIC-NEXT: "riscv.jalr"(%0) {"immediate" = 1 : si12, "rd" = !riscv.reg} : (!riscv.reg) -> () -// CHECK-GENERIC-NEXT: "riscv.jalr"(%0) {"immediate" = #riscv.label<"label">} : (!riscv.reg) -> () +// CHECK-GENERIC-NEXT: "riscv.jal"() {immediate = 1 : si20} : () -> () +// CHECK-GENERIC-NEXT: "riscv.jal"() {immediate = 1 : si20, rd = !riscv.reg} : () -> () +// CHECK-GENERIC-NEXT: "riscv.jal"() {immediate = #riscv.label<"label">} : () -> () +// CHECK-GENERIC-NEXT: "riscv.j"() {immediate = 1 : si20} : () -> () +// CHECK-GENERIC-NEXT: "riscv.j"() {immediate = #riscv.label<"label">} : () -> () +// CHECK-GENERIC-NEXT: "riscv.jalr"(%0) {immediate = 1 : si12} : (!riscv.reg) -> () +// CHECK-GENERIC-NEXT: "riscv.jalr"(%0) {immediate = 1 : si12, rd = !riscv.reg} : (!riscv.reg) -> () +// CHECK-GENERIC-NEXT: "riscv.jalr"(%0) {immediate = #riscv.label<"label">} : (!riscv.reg) -> () // CHECK-GENERIC-NEXT: "riscv.ret"() : () -> () // CHECK-GENERIC-NEXT: ^0(%2 : !riscv.reg, %3 : !riscv.reg): -// CHECK-GENERIC-NEXT: "riscv.beq"(%0, %1) {"offset" = 1 : si12} : (!riscv.reg, !riscv.reg) -> () -// CHECK-GENERIC-NEXT: "riscv.bne"(%0, %1) {"offset" = 1 : si12} : (!riscv.reg, !riscv.reg) -> () -// CHECK-GENERIC-NEXT: "riscv.blt"(%0, %1) {"offset" = 1 : si12} : (!riscv.reg, !riscv.reg) -> () -// CHECK-GENERIC-NEXT: "riscv.bge"(%0, %1) {"offset" = 1 : si12} : (!riscv.reg, !riscv.reg) -> () -// CHECK-GENERIC-NEXT: "riscv.bltu"(%0, %1) {"offset" = 1 : si12} : (!riscv.reg, !riscv.reg) -> () -// CHECK-GENERIC-NEXT: "riscv.bgeu"(%0, %1) {"offset" = 1 : si12} : (!riscv.reg, !riscv.reg) -> () -// CHECK-GENERIC-NEXT: %lb = "riscv.lb"(%0) {"immediate" = 1 : si12} : (!riscv.reg) -> !riscv.reg -// CHECK-GENERIC-NEXT: %lbu = "riscv.lbu"(%0) {"immediate" = 1 : si12} : (!riscv.reg) -> !riscv.reg -// CHECK-GENERIC-NEXT: %lh = "riscv.lh"(%0) {"immediate" = 1 : si12} : (!riscv.reg) -> !riscv.reg -// CHECK-GENERIC-NEXT: %lhu = "riscv.lhu"(%0) {"immediate" = 1 : si12} : (!riscv.reg) -> !riscv.reg -// CHECK-GENERIC-NEXT: %lw = "riscv.lw"(%0) {"immediate" = 1 : si12} : (!riscv.reg) -> !riscv.reg -// CHECK-GENERIC-NEXT: "riscv.sb"(%0, %1) {"immediate" = 1 : si12} : (!riscv.reg, !riscv.reg) -> () -// CHECK-GENERIC-NEXT: "riscv.sh"(%0, %1) {"immediate" = 1 : si12} : (!riscv.reg, !riscv.reg) -> () -// CHECK-GENERIC-NEXT: "riscv.sw"(%0, %1) {"immediate" = 1 : si12} : (!riscv.reg, !riscv.reg) -> () -// CHECK-GENERIC-NEXT: %csrrw_rw = "riscv.csrrw"(%0) {"csr" = 1024 : i32} : (!riscv.reg) -> !riscv.reg -// CHECK-GENERIC-NEXT: %csrrw_w = "riscv.csrrw"(%0) {"csr" = 1024 : i32, "writeonly"} : (!riscv.reg) -> !riscv.reg -// CHECK-GENERIC-NEXT: %csrrs_rw = "riscv.csrrs"(%0) {"csr" = 1024 : i32} : (!riscv.reg) -> !riscv.reg -// CHECK-GENERIC-NEXT: %csrrs_r = "riscv.csrrs"(%0) {"csr" = 1024 : i32, "readonly"} : (!riscv.reg) -> !riscv.reg -// CHECK-GENERIC-NEXT: %csrrc_rw = "riscv.csrrc"(%0) {"csr" = 1024 : i32} : (!riscv.reg) -> !riscv.reg -// CHECK-GENERIC-NEXT: %csrrc_r = "riscv.csrrc"(%0) {"csr" = 1024 : i32, "readonly"} : (!riscv.reg) -> !riscv.reg -// CHECK-GENERIC-NEXT: %csrrsi_rw = "riscv.csrrsi"() {"csr" = 1024 : i32, "immediate" = 8 : i32} : () -> !riscv.reg -// CHECK-GENERIC-NEXT: %csrrsi_r = "riscv.csrrsi"() {"csr" = 1024 : i32, "immediate" = 0 : i32} : () -> !riscv.reg -// CHECK-GENERIC-NEXT: %csrrci_rw = "riscv.csrrci"() {"csr" = 1024 : i32, "immediate" = 8 : i32} : () -> !riscv.reg -// CHECK-GENERIC-NEXT: %csrrci_r = "riscv.csrrci"() {"csr" = 1024 : i32, "immediate" = 0 : i32} : () -> !riscv.reg -// CHECK-GENERIC-NEXT: %csrrwi_rw = "riscv.csrrwi"() {"csr" = 1024 : i32, "immediate" = 1 : i32} : () -> !riscv.reg -// CHECK-GENERIC-NEXT: %csrrwi_w = "riscv.csrrwi"() {"csr" = 1024 : i32, "immediate" = 1 : i32, "writeonly"} : () -> !riscv.reg +// CHECK-GENERIC-NEXT: "riscv.beq"(%0, %1) {offset = 1 : si12} : (!riscv.reg, !riscv.reg) -> () +// CHECK-GENERIC-NEXT: "riscv.bne"(%0, %1) {offset = 1 : si12} : (!riscv.reg, !riscv.reg) -> () +// CHECK-GENERIC-NEXT: "riscv.blt"(%0, %1) {offset = 1 : si12} : (!riscv.reg, !riscv.reg) -> () +// CHECK-GENERIC-NEXT: "riscv.bge"(%0, %1) {offset = 1 : si12} : (!riscv.reg, !riscv.reg) -> () +// CHECK-GENERIC-NEXT: "riscv.bltu"(%0, %1) {offset = 1 : si12} : (!riscv.reg, !riscv.reg) -> () +// CHECK-GENERIC-NEXT: "riscv.bgeu"(%0, %1) {offset = 1 : si12} : (!riscv.reg, !riscv.reg) -> () +// CHECK-GENERIC-NEXT: %lb = "riscv.lb"(%0) {immediate = 1 : si12} : (!riscv.reg) -> !riscv.reg +// CHECK-GENERIC-NEXT: %lbu = "riscv.lbu"(%0) {immediate = 1 : si12} : (!riscv.reg) -> !riscv.reg +// CHECK-GENERIC-NEXT: %lh = "riscv.lh"(%0) {immediate = 1 : si12} : (!riscv.reg) -> !riscv.reg +// CHECK-GENERIC-NEXT: %lhu = "riscv.lhu"(%0) {immediate = 1 : si12} : (!riscv.reg) -> !riscv.reg +// CHECK-GENERIC-NEXT: %lw = "riscv.lw"(%0) {immediate = 1 : si12} : (!riscv.reg) -> !riscv.reg +// CHECK-GENERIC-NEXT: "riscv.sb"(%0, %1) {immediate = 1 : si12} : (!riscv.reg, !riscv.reg) -> () +// CHECK-GENERIC-NEXT: "riscv.sh"(%0, %1) {immediate = 1 : si12} : (!riscv.reg, !riscv.reg) -> () +// CHECK-GENERIC-NEXT: "riscv.sw"(%0, %1) {immediate = 1 : si12} : (!riscv.reg, !riscv.reg) -> () +// CHECK-GENERIC-NEXT: %csrrw_rw = "riscv.csrrw"(%0) {csr = 1024 : i32} : (!riscv.reg) -> !riscv.reg +// CHECK-GENERIC-NEXT: %csrrw_w = "riscv.csrrw"(%0) {csr = 1024 : i32, writeonly} : (!riscv.reg) -> !riscv.reg +// CHECK-GENERIC-NEXT: %csrrs_rw = "riscv.csrrs"(%0) {csr = 1024 : i32} : (!riscv.reg) -> !riscv.reg +// CHECK-GENERIC-NEXT: %csrrs_r = "riscv.csrrs"(%0) {csr = 1024 : i32, readonly} : (!riscv.reg) -> !riscv.reg +// CHECK-GENERIC-NEXT: %csrrc_rw = "riscv.csrrc"(%0) {csr = 1024 : i32} : (!riscv.reg) -> !riscv.reg +// CHECK-GENERIC-NEXT: %csrrc_r = "riscv.csrrc"(%0) {csr = 1024 : i32, readonly} : (!riscv.reg) -> !riscv.reg +// CHECK-GENERIC-NEXT: %csrrsi_rw = "riscv.csrrsi"() {csr = 1024 : i32, immediate = 8 : i32} : () -> !riscv.reg +// CHECK-GENERIC-NEXT: %csrrsi_r = "riscv.csrrsi"() {csr = 1024 : i32, immediate = 0 : i32} : () -> !riscv.reg +// CHECK-GENERIC-NEXT: %csrrci_rw = "riscv.csrrci"() {csr = 1024 : i32, immediate = 8 : i32} : () -> !riscv.reg +// CHECK-GENERIC-NEXT: %csrrci_r = "riscv.csrrci"() {csr = 1024 : i32, immediate = 0 : i32} : () -> !riscv.reg +// CHECK-GENERIC-NEXT: %csrrwi_rw = "riscv.csrrwi"() {csr = 1024 : i32, immediate = 1 : i32} : () -> !riscv.reg +// CHECK-GENERIC-NEXT: %csrrwi_w = "riscv.csrrwi"() {csr = 1024 : i32, immediate = 1 : i32, writeonly} : () -> !riscv.reg // CHECK-GENERIC-NEXT: "riscv.wfi"() : () -> () // CHECK-GENERIC-NEXT: %mul = "riscv.mul"(%0, %1) : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-GENERIC-NEXT: %mulh = "riscv.mulh"(%0, %1) : (!riscv.reg, !riscv.reg) -> !riscv.reg @@ -424,18 +424,18 @@ // CHECK-GENERIC-NEXT: %divu = "riscv.divu"(%0, %1) : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-GENERIC-NEXT: %rem = "riscv.rem"(%0, %1) : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-GENERIC-NEXT: %remu = "riscv.remu"(%0, %1) : (!riscv.reg, !riscv.reg) -> !riscv.reg -// CHECK-GENERIC-NEXT: %li = "riscv.li"() {"immediate" = 1 : i32} : () -> !riscv.reg +// CHECK-GENERIC-NEXT: %li = "riscv.li"() {immediate = 1 : i32} : () -> !riscv.reg // CHECK-GENERIC-NEXT: "riscv.ecall"() : () -> () // CHECK-GENERIC-NEXT: "riscv.ebreak"() : () -> () -// CHECK-GENERIC-NEXT: "riscv.directive"() {"directive" = ".bss"} : () -> () -// CHECK-GENERIC-NEXT: "riscv.directive"() {"directive" = ".align", "value" = "2"} : () -> () +// CHECK-GENERIC-NEXT: "riscv.directive"() {directive = ".bss"} : () -> () +// CHECK-GENERIC-NEXT: "riscv.directive"() {directive = ".align", value = "2"} : () -> () // CHECK-GENERIC-NEXT: "riscv.assembly_section"() ({ -// CHECK-GENERIC-NEXT: %nested_li = "riscv.li"() {"immediate" = 1 : i32} : () -> !riscv.reg -// CHECK-GENERIC-NEXT: }) {"directive" = ".text", "foo" = i32} : () -> () +// CHECK-GENERIC-NEXT: %nested_li = "riscv.li"() {immediate = 1 : i32} : () -> !riscv.reg +// CHECK-GENERIC-NEXT: }) {directive = ".text", foo = i32} : () -> () // CHECK-GENERIC-NEXT: "riscv.assembly_section"() ({ -// CHECK-GENERIC-NEXT: %nested_li = "riscv.li"() {"immediate" = 1 : i32} : () -> !riscv.reg -// CHECK-GENERIC-NEXT: }) {"directive" = ".text"} : () -> () -// CHECK-GENERIC-NEXT: %custom0, %custom1 = "riscv.custom_assembly_instruction"(%0, %1) {"instruction_name" = "hello"} : (!riscv.reg, !riscv.reg) -> (!riscv.reg, !riscv.reg) +// CHECK-GENERIC-NEXT: %nested_li = "riscv.li"() {immediate = 1 : i32} : () -> !riscv.reg +// CHECK-GENERIC-NEXT: }) {directive = ".text"} : () -> () +// CHECK-GENERIC-NEXT: %custom0, %custom1 = "riscv.custom_assembly_instruction"(%0, %1) {instruction_name = "hello"} : (!riscv.reg, !riscv.reg) -> (!riscv.reg, !riscv.reg) // CHECK-GENERIC-NEXT: %f0 = "riscv.get_float_register"() : () -> !riscv.freg // CHECK-GENERIC-NEXT: %f1 = "riscv.get_float_register"() : () -> !riscv.freg // CHECK-GENERIC-NEXT: %f2 = "riscv.get_float_register"() : () -> !riscv.freg @@ -444,58 +444,58 @@ // CHECK-GENERIC-NEXT: %fmsub_s = "riscv.fmsub.s"(%f0, %f1, %f2) : (!riscv.freg, !riscv.freg, !riscv.freg) -> !riscv.freg // CHECK-GENERIC-NEXT: %fnmsub_s = "riscv.fnmsub.s"(%f0, %f1, %f2) : (!riscv.freg, !riscv.freg, !riscv.freg) -> !riscv.freg // CHECK-GENERIC-NEXT: %fnmadd_s = "riscv.fnmadd.s"(%f0, %f1, %f2) : (!riscv.freg, !riscv.freg, !riscv.freg) -> !riscv.freg -// CHECK-GENERIC-NEXT: %fadd_s = "riscv.fadd.s"(%f0, %f1) {"fastmath" = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg -// CHECK-GENERIC-NEXT: %fsub_s = "riscv.fsub.s"(%f0, %f1) {"fastmath" = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg -// CHECK-GENERIC-NEXT: %fmul_s = "riscv.fmul.s"(%f0, %f1) {"fastmath" = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg -// CHECK-GENERIC-NEXT: %fdiv_s = "riscv.fdiv.s"(%f0, %f1) {"fastmath" = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg -// CHECK-GENERIC-NEXT: %fadd_s_fm = "riscv.fadd.s"(%f0, %f1) {"fastmath" = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg -// CHECK-GENERIC-NEXT: %fsub_s_fm = "riscv.fsub.s"(%f0, %f1) {"fastmath" = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg -// CHECK-GENERIC-NEXT: %fmul_s_fm = "riscv.fmul.s"(%f0, %f1) {"fastmath" = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg -// CHECK-GENERIC-NEXT: %fdiv_s_fm = "riscv.fdiv.s"(%f0, %f1) {"fastmath" = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg +// CHECK-GENERIC-NEXT: %fadd_s = "riscv.fadd.s"(%f0, %f1) {fastmath = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg +// CHECK-GENERIC-NEXT: %fsub_s = "riscv.fsub.s"(%f0, %f1) {fastmath = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg +// CHECK-GENERIC-NEXT: %fmul_s = "riscv.fmul.s"(%f0, %f1) {fastmath = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg +// CHECK-GENERIC-NEXT: %fdiv_s = "riscv.fdiv.s"(%f0, %f1) {fastmath = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg +// CHECK-GENERIC-NEXT: %fadd_s_fm = "riscv.fadd.s"(%f0, %f1) {fastmath = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg +// CHECK-GENERIC-NEXT: %fsub_s_fm = "riscv.fsub.s"(%f0, %f1) {fastmath = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg +// CHECK-GENERIC-NEXT: %fmul_s_fm = "riscv.fmul.s"(%f0, %f1) {fastmath = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg +// CHECK-GENERIC-NEXT: %fdiv_s_fm = "riscv.fdiv.s"(%f0, %f1) {fastmath = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg // CHECK-GENERIC-NEXT: %fsqrt_s = "riscv.fsqrt.s"(%f0) : (!riscv.freg) -> !riscv.freg // CHECK-GENERIC-NEXT: %fsgnj_s = "riscv.fsgnj.s"(%f0, %f1) : (!riscv.freg, !riscv.freg) -> !riscv.freg // CHECK-GENERIC-NEXT: %fsgnjn_s = "riscv.fsgnjn.s"(%f0, %f1) : (!riscv.freg, !riscv.freg) -> !riscv.freg // CHECK-GENERIC-NEXT: %fsgnjx_s = "riscv.fsgnjx.s"(%f0, %f1) : (!riscv.freg, !riscv.freg) -> !riscv.freg -// CHECK-GENERIC-NEXT: %fmin_s = "riscv.fmin.s"(%f0, %f1) {"fastmath" = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg -// CHECK-GENERIC-NEXT: %fmax_s = "riscv.fmax.s"(%f0, %f1) {"fastmath" = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg -// CHECK-GENERIC-NEXT: %fmin_s_fm = "riscv.fmin.s"(%f0, %f1) {"fastmath" = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg -// CHECK-GENERIC-NEXT: %fmax_s_fm = "riscv.fmax.s"(%f0, %f1) {"fastmath" = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg +// CHECK-GENERIC-NEXT: %fmin_s = "riscv.fmin.s"(%f0, %f1) {fastmath = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg +// CHECK-GENERIC-NEXT: %fmax_s = "riscv.fmax.s"(%f0, %f1) {fastmath = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg +// CHECK-GENERIC-NEXT: %fmin_s_fm = "riscv.fmin.s"(%f0, %f1) {fastmath = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg +// CHECK-GENERIC-NEXT: %fmax_s_fm = "riscv.fmax.s"(%f0, %f1) {fastmath = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg // CHECK-GENERIC-NEXT: %fcvt_w_s = "riscv.fcvt.w.s"(%f0) : (!riscv.freg) -> !riscv.reg // CHECK-GENERIC-NEXT: %fcvt_wu_s = "riscv.fcvt.wu.s"(%f0) : (!riscv.freg) -> !riscv.reg // CHECK-GENERIC-NEXT: %fmv_x_w = "riscv.fmv.x.w"(%f0) : (!riscv.freg) -> !riscv.reg -// CHECK-GENERIC-NEXT: %feq_s = "riscv.feq.s"(%f0, %f1) {"fastmath" = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.reg -// CHECK-GENERIC-NEXT: %flt_s = "riscv.flt.s"(%f0, %f1) {"fastmath" = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.reg -// CHECK-GENERIC-NEXT: %fle_s = "riscv.fle.s"(%f0, %f1) {"fastmath" = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.reg -// CHECK-GENERIC-NEXT: %feq_s_fm = "riscv.feq.s"(%f0, %f1) {"fastmath" = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.reg -// CHECK-GENERIC-NEXT: %flt_s_fm = "riscv.flt.s"(%f0, %f1) {"fastmath" = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.reg -// CHECK-GENERIC-NEXT: %fle_s_fm = "riscv.fle.s"(%f0, %f1) {"fastmath" = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.reg +// CHECK-GENERIC-NEXT: %feq_s = "riscv.feq.s"(%f0, %f1) {fastmath = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.reg +// CHECK-GENERIC-NEXT: %flt_s = "riscv.flt.s"(%f0, %f1) {fastmath = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.reg +// CHECK-GENERIC-NEXT: %fle_s = "riscv.fle.s"(%f0, %f1) {fastmath = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.reg +// CHECK-GENERIC-NEXT: %feq_s_fm = "riscv.feq.s"(%f0, %f1) {fastmath = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.reg +// CHECK-GENERIC-NEXT: %flt_s_fm = "riscv.flt.s"(%f0, %f1) {fastmath = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.reg +// CHECK-GENERIC-NEXT: %fle_s_fm = "riscv.fle.s"(%f0, %f1) {fastmath = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.reg // CHECK-GENERIC-NEXT: %fclass_s = "riscv.fclass.s"(%f0) : (!riscv.freg) -> !riscv.reg // CHECK-GENERIC-NEXT: %fcvt_s_w = "riscv.fcvt.s.w"(%0) : (!riscv.reg) -> !riscv.freg // CHECK-GENERIC-NEXT: %fcvt_s_wu = "riscv.fcvt.s.wu"(%0) : (!riscv.reg) -> !riscv.freg // CHECK-GENERIC-NEXT: %fmv_w_x = "riscv.fmv.w.x"(%0) : (!riscv.reg) -> !riscv.freg -// CHECK-GENERIC-NEXT: %flw = "riscv.flw"(%0) {"immediate" = 1 : i12} : (!riscv.reg) -> !riscv.freg -// CHECK-GENERIC-NEXT: "riscv.fsw"(%0, %f0) {"immediate" = 1 : i12} : (!riscv.reg, !riscv.freg) -> () -// CHECK-GENERIC-NEXT: %fld = "riscv.fld"(%0) {"immediate" = 1 : i12} : (!riscv.reg) -> !riscv.freg -// CHECK-GENERIC-NEXT: "riscv.fsd"(%0, %f0) {"immediate" = 1 : i12} : (!riscv.reg, !riscv.freg) -> () +// CHECK-GENERIC-NEXT: %flw = "riscv.flw"(%0) {immediate = 1 : i12} : (!riscv.reg) -> !riscv.freg +// CHECK-GENERIC-NEXT: "riscv.fsw"(%0, %f0) {immediate = 1 : i12} : (!riscv.reg, !riscv.freg) -> () +// CHECK-GENERIC-NEXT: %fld = "riscv.fld"(%0) {immediate = 1 : i12} : (!riscv.reg) -> !riscv.freg +// CHECK-GENERIC-NEXT: "riscv.fsd"(%0, %f0) {immediate = 1 : i12} : (!riscv.reg, !riscv.freg) -> () // CHECK-GENERIC-NEXT: %fmv_d = "riscv.fmv.d"(%f0) : (!riscv.freg) -> !riscv.freg // CHECK-GENERIC-NEXT: %vfadd_s = "riscv.vfadd.s"(%f0, %f1) : (!riscv.freg, !riscv.freg) -> !riscv.freg // CHECK-GENERIC-NEXT: %vfmul_s = "riscv.vfmul.s"(%f0, %f1) : (!riscv.freg, !riscv.freg) -> !riscv.freg -// CHECK-GENERIC-NEXT: %{{.*}} = "riscv.fadd.d"(%{{.*}}, %{{.*}}) {"fastmath" = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg -// CHECK-GENERIC-NEXT: %{{.*}} = "riscv.fsub.d"(%{{.*}}, %{{.*}}) {"fastmath" = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg -// CHECK-GENERIC-NEXT: %{{.*}} = "riscv.fmul.d"(%{{.*}}, %{{.*}}) {"fastmath" = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg -// CHECK-GENERIC-NEXT: %{{.*}} = "riscv.fdiv.d"(%{{.*}}, %{{.*}}) {"fastmath" = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg -// CHECK-GENERIC-NEXT: %{{.*}} = "riscv.fadd.d"(%{{.*}}, %{{.*}}) {"fastmath" = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg -// CHECK-GENERIC-NEXT: %{{.*}} = "riscv.fsub.d"(%{{.*}}, %{{.*}}) {"fastmath" = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg -// CHECK-GENERIC-NEXT: %{{.*}} = "riscv.fmul.d"(%{{.*}}, %{{.*}}) {"fastmath" = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg -// CHECK-GENERIC-NEXT: %{{.*}} = "riscv.fdiv.d"(%{{.*}}, %{{.*}}) {"fastmath" = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg +// CHECK-GENERIC-NEXT: %{{.*}} = "riscv.fadd.d"(%{{.*}}, %{{.*}}) {fastmath = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg +// CHECK-GENERIC-NEXT: %{{.*}} = "riscv.fsub.d"(%{{.*}}, %{{.*}}) {fastmath = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg +// CHECK-GENERIC-NEXT: %{{.*}} = "riscv.fmul.d"(%{{.*}}, %{{.*}}) {fastmath = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg +// CHECK-GENERIC-NEXT: %{{.*}} = "riscv.fdiv.d"(%{{.*}}, %{{.*}}) {fastmath = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg +// CHECK-GENERIC-NEXT: %{{.*}} = "riscv.fadd.d"(%{{.*}}, %{{.*}}) {fastmath = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg +// CHECK-GENERIC-NEXT: %{{.*}} = "riscv.fsub.d"(%{{.*}}, %{{.*}}) {fastmath = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg +// CHECK-GENERIC-NEXT: %{{.*}} = "riscv.fmul.d"(%{{.*}}, %{{.*}}) {fastmath = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg +// CHECK-GENERIC-NEXT: %{{.*}} = "riscv.fdiv.d"(%{{.*}}, %{{.*}}) {fastmath = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg // CHECK-GENERIC-NEXT: %fmadd_d = "riscv.fmadd.d"(%f0, %f1, %f2) : (!riscv.freg, !riscv.freg, !riscv.freg) -> !riscv.freg // CHECK-GENERIC-NEXT: %fmsub_d = "riscv.fmsub.d"(%f0, %f1, %f2) : (!riscv.freg, !riscv.freg, !riscv.freg) -> !riscv.freg -// CHECK-GENERIC-NEXT: %fmin_d = "riscv.fmin.d"(%f0, %f1) {"fastmath" = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg -// CHECK-GENERIC-NEXT: %fmax_d = "riscv.fmax.d"(%f0, %f1) {"fastmath" = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg -// CHECK-GENERIC-NEXT: %fmin_d_fm = "riscv.fmin.d"(%f0, %f1) {"fastmath" = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg -// CHECK-GENERIC-NEXT: %fmax_d_fm = "riscv.fmax.d"(%f0, %f1) {"fastmath" = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg +// CHECK-GENERIC-NEXT: %fmin_d = "riscv.fmin.d"(%f0, %f1) {fastmath = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg +// CHECK-GENERIC-NEXT: %fmax_d = "riscv.fmax.d"(%f0, %f1) {fastmath = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg +// CHECK-GENERIC-NEXT: %fmin_d_fm = "riscv.fmin.d"(%f0, %f1) {fastmath = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg +// CHECK-GENERIC-NEXT: %fmax_d_fm = "riscv.fmax.d"(%f0, %f1) {fastmath = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg // CHECK-GENERIC-NEXT: %{{.*}} = "riscv.fcvt.d.w"(%{{.*}}) : (!riscv.reg) -> !riscv.freg // CHECK-GENERIC-NEXT: %{{.*}} = "riscv.fcvt.d.wu"(%{{.*}}) : (!riscv.reg) -> !riscv.freg // CHECK-GENERIC-NEXT: "riscv_func.return"() : () -> () -// CHECK-GENERIC-NEXT: }) {"sym_name" = "main", "function_type" = () -> ()} : () -> () +// CHECK-GENERIC-NEXT: }) {sym_name = "main", function_type = () -> ()} : () -> () // CHECK-GENERIC-NEXT: }) : () -> () diff --git a/tests/filecheck/dialects/riscv_cf/canonicalize.mlir b/tests/filecheck/dialects/riscv_cf/canonicalize.mlir index ab603bbd0c..876583572e 100644 --- a/tests/filecheck/dialects/riscv_cf/canonicalize.mlir +++ b/tests/filecheck/dialects/riscv_cf/canonicalize.mlir @@ -33,7 +33,7 @@ riscv_func.func @never() { // CHECK-NEXT: %one = riscv.li 1 : !riscv.reg // CHECK-NEXT: %three = riscv.li 3 : !riscv.reg // CHECK-NEXT: %0 = riscv.mv %one : (!riscv.reg) -> !riscv.reg -// CHECK-NEXT: riscv_cf.branch ^0(%0 : !riscv.reg) attributes {"comment" = "Constant folded riscv_cf.bge"} +// CHECK-NEXT: riscv_cf.branch ^0(%0 : !riscv.reg) attributes {comment = "Constant folded riscv_cf.bge"} // CHECK-NEXT: ^0(%i : !riscv.reg): // CHECK-NEXT: riscv.label "scf_body_0_for" // CHECK-NEXT: "test.op"(%i) : (!riscv.reg) -> () @@ -47,7 +47,7 @@ riscv_func.func @never() { // CHECK-NEXT: riscv_func.func @never() { // CHECK-NEXT: %one = riscv.li 1 : !riscv.reg // CHECK-NEXT: %0 = riscv.mv %one : (!riscv.reg) -> !riscv.reg -// CHECK-NEXT: riscv_cf.j ^0(%0 : !riscv.reg) attributes {"comment" = "Constant folded riscv_cf.bge"} +// CHECK-NEXT: riscv_cf.j ^0(%0 : !riscv.reg) attributes {comment = "Constant folded riscv_cf.bge"} // CHECK-NEXT: ^0(%1 : !riscv.reg): // CHECK-NEXT: riscv.label "scf_body_end_0_for" // CHECK-NEXT: riscv_func.return diff --git a/tests/filecheck/dialects/riscv_cf/ops.mlir b/tests/filecheck/dialects/riscv_cf/ops.mlir index 7afe37508e..6b2ea0663c 100644 --- a/tests/filecheck/dialects/riscv_cf/ops.mlir +++ b/tests/filecheck/dialects/riscv_cf/ops.mlir @@ -33,10 +33,10 @@ // CHECK-NEXT: riscv_cf.beq %0 : !riscv.reg, %1 : !riscv.reg, ^{{.+}}(%2 : !riscv.reg, %3 : !riscv.reg), ^{{.+}}(%4 : !riscv.reg, %5 : !riscv.reg) // CHECK-NEXT: ^{{.+}}(%{{.+}} : !riscv.reg, %{{.+}} : !riscv.reg): // CHECK-NEXT: riscv.label "else0" -// CHECK-NEXT: riscv_cf.bne %0 : !riscv.reg, %1 : !riscv.reg, ^{{.+}}(%2 : !riscv.reg, %3 : !riscv.reg), ^{{.+}}(%4 : !riscv.reg, %5 : !riscv.reg) attributes {"comment" = "comment"} +// CHECK-NEXT: riscv_cf.bne %0 : !riscv.reg, %1 : !riscv.reg, ^{{.+}}(%2 : !riscv.reg, %3 : !riscv.reg), ^{{.+}}(%4 : !riscv.reg, %5 : !riscv.reg) attributes {comment = "comment"} // CHECK-NEXT: ^{{.+}}(%{{.+}} : !riscv.reg, %{{.+}} : !riscv.reg): // CHECK-NEXT: riscv.label "else1" -// CHECK-NEXT: riscv_cf.blt %0 : !riscv.reg, %1 : !riscv.reg, ^{{.+}}(%2 : !riscv.reg, %3 : !riscv.reg), ^{{.+}}(%4 : !riscv.reg, %5 : !riscv.reg) attributes {"hello" = "world"} +// CHECK-NEXT: riscv_cf.blt %0 : !riscv.reg, %1 : !riscv.reg, ^{{.+}}(%2 : !riscv.reg, %3 : !riscv.reg), ^{{.+}}(%4 : !riscv.reg, %5 : !riscv.reg) attributes {hello = "world"} // CHECK-NEXT: ^{{.+}}(%{{.+}} : !riscv.reg, %{{.+}} : !riscv.reg): // CHECK-NEXT: riscv.label "else2" // CHECK-NEXT: riscv_cf.bge %0 : !riscv.reg, %1 : !riscv.reg, ^{{.+}}(%2 : !riscv.reg, %3 : !riscv.reg), ^{{.+}}(%4 : !riscv.reg, %5 : !riscv.reg) @@ -48,10 +48,10 @@ // CHECK-NEXT: riscv_cf.bgeu %0 : !riscv.reg, %1 : !riscv.reg, ^{{.+}}(%2 : !riscv.reg, %3 : !riscv.reg), ^{{.+}}(%4 : !riscv.reg, %5 : !riscv.reg) // CHECK-NEXT: ^{{.+}}(%{{.+}} : !riscv.reg, %{{.+}} : !riscv.reg): // CHECK-NEXT: riscv.label "else5" -// CHECK-NEXT: riscv_cf.branch ^{{.+}}(%2 : !riscv.reg, %3 : !riscv.reg) attributes {"hello" = "world"} +// CHECK-NEXT: riscv_cf.branch ^{{.+}}(%2 : !riscv.reg, %3 : !riscv.reg) attributes {hello = "world"} // CHECK-NEXT: ^{{.+}}(%{{.+}} : !riscv.reg, %{{.+}} : !riscv.reg): // CHECK-NEXT: riscv.label "then" -// CHECK-NEXT: riscv_cf.j ^{{.+}}(%2 : !riscv.reg, %3 : !riscv.reg) attributes {"hello" = "world"} +// CHECK-NEXT: riscv_cf.j ^{{.+}}(%2 : !riscv.reg, %3 : !riscv.reg) attributes {hello = "world"} // CHECK-NEXT: }) : () -> () // CHECK-NEXT: } @@ -90,10 +90,10 @@ // CHECK-NEXT: riscv_cf.beq %0 : !riscv.reg, %1 : !riscv.reg, ^{{.+}}(%2 : !riscv.freg, %3 : !riscv.freg), ^{{.+}}(%4 : !riscv.freg, %5 : !riscv.freg) // CHECK-NEXT: ^{{.+}}(%{{.+}} : !riscv.freg, %{{.+}} : !riscv.freg): // CHECK-NEXT: riscv.label "else0" -// CHECK-NEXT: riscv_cf.bne %0 : !riscv.reg, %1 : !riscv.reg, ^{{.+}}(%2 : !riscv.freg, %3 : !riscv.freg), ^{{.+}}(%4 : !riscv.freg, %5 : !riscv.freg) attributes {"comment" = "comment"} +// CHECK-NEXT: riscv_cf.bne %0 : !riscv.reg, %1 : !riscv.reg, ^{{.+}}(%2 : !riscv.freg, %3 : !riscv.freg), ^{{.+}}(%4 : !riscv.freg, %5 : !riscv.freg) attributes {comment = "comment"} // CHECK-NEXT: ^{{.+}}(%{{.+}} : !riscv.freg, %{{.+}} : !riscv.freg): // CHECK-NEXT: riscv.label "else1" -// CHECK-NEXT: riscv_cf.blt %0 : !riscv.reg, %1 : !riscv.reg, ^{{.+}}(%2 : !riscv.freg, %3 : !riscv.freg), ^{{.+}}(%4 : !riscv.freg, %5 : !riscv.freg) attributes {"hello" = "world"} +// CHECK-NEXT: riscv_cf.blt %0 : !riscv.reg, %1 : !riscv.reg, ^{{.+}}(%2 : !riscv.freg, %3 : !riscv.freg), ^{{.+}}(%4 : !riscv.freg, %5 : !riscv.freg) attributes {hello = "world"} // CHECK-NEXT: ^{{.+}}(%{{.+}} : !riscv.freg, %{{.+}} : !riscv.freg): // CHECK-NEXT: riscv.label "else2" // CHECK-NEXT: riscv_cf.bge %0 : !riscv.reg, %1 : !riscv.reg, ^{{.+}}(%2 : !riscv.freg, %3 : !riscv.freg), ^{{.+}}(%4 : !riscv.freg, %5 : !riscv.freg) @@ -105,9 +105,9 @@ // CHECK-NEXT: riscv_cf.bgeu %0 : !riscv.reg, %1 : !riscv.reg, ^{{.+}}(%2 : !riscv.freg, %3 : !riscv.freg), ^{{.+}}(%4 : !riscv.freg, %5 : !riscv.freg) // CHECK-NEXT: ^{{.+}}(%{{.+}} : !riscv.freg, %{{.+}} : !riscv.freg): // CHECK-NEXT: riscv.label "else5" -// CHECK-NEXT: riscv_cf.branch ^{{.+}}(%2 : !riscv.freg, %3 : !riscv.freg) attributes {"hello" = "world"} +// CHECK-NEXT: riscv_cf.branch ^{{.+}}(%2 : !riscv.freg, %3 : !riscv.freg) attributes {hello = "world"} // CHECK-NEXT: ^{{.+}}(%{{.+}} : !riscv.freg, %{{.+}} : !riscv.freg): // CHECK-NEXT: riscv.label "then" -// CHECK-NEXT: riscv_cf.j ^{{.+}}(%2 : !riscv.freg, %3 : !riscv.freg) attributes {"hello" = "world"} +// CHECK-NEXT: riscv_cf.j ^{{.+}}(%2 : !riscv.freg, %3 : !riscv.freg) attributes {hello = "world"} // CHECK-NEXT: }) : () -> () // CHECK-NEXT: } diff --git a/tests/filecheck/dialects/riscv_debug/riscv_debug_ops.mlir b/tests/filecheck/dialects/riscv_debug/riscv_debug_ops.mlir index 665a14fe61..bd1637f729 100644 --- a/tests/filecheck/dialects/riscv_debug/riscv_debug_ops.mlir +++ b/tests/filecheck/dialects/riscv_debug/riscv_debug_ops.mlir @@ -21,7 +21,7 @@ riscv_debug.printf %0, %1, %2, "{}, {}, {}" : (!riscv.reg, !riscv.reg, ! // CHECK-GENERIC-NEXT: %0 = "riscv.get_register"() : () -> !riscv.reg // CHECK-GENERIC-NEXT: %1 = "riscv.get_register"() : () -> !riscv.reg // CHECK-GENERIC-NEXT: %2 = "riscv.get_float_register"() : () -> !riscv.freg -// CHECK-GENERIC-NEXT: "riscv_debug.printf"(%0, %1, %2) {"format_str" = "{}, {}, {}"} : (!riscv.reg, !riscv.reg, !riscv.freg) -> () +// CHECK-GENERIC-NEXT: "riscv_debug.printf"(%0, %1, %2) {format_str = "{}, {}, {}"} : (!riscv.reg, !riscv.reg, !riscv.freg) -> () // CHECK-GENERIC-NEXT: }) : () -> () // CHECK-ASM: printf "{}, {}, {}", a0, a1, fa0 diff --git a/tests/filecheck/dialects/riscv_func/riscv_func_ops.mlir b/tests/filecheck/dialects/riscv_func/riscv_func_ops.mlir index c9f78409f3..535b013503 100644 --- a/tests/filecheck/dialects/riscv_func/riscv_func_ops.mlir +++ b/tests/filecheck/dialects/riscv_func/riscv_func_ops.mlir @@ -26,7 +26,7 @@ builtin.module { } // CHECK: riscv_func.func @call_void_attributes() { - // CHECK-NEXT: riscv_func.call @call_void_attributes() {"hello" = "world"} : () -> () + // CHECK-NEXT: riscv_func.call @call_void_attributes() {hello = "world"} : () -> () // CHECK-NEXT: riscv_func.return // CHECK-NEXT: } diff --git a/tests/filecheck/dialects/riscv_snitch/ops.mlir b/tests/filecheck/dialects/riscv_snitch/ops.mlir index af3392eb0a..b8f0385bac 100644 --- a/tests/filecheck/dialects/riscv_snitch/ops.mlir +++ b/tests/filecheck/dialects/riscv_snitch/ops.mlir @@ -120,15 +120,15 @@ riscv_func.func @simd() { // CHECK-GENERIC-NEXT: %0 = "riscv.get_register"() : () -> !riscv.reg // CHECK-GENERIC-NEXT: %1 = "riscv.get_register"() : () -> !riscv.reg // CHECK-GENERIC-NEXT: "riscv_snitch.scfgw"(%0, %1) : (!riscv.reg, !riscv.reg) -> () -// CHECK-GENERIC-NEXT: "riscv_snitch.scfgwi"(%0) {"immediate" = 42 : si12} : (!riscv.reg) -> () +// CHECK-GENERIC-NEXT: "riscv_snitch.scfgwi"(%0) {immediate = 42 : si12} : (!riscv.reg) -> () // CHECK-GENERIC-NEXT: "riscv_snitch.frep_outer"(%{{.*}}) ({ // CHECK-GENERIC-NEXT: %{{.*}} = "riscv.add"(%{{.*}}, %{{.*}}) : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-GENERIC-NEXT: "riscv_snitch.frep_yield"() : () -> () -// CHECK-GENERIC-NEXT: }) {"stagger_mask" = #builtin.int<0>, "stagger_count" = #builtin.int<0>} : (!riscv.reg) -> () +// CHECK-GENERIC-NEXT: }) {stagger_mask = #builtin.int<0>, stagger_count = #builtin.int<0>} : (!riscv.reg) -> () // CHECK-GENERIC-NEXT: "riscv_snitch.frep_inner"(%{{.*}}) ({ // CHECK-GENERIC-NEXT: %{{.*}} = "riscv.add"(%{{.*}}, %{{.*}}) : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-GENERIC-NEXT: "riscv_snitch.frep_yield"() : () -> () -// CHECK-GENERIC-NEXT: }) {"stagger_mask" = #builtin.int<0>, "stagger_count" = #builtin.int<0>} : (!riscv.reg) -> () +// CHECK-GENERIC-NEXT: }) {stagger_mask = #builtin.int<0>, stagger_count = #builtin.int<0>} : (!riscv.reg) -> () // CHECK-GENERIC-NEXT: %readable = "riscv_snitch.get_stream"() : () -> !snitch.readable> // CHECK-GENERIC-NEXT: %writable = "riscv_snitch.get_stream"() : () -> !snitch.writable> // CHECK-GENERIC-NEXT: "riscv_snitch.frep_outer"(%0) ({ @@ -136,15 +136,15 @@ riscv_func.func @simd() { // CHECK-GENERIC-NEXT: %val1 = "riscv.fmv.d"(%val0) : (!riscv.freg) -> !riscv.freg // CHECK-GENERIC-NEXT: "riscv_snitch.write"(%val1, %writable) : (!riscv.freg, !snitch.writable>) -> () // CHECK-GENERIC-NEXT: "riscv_snitch.frep_yield"() : () -> () -// CHECK-GENERIC-NEXT: }) {"stagger_mask" = #builtin.int<0>, "stagger_count" = #builtin.int<0>} : (!riscv.reg) -> () +// CHECK-GENERIC-NEXT: }) {stagger_mask = #builtin.int<0>, stagger_count = #builtin.int<0>} : (!riscv.reg) -> () // CHECK-GENERIC-NEXT: %init = "test.op"() : () -> !riscv.freg // CHECK-GENERIC-NEXT: %z = "riscv_snitch.frep_outer"(%0, %init) ({ // CHECK-GENERIC-NEXT: ^0(%acc : !riscv.freg): -// CHECK-GENERIC-NEXT: %res = "riscv.fadd.d"(%acc, %acc) {"fastmath" = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg +// CHECK-GENERIC-NEXT: %res = "riscv.fadd.d"(%acc, %acc) {fastmath = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg // CHECK-GENERIC-NEXT: "riscv_snitch.frep_yield"(%res) : (!riscv.freg) -> () -// CHECK-GENERIC-NEXT: }) {"stagger_mask" = #builtin.int<0>, "stagger_count" = #builtin.int<0>} : (!riscv.reg, !riscv.freg) -> !riscv.freg +// CHECK-GENERIC-NEXT: }) {stagger_mask = #builtin.int<0>, stagger_count = #builtin.int<0>} : (!riscv.reg, !riscv.freg) -> !riscv.freg // CHECK-GENERIC-NEXT: "riscv_func.return"() : () -> () -// CHECK-GENERIC-NEXT: }) {"sym_name" = "xfrep", "function_type" = () -> ()} : () -> () +// CHECK-GENERIC-NEXT: }) {sym_name = "xfrep", function_type = () -> ()} : () -> () // CHECK-GENERIC-NEXT: "riscv_func.func"() ({ // CHECK-GENERIC-NEXT: %reg = "riscv.get_register"() : () -> !riscv.reg // CHECK-GENERIC-NEXT: "riscv_snitch.dmsrc"(%reg, %reg) : (!riscv.reg, !riscv.reg) -> () @@ -153,19 +153,19 @@ riscv_func.func @simd() { // CHECK-GENERIC-NEXT: "riscv_snitch.dmrep"(%reg) : (!riscv.reg) -> () // CHECK-GENERIC-NEXT: %{{.*}} = "riscv_snitch.dmcpy"(%reg, %reg) : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-GENERIC-NEXT: %{{.*}} = "riscv_snitch.dmstat"(%reg) : (!riscv.reg) -> !riscv.reg -// CHECK-GENERIC-NEXT: %{{.*}} = "riscv_snitch.dmcpyi"(%reg) <{"config" = 0 : ui5}> : (!riscv.reg) -> !riscv.reg -// CHECK-GENERIC-NEXT: %{{.*}} = "riscv_snitch.dmstati"() <{"status" = 0 : ui5}> : () -> !riscv.reg +// CHECK-GENERIC-NEXT: %{{.*}} = "riscv_snitch.dmcpyi"(%reg) <{config = 0 : ui5}> : (!riscv.reg) -> !riscv.reg +// CHECK-GENERIC-NEXT: %{{.*}} = "riscv_snitch.dmstati"() <{status = 0 : ui5}> : () -> !riscv.reg // CHECK-GENERIC-NEXT: "riscv_func.return"() : () -> () -// CHECK-GENERIC-NEXT: }) {"sym_name" = "xdma", "function_type" = () -> ()} : () -> () +// CHECK-GENERIC-NEXT: }) {sym_name = "xdma", function_type = () -> ()} : () -> () // CHECK-GENERIC-NEXT: "riscv_func.func"() ({ // CHECK-GENERIC-NEXT: %v = "riscv.get_float_register"() : () -> !riscv.freg -// CHECK-GENERIC-NEXT: %0 = "riscv_snitch.vfmul.s"(%v, %v) {"fastmath" = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg -// CHECK-GENERIC-NEXT: %1 = "riscv_snitch.vfadd.s"(%v, %v) {"fastmath" = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg +// CHECK-GENERIC-NEXT: %0 = "riscv_snitch.vfmul.s"(%v, %v) {fastmath = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg +// CHECK-GENERIC-NEXT: %1 = "riscv_snitch.vfadd.s"(%v, %v) {fastmath = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg // CHECK-GENERIC-NEXT: %2 = "riscv_snitch.vfcpka.s.s"(%v, %v) : (!riscv.freg, !riscv.freg) -> !riscv.freg -// CHECK-GENERIC-NEXT: %3 = "riscv_snitch.vfmac.s"(%v, %v, %v) {"fastmath" = #riscv.fastmath} : (!riscv.freg, !riscv.freg, !riscv.freg) -> !riscv.freg +// CHECK-GENERIC-NEXT: %3 = "riscv_snitch.vfmac.s"(%v, %v, %v) {fastmath = #riscv.fastmath} : (!riscv.freg, !riscv.freg, !riscv.freg) -> !riscv.freg // CHECK-GENERIC-NEXT: %4 = "riscv_snitch.vfsum.s"(%v, %v) : (!riscv.freg, !riscv.freg) -> !riscv.freg -// CHECK-GENERIC-NEXT: %5 = "riscv_snitch.vfadd.h"(%v, %v) {"fastmath" = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg -// CHECK-GENERIC-NEXT: %6 = "riscv_snitch.vfmax.s"(%v, %v) {"fastmath" = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg +// CHECK-GENERIC-NEXT: %5 = "riscv_snitch.vfadd.h"(%v, %v) {fastmath = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg +// CHECK-GENERIC-NEXT: %6 = "riscv_snitch.vfmax.s"(%v, %v) {fastmath = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg // CHECK-GENERIC-NEXT: "riscv_func.return"() : () -> () -// CHECK-GENERIC-NEXT: }) {"sym_name" = "simd", "function_type" = () -> ()} : () -> () +// CHECK-GENERIC-NEXT: }) {sym_name = "simd", function_type = () -> ()} : () -> () // CHECK-GENERIC-NEXT: }) : () -> () diff --git a/tests/filecheck/dialects/scf/reduce.mlir b/tests/filecheck/dialects/scf/reduce.mlir index c71bc71c3c..d3c7e81f8b 100644 --- a/tests/filecheck/dialects/scf/reduce.mlir +++ b/tests/filecheck/dialects/scf/reduce.mlir @@ -22,7 +22,7 @@ builtin.module { // CHECK-NEXT: %2 = arith.constant 3 : index // CHECK-NEXT: %3 = arith.constant 1.020000e+01 : f32 // CHECK-NEXT: %4 = arith.constant 1.810000e+01 : f32 -// CHECK-NEXT: %5 = "scf.parallel"(%0, %1, %2, %3) <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: %5 = "scf.parallel"(%0, %1, %2, %3) <{operandSegmentSizes = array}> ({ // CHECK-NEXT: ^0(%6 : index): // CHECK-NEXT: scf.reduce(%4 : f32) { // CHECK-NEXT: ^1(%7 : f32, %8 : f32): diff --git a/tests/filecheck/dialects/scf/scf_ops.mlir b/tests/filecheck/dialects/scf/scf_ops.mlir index f9b06b3e53..d0dcec15de 100644 --- a/tests/filecheck/dialects/scf/scf_ops.mlir +++ b/tests/filecheck/dialects/scf/scf_ops.mlir @@ -123,7 +123,7 @@ builtin.module { // CHECK-NEXT: %{{.*}}, %{{.*}} = scf.while (%{{.*}} = %{{.*}}, %{{.*}} = %{{.*}}) : (i32, f32) -> (i32, f32) { // CHECK-NEXT: %{{.*}} = arith.constant 0 : i32 // CHECK-NEXT: %{{.*}} = arith.cmpi eq, %{{.*}}, %{{.*}} : i32 - // CHECK-NEXT: scf.condition(%{{.*}}) {"hello" = "world"} %{{.*}}, %{{.*}} : i32, f32 + // CHECK-NEXT: scf.condition(%{{.*}}) {hello = "world"} %{{.*}}, %{{.*}} : i32, f32 // CHECK-NEXT: } do { // CHECK-NEXT: ^{{.*}}(%{{.*}} : i32, %{{.*}} : f32): // CHECK-NEXT: %{{.*}} = arith.constant 1.000000e+00 : f32 diff --git a/tests/filecheck/dialects/seq/seq_ops.mlir b/tests/filecheck/dialects/seq/seq_ops.mlir index 513bdd2f1d..c8e7c19d34 100644 --- a/tests/filecheck/dialects/seq/seq_ops.mlir +++ b/tests/filecheck/dialects/seq/seq_ops.mlir @@ -19,12 +19,12 @@ builtin.module { // CHECK: %compreg_all = seq.compreg %data, %clk reset %bool, %data powerOn %on : i14 %compreg_sym = seq.compreg sym @foo %data, %clk : i14 // CHECK: %compreg_sym = seq.compreg sym @foo %data, %clk : i14 - // CHECK-GENERIC: %compreg_sym = "seq.compreg"(%data, %clk) {"inner_sym" = #hw, "operandSegmentSizes" = array} : (i14, !seq.clock) -> i14 + // CHECK-GENERIC: %compreg_sym = "seq.compreg"(%data, %clk) {inner_sym = #hw, operandSegmentSizes = array} : (i14, !seq.clock) -> i14 %const_low = seq.const_clock low // CHECK: %const_low = seq.const_clock low %const_high = seq.const_clock high // CHECK: %const_high = seq.const_clock high %const_low_foo = seq.const_clock low {"foo"} - // CHECK: %const_low_foo = seq.const_clock low {"foo"} + // CHECK: %const_low_foo = seq.const_clock low {foo} } diff --git a/tests/filecheck/dialects/snitch/snitch_ops.mlir b/tests/filecheck/dialects/snitch/snitch_ops.mlir index f8561cce8a..ba28685e27 100644 --- a/tests/filecheck/dialects/snitch/snitch_ops.mlir +++ b/tests/filecheck/dialects/snitch/snitch_ops.mlir @@ -6,15 +6,15 @@ %rep = "test.op"() : () -> !riscv.reg // Usual SSR setup sequence: "snitch.ssr_set_dimension_bound"(%bound) {"dm" = #builtin.int<0>, "dimension" = #builtin.int<0>} : (!riscv.reg) -> () - // CHECK: "snitch.ssr_set_dimension_bound"(%bound) {"dm" = #builtin.int<0>, "dimension" = #builtin.int<0>} : (!riscv.reg) -> () + // CHECK: "snitch.ssr_set_dimension_bound"(%bound) {dm = #builtin.int<0>, dimension = #builtin.int<0>} : (!riscv.reg) -> () "snitch.ssr_set_dimension_stride"(%stride) {"dm" = #builtin.int<0>, "dimension" = #builtin.int<0>} : (!riscv.reg) -> () - // CHECK: "snitch.ssr_set_dimension_stride"(%stride) {"dm" = #builtin.int<0>, "dimension" = #builtin.int<0>} : (!riscv.reg) -> () + // CHECK: "snitch.ssr_set_dimension_stride"(%stride) {dm = #builtin.int<0>, dimension = #builtin.int<0>} : (!riscv.reg) -> () "snitch.ssr_set_dimension_source"(%addr) {"dm" = #builtin.int<0>, "dimension" = #builtin.int<0>} : (!riscv.reg) -> () - // CHECK: "snitch.ssr_set_dimension_source"(%addr) {"dm" = #builtin.int<0>, "dimension" = #builtin.int<0>} : (!riscv.reg) -> () + // CHECK: "snitch.ssr_set_dimension_source"(%addr) {dm = #builtin.int<0>, dimension = #builtin.int<0>} : (!riscv.reg) -> () "snitch.ssr_set_dimension_destination"(%addr) {"dm" = #builtin.int<0>, "dimension" = #builtin.int<0>} : (!riscv.reg) -> () - // CHECK: "snitch.ssr_set_dimension_destination"(%addr) {"dm" = #builtin.int<0>, "dimension" = #builtin.int<0>} : (!riscv.reg) -> () + // CHECK: "snitch.ssr_set_dimension_destination"(%addr) {dm = #builtin.int<0>, dimension = #builtin.int<0>} : (!riscv.reg) -> () "snitch.ssr_set_stream_repetition"(%rep) {"dm" = #builtin.int<0>} : (!riscv.reg) -> () - // CHECK: "snitch.ssr_set_stream_repetition"(%rep) {"dm" = #builtin.int<0>} : (!riscv.reg) -> () + // CHECK: "snitch.ssr_set_stream_repetition"(%rep) {dm = #builtin.int<0>} : (!riscv.reg) -> () "snitch.ssr_enable"() : () -> () // CHECK-NEXT: "snitch.ssr_enable"() : () -> () "snitch.ssr_disable"() : () -> () diff --git a/tests/filecheck/dialects/snitch/snitch_to_riscv_lowering.mlir b/tests/filecheck/dialects/snitch/snitch_to_riscv_lowering.mlir index ec534ca604..fc449e99e3 100644 --- a/tests/filecheck/dialects/snitch/snitch_to_riscv_lowering.mlir +++ b/tests/filecheck/dialects/snitch/snitch_to_riscv_lowering.mlir @@ -7,32 +7,32 @@ builtin.module { // SSR setup sequence for dimension 0 "snitch.ssr_set_dimension_bound"(%bound) {"dm" = #builtin.int<0>, "dimension" = #builtin.int<0>} : (!riscv.reg) -> () // CHECK: %{{.*}} = riscv.li 64 : !riscv.reg - // CHECK-NEXT: riscv_snitch.scfgw %bound, %{{.*}} {"comment" = "dm 0 dim 0 bound"} : (!riscv.reg, !riscv.reg) -> () + // CHECK-NEXT: riscv_snitch.scfgw %bound, %{{.*}} {comment = "dm 0 dim 0 bound"} : (!riscv.reg, !riscv.reg) -> () "snitch.ssr_set_dimension_stride"(%stride) {"dm" = #builtin.int<0>, "dimension" = #builtin.int<0>} : (!riscv.reg) -> () // CHECK: %{{.*}} = riscv.li 192 : !riscv.reg - // CHECK-NEXT: riscv_snitch.scfgw %stride, %{{.*}} {"comment" = "dm 0 dim 0 stride"} : (!riscv.reg, !riscv.reg) -> () + // CHECK-NEXT: riscv_snitch.scfgw %stride, %{{.*}} {comment = "dm 0 dim 0 stride"} : (!riscv.reg, !riscv.reg) -> () "snitch.ssr_set_dimension_source"(%addr) {"dm" = #builtin.int<0>, "dimension" = #builtin.int<0>} : (!riscv.reg) -> () // CHECK: %{{.*}} = riscv.li 768 : !riscv.reg - // riscv_snitch.scfgw %addr, %{{.*}} {"comment" = "dm 0 dim 0 source"} : (!riscv.reg, !riscv.reg) -> () + // riscv_snitch.scfgw %addr, %{{.*}} {comment = "dm 0 dim 0 source"} : (!riscv.reg, !riscv.reg) -> () "snitch.ssr_set_dimension_destination"(%addr) {"dm" = #builtin.int<0>, "dimension" = #builtin.int<0>} : (!riscv.reg) -> () // CHECK: %{{.*}} = riscv.li 896 : !riscv.reg - // CHECK-NEXT: riscv_snitch.scfgw %addr, %{{.*}} {"comment" = "dm 0 dim 0 destination"} : (!riscv.reg, !riscv.reg) -> () + // CHECK-NEXT: riscv_snitch.scfgw %addr, %{{.*}} {comment = "dm 0 dim 0 destination"} : (!riscv.reg, !riscv.reg) -> () // SSR setup sequence for dimension 3 "snitch.ssr_set_dimension_bound"(%bound) {"dm" = #builtin.int<0>, "dimension" = #builtin.int<3>} : (!riscv.reg) -> () // CHECK: %{{.*}} = riscv.li 160 : !riscv.reg - // CHECK-NEXT: riscv_snitch.scfgw %bound, %{{.*}} {"comment" = "dm 0 dim 3 bound"} : (!riscv.reg, !riscv.reg) -> () + // CHECK-NEXT: riscv_snitch.scfgw %bound, %{{.*}} {comment = "dm 0 dim 3 bound"} : (!riscv.reg, !riscv.reg) -> () "snitch.ssr_set_dimension_stride"(%stride) {"dm" = #builtin.int<0>, "dimension" = #builtin.int<3>} : (!riscv.reg) -> () // CHECK: %{{.*}} = riscv.li 288 : !riscv.reg - // CHECK-NEXT: riscv_snitch.scfgw %stride, %{{.*}} {"comment" = "dm 0 dim 3 stride"} : (!riscv.reg, !riscv.reg) -> () + // CHECK-NEXT: riscv_snitch.scfgw %stride, %{{.*}} {comment = "dm 0 dim 3 stride"} : (!riscv.reg, !riscv.reg) -> () "snitch.ssr_set_dimension_source"(%addr) {"dm" = #builtin.int<0>, "dimension" = #builtin.int<3>} : (!riscv.reg) -> () // CHECK: %{{.*}} = riscv.li 864 : !riscv.reg - // riscv_snitch.scfgw %addr, %{{.*}} {"comment" = "dm 0 dim 3 source"} : (!riscv.reg, !riscv.reg) -> () + // riscv_snitch.scfgw %addr, %{{.*}} {comment = "dm 0 dim 3 source"} : (!riscv.reg, !riscv.reg) -> () "snitch.ssr_set_dimension_destination"(%addr) {"dm" = #builtin.int<0>, "dimension" = #builtin.int<3>} : (!riscv.reg) -> () // CHECK: %{{.*}} = riscv.li 992 : !riscv.reg - // CHECK-NEXT: riscv_snitch.scfgw %addr, %{{.*}} {"comment" = "dm 0 dim 3 destination"} : (!riscv.reg, !riscv.reg) -> () + // CHECK-NEXT: riscv_snitch.scfgw %addr, %{{.*}} {comment = "dm 0 dim 3 destination"} : (!riscv.reg, !riscv.reg) -> () "snitch.ssr_set_stream_repetition"(%rep) {"dm" = #builtin.int<0>}: (!riscv.reg) -> () // CHECK: %{{.*}} = riscv.li 32 : !riscv.reg - // CHECK-NEXT: riscv_snitch.scfgw %rep, %{{.*}} {"comment" = "dm 0 repeat"} : (!riscv.reg, !riscv.reg) -> () + // CHECK-NEXT: riscv_snitch.scfgw %rep, %{{.*}} {comment = "dm 0 repeat"} : (!riscv.reg, !riscv.reg) -> () // On/Off switching sequence "snitch.ssr_enable"() : () -> () // CHECK: riscv.csrrsi 1984, 1 diff --git a/tests/filecheck/dialects/snitch_runtime/snitch_runtime_ops.mlir b/tests/filecheck/dialects/snitch_runtime/snitch_runtime_ops.mlir index 009dfde13a..e3500df2d8 100644 --- a/tests/filecheck/dialects/snitch_runtime/snitch_runtime_ops.mlir +++ b/tests/filecheck/dialects/snitch_runtime/snitch_runtime_ops.mlir @@ -80,22 +80,22 @@ builtin.module { %b0 = arith.constant 100 : index %i0 = arith.constant 101 : index "snrt.ssr_loop_1d"(%dm, %b0, %i0) {"operandSegmentSizes" = array} : (i32, index, index) -> () - // CHECK: "snrt.ssr_loop_1d"(%dm, %b0, %i0) {"operandSegmentSizes" = array} : (i32, index, index) -> () + // CHECK: "snrt.ssr_loop_1d"(%dm, %b0, %i0) {operandSegmentSizes = array} : (i32, index, index) -> () %b1 = arith.constant 102 : index %i1 = arith.constant 103 : index "snrt.ssr_loop_2d"(%dm, %b0, %b1, %i0, %i1) {"operandSegmentSizes" = array} : (i32, index, index, index, index) -> () - // CHECK: "snrt.ssr_loop_2d"(%dm, %b0, %b1, %i0, %i1) {"operandSegmentSizes" = array} : (i32, index, index, index, index) -> () + // CHECK: "snrt.ssr_loop_2d"(%dm, %b0, %b1, %i0, %i1) {operandSegmentSizes = array} : (i32, index, index, index, index) -> () %b2 = arith.constant 104 : index %i2 = arith.constant 105 : index "snrt.ssr_loop_3d"(%dm, %b0, %b1, %b2, %i0, %i1, %i2) {"operandSegmentSizes" = array} : (i32, index, index, index, index, index, index) -> () - // CHECK: "snrt.ssr_loop_3d"(%dm, %b0, %b1, %b2, %i0, %i1, %i2) {"operandSegmentSizes" = array} : (i32, index, index, index, index, index, index) -> () + // CHECK: "snrt.ssr_loop_3d"(%dm, %b0, %b1, %b2, %i0, %i1, %i2) {operandSegmentSizes = array} : (i32, index, index, index, index, index, index) -> () %b3 = arith.constant 106 : index %i3 = arith.constant 107 : index "snrt.ssr_loop_4d"(%dm, %b0, %b1, %b2, %b3, %i0, %i1, %i2, %i3) {"operandSegmentSizes" = array} : (i32, index, index, index, index, index, index, index, index) -> () - //CHECK: "snrt.ssr_loop_4d"(%dm, %b0, %b1, %b2, %b3, %i0, %i1, %i2, %i3) {"operandSegmentSizes" = array} : (i32, index, index, index, index, index, index, index, index) -> () + //CHECK: "snrt.ssr_loop_4d"(%dm, %b0, %b1, %b2, %b3, %i0, %i1, %i2, %i3) {operandSegmentSizes = array} : (i32, index, index, index, index, index, index, index, index) -> () %count = arith.constant 20 : i32 "snrt.ssr_repeat"(%count) <{dm = 3 : i32}> : (i32) -> () - // CHECK: "snrt.ssr_repeat"(%count) <{"dm" = 3 : i32}> : (i32) -> () + // CHECK: "snrt.ssr_repeat"(%count) <{dm = 3 : i32}> : (i32) -> () "snrt.ssr_enable"() : () -> () // CHECK: "snrt.ssr_enable"() : () -> () "snrt.ssr_disable"() : () -> () @@ -104,9 +104,9 @@ builtin.module { %dim = arith.constant 1 : i32 %ptr = arith.constant 0 : i32 "snrt.ssr_read"(%ptr) <{dm = 0 : i32, dim = 1 : i32}> : (i32) -> () - // CHECK: "snrt.ssr_read"(%ptr) <{"dm" = 0 : i32, "dim" = 1 : i32}> : (i32) -> () + // CHECK: "snrt.ssr_read"(%ptr) <{dm = 0 : i32, dim = 1 : i32}> : (i32) -> () "snrt.ssr_write"(%ptr) <{dm = 0 : i32, dim = 1 : i32}> : (i32) -> () - // CHECK: "snrt.ssr_write"(%ptr) <{"dm" = 0 : i32, "dim" = 1 : i32}> : (i32) -> () + // CHECK: "snrt.ssr_write"(%ptr) <{dm = 0 : i32, dim = 1 : i32}> : (i32) -> () "snrt.fpu_fence"() : () -> () // CHECK: "snrt.fpu_fence"() : () -> () diff --git a/tests/filecheck/dialects/snitch_stream/convert_snitch_stream_to_snitch.mlir b/tests/filecheck/dialects/snitch_stream/convert_snitch_stream_to_snitch.mlir index cfcd304d50..88ae335bab 100644 --- a/tests/filecheck/dialects/snitch_stream/convert_snitch_stream_to_snitch.mlir +++ b/tests/filecheck/dialects/snitch_stream/convert_snitch_stream_to_snitch.mlir @@ -18,63 +18,63 @@ }) : (!riscv.reg, !riscv.reg, !riscv.reg) -> () // CHECK-NEXT: %{{.*}} = riscv.li 2 : !riscv.reg // CHECK-NEXT: %{{.*}} = riscv.addi %{{.*}}, -1 : (!riscv.reg) -> !riscv.reg -// CHECK-NEXT: "snitch.ssr_set_dimension_bound"(%{{.*}}) {"dm" = #builtin.int<0>, "dimension" = #builtin.int<0>} : (!riscv.reg) -> () +// CHECK-NEXT: "snitch.ssr_set_dimension_bound"(%{{.*}}) {dm = #builtin.int<0>, dimension = #builtin.int<0>} : (!riscv.reg) -> () // CHECK-NEXT: %{{.*}} = riscv.li 8 : !riscv.reg -// CHECK-NEXT: "snitch.ssr_set_dimension_stride"(%{{.*}}) {"dm" = #builtin.int<0>, "dimension" = #builtin.int<0>} : (!riscv.reg) -> () +// CHECK-NEXT: "snitch.ssr_set_dimension_stride"(%{{.*}}) {dm = #builtin.int<0>, dimension = #builtin.int<0>} : (!riscv.reg) -> () // CHECK-NEXT: %{{.*}} = riscv.li 0 : !riscv.reg // CHECK-NEXT: %{{.*}} = riscv.li 0 : !riscv.reg -// CHECK-NEXT: "snitch.ssr_set_stream_repetition"(%{{.*}}) {"dm" = #builtin.int<0>} : (!riscv.reg) -> () +// CHECK-NEXT: "snitch.ssr_set_stream_repetition"(%{{.*}}) {dm = #builtin.int<0>} : (!riscv.reg) -> () // CHECK-NEXT: %{{.*}} = riscv.li 2 : !riscv.reg // CHECK-NEXT: %{{.*}} = riscv.li 3 : !riscv.reg // CHECK-NEXT: %{{.*}} = riscv.addi %{{.*}}, -1 : (!riscv.reg) -> !riscv.reg -// CHECK-NEXT: "snitch.ssr_set_dimension_bound"(%{{.*}}) {"dm" = #builtin.int<1>, "dimension" = #builtin.int<0>} : (!riscv.reg) -> () +// CHECK-NEXT: "snitch.ssr_set_dimension_bound"(%{{.*}}) {dm = #builtin.int<1>, dimension = #builtin.int<0>} : (!riscv.reg) -> () // CHECK-NEXT: %{{.*}} = riscv.addi %{{.*}}, -1 : (!riscv.reg) -> !riscv.reg -// CHECK-NEXT: "snitch.ssr_set_dimension_bound"(%{{.*}}) {"dm" = #builtin.int<1>, "dimension" = #builtin.int<1>} : (!riscv.reg) -> () +// CHECK-NEXT: "snitch.ssr_set_dimension_bound"(%{{.*}}) {dm = #builtin.int<1>, dimension = #builtin.int<1>} : (!riscv.reg) -> () // CHECK-NEXT: %{{.*}} = riscv.li 24 : !riscv.reg // CHECK-NEXT: %{{.*}} = riscv.li 8 : !riscv.reg -// CHECK-NEXT: "snitch.ssr_set_dimension_stride"(%{{.*}}) {"dm" = #builtin.int<1>, "dimension" = #builtin.int<0>} : (!riscv.reg) -> () +// CHECK-NEXT: "snitch.ssr_set_dimension_stride"(%{{.*}}) {dm = #builtin.int<1>, dimension = #builtin.int<0>} : (!riscv.reg) -> () // CHECK-NEXT: %{{.*}} = riscv.li 0 : !riscv.reg // CHECK-NEXT: %{{.*}} = riscv.mul %{{.*}}, %{{.*}} : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-NEXT: %{{.*}} = riscv.add %{{.*}}, %{{.*}} : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-NEXT: %{{.*}} = riscv.sub %{{.*}}, %{{.*}} : (!riscv.reg, !riscv.reg) -> !riscv.reg -// CHECK-NEXT: "snitch.ssr_set_dimension_stride"(%{{.*}}) {"dm" = #builtin.int<1>, "dimension" = #builtin.int<1>} : (!riscv.reg) -> () +// CHECK-NEXT: "snitch.ssr_set_dimension_stride"(%{{.*}}) {dm = #builtin.int<1>, dimension = #builtin.int<1>} : (!riscv.reg) -> () // CHECK-NEXT: %{{.*}} = riscv.li 0 : !riscv.reg -// CHECK-NEXT: "snitch.ssr_set_stream_repetition"(%{{.*}}) {"dm" = #builtin.int<1>} : (!riscv.reg) -> () +// CHECK-NEXT: "snitch.ssr_set_stream_repetition"(%{{.*}}) {dm = #builtin.int<1>} : (!riscv.reg) -> () // CHECK-NEXT: %{{.*}} = riscv.li 2 : !riscv.reg // CHECK-NEXT: %{{.*}} = riscv.li 3 : !riscv.reg // CHECK-NEXT: %{{.*}} = riscv.li 4 : !riscv.reg // CHECK-NEXT: %{{.*}} = riscv.li 5 : !riscv.reg // CHECK-NEXT: %{{.*}} = riscv.addi %{{.*}}, -1 : (!riscv.reg) -> !riscv.reg -// CHECK-NEXT: "snitch.ssr_set_dimension_bound"(%{{.*}}) {"dm" = #builtin.int<2>, "dimension" = #builtin.int<0>} : (!riscv.reg) -> () +// CHECK-NEXT: "snitch.ssr_set_dimension_bound"(%{{.*}}) {dm = #builtin.int<2>, dimension = #builtin.int<0>} : (!riscv.reg) -> () // CHECK-NEXT: %{{.*}} = riscv.addi %{{.*}}, -1 : (!riscv.reg) -> !riscv.reg -// CHECK-NEXT: "snitch.ssr_set_dimension_bound"(%{{.*}}) {"dm" = #builtin.int<2>, "dimension" = #builtin.int<1>} : (!riscv.reg) -> () +// CHECK-NEXT: "snitch.ssr_set_dimension_bound"(%{{.*}}) {dm = #builtin.int<2>, dimension = #builtin.int<1>} : (!riscv.reg) -> () // CHECK-NEXT: %{{.*}} = riscv.addi %{{.*}}, -1 : (!riscv.reg) -> !riscv.reg -// CHECK-NEXT: "snitch.ssr_set_dimension_bound"(%{{.*}}) {"dm" = #builtin.int<2>, "dimension" = #builtin.int<2>} : (!riscv.reg) -> () +// CHECK-NEXT: "snitch.ssr_set_dimension_bound"(%{{.*}}) {dm = #builtin.int<2>, dimension = #builtin.int<2>} : (!riscv.reg) -> () // CHECK-NEXT: %{{.*}} = riscv.addi %{{.*}}, -1 : (!riscv.reg) -> !riscv.reg -// CHECK-NEXT: "snitch.ssr_set_dimension_bound"(%{{.*}}) {"dm" = #builtin.int<2>, "dimension" = #builtin.int<3>} : (!riscv.reg) -> () +// CHECK-NEXT: "snitch.ssr_set_dimension_bound"(%{{.*}}) {dm = #builtin.int<2>, dimension = #builtin.int<3>} : (!riscv.reg) -> () // CHECK-NEXT: %{{.*}} = riscv.li 480 : !riscv.reg // CHECK-NEXT: %{{.*}} = riscv.li 160 : !riscv.reg // CHECK-NEXT: %{{.*}} = riscv.li 40 : !riscv.reg // CHECK-NEXT: %{{.*}} = riscv.li 8 : !riscv.reg -// CHECK-NEXT: "snitch.ssr_set_dimension_stride"(%{{.*}}) {"dm" = #builtin.int<2>, "dimension" = #builtin.int<0>} : (!riscv.reg) -> () +// CHECK-NEXT: "snitch.ssr_set_dimension_stride"(%{{.*}}) {dm = #builtin.int<2>, dimension = #builtin.int<0>} : (!riscv.reg) -> () // CHECK-NEXT: %{{.*}} = riscv.li 0 : !riscv.reg // CHECK-NEXT: %{{.*}} = riscv.mul %{{.*}}, %{{.*}} : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-NEXT: %{{.*}} = riscv.add %{{.*}}, %{{.*}} : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-NEXT: %{{.*}} = riscv.sub %{{.*}}, %{{.*}} : (!riscv.reg, !riscv.reg) -> !riscv.reg -// CHECK-NEXT: "snitch.ssr_set_dimension_stride"(%{{.*}}) {"dm" = #builtin.int<2>, "dimension" = #builtin.int<1>} : (!riscv.reg) -> () +// CHECK-NEXT: "snitch.ssr_set_dimension_stride"(%{{.*}}) {dm = #builtin.int<2>, dimension = #builtin.int<1>} : (!riscv.reg) -> () // CHECK-NEXT: %{{.*}} = riscv.mul %{{.*}}, %{{.*}} : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-NEXT: %{{.*}} = riscv.add %{{.*}}, %{{.*}} : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-NEXT: %{{.*}} = riscv.sub %{{.*}}, %{{.*}} : (!riscv.reg, !riscv.reg) -> !riscv.reg -// CHECK-NEXT: "snitch.ssr_set_dimension_stride"(%{{.*}}) {"dm" = #builtin.int<2>, "dimension" = #builtin.int<2>} : (!riscv.reg) -> () +// CHECK-NEXT: "snitch.ssr_set_dimension_stride"(%{{.*}}) {dm = #builtin.int<2>, dimension = #builtin.int<2>} : (!riscv.reg) -> () // CHECK-NEXT: %{{.*}} = riscv.mul %{{.*}}, %{{.*}} : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-NEXT: %{{.*}} = riscv.add %{{.*}}, %{{.*}} : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-NEXT: %{{.*}} = riscv.sub %{{.*}}, %{{.*}} : (!riscv.reg, !riscv.reg) -> !riscv.reg -// CHECK-NEXT: "snitch.ssr_set_dimension_stride"(%{{.*}}) {"dm" = #builtin.int<2>, "dimension" = #builtin.int<3>} : (!riscv.reg) -> () +// CHECK-NEXT: "snitch.ssr_set_dimension_stride"(%{{.*}}) {dm = #builtin.int<2>, dimension = #builtin.int<3>} : (!riscv.reg) -> () // CHECK-NEXT: %{{.*}} = riscv.li 0 : !riscv.reg -// CHECK-NEXT: "snitch.ssr_set_stream_repetition"(%{{.*}}) {"dm" = #builtin.int<2>} : (!riscv.reg) -> () -// CHECK-NEXT: "snitch.ssr_set_dimension_source"(%A) {"dm" = #builtin.int<0>, "dimension" = #builtin.int<0>} : (!riscv.reg) -> () -// CHECK-NEXT: "snitch.ssr_set_dimension_source"(%B) {"dm" = #builtin.int<1>, "dimension" = #builtin.int<1>} : (!riscv.reg) -> () -// CHECK-NEXT: "snitch.ssr_set_dimension_destination"(%C) {"dm" = #builtin.int<2>, "dimension" = #builtin.int<3>} : (!riscv.reg) -> () +// CHECK-NEXT: "snitch.ssr_set_stream_repetition"(%{{.*}}) {dm = #builtin.int<2>} : (!riscv.reg) -> () +// CHECK-NEXT: "snitch.ssr_set_dimension_source"(%A) {dm = #builtin.int<0>, dimension = #builtin.int<0>} : (!riscv.reg) -> () +// CHECK-NEXT: "snitch.ssr_set_dimension_source"(%B) {dm = #builtin.int<1>, dimension = #builtin.int<1>} : (!riscv.reg) -> () +// CHECK-NEXT: "snitch.ssr_set_dimension_destination"(%C) {dm = #builtin.int<2>, dimension = #builtin.int<3>} : (!riscv.reg) -> () // CHECK-NEXT: %a_stream, %b_stream, %c_stream = "snitch.ssr_enable"() : () -> (!snitch.readable>, !snitch.readable>, !snitch.writable>) // CHECK-NEXT: "test.op"() : () -> () // CHECK-NEXT: "snitch.ssr_disable"() : () -> () @@ -91,14 +91,14 @@ // CHECK-NEXT: %{{.*}} = riscv.li 2 : !riscv.reg // CHECK-NEXT: %{{.*}} = riscv.addi %{{.*}}, -1 : (!riscv.reg) -> !riscv.reg -// CHECK-NEXT: "snitch.ssr_set_dimension_bound"(%{{.*}}) {"dm" = #builtin.int<31>, "dimension" = #builtin.int<0>} : (!riscv.reg) -> () +// CHECK-NEXT: "snitch.ssr_set_dimension_bound"(%{{.*}}) {dm = #builtin.int<31>, dimension = #builtin.int<0>} : (!riscv.reg) -> () // CHECK-NEXT: %{{.*}} = riscv.li 8 : !riscv.reg -// CHECK-NEXT: "snitch.ssr_set_dimension_stride"(%{{.*}}) {"dm" = #builtin.int<31>, "dimension" = #builtin.int<0>} : (!riscv.reg) -> () +// CHECK-NEXT: "snitch.ssr_set_dimension_stride"(%{{.*}}) {dm = #builtin.int<31>, dimension = #builtin.int<0>} : (!riscv.reg) -> () // CHECK-NEXT: %{{.*}} = riscv.li 0 : !riscv.reg // CHECK-NEXT: %{{.*}} = riscv.li 0 : !riscv.reg -// CHECK-NEXT: "snitch.ssr_set_stream_repetition"(%{{.*}}) {"dm" = #builtin.int<31>} : (!riscv.reg) -> () -// CHECK-NEXT: "snitch.ssr_set_dimension_source"(%A) {"dm" = #builtin.int<0>, "dimension" = #builtin.int<0>} : (!riscv.reg) -> () -// CHECK-NEXT: "snitch.ssr_set_dimension_destination"(%B) {"dm" = #builtin.int<1>, "dimension" = #builtin.int<0>} : (!riscv.reg) -> () +// CHECK-NEXT: "snitch.ssr_set_stream_repetition"(%{{.*}}) {dm = #builtin.int<31>} : (!riscv.reg) -> () +// CHECK-NEXT: "snitch.ssr_set_dimension_source"(%A) {dm = #builtin.int<0>, dimension = #builtin.int<0>} : (!riscv.reg) -> () +// CHECK-NEXT: "snitch.ssr_set_dimension_destination"(%B) {dm = #builtin.int<1>, dimension = #builtin.int<0>} : (!riscv.reg) -> () // CHECK-NEXT: %{{.*}}, %{{.*}} = "snitch.ssr_enable"() : () -> (!snitch.readable>, !snitch.writable>) // CHECK-NEXT: "test.op"() : () -> () // CHECK-NEXT: "snitch.ssr_disable"() : () -> () diff --git a/tests/filecheck/dialects/snitch_stream/ops.mlir b/tests/filecheck/dialects/snitch_stream/ops.mlir index 4989bd8d4a..51a9104863 100644 --- a/tests/filecheck/dialects/snitch_stream/ops.mlir +++ b/tests/filecheck/dialects/snitch_stream/ops.mlir @@ -37,14 +37,14 @@ snitch_stream.streaming_region { // CHECK-NEXT: } // CHECK-GENERIC: %X, %Y, %Z = "test.op"() : () -> (!riscv.reg, !riscv.reg, !riscv.reg) -// CHECK-GENERIC-NEXT: "snitch_stream.streaming_region"(%X, %Y, %Z) <{"stride_patterns" = [#snitch_stream.stride_pattern], "operandSegmentSizes" = array}> ({ +// CHECK-GENERIC-NEXT: "snitch_stream.streaming_region"(%X, %Y, %Z) <{stride_patterns = [#snitch_stream.stride_pattern], operandSegmentSizes = array}> ({ // CHECK-GENERIC-NEXT: ^0(%a_stream : !snitch.readable>, %b_stream : !snitch.readable>, %c_stream : !snitch.writable>): -// CHECK-GENERIC-NEXT: %c5 = "riscv.li"() {"immediate" = 5 : i32} : () -> !riscv.reg +// CHECK-GENERIC-NEXT: %c5 = "riscv.li"() {immediate = 5 : i32} : () -> !riscv.reg // CHECK-GENERIC-NEXT: "riscv_snitch.frep_outer"(%c5) ({ // CHECK-GENERIC-NEXT: %a = "riscv_snitch.read"(%a_stream) : (!snitch.readable>) -> !riscv.freg // CHECK-GENERIC-NEXT: %b = "riscv_snitch.read"(%b_stream) : (!snitch.readable>) -> !riscv.freg -// CHECK-GENERIC-NEXT: %c = "riscv.fadd.d"(%a, %b) {"fastmath" = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg +// CHECK-GENERIC-NEXT: %c = "riscv.fadd.d"(%a, %b) {fastmath = #riscv.fastmath} : (!riscv.freg, !riscv.freg) -> !riscv.freg // CHECK-GENERIC-NEXT: "riscv_snitch.write"(%c, %c_stream) : (!riscv.freg, !snitch.writable>) -> () // CHECK-GENERIC-NEXT: "riscv_snitch.frep_yield"() : () -> () -// CHECK-GENERIC-NEXT: }) {"stagger_mask" = #builtin.int<0>, "stagger_count" = #builtin.int<0>} : (!riscv.reg) -> () +// CHECK-GENERIC-NEXT: }) {stagger_mask = #builtin.int<0>, stagger_count = #builtin.int<0>} : (!riscv.reg) -> () // CHECK-GENERIC-NEXT: }) : (!riscv.reg, !riscv.reg, !riscv.reg) -> () diff --git a/tests/filecheck/dialects/stablehlo/attrs.mlir b/tests/filecheck/dialects/stablehlo/attrs.mlir index aa21766d2e..472c1f8b5c 100644 --- a/tests/filecheck/dialects/stablehlo/attrs.mlir +++ b/tests/filecheck/dialects/stablehlo/attrs.mlir @@ -2,7 +2,7 @@ // CHECK: builtin.module { -// CHECK-NEXT: "test.op"() {"default" = #stablehlo, "high" = #stablehlo, "highest" = #stablehlo} : () -> () +// CHECK-NEXT: "test.op"() {default = #stablehlo, high = #stablehlo, highest = #stablehlo} : () -> () "test.op"() { default = #stablehlo, high = #stablehlo, @@ -12,7 +12,7 @@ // CHECK-NEXT: %token = "test.op"() : () -> !stablehlo.token %token = "test.op"() : () -> (!stablehlo.token) -// CHECK-NEXT: "test.op"() {"dot" = #stablehlo.dot< +// CHECK-NEXT: "test.op"() {dot = #stablehlo.dot< // CHECK-NEXT: lhs_batching_dimensions = [0], // CHECK-NEXT: rhs_batching_dimensions = [1], // CHECK-NEXT: lhs_contracting_dimensions = [2], diff --git a/tests/filecheck/dialects/stablehlo/ops.mlir b/tests/filecheck/dialects/stablehlo/ops.mlir index 2028bc45d6..289b862cfc 100644 --- a/tests/filecheck/dialects/stablehlo/ops.mlir +++ b/tests/filecheck/dialects/stablehlo/ops.mlir @@ -24,7 +24,7 @@ // [[1,2], [3,4], [5,6]], // [[7,8], [9,10], [11,12]] // ] -// CHECK: %transpose_result = "stablehlo.transpose"(%transpose_operand) {"permutation" = array} : (tensor<2x3x2xi32>) -> tensor<2x3x2xi32> +// CHECK: %transpose_result = "stablehlo.transpose"(%transpose_operand) {permutation = array} : (tensor<2x3x2xi32>) -> tensor<2x3x2xi32> %transpose_result = "stablehlo.transpose"(%transpose_operand) { permutation = array } : (tensor<2x3x2xi32>) -> tensor<2x3x2xi32> diff --git a/tests/filecheck/dialects/stencil/oec-kernels/fvtp2d_qi.mlir b/tests/filecheck/dialects/stencil/oec-kernels/fvtp2d_qi.mlir index 7af8529faa..cc99168468 100644 --- a/tests/filecheck/dialects/stencil/oec-kernels/fvtp2d_qi.mlir +++ b/tests/filecheck/dialects/stencil/oec-kernels/fvtp2d_qi.mlir @@ -118,7 +118,7 @@ func.func @fvtp2d_qi(%arg0: !stencil.field, %arg1: !stencil.field, %arg1 : !stencil.field, %arg2 : !stencil.field, %arg3 : !stencil.field, %arg4 : !stencil.field, %arg5 : !stencil.field, %arg6 : !stencil.field) attributes {"stencil.program"}{ +// CHECK: func.func @fvtp2d_qi(%arg0 : !stencil.field, %arg1 : !stencil.field, %arg2 : !stencil.field, %arg3 : !stencil.field, %arg4 : !stencil.field, %arg5 : !stencil.field, %arg6 : !stencil.field) attributes {stencil.program}{ // CHECK-NEXT: %0 = stencil.cast %arg0 : !stencil.field -> !stencil.field<[-4,68]x[-4,68]x[-4,68]xf64> // CHECK-NEXT: %1 = stencil.cast %arg1 : !stencil.field -> !stencil.field<[-4,68]x[-4,68]x[-4,68]xf64> // CHECK-NEXT: %2 = stencil.cast %arg2 : !stencil.field -> !stencil.field<[-4,68]x[-4,68]x[-4,68]xf64> @@ -233,7 +233,7 @@ func.func @fvtp2d_qi(%arg0: !stencil.field, %arg1: !stencil.field, %arg1 : !stencil.field, %arg2 : !stencil.field, %arg3 : !stencil.field, %arg4 : !stencil.field, %arg5 : !stencil.field, %arg6 : !stencil.field) attributes {"stencil.program"}{ +// SHAPE: func.func @fvtp2d_qi(%arg0 : !stencil.field, %arg1 : !stencil.field, %arg2 : !stencil.field, %arg3 : !stencil.field, %arg4 : !stencil.field, %arg5 : !stencil.field, %arg6 : !stencil.field) attributes {stencil.program}{ // SHAPE-NEXT: %0 = stencil.cast %arg0 : !stencil.field -> !stencil.field<[-4,68]x[-4,68]x[-4,68]xf64> // SHAPE-NEXT: %1 = stencil.cast %arg1 : !stencil.field -> !stencil.field<[-4,68]x[-4,68]x[-4,68]xf64> // SHAPE-NEXT: %2 = stencil.cast %arg2 : !stencil.field -> !stencil.field<[-4,68]x[-4,68]x[-4,68]xf64> @@ -354,7 +354,7 @@ func.func @fvtp2d_qi(%arg0: !stencil.field, %arg1: !stencil.field, %arg1 : memref, %arg2 : memref, %arg3 : memref, %arg4 : memref, %arg5 : memref, %arg6 : memref) attributes {"stencil.program"}{ +// MLIR: func.func @fvtp2d_qi(%arg0 : memref, %arg1 : memref, %arg2 : memref, %arg3 : memref, %arg4 : memref, %arg5 : memref, %arg6 : memref) attributes {stencil.program}{ // MLIR-NEXT: %arg8 = memref.alloc() : memref<64x67x64xf64, strided<[4288, 64, 1], offset: 64>> // MLIR-NEXT: %arg9 = memref.alloc() : memref<64x66x64xf64, strided<[4224, 64, 1], offset: 64>> // MLIR-NEXT: %arg10 = memref.alloc() : memref<64x66x64xf64, strided<[4224, 64, 1], offset: 64>> @@ -384,7 +384,7 @@ func.func @fvtp2d_qi(%arg0: !stencil.field, %arg1: !stencil.field}> ({ +// MLIR-NEXT: "scf.parallel"(%13, %14, %15, %19, %20, %21, %16, %17, %18) <{operandSegmentSizes = array}> ({ // MLIR-NEXT: ^0(%22 : index, %23 : index, %24 : index): // MLIR-NEXT: %cst = arith.constant 1.000000e+00 : f64 // MLIR-NEXT: %cst_1 = arith.constant 7.000000e+00 : f64 @@ -418,7 +418,7 @@ func.func @fvtp2d_qi(%arg0: !stencil.field, %arg1: !stencil.field}> ({ +// MLIR-NEXT: "scf.parallel"(%42, %43, %44, %48, %49, %50, %45, %46, %47) <{operandSegmentSizes = array}> ({ // MLIR-NEXT: ^1(%51 : index, %52 : index, %53 : index): // MLIR-NEXT: %cst_3 = arith.constant 0.000000e+00 : f64 // MLIR-NEXT: %cst_4 = arith.constant 1.000000e+00 : f64 @@ -448,7 +448,7 @@ func.func @fvtp2d_qi(%arg0: !stencil.field, %arg1: !stencil.field}> ({ +// MLIR-NEXT: "scf.parallel"(%65, %66, %67, %71, %72, %73, %68, %69, %70) <{operandSegmentSizes = array}> ({ // MLIR-NEXT: ^2(%74 : index, %75 : index, %76 : index): // MLIR-NEXT: %cst_5 = arith.constant 0.000000e+00 : f64 // MLIR-NEXT: %cst_6 = arith.constant 1.000000e+00 : f64 @@ -507,7 +507,7 @@ func.func @fvtp2d_qi(%arg0: !stencil.field, %arg1: !stencil.field}> ({ +// MLIR-NEXT: "scf.parallel"(%112, %113, %114, %118, %119, %120, %115, %116, %117) <{operandSegmentSizes = array}> ({ // MLIR-NEXT: ^3(%121 : index, %122 : index, %123 : index): // MLIR-NEXT: %124 = memref.load %11[%121, %122, %123] : memref<64x65x64xf64, strided<[5184, 72, 1], offset: 21028>> // MLIR-NEXT: %125 = memref.load %arg8_1[%121, %122, %123] : memref<64x65x64xf64, strided<[5184, 72, 1], offset: 21028>> @@ -524,7 +524,7 @@ func.func @fvtp2d_qi(%arg0: !stencil.field, %arg1: !stencil.field}> ({ +// MLIR-NEXT: "scf.parallel"(%127, %128, %129, %133, %134, %135, %130, %131, %132) <{operandSegmentSizes = array}> ({ // MLIR-NEXT: ^4(%136 : index, %137 : index, %138 : index): // MLIR-NEXT: %139 = memref.load %8[%136, %137, %138] : memref<64x70x64xf64, strided<[5184, 72, 1], offset: 21028>> // MLIR-NEXT: %140 = memref.load %12[%136, %137, %138] : memref<64x64x64xf64, strided<[5184, 72, 1], offset: 21028>> @@ -549,7 +549,7 @@ func.func @fvtp2d_qi(%arg0: !stencil.field, %arg1: !stencil.field, %arg1 : !stencil.field, %arg2 : !stencil.field, %arg3 : !stencil.field, %arg4 : !stencil.field, %arg5 : !stencil.field, %arg6 : !stencil.field) attributes {"stencil.program"}{ +// BUFF: func.func @fvtp2d_qi(%arg0 : !stencil.field, %arg1 : !stencil.field, %arg2 : !stencil.field, %arg3 : !stencil.field, %arg4 : !stencil.field, %arg5 : !stencil.field, %arg6 : !stencil.field) attributes {stencil.program}{ // BUFF-NEXT: %0 = stencil.alloc : !stencil.field<[0,64]x[0,65]x[0,64]xf64> // BUFF-NEXT: %1 = stencil.alloc : !stencil.field<[0,64]x[-1,65]x[0,64]xf64> // BUFF-NEXT: %2 = stencil.alloc : !stencil.field<[0,64]x[-1,65]x[0,64]xf64> diff --git a/tests/filecheck/dialects/stencil/stencil_ops.mlir b/tests/filecheck/dialects/stencil/stencil_ops.mlir index 88618fb8f5..4b951dcbcb 100644 --- a/tests/filecheck/dialects/stencil/stencil_ops.mlir +++ b/tests/filecheck/dialects/stencil/stencil_ops.mlir @@ -51,7 +51,7 @@ builtin.module { } } -// CHECK: func.func private @myfunc(%b0 : !stencil.field, %b1 : !stencil.field) attributes {"param_names" = ["data"]}{ +// CHECK: func.func private @myfunc(%b0 : !stencil.field, %b1 : !stencil.field) attributes {param_names = ["data"]}{ // CHECK-NEXT: %f0 = stencil.cast %b0 : !stencil.field -> !stencil.field<[-4,54]x[-4,84]x[-4,44]xf32> // CHECK-NEXT: %f1 = stencil.cast %b1 : !stencil.field -> !stencil.field<[-4,54]x[-4,84]x[-4,44]xf32> // CHECK-NEXT: %time_m = arith.constant 0 : index diff --git a/tests/filecheck/dialects/stim/attrs.mlir b/tests/filecheck/dialects/stim/attrs.mlir index 35a632de22..2f746aeb71 100644 --- a/tests/filecheck/dialects/stim/attrs.mlir +++ b/tests/filecheck/dialects/stim/attrs.mlir @@ -8,6 +8,6 @@ %qubit0 = "test.op"() : () -> (!stim.qubit<0>) // CHECK: builtin.module { -// CHECK-NEXT: "test.op"() {"qubit" = !stim.qubit<0>, "qubitcoord" = #stim.qubit_coord<(0, 0), !stim.qubit<0>>} : () -> () +// CHECK-NEXT: "test.op"() {qubit = !stim.qubit<0>, qubitcoord = #stim.qubit_coord<(0, 0), !stim.qubit<0>>} : () -> () // CHECK-NEXT: %qubit0 = "test.op"() : () -> !stim.qubit<0> // CHECK-NEXT: } diff --git a/tests/filecheck/dialects/stim/stim_ops.mlir b/tests/filecheck/dialects/stim/stim_ops.mlir index a84c5dbfb1..99eba756fa 100644 --- a/tests/filecheck/dialects/stim/stim_ops.mlir +++ b/tests/filecheck/dialects/stim/stim_ops.mlir @@ -11,38 +11,38 @@ stim.circuit {} // CHECK-GENERIC-NEXT: }) : () -> () stim.circuit attributes {"hello" = "world"} {} -// CHECK-NEXT: stim.circuit attributes {"hello" = "world"} { +// CHECK-NEXT: stim.circuit attributes {hello = "world"} { // CHECK-NEXT: } // CHECK-GENERIC-NEXT: "stim.circuit"() ({ -// CHECK-GENERIC-NEXT: }) {"hello" = "world"} : () -> () +// CHECK-GENERIC-NEXT: }) {hello = "world"} : () -> () stim.assign_qubit_coord <(0, 0), !stim.qubit<0>> // CHECK-NEXT: stim.assign_qubit_coord <(0, 0), !stim.qubit<0>> -// CHECK-GENERIC-NEXT: "stim.assign_qubit_coord"() <{"qubitmapping" = #stim.qubit_coord<(0, 0), !stim.qubit<0>>}> : () -> () +// CHECK-GENERIC-NEXT: "stim.assign_qubit_coord"() <{qubitmapping = #stim.qubit_coord<(0, 0), !stim.qubit<0>>}> : () -> () stim.circuit {stim.assign_qubit_coord <(0, 0), !stim.qubit<0>>} // CHECK-NEXT: stim.circuit { // CHECK-NEXT: stim.assign_qubit_coord <(0, 0), !stim.qubit<0>> // CHECK-NEXT: } // CHECK-GENERIC-NEXT: "stim.circuit"() ({ -// CHECK-GENERIC-NEXT: "stim.assign_qubit_coord"() <{"qubitmapping" = #stim.qubit_coord<(0, 0), !stim.qubit<0>>}> +// CHECK-GENERIC-NEXT: "stim.assign_qubit_coord"() <{qubitmapping = #stim.qubit_coord<(0, 0), !stim.qubit<0>>}> // CHECK-GENERIC-NEXT: }) : () -> () stim.circuit attributes {"hello" = "world"} {stim.assign_qubit_coord <(0, 0), !stim.qubit<0>>} -// CHECK-NEXT: stim.circuit attributes {"hello" = "world"} { +// CHECK-NEXT: stim.circuit attributes {hello = "world"} { // CHECK-NEXT: stim.assign_qubit_coord <(0, 0), !stim.qubit<0>> // CHECK-NEXT: } // CHECK-GENERIC-NEXT: "stim.circuit"() ({ -// CHECK-GENERIC-NEXT: "stim.assign_qubit_coord"() <{"qubitmapping" = #stim.qubit_coord<(0, 0), !stim.qubit<0>>}> -// CHECK-GENERIC-NEXT: }) {"hello" = "world"} : () -> () +// CHECK-GENERIC-NEXT: "stim.assign_qubit_coord"() <{qubitmapping = #stim.qubit_coord<(0, 0), !stim.qubit<0>>}> +// CHECK-GENERIC-NEXT: }) {hello = "world"} : () -> () stim.circuit qubitlayout [#stim.qubit_coord<(0, 0), !stim.qubit<0>>] attributes {"hello" = "world"} {stim.assign_qubit_coord <(1, 2), !stim.qubit<2>>} -// CHECK-NEXT: stim.circuit qubitlayout [#stim.qubit_coord<(0, 0), !stim.qubit<0>>] attributes {"hello" = "world"} +// CHECK-NEXT: stim.circuit qubitlayout [#stim.qubit_coord<(0, 0), !stim.qubit<0>>] attributes {hello = "world"} // CHECK-NEXT: stim.assign_qubit_coord <(1, 2), !stim.qubit<2>> // CHECK-NEXT: } -// CHECK-GENERIC-NEXT: "stim.circuit"() <{"qubitlayout" = [#stim.qubit_coord<(0, 0), !stim.qubit<0>>]}> ({ -// CHECK-GENERIC-NEXT: "stim.assign_qubit_coord"() <{"qubitmapping" = #stim.qubit_coord<(1, 2), !stim.qubit<2>>}> -// CHECK-GENERIC-NEXT: }) {"hello" = "world"} : () -> () +// CHECK-GENERIC-NEXT: "stim.circuit"() <{qubitlayout = [#stim.qubit_coord<(0, 0), !stim.qubit<0>>]}> ({ +// CHECK-GENERIC-NEXT: "stim.assign_qubit_coord"() <{qubitmapping = #stim.qubit_coord<(1, 2), !stim.qubit<2>>}> +// CHECK-GENERIC-NEXT: }) {hello = "world"} : () -> () // CHECK-NEXT: } // CHECK-GENERIC-NEXT: }) : () -> () diff --git a/tests/filecheck/dialects/tensor/ops.mlir b/tests/filecheck/dialects/tensor/ops.mlir index 2136425a8b..e88bab030d 100644 --- a/tests/filecheck/dialects/tensor/ops.mlir +++ b/tests/filecheck/dialects/tensor/ops.mlir @@ -13,9 +13,9 @@ // CHECK: builtin.module { // CHECK-NEXT: %index, %index1, %tensor = "test.op"() : () -> (index, index, tensor) // CHECK-NEXT: %dim1 = tensor.dim %tensor, %index : tensor -// CHECK-NEXT: %dim2 = tensor.dim {"hello" = "world"} %tensor, %index : tensor +// CHECK-NEXT: %dim2 = tensor.dim {hello = "world"} %tensor, %index : tensor // CHECK-NEXT: %cast1 = tensor.cast %tensor : tensor to tensor<4x4xf32> -// CHECK-NEXT: %cast2 = tensor.cast %tensor {"hello" = "world"} : tensor to tensor<4x4xf32> +// CHECK-NEXT: %cast2 = tensor.cast %tensor {hello = "world"} : tensor to tensor<4x4xf32> // CHECK-NEXT: %extract1 = tensor.extract %tensor[%index, %index1] : tensor // CHECK-NEXT: %insert1 = tensor.insert %extract1 into %tensor[%index, %index1] : tensor // CHECK-NEXT: } @@ -23,9 +23,9 @@ // CHECK-GENERIC: "builtin.module"() ({ // CHECK-GENERIC-NEXT: %index, %index1, %tensor = "test.op"() : () -> (index, index, tensor) // CHECK-GENERIC-NEXT: %dim1 = "tensor.dim"(%tensor, %index) : (tensor, index) -> index -// CHECK-GENERIC-NEXT: %dim2 = "tensor.dim"(%tensor, %index) {"hello" = "world"} : (tensor, index) -> index +// CHECK-GENERIC-NEXT: %dim2 = "tensor.dim"(%tensor, %index) {hello = "world"} : (tensor, index) -> index // CHECK-GENERIC-NEXT: %cast1 = "tensor.cast"(%tensor) : (tensor) -> tensor<4x4xf32> -// CHECK-GENERIC-NEXT: %cast2 = "tensor.cast"(%tensor) {"hello" = "world"} : (tensor) -> tensor<4x4xf32> +// CHECK-GENERIC-NEXT: %cast2 = "tensor.cast"(%tensor) {hello = "world"} : (tensor) -> tensor<4x4xf32> // CHECK-GENERIC-NEXT: %extract1 = "tensor.extract"(%tensor, %index, %index1) : (tensor, index, index) -> f32 // CHECK-GENERIC-NEXT: %insert1 = "tensor.insert"(%extract1, %tensor, %index, %index1) : (f32, tensor, index, index) -> tensor // CHECK-GENERIC-NEXT: }) : () -> () diff --git a/tests/filecheck/dialects/tosa/ops.mlir b/tests/filecheck/dialects/tosa/ops.mlir index 22d7558fd4..f8b63a31cf 100644 --- a/tests/filecheck/dialects/tosa/ops.mlir +++ b/tests/filecheck/dialects/tosa/ops.mlir @@ -7,6 +7,6 @@ // CHECK: builtin.module { // CHECK-NEXT: %0 = "test.op"() : () -> tensor<12x34xi32> -// CHECK-NEXT: %1 = tosa.clamp %0 {"min_int" = 0 : i64, "max_int" = 1 : i64, "min_fp" = 0.000000e+00 : f32, "max_fp" = 1.000000e+00 : f32} : (tensor<12x34xi32>) -> tensor<12x34xi32> -// CHECK-NEXT: %2 = tosa.rescale %0 {"input_zp" = 127 : i32, "output_zp" = -1 : i32, "multiplier" = array, "shift" = array, "scale32" = true, "double_round" = false, "per_channel" = false} : (tensor<12x34xi32>) -> tensor<12x34xi32> +// CHECK-NEXT: %1 = tosa.clamp %0 {min_int = 0 : i64, max_int = 1 : i64, min_fp = 0.000000e+00 : f32, max_fp = 1.000000e+00 : f32} : (tensor<12x34xi32>) -> tensor<12x34xi32> +// CHECK-NEXT: %2 = tosa.rescale %0 {input_zp = 127 : i32, output_zp = -1 : i32, multiplier = array, shift = array, scale32 = true, double_round = false, per_channel = false} : (tensor<12x34xi32>) -> tensor<12x34xi32> // CHECK-NEXT: } diff --git a/tests/filecheck/dialects/transform/transform_types.mlir b/tests/filecheck/dialects/transform/transform_types.mlir index dd0c52c60e..478b4ec3d2 100644 --- a/tests/filecheck/dialects/transform/transform_types.mlir +++ b/tests/filecheck/dialects/transform/transform_types.mlir @@ -45,8 +45,8 @@ builtin.module attributes {"transform.with_named_sequence"} { -//CHECK: builtin.module attributes {"transform.with_named_sequence"} { -//CHECK-NEXT: "transform.named_sequence"() <{"arg_attrs" = [{"transform.readonly"}], "function_type" = (!transform.any_op) -> (), "sym_name" = "foo"}> ({ +//CHECK: builtin.module attributes {transform.with_named_sequence} { +//CHECK-NEXT: "transform.named_sequence"() <{arg_attrs = [{transform.readonly}], function_type = (!transform.any_op) -> (), sym_name = "foo"}> ({ //CHECK-NEXT: ^0(%arg0 : !transform.any_op): //CHECK-NEXT: "transform.yield"() : () -> () //CHECK-NEXT: }) : () -> () @@ -56,34 +56,34 @@ builtin.module attributes {"transform.with_named_sequence"} { //CHECK-NEXT: %3 = "test.op"() : () -> !transform.op<"linalg.quantized_matmul"> //CHECK-NEXT: %4 = "test.op"() : () -> !transform.param //CHECK-NEXT: %5 = "test.op"() : () -> !transform.type -//CHECK-NEXT: "transform.named_sequence"() <{"function_type" = (!transform.any_op, !transform.op<"linalg.quantized_matmul">, !transform.op<"linalg.elemwise_binary">) -> (), "sym_name" = "__transform_main"}> ({ +//CHECK-NEXT: "transform.named_sequence"() <{function_type = (!transform.any_op, !transform.op<"linalg.quantized_matmul">, !transform.op<"linalg.elemwise_binary">) -> (), sym_name = "__transform_main"}> ({ //CHECK-NEXT: ^1(%arg0_1 : !transform.any_op, %arg1 : !transform.op<"linalg.quantized_matmul">, %arg2 : !transform.op<"linalg.elemwise_binary">): //CHECK-NEXT: %6 = "transform.cast"(%arg1) : (!transform.op<"linalg.quantized_matmul">) -> !transform.any_op -//CHECK-NEXT: %7, %8 = "transform.structured.tile_using_forall"(%arg1) <{"operandSegmentSizes" = array, "static_tile_sizes" = array}> : (!transform.op<"linalg.quantized_matmul">) -> (!transform.any_op, !transform.any_op) -//CHECK-NEXT: %9, %10, %11 = "transform.structured.tile_using_for"(%arg1) <{"scalable_sizes" = array, "static_sizes" = array}> : (!transform.op<"linalg.quantized_matmul">) -> (!transform.any_op, !transform.any_op, !transform.any_op) +//CHECK-NEXT: %7, %8 = "transform.structured.tile_using_forall"(%arg1) <{operandSegmentSizes = array, static_tile_sizes = array}> : (!transform.op<"linalg.quantized_matmul">) -> (!transform.any_op, !transform.any_op) +//CHECK-NEXT: %9, %10, %11 = "transform.structured.tile_using_for"(%arg1) <{scalable_sizes = array, static_sizes = array}> : (!transform.op<"linalg.quantized_matmul">) -> (!transform.any_op, !transform.any_op, !transform.any_op) //CHECK-NEXT: "transform.yield"() : () -> () //CHECK-NEXT: }) : () -> () -//CHECK-NEXT: "transform.sequence"() <{"failure_propagation_mode" = 1 : i32, "operandSegmentSizes" = array}> ({ +//CHECK-NEXT: "transform.sequence"() <{failure_propagation_mode = 1 : i32, operandSegmentSizes = array}> ({ //CHECK-NEXT: ^2(%arg0_2 : !transform.any_op): -//CHECK-NEXT: %arg1_1 = "transform.select"(%arg0_2) <{"op_name" = "linalg.quantized_matmul"}> : (!transform.any_op) -> !transform.op<"linalg.quantized_matmul"> -//CHECK-NEXT: %12, %13, %14 = "transform.structured.tile_using_for"(%arg1_1) <{"scalable_sizes" = array, "static_sizes" = array}> : (!transform.op<"linalg.quantized_matmul">) -> (!transform.any_op, !transform.any_op, !transform.any_op) +//CHECK-NEXT: %arg1_1 = "transform.select"(%arg0_2) <{op_name = "linalg.quantized_matmul"}> : (!transform.any_op) -> !transform.op<"linalg.quantized_matmul"> +//CHECK-NEXT: %12, %13, %14 = "transform.structured.tile_using_for"(%arg1_1) <{scalable_sizes = array, static_sizes = array}> : (!transform.op<"linalg.quantized_matmul">) -> (!transform.any_op, !transform.any_op, !transform.any_op) //CHECK-NEXT: "transform.yield"() : () -> () //CHECK-NEXT: }) : () -> () //CHECK-NEXT: %12 = "test.op"() : () -> !transform.any_op -//CHECK-NEXT: %13 = "transform.get_producer_of_operand"(%12) <{"operand_number" = 0 : i64}> : (!transform.any_op) -> !transform.any_op -//CHECK-NEXT: %14 = "transform.get_consumers_of_result"(%13) <{"result_number" = 0 : i64}> : (!transform.any_op) -> !transform.any_op +//CHECK-NEXT: %13 = "transform.get_producer_of_operand"(%12) <{operand_number = 0 : i64}> : (!transform.any_op) -> !transform.any_op +//CHECK-NEXT: %14 = "transform.get_consumers_of_result"(%13) <{result_number = 0 : i64}> : (!transform.any_op) -> !transform.any_op //CHECK-NEXT: %15 = "test.op"() : () -> !transform.any_value //CHECK-NEXT: %16 = "transform.get_defining_op"(%15) : (!transform.any_value) -> !transform.any_op -//CHECK-NEXT: %17 = "transform.get_parent_op"(%16) <{"isolated_from_above", "nth_parent" = 1 : i64}> : (!transform.any_op) -> !transform.any_op -//CHECK-NEXT: %18 = "transform.get_result"(%17) <{"raw_position_list" = array}> : (!transform.any_op) -> !transform.any_value -//CHECK-NEXT: %19 = "transform.get_type"(%18) <{"elemental"}> : (!transform.any_value) -> !transform.type -//CHECK-NEXT: "transform.include"(%17) <{"failure_propagation_mode" = 1 : i32, "target" = @foo}> : (!transform.any_op) -> () +//CHECK-NEXT: %17 = "transform.get_parent_op"(%16) <{isolated_from_above, nth_parent = 1 : i64}> : (!transform.any_op) -> !transform.any_op +//CHECK-NEXT: %18 = "transform.get_result"(%17) <{raw_position_list = array}> : (!transform.any_op) -> !transform.any_value +//CHECK-NEXT: %19 = "transform.get_type"(%18) <{elemental}> : (!transform.any_value) -> !transform.type +//CHECK-NEXT: "transform.include"(%17) <{failure_propagation_mode = 1 : i32, target = @foo}> : (!transform.any_op) -> () //CHECK-NEXT: "transform.match.operation_empty"(%17) : (!transform.any_op) -> () -//CHECK-NEXT: "transform.match.operation_name"(%17) <{"op_names" = ["scf.for"]}> : (!transform.any_op) -> () +//CHECK-NEXT: "transform.match.operation_name"(%17) <{op_names = ["scf.for"]}> : (!transform.any_op) -> () //CHECK-NEXT: %20 = "transform.merge_handles"(%16, %17) : (!transform.any_op, !transform.any_op) -> !transform.any_op //CHECK-NEXT: %21 = "test.op"() : () -> !transform.any_param //CHECK-NEXT: %22 = "test.op"() : () -> !transform.any_param -//CHECK-NEXT: "transform.match.param.cmpi"(%21, %22) <{"predicate" = 1 : i32}> : (!transform.any_param, !transform.any_param) -> () -//CHECK-NEXT: %23, %24 = "transform.split_handle"(%20) <{"fail_on_payload_too_small" = true, "pass_through_empty_handle" = true}> : (!transform.any_op) -> (!transform.any_op, !transform.any_op) -//CHECK-NEXT: %25 = "transform.structured.match"(%20) <{"op_attrs" = {"qmatmul_0"}}> : (!transform.any_op) -> !transform.any_op +//CHECK-NEXT: "transform.match.param.cmpi"(%21, %22) <{predicate = 1 : i32}> : (!transform.any_param, !transform.any_param) -> () +//CHECK-NEXT: %23, %24 = "transform.split_handle"(%20) <{fail_on_payload_too_small = true, pass_through_empty_handle = true}> : (!transform.any_op) -> (!transform.any_op, !transform.any_op) +//CHECK-NEXT: %25 = "transform.structured.match"(%20) <{op_attrs = {qmatmul_0}}> : (!transform.any_op) -> !transform.any_op //CHECK-NEXT:} diff --git a/tests/filecheck/dialects/varith/varith_ops.mlir b/tests/filecheck/dialects/varith/varith_ops.mlir index ef6cb4e2ba..398099b0c3 100644 --- a/tests/filecheck/dialects/varith/varith_ops.mlir +++ b/tests/filecheck/dialects/varith/varith_ops.mlir @@ -43,4 +43,4 @@ // CHECK-NEXT: 2: %fd // CHECK-NEXT: ] -// CHECK-GENERIC: %x7 = "varith.switch"(%ia, %fa, %fb, %fc, %fd) <{"case_values" = dense<[0, 1, 2]> : vector<3xi32>}> : (i32, f32, f32, f32, f32) -> f32 +// CHECK-GENERIC: %x7 = "varith.switch"(%ia, %fa, %fb, %fc, %fd) <{case_values = dense<[0, 1, 2]> : vector<3xi32>}> : (i32, f32, f32, f32, f32) -> f32 diff --git a/tests/filecheck/dialects/wasm/ops.mlir b/tests/filecheck/dialects/wasm/ops.mlir index 1e25ee3414..ee1e54bb67 100644 --- a/tests/filecheck/dialects/wasm/ops.mlir +++ b/tests/filecheck/dialects/wasm/ops.mlir @@ -10,8 +10,8 @@ wasm.module // CHECK-GENERIC-NEXT: "wasm.module"() : () -> () wasm.module attributes {"hello" = "world"} -// CHECK-NEXT: wasm.module attributes {"hello" = "world"} -// CHECK-GENERIC-NEXT: "wasm.module"() {"hello" = "world"} : () -> () +// CHECK-NEXT: wasm.module attributes {hello = "world"} +// CHECK-GENERIC-NEXT: "wasm.module"() {hello = "world"} : () -> () // CHECK-NEXT: } diff --git a/tests/filecheck/frontend/dialects/affine.py b/tests/filecheck/frontend/dialects/affine.py index f2b5541518..2a302ef430 100644 --- a/tests/filecheck/frontend/dialects/affine.py +++ b/tests/filecheck/frontend/dialects/affine.py @@ -7,7 +7,7 @@ p = FrontendProgram() with CodeContext(p): # CHECK: func.func @test_affine_for_I() { - # CHECK-NEXT: "affine.for"() <{"lowerBoundMap" = affine_map<() -> (0)>, "upperBoundMap" = affine_map<() -> (100)>, "step" = 1 : index, "operandSegmentSizes" = array}> ({ + # CHECK-NEXT: "affine.for"() <{lowerBoundMap = affine_map<() -> (0)>, upperBoundMap = affine_map<() -> (100)>, step = 1 : index, operandSegmentSizes = array}> ({ # CHECK-NEXT: ^0(%0 : index): # CHECK-NEXT: "affine.yield"() : () -> () # CHECK-NEXT: }) : () -> () @@ -20,7 +20,7 @@ def test_affine_for_I(): return # CHECK: func.func @test_affine_for_II() { - # CHECK-NEXT: "affine.for"() <{"lowerBoundMap" = affine_map<() -> (10)>, "upperBoundMap" = affine_map<() -> (30)>, "step" = 1 : index, "operandSegmentSizes" = array}> ({ + # CHECK-NEXT: "affine.for"() <{lowerBoundMap = affine_map<() -> (10)>, upperBoundMap = affine_map<() -> (30)>, step = 1 : index, operandSegmentSizes = array}> ({ # CHECK-NEXT: ^0(%0 : index): # CHECK-NEXT: "affine.yield"() : () -> () # CHECK-NEXT: }) : () -> () @@ -32,7 +32,7 @@ def test_affine_for_II(): return # CHECK: func.func @test_affine_for_III() { - # CHECK-NEXT: "affine.for"() <{"lowerBoundMap" = affine_map<() -> (1)>, "upperBoundMap" = affine_map<() -> (20)>, "step" = 5 : index, "operandSegmentSizes" = array}> ({ + # CHECK-NEXT: "affine.for"() <{lowerBoundMap = affine_map<() -> (1)>, upperBoundMap = affine_map<() -> (20)>, step = 5 : index, operandSegmentSizes = array}> ({ # CHECK-NEXT: ^0(%0 : index): # CHECK-NEXT: "affine.yield"() : () -> () # CHECK-NEXT: }) : () -> () @@ -44,11 +44,11 @@ def test_affine_for_III(): return # CHECK: func.func @test_affine_for_IV() { - # CHECK-NEXT: "affine.for"() <{"lowerBoundMap" = affine_map<() -> (0)>, "upperBoundMap" = affine_map<() -> (10)>, "step" = 1 : index, "operandSegmentSizes" = array}> ({ + # CHECK-NEXT: "affine.for"() <{lowerBoundMap = affine_map<() -> (0)>, upperBoundMap = affine_map<() -> (10)>, step = 1 : index, operandSegmentSizes = array}> ({ # CHECK-NEXT: ^0(%0 : index): - # CHECK-NEXT: "affine.for"() <{"lowerBoundMap" = affine_map<() -> (0)>, "upperBoundMap" = affine_map<() -> (20)>, "step" = 1 : index, "operandSegmentSizes" = array}> ({ + # CHECK-NEXT: "affine.for"() <{lowerBoundMap = affine_map<() -> (0)>, upperBoundMap = affine_map<() -> (20)>, step = 1 : index, operandSegmentSizes = array}> ({ # CHECK-NEXT: ^1(%1 : index): - # CHECK-NEXT: "affine.for"() <{"lowerBoundMap" = affine_map<() -> (0)>, "upperBoundMap" = affine_map<() -> (30)>, "step" = 1 : index, "operandSegmentSizes" = array}> ({ + # CHECK-NEXT: "affine.for"() <{lowerBoundMap = affine_map<() -> (0)>, upperBoundMap = affine_map<() -> (30)>, step = 1 : index, operandSegmentSizes = array}> ({ # CHECK-NEXT: ^2(%2 : index): # CHECK-NEXT: "affine.yield"() : () -> () # CHECK-NEXT: }) : () -> () diff --git a/tests/filecheck/frontend/dialects/func.py b/tests/filecheck/frontend/dialects/func.py index d97e3b00bd..899cce0b3b 100644 --- a/tests/filecheck/frontend/dialects/func.py +++ b/tests/filecheck/frontend/dialects/func.py @@ -7,8 +7,8 @@ p = FrontendProgram() with CodeContext(p): # CHECK: func.func @f1(%{{.*}} : i32) { - # CHECK-NEXT: "symref.declare"() {"sym_name" = "x"} : () -> () - # CHECK-NEXT: "symref.update"(%{{.*}}) {"symbol" = @x} : (i32) -> () + # CHECK-NEXT: "symref.declare"() {sym_name = "x"} : () -> () + # CHECK-NEXT: "symref.update"(%{{.*}}) {symbol = @x} : (i32) -> () # CHECK-NEXT: func.return # CHECK-NEXT: } def f1(x: i32): @@ -21,9 +21,9 @@ def f2(): return # CHECK: func.func @f3(%{{.*}} : i32) -> i32 { - # CHECK-NEXT: "symref.declare"() {"sym_name" = "x"} : () -> () - # CHECK-NEXT: "symref.update"(%{{.*}}) {"symbol" = @x} : (i32) -> () - # CHECK-NEXT: %{{.*}} = "symref.fetch"() {"symbol" = @x} : () -> i32 + # CHECK-NEXT: "symref.declare"() {sym_name = "x"} : () -> () + # CHECK-NEXT: "symref.update"(%{{.*}}) {symbol = @x} : (i32) -> () + # CHECK-NEXT: %{{.*}} = "symref.fetch"() {symbol = @x} : () -> i32 # CHECK-NEXT: func.return %{{.*}} : i32 # CHECK-NEXT: } def f3(x: i32) -> i32: diff --git a/tests/filecheck/frontend/dialects/scf.py b/tests/filecheck/frontend/dialects/scf.py index 4e45d47acf..b2d27b3641 100644 --- a/tests/filecheck/frontend/dialects/scf.py +++ b/tests/filecheck/frontend/dialects/scf.py @@ -9,7 +9,7 @@ with CodeContext(p): # CHECK: func.func @test_for_I(%{{.*}} : index) { # CHECK: %{{.*}} = arith.constant 0 : index - # CHECK-NEXT: %{{.*}} = "symref.fetch"() {"symbol" = @end} : () -> index + # CHECK-NEXT: %{{.*}} = "symref.fetch"() {symbol = @end} : () -> index # CHECK-NEXT: %{{.*}} = arith.constant 1 : index # CHECK-NEXT: scf.for %{{.*}} = %{{.*}} to %{{.*}} step %{{.*}} { # CHECK-NEXT: } @@ -24,8 +24,8 @@ def test_for_I(end: index): return # CHECK: func.func @test_for_II(%{{.*}} : index, %{{.*}} : index) { - # CHECK: %{{.*}} = "symref.fetch"() {"symbol" = @start} : () -> index - # CHECK-NEXT: %{{.*}} = "symref.fetch"() {"symbol" = @end} : () -> index + # CHECK: %{{.*}} = "symref.fetch"() {symbol = @start} : () -> index + # CHECK-NEXT: %{{.*}} = "symref.fetch"() {symbol = @end} : () -> index # CHECK-NEXT: %{{.*}} = arith.constant 1 : index # CHECK-NEXT: scf.for %{{.*}} = %{{.*}} to %{{.*}} step %{{.*}} { # CHECK-NEXT: } @@ -40,9 +40,9 @@ def test_for_II(start: index, end: index): return # CHECK: func.func @test_for_III(%{{.*}} : index, %{{.*}} : index, %{{.*}} : index) { - # CHECK: %{{.*}} = "symref.fetch"() {"symbol" = @start} : () -> index - # CHECK-NEXT: %{{.*}} = "symref.fetch"() {"symbol" = @end} : () -> index - # CHECK-NEXT: %{{.*}} = "symref.fetch"() {"symbol" = @step} : () -> index + # CHECK: %{{.*}} = "symref.fetch"() {symbol = @start} : () -> index + # CHECK-NEXT: %{{.*}} = "symref.fetch"() {symbol = @end} : () -> index + # CHECK-NEXT: %{{.*}} = "symref.fetch"() {symbol = @step} : () -> index # CHECK-NEXT: scf.for %{{.*}} = %{{.*}} to %{{.*}} step %{{.*}} { # CHECK-NEXT: } # CHECK-NEXT: func.return @@ -58,15 +58,15 @@ def test_for_III(start: index, end: index, step: index): # CHECK: func.func @test_for_IV(%{{.*}} : index, %{{.*}} : index, %{{.*}} : index) { # CHECK: %{{.*}} = arith.constant 0 : index - # CHECK-NEXT: %{{.*}} = "symref.fetch"() {"symbol" = @a} : () -> index + # CHECK-NEXT: %{{.*}} = "symref.fetch"() {symbol = @a} : () -> index # CHECK-NEXT: %{{.*}} = arith.constant 1 : index # CHECK-NEXT: scf.for %{{.*}} = %{{.*}} to %{{.*}} step %{{.*}} { # CHECK-NEXT: %{{.*}} = arith.constant 0 : index - # CHECK-NEXT: %{{.*}} = "symref.fetch"() {"symbol" = @b} : () -> index + # CHECK-NEXT: %{{.*}} = "symref.fetch"() {symbol = @b} : () -> index # CHECK-NEXT: %{{.*}} = arith.constant 1 : index # CHECK-NEXT: scf.for %{{.*}} = %{{.*}} to %{{.*}} step %{{.*}} { # CHECK-NEXT: %{{.*}} = arith.constant 0 : index - # CHECK-NEXT: %{{.*}} = "symref.fetch"() {"symbol" = @c} : () -> index + # CHECK-NEXT: %{{.*}} = "symref.fetch"() {symbol = @c} : () -> index # CHECK-NEXT: %{{.*}} = arith.constant 1 : index # CHECK-NEXT: scf.for %{{.*}} = %{{.*}} to %{{.*}} step %{{.*}} { # CHECK-NEXT: } @@ -134,10 +134,10 @@ def test_not_supported_loop_III(start: index, end: index, step: f32): p = FrontendProgram() with CodeContext(p): # CHECK: %{{.*}} = scf.if %{{.*}} -> (i32) { - # CHECK-NEXT: %{{.*}} = "symref.fetch"() {"symbol" = @x} : () -> i32 + # CHECK-NEXT: %{{.*}} = "symref.fetch"() {symbol = @x} : () -> i32 # CHECK-NEXT: scf.yield %{{.*}} : i32 # CHECK-NEXT: } else { - # CHECK-NEXT: %{{.*}} = "symref.fetch"() {"symbol" = @y} : () -> i32 + # CHECK-NEXT: %{{.*}} = "symref.fetch"() {symbol = @y} : () -> i32 # CHECK-NEXT: scf.yield %{{.*}} : i32 # CHECK-NEXT: } def test_if_expr(cond: i1, x: i32, y: i32) -> i32: @@ -152,13 +152,13 @@ def test_if_I(cond: i1): pass return - # CHECK: %{{.*}} = "symref.fetch"() {"symbol" = @a} : () -> i1 + # CHECK: %{{.*}} = "symref.fetch"() {symbol = @a} : () -> i1 # CHECK-NEXT: scf.if %{{.*}} { # CHECK-NEXT: } else { - # CHECK-NEXT: %{{.*}} = "symref.fetch"() {"symbol" = @b} : () -> i1 + # CHECK-NEXT: %{{.*}} = "symref.fetch"() {symbol = @b} : () -> i1 # CHECK-NEXT: scf.if %{{.*}} { # CHECK-NEXT: } else { - # CHECK-NEXT: %{{.*}} = "symref.fetch"() {"symbol" = @c} : () -> i1 + # CHECK-NEXT: %{{.*}} = "symref.fetch"() {symbol = @c} : () -> i1 # CHECK-NEXT: scf.if %{{.*}} { # CHECK-NEXT: } # CHECK-NEXT: } @@ -172,7 +172,7 @@ def test_if_II(a: i1, b: i1, c: i1): pass return - # CHECK: %{{.*}} = "symref.fetch"() {"symbol" = @cond} : () -> i1 + # CHECK: %{{.*}} = "symref.fetch"() {symbol = @cond} : () -> i1 # CHECK-NEXT: scf.if %{{.*}} { # CHECK-NEXT: } def test_if_III(cond: i1): diff --git a/tests/filecheck/frontend/passes/desymref.mlir b/tests/filecheck/frontend/passes/desymref.mlir index add46edcfb..71ab1daaac 100644 --- a/tests/filecheck/frontend/passes/desymref.mlir +++ b/tests/filecheck/frontend/passes/desymref.mlir @@ -8,7 +8,7 @@ }) : () -> () // CHECK-NEXT: "builtin.module"() ({ - // CHECK-NEXT: %{{.*}} = "arith.constant"() <{"value" = 5 : i32}> : () -> i32 + // CHECK-NEXT: %{{.*}} = "arith.constant"() <{value = 5 : i32}> : () -> i32 // CHECK-NEXT: }) : () -> () @@ -32,9 +32,9 @@ }) : () -> () // CHECK-NEXT: "builtin.module"() ({ - // CHECK-NEXT: %{{.*}} = "arith.constant"() <{"value" = 42 : i32}> : () -> i32 - // CHECK-NEXT: %{{.*}} = "arith.constant"() <{"value" = 11 : i32}> : () -> i32 - // CHECK-NEXT: %{{.*}} = "arith.constant"() <{"value" = 23 : i32}> : () -> i32 + // CHECK-NEXT: %{{.*}} = "arith.constant"() <{value = 42 : i32}> : () -> i32 + // CHECK-NEXT: %{{.*}} = "arith.constant"() <{value = 11 : i32}> : () -> i32 + // CHECK-NEXT: %{{.*}} = "arith.constant"() <{value = 23 : i32}> : () -> i32 // CHECK-NEXT: }) : () -> () @@ -49,10 +49,10 @@ }) : () -> () // CHECK-NEXT: "builtin.module"() ({ - // CHECK-NEXT: %{{.*}} = "arith.constant"() <{"value" = 42 : i32}> : () -> i32 - // CHECK-NEXT: %{{.*}} = "arith.addi"(%{{.*}}, %{{.*}}) <{"overflowFlags" = #arith.overflow}> : (i32, i32) -> i32 - // CHECK-NEXT: %{{.*}} = "arith.constant"() <{"value" = 7 : i32}> : () -> i32 - // CHECK-NEXT: %{{.*}} = "arith.muli"(%{{.*}}, %{{.*}}) <{"overflowFlags" = #arith.overflow}> : (i32, i32) -> i32 + // CHECK-NEXT: %{{.*}} = "arith.constant"() <{value = 42 : i32}> : () -> i32 + // CHECK-NEXT: %{{.*}} = "arith.addi"(%{{.*}}, %{{.*}}) <{overflowFlags = #arith.overflow}> : (i32, i32) -> i32 + // CHECK-NEXT: %{{.*}} = "arith.constant"() <{value = 7 : i32}> : () -> i32 + // CHECK-NEXT: %{{.*}} = "arith.muli"(%{{.*}}, %{{.*}}) <{overflowFlags = #arith.overflow}> : (i32, i32) -> i32 // CHECK-NEXT: }) : () -> () @@ -69,9 +69,9 @@ }) : () -> () // CHECK-NEXT: "builtin.module"() ({ - // CHECK-NEXT: %{{.*}} = "arith.constant"() <{"value" = 11 : i32}> : () -> i32 - // CHECK-NEXT: %{{.*}} = "arith.constant"() <{"value" = 22 : i32}> : () -> i32 - // CHECK-NEXT: %{{.*}} = "arith.addi"(%{{.*}}, %{{.*}}) <{"overflowFlags" = #arith.overflow}> : (i32, i32) -> i32 + // CHECK-NEXT: %{{.*}} = "arith.constant"() <{value = 11 : i32}> : () -> i32 + // CHECK-NEXT: %{{.*}} = "arith.constant"() <{value = 22 : i32}> : () -> i32 + // CHECK-NEXT: %{{.*}} = "arith.addi"(%{{.*}}, %{{.*}}) <{overflowFlags = #arith.overflow}> : (i32, i32) -> i32 // CHECK-NEXT: }) : () -> () @@ -99,12 +99,12 @@ }) : () -> () // CHECK-NEXT: "builtin.module"() ({ - // CHECK-NEXT: %{{.*}} = "arith.constant"() <{"value" = 0 : i32}> : () -> i32 - // CHECK-NEXT: %{{.*}} = "arith.constant"() <{"value" = 1 : i32}> : () -> i32 - // CHECK-NEXT: %{{.*}} = "arith.constant"() <{"value" = 2 : i32}> : () -> i32 - // CHECK-NEXT: %{{.*}} = "arith.addi"(%{{.*}}, %{{.*}}) <{"overflowFlags" = #arith.overflow}> : (i32, i32) -> i32 - // CHECK-NEXT: %{{.*}} = "arith.muli"(%{{.*}}, %{{.*}}) <{"overflowFlags" = #arith.overflow}> : (i32, i32) -> i32 - // CHECK-NEXT: %{{.*}} = "arith.addi"(%{{.*}}, %{{.*}}) <{"overflowFlags" = #arith.overflow}> : (i32, i32) -> i32 + // CHECK-NEXT: %{{.*}} = "arith.constant"() <{value = 0 : i32}> : () -> i32 + // CHECK-NEXT: %{{.*}} = "arith.constant"() <{value = 1 : i32}> : () -> i32 + // CHECK-NEXT: %{{.*}} = "arith.constant"() <{value = 2 : i32}> : () -> i32 + // CHECK-NEXT: %{{.*}} = "arith.addi"(%{{.*}}, %{{.*}}) <{overflowFlags = #arith.overflow}> : (i32, i32) -> i32 + // CHECK-NEXT: %{{.*}} = "arith.muli"(%{{.*}}, %{{.*}}) <{overflowFlags = #arith.overflow}> : (i32, i32) -> i32 + // CHECK-NEXT: %{{.*}} = "arith.addi"(%{{.*}}, %{{.*}}) <{overflowFlags = #arith.overflow}> : (i32, i32) -> i32 // CHECK-NEXT: }) : () -> () @@ -131,11 +131,11 @@ }) : () -> () // CHECK-NEXT: "builtin.module"() ({ - // CHECK-NEXT: %{{.*}} = "arith.constant"() <{"value" = 0 : i32}> : () -> i32 - // CHECK-NEXT: %{{.*}} = "arith.constant"() <{"value" = 1 : i32}> : () -> i32 - // CHECK-NEXT: "symref.update"(%{{.*}}) {"symbol" = @a} : (i32) -> () - // CHECK-NEXT: %{{.*}} = "arith.constant"() <{"value" = 2 : i32}> : () -> i32 - // CHECK-NEXT: "symref.update"(%{{.*}}) {"symbol" = @c} : (i32) -> () + // CHECK-NEXT: %{{.*}} = "arith.constant"() <{value = 0 : i32}> : () -> i32 + // CHECK-NEXT: %{{.*}} = "arith.constant"() <{value = 1 : i32}> : () -> i32 + // CHECK-NEXT: "symref.update"(%{{.*}}) {symbol = @a} : (i32) -> () + // CHECK-NEXT: %{{.*}} = "arith.constant"() <{value = 2 : i32}> : () -> i32 + // CHECK-NEXT: "symref.update"(%{{.*}}) {symbol = @c} : (i32) -> () // CHECK-NEXT: }) : () -> () @@ -148,10 +148,10 @@ }) : () -> () // CHECK-NEXT: "builtin.module"() ({ - // CHECK-NEXT: %{{.*}} = "symref.fetch"() {"symbol" = @b} : () -> i32 - // CHECK-NEXT: %{{.*}} = "arith.muli"(%{{.*}}, %{{.*}}) <{"overflowFlags" = #arith.overflow}> : (i32, i32) -> i32 - // CHECK-NEXT: %{{.*}} = "arith.constant"() <{"value" = 5 : i32}> : () -> i32 - // CHECK-NEXT: %{{.*}} = "arith.addi"(%{{.*}}, %{{.*}}) <{"overflowFlags" = #arith.overflow}> : (i32, i32) -> i32 + // CHECK-NEXT: %{{.*}} = "symref.fetch"() {symbol = @b} : () -> i32 + // CHECK-NEXT: %{{.*}} = "arith.muli"(%{{.*}}, %{{.*}}) <{overflowFlags = #arith.overflow}> : (i32, i32) -> i32 + // CHECK-NEXT: %{{.*}} = "arith.constant"() <{value = 5 : i32}> : () -> i32 + // CHECK-NEXT: %{{.*}} = "arith.addi"(%{{.*}}, %{{.*}}) <{overflowFlags = #arith.overflow}> : (i32, i32) -> i32 // CHECK-NEXT: }) : () -> () @@ -177,14 +177,14 @@ }) : () -> () // CHECK-NEXT: "builtin.module"() ({ - // CHECK-NEXT: %{{.*}} = "arith.constant"() <{"value" = 1 : i32}> : () -> i32 - // CHECK-NEXT: "symref.update"(%{{.*}}) {"symbol" = @b} : (i32) -> () - // CHECK-NEXT: %{{.*}} = "arith.constant"() <{"value" = 2 : i32}> : () -> i32 - // CHECK-NEXT: %{{.*}} = "arith.addi"(%{{.*}}, %{{.*}}) <{"overflowFlags" = #arith.overflow}> : (i32, i32) -> i32 - // CHECK-NEXT: %{{.*}} = "arith.muli"(%{{.*}}, %{{.*}}) <{"overflowFlags" = #arith.overflow}> : (i32, i32) -> i32 - // CHECK-NEXT: "symref.update"(%{{.*}}) {"symbol" = @a} : (i32) -> () - // CHECK-NEXT: %{{.*}} = "arith.addi"(%{{.*}}, %{{.*}}) <{"overflowFlags" = #arith.overflow}> : (i32, i32) -> i32 - // CHECK-NEXT: "symref.update"(%{{.*}}) {"symbol" = @c} : (i32) -> () + // CHECK-NEXT: %{{.*}} = "arith.constant"() <{value = 1 : i32}> : () -> i32 + // CHECK-NEXT: "symref.update"(%{{.*}}) {symbol = @b} : (i32) -> () + // CHECK-NEXT: %{{.*}} = "arith.constant"() <{value = 2 : i32}> : () -> i32 + // CHECK-NEXT: %{{.*}} = "arith.addi"(%{{.*}}, %{{.*}}) <{overflowFlags = #arith.overflow}> : (i32, i32) -> i32 + // CHECK-NEXT: %{{.*}} = "arith.muli"(%{{.*}}, %{{.*}}) <{overflowFlags = #arith.overflow}> : (i32, i32) -> i32 + // CHECK-NEXT: "symref.update"(%{{.*}}) {symbol = @a} : (i32) -> () + // CHECK-NEXT: %{{.*}} = "arith.addi"(%{{.*}}, %{{.*}}) <{overflowFlags = #arith.overflow}> : (i32, i32) -> i32 + // CHECK-NEXT: "symref.update"(%{{.*}}) {symbol = @c} : (i32) -> () // CHECK-NEXT: }) : () -> () }) : () -> () diff --git a/tests/filecheck/mlir-conversion/with-mlir/affinet_set.mlir b/tests/filecheck/mlir-conversion/with-mlir/affinet_set.mlir index f34044f5d4..c396a3a5e3 100644 --- a/tests/filecheck/mlir-conversion/with-mlir/affinet_set.mlir +++ b/tests/filecheck/mlir-conversion/with-mlir/affinet_set.mlir @@ -1,19 +1,19 @@ // RUN: xdsl-opt --allow-unregistered-dialect %s | mlir-opt --allow-unregistered-dialect --mlir-print-local-scope -mlir-print-op-generic | xdsl-opt --allow-unregistered-dialect | filecheck %s builtin.module { - // CHECK: "f1"() {"set" = affine_set<(d0) : ((d0 + -10) >= 0)>} : () -> () + // CHECK: "f1"() {set = affine_set<(d0) : ((d0 + -10) >= 0)>} : () -> () "f1"() {set = affine_set<(d0): ((d0 - 10) >= 0)>} : () -> () - // CHECK-NEXT: "f2"() {"set" = affine_set<(d0)[s0] : (d0 >= 0, ((d0 * -1) + s0) >= 0)>} : () -> () + // CHECK-NEXT: "f2"() {set = affine_set<(d0)[s0] : (d0 >= 0, ((d0 * -1) + s0) >= 0)>} : () -> () "f2"() {set = affine_set<(i)[N] : (i >= 0, N - i >= 0)>} : () -> () - // CHECK-NEXT: "f3"() {"set" = affine_set<(d0) : (1 == 0)>} : () -> () + // CHECK-NEXT: "f3"() {set = affine_set<(d0) : (1 == 0)>} : () -> () "f3"() {set = affine_set<(d0) : (1 == 0)>} : () -> () - // CHECK-NEXT: "f4"() {"set" = affine_set<(d0, d1)[s0, s1] : ((((d0 * -16) + s0) + -16) >= 0, (((d1 * -3) + s1) + -3) >= 0)>} : () -> () + // CHECK-NEXT: "f4"() {set = affine_set<(d0, d1)[s0, s1] : ((((d0 * -16) + s0) + -16) >= 0, (((d1 * -3) + s1) + -3) >= 0)>} : () -> () "f4"() {set = affine_set<(d0, d1)[s0, s1] : (d0 * -16 + s0 - 16 >= 0, d1 * -3 + s1 - 3 >= 0)>} : () -> () - // CHECK-NEXT: "f5"() {"set" = affine_set<(d0, d1)[s0, s1] : (((((d0 * 7) + (d1 * 5)) + (s0 * 11)) + s1) == 0, ((((d0 * 5) + (d1 * -11)) + (s0 * 7)) + s1) == 0, ((((d0 * 11) + (d1 * 7)) + (s0 * -5)) + s1) == 0, ((((d0 * 7) + (d1 * 5)) + (s0 * 11)) + s1) == 0)>} : () -> () + // CHECK-NEXT: "f5"() {set = affine_set<(d0, d1)[s0, s1] : (((((d0 * 7) + (d1 * 5)) + (s0 * 11)) + s1) == 0, ((((d0 * 5) + (d1 * -11)) + (s0 * 7)) + s1) == 0, ((((d0 * 11) + (d1 * 7)) + (s0 * -5)) + s1) == 0, ((((d0 * 7) + (d1 * 5)) + (s0 * 11)) + s1) == 0)>} : () -> () "f5"() {set = affine_set<(d0, d1)[s0, s1] : (d0 * 7 + d1 * 5 + s0 * 11 + s1 == 0, d0 * 5 - d1 * 11 + s0 * 7 + s1 == 0, d0 * 11 + d1 * 7 - s0 * 5 + s1 == 0, diff --git a/tests/filecheck/mlir-conversion/with-mlir/dialects/affine/affine_ops.mlir b/tests/filecheck/mlir-conversion/with-mlir/dialects/affine/affine_ops.mlir index aaf929c9bd..1cdc0f563e 100644 --- a/tests/filecheck/mlir-conversion/with-mlir/dialects/affine/affine_ops.mlir +++ b/tests/filecheck/mlir-conversion/with-mlir/dialects/affine/affine_ops.mlir @@ -9,7 +9,7 @@ "affine.yield"() : () -> () }) : () -> () - // CHECK: "affine.for"() <{"lowerBoundMap" = affine_map<() -> (0)>, "operandSegmentSizes" = array, "step" = 1 : index, "upperBoundMap" = affine_map<() -> (256)>}> ({ + // CHECK: "affine.for"() <{lowerBoundMap = affine_map<() -> (0)>, operandSegmentSizes = array, step = 1 : index, upperBoundMap = affine_map<() -> (256)>}> ({ // CHECK-NEXT: ^0(%{{.*}} : index): // CHECK-NEXT: "affine.yield"() : () -> () // CHECK-NEXT: }) : () -> () @@ -35,12 +35,12 @@ "affine.yield"() : () -> () }) : (index) -> () - // CHECK: %{{.*}} = "affine.for"(%{{.*}}) <{"lowerBoundMap" = affine_map<() -> (-10)>, "operandSegmentSizes" = array, "step" = 1 : index, "upperBoundMap" = affine_map<() -> (10)>}> ({ + // CHECK: %{{.*}} = "affine.for"(%{{.*}}) <{lowerBoundMap = affine_map<() -> (-10)>, operandSegmentSizes = array, step = 1 : index, upperBoundMap = affine_map<() -> (10)>}> ({ // CHECK-NEXT: ^1(%{{.*}} : index, %{{.*}} : i32): // CHECK-NEXT: %{{.*}} = "test.op"() : () -> i32 // CHECK-NEXT: "affine.yield"(%{{.*}}) : (i32) -> () // CHECK-NEXT: }) : (i32) -> i32 - // CHECK: "affine.parallel"(%{{.*}}) <{"lowerBoundsGroups" = dense<1> : vector<1xi32>, "lowerBoundsMap" = affine_map<() -> (0)>, "reductions" = [], "steps" = [1 : i64], "upperBoundsGroups" = dense<1> : vector<1xi32>, "upperBoundsMap" = affine_map<()[s0] -> (s0)>}> ({ + // CHECK: "affine.parallel"(%{{.*}}) <{lowerBoundsGroups = dense<1> : vector<1xi32>, lowerBoundsMap = affine_map<() -> (0)>, reductions = [], steps = [1 : i64], upperBoundsGroups = dense<1> : vector<1xi32>, upperBoundsMap = affine_map<()[s0] -> (s0)>}> ({ // CHECK-NEXT: ^{{.*}}(%{{.*}} : index): // CHECK-NEXT: "affine.yield"() : () -> () // CHECK-NEXT: }) : (index) -> () @@ -52,7 +52,7 @@ // CHECK: %{{.*}} = "test.op"() : () -> memref<2x3xf64> // CHECK-NEXT: %{{.*}} = "test.op"() : () -> f64 - // CHECK-NEXT: "affine.store"(%{{.*}}, %{{.*}}) <{"map" = affine_map<() -> (0, 0)>}> : (f64, memref<2x3xf64>) -> () + // CHECK-NEXT: "affine.store"(%{{.*}}, %{{.*}}) <{map = affine_map<() -> (0, 0)>}> : (f64, memref<2x3xf64>) -> () %zero = "test.op"() : () -> index %2 = affine.apply affine_map<(d0)[s0] -> (((d0 + (s0 * 42)) + -1))> (%zero)[%zero] @@ -61,8 +61,8 @@ // CHECK: %{{.*}} = "test.op"() : () -> index // CHECK-NEXT: %{{.*}} = affine.apply affine_map<(d0)[s0] -> (((d0 + (s0 * 42)) + -1))> (%{{.*}})[%{{.*}}] - // CHECK-NEXT: %{{.*}} = "affine.min"(%{{.*}}) <{"map" = affine_map<(d0) -> ((d0 + 41), d0)>}> : (index) -> index - // CHECK-NEXT: %{{.*}} = "affine.load"(%{{.*}}, %{{.*}}, %{{.*}}) <{"map" = affine_map<(d0, d1) -> (d0, d1)>}> : (memref<2x3xf64>, index, index) -> f64 + // CHECK-NEXT: %{{.*}} = "affine.min"(%{{.*}}) <{map = affine_map<(d0) -> ((d0 + 41), d0)>}> : (index) -> index + // CHECK-NEXT: %{{.*}} = "affine.load"(%{{.*}}, %{{.*}}, %{{.*}}) <{map = affine_map<(d0, d1) -> (d0, d1)>}> : (memref<2x3xf64>, index, index) -> f64 func.func @empty() { "affine.for"() <{"lowerBoundMap" = affine_map<() -> (0)>, "step" = 1 : index, "upperBoundMap" = affine_map<() -> (10)>, "operandSegmentSizes" = array}> ({ @@ -82,19 +82,19 @@ func.return } // CHECK: func.func @empty() { -// CHECK-NEXT: "affine.for"() <{"lowerBoundMap" = affine_map<() -> (0)>, "operandSegmentSizes" = array, "step" = 1 : index, "upperBoundMap" = affine_map<() -> (10)>}> ({ +// CHECK-NEXT: "affine.for"() <{lowerBoundMap = affine_map<() -> (0)>, operandSegmentSizes = array, step = 1 : index, upperBoundMap = affine_map<() -> (10)>}> ({ // CHECK-NEXT: ^{{.*}}(%{{.*}} : index): // CHECK-NEXT: "affine.yield"() : () -> () // CHECK-NEXT: }) : () -> () // CHECK-NEXT: "affine.if"() ({ // CHECK-NEXT: "affine.yield"() : () -> () // CHECK-NEXT: }, { -// CHECK-NEXT: }) {"condition" = affine_set<() : (0 == 0)>} : () -> () +// CHECK-NEXT: }) {condition = affine_set<() : (0 == 0)>} : () -> () // CHECK-NEXT: "affine.if"() ({ // CHECK-NEXT: "affine.yield"() : () -> () // CHECK-NEXT: }, { // CHECK-NEXT: "affine.yield"() : () -> () -// CHECK-NEXT: }) {"condition" = affine_set<() : (0 == 0)>} : () -> () +// CHECK-NEXT: }) {condition = affine_set<() : (0 == 0)>} : () -> () // CHECK-NEXT: func.return // CHECK-NEXT: } @@ -113,7 +113,7 @@ // CHECK-NEXT: "affine.yield"(%{{.*}}) : (f32) -> () // CHECK-NEXT: }, { // CHECK-NEXT: "affine.yield"(%{{.*}}) : (f32) -> () -// CHECK-NEXT: }) {"condition" = affine_set<() : (0 == 0)>} : () -> f32 +// CHECK-NEXT: }) {condition = affine_set<() : (0 == 0)>} : () -> f32 // CHECK-NEXT: func.return %{{.*}} : f32 // CHECK-NEXT: } diff --git a/tests/filecheck/mlir-conversion/with-mlir/dialects/arith/arith_attrs.mlir b/tests/filecheck/mlir-conversion/with-mlir/dialects/arith/arith_attrs.mlir index a51c7bf82a..3fbfeb326d 100644 --- a/tests/filecheck/mlir-conversion/with-mlir/dialects/arith/arith_attrs.mlir +++ b/tests/filecheck/mlir-conversion/with-mlir/dialects/arith/arith_attrs.mlir @@ -4,28 +4,28 @@ // fastmath single flag set "test.op"() {"fastmath" = #arith.fastmath}: ()->() -// CHECK: "fastmath" = #arith.fastmath +// CHECK: fastmath = #arith.fastmath "test.op"() {"fastmath" = #arith.fastmath}: ()->() -// CHECK: "fastmath" = #arith.fastmath +// CHECK: fastmath = #arith.fastmath "test.op"() {"fastmath" = #arith.fastmath}: ()->() -// CHECK: "fastmath" = #arith.fastmath +// CHECK: fastmath = #arith.fastmath "test.op"() {"fastmath" = #arith.fastmath}: ()->() -// CHECK: "fastmath" = #arith.fastmath +// CHECK: fastmath = #arith.fastmath "test.op"() {"fastmath" = #arith.fastmath}: ()->() -// CHECK: "fastmath" = #arith.fastmath +// CHECK: fastmath = #arith.fastmath "test.op"() {"fastmath" = #arith.fastmath}: ()->() -// CHECK: "fastmath" = #arith.fastmath +// CHECK: fastmath = #arith.fastmath "test.op"() {"fastmath" = #arith.fastmath}: ()->() -// CHECK: "fastmath" = #arith.fastmath +// CHECK: fastmath = #arith.fastmath // fastmath special cases "test.op"() {"fastmath" = #arith.fastmath}: ()->() -// CHECK: "fastmath" = #arith.fastmath +// CHECK: fastmath = #arith.fastmath "test.op"() {"fastmath" = #arith.fastmath}: ()->() -// CHECK: "fastmath" = #arith.fastmath +// CHECK: fastmath = #arith.fastmath // fastmath combination "test.op"() {"fastmath" = #arith.fastmath}: ()->() -// CHECK: "fastmath" = #arith.fastmath +// CHECK: fastmath = #arith.fastmath }) : ()->() diff --git a/tests/filecheck/mlir-conversion/with-mlir/dialects/arith/arith_cmp.mlir b/tests/filecheck/mlir-conversion/with-mlir/dialects/arith/arith_cmp.mlir index 18e56d36b2..4e029dd147 100644 --- a/tests/filecheck/mlir-conversion/with-mlir/dialects/arith/arith_cmp.mlir +++ b/tests/filecheck/mlir-conversion/with-mlir/dialects/arith/arith_cmp.mlir @@ -15,9 +15,9 @@ %11 = "arith.select"(%10, %4, %5) : (i1, index, index) -> index }) : ()->() -// CHECK: "arith.cmpf"(%0, %1) <{"fastmath" = #arith.fastmath, "predicate" = 2 : i64}> : (f64, f64) -> i1 +// CHECK: "arith.cmpf"(%0, %1) <{fastmath = #arith.fastmath, predicate = 2 : i64}> : (f64, f64) -> i1 // CHECK: "arith.select"(%6, %0, %1) : (i1, f64, f64) -> f64 -// CHECK: "arith.cmpi"(%2, %3) <{"predicate" = 1 : i64}> : (i32, i32) -> i1 +// CHECK: "arith.cmpi"(%2, %3) <{predicate = 1 : i64}> : (i32, i32) -> i1 // CHECK: "arith.select"(%8, %2, %3) : (i1, i32, i32) -> i32 -// CHECK: "arith.cmpi"(%4, %5) <{"predicate" = 1 : i64}> : (index, index) -> i1 +// CHECK: "arith.cmpi"(%4, %5) <{predicate = 1 : i64}> : (index, index) -> i1 // CHECK: "arith.select"(%10, %4, %5) : (i1, index, index) -> index diff --git a/tests/filecheck/mlir-conversion/with-mlir/dialects/arith/arith_fp_ops.mlir b/tests/filecheck/mlir-conversion/with-mlir/dialects/arith/arith_fp_ops.mlir index 240edd7554..21d175af51 100644 --- a/tests/filecheck/mlir-conversion/with-mlir/dialects/arith/arith_fp_ops.mlir +++ b/tests/filecheck/mlir-conversion/with-mlir/dialects/arith/arith_fp_ops.mlir @@ -7,5 +7,5 @@ %4 = "arith.negf"(%1) : (f64) -> f64 }) : ()->() -// CHECK: "arith.negf"(%0) <{"fastmath" = #arith.fastmath}> : (f32) -> f32 -// CHECK: "arith.negf"(%1) <{"fastmath" = #arith.fastmath}> : (f64) -> f64 +// CHECK: "arith.negf"(%0) <{fastmath = #arith.fastmath}> : (f32) -> f32 +// CHECK: "arith.negf"(%1) <{fastmath = #arith.fastmath}> : (f64) -> f64 diff --git a/tests/filecheck/mlir-conversion/with-mlir/dialects/builtin/unrealized_conversion_cast.mlir b/tests/filecheck/mlir-conversion/with-mlir/dialects/builtin/unrealized_conversion_cast.mlir index 0d202f8712..0f18c3467c 100644 --- a/tests/filecheck/mlir-conversion/with-mlir/dialects/builtin/unrealized_conversion_cast.mlir +++ b/tests/filecheck/mlir-conversion/with-mlir/dialects/builtin/unrealized_conversion_cast.mlir @@ -11,10 +11,10 @@ }) : () -> () // CHECK: "builtin.module"() ({ -// CHECK-NEXT: %0 = "arith.constant"() <{"value" = 0 : i64}> : () -> i64 +// CHECK-NEXT: %0 = "arith.constant"() <{value = 0 : i64}> : () -> i64 // CHECK-NEXT: %1 = "builtin.unrealized_conversion_cast"(%0) : (i64) -> i32 // CHECK-NEXT: %2 = "builtin.unrealized_conversion_cast"(%0) : (i64) -> f32 -// CHECK-NEXT: %3 = "arith.constant"() <{"value" = 1.040000e+01 : f32}> : () -> f32 +// CHECK-NEXT: %3 = "arith.constant"() <{value = 1.040000e+01 : f32}> : () -> f32 // CHECK-NEXT: %4 = "builtin.unrealized_conversion_cast"(%3) : (f32) -> f64 // CHECK-NEXT: %5 = "builtin.unrealized_conversion_cast"(%3) : (f32) -> i64 // CHECK-NEXT: %6 = "builtin.unrealized_conversion_cast"(%3) : (f32) -> f32 diff --git a/tests/filecheck/mlir-conversion/with-mlir/dialects/func/func_ops.mlir b/tests/filecheck/mlir-conversion/with-mlir/dialects/func/func_ops.mlir index c916c0bc72..ab5ce64ac6 100644 --- a/tests/filecheck/mlir-conversion/with-mlir/dialects/func/func_ops.mlir +++ b/tests/filecheck/mlir-conversion/with-mlir/dialects/func/func_ops.mlir @@ -27,7 +27,7 @@ builtin.module { } // CHECK: func.func @call_void_attributes() { - // CHECK-NEXT: func.call @call_void_attributes() {"hello" = "world"} : () -> () + // CHECK-NEXT: func.call @call_void_attributes() {hello = "world"} : () -> () // CHECK-NEXT: func.return // CHECK-NEXT: } @@ -69,7 +69,7 @@ builtin.module { func.return %X : tensor<8x8xf64> } - // CHECK: func.func public @arg_attrs(%{{.*}}: tensor<8x8xf64> {"llvm.noalias"}, %{{.*}}: tensor<8x8xf64> {"llvm.noalias"}, %{{.*}}: tensor<8x8xf64> {"llvm.noalias"}) -> tensor<8x8xf64> { + // CHECK: func.func public @arg_attrs(%{{.*}}: tensor<8x8xf64> {llvm.noalias}, %{{.*}}: tensor<8x8xf64> {llvm.noalias}, %{{.*}}: tensor<8x8xf64> {llvm.noalias}) -> tensor<8x8xf64> { // CHECK-NEXT: func.return %{{.*}} : tensor<8x8xf64> // CHECK-NEXT: } @@ -78,7 +78,7 @@ builtin.module { return %r1, %r2 : f32, f32 } - // CHECK: func.func @output_attributes() -> (f32 {"dialect.a" = 0 : i32}, f32 {"dialect.b" = 0 : i32, "dialect.c" = 1 : i64}) { + // CHECK: func.func @output_attributes() -> (f32 {dialect.a = 0 : i32}, f32 {dialect.b = 0 : i32, dialect.c = 1 : i64}) { // CHECK-NEXT: %0, %1 = "test.op"() : () -> (f32, f32) // CHECK-NEXT: func.return %0, %1 : f32, f32 // CHECK-NEXT: } diff --git a/tests/filecheck/mlir-conversion/with-mlir/dialects/func/func_ops_generic.mlir b/tests/filecheck/mlir-conversion/with-mlir/dialects/func/func_ops_generic.mlir index c7c49c3e65..c3bb066dec 100644 --- a/tests/filecheck/mlir-conversion/with-mlir/dialects/func/func_ops_generic.mlir +++ b/tests/filecheck/mlir-conversion/with-mlir/dialects/func/func_ops_generic.mlir @@ -6,7 +6,7 @@ }) : () -> () -// CHECK: "func.func"() <{"function_type" = (tensor<8x8xf64>, tensor<8x8xf64>) -> (tensor<8x8xf64>, tensor<8x8xf64>), "res_attrs" = [{"llvm.noalias"}, {"llvm.noalias"}], "sym_name" = "arg_attrs", "sym_visibility" = "public"}> ({ +// CHECK: "func.func"() <{function_type = (tensor<8x8xf64>, tensor<8x8xf64>) -> (tensor<8x8xf64>, tensor<8x8xf64>), res_attrs = [{llvm.noalias}, {llvm.noalias}], sym_name = "arg_attrs", sym_visibility = "public"}> ({ // CHECK-NEXT: ^0(%arg0 : tensor<8x8xf64>, %arg1 : tensor<8x8xf64>): // CHECK-NEXT: "func.return"(%arg0, %arg1) : (tensor<8x8xf64>, tensor<8x8xf64>) -> () // CHECK-NEXT: }) : () -> () diff --git a/tests/filecheck/mlir-conversion/with-mlir/dialects/gpu/ops.mlir b/tests/filecheck/mlir-conversion/with-mlir/dialects/gpu/ops.mlir index 098e1cd0f1..7a24484ca9 100644 --- a/tests/filecheck/mlir-conversion/with-mlir/dialects/gpu/ops.mlir +++ b/tests/filecheck/mlir-conversion/with-mlir/dialects/gpu/ops.mlir @@ -81,81 +81,81 @@ }) {"gpu.container_module"} : () -> () // CHECK: "builtin.module"() ({ -// CHECK-NEXT: "gpu.module"() <{"sym_name" = "gpu"}> ({ -// CHECK-NEXT: "func.func"() <{"function_type" = () -> (), "sym_name" = "kernel"}> ({ -// CHECK-NEXT: %{{.*}} = "arith.constant"() <{"value" = 13 : index}> : () -> index -// CHECK-NEXT: %{{.*}} = "arith.constant"() <{"value" = 1 : index}> {"loopdim" = #gpu.loop_dim_map (d0), bound = (d0) -> (d0)>} : () -> index -// CHECK-NEXT: %{{.*}} = "memref.alloc"() <{"alignment" = 0 : i64, "operandSegmentSizes" = array}> : () -> memref<10x10xi32> +// CHECK-NEXT: "gpu.module"() <{sym_name = "gpu"}> ({ +// CHECK-NEXT: "func.func"() <{function_type = () -> (), sym_name = "kernel"}> ({ +// CHECK-NEXT: %{{.*}} = "arith.constant"() <{value = 13 : index}> : () -> index +// CHECK-NEXT: %{{.*}} = "arith.constant"() <{value = 1 : index}> {loopdim = #gpu.loop_dim_map (d0), bound = (d0) -> (d0)>} : () -> index +// CHECK-NEXT: %{{.*}} = "memref.alloc"() <{alignment = 0 : i64, operandSegmentSizes = array}> : () -> memref<10x10xi32> // CHECK-NEXT: %{{.*}} = "memref.cast"(%{{.*}}) : (memref<10x10xi32>) -> memref<*xi32> // CHECK-NEXT: "gpu.host_register"(%{{.*}}) : (memref<*xi32>) -> () // CHECK-NEXT: "gpu.host_unregister"(%{{.*}}) : (memref<*xi32>) -> () // CHECK-NEXT: %{{.*}} = "gpu.wait"() : () -> !gpu.async.token -// CHECK-NEXT: %{{.*}} = "gpu.thread_id"() <{"dimension" = #gpu}> : () -> index -// CHECK-NEXT: %{{.*}} = "gpu.thread_id"() <{"dimension" = #gpu}> : () -> index -// CHECK-NEXT: %{{.*}} = "gpu.thread_id"() <{"dimension" = #gpu}> : () -> index +// CHECK-NEXT: %{{.*}} = "gpu.thread_id"() <{dimension = #gpu}> : () -> index +// CHECK-NEXT: %{{.*}} = "gpu.thread_id"() <{dimension = #gpu}> : () -> index +// CHECK-NEXT: %{{.*}} = "gpu.thread_id"() <{dimension = #gpu}> : () -> index -// CHECK-NEXT: %{{.*}} = "gpu.block_dim"() <{"dimension" = #gpu}> : () -> index -// CHECK-NEXT: %{{.*}} = "gpu.block_dim"() <{"dimension" = #gpu}> : () -> index -// CHECK-NEXT: %{{.*}} = "gpu.block_dim"() <{"dimension" = #gpu}> : () -> index +// CHECK-NEXT: %{{.*}} = "gpu.block_dim"() <{dimension = #gpu}> : () -> index +// CHECK-NEXT: %{{.*}} = "gpu.block_dim"() <{dimension = #gpu}> : () -> index +// CHECK-NEXT: %{{.*}} = "gpu.block_dim"() <{dimension = #gpu}> : () -> index -// CHECK-NEXT: %{{.*}} = "gpu.block_id"() <{"dimension" = #gpu}> : () -> index -// CHECK-NEXT: %{{.*}} = "gpu.block_id"() <{"dimension" = #gpu}> : () -> index -// CHECK-NEXT: %{{.*}} = "gpu.block_id"() <{"dimension" = #gpu}> : () -> index +// CHECK-NEXT: %{{.*}} = "gpu.block_id"() <{dimension = #gpu}> : () -> index +// CHECK-NEXT: %{{.*}} = "gpu.block_id"() <{dimension = #gpu}> : () -> index +// CHECK-NEXT: %{{.*}} = "gpu.block_id"() <{dimension = #gpu}> : () -> index -// CHECK-NEXT: %{{.*}} = "gpu.global_id"() <{"dimension" = #gpu}> : () -> index -// CHECK-NEXT: %{{.*}} = "gpu.global_id"() <{"dimension" = #gpu}> : () -> index -// CHECK-NEXT: %{{.*}} = "gpu.global_id"() <{"dimension" = #gpu}> : () -> index +// CHECK-NEXT: %{{.*}} = "gpu.global_id"() <{dimension = #gpu}> : () -> index +// CHECK-NEXT: %{{.*}} = "gpu.global_id"() <{dimension = #gpu}> : () -> index +// CHECK-NEXT: %{{.*}} = "gpu.global_id"() <{dimension = #gpu}> : () -> index -// CHECK-NEXT: %{{.*}} = "gpu.grid_dim"() <{"dimension" = #gpu}> : () -> index -// CHECK-NEXT: %{{.*}} = "gpu.grid_dim"() <{"dimension" = #gpu}> : () -> index -// CHECK-NEXT: %{{.*}} = "gpu.grid_dim"() <{"dimension" = #gpu}> : () -> index +// CHECK-NEXT: %{{.*}} = "gpu.grid_dim"() <{dimension = #gpu}> : () -> index +// CHECK-NEXT: %{{.*}} = "gpu.grid_dim"() <{dimension = #gpu}> : () -> index +// CHECK-NEXT: %{{.*}} = "gpu.grid_dim"() <{dimension = #gpu}> : () -> index -// CHECK-NEXT: %{{.*}} = "gpu.alloc"() <{"operandSegmentSizes" = array}> : () -> memref<10x10xi32> -// CHECK-NEXT: %{{.*}} = "gpu.alloc"(%{{.*}}, %{{.*}}, %{{.*}}) <{"operandSegmentSizes" = array}> : (index, index, index) -> memref +// CHECK-NEXT: %{{.*}} = "gpu.alloc"() <{operandSegmentSizes = array}> : () -> memref<10x10xi32> +// CHECK-NEXT: %{{.*}} = "gpu.alloc"(%{{.*}}, %{{.*}}, %{{.*}}) <{operandSegmentSizes = array}> : (index, index, index) -> memref -// CHECK-NEXT: "gpu.memcpy"(%{{.*}}, %{{.*}}) {"operandSegmentSizes" = array} : (memref<10x10xi32>, memref<10x10xi32>) -> () +// CHECK-NEXT: "gpu.memcpy"(%{{.*}}, %{{.*}}) {operandSegmentSizes = array} : (memref<10x10xi32>, memref<10x10xi32>) -> () -// CHECK-NEXT: "gpu.dealloc"(%{{.*}}) {"operandSegmentSizes" = array} : (memref) -> () +// CHECK-NEXT: "gpu.dealloc"(%{{.*}}) {operandSegmentSizes = array} : (memref) -> () // CHECK-NEXT: %{{.*}} = "gpu.lane_id"() : () -> index // CHECK-NEXT: %{{.*}} = "gpu.num_subgroups"() : () -> index -// CHECK-NEXT: %{{.*}} = "arith.constant"() <{"value" = 0 : i32}> : () -> i32 +// CHECK-NEXT: %{{.*}} = "arith.constant"() <{value = 0 : i32}> : () -> i32 // CHECK-NEXT: "gpu.set_default_device"(%{{.*}}) : (i32) -> () // CHECK-NEXT: %{{.*}} = "gpu.subgroup_id"() : () -> index // CHECK-NEXT: %{{.*}} = "gpu.subgroup_size"() : () -> index -// CHECK-NEXT: %{{.*}} = "gpu.all_reduce"(%{{.*}}) <{"op" = #gpu}> ({ +// CHECK-NEXT: %{{.*}} = "gpu.all_reduce"(%{{.*}}) <{op = #gpu}> ({ // CHECK-NEXT: }) : (i32) -> i32 // CHECK-NEXT: %{{.*}} = "gpu.all_reduce"(%{{.*}}) ({ // CHECK-NEXT: ^{{.*}}(%{{.*}} : i32, %{{.*}} : i32): -// CHECK-NEXT: %{{.*}} = "arith.addi"(%{{.*}}, %{{.*}}) <{"overflowFlags" = #arith.overflow}> : (i32, i32) -> i32 +// CHECK-NEXT: %{{.*}} = "arith.addi"(%{{.*}}, %{{.*}}) <{overflowFlags = #arith.overflow}> : (i32, i32) -> i32 // CHECK-NEXT: "gpu.yield"(%{{.*}}) : (i32) -> () // CHECK-NEXT: }) : (i32) -> i32 -// CHECK-NEXT: "gpu.launch"(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "gpu.launch"(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) <{operandSegmentSizes = array}> ({ // CHECK-NEXT: ^{{\S+}}(%{{\S+}} : index, %{{\S+}} : index, %{{\S+}} : index, // CHECK-SAME: %{{\S+}} : index, %{{\S+}} : index, %{{\S+}} : index, // CHECK-SAME: %{{\S+}} : index, %{{\S+}} : index, %{{\S+}} : index, // CHECK-SAME: %{{\S+}} : index, %{{\S+}} : index, %{{\S+}} : index): -// CHECK-NEXT: %{{.*}} = "gpu.all_reduce"(%{{.*}}) <{"op" = #gpu}> ({ +// CHECK-NEXT: %{{.*}} = "gpu.all_reduce"(%{{.*}}) <{op = #gpu}> ({ // CHECK-NEXT: }) : (i32) -> i32 -// CHECK-NEXT: %{{.*}} = "arith.muli"(%{{.*}}, %{{.*}}) <{"overflowFlags" = #arith.overflow}> : (i32, i32) -> i32 +// CHECK-NEXT: %{{.*}} = "arith.muli"(%{{.*}}, %{{.*}}) <{overflowFlags = #arith.overflow}> : (i32, i32) -> i32 // CHECK-NEXT: "gpu.terminator"() : () -> () // CHECK-NEXT: }) : (index, index, index, index, index, index) -> () -// CHECK-NEXT: "gpu.launch_func"(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) <{"kernel" = @gpu::@foo, "operandSegmentSizes" = array}> : (index, index, index, index, index, index, index, index, index, i32, index) -> () +// CHECK-NEXT: "gpu.launch_func"(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) <{kernel = @gpu::@foo, operandSegmentSizes = array}> : (index, index, index, index, index, index, index, index, index, i32, index) -> () // CHECK-NEXT: "func.return"() : () -> () // CHECK-NEXT: }) : () -> () -// CHECK-NEXT: "gpu.func"() <{"function_type" = (index) -> ()}> ({ +// CHECK-NEXT: "gpu.func"() <{function_type = (index) -> ()}> ({ // CHECK-NEXT: ^{{.*}}(%{{.*}}: index): // CHECK-NEXT: "gpu.return"() : () -> () -// CHECK-NEXT: }) {"gpu.known_block_size" = array, "gpu.known_grid_size" = array, "sym_name" = "foo"} : () -> () +// CHECK-NEXT: }) {gpu.known_block_size = array, gpu.known_grid_size = array, sym_name = "foo"} : () -> () // CHECK-NEXT: "gpu.module_end"() : () -> () // CHECK-NEXT: }) : () -> () -// CHECK-NEXT: }) {"gpu.container_module"} : () -> () +// CHECK-NEXT: }) {gpu.container_module} : () -> () diff --git a/tests/filecheck/mlir-conversion/with-mlir/dialects/llvm/arithmetic.mlir b/tests/filecheck/mlir-conversion/with-mlir/dialects/llvm/arithmetic.mlir index f60f1dd57c..e757fc08ac 100644 --- a/tests/filecheck/mlir-conversion/with-mlir/dialects/llvm/arithmetic.mlir +++ b/tests/filecheck/mlir-conversion/with-mlir/dialects/llvm/arithmetic.mlir @@ -7,7 +7,7 @@ builtin.module { // CHECK: %{{.*}} = llvm.add %{{.*}}, %{{.*}} : i32 %add2 = llvm.add %arg0, %arg1 {"nsw"} : i32 - // CHECK: %{{.*}} = llvm.add %{{.*}}, %{{.*}} {"nsw"} : i32 + // CHECK: %{{.*}} = llvm.add %{{.*}}, %{{.*}} {nsw} : i32 %sub = llvm.sub %arg0, %arg1 : i32 // CHECK: %{{.*}} = llvm.sub %{{.*}}, %{{.*}} : i32 diff --git a/tests/filecheck/mlir-conversion/with-mlir/dialects/llvm/global.mlir b/tests/filecheck/mlir-conversion/with-mlir/dialects/llvm/global.mlir index c9b7a8c8bd..6fb5f5fbdb 100644 --- a/tests/filecheck/mlir-conversion/with-mlir/dialects/llvm/global.mlir +++ b/tests/filecheck/mlir-conversion/with-mlir/dialects/llvm/global.mlir @@ -9,10 +9,10 @@ }) : () -> () // CHECK: "builtin.module"() ({ -// CHECK-NEXT: "llvm.mlir.global"() <{"addr_space" = 0 : i32, "constant", "global_type" = !llvm.array<12 x i8>, "linkage" = #llvm.linkage<"internal">, "sym_name" = "str0", "unnamed_addr" = 0 : i64, "value" = "Hello world!", "visibility_" = 0 : i64}> ({ +// CHECK-NEXT: "llvm.mlir.global"() <{addr_space = 0 : i32, constant, global_type = !llvm.array<12 x i8>, linkage = #llvm.linkage<"internal">, sym_name = "str0", unnamed_addr = 0 : i64, value = "Hello world!", visibility_ = 0 : i64}> ({ // CHECK-NEXT: }) : () -> () -// CHECK-NEXT: "llvm.mlir.global"() <{"addr_space" = 0 : i32, "constant", "global_type" = i32, "linkage" = #llvm.linkage<"internal">, "sym_name" = "data", "unnamed_addr" = 0 : i64, "value" = 0 : i32, "visibility_" = 0 : i64}> ({ +// CHECK-NEXT: "llvm.mlir.global"() <{addr_space = 0 : i32, constant, global_type = i32, linkage = #llvm.linkage<"internal">, sym_name = "data", unnamed_addr = 0 : i64, value = 0 : i32, visibility_ = 0 : i64}> ({ // CHECK-NEXT: }) : () -> () -// CHECK-NEXT: %0 = "llvm.mlir.addressof"() <{"global_name" = @data}> : () -> !llvm.ptr +// CHECK-NEXT: %0 = "llvm.mlir.addressof"() <{global_name = @data}> : () -> !llvm.ptr // CHECK-NEXT: %1 = "llvm.mlir.zero"() : () -> !llvm.struct<(i32, f32)> // CHECK-NEXT: }) : () -> () diff --git a/tests/filecheck/mlir-conversion/with-mlir/dialects/llvm/icmp.mlir b/tests/filecheck/mlir-conversion/with-mlir/dialects/llvm/icmp.mlir index 85982700d7..14d71c4821 100644 --- a/tests/filecheck/mlir-conversion/with-mlir/dialects/llvm/icmp.mlir +++ b/tests/filecheck/mlir-conversion/with-mlir/dialects/llvm/icmp.mlir @@ -4,50 +4,50 @@ %2 = "llvm.icmp"(%0, %1) <{predicate = 0 : i64}> : (i32, i32) -> i1 %3 = llvm.icmp "eq" %0, %1 : i32 -// CHECK: %2 = "llvm.icmp"(%0, %1) <{"predicate" = 0 : i64}> : (i32, i32) -> i1 -// CHECK-NEXT: %3 = "llvm.icmp"(%0, %1) <{"predicate" = 0 : i64}> : (i32, i32) -> i1 +// CHECK: %2 = "llvm.icmp"(%0, %1) <{predicate = 0 : i64}> : (i32, i32) -> i1 +// CHECK-NEXT: %3 = "llvm.icmp"(%0, %1) <{predicate = 0 : i64}> : (i32, i32) -> i1 %4 = "llvm.icmp"(%0, %1) <{predicate = 1 : i64}> : (i32, i32) -> i1 %5 = llvm.icmp "ne" %0, %1 : i32 -// CHECK-NEXT: %4 = "llvm.icmp"(%0, %1) <{"predicate" = 1 : i64}> : (i32, i32) -> i1 -// CHECK-NEXT: %5 = "llvm.icmp"(%0, %1) <{"predicate" = 1 : i64}> : (i32, i32) -> i1 +// CHECK-NEXT: %4 = "llvm.icmp"(%0, %1) <{predicate = 1 : i64}> : (i32, i32) -> i1 +// CHECK-NEXT: %5 = "llvm.icmp"(%0, %1) <{predicate = 1 : i64}> : (i32, i32) -> i1 %6 = "llvm.icmp"(%0, %1) <{predicate = 2 : i64}> : (i32, i32) -> i1 %7 = llvm.icmp "slt" %0, %1 : i32 -// CHECK-NEXT: %6 = "llvm.icmp"(%0, %1) <{"predicate" = 2 : i64}> : (i32, i32) -> i1 -// CHECK-NEXT: %7 = "llvm.icmp"(%0, %1) <{"predicate" = 2 : i64}> : (i32, i32) -> i1 +// CHECK-NEXT: %6 = "llvm.icmp"(%0, %1) <{predicate = 2 : i64}> : (i32, i32) -> i1 +// CHECK-NEXT: %7 = "llvm.icmp"(%0, %1) <{predicate = 2 : i64}> : (i32, i32) -> i1 %8 = "llvm.icmp"(%0, %1) <{predicate = 3 : i64}> : (i32, i32) -> i1 %9 = llvm.icmp "sle" %0, %1 : i32 -// CHECK-NEXT: %8 = "llvm.icmp"(%0, %1) <{"predicate" = 3 : i64}> : (i32, i32) -> i1 -// CHECK-NEXT: %9 = "llvm.icmp"(%0, %1) <{"predicate" = 3 : i64}> : (i32, i32) -> i1 +// CHECK-NEXT: %8 = "llvm.icmp"(%0, %1) <{predicate = 3 : i64}> : (i32, i32) -> i1 +// CHECK-NEXT: %9 = "llvm.icmp"(%0, %1) <{predicate = 3 : i64}> : (i32, i32) -> i1 %10 = "llvm.icmp"(%0, %1) <{predicate = 4 : i64}> : (i32, i32) -> i1 %11 = llvm.icmp "sgt" %0, %1 : i32 -// CHECK-NEXT: %10 = "llvm.icmp"(%0, %1) <{"predicate" = 4 : i64}> : (i32, i32) -> i1 -// CHECK-NEXT: %11 = "llvm.icmp"(%0, %1) <{"predicate" = 4 : i64}> : (i32, i32) -> i1 +// CHECK-NEXT: %10 = "llvm.icmp"(%0, %1) <{predicate = 4 : i64}> : (i32, i32) -> i1 +// CHECK-NEXT: %11 = "llvm.icmp"(%0, %1) <{predicate = 4 : i64}> : (i32, i32) -> i1 %12 = "llvm.icmp"(%0, %1) <{predicate = 5 : i64}> : (i32, i32) -> i1 %13 = llvm.icmp "sge" %0, %1 : i32 -// CHECK-NEXT: %12 = "llvm.icmp"(%0, %1) <{"predicate" = 5 : i64}> : (i32, i32) -> i1 -// CHECK-NEXT: %13 = "llvm.icmp"(%0, %1) <{"predicate" = 5 : i64}> : (i32, i32) -> i1 +// CHECK-NEXT: %12 = "llvm.icmp"(%0, %1) <{predicate = 5 : i64}> : (i32, i32) -> i1 +// CHECK-NEXT: %13 = "llvm.icmp"(%0, %1) <{predicate = 5 : i64}> : (i32, i32) -> i1 %14 = "llvm.icmp"(%0, %1) <{predicate = 6 : i64}> : (i32, i32) -> i1 %15 = llvm.icmp "ult" %0, %1 : i32 -// CHECK-NEXT: %14 = "llvm.icmp"(%0, %1) <{"predicate" = 6 : i64}> : (i32, i32) -> i1 -// CHECK-NEXT: %15 = "llvm.icmp"(%0, %1) <{"predicate" = 6 : i64}> : (i32, i32) -> i1 +// CHECK-NEXT: %14 = "llvm.icmp"(%0, %1) <{predicate = 6 : i64}> : (i32, i32) -> i1 +// CHECK-NEXT: %15 = "llvm.icmp"(%0, %1) <{predicate = 6 : i64}> : (i32, i32) -> i1 %16 = "llvm.icmp"(%0, %1) <{predicate = 7 : i64}> : (i32, i32) -> i1 %17 = llvm.icmp "ule" %0, %1 : i32 -// CHECK-NEXT: %16 = "llvm.icmp"(%0, %1) <{"predicate" = 7 : i64}> : (i32, i32) -> i1 -// CHECK-NEXT: %17 = "llvm.icmp"(%0, %1) <{"predicate" = 7 : i64}> : (i32, i32) -> i1 +// CHECK-NEXT: %16 = "llvm.icmp"(%0, %1) <{predicate = 7 : i64}> : (i32, i32) -> i1 +// CHECK-NEXT: %17 = "llvm.icmp"(%0, %1) <{predicate = 7 : i64}> : (i32, i32) -> i1 %18 = "llvm.icmp"(%0, %1) <{predicate = 8 : i64}> : (i32, i32) -> i1 %19 = llvm.icmp "ugt" %0, %1 : i32 -// CHECK-NEXT: %18 = "llvm.icmp"(%0, %1) <{"predicate" = 8 : i64}> : (i32, i32) -> i1 -// CHECK-NEXT: %19 = "llvm.icmp"(%0, %1) <{"predicate" = 8 : i64}> : (i32, i32) -> i1 +// CHECK-NEXT: %18 = "llvm.icmp"(%0, %1) <{predicate = 8 : i64}> : (i32, i32) -> i1 +// CHECK-NEXT: %19 = "llvm.icmp"(%0, %1) <{predicate = 8 : i64}> : (i32, i32) -> i1 %20 = "llvm.icmp"(%0, %1) <{predicate = 9 : i64}> : (i32, i32) -> i1 %21 = llvm.icmp "uge" %0, %1 : i32 -// CHECK-NEXT: %20 = "llvm.icmp"(%0, %1) <{"predicate" = 9 : i64}> : (i32, i32) -> i1 -// CHECK-NEXT: %21 = "llvm.icmp"(%0, %1) <{"predicate" = 9 : i64}> : (i32, i32) -> i1 +// CHECK-NEXT: %20 = "llvm.icmp"(%0, %1) <{predicate = 9 : i64}> : (i32, i32) -> i1 +// CHECK-NEXT: %21 = "llvm.icmp"(%0, %1) <{predicate = 9 : i64}> : (i32, i32) -> i1 diff --git a/tests/filecheck/mlir-conversion/with-mlir/dialects/llvm/inline_asm.mlir b/tests/filecheck/mlir-conversion/with-mlir/dialects/llvm/inline_asm.mlir index b734952ad8..dff1ac13f7 100644 --- a/tests/filecheck/mlir-conversion/with-mlir/dialects/llvm/inline_asm.mlir +++ b/tests/filecheck/mlir-conversion/with-mlir/dialects/llvm/inline_asm.mlir @@ -10,7 +10,7 @@ // CHECK: %0 = "test.op"() : () -> i32 // CHECK-NEXT: 1 = "test.op"() : () -> i32 -// CHECK-NEXT: "llvm.inline_asm"(%0, %1) <{"asm_string" = "csrw $0, $1", "constraints" = "i, r", "has_side_effects"}> : (i32, i32) -> () +// CHECK-NEXT: "llvm.inline_asm"(%0, %1) <{asm_string = "csrw $0, $1", constraints = "i, r", has_side_effects}> : (i32, i32) -> () // CHECK-NEXT: %2 = "test.op"() : () -> vector<8xf32> // CHECK-NEXT: %3 = "test.op"() : () -> vector<8xf32> -// CHECK-NEXT: %4 = "llvm.inline_asm"(%2, %3) <{"asm_dialect" = 1 : i64, "asm_string" = "vaddps $0, $1, $2", "constraints" = "=x,x,x"}> : (vector<8xf32>, vector<8xf32>) -> vector<8xf32> +// CHECK-NEXT: %4 = "llvm.inline_asm"(%2, %3) <{asm_dialect = 1 : i64, asm_string = "vaddps $0, $1, $2", constraints = "=x,x,x"}> : (vector<8xf32>, vector<8xf32>) -> vector<8xf32> diff --git a/tests/filecheck/mlir-conversion/with-mlir/dialects/llvm/llvm_func.mlir b/tests/filecheck/mlir-conversion/with-mlir/dialects/llvm/llvm_func.mlir index dccfc67ca8..d0b5321893 100644 --- a/tests/filecheck/mlir-conversion/with-mlir/dialects/llvm/llvm_func.mlir +++ b/tests/filecheck/mlir-conversion/with-mlir/dialects/llvm/llvm_func.mlir @@ -17,14 +17,14 @@ }) : () -> () -// CHECK: %{{.*}} = "arith.constant"() <{"value" = 10 : i32}> {"truc" = !llvm.func (i64, ...)>} : () -> i32 -// CHECK-NEXT: "llvm.func"() <{"CConv" = #llvm.cconv, "function_type" = !llvm.func, "linkage" = #llvm.linkage<"external">, "sym_name" = "printf", "visibility_" = 0 : i64}> ({ +// CHECK: %{{.*}} = "arith.constant"() <{value = 10 : i32}> {truc = !llvm.func (i64, ...)>} : () -> i32 +// CHECK-NEXT: "llvm.func"() <{CConv = #llvm.cconv, function_type = !llvm.func, linkage = #llvm.linkage<"external">, sym_name = "printf", visibility_ = 0 : i64}> ({ // CHECK-NEXT: ^{{.*}}(%arg0 : i64): -// CHECK-NEXT: %{{.*}} = "llvm.mlir.constant"() <{"value" = 1 : i64}> : () -> i64 -// CHECK-NEXT: %{{.*}} = "llvm.call_intrinsic"(%arg0, %{{.*}}) <{"fastmathFlags" = #llvm.fastmath, "intrin" = "llvm.my_intrin"}> : (i64, i64) -> i64 +// CHECK-NEXT: %{{.*}} = "llvm.mlir.constant"() <{value = 1 : i64}> : () -> i64 +// CHECK-NEXT: %{{.*}} = "llvm.call_intrinsic"(%arg0, %{{.*}}) <{fastmathFlags = #llvm.fastmath, intrin = "llvm.my_intrin"}> : (i64, i64) -> i64 // CHECK-NEXT: "llvm.return"() : () -> () // CHECK-NEXT: }) : () -> () // CHECK-NEXT: %{{.*}} = "test.op"() : () -> i64 -// CHECK-NEXT: "llvm.call"(%{{.*}}) <{"CConv" = #llvm.cconv, "TailCallKind" = #llvm.tailcallkind, "callee" = @printf, "fastmathFlags" = #llvm.fastmath, "var_callee_type" = !llvm.func}> : (i64) -> () -// CHECK-NEXT: "llvm.func"() <{"CConv" = #llvm.cconv, "function_type" = !llvm.func, "linkage" = #llvm.linkage<"external">, "sym_name" = "nop", "visibility_" = 1 : i64}> ({ +// CHECK-NEXT: "llvm.call"(%{{.*}}) <{CConv = #llvm.cconv, TailCallKind = #llvm.tailcallkind, callee = @printf, fastmathFlags = #llvm.fastmath, var_callee_type = !llvm.func}> : (i64) -> () +// CHECK-NEXT: "llvm.func"() <{CConv = #llvm.cconv, function_type = !llvm.func, linkage = #llvm.linkage<"external">, sym_name = "nop", visibility_ = 1 : i64}> ({ // CHECK-NEXT: }) : () -> () diff --git a/tests/filecheck/mlir-conversion/with-mlir/dialects/llvm/llvm_pointer_ops.mlir b/tests/filecheck/mlir-conversion/with-mlir/dialects/llvm/llvm_pointer_ops.mlir index b97168c945..2b49bf2eba 100644 --- a/tests/filecheck/mlir-conversion/with-mlir/dialects/llvm/llvm_pointer_ops.mlir +++ b/tests/filecheck/mlir-conversion/with-mlir/dialects/llvm/llvm_pointer_ops.mlir @@ -14,14 +14,14 @@ }) : () -> () // CHECK: "builtin.module"() ({ -// CHECK-NEXT: %0 = "arith.constant"() <{"value" = 0 : i64}> : () -> i64 +// CHECK-NEXT: %0 = "arith.constant"() <{value = 0 : i64}> : () -> i64 // CHECK-NEXT: %1 = "llvm.inttoptr"(%0) : (i64) -> !llvm.ptr -// CHECK-NEXT: %2 = "llvm.load"(%1) <{"ordering" = 0 : i64}> : (!llvm.ptr) -> i32 +// CHECK-NEXT: %2 = "llvm.load"(%1) <{ordering = 0 : i64}> : (!llvm.ptr) -> i32 // CHECK-NEXT: %3 = "llvm.mlir.null"() : () -> !llvm.ptr -// CHECK-NEXT: %4 = "llvm.alloca"(%0) <{"alignment" = 32 : i64, "elem_type" = i64}> : (i64) -> !llvm.ptr -// CHECK-NEXT: %5 = "llvm.load"(%4) <{"ordering" = 0 : i64}> : (!llvm.ptr) -> i64 -// CHECK-NEXT: %6 = "llvm.alloca"(%0) <{"alignment" = 32 : i64, "elem_type" = i32}> : (i64) -> !llvm.ptr -// CHECK-NEXT: %7 = "llvm.getelementptr"(%6, %0) <{"elem_type" = i64, "rawConstantIndices" = array}> : (!llvm.ptr, i64) -> !llvm.ptr -// CHECK-NEXT: %8 = "llvm.getelementptr"(%4, %0) <{"elem_type" = i32, "rawConstantIndices" = array}> : (!llvm.ptr, i64) -> !llvm.ptr -// CHECK-NEXT: "llvm.store"(%5, %6) <{"alignment" = 32 : i64, "nontemporal", "ordering" = 0 : i64, "volatile_"}> : (i64, !llvm.ptr) -> () +// CHECK-NEXT: %4 = "llvm.alloca"(%0) <{alignment = 32 : i64, elem_type = i64}> : (i64) -> !llvm.ptr +// CHECK-NEXT: %5 = "llvm.load"(%4) <{ordering = 0 : i64}> : (!llvm.ptr) -> i64 +// CHECK-NEXT: %6 = "llvm.alloca"(%0) <{alignment = 32 : i64, elem_type = i32}> : (i64) -> !llvm.ptr +// CHECK-NEXT: %7 = "llvm.getelementptr"(%6, %0) <{elem_type = i64, rawConstantIndices = array}> : (!llvm.ptr, i64) -> !llvm.ptr +// CHECK-NEXT: %8 = "llvm.getelementptr"(%4, %0) <{elem_type = i32, rawConstantIndices = array}> : (!llvm.ptr, i64) -> !llvm.ptr +// CHECK-NEXT: "llvm.store"(%5, %6) <{alignment = 32 : i64, nontemporal, ordering = 0 : i64, volatile_}> : (i64, !llvm.ptr) -> () // CHECK-NEXT: }) : () -> () diff --git a/tests/filecheck/mlir-conversion/with-mlir/dialects/memref/matmul.mlir b/tests/filecheck/mlir-conversion/with-mlir/dialects/memref/matmul.mlir index e367162f82..0d7a788ab3 100644 --- a/tests/filecheck/mlir-conversion/with-mlir/dialects/memref/matmul.mlir +++ b/tests/filecheck/mlir-conversion/with-mlir/dialects/memref/matmul.mlir @@ -35,16 +35,16 @@ }) : () -> () // CHECK: "builtin.module"() ({ -// CHECK-NEXT: "func.func"() <{"function_type" = (memref, memref) -> memref, "sym_name" = "matmul", "sym_visibility" = "private"}> ({ +// CHECK-NEXT: "func.func"() <{function_type = (memref, memref) -> memref, sym_name = "matmul", sym_visibility = "private"}> ({ // CHECK-NEXT: ^0(%{{.+}} : memref, %{{.+}} : memref): -// CHECK-NEXT: %{{.+}} = "arith.constant"() <{"value" = 0 : index}> : () -> index -// CHECK-NEXT: %{{.+}} = "arith.constant"() <{"value" = 1 : index}> : () -> index +// CHECK-NEXT: %{{.+}} = "arith.constant"() <{value = 0 : index}> : () -> index +// CHECK-NEXT: %{{.+}} = "arith.constant"() <{value = 1 : index}> : () -> index // CHECK-NEXT: %{{.+}} = "memref.dim"(%{{.+}}, %{{.+}}) : (memref, index) -> index // CHECK-NEXT: %{{.+}} = "memref.dim"(%{{.+}}, %{{.+}}) : (memref, index) -> index // CHECK-NEXT: %{{.+}} = "memref.dim"(%{{.+}}, %{{.+}}) : (memref, index) -> index // CHECK-NEXT: %{{.+}} = "memref.dim"(%{{.+}}, %{{.+}}) : (memref, index) -> index -// CHECK-NEXT: %{{.+}} = "memref.alloca"(%{{.+}}, %{{.+}}) <{"alignment" = 0 : i64, "operandSegmentSizes" = array}> : (index, index) -> memref -// CHECK-NEXT: %{{.+}} = "arith.constant"() <{"value" = 0 : i64}> : () -> i64 +// CHECK-NEXT: %{{.+}} = "memref.alloca"(%{{.+}}, %{{.+}}) <{alignment = 0 : i64, operandSegmentSizes = array}> : (index, index) -> memref +// CHECK-NEXT: %{{.+}} = "arith.constant"() <{value = 0 : i64}> : () -> i64 // CHECK-NEXT: "scf.for"(%{{.+}}, %{{.+}}, %{{.+}}) ({ // CHECK-NEXT: ^1(%{{.+}} : index): // CHECK-NEXT: "scf.for"(%{{.+}}, %{{.+}}, %{{.+}}) ({ @@ -55,8 +55,8 @@ // CHECK-NEXT: %{{.+}} = "memref.load"(%{{.+}}, %{{.+}}, %{{.+}}) : (memref, index, index) -> i64 // CHECK-NEXT: %{{.+}} = "memref.load"(%{{.+}}, %{{.+}}, %{{.+}}) : (memref, index, index) -> i64 // CHECK-NEXT: %{{.+}} = "memref.load"(%{{.+}}, %{{.+}}, %{{.+}}) : (memref, index, index) -> i64 -// CHECK-NEXT: %{{.+}} = "arith.muli"(%{{.+}}, %{{.+}}) <{"overflowFlags" = #arith.overflow}> : (i64, i64) -> i64 -// CHECK-NEXT: %{{.+}} = "arith.addi"(%{{.+}}, %{{.+}}) <{"overflowFlags" = #arith.overflow}> : (i64, i64) -> i64 +// CHECK-NEXT: %{{.+}} = "arith.muli"(%{{.+}}, %{{.+}}) <{overflowFlags = #arith.overflow}> : (i64, i64) -> i64 +// CHECK-NEXT: %{{.+}} = "arith.addi"(%{{.+}}, %{{.+}}) <{overflowFlags = #arith.overflow}> : (i64, i64) -> i64 // CHECK-NEXT: "memref.store"(%{{.+}}, %{{.+}}, %{{.+}}, %{{.+}}) : (i64, memref, index, index) -> () // CHECK-NEXT: "scf.yield"() : () -> () // CHECK-NEXT: }) : (index, index, index) -> () diff --git a/tests/filecheck/mlir-conversion/with-mlir/dialects/memref/memref_ops_mlir_conversion.mlir b/tests/filecheck/mlir-conversion/with-mlir/dialects/memref/memref_ops_mlir_conversion.mlir index 10a88fa375..b256069ccd 100644 --- a/tests/filecheck/mlir-conversion/with-mlir/dialects/memref/memref_ops_mlir_conversion.mlir +++ b/tests/filecheck/mlir-conversion/with-mlir/dialects/memref/memref_ops_mlir_conversion.mlir @@ -44,33 +44,33 @@ // CHECK: "builtin.module"() ({ -// CHECK-NEXT: "memref.global"() <{"alignment" = 64 : i64, "initial_value" = dense<0> : tensor<1xindex>, "sym_name" = "g_with_alignment", "sym_visibility" = "public", "type" = memref<1xindex>}> : () -> () -// CHECK-NEXT: "memref.global"() <{"initial_value" = dense<0> : tensor<1xindex>, "sym_name" = "g", "sym_visibility" = "public", "type" = memref<1xindex>}> : () -> () -// CHECK-NEXT: "memref.global"() <{"constant", "initial_value" = dense<0> : tensor<1xindex>, "sym_name" = "g_constant", "sym_visibility" = "public", "type" = memref<1xindex>}> : () -> () -// CHECK-NEXT: "func.func"() <{"function_type" = () -> (), "sym_name" = "memref_test", "sym_visibility" = "private"}> ({ -// CHECK-NEXT: %0 = "memref.get_global"() <{"name" = @g}> : () -> memref<1xindex> -// CHECK-NEXT: %1 = "arith.constant"() <{"value" = 0 : index}> : () -> index -// CHECK-NEXT: %2 = "memref.alloca"() <{"alignment" = 0 : i64, "operandSegmentSizes" = array}> : () -> memref<1xindex> -// CHECK-NEXT: %3 = "arith.constant"() <{"value" = 42 : index}> : () -> index +// CHECK-NEXT: "memref.global"() <{alignment = 64 : i64, initial_value = dense<0> : tensor<1xindex>, sym_name = "g_with_alignment", sym_visibility = "public", type = memref<1xindex>}> : () -> () +// CHECK-NEXT: "memref.global"() <{initial_value = dense<0> : tensor<1xindex>, sym_name = "g", sym_visibility = "public", type = memref<1xindex>}> : () -> () +// CHECK-NEXT: "memref.global"() <{constant, initial_value = dense<0> : tensor<1xindex>, sym_name = "g_constant", sym_visibility = "public", type = memref<1xindex>}> : () -> () +// CHECK-NEXT: "func.func"() <{function_type = () -> (), sym_name = "memref_test", sym_visibility = "private"}> ({ +// CHECK-NEXT: %0 = "memref.get_global"() <{name = @g}> : () -> memref<1xindex> +// CHECK-NEXT: %1 = "arith.constant"() <{value = 0 : index}> : () -> index +// CHECK-NEXT: %2 = "memref.alloca"() <{alignment = 0 : i64, operandSegmentSizes = array}> : () -> memref<1xindex> +// CHECK-NEXT: %3 = "arith.constant"() <{value = 42 : index}> : () -> index // CHECK-NEXT: "memref.store"(%3, %2, %1) : (index, memref<1xindex>, index) -> () // CHECK-NEXT: %4 = "memref.load"(%2, %1) : (memref<1xindex>, index) -> index -// CHECK-NEXT: %5 = "memref.alloc"() <{"alignment" = 0 : i64, "operandSegmentSizes" = array}> : () -> memref<10x2xindex> +// CHECK-NEXT: %5 = "memref.alloc"() <{alignment = 0 : i64, operandSegmentSizes = array}> : () -> memref<10x2xindex> // CHECK-NEXT: "memref.store"(%3, %5, %3, %4) : (index, memref<10x2xindex>, index, index) -> () -// CHECK-NEXT: %6 = "memref.subview"(%5) <{"operandSegmentSizes" = array, "static_offsets" = array, "static_sizes" = array, "static_strides" = array}> : (memref<10x2xindex>) -> memref<1x1xindex, strided<[2, 1]>> +// CHECK-NEXT: %6 = "memref.subview"(%5) <{operandSegmentSizes = array, static_offsets = array, static_sizes = array, static_strides = array}> : (memref<10x2xindex>) -> memref<1x1xindex, strided<[2, 1]>> // CHECK-NEXT: %7 = "memref.cast"(%5) : (memref<10x2xindex>) -> memref -// CHECK-NEXT: %8 = "memref.alloca"() <{"operandSegmentSizes" = array}> {"i64"} : () -> memref<1xindex> +// CHECK-NEXT: %8 = "memref.alloca"() <{operandSegmentSizes = array}> {i64} : () -> memref<1xindex> // CHECK-NEXT: "memref.copy"(%8, %2) : (memref<1xindex>, memref<1xindex>) -> () // CHECK-NEXT: "memref.dealloc"(%8) : (memref<1xindex>) -> () // CHECK-NEXT: "memref.dealloc"(%2) : (memref<1xindex>) -> () // CHECK-NEXT: "memref.dealloc"(%5) : (memref<10x2xindex>) -> () -// CHECK: "memref.dma_start"(%9, %1, %10, %1, %3, %11, %1) {"operandSegmentSizes" = array} : (memref<100xi32, 10 : i64>, index, memref<100xi32, 9 : i64>, index, index, memref<100xi32>, index) -> () -// CHECK-NEXT: "memref.dma_wait"(%11, %1, %3) {"operandSegmentSizes" = array} : (memref<100xi32>, index, index) -> () -// CHECK-NEXT: %12 = "memref.alloc"() <{"operandSegmentSizes" = array}> : () -> memref<32x32xf32> -// CHECK-NEXT: %13 = "arith.constant"() <{"value" = 1.000000e+00 : f32}> : () -> f32 -// CHECK-NEXT: %14 = "memref.atomic_rmw"(%13, %12, %1, %1) <{"kind" = 0 : i64}> : (f32, memref<32x32xf32>, index, index) -> f32 -// CHECK-NEXT: %15 = "arith.constant"() <{"value" = 2 : index}> : () -> index -// CHECK-NEXT: %16 = "memref.subview"(%5, %15) <{"operandSegmentSizes" = array, "static_offsets" = array, "static_sizes" = array, "static_strides" = array}> : (memref<10x2xindex>, index) -> memref<2xindex, strided<[1], offset: ?>> +// CHECK: "memref.dma_start"(%9, %1, %10, %1, %3, %11, %1) {operandSegmentSizes = array} : (memref<100xi32, 10 : i64>, index, memref<100xi32, 9 : i64>, index, index, memref<100xi32>, index) -> () +// CHECK-NEXT: "memref.dma_wait"(%11, %1, %3) {operandSegmentSizes = array} : (memref<100xi32>, index, index) -> () +// CHECK-NEXT: %12 = "memref.alloc"() <{operandSegmentSizes = array}> : () -> memref<32x32xf32> +// CHECK-NEXT: %13 = "arith.constant"() <{value = 1.000000e+00 : f32}> : () -> f32 +// CHECK-NEXT: %14 = "memref.atomic_rmw"(%13, %12, %1, %1) <{kind = 0 : i64}> : (f32, memref<32x32xf32>, index, index) -> f32 +// CHECK-NEXT: %15 = "arith.constant"() <{value = 2 : index}> : () -> index +// CHECK-NEXT: %16 = "memref.subview"(%5, %15) <{operandSegmentSizes = array, static_offsets = array, static_sizes = array, static_strides = array}> : (memref<10x2xindex>, index) -> memref<2xindex, strided<[1], offset: ?>> // CHECK-NEXT: "func.return"() : () -> () // CHECK-NEXT: }) : () -> () // CHECK-NEXT: }) : () -> () diff --git a/tests/filecheck/mlir-conversion/with-mlir/dialects/omp/ops.mlir b/tests/filecheck/mlir-conversion/with-mlir/dialects/omp/ops.mlir index 25232a5847..2518d9bd5e 100644 --- a/tests/filecheck/mlir-conversion/with-mlir/dialects/omp/ops.mlir +++ b/tests/filecheck/mlir-conversion/with-mlir/dialects/omp/ops.mlir @@ -151,39 +151,39 @@ builtin.module { // CHECK: builtin.module { // CHECK-NEXT: func.func @omp_parallel(%arg37 : memref<1xi32>, %arg38 : i1, %arg39 : i32) { -// CHECK-NEXT: "omp.parallel"(%arg38, %arg39, %arg37, %arg37) <{"operandSegmentSizes" = array, "proc_bind_val" = #omp}> ({ -// CHECK-NEXT: "omp.parallel"(%arg39, %arg37, %arg37) <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "omp.parallel"(%arg38, %arg39, %arg37, %arg37) <{operandSegmentSizes = array, proc_bind_val = #omp}> ({ +// CHECK-NEXT: "omp.parallel"(%arg39, %arg37, %arg37) <{operandSegmentSizes = array}> ({ // CHECK-NEXT: "omp.terminator"() : () -> () // CHECK-NEXT: }) : (i32, memref<1xi32>, memref<1xi32>) -> () -// CHECK-NEXT: "omp.parallel"(%arg38, %arg37, %arg37) <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "omp.parallel"(%arg38, %arg37, %arg37) <{operandSegmentSizes = array}> ({ // CHECK-NEXT: "omp.terminator"() : () -> () // CHECK-NEXT: }) : (i1, memref<1xi32>, memref<1xi32>) -> () -// CHECK-NEXT: "omp.parallel"(%arg38, %arg39) <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "omp.parallel"(%arg38, %arg39) <{operandSegmentSizes = array}> ({ // CHECK-NEXT: "omp.terminator"() : () -> () // CHECK-NEXT: }) : (i1, i32) -> () // CHECK-NEXT: "omp.terminator"() : () -> () // CHECK-NEXT: }) : (i1, i32, memref<1xi32>, memref<1xi32>) -> () -// CHECK-NEXT: "omp.parallel"(%arg37, %arg37) <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "omp.parallel"(%arg37, %arg37) <{operandSegmentSizes = array}> ({ // CHECK-NEXT: "omp.terminator"() : () -> () // CHECK-NEXT: }) : (memref<1xi32>, memref<1xi32>) -> () // CHECK-NEXT: func.return // CHECK-NEXT: } // CHECK-NEXT: func.func @omp_ordered(%arg27 : i32, %arg28 : i32, %arg29 : i32, %arg30 : i64, %arg31 : i64, %arg32 : i64, %arg33 : i64) { -// CHECK-NEXT: "omp.wsloop"() <{"operandSegmentSizes" = array, "ordered_val" = 0 : i64}> ({ +// CHECK-NEXT: "omp.wsloop"() <{operandSegmentSizes = array, ordered_val = 0 : i64}> ({ // CHECK-NEXT: "omp.loop_nest"(%arg27, %arg28, %arg29) ({ // CHECK-NEXT: ^0(%arg36 : i32): // CHECK-NEXT: omp.yield // CHECK-NEXT: }) : (i32, i32, i32) -> () // CHECK-NEXT: "omp.terminator"() : () -> () // CHECK-NEXT: }) : () -> () -// CHECK-NEXT: "omp.wsloop"() <{"operandSegmentSizes" = array, "ordered_val" = 1 : i64}> ({ +// CHECK-NEXT: "omp.wsloop"() <{operandSegmentSizes = array, ordered_val = 1 : i64}> ({ // CHECK-NEXT: "omp.loop_nest"(%arg27, %arg28, %arg29) ({ // CHECK-NEXT: ^1(%arg35 : i32): // CHECK-NEXT: omp.yield // CHECK-NEXT: }) : (i32, i32, i32) -> () // CHECK-NEXT: "omp.terminator"() : () -> () // CHECK-NEXT: }) : () -> () -// CHECK-NEXT: "omp.wsloop"() <{"operandSegmentSizes" = array, "ordered_val" = 2 : i64}> ({ +// CHECK-NEXT: "omp.wsloop"() <{operandSegmentSizes = array, ordered_val = 2 : i64}> ({ // CHECK-NEXT: "omp.loop_nest"(%arg27, %arg28, %arg29) ({ // CHECK-NEXT: ^2(%arg34 : i32): // CHECK-NEXT: omp.yield @@ -193,35 +193,35 @@ builtin.module { // CHECK-NEXT: func.return // CHECK-NEXT: } // CHECK-NEXT: func.func @omp_wsloop(%arg16 : index, %arg17 : index, %arg18 : index, %arg19 : memref<1xi32>, %arg20 : i32, %arg21 : i32) { -// CHECK-NEXT: "omp.wsloop"() <{"operandSegmentSizes" = array, "ordered_val" = 1 : i64}> ({ +// CHECK-NEXT: "omp.wsloop"() <{operandSegmentSizes = array, ordered_val = 1 : i64}> ({ // CHECK-NEXT: "omp.loop_nest"(%arg16, %arg17, %arg18) ({ // CHECK-NEXT: ^0(%arg26 : index): // CHECK-NEXT: omp.yield // CHECK-NEXT: }) : (index, index, index) -> () // CHECK-NEXT: "omp.terminator"() : () -> () // CHECK-NEXT: }) : () -> () -// CHECK-NEXT: "omp.wsloop"(%arg19, %arg20) <{"operandSegmentSizes" = array, "schedule_val" = #omp}> ({ +// CHECK-NEXT: "omp.wsloop"(%arg19, %arg20) <{operandSegmentSizes = array, schedule_val = #omp}> ({ // CHECK-NEXT: "omp.loop_nest"(%arg16, %arg17, %arg18) ({ // CHECK-NEXT: ^1(%arg25 : index): // CHECK-NEXT: omp.yield // CHECK-NEXT: }) : (index, index, index) -> () // CHECK-NEXT: "omp.terminator"() : () -> () // CHECK-NEXT: }) : (memref<1xi32>, i32) -> () -// CHECK-NEXT: "omp.wsloop"(%arg19, %arg19, %arg20, %arg20) <{"operandSegmentSizes" = array, "schedule_val" = #omp}> ({ +// CHECK-NEXT: "omp.wsloop"(%arg19, %arg19, %arg20, %arg20) <{operandSegmentSizes = array, schedule_val = #omp}> ({ // CHECK-NEXT: "omp.loop_nest"(%arg16, %arg17, %arg18) ({ // CHECK-NEXT: ^2(%arg24 : index): // CHECK-NEXT: omp.yield // CHECK-NEXT: }) : (index, index, index) -> () // CHECK-NEXT: "omp.terminator"() : () -> () // CHECK-NEXT: }) : (memref<1xi32>, memref<1xi32>, i32, i32) -> () -// CHECK-NEXT: "omp.wsloop"(%arg19, %arg20, %arg21) <{"operandSegmentSizes" = array, "ordered_val" = 2 : i64, "schedule_val" = #omp}> ({ +// CHECK-NEXT: "omp.wsloop"(%arg19, %arg20, %arg21) <{operandSegmentSizes = array, ordered_val = 2 : i64, schedule_val = #omp}> ({ // CHECK-NEXT: "omp.loop_nest"(%arg16, %arg17, %arg18) ({ // CHECK-NEXT: ^3(%arg23 : index): // CHECK-NEXT: omp.yield // CHECK-NEXT: }) : (index, index, index) -> () // CHECK-NEXT: "omp.terminator"() : () -> () // CHECK-NEXT: }) : (memref<1xi32>, i32, i32) -> () -// CHECK-NEXT: "omp.wsloop"() <{"nowait", "operandSegmentSizes" = array, "schedule_val" = #omp}> ({ +// CHECK-NEXT: "omp.wsloop"() <{nowait, operandSegmentSizes = array, schedule_val = #omp}> ({ // CHECK-NEXT: "omp.loop_nest"(%arg16, %arg17, %arg18) ({ // CHECK-NEXT: ^4(%arg22 : index): // CHECK-NEXT: omp.yield @@ -231,63 +231,63 @@ builtin.module { // CHECK-NEXT: func.return // CHECK-NEXT: } // CHECK-NEXT: func.func @omp_wsloop_pretty(%arg0 : index, %arg1 : index, %arg2 : index, %arg3 : memref<1xi32>, %arg4 : i32, %arg5 : i32, %arg6 : i16) { -// CHECK-NEXT: "omp.wsloop"() <{"operandSegmentSizes" = array, "ordered_val" = 2 : i64}> ({ +// CHECK-NEXT: "omp.wsloop"() <{operandSegmentSizes = array, ordered_val = 2 : i64}> ({ // CHECK-NEXT: "omp.loop_nest"(%arg0, %arg1, %arg2) ({ // CHECK-NEXT: ^0(%arg15 : index): // CHECK-NEXT: omp.yield // CHECK-NEXT: }) : (index, index, index) -> () // CHECK-NEXT: "omp.terminator"() : () -> () // CHECK-NEXT: }) : () -> () -// CHECK-NEXT: "omp.wsloop"(%arg3, %arg4) <{"operandSegmentSizes" = array, "schedule_val" = #omp}> ({ +// CHECK-NEXT: "omp.wsloop"(%arg3, %arg4) <{operandSegmentSizes = array, schedule_val = #omp}> ({ // CHECK-NEXT: "omp.loop_nest"(%arg0, %arg1, %arg2) ({ // CHECK-NEXT: ^1(%arg14 : index): // CHECK-NEXT: omp.yield // CHECK-NEXT: }) : (index, index, index) -> () // CHECK-NEXT: "omp.terminator"() : () -> () // CHECK-NEXT: }) : (memref<1xi32>, i32) -> () -// CHECK-NEXT: "omp.wsloop"(%arg3, %arg4, %arg5) <{"operandSegmentSizes" = array, "ordered_val" = 2 : i64, "schedule_val" = #omp}> ({ +// CHECK-NEXT: "omp.wsloop"(%arg3, %arg4, %arg5) <{operandSegmentSizes = array, ordered_val = 2 : i64, schedule_val = #omp}> ({ // CHECK-NEXT: "omp.loop_nest"(%arg0, %arg1, %arg2) ({ // CHECK-NEXT: ^2(%arg13 : index): // CHECK-NEXT: omp.yield // CHECK-NEXT: }) : (index, index, index) -> () // CHECK-NEXT: "omp.terminator"() : () -> () // CHECK-NEXT: }) : (memref<1xi32>, i32, i32) -> () -// CHECK-NEXT: "omp.wsloop"(%arg3, %arg4, %arg5) <{"operandSegmentSizes" = array, "ordered_val" = 2 : i64, "schedule_modifier" = #omp, "schedule_val" = #omp}> ({ +// CHECK-NEXT: "omp.wsloop"(%arg3, %arg4, %arg5) <{operandSegmentSizes = array, ordered_val = 2 : i64, schedule_modifier = #omp, schedule_val = #omp}> ({ // CHECK-NEXT: "omp.loop_nest"(%arg0, %arg1, %arg2) ({ // CHECK-NEXT: ^3(%arg12 : index): // CHECK-NEXT: omp.yield // CHECK-NEXT: }) : (index, index, index) -> () // CHECK-NEXT: "omp.terminator"() : () -> () // CHECK-NEXT: }) : (memref<1xi32>, i32, i32) -> () -// CHECK-NEXT: "omp.wsloop"(%arg3, %arg4, %arg6) <{"operandSegmentSizes" = array, "ordered_val" = 2 : i64, "schedule_modifier" = #omp, "schedule_val" = #omp}> ({ +// CHECK-NEXT: "omp.wsloop"(%arg3, %arg4, %arg6) <{operandSegmentSizes = array, ordered_val = 2 : i64, schedule_modifier = #omp, schedule_val = #omp}> ({ // CHECK-NEXT: "omp.loop_nest"(%arg0, %arg1, %arg2) ({ // CHECK-NEXT: ^4(%arg11 : index): // CHECK-NEXT: omp.yield // CHECK-NEXT: }) : (index, index, index) -> () // CHECK-NEXT: "omp.terminator"() : () -> () // CHECK-NEXT: }) : (memref<1xi32>, i32, i16) -> () -// CHECK-NEXT: "omp.wsloop"() <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "omp.wsloop"() <{operandSegmentSizes = array}> ({ // CHECK-NEXT: "omp.loop_nest"(%arg0, %arg1, %arg2) ({ // CHECK-NEXT: ^5(%arg10 : index): // CHECK-NEXT: omp.yield // CHECK-NEXT: }) : (index, index, index) -> () // CHECK-NEXT: "omp.terminator"() : () -> () // CHECK-NEXT: }) : () -> () -// CHECK-NEXT: "omp.wsloop"() <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "omp.wsloop"() <{operandSegmentSizes = array}> ({ // CHECK-NEXT: "omp.loop_nest"(%arg0, %arg1, %arg2) ({ // CHECK-NEXT: ^6(%arg9 : index): // CHECK-NEXT: omp.yield // CHECK-NEXT: }) : (index, index, index) -> () // CHECK-NEXT: "omp.terminator"() : () -> () // CHECK-NEXT: }) : () -> () -// CHECK-NEXT: "omp.wsloop"() <{"nowait", "operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "omp.wsloop"() <{nowait, operandSegmentSizes = array}> ({ // CHECK-NEXT: "omp.loop_nest"(%arg0, %arg1, %arg2) ({ // CHECK-NEXT: ^7(%arg8 : index): // CHECK-NEXT: omp.yield // CHECK-NEXT: }) : (index, index, index) -> () // CHECK-NEXT: "omp.terminator"() : () -> () // CHECK-NEXT: }) : () -> () -// CHECK-NEXT: "omp.wsloop"() <{"nowait", "operandSegmentSizes" = array, "order_val" = #omp}> ({ +// CHECK-NEXT: "omp.wsloop"() <{nowait, operandSegmentSizes = array, order_val = #omp}> ({ // CHECK-NEXT: "omp.loop_nest"(%arg0, %arg1, %arg2) ({ // CHECK-NEXT: ^8(%arg7 : index): // CHECK-NEXT: omp.yield diff --git a/tests/filecheck/mlir-conversion/with-mlir/dialects/scf/for_generic.mlir b/tests/filecheck/mlir-conversion/with-mlir/dialects/scf/for_generic.mlir index 4df216e8d1..2a71462d77 100644 --- a/tests/filecheck/mlir-conversion/with-mlir/dialects/scf/for_generic.mlir +++ b/tests/filecheck/mlir-conversion/with-mlir/dialects/scf/for_generic.mlir @@ -17,13 +17,13 @@ }) : () -> () // CHECK: "builtin.module"() ({ -// CHECK-NEXT: %{{.*}} = "arith.constant"() <{"value" = 0 : index}> : () -> index -// CHECK-NEXT: %{{.*}} = "arith.constant"() <{"value" = 42 : index}> : () -> index -// CHECK-NEXT: %{{.*}} = "arith.constant"() <{"value" = 7 : index}> : () -> index -// CHECK-NEXT: %{{.*}} = "arith.constant"() <{"value" = 36 : index}> : () -> index +// CHECK-NEXT: %{{.*}} = "arith.constant"() <{value = 0 : index}> : () -> index +// CHECK-NEXT: %{{.*}} = "arith.constant"() <{value = 42 : index}> : () -> index +// CHECK-NEXT: %{{.*}} = "arith.constant"() <{value = 7 : index}> : () -> index +// CHECK-NEXT: %{{.*}} = "arith.constant"() <{value = 36 : index}> : () -> index // CHECK-NEXT: %{{.*}} = "scf.for"(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) ({ // CHECK-NEXT: ^0(%{{.*}} : index, %{{.*}} : index): -// CHECK-NEXT: %{{.*}} = "arith.addi"(%{{.*}}, %{{.*}}) <{"overflowFlags" = #arith.overflow}> : (index, index) -> index +// CHECK-NEXT: %{{.*}} = "arith.addi"(%{{.*}}, %{{.*}}) <{overflowFlags = #arith.overflow}> : (index, index) -> index // CHECK-NEXT: "scf.yield"(%{{.*}}) : (index) -> () // CHECK-NEXT: }) : (index, index, index, index) -> index // CHECK-NEXT: "scf.for"(%{{.*}}, %{{.*}}, %{{.*}}) ({ diff --git a/tests/filecheck/mlir-conversion/with-mlir/dialects/scf/for_generic_non_index_iv.mlir b/tests/filecheck/mlir-conversion/with-mlir/dialects/scf/for_generic_non_index_iv.mlir index 16c60c9f7b..9ebdc574f3 100644 --- a/tests/filecheck/mlir-conversion/with-mlir/dialects/scf/for_generic_non_index_iv.mlir +++ b/tests/filecheck/mlir-conversion/with-mlir/dialects/scf/for_generic_non_index_iv.mlir @@ -17,13 +17,13 @@ }) : () -> () // CHECK: "builtin.module"() ({ -// CHECK-NEXT: %{{.*}} = "arith.constant"() <{"value" = 0 : i32}> : () -> i32 -// CHECK-NEXT: %{{.*}} = "arith.constant"() <{"value" = 42 : i32}> : () -> i32 -// CHECK-NEXT: %{{.*}} = "arith.constant"() <{"value" = 7 : i32}> : () -> i32 -// CHECK-NEXT: %{{.*}} = "arith.constant"() <{"value" = 36 : i32}> : () -> i32 +// CHECK-NEXT: %{{.*}} = "arith.constant"() <{value = 0 : i32}> : () -> i32 +// CHECK-NEXT: %{{.*}} = "arith.constant"() <{value = 42 : i32}> : () -> i32 +// CHECK-NEXT: %{{.*}} = "arith.constant"() <{value = 7 : i32}> : () -> i32 +// CHECK-NEXT: %{{.*}} = "arith.constant"() <{value = 36 : i32}> : () -> i32 // CHECK-NEXT: %{{.*}} = "scf.for"(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) ({ // CHECK-NEXT: ^0(%{{.*}} : i32, %{{.*}} : i32): -// CHECK-NEXT: %{{.*}} = "arith.addi"(%{{.*}}, %{{.*}}) <{"overflowFlags" = #arith.overflow}> : (i32, i32) -> i32 +// CHECK-NEXT: %{{.*}} = "arith.addi"(%{{.*}}, %{{.*}}) <{overflowFlags = #arith.overflow}> : (i32, i32) -> i32 // CHECK-NEXT: "scf.yield"(%{{.*}}) : (i32) -> () // CHECK-NEXT: }) : (i32, i32, i32, i32) -> i32 // CHECK-NEXT: "scf.for"(%{{.*}}, %{{.*}}, %{{.*}}) ({ diff --git a/tests/filecheck/mlir-conversion/with-mlir/dialects/scf/parallel.mlir b/tests/filecheck/mlir-conversion/with-mlir/dialects/scf/parallel.mlir index 281e349c4d..c213185cd3 100644 --- a/tests/filecheck/mlir-conversion/with-mlir/dialects/scf/parallel.mlir +++ b/tests/filecheck/mlir-conversion/with-mlir/dialects/scf/parallel.mlir @@ -11,10 +11,10 @@ }) : () -> () // CHECK-NEXT: "builtin.module"() ({ -// CHECK-NEXT: %{{.*}} = "arith.constant"() <{"value" = 0 : index}> : () -> index -// CHECK-NEXT: %{{.*}} = "arith.constant"() <{"value" = 1000 : index}> : () -> index -// CHECK-NEXT: %{{.*}} = "arith.constant"() <{"value" = 3 : index}> : () -> index -// CHECK-NEXT: "scf.parallel"(%{{.*}}, %{{.*}}, %{{.*}}) <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: %{{.*}} = "arith.constant"() <{value = 0 : index}> : () -> index +// CHECK-NEXT: %{{.*}} = "arith.constant"() <{value = 1000 : index}> : () -> index +// CHECK-NEXT: %{{.*}} = "arith.constant"() <{value = 3 : index}> : () -> index +// CHECK-NEXT: "scf.parallel"(%{{.*}}, %{{.*}}, %{{.*}}) <{operandSegmentSizes = array}> ({ // CHECK-NEXT: ^{{.*}}(%{{.*}}: index): // CHECK-NEXT: scf.reduce // CHECK-NEXT: }) : (index, index, index) -> () diff --git a/tests/filecheck/mlir-conversion/with-mlir/dialects/scf/parallel_with_reduce.mlir b/tests/filecheck/mlir-conversion/with-mlir/dialects/scf/parallel_with_reduce.mlir index 507db4231a..55c7f71573 100644 --- a/tests/filecheck/mlir-conversion/with-mlir/dialects/scf/parallel_with_reduce.mlir +++ b/tests/filecheck/mlir-conversion/with-mlir/dialects/scf/parallel_with_reduce.mlir @@ -18,16 +18,16 @@ // CHECK: "builtin.module"() ({ -// CHECK-NEXT: %0 = "arith.constant"() <{"value" = 0 : index}> : () -> index -// CHECK-NEXT: %1 = "arith.constant"() <{"value" = 1000 : index}> : () -> index -// CHECK-NEXT: %2 = "arith.constant"() <{"value" = 3 : index}> : () -> index -// CHECK-NEXT: %3 = "arith.constant"() <{"value" = 10 : i32}> : () -> i32 -// CHECK-NEXT: %4 = "arith.constant"() <{"value" = 100 : i32}> : () -> i32 -// CHECK-NEXT: %5 = "scf.parallel"(%0, %1, %2, %3) <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: %0 = "arith.constant"() <{value = 0 : index}> : () -> index +// CHECK-NEXT: %1 = "arith.constant"() <{value = 1000 : index}> : () -> index +// CHECK-NEXT: %2 = "arith.constant"() <{value = 3 : index}> : () -> index +// CHECK-NEXT: %3 = "arith.constant"() <{value = 10 : i32}> : () -> i32 +// CHECK-NEXT: %4 = "arith.constant"() <{value = 100 : i32}> : () -> i32 +// CHECK-NEXT: %5 = "scf.parallel"(%0, %1, %2, %3) <{operandSegmentSizes = array}> ({ // CHECK-NEXT: ^0(%arg0 : index): // CHECK-NEXT: "scf.reduce"(%4) ({ // CHECK-NEXT: ^1(%arg1 : i32, %arg2 : i32): -// CHECK-NEXT: %6 = "arith.addi"(%arg1, %arg2) <{"overflowFlags" = #arith.overflow}> : (i32, i32) -> i32 +// CHECK-NEXT: %6 = "arith.addi"(%arg1, %arg2) <{overflowFlags = #arith.overflow}> : (i32, i32) -> i32 // CHECK-NEXT: "scf.reduce.return"(%6) : (i32) -> () // CHECK-NEXT: }) : (i32) -> () // CHECK-NEXT: }) : (index, index, index, i32) -> i32 diff --git a/tests/filecheck/mlir-conversion/with-mlir/dialects/stencil/ops.mlir b/tests/filecheck/mlir-conversion/with-mlir/dialects/stencil/ops.mlir index eb30f4b01a..bcc46060fd 100644 --- a/tests/filecheck/mlir-conversion/with-mlir/dialects/stencil/ops.mlir +++ b/tests/filecheck/mlir-conversion/with-mlir/dialects/stencil/ops.mlir @@ -53,7 +53,7 @@ func.func @dyn_access(%arg0 : !stencil.field, %arg1 : !stencil.field< } // CHECK: builtin.module { -// CHECK-NEXT: func.func @dyn_access(%arg0 : !stencil.field, %arg1 : !stencil.field) attributes {"stencil.program"}{ +// CHECK-NEXT: func.func @dyn_access(%arg0 : !stencil.field, %arg1 : !stencil.field) attributes {stencil.program}{ // CHECK-NEXT: %0 = stencil.cast %arg0 : !stencil.field -> !stencil.field<[-3,67]x[-3,67]x[0,60]xf64> // CHECK-NEXT: %1 = stencil.cast %arg1 : !stencil.field -> !stencil.field<[-3,67]x[-3,67]x[0,60]xf64> // CHECK-NEXT: %2 = stencil.load %0 : !stencil.field<[-3,67]x[-3,67]x[0,60]xf64> -> !stencil.temp<[0,64]x[0,64]x[0,60]xf64> diff --git a/tests/filecheck/mlir-conversion/with-mlir/mlir_opt.mlir b/tests/filecheck/mlir-conversion/with-mlir/mlir_opt.mlir index cb3d4a695b..f9a7fe24e4 100644 --- a/tests/filecheck/mlir-conversion/with-mlir/mlir_opt.mlir +++ b/tests/filecheck/mlir-conversion/with-mlir/mlir_opt.mlir @@ -23,7 +23,7 @@ // CHECK: "builtin.module"() ({ -// CHECK-NEXT: "func.func"() <{"function_type" = () -> (), "sym_name" = "do_nothing"}> ({ +// CHECK-NEXT: "func.func"() <{function_type = () -> (), sym_name = "do_nothing"}> ({ // CHECK-NEXT: "func.return"() : () -> () // CHECK-NEXT: }) : () -> () -// CHECK-NEXT: }) {"gpu.container_module"} : () -> () +// CHECK-NEXT: }) {gpu.container_module} : () -> () diff --git a/tests/filecheck/mlir-conversion/with-mlir/parser-printer/attribute_names.mlir b/tests/filecheck/mlir-conversion/with-mlir/parser-printer/attribute_names.mlir new file mode 100644 index 0000000000..48ab0062af --- /dev/null +++ b/tests/filecheck/mlir-conversion/with-mlir/parser-printer/attribute_names.mlir @@ -0,0 +1,52 @@ +// RUN: xdsl-opt %s | filecheck %s +// RUN: mlir-opt %s | filecheck %s + +// CHECK: module + +"test.op"() {"name"} : () -> () +// CHECK-NEXT: {name} + +"test.op"() {"_name"} : () -> () +// CHECK-NEXT: {_name} + +"test.op"() {"Name"} : () -> () +// CHECK-NEXT: {Name} + +"test.op"() {"name$"} : () -> () +// CHECK-NEXT: {name$} + +"test.op"() {"name_"} : () -> () +// CHECK-NEXT: {name_} + +"test.op"() {"name."} : () -> () +// CHECK-NEXT: {name.} + +"test.op"() {"name0"} : () -> () +// CHECK-NEXT: {name0} + +"test.op"() {"name%"} : () -> () +// CHECK-NEXT: {"name%"} + +"test.op"() {"name#"} : () -> () +// CHECK-NEXT: {"name#"} + +"test.op"() {"name-"} : () -> () +// CHECK-NEXT: {"name-"} + +"test.op"() {"name{"} : () -> () +// CHECK-NEXT: {"name{"} + +"test.op"() {"name("} : () -> () +// CHECK-NEXT: {"name("} + +"test.op"() {"0name"} : () -> () +// CHECK-NEXT: {"0name"} + +"test.op"() {"name" = "name"} : () -> () +// CHECK-NEXT: {name = "name"} + +"test.op"() {"name#" = "name#"} : () -> () +// CHECK-NEXT: {"name#" = "name#"} + +"test.op"() {"name0" = "name0"} : () -> () +// CHECK-NEXT: {name0 = "name0"} diff --git a/tests/filecheck/parser-printer/affine_map.mlir b/tests/filecheck/parser-printer/affine_map.mlir index 3f257f2916..72d93f9428 100644 --- a/tests/filecheck/parser-printer/affine_map.mlir +++ b/tests/filecheck/parser-printer/affine_map.mlir @@ -1,51 +1,51 @@ // RUN: xdsl-opt --allow-unregistered-dialect %s | filecheck %s builtin.module { - // CHECK: "f1"() {"map" = affine_map<(d0) -> (d0)>} : () -> () + // CHECK: "f1"() {map = affine_map<(d0) -> (d0)>} : () -> () "f1"() {map = affine_map<(d0) -> (d0)>} : () -> () - // CHECK: "f2"() {"map" = affine_map<(d0) -> ((d0 + 1))>} : () -> () + // CHECK: "f2"() {map = affine_map<(d0) -> ((d0 + 1))>} : () -> () "f2"() {map = affine_map<(d0) -> (d0 + 1)>} : () -> () - // CHECK: "f3"() {"map" = affine_map<(d0) -> ((d0 + -1))>} : () -> () + // CHECK: "f3"() {map = affine_map<(d0) -> ((d0 + -1))>} : () -> () "f3"() {map = affine_map<(d0) -> (d0 - 1)>} : () -> () - // CHECK: "f4"() {"map" = affine_map<(d0) -> ((d0 floordiv 2))>} : () -> () + // CHECK: "f4"() {map = affine_map<(d0) -> ((d0 floordiv 2))>} : () -> () "f4"() {map = affine_map<(d0) -> (d0 floordiv 2)>} : () -> () - // CHECK: "f5"() {"map" = affine_map<(d0) -> ((d0 ceildiv 2))>} : () -> () + // CHECK: "f5"() {map = affine_map<(d0) -> ((d0 ceildiv 2))>} : () -> () "f5"() {map = affine_map<(d0) -> (d0 ceildiv 2)>} : () -> () - // CHECK: "f6"() {"map" = affine_map<(d0) -> ((d0 mod 2))>} : () -> () + // CHECK: "f6"() {map = affine_map<(d0) -> ((d0 mod 2))>} : () -> () "f6"() {map = affine_map<(d0) -> (d0 mod 2)>} : () -> () - // CHECK: "f7"() {"map" = affine_map<(d0) -> ((d0 * 2))>} : () -> () + // CHECK: "f7"() {map = affine_map<(d0) -> ((d0 * 2))>} : () -> () "f7"() {map = affine_map<(d0) -> (d0 * 2)>} : () -> () - // CHECK: "f8"() {"map" = affine_map<(d0, d1, d2) -> (d0, d1, d2)>} : () -> () + // CHECK: "f8"() {map = affine_map<(d0, d1, d2) -> (d0, d1, d2)>} : () -> () "f8"() {map = affine_map<(d0, d1, d2) -> (d0, d1, d2)>} : () -> () - // CHECK: "f9"() {"map" = affine_map<(d0, d1, d2)[s0, s1, s2] -> ((d0 + s0), (d1 + s1), (d2 + s2))>} : () -> () + // CHECK: "f9"() {map = affine_map<(d0, d1, d2)[s0, s1, s2] -> ((d0 + s0), (d1 + s1), (d2 + s2))>} : () -> () "f9"() {map = affine_map<(d0, d1, d2)[s0, s1, s2] -> (d0 + s0, d1 + s1, d2 + s2)>} : () -> () - // CHECK: "f10"() {"map" = affine_map<(d0, d1) -> (((((d1 + (d0 * -1)) + (((d0 * 2) + (d1 * -2)) + 2)) + d1) + -1), (((((d1 + d1) + 1) + d1) + d1) + 1))>} : () -> () + // CHECK: "f10"() {map = affine_map<(d0, d1) -> (((((d1 + (d0 * -1)) + (((d0 * 2) + (d1 * -2)) + 2)) + d1) + -1), (((((d1 + d1) + 1) + d1) + d1) + 1))>} : () -> () "f10"() { map = affine_map<(i, j) -> ((j - i) + 2*(i - j + 1) + j - 1 + 0, j + j + 1 + j + j + 1)> } : () -> () - // CHECK: "f11"() {"map" = affine_map<(d0, d1) -> ((d0 + 2), d1)>} : () -> () + // CHECK: "f11"() {map = affine_map<(d0, d1) -> ((d0 + 2), d1)>} : () -> () "f11"() { map = affine_map<(i, j) -> (3+3-2*2+i, j)>} : () -> () - // CHECK: "f12"() {"map" = affine_map<(d0, d1) -> ((d0 + 1), ((d1 * 4) + 2))>} : () -> () + // CHECK: "f12"() {map = affine_map<(d0, d1) -> ((d0 + 1), ((d1 * 4) + 2))>} : () -> () "f12"() { map = affine_map<(i, j) -> (1*i+3*2-2*2-1, 4*j + 2)>} : () -> () - // CHECK: "f13"() {"map" = affine_map<(d0, d1)[s0, s1] -> ((d0 * -5), (d1 * -3), -2, ((d0 * -1) + (d1 * -1)), (s0 * -1))>} : () -> () + // CHECK: "f13"() {map = affine_map<(d0, d1)[s0, s1] -> ((d0 * -5), (d1 * -3), -2, ((d0 * -1) + (d1 * -1)), (s0 * -1))>} : () -> () "f13"() { map = affine_map<(i, j)[s0, s1] -> (-5*i, -3*j, -2, -1*(i+j), -1*s0)>} : () -> () - // CHECK: "f14"() {"map" = affine_map<(d0, d1, d2) -> ((d0 * 0), d1, ((d0 * 128) floordiv 64), ((d1 * 0) floordiv 64))>} : () -> () + // CHECK: "f14"() {map = affine_map<(d0, d1, d2) -> ((d0 * 0), d1, ((d0 * 128) floordiv 64), ((d1 * 0) floordiv 64))>} : () -> () "f14"() { map = affine_map<(i, j, k) -> (i*0, 1*j, i * 128 floordiv 64, j * 0 floordiv 64)>} : () -> () - // CHECK: "f15"() {"map" = affine_map<(d0, d1) -> (8, 4, 1, 3, 2, 4)>} : () -> () + // CHECK: "f15"() {map = affine_map<(d0, d1) -> (8, 4, 1, 3, 2, 4)>} : () -> () "f15"() { map = affine_map<(i, j) -> (5+3, 2*2, 8-7, 100 floordiv 32, 5 mod 3, 10 ceildiv 3)>} : () -> () - // CHECK: "f16"() {"map" = affine_map<(d0, d1) -> (4, 11, 512, 15)>} : () -> () + // CHECK: "f16"() {map = affine_map<(d0, d1) -> (4, 11, 512, 15)>} : () -> () "f16"() { map = affine_map<(i, j) -> (5 mod 3 + 2, 5*3 - 4, 128 * (500 ceildiv 128), 40 floordiv 7 * 3)>} : () -> () } diff --git a/tests/filecheck/parser-printer/affine_set.mlir b/tests/filecheck/parser-printer/affine_set.mlir index d4304a4d64..948f7c91a4 100644 --- a/tests/filecheck/parser-printer/affine_set.mlir +++ b/tests/filecheck/parser-printer/affine_set.mlir @@ -1,22 +1,22 @@ // RUN: xdsl-opt --allow-unregistered-dialect %s | filecheck %s builtin.module { - // CHECK: "f0"() {"set" = affine_set<(d0) : ((d0 + -10) >= 0)>} : () -> () + // CHECK: "f0"() {set = affine_set<(d0) : ((d0 + -10) >= 0)>} : () -> () "f0"() {set = affine_set<(d0): (d0 >= 10)>} : () -> () - // CHECK: "f1"() {"set" = affine_set<(d0) : ((d0 + -10) >= 0)>} : () -> () + // CHECK: "f1"() {set = affine_set<(d0) : ((d0 + -10) >= 0)>} : () -> () "f1"() {set = affine_set<(d0): (d0 - 10 >= 0)>} : () -> () - // CHECK: "f2"() {"set" = affine_set<(d0)[s0] : (d0 >= 0, (s0 + (d0 * -1)) >= 0)>} : () -> () + // CHECK: "f2"() {set = affine_set<(d0)[s0] : (d0 >= 0, (s0 + (d0 * -1)) >= 0)>} : () -> () "f2"() {set = affine_set<(i)[N] : (i >= 0, N - i >= 0)>} : () -> () - // CHECK: "f3"() {"set" = affine_set<(d0) : (1 == 0)>} : () -> () + // CHECK: "f3"() {set = affine_set<(d0) : (1 == 0)>} : () -> () "f3"() {set = affine_set<(d0) : (1 == 0)>} : () -> () - // CHECK: "f4"() {"set" = affine_set<(d0, d1)[s0, s1] : ((((d0 * -16) + s0) + -16) >= 0, (((d1 * -3) + s1) + -3) >= 0)>} : () -> () + // CHECK: "f4"() {set = affine_set<(d0, d1)[s0, s1] : ((((d0 * -16) + s0) + -16) >= 0, (((d1 * -3) + s1) + -3) >= 0)>} : () -> () "f4"() {set = affine_set<(d0, d1)[s0, s1] : (d0 * -16 + s0 - 16 >= 0, d1 * -3 + s1 - 3 >= 0)>} : () -> () - // CHECK: "f5"() {"set" = affine_set<(d0, d1)[s0, s1] : (((((d0 * 7) + (d1 * 5)) + (s0 * 11)) + s1) == 0, ((((d0 * 5) + (d1 * -11)) + (s0 * 7)) + s1) == 0, ((((d0 * 11) + (d1 * 7)) + (s0 * -5)) + s1) == 0, ((((d0 * 7) + (d1 * 5)) + (s0 * 11)) + s1) == 0)>} : () -> () + // CHECK: "f5"() {set = affine_set<(d0, d1)[s0, s1] : (((((d0 * 7) + (d1 * 5)) + (s0 * 11)) + s1) == 0, ((((d0 * 5) + (d1 * -11)) + (s0 * 7)) + s1) == 0, ((((d0 * 11) + (d1 * 7)) + (s0 * -5)) + s1) == 0, ((((d0 * 7) + (d1 * 5)) + (s0 * 11)) + s1) == 0)>} : () -> () "f5"() {set = affine_set<(d0, d1)[s0, s1] : (d0 * 7 + d1 * 5 + s0 * 11 + s1 == 0, d0 * 5 - d1 * 11 + s0 * 7 + s1 == 0, d0 * 11 + d1 * 7 - s0 * 5 + s1 == 0, diff --git a/tests/filecheck/parser-printer/aliases.mlir b/tests/filecheck/parser-printer/aliases.mlir index 936465bdc6..6d47717880 100644 --- a/tests/filecheck/parser-printer/aliases.mlir +++ b/tests/filecheck/parser-printer/aliases.mlir @@ -9,11 +9,11 @@ "test.op"() {"attr" = #attr} : () -> () "test.op"() {"attr" = [#attr]} : () -> () -// CHECK: "test.op"() {"attr" = 1 : i32} : () -> () -// CHECK-NEXT: "test.op"() {"attr" = i32} : () -> () -// CHECK-NEXT: "test.op"() {"attr" = vector<1xi32>} : () -> () -// CHECK-NEXT: "test.op"() {"attr" = 0 : i32} : () -> () -// CHECK-NEXT: "test.op"() {"attr" = [0 : i32]} : () -> () +// CHECK: "test.op"() {attr = 1 : i32} : () -> () +// CHECK-NEXT: "test.op"() {attr = i32} : () -> () +// CHECK-NEXT: "test.op"() {attr = vector<1xi32>} : () -> () +// CHECK-NEXT: "test.op"() {attr = 0 : i32} : () -> () +// CHECK-NEXT: "test.op"() {attr = [0 : i32]} : () -> () // ----- diff --git a/tests/filecheck/parser-printer/attribute_names.mlir b/tests/filecheck/parser-printer/attribute_names.mlir index bf413017a5..de45d98865 100644 --- a/tests/filecheck/parser-printer/attribute_names.mlir +++ b/tests/filecheck/parser-printer/attribute_names.mlir @@ -3,4 +3,4 @@ builtin.module attributes {a = i32, _e = i32, "foo" = i32, _ = i32} { } -// CHECK: "a" = i32, "_e" = i32, "foo" = i32, "_" = i32 +// CHECK: a = i32, _e = i32, foo = i32, _ = i32 diff --git a/tests/filecheck/parser-printer/builtin_attrs.mlir b/tests/filecheck/parser-printer/builtin_attrs.mlir index 02488547f9..f67312a114 100644 --- a/tests/filecheck/parser-printer/builtin_attrs.mlir +++ b/tests/filecheck/parser-printer/builtin_attrs.mlir @@ -10,7 +10,7 @@ "func.func"() ({}) {function_type = () -> (), sym_name = "unit_attr_func", unitarray = [unit]} : () -> () - // CHECK: "unitarray" = [unit] + // CHECK: unitarray = [unit] "func.func"() ({ ^bb0(%arg0: i32, %arg1: i64, %arg2: i1): @@ -47,19 +47,19 @@ "func.func"() ({}) {function_type = () -> (), value = true, sym_name = "true_attr"} : () -> () - // CHECK: "value" = true + // CHECK: value = true "func.func"() ({}) {function_type = () -> (), value = 1 : i1, sym_name = "true_explicit_attr"} : () -> () - // CHECK: "value" = true + // CHECK: value = true "func.func"() ({}) {function_type = () -> (), value = false, sym_name = "false_attr"} : () -> () - // CHECK: "value" = false + // CHECK: value = false "func.func"() ({}) {function_type = () -> (), value = 0 : i1, sym_name = "false_explicit_attr"} : () -> () - // CHECK: "value" = false + // CHECK: value = false "func.func"() ({}) {function_type = () -> (), value = 42 : i32, sym_name = "int_attr"} : () -> () @@ -84,7 +84,7 @@ "func.func"() ({}) {function_type = () -> (), value = dense<"0xEEA7CC3DF47612BE2BA4173E8B75E8BDE0B915BDA3191CBE8388E0BDC826DB3DFE78273E6B037E3DEF140D3EF0B5803D4026693CD6B6E1BCE08B4DBDC3A9E63D943B163EE64E46BD808C253EB8F4893D30270CBE36696C3D045E1DBED06A703DA33EBBBD66D646BD36507BBD764D8FBD7010FA3DB6E1B53D9B83C8BDD33FA73D58AD293EB0A6123EAB2627BA40B4CB3C20E9B6BD805AB2BDE047BDBC809A743DE01ADD3D9B77D5BDCEE7043E00B8C1BDCBA80A3DBB03DA3D787C993D163968BC208510BDABFDB1BD8C07213EA34614BEAB06B73A0091413B8013B3BD768F193E7B6515BE7306833D363183BC36BC8B3CA016B7BD3E05D33DE67C28BDCABB0EBEDA2A013EA67DF6BD007EB5BA782A04BEAB69F73D16DD703D3B93A43D1BE45B3DEBAEE8BD8891F1BDF8B18F3D20EC923CE67101BE8382A8BDAB9EE7BA0006CA3AA3F224BE1B56A5BDC06B8A3DC3E6BE3D562310BB964B713C2CC11FBE4BC68F3DAEACD7BDFB093A3D00070F3EC3E4C93D5BCF0D3D1B01E13D9B7D7F3D537CD43D6BEDFBBC4BD9AEBD17BA023E569906BB86599CBD4E28073E1639F5BDF60909BE8B4727BEE4AD153EDF3C05BEB01913BEEB1A59BD03E8D4BD4BD3123D9EA381BD6058F03CD0EFF73D00747FBADBC5AEBD5054273E204DB4BD00CA683B1E28C93D3BCC2A3D9B0E683D4302923D9A3408BEABC89D3A565336BCC0A7F3BD76D1F93D68A3B93D44891C3E1685243E1B3FDBBD5E06A4BD2B4192BD2B19983C50C97B3D40A808BEC0994C3D4B3435BD0B88293D506749BDFC13063E2B7ADF3CF3B013BE"> : tensor<4x4x3x3xf32>, sym_name = "hex_f32_large_attr"} : () -> () - // CHECK: "value" = dense<[[[[0.0999296755, -0.143031895, 0.148087189], [-0.113505445, -0.0365542173, -0.152441546], [-0.109635375, 0.107007563, 0.163547486]], [[0.0620149784, 0.137775168, 0.0628470182], [0.0142303109, -0.0275530033, -0.0501822233], [0.112628482, 0.146711648, -0.0484150872]], [[0.161668777, 0.0673612952, -0.136868238], [0.0577175245, -0.153678954, 0.0586956143], [-0.0914280638, -0.04854431, -0.061355792]], [[-0.0699719638, 0.122101665, 0.0888094157], [-0.0979072675, 0.0816647038, 0.165700316], [0.143213987, -0.000637630641, 0.0248662233]]], [[[-0.0893118382, -0.0870866776, -0.0231055617], [0.0597176552, 0.107961416, -0.104232036], [0.129790515, -0.0945892334, 0.0338523798]], [[0.106452428, 0.0749444366, -0.0141737666], [-0.0352832079, -0.0869096145, 0.157255352], [-0.144800708, 0.00139637792, 0.00295358896]], [[-0.087439537, 0.149961323, -0.14589493], [0.0639771447, -0.0160146765, 0.0170575194], [-0.0893986225, 0.103037342, -0.0411347374]], [[-0.139388233, 0.126140028, -0.120356843], [-0.0013846755, -0.129068255, 0.120807014], [0.058804594, 0.0803589448, 0.0536843352]]], [[[-0.11361488, -0.11795336, 0.0701636672], [0.0179348588, -0.126411051, -0.0822801813], [-0.00176711881, 0.00154131651, -0.161081836]], [[-0.0807306394, 0.0675883293, 0.0932135805], [-0.00219937181, 0.0147274937, -0.15601033], [0.0702024326, -0.105309829, 0.0454196744]], [[0.13967514, 0.0985808596, 0.0346215777], [0.10986539, 0.0623756461, 0.103752755], [-0.0307528581, -0.0853753909, 1.276630e-01]], [[-0.00205381727, -0.0763426274, 0.131989688], [-0.119737789, -0.13382706, -0.163358852], [0.146171153, -0.130115017, -0.143652678]]], [[[-0.0530041866, -0.103958152, 0.0358460359], [-0.0633003563, 0.0293390155, 0.121062875], [-0.000974476337, -0.0853383169, 0.163407564]], [[-0.0880377293, 0.0035520792, 0.0982210487], [0.0416986756, 0.0566545539, 0.0712933764], [-0.133013159, 0.00120379531, -0.0111282673]], [[-0.118972301, 0.121981546, 0.0906437039], [0.152867377, 0.160663933, -0.107053958], [-0.0800902694, -0.0714133605, 0.0185666885]], [[0.0614712834, -0.133454323, 0.0499513149], [-0.0442393236, 0.0413895063, -0.0491707921], [0.130935609, 0.0272799339, -0.144229695]]]]> : tensor<4x4x3x3xf32>} : () -> () + // CHECK: value = dense<[[[[0.0999296755, -0.143031895, 0.148087189], [-0.113505445, -0.0365542173, -0.152441546], [-0.109635375, 0.107007563, 0.163547486]], [[0.0620149784, 0.137775168, 0.0628470182], [0.0142303109, -0.0275530033, -0.0501822233], [0.112628482, 0.146711648, -0.0484150872]], [[0.161668777, 0.0673612952, -0.136868238], [0.0577175245, -0.153678954, 0.0586956143], [-0.0914280638, -0.04854431, -0.061355792]], [[-0.0699719638, 0.122101665, 0.0888094157], [-0.0979072675, 0.0816647038, 0.165700316], [0.143213987, -0.000637630641, 0.0248662233]]], [[[-0.0893118382, -0.0870866776, -0.0231055617], [0.0597176552, 0.107961416, -0.104232036], [0.129790515, -0.0945892334, 0.0338523798]], [[0.106452428, 0.0749444366, -0.0141737666], [-0.0352832079, -0.0869096145, 0.157255352], [-0.144800708, 0.00139637792, 0.00295358896]], [[-0.087439537, 0.149961323, -0.14589493], [0.0639771447, -0.0160146765, 0.0170575194], [-0.0893986225, 0.103037342, -0.0411347374]], [[-0.139388233, 0.126140028, -0.120356843], [-0.0013846755, -0.129068255, 0.120807014], [0.058804594, 0.0803589448, 0.0536843352]]], [[[-0.11361488, -0.11795336, 0.0701636672], [0.0179348588, -0.126411051, -0.0822801813], [-0.00176711881, 0.00154131651, -0.161081836]], [[-0.0807306394, 0.0675883293, 0.0932135805], [-0.00219937181, 0.0147274937, -0.15601033], [0.0702024326, -0.105309829, 0.0454196744]], [[0.13967514, 0.0985808596, 0.0346215777], [0.10986539, 0.0623756461, 0.103752755], [-0.0307528581, -0.0853753909, 1.276630e-01]], [[-0.00205381727, -0.0763426274, 0.131989688], [-0.119737789, -0.13382706, -0.163358852], [0.146171153, -0.130115017, -0.143652678]]], [[[-0.0530041866, -0.103958152, 0.0358460359], [-0.0633003563, 0.0293390155, 0.121062875], [-0.000974476337, -0.0853383169, 0.163407564]], [[-0.0880377293, 0.0035520792, 0.0982210487], [0.0416986756, 0.0566545539, 0.0712933764], [-0.133013159, 0.00120379531, -0.0111282673]], [[-0.118972301, 0.121981546, 0.0906437039], [0.152867377, 0.160663933, -0.107053958], [-0.0800902694, -0.0714133605, 0.0185666885]], [[0.0614712834, -0.133454323, 0.0499513149], [-0.0442393236, 0.0413895063, -0.0491707921], [0.130935609, 0.0272799339, -0.144229695]]]]> : tensor<4x4x3x3xf32>} : () -> () "func.func"() ({}) {function_type = () -> (), value = "foo", sym_name = "string_attr"} : () -> () @@ -123,77 +123,77 @@ value2 = dense<[0.0, 1.0]> : tensor<2xf64>, sym_name = "dense_tensor_attr"} : () -> () - // CHECK{LITERAL}: "value1" = dense<[[2, 3]]> : tensor<1x2xi32>, "value2" = dense<[0.000000e+00, 1.000000e+00]> : tensor<2xf64> + // CHECK{LITERAL}: value1 = dense<[[2, 3]]> : tensor<1x2xi32>, value2 = dense<[0.000000e+00, 1.000000e+00]> : tensor<2xf64> "func.func"() ({}) {function_type = () -> (), value1 = dense<"0xFF00"> : tensor<2xi8>, value2 = dense<"0xFF00FF00"> : tensor<1xi32>, value3 = dense<"0xCAFEBABE"> : tensor<2xi32>, sym_name = "dense_tensor_attr_hex"} : () -> () - // CHECK: "value1" = dense<[-1, 0]> : tensor<2xi8>, "value2" = dense<16711935> : tensor<1xi32>, "value3" = dense<-1095041334> : tensor<2xi32> + // CHECK: value1 = dense<[-1, 0]> : tensor<2xi8>, value2 = dense<16711935> : tensor<1xi32>, value3 = dense<-1095041334> : tensor<2xi32> "func.func"() ({}) {function_type = () -> (), value1 = dense<"0xFF00CAFE"> : tensor<2x2xi8>, sym_name = "dense_tensor_attr_hex_long"} : () -> () - // CHECK{LITERAL}: "value1" = dense<[[-1, 0], [-54, -2]]> : tensor<2x2xi8> + // CHECK{LITERAL}: value1 = dense<[[-1, 0], [-54, -2]]> : tensor<2x2xi8> "func.func"() ({}) {function_type = () -> (), value1 = dense<"0xCAFEBABE"> : tensor<2xf32>, value2 = dense<"0xCAFEBABEB00BAABE"> : tensor<1xf64>, sym_name = "dense_tensor_attr_hex_float"} : () -> () - // CHECK: "value1" = dense<-0.365225136> : tensor<2xf32>, "value2" = dense<-7.7622132495927025e-07> : tensor<1xf64> + // CHECK: value1 = dense<-0.365225136> : tensor<2xf32>, value2 = dense<-7.7622132495927025e-07> : tensor<1xf64> "func.func"() ({}) {function_type = () -> (), value1 = dense<[0]> : vector<1xi32>, value2 = dense<[0.0, 1.0]> : vector<2xf64>, sym_name = "dense_vector_attr"} : () -> () - // CHECK: "value1" = dense<0> : vector<1xi32>, "value2" = dense<[0.000000e+00, 1.000000e+00]> : vector<2xf64> + // CHECK: value1 = dense<0> : vector<1xi32>, value2 = dense<[0.000000e+00, 1.000000e+00]> : vector<2xf64> "func.func"() ({}) {function_type = () -> (), value1 = dense<> : tensor<1x23x0x4xi32>, value2 = dense<[[0.0], [1.0]]> : tensor<2x1xf64>, sym_name = "dense_corner_attr"} : () -> () - // CHECK{LITERAL}: "value1" = dense<> : tensor<1x23x0x4xi32>, "value2" = dense<[[0.000000e+00], [1.000000e+00]]> : tensor<2x1xf64> + // CHECK{LITERAL}: value1 = dense<> : tensor<1x23x0x4xi32>, value2 = dense<[[0.000000e+00], [1.000000e+00]]> : tensor<2x1xf64> "func.func"() ({}) {function_type = () -> (), value1 = dense<> : tensor<1x23x0x4xi32>, value2 = dense<[[-0.0], [-1.0]]> : tensor<2x1xf64>, sym_name = "dense_negative_attr"} : () -> () - // CHECK{LITERAL}: "value1" = dense<> : tensor<1x23x0x4xi32>, "value2" = dense<[[-0.000000e+00], [-1.000000e+00]]> : tensor<2x1xf64> + // CHECK{LITERAL}: value1 = dense<> : tensor<1x23x0x4xi32>, value2 = dense<[[-0.000000e+00], [-1.000000e+00]]> : tensor<2x1xf64> "func.func"() ({}) {function_type = () -> (), value1 = dense<12> : tensor<2x3xi32>, sym_name = "dense_trivial_attr"} : () -> () - // CHECK: "value1" = dense<12> : tensor<2x3xi32> + // CHECK: value1 = dense<12> : tensor<2x3xi32> "func.func"() ({}) {function_type = () -> (), value1 = dense<[true, false]> : tensor<2xi1>, sym_name = "dense_bool_attr"} : () -> () - // CHECK: "value1" = dense<[true, false]> : tensor<2xi1> + // CHECK: value1 = dense<[true, false]> : tensor<2xi1> "func.func"() ({}) {function_type = () -> (), value1 = opaque<"test", "contents">, value2 = opaque<"test", "contents"> : tensor<2xf64>, sym_name = "opaque_attr"} : () -> () - // CHECK: "value1" = opaque<"test", "contents">, "value2" = opaque<"test", "contents"> : tensor<2xf64> + // CHECK: value1 = opaque<"test", "contents">, value2 = opaque<"test", "contents"> : tensor<2xf64> "func.func"() ({}) {function_type = () -> (), value = {"one"=1, "two"=2, "three"="three"}, sym_name = "dict_attr"} : () -> () - // CHECK: "one" = 1 : i64, "two" = 2 : i64, "three" = "three" + // CHECK: one = 1 : i64, two = 2 : i64, three = "three" "func.func"() ({}) {function_type = () -> (), symbol = @some_symbol, sym_name = "symbol_ref_attr"} : () -> () - // CHECK: "symbol" = @some_symbol + // CHECK: symbol = @some_symbol "func.func"() ({}) {function_type = () -> (), value1 = tensor, @@ -266,56 +266,56 @@ type_attr = index, sym_name = "index_type"} : () -> () - // CHECK: "type_attr" = index + // CHECK: type_attr = index "func.func"() ({}) {function_type = () -> (), strided = strided<[1, 0x23, -23, -0x21, ?], offset: -3>, sym_name = "strided"} : () -> () - // CHECK: "strided" = strided<[1, 35, -23, -33, ?], offset: -3> + // CHECK: strided = strided<[1, 35, -23, -33, ?], offset: -3> "func.func"() ({}) {function_type = () -> (), strided = strided<[], offset: ?>, sym_name = "what_strided"} : () -> () - // CHECK: "strided" = strided<[], offset: ?> + // CHECK: strided = strided<[], offset: ?> "func.func"() ({}) {function_type = () -> (), strided = strided<[], offset: 0>, sym_name = "trivial_strided"} : () -> () - // CHECK: "strided" = strided<[]> + // CHECK: strided = strided<[]> "func.func"() ({}) {function_type = () -> (), strided = strided<[]>, sym_name = "empty_strided"} : () -> () - // CHECK: "strided" = strided<[]> + // CHECK: strided = strided<[]> "func.func"() ({}) {function_type = () -> (), complex = complex, sym_name = "complex_i32"} : () -> () - // CHECK: "complex" = complex + // CHECK: complex = complex "func.func"() ({}) {function_type = () -> (), complex = complex, sym_name = "complex_f32"} : () -> () - // CHECK: "complex" = complex + // CHECK: complex = complex "func.func"() ({}) {function_type = () -> (), function = () -> i32, sym_name = "one_to_one_func"} : () -> () - // CHECK: "function" = () -> i32 + // CHECK: function = () -> i32 "func.func"() ({}) {function_type = () -> (), function = (i1) -> (i32), sym_name = "one_to_one_func_paren"} : () -> () - // CHECK: "function" = (i1) -> i32 + // CHECK: function = (i1) -> i32 "func.func"() ({}) {function_type = () -> (), function = (i1, i2) -> (i32, i64), sym_name = "two_to_two_func"} : () -> () - // CHECK: "function" = (i1, i2) -> (i32, i64) + // CHECK: function = (i1, i2) -> (i32, i64) "func.func"() ({}) {function_type = () -> (), function = () -> (() -> i32), sym_name = "higher_order_func"} : () -> () - // CHECK: "function" = () -> (() -> i32) + // CHECK: function = () -> (() -> i32) }) : () -> () diff --git a/tests/filecheck/parser-printer/escaped_characters.mlir b/tests/filecheck/parser-printer/escaped_characters.mlir index fe7654daab..343eb596fd 100644 --- a/tests/filecheck/parser-printer/escaped_characters.mlir +++ b/tests/filecheck/parser-printer/escaped_characters.mlir @@ -3,18 +3,18 @@ "builtin.module"() ({ "test.op"() {name = "\""} : () -> () - // CHECK: "test.op"() {"name" = "\""} : () -> () + // CHECK: "test.op"() {name = "\""} : () -> () "test.op"() {name = "\n"} : () -> () - // CHECK-NEXT: "test.op"() {"name" = "\n"} : () -> () + // CHECK-NEXT: "test.op"() {name = "\n"} : () -> () "test.op"() {name = "\t"} : () -> () - // CHECK-NEXT: "test.op"() {"name" = "\t"} : () -> () + // CHECK-NEXT: "test.op"() {name = "\t"} : () -> () "test.op"() {name = "\\"} : () -> () - // CHECK-NEXT: "test.op"() {"name" = "\\"} : () -> () + // CHECK-NEXT: "test.op"() {name = "\\"} : () -> () "test.op"() {name = "\E2\9A\A0\EF\B8\8FP\EDU\BA\01\00\10\00\A0\11\00\00\00\00\00\00\02\00\01\01@\00\00\00(\0D\00\00\00\00\00\00\00\00\00\00\00\00\00\00\07\00\01\00=\00\00\00\00\00\00\00\00\00\00\00\11\00\00\00\00\00\00\00\00\00\00\00\00\00\00\00\00\00\00\00\00\00\00\00\7FELF\02\01\013\07\00\00\00\00\00\00\00\02\00\BE\00x\00\00\00\00\00\00\00\00\00\00\00\80\0C\00\00\00\00\00\00\80\0A\00\00\00\00\00\00=\05=\00@\008\00\03\00@\00\08\00\01\00\00.shstrtab\00.strtab\00.symtab\00.symtab_shndx\00.nv.uft.entry\00.nv.info\00.text.gpu_kernel_kernel"} : () -> () - // CHECK-NEXT: "test.op"() {"name" = "\E2\9A\A0\EF\B8\8FP\EDU\BA\01\00\10\00\A0\11\00\00\00\00\00\00\02\00\01\01@\00\00\00(\0D\00\00\00\00\00\00\00\00\00\00\00\00\00\00\07\00\01\00=\00\00\00\00\00\00\00\00\00\00\00\11\00\00\00\00\00\00\00\00\00\00\00\00\00\00\00\00\00\00\00\00\00\00\00\7FELF\02\01\013\07\00\00\00\00\00\00\00\02\00\BE\00x\00\00\00\00\00\00\00\00\00\00\00\80\0C\00\00\00\00\00\00\80\0A\00\00\00\00\00\00=\05=\00@\008\00\03\00@\00\08\00\01\00\00.shstrtab\00.strtab\00.symtab\00.symtab_shndx\00.nv.uft.entry\00.nv.info\00.text.gpu_kernel_kernel"} : () -> () + // CHECK-NEXT: "test.op"() {name = "\E2\9A\A0\EF\B8\8FP\EDU\BA\01\00\10\00\A0\11\00\00\00\00\00\00\02\00\01\01@\00\00\00(\0D\00\00\00\00\00\00\00\00\00\00\00\00\00\00\07\00\01\00=\00\00\00\00\00\00\00\00\00\00\00\11\00\00\00\00\00\00\00\00\00\00\00\00\00\00\00\00\00\00\00\00\00\00\00\7FELF\02\01\013\07\00\00\00\00\00\00\00\02\00\BE\00x\00\00\00\00\00\00\00\00\00\00\00\80\0C\00\00\00\00\00\00\80\0A\00\00\00\00\00\00=\05=\00@\008\00\03\00@\00\08\00\01\00\00.shstrtab\00.strtab\00.symtab\00.symtab_shndx\00.nv.uft.entry\00.nv.info\00.text.gpu_kernel_kernel"} : () -> () }) : () -> () diff --git a/tests/filecheck/parser-printer/float_parsing.mlir b/tests/filecheck/parser-printer/float_parsing.mlir index 7db0ec4700..58d7c26e0e 100644 --- a/tests/filecheck/parser-printer/float_parsing.mlir +++ b/tests/filecheck/parser-printer/float_parsing.mlir @@ -3,24 +3,24 @@ "builtin.module"() ({ "test.op"() {"value" = 42.0 : f32} : () -> () - // CHECK: "test.op"() {"value" = 4.200000e+01 : f32} : () -> () + // CHECK: "test.op"() {value = 4.200000e+01 : f32} : () -> () "test.op"() {"value" = -42.0 : f32} : () -> () - // CHECK-NEXT: "test.op"() {"value" = -4.200000e+01 : f32} : () -> () + // CHECK-NEXT: "test.op"() {value = -4.200000e+01 : f32} : () -> () "test.op"() {"value" = 34.e0 : f32} : () -> () - // CHECK-NEXT: "test.op"() {"value" = 3.400000e+01 : f32} : () -> () + // CHECK-NEXT: "test.op"() {value = 3.400000e+01 : f32} : () -> () "test.op"() {"value" = 34.e-23 : f32} : () -> () - // CHECK-NEXT: "test.op"() {"value" = 3.400000e-22 : f32} : () -> () + // CHECK-NEXT: "test.op"() {value = 3.400000e-22 : f32} : () -> () "test.op"() {"value" = 34.e12 : f32} : () -> () - // CHECK-NEXT: "test.op"() {"value" = 3.400000e+13 : f32} : () -> () + // CHECK-NEXT: "test.op"() {value = 3.400000e+13 : f32} : () -> () "test.op"() {"value" = -34.e-12 : f32} : () -> () - // CHECK-NEXT: "test.op"() {"value" = -3.400000e-11 : f32} : () -> () + // CHECK-NEXT: "test.op"() {value = -3.400000e-11 : f32} : () -> () // this should print in full precision "test.op"() {"value" = 3.141592653589793 : f64} : () -> () - // CHECK-NEXT: "test.op"() {"value" = 3.1415926535897931 : f64} : () -> () + // CHECK-NEXT: "test.op"() {value = 3.1415926535897931 : f64} : () -> () }) : () -> () diff --git a/tests/filecheck/parser-printer/operation_with_properties.mlir b/tests/filecheck/parser-printer/operation_with_properties.mlir index 1b1db6e8a1..9d24bf917a 100644 --- a/tests/filecheck/parser-printer/operation_with_properties.mlir +++ b/tests/filecheck/parser-printer/operation_with_properties.mlir @@ -4,15 +4,15 @@ builtin.module { // An operation with a property "unregistered.op"() <{"test" = 2 : i32}> : () -> () - // CHECK: "unregistered.op"() <{"test" = 2 : i32}> : () -> () + // CHECK: "unregistered.op"() <{test = 2 : i32}> : () -> () // An operation with a property, a region, and an attribute "unregistered.op"() <{"test" = 42 : i64, "test2" = 71 : i32}> ({}) {"test3" = "foo"} : () -> () - // CHECK-NEXT: "unregistered.op"() <{"test" = 42 : i64, "test2" = 71 : i32}> ({ - // CHECK-NEXT: }) {"test3" = "foo"} : () -> () + // CHECK-NEXT: "unregistered.op"() <{test = 42 : i64, test2 = 71 : i32}> ({ + // CHECK-NEXT: }) {test3 = "foo"} : () -> () // An operation with a property and an attribute with the same name "unregistered.op"() <{"test" = 42 : i64}> {"test" = "foo"} : () -> () - // CHECK-NEXT: "unregistered.op"() <{"test" = 42 : i64}> {"test" = "foo"} : () -> () + // CHECK-NEXT: "unregistered.op"() <{test = 42 : i64}> {test = "foo"} : () -> () } diff --git a/tests/filecheck/parser-printer/scope.mlir b/tests/filecheck/parser-printer/scope.mlir index 53c796f120..352f76ff60 100644 --- a/tests/filecheck/parser-printer/scope.mlir +++ b/tests/filecheck/parser-printer/scope.mlir @@ -14,7 +14,7 @@ module { // CHECK-NEXT: } // CHECK-GENERIC: "builtin.module"() ({ -// CHECK-GENERIC-NEXT: "func.func"() <{"sym_name" = "my_func", "function_type" = () -> (), "sym_visibility" = "public"}> ({ +// CHECK-GENERIC-NEXT: "func.func"() <{sym_name = "my_func", function_type = () -> (), sym_visibility = "public"}> ({ // CHECK-GENERIC-NEXT: "func.return"() : () -> () // CHECK-GENERIC-NEXT: }) : () -> () // CHECK-GENERIC-NEXT: }) : () -> () diff --git a/tests/filecheck/parser-printer/unregistered_dialect.mlir b/tests/filecheck/parser-printer/unregistered_dialect.mlir index 318ea5d062..7d776654e6 100644 --- a/tests/filecheck/parser-printer/unregistered_dialect.mlir +++ b/tests/filecheck/parser-printer/unregistered_dialect.mlir @@ -12,11 +12,11 @@ // CHECK: %{{.*}} = "region_op"() ({ - // CHECK-NEXT: %{{.*}} = "op_with_res"() {"otherattr" = #unknowndialect.unknown_attr>>} : () -> i32 - // CHECK-NEXT: %{{.*}} = "op_with_res"() {"otherattr" = #unknowndialect>>} : () -> i32 + // CHECK-NEXT: %{{.*}} = "op_with_res"() {otherattr = #unknowndialect.unknown_attr>>} : () -> i32 + // CHECK-NEXT: %{{.*}} = "op_with_res"() {otherattr = #unknowndialect>>} : () -> i32 // CHECK-NEXT: %{{.*}} = "op_with_operands"(%{{.*}}, %{{.*}}) : (i32, i32) -> !unknowndialect.unknown_type<{[<()>]}> - // CHECK-NEXT: "op"() {"ab" = !unknowndialect.unknown_singleton_type} : () -> () - // CHECK-NEXT: }) {"testattr" = "foo"} : () -> i32 - // CHECK-NEXT: "builtin.unimplemented_op"() {"attr" = #builtin.unimplemented_attr} : () -> () + // CHECK-NEXT: "op"() {ab = !unknowndialect.unknown_singleton_type} : () -> () + // CHECK-NEXT: }) {testattr = "foo"} : () -> i32 + // CHECK-NEXT: "builtin.unimplemented_op"() {attr = #builtin.unimplemented_attr} : () -> () }) : () -> () diff --git a/tests/filecheck/transforms/apply-pdl/apply_pdl_extra_file.mlir b/tests/filecheck/transforms/apply-pdl/apply_pdl_extra_file.mlir index b9ff3f6229..d17227413a 100644 --- a/tests/filecheck/transforms/apply-pdl/apply_pdl_extra_file.mlir +++ b/tests/filecheck/transforms/apply-pdl/apply_pdl_extra_file.mlir @@ -3,5 +3,5 @@ "test.op"() {attr = 0} : () -> () //CHECK: builtin.module { -// CHECK-NEXT: "test.op"() {"attr" = 1 : i64} : () -> () +// CHECK-NEXT: "test.op"() {attr = 1 : i64} : () -> () // CHECK-NEXT: } diff --git a/tests/filecheck/transforms/apply-pdl/apply_pdl_property_rewrite.mlir b/tests/filecheck/transforms/apply-pdl/apply_pdl_property_rewrite.mlir index 14bc3b9e65..f25d7e6fcb 100644 --- a/tests/filecheck/transforms/apply-pdl/apply_pdl_property_rewrite.mlir +++ b/tests/filecheck/transforms/apply-pdl/apply_pdl_property_rewrite.mlir @@ -1,6 +1,6 @@ // RUN: xdsl-opt %s -p apply-pdl | filecheck %s -// CHECK: "test.op"() <{"prop1" = i64}> : () -> () +// CHECK: "test.op"() <{prop1 = i64}> : () -> () "test.op"() <{"prop1" = i32}> : () -> () pdl.pattern : benefit(42) { diff --git a/tests/filecheck/transforms/apply-pdl/apply_pdl_simple.mlir b/tests/filecheck/transforms/apply-pdl/apply_pdl_simple.mlir index c88535b705..118aa5e152 100644 --- a/tests/filecheck/transforms/apply-pdl/apply_pdl_simple.mlir +++ b/tests/filecheck/transforms/apply-pdl/apply_pdl_simple.mlir @@ -13,7 +13,7 @@ pdl.pattern : benefit(1) { } //CHECK: builtin.module { -// CHECK-NEXT: "test.op"() {"attr" = 1 : i64} : () -> () +// CHECK-NEXT: "test.op"() {attr = 1 : i64} : () -> () // CHECK-NEXT: pdl.pattern : benefit(1) { // CHECK-NEXT: %zero_attr = pdl.attribute = 0 : i64 // CHECK-NEXT: %root = pdl.operation "test.op" {"attr" = %zero_attr} diff --git a/tests/filecheck/transforms/convert-scf-to-openmp.mlir b/tests/filecheck/transforms/convert-scf-to-openmp.mlir index 1c9c1c18c8..1d3b570d0f 100644 --- a/tests/filecheck/transforms/convert-scf-to-openmp.mlir +++ b/tests/filecheck/transforms/convert-scf-to-openmp.mlir @@ -18,8 +18,8 @@ builtin.module { // Check the default lowering. // CHECK: func.func @parallel(%{{.*}} : index, %{{.*}} : index, %{{.*}} : index, %{{.*}} : index, %{{.*}} : index, %{{.*}} : index) { -// CHECK-NEXT: "omp.parallel"() <{"operandSegmentSizes" = array}> ({ -// CHECK-NEXT: "omp.wsloop"() <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "omp.parallel"() <{operandSegmentSizes = array}> ({ +// CHECK-NEXT: "omp.wsloop"() <{operandSegmentSizes = array}> ({ // CHECK-NEXT: "omp.loop_nest"(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) ({ // CHECK-NEXT: ^{{.*}}(%{{.*}} : index, %{{.*}} : index): // CHECK-NEXT: "memref.alloca_scope"() ({ @@ -38,12 +38,12 @@ builtin.module { // Check that using `collapse=1` converts only the first dimension to OpenMP, and keeps the // inner one(s) as an `scf.parallel` for any other further conversion. // COLLAPSE: func.func @parallel(%{{.*}} : index, %{{.*}} : index, %{{.*}} : index, %{{.*}} : index, %{{.*}} : index, %{{.*}} : index) { -// COLLAPSE-NEXT: "omp.parallel"() <{"operandSegmentSizes" = array}> ({ -// COLLAPSE-NEXT: "omp.wsloop"() <{"operandSegmentSizes" = array}> ({ +// COLLAPSE-NEXT: "omp.parallel"() <{operandSegmentSizes = array}> ({ +// COLLAPSE-NEXT: "omp.wsloop"() <{operandSegmentSizes = array}> ({ // COLLAPSE-NEXT: "omp.loop_nest"(%{{.*}}, %{{.*}}, %{{.*}}) ({ // COLLAPSE-NEXT: ^{{.*}}(%{{.*}} : index): // COLLAPSE-NEXT: "memref.alloca_scope"() ({ -// COLLAPSE-NEXT: "scf.parallel"(%{{.*}}, %{{.*}}, %{{.*}}) <{"operandSegmentSizes" = array}> ({ +// COLLAPSE-NEXT: "scf.parallel"(%{{.*}}, %{{.*}}, %{{.*}}) <{operandSegmentSizes = array}> ({ // COLLAPSE-NEXT: ^{{.*}}(%{{.*}} : index): // COLLAPSE-NEXT: "test.op"(%{{.*}}, %{{.*}}) : (index, index) -> () // COLLAPSE-NEXT: scf.reduce @@ -61,8 +61,8 @@ builtin.module { // Check that using `schedule` does set the OpenMP loop's schedule // DYNAMIC: func.func @parallel(%{{.*}} : index, %{{.*}} : index, %{{.*}} : index, %{{.*}} : index, %{{.*}} : index, %{{.*}} : index) { -// DYNAMIC-NEXT: "omp.parallel"() <{"operandSegmentSizes" = array}> ({ -// DYNAMIC-NEXT: "omp.wsloop"() <{"operandSegmentSizes" = array, "schedule_val" = #omp}> ({ +// DYNAMIC-NEXT: "omp.parallel"() <{operandSegmentSizes = array}> ({ +// DYNAMIC-NEXT: "omp.wsloop"() <{operandSegmentSizes = array, schedule_val = #omp}> ({ // DYNAMIC-NEXT: "omp.loop_nest"(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) ({ // DYNAMIC-NEXT: ^{{.*}}(%{{.*}} : index, %{{.*}} : index): // DYNAMIC-NEXT: "memref.alloca_scope"() ({ @@ -82,9 +82,9 @@ builtin.module { // Also, check that doing so without selecting a scheule sets it to static. // (It is invalid to set a chunk size without setting a schedule) // CHUNK: func.func @parallel(%{{.*}} : index, %{{.*}} : index, %{{.*}} : index, %{{.*}} : index, %{{.*}} : index, %{{.*}} : index) { -// CHUNK-NEXT: "omp.parallel"() <{"operandSegmentSizes" = array}> ({ +// CHUNK-NEXT: "omp.parallel"() <{operandSegmentSizes = array}> ({ // CHUNK-NEXT: %{{.*}} = arith.constant 4 : index -// CHUNK-NEXT: "omp.wsloop"(%{{.*}}) <{"operandSegmentSizes" = array, "schedule_val" = #omp}> ({ +// CHUNK-NEXT: "omp.wsloop"(%{{.*}}) <{operandSegmentSizes = array, schedule_val = #omp}> ({ // CHUNK-NEXT: "omp.loop_nest"(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) ({ // CHUNK-NEXT: ^{{.*}}(%{{.*}} : index, %{{.*}} : index): // CHUNK-NEXT: "memref.alloca_scope"() ({ @@ -115,12 +115,12 @@ builtin.module { // Check that the default conversion does not convert the nested loop. // CHECK-NEXT: func.func @nested_loops(%{{.*}} : index, %{{.*}} : index, %{{.*}} : index, %{{.*}} : index, %{{.*}} : index, %{{.*}} : index) { -// CHECK-NEXT: "omp.parallel"() <{"operandSegmentSizes" = array}> ({ -// CHECK-NEXT: "omp.wsloop"() <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "omp.parallel"() <{operandSegmentSizes = array}> ({ +// CHECK-NEXT: "omp.wsloop"() <{operandSegmentSizes = array}> ({ // CHECK-NEXT: "omp.loop_nest"(%{{.*}}, %{{.*}}, %{{.*}}) ({ // CHECK-NEXT: ^{{.*}}(%{{.*}} : index): // CHECK-NEXT: "memref.alloca_scope"() ({ -// CHECK-NEXT: "scf.parallel"(%{{.*}}, %{{.*}}, %{{.*}}) <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "scf.parallel"(%{{.*}}, %{{.*}}, %{{.*}}) <{operandSegmentSizes = array}> ({ // CHECK-NEXT: ^{{.*}}(%{{.*}} : index): // CHECK-NEXT: "test.op"(%{{.*}}, %{{.*}}) : (index, index) -> () // CHECK-NEXT: scf.reduce @@ -138,13 +138,13 @@ builtin.module { // Check that using `nested=true` allows to lower the nested loop. // NESTED: func.func @nested_loops(%{{.*}} : index, %{{.*}} : index, %{{.*}} : index, %{{.*}} : index, %{{.*}} : index, %{{.*}} : index) { -// NESTED-NEXT: "omp.parallel"() <{"operandSegmentSizes" = array}> ({ -// NESTED-NEXT: "omp.wsloop"() <{"operandSegmentSizes" = array}> ({ +// NESTED-NEXT: "omp.parallel"() <{operandSegmentSizes = array}> ({ +// NESTED-NEXT: "omp.wsloop"() <{operandSegmentSizes = array}> ({ // NESTED-NEXT: "omp.loop_nest"(%{{.*}}, %{{.*}}, %{{.*}}) ({ // NESTED-NEXT: ^{{.*}}(%{{.*}} : index): // NESTED-NEXT: "memref.alloca_scope"() ({ -// NESTED-NEXT: "omp.parallel"() <{"operandSegmentSizes" = array}> ({ -// NESTED-NEXT: "omp.wsloop"() <{"operandSegmentSizes" = array}> ({ +// NESTED-NEXT: "omp.parallel"() <{operandSegmentSizes = array}> ({ +// NESTED-NEXT: "omp.wsloop"() <{operandSegmentSizes = array}> ({ // NESTED-NEXT: "omp.loop_nest"(%{{.*}}, %{{.*}}, %{{.*}}) ({ // NESTED-NEXT: ^{{.*}}(%{{.*}} : index): // NESTED-NEXT: "memref.alloca_scope"() ({ @@ -184,8 +184,8 @@ builtin.module { // Just another example, copied from MLIR's filecheck. // CHECK: func.func @adjacent_loops(%{{.*}} : index, %{{.*}} : index, %{{.*}} : index, %{{.*}} : index, %{{.*}} : index, %{{.*}} : index) { -// CHECK-NEXT: "omp.parallel"() <{"operandSegmentSizes" = array}> ({ -// CHECK-NEXT: "omp.wsloop"() <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "omp.parallel"() <{operandSegmentSizes = array}> ({ +// CHECK-NEXT: "omp.wsloop"() <{operandSegmentSizes = array}> ({ // CHECK-NEXT: "omp.loop_nest"(%{{.*}}, %{{.*}}, %{{.*}}) ({ // CHECK-NEXT: ^{{.*}}(%{{.*}} : index): // CHECK-NEXT: "memref.alloca_scope"() ({ @@ -198,8 +198,8 @@ builtin.module { // CHECK-NEXT: }) : () -> () // CHECK-NEXT: "omp.terminator"() : () -> () // CHECK-NEXT: }) : () -> () -// CHECK-NEXT: "omp.parallel"() <{"operandSegmentSizes" = array}> ({ -// CHECK-NEXT: "omp.wsloop"() <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "omp.parallel"() <{operandSegmentSizes = array}> ({ +// CHECK-NEXT: "omp.wsloop"() <{operandSegmentSizes = array}> ({ // CHECK-NEXT: "omp.loop_nest"(%{{.*}}, %{{.*}}, %{{.*}}) ({ // CHECK-NEXT: ^{{.*}}(%{{.*}} : index): // CHECK-NEXT: "memref.alloca_scope"() ({ @@ -234,7 +234,7 @@ builtin.module { // CHECK: func.func @reduction1(%{{.*}} : index, %{{.*}} : index, %{{.*}} : index, %{{.*}} : index, %{{.*}} : index) { // CHECK-NEXT: %{{.*}} = arith.constant 1 : index // CHECK-NEXT: %{{.*}} = arith.constant 0.000000e+00 : f32 -// CHECK-NEXT: %{{.*}} = "scf.parallel"(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: %{{.*}} = "scf.parallel"(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) <{operandSegmentSizes = array}> ({ // CHECK-NEXT: ^{{.*}}(%{{.*}} : index, %{{.*}} : index): // CHECK-NEXT: %{{.*}} = arith.constant 1.000000e+00 : f32 // CHECK-NEXT: scf.reduce(%{{.*}} : f32) { diff --git a/tests/filecheck/transforms/convert-stencil-to-csl-stencil.mlir b/tests/filecheck/transforms/convert-stencil-to-csl-stencil.mlir index e550a04561..8126f45bbe 100644 --- a/tests/filecheck/transforms/convert-stencil-to-csl-stencil.mlir +++ b/tests/filecheck/transforms/convert-stencil-to-csl-stencil.mlir @@ -37,7 +37,7 @@ builtin.module { // CHECK-NEXT: func.func @gauss_seidel(%a : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>, %b : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>) { // CHECK-NEXT: %0 = stencil.load %a : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>> -> !stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>> // CHECK-NEXT: %1 = tensor.empty() : tensor<510xf32> -// CHECK-NEXT: %2 = csl_stencil.apply(%0 : !stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>>, %1 : tensor<510xf32>) -> (!stencil.temp<[0,1]x[0,1]xtensor<510xf32>>) <{"swaps" = [#csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange], "topo" = #dmp.topo<1022x510>, "num_chunks" = 2 : i64, "operandSegmentSizes" = array}> ({ +// CHECK-NEXT: %2 = csl_stencil.apply(%0 : !stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>>, %1 : tensor<510xf32>) -> (!stencil.temp<[0,1]x[0,1]xtensor<510xf32>>) <{swaps = [#csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange], topo = #dmp.topo<1022x510>, num_chunks = 2 : i64, operandSegmentSizes = array}> ({ // CHECK-NEXT: ^0(%3 : tensor<4x255xf32>, %4 : index, %5 : tensor<510xf32>): // CHECK-NEXT: %6 = arith.constant 1.666600e-01 : f32 // CHECK-NEXT: %7 = csl_stencil.access %3[1, 0] : tensor<4x255xf32> @@ -47,15 +47,15 @@ builtin.module { // CHECK-NEXT: %11 = arith.addf %10, %9 : tensor<255xf32> // CHECK-NEXT: %12 = arith.addf %11, %8 : tensor<255xf32> // CHECK-NEXT: %13 = arith.addf %12, %7 : tensor<255xf32> -// CHECK-NEXT: %14 = "tensor.insert_slice"(%13, %5, %4) <{"static_offsets" = array, "static_sizes" = array, "static_strides" = array, "operandSegmentSizes" = array}> : (tensor<255xf32>, tensor<510xf32>, index) -> tensor<510xf32> +// CHECK-NEXT: %14 = "tensor.insert_slice"(%13, %5, %4) <{static_offsets = array, static_sizes = array, static_strides = array, operandSegmentSizes = array}> : (tensor<255xf32>, tensor<510xf32>, index) -> tensor<510xf32> // CHECK-NEXT: csl_stencil.yield %14 : tensor<510xf32> // CHECK-NEXT: }, { // CHECK-NEXT: ^1(%15 : !stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>>, %16 : tensor<510xf32>): // CHECK-NEXT: %17 = arith.constant 1.666600e-01 : f32 // CHECK-NEXT: %18 = csl_stencil.access %15[0, 0] : !stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>> -// CHECK-NEXT: %19 = "tensor.extract_slice"(%18) <{"static_offsets" = array, "static_sizes" = array, "static_strides" = array, "operandSegmentSizes" = array}> : (tensor<512xf32>) -> tensor<510xf32> +// CHECK-NEXT: %19 = "tensor.extract_slice"(%18) <{static_offsets = array, static_sizes = array, static_strides = array, operandSegmentSizes = array}> : (tensor<512xf32>) -> tensor<510xf32> // CHECK-NEXT: %20 = csl_stencil.access %15[0, 0] : !stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>> -// CHECK-NEXT: %21 = "tensor.extract_slice"(%20) <{"static_offsets" = array, "static_sizes" = array, "static_strides" = array, "operandSegmentSizes" = array}> : (tensor<512xf32>) -> tensor<510xf32> +// CHECK-NEXT: %21 = "tensor.extract_slice"(%20) <{static_offsets = array, static_sizes = array, static_strides = array, operandSegmentSizes = array}> : (tensor<512xf32>) -> tensor<510xf32> // CHECK-NEXT: %22 = arith.addf %16, %21 : tensor<510xf32> // CHECK-NEXT: %23 = arith.addf %22, %19 : tensor<510xf32> // CHECK-NEXT: %24 = tensor.empty() : tensor<510xf32> @@ -89,19 +89,19 @@ builtin.module { // CHECK-NEXT: func.func @bufferized(%a : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>, %b : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>) { // CHECK-NEXT: %0 = tensor.empty() : tensor<510xf32> -// CHECK-NEXT: %1 = csl_stencil.apply(%a : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>, %0 : tensor<510xf32>) -> (!stencil.temp<[0,1]x[0,1]xtensor<510xf32>>) <{"swaps" = [#csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange], "topo" = #dmp.topo<1022x510>, "num_chunks" = 2 : i64, "bounds" = #stencil.bounds<[0, 0], [1, 1]>, "operandSegmentSizes" = array}> ({ +// CHECK-NEXT: %1 = csl_stencil.apply(%a : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>, %0 : tensor<510xf32>) -> (!stencil.temp<[0,1]x[0,1]xtensor<510xf32>>) <{swaps = [#csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange], topo = #dmp.topo<1022x510>, num_chunks = 2 : i64, bounds = #stencil.bounds<[0, 0], [1, 1]>, operandSegmentSizes = array}> ({ // CHECK-NEXT: ^0(%2 : tensor<4x255xf32>, %3 : index, %4 : tensor<510xf32>): // CHECK-NEXT: %5 = arith.constant dense<1.666600e-01> : tensor<510xf32> // CHECK-NEXT: %6 = csl_stencil.access %2[1, 0] : tensor<4x255xf32> // CHECK-NEXT: %7 = csl_stencil.access %2[0, -1] : tensor<4x255xf32> // CHECK-NEXT: %8 = arith.addf %7, %6 : tensor<255xf32> -// CHECK-NEXT: %9 = "tensor.insert_slice"(%8, %4, %3) <{"static_offsets" = array, "static_sizes" = array, "static_strides" = array, "operandSegmentSizes" = array}> : (tensor<255xf32>, tensor<510xf32>, index) -> tensor<510xf32> +// CHECK-NEXT: %9 = "tensor.insert_slice"(%8, %4, %3) <{static_offsets = array, static_sizes = array, static_strides = array, operandSegmentSizes = array}> : (tensor<255xf32>, tensor<510xf32>, index) -> tensor<510xf32> // CHECK-NEXT: csl_stencil.yield %9 : tensor<510xf32> // CHECK-NEXT: }, { // CHECK-NEXT: ^1(%10 : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>, %11 : tensor<510xf32>): // CHECK-NEXT: %12 = arith.constant dense<1.666600e-01> : tensor<510xf32> // CHECK-NEXT: %13 = csl_stencil.access %10[0, 0] : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>> -// CHECK-NEXT: %14 = "tensor.extract_slice"(%13) <{"static_offsets" = array, "static_sizes" = array, "static_strides" = array, "operandSegmentSizes" = array}> : (tensor<512xf32>) -> tensor<510xf32> +// CHECK-NEXT: %14 = "tensor.extract_slice"(%13) <{static_offsets = array, static_sizes = array, static_strides = array, operandSegmentSizes = array}> : (tensor<512xf32>) -> tensor<510xf32> // CHECK-NEXT: %15 = arith.addf %11, %14 : tensor<510xf32> // CHECK-NEXT: %16 = arith.mulf %15, %12 : tensor<510xf32> // CHECK-NEXT: csl_stencil.yield %16 : tensor<510xf32> @@ -135,7 +135,7 @@ builtin.module { // CHECK-NEXT: func.func @coefficients(%a : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>, %b : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>) { // CHECK-NEXT: %0 = tensor.empty() : tensor<510xf32> -// CHECK-NEXT: %1 = csl_stencil.apply(%a : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>, %0 : tensor<510xf32>) -> (!stencil.temp<[0,1]x[0,1]xtensor<510xf32>>) <{"swaps" = [#csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange], "topo" = #dmp.topo<1022x510>, "num_chunks" = 2 : i64, "bounds" = #stencil.bounds<[0, 0], [1, 1]>, "operandSegmentSizes" = array, "coeffs" = [#csl_stencil.coeff<#stencil.index<[0, -1]>, 3.141500e-01 : f32>, #csl_stencil.coeff<#stencil.index<[1, 0]>, 0.234567806 : f32>]}> ({ +// CHECK-NEXT: %1 = csl_stencil.apply(%a : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>, %0 : tensor<510xf32>) -> (!stencil.temp<[0,1]x[0,1]xtensor<510xf32>>) <{swaps = [#csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange], topo = #dmp.topo<1022x510>, num_chunks = 2 : i64, bounds = #stencil.bounds<[0, 0], [1, 1]>, operandSegmentSizes = array, coeffs = [#csl_stencil.coeff<#stencil.index<[0, -1]>, 3.141500e-01 : f32>, #csl_stencil.coeff<#stencil.index<[1, 0]>, 0.234567806 : f32>]}> ({ // CHECK-NEXT: ^0(%2 : tensor<4x255xf32>, %3 : index, %4 : tensor<510xf32>): // CHECK-NEXT: %5 = arith.constant dense<1.234500e-01> : tensor<510xf32> // CHECK-NEXT: %6 = arith.constant dense<0.234567806> : tensor<510xf32> @@ -143,13 +143,13 @@ builtin.module { // CHECK-NEXT: %8 = csl_stencil.access %2[1, 0] : tensor<4x255xf32> // CHECK-NEXT: %9 = csl_stencil.access %2[0, -1] : tensor<4x255xf32> // CHECK-NEXT: %10 = arith.addf %9, %8 : tensor<255xf32> -// CHECK-NEXT: %11 = "tensor.insert_slice"(%10, %4, %3) <{"static_offsets" = array, "static_sizes" = array, "static_strides" = array, "operandSegmentSizes" = array}> : (tensor<255xf32>, tensor<510xf32>, index) -> tensor<510xf32> +// CHECK-NEXT: %11 = "tensor.insert_slice"(%10, %4, %3) <{static_offsets = array, static_sizes = array, static_strides = array, operandSegmentSizes = array}> : (tensor<255xf32>, tensor<510xf32>, index) -> tensor<510xf32> // CHECK-NEXT: csl_stencil.yield %11 : tensor<510xf32> // CHECK-NEXT: }, { // CHECK-NEXT: ^1(%12 : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>, %13 : tensor<510xf32>): // CHECK-NEXT: %14 = arith.constant dense<1.234500e-01> : tensor<510xf32> // CHECK-NEXT: %15 = csl_stencil.access %12[0, 0] : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>> -// CHECK-NEXT: %16 = "tensor.extract_slice"(%15) <{"static_offsets" = array, "static_sizes" = array, "static_strides" = array, "operandSegmentSizes" = array}> : (tensor<512xf32>) -> tensor<510xf32> +// CHECK-NEXT: %16 = "tensor.extract_slice"(%15) <{static_offsets = array, static_sizes = array, static_strides = array, operandSegmentSizes = array}> : (tensor<512xf32>) -> tensor<510xf32> // CHECK-NEXT: %17 = arith.addf %13, %16 : tensor<510xf32> // CHECK-NEXT: %18 = arith.mulf %17, %14 : tensor<510xf32> // CHECK-NEXT: csl_stencil.yield %18 : tensor<510xf32> @@ -191,20 +191,20 @@ builtin.module { // CHECK-NEXT: %2 = arith.constant 1 : index // CHECK-NEXT: %3, %4 = scf.for %arg2 = %1 to %0 step %2 iter_args(%arg3 = %arg0, %arg4 = %arg1) -> (!stencil.field<[-2,3]x[-2,3]xtensor<604xf32>>, !stencil.field<[-2,3]x[-2,3]xtensor<604xf32>>) { // CHECK-NEXT: %5 = tensor.empty() : tensor<600xf32> -// CHECK-NEXT: csl_stencil.apply(%arg3 : !stencil.field<[-2,3]x[-2,3]xtensor<604xf32>>, %5 : tensor<600xf32>) outs (%arg4 : !stencil.field<[-2,3]x[-2,3]xtensor<604xf32>>) <{"swaps" = [#csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange], "topo" = #dmp.topo<600x600>, "num_chunks" = 2 : i64, "bounds" = #stencil.bounds<[0, 0], [1, 1]>, "operandSegmentSizes" = array, "coeffs" = [#csl_stencil.coeff<#stencil.index<[1, 0]>, 119600.297 : f32>, #csl_stencil.coeff<#stencil.index<[-1, 0]>, 119600.297 : f32>]}> ({ +// CHECK-NEXT: csl_stencil.apply(%arg3 : !stencil.field<[-2,3]x[-2,3]xtensor<604xf32>>, %5 : tensor<600xf32>) outs (%arg4 : !stencil.field<[-2,3]x[-2,3]xtensor<604xf32>>) <{swaps = [#csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange], topo = #dmp.topo<600x600>, num_chunks = 2 : i64, bounds = #stencil.bounds<[0, 0], [1, 1]>, operandSegmentSizes = array, coeffs = [#csl_stencil.coeff<#stencil.index<[1, 0]>, 119600.297 : f32>, #csl_stencil.coeff<#stencil.index<[-1, 0]>, 119600.297 : f32>]}> ({ // CHECK-NEXT: ^0(%6 : tensor<8x300xf32>, %7 : index, %8 : tensor<600xf32>): // CHECK-NEXT: %9 = arith.constant dense<1.28715802e+09> : tensor<600xf32> // CHECK-NEXT: %10 = arith.constant dense<119600.297> : tensor<600xf32> // CHECK-NEXT: %11 = csl_stencil.access %6[-1, 0] : tensor<8x300xf32> // CHECK-NEXT: %12 = csl_stencil.access %6[1, 0] : tensor<8x300xf32> // CHECK-NEXT: %13 = arith.addf %11, %12 : tensor<300xf32> -// CHECK-NEXT: %14 = "tensor.insert_slice"(%13, %8, %7) <{"static_offsets" = array, "static_sizes" = array, "static_strides" = array, "operandSegmentSizes" = array}> : (tensor<300xf32>, tensor<600xf32>, index) -> tensor<600xf32> +// CHECK-NEXT: %14 = "tensor.insert_slice"(%13, %8, %7) <{static_offsets = array, static_sizes = array, static_strides = array, operandSegmentSizes = array}> : (tensor<300xf32>, tensor<600xf32>, index) -> tensor<600xf32> // CHECK-NEXT: csl_stencil.yield %14 : tensor<600xf32> // CHECK-NEXT: }, { // CHECK-NEXT: ^1(%15 : !stencil.field<[-2,3]x[-2,3]xtensor<604xf32>>, %16 : tensor<600xf32>): // CHECK-NEXT: %17 = arith.constant dense<1.28715802e+09> : tensor<600xf32> // CHECK-NEXT: %18 = csl_stencil.access %15[0, 0] : !stencil.field<[-2,3]x[-2,3]xtensor<604xf32>> -// CHECK-NEXT: %19 = "tensor.extract_slice"(%18) <{"static_offsets" = array, "static_sizes" = array, "static_strides" = array, "operandSegmentSizes" = array}> : (tensor<604xf32>) -> tensor<600xf32> +// CHECK-NEXT: %19 = "tensor.extract_slice"(%18) <{static_offsets = array, static_sizes = array, static_strides = array, operandSegmentSizes = array}> : (tensor<604xf32>) -> tensor<600xf32> // CHECK-NEXT: %20 = arith.mulf %19, %17 : tensor<600xf32> // CHECK-NEXT: %21 = arith.addf %16, %20 : tensor<600xf32> // CHECK-NEXT: csl_stencil.yield %21 : tensor<600xf32> @@ -236,20 +236,20 @@ builtin.module { // CHECK-NEXT: func.func @uvbke(%arg0 : !stencil.field<[-1,1]x[-1,1]xtensor<64xf32>>, %arg1 : !stencil.field<[-1,1]x[-1,1]xtensor<64xf32>>, %arg4 : !stencil.field<[-1,1]x[-1,1]xtensor<64xf32>>) { // CHECK-NEXT: %0 = tensor.empty() : tensor<1x64xf32> -// CHECK-NEXT: csl_stencil.apply(%arg1 : !stencil.field<[-1,1]x[-1,1]xtensor<64xf32>>, %0 : tensor<1x64xf32>) -> () <{"swaps" = [#csl_stencil.exchange], "topo" = #dmp.topo<64x64>, "num_chunks" = 2 : i64, "operandSegmentSizes" = array}> ({ +// CHECK-NEXT: csl_stencil.apply(%arg1 : !stencil.field<[-1,1]x[-1,1]xtensor<64xf32>>, %0 : tensor<1x64xf32>) -> () <{swaps = [#csl_stencil.exchange], topo = #dmp.topo<64x64>, num_chunks = 2 : i64, operandSegmentSizes = array}> ({ // CHECK-NEXT: ^0(%1 : tensor<1x32xf32>, %2 : index, %3 : tensor<1x64xf32>): // CHECK-NEXT: %4 = csl_stencil.access %1[-1, 0] : tensor<1x32xf32> -// CHECK-NEXT: %5 = "tensor.insert_slice"(%4, %3, %2) <{"static_offsets" = array, "static_sizes" = array, "static_strides" = array, "operandSegmentSizes" = array}> : (tensor<32xf32>, tensor<1x64xf32>, index) -> tensor<1x64xf32> +// CHECK-NEXT: %5 = "tensor.insert_slice"(%4, %3, %2) <{static_offsets = array, static_sizes = array, static_strides = array, operandSegmentSizes = array}> : (tensor<32xf32>, tensor<1x64xf32>, index) -> tensor<1x64xf32> // CHECK-NEXT: csl_stencil.yield %5 : tensor<1x64xf32> // CHECK-NEXT: }, { // CHECK-NEXT: ^1(%6 : !stencil.field<[-1,1]x[-1,1]xtensor<64xf32>>, %7 : tensor<1x64xf32>): // CHECK-NEXT: csl_stencil.yield // CHECK-NEXT: }) // CHECK-NEXT: %1 = tensor.empty() : tensor<64xf32> -// CHECK-NEXT: csl_stencil.apply(%arg0 : !stencil.field<[-1,1]x[-1,1]xtensor<64xf32>>, %1 : tensor<64xf32>, %arg1 : !stencil.field<[-1,1]x[-1,1]xtensor<64xf32>>, %0 : tensor<1x64xf32>) outs (%arg4 : !stencil.field<[-1,1]x[-1,1]xtensor<64xf32>>) <{"swaps" = [#csl_stencil.exchange], "topo" = #dmp.topo<64x64>, "num_chunks" = 2 : i64, "bounds" = #stencil.bounds<[0, 0], [1, 1]>, "operandSegmentSizes" = array}> ({ +// CHECK-NEXT: csl_stencil.apply(%arg0 : !stencil.field<[-1,1]x[-1,1]xtensor<64xf32>>, %1 : tensor<64xf32>, %arg1 : !stencil.field<[-1,1]x[-1,1]xtensor<64xf32>>, %0 : tensor<1x64xf32>) outs (%arg4 : !stencil.field<[-1,1]x[-1,1]xtensor<64xf32>>) <{swaps = [#csl_stencil.exchange], topo = #dmp.topo<64x64>, num_chunks = 2 : i64, bounds = #stencil.bounds<[0, 0], [1, 1]>, operandSegmentSizes = array}> ({ // CHECK-NEXT: ^0(%2 : tensor<1x32xf32>, %3 : index, %4 : tensor<64xf32>): // CHECK-NEXT: %5 = csl_stencil.access %2[0, -1] : tensor<1x32xf32> -// CHECK-NEXT: %6 = "tensor.insert_slice"(%5, %4, %3) <{"static_offsets" = array, "static_sizes" = array, "static_strides" = array, "operandSegmentSizes" = array}> : (tensor<32xf32>, tensor<64xf32>, index) -> tensor<64xf32> +// CHECK-NEXT: %6 = "tensor.insert_slice"(%5, %4, %3) <{static_offsets = array, static_sizes = array, static_strides = array, operandSegmentSizes = array}> : (tensor<32xf32>, tensor<64xf32>, index) -> tensor<64xf32> // CHECK-NEXT: csl_stencil.yield %6 : tensor<64xf32> // CHECK-NEXT: }, { // CHECK-NEXT: ^1(%7 : !stencil.field<[-1,1]x[-1,1]xtensor<64xf32>>, %8 : tensor<64xf32>, %9 : !stencil.field<[-1,1]x[-1,1]xtensor<64xf32>>, %10 : tensor<1x64xf32>): diff --git a/tests/filecheck/transforms/convert-stencil-to-ll-mlir.mlir b/tests/filecheck/transforms/convert-stencil-to-ll-mlir.mlir index 61f4e2788d..41733edc4c 100644 --- a/tests/filecheck/transforms/convert-stencil-to-ll-mlir.mlir +++ b/tests/filecheck/transforms/convert-stencil-to-ll-mlir.mlir @@ -30,7 +30,7 @@ builtin.module { // CHECK-NEXT: %10 = arith.constant 65 : index // CHECK-NEXT: %11 = arith.constant 66 : index // CHECK-NEXT: %12 = arith.constant 63 : index -// CHECK-NEXT: "scf.parallel"(%4, %5, %6, %10, %11, %12, %7, %8, %9) <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "scf.parallel"(%4, %5, %6, %10, %11, %12, %7, %8, %9) <{operandSegmentSizes = array}> ({ // CHECK-NEXT: ^0(%13 : index, %14 : index, %15 : index): // CHECK-NEXT: %16 = arith.constant 1.000000e+00 : f64 // CHECK-NEXT: %17 = arith.addf %0, %16 : f64 @@ -69,7 +69,7 @@ builtin.module { // CHECK-NEXT: %3 = arith.constant 1 : index // CHECK-NEXT: %4 = arith.constant 2000 : index // CHECK-NEXT: %5 = arith.constant 2000 : index -// CHECK-NEXT: "scf.parallel"(%0, %1, %4, %5, %2, %3) <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "scf.parallel"(%0, %1, %4, %5, %2, %3) <{operandSegmentSizes = array}> ({ // CHECK-NEXT: ^0(%6 : index, %7 : index): // CHECK-NEXT: %i = memref.load %fim1_loadview[%6, %7] : memref<2000x2000xf32, strided<[2004, 1], offset: 4010>> // CHECK-NEXT: memref.store %i, %fi_storeview[%6, %7] : memref<2000x2000xf32, strided<[2004, 1], offset: 4010>> @@ -100,7 +100,7 @@ builtin.module { // CHECK-NEXT: %3 = arith.constant 0 : index // CHECK-NEXT: %4 = arith.constant 1 : index // CHECK-NEXT: %5 = arith.constant 68 : index -// CHECK-NEXT: "scf.parallel"(%3, %5, %4) <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "scf.parallel"(%3, %5, %4) <{operandSegmentSizes = array}> ({ // CHECK-NEXT: ^0(%6 : index): // CHECK-NEXT: %7 = arith.constant -1 : index // CHECK-NEXT: %8 = arith.addi %6, %7 : index @@ -130,7 +130,7 @@ builtin.module { // CHECK-NEXT: %6 = arith.constant 1 : index // CHECK-NEXT: %7 = arith.constant 64 : index // CHECK-NEXT: %8 = arith.constant 68 : index -// CHECK-NEXT: "scf.parallel"(%3, %4, %7, %8, %5, %6) <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "scf.parallel"(%3, %4, %7, %8, %5, %6) <{operandSegmentSizes = array}> ({ // CHECK-NEXT: ^0(%9 : index, %10 : index): // CHECK-NEXT: %11 = arith.constant -1 : index // CHECK-NEXT: %12 = arith.addi %9, %11 : index @@ -163,7 +163,7 @@ builtin.module { // CHECK-NEXT: %9 = arith.constant 64 : index // CHECK-NEXT: %10 = arith.constant 64 : index // CHECK-NEXT: %11 = arith.constant 68 : index -// CHECK-NEXT: "scf.parallel"(%3, %4, %5, %9, %10, %11, %6, %7, %8) <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "scf.parallel"(%3, %4, %5, %9, %10, %11, %6, %7, %8) <{operandSegmentSizes = array}> ({ // CHECK-NEXT: ^0(%12 : index, %13 : index, %14 : index): // CHECK-NEXT: %15 = arith.constant -1 : index // CHECK-NEXT: %16 = arith.addi %12, %15 : index @@ -226,7 +226,7 @@ builtin.module { // CHECK-NEXT: %14 = arith.constant 64 : index // CHECK-NEXT: %15 = arith.constant 64 : index // CHECK-NEXT: %16 = arith.constant 64 : index -// CHECK-NEXT: "scf.parallel"(%8, %9, %10, %14, %15, %16, %11, %12, %13) <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "scf.parallel"(%8, %9, %10, %14, %15, %16, %11, %12, %13) <{operandSegmentSizes = array}> ({ // CHECK-NEXT: ^0(%17 : index, %18 : index, %19 : index): // CHECK-NEXT: %20 = arith.constant -1 : index // CHECK-NEXT: %21 = arith.addi %17, %20 : index @@ -283,7 +283,7 @@ builtin.module { // CHECK-NEXT: %0 = arith.constant -16 : index // CHECK-NEXT: %1 = arith.constant 1 : index // CHECK-NEXT: %2 = arith.constant 16 : index -// CHECK-NEXT: "scf.parallel"(%0, %2, %1) <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "scf.parallel"(%0, %2, %1) <{operandSegmentSizes = array}> ({ // CHECK-NEXT: ^0(%3 : index): // CHECK-NEXT: %val = memref.load %in_loadview[%3] : memref<32xf64, strided<[1], offset: 32>> // CHECK-NEXT: memref.store %val, %out_storeview[%3] : memref<32xf64, strided<[1], offset: 32>> @@ -314,7 +314,7 @@ builtin.module { // CHECK-NEXT: %5 = arith.constant 1 : index // CHECK-NEXT: %6 = arith.constant 1 : index // CHECK-NEXT: %7 = arith.constant 65 : index -// CHECK-NEXT: "scf.parallel"(%5, %7, %6) <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "scf.parallel"(%5, %7, %6) <{operandSegmentSizes = array}> ({ // CHECK-NEXT: ^0(%8 : index): // CHECK-NEXT: %9 = arith.constant -1 : index // CHECK-NEXT: %10 = arith.addi %8, %9 : index @@ -325,7 +325,7 @@ builtin.module { // CHECK-NEXT: %12 = arith.constant 0 : index // CHECK-NEXT: %13 = arith.constant 1 : index // CHECK-NEXT: %14 = arith.constant 64 : index -// CHECK-NEXT: "scf.parallel"(%12, %14, %13) <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "scf.parallel"(%12, %14, %13) <{operandSegmentSizes = array}> ({ // CHECK-NEXT: ^1(%15 : index): // CHECK-NEXT: %16 = arith.constant 1 : index // CHECK-NEXT: %17 = arith.addi %15, %16 : index @@ -359,7 +359,7 @@ builtin.module { // CHECK-NEXT: %6 = arith.constant 1 : index // CHECK-NEXT: %7 = arith.constant 1 : index // CHECK-NEXT: %8 = arith.constant 65 : index -// CHECK-NEXT: "scf.parallel"(%6, %8, %7) <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "scf.parallel"(%6, %8, %7) <{operandSegmentSizes = array}> ({ // CHECK-NEXT: ^0(%9 : index): // CHECK-NEXT: %10 = arith.constant -1 : index // CHECK-NEXT: %11 = arith.addi %9, %10 : index @@ -370,7 +370,7 @@ builtin.module { // CHECK-NEXT: %13 = arith.constant 0 : index // CHECK-NEXT: %14 = arith.constant 1 : index // CHECK-NEXT: %15 = arith.constant 64 : index -// CHECK-NEXT: "scf.parallel"(%13, %15, %14) <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "scf.parallel"(%13, %15, %14) <{operandSegmentSizes = array}> ({ // CHECK-NEXT: ^1(%16 : index): // CHECK-NEXT: %17 = arith.constant 1 : index // CHECK-NEXT: %18 = arith.addi %16, %17 : index @@ -405,15 +405,15 @@ builtin.module { func.return } -// CHECK: func.func @apply_kernel(%0 : memref<15x15xf32>, %1 : memref<15x15xf32>, %timers : !llvm.ptr) attributes {"param_names" = ["u_vec_1", "u_vec", "timers"]}{ -// CHECK-NEXT: %2 = "gpu.alloc"() <{"operandSegmentSizes" = array}> : () -> memref<15x15xf32> +// CHECK: func.func @apply_kernel(%0 : memref<15x15xf32>, %1 : memref<15x15xf32>, %timers : !llvm.ptr) attributes {param_names = ["u_vec_1", "u_vec", "timers"]}{ +// CHECK-NEXT: %2 = "gpu.alloc"() <{operandSegmentSizes = array}> : () -> memref<15x15xf32> // CHECK-NEXT: %u_vec = builtin.unrealized_conversion_cast %2 : memref<15x15xf32> to memref<15x15xf32> // CHECK-NEXT: %3 = builtin.unrealized_conversion_cast %1 : memref<15x15xf32> to memref<15x15xf32> -// CHECK-NEXT: "gpu.memcpy"(%2, %3) {"operandSegmentSizes" = array} : (memref<15x15xf32>, memref<15x15xf32>) -> () -// CHECK-NEXT: %4 = "gpu.alloc"() <{"operandSegmentSizes" = array}> : () -> memref<15x15xf32> +// CHECK-NEXT: "gpu.memcpy"(%2, %3) {operandSegmentSizes = array} : (memref<15x15xf32>, memref<15x15xf32>) -> () +// CHECK-NEXT: %4 = "gpu.alloc"() <{operandSegmentSizes = array}> : () -> memref<15x15xf32> // CHECK-NEXT: %u_vec_1 = builtin.unrealized_conversion_cast %4 : memref<15x15xf32> to memref<15x15xf32> // CHECK-NEXT: %5 = builtin.unrealized_conversion_cast %0 : memref<15x15xf32> to memref<15x15xf32> -// CHECK-NEXT: "gpu.memcpy"(%4, %5) {"operandSegmentSizes" = array} : (memref<15x15xf32>, memref<15x15xf32>) -> () +// CHECK-NEXT: "gpu.memcpy"(%4, %5) {operandSegmentSizes = array} : (memref<15x15xf32>, memref<15x15xf32>) -> () // CHECK-NEXT: %time_m = arith.constant 0 : index // CHECK-NEXT: %time_M = arith.constant 10 : index // CHECK-NEXT: %step = arith.constant 1 : index @@ -426,7 +426,7 @@ builtin.module { // CHECK-NEXT: %11 = arith.constant 1 : index // CHECK-NEXT: %12 = arith.constant 11 : index // CHECK-NEXT: %13 = arith.constant 11 : index -// CHECK-NEXT: "scf.parallel"(%8, %9, %12, %13, %10, %11) <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "scf.parallel"(%8, %9, %12, %13, %10, %11) <{operandSegmentSizes = array}> ({ // CHECK-NEXT: ^0(%14 : index, %15 : index): // CHECK-NEXT: %16 = memref.load %t0_loadview[%14, %15] : memref<11x11xf32, strided<[15, 1], offset: 32>> // CHECK-NEXT: memref.store %16, %t1_storeview[%14, %15] : memref<11x11xf32, strided<[15, 1], offset: 32>> @@ -466,7 +466,7 @@ builtin.module { // CHECK-NEXT: %10 = arith.constant 65 : index // CHECK-NEXT: %11 = arith.constant 66 : index // CHECK-NEXT: %12 = arith.constant 63 : index -// CHECK-NEXT: "scf.parallel"(%4, %5, %6, %10, %11, %12, %7, %8, %9) <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "scf.parallel"(%4, %5, %6, %10, %11, %12, %7, %8, %9) <{operandSegmentSizes = array}> ({ // CHECK-NEXT: ^0(%13 : index, %14 : index, %15 : index): // CHECK-NEXT: %16 = arith.constant 1.000000e+00 : f64 // CHECK-NEXT: %17 = arith.constant 2.000000e+00 : f64 @@ -537,7 +537,7 @@ builtin.module { // CHECK-NEXT: %8 = arith.constant 64 : index // CHECK-NEXT: %9 = arith.constant 64 : index // CHECK-NEXT: %10 = arith.constant 64 : index -// CHECK-NEXT: "scf.parallel"(%2, %3, %4, %8, %9, %10, %5, %6, %7) <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "scf.parallel"(%2, %3, %4, %8, %9, %10, %5, %6, %7) <{operandSegmentSizes = array}> ({ // CHECK-NEXT: ^0(%x : index, %y : index, %z : index): // CHECK-NEXT: %xy = arith.addi %x, %y : index // CHECK-NEXT: %xyz = arith.addi %xy, %z : index @@ -571,7 +571,7 @@ builtin.module { // CHECK-NEXT: %8 = arith.constant 64 : index // CHECK-NEXT: %9 = arith.constant 64 : index // CHECK-NEXT: %10 = arith.constant 64 : index -// CHECK-NEXT: "scf.parallel"(%2, %3, %4, %8, %9, %10, %5, %6, %7) <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "scf.parallel"(%2, %3, %4, %8, %9, %10, %5, %6, %7) <{operandSegmentSizes = array}> ({ // CHECK-NEXT: ^0(%11 : index, %12 : index, %z : index): // CHECK-NEXT: %x = arith.constant 1 : index // CHECK-NEXT: %x_1 = arith.addi %11, %x : index @@ -608,7 +608,7 @@ func.func @store_result_lowering(%arg0 : f64) { // CHECK-NEXT: %8 = arith.constant 7 : index // CHECK-NEXT: %9 = arith.constant 7 : index // CHECK-NEXT: %10 = arith.constant 7 : index -// CHECK-NEXT: "scf.parallel"(%2, %3, %4, %8, %9, %10, %5, %6, %7) <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "scf.parallel"(%2, %3, %4, %8, %9, %10, %5, %6, %7) <{operandSegmentSizes = array}> ({ // CHECK-NEXT: ^0(%11 : index, %12 : index, %13 : index): // CHECK-NEXT: memref.store %arg0, %1[%11, %12, %13] : memref<7x7x7xf64, strided<[49, 7, 1]>> // CHECK-NEXT: memref.store %arg0, %0[%11, %12, %13] : memref<7x7x7xf64, strided<[49, 7, 1]>> @@ -637,7 +637,7 @@ func.func @if_lowering(%arg0_1 : f64, %b0 : !stencil.field<[0,7]x[0,7]x[0,7]xf64 func.return } -// CHECK: func.func @if_lowering(%arg0 : f64, %b0 : memref<7x7x7xf64>, %b1 : memref<7x7x7xf64>) attributes {"stencil.program"}{ +// CHECK: func.func @if_lowering(%arg0 : f64, %b0 : memref<7x7x7xf64>, %b1 : memref<7x7x7xf64>) attributes {stencil.program}{ // CHECK-NEXT: %b0_storeview = memref.subview %b0[0, 0, 0] [7, 7, 7] [1, 1, 1] : memref<7x7x7xf64> to memref<7x7x7xf64, strided<[49, 7, 1]>> // CHECK-NEXT: %b1_storeview = memref.subview %b1[0, 0, 0] [7, 7, 7] [1, 1, 1] : memref<7x7x7xf64> to memref<7x7x7xf64, strided<[49, 7, 1]>> // CHECK-NEXT: %0 = arith.constant 0 : index @@ -649,7 +649,7 @@ func.func @if_lowering(%arg0_1 : f64, %b0 : !stencil.field<[0,7]x[0,7]x[0,7]xf64 // CHECK-NEXT: %6 = arith.constant 7 : index // CHECK-NEXT: %7 = arith.constant 7 : index // CHECK-NEXT: %8 = arith.constant 7 : index -// CHECK-NEXT: "scf.parallel"(%0, %1, %2, %6, %7, %8, %3, %4, %5) <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "scf.parallel"(%0, %1, %2, %6, %7, %8, %3, %4, %5) <{operandSegmentSizes = array}> ({ // CHECK-NEXT: ^0(%9 : index, %10 : index, %11 : index): // CHECK-NEXT: %true = "test.op"() : () -> i1 // CHECK-NEXT: %12, %13 = scf.if %true -> (f64, f64) { @@ -689,7 +689,7 @@ func.func @if_lowering(%arg0_1 : f64, %b0 : !stencil.field<[0,7]x[0,7]x[0,7]xf64 // CHECK-NEXT: %6 = arith.constant 1 : index // CHECK-NEXT: %7 = arith.constant 33 : index // CHECK-NEXT: %8 = arith.constant 66 : index -// CHECK-NEXT: "scf.parallel"(%3, %4, %7, %8, %5, %6) <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "scf.parallel"(%3, %4, %7, %8, %5, %6) <{operandSegmentSizes = array}> ({ // CHECK-NEXT: ^0(%9 : index, %10 : index): // CHECK-NEXT: %11 = arith.constant 1.000000e+00 : f64 // CHECK-NEXT: memref.store %11, %2[%9, %10] : memref<64x64xf64, strided<[70, 1], offset: 213>> @@ -701,7 +701,7 @@ func.func @if_lowering(%arg0_1 : f64, %b0 : !stencil.field<[0,7]x[0,7]x[0,7]xf64 // CHECK-NEXT: %15 = arith.constant 1 : index // CHECK-NEXT: %16 = arith.constant 65 : index // CHECK-NEXT: %17 = arith.constant 66 : index -// CHECK-NEXT: "scf.parallel"(%12, %13, %16, %17, %14, %15) <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "scf.parallel"(%12, %13, %16, %17, %14, %15) <{operandSegmentSizes = array}> ({ // CHECK-NEXT: ^1(%18 : index, %19 : index): // CHECK-NEXT: %20 = arith.constant 2.000000e+00 : f64 // CHECK-NEXT: memref.store %20, %2[%18, %19] : memref<64x64xf64, strided<[70, 1], offset: 213>> @@ -742,7 +742,7 @@ func.func @buffered_combine(%115 : !stencil.field) { // CHECK-NEXT: %7 = arith.constant 1 : index // CHECK-NEXT: %8 = arith.constant 33 : index // CHECK-NEXT: %9 = arith.constant 66 : index -// CHECK-NEXT: "scf.parallel"(%4, %5, %8, %9, %6, %7) <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "scf.parallel"(%4, %5, %8, %9, %6, %7) <{operandSegmentSizes = array}> ({ // CHECK-NEXT: ^0(%10 : index, %11 : index): // CHECK-NEXT: %12 = arith.constant 1.000000e+00 : f64 // CHECK-NEXT: memref.store %12, %1[%10, %11] : memref<64x64xf64, strided<[64, 1], offset: -66>> @@ -754,7 +754,7 @@ func.func @buffered_combine(%115 : !stencil.field) { // CHECK-NEXT: %16 = arith.constant 1 : index // CHECK-NEXT: %17 = arith.constant 65 : index // CHECK-NEXT: %18 = arith.constant 66 : index -// CHECK-NEXT: "scf.parallel"(%13, %14, %17, %18, %15, %16) <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "scf.parallel"(%13, %14, %17, %18, %15, %16) <{operandSegmentSizes = array}> ({ // CHECK-NEXT: ^1(%19 : index, %20 : index): // CHECK-NEXT: %21 = arith.constant 2.000000e+00 : f64 // CHECK-NEXT: memref.store %21, %1[%19, %20] : memref<64x64xf64, strided<[64, 1], offset: -66>> @@ -766,7 +766,7 @@ func.func @buffered_combine(%115 : !stencil.field) { // CHECK-NEXT: %25 = arith.constant 1 : index // CHECK-NEXT: %26 = arith.constant 65 : index // CHECK-NEXT: %27 = arith.constant 66 : index -// CHECK-NEXT: "scf.parallel"(%22, %23, %26, %27, %24, %25) <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "scf.parallel"(%22, %23, %26, %27, %24, %25) <{operandSegmentSizes = array}> ({ // CHECK-NEXT: ^2(%28 : index, %29 : index): // CHECK-NEXT: %30 = arith.constant 1.000000e+00 : f64 // CHECK-NEXT: %31 = memref.load %1[%28, %29] : memref<64x64xf64, strided<[64, 1], offset: -66>> @@ -801,7 +801,7 @@ func.func @buffered_combine(%115 : !stencil.field) { // CHECK-NEXT: %9 = arith.constant 1 : index // CHECK-NEXT: %10 = arith.constant 8 : index // CHECK-NEXT: %11 = arith.constant 8 : index -// CHECK-NEXT: "scf.parallel"(%6, %7, %10, %11, %8, %9) <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "scf.parallel"(%6, %7, %10, %11, %8, %9) <{operandSegmentSizes = array}> ({ // CHECK-NEXT: ^0(%12 : index, %13 : index): // CHECK-NEXT: %14 = memref.load %4[%12] : memref<8xf64, strided<[1]>> // CHECK-NEXT: %15 = memref.load %5[%13] : memref<8xf64, strided<[1]>> @@ -833,7 +833,7 @@ func.func @buffered_combine(%115 : !stencil.field) { // CHECK-NEXT: %10 = arith.constant 64 : index // CHECK-NEXT: %11 = arith.constant 64 : index // CHECK-NEXT: %12 = arith.constant 64 : index -// CHECK-NEXT: "scf.parallel"(%4, %5, %6, %10, %11, %12, %7, %8, %9) <{"operandSegmentSizes" = array}> ({ +// CHECK-NEXT: "scf.parallel"(%4, %5, %6, %10, %11, %12, %7, %8, %9) <{operandSegmentSizes = array}> ({ // CHECK-NEXT: ^0(%13 : index, %14 : index, %15 : index): // CHECK-NEXT: %16 = memref.load %3[%13, %14, %15] : memref<72x72x72xf64, strided<[5184, 72, 1], offset: 21028>> // CHECK-NEXT: memref.store %16, %2[%13, %14, %15] : memref<72x72x72xf64, strided<[5184, 72, 1], offset: 21028>> @@ -869,7 +869,7 @@ func.func @buffered_combine(%115 : !stencil.field) { func.return } -// CHECK: "stencil.store"(%2, %0) {"bounds" = #stencil.bounds<[0, 0, 0], [64, 64, 64]>} : (!stencil.temp<[0,64]x[0,64]x[0,64]xindex>, !stencil.field<[0,64]x[0,64]x[0,64]xindex>) -> () +// CHECK: "stencil.store"(%2, %0) {bounds = #stencil.bounds<[0, 0, 0], [64, 64, 64]>} : (!stencil.temp<[0,64]x[0,64]x[0,64]xindex>, !stencil.field<[0,64]x[0,64]x[0,64]xindex>) -> () // CHECK-NEXT: ^^^^^^^^^^^^^^^------------------------------------------------------------------------------------------------------------------------- // CHECK-NEXT: | Error while applying pattern: Cannot lower directly if storing to the same field multiple times! Try running `stencil-bufferize` before. // CHECK-NEXT: ---------------------------------------------------------------------------------------------------------------------------------------- diff --git a/tests/filecheck/transforms/convert_ml_program_to_memref.mlir b/tests/filecheck/transforms/convert_ml_program_to_memref.mlir index 1c33de3de7..e8ddf91040 100644 --- a/tests/filecheck/transforms/convert_ml_program_to_memref.mlir +++ b/tests/filecheck/transforms/convert_ml_program_to_memref.mlir @@ -5,7 +5,7 @@ ml_program.global private @global_same_type(dense<4> : tensor<4xi32>) : tensor<4 %0 = ml_program.global_load_const @global_same_type : tensor<4xi32> // CHECK: builtin.module { -// CHECK-NEXT: "memref.global"() <{"sym_name" = "global_same_type", "type" = memref<4xi32>, "initial_value" = dense<4> : tensor<4xi32>, "sym_visibility" = "private", "constant"}> : () -> () +// CHECK-NEXT: "memref.global"() <{sym_name = "global_same_type", type = memref<4xi32>, initial_value = dense<4> : tensor<4xi32>, sym_visibility = "private", constant}> : () -> () // CHECK-NEXT: %0 = memref.get_global @global_same_type : memref<4xi32> // CHECK-NEXT: %1 = bufferization.to_tensor %0 : memref<4xi32> // CHECK-NEXT: } diff --git a/tests/filecheck/transforms/convert_onnx_to_linalg.mlir b/tests/filecheck/transforms/convert_onnx_to_linalg.mlir index 74fd0cec7b..76ff2fef9a 100644 --- a/tests/filecheck/transforms/convert_onnx_to_linalg.mlir +++ b/tests/filecheck/transforms/convert_onnx_to_linalg.mlir @@ -89,21 +89,21 @@ // CHECK-NEXT: %res_max_pool_single_out_1 = tensor.empty() : tensor<1x16x4x4xf32> // CHECK-NEXT: %res_max_pool_single_out_2 = arith.constant -1.000000e+308 : f64 // CHECK-NEXT: %res_max_pool_single_out_3 = linalg.fill ins(%res_max_pool_single_out_2 : f64) outs(%res_max_pool_single_out_1 : tensor<1x16x4x4xf32>) -> tensor<1x16x4x4xf32> -// CHECK-NEXT: %res_max_pool_single_out_4 = linalg.pooling_nchw_max {"dilations" = dense<1> : tensor<2xi64>, "strides" = dense<3> : tensor<2xi64>} ins(%t26, %res_max_pool_single_out : tensor<1x16x14x14xf32>, tensor<3x3xf32>) outs(%res_max_pool_single_out_3 : tensor<1x16x4x4xf32>) -> tensor<1x16x4x4xf32> +// CHECK-NEXT: %res_max_pool_single_out_4 = linalg.pooling_nchw_max {dilations = dense<1> : tensor<2xi64>, strides = dense<3> : tensor<2xi64>} ins(%t26, %res_max_pool_single_out : tensor<1x16x14x14xf32>, tensor<3x3xf32>) outs(%res_max_pool_single_out_3 : tensor<1x16x4x4xf32>) -> tensor<1x16x4x4xf32> %t20, %t21, %t22 = "test.op"() : () -> (tensor<1x1x5x5xf32>, tensor<1x1x3x3xf32>, none) %res_conv_2 = "onnx.Conv"(%t20, %t21, %t22) {onnx_node_name = "/Conv", "auto_pad" = "NOTSET", "group" = 1 : i64, "kernel_shape" = [3 : i64, 3 : i64], "dilations" = [1 : i64, 1 : i64], "strides" = [1 : i64, 1 : i64], "pads" = [0 : i64, 0 : i64, 0: i64, 0 : i64]}: (tensor<1x1x5x5xf32>, tensor<1x1x3x3xf32>, none) -> tensor<1x1x3x3xf32> // CHECK-NEXT: %t20, %t21, %t22 = "test.op"() : () -> (tensor<1x1x5x5xf32>, tensor<1x1x3x3xf32>, none) // CHECK-NEXT: %res_conv = tensor.empty() : tensor<1x1x3x3xf32> -// CHECK-NEXT: %res_conv_1 = linalg.conv_2d_nchw_fchw {"dilations" = dense<1> : tensor<2xi64>, "strides" = dense<1> : tensor<2xi64>} ins(%t20, %t21 : tensor<1x1x5x5xf32>, tensor<1x1x3x3xf32>) outs(%res_conv : tensor<1x1x3x3xf32>) -> tensor<1x1x3x3xf32> +// CHECK-NEXT: %res_conv_1 = linalg.conv_2d_nchw_fchw {dilations = dense<1> : tensor<2xi64>, strides = dense<1> : tensor<2xi64>} ins(%t20, %t21 : tensor<1x1x5x5xf32>, tensor<1x1x3x3xf32>) outs(%res_conv : tensor<1x1x3x3xf32>) -> tensor<1x1x3x3xf32> %t23, %t24, %t25 = "test.op"() : () -> (tensor<1x8x14x14xf32>, tensor<16x8x5x5xf32>, tensor<16xf32>) %res_conv_3 = "onnx.Conv"(%t23, %t24, %t25) {onnx_node_name = "/Conv", "auto_pad" = "SAME_UPPER", "group" = 1 : i64, "kernel_shape" = [5 : i64, 5 : i64], "dilations" = [1 : i64, 1 : i64], "strides" = [1 : i64, 1 : i64], "pads" = [0 : i64, 0 : i64, 0: i64, 0 : i64]} : (tensor<1x8x14x14xf32>, tensor<16x8x5x5xf32>, tensor<16xf32>) -> tensor<1x16x14x14xf32> // CHECK-NEXT: %t23, %t24, %t25 = "test.op"() : () -> (tensor<1x8x14x14xf32>, tensor<16x8x5x5xf32>, tensor<16xf32>) // CHECK-NEXT: %res_conv_2 = tensor.empty() : tensor<1x16x14x14xf32> -// CHECK-NEXT: %res_conv_3 = linalg.conv_2d_nchw_fchw {"dilations" = dense<1> : tensor<2xi64>, "strides" = dense<1> : tensor<2xi64>} ins(%t23, %t24 : tensor<1x8x14x14xf32>, tensor<16x8x5x5xf32>) outs(%res_conv_2 : tensor<1x16x14x14xf32>) -> tensor<1x16x14x14xf32> +// CHECK-NEXT: %res_conv_3 = linalg.conv_2d_nchw_fchw {dilations = dense<1> : tensor<2xi64>, strides = dense<1> : tensor<2xi64>} ins(%t23, %t24 : tensor<1x8x14x14xf32>, tensor<16x8x5x5xf32>) outs(%res_conv_2 : tensor<1x16x14x14xf32>) -> tensor<1x16x14x14xf32> // CHECK-NEXT: %res_conv_4 = linalg.add ins(%t25 : tensor<16xf32>) outs(%res_conv_3 : tensor<1x16x14x14xf32>) -> tensor<1x16x14x14xf32> %res_constant = "onnx.Constant"() {onnx_node_name = "/Constant", "value" = dense<1> : tensor<1xi64>}: () -> tensor<1xi64> diff --git a/tests/filecheck/transforms/convert_riscv_to_llvm.mlir b/tests/filecheck/transforms/convert_riscv_to_llvm.mlir index 401a0487ee..a4a60f657f 100644 --- a/tests/filecheck/transforms/convert_riscv_to_llvm.mlir +++ b/tests/filecheck/transforms/convert_riscv_to_llvm.mlir @@ -7,61 +7,61 @@ %x0 = riscv.get_register : !riscv.reg // CHECK: builtin.module { -// CHECK-NEXT: %reg = "llvm.inline_asm"() <{"asm_string" = "li $0, 0", "constraints" = "=r", "asm_dialect" = 0 : i64}> : () -> i32 +// CHECK-NEXT: %reg = "llvm.inline_asm"() <{asm_string = "li $0, 0", constraints = "=r", asm_dialect = 0 : i64}> : () -> i32 // CHECK-NEXT: %reg_1 = builtin.unrealized_conversion_cast %reg : i32 to !riscv.reg -// CHECK-NEXT: %a0 = "llvm.inline_asm"() <{"asm_string" = "li $0, 0", "constraints" = "=r", "asm_dialect" = 0 : i64}> : () -> i32 +// CHECK-NEXT: %a0 = "llvm.inline_asm"() <{asm_string = "li $0, 0", constraints = "=r", asm_dialect = 0 : i64}> : () -> i32 // CHECK-NEXT: %a0_1 = builtin.unrealized_conversion_cast %a0 : i32 to !riscv.reg // CHECK-NEXT: %x0 = riscv.get_register : !riscv.reg // standard risc-v instructions %li = riscv.li 0 : !riscv.reg -// CHECK-NEXT: %li = "llvm.inline_asm"() <{"asm_string" = "li $0, 0", "constraints" = "=r", "asm_dialect" = 0 : i64}> : () -> i32 +// CHECK-NEXT: %li = "llvm.inline_asm"() <{asm_string = "li $0, 0", constraints = "=r", asm_dialect = 0 : i64}> : () -> i32 // CHECK-NEXT: %li_1 = builtin.unrealized_conversion_cast %li : i32 to !riscv.reg %sub = riscv.sub %reg, %reg : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-NEXT: %sub = builtin.unrealized_conversion_cast %reg_1 : !riscv.reg to i32 // CHECK-NEXT: %sub_1 = builtin.unrealized_conversion_cast %reg_1 : !riscv.reg to i32 -// CHECK-NEXT: %sub_2 = "llvm.inline_asm"(%sub, %sub_1) <{"asm_string" = "sub $0, $1, $2", "constraints" = "=r,rI,rI", "asm_dialect" = 0 : i64}> : (i32, i32) -> i32 +// CHECK-NEXT: %sub_2 = "llvm.inline_asm"(%sub, %sub_1) <{asm_string = "sub $0, $1, $2", constraints = "=r,rI,rI", asm_dialect = 0 : i64}> : (i32, i32) -> i32 // CHECK-NEXT: %sub_3 = builtin.unrealized_conversion_cast %sub_2 : i32 to !riscv.reg %div = riscv.div %reg, %reg : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-NEXT: %div = builtin.unrealized_conversion_cast %reg_1 : !riscv.reg to i32 // CHECK-NEXT: %div_1 = builtin.unrealized_conversion_cast %reg_1 : !riscv.reg to i32 -// CHECK-NEXT: %div_2 = "llvm.inline_asm"(%div, %div_1) <{"asm_string" = "div $0, $1, $2", "constraints" = "=r,rI,rI", "asm_dialect" = 0 : i64}> : (i32, i32) -> i32 +// CHECK-NEXT: %div_2 = "llvm.inline_asm"(%div, %div_1) <{asm_string = "div $0, $1, $2", constraints = "=r,rI,rI", asm_dialect = 0 : i64}> : (i32, i32) -> i32 // CHECK-NEXT: %div_3 = builtin.unrealized_conversion_cast %div_2 : i32 to !riscv.reg // named riscv registers: %li_named = riscv.li 0 : !riscv.reg -// CHECK-NEXT: %li_named = "llvm.inline_asm"() <{"asm_string" = "li $0, 0", "constraints" = "=r", "asm_dialect" = 0 : i64}> : () -> i32 +// CHECK-NEXT: %li_named = "llvm.inline_asm"() <{asm_string = "li $0, 0", constraints = "=r", asm_dialect = 0 : i64}> : () -> i32 // CHECK-NEXT: %li_named_1 = builtin.unrealized_conversion_cast %li_named : i32 to !riscv.reg %sub_named = riscv.sub %a0, %a0 : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-NEXT: %sub_named = builtin.unrealized_conversion_cast %a0_1 : !riscv.reg to i32 // CHECK-NEXT: %sub_named_1 = builtin.unrealized_conversion_cast %a0_1 : !riscv.reg to i32 -// CHECK-NEXT: %sub_named_2 = "llvm.inline_asm"(%sub_named, %sub_named_1) <{"asm_string" = "sub $0, $1, $2", "constraints" = "=r,rI,rI", "asm_dialect" = 0 : i64}> : (i32, i32) -> i32 +// CHECK-NEXT: %sub_named_2 = "llvm.inline_asm"(%sub_named, %sub_named_1) <{asm_string = "sub $0, $1, $2", constraints = "=r,rI,rI", asm_dialect = 0 : i64}> : (i32, i32) -> i32 // CHECK-NEXT: %sub_named_3 = builtin.unrealized_conversion_cast %sub_named_2 : i32 to !riscv.reg %div_named = riscv.div %a0, %a0 : (!riscv.reg, !riscv.reg) -> !riscv.reg // CHECK-NEXT: %div_named = builtin.unrealized_conversion_cast %a0_1 : !riscv.reg to i32 // CHECK-NEXT: %div_named_1 = builtin.unrealized_conversion_cast %a0_1 : !riscv.reg to i32 -// CHECK-NEXT: %div_named_2 = "llvm.inline_asm"(%div_named, %div_named_1) <{"asm_string" = "div $0, $1, $2", "constraints" = "=r,rI,rI", "asm_dialect" = 0 : i64}> : (i32, i32) -> i32 +// CHECK-NEXT: %div_named_2 = "llvm.inline_asm"(%div_named, %div_named_1) <{asm_string = "div $0, $1, $2", constraints = "=r,rI,rI", asm_dialect = 0 : i64}> : (i32, i32) -> i32 // CHECK-NEXT: %div_named_3 = builtin.unrealized_conversion_cast %div_named_2 : i32 to !riscv.reg // csr instructions %csrss = riscv.csrrs %x0, 3860, "r" : (!riscv.reg) -> !riscv.reg -// CHECK-NEXT: %csrss = "llvm.inline_asm"() <{"asm_string" = "csrrs $0, 3860, x0", "constraints" = "=r", "asm_dialect" = 0 : i64}> : () -> i32 +// CHECK-NEXT: %csrss = "llvm.inline_asm"() <{asm_string = "csrrs $0, 3860, x0", constraints = "=r", asm_dialect = 0 : i64}> : () -> i32 // CHECK-NEXT: %csrss_1 = builtin.unrealized_conversion_cast %csrss : i32 to !riscv.reg %csrrs = riscv.csrrs %x0, 1986 : (!riscv.reg) -> !riscv.reg // CHECK-NEXT: %csrrs = riscv.get_register : !riscv.reg -// CHECK-NEXT: "llvm.inline_asm"() <{"asm_string" = "csrrs x0, 1986, x0", "constraints" = "", "asm_dialect" = 0 : i64}> : () -> () +// CHECK-NEXT: "llvm.inline_asm"() <{asm_string = "csrrs x0, 1986, x0", constraints = "", asm_dialect = 0 : i64}> : () -> () %csrrci = riscv.csrrci 1984, 1 : () -> !riscv.reg -// CHECK-NEXT: %csrrci = "llvm.inline_asm"() <{"asm_string" = "csrrci $0, 1984, 1", "constraints" = "=r", "asm_dialect" = 0 : i64}> : () -> i32 +// CHECK-NEXT: %csrrci = "llvm.inline_asm"() <{asm_string = "csrrci $0, 1984, 1", constraints = "=r", asm_dialect = 0 : i64}> : () -> i32 // CHECK-NEXT: %csrrci_1 = builtin.unrealized_conversion_cast %csrrci : i32 to !riscv.reg @@ -70,25 +70,25 @@ riscv_snitch.dmsrc %reg, %reg : (!riscv.reg, !riscv.reg) -> () // CHECK-NEXT: %0 = builtin.unrealized_conversion_cast %reg_1 : !riscv.reg to i32 // CHECK-NEXT: %1 = builtin.unrealized_conversion_cast %reg_1 : !riscv.reg to i32 -// CHECK-NEXT: "llvm.inline_asm"(%0, %1) <{"asm_string" = ".insn r 0x2b, 0, 0, x0, $0, $1", "constraints" = "rI,rI", "asm_dialect" = 0 : i64}> : (i32, i32) -> () +// CHECK-NEXT: "llvm.inline_asm"(%0, %1) <{asm_string = ".insn r 0x2b, 0, 0, x0, $0, $1", constraints = "rI,rI", asm_dialect = 0 : i64}> : (i32, i32) -> () riscv_snitch.dmdst %reg, %reg : (!riscv.reg, !riscv.reg) -> () // CHECK-NEXT: %2 = builtin.unrealized_conversion_cast %reg_1 : !riscv.reg to i32 // CHECK-NEXT: %3 = builtin.unrealized_conversion_cast %reg_1 : !riscv.reg to i32 -// CHECK-NEXT: "llvm.inline_asm"(%2, %3) <{"asm_string" = ".insn r 0x2b, 0, 1, x0, $0, $1", "constraints" = "rI,rI", "asm_dialect" = 0 : i64}> : (i32, i32) -> () +// CHECK-NEXT: "llvm.inline_asm"(%2, %3) <{asm_string = ".insn r 0x2b, 0, 1, x0, $0, $1", constraints = "rI,rI", asm_dialect = 0 : i64}> : (i32, i32) -> () riscv_snitch.dmstr %reg, %reg : (!riscv.reg, !riscv.reg) -> () // CHECK-NEXT: %4 = builtin.unrealized_conversion_cast %reg_1 : !riscv.reg to i32 // CHECK-NEXT: %5 = builtin.unrealized_conversion_cast %reg_1 : !riscv.reg to i32 -// CHECK-NEXT: "llvm.inline_asm"(%4, %5) <{"asm_string" = ".insn r 0x2b, 0, 6, x0, $0, $1", "constraints" = "rI,rI", "asm_dialect" = 0 : i64}> : (i32, i32) -> () +// CHECK-NEXT: "llvm.inline_asm"(%4, %5) <{asm_string = ".insn r 0x2b, 0, 6, x0, $0, $1", constraints = "rI,rI", asm_dialect = 0 : i64}> : (i32, i32) -> () riscv_snitch.dmrep %reg : (!riscv.reg) -> () // CHECK-NEXT: %6 = builtin.unrealized_conversion_cast %reg_1 : !riscv.reg to i32 -// CHECK-NEXT: "llvm.inline_asm"(%6) <{"asm_string" = ".insn r 0x2b, 0, 7, x0, $0, x0", "constraints" = "rI", "asm_dialect" = 0 : i64}> : (i32) -> () +// CHECK-NEXT: "llvm.inline_asm"(%6) <{asm_string = ".insn r 0x2b, 0, 7, x0, $0, x0", constraints = "rI", asm_dialect = 0 : i64}> : (i32) -> () %dmcpyi = riscv_snitch.dmcpyi %reg, 2 : (!riscv.reg) -> !riscv.reg // CHECK-NEXT: %dmcpyi = builtin.unrealized_conversion_cast %reg_1 : !riscv.reg to i32 -// CHECK-NEXT: %dmcpyi_1 = "llvm.inline_asm"(%dmcpyi) <{"asm_string" = ".insn r 0x2b, 0, 2, $0, $1, 2", "constraints" = "=r,rI", "asm_dialect" = 0 : i64}> : (i32) -> i32 +// CHECK-NEXT: %dmcpyi_1 = "llvm.inline_asm"(%dmcpyi) <{asm_string = ".insn r 0x2b, 0, 2, $0, $1, 2", constraints = "=r,rI", asm_dialect = 0 : i64}> : (i32) -> i32 // CHECK-NEXT: %dmcpyi_2 = builtin.unrealized_conversion_cast %dmcpyi_1 : i32 to !riscv.reg @@ -97,20 +97,20 @@ riscv_snitch.dmrep %reg : (!riscv.reg) -> () // ------------------------------------------------------- // // COMPACT: builtin.module { -// COMPACT-NEXT: %reg = "llvm.inline_asm"() <{"asm_string" = "li $0, 0", "constraints" = "=r", "asm_dialect" = 0 : i64}> : () -> i32 -// COMPACT-NEXT: %a0 = "llvm.inline_asm"() <{"asm_string" = "li $0, 0", "constraints" = "=r", "asm_dialect" = 0 : i64}> : () -> i32 -// COMPACT-NEXT: %li = "llvm.inline_asm"() <{"asm_string" = "li $0, 0", "constraints" = "=r", "asm_dialect" = 0 : i64}> : () -> i32 -// COMPACT-NEXT: %sub = "llvm.inline_asm"(%reg, %reg) <{"asm_string" = "sub $0, $1, $2", "constraints" = "=r,rI,rI", "asm_dialect" = 0 : i64}> : (i32, i32) -> i32 -// COMPACT-NEXT: %div = "llvm.inline_asm"(%reg, %reg) <{"asm_string" = "div $0, $1, $2", "constraints" = "=r,rI,rI", "asm_dialect" = 0 : i64}> : (i32, i32) -> i32 -// COMPACT-NEXT: %li_named = "llvm.inline_asm"() <{"asm_string" = "li $0, 0", "constraints" = "=r", "asm_dialect" = 0 : i64}> : () -> i32 -// COMPACT-NEXT: %sub_named = "llvm.inline_asm"(%a0, %a0) <{"asm_string" = "sub $0, $1, $2", "constraints" = "=r,rI,rI", "asm_dialect" = 0 : i64}> : (i32, i32) -> i32 -// COMPACT-NEXT: %div_named = "llvm.inline_asm"(%a0, %a0) <{"asm_string" = "div $0, $1, $2", "constraints" = "=r,rI,rI", "asm_dialect" = 0 : i64}> : (i32, i32) -> i32 -// COMPACT-NEXT: %csrss = "llvm.inline_asm"() <{"asm_string" = "csrrs $0, 3860, x0", "constraints" = "=r", "asm_dialect" = 0 : i64}> : () -> i32 -// COMPACT-NEXT: "llvm.inline_asm"() <{"asm_string" = "csrrs x0, 1986, x0", "constraints" = "", "asm_dialect" = 0 : i64}> : () -> () -// COMPACT-NEXT: %csrrci = "llvm.inline_asm"() <{"asm_string" = "csrrci $0, 1984, 1", "constraints" = "=r", "asm_dialect" = 0 : i64}> : () -> i32 -// COMPACT-NEXT: "llvm.inline_asm"(%reg, %reg) <{"asm_string" = ".insn r 0x2b, 0, 0, x0, $0, $1", "constraints" = "rI,rI", "asm_dialect" = 0 : i64}> : (i32, i32) -> () -// COMPACT-NEXT: "llvm.inline_asm"(%reg, %reg) <{"asm_string" = ".insn r 0x2b, 0, 1, x0, $0, $1", "constraints" = "rI,rI", "asm_dialect" = 0 : i64}> : (i32, i32) -> () -// COMPACT-NEXT: "llvm.inline_asm"(%reg, %reg) <{"asm_string" = ".insn r 0x2b, 0, 6, x0, $0, $1", "constraints" = "rI,rI", "asm_dialect" = 0 : i64}> : (i32, i32) -> () -// COMPACT-NEXT: "llvm.inline_asm"(%reg) <{"asm_string" = ".insn r 0x2b, 0, 7, x0, $0, x0", "constraints" = "rI", "asm_dialect" = 0 : i64}> : (i32) -> () -// COMPACT-NEXT: %dmcpyi = "llvm.inline_asm"(%reg) <{"asm_string" = ".insn r 0x2b, 0, 2, $0, $1, 2", "constraints" = "=r,rI", "asm_dialect" = 0 : i64}> : (i32) -> i32 +// COMPACT-NEXT: %reg = "llvm.inline_asm"() <{asm_string = "li $0, 0", constraints = "=r", asm_dialect = 0 : i64}> : () -> i32 +// COMPACT-NEXT: %a0 = "llvm.inline_asm"() <{asm_string = "li $0, 0", constraints = "=r", asm_dialect = 0 : i64}> : () -> i32 +// COMPACT-NEXT: %li = "llvm.inline_asm"() <{asm_string = "li $0, 0", constraints = "=r", asm_dialect = 0 : i64}> : () -> i32 +// COMPACT-NEXT: %sub = "llvm.inline_asm"(%reg, %reg) <{asm_string = "sub $0, $1, $2", constraints = "=r,rI,rI", asm_dialect = 0 : i64}> : (i32, i32) -> i32 +// COMPACT-NEXT: %div = "llvm.inline_asm"(%reg, %reg) <{asm_string = "div $0, $1, $2", constraints = "=r,rI,rI", asm_dialect = 0 : i64}> : (i32, i32) -> i32 +// COMPACT-NEXT: %li_named = "llvm.inline_asm"() <{asm_string = "li $0, 0", constraints = "=r", asm_dialect = 0 : i64}> : () -> i32 +// COMPACT-NEXT: %sub_named = "llvm.inline_asm"(%a0, %a0) <{asm_string = "sub $0, $1, $2", constraints = "=r,rI,rI", asm_dialect = 0 : i64}> : (i32, i32) -> i32 +// COMPACT-NEXT: %div_named = "llvm.inline_asm"(%a0, %a0) <{asm_string = "div $0, $1, $2", constraints = "=r,rI,rI", asm_dialect = 0 : i64}> : (i32, i32) -> i32 +// COMPACT-NEXT: %csrss = "llvm.inline_asm"() <{asm_string = "csrrs $0, 3860, x0", constraints = "=r", asm_dialect = 0 : i64}> : () -> i32 +// COMPACT-NEXT: "llvm.inline_asm"() <{asm_string = "csrrs x0, 1986, x0", constraints = "", asm_dialect = 0 : i64}> : () -> () +// COMPACT-NEXT: %csrrci = "llvm.inline_asm"() <{asm_string = "csrrci $0, 1984, 1", constraints = "=r", asm_dialect = 0 : i64}> : () -> i32 +// COMPACT-NEXT: "llvm.inline_asm"(%reg, %reg) <{asm_string = ".insn r 0x2b, 0, 0, x0, $0, $1", constraints = "rI,rI", asm_dialect = 0 : i64}> : (i32, i32) -> () +// COMPACT-NEXT: "llvm.inline_asm"(%reg, %reg) <{asm_string = ".insn r 0x2b, 0, 1, x0, $0, $1", constraints = "rI,rI", asm_dialect = 0 : i64}> : (i32, i32) -> () +// COMPACT-NEXT: "llvm.inline_asm"(%reg, %reg) <{asm_string = ".insn r 0x2b, 0, 6, x0, $0, $1", constraints = "rI,rI", asm_dialect = 0 : i64}> : (i32, i32) -> () +// COMPACT-NEXT: "llvm.inline_asm"(%reg) <{asm_string = ".insn r 0x2b, 0, 7, x0, $0, x0", constraints = "rI", asm_dialect = 0 : i64}> : (i32) -> () +// COMPACT-NEXT: %dmcpyi = "llvm.inline_asm"(%reg) <{asm_string = ".insn r 0x2b, 0, 2, $0, $1, 2", constraints = "=r,rI", asm_dialect = 0 : i64}> : (i32) -> i32 // COMPACT-NEXT: } diff --git a/tests/filecheck/transforms/cse.mlir b/tests/filecheck/transforms/cse.mlir index 68da590f19..e9997c06a1 100644 --- a/tests/filecheck/transforms/cse.mlir +++ b/tests/filecheck/transforms/cse.mlir @@ -127,7 +127,7 @@ func.func @different_ops() -> (i32, i32) { } // CHECK: %0 = arith.constant 1 : i32 -// CHECK-NEXT: "affine.for"() <{"lowerBoundMap" = affine_map<() -> (0)>, "operandSegmentSizes" = array, "step" = 1 : index, "upperBoundMap" = affine_map<() -> (4)>}> ({ +// CHECK-NEXT: "affine.for"() <{lowerBoundMap = affine_map<() -> (0)>, operandSegmentSizes = array, step = 1 : index, upperBoundMap = affine_map<() -> (4)>}> ({ // CHECK-NEXT: ^0(%arg0 : index): // CHECK-NEXT: "foo"(%0, %0) : (i32, i32) -> () // CHECK-NEXT: "affine.yield"() : () -> () @@ -176,7 +176,7 @@ func.func @down_propagate() -> i32 { func.return %32 : i32 } -// CHECK: "affine.for"() <{"lowerBoundMap" = affine_map<() -> (0)>, "operandSegmentSizes" = array, "step" = 1 : index, "upperBoundMap" = affine_map<() -> (4)>}> ({ +// CHECK: "affine.for"() <{lowerBoundMap = affine_map<() -> (0)>, operandSegmentSizes = array, step = 1 : index, upperBoundMap = affine_map<() -> (4)>}> ({ // CHECK-NEXT: ^0(%arg0 : index): // CHECK-NEXT: %0 = arith.constant 1 : i32 // CHECK-NEXT: "foo"(%0) : (i32) -> () @@ -445,7 +445,7 @@ func.func @cse_single_block_ops_identical_bodies(%arg0_8 : tensor, %arg // CHECK-NEXT: %0 = "test.pureop"(%arg0, %arg1) ({ // CHECK-NEXT: ^0(%arg4 : f32, %arg5 : f32): // CHECK-NEXT: %1 = arith.divf %arg4, %arg5 : f32 -// CHECK-NEXT: %2 = "arith.remf"(%arg4, %arg2) <{"fastmath" = #arith.fastmath}> : (f32, f32) -> f32 +// CHECK-NEXT: %2 = "arith.remf"(%arg4, %arg2) <{fastmath = #arith.fastmath}> : (f32, f32) -> f32 // CHECK-NEXT: %3 = arith.select %arg3, %1, %2 : f32 // CHECK-NEXT: "test.region_yield"(%3) : (f32) -> () // CHECK-NEXT: }) : (tensor, tensor) -> tensor @@ -475,14 +475,14 @@ func.func @no_cse_single_block_ops_different_bodies(%arg0_9 : tensor, % // CHECK-NEXT: %0 = "test.pureop"(%arg0, %arg1) ({ // CHECK-NEXT: ^0(%arg4 : f32, %arg5 : f32): // CHECK-NEXT: %1 = arith.divf %arg4, %arg5 : f32 -// CHECK-NEXT: %2 = "arith.remf"(%arg4, %arg2) <{"fastmath" = #arith.fastmath}> : (f32, f32) -> f32 +// CHECK-NEXT: %2 = "arith.remf"(%arg4, %arg2) <{fastmath = #arith.fastmath}> : (f32, f32) -> f32 // CHECK-NEXT: %3 = arith.select %arg3, %1, %2 : f32 // CHECK-NEXT: "test.region_yield"(%3) : (f32) -> () // CHECK-NEXT: }) : (tensor, tensor) -> tensor // CHECK-NEXT: %4 = "test.pureop"(%arg0, %arg1) ({ // CHECK-NEXT: ^1(%arg4_1 : f32, %arg5_1 : f32): // CHECK-NEXT: %5 = arith.divf %arg4_1, %arg5_1 : f32 -// CHECK-NEXT: %6 = "arith.remf"(%arg4_1, %arg2) <{"fastmath" = #arith.fastmath}> : (f32, f32) -> f32 +// CHECK-NEXT: %6 = "arith.remf"(%arg4_1, %arg2) <{fastmath = #arith.fastmath}> : (f32, f32) -> f32 // CHECK-NEXT: %7 = arith.select %arg3, %6, %5 : f32 // CHECK-NEXT: "test.region_yield"(%7) : (f32) -> () // CHECK-NEXT: }) : (tensor, tensor) -> tensor diff --git a/tests/filecheck/transforms/csl-stencil-handle-async-flow.mlir b/tests/filecheck/transforms/csl-stencil-handle-async-flow.mlir index b45ba3d5bb..3982096f37 100644 --- a/tests/filecheck/transforms/csl-stencil-handle-async-flow.mlir +++ b/tests/filecheck/transforms/csl-stencil-handle-async-flow.mlir @@ -67,14 +67,14 @@ }) : () -> () -// CHECK: "csl_wrapper.module"() <{"width" = 1022 : i16, "height" = 510 : i16, "params" = [#csl_wrapper.param<"z_dim" default=512 : i16>, #csl_wrapper.param<"pattern" default=2 : i16>, #csl_wrapper.param<"num_chunks" default=2 : i16>, #csl_wrapper.param<"chunk_size" default=255 : i16>, #csl_wrapper.param<"padded_z_dim" default=510 : i16>], "program_name" = "gauss_seidel_func"}> ({ +// CHECK: "csl_wrapper.module"() <{width = 1022 : i16, height = 510 : i16, params = [#csl_wrapper.param<"z_dim" default=512 : i16>, #csl_wrapper.param<"pattern" default=2 : i16>, #csl_wrapper.param<"num_chunks" default=2 : i16>, #csl_wrapper.param<"chunk_size" default=255 : i16>, #csl_wrapper.param<"padded_z_dim" default=510 : i16>], program_name = "gauss_seidel_func"}> ({ // CHECK-NEXT: ^0(%0 : i16, %1 : i16, %2 : i16, %3 : i16, %4 : i16, %5 : i16, %6 : i16, %7 : i16, %8 : i16): // CHECK-NEXT: %9 = arith.constant 0 : i16 // CHECK-NEXT: %10 = "csl.get_color"(%9) : (i16) -> !csl.color -// CHECK-NEXT: %11 = "csl_wrapper.import"(%2, %3, %10) <{"module" = "", "fields" = ["width", "height", "LAUNCH"]}> : (i16, i16, !csl.color) -> !csl.imported_module -// CHECK-NEXT: %12 = "csl_wrapper.import"(%5, %2, %3) <{"module" = "routes.csl", "fields" = ["pattern", "peWidth", "peHeight"]}> : (i16, i16, i16) -> !csl.imported_module -// CHECK-NEXT: %13 = "csl.member_call"(%12, %0, %1, %2, %3, %5) <{"field" = "computeAllRoutes"}> : (!csl.imported_module, i16, i16, i16, i16, i16) -> !csl.comptime_struct -// CHECK-NEXT: %14 = "csl.member_call"(%11, %0) <{"field" = "get_params"}> : (!csl.imported_module, i16) -> !csl.comptime_struct +// CHECK-NEXT: %11 = "csl_wrapper.import"(%2, %3, %10) <{module = "", fields = ["width", "height", "LAUNCH"]}> : (i16, i16, !csl.color) -> !csl.imported_module +// CHECK-NEXT: %12 = "csl_wrapper.import"(%5, %2, %3) <{module = "routes.csl", fields = ["pattern", "peWidth", "peHeight"]}> : (i16, i16, i16) -> !csl.imported_module +// CHECK-NEXT: %13 = "csl.member_call"(%12, %0, %1, %2, %3, %5) <{field = "computeAllRoutes"}> : (!csl.imported_module, i16, i16, i16, i16, i16) -> !csl.comptime_struct +// CHECK-NEXT: %14 = "csl.member_call"(%11, %0) <{field = "get_params"}> : (!csl.imported_module, i16) -> !csl.comptime_struct // CHECK-NEXT: %15 = arith.constant 1 : i16 // CHECK-NEXT: %16 = arith.subi %15, %5 : i16 // CHECK-NEXT: %17 = arith.subi %2, %0 : i16 @@ -86,20 +86,20 @@ // CHECK-NEXT: %23 = arith.ori %19, %20 : i1 // CHECK-NEXT: %24 = arith.ori %23, %21 : i1 // CHECK-NEXT: %25 = arith.ori %24, %22 : i1 -// CHECK-NEXT: "csl_wrapper.yield"(%14, %13, %25) <{"fields" = ["memcpy_params", "stencil_comms_params", "isBorderRegionPE"]}> : (!csl.comptime_struct, !csl.comptime_struct, i1) -> () +// CHECK-NEXT: "csl_wrapper.yield"(%14, %13, %25) <{fields = ["memcpy_params", "stencil_comms_params", "isBorderRegionPE"]}> : (!csl.comptime_struct, !csl.comptime_struct, i1) -> () // CHECK-NEXT: }, { // CHECK-NEXT: ^1(%26 : i16, %27 : i16, %28 : i16, %29 : i16, %30 : i16, %31 : i16, %32 : i16, %memcpy_params : !csl.comptime_struct, %stencil_comms_params : !csl.comptime_struct, %isBorderRegionPE : i1): -// CHECK-NEXT: %33 = "csl_wrapper.import"(%memcpy_params) <{"module" = "", "fields" = [""]}> : (!csl.comptime_struct) -> !csl.imported_module -// CHECK-NEXT: %34 = "csl_wrapper.import"(%29, %31, %stencil_comms_params) <{"module" = "stencil_comms.csl", "fields" = ["pattern", "chunkSize", ""]}> : (i16, i16, !csl.comptime_struct) -> !csl.imported_module +// CHECK-NEXT: %33 = "csl_wrapper.import"(%memcpy_params) <{module = "", fields = [""]}> : (!csl.comptime_struct) -> !csl.imported_module +// CHECK-NEXT: %34 = "csl_wrapper.import"(%29, %31, %stencil_comms_params) <{module = "stencil_comms.csl", fields = ["pattern", "chunkSize", ""]}> : (i16, i16, !csl.comptime_struct) -> !csl.imported_module // CHECK-NEXT: %arg0 = memref.alloc() : memref<512xf32> // CHECK-NEXT: %arg1 = memref.alloc() : memref<512xf32> // CHECK-NEXT: %35 = "csl.addressof"(%arg0) : (memref<512xf32>) -> !csl.ptr, #csl> // CHECK-NEXT: %36 = "csl.addressof"(%arg1) : (memref<512xf32>) -> !csl.ptr, #csl> -// CHECK-NEXT: "csl.export"(%35) <{"var_name" = "arg0", "type" = !csl.ptr, #csl>}> : (!csl.ptr, #csl>) -> () -// CHECK-NEXT: "csl.export"(%36) <{"var_name" = "arg1", "type" = !csl.ptr, #csl>}> : (!csl.ptr, #csl>) -> () -// CHECK-NEXT: "csl.export"() <{"var_name" = @gauss_seidel_func, "type" = () -> ()}> : () -> () -// CHECK-NEXT: %37 = memref.alloc() {"alignment" = 64 : i64} : memref<510xf32> -// CHECK-NEXT: %iteration = "csl.variable"() <{"default" = 0 : ui32}> : () -> !csl.var +// CHECK-NEXT: "csl.export"(%35) <{var_name = "arg0", type = !csl.ptr, #csl>}> : (!csl.ptr, #csl>) -> () +// CHECK-NEXT: "csl.export"(%36) <{var_name = "arg1", type = !csl.ptr, #csl>}> : (!csl.ptr, #csl>) -> () +// CHECK-NEXT: "csl.export"() <{var_name = @gauss_seidel_func, type = () -> ()}> : () -> () +// CHECK-NEXT: %37 = memref.alloc() {alignment = 64 : i64} : memref<510xf32> +// CHECK-NEXT: %iteration = "csl.variable"() <{default = 0 : ui32}> : () -> !csl.var // CHECK-NEXT: %var0 = "csl.variable"() : () -> !csl.var> // CHECK-NEXT: %var1 = "csl.variable"() : () -> !csl.var> // CHECK-NEXT: csl.func @gauss_seidel_func() { @@ -111,14 +111,14 @@ // CHECK-NEXT: csl.activate local, 1 : ui5 // CHECK-NEXT: csl.return // CHECK-NEXT: } -// CHECK-NEXT: csl.task @for_cond0() attributes {"kind" = #csl, "id" = 1 : ui5}{ +// CHECK-NEXT: csl.task @for_cond0() attributes {kind = #csl, id = 1 : ui5}{ // CHECK-NEXT: %41 = arith.constant 1000 : ui32 // CHECK-NEXT: %iteration_cond = "csl.load_var"(%iteration) : (!csl.var) -> ui32 // CHECK-NEXT: %42 = arith.cmpi slt, %iteration_cond, %41 : ui32 // CHECK-NEXT: scf.if %42 { -// CHECK-NEXT: "csl.call"() <{"callee" = @for_body0}> : () -> () +// CHECK-NEXT: "csl.call"() <{callee = @for_body0}> : () -> () // CHECK-NEXT: } else { -// CHECK-NEXT: "csl.call"() <{"callee" = @for_post0}> : () -> () +// CHECK-NEXT: "csl.call"() <{callee = @for_post0}> : () -> () // CHECK-NEXT: } // CHECK-NEXT: csl.return // CHECK-NEXT: } @@ -126,7 +126,7 @@ // CHECK-NEXT: %iteration_bdy = "csl.load_var"(%iteration) : (!csl.var) -> ui32 // CHECK-NEXT: %var0_bdy = "csl.load_var"(%var0) : (!csl.var>) -> memref<512xf32> // CHECK-NEXT: %var1_bdy = "csl.load_var"(%var1) : (!csl.var>) -> memref<512xf32> -// CHECK-NEXT: csl_stencil.apply(%var0_bdy : memref<512xf32>, %37 : memref<510xf32>) outs (%var1_bdy : memref<512xf32>) <{"bounds" = #stencil.bounds<[0, 0], [1, 1]>, "num_chunks" = 2 : i64, "operandSegmentSizes" = array, "swaps" = [#csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange], "topo" = #dmp.topo<1022x510>}> ({ +// CHECK-NEXT: csl_stencil.apply(%var0_bdy : memref<512xf32>, %37 : memref<510xf32>) outs (%var1_bdy : memref<512xf32>) <{bounds = #stencil.bounds<[0, 0], [1, 1]>, num_chunks = 2 : i64, operandSegmentSizes = array, swaps = [#csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange], topo = #dmp.topo<1022x510>}> ({ // CHECK-NEXT: ^2(%arg5 : memref<4x255xf32>, %arg6 : index, %arg7 : memref<510xf32>): // CHECK-NEXT: %43 = csl_stencil.access %arg5[1, 0] : memref<4x255xf32> // CHECK-NEXT: %44 = csl_stencil.access %arg5[-1, 0] : memref<4x255xf32> @@ -145,7 +145,7 @@ // CHECK-NEXT: "csl.fadds"(%arg6_1, %arg6_1, %48) : (memref<510xf32>, memref<510xf32>, memref<510xf32, strided<[1], offset: 2>>) -> () // CHECK-NEXT: %50 = arith.constant 1.666600e-01 : f32 // CHECK-NEXT: "csl.fmuls"(%arg6_1, %arg6_1, %50) : (memref<510xf32>, memref<510xf32>, f32) -> () -// CHECK-NEXT: "csl.call"() <{"callee" = @for_inc0}> : () -> () +// CHECK-NEXT: "csl.call"() <{callee = @for_inc0}> : () -> () // CHECK-NEXT: csl_stencil.yield %arg6_1 : memref<510xf32> // CHECK-NEXT: }) to <[0, 0], [1, 1]> // CHECK-NEXT: csl.return @@ -163,10 +163,10 @@ // CHECK-NEXT: csl.return // CHECK-NEXT: } // CHECK-NEXT: csl.func @for_post0() { -// CHECK-NEXT: "csl.member_call"(%33) <{"field" = "unblock_cmd_stream"}> : (!csl.imported_module) -> () +// CHECK-NEXT: "csl.member_call"(%33) <{field = "unblock_cmd_stream"}> : (!csl.imported_module) -> () // CHECK-NEXT: csl.return // CHECK-NEXT: } -// CHECK-NEXT: "csl_wrapper.yield"() <{"fields" = []}> : () -> () +// CHECK-NEXT: "csl_wrapper.yield"() <{fields = []}> : () -> () // CHECK-NEXT: }) : () -> () "csl_wrapper.module"() <{"width" = 1022 : i16, "height" = 510 : i16, "params" = [], "program_name" = "sequential_kernels_func"}> ({ @@ -207,47 +207,47 @@ }) : () -> () -// CHECK: "csl_wrapper.module"() <{"width" = 1022 : i16, "height" = 510 : i16, "params" = [], "program_name" = "sequential_kernels_func"}> ({ +// CHECK: "csl_wrapper.module"() <{width = 1022 : i16, height = 510 : i16, params = [], program_name = "sequential_kernels_func"}> ({ // CHECK-NEXT: ^2(%45 : i16, %46 : i16, %47 : i16, %48 : i16): -// CHECK-NEXT: "csl_wrapper.yield"() <{"fields" = []}> : () -> () +// CHECK-NEXT: "csl_wrapper.yield"() <{fields = []}> : () -> () // CHECK-NEXT: }, { // CHECK-NEXT: ^3(%49 : i16, %50 : i16): -// CHECK-NEXT: %51 = "csl_wrapper.import"() <{"module" = "", "fields" = []}> : () -> !csl.imported_module +// CHECK-NEXT: %51 = "csl_wrapper.import"() <{module = "", fields = []}> : () -> !csl.imported_module // CHECK-NEXT: %arg0_1 = memref.alloc() : memref<512xf32> // CHECK-NEXT: %arg1_1 = memref.alloc() : memref<512xf32> // CHECK-NEXT: %52 = "csl.addressof"(%arg0_1) : (memref<512xf32>) -> !csl.ptr, #csl> // CHECK-NEXT: %53 = "csl.addressof"(%arg1_1) : (memref<512xf32>) -> !csl.ptr, #csl> -// CHECK-NEXT: "csl.export"(%52) <{"var_name" = "arg0", "type" = !csl.ptr, #csl>}> : (!csl.ptr, #csl>) -> () -// CHECK-NEXT: "csl.export"(%53) <{"var_name" = "arg1", "type" = !csl.ptr, #csl>}> : (!csl.ptr, #csl>) -> () -// CHECK-NEXT: "csl.export"() <{"var_name" = @gauss_seidel_func, "type" = () -> ()}> : () -> () -// CHECK-NEXT: %54 = memref.alloc() {"alignment" = 64 : i64} : memref<510xf32> +// CHECK-NEXT: "csl.export"(%52) <{var_name = "arg0", type = !csl.ptr, #csl>}> : (!csl.ptr, #csl>) -> () +// CHECK-NEXT: "csl.export"(%53) <{var_name = "arg1", type = !csl.ptr, #csl>}> : (!csl.ptr, #csl>) -> () +// CHECK-NEXT: "csl.export"() <{var_name = @gauss_seidel_func, type = () -> ()}> : () -> () +// CHECK-NEXT: %54 = memref.alloc() {alignment = 64 : i64} : memref<510xf32> // CHECK-NEXT: csl.func @sequential_kernels_func() { -// CHECK-NEXT: csl_stencil.apply(%arg0_1 : memref<512xf32>, %54 : memref<510xf32>) outs (%arg1_1 : memref<512xf32>) <{"bounds" = #stencil.bounds<[0, 0], [1, 1]>, "num_chunks" = 2 : i64, "operandSegmentSizes" = array, "swaps" = [#csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange], "topo" = #dmp.topo<1022x510>}> ({ +// CHECK-NEXT: csl_stencil.apply(%arg0_1 : memref<512xf32>, %54 : memref<510xf32>) outs (%arg1_1 : memref<512xf32>) <{bounds = #stencil.bounds<[0, 0], [1, 1]>, num_chunks = 2 : i64, operandSegmentSizes = array, swaps = [#csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange], topo = #dmp.topo<1022x510>}> ({ // CHECK-NEXT: ^4(%arg5 : memref<4x255xf32>, %arg6 : index, %arg7 : memref<510xf32>): // CHECK-NEXT: csl_stencil.yield %arg7 : memref<510xf32> // CHECK-NEXT: }, { // CHECK-NEXT: ^5(%arg5_1 : memref<512xf32>, %arg6_1 : memref<510xf32>): // CHECK-NEXT: %55 = arith.constant 1.666600e-01 : f32 -// CHECK-NEXT: "csl.call"() <{"callee" = @step0}> : () -> () +// CHECK-NEXT: "csl.call"() <{callee = @step0}> : () -> () // CHECK-NEXT: csl_stencil.yield %arg6_1 : memref<510xf32> // CHECK-NEXT: }) to <[0, 0], [1, 1]> // CHECK-NEXT: csl.return // CHECK-NEXT: } // CHECK-NEXT: csl.func @step0() { -// CHECK-NEXT: csl_stencil.apply(%arg1_1 : memref<512xf32>, %54 : memref<510xf32>) outs (%arg0_1 : memref<512xf32>) <{"bounds" = #stencil.bounds<[0, 0], [1, 1]>, "num_chunks" = 2 : i64, "operandSegmentSizes" = array, "swaps" = [#csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange], "topo" = #dmp.topo<1022x510>}> ({ +// CHECK-NEXT: csl_stencil.apply(%arg1_1 : memref<512xf32>, %54 : memref<510xf32>) outs (%arg0_1 : memref<512xf32>) <{bounds = #stencil.bounds<[0, 0], [1, 1]>, num_chunks = 2 : i64, operandSegmentSizes = array, swaps = [#csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange], topo = #dmp.topo<1022x510>}> ({ // CHECK-NEXT: ^4(%arg5 : memref<4x255xf32>, %arg6 : index, %arg7 : memref<510xf32>): // CHECK-NEXT: csl_stencil.yield %arg7 : memref<510xf32> // CHECK-NEXT: }, { // CHECK-NEXT: ^5(%arg5_1 : memref<512xf32>, %arg6_1 : memref<510xf32>): // CHECK-NEXT: %55 = arith.constant 0.123456702 : f32 -// CHECK-NEXT: "csl.call"() <{"callee" = @step1}> : () -> () +// CHECK-NEXT: "csl.call"() <{callee = @step1}> : () -> () // CHECK-NEXT: csl_stencil.yield %arg6_1 : memref<510xf32> // CHECK-NEXT: }) to <[0, 0], [1, 1]> // CHECK-NEXT: csl.return // CHECK-NEXT: } // CHECK-NEXT: csl.func @step1() { -// CHECK-NEXT: "csl.member_call"(%51) <{"field" = "unblock_cmd_stream"}> : (!csl.imported_module) -> () +// CHECK-NEXT: "csl.member_call"(%51) <{field = "unblock_cmd_stream"}> : (!csl.imported_module) -> () // CHECK-NEXT: csl.return // CHECK-NEXT: } -// CHECK-NEXT: "csl_wrapper.yield"() <{"fields" = []}> : () -> () +// CHECK-NEXT: "csl_wrapper.yield"() <{fields = []}> : () -> () // CHECK-NEXT: }) : () -> () diff --git a/tests/filecheck/transforms/csl-stencil-materialize-stores.mlir b/tests/filecheck/transforms/csl-stencil-materialize-stores.mlir index d0065b95fa..133f14c6dd 100644 --- a/tests/filecheck/transforms/csl-stencil-materialize-stores.mlir +++ b/tests/filecheck/transforms/csl-stencil-materialize-stores.mlir @@ -65,14 +65,14 @@ builtin.module { } // CHECK-NEXT: builtin.module { -// CHECK-NEXT: "csl_wrapper.module"() <{"height" = 512 : i16, "params" = [#csl_wrapper.param<"z_dim" default=512 : i16>, #csl_wrapper.param<"pattern" default=2 : i16>, #csl_wrapper.param<"num_chunks" default=1 : i16>, #csl_wrapper.param<"chunk_size" default=510 : i16>, #csl_wrapper.param<"padded_z_dim" default=510 : i16>], "program_name" = "gauss_seidel", "width" = 1024 : i16}> ({ +// CHECK-NEXT: "csl_wrapper.module"() <{height = 512 : i16, params = [#csl_wrapper.param<"z_dim" default=512 : i16>, #csl_wrapper.param<"pattern" default=2 : i16>, #csl_wrapper.param<"num_chunks" default=1 : i16>, #csl_wrapper.param<"chunk_size" default=510 : i16>, #csl_wrapper.param<"padded_z_dim" default=510 : i16>], program_name = "gauss_seidel", width = 1024 : i16}> ({ // CHECK-NEXT: ^0(%arg0 : i16, %arg1 : i16, %arg2 : i16, %arg3 : i16, %arg4 : i16, %arg5 : i16, %arg6 : i16, %arg7 : i16, %arg8 : i16): // CHECK-NEXT: %0 = arith.constant 0 : i16 // CHECK-NEXT: %1 = "csl.get_color"(%0) : (i16) -> !csl.color -// CHECK-NEXT: %2 = "csl_wrapper.import"(%arg2, %arg3, %1) <{"fields" = ["width", "height", "LAUNCH"], "module" = ""}> : (i16, i16, !csl.color) -> !csl.imported_module -// CHECK-NEXT: %3 = "csl_wrapper.import"(%arg5, %arg2, %arg3) <{"fields" = ["pattern", "peWidth", "peHeight"], "module" = "routes.csl"}> : (i16, i16, i16) -> !csl.imported_module -// CHECK-NEXT: %4 = "csl.member_call"(%3, %arg0, %arg1, %arg2, %arg3, %arg5) <{"field" = "computeAllRoutes"}> : (!csl.imported_module, i16, i16, i16, i16, i16) -> !csl.comptime_struct -// CHECK-NEXT: %5 = "csl.member_call"(%2, %arg0) <{"field" = "get_params"}> : (!csl.imported_module, i16) -> !csl.comptime_struct +// CHECK-NEXT: %2 = "csl_wrapper.import"(%arg2, %arg3, %1) <{fields = ["width", "height", "LAUNCH"], module = ""}> : (i16, i16, !csl.color) -> !csl.imported_module +// CHECK-NEXT: %3 = "csl_wrapper.import"(%arg5, %arg2, %arg3) <{fields = ["pattern", "peWidth", "peHeight"], module = "routes.csl"}> : (i16, i16, i16) -> !csl.imported_module +// CHECK-NEXT: %4 = "csl.member_call"(%3, %arg0, %arg1, %arg2, %arg3, %arg5) <{field = "computeAllRoutes"}> : (!csl.imported_module, i16, i16, i16, i16, i16) -> !csl.comptime_struct +// CHECK-NEXT: %5 = "csl.member_call"(%2, %arg0) <{field = "get_params"}> : (!csl.imported_module, i16) -> !csl.comptime_struct // CHECK-NEXT: %6 = arith.constant 1 : i16 // CHECK-NEXT: %7 = arith.subi %arg5, %6 : i16 // CHECK-NEXT: %8 = arith.subi %arg2, %arg0 : i16 @@ -84,21 +84,21 @@ builtin.module { // CHECK-NEXT: %14 = arith.ori %10, %11 : i1 // CHECK-NEXT: %15 = arith.ori %14, %12 : i1 // CHECK-NEXT: %16 = arith.ori %15, %13 : i1 -// CHECK-NEXT: "csl_wrapper.yield"(%5, %4, %16) <{"fields" = ["memcpy_params", "stencil_comms_params", "isBorderRegionPE"]}> : (!csl.comptime_struct, !csl.comptime_struct, i1) -> () +// CHECK-NEXT: "csl_wrapper.yield"(%5, %4, %16) <{fields = ["memcpy_params", "stencil_comms_params", "isBorderRegionPE"]}> : (!csl.comptime_struct, !csl.comptime_struct, i1) -> () // CHECK-NEXT: }, { // CHECK-NEXT: ^1(%arg0_1 : i16, %arg1_1 : i16, %arg2_1 : i16, %arg3_1 : i16, %arg4_1 : i16, %arg5_1 : i16, %arg6_1 : i16, %arg7_1 : !csl.comptime_struct, %arg8_1 : !csl.comptime_struct, %arg9 : i1): -// CHECK-NEXT: %17 = "csl_wrapper.import"(%arg7_1) <{"fields" = [""], "module" = ""}> : (!csl.comptime_struct) -> !csl.imported_module -// CHECK-NEXT: %18 = "csl_wrapper.import"(%arg3_1, %arg5_1, %arg8_1) <{"fields" = ["pattern", "chunkSize", ""], "module" = "stencil_comms.csl"}> : (i16, i16, !csl.comptime_struct) -> !csl.imported_module +// CHECK-NEXT: %17 = "csl_wrapper.import"(%arg7_1) <{fields = [""], module = ""}> : (!csl.comptime_struct) -> !csl.imported_module +// CHECK-NEXT: %18 = "csl_wrapper.import"(%arg3_1, %arg5_1, %arg8_1) <{fields = ["pattern", "chunkSize", ""], module = "stencil_comms.csl"}> : (i16, i16, !csl.comptime_struct) -> !csl.imported_module // CHECK-NEXT: %19 = memref.alloc() : memref<512xf32> // CHECK-NEXT: %20 = memref.alloc() : memref<512xf32> // CHECK-NEXT: %21 = "csl.addressof"(%19) : (memref<512xf32>) -> !csl.ptr, #csl> // CHECK-NEXT: %22 = "csl.addressof"(%20) : (memref<512xf32>) -> !csl.ptr, #csl> -// CHECK-NEXT: "csl.export"(%21) <{"type" = !csl.ptr, #csl>, "var_name" = "a"}> : (!csl.ptr, #csl>) -> () -// CHECK-NEXT: "csl.export"(%22) <{"type" = !csl.ptr, #csl>, "var_name" = "b"}> : (!csl.ptr, #csl>) -> () -// CHECK-NEXT: "csl.export"() <{"type" = () -> (), "var_name" = @gauss_seidel}> : () -> () +// CHECK-NEXT: "csl.export"(%21) <{type = !csl.ptr, #csl>, var_name = "a"}> : (!csl.ptr, #csl>) -> () +// CHECK-NEXT: "csl.export"(%22) <{type = !csl.ptr, #csl>, var_name = "b"}> : (!csl.ptr, #csl>) -> () +// CHECK-NEXT: "csl.export"() <{type = () -> (), var_name = @gauss_seidel}> : () -> () // CHECK-NEXT: csl.func @gauss_seidel() { -// CHECK-NEXT: %23 = memref.alloc() {"alignment" = 64 : i64} : memref<510xf32> -// CHECK-NEXT: csl_stencil.apply(%19 : memref<512xf32>, %23 : memref<510xf32>, %20 : memref<512xf32>, %arg9 : i1) outs (%20 : memref<512xf32>) <{"bounds" = #stencil.bounds<[0, 0], [1, 1]>, "num_chunks" = 1 : i64, "operandSegmentSizes" = array, "swaps" = [#csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange], "topo" = #dmp.topo<1022x510>}> ({ +// CHECK-NEXT: %23 = memref.alloc() {alignment = 64 : i64} : memref<510xf32> +// CHECK-NEXT: csl_stencil.apply(%19 : memref<512xf32>, %23 : memref<510xf32>, %20 : memref<512xf32>, %arg9 : i1) outs (%20 : memref<512xf32>) <{bounds = #stencil.bounds<[0, 0], [1, 1]>, num_chunks = 1 : i64, operandSegmentSizes = array, swaps = [#csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange], topo = #dmp.topo<1022x510>}> ({ // CHECK-NEXT: ^2(%arg10 : memref<4x510xf32>, %arg11 : index, %arg12 : memref<510xf32>): // CHECK-NEXT: %24 = csl_stencil.access %arg10[1, 0] : memref<4x510xf32> // CHECK-NEXT: %25 = csl_stencil.access %arg10[-1, 0] : memref<4x510xf32> @@ -126,9 +126,9 @@ builtin.module { // CHECK-NEXT: } // CHECK-NEXT: csl_stencil.yield // CHECK-NEXT: }) to <[0, 0], [1, 1]> -// CHECK-NEXT: "csl.member_call"(%17) <{"field" = "unblock_cmd_stream"}> : (!csl.imported_module) -> () +// CHECK-NEXT: "csl.member_call"(%17) <{field = "unblock_cmd_stream"}> : (!csl.imported_module) -> () // CHECK-NEXT: csl.return // CHECK-NEXT: } -// CHECK-NEXT: "csl_wrapper.yield"() <{"fields" = []}> : () -> () +// CHECK-NEXT: "csl_wrapper.yield"() <{fields = []}> : () -> () // CHECK-NEXT: }) : () -> () // CHECK-NEXT: } diff --git a/tests/filecheck/transforms/csl-stencil-to-csl-wrapper.mlir b/tests/filecheck/transforms/csl-stencil-to-csl-wrapper.mlir index 32012aa3bc..584d8bc55f 100644 --- a/tests/filecheck/transforms/csl-stencil-to-csl-wrapper.mlir +++ b/tests/filecheck/transforms/csl-stencil-to-csl-wrapper.mlir @@ -33,14 +33,14 @@ func.func @gauss_seidel(%a : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>, func.return } -// CHECK: "csl_wrapper.module"() <{"width" = 1024 : i16, "height" = 512 : i16, "params" = [#csl_wrapper.param<"z_dim" default=512 : i16>, #csl_wrapper.param<"pattern" default=2 : i16>, #csl_wrapper.param<"num_chunks" default=2 : i16>, #csl_wrapper.param<"chunk_size" default=255 : i16>, #csl_wrapper.param<"padded_z_dim" default=510 : i16>], "program_name" = "gauss_seidel"}> ({ +// CHECK: "csl_wrapper.module"() <{width = 1024 : i16, height = 512 : i16, params = [#csl_wrapper.param<"z_dim" default=512 : i16>, #csl_wrapper.param<"pattern" default=2 : i16>, #csl_wrapper.param<"num_chunks" default=2 : i16>, #csl_wrapper.param<"chunk_size" default=255 : i16>, #csl_wrapper.param<"padded_z_dim" default=510 : i16>], program_name = "gauss_seidel"}> ({ // CHECK-NEXT: ^0(%0 : i16, %1 : i16, %2 : i16, %3 : i16, %4 : i16, %5 : i16, %6 : i16, %7 : i16, %8 : i16): // CHECK-NEXT: %9 = arith.constant 0 : i16 // CHECK-NEXT: %10 = "csl.get_color"(%9) : (i16) -> !csl.color -// CHECK-NEXT: %11 = "csl_wrapper.import"(%2, %3, %10) <{"module" = "", "fields" = ["width", "height", "LAUNCH"]}> : (i16, i16, !csl.color) -> !csl.imported_module -// CHECK-NEXT: %12 = "csl_wrapper.import"(%5, %2, %3) <{"module" = "routes.csl", "fields" = ["pattern", "peWidth", "peHeight"]}> : (i16, i16, i16) -> !csl.imported_module -// CHECK-NEXT: %13 = "csl.member_call"(%12, %0, %1, %2, %3, %5) <{"field" = "computeAllRoutes"}> : (!csl.imported_module, i16, i16, i16, i16, i16) -> !csl.comptime_struct -// CHECK-NEXT: %14 = "csl.member_call"(%11, %0) <{"field" = "get_params"}> : (!csl.imported_module, i16) -> !csl.comptime_struct +// CHECK-NEXT: %11 = "csl_wrapper.import"(%2, %3, %10) <{module = "", fields = ["width", "height", "LAUNCH"]}> : (i16, i16, !csl.color) -> !csl.imported_module +// CHECK-NEXT: %12 = "csl_wrapper.import"(%5, %2, %3) <{module = "routes.csl", fields = ["pattern", "peWidth", "peHeight"]}> : (i16, i16, i16) -> !csl.imported_module +// CHECK-NEXT: %13 = "csl.member_call"(%12, %0, %1, %2, %3, %5) <{field = "computeAllRoutes"}> : (!csl.imported_module, i16, i16, i16, i16, i16) -> !csl.comptime_struct +// CHECK-NEXT: %14 = "csl.member_call"(%11, %0) <{field = "get_params"}> : (!csl.imported_module, i16) -> !csl.comptime_struct // CHECK-NEXT: %15 = arith.constant 1 : i16 // CHECK-NEXT: %16 = arith.subi %5, %15 : i16 // CHECK-NEXT: %17 = arith.subi %2, %0 : i16 @@ -52,11 +52,11 @@ func.func @gauss_seidel(%a : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>, // CHECK-NEXT: %23 = arith.ori %19, %20 : i1 // CHECK-NEXT: %24 = arith.ori %23, %21 : i1 // CHECK-NEXT: %25 = arith.ori %24, %22 : i1 -// CHECK-NEXT: "csl_wrapper.yield"(%14, %13, %25) <{"fields" = ["memcpy_params", "stencil_comms_params", "isBorderRegionPE"]}> : (!csl.comptime_struct, !csl.comptime_struct, i1) -> () +// CHECK-NEXT: "csl_wrapper.yield"(%14, %13, %25) <{fields = ["memcpy_params", "stencil_comms_params", "isBorderRegionPE"]}> : (!csl.comptime_struct, !csl.comptime_struct, i1) -> () // CHECK-NEXT: }, { // CHECK-NEXT: ^1(%26 : i16, %27 : i16, %28 : i16, %29 : i16, %30 : i16, %31 : i16, %32 : i16, %memcpy_params : !csl.comptime_struct, %stencil_comms_params : !csl.comptime_struct, %isBorderRegionPE : i1): -// CHECK-NEXT: %33 = "csl_wrapper.import"(%memcpy_params) <{"module" = "", "fields" = [""]}> : (!csl.comptime_struct) -> !csl.imported_module -// CHECK-NEXT: %34 = "csl_wrapper.import"(%29, %31, %stencil_comms_params) <{"module" = "stencil_comms.csl", "fields" = ["pattern", "chunkSize", ""]}> : (i16, i16, !csl.comptime_struct) -> !csl.imported_module +// CHECK-NEXT: %33 = "csl_wrapper.import"(%memcpy_params) <{module = "", fields = [""]}> : (!csl.comptime_struct) -> !csl.imported_module +// CHECK-NEXT: %34 = "csl_wrapper.import"(%29, %31, %stencil_comms_params) <{module = "stencil_comms.csl", fields = ["pattern", "chunkSize", ""]}> : (i16, i16, !csl.comptime_struct) -> !csl.imported_module // CHECK-NEXT: %35 = memref.alloc() : memref<512xf32> // CHECK-NEXT: %36 = memref.alloc() : memref<512xf32> // CHECK-NEXT: %c = memref.alloc() : memref<255xf32> @@ -65,14 +65,14 @@ func.func @gauss_seidel(%a : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>, // CHECK-NEXT: %37 = "csl.addressof"(%35) : (memref<512xf32>) -> !csl.ptr, #csl> // CHECK-NEXT: %38 = "csl.addressof"(%36) : (memref<512xf32>) -> !csl.ptr, #csl> // CHECK-NEXT: %39 = "csl.addressof"(%c) : (memref<255xf32>) -> !csl.ptr, #csl> -// CHECK-NEXT: "csl.export"(%37) <{"var_name" = "a", "type" = !csl.ptr, #csl>}> : (!csl.ptr, #csl>) -> () -// CHECK-NEXT: "csl.export"(%38) <{"var_name" = "b", "type" = !csl.ptr, #csl>}> : (!csl.ptr, #csl>) -> () -// CHECK-NEXT: "csl.export"(%39) <{"var_name" = "c", "type" = !csl.ptr, #csl>}> : (!csl.ptr, #csl>) -> () -// CHECK-NEXT: "csl.export"() <{"var_name" = @gauss_seidel, "type" = () -> ()}> : () -> () +// CHECK-NEXT: "csl.export"(%37) <{var_name = "a", type = !csl.ptr, #csl>}> : (!csl.ptr, #csl>) -> () +// CHECK-NEXT: "csl.export"(%38) <{var_name = "b", type = !csl.ptr, #csl>}> : (!csl.ptr, #csl>) -> () +// CHECK-NEXT: "csl.export"(%39) <{var_name = "c", type = !csl.ptr, #csl>}> : (!csl.ptr, #csl>) -> () +// CHECK-NEXT: "csl.export"() <{var_name = @gauss_seidel, type = () -> ()}> : () -> () // CHECK-NEXT: csl.func @gauss_seidel() { // CHECK-NEXT: %40 = stencil.load %a : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>> -> !stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>> // CHECK-NEXT: %41 = tensor.empty() : tensor<510xf32> -// CHECK-NEXT: %42 = csl_stencil.apply(%40 : !stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>>, %41 : tensor<510xf32>, %c : memref<255xf32>) -> (!stencil.temp<[0,1]x[0,1]xtensor<510xf32>>) <{"swaps" = [#csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange], "topo" = #dmp.topo<1022x510>, "num_chunks" = 2 : i64, "operandSegmentSizes" = array}> ({ +// CHECK-NEXT: %42 = csl_stencil.apply(%40 : !stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>>, %41 : tensor<510xf32>, %c : memref<255xf32>) -> (!stencil.temp<[0,1]x[0,1]xtensor<510xf32>>) <{swaps = [#csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange, #csl_stencil.exchange], topo = #dmp.topo<1022x510>, num_chunks = 2 : i64, operandSegmentSizes = array}> ({ // CHECK-NEXT: ^2(%43 : tensor<4x255xf32>, %44 : index, %45 : tensor<510xf32>, %46 : memref<255xf32>): // CHECK-NEXT: %47 = csl_stencil.access %43[1, 0] : tensor<4x255xf32> // CHECK-NEXT: %48 = csl_stencil.access %43[-1, 0] : tensor<4x255xf32> @@ -82,15 +82,15 @@ func.func @gauss_seidel(%a : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>, // CHECK-NEXT: %52 = arith.addf %50, %49 : tensor<255xf32> // CHECK-NEXT: %53 = arith.addf %52, %48 : tensor<255xf32> // CHECK-NEXT: %54 = arith.addf %53, %47 : tensor<255xf32> -// CHECK-NEXT: %55 = "tensor.insert_slice"(%54, %45, %44) <{"static_offsets" = array, "static_sizes" = array, "static_strides" = array, "operandSegmentSizes" = array}> : (tensor<255xf32>, tensor<510xf32>, index) -> tensor<510xf32> +// CHECK-NEXT: %55 = "tensor.insert_slice"(%54, %45, %44) <{static_offsets = array, static_sizes = array, static_strides = array, operandSegmentSizes = array}> : (tensor<255xf32>, tensor<510xf32>, index) -> tensor<510xf32> // CHECK-NEXT: csl_stencil.yield %55 : tensor<510xf32> // CHECK-NEXT: }, { // CHECK-NEXT: ^3(%56 : !stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>>, %57 : tensor<510xf32>): // CHECK-NEXT: %58 = csl_stencil.access %56[0, 0] : !stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>> // CHECK-NEXT: %59 = csl_stencil.access %56[0, 0] : !stencil.temp<[-1,2]x[-1,2]xtensor<512xf32>> // CHECK-NEXT: %60 = arith.constant 1.666600e-01 : f32 -// CHECK-NEXT: %61 = "tensor.extract_slice"(%58) <{"static_offsets" = array, "static_sizes" = array, "static_strides" = array, "operandSegmentSizes" = array}> : (tensor<512xf32>) -> tensor<510xf32> -// CHECK-NEXT: %62 = "tensor.extract_slice"(%59) <{"static_offsets" = array, "static_sizes" = array, "static_strides" = array, "operandSegmentSizes" = array}> : (tensor<512xf32>) -> tensor<510xf32> +// CHECK-NEXT: %61 = "tensor.extract_slice"(%58) <{static_offsets = array, static_sizes = array, static_strides = array, operandSegmentSizes = array}> : (tensor<512xf32>) -> tensor<510xf32> +// CHECK-NEXT: %62 = "tensor.extract_slice"(%59) <{static_offsets = array, static_sizes = array, static_strides = array, operandSegmentSizes = array}> : (tensor<512xf32>) -> tensor<510xf32> // CHECK-NEXT: %63 = arith.addf %57, %62 : tensor<510xf32> // CHECK-NEXT: %64 = arith.addf %63, %61 : tensor<510xf32> // CHECK-NEXT: %65 = tensor.empty() : tensor<510xf32> @@ -99,10 +99,10 @@ func.func @gauss_seidel(%a : !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>>, // CHECK-NEXT: csl_stencil.yield %67 : tensor<510xf32> // CHECK-NEXT: }) // CHECK-NEXT: stencil.store %42 to %b(<[0, 0], [1, 1]>) : !stencil.temp<[0,1]x[0,1]xtensor<510xf32>> to !stencil.field<[-1,1023]x[-1,511]xtensor<512xf32>> -// CHECK-NEXT: "csl.member_call"(%33) <{"field" = "unblock_cmd_stream"}> : (!csl.imported_module) -> () +// CHECK-NEXT: "csl.member_call"(%33) <{field = "unblock_cmd_stream"}> : (!csl.imported_module) -> () // CHECK-NEXT: csl.return // CHECK-NEXT: } -// CHECK-NEXT: "csl_wrapper.yield"() <{"fields" = []}> : () -> () +// CHECK-NEXT: "csl_wrapper.yield"() <{fields = []}> : () -> () // CHECK-NEXT: }) : () -> () @@ -129,14 +129,14 @@ func.func private @timer_start() -> f64 func.func private @timer_end(f64) -> f64 -// CHECK: "csl_wrapper.module"() <{"width" = 1024 : i16, "height" = 512 : i16, "params" = [#csl_wrapper.param<"z_dim" default=512 : i16>, #csl_wrapper.param<"pattern" default=2 : i16>, #csl_wrapper.param<"num_chunks" default=2 : i16>, #csl_wrapper.param<"chunk_size" default=255 : i16>, #csl_wrapper.param<"padded_z_dim" default=510 : i16>], "program_name" = "bufferized"}> ({ +// CHECK: "csl_wrapper.module"() <{width = 1024 : i16, height = 512 : i16, params = [#csl_wrapper.param<"z_dim" default=512 : i16>, #csl_wrapper.param<"pattern" default=2 : i16>, #csl_wrapper.param<"num_chunks" default=2 : i16>, #csl_wrapper.param<"chunk_size" default=255 : i16>, #csl_wrapper.param<"padded_z_dim" default=510 : i16>], program_name = "bufferized"}> ({ // CHECK-NEXT: ^2(%43 : i16, %44 : i16, %45 : i16, %46 : i16, %47 : i16, %48 : i16, %49 : i16, %50 : i16, %51 : i16): // CHECK-NEXT: %52 = arith.constant 0 : i16 // CHECK-NEXT: %53 = "csl.get_color"(%52) : (i16) -> !csl.color -// CHECK-NEXT: %54 = "csl_wrapper.import"(%45, %46, %53) <{"module" = "", "fields" = ["width", "height", "LAUNCH"]}> : (i16, i16, !csl.color) -> !csl.imported_module -// CHECK-NEXT: %55 = "csl_wrapper.import"(%48, %45, %46) <{"module" = "routes.csl", "fields" = ["pattern", "peWidth", "peHeight"]}> : (i16, i16, i16) -> !csl.imported_module -// CHECK-NEXT: %56 = "csl.member_call"(%55, %43, %44, %45, %46, %48) <{"field" = "computeAllRoutes"}> : (!csl.imported_module, i16, i16, i16, i16, i16) -> !csl.comptime_struct -// CHECK-NEXT: %57 = "csl.member_call"(%54, %43) <{"field" = "get_params"}> : (!csl.imported_module, i16) -> !csl.comptime_struct +// CHECK-NEXT: %54 = "csl_wrapper.import"(%45, %46, %53) <{module = "", fields = ["width", "height", "LAUNCH"]}> : (i16, i16, !csl.color) -> !csl.imported_module +// CHECK-NEXT: %55 = "csl_wrapper.import"(%48, %45, %46) <{module = "routes.csl", fields = ["pattern", "peWidth", "peHeight"]}> : (i16, i16, i16) -> !csl.imported_module +// CHECK-NEXT: %56 = "csl.member_call"(%55, %43, %44, %45, %46, %48) <{field = "computeAllRoutes"}> : (!csl.imported_module, i16, i16, i16, i16, i16) -> !csl.comptime_struct +// CHECK-NEXT: %57 = "csl.member_call"(%54, %43) <{field = "get_params"}> : (!csl.imported_module, i16) -> !csl.comptime_struct // CHECK-NEXT: %58 = arith.constant 1 : i16 // CHECK-NEXT: %59 = arith.subi %48, %58 : i16 // CHECK-NEXT: %60 = arith.subi %45, %43 : i16 @@ -148,29 +148,29 @@ func.func private @timer_end(f64) -> f64 // CHECK-NEXT: %66 = arith.ori %62, %63 : i1 // CHECK-NEXT: %67 = arith.ori %66, %64 : i1 // CHECK-NEXT: %68 = arith.ori %67, %65 : i1 -// CHECK-NEXT: "csl_wrapper.yield"(%57, %56, %68) <{"fields" = ["memcpy_params", "stencil_comms_params", "isBorderRegionPE"]}> : (!csl.comptime_struct, !csl.comptime_struct, i1) -> () +// CHECK-NEXT: "csl_wrapper.yield"(%57, %56, %68) <{fields = ["memcpy_params", "stencil_comms_params", "isBorderRegionPE"]}> : (!csl.comptime_struct, !csl.comptime_struct, i1) -> () // CHECK-NEXT: }, { // CHECK-NEXT: ^3(%69 : i16, %70 : i16, %71 : i16, %72 : i16, %73 : i16, %74 : i16, %75 : i16, %memcpy_params_1 : !csl.comptime_struct, %stencil_comms_params_1 : !csl.comptime_struct, %isBorderRegionPE_1 : i1): -// CHECK-NEXT: %76 = "csl_wrapper.import"(%memcpy_params_1) <{"module" = "", "fields" = [""]}> : (!csl.comptime_struct) -> !csl.imported_module -// CHECK-NEXT: %77 = "csl_wrapper.import"(%72, %74, %stencil_comms_params_1) <{"module" = "stencil_comms.csl", "fields" = ["pattern", "chunkSize", ""]}> : (i16, i16, !csl.comptime_struct) -> !csl.imported_module +// CHECK-NEXT: %76 = "csl_wrapper.import"(%memcpy_params_1) <{module = "", fields = [""]}> : (!csl.comptime_struct) -> !csl.imported_module +// CHECK-NEXT: %77 = "csl_wrapper.import"(%72, %74, %stencil_comms_params_1) <{module = "stencil_comms.csl", fields = ["pattern", "chunkSize", ""]}> : (i16, i16, !csl.comptime_struct) -> !csl.imported_module // CHECK-NEXT: %in = memref.alloc() : memref<512xf32> // CHECK-NEXT: %arg1 = memref.alloc() : memref<512xf32> // CHECK-NEXT: %timers = memref.alloc() : memref<6xui16> // CHECK-NEXT: %78 = "csl.addressof"(%in) : (memref<512xf32>) -> !csl.ptr, #csl> // CHECK-NEXT: %79 = "csl.addressof"(%arg1) : (memref<512xf32>) -> !csl.ptr, #csl> // CHECK-NEXT: %80 = "csl.addressof"(%timers) : (memref<6xui16>) -> !csl.ptr, #csl> -// CHECK-NEXT: "csl.export"(%78) <{"var_name" = "in", "type" = !csl.ptr, #csl>}> : (!csl.ptr, #csl>) -> () -// CHECK-NEXT: "csl.export"(%79) <{"var_name" = "arg1", "type" = !csl.ptr, #csl>}> : (!csl.ptr, #csl>) -> () -// CHECK-NEXT: "csl.export"(%80) <{"var_name" = "timers", "type" = !csl.ptr, #csl>}> : (!csl.ptr, #csl>) -> () -// CHECK-NEXT: %81 = "csl_wrapper.import"() <{"module" = "