================
@@ -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"),
+ *name);
+ if (!parsed && triple.isSPIRV())
+ parsed = llvm::parseAtomicScopeIRString(triple, *name);
+ }
+ // SPIR-V drops one-as, AMDGPU needs a CIR sync scope for it.
+ if (!parsed || (parsed->second && !triple.isSPIRV())) {
+ cgm.errorNYI(expr->getSourceRange(), "__builtin_amdgcn_fence scope");
+ return mlir::Value{};
+ }
+ cir::SyncScopeKind syncScope;
+ switch (parsed->first) {
+ case llvm::AtomicScope::Single:
+ syncScope = cir::SyncScopeKind::SingleThread;
+ break;
+ case llvm::AtomicScope::Wavefront:
+ syncScope = cir::SyncScopeKind::Wavefront;
+ break;
+ case llvm::AtomicScope::Workgroup:
+ syncScope = cir::SyncScopeKind::Workgroup;
+ break;
+ case llvm::AtomicScope::Cluster:
+ syncScope = cir::SyncScopeKind::Cluster;
+ break;
+ case llvm::AtomicScope::Device:
+ syncScope = cir::SyncScopeKind::Device;
+ break;
+ case llvm::AtomicScope::System:
+ syncScope = cir::SyncScopeKind::System;
+ break;
+ }
+
+ // Sorted and unique, like MMRAMetadata::appendTags.
+ llvm::SmallVector<std::string> addrSpaces;
+ for (unsigned i = 2, e = expr->getNumArgs(); i < e; ++i) {
+ if (std::optional<std::string> as =
+ expr->getArg(i)->tryEvaluateString(getContext()))
+ addrSpaces.push_back(std::move(*as));
+ else
+ cgm.error(expr->getExprLoc(),
+ "expected an address space name as a string literal");
+ }
+ llvm::sort(addrSpaces);
+ addrSpaces.erase(llvm::unique(addrSpaces), addrSpaces.end());
+ llvm::SmallVector<mlir::Attribute> mmras;
----------------
steffenlarsen wrote:
```suggestion
llvm::SmallVector<mlir::Attribute> mmras;
mmras.reserve(addrSpaces.size());
```
https://github.com/llvm/llvm-project/pull/230829
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
