r271001 - [OpenMP] Fix SEMA bug in the capture of global variables in template functions.
Samuel Antao via cfe-commits
cfe-commits at lists.llvm.org
Fri May 27 08:21:35 PDT 2016
Author: sfantao
Date: Fri May 27 10:21:27 2016
New Revision: 271001
URL: http://llvm.org/viewvc/llvm-project?rev=271001&view=rev
Log:
[OpenMP] Fix SEMA bug in the capture of global variables in template functions.
Summary:
Target regions require globals to be captured. This patch fixes a bug exposed when that happens in a template function.
Reviewers: hfinkel, arpith-jacob, kkwli0, carlo.bertolli, ABataev
Subscribers: guansong, ABataev, cfe-commits, caomhin, fraggamuffin
Differential Revision: http://reviews.llvm.org/D18110
Modified:
cfe/trunk/lib/Sema/SemaOpenMP.cpp
cfe/trunk/test/OpenMP/target_codegen_global_capture.cpp
Modified: cfe/trunk/lib/Sema/SemaOpenMP.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaOpenMP.cpp?rev=271001&r1=271000&r2=271001&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaOpenMP.cpp (original)
+++ cfe/trunk/lib/Sema/SemaOpenMP.cpp Fri May 27 10:21:27 2016
@@ -818,6 +818,9 @@ bool DSAStackTy::hasDirective(
const DeclarationNameInfo &, SourceLocation)>
&DPred,
bool FromParent) {
+ // We look only in the enclosing region.
+ if (Stack.size() < 2)
+ return false;
auto StartI = std::next(Stack.rbegin());
auto EndI = std::prev(Stack.rend());
if (FromParent && StartI != EndI) {
@@ -990,8 +993,7 @@ VarDecl *Sema::IsOpenMPCapturedDecl(Valu
if (DSAStack->getCurrentDirective() == OMPD_target &&
!DSAStack->isClauseParsingMode())
return VD;
- if (DSAStack->getCurScope() &&
- DSAStack->hasDirective(
+ if (DSAStack->hasDirective(
[](OpenMPDirectiveKind K, const DeclarationNameInfo &,
SourceLocation) -> bool {
return isOpenMPTargetExecutionDirective(K);
Modified: cfe/trunk/test/OpenMP/target_codegen_global_capture.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/target_codegen_global_capture.cpp?rev=271001&r1=271000&r2=271001&view=diff
==============================================================================
--- cfe/trunk/test/OpenMP/target_codegen_global_capture.cpp (original)
+++ cfe/trunk/test/OpenMP/target_codegen_global_capture.cpp Fri May 27 10:21:27 2016
@@ -21,6 +21,11 @@
// CHECK-DAG: [[BB:@.+]] = internal global float 1.000000e+01
// CHECK-DAG: [[BC:@.+]] = internal global float 1.100000e+01
// CHECK-DAG: [[BD:@.+]] = internal global float 1.200000e+01
+// CHECK-DAG: [[TBA:@.+]] = {{.*}}global float 1.700000e+01
+// CHECK-DAG: [[TBB:@.+]] = {{.*}}global float 1.800000e+01
+// CHECK-DAG: [[TBC:@.+]] = {{.*}}global float 1.900000e+01
+// CHECK-DAG: [[TBD:@.+]] = {{.*}}global float 2.000000e+01
+
double Ga = 1.0;
double Gb = 2.0;
double Gc = 3.0;
@@ -42,14 +47,14 @@ int foo(short a, short b, short c, short
static float Sd = 8.0;
// CHECK-DAG: [[VALLB:%.+]] = load i16, i16* [[LB]],
- // CHECK-64-DAG: [[VALGB:%.+]] = load double, double* @Gb,
- // CHECK-DAG: [[VALFB:%.+]] = load float, float* @_ZZ3foossssE2Sb,
- // CHECK-64-DAG: [[VALGC:%.+]] = load double, double* @Gc,
+ // CHECK-64-DAG: [[VALGB:%.+]] = load double, double* [[GB]],
+ // CHECK-DAG: [[VALFB:%.+]] = load float, float* [[FB]],
+ // CHECK-64-DAG: [[VALGC:%.+]] = load double, double* [[GC]],
// CHECK-DAG: [[VALLC:%.+]] = load i16, i16* [[LC]],
- // CHECK-DAG: [[VALFC:%.+]] = load float, float* @_ZZ3foossssE2Sc,
+ // CHECK-DAG: [[VALFC:%.+]] = load float, float* [[FC]],
// CHECK-DAG: [[VALLD:%.+]] = load i16, i16* [[LD]],
- // CHECK-64-DAG: [[VALGD:%.+]] = load double, double* @Gd,
- // CHECK-DAG: [[VALFD:%.+]] = load float, float* @_ZZ3foossssE2Sd,
+ // CHECK-64-DAG: [[VALGD:%.+]] = load double, double* [[GD]],
+ // CHECK-DAG: [[VALFD:%.+]] = load float, float* [[FD]],
// 3 local vars being captured.
@@ -178,14 +183,156 @@ int bar(short a, short b, short c, short
#pragma omp parallel
{
// CHECK-DAG: [[VALLB:%.+]] = load i16, i16* [[LLB]],
- // CHECK-64-DAG: [[VALGB:%.+]] = load double, double* @Gb,
- // CHECK-DAG: [[VALFB:%.+]] = load float, float* @_ZZ3barssssE2Sb,
- // CHECK-64-DAG: [[VALGC:%.+]] = load double, double* @Gc,
+ // CHECK-64-DAG: [[VALGB:%.+]] = load double, double* [[GB]],
+ // CHECK-DAG: [[VALFB:%.+]] = load float, float* [[BB]],
+ // CHECK-64-DAG: [[VALGC:%.+]] = load double, double* [[GC]],
+ // CHECK-DAG: [[VALLC:%.+]] = load i16, i16* [[LLC]],
+ // CHECK-DAG: [[VALFC:%.+]] = load float, float* [[BC]],
+ // CHECK-DAG: [[VALLD:%.+]] = load i16, i16* [[LLD]],
+ // CHECK-64-DAG: [[VALGD:%.+]] = load double, double* [[GD]],
+ // CHECK-DAG: [[VALFD:%.+]] = load float, float* [[BD]],
+
+ // 3 local vars being captured.
+
+ // CHECK-DAG: store i16 [[VALLB]], i16* [[CONVLB:%.+]],
+ // CHECK-DAG: [[CONVLB]] = bitcast i[[sz:64|32]]* [[CADDRLB:%.+]] to i16*
+ // CHECK-DAG: [[CVALLB:%.+]] = load i[[sz]], i[[sz]]* [[CADDRLB]],
+ // CHECK-DAG: [[CPTRLB:%.+]] = inttoptr i[[sz]] [[CVALLB]] to i8*
+ // CHECK-DAG: store i8* [[CPTRLB]], i8** [[GEPLB:%.+]],
+ // CHECK-DAG: [[GEPLB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
+
+ // CHECK-DAG: store i16 [[VALLC]], i16* [[CONVLC:%.+]],
+ // CHECK-DAG: [[CONVLC]] = bitcast i[[sz]]* [[CADDRLC:%.+]] to i16*
+ // CHECK-DAG: [[CVALLC:%.+]] = load i[[sz]], i[[sz]]* [[CADDRLC]],
+ // CHECK-DAG: [[CPTRLC:%.+]] = inttoptr i[[sz]] [[CVALLC]] to i8*
+ // CHECK-DAG: store i8* [[CPTRLC]], i8** [[GEPLC:%.+]],
+ // CHECK-DAG: [[GEPLC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
+
+ // CHECK-DAG: store i16 [[VALLD]], i16* [[CONVLD:%.+]],
+ // CHECK-DAG: [[CONVLD]] = bitcast i[[sz]]* [[CADDRLD:%.+]] to i16*
+ // CHECK-DAG: [[CVALLD:%.+]] = load i[[sz]], i[[sz]]* [[CADDRLD]],
+ // CHECK-DAG: [[CPTRLD:%.+]] = inttoptr i[[sz]] [[CVALLD]] to i8*
+ // CHECK-DAG: store i8* [[CPTRLD]], i8** [[GEPLD:%.+]],
+ // CHECK-DAG: [[GEPLD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
+
+ // 3 static vars being captured.
+
+ // CHECK-DAG: store float [[VALFB]], float* [[CONVFB:%.+]],
+ // CHECK-DAG: [[CONVFB]] = bitcast i[[sz]]* [[CADDRFB:%.+]] to float*
+ // CHECK-DAG: [[CVALFB:%.+]] = load i[[sz]], i[[sz]]* [[CADDRFB]],
+ // CHECK-DAG: [[CPTRFB:%.+]] = inttoptr i[[sz]] [[CVALFB]] to i8*
+ // CHECK-DAG: store i8* [[CPTRFB]], i8** [[GEPFB:%.+]],
+ // CHECK-DAG: [[GEPFB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
+
+ // CHECK-DAG: store float [[VALFC]], float* [[CONVFC:%.+]],
+ // CHECK-DAG: [[CONVFC]] = bitcast i[[sz]]* [[CADDRFC:%.+]] to float*
+ // CHECK-DAG: [[CVALFC:%.+]] = load i[[sz]], i[[sz]]* [[CADDRFC]],
+ // CHECK-DAG: [[CPTRFC:%.+]] = inttoptr i[[sz]] [[CVALFC]] to i8*
+ // CHECK-DAG: store i8* [[CPTRFC]], i8** [[GEPFC:%.+]],
+ // CHECK-DAG: [[GEPFC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
+
+ // CHECK-DAG: store float [[VALFD]], float* [[CONVFD:%.+]],
+ // CHECK-DAG: [[CONVFD]] = bitcast i[[sz]]* [[CADDRFD:%.+]] to float*
+ // CHECK-DAG: [[CVALFD:%.+]] = load i[[sz]], i[[sz]]* [[CADDRFD]],
+ // CHECK-DAG: [[CPTRFD:%.+]] = inttoptr i[[sz]] [[CVALFD]] to i8*
+ // CHECK-DAG: store i8* [[CPTRFD]], i8** [[GEPFD:%.+]],
+ // CHECK-DAG: [[GEPFD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
+
+ // 3 static global vars being captured.
+
+ // CHECK-64-DAG: store double [[VALGB]], double* [[CONVGB:%.+]],
+ // CHECK-64-DAG: [[CONVGB]] = bitcast i[[sz]]* [[CADDRGB:%.+]] to double*
+ // CHECK-64-DAG: [[CVALGB:%.+]] = load i[[sz]], i[[sz]]* [[CADDRGB]],
+ // CHECK-64-DAG: [[CPTRGB:%.+]] = inttoptr i[[sz]] [[CVALGB]] to i8*
+ // CHECK-64-DAG: store i8* [[CPTRGB]], i8** [[GEPGB:%.+]],
+ // CHECK-32-DAG: store i8* bitcast (double* @Gb to i8*), i8** [[GEPGB:%.+]],
+ // CHECK-DAG: [[GEPGB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
+
+ // CHECK-64-DAG: store double [[VALGC]], double* [[CONVGC:%.+]],
+ // CHECK-64-DAG: [[CONVGC]] = bitcast i[[sz]]* [[CADDRGC:%.+]] to double*
+ // CHECK-64-DAG: [[CVALGC:%.+]] = load i[[sz]], i[[sz]]* [[CADDRGC]],
+ // CHECK-64-DAG: [[CPTRGC:%.+]] = inttoptr i[[sz]] [[CVALGC]] to i8*
+ // CHECK-64-DAG: store i8* [[CPTRGC]], i8** [[GEPGC:%.+]],
+ // CHECK-32-DAG: store i8* bitcast (double* @Gc to i8*), i8** [[GEPGC:%.+]],
+ // CHECK-DAG: [[GEPGC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
+
+ // CHECK-64-DAG: store double [[VALGD]], double* [[CONVGD:%.+]],
+ // CHECK-64-DAG: [[CONVGD]] = bitcast i[[sz]]* [[CADDRGD:%.+]] to double*
+ // CHECK-64-DAG: [[CVALGD:%.+]] = load i[[sz]], i[[sz]]* [[CADDRGD]],
+ // CHECK-64-DAG: [[CPTRGD:%.+]] = inttoptr i[[sz]] [[CVALGD]] to i8*
+ // CHECK-64-DAG: store i8* [[CPTRGD]], i8** [[GEPGD:%.+]],
+ // CHECK-32-DAG: store i8* bitcast (double* @Gd to i8*), i8** [[GEPGD:%.+]],
+ // CHECK-DAG: [[GEPGD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
+
+ // CHECK: call i32 @__tgt_target
+ // CHECK: call void [[OFFLOADF:@.+]](
+ // Capture b, Gb, Sb, Gc, c, Sc, d, Gd, Sd
+ #pragma omp target if(Ga>0.0 && a>0 && Sa>0.0)
+ {
+ b += 1;
+ Gb += 1.0;
+ Sb += 1.0;
+
+ // CHECK: define internal void [[OFFLOADF]]({{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}})
+ // CHECK: call void {{.*}}@__kmpc_fork_call(%ident_t* {{.+}}, i32 {{.+}}, void (i32*, i32*, ...)* bitcast ({{.*}}[[PARF:@.+]] to {{.*}})
+
+ // CHECK: define internal void [[PARF]](i32* noalias %{{.*}}, i32* noalias %{{.*}}, {{.+}}* dereferenceable({{.+}}) %{{.+}}, {{.+}}* dereferenceable({{.+}}) %{{.+}}, {{.+}}* dereferenceable({{.+}}) %{{.+}})
+ // Capture d, Gd, Sd
+ #pragma omp parallel if(Gc>0.0 && c>0 && Sc>0.0)
+ {
+ d += 1;
+ Gd += 1.0;
+ Sd += 1.0;
+ }
+ }
+ }
+ return a + b + c + d + (int)Sa + (int)Sb + (int)Sc + (int)Sd;
+}
+
+///
+/// Tests with template functions.
+///
+
+// CHECK: define {{.*}} @{{.*}}tbar2{{.*}}(
+
+// CHECK: define {{.*}} @{{.*}}tbar{{.*}}(
+// CHECK-SAME: i16 {{[^,]*}}[[A:%[^,]+]],
+// CHECK-SAME: i16 {{[^,]*}}[[B:%[^,]+]],
+// CHECK-SAME: i16 {{[^,]*}}[[C:%[^,]+]],
+// CHECK-SAME: i16 {{[^,]*}}[[D:%[^,]+]])
+// CHECK: [[LA:%.+]] = alloca i16
+// CHECK: [[LB:%.+]] = alloca i16
+// CHECK: [[LC:%.+]] = alloca i16
+// CHECK: [[LD:%.+]] = alloca i16
+template<typename T>
+int tbar(T a, T b, T c, T d){
+ static float Sa = 17.0;
+ static float Sb = 18.0;
+ static float Sc = 19.0;
+ static float Sd = 20.0;
+
+ // CHECK: call void {{.*}}@__kmpc_fork_call(%ident_t* {{.+}}, i32 {{.+}}, void (i32*, i32*, ...)* bitcast ({{.*}}[[PARF:@.+]] to {{.*}}), i16* %{{.+}}, i16* %{{.+}}, i16* %{{.+}}, i16* %{{.+}})
+ // CHECK: define internal void [[PARF]](i32* noalias %{{.*}}, i32* noalias %{{.*}}, i16* dereferenceable(2) [[A:%.+]], i16* dereferenceable(2) [[B:%.+]], i16* dereferenceable(2) [[C:%.+]], i16* dereferenceable(2) [[D:%.+]])
+ // Capture a, b, c, d
+ // CHECK: [[ALLOCLA:%.+]] = alloca i16
+ // CHECK: [[ALLOCLB:%.+]] = alloca i16
+ // CHECK: [[ALLOCLC:%.+]] = alloca i16
+ // CHECK: [[ALLOCLD:%.+]] = alloca i16
+ // CHECK: [[LLA:%.+]] = load i16*, i16** [[ALLOCLA]],
+ // CHECK: [[LLB:%.+]] = load i16*, i16** [[ALLOCLB]],
+ // CHECK: [[LLC:%.+]] = load i16*, i16** [[ALLOCLC]],
+ // CHECK: [[LLD:%.+]] = load i16*, i16** [[ALLOCLD]],
+ #pragma omp parallel
+ {
+ // CHECK-DAG: [[VALLB:%.+]] = load i16, i16* [[LLB]],
+ // CHECK-64-DAG: [[VALGB:%.+]] = load double, double* [[GB]],
+ // CHECK-DAG: [[VALFB:%.+]] = load float, float* [[TBB]],
+ // CHECK-64-DAG: [[VALGC:%.+]] = load double, double* [[GC]],
// CHECK-DAG: [[VALLC:%.+]] = load i16, i16* [[LLC]],
- // CHECK-DAG: [[VALFC:%.+]] = load float, float* @_ZZ3barssssE2Sc,
+ // CHECK-DAG: [[VALFC:%.+]] = load float, float* [[TBC]],
// CHECK-DAG: [[VALLD:%.+]] = load i16, i16* [[LLD]],
- // CHECK-64-DAG: [[VALGD:%.+]] = load double, double* @Gd,
- // CHECK-DAG: [[VALFD:%.+]] = load float, float* @_ZZ3barssssE2Sd,
+ // CHECK-64-DAG: [[VALGD:%.+]] = load double, double* [[GD]],
+ // CHECK-DAG: [[VALFD:%.+]] = load float, float* [[TBD]],
// 3 local vars being captured.
@@ -284,4 +431,8 @@ int bar(short a, short b, short c, short
return a + b + c + d + (int)Sa + (int)Sb + (int)Sc + (int)Sd;
}
+int tbar2(short a, short b, short c, short d){
+ return tbar(a, b, c, d);
+}
+
#endif
More information about the cfe-commits
mailing list