Merge branch 'main' into main

xdslproject · Mar 13, 2023 · 19aec4b · 19aec4b
2 parents bc05f99 + 6bc3779
commit 19aec4b
Show file tree

Hide file tree

Showing 43 changed files with 1,025 additions and 444 deletions.
diff --git a/.github/workflows/ci-mlir.yml b/.github/workflows/ci-mlir.yml
@@ -54,7 +54,6 @@ jobs:
         path: llvm-project/build
         key: binaries-${{ runner.os }}-${{ env.MLIR-Version }}
         restore-keys: binaries-${{ runner.os }}-${{ env.MLIR-Version }}
-        append-timestamp: false
 
     - name: Checkout MLIR
       if: steps.cache-binary.outputs.cache-hit != 'true'

diff --git a/.github/workflows/ci-pyright.yml b/.github/workflows/ci-pyright.yml
@@ -1,7 +1,7 @@
 # This workflow will install Python dependencies, run tests and lint with a single version of Python
 # For more information see: https://help.github.com/actions/language-and-framework-guides/using-python-with-github-actions
 
-name: CI - Pyright
+name: CI - Pyright - Reviewdog
 
 on:
   # Trigger the workflow on push or pull request,

diff --git a/README.md b/README.md
@@ -88,3 +88,7 @@ format the code in a uniform manner.
 To automate the formatting within vim, one can use
 https://github.com/vim-autoformat/vim-autoformat and trigger a `:Autoformat` on
 save.
+
+### Discussion
+
+You can also join the discussion at our [Zulip chat room](https://xdsl.zulipchat.com), kindly supported by community hosting from [Zulip](https://zulip.com/).
diff --git a/pyright-ci.json b/pyright-ci.json
@@ -4,7 +4,6 @@
   "reportUnnecessaryIsInstance": false,
   "typeCheckingMode": "strict",
   "exclude": [
-    "xdsl/parser.py",
     "xdsl/irdl_mlir_printer.py",
     "xdsl/irdl.py",
     "xdsl/ir.py"

diff --git a/tests/dialects/test_mpi.py b/tests/dialects/test_mpi.py
@@ -4,15 +4,32 @@
 
 
 def test_mpi_baseop():
+    """
+    This test is used to track changes in `.get` and other accessors
+    """
     alloc0 = memref.Alloc.get(f64, 32, [100, 14, 14])
     dest = Constant.from_int_and_width(1, i32)
-    send = mpi.ISend.get(alloc0, dest, 1)
-    recv = mpi.IRecv.get(dest, alloc0.memref, 1)
+    unwrap = mpi.UnwrapMemrefOp.get(alloc0)
+    tag = Constant.from_int_and_width(1, i32)
+    send = mpi.ISend.get(unwrap.ptr, unwrap.len, unwrap.typ, dest, tag)
+    wait = mpi.Wait.get(send.request, ignore_status=False)
+    recv = mpi.IRecv.get(unwrap.ptr, unwrap.len, unwrap.typ, dest, tag)
     test_res = mpi.Test.get(recv.request)
-    code2 = mpi.Wait.get(recv.request)
+    source = mpi.GetStatusField.get(wait.status,
+                                    mpi.StatusTypeField.MPI_SOURCE)
 
-    assert send.operands[0] is alloc0.results[0]
-    assert send.operands[1] is dest.results[0]
-    assert recv.operands[0] is send.operands[1]
-    assert code2.operands[0] is recv.results[0]
-    assert test_res.operands[0] is recv.results[0]
+    assert unwrap.ref == alloc0.memref
+    assert send.buffer == unwrap.ptr
+    assert send.count == unwrap.len
+    assert send.datatype == unwrap.typ
+    assert send.dest == dest.result
+    assert send.tag == tag.result
+    assert wait.request == send.request
+    assert recv.buffer == unwrap.ptr
+    assert recv.count == unwrap.len
+    assert recv.datatype == unwrap.typ
+    assert recv.source == dest.result
+    assert recv.tag == tag.result
+    assert test_res.request == recv.request
+    assert source.status == wait.status
+    assert source.field.data == mpi.StatusTypeField.MPI_SOURCE.value
diff --git a/tests/dialects/test_mpi_lowering.py b/tests/dialects/test_mpi_lowering.py
@@ -153,12 +153,11 @@ def test_lower_mpi_send():
 
 
 def test_lower_mpi_isend():
-    buff, dest = CreateTestValsOp.get(
-        mpi.MemRefType.from_element_type_and_shape(builtin.f64, [32, 32, 32]),
-        i32).results
+    ptr, count, dtype, dest, tag = CreateTestValsOp.get(
+        llvm.LLVMPointerType.opaque(), i32, mpi.DataType(), i32, i32).results
 
     ops, result = lower_mpi.LowerMpiISend(info).lower(
-        mpi.ISend.get(buff, dest, 1))
+        mpi.ISend.get(ptr, count, dtype, dest, tag))
     """
     Check for function with signature like:
     int MPI_Isend(const void *buf, int count, MPI_Datatype datatype, int dest,
@@ -219,16 +218,15 @@ def test_lower_mpi_recv_with_status():
 
 
 def test_lower_mpi_irecv():
-    buff, source = CreateTestValsOp.get(
-        mpi.MemRefType.from_element_type_and_shape(builtin.f64, [32, 32, 32]),
-        i32).results
+    ptr, count, dtype, source, tag = CreateTestValsOp.get(
+        llvm.LLVMPointerType.opaque(), i32, mpi.DataType(), i32, i32).results
     """
     int MPI_Irecv(void *buf, int count, MPI_Datatype datatype,
             int source, int tag, MPI_Comm comm, MPI_Request *request)
     """
 
     ops, result = lower_mpi.LowerMpiIRecv(info).lower(
-        mpi.IRecv.get(source, buff, tag=3))
+        mpi.IRecv.get(ptr, count, dtype, source, tag))
 
     assert len(result) == 1
 

diff --git a/tests/dialects/test_scf.py b/tests/dialects/test_scf.py
@@ -1,5 +1,5 @@
-from xdsl.dialects.arith import Constant, IndexType
-from xdsl.dialects.builtin import Region
+from xdsl.dialects.arith import Constant
+from xdsl.dialects.builtin import Region, IndexType
 from xdsl.dialects.cf import Block
 from xdsl.dialects.scf import For
 

diff --git a/tests/filecheck/dialects/gpu/all_reduce_types.mlir b/tests/filecheck/dialects/gpu/all_reduce_types.mlir
@@ -3,7 +3,8 @@
 "builtin.module"()({
     "gpu.module"()({
         %init = "arith.constant"() {"value" = 42 : index} : () -> index
-        %sum = "gpu.all_reduce"(%init) ({}) {"op" = #gpu<all_reduce_op add>} : (index) -> f32
+        %sum = "gpu.all_reduce"(%init) ({
+        }) {"op" = #gpu<all_reduce_op add>} : (index) -> f32
         "gpu.module_end"() : () -> ()
     }) {"sym_name" = "gpu"} : () -> ()
 }) {} : () -> ()

diff --git a/tests/filecheck/dialects/gpu/example.mlir b/tests/filecheck/dialects/gpu/example.mlir
@@ -39,7 +39,8 @@
             %subgroupid = "gpu.subgroup_id"() : () -> index
             %subgroupsize = "gpu.subgroup_size"() : () -> index
 
-            %globalprodx = "gpu.all_reduce"(%globalidx) ({}) {"op" = #gpu<all_reduce_op mul>} : (index) -> index
+            %globalprodx = "gpu.all_reduce"(%globalidx) ({
+            }) {"op" = #gpu<all_reduce_op mul>} : (index) -> index
 
             %globalsumy = "gpu.all_reduce"(%globalidy) ({
             ^bb(%lhs : index, %rhs : index):
@@ -52,7 +53,8 @@
                 %tx : index, %ty : index, %tz : index,
                 %num_bx : index, %num_by : index, %num_bz : index,
                 %num_tx : index, %num_ty : index, %num_tz : index):
-                %sum = "gpu.all_reduce"(%tx) ({}) {"op" = #gpu<all_reduce_op add>} : (index) -> index
+                %sum = "gpu.all_reduce"(%tx) ({
+                }) {"op" = #gpu<all_reduce_op add>} : (index) -> index
                 %final = "arith.muli"(%sum, %one) : (index, index) -> index    
                 "gpu.terminator"() : () -> ()
             }) {"operand_segment_sizes" = array<i32: 0, 1, 1, 1, 1, 1, 1, 0>} : (index, index, index, index, index, index) -> () 
@@ -102,7 +104,8 @@
 // CHECK-NEXT:             %{{.*}} = "gpu.subgroup_id"() : () -> index
 // CHECK-NEXT:             %{{.*}} = "gpu.subgroup_size"() : () -> index
 
-// CHECK-NEXT:             %{{.*}} = "gpu.all_reduce"(%{{.*}}) ({}) {"op" = #gpu<all_reduce_op mul>} : (index) -> index
+// CHECK-NEXT:             %{{.*}} = "gpu.all_reduce"(%{{.*}}) ({
+// CHECK-NEXT:             }) {"op" = #gpu<all_reduce_op mul>} : (index) -> index
 
 // CHECK-NEXT:             %{{.*}} = "gpu.all_reduce"(%{{.*}}) ({
 // CHECK-NEXT:             ^{{.*}}(%{{.*}} : index, %{{.*}} : index):
@@ -115,7 +118,8 @@
 // CHECK-SAME:                 %{{.*}} : index, %{{.*}} : index, %{{.*}} : index,
 // CHECK-SAME:                 %{{.*}} : index, %{{.*}} : index, %{{.*}} : index,
 // CHECK-SAME:                 %{{.*}} : index, %{{.*}} : index, %{{.*}} : index):
-// CHECK-NEXT:                 %{{.*}} = "gpu.all_reduce"(%{{.*}}) ({}) {"op" = #gpu<all_reduce_op add>} : (index) -> index
+// CHECK-NEXT:                 %{{.*}} = "gpu.all_reduce"(%{{.*}}) ({
+// CHECK-NEXT:                 }) {"op" = #gpu<all_reduce_op add>} : (index) -> index
 // CHECK-NEXT:                 %{{.*}} = "arith.muli"(%{{.*}}, %{{.*}}) : (index, index) -> index    
 // CHECK-NEXT:                 "gpu.terminator"() : () -> ()
 // CHECK-NEXT:             }) {"operand_segment_sizes" = array<i32: 0, 1, 1, 1, 1, 1, 1, 0>} : (index, index, index, index, index, index) -> () 

diff --git a/tests/filecheck/dialects/gpu/launch_args.mlir b/tests/filecheck/dialects/gpu/launch_args.mlir
@@ -8,7 +8,8 @@
         "gpu.launch"(%one, %one, %n, %one, %one, %one) ({
         ^bb0(%bx : index, %by : index, %bz : index,
             %tx : index, %ty : index, %tz : index):
-            %sum = "gpu.all_reduce"(%tx) ({}) {"op" = #gpu<all_reduce_op add>} : (index) -> index
+            %sum = "gpu.all_reduce"(%tx) ({
+            }) {"op" = #gpu<all_reduce_op add>} : (index) -> index
             %final = "arith.muli"(%sum, %one) : (index, index) -> index    
             "gpu.terminator"() : () -> ()
         }) {"operand_segment_sizes" = array<i32: 0, 1, 1, 1, 1, 1, 1, 0>} : (index, index, index, index, index, index) -> () 

diff --git a/tests/filecheck/dialects/gpu/terminator_launch.mlir b/tests/filecheck/dialects/gpu/terminator_launch.mlir
@@ -10,7 +10,8 @@
             %tx : index, %ty : index, %tz : index,
             %num_bx : index, %num_by : index, %num_bz : index,
             %num_tx : index, %num_ty : index, %num_tz : index):
-            %sum = "gpu.all_reduce"(%tx) ({}) {"op" = #gpu<all_reduce_op add>} : (index) -> index
+            %sum = "gpu.all_reduce"(%tx) ({
+            }) {"op" = #gpu<all_reduce_op add>} : (index) -> index
             %final = "arith.muli"(%sum, %one) : (index, index) -> index    
         }) {"operand_segment_sizes" = array<i32: 0, 1, 1, 1, 1, 1, 1, 0>} : (index, index, index, index, index, index) -> () 
 

diff --git a/tests/filecheck/dialects/gpu/terminator_terminate.mlir b/tests/filecheck/dialects/gpu/terminator_terminate.mlir
@@ -10,7 +10,8 @@
             %tx : index, %ty : index, %tz : index,
             %num_bx : index, %num_by : index, %num_bz : index,
             %num_tx : index, %num_ty : index, %num_tz : index):
-            %sum = "gpu.all_reduce"(%tx) ({}) {"op" = #gpu<all_reduce_op add>} : (index) -> index
+            %sum = "gpu.all_reduce"(%tx) ({
+            }) {"op" = #gpu<all_reduce_op add>} : (index) -> index
             "gpu.terminator"() : () -> ()
             %final = "arith.muli"(%sum, %one) : (index, index) -> index    
         }) {"operand_segment_sizes" = array<i32: 0, 1, 1, 1, 1, 1, 1, 0>} : (index, index, index, index, index, index) -> () 

diff --git a/tests/filecheck/dialects/stencil/copy.mlir b/tests/filecheck/dialects/stencil/copy.mlir
@@ -0,0 +1,35 @@
+// RUN: xdsl-opt %s -t mlir | filecheck %s
+
+"builtin.module"() ({
+  "func.func"() ({
+  ^0(%0 : !stencil.field<[-1 : i32, -1 : i32, -1 : i32], f64>, %1 : !stencil.field<[-1 : i32, -1 : i32, -1 : i32], f64>):
+    %2 = "stencil.cast"(%0) {"lb" = #stencil.index<[-4 : i32, -4 : i32, -4 : i32]>, "ub" = #stencil.index<[68 : i32, 68 : i32, 68 : i32]>} : (!stencil.field<[-1 : i32, -1 : i32, -1 : i32], f64>) -> !stencil.field<[72 : i32, 72 : i32, 72 : i32], f64>
+    %3 = "stencil.cast"(%1) {"lb" = #stencil.index<[-4 : i32, -4 : i32, -4 : i32]>, "ub" = #stencil.index<[68 : i32, 68 : i32, 68 : i32]>} : (!stencil.field<[-1 : i32, -1 : i32, -1 : i32], f64>) -> !stencil.field<[72 : i32, 72 : i32, 72 : i32], f64>
+    %4 = "stencil.load"(%2) {"lb" = #stencil.index<[0 : i32, 0 : i32, 0 : i32]>, "ub" = #stencil.index<[64 : i32, 64 : i32, 64 : i32]>} : (!stencil.field<[72 : i32, 72 : i32, 72 : i32], f64>) -> !stencil.temp<[64 : i32, 64 : i32, 64 : i32], f64>
+    %5 = "stencil.apply"(%4) ({
+    ^1(%6 : !stencil.temp<[64 : i32, 64 : i32, 64 : i32], f64>):
+      %7 = "stencil.access"(%6) {"offset" = #stencil.index<[0 : i32, 0 : i32, 0 : i32]>} : (!stencil.temp<[64 : i32, 64 : i32, 64 : i32], f64>) -> f64
+      %8 = "stencil.store_result"(%7) : (f64) -> !stencil.result<f64>
+      "stencil.return"(%8) : (!stencil.result<f64>) -> ()
+    }) {"lb" = #stencil.index<[0 : i32, 0 : i32, 0 : i32]>, "ub" = #stencil.index<[64 : i32, 64 : i32, 64 : i32]>} : (!stencil.temp<[64 : i32, 64 : i32, 64 : i32], f64>) -> !stencil.temp<[64 : i32, 64 : i32, 64 : i32], f64>
+    "stencil.store"(%5, %3) {"lb" = #stencil.index<[0 : i32, 0 : i32, 0 : i32]>, "ub" = #stencil.index<[64 : i32, 64 : i32, 64 : i32]>} : (!stencil.temp<[64 : i32, 64 : i32, 64 : i32], f64>, !stencil.field<[72 : i32, 72 : i32, 72 : i32], f64>) -> ()
+    "func.return"() : () -> ()
+  }) {"sym_name" = "stencil_copy", "function_type" = (!stencil.field<[-1 : i32, -1 : i32, -1 : i32], f64>, !stencil.field<[-1 : i32, -1 : i32, -1 : i32], f64>) -> (), "sym_visibility" = "private"} : () -> ()
+}) : () -> ()
+
+// CHECK-NEXT: "builtin.module"() ({
+// CHECK-NEXT:   "func.func"() ({
+// CHECK-NEXT:   ^0(%0 : !stencil.field<[-1 : i32, -1 : i32, -1 : i32], f64>, %1 : !stencil.field<[-1 : i32, -1 : i32, -1 : i32], f64>):
+// CHECK-NEXT:     %2 = "stencil.cast"(%0) {"lb" = #stencil.index<[-4 : i32, -4 : i32, -4 : i32]>, "ub" = #stencil.index<[68 : i32, 68 : i32, 68 : i32]>} : (!stencil.field<[-1 : i32, -1 : i32, -1 : i32], f64>) -> !stencil.field<[72 : i32, 72 : i32, 72 : i32], f64>
+// CHECK-NEXT:     %3 = "stencil.cast"(%1) {"lb" = #stencil.index<[-4 : i32, -4 : i32, -4 : i32]>, "ub" = #stencil.index<[68 : i32, 68 : i32, 68 : i32]>} : (!stencil.field<[-1 : i32, -1 : i32, -1 : i32], f64>) -> !stencil.field<[72 : i32, 72 : i32, 72 : i32], f64>
+// CHECK-NEXT:     %4 = "stencil.load"(%2) {"lb" = #stencil.index<[0 : i32, 0 : i32, 0 : i32]>, "ub" = #stencil.index<[64 : i32, 64 : i32, 64 : i32]>} : (!stencil.field<[72 : i32, 72 : i32, 72 : i32], f64>) -> !stencil.temp<[64 : i32, 64 : i32, 64 : i32], f64>
+// CHECK-NEXT:     %5 = "stencil.apply"(%4) ({
+// CHECK-NEXT:     ^1(%6 : !stencil.temp<[64 : i32, 64 : i32, 64 : i32], f64>):
+// CHECK-NEXT:       %7 = "stencil.access"(%6) {"offset" = #stencil.index<[0 : i32, 0 : i32, 0 : i32]>} : (!stencil.temp<[64 : i32, 64 : i32, 64 : i32], f64>) -> f64
+// CHECK-NEXT:       %8 = "stencil.store_result"(%7) : (f64) -> !stencil.result<f64>
+// CHECK-NEXT:       "stencil.return"(%8) : (!stencil.result<f64>) -> ()
+// CHECK-NEXT:     }) {"lb" = #stencil.index<[0 : i32, 0 : i32, 0 : i32]>, "ub" = #stencil.index<[64 : i32, 64 : i32, 64 : i32]>} : (!stencil.temp<[64 : i32, 64 : i32, 64 : i32], f64>) -> !stencil.temp<[64 : i32, 64 : i32, 64 : i32], f64>
+// CHECK-NEXT:     "stencil.store"(%5, %3) {"lb" = #stencil.index<[0 : i32, 0 : i32, 0 : i32]>, "ub" = #stencil.index<[64 : i32, 64 : i32, 64 : i32]>} : (!stencil.temp<[64 : i32, 64 : i32, 64 : i32], f64>, !stencil.field<[72 : i32, 72 : i32, 72 : i32], f64>) -> ()
+// CHECK-NEXT:     "func.return"() : () -> ()
+// CHECK-NEXT:   }) {"sym_name" = "stencil_copy", "function_type" = (!stencil.field<[-1 : i32, -1 : i32, -1 : i32], f64>, !stencil.field<[-1 : i32, -1 : i32, -1 : i32], f64>) -> (), "sym_visibility" = "private"} : () -> ()
+// CHECK-NEXT: }) : () -> ()
diff --git a/tests/filecheck/dialects/stencil/hdiff.mlir b/tests/filecheck/dialects/stencil/hdiff.mlir
@@ -0,0 +1,57 @@
+// RUN: xdsl-opt %s -t mlir | filecheck %s
+
+"builtin.module"() ({
+  "func.func"() ({
+  ^0(%0 : !stencil.field<[-1 : i64, -1 : i64, -1 : i64], f64>, %1 : !stencil.field<[-1 : i64, -1 : i64, -1 : i64], f64>, %2 : !stencil.field<[-1 : i64, -1 : i64, -1 : i64], f64>):
+    %3 = "stencil.cast"(%0) {"lb" = #stencil.index<[-4 : i64, -4 : i64, -4 : i64]>, "ub" = #stencil.index<[68 : i64, 68 : i64, 68 : i64]>} : (!stencil.field<[-1 : i64, -1 : i64, -1 : i64], f64>) -> !stencil.field<[72 : i64, 72 : i64, 72 : i64], f64>
+    %4 = "stencil.cast"(%1) {"lb" = #stencil.index<[-4 : i64, -4 : i64, -4 : i64]>, "ub" = #stencil.index<[68 : i64, 68 : i64, 68 : i64]>} : (!stencil.field<[-1 : i64, -1 : i64, -1 : i64], f64>) -> !stencil.field<[72 : i64, 72 : i64, 72 : i64], f64>
+    %5 = "stencil.cast"(%2) {"lb" = #stencil.index<[-4 : i64, -4 : i64, -4 : i64]>, "ub" = #stencil.index<[68 : i64, 68 : i64, 68 : i64]>} : (!stencil.field<[-1 : i64, -1 : i64, -1 : i64], f64>) -> !stencil.field<[72 : i64, 72 : i64, 72 : i64], f64>
+    %6 = "stencil.load"(%3) : (!stencil.field<[72 : i64, 72 : i64, 72 : i64], f64>) -> !stencil.temp<[-1 : i64, -1 : i64, -1 : i64], f64>
+    %7 = "stencil.load"(%4) : (!stencil.field<[72 : i64, 72 : i64, 72 : i64], f64>) -> !stencil.temp<[-1 : i64, -1 : i64, -1 : i64], f64>
+    %8 = "stencil.apply"(%6) ({
+    ^1(%9 : !stencil.temp<[-1 : i64, -1 : i64, -1 : i64], f64>):
+      %10 = "stencil.access"(%9) {"offset" = #stencil.index<[-1 : i64, 0 : i64, 0 : i64]>} : (!stencil.temp<[-1 : i64, -1 : i64, -1 : i64], f64>) -> f64
+      %11 = "stencil.access"(%9) {"offset" = #stencil.index<[1 : i64, 0 : i64, 0 : i64]>} : (!stencil.temp<[-1 : i64, -1 : i64, -1 : i64], f64>) -> f64
+      %12 = "stencil.access"(%9) {"offset" = #stencil.index<[0 : i64, 1 : i64, 0 : i64]>} : (!stencil.temp<[-1 : i64, -1 : i64, -1 : i64], f64>) -> f64
+      %13 = "stencil.access"(%9) {"offset" = #stencil.index<[0 : i64, -1 : i64, 0 : i64]>} : (!stencil.temp<[-1 : i64, -1 : i64, -1 : i64], f64>) -> f64
+      %14 = "stencil.access"(%9) {"offset" = #stencil.index<[0 : i64, 0 : i64, 0 : i64]>} : (!stencil.temp<[-1 : i64, -1 : i64, -1 : i64], f64>) -> f64
+      %15 = "arith.addf"(%10, %11) : (f64, f64) -> f64
+      %16 = "arith.addf"(%12, %13) : (f64, f64) -> f64
+      %17 = "arith.addf"(%15, %16) : (f64, f64) -> f64
+      %cst = "arith.constant"() {"value" = -4.0 : f32} : () -> f64
+      %18 = "arith.mulf"(%14, %cst) : (f64, f64) -> f64
+      %19 = "arith.addf"(%18, %17) : (f64, f64) -> f64
+      %20 = "stencil.store_result"(%19) : (f64) -> !stencil.result<f64>
+      "stencil.return"(%20) : (!stencil.result<f64>) -> ()
+    }) : (!stencil.temp<[-1 : i64, -1 : i64, -1 : i64], f64>) -> !stencil.temp<[-1 : i64, -1 : i64, -1 : i64], f64>
+  }) {"function_type" = (!stencil.field<[-1 : i64, -1 : i64, -1 : i64], f64>, !stencil.field<[-1 : i64, -1 : i64, -1 : i64], f64>, !stencil.field<[-1 : i64, -1 : i64, -1 : i64], f64>) -> !stencil.field<[-1 : i64, -1 : i64, -1 : i64], f64>, "sym_name" = "stencil_hdiff"} : () -> ()
+}) : () -> ()
+
+
+// CHECK-NEXT: "builtin.module"() ({
+// CHECK-NEXT:   "func.func"() ({
+// CHECK-NEXT:   ^0(%0 : !stencil.field<[-1 : i64, -1 : i64, -1 : i64], f64>, %1 : !stencil.field<[-1 : i64, -1 : i64, -1 : i64], f64>, %2 : !stencil.field<[-1 : i64, -1 : i64, -1 : i64], f64>):
+// CHECK-NEXT:     %3 = "stencil.cast"(%0) {"lb" = #stencil.index<[-4 : i64, -4 : i64, -4 : i64]>, "ub" = #stencil.index<[68 : i64, 68 : i64, 68 : i64]>} : (!stencil.field<[-1 : i64, -1 : i64, -1 : i64], f64>) -> !stencil.field<[72 : i64, 72 : i64, 72 : i64], f64>
+// CHECK-NEXT:     %4 = "stencil.cast"(%1) {"lb" = #stencil.index<[-4 : i64, -4 : i64, -4 : i64]>, "ub" = #stencil.index<[68 : i64, 68 : i64, 68 : i64]>} : (!stencil.field<[-1 : i64, -1 : i64, -1 : i64], f64>) -> !stencil.field<[72 : i64, 72 : i64, 72 : i64], f64>
+// CHECK-NEXT:     %5 = "stencil.cast"(%2) {"lb" = #stencil.index<[-4 : i64, -4 : i64, -4 : i64]>, "ub" = #stencil.index<[68 : i64, 68 : i64, 68 : i64]>} : (!stencil.field<[-1 : i64, -1 : i64, -1 : i64], f64>) -> !stencil.field<[72 : i64, 72 : i64, 72 : i64], f64>
+// CHECK-NEXT:     %6 = "stencil.load"(%3) : (!stencil.field<[72 : i64, 72 : i64, 72 : i64], f64>) -> !stencil.temp<[-1 : i64, -1 : i64, -1 : i64], f64>
+// CHECK-NEXT:     %7 = "stencil.load"(%4) : (!stencil.field<[72 : i64, 72 : i64, 72 : i64], f64>) -> !stencil.temp<[-1 : i64, -1 : i64, -1 : i64], f64>
+// CHECK-NEXT:     %8 = "stencil.apply"(%6) ({
+// CHECK-NEXT:     ^1(%9 : !stencil.temp<[-1 : i64, -1 : i64, -1 : i64], f64>):
+// CHECK-NEXT:       %10 = "stencil.access"(%9) {"offset" = #stencil.index<[-1 : i64, 0 : i64, 0 : i64]>} : (!stencil.temp<[-1 : i64, -1 : i64, -1 : i64], f64>) -> f64
+// CHECK-NEXT:       %11 = "stencil.access"(%9) {"offset" = #stencil.index<[1 : i64, 0 : i64, 0 : i64]>} : (!stencil.temp<[-1 : i64, -1 : i64, -1 : i64], f64>) -> f64
+// CHECK-NEXT:       %12 = "stencil.access"(%9) {"offset" = #stencil.index<[0 : i64, 1 : i64, 0 : i64]>} : (!stencil.temp<[-1 : i64, -1 : i64, -1 : i64], f64>) -> f64
+// CHECK-NEXT:       %13 = "stencil.access"(%9) {"offset" = #stencil.index<[0 : i64, -1 : i64, 0 : i64]>} : (!stencil.temp<[-1 : i64, -1 : i64, -1 : i64], f64>) -> f64
+// CHECK-NEXT:       %14 = "stencil.access"(%9) {"offset" = #stencil.index<[0 : i64, 0 : i64, 0 : i64]>} : (!stencil.temp<[-1 : i64, -1 : i64, -1 : i64], f64>) -> f64
+// CHECK-NEXT:       %15 = "arith.addf"(%10, %11) : (f64, f64) -> f64
+// CHECK-NEXT:       %16 = "arith.addf"(%12, %13) : (f64, f64) -> f64
+// CHECK-NEXT:       %17 = "arith.addf"(%15, %16) : (f64, f64) -> f64
+// CHECK-NEXT:       %cst = "arith.constant"() {"value" = -4.0 : f32} : () -> f64
+// CHECK-NEXT:       %18 = "arith.mulf"(%14, %cst) : (f64, f64) -> f64
+// CHECK-NEXT:       %19 = "arith.addf"(%18, %17) : (f64, f64) -> f64
+// CHECK-NEXT:       %20 = "stencil.store_result"(%19) : (f64) -> !stencil.result<f64>
+// CHECK-NEXT:       "stencil.return"(%20) : (!stencil.result<f64>) -> ()
+// CHECK-NEXT:     }) : (!stencil.temp<[-1 : i64, -1 : i64, -1 : i64], f64>) -> !stencil.temp<[-1 : i64, -1 : i64, -1 : i64], f64>
+// CHECK-NEXT:   }) {"function_type" = (!stencil.field<[-1 : i64, -1 : i64, -1 : i64], f64>, !stencil.field<[-1 : i64, -1 : i64, -1 : i64], f64>, !stencil.field<[-1 : i64, -1 : i64, -1 : i64], f64>) -> !stencil.field<[-1 : i64, -1 : i64, -1 : i64], f64>, "sym_name" = "stencil_hdiff"} : () -> ()
+// CHECK-NEXT: }) : () -> ()
+