Create a block that reads a buffer region into a write cache. It requires: 1) There is only one block who write the buffer in the scope. 2) The scope block have stage-pipeline property. Parameters ---------- block : SBlockRV | str The producer
(
self,
block: SBlockRV | str,
write_buffer_index: int | str | Buffer,
storage_scope: str,
consumer_blocks: list[SBlockRV | str] | None = None,
)
| 1450 | |
| 1451 | @type_checked |
| 1452 | def cache_write( |
| 1453 | self, |
| 1454 | block: SBlockRV | str, |
| 1455 | write_buffer_index: int | str | Buffer, |
| 1456 | storage_scope: str, |
| 1457 | consumer_blocks: list[SBlockRV | str] | None = None, |
| 1458 | ) -> SBlockRV: |
| 1459 | """Create a block that reads a buffer region into a write cache. It requires: |
| 1460 | |
| 1461 | |
| 1462 | 1) There is only one block who write the buffer in the scope. |
| 1463 | |
| 1464 | 2) The scope block have stage-pipeline property. |
| 1465 | |
| 1466 | Parameters |
| 1467 | ---------- |
| 1468 | block : SBlockRV | str |
| 1469 | The producer block of the target buffer. |
| 1470 | |
| 1471 | write_buffer_index: int |
| 1472 | The index of the buffer in block's write region, the unique |
| 1473 | name of a write buffer in the block, or a Buffer object |
| 1474 | that is within the blocks write region. |
| 1475 | |
| 1476 | storage_scope: str |
| 1477 | The target storage scope. |
| 1478 | |
| 1479 | consumer_blocks: Optional[List[SBlockRV | str]] |
| 1480 | An optional list of consumers that should read directly from the cache. |
| 1481 | If not specified, all consumers will read from the original buffer. |
| 1482 | |
| 1483 | Returns |
| 1484 | ------- |
| 1485 | cached_block : SBlockRV |
| 1486 | The block of the cache stage |
| 1487 | |
| 1488 | Examples |
| 1489 | -------- |
| 1490 | Before cache_write, in TensorIR, the IR is: |
| 1491 | |
| 1492 | .. code-block:: python |
| 1493 | |
| 1494 | @T.prim_func(s_tir=True) |
| 1495 | def before_cache_write(a: T.handle, b: T.handle) -> None: |
| 1496 | A = T.match_buffer(a, (128, 128)) |
| 1497 | B = T.match_buffer(b, (128, 128)) |
| 1498 | for i, j in T.grid(128, 128): |
| 1499 | with T.sblock("B"): |
| 1500 | vi, vj = T.axis.remap("SS", [i, j]) |
| 1501 | B[vi, vj] = A[vi, vj] * 2.0 |
| 1502 | |
| 1503 | Create the schedule and cache_write: |
| 1504 | |
| 1505 | .. code-block:: python |
| 1506 | |
| 1507 | sch = tvm.s_tir.Schedule(before_cache_write) |
| 1508 | block_b = sch.get_sblock("B") |
| 1509 | sch.cache_write(block_b, 0, "local") |