sfantao updated this revision to Diff 58704.
sfantao updated the summary for this revision.
sfantao added a comment.
Remove most of the logic in the first diff. It is no longer necessary given
that all firstprivate
captures are now passed by value no matter the directive they are captured in.
Still, there is a small fix related with template function and capture of
globals.
http://reviews.llvm.org/D18110
Files:
lib/Sema/SemaOpenMP.cpp
test/OpenMP/target_codegen_global_capture.cpp
Index: test/OpenMP/target_codegen_global_capture.cpp
===================================================================
--- test/OpenMP/target_codegen_global_capture.cpp
+++ test/OpenMP/target_codegen_global_capture.cpp
@@ -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 @@
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 @@
#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 @@
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
Index: lib/Sema/SemaOpenMP.cpp
===================================================================
--- lib/Sema/SemaOpenMP.cpp
+++ lib/Sema/SemaOpenMP.cpp
@@ -818,6 +818,9 @@
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 @@
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);
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
