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

[chenglin.bi via Phabricator via llvm-commits]

bcl5980 added a comment.

In D132511#3752207 <https://reviews.llvm.org/D132511#3752207>, @rampitec wrote:

> 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.

What if the blockDim.x is not 64, like 65, blockDim.y is not 1
for example:
the workgroup shape is <65, 2, 1>
warp 0 should be [0,0] to [63, 0]
warp 1 should be [64,0] to [62, 1]
warp1's threadIdx.x / 64 should be still divergency


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D132511