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/12/22 06:17:08 UTC
[GitHub] [tvm] ANSHUMAN87 opened a new pull request #7148: [Frontend][Tensorflow] Sparse_Dense Op CSR scheduling issue resolved for Cuda & X86
ANSHUMAN87 opened a new pull request #7148:
URL: https://github.com/apache/tvm/pull/7148
This is a follow up PR.
1. It has resolved the issue in CSR scheduling for both Cuda & X86.
2. Also the test cases in Tensorflow frontends are enabled for the same.
cc @tkonolige , @FrozenGene !
----------------------------------------------------------------
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] [tvm] ANSHUMAN87 commented on pull request #7148: [Frontend][Tensorflow] Sparse_Dense Op CSR scheduling issue resolved for Cuda & X86
Posted by GitBox <gi...@apache.org>.
ANSHUMAN87 commented on pull request #7148:
URL: https://github.com/apache/tvm/pull/7148#issuecomment-750325140
> What was the issue holding up CSR scheduling?
>
> We already talked about performance for this right?
Cuda scheduling for Sparse_dense Op is internally changed to Sparse_dense_padded. But it works only when multiple of warp_size, but if it is lower than that, there is no fallback scheduling for CSR, so i have resolved that part here. Please let me know in case i am not clear. 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] [tvm] tkonolige commented on a change in pull request #7148: [Frontend][Tensorflow] Sparse_Dense Op CSR scheduling issue resolved for Cuda & X86
Posted by GitBox <gi...@apache.org>.
tkonolige commented on a change in pull request #7148:
URL: https://github.com/apache/tvm/pull/7148#discussion_r557524478
##########
File path: tests/python/topi/python/test_topi_sparse.py
##########
@@ -530,5 +557,6 @@ def test_sparse_dense_padded_alter_op():
test_sparse_transpose_csr()
test_sparse_dense_padded_cuda()
test_sparse_dense_padded_alter_op()
Review comment:
the change from constant to var is so that the input variable does not get inlined in build.
----------------------------------------------------------------
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] [tvm] tkonolige commented on a change in pull request #7148: [Frontend][Tensorflow] Sparse_Dense Op CSR scheduling issue resolved for Cuda & X86
Posted by GitBox <gi...@apache.org>.
tkonolige commented on a change in pull request #7148:
URL: https://github.com/apache/tvm/pull/7148#discussion_r547395050
##########
File path: python/tvm/topi/cuda/sparse.py
##########
@@ -311,6 +339,8 @@ def sparse_dense_padded(data, weight_data, weight_indices, weight_indptr):
output : tvm.te.Tensor
2-D with shape [M, N]
"""
+ # TODO(ANSHUMAN87): Handle for sparse_lhs case too
+ assert not sparse_lhs
Review comment:
Can you add a message to this one.
##########
File path: python/tvm/topi/x86/sparse.py
##########
@@ -28,15 +28,17 @@ def schedule_sparse_dense(outs):
def _callback(op):
simd_width = get_fp32_len()
- if op.tag == "sparse_dense_csrmm" and op != outs[0].op:
- (_, v_i) = s[op].op.axis
- s[op].vectorize(v_i)
- (y_o, y_i) = s[outs[0].op].split(s[outs[0].op].op.axis[1], 2 * simd_width)
- s[op].compute_at(s[outs[0]], y_o)
- s[outs[0].op].vectorize(y_i)
- if op.tag == "sparse_dense_bsrmm":
+ if op.tag == "sparse_dense_csrmm_v2" or op.tag == "sparse_dense_csrmm_v1":
+ (y_o, y_i) = s[op].split(s[op].op.axis[1], 2)
+ fused = s[op].fuse(s[op].op.axis[0], y_o)
+ s[op].parallel(fused)
+ s[op].vectorize(y_i)
+ elif op.tag == "sparse_dense_bsrmm_v2" or op.tag == "sparse_dense_bsrmm_v1":
Review comment:
What is the difference between v1 and v2? Can we consolidate tags? If not, maybe we should rename them to be more clear about what they are doing.
----------------------------------------------------------------
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] [tvm] tkonolige commented on a change in pull request #7148: [Frontend][Tensorflow] Sparse_Dense Op CSR scheduling issue resolved for Cuda & X86
Posted by GitBox <gi...@apache.org>.
tkonolige commented on a change in pull request #7148:
URL: https://github.com/apache/tvm/pull/7148#discussion_r557512782
##########
File path: tests/python/topi/python/test_topi_sparse.py
##########
@@ -530,5 +557,6 @@ def test_sparse_dense_padded_alter_op():
test_sparse_transpose_csr()
test_sparse_dense_padded_cuda()
test_sparse_dense_padded_alter_op()
Review comment:
You can just remove the old test. I wrote the new one as a replacement for 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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] ANSHUMAN87 commented on a change in pull request #7148: [Frontend][Tensorflow] Sparse_Dense Op CSR scheduling issue resolved for Cuda & X86
Posted by GitBox <gi...@apache.org>.
ANSHUMAN87 commented on a change in pull request #7148:
URL: https://github.com/apache/tvm/pull/7148#discussion_r557532462
##########
File path: tests/python/topi/python/test_topi_sparse.py
##########
@@ -530,5 +557,6 @@ def test_sparse_dense_padded_alter_op():
test_sparse_transpose_csr()
test_sparse_dense_padded_cuda()
test_sparse_dense_padded_alter_op()
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] ANSHUMAN87 commented on a change in pull request #7148: [Frontend][Tensorflow] Sparse_Dense Op CSR scheduling issue resolved for Cuda & X86
Posted by GitBox <gi...@apache.org>.
ANSHUMAN87 commented on a change in pull request #7148:
URL: https://github.com/apache/tvm/pull/7148#discussion_r557529148
##########
File path: tests/python/topi/python/test_topi_sparse.py
##########
@@ -530,5 +557,6 @@ def test_sparse_dense_padded_alter_op():
test_sparse_transpose_csr()
test_sparse_dense_padded_cuda()
test_sparse_dense_padded_alter_op()
Review comment:
Right. Anyway does not give much benefit also. Lets remove it then!
----------------------------------------------------------------
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] [tvm] ANSHUMAN87 commented on pull request #7148: [Frontend][Tensorflow] Sparse_Dense Op CSR scheduling issue resolved for Cuda & X86
Posted by GitBox <gi...@apache.org>.
ANSHUMAN87 commented on pull request #7148:
URL: https://github.com/apache/tvm/pull/7148#issuecomment-759136347
Gentle ping @tkonolige !
cc @junrushao1994 , @comaniac too :)
----------------------------------------------------------------
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] [tvm] comaniac commented on pull request #7148: [Frontend][Tensorflow] Sparse_Dense Op CSR scheduling issue resolved for Cuda & X86
Posted by GitBox <gi...@apache.org>.
comaniac commented on pull request #7148:
URL: https://github.com/apache/tvm/pull/7148#issuecomment-760603648
Thanks @ANSHUMAN87 @tkonolige
----------------------------------------------------------------
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] [tvm] ANSHUMAN87 commented on a change in pull request #7148: [Frontend][Tensorflow] Sparse_Dense Op CSR scheduling issue resolved for Cuda & X86
Posted by GitBox <gi...@apache.org>.
ANSHUMAN87 commented on a change in pull request #7148:
URL: https://github.com/apache/tvm/pull/7148#discussion_r557523604
##########
File path: tests/python/topi/python/test_topi_sparse.py
##########
@@ -530,5 +557,6 @@ def test_sparse_dense_padded_alter_op():
test_sparse_transpose_csr()
test_sparse_dense_padded_cuda()
test_sparse_dense_padded_alter_op()
Review comment:
Yes i agree on that. I was referring the var input & constant input diff :)
----------------------------------------------------------------
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] [tvm] tkonolige commented on a change in pull request #7148: [Frontend][Tensorflow] Sparse_Dense Op CSR scheduling issue resolved for Cuda & X86
Posted by GitBox <gi...@apache.org>.
tkonolige commented on a change in pull request #7148:
URL: https://github.com/apache/tvm/pull/7148#discussion_r548054955
##########
File path: python/tvm/topi/x86/sparse.py
##########
@@ -28,15 +28,17 @@ def schedule_sparse_dense(outs):
def _callback(op):
simd_width = get_fp32_len()
- if op.tag == "sparse_dense_csrmm" and op != outs[0].op:
- (_, v_i) = s[op].op.axis
- s[op].vectorize(v_i)
- (y_o, y_i) = s[outs[0].op].split(s[outs[0].op].op.axis[1], 2 * simd_width)
- s[op].compute_at(s[outs[0]], y_o)
- s[outs[0].op].vectorize(y_i)
- if op.tag == "sparse_dense_bsrmm":
+ if op.tag == "sparse_dense_csrmm_v2" or op.tag == "sparse_dense_csrmm_v1":
+ (y_o, y_i) = s[op].split(s[op].op.axis[1], 2)
+ fused = s[op].fuse(s[op].op.axis[0], y_o)
+ s[op].parallel(fused)
+ s[op].vectorize(y_i)
+ elif op.tag == "sparse_dense_bsrmm_v2" or op.tag == "sparse_dense_bsrmm_v1":
Review comment:
Yeah, lets do "sparse_dense_bsrmm_v2" --> "sparse_dense_bsrmm_sparse_rhs".
----------------------------------------------------------------
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] [tvm] ANSHUMAN87 commented on pull request #7148: [Frontend][Tensorflow] Sparse_Dense Op CSR scheduling issue resolved for Cuda & X86
Posted by GitBox <gi...@apache.org>.
ANSHUMAN87 commented on pull request #7148:
URL: https://github.com/apache/tvm/pull/7148#issuecomment-760296703
> So I just hit the bug that this fixes. Can we add a test to make sure we don't hit it again in the future. Here is the test I wrote:
>
> ```python
> @tvm.testing.requires_cuda
> def test_sparse_dense_padded_alter_op():
> with tvm.target.Target("cuda"):
> M = 128
> N = 16
> K = 128
> X_np = np.random.randn(M, K).astype("float32")
> W_sp_np = random_bsr_matrix(N, K, 2, 2, density=0.01, dtype="float32")
> x = relay.var("x", relay.TensorType(X_np.shape,"float32"))
> mult = relay.op.nn.sparse_dense(
> x,
> (
> relay.Constant(tvm.nd.array(W_sp_np.data)),
> relay.Constant(tvm.nd.array(W_sp_np.indices)),
> relay.Constant(tvm.nd.array(W_sp_np.indptr)),
> ),
> )
> f = relay.Function([x], mult)
> f_ = relay.transform.InferType()(tvm.IRModule.from_expr(f))
> f_ = relay.transform.AlterOpLayout()(f_)
> assert f_["main"].body.op.name == "nn.internal.sparse_dense_padded"
>
> # build with cuda and AlterOpLayout to ensure that sparse_dense_padded is has an implementation
> with tvm.transform.PassContext(opt_level=3, required_pass="AlterOpLayout"):
> x = relay.build(tvm.IRModule.from_expr(f), target=tvm.target.Target("cuda"))
> ```
>
> in `tests/python/topi/python/test_topi_sparse.py`
Thanks @tkonolige ! The test case is added 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] [tvm] tkonolige commented on a change in pull request #7148: [Frontend][Tensorflow] Sparse_Dense Op CSR scheduling issue resolved for Cuda & X86
Posted by GitBox <gi...@apache.org>.
tkonolige commented on a change in pull request #7148:
URL: https://github.com/apache/tvm/pull/7148#discussion_r557521867
##########
File path: tests/python/topi/python/test_topi_sparse.py
##########
@@ -530,5 +557,6 @@ def test_sparse_dense_padded_alter_op():
test_sparse_transpose_csr()
test_sparse_dense_padded_cuda()
test_sparse_dense_padded_alter_op()
Review comment:
The new tests is a superset of the old one. They both test the same things, this just tests a little more.
----------------------------------------------------------------
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] [tvm] comaniac merged pull request #7148: [Frontend][Tensorflow] Sparse_Dense Op CSR scheduling issue resolved for Cuda & X86
Posted by GitBox <gi...@apache.org>.
comaniac merged pull request #7148:
URL: https://github.com/apache/tvm/pull/7148
----------------------------------------------------------------
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] [tvm] ANSHUMAN87 commented on a change in pull request #7148: [Frontend][Tensorflow] Sparse_Dense Op CSR scheduling issue resolved for Cuda & X86
Posted by GitBox <gi...@apache.org>.
ANSHUMAN87 commented on a change in pull request #7148:
URL: https://github.com/apache/tvm/pull/7148#discussion_r547985023
##########
File path: python/tvm/topi/x86/sparse.py
##########
@@ -28,15 +28,17 @@ def schedule_sparse_dense(outs):
def _callback(op):
simd_width = get_fp32_len()
- if op.tag == "sparse_dense_csrmm" and op != outs[0].op:
- (_, v_i) = s[op].op.axis
- s[op].vectorize(v_i)
- (y_o, y_i) = s[outs[0].op].split(s[outs[0].op].op.axis[1], 2 * simd_width)
- s[op].compute_at(s[outs[0]], y_o)
- s[outs[0].op].vectorize(y_i)
- if op.tag == "sparse_dense_bsrmm":
+ if op.tag == "sparse_dense_csrmm_v2" or op.tag == "sparse_dense_csrmm_v1":
+ (y_o, y_i) = s[op].split(s[op].op.axis[1], 2)
+ fused = s[op].fuse(s[op].op.axis[0], y_o)
+ s[op].parallel(fused)
+ s[op].vectorize(y_i)
+ elif op.tag == "sparse_dense_bsrmm_v2" or op.tag == "sparse_dense_bsrmm_v1":
Review comment:
Basically the difference is whether lhs is sparse or rhs is sparse.
May be we change name like "sparse_dense_bsrmm_v2" --> "sparse_dense_bsrmm_sparse_rhs" ?
----------------------------------------------------------------
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] [tvm] ANSHUMAN87 commented on a change in pull request #7148: [Frontend][Tensorflow] Sparse_Dense Op CSR scheduling issue resolved for Cuda & X86
Posted by GitBox <gi...@apache.org>.
ANSHUMAN87 commented on a change in pull request #7148:
URL: https://github.com/apache/tvm/pull/7148#discussion_r557520145
##########
File path: tests/python/topi/python/test_topi_sparse.py
##########
@@ -530,5 +557,6 @@ def test_sparse_dense_padded_alter_op():
test_sparse_transpose_csr()
test_sparse_dense_padded_cuda()
test_sparse_dense_padded_alter_op()
Review comment:
Yes right. But i was thinking to keep that as well, so that covers one plus scenario, what do u think?
----------------------------------------------------------------
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