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

    <tr>
        <th>Summary</th>
        <td>
            [NVPTX] performance regression caused by register splitting/merging combined with `wgmma` instructions.
        </td>
    </tr>

    <tr>
      <th>Labels</th>
      <td>
            backend:NVPTX
      </td>
    </tr>

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

    <tr>
      <th>Reporter</th>
      <td>
          Artem-B
      </td>
    </tr>
</table>

<pre>
    Some users reported substantial performance regression (apparently in nvptx) attributed to register splitting/merging after recent LLVM changes.

> [kernel] keeps the WGMMA accumulator registers as <2 x float> and unpacks them right before passing them into WGMMA:
> ```
> ...
>   %571 = extractelement <2 x float> %458, i64 0, !dbg !10
>   %572 = extractelement <2 x float> %458, i64 1, !dbg !10
>   %573 = extractelement <2 x float> %457, i64 0, !dbg !10
>   %574 = extractelement <2 x float> %457, i64 1, !dbg !10
>   %575 = call { float, float, ... } asm sideeffect "{ .reg .pred p; setp.ne.b32 p, $130, 0; wgmma.mma_async.sync.aligned.m64n128k16.f32.f16.f16 {$0,$1,$2,$3,$4,...}, $128, $129, p, $131, $132, $133, $134; }\0A", "=f,=f,=f,=f,=f,...,0,1,2,3,4,...,l,l,n,n,n,n,n"(float %511, float %512, float %513, float %514, ..., float %568, float %569, float %570, float %571, float %572, float %573, float %574, i64 %506, i64 %510, i32 1, i32 1, i32 1, i32 0, i32 1) #5, !dbg !10
>   ...
> ```
> 

> Before those changes, NVPTX would always represent the <2 x float> as 2 32-bit registers and the extractelement ops would not be visible in the generated ptx. However, if I look at what happens now, I see this:
> ```
>         ...
>    mov.b64         {%r1533, %r1534}, %rd281;
>    mov.b64         {%r1535, %r1536}, %rd280;
>    mov.b64         {%r1537, %r1538}, %rd279;
>    mov.b64         {%r1539, %r1540}, %rd278;
>    mov.b64         {%r1541, %r1542}, %rd277;
>    // begin inline asm
>    { .reg .pred p; setp.ne.b32 p, 1, 0; wgmma.mma_async.sync.aligned.m64n128k16.f32.f16.f16 {%r1479,%r1480,%r1481,%r1482,%r1483,%r1484,%r1485,%r1486,%r1487,%r1488,%r1489,%r1490,%r1491,%r1492,%r1493,%r1494,%r1495,%r1496,%r1497,%r1498,%r1499,%r1500,%r1501,%r1502,%r1503,%r1504,%r1505,%r1506,%r1507,%r1508,%r1509,%r1510,%r1511,%r1512,%r1513,%r1514,%r1515,%r1516,%r1517,%r1518,%r1519,%r1520,%r1521,%r1522,%r1523,%r1524,%r1525,%r1526,%r1527,%r1528,%r1529,%r1530,%r1531,%r1532,%r1533,%r1534,%r1535,%r1536,%r1537,%r1538,%r1539,%r1540,%r1541,%r1542}, %rd235, %rd220, p, 1, 1, 0, 1; }
>    ...
> ```
> and this is a problem, because ptxas really doesn't like it when you touch registers that end up as WGMMA accumulators between a wgmma.fence and the wgmma instruction, even if it's for something as trivial as this. It emits a warning saying that it will insert additional waits to make sure this is safe, and this slows down the whole kernel. 
> 

</pre>
<img width="1" height="1" alt="" src="http://email.email.llvm.org/o/eJykl1-P46wVxj8NuTlay4Ad2xe5mD-ddqV3q0qt3vauwvaxQweDBXiy-fYVTiYnyXZ29lVHyfgh5pwfYHjAKgQ9WsQdKx9Z-bxRS9w7v3vwEacvj5vW9cfd392EsAT0ATzOzkfsISxtiMpGrQzM6AfnJ2U7BI-jxxC0s8BEreZZebTRHEFbsG9z_M5EAypGr9sl5YkuhegQ0UOYjY5R25GJlwn9qO0Iakh3PHZoI_z22-_foNsrO2LIWP6QPvJPwMrHV_QWDSuf4RVxDhD3CP_887dvD6C6bpkWo6LzF1IAFYDJJwHfYTBOxZRF2R4WO6vudQ2fwOtxH6HFwXmEOQ2UHU93tI3ulJ7JSyO2-flzKmdZdlYATJRlxYHJZ8Dv0asuosEp9ei-EUyURVkz8QR6W0CeBBO8b8d04fltRvHHM_KfZ5S_nLH6xTYWfzzjJ20s14ydMgZY9XhOIp5IZFkGrHoGFSYIukccBuwiMCFS_czjCNnssYeZyUcIGOfMYtZKAfOJXHC59ipP9w_jNKlsmtS_VTjaLlv_KZMWTZ9N28JyUb_ybTZIkQ3pyrepXUwUKUdKdrqI00WeLgUTT2mGVM_vSFFfVJMUtYVflLgoeVFFamTKUz7lD6mP6w3B5POQUD-5JL54Sq1MhBSYshaXO-b8tT9-BRP1Ot7rI-H8Mv6nsrgry7tycX5Mtz9v67tyc1uu8rvyHba6w1Z32Kp4n2GplG-vS3zNraU4zb7_La7rNMCELD-cqbT6732BbOvx5Cxx7wK-u1pK-Nff__aPf8HBLaYHZQ7quNqux5CWTjK2H5wrgAApvrQ6Xluc7dfad2vPzeGc27rkbvCmg24NJoNO1Ue06FWy5jl-z-Av7oBv6NeeD_AVjHOvoCIc9irCXs0z2gDWHVKFrxAw9UeHj23x_e9qgPJmcm9Zm55E3qxLp_S8fJ_iqy4uC6X0vag5k4-fRJdX0dvb6Pzz6Ooqur6JrprPoxuKLvLb6PrT6IJfRYvb6OommokXJl6gxVFb0NZoi8nzrir8it3x_9_pSs-LqlmNLck6J8lJCpKSZEGyJLklWZGsSRKtIVpDtIZoDdEaojVEa4jWEK0hWnOhlXlOkpMUJCXJgmRJckuyIlmTJBonGicaJxonGicaJxonGicaJxonmiCaIJogmiCaIJogmiCaIJogmiCaJJokmiSaJJokmiSaJJokmiSaJFpBtIJod-uLTKMX62BcrZDzMlnVecO9rLOPzf7kwTqADqBg9q41OKUkLXZqCZg8ViV7V8YcoXcYLBNVBKNfEXRyWbRwdAtEt3T7K3OPyX4xnVjnZP8_HHUDtBgPiBbUeV0PmI7n75vC-htoG6Jfuqhd2tcB39Aml9eRiSrA4DwEN2HcrwfxANHrt3TeT3KvQwZfI-CkY-rbQXmbqgV1PB2TVVw7oI1JGPQRVN_rhFIGDipFRQeTekUIi8fLMAU1YGrMZeiCcYcAvTuc9qfD3hmE03k_A9pUN_1O9o1s1AZ3vCplUzdFUWz2u7rpt00pq6KralTD0Bd5ybHFNi8EVyg2eidyUeaV5LzhlSgyUQ5tPpTNIPOtFKpnRY6T0iYz5m3KnB83OoQFd2ml1fnGqBZNWF-ghGhV94q2Z_Jh3cbTUal83vhdCv3SLmNgRW50iIGSRR3N-v51iiifP3qhWqdMD-3x569MnZtabbGHg477NCXXh822-fXzDtlm8Wa3j3Fed-t1Lxl13C9t1rmJiZfUvvPly-zdf7CLTLysPQ9MvJw7_7YT_w0AAP__KYbjeA">