(register_te_grads)
| 146 | |
| 147 | |
| 148 | def test_call_tir(register_te_grads): |
| 149 | # fmt: off |
| 150 | @I.ir_module(s_tir=True) |
| 151 | class Before: |
| 152 | @T.prim_func(private=True, s_tir=True) |
| 153 | def f_mul(A: T.Buffer((T.int64(5), T.int64(5)), "float32"), B: T.Buffer((T.int64(5), T.int64(5)), "float32"), f_mul_1: T.Buffer((T.int64(5), T.int64(5)), "float32")): |
| 154 | T.func_attr({"tirx.noalias": True}) |
| 155 | # with T.sblock("root"): |
| 156 | for i0, i1 in T.grid(T.int64(5), T.int64(5)): |
| 157 | with T.sblock("f_mul"): |
| 158 | v_i0, v_i1 = T.axis.remap("SS", [i0, i1]) |
| 159 | T.reads(A[v_i0, v_i1], B[v_i0, v_i1]) |
| 160 | T.writes(f_mul_1[v_i0, v_i1]) |
| 161 | f_mul_1[v_i0, v_i1] = A[v_i0, v_i1] * B[v_i0, v_i1] |
| 162 | |
| 163 | @R.function |
| 164 | def main(a: R.Tensor((5, 5), dtype="float32"), b: R.Tensor((5, 5), dtype="float32")) -> R.Tensor((), dtype="float32"): |
| 165 | cls = Before |
| 166 | with R.dataflow(): |
| 167 | lv = R.call_tir_with_grad(cls.f_mul, (a, b), out_sinfo=R.Tensor((5, 5), dtype="float32"), te_grad_name="f_mul_grad") |
| 168 | gv: R.Tensor((), dtype="float32") = R.sum(lv, axis=None, keepdims=False) |
| 169 | R.output(gv) |
| 170 | return gv |
| 171 | # fmt: off |
| 172 | |
| 173 | After = Gradient("main")(Before) |
| 174 | assert_structural_equal(After, get_expected_1()) |
| 175 | |
| 176 | |
| 177 | def get_expected_2(): |
nothing calls this directly
no test coverage detected
searching dependent graphs…