Skip to content

[Fix] Tensor core type issue for dense #7187

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom
Merged
junrushao merged 3 commits into from
Jan 2, 2021
Merged

[Fix] Tensor core type issue for dense #7187

junrushao merged 3 commits into from
Jan 2, 2021

Conversation

@Laurawly
Copy link
Contributor

@Laurawly Laurawly commented Dec 31, 2020

As mentioned in PR #7146, previous tensor core code lacks type check and this PR is trying to fix that.
cc @jcf94 @merrymercy @Meteorix

junrushao
junrushao approved these changes Dec 31, 2020
View reviewed changes
Copy link
Member

@junrushao junrushao left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM

@merrymercy
Copy link
Member

merrymercy commented Dec 31, 2020

Do we allow data.dtype == "float32"?
The problem comes from this downcast, which may cause accuracy problem.

tvm/python/tvm/topi/cuda/dense_tensorcore.py

Lines 72 to 73 in c02c9c5

data_16 = te.compute((batch, in_dim), lambda b, i: data[b, i].astype("float16"))
weight_16 = te.compute((out_dim, in_dim), lambda o, i: weight[o, i].astype("float16"))

@Laurawly
Copy link
Contributor Author

Laurawly commented Dec 31, 2020

Do we allow data.dtype == "float32"?
The problem comes from this downcast, which may cause accuracy problem.

tvm/python/tvm/topi/cuda/dense_tensorcore.py

Lines 72 to 73 in c02c9c5

data_16 = te.compute((batch, in_dim), lambda b, i: data[b, i].astype("float16"))
weight_16 = te.compute((out_dim, in_dim), lambda o, i: weight[o, i].astype("float16"))

Good catch, just removed it.

@junrushao junrushao merged commit 6258fae into apache:main Jan 2, 2021
@Meteorix Meteorix mentioned this pull request Jan 11, 2021
Merged
tkonolige pushed a commit to tkonolige/incubator-tvm that referenced this pull request Jan 11, 2021
@Laurawly @tkonolige
[Fix] Tensor core type issue for dense (apache#7187)
3892bd9 
* fix tc type issue for dense

* fix lint

* rm float 32

Co-authored-by: Leyuan Wang <ziyu.guo@bytedance.com>
TusharKanekiDey pushed a commit to TusharKanekiDey/tvm that referenced this pull request Jan 20, 2021
@Laurawly
[Fix] Tensor core type issue for dense (apache#7187)
105a9f0 
* fix tc type issue for dense

* fix lint

* rm float 32

Co-authored-by: Leyuan Wang <ziyu.guo@bytedance.com>
trevor-m pushed a commit to neo-ai/tvm that referenced this pull request Jan 21, 2021
@Laurawly @trevor-m
[Fix] Tensor core type issue for dense (apache#7187)
eaa2cae 
* fix tc type issue for dense

* fix lint

* rm float 32

Co-authored-by: Leyuan Wang <ziyu.guo@bytedance.com>
electriclilies pushed a commit to electriclilies/tvm that referenced this pull request Feb 18, 2021
@Laurawly @electriclilies
[Fix] Tensor core type issue for dense (apache#7187)
4349f83 
* fix tc type issue for dense

* fix lint

* rm float 32

Co-authored-by: Leyuan Wang <ziyu.guo@bytedance.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

@junrushao junrushao junrushao approved these changes

Assignees

No one assigned

Labels

None yet

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

3 participants

@Laurawly @merrymercy @junrushao