| Issue |
163835
|
| Summary |
[AMDGPU] Add cases for or emulate wave_id
|
| Labels |
new issue
|
| Assignees |
|
| Reporter |
krzysz00
|
The `amdgcn.wave.id` intrinsic reads a special-purpose register that holds the wave ID (that is `(threadIdx.x + blockDim.x * (threadId.y + blockDim.y * threadId.z)) / waveSize`) on gfx12. It's undefined on other architectures.
However
1. This feature is still available on other architectures (ex. gfx942 has the wave ID in ttmp11)
2. In cases where it isn't supported, the wave ID intrinsic can still be expanded to a computation of the wave ID instead of failing compilation
![]()
_______________________________________________
llvm-bugs mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-bugs
