llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-backend-amdgpu @llvm/pr-subscribers-clangir Author: Steffen Larsen (steffenlarsen) <details> <summary>Changes</summary> This patch implements support for the clang::atomic attribute for the AMDGPU target. This attribute adjusts the atomic options in effect for the statement it is attached to, which decides the AMDGPU metadata on the resulting atomicrmw. That also closes a gap the attribute exposed, as __hip_atomic_* read-modify-writes carried no AMDGPU metadata at all. The "no.X" metadata asserts the absence of a memory kind, so it is emitted when the corresponding option is off, and amdgpu.no.fine.grained.memory, amdgpu.no.remote.memory and amdgpu.ignore.denormal.mode now match classic CodeGen. Large parts of these changes correspond to CGAtomicOptionsRAII and AMDGPUTargetCodeGenInfo::setTargetAtomicMetadata from OGCG. Assisted-by: Claude Code Sonnet 5 --- Patch is 44.72 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/225364.diff 11 Files Affected: - (modified) clang/include/clang/CIR/MissingFeatures.h (+2) - (modified) clang/lib/CIR/CodeGen/CIRGenAtomic.cpp (+38-20) - (modified) clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp (+114-14) - (modified) clang/lib/CIR/CodeGen/CIRGenModule.cpp (+2-1) - (modified) clang/lib/CIR/CodeGen/CIRGenModule.h (+9) - (modified) clang/lib/CIR/CodeGen/CIRGenStmt.cpp (+46-2) - (modified) clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/AMDGPU.cpp (+5) - (modified) clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp (+69-13) - (modified) clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp (+29) - (added) clang/test/CIR/CodeGenHIP/atomic-options.hip (+70) - (added) clang/test/CIR/CodeGenHIP/builtins-amdgcn-raw-atomic.hip (+183) ``````````diff diff --git a/clang/include/clang/CIR/MissingFeatures.h b/clang/include/clang/CIR/MissingFeatures.h index 16d97871ec7b6..9517ed098d9ba 100644 --- a/clang/include/clang/CIR/MissingFeatures.h +++ b/clang/include/clang/CIR/MissingFeatures.h @@ -151,6 +151,8 @@ struct MissingFeatures { static bool atomicUseLibCall() { return false; } static bool atomicMicrosoftVolatile() { return false; } static bool atomicOpenMP() { return false; } + static bool atomicAMDGPUNoaliasAddrspace() { return false; } + static bool atomicAMDGPUAvailableVisibleMMRA() { return false; } // Global ctor handling static bool globalCtorLexOrder() { return false; } diff --git a/clang/lib/CIR/CodeGen/CIRGenAtomic.cpp b/clang/lib/CIR/CodeGen/CIRGenAtomic.cpp index c4c3b455bf11c..de5be38090cb4 100644 --- a/clang/lib/CIR/CodeGen/CIRGenAtomic.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenAtomic.cpp @@ -646,6 +646,41 @@ static void emitAtomicCmpXchgFailureSetCheckWeak( }); } +/// Attach the AMDGPU atomic metadata markers that the current atomic options +/// call for. The "no.X" metadata is emitted when the corresponding option is +/// off, since it asserts the absence of that memory kind. The clang::atomic +/// attribute is what turns the options on and off. +static void setAMDGPUAtomicMetadata(CIRGenFunction &cgf, mlir::Operation *op) { + // TODO: AMDGPUTargetCodeGenInfo::setTargetAtomicMetadata also emits + // !noalias.addrspace on a flat-pointer atomic when the source atomic + // expression's memory is thread-private-undefined (OpenCL / old-style HIP + // atomics), regardless of whether it is a read-modify-write or cmpxchg. + assert(!cir::MissingFeatures::atomicAMDGPUNoaliasAddrspace()); + // TODO: AMDGPUTargetCodeGenInfo::setTargetAtomicMetadata also calls + // CGF.AddAMDGPUAvailableVisibleMMRA on every atomic instruction; this is + // tied to the AMDGPUAvailableVisible statement attribute, which + // CIRGenStmt.cpp's emitAttributedStmt currently rejects via errorNYI. + assert(!cir::MissingFeatures::atomicAMDGPUAvailableVisibleMMRA()); + + // Only a read-modify-write instruction carries these; a plain load, store or + // cmpxchg does not. + auto fetchOp = mlir::dyn_cast<cir::AtomicFetchOp>(op); + if (!fetchOp) + return; + + clang::AtomicOptions atomicOpts = cgf.cgm.getAtomicOpts(); + mlir::UnitAttr unit = cgf.getBuilder().getUnitAttr(); + if (!atomicOpts.getOption(clang::AtomicOptionKind::FineGrainedMemory)) + op->setAttr("cir.amdgpu_no_fine_grained_memory", unit); + if (!atomicOpts.getOption(clang::AtomicOptionKind::RemoteMemory)) + op->setAttr("cir.amdgpu_no_remote_memory", unit); + // Denormal flushing only matters for a float add. + if (atomicOpts.getOption(clang::AtomicOptionKind::IgnoreDenormalMode) && + fetchOp.getBinop() == cir::AtomicFetchKind::Add && + mlir::isa<cir::SingleType>(fetchOp.getVal().getType())) + op->setAttr("cir.amdgpu_ignore_denormal_mode", unit); +} + static void emitAtomicOp(CIRGenFunction &cgf, AtomicExpr *expr, Address dest, Address ptr, Address val1, Address val2, Expr *isWeakExpr, Expr *failureOrderExpr, int64_t size, @@ -914,6 +949,9 @@ static void emitAtomicOp(CIRGenFunction &cgf, AtomicExpr *expr, Address dest, if (fetchFirst && opName == cir::AtomicFetchOp::getOperationName()) rmwOp->setAttr("fetch_first", builder.getUnitAttr()); + if (cgf.cgm.getTriple().isAMDGCN()) + setAMDGPUAtomicMetadata(cgf, rmwOp); + mlir::Value result = rmwOp->getResult(0); builder.createStore(loc, result, dest); @@ -1209,10 +1247,6 @@ static RValue emitLibCallForAtomicExpr(CIRGenFunction &cgf, AtomicExpr *e, case AtomicExpr::AO__atomic_compare_exchange_n: case AtomicExpr::AO__c11_atomic_compare_exchange_weak: case AtomicExpr::AO__c11_atomic_compare_exchange_strong: - case AtomicExpr::AO__hip_atomic_compare_exchange_weak: - case AtomicExpr::AO__hip_atomic_compare_exchange_strong: - case AtomicExpr::AO__opencl_atomic_compare_exchange_weak: - case AtomicExpr::AO__opencl_atomic_compare_exchange_strong: case AtomicExpr::AO__scoped_atomic_compare_exchange: case AtomicExpr::AO__scoped_atomic_compare_exchange_n: { calleeName = "__atomic_compare_exchange"; @@ -1235,8 +1269,6 @@ static RValue emitLibCallForAtomicExpr(CIRGenFunction &cgf, AtomicExpr *e, case AtomicExpr::AO__atomic_exchange: case AtomicExpr::AO__atomic_exchange_n: case AtomicExpr::AO__c11_atomic_exchange: - case AtomicExpr::AO__hip_atomic_exchange: - case AtomicExpr::AO__opencl_atomic_exchange: case AtomicExpr::AO__scoped_atomic_exchange: case AtomicExpr::AO__scoped_atomic_exchange_n: calleeName = "__atomic_exchange"; @@ -1276,36 +1308,26 @@ static RValue emitLibCallForAtomicExpr(CIRGenFunction &cgf, AtomicExpr *e, case AtomicExpr::AO__scoped_atomic_add_fetch: case AtomicExpr::AO__atomic_fetch_add: case AtomicExpr::AO__c11_atomic_fetch_add: - case AtomicExpr::AO__hip_atomic_fetch_add: - case AtomicExpr::AO__opencl_atomic_fetch_add: case AtomicExpr::AO__scoped_atomic_fetch_add: case AtomicExpr::AO__atomic_and_fetch: case AtomicExpr::AO__scoped_atomic_and_fetch: case AtomicExpr::AO__atomic_fetch_and: case AtomicExpr::AO__c11_atomic_fetch_and: - case AtomicExpr::AO__hip_atomic_fetch_and: - case AtomicExpr::AO__opencl_atomic_fetch_and: case AtomicExpr::AO__scoped_atomic_fetch_and: case AtomicExpr::AO__atomic_or_fetch: case AtomicExpr::AO__scoped_atomic_or_fetch: case AtomicExpr::AO__atomic_fetch_or: case AtomicExpr::AO__c11_atomic_fetch_or: - case AtomicExpr::AO__hip_atomic_fetch_or: - case AtomicExpr::AO__opencl_atomic_fetch_or: case AtomicExpr::AO__scoped_atomic_fetch_or: case AtomicExpr::AO__atomic_sub_fetch: case AtomicExpr::AO__scoped_atomic_sub_fetch: case AtomicExpr::AO__atomic_fetch_sub: case AtomicExpr::AO__c11_atomic_fetch_sub: - case AtomicExpr::AO__hip_atomic_fetch_sub: - case AtomicExpr::AO__opencl_atomic_fetch_sub: case AtomicExpr::AO__scoped_atomic_fetch_sub: case AtomicExpr::AO__atomic_xor_fetch: case AtomicExpr::AO__scoped_atomic_xor_fetch: case AtomicExpr::AO__atomic_fetch_xor: case AtomicExpr::AO__c11_atomic_fetch_xor: - case AtomicExpr::AO__hip_atomic_fetch_xor: - case AtomicExpr::AO__opencl_atomic_fetch_xor: case AtomicExpr::AO__scoped_atomic_fetch_xor: case AtomicExpr::AO__atomic_nand_fetch: case AtomicExpr::AO__atomic_fetch_nand: @@ -1315,15 +1337,11 @@ static RValue emitLibCallForAtomicExpr(CIRGenFunction &cgf, AtomicExpr *e, case AtomicExpr::AO__atomic_min_fetch: case AtomicExpr::AO__atomic_fetch_min: case AtomicExpr::AO__c11_atomic_fetch_min: - case AtomicExpr::AO__hip_atomic_fetch_min: - case AtomicExpr::AO__opencl_atomic_fetch_min: case AtomicExpr::AO__scoped_atomic_fetch_min: case AtomicExpr::AO__scoped_atomic_min_fetch: case AtomicExpr::AO__atomic_max_fetch: case AtomicExpr::AO__atomic_fetch_max: case AtomicExpr::AO__c11_atomic_fetch_max: - case AtomicExpr::AO__hip_atomic_fetch_max: - case AtomicExpr::AO__opencl_atomic_fetch_max: case AtomicExpr::AO__scoped_atomic_fetch_max: case AtomicExpr::AO__scoped_atomic_max_fetch: case AtomicExpr::AO__scoped_atomic_fetch_uinc: diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index 8d23951dd64ba..2e6aa47b59535 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -21,6 +21,84 @@ using namespace clang; using namespace clang::CIRGen; using namespace cir; +/// Map a constant integeral to memory order. +static cir::MemOrder decodeAtomicOrder(const Expr *arg, ASTContext &ctx) { + Expr::EvalResult orderRes; + if (!arg->EvaluateAsInt(orderRes, ctx)) + return cir::MemOrder::SequentiallyConsistent; + switch (orderRes.Val.getInt().getZExtValue()) { + case 0: + return cir::MemOrder::Relaxed; + case 1: // consume -> acquire + case 2: + return cir::MemOrder::Acquire; + case 3: + return cir::MemOrder::Release; + case 4: + return cir::MemOrder::AcquireRelease; + default: + return cir::MemOrder::SequentiallyConsistent; + } +} + +/// Map an AMDGPU sync-scope string-literal argument onto cir::SyncScopeKind. +/// An absent or unrecognized scope is system scope, which is the conservative +/// choice and matches what an empty syncscope string means in LLVM. +static cir::SyncScopeKind decodeAMDGPUSyncScope(const Expr *arg) { + const auto *sl = + llvm::dyn_cast<clang::StringLiteral>(arg->IgnoreParenCasts()); + if (!sl) + return cir::SyncScopeKind::System; + return llvm::StringSwitch<cir::SyncScopeKind>(sl->getString()) + .Case("singlethread", cir::SyncScopeKind::SingleThread) + .Case("wavefront", cir::SyncScopeKind::Wavefront) + .Case("workgroup", cir::SyncScopeKind::Workgroup) + .Case("agent", cir::SyncScopeKind::Device) + .Default(cir::SyncScopeKind::System); +} + +/// Emit one of the AMDGPU raw hardware atomic builtins as a cir.atomic.fetch. +static mlir::Value emitAMDGPUAtomicRMW(CIRGenFunction &cgf, + const CallExpr *expr, + cir::AtomicFetchKind binOp, + bool hasVolatileArg) { + CIRGenBuilderTy &builder = cgf.getBuilder(); + mlir::Location loc = cgf.getLoc(expr->getExprLoc()); + + Address ptr = cgf.emitPointerWithAlignment(expr->getArg(0)); + mlir::Value val = cgf.emitScalarExpr(expr->getArg(1)); + + bool isVolatile; + if (hasVolatileArg) { + assert(expr->getNumArgs() >= 5); + // ds_faddf/fminf/fmaxf spell the volatile flag out as a constant argument. + Expr::EvalResult volRes; + isVolatile = expr->getArg(4)->EvaluateAsInt(volRes, cgf.getContext()) && + volRes.Val.getInt().getBoolValue(); + } else { + // Everything else infers it from the pointee type. + QualType argTy = expr->getArg(0)->IgnoreImpCasts()->getType(); + isVolatile = argTy->castAs<clang::PointerType>() + ->getPointeeType() + .isVolatileQualified(); + } + + // Some of these builtins spell out the ordering and scope; the rest take the + // monotonic/agent default described above. + cir::MemOrder order = cir::MemOrder::Relaxed; + cir::SyncScopeKind scope = cir::SyncScopeKind::Device; + if (expr->getNumArgs() >= 4) { + order = decodeAtomicOrder(expr->getArg(2), cgf.getContext()); + scope = decodeAMDGPUSyncScope(expr->getArg(3)); + } + + auto rmw = cir::AtomicFetchOp::create(builder, loc, ptr.emitRawPointer(), val, + binOp, order, scope, isVolatile, + /*fetch_first=*/true); + rmw->setAttr("cir.amdgpu_raw_atomic", builder.getUnitAttr()); + return rmw->getResult(0); +} + // Emit the `amdgcn.dispatch.ptr` intrinsic, address-space-casting the // result to match \p e's return type when needed. // If \p e is null, returns the raw AS-4 pointer. @@ -965,34 +1043,56 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, return mlir::Value{}; } case AMDGPU::BI__builtin_amdgcn_fence: { - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented AMDGPU builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); + 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); return mlir::Value{}; } case AMDGPU::BI__builtin_amdgcn_atomic_inc32: case AMDGPU::BI__builtin_amdgcn_atomic_inc64: + return emitAMDGPUAtomicRMW(*this, expr, cir::AtomicFetchKind::UIncWrap, + /*hasVolatileArg=*/false); case AMDGPU::BI__builtin_amdgcn_atomic_dec32: case AMDGPU::BI__builtin_amdgcn_atomic_dec64: + return emitAMDGPUAtomicRMW(*this, expr, cir::AtomicFetchKind::UDecWrap, + /*hasVolatileArg=*/false); case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_f64: case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_f32: - case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_v2f16: - case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_v2bf16: - case AMDGPU::BI__builtin_amdgcn_ds_faddf: - case AMDGPU::BI__builtin_amdgcn_ds_fminf: - case AMDGPU::BI__builtin_amdgcn_ds_fmaxf: case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_f32: case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_f64: - case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2f16: - case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2f16: case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_f32: case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_f64: - case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2bf16: - case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2bf16: + return emitAMDGPUAtomicRMW(*this, expr, cir::AtomicFetchKind::Add, + /*hasVolatileArg=*/false); case AMDGPU::BI__builtin_amdgcn_global_atomic_fmin_f64: - case AMDGPU::BI__builtin_amdgcn_global_atomic_fmax_f64: case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmin_f64: - case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmax_f64: { + return emitAMDGPUAtomicRMW(*this, expr, cir::AtomicFetchKind::Min, + /*hasVolatileArg=*/false); + case AMDGPU::BI__builtin_amdgcn_global_atomic_fmax_f64: + case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmax_f64: + return emitAMDGPUAtomicRMW(*this, expr, cir::AtomicFetchKind::Max, + /*hasVolatileArg=*/false); + // The ds_ float forms are the same operations with the volatile flag spelled + // out as an argument. + case AMDGPU::BI__builtin_amdgcn_ds_faddf: + return emitAMDGPUAtomicRMW(*this, expr, cir::AtomicFetchKind::Add, + /*hasVolatileArg=*/true); + case AMDGPU::BI__builtin_amdgcn_ds_fminf: + return emitAMDGPUAtomicRMW(*this, expr, cir::AtomicFetchKind::Min, + /*hasVolatileArg=*/true); + case AMDGPU::BI__builtin_amdgcn_ds_fmaxf: + return emitAMDGPUAtomicRMW(*this, expr, cir::AtomicFetchKind::Max, + /*hasVolatileArg=*/true); + case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_v2f16: + case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_v2bf16: + case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2f16: + case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2f16: + case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2bf16: + case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2bf16: { cgm.errorNYI(expr->getSourceRange(), std::string("unimplemented AMDGPU builtin call: ") + getContext().BuiltinInfo.getName(builtinId)); diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp index 3af6ce4ce6e94..47c5b1e5af457 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp @@ -85,7 +85,8 @@ CIRGenModule::CIRGenModule(mlir::MLIRContext &mlirContext, const clang::CodeGenOptions &cgo, DiagnosticsEngine &diags) : builder(mlirContext, *this), astContext(astContext), - langOpts(astContext.getLangOpts()), codeGenOpts(cgo), + langOpts(astContext.getLangOpts()), atomicOpts(astContext.getLangOpts()), + codeGenOpts(cgo), theModule{mlir::ModuleOp::create(mlir::UnknownLoc::get(&mlirContext))}, diags(diags), target(astContext.getTargetInfo()), abi(createCXXABI(*this)), genTypes(*this), vtables(*this) { diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.h b/clang/lib/CIR/CodeGen/CIRGenModule.h index 51b9c420c94be..93e007381d292 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.h +++ b/clang/lib/CIR/CodeGen/CIRGenModule.h @@ -82,6 +82,9 @@ class CIRGenModule : public CIRGenTypeCache { const clang::LangOptions &langOpts; + /// Seeded from langOpts and then pushed and popped by clang::atomic. + clang::AtomicOptions atomicOpts; + const clang::CodeGenOptions &codeGenOpts; /// A "module" matches a c/cpp source file: containing a list of functions. @@ -176,6 +179,12 @@ class CIRGenModule : public CIRGenTypeCache { CIRGenTypes &getTypes() { return genTypes; } const clang::LangOptions &getLangOpts() const { return langOpts; } + /// The atomic options in effect at the point currently being emitted. + /// clang::atomic adjusts these for the extent of the statement it is attached + /// to, so they are module state rather than a fixed language option. + clang::AtomicOptions getAtomicOpts() const { return atomicOpts; } + void setAtomicOpts(clang::AtomicOptions opts) { atomicOpts = opts; } + CIRGenCXXABI &getCXXABI() const { return *abi; } mlir::MLIRContext &getMLIRContext() { return *builder.getContext(); } diff --git a/clang/lib/CIR/CodeGen/CIRGenStmt.cpp b/clang/lib/CIR/CodeGen/CIRGenStmt.cpp index e9f5e466c63d2..864998f231619 100644 --- a/clang/lib/CIR/CodeGen/CIRGenStmt.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenStmt.cpp @@ -28,6 +28,38 @@ using namespace clang; using namespace clang::CIRGen; using namespace cir; +/// Compute the atomic options that apply for the extent of a clang::atomic +/// attributed statement, starting from the enclosing options adjustments on top +/// of them. +static clang::AtomicOptions getAdjustedAtomicOptions(clang::AtomicOptions ao, + const AtomicAttr *aa) { + if (!aa) + return ao; + for (auto option : aa->atomicOptions()) { + switch (option) { + case AtomicAttr::remote_memory: + ao.remote_memory = true; + break; + case AtomicAttr::no_remote_memory: + ao.remote_memory = false; + break; + case AtomicAttr::fine_grained_memory: + ao.fine_grained_memory = true; + break; + case AtomicAttr::no_fine_grained_memory: + ao.fine_grained_memory = false; + break; + case AtomicAttr::ignore_denormal_mode: + ao.ignore_denormal_mode = true; + break; + case AtomicAttr::no_ignore_denormal_mode: + ao.ignore_denormal_mode = false; + break; + } + } + return ao; +} + static mlir::LogicalResult emitStmtWithResult(CIRGenFunction &cgf, const Stmt *exprResult, AggValueSlot slot, @@ -93,14 +125,17 @@ CIRGenFunction::emitAttributedStmt(const AttributedStmt &s) { bool noinline = inNoInlineAttributedStmt; bool alwaysinline = inAlwaysInlineAttributedStmt; const CallExpr *musttail = mustTailCall; + const AtomicAttr *atomicAttr = nullptr; for (const Attr *attr : s.getAttrs()) { switch (attr->getKind()) { default: break; + case attr::Atomic: + atomicAttr = cast<AtomicAttr>(attr); + break; case attr::NoMerge: case attr::NoConvergent: - case attr::Atomic: case attr::AMDGPUAvailableVisible: case attr::HLSLControlFlowHint: cgm.errorNYI(s.getSourceRange(), @@ -141,7 +176,16 @@ CIRGenFunction::emitAttributedStmt(const AttributedStmt &s) { SaveAndRestore save_musttail(mustTailCall, musttail); - return emitStmt(s.getSubStmt(), /*useCurrentScope=*/true, s.getAttrs()); + // clang::atomic adjusts the atomic options for the extent of the statement it + // is attached to, so they are saved and restored around it. + clang::AtomicOptions savedAtomicOpts = cgm.getAtomicOpts(); + cgm.setAtomicOpts(getAdjustedAtomicOptions(savedAtomicOpts, atomicAttr)); + + mlir::LogicalResult result = + emitStmt(s.getSubStmt(), /*useCurrentScope=*/true, s.getAttrs()); + + cgm.setAtomicOpts(savedAtomicOpts); + return result; } mlir::LogicalResult CIRGenFunction::emitCompoundStmt(const CompoundStmt &s, diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/AMDGPU.cpp b/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/AMDGPU.cpp index f4cdc88d60cf4..9282b91d048c5 100644 --- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/AMDGPU.cpp +++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/AMDGPU.cpp @@ -38,6 +38,11 @@ class AMDGPUTargetLoweringIn... [truncated] `````````` </details> https://github.com/llvm/llvm-project/pull/225364 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
