Skip to content
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.

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

[PASS] Enhance gpu verify pass #1660

Merged
merged 1 commit into from
Aug 30, 2018
Merged

[PASS] Enhance gpu verify pass #1660

merged 1 commit into from
Aug 30, 2018

Conversation

merrymercy
Copy link
Member

Check the mismatch of axis lengths.
This error is easy to meet when we schedule cooperative fetching

cc @eqy

@eqy
Copy link
Contributor

eqy commented Aug 26, 2018

Do we have another way of supporting multi-axis bind with different lengths?
I think the typical CUDA style for doing this if the axes are independent is to just do

if (threadIdx.x ... < axis_size)
     do_something;

@merrymercy
Copy link
Member Author

merrymercy commented Aug 27, 2018

It is a problem of the compiler. But this pattern typically cannot make the performance better. So we can leave this problem later and use this verification pass to ensure correctness.

@yzhliu yzhliu self-assigned this Aug 27, 2018
@tqchen
Copy link
Member

tqchen commented Aug 30, 2018

I am going to merge this for now as this is an improvement over previous one. However, note that if there are two thread extent, there can actually be the different thread bound if a schedule contains multiple stages of kernels, a better way is to not check the string, but check by address IterVar

@tqchen tqchen merged commit 0c52378 into apache:master Aug 30, 2018
@merrymercy merrymercy deleted the gpu-verify branch August 30, 2018 09:06
@merrymercy
Copy link
Member Author

@tqchen My code can handle this case. I only do check inside a kernel

ZhennanQin pushed a commit to ZhennanQin/tvm that referenced this pull request Sep 6, 2018
FrozenGene pushed a commit to FrozenGene/tvm that referenced this pull request Dec 27, 2018
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Projects
None yet
Development

Successfully merging this pull request may close these issues.

4 participants