Author: Arseniy Obolenskiy Date: 2026-10-16T08:01:15+02:00 New Revision: 95c78f20cc7523ba47ee9e81c877b46a9873775f
URL: https://github.com/llvm/llvm-project/commit/95c78f20cc7523ba47ee9e81c877b46a9873775f DIFF: https://github.com/llvm/llvm-project/commit/95c78f20cc7523ba47ee9e81c877b46a9873775f.diff LOG: [CIR][AMDGPU] Emit !noalias.addrspace on HIP and OpenCL atomics (#232350) Added: Modified: clang/include/clang/CIR/Dialect/IR/CIRDialect.td clang/include/clang/CIR/MissingFeatures.h clang/lib/CIR/CodeGen/CIRGenAtomic.cpp clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp clang/test/CIR/CodeGenHIP/atomic-options.hip Removed: ################################################################################ diff --git a/clang/include/clang/CIR/Dialect/IR/CIRDialect.td b/clang/include/clang/CIR/Dialect/IR/CIRDialect.td index 429af17eb194b..cd6601c213ba9 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIRDialect.td +++ b/clang/include/clang/CIR/Dialect/IR/CIRDialect.td @@ -133,6 +133,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 4afa6c3a64159..0ff6c6d44d191 100644 --- a/clang/include/clang/CIR/MissingFeatures.h +++ b/clang/include/clang/CIR/MissingFeatures.h @@ -151,7 +151,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..a758057f78167 100644 --- a/clang/lib/CIR/CodeGen/CIRGenAtomic.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenAtomic.cpp @@ -159,6 +159,41 @@ class AtomicInfo { }; } // namespace +/// 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); +} + // This function emits any expression (scalar, complex, or aggregate) // into a temporary alloca. static Address emitValToTemp(CIRGenFunction &cgf, Expr *e) { @@ -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 c745c1864c0cc..a6e48b2985c67 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp @@ -1427,6 +1427,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()); + mlir::Attribute a = op->getAttr(name); + if (a && ptrTy.getAddressSpace() == llvm::AMDGPUAS::FLAT_ADDRESS) + llvmOp->setAttr(name, a); +} + mlir::LogicalResult CIRToLLVMAtomicCmpXchgOpLowering::matchAndRewrite( cir::AtomicCmpXchgOp op, OpAdaptor adaptor, mlir::ConversionPatternRewriter &rewriter) const { @@ -1443,6 +1453,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(), @@ -1461,9 +1472,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(); } @@ -1676,6 +1688,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..2b80606fea613 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; @@ -73,7 +75,9 @@ class CIRDialectLLVMIRTranslationInterface attribute.getName() == cir::CIRDialect::getAMDGPUNoRemoteMemoryAttrName() || attribute.getName() == - cir::CIRDialect::getAMDGPUIgnoreDenormalModeAttrName()) { + cir::CIRDialect::getAMDGPUIgnoreDenormalModeAttrName() || + attribute.getName() == + cir::CIRDialect::getAMDGPUNoPrivateAddrSpaceAttrName()) { amendAMDGPUAtomicMetadata(instructions, attribute, moduleTranslation); } return mlir::success(); @@ -106,6 +110,16 @@ class CIRDialectLLVMIRTranslationInterface llvm::ArrayRef<llvm::Instruction *> instructions, mlir::NamedAttribute attribute, mlir::LLVM::ModuleTranslation &moduleTranslation) const { + if (attribute.getName() == + cir::CIRDialect::getAMDGPUNoPrivateAddrSpaceAttrName()) { + 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); + return; + } llvm::MDNode *empty = llvm::MDNode::get(moduleTranslation.getLLVMContext(), {}); // !atomic.ignore.denormal.mode is a fixed metadata kind so it has to be diff --git a/clang/test/CIR/CodeGenHIP/atomic-options.hip b/clang/test/CIR/CodeGenHIP/atomic-options.hip index 7fa524f23de51..c88f4c7c687f5 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,45 @@ __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); +} + +// CIR-LABEL: @_Z10global_addPU3AS1ii +// CIR: cir.atomic.fetch add relaxed syncscope(hip_agent) fetch_first %{{.+}}, %{{.+}} : (!cir.ptr<!s32i, target_address_space(1)>, !s32i) -> !s32i {cir.amdgpu_no_fine_grained_memory, cir.amdgpu_no_private_addrspace, cir.amdgpu_no_remote_memory} +// 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} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
