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

    <tr>
        <th>Summary</th>
        <td>
            Invalid register allocated for v_mad_u64_u32 on gfx1100
        </td>
    </tr>

    <tr>
      <th>Labels</th>
      <td>
      </td>
    </tr>

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

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

<pre>
    Given the following source code:
```cpp
__global__ void test(unsigned long long* out) {
    *out = __builtin_amdgcn_workgroup_id_x() * __builtin_amdgcn_workgroup_size_x() + __builtin_amdgcn_workitem_id_x();
}
```

When generating assembly for this snippet for gfx1100 and assembling it to machine code, an error is returned:
```shell
$ /opt/rocm/llvm/bin/clang++ -x hip test.cpp -std=c++17 --offload-arch=gfx1100 --cuda-device-only -S -o test.s -O3
$ /opt/rocm/llvm/bin/clang++ -otest test.s -target amdgcn-amd-amdhsa -mcpu=gfx1100
test.s:15:39: error: destination must be different than all sources
        v_mad_u64_u32 v[1:2], null, s15, s2, v[0:1]
$
```

Note that the same is not the case for gfx10:
```
$ /opt/rocm/llvm/bin/clang++ -x hip test.cpp -std=c++17 --offload-arch=gfx1030 --cuda-device-only -S -o test.s -O3
$ /opt/rocm/llvm/bin/clang++ -otest test.s -target amdgcn-amd-amdhsa -mcpu=gfx1030
$
```
It seems that, at least in disassembly, a different register is allocated as the result:
```nasm
v_mad_u64_u32 v[0:1], null, s8, s2, v[0:1]
```

It looks like the final instruction is also different, even apart from the s8/s15 (the last 4 bits of the encoded instruction)
```
v_mad_u64_u32 v[0:1], null, s8, s2, v[0:1] // gfx1030: D5767D00 04000408
v_mad_u64_u32 v[1:2], null, s15, s2, v[0:1]  // gfx1100: D6FE7C01 0400040F
```

It looks like this is either an invalid register being assigned, or a problem with the disassembler.
</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJzMVsFu4zYQ_Rr6MpBBkZItH3xw1nWxl_awhx4FShxLbChSIEfeTb--oGTH7iYNsC1QNIhJk5rhDN97M7KK0XQOcc_KJ1YeV2qi3of9F4fPFLxbNV6_7H82F3RAPcLZW-u_GtdB9FNoEVqvkckD40fGD2zDl_92HJeduu6sb5Sta7h4o4EwEhPV5OaoGqx33TwwcQA_ERM7YNunxRkAgImDnwiYPEJdN5OxZFytBt21rv7qw3MX_DTWRtffmKhmb3H4yDKaP_DB9ul9W0M4PBzK5DUjtj1-d9Prch5_69FBhw6DogSRihGHxr7A2Qeg3kSIzowj0rzRnb_lOeegnL5ZJidDQB4G1fbGXeEVn0A5wBB8ABMhIE3BoX4Le-zR2uueKICJkx-JiVPw7cDEydpLmhrjmDi1ViXUnxIG2TfozTiTs27HEbJImsljuzzOt5Bl_ny2XulMhbZn8nhLPsvaSatM48W0mHlnXyD7AplfzoqQ_Sr_ST4-ub-eQSp0SLDwk6lBp08fFWRDO073bJZIixeTh7xk8iB3TB4W7NIXjZGMU2S8g2GKBA2CNuczBnQE1CsHytqruONdhunvUg9K19OmqCcp4MLKp5zJg2DlMTHkJmvTHPNynkQakxFPqSSjGw4fKOgXT5iyoLnYohowEe78sm5VxFfp8Lf0_2fMc_l_Yp5L_jG2nwki4hBnZOdqIrCoIoFxoE281en86EEOATsTCeeiU9b6VhGmWp3JCBgnS29JcCoOy9ZbvbxK4VEv1YdyeU8lnwms988RrHnGpS0bpywYFylM7SzuOeXo77dJ52Nq42pUgeAc_LCIrGLiFPMSmKjS2iZcCmgMRfDn2QRdakP68fzUFN9L8N_dOemEiRPcWJUHOJbbzfbIOfCCc17w6u_i_FAtwmOk1DlSpM3pp-0nnt8inX6AARMT4Giox5B6tXEXZY2-K6jB6wthfu2ldHwABWPwjcUBvhrqZ6jvasSwXum91Du5Uyvc55stl7LMpVz1-7Io1K6sSqzOkkspeMWbHLFt1ea82270yuwFFyIXfJPLsip2a13gdtu0G2z0FosKWcFxUMauUyWufehWJsYJ9-VOltXKqgZtvP0kCPtklDVTF1nBrYkU725kyOL-8_f3vZdL6ld_Zcu7G-yrKdh9TzSmdr0w0hnqp2bd-oc2MYcfg_8d29RD5kwjE6c52T8DAAD__wT3p9w">