<table border="1" cellspacing="0" cellpadding="8">
    <tr>
        <th>Issue</th>
        <td>
            <a href=https://github.com/llvm/llvm-project/issues/136727>136727</a>
        </td>
    </tr>

    <tr>
        <th>Summary</th>
        <td>
            AMDGPU misses optimization on check-all-workitem-ids are 0 pattern
        </td>
    </tr>

    <tr>
      <th>Labels</th>
      <td>
            backend:AMDGPU,
            missed-optimization
      </td>
    </tr>

    <tr>
      <th>Assignees</th>
      <td>
      </td>
    </tr>

    <tr>
      <th>Reporter</th>
      <td>
          arsenm
      </td>
    </tr>
</table>

<pre>
    The device libraries include this pattern to check if all workitem IDs are 0.

```
// RUN: clang -target amdgcn-amd-amdhsa -S -O3 -mcpu=gfx900 -nogpulib < %s
bool
choose_one_workgroup_workitem(void)
{
    return (__builtin_amdgcn_workitem_id_x() | __builtin_amdgcn_workitem_id_y() | __builtin_amdgcn_workitem_id_z()) == 0;
}
```

https://github.com/ROCm/llvm-project/blob/662bae8d56ae5ba900a81b468936f47769b0fc2d/amd/device-libs/ockl/src/cg.cl#L46


This is equivalent to checking x == 0 && y == 0 && z == 0. If we codegen this, we see:

```
        v_and_b32_e32 v0, 0x3ff, v31
        v_bfe_u32 v1, v31, 20, 10
        v_bfe_u32 v2, v31, 10, 10
        v_or3_b32 v0, v0, v2, v1
        v_cmp_eq_u32_e32 vcc, 0, v0
        v_cndmask_b32_e64 v0, 0, 1, vcc
        s_setpc_b64 s[30:31]
```

In the function ABI, the work item IDs are packed into v31. We should be able to just check v31 == 0, so this would shrink to

```
        v_cmp_eq_u32_e32 vcc, 0, v31
        v_cndmask_b32_e64 v0, 0, 1, vcc
        s_setpc_b64 s[30:31]

```

</pre>
<img width="1" height="1" alt="" src="http://email.email.llvm.org/o/eJysVE2P2zYQ_TX0ZSCDGkqyddBBa8PFAm1TpAl6FChyLDGmREeknE1-fUHZ3k2y7aKHArIskG9m3nw96b3pRqKK5Q8s36_kHHo3VXLyNA6r1umv1YeeQNPFKAJr2klOhjyYUdlZE4TeeDjLEGgaIThQPakTmCNIa-GLm04m0ACPew9yIuBrxuv4FPz28JrhgeEB3n_8nYkalJVjB0mQU0cB5KA7NSZy0PHXewnJn5C8E5AM6jwzse-OTyXnkIyuO8_WtMDEDhjmnvG6dc4yXqveOU-NG6mJfLrJzefmzozh9uKMZlhGJpsHxmsAgInCPI3AcNs07WxsMGNz5fJs2RjdPDHcMiyBbXbwJvDrfwV-uwIXrNgzsQfOxMPCbf9z3Xjdh3D2TNxK2JnQz-1auYHh4f27Xfyz9jIk58l9IhUYHlrrWoaHosBW0lbnhaS8lSXncpu2WbEtRXHMNpuibPlRoWZ4kEN8X9ufWNN6hgenTpbhwU-K4UF1a2UZil-z4tZbXn-IQ2E80OfZXKSlMTyPhhk7eHrODRgWDAv4-urk2_PJGh6P8IVAOU0djcvEMdzFI08Us389Ury8NHLUTSuwIYFw4dGCP4njMX5cRHpHtUdq5ohI7xe4A1zgKX8Fwu9A6U8gN4kY7xbr9r4aPEdTw7mhz9HXlZZSC6-bwR006kH605V8kd3JL9EWpFJXqG88hbNq2iIDz_IHwZmoRcryfxiVx1g4guM8qmDcCPXDY_QVz-L8wQ9repbqRBrMGFzMdg1_EfjezVZDSyBbS7Gfn2Yfbvt-EenLvOIOvLsKw5fFxveTGU8Q3L906q2qvHTqfyrLj_FXuhK6FKVcUZVuslyIUhTZqq_KfCukFFxtNAq9KdtUIJeSciyPeSaLlamQY84zxLTAErO1LDXJQmc83yil0oJlnAZp7Dru4NpN3cp4P1OVimKDm5WVLVm_yC5iGws-aibq-rf9L398ZBgHhyEOxnvSiTsHM5hvMrYu3uX71VQtu93OnWcZt8YH_xIpmGCpuvqCxYeH732AG6-dS6S1yV2AEqNvMn1X9NU82eoNlYkBX6vMkmcUiluqlwr_DgAA__-utOtK">