[Vulkan] Backport SPIR-V fixes from apache/tvm main#35
Open
kabu1204 wants to merge 2 commits intotile-ai:tilelang_mainfrom
Open
[Vulkan] Backport SPIR-V fixes from apache/tvm main#35kabu1204 wants to merge 2 commits intotile-ai:tilelang_mainfrom
kabu1204 wants to merge 2 commits intotile-ai:tilelang_mainfrom
Conversation
This PR fixes a regression in vulkan codegen where we need to create phi node after the start label.
apache#18914) SPIR-V codegen currently emits `ArrayStride` and `Offset` decorations for non-interface allocations in `GetStructArrayType()`. That is correct for descriptor-backed interface blocks, but not for static workgroup allocations. I hit this while bringing up tilelang vulkan shared memory allocation path: the vulkan validation rejected shaders that used shared memory lowered through this path: ``` tvm.error.InternalError: Check failed: res == SPV_SUCCESS (-10 vs. 0) : index=44 error:[VUID-StandaloneSpirv-None-10684] Invalid explicit layout decorations on type for operand '25[%_ptr_Workgroup__struct_24]' %A_shared = OpVariable %_ptr_Workgroup__struct_24 Workgroup ``` FIX: This PR keeps layout decoration for interface blocks, and skips for non-interface allocations such as static shared/workgroup memory. A new compile-only test is added for this. One possible concern is that there's already a pre-existing test using `fetch_to_shared`.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Necessary fixes for preparation of vulkan backend.
This PR cherry-pick two fixes from apache/tvm main:
739027d [SPIRV] Fix forloop codegen in vulkan (#18715) — fixes OpPhi placement in loop headers so the OpPhi instruction appears immediately after OpLabel per SPIR-V spec. Without this, any kernel with a serial loop triggers SPIR-V validation error: "OpPhi's number of incoming blocks (2) does not match block's predecessor count (0)".
d33630c [Vulkan] Avoid explicit layout decoration on non-interface allocations (#18914) — restricts ArrayStride/Offset decorations to descriptor-backed interface blocks only, avoiding validation error "Invalid explicit layout decorations on type" for Workgroup storage class allocations (e.g. shared memory).