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 2020/11/09 18:05:29 UTC
[GitHub] [incubator-tvm] ANSHUMAN87 opened a new pull request #6889: [TOPI] sparse_dense Op sparse_data input added
ANSHUMAN87 opened a new pull request #6889:
URL: https://github.com/apache/incubator-tvm/pull/6889
Change Summary:
Current sparse_dense Op in Topi support only weight as sparse input.
This PR add support for data also to be as sparse input.
So, in either case any one can be a sparse and other one will be dense.
This is a follow up PR from #6685
cc @tkonolige , @siju-samuel !
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] ANSHUMAN87 edited a comment on pull request #6889: [TOPI] sparse_dense Op sparse_data input added
Posted by GitBox <gi...@apache.org>.
ANSHUMAN87 edited a comment on pull request #6889:
URL: https://github.com/apache/incubator-tvm/pull/6889#issuecomment-724198770
> I'm not sure that the correct course of action is to add a flag to `sparse_dense` to support AB^T with B sparse. This makes all the implementations of `sparse_dense` confusing because they now have to check this flag and use a compute/schedule depending on if it is enabled or not. I'd instead prefer to make a new op called `dense_sparse` that does AB^T with B sparse.
>
> Alternatively, I don't really see a reason for supporting AB^T with B sparse directly. Instead, when we convert from a frontend to tvm, we can just insert the correct transposes. In a lot of ways this is a better choice because we do not need to reimplement the same operators and the operators prefer the data to be in this format. I think this is the best choice.
Thanks @tkonolige for response. I also had similar dilemma at the front. But later i resolved to re-utilize the existing Op, in order to enable reuse of small portion of code and it bind the concept of Op together too. But i am in favor of 'dense_sparse' Op as well(But that would be the name for existing Op i think :slightly_smiling_face: ). However i think we do not have much impact on the schedule side, in between these 2 flavor of Ops.
But the utilizing Op with transpose, i felt we are adding additional overhead every-time may not be much in smaller matrix.
Please let me know your thoughts on this. Thanks!
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] ANSHUMAN87 commented on pull request #6889: [TOPI] sparse_dense Op sparse_data input added
Posted by GitBox <gi...@apache.org>.
ANSHUMAN87 commented on pull request #6889:
URL: https://github.com/apache/incubator-tvm/pull/6889#issuecomment-724235751
Sure I will check on how much overhead added in case of transpose with existing Op case.
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] ANSHUMAN87 commented on pull request #6889: [TOPI] sparse_dense Op sparse_data input added
Posted by GitBox <gi...@apache.org>.
ANSHUMAN87 commented on pull request #6889:
URL: https://github.com/apache/incubator-tvm/pull/6889#issuecomment-728244556
> But maybe someone else can chime in.
Thanks @tkonolige for your feedback. I believe the performance stats are quite clear to opt for a new Op in the case.
However if you have any kind of concern, you can always let me know either here or offline too. We can discuss more in detail to reach to a more convincing point.
@tqchen : Tristan is helping here with his valuable review efforts. But i think we need a third opinion here (possibly an official reviewer or committer) to help proceed with the PR to next level. As i am not very sure who might be interested in Sparse changes, so if you can help tag few people here. TIA!
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] tkonolige commented on a change in pull request #6889: [TOPI] sparse_dense Op sparse_data input added
Posted by GitBox <gi...@apache.org>.
tkonolige commented on a change in pull request #6889:
URL: https://github.com/apache/incubator-tvm/pull/6889#discussion_r524495292
##########
File path: src/relay/op/nn/sparse.cc
##########
@@ -85,10 +117,11 @@ RELAY_REGISTER_OP("nn.sparse_dense")
)code" TVM_ADD_FILELINE)
.set_attrs_type<SparseDenseAttrs>()
.set_num_inputs(4)
- .add_argument("data", "nD Tensor", "Input data.")
- .add_argument("weight_data", "1D Tensor", "Weight data matrix.")
- .add_argument("weight_indices", "1D Tensor", "Weight indices matrix.")
- .add_argument("weight_indptr", "1D Tensor", "Weight indptr matrix.")
+ .add_argument("input_tensor1", "nD Tensor",
Review comment:
I'd prefer `dense_data`, `sparse_data`, `sparse_indices`, `sparse_indptr` as it is less verbose.
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] tkonolige commented on a change in pull request #6889: [TOPI] sparse_dense Op sparse_data input added
Posted by GitBox <gi...@apache.org>.
tkonolige commented on a change in pull request #6889:
URL: https://github.com/apache/incubator-tvm/pull/6889#discussion_r524456070
##########
File path: src/relay/op/nn/sparse.cc
##########
@@ -85,10 +117,11 @@ RELAY_REGISTER_OP("nn.sparse_dense")
)code" TVM_ADD_FILELINE)
.set_attrs_type<SparseDenseAttrs>()
.set_num_inputs(4)
- .add_argument("data", "nD Tensor", "Input data.")
- .add_argument("weight_data", "1D Tensor", "Weight data matrix.")
- .add_argument("weight_indices", "1D Tensor", "Weight indices matrix.")
- .add_argument("weight_indptr", "1D Tensor", "Weight indptr matrix.")
+ .add_argument("input_tensor1", "nD Tensor",
Review comment:
Please switch these back to their original names (data, weight_data, etc) or something descriptive.
##########
File path: python/tvm/topi/nn/sparse.py
##########
@@ -52,13 +52,120 @@ def sparse_dense(data, weight_data, weight_indices, weight_indptr):
"""
assert len(weight_data.shape) in (1, 3)
if len(weight_data.shape) == 1:
- func = _sparse_dense_csrmm
+ func = _sparse_dense_csrmm_v2
if len(weight_data.shape) == 3:
- func = _sparse_dense_bsrmm
+ func = _sparse_dense_bsrmm_v2
return func(data, weight_data, weight_indices, weight_indptr)
-def _sparse_dense_csrmm(data, weight_data, weight_indices, weight_indptr):
+def sparse_dense_v1(data_data, data_indices, data_indptr, weight):
+ """
+ Computes sparse-dense matrix multiplication of
+ `(data_data, data_indices, data_indptr)` and `weight.T`
+
+ Parameters
+ ----------
+ data_data:
+ 1-D with shape [nnz] (CSR) or
+ 3-D with shape [num_blocks, bs_r, bs_c] (BSR)
+
+ data_indices:
+ 1-D with shape [nnz] (CSR) or
+ 1-D with shape [num_blocks] (BSR)
+
+ data_indptr:
+ 1-D with shape [M + 1] (CSR) or
+ 1-D with shape [(M + 1) // bs_r] (BSR)
+
+ weight:
+ 2-D with shape [N, K], float32 when weight is dense
+
+ Returns
+ -------
+ output : tvm.te.Tensor
+ 2-D with shape [M, N]
+ """
+ assert len(data_data.shape) in (1, 3)
+ if len(data_data.shape) == 1:
+ func = _sparse_dense_csrmm_v1
+ if len(data_data.shape) == 3:
+ func = _sparse_dense_bsrmm_v1
+ return func(data_data, data_indices, data_indptr, weight)
+
+
+# pylint: disable=no-else-return,inconsistent-return-statements
+def sparse_dense(input_1, input_2, input_3, input_4, sparse_data=False):
+ """
+ Computes sparse-dense matrix multiplication of `data` and
+ `(weight_data, weight_indices, weight_indptr).T`
+ or
+ Computes sparse-dense matrix multiplication of
+ `(data_data, data_indices, data_indptr)` and `weight.T`
+
+ Parameters
+ ----------
+ input_1 : tvm.te.Tensor
Review comment:
Please name these something more descriptive. Maybe sparse_data, sparse_indices, sparse_indptr?
##########
File path: python/tvm/relay/op/nn/nn.py
##########
@@ -1993,7 +1993,8 @@ def batch_matmul(x, y):
return _make.batch_matmul(x, y)
-def sparse_dense(data, weight):
+# pylint: disable=no-else-return,inconsistent-return-statements
+def sparse_dense(data, weight, sparse_data=False):
Review comment:
Please document `sparse_data` here. Also update the doc comment with the new input types.
##########
File path: python/tvm/topi/nn/sparse.py
##########
@@ -52,13 +52,120 @@ def sparse_dense(data, weight_data, weight_indices, weight_indptr):
"""
assert len(weight_data.shape) in (1, 3)
if len(weight_data.shape) == 1:
- func = _sparse_dense_csrmm
+ func = _sparse_dense_csrmm_v2
if len(weight_data.shape) == 3:
- func = _sparse_dense_bsrmm
+ func = _sparse_dense_bsrmm_v2
return func(data, weight_data, weight_indices, weight_indptr)
-def _sparse_dense_csrmm(data, weight_data, weight_indices, weight_indptr):
+def sparse_dense_v1(data_data, data_indices, data_indptr, weight):
+ """
+ Computes sparse-dense matrix multiplication of
+ `(data_data, data_indices, data_indptr)` and `weight.T`
+
+ Parameters
+ ----------
+ data_data:
+ 1-D with shape [nnz] (CSR) or
+ 3-D with shape [num_blocks, bs_r, bs_c] (BSR)
+
+ data_indices:
+ 1-D with shape [nnz] (CSR) or
+ 1-D with shape [num_blocks] (BSR)
+
+ data_indptr:
+ 1-D with shape [M + 1] (CSR) or
+ 1-D with shape [(M + 1) // bs_r] (BSR)
+
+ weight:
+ 2-D with shape [N, K], float32 when weight is dense
+
+ Returns
+ -------
+ output : tvm.te.Tensor
+ 2-D with shape [M, N]
+ """
+ assert len(data_data.shape) in (1, 3)
+ if len(data_data.shape) == 1:
+ func = _sparse_dense_csrmm_v1
+ if len(data_data.shape) == 3:
+ func = _sparse_dense_bsrmm_v1
+ return func(data_data, data_indices, data_indptr, weight)
+
+
+# pylint: disable=no-else-return,inconsistent-return-statements
+def sparse_dense(input_1, input_2, input_3, input_4, sparse_data=False):
Review comment:
Please document `sparse_data`.
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] tkonolige commented on a change in pull request #6889: [TOPI] sparse_dense Op sparse_data input added
Posted by GitBox <gi...@apache.org>.
tkonolige commented on a change in pull request #6889:
URL: https://github.com/apache/incubator-tvm/pull/6889#discussion_r527281038
##########
File path: python/tvm/relay/op/nn/nn.py
##########
@@ -1994,7 +1994,7 @@ def batch_matmul(x, y):
# pylint: disable=no-else-return,inconsistent-return-statements
-def sparse_dense(data, weight, sparse_data=False):
+def sparse_dense(data, weight, sparse_lhs=False):
r"""
Computes the matrix multiplication of `data` and `weight`, where `data` is
Review comment:
Yeah, that might make more sense.
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] ANSHUMAN87 commented on a change in pull request #6889: [TOPI] sparse_dense Op sparse_data input added
Posted by GitBox <gi...@apache.org>.
ANSHUMAN87 commented on a change in pull request #6889:
URL: https://github.com/apache/incubator-tvm/pull/6889#discussion_r527208388
##########
File path: python/tvm/relay/op/nn/nn.py
##########
@@ -1994,7 +1994,7 @@ def batch_matmul(x, y):
# pylint: disable=no-else-return,inconsistent-return-statements
-def sparse_dense(data, weight, sparse_data=False):
+def sparse_dense(data, weight, sparse_lhs=False):
r"""
Computes the matrix multiplication of `data` and `weight`, where `data` is
Review comment:
Should we change the name data & weight to dense_mat & sparse_mat ?
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] ANSHUMAN87 edited a comment on pull request #6889: [TOPI] sparse_dense Op sparse_data input added
Posted by GitBox <gi...@apache.org>.
ANSHUMAN87 edited a comment on pull request #6889:
URL: https://github.com/apache/incubator-tvm/pull/6889#issuecomment-726801407
> Just to check, you're only transposing the dense matrix? Also, what is the density of the sparse matrix?
>
> I'm curious, could you do a benchmark with a smaller batch size (dense_input=[16, 1024])? And could you also test on cuda? And if you have time, could you give a breakdown of time spent in the matrix multiply vs time spent in the transpose?
>
> From these numbers, it does look like it is worthwhile to do this. It just means a lot more sparse kernels to write.
Yes , the additional transpose is done on final output from sparse_dense Op which is a dense tensor.
Cuda scheduling has issues currently which i am working on, maybe next PR i will handle it.
However i think cuda benchmark is not needed to conclude the benifit of new Op added in this PR.
A lower dimension output without transpose{Sparse_input = [4096, 1024], dense_input=[1024, 16]}::
With transpose:
With this i think we can start the code review.
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] ANSHUMAN87 commented on a change in pull request #6889: [TOPI] sparse_dense Op sparse_data input added
Posted by GitBox <gi...@apache.org>.
ANSHUMAN87 commented on a change in pull request #6889:
URL: https://github.com/apache/incubator-tvm/pull/6889#discussion_r526368800
##########
File path: src/relay/op/nn/sparse.cc
##########
@@ -85,10 +117,11 @@ RELAY_REGISTER_OP("nn.sparse_dense")
)code" TVM_ADD_FILELINE)
.set_attrs_type<SparseDenseAttrs>()
.set_num_inputs(4)
- .add_argument("data", "nD Tensor", "Input data.")
- .add_argument("weight_data", "1D Tensor", "Weight data matrix.")
- .add_argument("weight_indices", "1D Tensor", "Weight indices matrix.")
- .add_argument("weight_indptr", "1D Tensor", "Weight indptr matrix.")
+ .add_argument("input_tensor1", "nD Tensor",
Review comment:
Okay thanks for details, I got your point clearly now!
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] ANSHUMAN87 commented on pull request #6889: [TOPI] sparse_dense Op sparse_data input added
Posted by GitBox <gi...@apache.org>.
ANSHUMAN87 commented on pull request #6889:
URL: https://github.com/apache/incubator-tvm/pull/6889#issuecomment-724198770
> I'm not sure that the correct course of action is to add a flag to `sparse_dense` to support AB^T with B sparse. This makes all the implementations of `sparse_dense` confusing because they now have to check this flag and use a compute/schedule depending on if it is enabled or not. I'd instead prefer to make a new op called `dense_sparse` that does AB^T with B sparse.
>
> Alternatively, I don't really see a reason for supporting AB^T with B sparse directly. Instead, when we convert from a frontend to tvm, we can just insert the correct transposes. In a lot of ways this is a better choice because we do not need to reimplement the same operators and the operators prefer the data to be in this format. I think this is the best choice.
Thanks @tkonolige for response. I also had similar dilemma at the front. But later i resolved to re-utilize the existing Op, in order to enable reuse of small portion of code and it bind the concept of Op together too. But i am in favor of 'dense_sparse' Op as well(But that would be the name for existing Op i think :slightly_smiling_face: ).
But the utilizing Op with transpose, i felt we are adding additional overhead every-time may not be much in smaller matrix.
Please let me know your thoughts on this. Thanks!
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] ANSHUMAN87 commented on pull request #6889: [TOPI] sparse_dense Op sparse_data input added
Posted by GitBox <gi...@apache.org>.
ANSHUMAN87 commented on pull request #6889:
URL: https://github.com/apache/incubator-tvm/pull/6889#issuecomment-726801407
> Just to check, you're only transposing the dense matrix? Also, what is the density of the sparse matrix?
>
> I'm curious, could you do a benchmark with a smaller batch size (dense_input=[16, 1024])? And could you also test on cuda? And if you have time, could you give a breakdown of time spent in the matrix multiply vs time spent in the transpose?
>
> From these numbers, it does look like it is worthwhile to do this. It just means a lot more sparse kernels to write.
Yes , the additional transpose is done on final output from sparse_dense Op which is a dense tensor.
Cuda scheduling has issues currently which i am working on, maybe next PR i will handle it.
However i think cuda benchmark is not needed to conclude the benifit of new Op added in this PR.
A lower dimension output without transpose:
With transpose:
With this i think we can start the code review.
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] tkonolige commented on pull request #6889: [TOPI] sparse_dense Op sparse_data input added
Posted by GitBox <gi...@apache.org>.
tkonolige commented on pull request #6889:
URL: https://github.com/apache/incubator-tvm/pull/6889#issuecomment-725585407
Just to check, you're only transposing the dense matrix? Also, what is the density of the sparse matrix?
I'm curious, could you do a benchmark with a smaller batch size (dense_input=[16, 1024])? And could you also test on cuda? And if you have time, could you give a breakdown of time spent in the matrix multiply vs time spent in the transpose?
From these numbers, it does look like it is worthwhile to do this. It just means a lot more sparse kernels to write.
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] ANSHUMAN87 commented on a change in pull request #6889: [TOPI] sparse_dense Op sparse_data input added
Posted by GitBox <gi...@apache.org>.
ANSHUMAN87 commented on a change in pull request #6889:
URL: https://github.com/apache/incubator-tvm/pull/6889#discussion_r524490253
##########
File path: src/relay/op/nn/sparse.cc
##########
@@ -85,10 +117,11 @@ RELAY_REGISTER_OP("nn.sparse_dense")
)code" TVM_ADD_FILELINE)
.set_attrs_type<SparseDenseAttrs>()
.set_num_inputs(4)
- .add_argument("data", "nD Tensor", "Input data.")
- .add_argument("weight_data", "1D Tensor", "Weight data matrix.")
- .add_argument("weight_indices", "1D Tensor", "Weight indices matrix.")
- .add_argument("weight_indptr", "1D Tensor", "Weight indptr matrix.")
+ .add_argument("input_tensor1", "nD Tensor",
Review comment:
I think we can not go back to their original names, as the tensor will be used for multiple purposes as given in the description string. I too am not impressed by the name "input_tensor1" :slightly_smiling_face:
Would appreciate if you can suggest some names.
How about merger of both use case as below?
"input_tensor1" => "data_dense_or_sparse"
"input_tensor2" => "weight_data_or_data_indices"
"input_tensor3" => "weight_indices_or_data_indptr"
"input_tensor3" => "weight_indptr_or_dense_weight"
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] tkonolige commented on a change in pull request #6889: [TOPI] sparse_dense Op sparse_data input added
Posted by GitBox <gi...@apache.org>.
tkonolige commented on a change in pull request #6889:
URL: https://github.com/apache/incubator-tvm/pull/6889#discussion_r526311567
##########
File path: src/relay/op/nn/sparse.cc
##########
@@ -85,10 +117,11 @@ RELAY_REGISTER_OP("nn.sparse_dense")
)code" TVM_ADD_FILELINE)
.set_attrs_type<SparseDenseAttrs>()
.set_num_inputs(4)
- .add_argument("data", "nD Tensor", "Input data.")
- .add_argument("weight_data", "1D Tensor", "Weight data matrix.")
- .add_argument("weight_indices", "1D Tensor", "Weight indices matrix.")
- .add_argument("weight_indptr", "1D Tensor", "Weight indptr matrix.")
+ .add_argument("input_tensor1", "nD Tensor",
Review comment:
I would keep the inputs in the original order (no matter what the value of the sparse_data flag is) . That is data, then sparse_data, sparse_indices, sparse_indptr. The order always remains the same, and the flag just changes which computation happens.
Also please change the name of the sparse_data flag. It will conflict with the new names of the parameters.
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] tkonolige commented on a change in pull request #6889: [TOPI] sparse_dense Op sparse_data input added
Posted by GitBox <gi...@apache.org>.
tkonolige commented on a change in pull request #6889:
URL: https://github.com/apache/incubator-tvm/pull/6889#discussion_r527200792
##########
File path: python/tvm/topi/nn/sparse.py
##########
@@ -52,13 +52,105 @@ def sparse_dense(data, weight_data, weight_indices, weight_indptr):
"""
assert len(weight_data.shape) in (1, 3)
if len(weight_data.shape) == 1:
- func = _sparse_dense_csrmm
+ func = _sparse_dense_csrmm_v2
if len(weight_data.shape) == 3:
- func = _sparse_dense_bsrmm
+ func = _sparse_dense_bsrmm_v2
return func(data, weight_data, weight_indices, weight_indptr)
-def _sparse_dense_csrmm(data, weight_data, weight_indices, weight_indptr):
+def sparse_dense_v1(data_data, data_indices, data_indptr, weight):
+ """
+ Computes sparse-dense matrix multiplication of
+ `(data_data, data_indices, data_indptr)` and `weight.T`
+
+ Parameters
+ ----------
+ data_data:
+ 1-D with shape [nnz] (CSR) or
+ 3-D with shape [num_blocks, bs_r, bs_c] (BSR)
+
+ data_indices:
+ 1-D with shape [nnz] (CSR) or
+ 1-D with shape [num_blocks] (BSR)
+
+ data_indptr:
+ 1-D with shape [M + 1] (CSR) or
+ 1-D with shape [(M + 1) // bs_r] (BSR)
+
+ weight:
+ 2-D with shape [N, K], float32 when weight is dense
+
+ Returns
+ -------
+ output : tvm.te.Tensor
+ 2-D with shape [M, N]
+ """
+ assert len(data_data.shape) in (1, 3)
+ if len(data_data.shape) == 1:
+ func = _sparse_dense_csrmm_v1
+ if len(data_data.shape) == 3:
+ func = _sparse_dense_bsrmm_v1
+ return func(data_data, data_indices, data_indptr, weight)
+
+
+# pylint: disable=no-else-return,inconsistent-return-statements
+def sparse_dense(dense_data, sparse_data, sparse_indices, sparse_indptr, sparse_lhs=False):
+ """
+ Computes sparse-dense matrix multiplication of `data` and
+ `(weight_data, weight_indices, weight_indptr).T`
+ or
+ Computes sparse-dense matrix multiplication of
+ `(data_data, data_indices, data_indptr)` and `weight.T`
+
+ Parameters
+ ----------
+ dense_data : tvm.te.Tensor
+ data:
+ 2-D with shape [M, K], float32 when input data is dense or
+
+ sparse_data : tvm.te.Tensor
+ 1-D with shape [nnz] (CSR) or
+ 3-D with shape [num_blocks, bs_r, bs_c] (BSR) or
Review comment:
remove the trailing or.
##########
File path: python/tvm/topi/nn/sparse.py
##########
@@ -52,13 +52,105 @@ def sparse_dense(data, weight_data, weight_indices, weight_indptr):
"""
assert len(weight_data.shape) in (1, 3)
if len(weight_data.shape) == 1:
- func = _sparse_dense_csrmm
+ func = _sparse_dense_csrmm_v2
if len(weight_data.shape) == 3:
- func = _sparse_dense_bsrmm
+ func = _sparse_dense_bsrmm_v2
return func(data, weight_data, weight_indices, weight_indptr)
-def _sparse_dense_csrmm(data, weight_data, weight_indices, weight_indptr):
+def sparse_dense_v1(data_data, data_indices, data_indptr, weight):
+ """
+ Computes sparse-dense matrix multiplication of
+ `(data_data, data_indices, data_indptr)` and `weight.T`
+
+ Parameters
+ ----------
+ data_data:
+ 1-D with shape [nnz] (CSR) or
+ 3-D with shape [num_blocks, bs_r, bs_c] (BSR)
+
+ data_indices:
+ 1-D with shape [nnz] (CSR) or
+ 1-D with shape [num_blocks] (BSR)
+
+ data_indptr:
+ 1-D with shape [M + 1] (CSR) or
+ 1-D with shape [(M + 1) // bs_r] (BSR)
+
+ weight:
+ 2-D with shape [N, K], float32 when weight is dense
+
+ Returns
+ -------
+ output : tvm.te.Tensor
+ 2-D with shape [M, N]
+ """
+ assert len(data_data.shape) in (1, 3)
+ if len(data_data.shape) == 1:
+ func = _sparse_dense_csrmm_v1
+ if len(data_data.shape) == 3:
+ func = _sparse_dense_bsrmm_v1
+ return func(data_data, data_indices, data_indptr, weight)
+
+
+# pylint: disable=no-else-return,inconsistent-return-statements
+def sparse_dense(dense_data, sparse_data, sparse_indices, sparse_indptr, sparse_lhs=False):
+ """
+ Computes sparse-dense matrix multiplication of `data` and
+ `(weight_data, weight_indices, weight_indptr).T`
+ or
+ Computes sparse-dense matrix multiplication of
+ `(data_data, data_indices, data_indptr)` and `weight.T`
Review comment:
Indicate what multiplication happens with the `sparse_lhs` flag.
##########
File path: python/tvm/relay/op/nn/nn.py
##########
@@ -1994,7 +1994,7 @@ def batch_matmul(x, y):
# pylint: disable=no-else-return,inconsistent-return-statements
-def sparse_dense(data, weight, sparse_data=False):
+def sparse_dense(data, weight, sparse_lhs=False):
r"""
Computes the matrix multiplication of `data` and `weight`, where `data` is
Review comment:
Please clarify what operation is occurring here. S * D^T if sparse_lhs, D * S^T otherwise.
##########
File path: python/tvm/relay/op/nn/nn.py
##########
@@ -2020,19 +2020,18 @@ def sparse_dense(data, weight, sparse_data=False):
weight : Union[namedtuple, Tuple[ndarray, ndarray, ndarray]].
The sparse weight matrix for the matrix multiplication.
+ sparse_lhs : bool, optional
+ Indicates whether lhs or rhs matrix is sparse.
Review comment:
Indicate what the default value is.
##########
File path: python/tvm/topi/nn/sparse.py
##########
@@ -52,13 +52,105 @@ def sparse_dense(data, weight_data, weight_indices, weight_indptr):
"""
assert len(weight_data.shape) in (1, 3)
if len(weight_data.shape) == 1:
- func = _sparse_dense_csrmm
+ func = _sparse_dense_csrmm_v2
if len(weight_data.shape) == 3:
- func = _sparse_dense_bsrmm
+ func = _sparse_dense_bsrmm_v2
return func(data, weight_data, weight_indices, weight_indptr)
-def _sparse_dense_csrmm(data, weight_data, weight_indices, weight_indptr):
+def sparse_dense_v1(data_data, data_indices, data_indptr, weight):
+ """
+ Computes sparse-dense matrix multiplication of
+ `(data_data, data_indices, data_indptr)` and `weight.T`
+
+ Parameters
+ ----------
+ data_data:
+ 1-D with shape [nnz] (CSR) or
+ 3-D with shape [num_blocks, bs_r, bs_c] (BSR)
+
+ data_indices:
+ 1-D with shape [nnz] (CSR) or
+ 1-D with shape [num_blocks] (BSR)
+
+ data_indptr:
+ 1-D with shape [M + 1] (CSR) or
+ 1-D with shape [(M + 1) // bs_r] (BSR)
+
+ weight:
+ 2-D with shape [N, K], float32 when weight is dense
+
+ Returns
+ -------
+ output : tvm.te.Tensor
+ 2-D with shape [M, N]
+ """
+ assert len(data_data.shape) in (1, 3)
+ if len(data_data.shape) == 1:
+ func = _sparse_dense_csrmm_v1
+ if len(data_data.shape) == 3:
+ func = _sparse_dense_bsrmm_v1
+ return func(data_data, data_indices, data_indptr, weight)
+
+
+# pylint: disable=no-else-return,inconsistent-return-statements
+def sparse_dense(dense_data, sparse_data, sparse_indices, sparse_indptr, sparse_lhs=False):
+ """
+ Computes sparse-dense matrix multiplication of `data` and
+ `(weight_data, weight_indices, weight_indptr).T`
+ or
+ Computes sparse-dense matrix multiplication of
+ `(data_data, data_indices, data_indptr)` and `weight.T`
+
+ Parameters
+ ----------
+ dense_data : tvm.te.Tensor
+ data:
+ 2-D with shape [M, K], float32 when input data is dense or
Review comment:
Update this description.
##########
File path: python/tvm/topi/nn/sparse.py
##########
@@ -52,13 +52,105 @@ def sparse_dense(data, weight_data, weight_indices, weight_indptr):
"""
assert len(weight_data.shape) in (1, 3)
if len(weight_data.shape) == 1:
- func = _sparse_dense_csrmm
+ func = _sparse_dense_csrmm_v2
if len(weight_data.shape) == 3:
- func = _sparse_dense_bsrmm
+ func = _sparse_dense_bsrmm_v2
return func(data, weight_data, weight_indices, weight_indptr)
-def _sparse_dense_csrmm(data, weight_data, weight_indices, weight_indptr):
+def sparse_dense_v1(data_data, data_indices, data_indptr, weight):
+ """
+ Computes sparse-dense matrix multiplication of
+ `(data_data, data_indices, data_indptr)` and `weight.T`
+
+ Parameters
+ ----------
+ data_data:
+ 1-D with shape [nnz] (CSR) or
+ 3-D with shape [num_blocks, bs_r, bs_c] (BSR)
+
+ data_indices:
+ 1-D with shape [nnz] (CSR) or
+ 1-D with shape [num_blocks] (BSR)
+
+ data_indptr:
+ 1-D with shape [M + 1] (CSR) or
+ 1-D with shape [(M + 1) // bs_r] (BSR)
+
+ weight:
+ 2-D with shape [N, K], float32 when weight is dense
+
+ Returns
+ -------
+ output : tvm.te.Tensor
+ 2-D with shape [M, N]
+ """
+ assert len(data_data.shape) in (1, 3)
+ if len(data_data.shape) == 1:
+ func = _sparse_dense_csrmm_v1
+ if len(data_data.shape) == 3:
+ func = _sparse_dense_bsrmm_v1
+ return func(data_data, data_indices, data_indptr, weight)
+
+
+# pylint: disable=no-else-return,inconsistent-return-statements
+def sparse_dense(dense_data, sparse_data, sparse_indices, sparse_indptr, sparse_lhs=False):
+ """
+ Computes sparse-dense matrix multiplication of `data` and
+ `(weight_data, weight_indices, weight_indptr).T`
+ or
+ Computes sparse-dense matrix multiplication of
+ `(data_data, data_indices, data_indptr)` and `weight.T`
+
+ Parameters
+ ----------
+ dense_data : tvm.te.Tensor
+ data:
+ 2-D with shape [M, K], float32 when input data is dense or
+
+ sparse_data : tvm.te.Tensor
+ 1-D with shape [nnz] (CSR) or
+ 3-D with shape [num_blocks, bs_r, bs_c] (BSR) or
+
+ sparse_indices : tvm.te.Tensor
+ 1-D with shape [nnz] (CSR) or
+ 1-D with shape [num_blocks] (BSR) or
Review comment:
remove trailing or
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] tkonolige commented on pull request #6889: [TOPI] sparse_dense Op sparse_data input added
Posted by GitBox <gi...@apache.org>.
tkonolige commented on pull request #6889:
URL: https://github.com/apache/incubator-tvm/pull/6889#issuecomment-724214412
I think it is still possible to reuse code with a separate op, but maybe its to a lesser extent.
I'm not sure how much overhead the transposes will add to the code. We don't have to pay the cost of transposing the sparse matrix because that can be done at compile time (the sparse matrix is usually a weight). Could you maybe do some benchmarking and see if the duplication of code is worth it from a performance standpoint?
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] ANSHUMAN87 commented on pull request #6889: [TOPI] sparse_dense Op sparse_data input added
Posted by GitBox <gi...@apache.org>.
ANSHUMAN87 commented on pull request #6889:
URL: https://github.com/apache/incubator-tvm/pull/6889#issuecomment-727237614
Gentle ping @tkonolige !!!
I am not too sure who else from TVM official reviewer or committer interested in sparse. If you are aware of anyone please feel free to invite. TIA!
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] tkonolige commented on pull request #6889: [TOPI] sparse_dense Op sparse_data input added
Posted by GitBox <gi...@apache.org>.
tkonolige commented on pull request #6889:
URL: https://github.com/apache/incubator-tvm/pull/6889#issuecomment-724191460
I'm not sure that the correct course of action is to add a flag to `sparse_dense` to support AB^T with B sparse. This makes all the implementations of `sparse_dense` confusing because they now have to check this flag and use a compute/schedule depending on if it is enabled or not. I'd instead prefer to make a new op called `dense_sparse` that does AB^T with B sparse.
Alternatively, I don't really see a reason for supporting AB^T with B sparse directly. Instead, when we convert from a frontend to tvm, we can just insert the correct transposes. In a lot of ways this is a better choice because we do not need to reimplement the same operators and the operators prefer the data to be in this format. I think this is the best choice.
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] ANSHUMAN87 commented on a change in pull request #6889: [TOPI] sparse_dense Op sparse_data input added
Posted by GitBox <gi...@apache.org>.
ANSHUMAN87 commented on a change in pull request #6889:
URL: https://github.com/apache/incubator-tvm/pull/6889#discussion_r526079596
##########
File path: src/relay/op/nn/sparse.cc
##########
@@ -85,10 +117,11 @@ RELAY_REGISTER_OP("nn.sparse_dense")
)code" TVM_ADD_FILELINE)
.set_attrs_type<SparseDenseAttrs>()
.set_num_inputs(4)
- .add_argument("data", "nD Tensor", "Input data.")
- .add_argument("weight_data", "1D Tensor", "Weight data matrix.")
- .add_argument("weight_indices", "1D Tensor", "Weight indices matrix.")
- .add_argument("weight_indptr", "1D Tensor", "Weight indptr matrix.")
+ .add_argument("input_tensor1", "nD Tensor",
Review comment:
@tkonolige : Sorry for late response. I think somehow i missed your last reply.
I agree with the naming you suggested, but the issue here is one tensor represents 2 different use case altogether.
For example:
"input_tensor2" => "weight_data_or_data_indices"
So it can be sparse_data in one case and sparse_indices in another case.
So i am wondering what is the best way to merge both representation ?
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] ANSHUMAN87 commented on pull request #6889: [TOPI] sparse_dense Op sparse_data input added
Posted by GitBox <gi...@apache.org>.
ANSHUMAN87 commented on pull request #6889:
URL: https://github.com/apache/incubator-tvm/pull/6889#issuecomment-729671304
> @tqchen : Tristan is helping here with his valuable review efforts. But i think we need a third opinion here (possibly an official reviewer or committer) to help proceed with the PR to next level. As i am not very sure who might be interested in Sparse changes, so if you can help tag few people here. TIA!
Gentle ping @tqchen !
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] tkonolige commented on a change in pull request #6889: [TOPI] sparse_dense Op sparse_data input added
Posted by GitBox <gi...@apache.org>.
tkonolige commented on a change in pull request #6889:
URL: https://github.com/apache/incubator-tvm/pull/6889#discussion_r527198571
##########
File path: include/tvm/relay/attrs/nn.h
##########
@@ -937,12 +937,12 @@ struct DenseAttrs : public tvm::AttrsNode<DenseAttrs> {
/*! \brief Attributes for sparse_dense operator */
struct SparseDenseAttrs : public tvm::AttrsNode<SparseDenseAttrs> {
- bool sparse_data;
+ bool sparse_lhs;
TVM_DECLARE_ATTRS(SparseDenseAttrs, "relay.attrs.SparseDenseAttrs") {
- TVM_ATTR_FIELD(sparse_data)
+ TVM_ATTR_FIELD(sparse_lhs)
.set_default(false)
- .describe("Indicate whether data or weight is sparse. True if data is sparse");
+ .describe("Indicate whether lhs or rhs matrix is sparse. True if lhs matrix is sparse");
Review comment:
I think this could use a little more clarification. Maybe "Indicate whether sparse matrix is multiplied on the right or the left. If true, then the operation is S * D^T (D dense, S sparse). If false, the operation is D * S^T"
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] ANSHUMAN87 commented on pull request #6889: [TOPI] sparse_dense Op sparse_data input added
Posted by GitBox <gi...@apache.org>.
ANSHUMAN87 commented on pull request #6889:
URL: https://github.com/apache/incubator-tvm/pull/6889#issuecomment-725492811
Hi @tkonolige , below is the benchmark data i have obtained for 4 different input dimensions.
NOTE: Here with Transpose means using existing sparse_dense Op with additional Transpose layer.
Without Transpose means using the new flavor of the Op implemented in the current PR.
And all the data is collected with iteration=3000 and repeat=3 .
**Case 1{Sparse_input = [1024, 512], dense_input=[512, 4096]}:**
**Case 2{Sparse_input = [1024, 512], dense_input=[512, 1024]}:**
**Case 3{Sparse_input = [2048, 512], dense_input=[512, 2048]}:**
**Case 4{Sparse_input = [4096, 512], dense_input=[512, 4096]}:**
Clearly as the dimension increases the amount of improvement is significant. So i think we should keep the new flavor of the Op implemented in current 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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] tkonolige commented on pull request #6889: [TOPI] sparse_dense Op sparse_data input added
Posted by GitBox <gi...@apache.org>.
tkonolige commented on pull request #6889:
URL: https://github.com/apache/incubator-tvm/pull/6889#issuecomment-728249449
Sorry, I meant that I'm not sure about the using flags approach vs having a separate operator for sparse x dense vs dense x sparse. From your benchmarks, it does look like we need some way of doing dense x sparse directly.
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org