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

    <tr>
        <th>Summary</th>
        <td>
            [AMDGPU][MC][GFX940] Incorrect disassembly of MFMA opcodes
        </td>
    </tr>

    <tr>
      <th>Labels</th>
      <td>
            bug,
            backend:AMDGPU,
            mc
      </td>
    </tr>

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

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

<pre>
    MFMA opcodes common to all MI GPUs are disassembled incorrectly. Each GPU has its own version of these opcodes, but disassembler selects a wrong table and decodes GFX940 opcodes as if they were GFX908 ops.

An example of failed test:

    0x0a,0x00,0xcd,0xd3,0x02,0x09,0x1a,0x0c

Expected output:

    v_mfma_f32_16x16x16f16 v[10:13], a[2:3], v[4:5], v[6:9]    // V_MFMA_F32_16X16X16F16_gfx940_vcd

Actual result:

    v_mfma_f32_16x16x16f16 a[10:13], a[2:3], v[4:5], a[6:9]    // V_MFMA_F32_16X16X16F16_vi

A fix is pending.

</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJyVU0uPmzAQ_jVwGTUy5hE4cMhuNqs9ROql1d6QwQbcGoxsk8e_75gkWyr1EmQxzEPf980wrjW_lsfDcQd6ajQXFho9DHoEp4EpBccPeP_-wwIzAri0zFox1EpwkGOjjRGNU9cNvLGm93XQMwvSWdDnEU7CWIlIugXXCyseDAF9hXp2azgDVijEQh44Gz124BiGgY0cuLjpej98Fgn5kumJFuArnAWK82mSY9puArIPyO723o0gLmyYEAx1tEx67U5YF8S7dR3gQy6EoTg0ZDENXwyPb0F6M8ViontlswZ5u0zYBBLo2U3z_yhO1dAOrGpjWkXZZTltlMEpSF8igvVRHKR7PyCGEYqBh-8rEvTTlZ-hX6DvgQN6wAM_K_8vq8OC_7mcQ5RVXXvB2VUn7Gg9m8bNTIERdlZPiGVPi2VPiT3Jf0RCKy8gLUxi5HLs7n83FGWUpQWJoyKPQ17GvIgLFjrplCiRbnfc4z56_vTl-Hqztw3yGj4ey7tawqtfkPVNCGejyt65yfrZLJI76fq53uAVQUep08N8m4z-hXDoSmtnv-OHNNvmSdiXGWMZjbYR5XmWbmNS1zXbZi1LSEFJlIpQsVoo60UHlNZzh28_NO-w5jd2jfT3dr4yQ-O_030oS0ooJVuaR3kUE7ppSJ5skbmuE1q0RRMkRAy49huvc6NNF5pykYxMFpNKWmf_JnEYshuFuMvhkxG6NuzOxmbXa1M-ouHSbLl0-gcACz8u">