KhronosGroup / KhronosGroup/SPIRV-LLVM-Translator
Bugs when there are two consecutive get_local_id
- Dominant language
- LLVM
- Stars
- 625
- Forks
- 279
- Avg merge
- 3d 5h
- Merged PRs (30d)
- 34
Description
Hi, I have tried to support converting NVPTX into SPIRV and then convert into LLVM IR. And when I convert the following code
`int row = threadIdx.y;`
, it can generate the corret LLVM bitcode
`%0 = call spir_func i64 @_Z12get_local_idj(i32 1) #1`
However, when the input code is
`int row = threadIdx.y;
int col = threadIdx.x;`
The generated LLVM will be
`%0 = call spir_func i64 @_Z12get_local_idj(i32 1)`
`%2 = load <3 x i64>, <3 x i64> addrspace(5)* @__spirv_BuiltInLocalInvocationId1`
`%3 = extractelement <3 x i64> %2, i32 0 `
So I think when converting a SPIRV file into LLVM IR, when there are two consecutive get_local_id, it will raise error.
Would you please give me some suggestions?
Thanks in advance!
Contributor guide
Research direction
Reproduce the NVPTX-to-SPIR-V-to-LLVM IR conversion for the two snippets in the issue and compare it with the single get_local_id case. Inspect the generated calls, load, and extractelement instructions to determine why consecutive calls produce different IR; done when the two-variable example translates without the reported error or malformed output.
Written by the indexing model from the issue text.
Assessment
- Domain
- compilers
- Issue type
- Bug
- Difficulty
- 4/5
- Estimated time
- 3-5 days
- Activity status
- Stale
- Clarity
- Needs clarification
- Newbie friendliness
- 25/100