Repository: exo-lang/exo

Issue:
We should check Check_Aliasing from the frontend typecheck

Discussion:
yamaguchi1024: @rachitnigam It looks like some tests are failing because their algorithm has aliasing. Could you modify the tests and golden output so that they don't have aliasing?
rachitnigam: Hey @yamaguchi1024! I’ll get this done early next week. I got pulled into a deadline so a little busy right now
yamaguchi1024: @rachitnigam No worries & no rush, good luck for the deadline!
yamaguchi1024: Actually, it seems what we want is to allow aliasing in general, but check if the argument can be aliased in the backend and add `__restrict__` if not. 
gilbo: If you want to allow aliasing you need to account for it everywhere in the program analysis…On May 9, 2023, at 8:05 AM, Yuka Ikarashi ***@***.***> wrote:﻿
Actually, it seems what we want is to allow aliasing in general, but check if the argument can be aliased in the backend and add __restrict__ if not.

—Reply to this email directly, view it on GitHub, or unsubscribe.You are receiving this because you are subscribed to this thread.Message ID: ***@***.***>
codecov-commenter: ## [Codecov](https://app.codecov.io/gh/exo-lang/exo/pull/427?src=pr&el=h1&utm_medium=referral&utm_source=github&utm_content=comment&utm_campaign=pr+comments&utm_term=exo-lang) Report
> Merging [#427](https://app.codecov.io/gh/exo-lang/exo/pull/427?src=pr&el=desc&utm_medium=referral&utm_source=github&utm_content=comment&utm_campaign=pr+comments&utm_term=exo-lang) (218e2cc) into [master](https://app.codecov.io/gh/exo-lang/exo/commit/acff39a7b13d9170c2de2599787130f7bfddbf45?el=desc&utm_medium=referral&utm_source=github&utm_content=comment&utm_campaign=pr+comments&utm_term=exo-lang) (acff39a) will **decrease** coverage by `0.01%`.
> The diff coverage is `100.00%`.

> :exclamation: Current head 218e2cc differs from pull request most recent head f118619. Consider uploading reports for the commit f118619 to get more accurate results

:exclamation: Your organization is not using the GitHub App Integration. As a result you may experience degraded service beginning May 15th. Please [install the Github App Integration](https://github.com/apps/codecov) for your organization. [Read more](https://about.codecov.io/blog/codecov-is-updating-its-github-integration/?utm_medium=referral&utm_source=github&utm_content=comment&utm_campaign=pr+comments&utm_term=exo-lang).

```diff
@@            Coverage Diff             @@
##           master     #427      +/-   ##
==========================================
- Coverage   86.76%   86.76%   -0.01%     
==========================================
  Files          76       76              
  Lines       17600    17605       +5     
==========================================
+ Hits        15271    15275       +4     
- Misses       2329     2330       +1     
```


| [Impacted Files](https://app.codecov.io/gh/exo-lang/exo/pull/427?src=pr&el=tree&utm_medium=referral&utm_source=github&utm_content=comment&utm_campaign=pr+comments&utm_term=exo-lang) | Coverage Δ | |
|---|---|---|
| [tests/test\_window.py](https://app.codecov.io/gh/exo-lang/exo/pull/427?src=pr&el=tree&utm_medium=referral&utm_source=github&utm_content=comment&utm_campaign=pr+comments&utm_term=exo-lang#diff-dGVzdHMvdGVzdF93aW5kb3cucHk=) | `100.00% <ø> (ø)` | |
| [tests/test\_x86.py](https://app.codecov.io/gh/exo-lang/exo/pull/427?src=pr&el=tree&utm_medium=referral&utm_source=github&utm_content=comment&utm_campaign=pr+comments&utm_term=exo-lang#diff-dGVzdHMvdGVzdF94ODYucHk=) | `100.00% <ø> (ø)` | |
| [src/exo/API.py](https://app.codecov.io/gh/exo-lang/exo/pull/427?src=pr&el=tree&utm_medium=referral&utm_source=github&utm_content=comment&utm_campaign=pr+comments&utm_term=exo-lang#diff-c3JjL2V4by9BUEkucHk=) | `80.00% <100.00%> (+0.16%)` | :arrow_up: |
| [tests/test\_effects.py](https://app.codecov.io/gh/exo-lang/exo/pull/427?src=pr&el=tree&utm_medium=referral&utm_source=github&utm_content=comment&utm_campaign=pr+comments&utm_term=exo-lang#diff-dGVzdHMvdGVzdF9lZmZlY3RzLnB5) | `100.00% <100.00%> (ø)` | |
| [tests/test\_new\_eff.py](https://app.codecov.io/gh/exo-lang/exo/pull/427?src=pr&el=tree&utm_medium=referral&utm_source=github&utm_content=comment&utm_campaign=pr+comments&utm_term=exo-lang#diff-dGVzdHMvdGVzdF9uZXdfZWZmLnB5) | `100.00% <100.00%> (ø)` | |
| [tests/test\_schedules.py](https://app.codecov.io/gh/exo-lang/exo/pull/427?src=pr&el=tree&utm_medium=referral&utm_source=github&utm_content=comment&utm_campaign=pr+comments&utm_term=exo-lang#diff-dGVzdHMvdGVzdF9zY2hlZHVsZXMucHk=) | `99.30% <100.00%> (ø)` | |

... and [2 files with indirect coverage changes](https://app.codecov.io/gh/exo-lang/exo/pull/427/indirect-changes?src=pr&el=tree-more&utm_medium=referral&utm_source=github&utm_content=comment&utm_campaign=pr+comments&utm_term=exo-lang)



Patch:
diff --git a/src/exo/API.py b/src/exo/API.py
index 2977369..8ad02cb 100644
--- a/src/exo/API.py
+++ b/src/exo/API.py
@@ -20,6 +20,7 @@ from .memory import Memory
 from .parse_fragment import parse_fragment
 from .pattern_match import match_pattern, get_match_no
 from .prelude import *
+from .new_eff import Check_Aliasing
 
 # Moved to new file
 from .proc_eqv import decl_new_proc, derive_proc, assert_eqv_proc, check_eqv_proc
@@ -174,6 +175,7 @@ class Procedure(ProcedureBase):
             proc = TypeChecker(proc).get_loopir()
             proc = InferEffects(proc).result()
             CheckEffects(proc)
+            Check_Aliasing(proc)
 
         assert isinstance(proc, LoopIR.LoopIR.proc)
 
diff --git a/src/exo/LoopIR_compiler.py b/src/exo/LoopIR_compiler.py
index 9c9ba59..458ce01 100644
--- a/src/exo/LoopIR_compiler.py
+++ b/src/exo/LoopIR_compiler.py
@@ -382,11 +382,11 @@ def compile_to_strings(lib_name, proc_list):
             )
         else:
             is_public_decl = id(p) in orig_procs
-            
+
             p = PrecisionAnalysis().run(p)
             p = WindowAnalysis().apply_proc(p)
             p = MemoryAnalysis().run(p)
-            
+
             comp = Compiler(p, ctxt_name, is_public_decl=is_public_decl)
             d, b = comp.comp_top()
             struct_defns |= comp.struct_defns()
diff --git a/src/exo/platforms/neon.py b/src/exo/platforms/neon.py
index a13c9a4..0b449af 100644
--- a/src/exo/platforms/neon.py
+++ b/src/exo/platforms/neon.py
@@ -30,7 +30,11 @@ class Neon(Memory):
         if not shape:
             raise MemGenError(f"{srcinfo}: Neon vectors are not scalar values")
 
-        vec_types = {"float": (4, "float32x4_t"), "double": (2, "float64x2_t"), "_Float16" : (8, "float16x8_t")}
+        vec_types = {
+            "float": (4, "float32x4_t"),
+            "double": (2, "float64x2_t"),
+            "_Float16": (8, "float16x8_t"),
+        }
 
         if not prim_type in vec_types.keys():
             raise MemGenError(f"{srcinfo}: Neon vectors must be f32/f64 (for now)")
@@ -64,6 +68,8 @@ class Neon(Memory):
         if idxs:
             idxs = "[" + "][".join(idxs) + "]"
         return f"{baseptr}{idxs}"
+
+
 # --------------------------------------------------------------------------- #
 #   f32 Neon intrinsics
 # --------------------------------------------------------------------------- #
@@ -150,6 +156,7 @@ def neon_vmul_4xf32(dst: [f32][4] @ Neon, lhs: [f32][4] @ Neon, rhs: [f32][4] @
     for i in seq(0, 4):
         dst[i] = lhs[i] * rhs[i]
 
+
 @instr("{dst_data} = vmulq_f32({dst_data}, {rhs_data});")
 def neon_vmul2_4xf32(dst: [f32][4] @ Neon, lhs: [f32][4] @ Neon, rhs: [f32][4] @ Neon):
     assert stride(dst, 0) == 1
@@ -172,10 +179,15 @@ def neon_vfmla_4xf32_4xf32(
     for i in seq(0, 4):
         dst[i] += lhs[i] * rhs[lane]
 
-#This function uses an extra buffer for a beta=0 approach
+
+# This function uses an extra buffer for a beta=0 approach
 @instr("{dst_data} = vfmaq_laneq_f32({b_data}, {lhs_data}, {rhs_data}, {lane});")
 def neon_vfmla2_4xf32_4xf32(
-        dst: [f32][4] @ Neon, b: [f32][4] @ Neon, lhs: [f32][4] @ Neon, rhs: [f32][4] @ Neon, lane: index
+    dst: [f32][4] @ Neon,
+    b: [f32][4] @ Neon,
+    lhs: [f32][4] @ Neon,
+    rhs: [f32][4] @ Neon,
+    lane: index,
 ):
     assert stride(dst, 0) == 1
     assert stride(lhs, 0) == 1
@@ -186,6 +198,7 @@ def neon_vfmla2_4xf32_4xf32(
     for i in seq(0, 4):
         dst[i] = b[i] + lhs[i] * rhs[lane]
 
+
 @instr("{dst_data} = vmlaq_f32({dst_data}, {lhs_data}, {rhs_data});")
 def neon_vfmadd_4xf32_4xf32(
     dst: [f32][4] @ Neon, lhs: [f32][4] @ Neon, rhs: [f32][4] @ Neon
@@ -197,9 +210,13 @@ def neon_vfmadd_4xf32_4xf32(
     for i in seq(0, 4):
         dst[i] += lhs[i] * rhs[i]
 
+
 @instr("{dst_data} = vmlaq_f32({res_data}, {lhs_data}, {rhs_data});")
 def neon_vfmadd_ex_4xf32_4xf32(
-        dst: [f32][4] @ Neon, res: [f32][4] @ Neon, lhs: [f32][4] @ Neon, rhs: [f32][4] @ Neon
+    dst: [f32][4] @ Neon,
+    res: [f32][4] @ Neon,
+    lhs: [f32][4] @ Neon,
+    rhs: [f32][4] @ Neon,
 ):
     assert stride(dst, 0) == 1
     assert stride(res, 0) == 1
@@ -232,11 +249,12 @@ def neon_vfmadd_1xf32_4xf32(
         dst[i] += lhs[0] * rhs[i]
 
 
-#-----------------------------------------------
+# -----------------------------------------------
 # Load, Store, Broadcast, FMAdd, Mul, Add?
 #
 # float16
 
+
 @instr("{dst_data} = vld1q_f16((float16_t *)&{src_data});")
 def neon_vld_8xf16(dst: [f16][8] @ Neon, src: [f16][8] @ DRAM):
     assert stride(src, 0) == 1
@@ -272,9 +290,7 @@ def neon_zero_8xf16(dst: [f16][8] @ Neon):
 
 
 @instr("{dst_data} = vaddq_f16({lhs_data}, {rhs_data});")
-def neon_vadd_8xf16(
-    dst: [f16][8] @ Neon, lhs: [f16][8] @ Neon, rhs: [f16][8] @ Neon
-):
+def neon_vadd_8xf16(dst: [f16][8] @ Neon, lhs: [f16][8] @ Neon, rhs: [f16][8] @ Neon):
     assert stride(dst, 0) == 1
     assert stride(lhs, 0) == 1
     assert stride(rhs, 0) == 1
@@ -284,9 +300,7 @@ def neon_vadd_8xf16(
 
 
 @instr("{dst_data} = vmulq_f16({lhs_data}, {rhs_data});")
-def neon_vmul_8xf16(
-    dst: [f16][8] @ Neon, lhs: [f16][8] @ Neon, rhs: [f16][8] @ Neon
-):
+def neon_vmul_8xf16(dst: [f16][8] @ Neon, lhs: [f16][8] @ Neon, rhs: [f16][8] @ Neon):
     assert stride(dst, 0) == 1
     assert stride(lhs, 0) == 1
     assert stride(rhs, 0) == 1
@@ -294,9 +308,10 @@ def neon_vmul_8xf16(
     for i in seq(0, 8):
         dst[i] = lhs[i] * rhs[i]
 
+
 @instr("{dst_data} = vfmaq_laneq_f16({dst_data}, {lhs_data}, {rhs_data}, {lane});")
 def neon_vfmla_8xf16_8xf16(
-        dst: [f16][8] @ Neon, lhs: [f16][8] @ Neon, rhs: [f16][8] @ Neon, lane: index
+    dst: [f16][8] @ Neon, lhs: [f16][8] @ Neon, rhs: [f16][8] @ Neon, lane: index
 ):
     assert stride(dst, 0) == 1
     assert stride(lhs, 0) == 1
@@ -305,7 +320,7 @@ def neon_vfmla_8xf16_8xf16(
     assert lane < 8
     for i in seq(0, 8):
         dst[i] += lhs[i] * rhs[lane]
-        
+
 
 @instr("{dst_data} = vfmaq_f16({dst_data}, {lhs_data}, {rhs_data});")
 def neon_vfmadd_8xf16_8xf16(
@@ -318,9 +333,13 @@ def neon_vfmadd_8xf16_8xf16(
     for i in seq(0, 8):
         dst[i] += lhs[i] * rhs[i]
 
+
 @instr("{dst_data} = vfmaq_f16({res_data}, {lhs_data}, {rhs_data});")
 def neon_vfmadd_ex_8xf16_8xf16(
-        dst: [f16][8] @ Neon, res: [f16][8] @ Neon, lhs: [f16][8] @ Neon, rhs: [f16][8] @ Neon
+    dst: [f16][8] @ Neon,
+    res: [f16][8] @ Neon,
+    lhs: [f16][8] @ Neon,
+    rhs: [f16][8] @ Neon,
 ):
     assert stride(dst, 0) == 1
     assert stride(res, 0) == 1
diff --git a/src/exo/query_asts.py b/src/exo/query_asts.py
index 9edee3e..1b69e27 100644
--- a/src/exo/query_asts.py
+++ b/src/exo/query_asts.py
@@ -97,10 +97,12 @@ class Stmt(QueryAST):
 class R(Type):
     pass
 
+
 @_dataclass
 class f16(Type):
     pass
 
+
 @_dataclass
 class f32(Type):
     pass


Test patch:
diff --git a/tests/golden/test_schedules/test_simple_inline.txt b/tests/golden/test_schedules/test_simple_inline.txt
index 5d6a4d9..34034d4 100644
--- a/tests/golden/test_schedules/test_simple_inline.txt
+++ b/tests/golden/test_schedules/test_simple_inline.txt
@@ -1,7 +1,9 @@
 def bar(n: size, src: i8[n] @ DRAM, dst: i8[n] @ DRAM):
     for i in seq(0, n):
-        tmp_src: i8 @ DRAM
+        tmp_src1: i8 @ DRAM
+        tmp_src2: i8 @ DRAM
+        tmp_src1 = src[i]
+        tmp_src2 = src[i]
         tmp_dst: i8 @ DRAM
-        tmp_src = src[i]
         tmp_dst = dst[i]
-        tmp_dst = tmp_src + tmp_src
\ No newline at end of file
+        tmp_dst = tmp_src1 + tmp_src2
\ No newline at end of file
diff --git a/tests/golden/test_window/test_normalize.txt b/tests/golden/test_window/test_normalize.txt
index 2794e56..f309d51 100644
--- a/tests/golden/test_window/test_normalize.txt
+++ b/tests/golden/test_window/test_normalize.txt
@@ -73,6 +73,6 @@ EXO_ASSUME(m > 4);
 float xy;
 float y2;
 dot(ctxt,m,(struct exo_win_1f32c){ &x[m], { 1 } },(struct exo_win_1f32c){ &y[2], { n } },&xy);
-dot(ctxt,m,(struct exo_win_1f32c){ &y[3], { n } },(struct exo_win_1f32c){ &y[3], { n } },&y2);
+dot(ctxt,m,(struct exo_win_1f32c){ &y[3], { n } },(struct exo_win_1f32c){ &x[(2) * m], { 1 } },&y2);
 }
 
diff --git a/tests/test_effects.py b/tests/test_effects.py
index a23c925..c39e754 100644
--- a/tests/test_effects.py
+++ b/tests/test_effects.py
@@ -4,6 +4,7 @@ import pytest
 
 from exo import proc, DRAM
 from exo.libs.memories import GEMM_SCRATCH
+from exo.stdlib.scheduling import SchedulingError
 
 
 # ------- Effect check tests ---------
@@ -288,19 +289,18 @@ def test_race3(golden):
 
 
 # one big issue is aliasing in sub-procedure arguments
-# TODO: Think about this behaviour
-def test_race4(golden):
+def test_race4():
     @proc
     def foo(n: size, x: [R][n, n], y: [R][n, n]):
         for i in seq(0, n):
             if i + 1 < n:
                 x[i, i] = y[i, i]
 
-    @proc
-    def bar(n: size, z: R[n, n]):
-        foo(n, z, z)
+    with pytest.raises(SchedulingError, match="Cannot Pass the same buffer"):
 
-    assert bar.show_effects() == golden
+        @proc
+        def bar(n: size, z: R[n, n]):
+            foo(n, z, z)
 
 
 def test_div1(golden):
diff --git a/tests/test_neon.py b/tests/test_neon.py
index 2d0d37e..8a3cb19 100644
--- a/tests/test_neon.py
+++ b/tests/test_neon.py
@@ -131,10 +131,12 @@ def test_neon_vfmla():
 
     return simplify(simple_neon_vfmla)
 
+
 @pytest.mark.isa("neon")
 def test_gen_neon_vfmla(golden, test_neon_vfmla):
     assert str(test_neon_vfmla) == golden
 
+
 @pytest.fixture
 def test_neon_vfmla_f16():
     """
diff --git a/tests/test_new_eff.py b/tests/test_new_eff.py
index f8e0161..31f0736 100644
--- a/tests/test_new_eff.py
+++ b/tests/test_new_eff.py
@@ -321,3 +321,18 @@ def test_scal():
 
     with pytest.raises(SchedulingError):
         sscal = reorder_stmts(sscal, stmt_c.expand(1, 0))
+
+
+def test_alias_check():
+    """This fails when we have an aliasing in a proc"""
+
+    @proc
+    def foo(N: size, x: [f32][N], y: [f32][N]):
+        for i in seq(0, N):
+            x[i] = y[i]
+
+    with pytest.raises(SchedulingError, match="Cannot Pass the same buffer"):
+
+        @proc
+        def bar(N: size, x: [f32][N]):
+            foo(N, x, x)
diff --git a/tests/test_schedules.py b/tests/test_schedules.py
index 92c2bb6..1e33b1e 100644
--- a/tests/test_schedules.py
+++ b/tests/test_schedules.py
@@ -1147,20 +1147,15 @@ def test_simple_inline(golden):
     @proc
     def bar(n: size, src: i8[n], dst: i8[n]):
         for i in seq(0, n):
-            tmp_src: i8
+            tmp_src1: i8
+            tmp_src2: i8
+            tmp_src1 = src[i]
+            tmp_src2 = src[i]
             tmp_dst: i8
-            tmp_src = src[i]
             tmp_dst = dst[i]
-            foo(tmp_src, tmp_src, tmp_dst)
+            foo(tmp_src1, tmp_src2, tmp_dst)
 
-    # TODO: these should fail
-    # with pytest.raises(SchedulingError, match='blah'):
-    #     inline(bar, 'foo(_)')
-    #
-    # with pytest.raises(SchedulingError, match='blah'):
-    #     inline(bar, 'foo(io, i1, i2)')
-
-    bar = inline(bar, "foo(_, _, _)")
+    bar = inline(bar, "foo(_)")
     assert str(bar) == golden
 
 
diff --git a/tests/test_window.py b/tests/test_window.py
index d6815ce..5d0b32a 100644
--- a/tests/test_window.py
+++ b/tests/test_window.py
@@ -54,9 +54,10 @@ def test_input4():
             dst2[i] = src2[i]
 
     @proc
-    def bar(x: [f32][8]):
+    def bar(x: [f32][8], y: [f32][8]):
         assert stride(x, 0) == 1
-        foo(x, x)
+        assert stride(y, 0) == 1
+        foo(x, y)
 
 
 def test_window(golden):
@@ -123,6 +124,6 @@ def test_normalize(golden):
         xy: f32
         y2: f32
         dot(m, x[1, :], y[:, 2], xy)
-        dot(m, y[:, 3], y[:, 3], y2)
+        dot(m, y[:, 3], x[2, :], y2)
 
     assert proj.c_code_str() == golden
diff --git a/tests/test_x86.py b/tests/test_x86.py
index b193dad..f762701 100644
--- a/tests/test_x86.py
+++ b/tests/test_x86.py
@@ -65,11 +65,12 @@ def test_avx2_simple_math(compiler):
         assert n % 8 == 0
         for i in seq(0, n / 8):
             xVec: f32[8] @ AVX2
+            tmp: f32[8] @ AVX2
             yVec: f32[8] @ AVX2
             mm256_loadu_ps(xVec, x[8 * i : 8 * i + 8])
             mm256_loadu_ps(yVec, y[8 * i : 8 * i + 8])
-            mm256_mul_ps(xVec, xVec, yVec)
-            mm256_mul_ps(xVec, xVec, yVec)
+            mm256_mul_ps(tmp, xVec, yVec)
+            mm256_mul_ps(xVec, tmp, yVec)
             mm256_storeu_ps(x[8 * i : 8 * i + 8], xVec)
 
     fn = compiler.compile(