Lunderberg commented on code in PR #16660:
URL: https://github.com/apache/tvm/pull/16660#discussion_r1509063367
##########
tests/python/tir-schedule/test_tir_schedule_cache_read_write.py:
##########
@@ -1379,6 +1379,54 @@ def
test_cache_read_fail_invalid_storage_scope(use_block_name):
sch.cache_read(block_b, 0, "test_scope")
+def test_cache_read_allocate_const():
+ @T.prim_func
+ def before(A: T.Buffer((8), "float32"), C: T.Buffer((8), "float32")):
+ B = T.allocate_const([0.0, 0.1, 0.2, 0.3, 0.4, 0.5, 0.6, 0.7],
"float32", [8])
+ B_buf = T.decl_buffer((8), dtype="float32", data=B)
+ for i in T.serial(128):
+ with T.block("C"):
+ vi = T.axis.remap("S", [i])
+ T.reads(A[vi], B_buf[vi])
Review Comment:
Thank you for cleaning it up, and that's an iteresting failure. It looks
like the `T.reads()` annotations are only generated for buffers that can be
manipulated withing schedulable TIR. The filtering is applied
[here](https://github.com/apache/tvm/blob/main/src/tir/analysis/block_access_region_detector.cc#L300),
and only contains buffers that were passed as arguments ([added
here](https://github.com/apache/tvm/blob/main/src/tir/ir/script/script_complete.cc#L107))
or that were part of a `tir::Block` node ([added
here](https://github.com/apache/tvm/blob/main/src/tir/ir/script/script_complete.cc#L52)).
I think this is a bug in the `GetBlockAccessRegion`, that it should track
reads/writes in buffers declared by `DeclBuffer` nodes as well.
But that's for a later clean-up on my side, and this PR looks good! Thank
you!
--
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: [email protected]
For queries about this service, please contact Infrastructure at:
[email protected]
