# Codegen support for relax.layout\_transform **URL:** \ **Category:** unity\ **Created:** [June 26, 2023, 11:09am UTC](https://discuss.tvm.apache.org/t/codegen-support-for-relax-layout-transform/15187 "2023-06-26T11:09:08Z")\ **Posts on this page:** 7\ **Page:** 1
**Author:** ![rutkoor](https://discuss.tvm.apache.org/letter_avatar_proxy/v4/letter/r/90ced4/32.png) [@rutkoor](https://discuss.tvm.apache.org/u/rutkoor)\ **Post date:** [June 26, 2023, 11:09am UTC](https://discuss.tvm.apache.org/t/codegen-support-for-relax-layout-transform/15187/1 "2023-06-26T11:09:08Z")
Hello, When I am trying to do relax.build on the IRModule from _tests/python/relax/test\_transform\_alter\_op\_impl.py_ example, I am noticing build errors saying that the codegen support for the intrinsic(**Op(relax.layout\_transform)**) is not handled. Reduced test case: > ``` > from tvm import relax > from tvm.script import tir as T, ir as I, relax as R > > @I.ir_module > class AddTest: > @T.prim_func > def relax_add_replacement(arg0: T.Buffer((4, 4), "float32"), arg1: T.Buffer((4, 4), "float32"), output: T.Buffer((4, 4), "float32")): > T.func_attr({"operator_name": "relax.add"}) > for ax0, ax1 in T.grid(4, 4): > with T.block("T_add"): > v_ax0, v_ax1 = T.axis.remap("SS", [ax0, ax1]) > T.reads(arg0[v_ax0, v_ax1], arg1[v_ax0, v_ax1]) > T.writes(output[v_ax0, v_ax1]) > output[v_ax0, v_ax1] = arg0[v_ax0, v_ax1] + arg1[v_ax0, v_ax1] > > @R.function > def main(x: R.Tensor((16,), dtype="float32"), y: R.Tensor((16,), dtype="float32")) -> R.Tensor((16,), dtype="float32"): > with R.dataflow(): > lv: R.Tensor((4, 4), dtype="float32") = R.layout_transform(x, index_map=lambda i: (i // 4, i % 4), pad_value=None) > lv1: R.Tensor((4, 4), dtype="float32") = R.layout_transform(y, index_map=lambda i: (i // 4, i % 4), pad_value=None) > lv2 = R.call_tir(AddTest.relax_add_replacement, (lv, lv1), out_sinfo=R.Tensor((4, 4), dtype="float32")) > lv_1: R.Tensor((16,), dtype="float32") = R.layout_transform(lv2, index_map=lambda axis0, axis1: (axis0 * 4 + axis1,), pad_value=None) > gv: R.Tensor((16,), dtype="float32") = lv_1 > R.output(gv) > return gv > > lib = relax.build(mod=AddTest, target="llvm", exec_mode="compiled") > > ``` Upon compiling above test case, I am seeing this error, > _\>\>\>\>\>\> File “tvm/src/relax/backend/vm/codegen\_vm\_tir.cc”, line 240 ::_ _TVMError: CodeGenVMTIR cannot handle this intrinsic now: Op(relax.layout\_transform)_ Can someone please let me know why the support is not being added or point me to any existing issue where they are tracking **adding codegen support for relax.layout\_transform**? Thanks. ---
**Author:** ![Hzfengsy](https://discuss.tvm.apache.org/user_avatar/discuss.tvm.apache.org/hzfengsy/32/2961_2.png) [@Hzfengsy](https://discuss.tvm.apache.org/u/Hzfengsy)\ **Post date:** [June 26, 2023, 12:34pm UTC](https://discuss.tvm.apache.org/t/codegen-support-for-relax-layout-transform/15187/2 "2023-06-26T12:34:34Z")
I guess there is no op legalization support for `relax.layout_transform`. cc @MasterJH5574 ---
**Author:** ![MasterJH5574](https://discuss.tvm.apache.org/user_avatar/discuss.tvm.apache.org/masterjh5574/32/3449_2.png) [@MasterJH5574](https://discuss.tvm.apache.org/u/MasterJH5574)\ **Post date:** [June 26, 2023, 4:35pm UTC](https://discuss.tvm.apache.org/t/codegen-support-for-relax-layout-transform/15187/3 "2023-06-26T16:35:34Z")
Here is the tracking issue for the current layout planning [https://github.com/apache/tvm/issues/14070](https://github.com/apache/tvm/issues/14070). Going through the list of points, it looks to me that the `relax.layout_transform` op is intendedly used by some other passes and is not directly codegened (lowered to TIR). So at this moment, a Relax function with `relax.layout_transform` cannot be build. cc @psrivas2 who manages the tracking thread for more followups ---
**Author:** ![psrivas2](https://discuss.tvm.apache.org/user_avatar/discuss.tvm.apache.org/psrivas2/32/3767_2.png) [@psrivas2](https://discuss.tvm.apache.org/u/psrivas2)\ **Post date:** [June 26, 2023, 4:57pm UTC](https://discuss.tvm.apache.org/t/codegen-support-for-relax-layout-transform/15187/4 "2023-06-26T16:57:27Z")
@rutkoor yes, the layout transform op is missing legalization. Let me send a PR that implements layout\_transform legalization. ---
**Author:** ![rutkoor](https://discuss.tvm.apache.org/letter_avatar_proxy/v4/letter/r/90ced4/32.png) [@rutkoor](https://discuss.tvm.apache.org/u/rutkoor)\ **Post date:** [June 27, 2023, 6:15am UTC](https://discuss.tvm.apache.org/t/codegen-support-for-relax-layout-transform/15187/5 "2023-06-27T06:15:11Z")
Thanks @Hzfengsy, @MasterJH5574, @psrivas2 for quick reply. I will wait for the PR that implements layout\_transform legalization. ---
**Author:** ![psrivas2](https://discuss.tvm.apache.org/user_avatar/discuss.tvm.apache.org/psrivas2/32/3767_2.png) [@psrivas2](https://discuss.tvm.apache.org/u/psrivas2)\ **Post date:** [June 29, 2023, 6:19pm UTC](https://discuss.tvm.apache.org/t/codegen-support-for-relax-layout-transform/15187/6 "2023-06-29T18:19:55Z")
PR to add legalization for layout\_transform was sent out today [https://github.com/apache/tvm/pull/15184](https://github.com/apache/tvm/pull/15184) Thanks @leshenj15! ---
**Author:** ![rutkoor](https://discuss.tvm.apache.org/letter_avatar_proxy/v4/letter/r/90ced4/32.png) [@rutkoor](https://discuss.tvm.apache.org/u/rutkoor)\ **Post date:** [June 30, 2023, 5:23am UTC](https://discuss.tvm.apache.org/t/codegen-support-for-relax-layout-transform/15187/7 "2023-06-30T05:23:56Z")
Thanks, @leshenj15 and @psrivas2 for adding legalization support for layout\_transform.