[TIR] Get read/write access precisely for opaque access.#11110
Conversation
| 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" |
There was a problem hiding this comment.
It would be great if we could also handle "r" and "w" for better readability
There was a problem hiding this comment.
OK, I use access_prt in this case.:
T.evaluate(B.access_ptr("r", extent=128))
qsqqsqqsq-intellif
force-pushed
the
get_read_write_region_precisely
branch
from
f1f152a to
42dc437
Compare
| 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()) |
qsqqsqqsq-intellif
force-pushed
the
get_read_write_region_precisely
branch
from
bb730de to
2998158
Compare
| 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; |
There was a problem hiding this comment.
if extent can be nullptr, please use Optional<PrimExpr> instead of PrimExpr.
| } | ||
|
|
||
| void BlockReadWriteDetector::VisitExpr_(const CallNode* op) { | ||
| if (op->op.same_as(builtin::tvm_access_ptr())) { |
There was a problem hiding this comment.
Please add testcase for BlockReadWriteDetector directly in tests/python/unittest/test_tir_analysis_get_block_access_region.py
There was a problem hiding this comment.
Should I keep the testcase in tests/python/unittest/test_tir_schedule_compute_inline.py?
There was a problem hiding this comment.
Of course, the more test cases the better
qsqqsqqsq-intellif
force-pushed
the
get_read_write_region_precisely
branch
4 times, most recently
from
e92d012 to
3445935
Compare
| 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")) |
There was a problem hiding this comment.
Would be great if we can test different start point and different extent
There was a problem hiding this comment.
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
force-pushed
the
get_read_write_region_precisely
branch
2 times, most recently
from
cc04d7b to
756c16c
Compare
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.
qsqqsqqsq-intellif
force-pushed
the
get_read_write_region_precisely
branch
from
756c16c to
4e28131
Compare
* [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>
* [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>
3 participants
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:
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.