[PATCH] D132511: [AMDGPU] Detect uniformness of TID / wavefrontsize

Stanislav Mekhanoshin via Phabricator via llvm-commits llvm-commits at lists.llvm.org
Fri Aug 26 10:16:18 PDT 2022


rampitec added a comment.

In D132511#3749004 <https://reviews.llvm.org/D132511#3749004>, @yaxunl wrote:

> what if the operation happens in divergent control flow, e.g
>
>   int x = 0;
>   x = threadIdx.x > 32 : threadIdx/64 : 0;
>
> will this patch still work?

Regardless of the CFG the value of TID / 64 is always uniform. The value of 'x' here is another value, derived from that uniform value. It is no different from `x = cc ? sgpr0 : sgpr1;` LHS and RHS are uniform, but 'x' is not.


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

https://reviews.llvm.org/D132511



More information about the llvm-commits mailing list