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

    <tr>
        <th>Summary</th>
        <td>
            [AMDGPU] CGBuiltin - bad handling of builtin_amdgcn_fence
        </td>
    </tr>

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

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

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

<pre>
    As mentioned here: https://reviews.llvm.org/D80804#inline-1081007

The LLVM_FALLTHROUGH handling of builtin_amdgcn_fence is missing:
```
  case AMDGPU::BI__builtin_amdgcn_fence: {
    if (ProcessOrderScopeAMDGCN(EmitScalarExpr(E->getArg(0)),
                                EmitScalarExpr(E->getArg(1)), AO, SSID))
      return Builder.CreateFence(AO, SSID);
    LLVM_FALLTHROUGH;
  }
  case AMDGPU::BI__builtin_amdgcn_atomic_inc32:
  case AMDGPU::BI__builtin_amdgcn_atomic_inc64:
  case AMDGPU::BI__builtin_amdgcn_atomic_dec32:
  case AMDGPU::BI__builtin_amdgcn_atomic_dec64: {
    unsigned BuiltinAtomicOp;
    llvm::Type *ResultType = ConvertType(E->getType());

    switch (BuiltinID) {
    case AMDGPU::BI__builtin_amdgcn_atomic_inc32:
    case AMDGPU::BI__builtin_amdgcn_atomic_inc64:
      BuiltinAtomicOp = Intrinsic::amdgcn_atomic_inc;
      break;
    case AMDGPU::BI__builtin_amdgcn_atomic_dec32:
    case AMDGPU::BI__builtin_amdgcn_atomic_dec64:
      BuiltinAtomicOp = Intrinsic::amdgcn_atomic_dec;
      break;
    }
```
</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJylVd9zoyAQ_mv0hUkGIUbz4INJ-msuvdw07b1mUFflqpgBbK___SHaNPZurtOGYZQF9tv9dhdImuwlihWqQWjeCMhQCRIcGqNS64MyA4dcmi7hicOzmlbVUz1tZGGm1iEO8cwhlIuKC5h4OPQwDhy8dnDcf-9LQJvNz9v9ZbzZ3F_fbR-urlHJRGYUCtTkKGl5pbnYszorUrHPQaSAuHGHK2W2dOZ7uDkeuhURSpkCFN-ur348dJtovLzZ7_-F1lFxguWrHkI8Rw4Jf8gmBaW2MgO5S5sDdFir72blouZ6l7KKyYvfB9lNTBx6UYCOO9YhdsjC9tUb5P_aB3DeEQ7F2-67292sh7kTfAm6lQItDUHj8HQlgWm4tPxIOFakJ1zfx_5k0QnWn4kl003N0z0XKSXHrHxWdT77mmoGX7ZqVK3VcQ20QvGiK_ZlrxTbzdvDKHhdqffw9y8HMEUT34FqK91LdI1WjXgCaeWTvA7ikMNXwDdY9cx1WnY1OBi3aRv7d0ZGzspJ196FxDK9EVpyE7O0B_wLZBQ3hBJTno-juTOS_aV0n0XIgHxI6Hh8jjeTC5E39xckpNgL3Syi2YIumKu5riBy_OVAwDd1czW4hCYoYdmHF6LbyioaX8cF12WbTNOmNoIt1P43OcjmF6TaiOYGbUGZgT8PKXHLKM-z2SLAHgshnOeMZF5GaOpRnOCQBSx3K5ZApTpfHUISlj6CyIzFwXFCjO8ujwgmBAfU83x_7gdTPPNYkAQ59ryQBECdGYaa8er4Urgysp4lbaHMYsWVfntGXKbsSbQR6vBZq8tGRnffdrxuhGtJRJbBH0OK_hA">