You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by "echuraev (via GitHub)" <gi...@apache.org> on 2023/06/21 15:33:48 UTC
[GitHub] [tvm] echuraev opened a new pull request, #15137: [Relay] Introduce arguments limit to FuseOps pass
echuraev opened a new pull request, #15137:
URL: https://github.com/apache/tvm/pull/15137
In PR #8313 a parameter `max_function_args` was introduced. It leads to limit number of function argument and in case when this value is exceeded then concatenation layer is split to a several concat operations.
I faced a problem on Adreno GPU that for kernel with big number of arguments the enqueueNDRange was crashed without any errors. The problem appeared because of the huge number of arguments. But in this case not only concat layer was a root cause of the problem. Also after fusing several operations the final functions had a big number of arguments.
As it was discussed in #8313, adding a limitation on the number of function arguments to the FuseOps pass might be a good improvement. In this PR I introduced such mechanism for limitation number of function arguments for FuseOps pass and add an arguments limit to OpenCL devices at 128 parameters.
cc: @Hzfengsy, @masahi, @csullivan
--
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] tvm-bot commented on pull request #15137: [Relay] Introduce arguments limit to FuseOps pass
Posted by "tvm-bot (via GitHub)" <gi...@apache.org>.
tvm-bot commented on PR #15137:
URL: https://github.com/apache/tvm/pull/15137#issuecomment-1601066405
<!---bot-comment-->
Thanks for contributing to TVM! Please refer to the contributing guidelines https://tvm.apache.org/docs/contribute/ for useful information and tips. Please request code reviews from [Reviewers](https://github.com/apache/incubator-tvm/blob/master/CONTRIBUTORS.md#reviewers) by @-ing them in a comment.
<!--bot-comment-ccs-start-->
* cc @shingjan <sub>See [#10317](https://github.com/apache/tvm/issues/10317) for details</sub><!--bot-comment-ccs-end-->
<sub>Generated by [tvm-bot](https://github.com/apache/tvm/blob/main/ci/README.md#github-actions)</sub>
--
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] masahi merged pull request #15137: [Relay] Introduce arguments limit to FuseOps pass
Posted by "masahi (via GitHub)" <gi...@apache.org>.
masahi merged PR #15137:
URL: https://github.com/apache/tvm/pull/15137
--
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] echuraev commented on pull request #15137: [Relay] Introduce arguments limit to FuseOps pass
Posted by "echuraev (via GitHub)" <gi...@apache.org>.
echuraev commented on PR #15137:
URL: https://github.com/apache/tvm/pull/15137#issuecomment-1643434216
@masahi thank you for your review and help! Applied all comments.
--
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] csullivan commented on a diff in pull request #15137: [Relay] Introduce arguments limit to FuseOps pass
Posted by "csullivan (via GitHub)" <gi...@apache.org>.
csullivan commented on code in PR #15137:
URL: https://github.com/apache/tvm/pull/15137#discussion_r1266998446
##########
src/target/target_kind.cc:
##########
@@ -365,6 +365,11 @@ TVM_REGISTER_TARGET_KIND("opencl", kDLOpenCL)
.add_attr_option<Integer>("max_num_threads", Integer(256))
.add_attr_option<Integer>("thread_warp_size", Integer(1))
.add_attr_option<Integer>("texture_spatial_limit", Integer(16384))
+ // Faced that Qualcomm OpenCL runtime was crashed without any error message in
+ // case when the number of kernel arguments was pretty big. OpenCL doesn't
+ // specify any limitations on the number of kernel arguments. max_function_args
+ // equals to 128 looks like a reasonable number of kernel arguments.
+ .add_attr_option<Integer>("max_function_args", Integer(128))
Review Comment:
We discussed offline considering making this only the case for Adreno instead of any opencl. The issue is there appears to be a limitation in FuseOps exposed by this PR where the remainder of ops (mod max_function_args) not fused do not fuse together and remain individual primfuncs. @echuraev will follow up with an issue to track this exposed limitation of FuseOps.
--
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] echuraev commented on pull request #15137: [Relay] Introduce arguments limit to FuseOps pass
Posted by "echuraev (via GitHub)" <gi...@apache.org>.
echuraev commented on PR #15137:
URL: https://github.com/apache/tvm/pull/15137#issuecomment-1623840171
@csullivan, @masahi, @Hzfengsy now all functional and performance issues are fixed. Could you please review this PR?
--
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] Hzfengsy commented on pull request #15137: [Relay] Introduce arguments limit to FuseOps pass
Posted by "Hzfengsy (via GitHub)" <gi...@apache.org>.
Hzfengsy commented on PR #15137:
URL: https://github.com/apache/tvm/pull/15137#issuecomment-1628779796
@masahi Could you please review this one?
--
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] masahi commented on a diff in pull request #15137: [Relay] Introduce arguments limit to FuseOps pass
Posted by "masahi (via GitHub)" <gi...@apache.org>.
masahi commented on code in PR #15137:
URL: https://github.com/apache/tvm/pull/15137#discussion_r1258806135
##########
include/tvm/relay/transform.h:
##########
@@ -120,9 +120,12 @@ TVM_DLL Pass FoldConstant(bool fold_qnn = false);
/*!
* \brief Split function with huge number of arguments to smaller pieces.
*
+ * \param max_function_args Maximum number of function arguments. If it is 0 then SplitArgs won't
+ * split funciton.
Review Comment:
typo
##########
src/relay/analysis/graph_partitioner.cc:
##########
@@ -238,6 +320,8 @@ void GraphPartitioner::InitGroups(const IndexedForwardGraph& graph) {
void GraphPartitioner::RunFuse(const IndexedForwardGraph& graph, //
const DominatorTree& post_dom_tree, //
int phase) {
+ IndexedForwardGraph::Node* prev_node = nullptr;
Review Comment:
Describe what is meant by "previous". It is only used for counting args but the name is too generic.
##########
src/relay/analysis/graph_partitioner.h:
##########
@@ -247,6 +255,9 @@ class GraphPartitioner {
void CommitFuse(IndexedForwardGraph::Node* src, IndexedForwardGraph::Node* sink);
size_t CountNodesUptoSink_(IndexedForwardGraph::Node* src, IndexedForwardGraph::Node* sink);
+ size_t CountAdditionalArgs_(const TensorTypeNode* ttype, bool with_strides = true);
Review Comment:
Describe what "Additional" means
##########
src/relay/analysis/graph_partitioner.h:
##########
@@ -247,6 +255,9 @@ class GraphPartitioner {
void CommitFuse(IndexedForwardGraph::Node* src, IndexedForwardGraph::Node* sink);
size_t CountNodesUptoSink_(IndexedForwardGraph::Node* src, IndexedForwardGraph::Node* sink);
+ size_t CountAdditionalArgs_(const TensorTypeNode* ttype, bool with_strides = true);
+ size_t CountArgs_(const tvm::Object* child, const tvm::Object* till_node);
+ size_t CountArgsLimit_(const IndexedForwardGraph::Node* child);
Review Comment:
Describe what this function is for, even though there is already `max_function_args_`.
##########
src/relay/analysis/graph_partitioner.h:
##########
@@ -205,10 +209,14 @@ class GraphPartitioner {
int opt_level_;
/*! \brief The maximum number of operations in one fused function */
size_t max_fuse_depth_;
+ /*! \brief The maximum number of arguments in one fused function */
+ size_t max_function_args_;
/*! \brief The internal groups. */
std::vector<Group*> groups_;
/*! \brief internal field used for deduplication */
std::unordered_set<IndexedForwardGraph::Node*> visited_;
+ /*! \brief internal field used for hashing arguments number for a node */
+ std::unordered_map<const tvm::Object*, size_t> argsMap_;
Review Comment:
args_map_
"arguments number" is a weird phrase
##########
src/relay/analysis/graph_partitioner.h:
##########
@@ -205,10 +209,14 @@ class GraphPartitioner {
int opt_level_;
/*! \brief The maximum number of operations in one fused function */
size_t max_fuse_depth_;
+ /*! \brief The maximum number of arguments in one fused function */
+ size_t max_function_args_;
/*! \brief The internal groups. */
std::vector<Group*> groups_;
/*! \brief internal field used for deduplication */
std::unordered_set<IndexedForwardGraph::Node*> visited_;
+ /*! \brief internal field used for hashing arguments number for a node */
+ std::unordered_map<const tvm::Object*, size_t> argsMap_;
Review Comment:
Also the name `args_map_` doesn't tell what it is for.
--
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] echuraev commented on pull request #15137: [Relay] Introduce arguments limit to FuseOps pass
Posted by "echuraev (via GitHub)" <gi...@apache.org>.
echuraev commented on PR #15137:
URL: https://github.com/apache/tvm/pull/15137#issuecomment-1641944943
@csullivan In the bug #15358 I have described a problem related to the FuseOps pass algorithm. The same problem applicable for the functionality which was introduced in this PR. When we stop fusing because the number of the function arguments was exceeded, then the first PrimFunc will contain the maximum possible number of base functions (if one more base function is added, then the `max_function_args` is exceeded). But other PrimFuncs will contain only one base function in it.
I suppose it is a generic problem in current fusion algorithm and it should be fixed separately in another PR.
--
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] masahi commented on a diff in pull request #15137: [Relay] Introduce arguments limit to FuseOps pass
Posted by "masahi (via GitHub)" <gi...@apache.org>.
masahi commented on code in PR #15137:
URL: https://github.com/apache/tvm/pull/15137#discussion_r1269037848
##########
src/relay/analysis/graph_partitioner.h:
##########
@@ -247,6 +266,23 @@ class GraphPartitioner {
void CommitFuse(IndexedForwardGraph::Node* src, IndexedForwardGraph::Node* sink);
size_t CountNodesUptoSink_(IndexedForwardGraph::Node* src, IndexedForwardGraph::Node* sink);
+ // Count the number of additional arguments. In the case of dynamic shape,
+ // generated function takes several additional arguments, such as the sizes of
+ // the dynamic dimensions and strides.
+ // This function calculates the number of such additional arguments.
+ size_t CountAdditionalArgs_(const TensorTypeNode* ttype, bool with_strides = true);
+ // Calculate the number of arguments for the node.
+ size_t CountArgs_(IndexedForwardGraph::Node* src, const IndexedForwardGraph& graph,
+ bool update_postpone = true);
+ // Count the actual limit of arguments for a generated function.
+ // max_function_args_ specifies the number of maximum function arguments. But
+ // usually, output tensors are also passed to the function as arguments.
+ // Additionally, in the case of dynamic shape, it is necessary to take into
+ // account the number of parameters which specifies the sizes of the dynamic
+ // dimensions.
+ // This function computes the limit of arguments by the following formula:
Review Comment:
the limit of arguments -> the maximum number of arguments
##########
src/relay/analysis/graph_partitioner.h:
##########
@@ -205,10 +213,21 @@ class GraphPartitioner {
int opt_level_;
/*! \brief The maximum number of operations in one fused function */
size_t max_fuse_depth_;
+ /*! \brief The maximum number of arguments in one fused function */
+ size_t max_function_args_;
/*! \brief The internal groups. */
std::vector<Group*> groups_;
/*! \brief internal field used for deduplication */
std::unordered_set<IndexedForwardGraph::Node*> visited_;
+ /*! \brief The map with nodes which were postponed for fusing. */
+ std::unordered_multimap<const IndexedForwardGraph::Node*, IndexedForwardGraph::Node*>
+ postponed_fusing_map_;
+ /*!
+ * \brief Fusing of this node should be postponed till all child nodes will be evaluated.
Review Comment:
till all child nodes **are** evaluated.
--
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] echuraev commented on pull request #15137: [Relay] Introduce arguments limit to FuseOps pass
Posted by "echuraev (via GitHub)" <gi...@apache.org>.
echuraev commented on PR #15137:
URL: https://github.com/apache/tvm/pull/15137#issuecomment-1639802819
Hi @masahi!
Thank you for your review, and sorry for the delay in the response. I found one accuracy problem in previous implementation and it was necessary to reimplement this algorithm almost from scratch. I have added a few new tests and also checked the accuracy on several networks. Now everything is ok. Also, I have applied some of your comments in the new commit.
Could you please take a look at this PR once again?
--
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] echuraev commented on pull request #15137: [Relay] Introduce arguments limit to FuseOps pass
Posted by "echuraev (via GitHub)" <gi...@apache.org>.
echuraev commented on PR #15137:
URL: https://github.com/apache/tvm/pull/15137#issuecomment-1609406063
@tqchen During implementing a limitation on the number of function arguments, I met a problem. In case of tensor with dynamic shape, it is possible that in addition to the parameters of the kernel and indices by dynamic dimensions, strides will be passed to the function. For example, below an example of the OpenCL code which was generated for one of concats from the test on split_args:
```opencl
__kernel void vm_mod_fused_concatenate_kernel(__global float* restrict T_concat,
__global float* restrict p0,
__global float* restrict p1,
__global float* restrict p2,
__global float* restrict p3,
__global float* restrict p4,
int any_dim,
int any_dim_1,
int any_dim_2,
int any_dim_3,
int any_dim_4,
int stride,
int stride_1,
int stride_2,
int stride_3,
int stride_4,
int stride_5,
int stride_6,
int stride_7,
int stride_8,
int stride_9,
int stride_10,
int stride_11,
int stride_12,
int stride_13,
int stride_14,
int stride_15,
int stride_16) {
// ...
}
```
As you can see, OpenCL kernel has three types of arguments: buffers, indices of dynamic dimensions and strides. Unfortunately, I'm not able to get the number of strides in `SplitArgs` and `FuseOps` passes. Strides appear only after the stage of `LowerTE`.
To handle this problem, I added a new attribute to target `max_function_args_dyn` which is used in case if at least one of the tensors has dynamic shape. This attribute should have less value than `max_function_args` and it helps to run topologies, which previously didn't work in TVM. But I have some doubts about this solution. Probably you know the better way, how can I calculate the number of strides which will be passed to the kernel in case of dynamic shapes?
--
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] masahi commented on a diff in pull request #15137: [Relay] Introduce arguments limit to FuseOps pass
Posted by "masahi (via GitHub)" <gi...@apache.org>.
masahi commented on code in PR #15137:
URL: https://github.com/apache/tvm/pull/15137#discussion_r1268853059
##########
src/relay/analysis/graph_partitioner.cc:
##########
@@ -220,7 +225,113 @@ size_t GraphPartitioner::CountFusedNodesWithNewChild(IndexedForwardGraph::Node*
return target->FindRoot()->num_nodes + CountNodesUptoSink_(child, dom_parent);
}
+size_t GraphPartitioner::CountAdditionalArgs_(const TensorTypeNode* ttype, bool with_strides) {
+ size_t any_dims = 0;
+ for (const auto& dim : ttype->shape) {
+ if (dim.as<AnyNode>()) {
+ any_dims++;
+ }
+ }
+ if (with_strides && any_dims > 0) any_dims += ttype->shape.size();
+ return any_dims;
+}
+
+size_t GraphPartitioner::CountArgs_(IndexedForwardGraph::Node* src,
+ const IndexedForwardGraph& graph, bool update_postpone) {
+ std::unordered_set<Group*> visited_groups;
+ Group* gnode = groups_[src->index];
+ ICHECK(gnode != nullptr);
+ auto sum = gnode->args_num;
+ visited_groups.insert(gnode->FindRoot());
+ auto calcArgs = [this, src, &graph, &visited_groups,
Review Comment:
code style
--
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] echuraev commented on pull request #15137: [Relay] Introduce arguments limit to FuseOps pass
Posted by "echuraev (via GitHub)" <gi...@apache.org>.
echuraev commented on PR #15137:
URL: https://github.com/apache/tvm/pull/15137#issuecomment-1645271618
@tvm-bot rerun
--
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] echuraev commented on pull request #15137: [Relay] Introduce arguments limit to FuseOps pass
Posted by "echuraev (via GitHub)" <gi...@apache.org>.
echuraev commented on PR #15137:
URL: https://github.com/apache/tvm/pull/15137#issuecomment-1643300020
@masahi Thank you for your review! I have corrected all mistakes that I have found. Please, let me know if I missed something.
--
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