()
| 59 | |
| 60 | |
| 61 | def get_expected_1(): |
| 62 | # fmt: off |
| 63 | @I.ir_module(s_tir=True) |
| 64 | class Expected: |
| 65 | @T.prim_func(private=True, s_tir=True) |
| 66 | 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")): |
| 67 | T.func_attr({"tirx.noalias": True}) |
| 68 | # with T.sblock("root"): |
| 69 | for i0, i1 in T.grid(T.int64(5), T.int64(5)): |
| 70 | with T.sblock("f_mul"): |
| 71 | v_i0, v_i1 = T.axis.remap("SS", [i0, i1]) |
| 72 | T.reads(A[v_i0, v_i1], B[v_i0, v_i1]) |
| 73 | T.writes(f_mul_1[v_i0, v_i1]) |
| 74 | f_mul_1[v_i0, v_i1] = A[v_i0, v_i1] * B[v_i0, v_i1] |
| 75 | |
| 76 | @T.prim_func(private=True, s_tir=True) |
| 77 | def f_mul_grad(A: T.Buffer((T.int64(5), T.int64(5)), "float32"), B: T.Buffer((T.int64(5), T.int64(5)), "float32"), C: T.Buffer((T.int64(5), T.int64(5)), "float32"), f_mul_grad_1: T.Buffer((T.int64(5), T.int64(5)), "float32"), f_mul_grad_2: T.Buffer((T.int64(5), T.int64(5)), "float32")): |
| 78 | T.func_attr({"tirx.noalias": True}) |
| 79 | # with T.sblock("root"): |
| 80 | for i0, i1 in T.grid(T.int64(5), T.int64(5)): |
| 81 | with T.sblock("f_mul_grad_1"): |
| 82 | v_i0, v_i1 = T.axis.remap("SS", [i0, i1]) |
| 83 | T.reads(C[v_i0, v_i1], A[v_i0, v_i1]) |
| 84 | T.writes(f_mul_grad_1[v_i0, v_i1]) |
| 85 | f_mul_grad_1[v_i0, v_i1] = C[v_i0, v_i1] * A[v_i0, v_i1] |
| 86 | for i0, i1 in T.grid(T.int64(5), T.int64(5)): |
| 87 | with T.sblock("f_mul_grad_2"): |
| 88 | v_i0, v_i1 = T.axis.remap("SS", [i0, i1]) |
| 89 | T.reads(B[v_i0, v_i1], A[v_i0, v_i1]) |
| 90 | T.writes(f_mul_grad_2[v_i0, v_i1]) |
| 91 | f_mul_grad_2[v_i0, v_i1] = B[v_i0, v_i1] * A[v_i0, v_i1] |
| 92 | |
| 93 | @R.function |
| 94 | def main_adjoint(a: R.Tensor((5, 5), dtype="float32"), b: R.Tensor((5, 5), dtype="float32")) -> R.Tuple(R.Tensor((), dtype="float32"), R.Tuple(R.Tensor((5, 5), dtype="float32"), R.Tensor((5, 5), dtype="float32"))): |
| 95 | cls = Expected |
| 96 | with R.dataflow(): |
| 97 | lv = R.call_tir(cls.f_mul, (a, b), out_sinfo=R.Tensor((5, 5), dtype="float32")) |
| 98 | gv: R.Tensor((), dtype="float32") = R.sum(lv, axis=None, keepdims=False) |
| 99 | gv_adjoint: R.Tensor((), dtype="float32") = R.ones(R.shape([]), dtype="float32") |
| 100 | lv_adjoint: R.Tensor((5, 5), dtype="float32") = R.broadcast_to(gv_adjoint, R.shape([5, 5])) |
| 101 | lv_1 = R.call_tir(cls.f_mul_grad, (lv_adjoint, a, b), out_sinfo=[R.Tensor((5, 5), dtype="float32"), R.Tensor((5, 5), dtype="float32")]) |
| 102 | a_adjoint: R.Tensor((5, 5), dtype="float32") = lv_1[0] |
| 103 | b_adjoint: R.Tensor((5, 5), dtype="float32") = lv_1[1] |
| 104 | a_adjoint_out: R.Tensor((5, 5), dtype="float32") = a_adjoint |
| 105 | b_adjoint_out: R.Tensor((5, 5), dtype="float32") = b_adjoint |
| 106 | R.output(gv, a_adjoint_out, b_adjoint_out) |
| 107 | return (gv, (a_adjoint_out, b_adjoint_out)) |
| 108 | |
| 109 | @R.function |
| 110 | def main(a: R.Tensor((5, 5), dtype="float32"), b: R.Tensor((5, 5), dtype="float32")) -> R.Tensor((), dtype="float32"): |
| 111 | cls = Expected |
| 112 | with R.dataflow(): |
| 113 | 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") |
| 114 | gv: R.Tensor((), dtype="float32") = R.sum(lv, axis=None, keepdims=False) |
| 115 | R.output(gv) |
| 116 | return gv |
| 117 | # fmt: on |
| 118 | return Expected |
no outgoing calls
no test coverage detected
searching dependent graphs…