yelite commented on code in PR #79: URL: https://github.com/apache/tvm-rfcs/pull/79#discussion_r915141600
########## rfcs/0079-tvmscript-metaprogramming.md: ########## @@ -0,0 +1,398 @@ +- Feature Name: tvmscript-metaprogramming +- Start Date: 2022-06-16 +- RFC PR: [apache/tvm-rfcs#79](https://github.com/apache/tvm-rfcs/pull/79) +- GitHub Issue: [apache/tvm#0000](https://github.com/apache/tvm/issues/0000) +- Co-Authors: Yaxing Cai ([**@cyx-6**](https://github.com/cyx-6), main implementation), Lite Ye + ([**@yelite**](https://github.com/yelite)), Yong Wu + ([**@yongwww**](https://github.com/yongwww)), Yuchen Jin + ([**@YuchenJin**](https://github.com/YuchenJin)), Eric Lunderberg + ([**@Lunderberg**](https://github.com/Lunderberg)), Masahiro Masuda + ([**@masahi**](https://github.com/masahi)), Junru Shao + ([**@junrushao1994**](https://github.com/junrushao1994), main designer) + +# Summary +[summary]: #summary + +This RFC proposes a new TVMScript parser infrastructure, supporting extensive +metaprogramming and syntactic sugars. The new infrastructure is IR-agnostic, +treating TIR just as one of dialects. Additionally, the new infrastructure will +provide better tooling around Python ecosystem (pylint, mypy, etc.). + +# Motivation +[motivation]: #motivation + +**What is TVMScript**. +Check [Blitz Course to TensorIR](https://tvm.apache.org/docs/tutorial/tensor_ir_blitz_course.html) and +[TVMScript Unified Printer RFC](https://github.com/apache/tvm-rfcs/pull/74/files#diff-6965a40ad8df7618ae68e11c88f924542a506c74a931cc3011ae9f99989b5f51R20-R26) +for an introduction into TVMScript. + +**What is metaprogramming.** In the context of TVMScript, metaprogramming means +a programmable way to control IR generation. For example, in +https://github.com/apache/tvm/pull/11097, a metaprogramming feature was added +to the TVMScript parser, allows users to programmably control the shapes of the +input buffers of a `PrimFunc`. + +### Limitation of current design + +The current parser lacks capability on generic metaprogramming that allows user +to have more control on IR construction. This makes it challenging to support +operators like NMS (non-maximum suppression, which is crucial to object +detection model). There is an implementation of NMS at +[python/tvm/topi/cuda/nms.py#L367-L386](https://github.com/apache/tvm/blob/d0650bad66d0ff89a01347537021bc442a98c223/python/tvm/topi/cuda/nms.py#L367-L386). +The implementation of NMS-like operators requires rank-polymorphism and the +ability to interleave host program with TVMScript, which is difficult to be +implemented under the current design. + +TVMScript also needs reasonable support on Python tooling. Currently it doesn’t +play nicely with pylint and mypy. For example, +[test_meta_schedule_postproc_rewrite_tensorize.py](https://github.com/apache/tvm/blob/d0650bad66d0ff89a01347537021bc442a98c223/tests/python/unittest/test_meta_schedule_postproc_rewrite_tensorize.py) +has 100+ warnings from pylint within only 500 hundred lines of code. This +creates confusion to the user and leaves an impression that TVMScript isn’t a +mature product and not production-ready. Even though it’s something that can be +incrementally improved under the current design, we believe it’s easier to get +an ideal result if we have a design with the tooling support in mind. + +The current design also lacks of unified approach for different IRs. At +[https://github.com/tlc-pack/relax/tree/relax/python/tvm/script/relax](https://github.com/tlc-pack/relax/tree/relax/python/tvm/script/relax), +a mature implementation of TVMScript parser is maintained for Relax. But it’s +hard to extend if we want to support more IRs for TVM unity. + +To conclude, with this RFC, we want to: + +1. Add more metaprogramming features to TVMScript, making it easier for TVM + developers to write complicated operators. +2. Improve tooling and documentation of TVMScript, reducing the friction for an + average machine learning practitioner to use TVMScript. +3. Modularize and infrastructuralize the TVMScript parser, lowering the cost to + implement parser for new IR. + + +# Guide-level explanation +[guide-level-explanation]: #guide-level-explanation + +## Metaprogramming features to support + +### (F1) Template Metaprogramming + +Users should be able to use variables from outer scope in the TVMScript +function/class. The parsed result should be identical to function/class with +the variable replaced by its value. For instance, + +```python [email protected]_func +def matmul( + A: T.Buffer[(128, 128)], +) -> None: + ... + +def gen_matmul(n, m) -> None: + @T.prim_func + def f(A: T.Buffer[(n, m)]): + ... + return f + +f = matmul(n=128, m=128) # `f` should be identical to `matmul` +``` + +This is already partially supported by https://github.com/apache/tvm/pull/11097 +for using `PrimExpr` captured by outer function. With the new parser, we want +to support this feature in more places and with more variable types. + +### (F2) Rank-polymorphism + +Users should be able to write a single function to handle different ranks of +input buffers (different numbers of dimensions). For example, user should be +able to write a generic function to do broadcast add, + +```python +def broadcast_add(a, b, c): + @T.prim_func + def f( + A: T.BufferFrom(a), + B: T.BufferFrom(b), + C: T.BufferFrom(c), + ) -> None: + for i, i_a, i_b in T.some_broadcast_method(A.shape, B.shape): + with T.block(): + C[*i] = A[*i_a] + B[*i_b] + +broadcast_add( + a = Buffer((128, 1), "float32"), + b = Buffer((1, 128), "float32"), + c = Buffer((128, 128), "float32"), +) +``` + +### (F3) Sugar: TE Compute in TIR + +Users should be able to replace boilerplate code with a function call, which’s +expanded to large chunk of code during parsing. For example, we may want to use +TE’s compute-like syntax to replace nested loop, + +```python [email protected]_func +def te_compute_sugar( + A: T.Buffer[(128, 128)], + B: T.Buffer[(128, 128)], +) -> None: + ... + C = T.compute((128, 128), lambda i, j: A[i, j] + B[i, j]) + ... + +## expands to ====> + [email protected]_func +def te_compute_expanded( + A: T.Buffer[(128, 128)], + B: T.Buffer[(128, 128)], +) -> None: + ... + for i in range(128): + for j in range(128): + with T.block("..."): + C[i, j] = A[i, j] + B[i, j] + ... +``` + +### (F4) Interleave host program and TVMScript program to customize metaprogramming + +As an escape hatch from writing code to be parsed by the TVMScript +parser, users should be able to write imperative code to construct IR nodes Review Comment: It's actually more flexible than what you described. We will not place extra restriction on the signature of user-provided function. More precisely, the argument types can be anything because TVMScript is able to capture outter variables, and the restriction on returned type depends on the language rules of the target IR. For example, for TIR, user can write `A[i] = ...` for BufferStore with a PrimExpr rhs, then the user-provided function that's called in the rhs of BufferStore needs to return `PrimExpr`. I updated the RFC to better clarify this. -- This is an automated message from the Apache Git Service. To respond to the message, please log on to GitHub and use the URL above to go to the specific comment. To unsubscribe, e-mail: [email protected] For queries about this service, please contact Infrastructure at: [email protected]
