https://github.com/ssahasra updated https://github.com/llvm/llvm-project/pull/198250
>From 39e36ba0048cc6ab4fc0e40760c4ae1231a11b19 Mon Sep 17 00:00:00 2001 From: Sameer Sahasrabuddhe <[email protected]> Date: Tue, 12 May 2026 19:36:09 +0530 Subject: [PATCH 1/2] [AMDGPU][Clang] use a ScopeModel when emitting load_monitor [NFC] Assisted-By: Claude Opus 4.6 --- clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp | 25 ++++++++++------- .../builtins-amdgcn-gfx1250-load-monitor.cl | 27 +++++++++++++++++++ clang/test/SemaHIP/incorrect-atomic-scope.hip | 2 +- 3 files changed, 44 insertions(+), 10 deletions(-) diff --git a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp index 751cd9847bd31..1d674803931d2 100644 --- a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp +++ b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp @@ -383,6 +383,21 @@ static llvm::AtomicOrdering mapCABIAtomicOrdering(unsigned AO) { llvm_unreachable("Unknown AtomicOrderingCABI enum"); } +/// Convert a __MEMORY_SCOPE_* integer constant to a metadata node containing +/// the target-specific sync scope string. +static llvm::MetadataAsValue *emitScopeMD( + CodeGenFunction &CGF, unsigned ScopeInt, + llvm::AtomicOrdering AO = llvm::AtomicOrdering::SequentiallyConsistent) { + AtomicScopeGenericModel ScopeModel; + clang::SyncScope Scope = ScopeModel.map(ScopeInt); + StringRef ScopeStr = CGF.CGM.getTargetCodeGenInfo().getLLVMSyncScopeStr( + CGF.CGM.getLangOpts(), Scope, AO); + llvm::LLVMContext &Ctx = CGF.CGM.getLLVMContext(); + llvm::MDNode *MD = + llvm::MDNode::get(Ctx, {llvm::MDString::get(Ctx, ScopeStr)}); + return llvm::MetadataAsValue::get(Ctx, MD); +} + // For processing memory ordering and memory scope arguments of various // amdgcn builtins. // \p Order takes a C++11 compatible memory-ordering specifier and converts @@ -927,22 +942,14 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID, break; } - LLVMContext &Ctx = CGM.getLLVMContext(); llvm::Type *LoadTy = ConvertType(E->getType()); llvm::Value *Addr = EmitScalarExpr(E->getArg(0)); auto *AOExpr = cast<llvm::ConstantInt>(EmitScalarExpr(E->getArg(1))); auto *ScopeExpr = cast<llvm::ConstantInt>(EmitScalarExpr(E->getArg(2))); - - auto Scope = static_cast<SyncScope>(ScopeExpr->getZExtValue()); llvm::AtomicOrdering AO = mapCABIAtomicOrdering(AOExpr->getZExtValue()); - StringRef ScopeStr = CGM.getTargetCodeGenInfo().getLLVMSyncScopeStr( - CGM.getLangOpts(), Scope, AO); - - llvm::MDNode *MD = - llvm::MDNode::get(Ctx, {llvm::MDString::get(Ctx, ScopeStr)}); - llvm::Value *ScopeMD = llvm::MetadataAsValue::get(Ctx, MD); + llvm::Value *ScopeMD = emitScopeMD(*this, ScopeExpr->getZExtValue(), AO); llvm::Function *F = CGM.getIntrinsic(IID, {LoadTy}); return Builder.CreateCall(F, {Addr, AOExpr, ScopeMD}); } diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250-load-monitor.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250-load-monitor.cl index 8ecd6ba61a03e..4e285fcb217e2 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250-load-monitor.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250-load-monitor.cl @@ -64,3 +64,30 @@ v4i test_amdgcn_flat_load_monitor_b128(v4i* inptr) { return __builtin_amdgcn_flat_load_monitor_b128(inptr, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM); } + +// CHECK-GFX1250-LABEL: @test_amdgcn_global_load_monitor_b32_wavefront( +// CHECK-GFX1250-NEXT: entry: +// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.global.load.monitor.b32.i32(ptr addrspace(1) [[INPTR:%.*]], i32 0, metadata [[META12:![0-9]+]]) +// CHECK-GFX1250-NEXT: ret i32 [[TMP0]] +// +int test_amdgcn_global_load_monitor_b32_wavefront(global int* inptr) +{ + return __builtin_amdgcn_global_load_monitor_b32(inptr, __ATOMIC_RELAXED, __MEMORY_SCOPE_WVFRNT); +} + +// CHECK-GFX1250-LABEL: @test_amdgcn_global_load_monitor_b32_single( +// CHECK-GFX1250-NEXT: entry: +// CHECK-GFX1250-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.global.load.monitor.b32.i32(ptr addrspace(1) [[INPTR:%.*]], i32 0, metadata [[META13:![0-9]+]]) +// CHECK-GFX1250-NEXT: ret i32 [[TMP0]] +// +int test_amdgcn_global_load_monitor_b32_single(global int* inptr) +{ + return __builtin_amdgcn_global_load_monitor_b32(inptr, __ATOMIC_RELAXED, __MEMORY_SCOPE_SINGLE); +} + +// CHECK-GFX1250: [[META8]] = !{!""} +// CHECK-GFX1250: [[META9]] = !{!"agent"} +// CHECK-GFX1250: [[META10]] = !{!"workgroup"} +// CHECK-GFX1250: [[META11]] = !{!"cluster"} +// CHECK-GFX1250: [[META12]] = !{!"wavefront"} +// CHECK-GFX1250: [[META13]] = !{!"singlethread"} diff --git a/clang/test/SemaHIP/incorrect-atomic-scope.hip b/clang/test/SemaHIP/incorrect-atomic-scope.hip index 07499b412aa30..1c5aaee710051 100644 --- a/clang/test/SemaHIP/incorrect-atomic-scope.hip +++ b/clang/test/SemaHIP/incorrect-atomic-scope.hip @@ -14,7 +14,7 @@ // // CHECK-LABEL: test_intrinsic_metadata // CHECK: call i32 @llvm.amdgcn.flat.load.monitor{{.*}} metadata [[SCOPE:![0-9]+]] -// CHECK: [[SCOPE]] = !{!"wavefront"} +// CHECK: [[SCOPE]] = !{!"singlethread"} __device__ void test_builtin_rmw(__attribute__((address_space(3))) float *out, float src) { *out = __builtin_amdgcn_ds_fmaxf(out, src, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT, false); >From fea0bd021e8ec3bcac3d996ec3f46fa37a218149 Mon Sep 17 00:00:00 2001 From: Sameer Sahasrabuddhe <[email protected]> Date: Tue, 19 May 2026 11:12:52 +0530 Subject: [PATCH 2/2] [AMDGPU][Clang] Use ScopeModel in ProcessOrderScopeAMDGCN [NFC] --- clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp | 47 +++++++-------------- 1 file changed, 16 insertions(+), 31 deletions(-) diff --git a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp index 1d674803931d2..cb883e8780e59 100644 --- a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp +++ b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp @@ -383,15 +383,24 @@ static llvm::AtomicOrdering mapCABIAtomicOrdering(unsigned AO) { llvm_unreachable("Unknown AtomicOrderingCABI enum"); } +// Map a __MEMORY_SCOPE_* integer constant to the AMDGPU-specific syncscope. +// Invalid scope values are mapped to system scope (empty string). +static StringRef getAMDGPUSyncScopeStr(CodeGenModule &CGM, unsigned ScopeInt, + llvm::AtomicOrdering AO) { + AtomicScopeGenericModel ScopeModel; + if (!ScopeModel.isValid(ScopeInt)) + return ""; + clang::SyncScope Scope = ScopeModel.map(ScopeInt); + return CGM.getTargetCodeGenInfo().getLLVMSyncScopeStr(CGM.getLangOpts(), + Scope, AO); +} + /// Convert a __MEMORY_SCOPE_* integer constant to a metadata node containing /// the target-specific sync scope string. static llvm::MetadataAsValue *emitScopeMD( CodeGenFunction &CGF, unsigned ScopeInt, llvm::AtomicOrdering AO = llvm::AtomicOrdering::SequentiallyConsistent) { - AtomicScopeGenericModel ScopeModel; - clang::SyncScope Scope = ScopeModel.map(ScopeInt); - StringRef ScopeStr = CGF.CGM.getTargetCodeGenInfo().getLLVMSyncScopeStr( - CGF.CGM.getLangOpts(), Scope, AO); + StringRef ScopeStr = getAMDGPUSyncScopeStr(CGF.CGM, ScopeInt, AO); llvm::LLVMContext &Ctx = CGF.CGM.getLLVMContext(); llvm::MDNode *MD = llvm::MDNode::get(Ctx, {llvm::MDString::get(Ctx, ScopeStr)}); @@ -422,33 +431,9 @@ void CodeGenFunction::ProcessOrderScopeAMDGCN(Value *Order, Value *Scope, } // Older builtins had an enum argument for the memory scope. - const char *SSN = nullptr; - int scope = cast<llvm::ConstantInt>(Scope)->getZExtValue(); - switch (scope) { - case AtomicScopeGenericModel::System: // __MEMORY_SCOPE_SYSTEM - SSID = llvm::SyncScope::System; - break; - case AtomicScopeGenericModel::Device: // __MEMORY_SCOPE_DEVICE - SSN = getTarget().getTriple().isSPIRV() ? "device" : "agent"; - break; - case AtomicScopeGenericModel::Workgroup: // __MEMORY_SCOPE_WRKGRP - SSN = "workgroup"; - break; - case AtomicScopeGenericModel::Cluster: // __MEMORY_SCOPE_CLUSTR - SSN = getTarget().getTriple().isSPIRV() ? "workgroup" : "cluster"; - break; - case AtomicScopeGenericModel::Wavefront: // __MEMORY_SCOPE_WVFRNT - SSN = getTarget().getTriple().isSPIRV() ? "subgroup" : "wavefront"; - break; - case AtomicScopeGenericModel::Single: // __MEMORY_SCOPE_SINGLE - SSID = llvm::SyncScope::SingleThread; - break; - default: - SSID = llvm::SyncScope::System; - break; - } - if (SSN) - SSID = getLLVMContext().getOrInsertSyncScopeID(SSN); + unsigned scope = cast<llvm::ConstantInt>(Scope)->getZExtValue(); + StringRef SSN = getAMDGPUSyncScopeStr(CGM, scope, AO); + SSID = getLLVMContext().getOrInsertSyncScopeID(SSN); } void CodeGenFunction::AddAMDGPUFenceAddressSpaceMMRA(llvm::Instruction *Inst, _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
