Author: abataev
Date: Fri Oct 6 10:00:28 2017
New Revision: 315076
URL: http://llvm.org/viewvc/llvm-project?rev=315076&view=rev
Log:
[OPENMP] Do not capture local static variables.
Previously we may erroneously try to capture locally declared static
variables, which will lead to crash for target-based constructs.
Patch fixes this problem.
Modified:
cfe/trunk/lib/Sema/SemaOpenMP.cpp
cfe/trunk/test/OpenMP/target_codegen.cpp
Modified: cfe/trunk/lib/Sema/SemaOpenMP.cpp
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaOpenMP.cpp?rev=315076&r1=315075&r2=315076&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaOpenMP.cpp (original)
+++ cfe/trunk/lib/Sema/SemaOpenMP.cpp Fri Oct 6 10:00:28 2017
@@ -1838,6 +1838,10 @@ public:
if (DVar.RefExpr || !ImplicitDeclarations.insert(VD).second)
return;
+ // Skip internally declared static variables.
+ if (VD->hasGlobalStorage() && !CS->capturesVariable(VD))
+ return;
+
auto ELoc = E->getExprLoc();
auto DKind = Stack->getCurrentDirective();
// The default(none) clause requires that each variable that is
referenced
Modified: cfe/trunk/test/OpenMP/target_codegen.cpp
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/target_codegen.cpp?rev=315076&r1=315075&r2=315076&view=diff
==============================================================================
--- cfe/trunk/test/OpenMP/target_codegen.cpp (original)
+++ cfe/trunk/test/OpenMP/target_codegen.cpp Fri Oct 6 10:00:28 2017
@@ -34,7 +34,9 @@
// code, only 6 will have mapped arguments, and only 4 have all-constant map
// sizes.
-// CHECK-DAG: [[SIZET2:@.+]] = private unnamed_addr constant [1 x i{{32|64}}]
[i[[SZ:32|64]] 2]
+// CHECK-DAG: [[SIZET:@.+]] = private unnamed_addr constant [2 x i[[SZ]]]
[i[[SZ]] 0, i[[SZ]] 4]
+// CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [2 x i32] [i32 32,
i32 288]
+// CHECK-DAG: [[SIZET2:@.+]] = private unnamed_addr constant [1 x i{{32|64}}]
[i[[SZ]] 2]
// CHECK-DAG: [[MAPT2:@.+]] = private unnamed_addr constant [1 x i32] [i32 288]
// CHECK-DAG: [[SIZET3:@.+]] = private unnamed_addr constant [2 x i[[SZ]]]
[i[[SZ]] 4, i[[SZ]] 2]
// CHECK-DAG: [[MAPT3:@.+]] = private unnamed_addr constant [2 x i32] [i32
288, i32 288]
@@ -59,6 +61,7 @@
// TCHECK: @{{.+}} = constant [[ENTTY]]
// TCHECK: @{{.+}} = constant [[ENTTY]]
// TCHECK: @{{.+}} = constant [[ENTTY]]
+// TCHECK: @{{.+}} = {{.*}}constant [[ENTTY]]
// TCHECK-NOT: @{{.+}} = constant [[ENTTY]]
// Check if offloading descriptor is created.
@@ -91,6 +94,7 @@ int foo(int n) {
double c[5][10];
double cn[5][n];
TT<long long, char> d;
+ static long *plocal;
// CHECK: [[ADD:%.+]] = add nsw i32
// CHECK: store i32 [[ADD]], i32* [[CAPTURE:%.+]],
@@ -106,6 +110,39 @@ int foo(int n) {
{
}
+ // CHECK: [[ADD:%.+]] = add nsw i32
+ // CHECK: store i32 [[ADD]], i32* [[CAPTURE:%.+]],
+ // CHECK-DAG: [[LD:%.+]] = load i32, i32* [[CAPTURE]],
+ // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target(i32 [[LD]], i8*
@{{[^,]+}}, i32 2, i8** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i[[SZ]]*
getelementptr inbounds ([2 x i[[SZ]]], [2 x i[[SZ]]]* [[SIZET]], i32 0, i32 0),
i32* getelementptr inbounds ([2 x i32], [2 x i32]* [[MAPT]], i32 0, i32 0)
+ // CHECK-DAG: [[BPR]] = getelementptr inbounds [2 x i8*], [2 x i8*]*
[[BP:%[^,]+]], i32 0, i32 0
+ // CHECK-DAG: [[PR]] = getelementptr inbounds [2 x i8*], [2 x i8*]*
[[P:%[^,]+]], i32 0, i32 0
+
+ // CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [2 x i8*], [2 x
i8*]* [[BP]], i32 0, i32 0
+ // CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [2 x i8*], [2 x
i8*]* [[P]], i32 0, i32 0
+ // CHECK-DAG: [[CBPADDR0:%.+]] = bitcast i8** [[BPADDR0]] to i[[SZ]]**
+ // CHECK-DAG: [[CPADDR0:%.+]] = bitcast i8** [[PADDR0]] to i[[SZ]]**
+ // CHECK-DAG: store i[[SZ]]* [[BP0:%[^,]+]], i[[SZ]]** [[CBPADDR0]]
+ // CHECK-DAG: store i[[SZ]]* [[BP0]], i[[SZ]]** [[CPADDR0]]
+
+ // CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [2 x i8*], [2 x
i8*]* [[BP]], i32 0, i32 1
+ // CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [2 x i8*], [2 x
i8*]* [[P]], i32 0, i32 1
+ // CHECK-DAG: [[CBPADDR1:%.+]] = bitcast i8** [[BPADDR1]] to i[[SZ]]*
+ // CHECK-DAG: [[CPADDR1:%.+]] = bitcast i8** [[PADDR1]] to i[[SZ]]*
+ // CHECK-DAG: store i[[SZ]] [[BP1:%[^,]+]], i[[SZ]]* [[CBPADDR1]]
+ // CHECK-DAG: store i[[SZ]] [[BP1]], i[[SZ]]* [[CPADDR1]]
+ // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
+ // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
+ // CHECK: [[FAIL]]
+ // CHECK: call void [[HVT0_:@.+]](i[[SZ]]* [[BP0]], i[[SZ]] [[BP1]])
+ // CHECK-NEXT: br label %[[END]]
+ // CHECK: [[END]]
+ #pragma omp target device(global + a)
+ {
+ static int local1;
+ *plocal = global;
+ local1 = global;
+ }
+
// CHECK: call void [[HVT1:@.+]](i[[SZ]] {{[^,]+}})
#pragma omp target if(0) firstprivate(global)
{
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
