From a0932ddf9789be5bb8b26f9019487afbf94121f7 Mon Sep 17 00:00:00 2001 From: Eric Lunderberg Date: Wed, 27 Oct 2021 12:31:34 -0500 Subject: [PATCH 1/3] [TIR] Adding BufferPointer node, used by BufferLoad and BufferStore --- rfcs/XXXX-buffer-pointer.md | 109 ++++++++++++++++++++++++++++++++++++ 1 file changed, 109 insertions(+) create mode 100644 rfcs/XXXX-buffer-pointer.md diff --git a/rfcs/XXXX-buffer-pointer.md b/rfcs/XXXX-buffer-pointer.md new file mode 100644 index 00000000..f336a595 --- /dev/null +++ b/rfcs/XXXX-buffer-pointer.md @@ -0,0 +1,109 @@ +- Feature Name: (fill me in with a unique identifier, `my_awesome_feature`) +- Start Date: (fill me in with today's date, YYYY-MM-DD) +- RFC PR: [apache/tvm-rfcs#0000](https://github.com/apache/tvm-rfcs/pull/0000) +- GitHub Issue: [apache/tvm#0000](https://github.com/apache/tvm/issues/0000) + +# 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. + +# 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 + +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`. From 5bf206982ad61eb864491f1627c50d30ad579579 Mon Sep 17 00:00:00 2001 From: Eric Lunderberg Date: Thu, 28 Oct 2021 15:57:14 -0500 Subject: [PATCH 2/3] Updated with link to RFC PR --- rfcs/{XXXX-buffer-pointer.md => 0042-buffer-pointer.md} | 9 ++++++++- 1 file changed, 8 insertions(+), 1 deletion(-) rename rfcs/{XXXX-buffer-pointer.md => 0042-buffer-pointer.md} (92%) diff --git a/rfcs/XXXX-buffer-pointer.md b/rfcs/0042-buffer-pointer.md similarity index 92% rename from rfcs/XXXX-buffer-pointer.md rename to rfcs/0042-buffer-pointer.md index f336a595..6f349320 100644 --- a/rfcs/XXXX-buffer-pointer.md +++ b/rfcs/0042-buffer-pointer.md @@ -1,7 +1,8 @@ - Feature Name: (fill me in with a unique identifier, `my_awesome_feature`) - Start Date: (fill me in with today's date, YYYY-MM-DD) -- RFC PR: [apache/tvm-rfcs#0000](https://github.com/apache/tvm-rfcs/pull/0000) +- RFC PR: [apache/tvm-rfcs#0042](https://github.com/apache/tvm-rfcs/pull/0042) - GitHub Issue: [apache/tvm#0000](https://github.com/apache/tvm/issues/0000) +- Related RFCs: [RFC#0039](https://github.com/apache/tvm-rfcs/pull/0039) # Summary [summary]: #summary @@ -12,6 +13,8 @@ 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 @@ -24,6 +27,10 @@ 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 From 04719307bc33e8632603c4d24a17b6fd3d40cdf4 Mon Sep 17 00:00:00 2001 From: Eric Lunderberg Date: Mon, 1 Nov 2021 12:30:41 -0500 Subject: [PATCH 3/3] Finished an unfinished sentence, filled in additional header details. --- rfcs/0042-buffer-pointer.md | 10 ++++++---- 1 file changed, 6 insertions(+), 4 deletions(-) diff --git a/rfcs/0042-buffer-pointer.md b/rfcs/0042-buffer-pointer.md index 6f349320..9088a35f 100644 --- a/rfcs/0042-buffer-pointer.md +++ b/rfcs/0042-buffer-pointer.md @@ -1,7 +1,7 @@ -- Feature Name: (fill me in with a unique identifier, `my_awesome_feature`) -- Start Date: (fill me in with today's date, YYYY-MM-DD) +- Feature Name: BufferPointer +- Start Date: 2021-10-28 - RFC PR: [apache/tvm-rfcs#0042](https://github.com/apache/tvm-rfcs/pull/0042) -- GitHub Issue: [apache/tvm#0000](https://github.com/apache/tvm/issues/0000) +- GitHub PR: [PR#9391](https://github.com/apache/tvm/pull/9391) - Related RFCs: [RFC#0039](https://github.com/apache/tvm-rfcs/pull/0039) # Summary @@ -77,7 +77,9 @@ become more verbose. 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 +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