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

    <tr>
        <th>Summary</th>
        <td>
            CUDA code compile error for asm code: expected ')'
        </td>
    </tr>

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

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

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

<pre>
    ```
inline __device__ float4 hmma_fp32(const uint2& a, unsigned int b)
{
    float4 c;
    float  zero = 0.f;
 asm volatile("mma.sync.aligned.m16n8k8.row.col.f32.f16.f16.f32 \n"
 "    {%0, %1, %2, %3}, \n"
                 "    {%4, %5}, \n"
                 "    {%6}, \n"
                 " {%7, %7, %7, %7}; \n"

                 : "=f"(c.x), "=f"(c.y), "=f"(c.z), "=f"(c.w)
                 : "r"(a.x) "r"(a.y), "r"(b), "f"(zero));
    return c;
}
```
https://godbolt.org/z/jxd483a8j

This code can be built by nvcc: https://godbolt.org/z/34Gaoz9r6

The original code comes from https://github.com/NVIDIA/FasterTransformer/blob/afdf9a9eb86f15363c0249117d166d6b45dbb371/src/fastertransformer/kernels/decoder_masked_multihead_attention/decoder_masked_multihead_attention_template.hpp#L636

It seems that nvcc does not depend on the `,` for separation.
Below code also can be built:
```
inline __device__ float4 hmma_fp32(const uint2& a, unsigned int b)
{
    float4 c;
    float  zero = 0.f;
    asm volatile("mma.sync.aligned.m16n8k8.row.col.f32.f16.f16.f32 \n"
 "    {%0, %1, %2, %3}, \n"
                 "    {%4, %5}, \n"
                 "    {%6}, \n"
                 "    {%7, %7, %7, %7}; \n"

                 : "=f"(c.x) "=f"(c.y) "=f"(c.z) "=f"(c.w)
                 : "r"(a.x) "r"(a.y) "r"(b) "f"(zero));
    return c;
}
```
https://godbolt.org/z/eqqveTqTx
</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJzsVl1v6zYM_TXKC3ENi4o_8uCHtoGHAsOeur0GskzXamUplZR-5NcPcrLbusvd7oBt2MOCBBRI8RyaZA4sQ9D3lqhhxTUrtit5iKPzzXHUL6Qf-Kpz_VvDyvz8zbcsv9LWaEuw2_X0rBXtdjAYJ-MaxmmSu2EvkGGtnA0RDtpGZFiCZHgDBzuT9aBthI7h5oTHquvTAQB-g1JMfHYCHMk7YGILeTa8x2WY4NkZGbUhhjVDnCaZhTerMmlmvmzipa0f68y7l0w5kw0Cs4GXp59AYMWNZYhnQIaYSFNZWOSpcIYFP1s8W8Gq7XxcpH7-LKDW59Tir6eW351yul-dqX5vqy0T10ucb6GJuRVMbIdksFbZa5rZjLRwv112Hy-7X74O_luU_nRVznwLxwems69795wZ0pbMzs1ihzzFg7cfFiu14nRYrvcY4z4wccWwZdjeu75zJmbO3zNsjwzbh9d-XQtZP3zs3t2oAyjXEyhpoSPoDtpE6N7APiuVnuvPYMX6B-mOG18uYQmc1_faSnOGdxMFGLybPiPqOB66TLmJYfvTL7fb2xRoZYjk77y0YXB-Is-w7YzrGLZy6IeN3FBXlwMvRClUjusN51XPy7Ivu3XRd52oOMM2eMWwHWasuMB6JG_JBIZtT6k-v5tkeKR-Nx1M1CPJfidjJBu1s991aRdp2hsZKRv3e4bix1IsOnIbIRBNAeIo49xc6B0FsC5CT3uyPTgLcSRIA8UbVuYwOA-B9tLLxJCdgK7JuJdTT6UJbjG31NRLq_EfUz6A_8Xvcso_qH-XxO-S8v3dsvdJ8_5FwaOnp2e6e7p7XfWN6DdiI1fU8HJT5TznBV-NDcm8Lsqiy7Go-EaWqsJSKNXltay5rIaVbjBHwXPO82qNgmc5DWpQPK_rXBJXa7bOaZLaZMY8T4l7pUM4UFPWVSVWRnZkwvyOgmjpBeZgevpiu_JNyvnSHe4DW-dGhxjeUaKOhpqbn7dXX-Vzrw0Bee_8LAzpH5RCaQj0uicVqQeG1dzRanXwpvkDoU1EZ_Nl790DqciwnctLojiX_2sAAAD__yhgfak">