[FEAT] Add the race detector's multipath and loaded-value reader facts

tests/unit/ir/reader_kernels.py

@@ -361,6 +361,70 @@ def poll_while(flag_ptr, out_ptr, n):
     tl.store(out_ptr + pid, 1)
 
 
+@triton.jit
+def mp_loops(x_ptr, n, m):
+    # nested loops (the inner pointer advances by the outer induction
+    # variable), then a sequential loop under a branch
+    for i in range(0, n):
+        q = x_ptr + i
+        for j in range(0, m, 2):
+            tl.store(q + j, 1)
+            q += i
+    if tl.program_id(0) < n:
+        p = x_ptr
+        for k in range(0, n):
+            tl.store(p, 2)
+            p += m
+
+
+@triton.jit
+def mp_loaded(idx_ptr, x_ptr, f_ptr, n, B: tl.constexpr):
+    # integer loads in an address, a mask and a branch; a float load and a
+    # dropped mask; a loaded tile expanded; masks with a modelable conjunct
+    # (kept, also expanded) or disjunct (dropped); a float atomic operand
+    offs = tl.arange(0, B)
+    j = tl.load(idx_ptr + offs, mask=offs < n, other=0)
+    tl.store(x_ptr + j, 1.0)
+    v = tl.load(f_ptr + offs)
+    tl.store(x_ptr + offs, v, mask=j > 0)
+    t = tl.load(idx_ptr + offs, mask=offs < n, other=0)
+    tl.store(x_ptr + t[:, None] + offs[None, :], 3.0)
+    w = tl.load(idx_ptr + offs, mask=v > 0)
+    m = (offs < n) & (w > 0)
+    tl.store(x_ptr + offs, 4.0, mask=m)
+    if tl.load(idx_ptr) > 0:
+        tl.store(x_ptr, 5.0)
+    tl.store(x_ptr + offs, 6.0, mask=(offs < n) | (w > 0))
+    tl.store(x_ptr + offs[:, None] * B + offs[None, :], 7.0, mask=m[:, None])
+    tl.atomic_add(f_ptr + offs, v)
+
+
+@triton.jit
+def mp_early_return(x_ptr, f_ptr, n):
+    # an early return under a float condition: the rest is guarded
+    pid = tl.program_id(0)
+    if tl.load(f_ptr) > 0.0:
+        return
+    tl.store(x_ptr + pid, 1.0)
+
+
+@triton.jit
+def mp_loaded_advance(idx_ptr, x_ptr, n):
+    # an advance by a value loaded in the loop: variant
+    p = x_ptr
+    for i in range(0, n):
+        tl.store(p, 1)
+        tl.store(idx_ptr, i)
+        p += tl.load(idx_ptr)
+
+
+@triton.jit
+def mp_for_in_while(flag_ptr, x_ptr, n):
+    while tl.load(flag_ptr) == 0:
+        for i in range(0, n):
+            tl.store(x_ptr + i, 1)
+
+
 _F32 = "*fp32"
 # corpus name (reader_ttir/<name>.ttir, the kernel of that name) ->
 # (signature, constexprs), compiled for sm80 with the default options
@@ -425,4 +489,12 @@ def poll_while(flag_ptr, out_ptr, n):
     "fences": ({"x_ptr": "*i32", "n": "i32"}, {}),
     "pid_in_combine": ({"x_ptr": "*i32", "o_ptr": "*i32", "B": "constexpr"}, {"B": 8}),
     "poll_while": ({"flag_ptr": "*i32", "out_ptr": "*i32", "n": "i32"}, {}),
+    "mp_loops": ({"x_ptr": "*i32", "n": "i32", "m": "i32"}, {}),
+    "mp_loaded": (
+        {"idx_ptr": "*i32", "x_ptr": _F32, "f_ptr": _F32, "n": "i32", "B": "constexpr"},
+        {"B": 8},
+    ),
+    "mp_early_return": ({"x_ptr": _F32, "f_ptr": _F32, "n": "i32"}, {}),
+    "mp_loaded_advance": ({"idx_ptr": "*i32", "x_ptr": "*i32", "n": "i32"}, {}),
+    "mp_for_in_while": ({"flag_ptr": "*i32", "x_ptr": "*i32", "n": "i32"}, {}),
 }

tests/unit/ir/test_lowering.py

@@ -27,6 +27,7 @@
     IntCast,
     IterArgInfo,
     IterArgOffset,
+    Loaded,
     LoopInfo,
     LoopVar,
     Not,
@@ -75,6 +76,9 @@ def iteration(self, loop_ssa):
     def observed(self, t):
         return self._var(f"observed_{t.access_index}")
 
+    def loaded(self, t):
+        return self._var(f"loaded_{t.access_index}")
+
     def data_dep(self, t):
         raise _Refusal(t.why)
 
@@ -201,7 +205,9 @@ def kids(t) -> list[int]:
     assert kids(LoopVar("%loop")) == [id(LOOP.lower), id(LOOP.step)]
     for leaf in (Const(1), Pid(0), NumPrograms(1), Arange("%r", 0, 4), X, Observed(0)):
         assert kids(leaf) == []
-    assert kids(DataDep()) == []
+    # a loaded value's terms are the client's; keep is no value dependency
+    assert kids(DataDep()) == kids(DataDep("x", keep=less)) == []
+    assert kids(Loaded(0, X, less, Y)) == []
 
 
 # ─────────────────────────── operator semantics ───────────────────────────
@@ -325,11 +331,30 @@ def test_loop_terms_read_the_clients_iteration():
     assert _proved(lowerer.lower(IterArgOffset(0)) == _v("o0") + k * _v("d0"))
     assert _proved(lowerer.lower(IterArgOffset(1)) == _v("o1") + other * _v("d1"))
     assert "hi" not in leaves.made  # the upper bound is not the variable's value
+    # with several loops, each variable reads its own loop by name
+    inner = LoopInfo(
+        "%loop1", lower=Param("lo1"), upper=Param("hi1"), step=Param("st1")
+    )
+    graph = AccessGraph("k", (), (), None, loops=(LOOP, inner))
+    lowerer = Lowerer(graph, _Leaves())
+    assert _proved(
+        lowerer.lower(LoopVar("%loop1")) == _v("lo1") + _v("k%loop1") * _v("st1")
+    )
+    assert _proved(lowerer.lower(LoopVar("%loop")) == _v("lo") + k * _v("st"))
+    with pytest.raises(ValueError, match="without a loop named '%loop2'"):
+        lowerer.lower(LoopVar("%loop2"))
 
 
 # ─────────────────────────── leaves and bugs ───────────────────────────
 
 
+def test_a_loaded_value_is_the_clients_leaf():
+    leaves = _Leaves()
+    term = Bin("+", X, Loaded(3, Param("o"), None, None))
+    assert _proved(Lowerer(_graph(), leaves).lower(term) == _v("x") + _v("loaded_3"))
+    assert "o" not in leaves.made
+
+
 def test_a_leaf_refusal_propagates_unchanged():
     with pytest.raises(_Refusal, match="loaded value"):
         Lowerer(_graph(), _Leaves()).lower(BoolBin("and", X, DataDep("loaded value")))

tests/unit/ir/test_mlir_walk.py

@@ -226,6 +226,21 @@ def test_cf_blocks_align():
     labels = [m.blocks[b].label for b in func.regions[0]]
     assert labels[0] is None and "^bb5" in labels and len(labels) == 6
     assert all(op.loc is not None for op in _ops(m, "cf.br") + _ops(m, "cf.cond_br"))
+    # a branch's successors are its target blocks, read from the text
+    by_label = {m.blocks[b].label: b for b in func.regions[0][1:]}
+    succ = [op.attrs["successors"] for op in _ops(m, "cf.br") + _ops(m, "cf.cond_br")]
+    assert succ == [
+        (by_label["^bb5"],),
+        (by_label["^bb1"], by_label["^bb2"]),
+        (by_label["^bb3"], by_label["^bb4"]),
+        (by_label["^bb1"], by_label["^bb5"]),
+    ]
+    same = W._walk(_text("crafted_same_dest.ttir"))
+    first = _ops(same, "cf.cond_br")[0]
+    assert first.attrs["successors"] == (first.block + 1,) * 2  # ^bb1(%a), ^bb1(%b)
+    bad = _text("crafted_same_dest.ttir").replace("^bb2(%y, %x : i32, i32)", "^bb2")
+    with pytest.raises(W.MisalignedModule, match="0 operands for its arguments"):
+        W._walk(_text("crafted_same_dest.ttir"), scan_text=bad)
 
 
 def test_parser_invented_locs_are_absent():