Skip to content

[Unity][Dlight] Rule matmul avoiding blockIdx.z#15333

Merged
yzh119 merged 1 commit intoapache:unityfrom
MasterJH5574:unity-dev/2023-07-16-matmul-blockidxz
Jul 17, 2023
Merged

[Unity][Dlight] Rule matmul avoiding blockIdx.z#15333
yzh119 merged 1 commit intoapache:unityfrom
MasterJH5574:unity-dev/2023-07-16-matmul-blockidxz

Conversation

@MasterJH5574
Copy link
Contributor

Prior to this PR, the matmul rule of dlight binds loops to blockIdx.z. However, not every device supports this blockIdx dimension (for example, WebGPU does not support blockIdx.z), which makes dlight fails to apply and build.

Therefore, this PR fuses the blockIdx.z loop with other blockIdx loop.

@tvm-bot
Copy link
Collaborator

tvm-bot commented Jul 16, 2023

Thanks for contributing to TVM! Please refer to the contributing guidelines https://tvm.apache.org/docs/contribute/ for useful information and tips. Please request code reviews from Reviewers by @-ing them in a comment.

Generated by tvm-bot

1 similar comment
@tvm-bot
Copy link
Collaborator

tvm-bot commented Jul 16, 2023

Thanks for contributing to TVM! Please refer to the contributing guidelines https://tvm.apache.org/docs/contribute/ for useful information and tips. Please request code reviews from Reviewers by @-ing them in a comment.

Generated by tvm-bot

@MasterJH5574
Copy link
Contributor Author

cc @Hzfengsy

Copy link
Member

@yzh119 yzh119 left a comment

Choose a reason for hiding this comment

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

LGTM, thank you for the patch.

Prior to this PR, the matmul rule of dlight binds loops to `blockIdx.z`.
However, not every device supports this blockIdx dimension (for example,
WebGPU does not support `blockIdx.z`), which makes dlight fails to
apply and build.

Therefore, this PR fuses the `blockIdx.z` loop with other `blockIdx`
loop.
@MasterJH5574 MasterJH5574 force-pushed the unity-dev/2023-07-16-matmul-blockidxz branch from 3ae8a6e to f0c62de Compare July 17, 2023 03:46
@yzh119 yzh119 merged commit 66d3957 into apache:unity Jul 17, 2023
junrushao pushed a commit that referenced this pull request Jul 18, 2023
Prior to this PR, the matmul rule of dlight binds loops to `blockIdx.z`.
However, not every device supports this blockIdx dimension (for example,
WebGPU does not support `blockIdx.z`), which makes dlight fails to
apply and build.

Therefore, this PR fuses the `blockIdx.z` loop with other `blockIdx`
loop.
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.

3 participants