[PATCH] D132511: [AMDGPU] Detect uniformness of TID / wavefrontsize
chenglin.bi via Phabricator via llvm-commits
llvm-commits at lists.llvm.org
Sat Aug 27 05:10:10 PDT 2022
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
More information about the llvm-commits
mailing list