KhronosGroup / KhronosGroup/SPIRV-LLVM-Translator

Bugs when there are two consecutive get_local_id

Open
#978 7 comments 0 reactions 0 assignees View on GitHub
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

Open the contributing 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

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.