[apache/tvm] [TIR][REFACTOR][RFC] ForNode -- Introduce Annotations and ThreadBinding to for_type (#7302)

[Tianqi Chen <no...@github.com.INVALID>]

This proposed refactor is part of TensorIR refactoring. https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872

To enable the TensorIR scheduling, we will need to introduce support for more annotations to the loop node, specifically two fields will be introduced:

- Optional<IterVar> thread_binding: Used for optionally bind the loop to a context thread such as GPU's blockIdx.x
- Map<String, Object> annotations: Used for additional hint information used in high level transformations

We will also remove the legacy `device_api` field as the field is no longer being used

The refactor will change the IR construction API of ForNode, which requires migration of the code, but won't results in major refactor of the existing code as most cases CopyOnWrite are used. 


-- 
You are receiving this because you are subscribed to this thread.
Reply to this email directly or view it on GitHub:
https://github.com/apache/tvm/issues/7302

Re: [apache/tvm] [TIR][REFACTOR][RFC] ForNode -- Introduce Annotations and ThreadBinding to for_type (#7302)

[Bohan Hou <no...@github.com.INVALID>]

The problem is that launch_thread was not designed to print those additional hints，since launch_thread was syntax for Attr  now.

-- 
You are receiving this because you are subscribed to this thread.
Reply to this email directly or view it on GitHub:
https://github.com/apache/tvm/issues/7302#issuecomment-762306580

Re: [apache/tvm] [TIR][REFACTOR][RFC] ForNode -- Introduce Annotations and ThreadBinding to for_type (#7302)

[Tianqi Chen <no...@github.com.INVALID>]

Closed #7302.

-- 
You are receiving this because you are subscribed to this thread.
Reply to this email directly or view it on GitHub:
https://github.com/apache/tvm/issues/7302#event-4222721789

Re: [apache/tvm] [TIR][REFACTOR][RFC] ForNode -- Introduce Annotations and ThreadBinding to for_type (#7302)

[Tianqi Chen <no...@github.com.INVALID>]

@spectrometerHBH it would be great if you can take a stab to get the first strawman syntax and parser

-- 
You are receiving this because you are subscribed to this thread.
Reply to this email directly or view it on GitHub:
https://github.com/apache/tvm/issues/7302#issuecomment-762330423

Re: [apache/tvm] [TIR][REFACTOR][RFC] ForNode -- Introduce Annotations and ThreadBinding to for_type (#7302)

[Junru Shao <no...@github.com.INVALID>]

I like the proposal, especially the design philosophy that separates `thread_binding` and other `annotations`, where `thread_binding` are required to be taken care of by each pass, and `annotations` are optional and are safe to ignore in irrelevant passes.

One question: what is the implication for the roundtrippable TIR parser? How do we print `thread_binding` in this case?


-- 
You are receiving this because you are subscribed to this thread.
Reply to this email directly or view it on GitHub:
https://github.com/apache/tvm/issues/7302#issuecomment-761948446

Re: [apache/tvm] [TIR][REFACTOR][RFC] ForNode -- Introduce Annotations and ThreadBinding to for_type (#7302)

[Tianqi Chen <no...@github.com.INVALID>]

I see, I don't think the thread binding can come with additional annotations, which might simplify the problem (if we enforce such check in IR validation). 

-- 
You are receiving this because you are subscribed to this thread.
Reply to this email directly or view it on GitHub:
https://github.com/apache/tvm/issues/7302#issuecomment-762308405

Re: [apache/tvm] [TIR][REFACTOR][RFC] ForNode -- Introduce Annotations and ThreadBinding to for_type (#7302)

[Tianqi Chen <no...@github.com.INVALID>]

https://github.com/apache/tvm/pull/7306

-- 
You are receiving this because you are subscribed to this thread.
Reply to this email directly or view it on GitHub:
https://github.com/apache/tvm/issues/7302#issuecomment-762863140

Re: [apache/tvm] [TIR][REFACTOR][RFC] ForNode -- Introduce Annotations and ThreadBinding to for_type (#7302)

[Tianqi Chen <no...@github.com.INVALID>]

Thanks @junrushao1994  We have already some mechanism to print the low level thread binding via attr (printed as launch_thread). When the threadIdx is shared in multiple launch point, the var is printed separately. Otherwise it is printed inline to the code. I believe we could adopt a similar strategy.

Note that because the initial refactor is only wrt to the data structure, we can defer the parsing syntax decision until the parser support gets in. but would be great to discuss it.

cc @spectrometerHBH @tkonolige 

-- 
You are receiving this because you are subscribed to this thread.
Reply to this email directly or view it on GitHub:
https://github.com/apache/tvm/issues/7302#issuecomment-762303027