[OMPIRBuilder] Fix addrspace of internal critical section lock #166459
+31
−17
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Signed-off-by: Nick Sarnie <nick.sarnie@intel.com>
llvmbot
added
clang
Member
|
@llvm/pr-subscribers-clang-codegen @llvm/pr-subscribers-clang Author: Nick Sarnie (sarnex) ChangesFirst, for internal variables, they are always global, so use the global AS by default unless specified otherwise. We can't really use Second, for the critical lock variable, add an addrspace cast if needed. Full diff: https://github.com/llvm/llvm-project/pull/166459.diff 5 Files Affected:
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 different versions of the function for different 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;
}
|
Member
|
@llvm/pr-subscribers-flang-openmp Author: Nick Sarnie (sarnex) ChangesFirst, for internal variables, they are always global, so use the global AS by default unless specified otherwise. We can't really use Second, for the critical lock variable, add an addrspace cast if needed. Full diff: https://github.com/llvm/llvm-project/pull/166459.diff 5 Files Affected:
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 different versions of the function for different 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;
}
|
tblah
approved these changes
Nov 10, 2025
Contributor
tblah
left a comment
tblah
left a comment
There was a problem hiding this comment.
LGTM. Thanks for your patch!
Member
Author
|
Thanks Tom! |
Contributor
|
Hi, I think this patch broke our flang OpenMP builder: https://lab.llvm.org/staging/#/builders/105/builds/36206 |
Member
Author
|
Thanks, I'll restrict the target for now to fix it, seems AMDGPU doesn't expect the global AS for globals in this case. |
ronlieb
added a commit
to ROCm/llvm-project
that referenced
this pull request
Nov 10, 2025
llvm#166459)" This reverts commit c17a839.
Kewen12
pushed a commit
that referenced
this pull request
Nov 11, 2025
We see some libomptarget test failures if we use the default global AS. See #166459 for more info. Signed-off-by: Nick Sarnie <nick.sarnie@intel.com>
ronlieb
pushed a commit
to ROCm/llvm-project
that referenced
this pull request
Nov 11, 2025
…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@intel.com>
llvm-sync bot
pushed a commit
to arm/arm-toolchain
that referenced
this pull request
Nov 11, 2025
…(#167377) We see some libomptarget test failures if we use the default global AS. See llvm/llvm-project#166459 for more info. Signed-off-by: Nick Sarnie <nick.sarnie@intel.com>
ronlieb
pushed a commit
to ROCm/llvm-project
that referenced
this pull request
Nov 12, 2025
…llvm#166459) Update revert_patches.txt
ronlieb
added a commit
to ROCm/llvm-project
that referenced
this pull request
Nov 13, 2025
#573) … (llvm#166459) Update revert_patches.txt Co-authored-by: Nick Sarnie <nick.sarnie@intel.com>
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
4 participants
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
First, for internal variables, they are always global, so use the global AS by default unless specified otherwise. We can't really use
0as a default like we do now because that has an actual meaning on some targets, so we really need specified vs unspecified, so I usedstd::optionalwhich is already used in many places in OMPIRBuilder.Second, for the critical lock variable, add an addrspace cast if needed.