Skip to content

Commit 3663563

Browse files
alexey-bataevalexey-bataev
committed
[OPENMP]Fix PR44578: crash in target construct with captured global.
Target regions have implicit outer region which may erroneously capture some globals when it should not. It may lead to a compiler crash at the compile time.
1 parent 4c9d691 commit 3663563

14 files changed

+58
-27
lines changed

‎clang/include/clang/Sema/Sema.h‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -9681,7 +9681,8 @@ class Sema final {
96819681
/// Check if the specified variable is captured by 'target' directive.
96829682
/// \param Level Relative level of nested OpenMP construct for that the check
96839683
/// is performed.
9684-
bool isOpenMPTargetCapturedDecl(const ValueDecl *D, unsigned Level) const;
9684+
bool isOpenMPTargetCapturedDecl(const ValueDecl *D, unsigned Level,
9685+
unsigned CaptureLevel) const;
96859686

96869687
ExprResult PerformOpenMPImplicitIntegerConversion(SourceLocation OpLoc,
96879688
Expr *Op);

‎clang/lib/Sema/SemaExpr.cpp‎

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -16280,8 +16280,10 @@ bool Sema::tryCaptureVariable(
1628016280
captureVariablyModifiedType(Context, QTy, OuterRSI);
1628116281
}
1628216282
}
16283-
bool IsTargetCap = !IsOpenMPPrivateDecl &&
16284-
isOpenMPTargetCapturedDecl(Var, RSI->OpenMPLevel);
16283+
bool IsTargetCap =
16284+
!IsOpenMPPrivateDecl &&
16285+
isOpenMPTargetCapturedDecl(Var, RSI->OpenMPLevel,
16286+
RSI->OpenMPCaptureLevel);
1628516287
// When we detect target captures we are looking from inside the
1628616288
// target region, therefore we need to propagate the capture from the
1628716289
// enclosing region. Therefore, the capture is not initially nested.

‎clang/lib/Sema/SemaOpenMP.cpp‎

Lines changed: 24 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -31,6 +31,7 @@
3131
#include "clang/Sema/SemaInternal.h"
3232
#include "llvm/ADT/IndexedMap.h"
3333
#include "llvm/ADT/PointerEmbeddedInt.h"
34+
#include "llvm/ADT/STLExtras.h"
3435
#include "llvm/Frontend/OpenMP/OMPConstants.h"
3536
using namespace clang;
3637
using namespace llvm::omp;
@@ -2010,7 +2011,23 @@ VarDecl *Sema::isOpenMPCapturedDecl(ValueDecl *D, bool CheckScopeInfo,
20102011
//
20112012
if (OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD))
20122013
return nullptr;
2013-
return VD;
2014+
CapturedRegionScopeInfo *CSI = nullptr;
2015+
for (FunctionScopeInfo *FSI : llvm::drop_begin(
2016+
llvm::reverse(FunctionScopes),
2017+
CheckScopeInfo ? (FunctionScopes.size() - (StopAt + 1)) : 0)) {
2018+
if (!isa<CapturingScopeInfo>(FSI))
2019+
return nullptr;
2020+
if (auto *RSI = dyn_cast<CapturedRegionScopeInfo>(FSI))
2021+
if (RSI->CapRegionKind == CR_OpenMP) {
2022+
CSI = RSI;
2023+
break;
2024+
}
2025+
}
2026+
SmallVector<OpenMPDirectiveKind, 4> Regions;
2027+
getOpenMPCaptureRegions(Regions,
2028+
DSAStack->getDirective(CSI->OpenMPLevel));
2029+
if (Regions[CSI->OpenMPCaptureLevel] != OMPD_task)
2030+
return VD;
20142031
}
20152032
}
20162033

@@ -2151,15 +2168,18 @@ void Sema::setOpenMPCaptureKind(FieldDecl *FD, const ValueDecl *D,
21512168
FD->addAttr(OMPCaptureKindAttr::CreateImplicit(Context, OMPC));
21522169
}
21532170

2154-
bool Sema::isOpenMPTargetCapturedDecl(const ValueDecl *D,
2155-
unsigned Level) const {
2171+
bool Sema::isOpenMPTargetCapturedDecl(const ValueDecl *D, unsigned Level,
2172+
unsigned CaptureLevel) const {
21562173
assert(LangOpts.OpenMP && "OpenMP is not allowed");
21572174
// Return true if the current level is no longer enclosed in a target region.
21582175

2176+
SmallVector<OpenMPDirectiveKind, 4> Regions;
2177+
getOpenMPCaptureRegions(Regions, DSAStack->getDirective(Level));
21592178
const auto *VD = dyn_cast<VarDecl>(D);
21602179
return VD && !VD->hasLocalStorage() &&
21612180
DSAStack->hasExplicitDirective(isOpenMPTargetExecutionDirective,
2162-
Level);
2181+
Level) &&
2182+
Regions[CaptureLevel] != OMPD_task;
21632183
}
21642184

21652185
void Sema::DestroyDataSharingAttributesStack() { delete DSAStack; }

‎clang/test/OpenMP/target_depend_codegen.cpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -209,7 +209,7 @@ int foo(int n) {
209209
// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
210210
// CHECK: [[FAIL]]
211211
// CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
212-
// CHECK: [[BP1_I32:%.+]] = load i32, i32* %
212+
// CHECK: [[BP1_I32:%.+]] = load i32, i32* @
213213
// CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
214214
// CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]],
215215
// CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],
@@ -223,7 +223,7 @@ int foo(int n) {
223223
// CHECK: call void (i8*, ...) %
224224
// CHECK: [[DEVICE_CAP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2
225225
// CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
226-
// CHECK: [[BP1_I32:%.+]] = load i32, i32* %
226+
// CHECK: [[BP1_I32:%.+]] = load i32, i32* @
227227
// CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
228228
// CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]],
229229
// CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],

‎clang/test/OpenMP/target_messages.cpp‎

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -112,4 +112,12 @@ int main(int argc, char **argv) {
112112

113113
return 0;
114114
}
115+
116+
template <class> struct a { static bool b; };
117+
template <class c, bool = a<c>::b> void e(c) { // expected-note {{candidate template ignored: substitution failure [with c = int]: non-type template argument is not a constant expression}}
118+
#pragma omp target
119+
{
120+
int d ; e(d); // expected-error {{no matching function for call to 'e'}}
121+
}
122+
}
115123
#endif

‎clang/test/OpenMP/target_parallel_depend_codegen.cpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -209,7 +209,7 @@ int foo(int n) {
209209
// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
210210
// CHECK: [[FAIL]]
211211
// CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
212-
// CHECK: [[BP1_I32:%.+]] = load i32, i32* %
212+
// CHECK: [[BP1_I32:%.+]] = load i32, i32* @
213213
// CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
214214
// CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]],
215215
// CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],
@@ -223,7 +223,7 @@ int foo(int n) {
223223
// CHECK: call void (i8*, ...) %
224224
// CHECK: [[DEVICE_CAP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2
225225
// CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
226-
// CHECK: [[BP1_I32:%.+]] = load i32, i32* %
226+
// CHECK: [[BP1_I32:%.+]] = load i32, i32* @
227227
// CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
228228
// CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]],
229229
// CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],

‎clang/test/OpenMP/target_parallel_for_depend_codegen.cpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -209,7 +209,7 @@ int foo(int n) {
209209
// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
210210
// CHECK: [[FAIL]]
211211
// CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
212-
// CHECK: [[BP1_I32:%.+]] = load i32, i32* %
212+
// CHECK: [[BP1_I32:%.+]] = load i32, i32* @
213213
// CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
214214
// CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]],
215215
// CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],
@@ -223,7 +223,7 @@ int foo(int n) {
223223
// CHECK: call void (i8*, ...) %
224224
// CHECK: [[DEVICE_CAP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2
225225
// CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
226-
// CHECK: [[BP1_I32:%.+]] = load i32, i32* %
226+
// CHECK: [[BP1_I32:%.+]] = load i32, i32* @
227227
// CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
228228
// CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]],
229229
// CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],

‎clang/test/OpenMP/target_parallel_for_simd_depend_codegen.cpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -209,7 +209,7 @@ int foo(int n) {
209209
// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
210210
// CHECK: [[FAIL]]
211211
// CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
212-
// CHECK: [[BP1_I32:%.+]] = load i32, i32* %
212+
// CHECK: [[BP1_I32:%.+]] = load i32, i32* @
213213
// CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
214214
// CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]],
215215
// CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],
@@ -223,7 +223,7 @@ int foo(int n) {
223223
// CHECK: call void (i8*, ...) %
224224
// CHECK: [[DEVICE_CAP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2
225225
// CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
226-
// CHECK: [[BP1_I32:%.+]] = load i32, i32* %
226+
// CHECK: [[BP1_I32:%.+]] = load i32, i32* @
227227
// CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
228228
// CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]],
229229
// CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],

‎clang/test/OpenMP/target_simd_depend_codegen.cpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -209,7 +209,7 @@ int foo(int n) {
209209
// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
210210
// CHECK: [[FAIL]]
211211
// CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
212-
// CHECK: [[BP1_I32:%.+]] = load i32, i32* %
212+
// CHECK: [[BP1_I32:%.+]] = load i32, i32* @
213213
// CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
214214
// CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]],
215215
// CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],
@@ -223,7 +223,7 @@ int foo(int n) {
223223
// CHECK: call void (i8*, ...) %
224224
// CHECK: [[DEVICE_CAP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2
225225
// CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
226-
// CHECK: [[BP1_I32:%.+]] = load i32, i32* %
226+
// CHECK: [[BP1_I32:%.+]] = load i32, i32* @
227227
// CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
228228
// CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]],
229229
// CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],

‎clang/test/OpenMP/target_teams_depend_codegen.cpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -209,7 +209,7 @@ int foo(int n) {
209209
// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
210210
// CHECK: [[FAIL]]
211211
// CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
212-
// CHECK: [[BP1_I32:%.+]] = load i32, i32* %
212+
// CHECK: [[BP1_I32:%.+]] = load i32, i32* @
213213
// CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
214214
// CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]],
215215
// CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],
@@ -223,7 +223,7 @@ int foo(int n) {
223223
// CHECK: call void (i8*, ...) %
224224
// CHECK: [[DEVICE_CAP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2
225225
// CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
226-
// CHECK: [[BP1_I32:%.+]] = load i32, i32* %
226+
// CHECK: [[BP1_I32:%.+]] = load i32, i32* @
227227
// CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
228228
// CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]],
229229
// CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],

0 commit comments

Comments
 (0)