[GitHub] [tvm-rfcs] Lunderberg commented on a change in pull request #42: [TIR] Adding BufferPointer node, used by BufferLoad and BufferStore

[GitBox <gi...@apache.org>]

Lunderberg commented on a change in pull request #42:
URL: https://github.com/apache/tvm-rfcs/pull/42#discussion_r740396412



##########
File path: rfcs/0042-buffer-pointer.md
##########
@@ -0,0 +1,116 @@
+- 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#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
+
+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 

Review comment:
       Thank you for that catch, and updated.  I tend to bounce between different sentences when editing, and I thought I had gone through it to catch all of these, but it looks like I missed one.




-- 
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: commits-unsubscribe@tvm.apache.org

For queries about this service, please contact Infrastructure at:
users@infra.apache.org