================
@@ -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
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
