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

    <tr>
        <th>Summary</th>
        <td>
            [AMDGPU] generates v_cndmask/lshlrev for uniform select between 0 and a power of 2
        </td>
    </tr>

    <tr>
      <th>Labels</th>
      <td>
            new issue
      </td>
    </tr>

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

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

<pre>
    Here is my test kernel:
```cpp
static __global__ void test(const int32_t input, int32_t * result) {
    const int data = input ? 32 : 0;
    int output;

    asm volatile("s_mov_b32 %0, %1" : "=s"(output) : "s"(data));

 result[threadIdx.x] = output;
}
```
and here's the asm it compiles to for gfx900 on rocm 6.0.0, where it fails to assemble because of an invalid operand.
```asm
        s_load_dword s2, s[4:5], 0x0
        s_load_dwordx2 s[0:1], s[4:5], 0x8
        v_lshlrev_b32_e32 v0, 2, v0
        s_waitcnt lgkmcnt(0)
        s_cmp_lg_u32 s2, 0
        s_cselect_b64 s[2:3], -1, 0
        v_cndmask_b32_e64 v1, 0, 1, s[2:3]
        v_lshlrev_b32_e32 v1, 5, v1
        ;;#ASMSTART
        s_mov_b32 s2, v1
        ;;#ASMEND
        v_mov_b32_e32 v1, s2
        global_store_dword v0, v1, s[0:1]
```
If either of the values being selected is changed so that either one isn't 0, or the other isn't a power of 2, instead it compiles to this, which works.
```asm
        s_load_dword s2, s[4:5], 0x0
        s_load_dwordx2 s[0:1], s[4:5], 0x8
        v_lshlrev_b32_e32 v0, 2, v0
        s_waitcnt lgkmcnt(0)
        s_cmp_eq_u32 s2, 0
        s_cselect_b32 s2, 0, 33
        ;;#ASMSTART
        s_mov_b32 s2, s2
        ;;#ASMEND
        v_mov_b32_e32 v1, s2
        global_store_dword v0, v1, s[0:1]
```

</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJzcVcGO2zYQ_Rr6MoghDyXZPvjgjeM2hxRFk54FShxJzFKkS1Ly5u8LUrLX3aYt0EuBAgIEcebNvEc-DYX3qjNEB1Y8seK0EmPorTt8MJ0yynwVuKqt_Hb4kRyB8jB8g0A-wDM5Q5rxI8tOLDuyMpuf5nKZV3wQQTVQVZ22tdBVBZNVMoEZ7hprfABlAscqvi9jYPj-vsDwCI78qAPDPbDt01wTAOCOBCmCAMZPMxwYPwNHYPwIGeMPiJhrxxBb3JZfg8IPMFktgtLEcMcQfTXYqapjKSyyyIphsWE4l2aIjJ98fOFuqRopzqFlPTJjuI_Pm46LqOIp9I6E_Chf1i-sOCUZbzluT2_2dv4URkJPjhhuPYSekgQVoLHDRWnyECy01kHXvuyzDKwBZ5sBynW2TmqufTrKAK1QOmUL72moNUFNjRg9gW1BGFBmElpJsBdywsj1GzbCDzdpe19pK2Qlr9ZJ8BjbeFY85YwfC1ac4nf2kn03_QVTasb4cbOk_hm6u0OnSvteO0onVBFHmJKq1HN6bHEVKjQmgO6eh8ZE02XxQF4TmuFS6a4aOS6UH9GNJ01NqOoyT3yQ8SNf-Lzb_DF7qhojB-GfZ05lDtOSge9hc1N0r_A3UlJykaRs7nnREfyJIT9-_vT5y_GXLw80b1adBfwF6MNPj00XyENDj_fw8rP6YB0txznv73TXcT-p75nzYwukQk8ueiiacxJ6JA81KdPBvKck4yBpemE6kuAthF6EO8zEMWMYbsO8f9alOjZFbxEBF3udm-A8N3wgId_-BqFXfra8anq4Wvfs_5cupt_-ycUPUXwPnP8Lcz245D8y10oeuNzzvVjRYbPdYIEbjuWqP1CZ5XXdbsu63eW7nZA5b_blnhpZ5A3J_UodMMM8y7MtZohZuS54yVssi7Ytc5TtjuUZDULptdbTsLauWynvRzrstnu-W2lRk_bphkQ0dIUUjLO-OK3cIWLe1WPnWZ5p5YN_rRJU0OlqPX46_fDzr3HUd2TIiUAe7nOD4XnxQ5rco1GtdcPyu0BN4UpkIIM4-h-tvxqdPvQhXHy8iPHM8Nyp0I_1urFDLKqn2-vdxdmv1ASG50TeMzwncb8HAAD__xJCTOE">