[clang] [CIR][AMDGPU] Emit !noalias.addrspace on HIP and OpenCL atomics (PR #232350)
via cfe-commits
cfe-commits at lists.llvm.org
Thu Oct 15 07:45:56 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-clang
Author: Arseniy Obolenskiy (aobolensk)
<details>
<summary>Changes</summary>
Add it to `__hip_*` and `__opencl_*` atomicrmw and cmpxchg on a flat pointer, so the backend can skip the private memory check
---
Full diff: https://github.com/llvm/llvm-project/pull/232350.diff
6 Files Affected:
- (modified) clang/include/clang/CIR/Dialect/IR/CIRDialect.td (+1)
- (modified) clang/include/clang/CIR/MissingFeatures.h (-1)
- (modified) clang/lib/CIR/CodeGen/CIRGenAtomic.cpp (+38-36)
- (modified) clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp (+14-1)
- (modified) clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp (+16)
- (modified) clang/test/CIR/CodeGenHIP/atomic-options.hip (+39-6)
``````````diff
diff --git a/clang/include/clang/CIR/Dialect/IR/CIRDialect.td b/clang/include/clang/CIR/Dialect/IR/CIRDialect.td
index 5fb3366d0c4e7..dff732ec40f74 100644
--- a/clang/include/clang/CIR/Dialect/IR/CIRDialect.td
+++ b/clang/include/clang/CIR/Dialect/IR/CIRDialect.td
@@ -127,6 +127,7 @@ def CIR_Dialect : Dialect {
static llvm::StringRef getAMDGPUNoRemoteMemoryAttrName() { return "cir.amdgpu_no_remote_memory"; }
static llvm::StringRef getAMDGPUIgnoreDenormalModeAttrName() { return "cir.amdgpu_ignore_denormal_mode"; }
static llvm::StringRef getAMDGPURawAtomicAttrName() { return "cir.amdgpu_raw_atomic"; }
+ static llvm::StringRef getAMDGPUNoPrivateAddrSpaceAttrName() { return "cir.amdgpu_no_private_addrspace"; }
static llvm::StringRef getOpenCLKernelArgMetadataAttrName() { return "cir.cl.kernel_arg_metadata"; }
static llvm::StringRef getMaxWorkGroupSizeAttrName() { return "cir.max_work_group_size"; }
static llvm::StringRef getOpenCLVersionAttrName() { return "cir.cl.version"; }
diff --git a/clang/include/clang/CIR/MissingFeatures.h b/clang/include/clang/CIR/MissingFeatures.h
index cd1e43b17ec3f..709e4a8239136 100644
--- a/clang/include/clang/CIR/MissingFeatures.h
+++ b/clang/include/clang/CIR/MissingFeatures.h
@@ -155,7 +155,6 @@ 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
diff --git a/clang/lib/CIR/CodeGen/CIRGenAtomic.cpp b/clang/lib/CIR/CodeGen/CIRGenAtomic.cpp
index 095ebdaa04ce1..ccbb453b13673 100644
--- a/clang/lib/CIR/CodeGen/CIRGenAtomic.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenAtomic.cpp
@@ -483,6 +483,41 @@ static void emitMemOrderCaseLabel(CIRGenBuilderTy &builder, mlir::Location loc,
builder.restoreInsertionPoint(insertPoint);
}
+/// 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,
+ const AtomicExpr *expr) {
+ // 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());
+
+ // Lowering keeps this only on a flat pointer.
+ mlir::UnitAttr unit = cgf.getBuilder().getUnitAttr();
+ if (expr->threadPrivateMemoryAtomicsAreUndefined())
+ op->setAttr(cir::CIRDialect::getAMDGPUNoPrivateAddrSpaceAttrName(), unit);
+
+ // 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();
+ if (!atomicOpts.getOption(clang::AtomicOptionKind::FineGrainedMemory))
+ op->setAttr(cir::CIRDialect::getAMDGPUNoFineGrainedMemoryAttrName(), unit);
+ if (!atomicOpts.getOption(clang::AtomicOptionKind::RemoteMemory))
+ op->setAttr(cir::CIRDialect::getAMDGPUNoRemoteMemoryAttrName(), 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::CIRDialect::getAMDGPUIgnoreDenormalModeAttrName(), unit);
+}
+
static void emitAtomicCmpXchg(CIRGenFunction &cgf, AtomicExpr *e, bool isWeak,
Address dest, Address ptr, Address val1,
Address val2, Address expectedResult,
@@ -505,6 +540,8 @@ static void emitAtomicCmpXchg(CIRGenFunction &cgf, AtomicExpr *e, bool isWeak,
cmpxchg.setIsVolatile(e->isVolatile());
cmpxchg.setWeak(isWeak);
+ if (cgf.cgm.getTriple().isAMDGCN())
+ setAMDGPUAtomicMetadata(cgf, cmpxchg, e);
mlir::Value failed = builder.createNot(cmpxchg.getSuccess());
cir::IfOp::create(
@@ -661,41 +698,6 @@ 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::CIRDialect::getAMDGPUNoFineGrainedMemoryAttrName(), unit);
- if (!atomicOpts.getOption(clang::AtomicOptionKind::RemoteMemory))
- op->setAttr(cir::CIRDialect::getAMDGPUNoRemoteMemoryAttrName(), 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::CIRDialect::getAMDGPUIgnoreDenormalModeAttrName(), unit);
-}
-
static void emitAtomicOp(CIRGenFunction &cgf, AtomicExpr *expr, Address dest,
Address ptr, Address val1, Address val2,
Address expectedResult, Expr *isWeakExpr,
@@ -969,7 +971,7 @@ static void emitAtomicOp(CIRGenFunction &cgf, AtomicExpr *expr, Address dest,
rmwOp->setAttr("fetch_first", builder.getUnitAttr());
if (cgf.cgm.getTriple().isAMDGCN())
- setAMDGPUAtomicMetadata(cgf, rmwOp);
+ setAMDGPUAtomicMetadata(cgf, rmwOp, expr);
mlir::Value result = rmwOp->getResult(0);
diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
index fb2be40281883..695deda3980b4 100644
--- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
+++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
@@ -1411,6 +1411,16 @@ getLLVMSyncScope(std::optional<cir::SyncScopeKind> syncScope,
return std::nullopt;
}
+static void forwardAMDGPUNoPrivateAddrSpace(mlir::Operation *op,
+ mlir::Operation *llvmOp,
+ mlir::Value llvmPtr) {
+ llvm::StringRef name = cir::CIRDialect::getAMDGPUNoPrivateAddrSpaceAttrName();
+ auto ptrTy = mlir::cast<mlir::LLVM::LLVMPointerType>(llvmPtr.getType());
+ if (mlir::Attribute a = op->getAttr(name);
+ a && ptrTy.getAddressSpace() == llvm::AMDGPUAS::FLAT_ADDRESS)
+ llvmOp->setAttr(name, a);
+}
+
mlir::LogicalResult CIRToLLVMAtomicCmpXchgOpLowering::matchAndRewrite(
cir::AtomicCmpXchgOp op, OpAdaptor adaptor,
mlir::ConversionPatternRewriter &rewriter) const {
@@ -1427,6 +1437,7 @@ mlir::LogicalResult CIRToLLVMAtomicCmpXchgOpLowering::matchAndRewrite(
cmpxchg.setAlignment(adaptor.getAlignment());
cmpxchg.setWeak(adaptor.getWeak());
cmpxchg.setVolatile_(adaptor.getIsVolatile());
+ forwardAMDGPUNoPrivateAddrSpace(op, cmpxchg, adaptor.getPtr());
// Check result and apply stores accordingly.
auto old = mlir::LLVM::ExtractValueOp::create(rewriter, op.getLoc(),
@@ -1445,9 +1456,10 @@ mlir::LogicalResult CIRToLLVMAtomicXchgOpLowering::matchAndRewrite(
mlir::LLVM::AtomicOrdering llvmOrder = getLLVMMemOrder(adaptor.getMemOrder());
llvm::StringRef llvmSyncScope = getLLVMSyncScope(
op.getSyncScope(), op->getParentOfType<mlir::ModuleOp>());
- rewriter.replaceOpWithNewOp<mlir::LLVM::AtomicRMWOp>(
+ auto xchg = rewriter.replaceOpWithNewOp<mlir::LLVM::AtomicRMWOp>(
op, mlir::LLVM::AtomicBinOp::xchg, adaptor.getPtr(), adaptor.getVal(),
llvmOrder, llvmSyncScope, /*alignment=*/0, op.getIsVolatile());
+ forwardAMDGPUNoPrivateAddrSpace(op, xchg, adaptor.getPtr());
return mlir::success();
}
@@ -1660,6 +1672,7 @@ mlir::LogicalResult CIRToLLVMAtomicFetchOpLowering::matchAndRewrite(
cir::CIRDialect::getAMDGPUIgnoreDenormalModeAttrName()})
if (mlir::Attribute a = op->getAttr(marker))
rmwVal->setAttr(marker, a);
+ forwardAMDGPUNoPrivateAddrSpace(op, rmwVal, adaptor.getPtr());
// The AMDGPU raw hardware atomic builtins need metadata for the backend to
// select the native instruction. LDS atomics are always native, so the
diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp
index 9fc12631d8ed5..0cd3f2a19747e 100644
--- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp
+++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp
@@ -20,7 +20,9 @@
#include "llvm/IR/Constant.h"
#include "llvm/IR/Constants.h"
#include "llvm/IR/GlobalVariable.h"
+#include "llvm/IR/MDBuilder.h"
#include "llvm/IR/Metadata.h"
+#include "llvm/Support/AMDGPUAddrSpace.h"
using namespace llvm;
@@ -75,6 +77,9 @@ class CIRDialectLLVMIRTranslationInterface
attribute.getName() ==
cir::CIRDialect::getAMDGPUIgnoreDenormalModeAttrName()) {
amendAMDGPUAtomicMetadata(instructions, attribute, moduleTranslation);
+ } else if (attribute.getName() ==
+ cir::CIRDialect::getAMDGPUNoPrivateAddrSpaceAttrName()) {
+ amendAMDGPUNoPrivateAddrSpace(instructions, moduleTranslation);
}
return mlir::success();
}
@@ -128,6 +133,17 @@ class CIRDialectLLVMIRTranslationInterface
inst->setMetadata(mdName, empty);
}
+ void amendAMDGPUNoPrivateAddrSpace(
+ llvm::ArrayRef<llvm::Instruction *> instructions,
+ mlir::LLVM::ModuleTranslation &moduleTranslation) const {
+ llvm::MDBuilder mdHelper(moduleTranslation.getLLVMContext());
+ llvm::MDNode *asRange = mdHelper.createRange(
+ llvm::APInt(32, llvm::AMDGPUAS::PRIVATE_ADDRESS),
+ llvm::APInt(32, llvm::AMDGPUAS::PRIVATE_ADDRESS + 1));
+ for (llvm::Instruction *inst : instructions)
+ inst->setMetadata(llvm::LLVMContext::MD_noalias_addrspace, asRange);
+ }
+
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
index 7fa524f23de51..5dc137ed7cfc5 100644
--- a/clang/test/CIR/CodeGenHIP/atomic-options.hip
+++ b/clang/test/CIR/CodeGenHIP/atomic-options.hip
@@ -10,7 +10,7 @@
#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}
+// 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_private_addrspace, 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) {
@@ -18,7 +18,7 @@ __device__ float plain_add(float *p, float v) {
}
// 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}
+// CIR: cir.atomic.fetch add relaxed syncscope(hip_agent) fetch_first %{{.+}}, %{{.+}} : (!cir.ptr<!cir.float>, !cir.float) -> !cir.float {cir.amdgpu_no_private_addrspace, cir.amdgpu_no_remote_memory}
// LLVM-LABEL: @_Z10unsafe_addPff
// LLVM: atomicrmw fadd ptr %{{.+}}, float %{{.+}} syncscope("agent") monotonic
// LLVM-NOT: amdgpu.no.fine.grained.memory
@@ -29,7 +29,7 @@ __device__ float unsafe_add(float *p, float v) {
}
// 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}
+// 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_private_addrspace, 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) {
@@ -39,7 +39,7 @@ __device__ float no_fg_add(float *p, float v) {
}
// 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}
+// 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_private_addrspace, cir.amdgpu_no_remote_memory}
// LLVM-LABEL: @_Z9no_fg_addPdd
// LLVM: atomicrmw fadd ptr %{{.+}}, double %{{.+}} syncscope("agent") monotonic
// LLVM-NOT: atomic.ignore.denormal.mode
@@ -50,7 +50,7 @@ __device__ double no_fg_add(double *p, double v) {
}
// 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}
+// 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_private_addrspace, 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) {
@@ -58,10 +58,43 @@ __device__ double workgroup_dd(double *p, double v) {
}
// 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}
+// 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_private_addrspace, 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);
}
+
+// CIR-LABEL: @_Z4xchgPii
+// CIR: cir.atomic.xchg relaxed syncscope(hip_agent) %{{.+}}, %{{.+}} : (!cir.ptr<!s32i>, !s32i) -> !s32i {cir.amdgpu_no_private_addrspace}
+// LLVM-LABEL: @_Z4xchgPii
+// LLVM: atomicrmw xchg ptr %{{.+}}, i32 %{{.+}} syncscope("agent") monotonic, align 4, !noalias.addrspace ![[#NOPRIV:]]
+__device__ int xchg(int *p, int v) {
+ return __hip_atomic_exchange(p, v, __ATOMIC_RELAXED, 4);
+}
+
+// CIR-LABEL: @_Z7cmpxchgPiii
+// CIR: cir.atomic.cmpxchg success(relaxed) failure(relaxed) syncscope(hip_agent) %{{.+}}, %{{.+}}, %{{.+}} align(4) : (!cir.ptr<!s32i>, !s32i, !s32i) -> (!s32i, !cir.bool) {cir.amdgpu_no_private_addrspace}
+// LLVM-LABEL: @_Z7cmpxchgPiii
+// LLVM: cmpxchg ptr %{{.+}}, i32 %{{.+}}, i32 %{{.+}} syncscope("agent") monotonic monotonic, align 4, !noalias.addrspace ![[#NOPRIV]]{{$}}
+__device__ bool cmpxchg(int *p, int e, int d) {
+ return __hip_atomic_compare_exchange_strong(p, &e, d, __ATOMIC_RELAXED,
+ __ATOMIC_RELAXED, 4);
+}
+
+// LLVM-LABEL: @_Z10global_addPU3AS1ii
+// LLVM: atomicrmw add ptr addrspace(1) %{{.+}}, i32 %{{.+}} syncscope("agent") monotonic, align 4, !amdgpu.no.fine.grained.memory
+__device__ int global_add(__attribute__((address_space(1))) int *p, int v) {
+ return __hip_atomic_fetch_add(p, v, __ATOMIC_RELAXED, 4);
+}
+
+// CIR-LABEL: @_Z7gnu_addPii
+// CIR: cir.atomic.fetch add relaxed syncscope(system) fetch_first %{{.+}}, %{{.+}} : (!cir.ptr<!s32i>, !s32i) -> !s32i {cir.amdgpu_no_fine_grained_memory, cir.amdgpu_no_remote_memory}
+// LLVM-LABEL: @_Z7gnu_addPii
+// LLVM: atomicrmw add ptr %{{.+}}, i32 %{{.+}} monotonic, align 4, !amdgpu.no.fine.grained.memory
+__device__ int gnu_add(int *p, int v) {
+ return __atomic_fetch_add(p, v, __ATOMIC_RELAXED);
+}
+
+// LLVM: ![[#NOPRIV]] = !{i32 5, i32 6}
``````````
</details>
https://github.com/llvm/llvm-project/pull/232350
More information about the cfe-commits
mailing list