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

Stanislav Mekhanoshin via Phabricator via llvm-commits llvm-commits at lists.llvm.org
Mon Aug 29 12:25:18 PDT 2022


rampitec added a comment.

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

> In D132511#3753491 <https://reviews.llvm.org/D132511#3753491>, @bcl5980 wrote:
>
>> 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]
>> warp 2 should be [63, 1] to [64, 1]
>> warp1&2's threadIdx.x / 64 should be still divergent
>
> Yes, you are right, thanks! Looks like I need to limit it to the case when there is `amdgpu-no-workitem-id-y` attribute on the function.

D132879 <https://reviews.llvm.org/D132879> limits it.


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