[AMDGPU] Add cases for or emulate wave_id

[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

[llvmbot]

@llvm/issue-subscribers-backend-amdgpu

Author: Krzysztof Drewniak (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

[llvmbot]

Hi!

This issue may be a good introductory issue for people new to working on LLVM. If you would like to work on this issue, your first steps are:

1. Check that no other contributor is working on this issue. If someone is assigned to the issue or claimed to be working on it, ping the person. After one week without a response, the assignee may be changed.
2. Leave a comment indicating that you are working on the issue, or just create a pull request after following the steps below. Mention this issue in the description of the pull request.
3. Fix the issue locally.
4. Run the test suite locally. Remember that the subdirectories under test/ create fine-grained testing targets, so you can e.g. use make check-clang-ast to only run Clang's AST tests.
5. Create a Git commit.
6. Run git clang-format HEAD~1 to format your changes.
7. Open a pull request to the upstream repository on GitHub. Detailed instructions can be found in GitHub's documentation. Mention this issue in the description of the pull request.

If you have any further questions about this issue, don't hesitate to ask via a comment in the thread below.

[llvmbot]

@llvm/issue-subscribers-good-first-issue

Author: Krzysztof Drewniak (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

[XChy]

I will give it a try.

[XChy]

I have implemented a local version. However, upon reading the ISA manual for CDNA3 and CDNA4, I found that there are two similar terms: wave_id_in_group and wave_id_in_workgroup. The former is a 5-bit integer in ttmp3, while the latter is a 6-bit integer in ttmp11. Is there any difference between them? Thanks for any information.

[arsenm]

For 1, yes we should just read the other target register. For 2, I'm not sure we should do that. The semantics of these intrinsics is supposed to be "give me this ABI register", not "give me the semantics of this value"

[shiltian]

Unlike gfx12, TTMPs are only initialized or loaded when the dispatch has global traps enabled on other architectures. That said, I'm not sure if we can reliably use TTMP11 for the lowering.

[zwu-2025]

First of all, I hope you don’t mind me asking — I’m just curious about this question about the PR:

  // With architected SGPRs, waveIDinGroup is in TTMP8[29:25].
  if (ST.hasArchitectedSGPRs()) {
    LLT S32 = LLT::scalar(32);
    auto TTMP8 = B.buildCopy(S32, Register(AMDGPU::TTMP8));
    auto LSB = B.buildConstant(S32, 25);
    auto Width = B.buildConstant(S32, 5);
    B.buildUbfx(DstReg, TTMP8, LSB, Width);
    MI.eraseFromParent();
    return true;
  }

  // GFX942/GFX950 has wave_id_in_workgroup in ttmp11
  if (ST.hasGFX940Insts()) {
    B.buildCopy(DstReg, Register(AMDGPU::TTMP11));
    MI.eraseFromParent();
    return true;
  }

  return false;

but from the document, is it only bits from 0 to 5 in TTMP11 is for the wave_id defined here?

[XChy]

but from the document, is it only bits from 0 to 5 in TTMP11 is for the wave_id defined here?

IIUC, as other bits are zeros, TTMP11 is zero-extended from the 6-bit wave_id to 32 bits.