Skip to content

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