Lunderberg commented on code in PR #70: URL: https://github.com/apache/tvm-rfcs/pull/70#discussion_r876343679
########## rfcs/0070-introducing-decl-buffer.md: ########## @@ -0,0 +1,210 @@ +- Feature Name: introducing-decl-buffer +- Author: Wuwei Lin (@vinx13), Eric Lunderberg (@Lunderberg) +- Start Date: 2022-05-04 +- RFC PR: [apache/tvm-rfcs#0000](https://github.com/apache/tvm-rfcs/pull/70) +- GitHub Issue: TBD + +# Summary +[summary]: #summary + +This is a follow-up of https://github.com/apache/tvm/pull/9727 and +[RFC#63](https://github.com/apache/tvm-rfcs/pull/63). Currently buffer can be implicitly +declared and then used. The implicit behavior can be error prone and makes analysis more difficult. +This RFC introduces `DeclBuffer`, a new IR construct as an explicit statement for buffer declaration. + +# Motivation +[motivation]: #motivation + +Currently a Buffer object can be created and then referenced in TIR, without explicit declaration +or allocation. For example, in TVM script, one can use `T.buffer_decl` to create a new buffer and +then use it in the rest of the program. +``` [email protected]_func +def buffer_alias(A: T.Buffer[(16,), "float"]): + A_vector = T.buffer_decl([4], "float32x4", data=A.data) + T.evaluate(A_vector[0]) # read from buffer alias +``` +However, `T.buffer_decl` doesn’t translate to a node in AST. The AST will be +``` +PrimFunc { + buffer_map: {A: Buffer[(16,), "float"}, + body: Evaluate { + BufferLoad { + buffer: Buffer(data = A.data, [4], "float32x4") # implicit creation of new buffer + index: [0] + } + } +} +``` +In this example, `BufferLoad` loads from an implicitly-created new buffer which aliases another +buffer. This example shows that a data variable can be used to create a buffer in arbitrary ways. +There are no guarantee that the created buffer and the underlying data variable have consistent +physical memory. This makes analysis in TIR difficult and error-prone as one should always check +whether a buffer in TIR is an implicitly-created one. + +By introducing explicit `DeclBuffer` statement, we can require that a buffer must always be declared +before any usages. This makes the creation and the usage of buffer better-managed within TIR. +Developers (e.g pass writers) can collect buffer information such as allocation, aliasing by +visiting `DeclBuffer` nodes. + +# Guide-level explanation +[guide-level-explanation]: #guide-level-explanation + +`DeclBuffer` will be defined as +``` +class DeclBuffer : public Stmt { + Buffer buffer; // the buffer declared + Stmt body; // the scope of the buffer +}; +``` + +In TVM script, `T.buffer_decl` will be renamed to `T.decl_buffer` to make the name a verb phase that +is consistent with the existing ones such as `T.alloc_buffer`, `T.match_buffer`. `T.decl_buffer` +will be translated to a `DeclBuffer` object in TIR. This only changes the way parser handles +`T.decl_buffer`, the user API of `T.decl_buffer` in TVM script will stay the same. + +In TIR, `DeclBuffer` will be handled in `StmtFunctor`. Visitors or mutators of `DeclBuffer` can be +override to handle `DeclBuffer` in TIR passes. + +# Reference-level explanation +[reference-level-explanation]: #reference-level-explanation + +## Allocation of intermediate buffer +The intermediate buffer inside `PrimFunc` can be declared and allocated in the following way: + +``` +Allocate { + data: A_data(Var(name=...)) + extent: ... + body: DeclBuffer { + buffer: A(data=A_data, dtype=..., shape=...), Review Comment: > The question here is whether this should be allowed I think I'd lean toward removing the alternative form and requiring the definition of the `data` variable to be defined prior to use in a `DeclBuffer` node in the PrimFunc's TIR. If a user gets an error message that states `A is undefined`, they may point at the `DeclBuffer` node and say that of course it is defined, conflating the definition of the buffer and the definition of the buffer's backing allocation. The confusion is especially likely, because the default behavior of `tvm.tir.decl_buffer` is to name the `BufferNode` and the `VarNode` with the same name. By making it impossible to have a buffer without a defined `data` variable, we avoid potential confusion about why a buffer is partially undefined. > that all backing buffers would come from tir.allocate nodes or from function arguments They could also come from return values of functions. After `MakePackedAPI`, the backing buffers are the return value of `@tir.tvm_struct_get`. It could also be an entirely separate function call, such as `data: T.Ptr[T.int32] = T.call_extern("device_specific_malloc", 1024, dtype="handle")` -- 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]
