(sch: s_tir.Schedule)
| 131 | B[i, j] = A[i, j] |
| 132 | |
| 133 | def schedule_texture_read(sch: s_tir.Schedule): |
| 134 | B_blk = sch.get_sblock("Copy") |
| 135 | Ai_block = sch.cache_read(B_blk, 0, "global.texture") |
| 136 | sch.transform_layout(Ai_block, ("write", 0), lambda i, j: (i, j // lanes, j % lanes)) |
| 137 | |
| 138 | def schedule_default(blk, lanes): |
| 139 | i, j = sch.get_loops(blk) |
| 140 | jo, jv = sch.split(j, [None, lanes]) |
| 141 | |
| 142 | b = sch.fuse(i, jo) |
| 143 | bx, tx = sch.split(b, [None, 256]) |
| 144 | sch.bind(bx, "blockIdx.x") |
| 145 | sch.bind(tx, "threadIdx.x") |
| 146 | |
| 147 | sch.vectorize(jv) |
| 148 | |
| 149 | schedule_default(Ai_block, lanes) |
| 150 | schedule_default(B_blk, read_width) |
| 151 | |
| 152 | mod = TextureCopy |
| 153 |
no test coverage detected
searching dependent graphs…