[PATCH] D153047: [AMDGPU] Stop replacing amdgcn.ballot(1) with amdgcn.s.getreg(exec)

Matt Arsenault via Phabricator via llvm-commits llvm-commits at lists.llvm.org
Thu Jun 15 09:56:25 PDT 2023


arsenm accepted this revision.
arsenm added a comment.
This revision is now accepted and ready to land.

> There is also a correctness problem that getreg can read from exec but it is currently not marked as convergent.

You can mark the call sites convergent, which this is doing. We still have __builtin_amdgcn_read_exec emitting read_register, but I was thinking we should move them to ballot(true)


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D153047/new/

https://reviews.llvm.org/D153047



More information about the llvm-commits mailing list