You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by GitBox <gi...@apache.org> on 2022/07/28 10:22:00 UTC
[GitHub] [tvm] ashutosh-arm opened a new pull request, #12215: Pass that removes reshapes post LowerTE
ashutosh-arm opened a new pull request, #12215:
URL: https://github.com/apache/tvm/pull/12215
Introduces a Pass for removing intermediate reshapes post
LowerTE() in AOT compiler. This commit adds pass specific
tests and updates usmp generated workspace pools due to
reduction in number of allocations post reshape removals.
Note: this pass at present does not support first reshape
appearing in the graph. If seen as a useful case, it can be
added in the future.
cc: @manupa-arm @grant-arm
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] ashutosh-arm commented on a diff in pull request #12215: Pass that removes reshapes post LowerTE
Posted by GitBox <gi...@apache.org>.
ashutosh-arm commented on code in PR #12215:
URL: https://github.com/apache/tvm/pull/12215#discussion_r935291033
##########
include/tvm/relay/transform.h:
##########
@@ -580,6 +580,14 @@ TVM_DLL Pass AnnotateUsedMemory();
*/
TVM_DLL Pass CapturePostDfsIndexInSpans();
+/*!
+ * \brief Remove reshapes after lowering the graph.
+ *
+ *
+ * \return The pass.
+ */
+TVM_DLL Pass RemoveReshapes();
Review Comment:
ACK
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] ashutosh-arm commented on a diff in pull request #12215: Pass that removes reshapes post LowerTE
Posted by GitBox <gi...@apache.org>.
ashutosh-arm commented on code in PR #12215:
URL: https://github.com/apache/tvm/pull/12215#discussion_r940011392
##########
tests/python/relay/backend/test_pass_remove_standalone_reshapes.py:
##########
@@ -0,0 +1,260 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License. You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied. See the License for the
+# specific language governing permissions and limitations
+# under the License.
+
+# Exercises the RemoveStandaloneReshapes pass.
+
+import tvm
+from tvm import relay
+from tvm.relay.expr_functor import ExprMutator
+import tvm.testing
+from tvm.script import tir as T
+
+
+HOST_DEVICE = tvm.device("cpu")
+HOST_TARGET = tvm.target.Target("llvm")
+
+CPU_DEVICE = tvm.device("cpu")
+CPU_TARGET = tvm.target.Target("llvm").with_host(HOST_TARGET)
+
+CPU = tvm.target.VirtualDevice(CPU_DEVICE, CPU_TARGET) # device_type=1
+
+
+RemoveStandaloneReshapes = tvm._ffi.get_global_func("relay._transform.RemoveStandaloneReshapes")
+
+
+class MarkReshapeOnlyMutator(ExprMutator):
+ """A pass for marking call_lowered as ReshapeOnly where reshapes exist unfused"""
+
+ def __init__(self):
+ ExprMutator.__init__(self)
+
+ def visit_call(self, call):
+ if isinstance(call.args[0], tvm.ir.GlobalVar) and "reshape" in call.args[0].name_hint:
+ # attrs = {"relay_attrs" : {"relay.reshape_only" : 1}}
+ dict_attrs = tvm.ir.make_node("DictAttrs", **{"relay.reshape_only": 1})
+ attrs = tvm.ir.make_node(
+ "relay.attrs.CallLoweredAttrs", **{"metadata": {"relay_attrs": dict_attrs}}
+ )
+ return relay.Call(call.op, call.args, attrs)
+ return super().visit_call(call)
+
+
+# Reshape should not be removed if its the first layer in the network
+def test_first_reshape():
+ mod = tvm.ir.IRModule()
+
+ @T.prim_func
+ def reshape_primfunc(a: T.handle, d: T.handle) -> None:
+ A = T.match_buffer(a, [128, 128])
+ D = T.match_buffer(d, [128, 128])
+
+ for i, j in T.grid(128, 128):
+ D[i, j] = A[i, j]
+
+ metatable = {"VirtualDevice": [CPU]}
+ reshape_ty = relay.FuncType(
+ [
+ relay.TensorType((128, 128), "float32"),
+ ],
+ relay.TensorType((128, 128), "float32"),
+ )
+
+ reshape_gv = relay.GlobalVar("reshape", type_annot=reshape_ty)
+ mod[reshape_gv] = reshape_primfunc
+ mod = tvm.parser.parse(
+ """
+ #[version = "0.0.5"]
+ def @main(%x {virtual_device=meta[VirtualDevice][0]}: Tensor[(128, 128), float32],
+ virtual_device=meta[VirtualDevice][0]) {
+ %1 = call_lowered(@reshape, (%x,) );
+ let %x_14: Tensor[(128, 128), float32] = on_device(%1, virtual_device=meta[VirtualDevice][0], constrain_result=True);
+ %x_14
+ }
+ """,
+ "from_string",
+ mod,
+ metatable,
+ )
+
+ mod["main"] = MarkReshapeOnlyMutator().visit(mod["main"])
+ mod = RemoveStandaloneReshapes()(mod)
+ reshapes_present = any(["reshape" in gv.name_hint for gv in mod.get_global_vars()])
+ assert reshapes_present, "Reshape should have been removed."
+ return
+
+
+# When reshape layer is the last one in the network
Review Comment:
Sure, I will make a note of it.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] ashutosh-arm commented on a diff in pull request #12215: Pass that removes reshapes post LowerTE
Posted by GitBox <gi...@apache.org>.
ashutosh-arm commented on code in PR #12215:
URL: https://github.com/apache/tvm/pull/12215#discussion_r933262523
##########
src/relay/transforms/remove_standalone_reshapes.cc:
##########
@@ -0,0 +1,122 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License. You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied. See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+/*!
+ * \file remove_standalone_reshapes.cc
+ * \brief This file contains the Relay pass for removing unfused reshapes from lowered graph.
+ * InferType() cannot be invoked after calling this pass as it removes reshapes from the call
+ * graph. Many targets only need buffer addresses irrespective of the shapes of them. This makes
+ * reshapes symbolic once the graph has been lowered. Reshape removal results into smaller code
+ * size and reduced buffer allocations. It opens up opportunities of operator fusion in the target
+ * backend. Thus, consequently, it improves the performance of the inference.
Review Comment:
I have moved this description to transform.h where it makes more sense.
##########
src/relay/transforms/remove_standalone_reshapes.cc:
##########
@@ -0,0 +1,122 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License. You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied. See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+/*!
+ * \file remove_standalone_reshapes.cc
Review Comment:
ACK
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] ashutosh-arm commented on a diff in pull request #12215: Pass that removes reshapes post LowerTE
Posted by GitBox <gi...@apache.org>.
ashutosh-arm commented on code in PR #12215:
URL: https://github.com/apache/tvm/pull/12215#discussion_r937519497
##########
src/relay/backend/aot_executor_codegen.cc:
##########
@@ -1096,6 +1096,12 @@ class AOTExecutorCodegen : public MixedModeVisitor {
tec::UpdateFunctionMetadata(func, this->function_metadata_, workspace_byte_alignment);
})(mod);
+ transform::PassContext pass_ctx = transform::PassContext::Current();
+ bool enable_remove_reshapes =
+ pass_ctx->GetConfig<Bool>("relay.RemoveStandaloneReshapes.enable", Bool(true)).value();
Review Comment:
ACK
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] manupa-arm commented on a diff in pull request #12215: Pass that removes reshapes post LowerTE
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on code in PR #12215:
URL: https://github.com/apache/tvm/pull/12215#discussion_r932045665
##########
include/tvm/relay/transform.h:
##########
@@ -580,6 +580,14 @@ TVM_DLL Pass AnnotateUsedMemory();
*/
TVM_DLL Pass CapturePostDfsIndexInSpans();
+/*!
+ * \brief Remove reshapes after lowering the graph.
+ *
+ *
+ * \return The pass.
+ */
+TVM_DLL Pass RemoveReshapes();
Review Comment:
Lets rename this to RemoveStandaloneReshapes
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] ashutosh-arm commented on a diff in pull request #12215: Pass that removes reshapes post LowerTE
Posted by GitBox <gi...@apache.org>.
ashutosh-arm commented on code in PR #12215:
URL: https://github.com/apache/tvm/pull/12215#discussion_r937512595
##########
src/relay/transforms/remove_standalone_reshapes.cc:
##########
@@ -0,0 +1,120 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License. You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied. See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+/*!
+ * \file src/relay/transforms/remove_standalone_reshapes.cc
+ * \brief This file contains the Relay pass for removing unfused reshapes from lowered graph.
+ */
+
+#include <tvm/relay/expr_functor.h>
+#include <tvm/relay/transform.h>
+
+#include "../op/call/call.h"
+#include "../op/memory/on_device.h"
+
+namespace tvm {
+namespace relay {
+
+TVM_REGISTER_PASS_CONFIG_OPTION("relay.RemoveStandaloneReshapes.enable", Bool);
+/*! Removes reshapes right after LowerTE. Removes preceding on_device calls
+ * while removing reshapes.
+ */
+class RemoveStandaloneReshapesMutator : public MixedModeMutator {
+ public:
+ explicit RemoveStandaloneReshapesMutator(IRModule& mod) : ir_module_(mod) {}
+
+ using MixedModeMutator::VisitExpr_;
+
+ /*! * \brief Generated map of let variables to preceding CallLowered */
+ Expr VisitExpr_(const LetNode* let) final {
+ Let ret_let;
+ Var var = Downcast<Var>(this->Mutate(let->var));
+ auto value = this->Mutate(let->value);
+ if (auto* on_device_call = value.as<CallNode>()) {
+ OnDeviceProps on_device_props = GetOnDeviceProps(on_device_call);
+ if (on_device_props.body.defined() && on_device_props.body->IsInstance<CallNode>()) {
+ const Call call_lowered = Downcast<Call>(on_device_props.body);
+ if (call_lowered.defined() && call_lowered->op.same_as(CallLoweredOp())) {
+ let_var_to_call_lowered_.Set(var, call_lowered);
+ }
+ }
+ }
+ auto body = this->Mutate(let->body);
+ return WithFields(GetRef<Let>(let), var, value, body);
+ }
+
+ /*! * \brief Returns preceding CallLowered when call is a CallLowered(Reshape) */
Review Comment:
Graph contains let nodes in between the call_lowered(). I've included the following piece as part of the Rewrite_() as well.
```
/*
%1 = call_lowered(@tvmgen_default_non_reshape_function, %input, ...);
let %x: = on_device(%1, ...);
%2 = (%x,);
%3 = call_lowered(@tvmgen_default_fused_reshape, %2, ...,
"relay_attrs"=__dict__="relay.reshape_only"=1, ...);
*/
```
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] manupa-arm commented on a diff in pull request #12215: Pass that removes reshapes post LowerTE
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on code in PR #12215:
URL: https://github.com/apache/tvm/pull/12215#discussion_r932054959
##########
src/relay/transforms/remove_reshapes.cc:
##########
@@ -0,0 +1,116 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License. You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied. See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+
+/*!
+ * \file remove_reshapes.cc
+ * \brief Relay pass for removing reshapes from lowered graph.
+ */
+
+#include <tvm/relay/expr_functor.h>
+#include <tvm/relay/transform.h>
+
+#include "../op/call/call.h"
+#include "../op/memory/on_device.h"
+
+namespace tvm {
+namespace relay {
+
+TVM_REGISTER_PASS_CONFIG_OPTION("relay.RemoveReshapes", Bool);
+/*! Removes reshapes right after LowerTE. Removes preceding on_device calls
+ * while removing reshapes.
+ */
+class RemoveReshapesMutator : public MixedModeMutator {
+ public:
+ explicit RemoveReshapesMutator(IRModule& mod) : ir_module_(mod) {}
+
+ using MixedModeMutator::VisitExpr_;
+
+ Expr VisitExpr_(const LetNode* let) final {
+ Let ret_let;
+ Var var = Downcast<Var>(this->Mutate(let->var));
+ auto value = this->Mutate(let->value);
+ if (auto* on_device_call = value.as<CallNode>()) {
+ OnDeviceProps on_device_props = GetOnDeviceProps(on_device_call);
+ if (on_device_props.body.defined() && on_device_props.body->IsInstance<CallNode>()) {
+ const Call call_lowered = Downcast<Call>(on_device_props.body);
+ if (call_lowered.defined() && call_lowered->op.same_as(CallLoweredOp())) {
+ let_var_to_call_lowered_.Set(var, call_lowered);
+ }
+ }
+ }
+ auto body = this->Mutate(let->body);
+ return WithFields(GetRef<Let>(let), var, value, body);
+ }
+
+ Expr Rewrite_(const CallNode* call, const Expr& post) final {
+ /*
Review Comment:
nit : might worth explaining what the reader should get out of this block
##########
src/relay/transforms/remove_reshapes.cc:
##########
@@ -0,0 +1,116 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License. You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied. See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+
+/*!
+ * \file remove_reshapes.cc
+ * \brief Relay pass for removing reshapes from lowered graph.
+ */
+
+#include <tvm/relay/expr_functor.h>
+#include <tvm/relay/transform.h>
+
+#include "../op/call/call.h"
+#include "../op/memory/on_device.h"
+
+namespace tvm {
+namespace relay {
+
+TVM_REGISTER_PASS_CONFIG_OPTION("relay.RemoveReshapes", Bool);
Review Comment:
Please use something like relay.RemoveStandaloneReshapes.enable or relay.use_remove_standalone_reshapes.
I personally prefer the former as it creates a namespace of any future options of the pass.
##########
include/tvm/relay/transform.h:
##########
@@ -580,6 +580,14 @@ TVM_DLL Pass AnnotateUsedMemory();
*/
TVM_DLL Pass CapturePostDfsIndexInSpans();
+/*!
+ * \brief Remove reshapes after lowering the graph.
+ *
+ *
+ * \return The pass.
+ */
+TVM_DLL Pass RemoveReshapes();
Review Comment:
Lets rename this to remove StandaloneReshapes
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] ashutosh-arm commented on a diff in pull request #12215: Pass that removes reshapes post LowerTE
Posted by GitBox <gi...@apache.org>.
ashutosh-arm commented on code in PR #12215:
URL: https://github.com/apache/tvm/pull/12215#discussion_r935299756
##########
src/relay/transforms/remove_standalone_reshapes.cc:
##########
@@ -0,0 +1,122 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License. You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied. See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+/*!
+ * \file remove_standalone_reshapes.cc
+ * \brief This file contains the Relay pass for removing unfused reshapes from lowered graph.
+ * InferType() cannot be invoked after calling this pass as it removes reshapes from the call
+ * graph. Many targets only need buffer addresses irrespective of the shapes of them. This makes
+ * reshapes symbolic once the graph has been lowered. Reshape removal results into smaller code
+ * size and reduced buffer allocations. It opens up opportunities of operator fusion in the target
+ * backend. Thus, consequently, it improves the performance of the inference.
+ */
+
+#include <tvm/relay/expr_functor.h>
+#include <tvm/relay/transform.h>
+
+#include "../op/call/call.h"
+#include "../op/memory/on_device.h"
+
+namespace tvm {
+namespace relay {
+
+TVM_REGISTER_PASS_CONFIG_OPTION("relay.RemoveStandaloneReshapes.enable", Bool);
+/*! Removes reshapes right after LowerTE. Removes preceding on_device calls
+ * while removing reshapes.
+ */
+class RemoveStandaloneReshapesMutator : public MixedModeMutator {
+ public:
+ explicit RemoveStandaloneReshapesMutator(IRModule& mod) : ir_module_(mod) {}
+
+ using MixedModeMutator::VisitExpr_;
+
+ /*! * \brief Generated map of let variables to preceding CallLowered */
+ Expr VisitExpr_(const LetNode* let) final {
+ Let ret_let;
+ Var var = Downcast<Var>(this->Mutate(let->var));
+ auto value = this->Mutate(let->value);
+ if (auto* on_device_call = value.as<CallNode>()) {
+ OnDeviceProps on_device_props = GetOnDeviceProps(on_device_call);
+ if (on_device_props.body.defined() && on_device_props.body->IsInstance<CallNode>()) {
+ const Call call_lowered = Downcast<Call>(on_device_props.body);
+ if (call_lowered.defined() && call_lowered->op.same_as(CallLoweredOp())) {
+ let_var_to_call_lowered_.Set(var, call_lowered);
+ }
+ }
+ }
+ auto body = this->Mutate(let->body);
+ return WithFields(GetRef<Let>(let), var, value, body);
Review Comment:
Discussed this offline with @lhutton1 to study the impact of not having ExpandANormalForm() post lowering. It seems the LowerTensorExprMutator takes care of expanding the network during lowering. So, introducing ExpandANormalForm() should not make any difference.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] lhutton1 commented on a diff in pull request #12215: Pass that removes reshapes post LowerTE
Posted by GitBox <gi...@apache.org>.
lhutton1 commented on code in PR #12215:
URL: https://github.com/apache/tvm/pull/12215#discussion_r932243749
##########
src/relay/transforms/remove_standalone_reshapes.cc:
##########
@@ -0,0 +1,122 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License. You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied. See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+/*!
+ * \file remove_standalone_reshapes.cc
+ * \brief This file contains the Relay pass for removing unfused reshapes from lowered graph.
+ * InferType() cannot be invoked after calling this pass as it removes reshapes from the call
+ * graph. Many targets only need buffer addresses irrespective of the shapes of them. This makes
+ * reshapes symbolic once the graph has been lowered. Reshape removal results into smaller code
+ * size and reduced buffer allocations. It opens up opportunities of operator fusion in the target
+ * backend. Thus, consequently, it improves the performance of the inference.
+ */
+
+#include <tvm/relay/expr_functor.h>
+#include <tvm/relay/transform.h>
+
+#include "../op/call/call.h"
+#include "../op/memory/on_device.h"
+
+namespace tvm {
+namespace relay {
+
+TVM_REGISTER_PASS_CONFIG_OPTION("relay.RemoveStandaloneReshapes.enable", Bool);
+/*! Removes reshapes right after LowerTE. Removes preceding on_device calls
+ * while removing reshapes.
+ */
+class RemoveStandaloneReshapesMutator : public MixedModeMutator {
+ public:
+ explicit RemoveStandaloneReshapesMutator(IRModule& mod) : ir_module_(mod) {}
+
+ using MixedModeMutator::VisitExpr_;
+
+ /*! * \brief Generated map of let variables to preceding CallLowered */
+ Expr VisitExpr_(const LetNode* let) final {
+ Let ret_let;
+ Var var = Downcast<Var>(this->Mutate(let->var));
+ auto value = this->Mutate(let->value);
+ if (auto* on_device_call = value.as<CallNode>()) {
+ OnDeviceProps on_device_props = GetOnDeviceProps(on_device_call);
+ if (on_device_props.body.defined() && on_device_props.body->IsInstance<CallNode>()) {
+ const Call call_lowered = Downcast<Call>(on_device_props.body);
+ if (call_lowered.defined() && call_lowered->op.same_as(CallLoweredOp())) {
+ let_var_to_call_lowered_.Set(var, call_lowered);
+ }
+ }
+ }
+ auto body = this->Mutate(let->body);
+ return WithFields(GetRef<Let>(let), var, value, body);
Review Comment:
We should use `ExpandANormalForm` here to avoid stack overflows. See similar expansion in: https://github.com/apache/tvm/blob/main/src/relay/backend/contrib/ethosu/codegen.cc#L60
##########
src/relay/transforms/remove_standalone_reshapes.cc:
##########
@@ -0,0 +1,122 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License. You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied. See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+/*!
+ * \file remove_standalone_reshapes.cc
Review Comment:
nit: `src/relay/transforms/remove_standalone_reshapes.cc`
##########
src/relay/transforms/remove_standalone_reshapes.cc:
##########
@@ -0,0 +1,122 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License. You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied. See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+/*!
+ * \file remove_standalone_reshapes.cc
+ * \brief This file contains the Relay pass for removing unfused reshapes from lowered graph.
+ * InferType() cannot be invoked after calling this pass as it removes reshapes from the call
+ * graph. Many targets only need buffer addresses irrespective of the shapes of them. This makes
+ * reshapes symbolic once the graph has been lowered. Reshape removal results into smaller code
+ * size and reduced buffer allocations. It opens up opportunities of operator fusion in the target
+ * backend. Thus, consequently, it improves the performance of the inference.
Review Comment:
Shall we use this description in transform.h as well?
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] lhutton1 merged pull request #12215: Pass that removes reshapes post LowerTE
Posted by GitBox <gi...@apache.org>.
lhutton1 merged PR #12215:
URL: https://github.com/apache/tvm/pull/12215
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] areusch commented on a diff in pull request #12215: Pass that removes reshapes post LowerTE
Posted by GitBox <gi...@apache.org>.
areusch commented on code in PR #12215:
URL: https://github.com/apache/tvm/pull/12215#discussion_r936109474
##########
src/relay/backend/aot_executor_codegen.cc:
##########
@@ -1096,6 +1096,12 @@ class AOTExecutorCodegen : public MixedModeVisitor {
tec::UpdateFunctionMetadata(func, this->function_metadata_, workspace_byte_alignment);
})(mod);
+ transform::PassContext pass_ctx = transform::PassContext::Current();
+ bool enable_remove_reshapes =
+ pass_ctx->GetConfig<Bool>("relay.RemoveStandaloneReshapes.enable", Bool(true)).value();
Review Comment:
i think can we use snake_case here to match the other pass options?
##########
src/relay/transforms/remove_standalone_reshapes.cc:
##########
@@ -0,0 +1,120 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License. You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied. See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+/*!
+ * \file src/relay/transforms/remove_standalone_reshapes.cc
+ * \brief This file contains the Relay pass for removing unfused reshapes from lowered graph.
+ */
+
+#include <tvm/relay/expr_functor.h>
+#include <tvm/relay/transform.h>
+
+#include "../op/call/call.h"
+#include "../op/memory/on_device.h"
+
+namespace tvm {
+namespace relay {
+
+TVM_REGISTER_PASS_CONFIG_OPTION("relay.RemoveStandaloneReshapes.enable", Bool);
+/*! Removes reshapes right after LowerTE. Removes preceding on_device calls
+ * while removing reshapes.
+ */
+class RemoveStandaloneReshapesMutator : public MixedModeMutator {
+ public:
+ explicit RemoveStandaloneReshapesMutator(IRModule& mod) : ir_module_(mod) {}
+
+ using MixedModeMutator::VisitExpr_;
+
+ /*! * \brief Generated map of let variables to preceding CallLowered */
+ Expr VisitExpr_(const LetNode* let) final {
+ Let ret_let;
+ Var var = Downcast<Var>(this->Mutate(let->var));
+ auto value = this->Mutate(let->value);
+ if (auto* on_device_call = value.as<CallNode>()) {
+ OnDeviceProps on_device_props = GetOnDeviceProps(on_device_call);
+ if (on_device_props.body.defined() && on_device_props.body->IsInstance<CallNode>()) {
+ const Call call_lowered = Downcast<Call>(on_device_props.body);
+ if (call_lowered.defined() && call_lowered->op.same_as(CallLoweredOp())) {
+ let_var_to_call_lowered_.Set(var, call_lowered);
+ }
+ }
+ }
+ auto body = this->Mutate(let->body);
+ return WithFields(GetRef<Let>(let), var, value, body);
+ }
+
+ /*! * \brief Returns preceding CallLowered when call is a CallLowered(Reshape) */
Review Comment:
i'm probably missing some context here, but what about just returning the args to reshape()?
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] ashutosh-arm commented on a diff in pull request #12215: Pass that removes reshapes post LowerTE
Posted by GitBox <gi...@apache.org>.
ashutosh-arm commented on code in PR #12215:
URL: https://github.com/apache/tvm/pull/12215#discussion_r935292005
##########
src/relay/transforms/remove_reshapes.cc:
##########
@@ -0,0 +1,116 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License. You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied. See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+
+/*!
+ * \file remove_reshapes.cc
+ * \brief Relay pass for removing reshapes from lowered graph.
+ */
+
+#include <tvm/relay/expr_functor.h>
+#include <tvm/relay/transform.h>
+
+#include "../op/call/call.h"
+#include "../op/memory/on_device.h"
+
+namespace tvm {
+namespace relay {
+
+TVM_REGISTER_PASS_CONFIG_OPTION("relay.RemoveReshapes", Bool);
+/*! Removes reshapes right after LowerTE. Removes preceding on_device calls
+ * while removing reshapes.
+ */
+class RemoveReshapesMutator : public MixedModeMutator {
+ public:
+ explicit RemoveReshapesMutator(IRModule& mod) : ir_module_(mod) {}
+
+ using MixedModeMutator::VisitExpr_;
+
+ Expr VisitExpr_(const LetNode* let) final {
+ Let ret_let;
+ Var var = Downcast<Var>(this->Mutate(let->var));
+ auto value = this->Mutate(let->value);
+ if (auto* on_device_call = value.as<CallNode>()) {
+ OnDeviceProps on_device_props = GetOnDeviceProps(on_device_call);
+ if (on_device_props.body.defined() && on_device_props.body->IsInstance<CallNode>()) {
+ const Call call_lowered = Downcast<Call>(on_device_props.body);
+ if (call_lowered.defined() && call_lowered->op.same_as(CallLoweredOp())) {
+ let_var_to_call_lowered_.Set(var, call_lowered);
+ }
+ }
+ }
+ auto body = this->Mutate(let->body);
+ return WithFields(GetRef<Let>(let), var, value, body);
+ }
+
+ Expr Rewrite_(const CallNode* call, const Expr& post) final {
+ /*
Review Comment:
Moved this description finally to transform.h
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] ashutosh-arm commented on a diff in pull request #12215: Pass that removes reshapes post LowerTE
Posted by GitBox <gi...@apache.org>.
ashutosh-arm commented on code in PR #12215:
URL: https://github.com/apache/tvm/pull/12215#discussion_r935291413
##########
src/relay/transforms/remove_reshapes.cc:
##########
@@ -0,0 +1,116 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License. You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied. See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+
+/*!
+ * \file remove_reshapes.cc
+ * \brief Relay pass for removing reshapes from lowered graph.
+ */
+
+#include <tvm/relay/expr_functor.h>
+#include <tvm/relay/transform.h>
+
+#include "../op/call/call.h"
+#include "../op/memory/on_device.h"
+
+namespace tvm {
+namespace relay {
+
+TVM_REGISTER_PASS_CONFIG_OPTION("relay.RemoveReshapes", Bool);
Review Comment:
Done
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] lhutton1 commented on a diff in pull request #12215: Pass that removes reshapes post LowerTE
Posted by GitBox <gi...@apache.org>.
lhutton1 commented on code in PR #12215:
URL: https://github.com/apache/tvm/pull/12215#discussion_r938922796
##########
tests/python/relay/backend/test_pass_remove_standalone_reshapes.py:
##########
@@ -0,0 +1,260 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License. You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied. See the License for the
+# specific language governing permissions and limitations
+# under the License.
+
+# Exercises the RemoveStandaloneReshapes pass.
+
+import tvm
+from tvm import relay
+from tvm.relay.expr_functor import ExprMutator
+import tvm.testing
+from tvm.script import tir as T
+
+
+HOST_DEVICE = tvm.device("cpu")
+HOST_TARGET = tvm.target.Target("llvm")
+
+CPU_DEVICE = tvm.device("cpu")
+CPU_TARGET = tvm.target.Target("llvm").with_host(HOST_TARGET)
+
+CPU = tvm.target.VirtualDevice(CPU_DEVICE, CPU_TARGET) # device_type=1
+
+
+RemoveStandaloneReshapes = tvm._ffi.get_global_func("relay._transform.RemoveStandaloneReshapes")
+
+
+class MarkReshapeOnlyMutator(ExprMutator):
+ """A pass for marking call_lowered as ReshapeOnly where reshapes exist unfused"""
+
+ def __init__(self):
+ ExprMutator.__init__(self)
+
+ def visit_call(self, call):
+ if isinstance(call.args[0], tvm.ir.GlobalVar) and "reshape" in call.args[0].name_hint:
+ # attrs = {"relay_attrs" : {"relay.reshape_only" : 1}}
+ dict_attrs = tvm.ir.make_node("DictAttrs", **{"relay.reshape_only": 1})
+ attrs = tvm.ir.make_node(
+ "relay.attrs.CallLoweredAttrs", **{"metadata": {"relay_attrs": dict_attrs}}
+ )
+ return relay.Call(call.op, call.args, attrs)
+ return super().visit_call(call)
+
+
+# Reshape should not be removed if its the first layer in the network
+def test_first_reshape():
+ mod = tvm.ir.IRModule()
+
+ @T.prim_func
+ def reshape_primfunc(a: T.handle, d: T.handle) -> None:
+ A = T.match_buffer(a, [128, 128])
+ D = T.match_buffer(d, [128, 128])
+
+ for i, j in T.grid(128, 128):
+ D[i, j] = A[i, j]
+
+ metatable = {"VirtualDevice": [CPU]}
+ reshape_ty = relay.FuncType(
+ [
+ relay.TensorType((128, 128), "float32"),
+ ],
+ relay.TensorType((128, 128), "float32"),
+ )
+
+ reshape_gv = relay.GlobalVar("reshape", type_annot=reshape_ty)
+ mod[reshape_gv] = reshape_primfunc
+ mod = tvm.parser.parse(
+ """
+ #[version = "0.0.5"]
+ def @main(%x {virtual_device=meta[VirtualDevice][0]}: Tensor[(128, 128), float32],
+ virtual_device=meta[VirtualDevice][0]) {
+ %1 = call_lowered(@reshape, (%x,) );
+ let %x_14: Tensor[(128, 128), float32] = on_device(%1, virtual_device=meta[VirtualDevice][0], constrain_result=True);
+ %x_14
+ }
+ """,
+ "from_string",
+ mod,
+ metatable,
+ )
+
+ mod["main"] = MarkReshapeOnlyMutator().visit(mod["main"])
+ mod = RemoveStandaloneReshapes()(mod)
+ reshapes_present = any(["reshape" in gv.name_hint for gv in mod.get_global_vars()])
+ assert reshapes_present, "Reshape should have been removed."
+ return
+
+
+# When reshape layer is the last one in the network
Review Comment:
Nit: we should probably lint new tests being added to help with the efforts in https://github.com/apache/tvm/issues/11414, perhaps in a follow up?
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] lhutton1 commented on pull request #12215: Pass that removes reshapes post LowerTE
Posted by GitBox <gi...@apache.org>.
lhutton1 commented on PR #12215:
URL: https://github.com/apache/tvm/pull/12215#issuecomment-1207872900
Thanks @ashutosh-arm @manupa-arm @areusch!
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org