diff --git a/rfcs/0042-buffer-pointer.md b/rfcs/0042-buffer-pointer.md new file mode 100644 index 00000000..9088a35f --- /dev/null +++ b/rfcs/0042-buffer-pointer.md @@ -0,0 +1,118 @@ +- Feature Name: BufferPointer +- Start Date: 2021-10-28 +- RFC PR: [apache/tvm-rfcs#0042](https://github.com/apache/tvm-rfcs/pull/0042) +- GitHub PR: [PR#9391](https://github.com/apache/tvm/pull/9391) +- Related RFCs: [RFC#0039](https://github.com/apache/tvm-rfcs/pull/0039) + +# Summary +[summary]: #summary + +A location being access in a buffer is represented as a +`BufferPointer` object, which holds a reference to the buffer being +accessed, and an array of indices to specify the location. +`BufferLoad` and `BufferStore` objects each hold a `BufferPointer` +object to specify where they operate. + + + +# Motivation +[motivation]: #motivation + +`BufferLoad` and `BufferStore` both act on a location within a buffer, +and must know where read/write their respective values. However, many +transformations are unconcerned with whether an access is a read or a +write. (e.g. `tir.transform.LowerMatchBuffer`, which rewrites access +of a matched buffers to instead be direct access of the backing +buffer.) By having a `BufferPointer` object to represent a pointer +into a buffer's memory, these transformations can be done without +repeated code. + +This is intended to make the layout transformations specified in +[RFC#0039](https://github.com/apache/tvm-rfcs/pull/0039) be more +straightforward to implement, but is not strictly required for it. + +# Guide-level explanation +[guide-level-explanation]: #guide-level-explanation + +The `BufferPointer` object contains a reference to the buffer that it +acts upon, and the indices at which the access is being performed. + +`BufferPointer` also provides a utility function +`BufferPointerNode::value_dtype()`, which returns the expected +datatype at the specified location. This will typically be the same +as the buffer's datatype, but may have a different number of lanes. +For example, a `BufferPointer` whose buffer's datatype is `float16`, +and whose index is `Ramp(pos, stride=1, lanes=4)` for vectorized +access will return `float16*4` for `value_dtype()`. + +Previously, `BufferLoad` and `BufferStore` held references to the +buffer and indices directly. To migrate these codes, replace +references to `BufferX::buffer` with `BufferX::pointer::buffer` and +replace references to `BufferX::indices` with +`BufferX::pointer::indices`. + + +# Reference-level explanation +[reference-level-explanation]: #reference-level-explanation + +This is the technical portion of the RFC. Explain the design in sufficient detail that: + +- Its interaction with other features is clear. +- It is reasonably clear how the feature would be implemented. +- Corner cases are dissected by example. + +The section should return to the examples given in the previous section, +and explain more fully how the detailed proposal makes those examples work. + +# Drawbacks +[drawbacks]: #drawbacks + +Requires an additional indirection to access the buffer and pointer, +so transformations that must distinguish between reads and writes +become more verbose. + +# Rationale and alternatives +[rationale-and-alternatives]: #rationale-and-alternatives + +The `BufferLoad::pointer` and `BufferStore::pointer` could be generic +`PrimExpr`, instead of being `BufferPointer` objects. This would +require the datatype to be a handle, with an additional parameter to +indicate what is being stored. However, this would require all +visitors that interact or modify a buffer to be duplicated across both +reads and writes. + +Currently, the `BufferLoad::pointer` and `BufferStore::pointer` +objects are visited by `ExprMutator` and `StmtMutator`, but are +required to be `BufferPointer` objects. This is implemented in +type-checking in `ExprMutator::VisitExpr_(BufferLoad*)` and +`StmtMutator::VisitStmt_(BufferStore*)`, but it isn't apparent at the +callsite that the returned `PrimExpr` must be a `BufferPointer`. + +Prior to this RFC's implementation, all transformations that modify a +buffer must have near-equivalent mutators for both the `BufferLoad` +and `BufferStore` nodes. + +# Prior art +[prior-art]: #prior-art + +This follows a C-style convention. Given an array `int x[100]`, the +location `int* ptr = (x+50)` represents the location of element 50 in +array `x`. A buffer load is then represented as `*ptr`, and buffer +store is represented as `*ptr = val;`. + +This also maps onto Vulkan semantics, where a `OpAccessChain` +instruction is used to generate a pointer into an array, which can +then be used with either `OpLoad` or `OpStore`. + +# Unresolved questions +[unresolved-questions]: #unresolved-questions + + + +# Future possibilities +[future-possibilities]: #future-possibilities + +The `BufferPointer` could represent a pointer to a specific element in +the generated C code. This can be used for generating pointers to +pass into hardware-specific intrinsics, rather than using +`BufferNode::elem_offset` or the built-in `tvm_access_ptr`.