[async] Add element-wise info into TaskMeta for fusion

[xumingkuan]

Related issue = #742

Changes:

• Add std::unordered_map<SNode *, bool> TaskMeta::element_wise
• Refactor AsyncStateHash into std::hash<AsyncState> because I think we don't need other hashes of AsyncState
• Fix typo "fusable" -> "fusible"

Question:

• How does the element-wise info affect mask/list states?

[Click here for the format server]

---

[codecov]

Codecov Report

Merging #1884 into master will not change coverage.
The diff coverage is n/a.

@@           Coverage Diff           @@
##           master    #1884   +/-   ##
=======================================
  Coverage   43.19%   43.19%           
=======================================
  Files          44       44           
  Lines        6330     6330           
  Branches     1092     1092           
=======================================
  Hits         2734     2734           
  Misses       3426     3426           
  Partials      170      170           

---

Continue to review full report at Codecov.

Legend - Click here to learn more
Δ = absolute <relative> (impact), ø = not affected, ? = missing data
Powered by Codecov. Last update f2a798d...427453b. Read the comment docs.

[yuanming-hu]

Awesome!! LGTM.

How does the element-wise info affect mask/list states?

Great question. The short answer is we don't need to worry about it.
If a write access is element-wise, the activation will be demoted here:

taichi/taichi/transforms/flag_access.cpp

Line 114 in f2a798d

void visit(GlobalPtrStmt *stmt) {

So by the time you are doing the is_element_wise test, whether a task needs list generation/gc/mask state output is already determined. We don't need to handle that once more.

(Sorry I didn't notice we already have a similar implementation of is_element_wise there - maybe we can unify two implementations sometime in the future.)

[xumingkuan]

If a write access is element-wise, the activation will be demoted here:

taichi/taichi/transforms/flag_access.cpp

Line 114 in f2a798d

void visit(GlobalPtrStmt *stmt) {

(Sorry I didn't notice we already have a similar implementation of is_element_wise there - maybe we can unify two implementations sometime in the future.)

Interesting! It looks like the element-wise check in flag_access also includes a sparsity check. For example, if x and y have different shapes on the parent sparse SNode (x is smaller than y) or they have the same shape but their parent sparse SNodes are different, y will not be considered element-wise in

for I in ti.grouped(x):
    y[I] = 1

(y[I] needs activation here). However, for fusion, we can consider the task as element-wise on y.

A fusible example:

x = ti.field(ti.i32)
y = ti.field(ti.i32)
ti.root.pointer(ti.i, 4).place(x)
ti.root.pointer(ti.i, 8).place(y)

@ti.kernel
def foo():
    for i in x:
        y[i] = 1

@ti.kernel
def bar():
    for i in x:
        y[i] = 2

foo()
bar()

ti.sync()

Now I think the mask/list states have the same criteria as value states here -- if task a has an edge to task b (struct for or range for tasks), no matter which type the edge's state is, only if both tasks are element-wise on the edge's SNode, can a and b be fused.
Now I think it's fine to not check mask/list state edges -- if there are these edges, there must be value state edges too. Mask state edges can be of parent SNodes, so it would be easier to check value state edges directly.

[xumingkuan]

Why Some checks haven’t completed yet -- Build and Test (CPU) (ubuntu-latest, 3.6, OFF) Expected — Waiting for status to be reported?