Skip to content

Scatter on Cuda - #6533

Merged
zhiics merged 4 commits into
apache:mainfrom
mbrookhart:mbrookhart/scatter_cuda_schedule
Oct 28, 2020
Merged

Scatter on Cuda#6533
zhiics merged 4 commits into
apache:mainfrom
mbrookhart:mbrookhart/scatter_cuda_schedule

Conversation

@mbrookhart

@mbrookhartmbrookhart commented Sep 22, 2020

Copy link
Copy Markdown
Contributor

I was unable to get scatter working with te schedules on cuda due to the two loops that update values in place, so I resorted to using ir_builder directly.

Attempts to better parallelize the algorithm have created some strange behavior. If I can get a correct and faster implementation to work, I'll submit another PR.

Thanks to @tkonolige for very helpful discussions. @notoraptor, this may be useful for your scatter_add work.

cc @zhiics

@zhiics

zhiics commented Sep 22, 2020•
edited
Loading

Copy link
Copy Markdown
Member

cc @Laurawly@vinx13 @icemelon9

zhiics
zhiics reviewed Sep 22, 2020
View reviewed changes
Comment threadpython/tvm/topi/cuda/scatter.py Outdated
@mbrookhart
mbrookhartforce-pushed the mbrookhart/scatter_cuda_schedule branch from 2e3c4d4 to 963147aCompareOctober 1, 2020 15:16
Laurawly
Laurawly reviewed Oct 2, 2020
View reviewed changes
Comment threadpython/tvm/topi/cuda/scatter.py

with ib.new_scope():
bx = te.thread_axis("blockIdx.x")
ib.scope_attr(bx, "thread_extent", 1)

@LaurawlyLaurawlyOct 2, 2020

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

if bx is not used in this context, we can remove the above two lines.

@mbrookhartmbrookhartOct 5, 2020•
edited
Loading

Copy link
Copy Markdown
ContributorAuthor

Choose a reason for hiding this comment

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

If I don't define at least one block in the scope, the generated cuda code fails to compile.

zhiics
zhiics reviewed Oct 2, 2020
View reviewed changes
Comment threadtests/python/relay/test_op_level3.py
@mbrookhart
mbrookhartforce-pushed the mbrookhart/scatter_cuda_schedule branch from 963147a to 488ad3dCompareOctober 8, 2020 15:31
zhiics
zhiics reviewed Oct 8, 2020
View reviewed changes
Comment threadpython/tvm/topi/cuda/scatter.py
@tqchen
tqchen changed the base branch from master to mainOctober 11, 2020 18:19
zhiics
zhiics approved these changes Oct 19, 2020
View reviewed changes

@zhiicszhiics left a comment

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

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

LGTM. just one minor nitpick. @Laurawly could you please take another look?

Comment threadpython/tvm/topi/cuda/scatter.py Outdated
@masahi

masahi commented Oct 27, 2020•
edited
Loading

Copy link
Copy Markdown
Member

@zhiics Can we merge this? I want to send a PR to enable scatter_add on GPU, building off of this PR

@zhiics

zhiics commented Oct 27, 2020

Copy link
Copy Markdown
Member

yeah, let's merge this. before that, we probably want to rebase again just in case bad things would happen

@mbrookhart
mbrookhartforce-pushed the mbrookhart/scatter_cuda_schedule branch from 7881b2e to 03de673CompareOctober 28, 2020 00:03
@mbrookhart

mbrookhart commented Oct 28, 2020

Copy link
Copy Markdown
ContributorAuthor

I rebased, I'll keep an eye on it this evening to make sure it still passes CI. Thanks!

@zhiics
zhiics merged commit 3d624ec into apache:mainOct 28, 2020
@zhiics

zhiics commented Oct 28, 2020

Copy link
Copy Markdown
Member

Thanks @mbrookhart@Laurawly@tkonolige@masahi

@mbrookhart
mbrookhart deleted the mbrookhart/scatter_cuda_schedule branch October 28, 2020 14:47
trevor-m pushed a commit to trevor-m/tvm that referenced this pull request Oct 29, 2020
Matthew BrookhartTrevor Morris
Scatter on Cuda (apache#6533)
1acdf12
* working cuda scatter
fix lint
fix pylint again
* cuda scatter with threading
* add dynamic shape tests
* remove unused variable
@masahimasahi mentioned this pull request Nov 5, 2020
Merged
trevor-m pushed a commit to trevor-m/tvm that referenced this pull request Dec 2, 2020
Matthew BrookhartTrevor Morris
Scatter on Cuda (apache#6533)
eb2e0dc
* working cuda scatter
fix lint
fix pylint again
* cuda scatter with threading
* add dynamic shape tests
* remove unused variable
trevor-m pushed a commit to trevor-m/tvm that referenced this pull request Dec 4, 2020
Matthew BrookhartTrevor Morris
Scatter on Cuda (apache#6533)
d75a7ca
* working cuda scatter
fix lint
fix pylint again
* cuda scatter with threading
* add dynamic shape tests
* remove unused variable
trevor-m pushed a commit to neo-ai/tvm that referenced this pull request Dec 4, 2020
Matthew Brookhart@trevor-m
Scatter on Cuda (apache#6533)
fb1bbb0
* working cuda scatter
fix lint
fix pylint again
* cuda scatter with threading
* add dynamic shape tests
* remove unused variable
@junrushaojunrushao mentioned this pull request Nov 1, 2021
Closed
Sign up for freeto join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

@tkonoligetkonoligetkonolige left review comments
@LaurawlyLaurawlyLaurawly left review comments
@zhiicszhiicszhiics approved these changes

Assignees

No one assigned

Labels

status: accepted

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

6 participants

@mbrookhart@zhiics@masahi@tkonolige@Laurawly@ZihengJiang