[llvm] c17a839 - [OMPIRBuilder] Fix addrspace of internal critical section lock (#166459)
via llvm-commits
llvm-commits at lists.llvm.org
Mon Nov 10 07:12:18 PST 2025
Author: Nick Sarnie
Date: 2025-11-10T15:12:13Z
New Revision: c17a839d8335ae75447221adf62f7993e575a913
URL: https://github.com/llvm/llvm-project/commit/c17a839d8335ae75447221adf62f7993e575a913
DIFF: https://github.com/llvm/llvm-project/commit/c17a839d8335ae75447221adf62f7993e575a913.diff
LOG: [OMPIRBuilder] Fix addrspace of internal critical section lock (#166459)
First, for internal variables, they are always global, so use the global
AS by default unless specified otherwise. We can't really use `0` as a
default like we do now because that has an actual meaning on some
targets, so we really need specified vs unspecified, so I used
`std::optional` which is already used in many places in OMPIRBuilder.
Second, for the critical lock variable, add an addrspace cast if needed.
Signed-off-by: Nick Sarnie <nick.sarnie at intel.com>
Added:
Modified:
clang/lib/CodeGen/CGOpenMPRuntime.cpp
clang/test/OpenMP/force-usm.c
clang/test/OpenMP/spirv_target_codegen_basic.cpp
llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
Removed:
################################################################################
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 121de42248e3b..44ba72c5c76c7 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -2000,22 +2000,29 @@ void CGOpenMPRuntime::emitCriticalRegion(CodeGenFunction &CGF,
// Prepare arguments and build a call to __kmpc_critical
if (!CGF.HaveInsertPoint())
return;
+ llvm::FunctionCallee RuntimeFcn = OMPBuilder.getOrCreateRuntimeFunction(
+ CGM.getModule(),
+ Hint ? OMPRTL___kmpc_critical_with_hint : OMPRTL___kmpc_critical);
+ llvm::Value *LockVar = getCriticalRegionLock(CriticalName);
+ unsigned LockVarArgIdx = 2;
+ if (cast<llvm::GlobalVariable>(LockVar)->getAddressSpace() !=
+ RuntimeFcn.getFunctionType()
+ ->getParamType(LockVarArgIdx)
+ ->getPointerAddressSpace())
+ LockVar = CGF.Builder.CreateAddrSpaceCast(
+ LockVar, RuntimeFcn.getFunctionType()->getParamType(LockVarArgIdx));
llvm::Value *Args[] = {emitUpdateLocation(CGF, Loc), getThreadID(CGF, Loc),
- getCriticalRegionLock(CriticalName)};
+ LockVar};
llvm::SmallVector<llvm::Value *, 4> EnterArgs(std::begin(Args),
std::end(Args));
if (Hint) {
EnterArgs.push_back(CGF.Builder.CreateIntCast(
CGF.EmitScalarExpr(Hint), CGM.Int32Ty, /*isSigned=*/false));
}
- CommonActionTy Action(
- OMPBuilder.getOrCreateRuntimeFunction(
- CGM.getModule(),
- Hint ? OMPRTL___kmpc_critical_with_hint : OMPRTL___kmpc_critical),
- EnterArgs,
- OMPBuilder.getOrCreateRuntimeFunction(CGM.getModule(),
- OMPRTL___kmpc_end_critical),
- Args);
+ CommonActionTy Action(RuntimeFcn, EnterArgs,
+ OMPBuilder.getOrCreateRuntimeFunction(
+ CGM.getModule(), OMPRTL___kmpc_end_critical),
+ Args);
CriticalOpGen.setAction(Action);
emitInlinedDirective(CGF, OMPD_critical, CriticalOpGen);
}
diff --git a/clang/test/OpenMP/force-usm.c b/clang/test/OpenMP/force-usm.c
index 5c63a9a5e7004..45c0e28b525da 100644
--- a/clang/test/OpenMP/force-usm.c
+++ b/clang/test/OpenMP/force-usm.c
@@ -46,7 +46,7 @@ int main(void) {
// CHECK-USM-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
// CHECK-USM: user_code.entry:
// CHECK-USM-NEXT: store i32 1, ptr [[TMP0]], align 4
-// CHECK-USM-NEXT: [[TMP2:%.*]] = load ptr, ptr @pGI_decl_tgt_ref_ptr, align 8
+// CHECK-USM-NEXT: [[TMP2:%.*]] = load ptr, ptr addrspace(1) @pGI_decl_tgt_ref_ptr, align 8
// CHECK-USM-NEXT: [[TMP3:%.*]] = load ptr, ptr [[TMP2]], align 8
// CHECK-USM-NEXT: store i32 2, ptr [[TMP3]], align 4
// CHECK-USM-NEXT: call void @__kmpc_target_deinit()
diff --git a/clang/test/OpenMP/spirv_target_codegen_basic.cpp b/clang/test/OpenMP/spirv_target_codegen_basic.cpp
index fb2810e88c063..6e029fb93644d 100644
--- a/clang/test/OpenMP/spirv_target_codegen_basic.cpp
+++ b/clang/test/OpenMP/spirv_target_codegen_basic.cpp
@@ -6,12 +6,18 @@
// CHECK: @__omp_offloading_{{.*}}_dynamic_environment = weak_odr protected addrspace(1) global %struct.DynamicEnvironmentTy zeroinitializer
// CHECK: @__omp_offloading_{{.*}}_kernel_environment = weak_odr protected addrspace(1) constant %struct.KernelEnvironmentTy
+// CHECK: @"_gomp_critical_user_$var" = common addrspace(1) global [8 x i32] zeroinitializer, align 8
+
// CHECK: define weak_odr protected spir_kernel void @__omp_offloading_{{.*}}
+// CHECK: call spir_func addrspace(9) void @__kmpc_critical(ptr addrspace(4) addrspacecast (ptr addrspace(1) @{{.*}} to ptr addrspace(4)), i32 %{{.*}}, ptr addrspace(4) addrspacecast (ptr addrspace(1) @"_gomp_critical_user_$var" to ptr addrspace(4)))
+// CHECK: call spir_func addrspace(9) void @__kmpc_end_critical(ptr addrspace(4) addrspacecast (ptr addrspace(1) @{{.*}} to ptr addrspace(4)), i32 %{{.*}}, ptr addrspace(4) addrspacecast (ptr addrspace(1) @"_gomp_critical_user_$var" to ptr addrspace(4)))
+
int main() {
int ret = 0;
#pragma omp target
for(int i = 0; i < 5; i++)
+ #pragma omp critical
ret++;
return ret;
}
diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
index b3d7ab4acf303..fd6b9729658c1 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
+++ b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
@@ -3654,7 +3654,7 @@ class OpenMPIRBuilder {
/// \param Name Name of the variable.
LLVM_ABI GlobalVariable *
getOrCreateInternalVariable(Type *Ty, const StringRef &Name,
- unsigned AddressSpace = 0);
+ std::optional<unsigned> AddressSpace = {});
};
/// Class to represented the control flow structure of an OpenMP canonical loop.
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index fff9a815e5368..7dc32fda0eed6 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -8460,9 +8460,8 @@ OpenMPIRBuilder::createPlatformSpecificName(ArrayRef<StringRef> Parts) const {
Config.separator());
}
-GlobalVariable *
-OpenMPIRBuilder::getOrCreateInternalVariable(Type *Ty, const StringRef &Name,
- unsigned AddressSpace) {
+GlobalVariable *OpenMPIRBuilder::getOrCreateInternalVariable(
+ Type *Ty, const StringRef &Name, std::optional<unsigned> AddressSpace) {
auto &Elem = *InternalVars.try_emplace(Name, nullptr).first;
if (Elem.second) {
assert(Elem.second->getValueType() == Ty &&
@@ -8472,16 +8471,18 @@ OpenMPIRBuilder::getOrCreateInternalVariable(Type *Ty, const StringRef &Name,
// variable for possibly changing that to internal or private, or maybe
// create
diff erent versions of the function for
diff erent OMP internal
// variables.
+ const DataLayout &DL = M.getDataLayout();
+ unsigned AddressSpaceVal =
+ AddressSpace ? *AddressSpace : DL.getDefaultGlobalsAddressSpace();
auto Linkage = this->M.getTargetTriple().getArch() == Triple::wasm32
? GlobalValue::InternalLinkage
: GlobalValue::CommonLinkage;
auto *GV = new GlobalVariable(M, Ty, /*IsConstant=*/false, Linkage,
Constant::getNullValue(Ty), Elem.first(),
/*InsertBefore=*/nullptr,
- GlobalValue::NotThreadLocal, AddressSpace);
- const DataLayout &DL = M.getDataLayout();
+ GlobalValue::NotThreadLocal, AddressSpaceVal);
const llvm::Align TypeAlign = DL.getABITypeAlign(Ty);
- const llvm::Align PtrAlign = DL.getPointerABIAlignment(AddressSpace);
+ const llvm::Align PtrAlign = DL.getPointerABIAlignment(AddressSpaceVal);
GV->setAlignment(std::max(TypeAlign, PtrAlign));
Elem.second = GV;
}
More information about the llvm-commits
mailing list