()
| 175 | |
| 176 | |
| 177 | def get_expected_2(): |
| 178 | # fmt: off |
| 179 | @I.ir_module(s_tir=True) |
| 180 | class Expected: |
| 181 | @T.prim_func(private=True, s_tir=True) |
| 182 | def f_mul(A: T.Buffer((T.int64(5), T.int64(5)), "float32"), f_mul2: T.Buffer((T.int64(5), T.int64(5)), "float32")): |
| 183 | T.func_attr({"tirx.noalias": True}) |
| 184 | # with T.sblock("root"): |
| 185 | for i0, i1 in T.grid(T.int64(5), T.int64(5)): |
| 186 | with T.sblock("f_mul2"): |
| 187 | v_i0, v_i1 = T.axis.remap("SS", [i0, i1]) |
| 188 | T.reads(A[v_i0, v_i1]) |
| 189 | T.writes(f_mul2[v_i0, v_i1]) |
| 190 | f_mul2[v_i0, v_i1] = A[v_i0, v_i1] * T.float32(2) |
| 191 | |
| 192 | @T.prim_func(private=True, s_tir=True) |
| 193 | def f_mulk_grad(A: T.Buffer((T.int64(5), T.int64(5)), "float32"), B: T.Buffer((T.int64(5), T.int64(5)), "float32"), f_mulk_grad_1: T.Buffer((T.int64(5), T.int64(5)), "float32")): |
| 194 | T.func_attr({"tirx.noalias": True}) |
| 195 | # with T.sblock("root"): |
| 196 | for i0, i1 in T.grid(T.int64(5), T.int64(5)): |
| 197 | with T.sblock("f_mulk_grad"): |
| 198 | v_i0, v_i1 = T.axis.remap("SS", [i0, i1]) |
| 199 | T.reads(A[v_i0, v_i1]) |
| 200 | T.writes(f_mulk_grad_1[v_i0, v_i1]) |
| 201 | f_mulk_grad_1[v_i0, v_i1] = A[v_i0, v_i1] * T.float32(2) |
| 202 | |
| 203 | @R.function |
| 204 | def main_adjoint(a: R.Tensor((5, 5), dtype="float32")) -> R.Tuple(R.Tensor((), dtype="float32"), R.Tuple(R.Tensor((5, 5), dtype="float32"))): |
| 205 | cls = Expected |
| 206 | with R.dataflow(): |
| 207 | lv = R.call_tir(cls.f_mul, (a,), out_sinfo=R.Tensor((5, 5), dtype="float32")) |
| 208 | gv: R.Tensor((), dtype="float32") = R.sum(lv, axis=None, keepdims=False) |
| 209 | gv_adjoint: R.Tensor((), dtype="float32") = R.ones(R.shape([]), dtype="float32") |
| 210 | lv_adjoint: R.Tensor((5, 5), dtype="float32") = R.broadcast_to(gv_adjoint, R.shape([5, 5])) |
| 211 | lv_1 = R.call_tir(cls.f_mulk_grad, (lv_adjoint, a), out_sinfo=R.Tensor((5, 5), dtype="float32")) |
| 212 | a_adjoint: R.Tensor((5, 5), dtype="float32") = lv_1 |
| 213 | a_adjoint_out: R.Tensor((5, 5), dtype="float32") = a_adjoint |
| 214 | R.output(gv, a_adjoint_out) |
| 215 | return (gv, (a_adjoint_out,)) |
| 216 | |
| 217 | @R.function |
| 218 | def main(a: R.Tensor((5, 5), dtype="float32")) -> R.Tensor((), dtype="float32"): |
| 219 | cls = Expected |
| 220 | with R.dataflow(): |
| 221 | lv = R.call_tir_with_grad(cls.f_mul, (a,), out_sinfo=R.Tensor((5, 5), dtype="float32"), te_grad_name="f_mulk_grad", te_grad_kwargs={"k": T.float32(2)}) |
| 222 | gv: R.Tensor((), dtype="float32") = R.sum(lv, axis=None, keepdims=False) |
| 223 | R.output(gv) |
| 224 | return gv |
| 225 | # fmt: on |
| 226 | return Expected |
| 227 | |
| 228 | |
| 229 | def test_emit_te_kwargs(register_te_grads): |
no outgoing calls
no test coverage detected
searching dependent graphs…