[clang] 3663563 - [OPENMP]Fix PR44578: crash in target construct with captured global.
Alexey Bataev via cfe-commits
cfe-commits at lists.llvm.org
Mon Jan 20 08:13:50 PST 2020
Author: Alexey Bataev
Date: 2020-01-20T11:10:17-05:00
New Revision: 366356361cf3abf97fbcfe9a0467b1ed1610401f
URL: https://github.com/llvm/llvm-project/commit/366356361cf3abf97fbcfe9a0467b1ed1610401f
DIFF: https://github.com/llvm/llvm-project/commit/366356361cf3abf97fbcfe9a0467b1ed1610401f.diff
LOG: [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.
Added:
Modified:
clang/include/clang/Sema/Sema.h
clang/lib/Sema/SemaExpr.cpp
clang/lib/Sema/SemaOpenMP.cpp
clang/test/OpenMP/target_depend_codegen.cpp
clang/test/OpenMP/target_messages.cpp
clang/test/OpenMP/target_parallel_depend_codegen.cpp
clang/test/OpenMP/target_parallel_for_depend_codegen.cpp
clang/test/OpenMP/target_parallel_for_simd_depend_codegen.cpp
clang/test/OpenMP/target_simd_depend_codegen.cpp
clang/test/OpenMP/target_teams_depend_codegen.cpp
clang/test/OpenMP/target_teams_distribute_depend_codegen.cpp
clang/test/OpenMP/target_teams_distribute_parallel_for_depend_codegen.cpp
clang/test/OpenMP/target_teams_distribute_parallel_for_simd_depend_codegen.cpp
clang/test/OpenMP/target_teams_distribute_simd_depend_codegen.cpp
Removed:
################################################################################
diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index 2fb758bc75d6..6e3ae96b3afc 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -9681,7 +9681,8 @@ class Sema final {
/// Check if the specified variable is captured by 'target' directive.
/// \param Level Relative level of nested OpenMP construct for that the check
/// is performed.
- bool isOpenMPTargetCapturedDecl(const ValueDecl *D, unsigned Level) const;
+ bool isOpenMPTargetCapturedDecl(const ValueDecl *D, unsigned Level,
+ unsigned CaptureLevel) const;
ExprResult PerformOpenMPImplicitIntegerConversion(SourceLocation OpLoc,
Expr *Op);
diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp
index ea4b93ee6a5a..67b68ffb5eda 100644
--- a/clang/lib/Sema/SemaExpr.cpp
+++ b/clang/lib/Sema/SemaExpr.cpp
@@ -16280,8 +16280,10 @@ bool Sema::tryCaptureVariable(
captureVariablyModifiedType(Context, QTy, OuterRSI);
}
}
- bool IsTargetCap = !IsOpenMPPrivateDecl &&
- isOpenMPTargetCapturedDecl(Var, RSI->OpenMPLevel);
+ bool IsTargetCap =
+ !IsOpenMPPrivateDecl &&
+ isOpenMPTargetCapturedDecl(Var, RSI->OpenMPLevel,
+ RSI->OpenMPCaptureLevel);
// When we detect target captures we are looking from inside the
// target region, therefore we need to propagate the capture from the
// enclosing region. Therefore, the capture is not initially nested.
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index 3fce0e27e9b3..72afe63749ac 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -31,6 +31,7 @@
#include "clang/Sema/SemaInternal.h"
#include "llvm/ADT/IndexedMap.h"
#include "llvm/ADT/PointerEmbeddedInt.h"
+#include "llvm/ADT/STLExtras.h"
#include "llvm/Frontend/OpenMP/OMPConstants.h"
using namespace clang;
using namespace llvm::omp;
@@ -2010,7 +2011,23 @@ VarDecl *Sema::isOpenMPCapturedDecl(ValueDecl *D, bool CheckScopeInfo,
//
if (OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD))
return nullptr;
- return VD;
+ CapturedRegionScopeInfo *CSI = nullptr;
+ for (FunctionScopeInfo *FSI : llvm::drop_begin(
+ llvm::reverse(FunctionScopes),
+ CheckScopeInfo ? (FunctionScopes.size() - (StopAt + 1)) : 0)) {
+ if (!isa<CapturingScopeInfo>(FSI))
+ return nullptr;
+ if (auto *RSI = dyn_cast<CapturedRegionScopeInfo>(FSI))
+ if (RSI->CapRegionKind == CR_OpenMP) {
+ CSI = RSI;
+ break;
+ }
+ }
+ SmallVector<OpenMPDirectiveKind, 4> Regions;
+ getOpenMPCaptureRegions(Regions,
+ DSAStack->getDirective(CSI->OpenMPLevel));
+ if (Regions[CSI->OpenMPCaptureLevel] != OMPD_task)
+ return VD;
}
}
@@ -2151,15 +2168,18 @@ void Sema::setOpenMPCaptureKind(FieldDecl *FD, const ValueDecl *D,
FD->addAttr(OMPCaptureKindAttr::CreateImplicit(Context, OMPC));
}
-bool Sema::isOpenMPTargetCapturedDecl(const ValueDecl *D,
- unsigned Level) const {
+bool Sema::isOpenMPTargetCapturedDecl(const ValueDecl *D, unsigned Level,
+ unsigned CaptureLevel) const {
assert(LangOpts.OpenMP && "OpenMP is not allowed");
// Return true if the current level is no longer enclosed in a target region.
+ SmallVector<OpenMPDirectiveKind, 4> Regions;
+ getOpenMPCaptureRegions(Regions, DSAStack->getDirective(Level));
const auto *VD = dyn_cast<VarDecl>(D);
return VD && !VD->hasLocalStorage() &&
DSAStack->hasExplicitDirective(isOpenMPTargetExecutionDirective,
- Level);
+ Level) &&
+ Regions[CaptureLevel] != OMPD_task;
}
void Sema::DestroyDataSharingAttributesStack() { delete DSAStack; }
diff --git a/clang/test/OpenMP/target_depend_codegen.cpp b/clang/test/OpenMP/target_depend_codegen.cpp
index e2810f946b04..e5473dad8055 100644
--- a/clang/test/OpenMP/target_depend_codegen.cpp
+++ b/clang/test/OpenMP/target_depend_codegen.cpp
@@ -209,7 +209,7 @@ int foo(int n) {
// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
// CHECK: [[FAIL]]
// CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
-// CHECK: [[BP1_I32:%.+]] = load i32, i32* %
+// CHECK: [[BP1_I32:%.+]] = load i32, i32* @
// CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
// CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]],
// CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],
@@ -223,7 +223,7 @@ int foo(int n) {
// CHECK: call void (i8*, ...) %
// CHECK: [[DEVICE_CAP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2
// CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
-// CHECK: [[BP1_I32:%.+]] = load i32, i32* %
+// CHECK: [[BP1_I32:%.+]] = load i32, i32* @
// CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
// CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]],
// CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],
diff --git a/clang/test/OpenMP/target_messages.cpp b/clang/test/OpenMP/target_messages.cpp
index 5f22fbcaf606..ad04e9306cb0 100644
--- a/clang/test/OpenMP/target_messages.cpp
+++ b/clang/test/OpenMP/target_messages.cpp
@@ -112,4 +112,12 @@ int main(int argc, char **argv) {
return 0;
}
+
+template <class> struct a { static bool b; };
+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}}
+#pragma omp target
+ {
+ int d ; e(d); // expected-error {{no matching function for call to 'e'}}
+ }
+}
#endif
diff --git a/clang/test/OpenMP/target_parallel_depend_codegen.cpp b/clang/test/OpenMP/target_parallel_depend_codegen.cpp
index dceb585b14c4..de7e37efc484 100644
--- a/clang/test/OpenMP/target_parallel_depend_codegen.cpp
+++ b/clang/test/OpenMP/target_parallel_depend_codegen.cpp
@@ -209,7 +209,7 @@ int foo(int n) {
// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
// CHECK: [[FAIL]]
// CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
-// CHECK: [[BP1_I32:%.+]] = load i32, i32* %
+// CHECK: [[BP1_I32:%.+]] = load i32, i32* @
// CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
// CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]],
// CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],
@@ -223,7 +223,7 @@ int foo(int n) {
// CHECK: call void (i8*, ...) %
// CHECK: [[DEVICE_CAP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2
// CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
-// CHECK: [[BP1_I32:%.+]] = load i32, i32* %
+// CHECK: [[BP1_I32:%.+]] = load i32, i32* @
// CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
// CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]],
// CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],
diff --git a/clang/test/OpenMP/target_parallel_for_depend_codegen.cpp b/clang/test/OpenMP/target_parallel_for_depend_codegen.cpp
index 6b2325592c6a..73a5cb1a8beb 100644
--- a/clang/test/OpenMP/target_parallel_for_depend_codegen.cpp
+++ b/clang/test/OpenMP/target_parallel_for_depend_codegen.cpp
@@ -209,7 +209,7 @@ int foo(int n) {
// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
// CHECK: [[FAIL]]
// CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
-// CHECK: [[BP1_I32:%.+]] = load i32, i32* %
+// CHECK: [[BP1_I32:%.+]] = load i32, i32* @
// CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
// CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]],
// CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],
@@ -223,7 +223,7 @@ int foo(int n) {
// CHECK: call void (i8*, ...) %
// CHECK: [[DEVICE_CAP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2
// CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
-// CHECK: [[BP1_I32:%.+]] = load i32, i32* %
+// CHECK: [[BP1_I32:%.+]] = load i32, i32* @
// CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
// CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]],
// CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],
diff --git a/clang/test/OpenMP/target_parallel_for_simd_depend_codegen.cpp b/clang/test/OpenMP/target_parallel_for_simd_depend_codegen.cpp
index 74ff678316d7..6b9950428dba 100644
--- a/clang/test/OpenMP/target_parallel_for_simd_depend_codegen.cpp
+++ b/clang/test/OpenMP/target_parallel_for_simd_depend_codegen.cpp
@@ -209,7 +209,7 @@ int foo(int n) {
// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
// CHECK: [[FAIL]]
// CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
-// CHECK: [[BP1_I32:%.+]] = load i32, i32* %
+// CHECK: [[BP1_I32:%.+]] = load i32, i32* @
// CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
// CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]],
// CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],
@@ -223,7 +223,7 @@ int foo(int n) {
// CHECK: call void (i8*, ...) %
// CHECK: [[DEVICE_CAP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2
// CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
-// CHECK: [[BP1_I32:%.+]] = load i32, i32* %
+// CHECK: [[BP1_I32:%.+]] = load i32, i32* @
// CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
// CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]],
// CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],
diff --git a/clang/test/OpenMP/target_simd_depend_codegen.cpp b/clang/test/OpenMP/target_simd_depend_codegen.cpp
index 0fb75b0b7fc4..0c28c43dc734 100644
--- a/clang/test/OpenMP/target_simd_depend_codegen.cpp
+++ b/clang/test/OpenMP/target_simd_depend_codegen.cpp
@@ -209,7 +209,7 @@ int foo(int n) {
// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
// CHECK: [[FAIL]]
// CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
-// CHECK: [[BP1_I32:%.+]] = load i32, i32* %
+// CHECK: [[BP1_I32:%.+]] = load i32, i32* @
// CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
// CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]],
// CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],
@@ -223,7 +223,7 @@ int foo(int n) {
// CHECK: call void (i8*, ...) %
// CHECK: [[DEVICE_CAP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2
// CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
-// CHECK: [[BP1_I32:%.+]] = load i32, i32* %
+// CHECK: [[BP1_I32:%.+]] = load i32, i32* @
// CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
// CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]],
// CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],
diff --git a/clang/test/OpenMP/target_teams_depend_codegen.cpp b/clang/test/OpenMP/target_teams_depend_codegen.cpp
index 85fc5e297ce8..4cc402c89344 100644
--- a/clang/test/OpenMP/target_teams_depend_codegen.cpp
+++ b/clang/test/OpenMP/target_teams_depend_codegen.cpp
@@ -209,7 +209,7 @@ int foo(int n) {
// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
// CHECK: [[FAIL]]
// CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
-// CHECK: [[BP1_I32:%.+]] = load i32, i32* %
+// CHECK: [[BP1_I32:%.+]] = load i32, i32* @
// CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
// CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]],
// CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],
@@ -223,7 +223,7 @@ int foo(int n) {
// CHECK: call void (i8*, ...) %
// CHECK: [[DEVICE_CAP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2
// CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
-// CHECK: [[BP1_I32:%.+]] = load i32, i32* %
+// CHECK: [[BP1_I32:%.+]] = load i32, i32* @
// CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
// CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]],
// CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],
diff --git a/clang/test/OpenMP/target_teams_distribute_depend_codegen.cpp b/clang/test/OpenMP/target_teams_distribute_depend_codegen.cpp
index 37b80b097224..c8694761bf7b 100644
--- a/clang/test/OpenMP/target_teams_distribute_depend_codegen.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_depend_codegen.cpp
@@ -209,7 +209,7 @@ int foo(int n) {
// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
// CHECK: [[FAIL]]
// CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
-// CHECK: [[BP1_I32:%.+]] = load i32, i32* %
+// CHECK: [[BP1_I32:%.+]] = load i32, i32* @
// CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
// CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]],
// CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],
@@ -223,7 +223,7 @@ int foo(int n) {
// CHECK: call void (i8*, ...) %
// CHECK: [[DEVICE_CAP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2
// CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
-// CHECK: [[BP1_I32:%.+]] = load i32, i32* %
+// CHECK: [[BP1_I32:%.+]] = load i32, i32* @
// CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
// CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]],
// CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],
diff --git a/clang/test/OpenMP/target_teams_distribute_parallel_for_depend_codegen.cpp b/clang/test/OpenMP/target_teams_distribute_parallel_for_depend_codegen.cpp
index b136e7b75e4e..df1ba2772afc 100644
--- a/clang/test/OpenMP/target_teams_distribute_parallel_for_depend_codegen.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_parallel_for_depend_codegen.cpp
@@ -209,7 +209,7 @@ int foo(int n) {
// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
// CHECK: [[FAIL]]
// CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
-// CHECK: [[BP1_I32:%.+]] = load i32, i32* %
+// CHECK: [[BP1_I32:%.+]] = load i32, i32* @
// CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
// CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]],
// CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],
@@ -223,7 +223,7 @@ int foo(int n) {
// CHECK: call void (i8*, ...) %
// CHECK: [[DEVICE_CAP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2
// CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
-// CHECK: [[BP1_I32:%.+]] = load i32, i32* %
+// CHECK: [[BP1_I32:%.+]] = load i32, i32* @
// CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
// CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]],
// CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],
diff --git a/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_depend_codegen.cpp b/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_depend_codegen.cpp
index 16c73e7406d2..69cdd99e3f46 100644
--- a/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_depend_codegen.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_depend_codegen.cpp
@@ -209,7 +209,7 @@ int foo(int n) {
// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
// CHECK: [[FAIL]]
// CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
-// CHECK: [[BP1_I32:%.+]] = load i32, i32* %
+// CHECK: [[BP1_I32:%.+]] = load i32, i32* @
// CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
// CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]],
// CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],
@@ -223,7 +223,7 @@ int foo(int n) {
// CHECK: call void (i8*, ...) %
// CHECK: [[DEVICE_CAP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2
// CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
-// CHECK: [[BP1_I32:%.+]] = load i32, i32* %
+// CHECK: [[BP1_I32:%.+]] = load i32, i32* @
// CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
// CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]],
// CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],
diff --git a/clang/test/OpenMP/target_teams_distribute_simd_depend_codegen.cpp b/clang/test/OpenMP/target_teams_distribute_simd_depend_codegen.cpp
index 4fffbc50a5eb..06fb3daa3240 100644
--- a/clang/test/OpenMP/target_teams_distribute_simd_depend_codegen.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_simd_depend_codegen.cpp
@@ -209,7 +209,7 @@ int foo(int n) {
// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
// CHECK: [[FAIL]]
// CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
-// CHECK: [[BP1_I32:%.+]] = load i32, i32* %
+// CHECK: [[BP1_I32:%.+]] = load i32, i32* @
// CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
// CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]],
// CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],
@@ -223,7 +223,7 @@ int foo(int n) {
// CHECK: call void (i8*, ...) %
// CHECK: [[DEVICE_CAP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2
// CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
-// CHECK: [[BP1_I32:%.+]] = load i32, i32* %
+// CHECK: [[BP1_I32:%.+]] = load i32, i32* @
// CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
// CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]],
// CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],
More information about the cfe-commits
mailing list