maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

[openmp] Add addrspacecast to getOrCreateIdent 

Fixes 51982. Adds a missing CreatePointerCast and allocates a global in
the correct address space.

Test case derived from https://github.com/ROCm-Developer-Tools/aomp/\
blob/aomp-dev/test/smoke/nest_call_par2/nest_call_par2.c by deleting
parts while checking the assertion failure still occurred.

Reviewed By: jdoerfert

Differential Revision: https://reviews.llvm.org/D110556
This commit is contained in:
Jon Chesterfield 2021-09-30 21:36:30 +01:00
parent b75a7481ba
commit 3247329107
2 changed files with 36 additions and 7 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

19
llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
View File

@ -276,15 +276,20 @@ Value *OpenMPIRBuilder::getOrCreateIdent(Constant *SrcLocStr,
for (GlobalVariable &GV : M.getGlobalList())
if (GV.getValueType() == OpenMPIRBuilder::Ident && GV.hasInitializer())
if (GV.getInitializer() == Initializer)
return Ident = &GV;
Ident = &GV;
auto *GV = new GlobalVariable(M, OpenMPIRBuilder::Ident,
/* isConstant = */ true,
GlobalValue::PrivateLinkage, Initializer);
GV->setUnnamedAddr(GlobalValue::UnnamedAddr::Global);
GV->setAlignment(Align(8));
Ident = GV;
if (!Ident) {
auto *GV = new GlobalVariable(
M, OpenMPIRBuilder::Ident,
/* isConstant = */ true, GlobalValue::PrivateLinkage, Initializer, "",
nullptr, GlobalValue::NotThreadLocal,
M.getDataLayout().getDefaultGlobalsAddressSpace());
GV->setUnnamedAddr(GlobalValue::UnnamedAddr::Global);
GV->setAlignment(Align(8));
Ident = GV;
}
}
return Builder.CreatePointerCast(Ident, IdentPtr);
}

24
openmp/libomptarget/test/offloading/bug51982.c Normal file
View File

@ -0,0 +1,24 @@
// RUN: %libomptarget-compile-generic -O1 && %libomptarget-run-generic
// -O1 to run openmp-opt
int main(void) {
long int aa = 0;
int ng = 12;
int nxyz = 5;
const long exp = ng * nxyz;
#pragma omp target map(tofrom : aa)
for (int gid = 0; gid < nxyz; gid++) {
#pragma omp parallel for
for (unsigned int g = 0; g < ng; g++) {
#pragma omp atomic
aa += 1;
}
}
if (aa != exp) {
return 1;
}
return 0;
}