https://github.com/steffenlarsen created https://github.com/llvm/llvm-project/pull/225364
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 >From 6c48efa1c17224559c912ee6650db84d2b99926d Mon Sep 17 00:00:00 2001 From: Steffen Holst Larsen <[email protected]> Date: Tue, 1 Sep 2026 10:01:40 -0500 Subject: [PATCH] [CIR][AMDGPU] Implement the clang::atomic statement attribute 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 --- clang/include/clang/CIR/MissingFeatures.h | 2 + clang/lib/CIR/CodeGen/CIRGenAtomic.cpp | 58 ++++-- clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 128 ++++++++++-- clang/lib/CIR/CodeGen/CIRGenModule.cpp | 3 +- clang/lib/CIR/CodeGen/CIRGenModule.h | 9 + clang/lib/CIR/CodeGen/CIRGenStmt.cpp | 48 ++++- .../TargetLowering/Targets/AMDGPU.cpp | 5 + .../CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp | 82 ++++++-- .../Lowering/DirectToLLVM/LowerToLLVMIR.cpp | 29 +++ clang/test/CIR/CodeGenHIP/atomic-options.hip | 70 +++++++ .../CodeGenHIP/builtins-amdgcn-raw-atomic.hip | 183 ++++++++++++++++++ 11 files changed, 567 insertions(+), 50 deletions(-) create mode 100644 clang/test/CIR/CodeGenHIP/atomic-options.hip create mode 100644 clang/test/CIR/CodeGenHIP/builtins-amdgcn-raw-atomic.hip 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 AMDGPUTargetLoweringInfo : public TargetLoweringInfo { "Unknown CIR address space for AMDGPU target"); return AMDGPUAddrSpaceMap[idx]; } + + cir::SyncScopeKind + convertSyncScope(cir::SyncScopeKind syncScope) const override { + return syncScope; + } }; } // namespace diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp index 596c0d72264a0..e6dd447c9e9d0 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp @@ -50,6 +50,7 @@ #include "llvm/ADT/StringMap.h" #include "llvm/ADT/TypeSwitch.h" #include "llvm/IR/Module.h" +#include "llvm/Support/AMDGPUAddrSpace.h" #include "llvm/Support/Casting.h" #include "llvm/Support/ErrorHandling.h" #include "llvm/Support/TimeProfiler.h" @@ -1215,21 +1216,52 @@ getLLVMMemOrder(std::optional<cir::MemOrder> memorder) { llvm_unreachable("unknown memory order"); } -static llvm::StringRef getLLVMSyncScope(cir::SyncScopeKind syncScope) { +static bool isNVPTXTriple(mlir::Operation *op) { + auto moduleOp = op->getParentOfType<mlir::ModuleOp>(); + if (!moduleOp) + return false; + if (auto tripleAttr = moduleOp->getAttrOfType<mlir::StringAttr>( + cir::CIRDialect::getTripleAttrName())) + return llvm::Triple(tripleAttr.getValue()).isNVPTX(); + return false; +} + +static llvm::StringRef getLLVMSyncScope(cir::SyncScopeKind syncScope, + mlir::Operation *op) { switch (syncScope) { case cir::SyncScopeKind::SingleThread: + case cir::SyncScopeKind::HIPSingleThread: return "singlethread"; - case cir::SyncScopeKind::Workgroup: - return "block"; - default: + case cir::SyncScopeKind::System: + case cir::SyncScopeKind::HIPSystem: + case cir::SyncScopeKind::OpenCLAllSVMDevices: return ""; - } + case cir::SyncScopeKind::Device: + case cir::SyncScopeKind::HIPAgent: + case cir::SyncScopeKind::OpenCLDevice: + return "agent"; + case cir::SyncScopeKind::Workgroup: + case cir::SyncScopeKind::HIPWorkgroup: + case cir::SyncScopeKind::OpenCLWorkGroup: + // NVPTX uses "block" for workgroup sync scope; AMDGPU uses "workgroup". + return isNVPTXTriple(op) ? "block" : "workgroup"; + case cir::SyncScopeKind::Wavefront: + case cir::SyncScopeKind::HIPWavefront: + return "wavefront"; + case cir::SyncScopeKind::Cluster: + case cir::SyncScopeKind::HIPCluster: + return "cluster"; + case cir::SyncScopeKind::OpenCLSubGroup: + return "sub_group"; + } + llvm_unreachable("unknown sync scope"); } static std::optional<llvm::StringRef> -getLLVMSyncScope(std::optional<cir::SyncScopeKind> syncScope) { +getLLVMSyncScope(std::optional<cir::SyncScopeKind> syncScope, + mlir::Operation *op) { if (syncScope.has_value()) - return getLLVMSyncScope(*syncScope); + return getLLVMSyncScope(*syncScope, op); return std::nullopt; } @@ -1243,7 +1275,7 @@ mlir::LogicalResult CIRToLLVMAtomicCmpXchgOpLowering::matchAndRewrite( rewriter, op.getLoc(), adaptor.getPtr(), expected, desired, getLLVMMemOrder(adaptor.getSuccOrder()), getLLVMMemOrder(adaptor.getFailOrder()), - getLLVMSyncScope(op.getSyncScope())); + getLLVMSyncScope(op.getSyncScope(), op)); cmpxchg.setAlignment(adaptor.getAlignment()); cmpxchg.setWeak(adaptor.getWeak()); @@ -1264,7 +1296,7 @@ mlir::LogicalResult CIRToLLVMAtomicXchgOpLowering::matchAndRewrite( mlir::ConversionPatternRewriter &rewriter) const { assert(!cir::MissingFeatures::atomicSyncScopeID()); mlir::LLVM::AtomicOrdering llvmOrder = getLLVMMemOrder(adaptor.getMemOrder()); - llvm::StringRef llvmSyncScope = getLLVMSyncScope(adaptor.getSyncScope()); + llvm::StringRef llvmSyncScope = getLLVMSyncScope(op.getSyncScope(), op); rewriter.replaceOpWithNewOp<mlir::LLVM::AtomicRMWOp>( op, mlir::LLVM::AtomicBinOp::xchg, adaptor.getPtr(), adaptor.getVal(), llvmOrder, llvmSyncScope, /*alignment=*/0, op.getIsVolatile()); @@ -1317,7 +1349,7 @@ mlir::LogicalResult CIRToLLVMAtomicFenceOpLowering::matchAndRewrite( mlir::LLVM::AtomicOrdering llvmOrder = getLLVMMemOrder(adaptor.getOrdering()); auto fence = mlir::LLVM::FenceOp::create(rewriter, op.getLoc(), llvmOrder); - fence.setSyncscope(getLLVMSyncScope(adaptor.getSyncscope())); + fence.setSyncscope(getLLVMSyncScope(op.getSyncscope(), op)); rewriter.replaceOp(op, fence); @@ -1460,13 +1492,37 @@ mlir::LogicalResult CIRToLLVMAtomicFetchOpLowering::matchAndRewrite( } mlir::LLVM::AtomicOrdering llvmOrder = getLLVMMemOrder(op.getMemOrder()); - llvm::StringRef llvmSyncScope = getLLVMSyncScope(op.getSyncScope()); + llvm::StringRef llvmSyncScope = getLLVMSyncScope(op.getSyncScope(), op); mlir::LLVM::AtomicBinOp llvmBinOp = getLLVMAtomicBinOp(op.getBinop(), isInt, isSignedInt); auto rmwVal = mlir::LLVM::AtomicRMWOp::create( rewriter, op.getLoc(), llvmBinOp, adaptor.getPtr(), adaptor.getVal(), llvmOrder, llvmSyncScope, /*alignment=*/0, op.getIsVolatile()); + // CIRGen decides the metadata for a C++/HIP atomic from the atomic options + // in effect, so those markers are simply carried across. + for (llvm::StringRef marker : + {"cir.amdgpu_no_fine_grained_memory", "cir.amdgpu_no_remote_memory", + "cir.amdgpu_ignore_denormal_mode"}) + if (mlir::Attribute a = op->getAttr(marker)) + rmwVal->setAttr(marker, a); + + // The AMDGPU raw hardware atomic builtins need metadata for the backend to + // select the native instruction. LDS atomics are always native, so the + // metadata is only needed for the global and flat address spaces. + if (op->hasAttr("cir.amdgpu_raw_atomic")) { + auto ptrTy = + mlir::cast<mlir::LLVM::LLVMPointerType>(adaptor.getPtr().getType()); + if (ptrTy.getAddressSpace() != llvm::AMDGPUAS::LOCAL_ADDRESS) { + mlir::UnitAttr unit = rewriter.getUnitAttr(); + rmwVal->setAttr("cir.amdgpu_no_fine_grained_memory", unit); + // Denormal flushing only matters for a float add. + if (llvmBinOp == mlir::LLVM::AtomicBinOp::fadd && + mlir::isa<cir::SingleType>(op.getVal().getType())) + rmwVal->setAttr("cir.amdgpu_ignore_denormal_mode", unit); + } + } + mlir::Value result = rmwVal.getResult(); if (!op.getFetchFirst()) { if (op.getBinop() == cir::AtomicFetchKind::Max || @@ -2367,7 +2423,7 @@ mlir::LogicalResult CIRToLLVMLoadOpLowering::matchAndRewrite( assert(!cir::MissingFeatures::lowerModeOptLevel()); std::optional<llvm::StringRef> llvmSyncScope = - getLLVMSyncScope(op.getSyncScope()); + getLLVMSyncScope(op.getSyncScope(), op); mlir::LLVM::LoadOp newLoad = mlir::LLVM::LoadOp::create( rewriter, op->getLoc(), llvmTy, adaptor.getAddr(), alignment, @@ -2430,7 +2486,7 @@ mlir::LogicalResult CIRToLLVMStoreOpLowering::matchAndRewrite( assert(!cir::MissingFeatures::opLoadStoreTbaa()); std::optional<llvm::StringRef> llvmSyncScope = - getLLVMSyncScope(op.getSyncScope()); + getLLVMSyncScope(op.getSyncScope(), op); mlir::LLVM::StoreOp storeOp = mlir::LLVM::StoreOp::create( rewriter, op->getLoc(), value, adaptor.getAddr(), alignment, diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp index dd95f09ee77ae..47abfe9faff43 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp @@ -65,11 +65,40 @@ class CIRDialectLLVMIRTranslationInterface if (mlir::failed(amendRISCVNontemporalDomain(op, instructions, attribute, moduleTranslation))) return mlir::failure(); + } else if (attribute.getName() == "cir.amdgpu_no_fine_grained_memory" || + attribute.getName() == "cir.amdgpu_no_remote_memory" || + attribute.getName() == "cir.amdgpu_ignore_denormal_mode") { + amendAMDGPUAtomicMetadata(instructions, attribute, moduleTranslation); } return mlir::success(); } private: + /// Attach the AMDGPU atomic metadata that lets the backend select a native + /// atomic instruction instead of expanding to a cmpxchg loop. + void amendAMDGPUAtomicMetadata( + llvm::ArrayRef<llvm::Instruction *> instructions, + mlir::NamedAttribute attribute, + mlir::LLVM::ModuleTranslation &moduleTranslation) const { + llvm::MDNode *empty = + llvm::MDNode::get(moduleTranslation.getLLVMContext(), {}); + // !atomic.ignore.denormal.mode is a fixed metadata kind so it has to be + // attached via its enum rather than by name. + if (attribute.getName() == "cir.amdgpu_ignore_denormal_mode") { + for (llvm::Instruction *inst : instructions) + inst->setMetadata(llvm::LLVMContext::MD_atomic_ignore_denormal_mode, + empty); + return; + } + llvm::StringRef mdName = + llvm::StringSwitch<llvm::StringRef>(attribute.getName().strref()) + .Case("cir.amdgpu_no_fine_grained_memory", + "amdgpu.no.fine.grained.memory") + .Case("cir.amdgpu_no_remote_memory", "amdgpu.no.remote.memory"); + for (llvm::Instruction *inst : instructions) + inst->setMetadata(mdName, empty); + } + mlir::LogicalResult amendRISCVNontemporalDomain( mlir::Operation *op, llvm::ArrayRef<llvm::Instruction *> instructions, mlir::NamedAttribute attribute, diff --git a/clang/test/CIR/CodeGenHIP/atomic-options.hip b/clang/test/CIR/CodeGenHIP/atomic-options.hip new file mode 100644 index 0000000000000..8c0509d85c8d6 --- /dev/null +++ b/clang/test/CIR/CodeGenHIP/atomic-options.hip @@ -0,0 +1,70 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx90a -x hip \ +// RUN: -std=c++11 -fclangir -fcuda-is-device -emit-cir %s -o - \ +// RUN: | FileCheck --check-prefix=CIR %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx90a -x hip \ +// RUN: -std=c++11 -fclangir -fcuda-is-device -emit-llvm %s -o - \ +// RUN: | FileCheck --check-prefix=LLVM %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx90a -x hip \ +// RUN: -std=c++11 -fcuda-is-device -emit-llvm %s -o - \ +// RUN: | FileCheck --check-prefix=LLVM %s + +#define __device__ __attribute__((device)) +#define __global__ __attribute__((global)) + +// CIR-LABEL: @_Z9plain_addPff +// CIR: cir.atomic.fetch add relaxed syncscope(hip_agent) fetch_first %{{.+}}, %{{.+}} : (!cir.ptr<!cir.float>, !cir.float) -> !cir.float {cir.amdgpu_no_fine_grained_memory, cir.amdgpu_no_remote_memory} +// LLVM-LABEL: @_Z9plain_addPff +// LLVM: atomicrmw fadd ptr %{{.+}}, float %{{.+}} syncscope("agent") monotonic, align 4{{.*}}!amdgpu.{{(no.remote.memory|no.fine.grained.memory)}} !{{[0-9]+}}{{.*}}!amdgpu.{{(no.remote.memory|no.fine.grained.memory)}} +__device__ float plain_add(float *p, float v) { + return __hip_atomic_fetch_add(p, v, __ATOMIC_RELAXED, 4); +} + +// CIR-LABEL: @_Z10unsafe_addPff +// CIR: cir.atomic.fetch add relaxed syncscope(hip_agent) fetch_first %{{.+}}, %{{.+}} : (!cir.ptr<!cir.float>, !cir.float) -> !cir.float {cir.amdgpu_no_remote_memory} +// LLVM-LABEL: @_Z10unsafe_addPff +// LLVM: atomicrmw fadd ptr %{{.+}}, float %{{.+}} syncscope("agent") monotonic +// LLVM-NOT: amdgpu.no.fine.grained.memory +__device__ float unsafe_add(float *p, float v) { + [[clang::atomic(fine_grained_memory)]] { + return __hip_atomic_fetch_add(p, v, __ATOMIC_RELAXED, 4); + } +} + +// CIR-LABEL: @_Z9no_fg_addPff +// CIR: cir.atomic.fetch add relaxed syncscope(hip_agent) fetch_first %{{.+}}, %{{.+}} : (!cir.ptr<!cir.float>, !cir.float) -> !cir.float {cir.amdgpu_ignore_denormal_mode, cir.amdgpu_no_fine_grained_memory, cir.amdgpu_no_remote_memory} +// LLVM-LABEL: @_Z9no_fg_addPff +// LLVM: atomicrmw fadd ptr %{{.+}}, float %{{.+}} syncscope("agent") monotonic, align 4{{.*}}!atomic.ignore.denormal.mode !{{[0-9]+}}, !amdgpu.no.fine.grained.memory !{{[0-9]+}}, !amdgpu.no.remote.memory +__device__ float no_fg_add(float *p, float v) { + [[clang::atomic(no_fine_grained_memory, ignore_denormal_mode)]] { + return __hip_atomic_fetch_add(p, v, __ATOMIC_RELAXED, 4); + } +} + +// CIR-LABEL: @_Z9no_fg_addPdd +// CIR: cir.atomic.fetch add relaxed syncscope(hip_agent) fetch_first %{{.+}}, %{{.+}} : (!cir.ptr<!cir.double>, !cir.double) -> !cir.double {cir.amdgpu_no_fine_grained_memory, cir.amdgpu_no_remote_memory} +// LLVM-LABEL: @_Z9no_fg_addPdd +// LLVM: atomicrmw fadd ptr %{{.+}}, double %{{.+}} syncscope("agent") monotonic +// LLVM-NOT: atomic.ignore.denormal.mode +__device__ double no_fg_add(double *p, double v) { + [[clang::atomic(no_fine_grained_memory, ignore_denormal_mode)]] { + return __hip_atomic_fetch_add(p, v, __ATOMIC_RELAXED, 4); + } +} + +// CIR-LABEL: @_Z12workgroup_ddPdd +// CIR: cir.atomic.fetch add relaxed syncscope(hip_workgroup) fetch_first %{{.+}}, %{{.+}} : (!cir.ptr<!cir.double>, !cir.double) -> !cir.double {cir.amdgpu_no_fine_grained_memory, cir.amdgpu_no_remote_memory} +// LLVM-LABEL: @_Z12workgroup_ddPdd +// LLVM: atomicrmw fadd ptr %{{.+}}, double %{{.+}} syncscope("workgroup") monotonic +__device__ double workgroup_dd(double *p, double v) { + return __hip_atomic_fetch_add(p, v, __ATOMIC_RELAXED, 3); +} + +// CIR-LABEL: @_Z12after_scopedPff +// CIR: cir.atomic.fetch add relaxed syncscope(hip_agent) fetch_first %{{.+}}, %{{.+}} : (!cir.ptr<!cir.float>, !cir.float) -> !cir.float {cir.amdgpu_no_fine_grained_memory, cir.amdgpu_no_remote_memory} +// LLVM-LABEL: @_Z12after_scopedPff +// LLVM: atomicrmw fadd ptr %{{.+}}, float %{{.+}} syncscope("agent") monotonic, align 4{{.*}}!amdgpu.no.fine.grained.memory +__device__ float after_scoped(float *p, float v) { + [[clang::atomic(fine_grained_memory)]] { (void)0; } + return __hip_atomic_fetch_add(p, v, __ATOMIC_RELAXED, 4); +} diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-raw-atomic.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-raw-atomic.hip new file mode 100644 index 0000000000000..17989f71c87eb --- /dev/null +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-raw-atomic.hip @@ -0,0 +1,183 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx942 -x hip \ +// RUN: -std=c++11 -fclangir -fcuda-is-device -emit-cir %s -o - \ +// RUN: | FileCheck --check-prefix=CIR %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx942 -x hip \ +// RUN: -std=c++11 -fclangir -fcuda-is-device -emit-llvm %s -o - \ +// RUN: | FileCheck --check-prefix=LLVM %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx942 -x hip \ +// RUN: -std=c++11 -fcuda-is-device -emit-llvm %s -o - \ +// RUN: | FileCheck --check-prefix=LLVM %s + +#define __device__ __attribute__((device)) + +typedef __attribute__((address_space(3))) float *float_lds_ptr; +typedef __attribute__((address_space(3))) double *double_lds_ptr; +typedef __attribute__((address_space(1))) float *float_global_ptr; +typedef __attribute__((address_space(1))) double *double_global_ptr; + +// CIR-LABEL: @_Z5inc32PVjj +// CIR: cir.atomic.fetch uinc_wrap seq_cst syncscope(device) fetch_first %{{.+}}, %{{.+}} volatile : (!cir.ptr<!u32i>, !u32i) -> !u32i {cir.amdgpu_raw_atomic} +// LLVM-LABEL: @_Z5inc32PVjj +// LLVM: atomicrmw volatile uinc_wrap ptr %{{.+}}, i32 %{{.+}} syncscope("agent") seq_cst, align 4, !amdgpu.no.fine.grained.memory +__device__ unsigned int inc32(unsigned int volatile *p, unsigned int v) { + return __builtin_amdgcn_atomic_inc32(p, v, __ATOMIC_SEQ_CST, "agent"); +} + +// CIR-LABEL: @_Z17inc32_nonvolatilePjj +// CIR: cir.atomic.fetch uinc_wrap seq_cst syncscope(device) fetch_first %{{.+}}, %{{.+}} : (!cir.ptr<!u32i>, !u32i) -> !u32i {cir.amdgpu_raw_atomic} +// LLVM-LABEL: @_Z17inc32_nonvolatilePjj +// LLVM: atomicrmw uinc_wrap ptr %{{.+}}, i32 %{{.+}} syncscope("agent") seq_cst, align 4, !amdgpu.no.fine.grained.memory +__device__ unsigned int inc32_nonvolatile(unsigned int *p, unsigned int v) { + return __builtin_amdgcn_atomic_inc32(p, v, __ATOMIC_SEQ_CST, "agent"); +} + +// CIR-LABEL: @_Z5inc64PVmm +// CIR: cir.atomic.fetch uinc_wrap seq_cst syncscope(device) fetch_first %{{.+}}, %{{.+}} volatile : (!cir.ptr<!u64i>, !u64i) -> !u64i {cir.amdgpu_raw_atomic} +// LLVM-LABEL: @_Z5inc64PVmm +// LLVM: atomicrmw volatile uinc_wrap ptr %{{.+}}, i64 %{{.+}} syncscope("agent") seq_cst, align 8, !amdgpu.no.fine.grained.memory +__device__ unsigned long inc64(unsigned long volatile *p, unsigned long v) { + return __builtin_amdgcn_atomic_inc64(p, v, __ATOMIC_SEQ_CST, "agent"); +} + +// CIR-LABEL: @_Z5dec32PVjj +// CIR: cir.atomic.fetch udec_wrap seq_cst syncscope(workgroup) fetch_first %{{.+}}, %{{.+}} volatile : (!cir.ptr<!u32i>, !u32i) -> !u32i {cir.amdgpu_raw_atomic} +// LLVM-LABEL: @_Z5dec32PVjj +// LLVM: atomicrmw volatile udec_wrap ptr %{{.+}}, i32 %{{.+}} syncscope("workgroup") seq_cst, align 4, !amdgpu.no.fine.grained.memory +__device__ unsigned int dec32(unsigned int volatile *p, unsigned int v) { + return __builtin_amdgcn_atomic_dec32(p, v, __ATOMIC_SEQ_CST, "workgroup"); +} + +// CIR-LABEL: @_Z5dec64PVmm +// CIR: cir.atomic.fetch udec_wrap seq_cst syncscope(system) fetch_first %{{.+}}, %{{.+}} volatile : (!cir.ptr<!u64i>, !u64i) -> !u64i {cir.amdgpu_raw_atomic} +// LLVM-LABEL: @_Z5dec64PVmm +// LLVM: atomicrmw volatile udec_wrap ptr %{{.+}}, i64 %{{.+}} seq_cst, align 8, !amdgpu.no.fine.grained.memory +__device__ unsigned long dec64(unsigned long volatile *p, unsigned long v) { + return __builtin_amdgcn_atomic_dec64(p, v, __ATOMIC_SEQ_CST, ""); +} + +// CIR-LABEL: @_Z8ds_faddfPU3AS3ff +// CIR: cir.atomic.fetch add relaxed syncscope(system) fetch_first %{{.+}}, %{{.+}} : (!cir.ptr<!cir.float, target_address_space(3)>, !cir.float) -> !cir.float {cir.amdgpu_raw_atomic} +// LLVM-LABEL: @_Z8ds_faddfPU3AS3ff +// LLVM: atomicrmw fadd ptr addrspace(3) %{{.+}}, float %{{.+}} monotonic, align 4 +// LLVM-NOT: amdgpu.no.fine.grained.memory +__device__ float ds_faddf(float_lds_ptr p, float v) { + return __builtin_amdgcn_ds_faddf(p, v, 0, 0, false); +} + +// CIR-LABEL: @_Z17ds_faddf_volatilePU3AS3ff +// CIR: cir.atomic.fetch add relaxed syncscope(system) fetch_first %{{.+}}, %{{.+}} volatile : (!cir.ptr<!cir.float, target_address_space(3)>, !cir.float) -> !cir.float {cir.amdgpu_raw_atomic} +// LLVM-LABEL: @_Z17ds_faddf_volatilePU3AS3ff +// LLVM: atomicrmw volatile fadd ptr addrspace(3) %{{.+}}, float %{{.+}} monotonic, align 4 +__device__ float ds_faddf_volatile(float_lds_ptr p, float v) { + return __builtin_amdgcn_ds_faddf(p, v, 0, 0, true); +} + +// CIR-LABEL: @_Z8ds_fminfPU3AS3ff +// CIR: cir.atomic.fetch min relaxed syncscope(system) fetch_first %{{.+}}, %{{.+}} : (!cir.ptr<!cir.float, target_address_space(3)>, !cir.float) -> !cir.float {cir.amdgpu_raw_atomic} +// LLVM-LABEL: @_Z8ds_fminfPU3AS3ff +// LLVM: atomicrmw fmin ptr addrspace(3) %{{.+}}, float %{{.+}} monotonic, align 4 +__device__ float ds_fminf(float_lds_ptr p, float v) { + return __builtin_amdgcn_ds_fminf(p, v, 0, 0, false); +} + +// CIR-LABEL: @_Z8ds_fmaxfPU3AS3ff +// CIR: cir.atomic.fetch max relaxed syncscope(system) fetch_first %{{.+}}, %{{.+}} : (!cir.ptr<!cir.float, target_address_space(3)>, !cir.float) -> !cir.float {cir.amdgpu_raw_atomic} +// LLVM-LABEL: @_Z8ds_fmaxfPU3AS3ff +// LLVM: atomicrmw fmax ptr addrspace(3) %{{.+}}, float %{{.+}} monotonic, align 4 +__device__ float ds_fmaxf(float_lds_ptr p, float v) { + return __builtin_amdgcn_ds_fmaxf(p, v, 0, 0, false); +} + +// CIR-LABEL: @_Z18ds_atomic_fadd_f32PU3AS3ff +// CIR: cir.atomic.fetch add relaxed syncscope(device) fetch_first %{{.+}}, %{{.+}} : (!cir.ptr<!cir.float, target_address_space(3)>, !cir.float) -> !cir.float {cir.amdgpu_raw_atomic} +// LLVM-LABEL: @_Z18ds_atomic_fadd_f32PU3AS3ff +// LLVM: atomicrmw fadd ptr addrspace(3) %{{.+}}, float %{{.+}} syncscope("agent") monotonic, align 4 +// LLVM-NOT: amdgpu.no.fine.grained.memory +__device__ float ds_atomic_fadd_f32(float_lds_ptr p, float v) { + return __builtin_amdgcn_ds_atomic_fadd_f32(p, v); +} + +// CIR-LABEL: @_Z18ds_atomic_fadd_f64PU3AS3dd +// CIR: cir.atomic.fetch add relaxed syncscope(device) fetch_first %{{.+}}, %{{.+}} : (!cir.ptr<!cir.double, target_address_space(3)>, !cir.double) -> !cir.double {cir.amdgpu_raw_atomic} +// LLVM-LABEL: @_Z18ds_atomic_fadd_f64PU3AS3dd +// LLVM: atomicrmw fadd ptr addrspace(3) %{{.+}}, double %{{.+}} syncscope("agent") monotonic, align 8 +// LLVM-NOT: amdgpu.no.fine.grained.memory +__device__ double ds_atomic_fadd_f64(double_lds_ptr p, double v) { + return __builtin_amdgcn_ds_atomic_fadd_f64(p, v); +} + +// CIR-LABEL: @_Z22global_atomic_fadd_f32PU3AS1ff +// CIR: cir.atomic.fetch add relaxed syncscope(device) fetch_first %{{.+}}, %{{.+}} : (!cir.ptr<!cir.float, target_address_space(1)>, !cir.float) -> !cir.float {cir.amdgpu_raw_atomic} +// LLVM-LABEL: @_Z22global_atomic_fadd_f32PU3AS1ff +// LLVM: atomicrmw fadd ptr addrspace(1) %{{.+}}, float %{{.+}} syncscope("agent") monotonic, align 4{{.*}}!atomic.ignore.denormal.mode{{.*}}!amdgpu.no.fine.grained.memory +__device__ float global_atomic_fadd_f32(float_global_ptr p, float v) { + return __builtin_amdgcn_global_atomic_fadd_f32(p, v); +} + +// CIR-LABEL: @_Z22global_atomic_fadd_f64PU3AS1dd +// CIR: cir.atomic.fetch add relaxed syncscope(device) fetch_first %{{.+}}, %{{.+}} : (!cir.ptr<!cir.double, target_address_space(1)>, !cir.double) -> !cir.double {cir.amdgpu_raw_atomic} +// LLVM-LABEL: @_Z22global_atomic_fadd_f64PU3AS1dd +// LLVM: atomicrmw fadd ptr addrspace(1) %{{.+}}, double %{{.+}} syncscope("agent") monotonic, align 8, !amdgpu.no.fine.grained.memory +// LLVM-NOT: atomic.ignore.denormal.mode +__device__ double global_atomic_fadd_f64(double_global_ptr p, double v) { + return __builtin_amdgcn_global_atomic_fadd_f64(p, v); +} + +// CIR-LABEL: @_Z20flat_atomic_fadd_f32Pff +// CIR: cir.atomic.fetch add relaxed syncscope(device) fetch_first %{{.+}}, %{{.+}} : (!cir.ptr<!cir.float>, !cir.float) -> !cir.float {cir.amdgpu_raw_atomic} +// LLVM-LABEL: @_Z20flat_atomic_fadd_f32Pff +// LLVM: atomicrmw fadd ptr %{{.+}}, float %{{.+}} syncscope("agent") monotonic, align 4{{.*}}!atomic.ignore.denormal.mode{{.*}}!amdgpu.no.fine.grained.memory +__device__ float flat_atomic_fadd_f32(float *p, float v) { + return __builtin_amdgcn_flat_atomic_fadd_f32(p, v); +} + +// CIR-LABEL: @_Z20flat_atomic_fadd_f64Pdd +// CIR: cir.atomic.fetch add relaxed syncscope(device) fetch_first %{{.+}}, %{{.+}} : (!cir.ptr<!cir.double>, !cir.double) -> !cir.double {cir.amdgpu_raw_atomic} +// LLVM-LABEL: @_Z20flat_atomic_fadd_f64Pdd +// LLVM: atomicrmw fadd ptr %{{.+}}, double %{{.+}} syncscope("agent") monotonic, align 8, !amdgpu.no.fine.grained.memory +// LLVM-NOT: atomic.ignore.denormal.mode +__device__ double flat_atomic_fadd_f64(double *p, double v) { + return __builtin_amdgcn_flat_atomic_fadd_f64(p, v); +} + +// CIR-LABEL: @_Z22global_atomic_fmin_f64PU3AS1dd +// CIR: cir.atomic.fetch min relaxed syncscope(device) fetch_first %{{.+}}, %{{.+}} : (!cir.ptr<!cir.double, target_address_space(1)>, !cir.double) -> !cir.double {cir.amdgpu_raw_atomic} +// LLVM-LABEL: @_Z22global_atomic_fmin_f64PU3AS1dd +// LLVM: atomicrmw fmin ptr addrspace(1) %{{.+}}, double %{{.+}} syncscope("agent") monotonic, align 8, !amdgpu.no.fine.grained.memory +// LLVM-NOT: atomic.ignore.denormal.mode +__device__ double global_atomic_fmin_f64(double_global_ptr p, double v) { + return __builtin_amdgcn_global_atomic_fmin_f64(p, v); +} + +// CIR-LABEL: @_Z22global_atomic_fmax_f64PU3AS1dd +// CIR: cir.atomic.fetch max relaxed syncscope(device) fetch_first %{{.+}}, %{{.+}} : (!cir.ptr<!cir.double, target_address_space(1)>, !cir.double) -> !cir.double {cir.amdgpu_raw_atomic} +// LLVM-LABEL: @_Z22global_atomic_fmax_f64PU3AS1dd +// LLVM: atomicrmw fmax ptr addrspace(1) %{{.+}}, double %{{.+}} syncscope("agent") monotonic, align 8, !amdgpu.no.fine.grained.memory +__device__ double global_atomic_fmax_f64(double_global_ptr p, double v) { + return __builtin_amdgcn_global_atomic_fmax_f64(p, v); +} + +// CIR-LABEL: @_Z20flat_atomic_fmin_f64Pdd +// CIR: cir.atomic.fetch min relaxed syncscope(device) fetch_first %{{.+}}, %{{.+}} : (!cir.ptr<!cir.double>, !cir.double) -> !cir.double {cir.amdgpu_raw_atomic} +// LLVM-LABEL: @_Z20flat_atomic_fmin_f64Pdd +// LLVM: atomicrmw fmin ptr %{{.+}}, double %{{.+}} syncscope("agent") monotonic, align 8, !amdgpu.no.fine.grained.memory +__device__ double flat_atomic_fmin_f64(double *p, double v) { + return __builtin_amdgcn_flat_atomic_fmin_f64(p, v); +} + +// CIR-LABEL: @_Z20flat_atomic_fmax_f64Pdd +// CIR: cir.atomic.fetch max relaxed syncscope(device) fetch_first %{{.+}}, %{{.+}} : (!cir.ptr<!cir.double>, !cir.double) -> !cir.double {cir.amdgpu_raw_atomic} +// LLVM-LABEL: @_Z20flat_atomic_fmax_f64Pdd +// LLVM: atomicrmw fmax ptr %{{.+}}, double %{{.+}} syncscope("agent") monotonic, align 8, !amdgpu.no.fine.grained.memory +__device__ double flat_atomic_fmax_f64(double *p, double v) { + return __builtin_amdgcn_flat_atomic_fmax_f64(p, v); +} + +// CIR-LABEL: @_Z5fencev +// CIR: cir.atomic.fence syncscope(device) seq_cst +// LLVM-LABEL: @_Z5fencev +// LLVM: fence syncscope("agent") seq_cst +__device__ void fence() { + __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "agent"); +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
