<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">