[clang] [CIR][AMDGPU][SPIR-V] Complete __builtin_amdgcn_fence lowering implementation (PR #230829)

[Steffen Larsen via cfe-commits]

================
@@ -1118,13 +1121,68 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
     return mlir::Value{};
   }
   case AMDGPU::BI__builtin_amdgcn_fence: {
-    CIRGenBuilderTy &b = getBuilder();
-    cir::MemOrder mo = decodeAtomicOrder(expr->getArg(0), getContext());
-    cir::SyncScopeKind syncScope = decodeAMDGPUSyncScope(expr->getArg(1));
-    cir::SyncScopeKindAttr syncScopeAttr =
-        cir::SyncScopeKindAttr::get(b.getContext(), syncScope);
-    cir::AtomicFenceOp::create(b, getLoc(expr->getExprLoc()), mo,
-                               syncScopeAttr);
+    // The scope uses AMDGPU spelling on every target, like mapScopeToSPIRV.
+    const llvm::Triple &triple = getTarget().getTriple();
+    std::optional<std::pair<llvm::AtomicScope, bool>> parsed;
+    if (std::optional<std::string> name =
+            expr->getArg(1)->tryEvaluateString(getContext())) {
+      parsed = llvm::parseAtomicScopeIRString(llvm::Triple("amdgcn-amd-amdhsa"),
----------------
steffenlarsen wrote:

Why do we generate the generic AMDGCN triple here? Could we use the actual triple or would SPIR-V mess it up? If the latter, could we amend it at the tooling level?

https://github.com/llvm/llvm-project/pull/230829