https://github.com/aobolensk updated https://github.com/llvm/llvm-project/pull/233440
>From f78e16a20f96a347fe1aaf8c7ff08dcba6e4c848 Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Mon, 19 Oct 2026 16:01:00 +0200 Subject: [PATCH] [CIR][AMDGPU] Emit `!amdgpu.no.{fine.grained,remote}.memory` metadata on atomic xchg Add it to xchg too (not only atomic fetch ops), so IR matches what classic codegen emits --- clang/lib/CIR/CodeGen/CIRGenAtomic.cpp | 7 +++--- .../CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp | 22 ++++++++++++------- clang/test/CIR/CodeGenHIP/atomic-options.hip | 4 ++-- 3 files changed, 20 insertions(+), 13 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenAtomic.cpp b/clang/lib/CIR/CodeGen/CIRGenAtomic.cpp index a758057f78167..c4ebf6f0afcec 100644 --- a/clang/lib/CIR/CodeGen/CIRGenAtomic.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenAtomic.cpp @@ -178,8 +178,7 @@ static void setAMDGPUAtomicMetadata(CIRGenFunction &cgf, mlir::Operation *op, // 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) + if (!mlir::isa<cir::AtomicFetchOp, cir::AtomicXchgOp>(op)) return; clang::AtomicOptions atomicOpts = cgf.cgm.getAtomicOpts(); @@ -188,7 +187,9 @@ static void setAMDGPUAtomicMetadata(CIRGenFunction &cgf, mlir::Operation *op, 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) && + auto fetchOp = mlir::dyn_cast<cir::AtomicFetchOp>(op); + if (fetchOp && + atomicOpts.getOption(clang::AtomicOptionKind::IgnoreDenormalMode) && fetchOp.getBinop() == cir::AtomicFetchKind::Add && mlir::isa<cir::SingleType>(fetchOp.getVal().getType())) op->setAttr(cir::CIRDialect::getAMDGPUIgnoreDenormalModeAttrName(), unit); diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp index 3b2e76fde1990..40e6a5624366c 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp @@ -1467,6 +1467,18 @@ static void forwardAMDGPUNoPrivateAddrSpace(mlir::Operation *op, llvmOp->setAttr(name, a); } +// CIRGen decides the metadata for a C++/HIP atomic from the atomic options +// in effect, so those markers are simply carried across. +static void forwardAMDGPUAtomicOptionMarkers(mlir::Operation *op, + mlir::Operation *llvmOp) { + for (llvm::StringRef marker : + {cir::CIRDialect::getAMDGPUNoFineGrainedMemoryAttrName(), + cir::CIRDialect::getAMDGPUNoRemoteMemoryAttrName(), + cir::CIRDialect::getAMDGPUIgnoreDenormalModeAttrName()}) + if (mlir::Attribute a = op->getAttr(marker)) + llvmOp->setAttr(marker, a); +} + mlir::LogicalResult CIRToLLVMAtomicCmpXchgOpLowering::matchAndRewrite( cir::AtomicCmpXchgOp op, OpAdaptor adaptor, mlir::ConversionPatternRewriter &rewriter) const { @@ -1505,6 +1517,7 @@ mlir::LogicalResult CIRToLLVMAtomicXchgOpLowering::matchAndRewrite( auto xchg = rewriter.replaceOpWithNewOp<mlir::LLVM::AtomicRMWOp>( op, mlir::LLVM::AtomicBinOp::xchg, adaptor.getPtr(), adaptor.getVal(), llvmOrder, llvmSyncScope, /*alignment=*/0, op.getIsVolatile()); + forwardAMDGPUAtomicOptionMarkers(op, xchg); forwardAMDGPUNoPrivateAddrSpace(op, xchg, adaptor.getPtr()); return mlir::success(); } @@ -1710,14 +1723,7 @@ mlir::LogicalResult CIRToLLVMAtomicFetchOpLowering::matchAndRewrite( 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::CIRDialect::getAMDGPUNoFineGrainedMemoryAttrName(), - cir::CIRDialect::getAMDGPUNoRemoteMemoryAttrName(), - cir::CIRDialect::getAMDGPUIgnoreDenormalModeAttrName()}) - if (mlir::Attribute a = op->getAttr(marker)) - rmwVal->setAttr(marker, a); + forwardAMDGPUAtomicOptionMarkers(op, rmwVal); forwardAMDGPUNoPrivateAddrSpace(op, rmwVal, adaptor.getPtr()); // The AMDGPU raw hardware atomic builtins need metadata for the backend to diff --git a/clang/test/CIR/CodeGenHIP/atomic-options.hip b/clang/test/CIR/CodeGenHIP/atomic-options.hip index c88f4c7c687f5..da0f2f2dcdb73 100644 --- a/clang/test/CIR/CodeGenHIP/atomic-options.hip +++ b/clang/test/CIR/CodeGenHIP/atomic-options.hip @@ -67,9 +67,9 @@ __device__ float after_scoped(float *p, float v) { } // CIR-LABEL: @_Z4xchgPii -// CIR: cir.atomic.xchg relaxed syncscope(hip_agent) %{{.+}}, %{{.+}} : (!cir.ptr<!s32i>, !s32i) -> !s32i {cir.amdgpu_no_private_addrspace} +// CIR: cir.atomic.xchg relaxed syncscope(hip_agent) %{{.+}}, %{{.+}} : (!cir.ptr<!s32i>, !s32i) -> !s32i {cir.amdgpu_no_fine_grained_memory, cir.amdgpu_no_private_addrspace, cir.amdgpu_no_remote_memory} // LLVM-LABEL: @_Z4xchgPii -// LLVM: atomicrmw xchg ptr %{{.+}}, i32 %{{.+}} syncscope("agent") monotonic, align 4, !noalias.addrspace ![[#NOPRIV:]] +// LLVM: atomicrmw xchg ptr %{{.+}}, i32 %{{.+}} syncscope("agent") monotonic, align 4, !noalias.addrspace ![[#NOPRIV:]], !amdgpu.no.fine.grained.memory !{{[0-9]+}}, !amdgpu.no.remote.memory !{{[0-9]+}}{{$}} __device__ int xchg(int *p, int v) { return __hip_atomic_exchange(p, v, __ATOMIC_RELAXED, 4); } _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
