[clang] 6255ecd - [AMDGPU][Clang] use ScopeModel to translate integer scopes [NFC] (#198250)
via cfe-commits
cfe-commits at lists.llvm.org
Thu May 21 04:33:51 PDT 2026
Author: Sameer Sahasrabuddhe
Date: 2026-05-21T17:03:45+05:30
New Revision: 6255ecd593f488bd99123771e9803ad2f1be49c5
URL: https://github.com/llvm/llvm-project/commit/6255ecd593f488bd99123771e9803ad2f1be49c5
DIFF: https://github.com/llvm/llvm-project/commit/6255ecd593f488bd99123771e9803ad2f1be49c5.diff
LOG: [AMDGPU][Clang] use ScopeModel to translate integer scopes [NFC] (#198250)
The assumption here is that AMDGPU builtins (typically suffixed with
`__builtin_amdgcn`) use the `__MEMORY_SCOPE_*` enumeration, and not the
`__HIP_MEMORY_SCOPE_*` enumeration (which is how it should be).
Assisted-By: Claude Opus 4.6
Added:
Modified:
clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250-load-monitor.cl
clang/test/SemaHIP/incorrect-atomic-scope.hip
Removed:
################################################################################
diff --git a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
index 751cd9847bd31..cb883e8780e59 100644
--- a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
+++ b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
@@ -383,6 +383,30 @@ 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) {
+ 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)});
+ 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
@@ -407,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,
@@ -927,22 +927,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);
More information about the cfe-commits
mailing list