Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

[TIR] Get read/write access precisely for opaque access. #11110

Merged
merged 2 commits into from
Apr 28, 2022
Merged

[TIR] Get read/write access precisely for opaque access. #11110

merged 2 commits into from
Apr 28, 2022

Conversation

qsqqsqqsq-intellif
Copy link
Contributor

Previously we assume opaque access as both read and write access in compute_inline. Here we present a way to tell the opaque access is whether a read or write access precisely. The opaque access is wrapped with tvm_access_ptr as following:

T.tvm_access_ptr(
    T.type_annotation(dtype="int8"),
    A.data,
    0,
    1024,
    1, # access_mask, 1:read access, 2:write access
    dtype="handle",
),

We can get the access_mask from tvm_access_ptr in BlockReadWriteDetector and put this opaque access to read_regions or write_regions according to access_mask.

Hzfengsy
Hzfengsy requested changes Apr 25, 2022
View reviewed changes
tests/python/unittest/test_tir_schedule_compute_inline.py Outdated
@@ -185,7 +185,7 @@ def opaque_access_load(a: T.handle, c: T.handle) -> None:
T.writes(C[0:128, 0:128])
T.evaluate(
T.tvm_access_ptr(
T.type_annotation(dtype="float32"), B.data, 0, 128, "r", dtype="handle"
T.type_annotation(dtype="float32"), B.data, 0, 128, 1, dtype="handle"
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

It would be great if we could also handle "r" and "w" for better readability

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

OK, I use access_prt in this case.:
T.evaluate(B.access_ptr("r", extent=128))

@qsqqsqqsq-intellif qsqqsqqsq-intellif force-pushed the get_read_write_region_precisely branch from f1f152a to 42dc437 Compare April 25, 2022 04:57
wrongtest-intellif
wrongtest-intellif reviewed Apr 25, 2022
View reviewed changes
tests/python/unittest/test_tir_schedule_compute_inline.py Outdated
sch = tir.Schedule(exp_exp_opaque_access_with_tvm_access_ptr, debug_mask="all")
compute = sch.get_block("compute")
sch.compute_inline(compute)
print(sch.mod.script())
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

remove print

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

removed

@qsqqsqqsq-intellif qsqqsqqsq-intellif force-pushed the get_read_write_region_precisely branch from bb730de to 2998158 Compare April 26, 2022 06:34
Hzfengsy
Hzfengsy requested changes Apr 26, 2022
View reviewed changes
include/tvm/tir/buffer.h Outdated
int content_lanes = 1,
PrimExpr offset = IntImm(DataType::Int(32), 0)) const;
int content_lanes = 1, PrimExpr offset = IntImm(DataType::Int(32), 0),
PrimExpr extent = PrimExpr(ObjectPtr<Object>(nullptr))) const;
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

if extent can be nullptr, please use Optional<PrimExpr> instead of PrimExpr.

src/tir/analysis/block_access_region_detector.cc
@@ -181,6 +181,34 @@ void BlockReadWriteDetector::VisitStmt_(const IfThenElseNode* op) {
}

void BlockReadWriteDetector::VisitExpr_(const CallNode* op) {
if (op->op.same_as(builtin::tvm_access_ptr())) {
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Please add testcase for BlockReadWriteDetector directly in tests/python/unittest/test_tir_analysis_get_block_access_region.py

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Should I keep the testcase in tests/python/unittest/test_tir_schedule_compute_inline.py?

Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Of course, the more test cases the better

@qsqqsqqsq-intellif qsqqsqqsq-intellif force-pushed the get_read_write_region_precisely branch 4 times, most recently from e92d012 to 3445935 Compare April 27, 2022 02:10
Hzfengsy
Hzfengsy reviewed Apr 27, 2022
View reviewed changes
tests/python/unittest/test_tir_analysis_get_block_access_region.py
with T.block("opaque"):
T.reads(A[0:1024], C[0:1024])
T.writes(B[0:1024], C[0:1024])
T.evaluate(A.access_ptr("r"))
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Would be great if we can test different start point and different extent

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Currently, we use FullRegion of the opaque access with tvm_access_ptr, so the offset and extent does not count. I try to use access_ptr's offset and extent as region, but access_ptr can only express the region as flattened which will be failed when the buffer dims > 1.

@qsqqsqqsq-intellif qsqqsqqsq-intellif force-pushed the get_read_write_region_precisely branch 2 times, most recently from cc04d7b to 756c16c Compare April 27, 2022 09:33
sqing added 2 commits April 28, 2022 17:05
[TIR] Get read/write access precisely for opaque access.
b202b76 
When the opaque access is wrapped with tvm_access_ptr, we can get the access_mask
from tvm_access_ptr in BlockReadWriteDetector and put this opaque access to read_regions
or write_regions according to access_mask.
[TIR] Add parameter extent for access_ptr.
4e28131
@qsqqsqqsq-intellif qsqqsqqsq-intellif force-pushed the get_read_write_region_precisely branch from 756c16c to 4e28131 Compare April 28, 2022 09:05
@Hzfengsy Hzfengsy merged commit 7710dfd into apache:main Apr 28, 2022
qsqqsqqsq-intellif added a commit to qsqqsqqsq-intellif/tvm that referenced this pull request Apr 29, 2022
@qsqqsqqsq-intellif
[TIR] Get read/write access precisely for opaque access. (apache#11110)
9766332 
* [TIR] Get read/write access precisely for opaque access.

When the opaque access is wrapped with tvm_access_ptr, we can get the access_mask
from tvm_access_ptr in BlockReadWriteDetector and put this opaque access to read_regions
or write_regions according to access_mask.

* [TIR] Add parameter extent for access_ptr.

Co-authored-by: sqing <qing.siqi@intellif.com>
shtinsa pushed a commit to Deelvin/tvm that referenced this pull request May 17, 2022
@qsqqsqqsq-intellif
[TIR] Get read/write access precisely for opaque access. (apache#11110)
96a274a 
* [TIR] Get read/write access precisely for opaque access.

When the opaque access is wrapped with tvm_access_ptr, we can get the access_mask
from tvm_access_ptr in BlockReadWriteDetector and put this opaque access to read_regions
or write_regions according to access_mask.

* [TIR] Add parameter extent for access_ptr.

Co-authored-by: sqing <qing.siqi@intellif.com>
@driazati driazati mentioned this pull request Jul 14, 2022
Closed
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Reviewers

@wrongtest-intellif wrongtest-intellif wrongtest-intellif left review comments

@Hzfengsy Hzfengsy Hzfengsy approved these changes

Assignees
No one assigned
Labels
None yet
Projects
None yet
Milestone
No milestone
Development

Successfully merging this pull request may close these issues.
3 participants
@qsqqsqqsq-intellif @wrongtest-intellif @Hzfengsy