| 102 | |
| 103 | @T.prim_func(private=True) |
| 104 | def before2(in_buf: T.Buffer((16, 16), "float32"), out: T.Buffer((16, 16), "float32")) -> None: |
| 105 | T.device_entry() |
| 106 | bx, by, bz = T.cta_id([1, 1, 1]) |
| 107 | T.warp_id([1]) |
| 108 | lane_id = T.lane_id([32]) |
| 109 | atom = T.TileLayout(T.S[(1, 2) : (2, 1)]) |
| 110 | tile = T.TileLayout(T.S[(2, 2) : (2, 1)]) |
| 111 | warp_atom = atom.tile(L_LANE, (8, 4), (1, 2)) |
| 112 | A = T.alloc_buffer( |
| 113 | [4, 2], dtype="float32", scope="local", layout=atom.tile(tile, (2, 2), (1, 2)) |
| 114 | ) |
| 115 | B_layout = warp_atom.tile(tile, (2, 2), (8, 8)) |
| 116 | B = A.view(16, 16, layout=B_layout) |
| 117 | A_local = B.local(2, 2, 2) |
| 118 | for i in T.unroll(4): |
| 119 | for j in T.vectorized(2): |
| 120 | A_local[i // 2, i % 2, j] = in_buf[ |
| 121 | i // 2 * 8 + lane_id // 4, i % 2 * 8 + lane_id % 4 + j |
| 122 | ] |
| 123 | B_1 = A.view(16, 16, layout=B_layout) |
| 124 | A_local_1 = B_1.local(8) |
| 125 | for i in T.vectorized(2): |
| 126 | out[lane_id // 4 * 8 + i // 2 * 8 + lane_id % 4, lane_id % 4 * 2 + i % 2] = A_local_1[i] |
| 127 | |
| 128 | @T.prim_func(private=True) |
| 129 | def after2(in_buf_handle: T.handle, out_handle: T.handle): |