You are viewing a plain text version of this content. The canonical link for it is here.
Posted to dev@tvm.apache.org by Siyuan Feng via Apache TVM Discuss <no...@discuss.tvm.ai> on 2020/09/10 16:52:23 UTC
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
## Background and Motivation
TVM is an end-to-end deep learning compiler with two levels of IR and optimization. TVM translates popular DL frameworks into Relay and optimizes the computation graph, after which it lowers each graph node into Tensor Expression(TE) and does another function-level optimization before finally lowering it into TIR and generating backend code. In brief, the current workflow is
```
TF/PyTorch/ONNX -> Relay -> TE (based on TE schedule) -> TIR -> C++/CUDA
```
Currently, low-level optimization is done through TE scheduling, which has several limitations:
- Based on an accessory data structure: schedule tree. Schedule primitives operate on schedule tree rather than TIR itself directly, which makes the scheduling result less intuitive.
- All primitives are coupled on the schedule tree representation, it is not easy to add new primitives and ensure the correctness of transformation at the same time.
- Limited high-dimension instruction and tensorization support. TE is not schedulable after tensorization. The description of tensor intrinsics is not friendly to users.
## Introduction
TensorIR is a brand new low-level IR with full scheduling support. Here are some key features and novel techniques.
### Core data structure: Block and Block Realize
To generalize the high-dimension tensor expression, we introduce a new statement structure *Block*. A block can wrap any part of the IR to provide isolation. A block is *the minimal unit* for scheduling and tensorization, which stores all the fundamental information of computation including block iteration vars, the region that the block reads and writes, the buffer which allocated inside the block, and the critical body statement. Block declare the expression computation with its iteration vars and types.
### Workload Example
1. First of all, we will define the workload(gemm in the example) with hybrid. Note that the hybrid scipt direct generate a **TIR** program rather than TE stages. (We can auto-complete the loop nesting with block definition by default)
```python
@tvm.hybrid.script
def matmul(a: ty.handle, b: ty.handle, c: ty.handle) -> None:
C = tir.match_buffer(c, (1024, 1024), "float32")
A = tir.match_buffer(a, (1024, 1024), "float32")
B = tir.match_buffer(b, (1024, 1024), "float32")
reducer = tir.comm_reducer(lambda x, y: x + y, tir.float32(0))
with tir.block([1024, 1024, tir.reduce_axis(0, 1024)], "C") as [vi, vj, vk]:
reducer.step(C[vi, vj], A[vi, vk] * B[vk, vj])
```
2. Then create the schedule from **TIR** and tile the loops using primitives.
```python
s = tir.create_schedule(matmul)
update = s.get_block("C")
i, j, k = s.get_axes(update)
i_o, i_i = s.split(i, bn)
j_o, j_i = s.split(j, bn)
k_o, k_i = s.split(k, 4)
s.reorder(i_o, j_o, k_o, k_i, i_i, j_i)
```
Result
```python
def func(a: ty.handle, b: ty.handle, c: ty.handle) -> None:
# function attr dict
B = tir.match_buffer(b, [1024, 1024])
A = tir.match_buffer(a, [1024, 1024])
C = tir.match_buffer(c, [1024, 1024])
reducer = tir.comm_reducer(lambda x, y: x + y, tir.float32(0))
# body
for i0_outer, i1_outer, i2_outer, i2_inner, i0_inner, i1_inner in tir.grid(32, 32, 256, 4, 32, 32):
with tir.block([1024, 1024, tir.reduce_axis(0, 1024)], "C") as [vi, vj, vk]:
tir.bind(vi, ((i0_outer*32) + i0_inner))
tir.bind(vj, ((i1_outer*32) + i1_inner))
tir.bind(vk, ((i2_outer*4) + i2_inner))
reducer.step(C[vi, vj], (A[vi, vk]*B[vk, vj]))
```
3. Vectorize and decompose reduction
```python
s.vectorize(j_i)
s.decompose_reduction(update, j_o)
```
Result
```python
@tvm.hybrid.script
def func(a: ty.handle, b: ty.handle, c: ty.handle) -> None:
# function attr dict
B = tir.match_buffer(b, [1024, 1024])
A = tir.match_buffer(a, [1024, 1024])
C = tir.match_buffer(c, [1024, 1024])
# body
for i0_outer in, i1_outer_init, i0_inner_init in tir.grid(32, 32, 32):
for i1_inner_init in range(0, 32, annotation = {"loop_type":"vectorize"}):
with tir.block([1024, 1024], "C_init") as [vi_init, vj_init]:
tir.bind(vi_init, ((i0_outer*32) + i0_inner_init))
tir.bind(vj_init, ((i1_outer_init*32) + i1_inner_init))
C[vi_init, vj_init] = tir.float32(0)
for i1_outer, i2_outer, i2_inner, i0_inner in tir.grid(32, 256, 4, 32):
for i1_inner in range(0, 32, annotation = {"loop_type":"vectorize"}):
with tir.block([1024, 1024, tir.reduce_axis(0, 1024)], "C_update") as [vi, vj, vk]:
tir.bind(vi, ((i0_outer*32) + i0_inner))
tir.bind(vj, ((i1_outer*32) + i1_inner))
tir.bind(vk, ((i2_outer*4) + i2_inner))
C[vi, vj] = (C[vi, vj] + (A[vi, vk]*B[vk, vj]))
```
4. Print it and run
```python
build_func = tvm.build(s.func, target=target)
build_func(a, b, c)
tvm.testing.assert_allclose(c.asnumpy(), np.matmul(a.asnumpy(), b.asnumpy()), rtol=1e-5)
evaluator = build_func.time_evaluator(build_func.entry_name, ctx, number=1)
```
In this example, we are imperatively changing the IR rather than waiting until the end to change(TE). It is very important for users and developers to directly see what happens during the scheduling. Also, at every stage during the schedule, we can get a *verifiable* IR, that's the major improvement for both user experience and correctness proof.
## Key Features
### Independent scheduling based on IR itself
Different from the TE schedule, TensorIR has a complete set of schedule algorithms, which does not need a schedule tree or any extra data structure. We will introduce a brand new set of schedule primitives and it has full backward compatibility for the TE schedule.
We simplify the compiling workload and conception.
```
TF/PyTorch/ONNX -> Relay -> TIR -> schedule -> TIR -> scheudle -> TIR -> C++/CUDA
```
Now, there is no stage during the schedule. Rather than lowering the schedule into TIR, we directly mutate the TIR itself. Also, it enables the sequential schedule (schedule several times for a single workload).
### Stronger Expressiveness and Optimization Ability
TE has limited expressiveness since each stage is defined by `Stage = te.compute(lambda expr)`, while TensorIR is a full c++-like IR. We can write any program with TensorIR as you want. Although not all programs can be scheduled, there are still more workloads that can be optimized by TensorIR.
One of the improved tasks is concatenating:
#### TE
```python
B = te.compute(i, tvm.tir.if_then_else(i < 10, A0[i], tvm.tir.if_then_else(i < 20, A1[i - 10], A2[i - 20])
```
#### TensorIR:
```python
with tir.block([10]) as vi:
B[vi] = A0[vi]
with tir.block([10]) as vi:
B[vi + 10] = A1[vi]
with tir.block([10]) as vi:
B[vi + 20] = A2[vi]
```
The critical improvement is performance. In TIR we optimize the program by deprecating the `if` branch, which is impossible in the TE schedule.
### Memory and Execution Scope
Hardware accelerators led by GPUs are increasingly using hierarchical architecture, including memory hierarchy(global, shared, local/wmma in NV-GPU) and execution hierarchy(SM, warp, thread in NV-GPU). TVM defines the memory hierarchy and TensorIR provides corresponding native hierarchical block and execution scope. Different execution scope can access different memory scope.
TensorIR natively supports hierarchy checks. We will check the memory access and thread binding, including warp level instruction(wmma) validation during the schedule. Following is an example of the GPU hierarchy.
```python
for bx in range(0, 32, annotation = {"loop_type":"blockIdx.x"}):
with block(exec_scope="gpu_block"):
for ty in range(0, 32, annotation = {"loop_type":"threadIdx.y"}):
with block(exec_scope="gpu_warp"):
for ty in range(0, 32, annotation = {"loop_type":"threadIdx.x"}):
with block(exec_scope="gpu_thread"):
A[i] = B[i] + 1
```
### High-dimension Scheduling and Tensorization.
With more and more backend provides tensor operator and instruction, the limitation of the TE schedule shows off. It is hard to tensorize a complex stage. TensorIR chooses a block (a wrap of high-dimension computation) as the minimum schedule unit. We can natively tensorize a sub-program and even do schedule after tensorization.
### Decouple Primitives and Correctness Proof
Since every primitive directly rewrite the AST (it works like a `StmtMutator`), we easily decouple primitives. Developers can add a schedule primitive as easy as add a pass. Also, it is easy to ensure the program has the same behavior between origin one and scheduled one by proving the correctness of each primitive and checking intermedia IR during the schedule.
### Round-trip Python Hybrid Syntax (Already upstream)
Hybrid enables developers to write TensorIR through python syntax, one of the most popular languages. Also, it provides a way to store the IR after or during the schedule. Please see the detail by https://discuss.tvm.apache.org/t/rfc-hybrid-script-support-for-tir/7516
## Migration Plan
1. Upstream TensorIR data structure.
1. Upstream TensorIR schedule primitive
1. Support AutoTVM/Ansor on TensorIR
Co-author @spectrometerHBH @tqchen @junrushao1994
I appreciate discuss and advise from @merrymercy
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/1) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/2fcebc91faa49032949e50d0909b466757f17f8503246f805d57d4a6782822d6).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Siyuan Feng via Apache TVM Discuss <no...@discuss.tvm.ai>.
Technically, it should support. However, due to time constraints, we have not yet supported.
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/25) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/a71731b7108ffbc73c1f296fb52a04700190c8932814f175baa0804269adbe2e).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Matt Barrett via Apache TVM Discuss <no...@discuss.tvm.ai>.
Thanks for this explanation. I'm interested if it might be possible to match tensor intrinsics with variable size? For example, Arm SVE introduces vector instructions of variable size.
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/24) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/7f7a92b4a5703274048fc74921bb6211fa4155d7785f5e10541c51e1cbf803c2).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Siyuan Feng via Apache TVM Discuss <no...@discuss.tvm.ai>.
Thank you for your interest.
Tensorize in TensorIR is completely different from the TE ones. In TensorIR, we use two functions (desc_func and intrin_func) to define an intrinsic. Here would be an example of intrinsic (Note that TensorIR is still WIP, so the API may be changed).
```python
@tvm.hybrid.script
def desc_func(a: ty.handle, b: ty.handle, c: ty.handle) -> None:
A = tir.match_buffer(a, [16, 16])
B = tir.match_buffer(b, [16, 16])
C = tir.match_buffer(c, [16, 16])
with tir.block([16, 16, tir.reduce_axis(0, 16)], "root") as [vi, vj, vk]:
for i, j, k in tir.grid(16, 16, 16):
with tir.block([16, 16, tir.reduce_axis(0, 16)], "update") as [vii, vjj, vkk]:
tir.bind(vii, vi + i)
tir.bind(vjj, vj + j)
tir.bind(vkk, vk + k)
C[vii, vjj] = C[vii, vjj] + A[vii, vkk] * B[vjj, vkk]
@tvm.hybrid.script
def intrin_func(a: ty.handle, b: ty.handle, c: ty.handle) -> None:
A = tir.match_buffer(a, [16, 16])
B = tir.match_buffer(b, [16, 16])
C = tir.match_buffer(c, [16, 16])
with tir.block([16, 16, tir.reduce_axis(0, 16)], "root") as [vi, vj, vk]:
tir.evaluate(tir.tvm_mma_sync(C.data, C.elem_offset // 256,
A.data, A.elem_offset // 256,
B.data, B.elem_offset // 256,
C.data, C.elem_offset // 256,
dtype="handle"))
```
Tensorize will match the sub-AST(usually is a block) with the `desc_func`, and then replace by `intrin_func`.
TensorIR is in the schedule level and has no coupling with low-level passes. However, we can directly schedule each loop directly and add primitives as you want. :slight_smile:
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/23) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/7938209de61c59e971405375dd4f261f54134d2cfc1c77eec98c83dbc9ea60ea).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Matt Barrett via Apache TVM Discuss <no...@discuss.tvm.ai>.
Thanks for this RFC, I think it's a great idea and will help solve a number of issues I've been facing recently. I'm particularly interested in what 'tensorize' will look like for this new IR. Could you give a snippet as an example?
I'm also interested in what the interaction of this will be with the loop partition pass. Will this mean that each partitioned loop will then be individually schedulable?
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/22) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/f077fc1b83f5b0d8820115fef9a0d6dbe5f25b0fd5488d8c73cd6f4b5f71aa48).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by tqchen via Apache TVM Discuss <no...@discuss.tvm.ai>.
TensorIR can be viewed as major feature enhancements(upgrade) to the TIR in master. That is why TensorIR and TIR are used interchangeably as they are supposed to be so.
Some of the elements like multidimensional buffer load and tvm script are already present as part of the unified IR effort.
The upgrade will happen quite naturally. With the TIR continue to support current code as as they are, and gains new scheduling capabilities with the new block constructs.
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/41) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/ae8056c981087ec9ceae54fef35fde4e0761b5af456a3db5512238ccfbdbbdd2).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Manupa Karunaratne via Apache TVM Discuss <no...@discuss.tvm.ai>.
Yes, the ambiguity is something I was struggling with too, when having a conversation. May I ask what does the "T" of old TIR stands for ? TVM ?
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/40) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/77baf2aff59764bbb237948b7030fdbccfd4f5fd085530df64a433da98edecc4).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Jason Knight via Apache TVM Discuss <no...@discuss.tvm.ai>.
When I read this RFC, I was confused because I read that the compilation flow **currently** goes from
> TF/PyTorch/ONNX -> Relay -> TE (based on TE schedule) -> TIR -> C++/CUDA
And then I read:
> TensorIR is a brand new low-level IR with full scheduling support. Here are some [...]
And then see references to `TIR` throughout the rest of the RFC, which seem unclear to me if these references are to the **new** `TensorIR`, or the **old** `TIR`.
Can we clarify both here and in the code base moving forward whether we are referring to `TensorIR` or the old `TIR`? I think there are two ways of doing this going forward and we should explicitly pick one:
1. Have `TIR` refer to the old version of TIR, and always use `TensorIR` when talking about the new IR.
2. Be clear that TensorIR moving forward will be replacing the old `TIR`, so that any reference to `TIR` could be referring to the new or old version (and should be clarified if not obvious from context).
Someone more knowledgeable than me should pick one of those (or correct me and/or point out other options if I've gotten anything wrong here).
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/39) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/793aaee1938382b9b1215945f8f86517563aa3bbf8683fc0f33e1b09fb523744).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Junru Shao via Apache TVM Discuss <no...@discuss.tvm.ai>.
@masahi Yeah. We should allow to embed IR fragments, functions that produce IR fragments, or replace some tokens with caller specified IR fragments into the script :-)
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/61) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/1317ecf60d421725a52f3ad01ec395a2c1a78dfd2fdbc95c14994552edaba6e9).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Siyuan Feng via Apache TVM Discuss <no...@discuss.tvm.ai>.
Thanks, @yzh119. Currently, we have not considered the cross-kernel schedule in TensorIR. But it may be possible if we make it as one large kernel. Could you please show an example? (e.g. the IR before and after the schedule)
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/57) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/40d9514623b75ff4a760bd4c91e43ce287a87de095b20b0c230597d1dfabd02a).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Zihao Ye via Apache TVM Discuss <no...@discuss.tvm.ai>.
[quote="Hzfengsy, post:1, topic:7872"]
#### TensorIR:
```
with tir.block([10]) as vi:
B[vi] = A0[vi]
with tir.block([10]) as vi:
B[vi + 10] = A1[vi]
with tir.block([10]) as vi:
B[vi + 20] = A2[vi]
```
The critical improvement is performance. In TIR we optimize the program by deprecating the `if` branch, which is impossible in the TE schedule.
[/quote]
This looks confusing to me because there is no example showing how to schedule the three blocks.
I'm curious if TIR's Block construct is strong enough to support holistic fusion like [Rammer](https://www.usenix.org/system/files/osdi20-ma.pdf) and [HFuse](https://arxiv.org/pdf/2007.01277.pdf), where different Block has heterogeneous workload but we can fuse them in a single kernel. If so we can make such fusion a primitive in TIR schedule and enlarge the auto-scheduling search space.
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/56) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/efec97d52d96bebf2d6c5966d951ac12d9a81dfec8fd6988ec19ddee04582b85).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Junru Shao via Apache TVM Discuss <no...@discuss.tvm.ai>.
@xqdan In Ansor, Fusion analysis is handled in TE with some straightforward heuristics, which I believe have covered our usecases. CC: @merrymercy @jcf94
Agree that ISL provides effective information about vectorization, and I believe there might be other competitive heuristics too. Tensorization is a more general topic that would be super interesting to explore :-)
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/33) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/781aa8cc2d490e9d898000241082a29a887d62556e043cba1c9e5b571e21c087).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Bohan Hou via Apache TVM Discuss <no...@discuss.tvm.ai>.
Thanks for your reply! @MinminSun
The cache_read/cache_write API accepts a Buffer and new scope as input, do some checks to ensure it brings no problem to read/write the Buffer into cache, and create new blocks to do the cache transfer.
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/21) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/64c91b3770f527ea95605e791b36970ef05e8e014849afa269e5235933e6ab01).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Minmin Sun via Apache TVM Discuss <no...@discuss.tvm.ai>.
Thanks for the proposal! Just courious about the schuedule primitives like cache_write and cache_read, since there are no stages in TensorIR.
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/20) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/995404f2b4f095d499dc4ba3d2a9a6d2714d01aba44cc2fe8cbea4bf26a72d36).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Manupa Karunaratne via Apache TVM Discuss <no...@discuss.tvm.ai>.
On a side note to this conversation about new primitives, would the new TensorIR will include primitive "store_at" -- the one present in Halide/Tiramisu ? -- I just want to know if thats something in the roadmap :slight_smile: .
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/52) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/ce1043115f9ec89872fb5d93956816c2245245268c083c3b5510913d43eecfe7).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Junru Shao via Apache TVM Discuss <no...@discuss.tvm.ai>.
Thanks @srush! I totally agree!
Simplifying the scheduling language is definitely an important issue, especially when we want to approach broader audience.
An idea to make it happen is to compose schedule "primitives" into more high-level "composite rules". For example, making "tiling + cache_read + compute_at" a single rule that user can directly call - it works like a single schedule instruction, but will be decomposed into many small "primitives" inside our system.
Let's take a step forward, if we push the idea to the extreme, i.e. our rule could work with no human interaction at all, then we can actually build an automatic scheduling framework on top of it :-)
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/51) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/07a70667a23f4d5a8cc468c4358377055cdfd754b816ae4d6778eba7156b27c8).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Sasha Rush via Apache TVM Discuss <no...@discuss.tvm.ai>.
Hi @Hzfengsy, Thanks for your answer, that makes a lot of sense.
However, what I am trying to communicate is that even with a new TIR and the ability to print it out at each step, the underlying scheduling language seems to be roughly the same in all of these examples (with the nice exception of removing the unnecessary `s[C]`). This is okay in all the examples in the thread that are very small.
What I am suggesting is that this might be a perfect time to think about simplifying the scheduling language as well. For instance the example here: https://tvm.d2l.ai/chapter_gpu_schedules/matmul.html#blocked-matrix-multiplication-on-gpu is in some sense the simplest scheduling example of TVM, and it is still extremely complicated. It has roughly 20 local variables in scope all with one letter names and many overlapping mutations (without even using autotvm).
I am curious about any ideas the new approach has to make this more robust, easier to extend, and understandable to non-experts?
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/50) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/cfa418fa84383166a0b9981a44ba3c1d8cf6a19879e34b4802d8fe6fbf590a19).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Siyuan Feng via Apache TVM Discuss <no...@discuss.tvm.ai>.
Thank you for such a valuable question.
Your understanding is correct. We still need a schedule language to schedule. That is because we need a simple API and abstraction for both human experts and automatical optimization (like AutoTVM, Ansor, and our new meta-schedule). Also, we try to keep user habbit, so we do not change API too much.
The critical challenge you are mentioned seems user experience especially loop axes. TensorIR is an eager schedule, which means every schedule primitive will change IR as soon as it exuecutes. Then, user can see axes and whole AST whenever they want, here is a simple example:
```Python
import tvm
from tvm import te, tir
A = te.placeholder((128, 128), name="A")
Update = te.compute((128, 128), lambda i, j: A[i, j] + 1, name="update")
"""Create PrimFunc from TE compute for further scheduling"""
func = te.create_func(Update)
"""Create TensorIR schedule"""
s = tir.create_schedule(func)
print(tvm.script.asscript(func))
"""Output
@tvm.script.tir
def func(var_A: ty.handle, var_update: ty.handle) -> None:
A = tir.match_buffer(var_A, [128, 128], elem_offset=0, align=128, offset_factor=1)
update = tir.match_buffer(var_update, [128, 128], elem_offset=0, align=128, offset_factor=1)
# body
with tir.block([], "root") as []:
tir.reads([])
tir.writes([])
for i0, i1 in tir.grid(128, 128):
with tir.block([128, 128], "update") as [i, j]:
tir.bind(i, i0)
tir.bind(j, i1)
tir.reads([A[i:(i + 1), j:(j + 1)]])
tir.writes([update[i:(i + 1), j:(j + 1)]])
update[i, j] = (A[i, j] + tir.float32(1))
"""
update = s.get_block("update")
x, y = s.get_axes(update)
print(x)
"""Output
for i0 = 0 to 128
"""
xo, xi = s.split(x, factor=32)
print(xo, xi, sep="\n")
"""Output
for i0_outer = 0 to 4
for i0_inner = 0 to 32
"""
print(x)
"""Output
(nullptr)
"""
```
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/49) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/e3b447f21e04a989a76c79504f7d65d0d000cba8a15cde190cd933c6b3987bc0).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Sasha Rush via Apache TVM Discuss <no...@discuss.tvm.ai>.
Yeah, I think your explanation is a good summary. I see what you mean about the TensorIR blocks
My understanding though is that the user doesn't actually write TensorIR (except maybe to start), they still schedule with a separate language? The blocks in TIR seem really nice, but I still worry that the scheduling code itself also needs some ability to abstract. For instance the example here, https://tvm.d2l.ai/chapter_gpu_schedules/matmul.html#blocked-matrix-multiplication-on-gpu . It doesn't seem like this changes that too much? There are so many axes in scope in this function at once, and it seems very hard to separate them all from each other.
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/48) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/c62aca6fd42a8ac1d6a0ef90d44c1ea2dcb3324f4501e7bf60ccc08a1f5234c2).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Manupa Karunaratne via Apache TVM Discuss <no...@discuss.tvm.ai>.
Thanks for the clarification! I concur that such a primitive should be useful and would allow more flexible compute movements.
Regarding the full graph, I agree that relay (along with optimization) being very useful. I was thinking whether there would be a benefit of lowering the full graph to tensorIR post relay optimization rather than lowering each primitive function. I guess this has to do with how AutoTVM/Ansor will allow the exploration of schedules but I got a feeling that could be scoped via the "blocks" that would otherwise lead to explosion of search space.
(Looking from an AoT angle here).
Moreover, may be that could lay a foundation to inter-primitive function optimizations later.
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/28) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/4a1848f587b78097f506c6c9f44163b53f14ed80842bc37f4825a6e29c2069b8).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by tqchen via Apache TVM Discuss <no...@discuss.tvm.ai>.
for concat, we could introduce a reverse inlining primitive that inlines elemenwise operations(after concat) back to the concat, which should be helpful in many cases.
While it is possible to represent a full graph, we would still imagine relay being super useful as a coarse grained repr for graph level opt. So that would suggest to have a continued effort on multi-level repr(relay and tir)
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/27) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/498c5892f36ef38a0efccec2d9ee5f428a7b1df1714058bfd6c7bc4b14544a90).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Manupa Karunaratne via Apache TVM Discuss <no...@discuss.tvm.ai>.
Thanks for the proposal! Looks quite interesting!
Out of curiosity,
1) The concat example you've shown where the original stage is represented in three blocks that seems to be assigning to the same buffer. I'm curious to know what if we want to move the concat (using compute_at, if possible ?) to a consumer of the concat's output (to some loop of the consumer), how could it be done ? Will it create multiple blocks there as well ?
2) Since the proposed TensorIR enables scoping of scheduling transformations in terms of blocks, will there be a prospect of representing a full relay graph in TensorIR ?
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/26) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/1b39f049a7d4c003332d3f27314e92dda432fb20f52341f91881ac343cbfecdf).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Sasha Rush via Apache TVM Discuss <no...@discuss.tvm.ai>.
Neat! This is a really good explanation. I think I get most everything you are explaining. (I'm still stuck on `reverse_compute_at` which seems like a long name, and is still a bit too magical for me to understand.)
In terms of splitting / reordering, my goofy thought is that my favorite construct for tensor abstractions is `vmap` in Jax. What makes programming tensors so hard is that keeping 6 tensor dimensions in your head is really hard. vmap lets you do something then "zoom" into the area, forget entirely about the outer dimension, and focus on that.
When writing tvm code for matrix multiply with double buffering. I would really like to 1) first decide on my tiling of the output, split, assign to blocks and threads, and then 2) write a separately scoped bit of code that doesn't even know about the outer construction at all. Ideally, I would create my outer scope, vmap in, then all my buffers are automatically "reverse_compute_at" / vmapped, and then create my inner setup.
I don't know if this totally works but this would be my ideal:
```python
l_o = split_out(C, l) # only exposes the outer
n_o = split_out(C, n)
with prefix(ll_o, nn_o, tensors=[A, B], threads=[]) as s2:
s2.cache_read(... ) # this cache read is now local computed at here
l, n = s2.axes(C) # these are now the inner splitted axes
m = s2.axes(A) # A's outer axes are now invisible
s2.reorder(n, l) # only touches the visible axes.
```
(Maybe instead of a `with` this is an inner function like in jax)
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/45) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/05e4fd027d9c9c1a9ca4be05a983d0df9f859ea1f43bfac0a64be8f03e0bf98f).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Junru Shao via Apache TVM Discuss <no...@discuss.tvm.ai>.
Hey @srush,
Thanks for your valuable feedback! Please allow me to try to explain the design rationale below:
> **Q1.** There are strings for get_block?
Yeah. One design principle we hold for TensorIR is that all needed for scheduling is contained in the TensorIR python syntax, so that there is no "mysteriously hidden" information any more. Given a TensorIR script in python, we can schedule on it.
In the particular case, the string here is the name of the block in the text format.
Here is the example:
```python
@tvm.script.tir
def matmul(x: ty.handle, y: ty.handle, z: ty.handle) -> None:
X = tir.match_buffer(x, [128, 128], "float32")
Y = tir.match_buffer(y, [128, 128], "float32")
Z = tir.match_buffer(z, [128, 128], "float32")
# ⬇️ name of the block is "Z"
with tir.block([128, 128, tir.reduce_axis(0, 128)], "Z") as [i, j, k]:
with tir.init():
Z[i, j] = tir.float32(0)
Z[i, j] = Z[i, j] + (X[i, k] * Y[k, j])
"""Print TensorIR in python syntax"""
print(tvm.script.asscript(matmul))
"""
Create a schedule
Scehdule is TensorIR => TensorIR
so we can print the latest TensorIR at any step of scheduling
"""
sch = tvm.tir.create_schedule(matmul)
ax0, ax1 = sch.split(...)
"""Print TensorIR at any step"""
print(tvm.script.asscript(sch.func))
sch.fuse(...)
print(tvm.script.asscript(sch.func))
"""
We can print the loops and blocks too into syntax like:
for i in range(0, 100)
"""
print(ax0)
```
> **Q2.** Why not have split go into a named out/in tuple to discourage this _ style naming. It gets so messy so quickly
Agreed. We do find the _ style naming annoying in our experiments, especially when scheduling gets complicated and it generates really horrible names like `i0_outer_outer_outer_outer_i1_outer_outer_outer_outer_fused_i2_outer_outer_outer_outer_fused_i3_outer_outer_outer_outer_fused`.
We have some internal strawman proposals for syntactic sugars, but converged to a perfect solution yet.
**Solution 1.** Allow splitting by multiple factors + accept customized naming of axes.
```python
a, b, c, d = s.split(axes, factors=[None, 2, 4, 8], names=["a", "b", "c", "d"])
```
Note that the `None` in the factors means letting the schedule to infer.
**Solution 2.** Allow splitting by multiple factors + einsum-like style API + get axes by name
```python
i0, _, _, _ = s.split("i => i0, i1, i2, i3", factors=[None, 2, 4, 8])
# or allow retrieve by name
i0 = s.get_axis(name="i0")
```
We are open to new proposals too :slight_smile:
> **Q3.** Does this propose to fix the issue of having to repeat the identical splits for things like shared and local buffers that need to be done later in the code. (In order for compute at to work)
In short, yes, and it is solved by introducing a new scheduling primitive `reverse_compute_at`.
**Definition of `compute_at`.** Given a producer and a consumer, `compute_at` allows to compute part of the producer's region under one of the consumer's loop.
**Definition of `reverse_compute_at`.** Given a producer and a consumer, `reverse_compute_at` allows to compute part of the consumer's region under one of the producer's loop.
**Our typical usecase.** We have a heavy producer, like conv2d, and a light consumer, like what is generated by `cache_write`, or a small ReLU, and we want to fuse them for better locality.
**Why bother duplicated splitting in `compute_at`.** With `compute_at`, we are moving the producer under a loop of the consumer. First, user has to split the consumer, otherwise the user doesn't even know which axis to be computed at; Second, user has to split the producer so that the other tiles are correctly positions - that is why it is so lengthy and tedious.
**Why `reverse_compute_at` avoids duplicated splitting.** In this case, we only need to split the producer, then put the consumer under a specific loop of the producer. Then we don't have to do any duplicate splitting :-)
CC: @tqchen @Hzfengsy @spectrometerHBH @vinx13 @masahi
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/44) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/2462cba174a272d5f70b53b64c084f41a5043bc4f07a0ffce9ee23ede404aa27).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Siyuan Feng via Apache TVM Discuss <no...@discuss.tvm.ai>.
Thank you for your interest.
A1: Current op fusing is based on `stage` but the critical point is fusing the injective computation. We can also inline injective computation by `traverse_inline`. So there is no doubt that FuseOps works. As for the philosophy, I think there are only few changes. TIR is not only an IR but only can be a computation declaration. We provide very user-friendly API (as easy as TE) to define compute rules.
A2: TIR is a general IR which can represent almost every program. It's really hard to schedule a general program, but we promise TIR can schedule all programs which TE can schedule.
A3: Most of the primitives are similar to the TE ones. For now, only two new primitives are `decompose_reduction` and `merge_reduction`
A4: Upstream on Oct and Nov. Ansor supporting is WIP
I hope I can answer your question.
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/8) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/04866ec0b8abc98692fbe14e00cb98f01a3ac3fd00abee2ab8289347a07afd28).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by aca88 via Apache TVM Discuss <no...@discuss.tvm.ai>.
Hi,
Even though I don't think I understood everything, I like the idea of solving some of the limitations of `te.compute`. Since the `te.compute` is in a central part of the TVM stack changing it requires a lot of work and understanding. So thank you all for continuing such development.
Q1: I was wondering how this fits in the Relay ->Topi->TE->TIR flow. In a more specific case take FuseOps. AFAIK the FuseOps pass creates the multi-stage operator based on the `te.compute`s. Since you mentioned that there would be no notion of a `stage` in the new TensorIR, how would FuseOps work? and more generally how TVM's philosophy of "defining a compute rule and a separate schedule" be changed?
[quote="Hzfengsy, post:1, topic:7872"]
TE has limited expressiveness since each stage is defined by `Stage = te.compute(lambda expr)` , while TensorIR is a full c+±like IR. We can write any program with TensorIR as you want. Although not all programs can be scheduled, there are still more workloads that can be optimized by TensorIR.
[/quote]
Q2: What exactly do you mean by "not all program can be scheduled"? maybe an example?
Q3: You mentioned "new scheduling primitives", could you maybe give a list?
EDIT:
Q4: Expected timeline of the steps?
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/6) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/4c0ef8be45bfe6cfef23f1dfe5e0407dbf9ced10d399080ab68a174431c19b49).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by "Cody H. Yu via Apache TVM Discuss" <no...@discuss.tvm.ai>.
Thanks for clarification. Make sense to me.
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/19) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/e1ab54264948a038da4df1022a21cfcf20ce788c2cbdb53e65b3cd437630319f).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by tqchen via Apache TVM Discuss <no...@discuss.tvm.ai>.
Because there is a 1-1 mapping between te.Stage and Block. It should actually not be hard to use tir schedule to schedule a te compute generated PrimFunc (either by getting block via name, or pragmatically traverse the blocks like we do pragmatically on stages). But i agree that we can keep te.schedule for a bit.
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/18) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/29fb2b8264094ac7249ced2921771e02e25948f5ee4d2e44383207f5aee76141).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by "Cody H. Yu via Apache TVM Discuss" <no...@discuss.tvm.ai>.
So the scenario is like you can choose to use TE or TIR to write a compute, but if you choose TE, you have to first lower it to TIR and then add schedule primitives?
IIUC, it seems to me that this is nontrivial, because TIR was not written by human and you may need to first print it out to figure out how to schedule it. It sounds more straightforward to keep TE schedule as syntactic sugar. At least you can get the sense about how the schedule looks like by tracing the Python code.
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/17) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/30b15ab093204382d89b4c5b1433f51527a645e6e74ab1bd61e2e2ac5c5f673a).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Yao Wang via Apache TVM Discuss <no...@discuss.tvm.ai>.
Would love to see dynamic shape supported otherwise a large set of models can't be backed by new TensorIR. :D
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/16) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/c2aa7dfadcc7a0b4ba60e463d1fddbcc779e4eef36d8389ccd905230a1cc4458).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Siyuan Feng via Apache TVM Discuss <no...@discuss.tvm.ai>.
Good questions!
1. As for as we know, we would like to let users use TensorIR schedule rather than TE schedule one we fully upstream the TensorIR. For three reasons:
1. Just as you have mentioned, TE is a fronted wrapper, and it directly generates TIR with blocks. Somehow, TE is more like a sugar to define TIR.
1. Most of the schedules and primitives in TensorIR are very similar to those in TE. The cost of learning TensorIR schedule is extremely low (maybe just one day).
2. All primitives are based on the block (no stage concept in TensorIR schedule). It's hard to keep the TE schedule with block
2. Dynamic shapes are not supported now. However, thanks to our de-coupled primitives, it's ok to support later.
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/15) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/999ff8136307130669c2195d9ca73065cd97e1b6b221aa1029a88630487e15e9).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by "Cody H. Yu via Apache TVM Discuss" <no...@discuss.tvm.ai>.
Thanks for the proposal! This definitely opens more opportunities for performance optimization. Two questions for clarification:
1. IIUC, based on the proposal and discussion, we will have both TE and TIR, but TE is more like a frontend wrapper of TIR to serve some users that prefer to write high-level DSL. Then, what will we do with the TE schedule primitives? Intuitively, we should still keep them; otherwise TE writers will have no way to schedule their computes, because they know nothing about TIR and blocks.
2. Does this proposal support dynamic shape (i.e., `Any`)? For example, can we have something like:
```python
@tvm.hybrid.script
def matmul(a: ty.handle, b: ty.handle, c: ty.handle) -> None:
C = tir.match_buffer(c, (1024, 1024), "float32")
A = tir.match_buffer(a, (1024, Any), "float32")
B = tir.match_buffer(b, (Any, 1024), "float32")
reducer = tir.comm_reducer(lambda x, y: x + y, tir.float32(0))
with tir.block([1024, 1024, tir.reduce_axis(0, 1024)], "C") as [vi, vj, vk]:
reducer.step(C[vi, vj], A[vi, vk] * B[vk, vj])
s = tir.create_schedule(matmul)
update = s.get_block("C")
i, j, k = s.get_axes(update)
i_o, i_i = s.split(i, bn)
j_o, j_i = s.split(j, bn)
k_o, k_i = s.split(k, 4)
```
In this case, the length of `vk` (or `k`) is `Any`. Can we still apply `split` to it with a fixed factor
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/12) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/6a5d6e8ca987262e53cb39c47d1a37292fa07c774c177061cc3f1cbfc6d87d01).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Xqdan via Apache TVM Discuss <no...@discuss.tvm.ai>.
One issue in old schedule ops is we can not get the accurate bouds with inferbound, what will it be like in new schedule system? thanks.
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/64) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/a665d101ee258a31107279867294b7a7dc7427c420169cc6f1fed9f459564cc6).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Zihao Ye via Apache TVM Discuss <no...@discuss.tvm.ai>.
[quote="Hzfengsy, post:1, topic:7872"]
TensorIR natively supports hierarchy checks. We will check the memory access and thread binding, including warp level instruction(wmma) validation during the schedule. Following is an example of the GPU hierarchy.
[/quote]
@Hzfengsy What do we mean by "check" here? Is it something like a boundary check? What if tensorized operators provided by vendor have some alignment requirements (e.g. the start pointer must be multiple of 8/16).
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/54) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/07452c616584e30afb169f93687d75eaafee6a261f0194aa342989934a40b82b).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Junru Shao via Apache TVM Discuss <no...@discuss.tvm.ai>.
> I’m still stuck on reverse_compute_at which seems like a long name, and is still a bit too magical for me to understand
Yeah I once had some discussion with @tqchen and @spectrometerHBH that reverse_compute_at is too long and tedious, and I agree it will be great to find some better names (we are always bad at it admittedly)...
Perhaps @spectrometerHBH could give a specific example of using `reverse_compute_at` to avoid duplicate splitting?
> vmap lets you do something then “zoom” into the area, forget entirely about the outer dimension, and focus on that
Yeah `vmap` is a powerful tool - it sounds more like "batching" (either static or dynamic), but its functionality can certainly go beyond that by being applied multiple times.
One problem of `vmap` if we want to introduce to TensorIR is the semantics: do we consider it as adding a data parallel loop on top of the producer block of each output buffer? I believe @tqchen has more thoughts on this.
> write a separately scoped bit of code that doesn’t even know about the outer construction at all
The idea sounds pretty much like our "block isolation" construct. A "block" in the TensorIR means doing isolated computation without having to know the outer scope.
The difference is that the `with` scope you mentioned does not effectively change the IR, hence it is more like syntactic sugar to hint the schedule class that some primitives should operative locally, while the `Block` in TensorIR is an IR construct that really enforces some restrictions.
CC @Hzfengsy @spectrometerHBH @tqchen would love to hear your opinions
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/47) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/448533130f112297da3575711cd06f390c7c5424baca65748389314641302c96).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Daniel Mu via Apache TVM Discuss <no...@discuss.tvm.ai>.
Well-received with thanks!
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/7) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/08ac0a4be5a8f19be521f57ad137d22e814a0fd16ba753bd90180deb0631206e).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Bohan Hou via Apache TVM Discuss <no...@discuss.tvm.ai>.
[quote="ds1231h, post:3, topic:7872"]
However, will this increase the coupling between the schedule and the lower pass, which may lead to an increase in the complexity of the lower pass?
[/quote]
Thanks for your reply! @ds1231h
At the moment, we at first transform TIR with block to TIR without block in one pass(the latter is current TIR in the main repo) and utilizes the current TIR lowering passes to do codegen. And we can gradually refactor this lowering process in the future.
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/5) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/38cd522ac0e3f1276f3367910d52587886b0c138b37e7f9c9da82528635125c7).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Daniel Mu via Apache TVM Discuss <no...@discuss.tvm.ai>.
Great work! I believe this will make "scheduling" more flexible and intuitive!
However, will this increase the coupling between the schedule and the lower pass, which may lead to an increase in the complexity of the lower pass?
By the way, I'm also looking forward to know how to auto-schedule based on the the TensorIR.
Thanks!
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/3) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/5925617787fa2ee2d343a58ba457d1aa6839fc66402c93e3b0c7ef63c225737c).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Xqdan via Apache TVM Discuss <no...@discuss.tvm.ai>.
Is Fusion in Ansor based on tir?
For other transforms, you may checkout here, that's what we've done in AKG. I can explain some if you are intrested.
https://github.com/mindspore-ai/akg/blob/master/src/codegen/build_module.cc#L439
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/31) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/a7f0d2bafba438aef187410bc3f676663b7b15309b2644f747d34a10d3bc45bd).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Xqdan via Apache TVM Discuss <no...@discuss.tvm.ai>.
@junrushao1994 It's better to know loops can be vectoried, permutable or distributied, isl can provide these information，so we can do loop optimization and tensorization/vectorization automatically.
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/32) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/9935959d85972017de17516f48d2c09e3a5b07c0857a9cdcdd3306e512945c9f).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Junru Shao via Apache TVM Discuss <no...@discuss.tvm.ai>.
@xqdan Thank you for the valuable feedback! Fusion can be done automatically with some analysis provided in Ansor.
Do you have any other kind of analysis in mind that might be potentially useful?
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/30) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/eaf9e71884e702c672d7ce255eee7c8e9a0f4d9157e54c5ba2539be5852022ad).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Zihao Ye via Apache TVM Discuss <no...@discuss.tvm.ai>.
I think the introduction about "Block Realize" is missing here.
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/55) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/5c8b472157586207459d5e7d8dd828c9798b9c2cdce0c8c3636c2d1696c32972).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Lianmin Zheng via Apache TVM Discuss <no...@discuss.tvm.ai>.
How is the compilation speed compared to the original TE?
In Ansor/Autotvm, we have to compile a lot of schedules for feature extraction, so the speed of schedule transformation matters.
Do you have any benchmark results? Intuitively, I think the original TE will be faster because it can do a batched bound inference and AST construction. If it is true, how can we fix this performance gap?
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/34) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/d44828c61231fffb080d241914c92520cd26123d675c2c0a96aac51d5d97bab2).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Sasha Rush via Apache TVM Discuss <no...@discuss.tvm.ai>.
Hi TVM,
This idea is so cool, I think it is going to make it possible for mortals to use TVM effectively.
I have a couple of questions about the snippet of the scheduling language.
The three issues I have when programming TVM are:
* Too many variables in scope with meaningless names, and forgetting where they can from.
* Losing track of which axes need to be split identically
* Not understanding the semantics of compute_at and how it tells which axes line up.
It seems like maybe this fixes a couple of these. However the following still bugs me a bit:
```
s = tir.create_schedule(matmul)
update = s.get_block("C")
i, j, k = s.get_axes(update)
i_o, i_i = s.split(i, bn)
j_o, j_i = s.split(j, bn)
k_o, k_i = s.split(k, 4)
s.reorder(i_o, j_o, k_o, k_i, i_i, j_i)
```
Curious why:
- There are strings for get_block?
- Why not have split go into a named out/in tuple to discourage this _ style naming. It gets so messy so quickly.
- does this propose to fix the issue of having to repeat the identical splits for things like shared and local buffers that need to be done later in the code. (In order for compute at to work)
Thanks so much! Love the library, want to get to a point where I can teach it to my students effectively.
/Sasha
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/43) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/d870aadc9f3242a2985168611f0f69f57f3c6578e3e89818abd6916cd4095775).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Chenfan via Apache TVM Discuss <no...@discuss.tvm.ai>.
E... @junrushao1994 I guess @merrymercy 's opinion is that doing analysis in TE is quicker than using the ISL.
ISL is sure a powerful tool for loop analyse, but in my understanding we should lower the schedule to C code first before using ISL? Which I think is more time consuming.
Currently, Ansor applies some simple but useful analyses based on TE. Though it may not be as accurate as ISL does, but it's cheap. Then we count on the tuning to try lots of uncertain schedules and find the best one by measuring.
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/36) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/ed7bad22d57cd3d30d99cfba8be5dd21289c41a85ad4e40e8ab21b989b3c9b4f).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Junru Shao via Apache TVM Discuss <no...@discuss.tvm.ai>.
@merrymercy I didn't get it about batched bound inference, doesn't Ansor use a pool of threads for massive bound inference?
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/35) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/125fc44665f3b014e07161454a28ad60c0dd07ce2dab0c12233aa9d79b4fb79c).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Bohan Hou via Apache TVM Discuss <no...@discuss.tvm.ai>.
Thanks for your reply! @jcf94
A1. We've tried to tensorize intrinsic using this new IR, and are working on the TensorCore demo. Our design is really close to the original tensorize programming logic, only differs in the declaration of description&implementation of HW intrinsic (we can use Hybrid Script to program them more conveniently)
A2. There is no concept of Stage in TensorIR. What you can see in Hybrid Script contains full IR information and our schedule transforms IR directly. The basic schedule unit is Block.
A3. Ansor + TensorIR is on our roadmap.
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/4) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/1742456eab27a02e24c4f5609000bf6a04da32e296d7aadbb5df1cf6779b8c54).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Chenfan via Apache TVM Discuss <no...@discuss.tvm.ai>.
Great work!
Have you tried tensorize intrinsic(e.g. TensorCore schedule) using this new IR? Since I remember that to support tensorize is also one of your initial motivations.
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/2) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/42302489d4b03ae5fb18b0022f87960b01432353c1bd08d016d27041060280a2).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Xqdan via Apache TVM Discuss <no...@discuss.tvm.ai>.
This is the right way to go. However I have two concern,
1) How to fuse ops as much as possible? Basically fusion is copy propagation optimization in compilers, which is based on data flow analysis, but still lack of programming analysis in TVM now.
2) TE tensorize can not handle some complex pattern matching, see https://github.com/apache/incubator-tvm/pull/1053, can we do 100% pattern matching in tir?
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/29) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/78c2de28cf50a3f0e21bd234a9fd975d7fd77c870c4627104dab67469571f219).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by masahi via Apache TVM Discuss <no...@discuss.tvm.ai>.
@junrushao1994 If by "meta programming" you mean an ability to call python function from script to generate other code, and embed the generated code into the calling context, then YES, we absolutely need this! I think it's called "splicing" or "unquote" in the literature.
The lack of such feature is what turned me away from hybrid script. Since there is basically no abstraction to compose multiple code to generate bigger code, I had to manually duplicate a lot of code. The resulting code is super ugly (just look at various shape functions). For example, https://github.com/apache/tvm/blob/8e23806d2d522b71979d0a2730b38cc5c3bf6185/python/tvm/relay/op/_transform.py#L1000-L1008 (this one I wrote)
For examples of kind of programming style we've developed for ir builder, see `topi/cuda/sort.py` or `topi/cuda/nms.py`. cc @mbrookhart
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/60) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/dd7c78090fae7b2639c16f042e9edd905248fe7ef098a74acb9eff11069b39e5).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Junru Shao via Apache TVM Discuss <no...@discuss.tvm.ai>.
Yeah, to fully replace IRBuilder, there are still some missing elements of TVM script: meta programming and hygiene macros. Let's consider the support after the upstreaming is done :-)
CC @tqchen
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/59) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/a4d097f489e59339a92e4b5699181813a7dd78c0e8aa34338ef0abed716e95d5).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Yao Wang via Apache TVM Discuss <no...@discuss.tvm.ai>.
Thanks for explanation. The relation between te and new tir is now more clear to me.
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/13) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/55dbd5407fa8c8f3a50c45fe5cb8b5c51bb644207975358d07945b5bde75be37).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Bohan Hou via Apache TVM Discuss <no...@discuss.tvm.ai>.
Thanks for your reply! @kevinthesun
[quote="kevinthesun, post:9, topic:7872"]
Thank you for this proposal! This work does make scheduling much easier. I have a concern about using this way to write a tensor expression. It looks like more complicated than tvm.compute when defining matmul. We need to define some buffers and creating block with corresponding shape dimension. It would be helpful if you can add a conv2d example which can replace existing topi.nn.conv2d definition to better understand what developer would need to write.
[/quote]
In original te programming, we also have to declare buffers and create lambda expression with iter_vars having the correct shape dimension. If we take this additional info into account, TE programming is close to Hybrid Script's programming in complexity.
Currently, we can not replace existing topi operators, since they are represented by Stage/Op and optimized by te schedule, while Hybrid Script will be parsed into TIR directly.
[quote="kevinthesun, post:9, topic:7872"]
Another question is about representing generic programming style ops such as shape functions. Since these programs don’t fit into tvm scheduling, I assume it would still be more convenient to use existing te hybrid script to create these ops?
[/quote]
If we don't have to schedule the PrimFunc, we don't have to declare blocks in TIR.
Actually, TE hybrid script is also a text representation of TIR to a large extent, with loop and condition statement directly representing the IR structure, using sugar like variable being translated to Array of size 1 to ease the format of Store&Load. At the moment, we haven't introduced sugars to simplify such Load&Store, but the rest writing are largely simlilar.
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/10) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/b56ad45dd9069fd09fb5ae1b434a3d4d598842a38b39e2b2aea236c55b1ccad3).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Yao Wang via Apache TVM Discuss <no...@discuss.tvm.ai>.
Thanks for clarification. It would be nice if we can use various methods to create tensor programs and use new tir to schedule them.
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/14) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/1a79719c8acbecadd7e9886a2e71139728fc881794bd3fc6c2b1b1182429d91b).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by tqchen via Apache TVM Discuss <no...@discuss.tvm.ai>.
TIR and TE does not conflict with each other. TE is still a useful DSL to stitch fragments of TIR together to form a PrimFunc.
We could still define TE based DSL(backed by TIR) that enables primitives like compute and hybrid calls to stitch together a dataflow graph to form a PrimFunc And then use the TIR for scheduling.
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/11) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/bcfe5c7a9a120a0b4b355ab0cef2b57e3c901087fcb0a186af2c149618e2b53a).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Yao Wang via Apache TVM Discuss <no...@discuss.tvm.ai>.
Thank you for this proposal! This work does make scheduling much easier. I have a concern about using this way to write a tensor expression. It looks like more complicated than tvm.compute when defining matmul. We need to define some buffers and creating block with corresponding shape dimension. It would be helpful if you can add a conv2d example which can replace existing topi.nn.conv2d definition to better understand what developer would need to write.
Another question is about representing generic programming style ops such as shape functions. Since these programs don't fit into tvm scheduling, I assume it would still be more convenient to use existing te hybrid script to create these ops?
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/9) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/368fa3a2cf5e4531886a799e6f4ed9355e153e65b0c9853c9a0d477a7b7dbe8e).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Bohan Hou via Apache TVM Discuss <no...@discuss.tvm.ai>.
[quote="merrymercy, post:37, topic:7872"]
I mean the original TE is a declarative language so it can know all transformation before it starts to generate low-level AST. But the new schedule primitives are done imperatively. In the original TE, we can share some analysis results (e.g. dependency analysis), so it is expected to be faster.
[/quote]
@merrymercy Good question! Here's an example of TIR's schedule.
```python
s = tir.create_schedule(original_func)
update = s.get_block("C")
i, j, k = s.get_axes(update)
i_o, i_i = s.split(i, bn)
j_o, j_i = s.split(j, bn)
k_o, k_i = s.split(k, 4)
s.reorder(i_o, j_o, k_o, k_i, i_i, j_i)
```
TIR's schedule is not totally stateless. Scope info, dependency graph info is actively maintained during the scheduling process in class Schedule. We don't calculate them each time we apply a new primitive. After lowering to TIR without blocks, we don't maintain these info any more since it is not schedulable.
All in all, it is good to run the benchmark to compare them in practice. I hope I understand your question correctly. :smile:
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/38) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/010365ba77175ff89f5b27d2ef08cfa25d6925e8d1de22308415b0f0a328715e).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Lianmin Zheng via Apache TVM Discuss <no...@discuss.tvm.ai>.
@jcf94 @junrushao1994 Sorry, both of you don't understand my question correctly.
I mean the original TE is a declarative language so it can know all transformation before it starts to generate low-level AST. But the new schedule primitives are done imperatively. In the original TE, we can share some analysis results (e.g. dependency analysis), so it is expected to be faster.
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/37) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/0253ff039aa350b4b524bc1200e2b1a0b902940762ecf5975720adafafe655da).
[Apache TVM Discuss] [Development/RFC] [RFC] TensorIR: A schedulable
IR for TVM
Posted by Siyuan Feng via Apache TVM Discuss <no...@discuss.tvm.ai>.
Thanks for such a great suggestion. Yes, we do support IRBuilder for TensorIR. However, it is not recommended. Because it is likely to generate illegal or opaque IR (which lacks some of the information). Besides, there are so many attributes/annotations (e.g block read/write regions and block iter_var) to be provided by users if you want to use native IRBuilder.
On the other hand, TVMScript can also represent ANY IR supported in TensorIR and TVM as far as I know. It provides extra syntax checks and sugar to make it easy to write a schedulable IR. So we strongly recommend you try TVMScript if possible. :)
---
[Visit Topic](https://discuss.tvm.apache.org/t/rfc-tensorir-a-schedulable-ir-for-tvm/7872/58) to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click here](https://discuss.tvm.apache.org/email/unsubscribe/6ae37e24a7579d941bad4814417354e37728bd52cdc79ab11f2fb52c12bda709).