lildmh created this revision.
Herald added subscribers: cfe-commits, guansong.
Herald added a reviewer: jdoerfert.
Herald added a project: clang.
This patch implements the code generation to use OpenMP 5.0 declare mapper
(e.g., user-defined mapper) constructs. It looks up the proper mapper function
for each map, to, or from clause that has a user-defined mapper associated, and
passes them to the OpenMP runtime function.
The design slides can be found at
https://github.com/lingda-li/public-sharing/blob/master/mapper_runtime_design.pptx
Repository:
rC Clang
https://reviews.llvm.org/D67833
Files:
include/clang/AST/OpenMPClause.h
lib/CodeGen/CGOpenMPRuntime.cpp
lib/CodeGen/CGOpenMPRuntime.h
lib/CodeGen/CGStmtOpenMP.cpp
lib/CodeGen/CodeGenFunction.h
test/OpenMP/capturing_in_templates.cpp
test/OpenMP/declare_mapper_codegen.cpp
test/OpenMP/declare_target_link_codegen.cpp
test/OpenMP/target_is_device_ptr_codegen.cpp
Index: test/OpenMP/target_is_device_ptr_codegen.cpp
===================================================================
--- test/OpenMP/target_is_device_ptr_codegen.cpp
+++ test/OpenMP/target_is_device_ptr_codegen.cpp
@@ -49,7 +49,7 @@
float *l;
T *t;
- // CK1-DAG: call i32 @__tgt_target(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES00]]{{.+}}, {{.+}}[[TYPES00]]{{.+}})
+ // CK1-DAG: call i32 @__tgt_target_mapper(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES00]]{{.+}}, {{.+}}[[TYPES00]]{{.+}})
// CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
// CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
// CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
@@ -66,7 +66,7 @@
++g;
}
- // CK1-DAG: call i32 @__tgt_target(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES01]]{{.+}}, {{.+}}[[TYPES01]]{{.+}})
+ // CK1-DAG: call i32 @__tgt_target_mapper(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES01]]{{.+}}, {{.+}}[[TYPES01]]{{.+}})
// CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
// CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
// CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
@@ -83,7 +83,7 @@
++l;
}
- // CK1-DAG: call i32 @__tgt_target(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES02]]{{.+}}, {{.+}}[[TYPES02]]{{.+}})
+ // CK1-DAG: call i32 @__tgt_target_mapper(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES02]]{{.+}}, {{.+}}[[TYPES02]]{{.+}})
// CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
// CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
// CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
@@ -100,7 +100,7 @@
++t;
}
- // CK1-DAG: call i32 @__tgt_target(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES03]]{{.+}}, {{.+}}[[TYPES03]]{{.+}})
+ // CK1-DAG: call i32 @__tgt_target_mapper(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES03]]{{.+}}, {{.+}}[[TYPES03]]{{.+}})
// CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
// CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
// CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
@@ -118,7 +118,7 @@
++lr;
}
- // CK1-DAG: call i32 @__tgt_target(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES04]]{{.+}}, {{.+}}[[TYPES04]]{{.+}})
+ // CK1-DAG: call i32 @__tgt_target_mapper(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES04]]{{.+}}, {{.+}}[[TYPES04]]{{.+}})
// CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
// CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
// CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
@@ -136,7 +136,7 @@
++tr;
}
- // CK1-DAG: call i32 @__tgt_target(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES05]]{{.+}}, {{.+}}[[TYPES05]]{{.+}})
+ // CK1-DAG: call i32 @__tgt_target_mapper(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES05]]{{.+}}, {{.+}}[[TYPES05]]{{.+}})
// CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
// CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
// CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
@@ -154,7 +154,7 @@
++tr;
}
- // CK1-DAG: call i32 @__tgt_target(i64 {{.+}}, i8* {{.+}}, i32 2, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES06]]{{.+}}, {{.+}}[[TYPES06]]{{.+}})
+ // CK1-DAG: call i32 @__tgt_target_mapper(i64 {{.+}}, i8* {{.+}}, i32 2, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES06]]{{.+}}, {{.+}}[[TYPES06]]{{.+}})
// CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
// CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
// CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
@@ -231,7 +231,7 @@
void foo(double *&arg) {
int *la = 0;
- // CK2-DAG: call i32 @__tgt_target(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}})
+ // CK2-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}})
// CK2-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK2-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
@@ -246,7 +246,7 @@
a++;
}
- // CK2-DAG: call i32 @__tgt_target(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE01]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}})
+ // CK2-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE01]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}})
// CK2-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK2-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
@@ -261,7 +261,7 @@
b++;
}
- // CK2-DAG: call i32 @__tgt_target(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}})
+ // CK2-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}})
// CK2-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK2-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
Index: test/OpenMP/declare_target_link_codegen.cpp
===================================================================
--- test/OpenMP/declare_target_link_codegen.cpp
+++ test/OpenMP/declare_target_link_codegen.cpp
@@ -77,9 +77,9 @@
// HOST: [[BP0:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BASEPTRS]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
// HOST: [[P0:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PTRS]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
-// HOST: call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 3, i8** [[BP0]], i8** [[P0]], i64* getelementptr inbounds ([3 x i64], [3 x i64]* [[SIZES]], i{{[0-9]+}} 0, i{{[0-9]+}} 0), i64* getelementptr inbounds ([3 x i64], [3 x i64]* [[MAPTYPES]], i{{[0-9]+}} 0, i{{[0-9]+}} 0))
+// HOST: call i32 @__tgt_target_mapper(i64 -1, i8* @{{[^,]+}}, i32 3, i8** [[BP0]], i8** [[P0]], i64* getelementptr inbounds ([3 x i64], [3 x i64]* [[SIZES]], i{{[0-9]+}} 0, i{{[0-9]+}} 0), i64* getelementptr inbounds ([3 x i64], [3 x i64]* [[MAPTYPES]], i{{[0-9]+}} 0, i{{[0-9]+}} 0))
// HOST: call void @__omp_offloading_{{.*}}_{{.*}}_{{.*}}maini1{{.*}}_l42(i32* %{{[^,]+}})
-// HOST: call i32 @__tgt_target_teams(i64 -1, i8* @.__omp_offloading_{{.+}}_l47.region_id, i32 2, {{.+}})
+// HOST: call i32 @__tgt_target_teams_mapper(i64 -1, i8* @.__omp_offloading_{{.+}}_l47.region_id, i32 2, {{.+}})
// HOST: define internal void @__omp_offloading_{{.*}}_{{.*}}maini1{{.*}}_l42(i32* dereferenceable{{.*}})
// HOST: [[C:%.*]] = load i32, i32* @c,
Index: test/OpenMP/declare_mapper_codegen.cpp
===================================================================
--- test/OpenMP/declare_mapper_codegen.cpp
+++ test/OpenMP/declare_mapper_codegen.cpp
@@ -22,14 +22,11 @@
#ifdef CK0
// CK0-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}.region_id = weak constant i8 0
-// CK0-64: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
-// CK0-32: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
+// CK0: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 1]
// CK0: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 35]
-// CK0-64: [[TSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
-// CK0-32: [[TSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
+// CK0: [[TSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 1]
// CK0: [[TTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 33]
-// CK0-64: [[FSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
-// CK0-32: [[FSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
+// CK0: [[FSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 1]
// CK0: [[FTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 34]
class C {
@@ -40,7 +37,7 @@
#pragma omp declare mapper(id: C s) map(s.a, s.b[0:2])
-// CK0-LABEL: define {{.*}}void @.omp_mapper.{{.*}}C.id{{.*}}(i8*{{.*}}, i8*{{.*}}, i8*{{.*}}, i64{{.*}}, i64{{.*}})
+// CK0: define {{.*}}void [[MPRFUNC:@[.]omp_mapper[.].*C[.]id]](i8*{{.*}}, i8*{{.*}}, i8*{{.*}}, i64{{.*}}, i64{{.*}})
// CK0: store i8* %{{[^,]+}}, i8** [[HANDLEADDR:%[^,]+]]
// CK0: store i8* %{{[^,]+}}, i8** [[BPTRADDR:%[^,]+]]
// CK0: store i8* %{{[^,]+}}, i8** [[VPTRADDR:%[^,]+]]
@@ -209,41 +206,53 @@
C c;
c.a = a;
- // CK0-DAG: call i32 @__tgt_target(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}})
+ // CK0-DAG: call i32 @__tgt_target_mapper(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}, i8** [[MPRGEP:%.+]])
// CK0-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
// CK0-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
+ // CK0-DAG: [[MPRGEP]] = getelementptr inbounds {{.+}}[[MPR:%[^,]+]], i32 0, i32 0
// CK0-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
// CK0-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
+ // CK0-DAG: [[MPR1:%.+]] = getelementptr inbounds {{.+}}[[MPR]], i32 0, i32 0
// CK0-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to %class.C**
// CK0-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to %class.C**
+ // CK0-DAG: [[CMPR1:%.+]] = bitcast i8** [[MPR1]] to void (i8*, i8*, i8*, i64, i64)**
// CK0-DAG: store %class.C* [[VAL:%[^,]+]], %class.C** [[CBP1]]
// CK0-DAG: store %class.C* [[VAL]], %class.C** [[CP1]]
+ // CK0-DAG: store void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]], void (i8*, i8*, i8*, i64, i64)** [[CMPR1]]
// CK0: call void [[KERNEL:@.+]](%class.C* [[VAL]])
#pragma omp target map(mapper(id),tofrom: c)
{
- ++c.a;
+ ++c.a;
}
- // CK0-DAG: call void @__tgt_target_data_update(i64 -1, i32 1, i8** [[TGEPBP:%.+]], i8** [[TGEPP:%.+]], i64* getelementptr {{.+}}[1 x i64]* [[TSIZES]], i32 0, i32 0), {{.+}}getelementptr {{.+}}[1 x i64]* [[TTYPES]]{{.+}})
+ // CK0-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[TGEPBP:%.+]], i8** [[TGEPP:%.+]], i64* getelementptr {{.+}}[1 x i64]* [[TSIZES]], i32 0, i32 0), {{.+}}getelementptr {{.+}}[1 x i64]* [[TTYPES]]{{.+}}, i8** [[TMPRGEP:%.+]])
// CK0-DAG: [[TGEPBP]] = getelementptr inbounds {{.+}}[[TBP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
// CK0-DAG: [[TGEPP]] = getelementptr inbounds {{.+}}[[TP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
+ // CK0-DAG: [[TMPRGEP]] = getelementptr inbounds {{.+}}[[TMPR:%[^,]+]], i32 0, i32 0
// CK0-DAG: [[TBP0:%.+]] = getelementptr inbounds {{.+}}[[TBP]], i{{.+}} 0, i{{.+}} 0
// CK0-DAG: [[TP0:%.+]] = getelementptr inbounds {{.+}}[[TP]], i{{.+}} 0, i{{.+}} 0
+ // CK0-DAG: [[TMPR1:%.+]] = getelementptr inbounds {{.+}}[[TMPR]], i32 0, i32 0
// CK0-DAG: [[TCBP0:%.+]] = bitcast i8** [[TBP0]] to %class.C**
// CK0-DAG: [[TCP0:%.+]] = bitcast i8** [[TP0]] to %class.C**
+ // CK0-DAG: [[TCMPR1:%.+]] = bitcast i8** [[TMPR1]] to void (i8*, i8*, i8*, i64, i64)**
// CK0-DAG: store %class.C* [[VAL]], %class.C** [[TCBP0]]
// CK0-DAG: store %class.C* [[VAL]], %class.C** [[TCP0]]
+ // CK0-DAG: store void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]], void (i8*, i8*, i8*, i64, i64)** [[TCMPR1]]
#pragma omp target update to(mapper(id): c)
- // CK0-DAG: call void @__tgt_target_data_update(i64 -1, i32 1, i8** [[FGEPBP:%.+]], i8** [[FGEPP:%.+]], i64* getelementptr {{.+}}[1 x i64]* [[FSIZES]], i32 0, i32 0), {{.+}}getelementptr {{.+}}[1 x i64]* [[FTYPES]]{{.+}})
+ // CK0-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[FGEPBP:%.+]], i8** [[FGEPP:%.+]], i64* getelementptr {{.+}}[1 x i64]* [[FSIZES]], i32 0, i32 0), {{.+}}getelementptr {{.+}}[1 x i64]* [[FTYPES]]{{.+}}, i8** [[FMPRGEP:%.+]])
// CK0-DAG: [[FGEPBP]] = getelementptr inbounds {{.+}}[[FBP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
// CK0-DAG: [[FGEPP]] = getelementptr inbounds {{.+}}[[FP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
+ // CK0-DAG: [[FMPRGEP]] = getelementptr inbounds {{.+}}[[FMPR:%[^,]+]], i32 0, i32 0
// CK0-DAG: [[FBP0:%.+]] = getelementptr inbounds {{.+}}[[FBP]], i{{.+}} 0, i{{.+}} 0
// CK0-DAG: [[FP0:%.+]] = getelementptr inbounds {{.+}}[[FP]], i{{.+}} 0, i{{.+}} 0
+ // CK0-DAG: [[FMPR1:%.+]] = getelementptr inbounds {{.+}}[[FMPR]], i32 0, i32 0
// CK0-DAG: [[FCBP0:%.+]] = bitcast i8** [[FBP0]] to %class.C**
// CK0-DAG: [[FCP0:%.+]] = bitcast i8** [[FP0]] to %class.C**
+ // CK0-DAG: [[FCMPR1:%.+]] = bitcast i8** [[FMPR1]] to void (i8*, i8*, i8*, i64, i64)**
// CK0-DAG: store %class.C* [[VAL]], %class.C** [[FCBP0]]
// CK0-DAG: store %class.C* [[VAL]], %class.C** [[FCP0]]
+ // CK0-DAG: store void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]], void (i8*, i8*, i8*, i64, i64)** [[FCMPR1]]
#pragma omp target update from(mapper(id): c)
}
@@ -257,7 +266,7 @@
// CK0: {{.+}} = add nsw i32 [[VAL]], 1
// CK0: }
-#endif
+#endif // CK0
///==========================================================================///
@@ -276,6 +285,7 @@
// RUN: %clang_cc1 -DCK1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
#ifdef CK1
+// C++ template
template <class T>
class C {
@@ -409,6 +419,168 @@
// CK1: [[DONE]]
// CK1: ret void
-#endif
+#endif // CK1
-#endif
+
+///==========================================================================///
+// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix CK2 --check-prefix CK2-64 %s
+// RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
+// RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix CK2 --check-prefix CK2-64 %s
+// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix CK2 --check-prefix CK2-32 %s
+// RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
+// RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix CK2 --check-prefix CK2-32 %s
+
+// RUN: %clang_cc1 -DCK2 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
+// RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK2 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
+// RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+
+#ifdef CK2
+// Nested mappers.
+
+class B {
+public:
+ double a;
+};
+
+class C {
+public:
+ double a;
+ B b;
+};
+
+#pragma omp declare mapper(B s) map(s.a)
+
+#pragma omp declare mapper(id: C s) map(s.b)
+
+// CK2: define {{.*}}void [[BMPRFUNC:@[.]omp_mapper[.].*B[.]default]](i8*{{.*}}, i8*{{.*}}, i8*{{.*}}, i64{{.*}}, i64{{.*}})
+
+// CK2-LABEL: define {{.*}}void @.omp_mapper.{{.*}}C{{.*}}.id(i8*{{.*}}, i8*{{.*}}, i8*{{.*}}, i64{{.*}}, i64{{.*}})
+// CK2: store i8* %{{[^,]+}}, i8** [[HANDLEADDR:%[^,]+]]
+// CK2: store i8* %{{[^,]+}}, i8** [[BPTRADDR:%[^,]+]]
+// CK2: store i8* %{{[^,]+}}, i8** [[VPTRADDR:%[^,]+]]
+// CK2: store i64 %{{[^,]+}}, i{{64|32}}* [[SIZEADDR:%[^,]+]]
+// CK2: store i64 %{{[^,]+}}, i64* [[TYPEADDR:%[^,]+]]
+// CK2-DAG: [[SIZE:%.+]] = load i64, i64* [[SIZEADDR]]
+// CK2-DAG: [[TYPE:%.+]] = load i64, i64* [[TYPEADDR]]
+// CK2-DAG: [[HANDLE:%.+]] = load i8*, i8** [[HANDLEADDR]]
+// CK2-DAG: [[PTRBEGIN:%.+]] = bitcast i8** [[VPTRADDR]] to %class.C**
+// CK2-DAG: [[PTREND:%.+]] = getelementptr %class.C*, %class.C** [[PTRBEGIN]], i64 [[SIZE]]
+// CK2-DAG: [[BPTR:%.+]] = load i8*, i8** [[BPTRADDR]]
+// CK2-DAG: [[BEGIN:%.+]] = load i8*, i8** [[VPTRADDR]]
+// CK2: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1
+// CK2: br i1 [[ISARRAY]], label %[[INITEVALDEL:[^,]+]], label %[[LHEAD:[^,]+]]
+
+// CK2: [[INITEVALDEL]]
+// CK2: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
+// CK2: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0
+// CK2: br i1 [[ISNOTDEL]], label %[[INIT:[^,]+]], label %[[LHEAD:[^,]+]]
+// CK2: [[INIT]]
+// CK2-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
+// CK2-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4
+// CK2: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]])
+// CK2: br label %[[LHEAD:[^,]+]]
+
+// CK2: [[LHEAD]]
+// CK2: [[ISEMPTY:%.+]] = icmp eq %class.C** [[PTRBEGIN]], [[PTREND]]
+// CK2: br i1 [[ISEMPTY]], label %[[DONE:[^,]+]], label %[[LBODY:[^,]+]]
+// CK2: [[LBODY]]
+// CK2: [[PTR:%.+]] = phi %class.C** [ [[PTRBEGIN]], %[[LHEAD]] ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ]
+// CK2: [[OBJ:%.+]] = load %class.C*, %class.C** [[PTR]]
+// CK2-DAG: [[BBEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 1
+// CK2-DAG: [[BEND:%.+]] = getelementptr %class.B, %class.B* [[BBEGIN]], i32 1
+// CK2-DAG: [[BBEGINV:%.+]] = bitcast %class.B* [[BBEGIN]] to i8*
+// CK2-DAG: [[BENDV:%.+]] = bitcast %class.B* [[BEND]] to i8*
+// CK2-DAG: [[BBEGINI:%.+]] = ptrtoint i8* [[BBEGINV]] to i64
+// CK2-DAG: [[BENDI:%.+]] = ptrtoint i8* [[BENDV]] to i64
+// CK2-DAG: [[BSIZE:%.+]] = sub i64 [[BENDI]], [[BBEGINI]]
+// CK2-DAG: [[BUSIZE:%.+]] = sdiv exact i64 [[BSIZE]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
+// CK2-DAG: [[BPTRADDR0BC:%.+]] = bitcast %class.C* [[OBJ]] to i8*
+// CK2-DAG: [[PTRADDR0BC:%.+]] = bitcast %class.B* [[BBEGIN]] to i8*
+// CK2-DAG: [[PRESIZE:%.+]] = call i64 @__tgt_mapper_num_components(i8* [[HANDLE]])
+// CK2-DAG: [[SHIPRESIZE:%.+]] = shl i64 [[PRESIZE]], 48
+// CK2-DAG: br label %[[MEMBER:[^,]+]]
+// CK2-DAG: [[MEMBER]]
+// CK2-DAG: br i1 true, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]]
+// CK2-DAG: [[MEMBERCOM]]
+// CK2-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 32, [[SHIPRESIZE]]
+// CK2-DAG: br label %[[LTYPE]]
+// CK2-DAG: [[LTYPE]]
+// CK2-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 32, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ]
+// CK2-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
+// CK2-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
+// CK2-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
+// CK2-DAG: [[ALLOC]]
+// CK2-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
+// CK2-DAG: br label %[[TYEND:[^,]+]]
+// CK2-DAG: [[ALLOCELSE]]
+// CK2-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
+// CK2-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
+// CK2-DAG: [[TO]]
+// CK2-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3
+// CK2-DAG: br label %[[TYEND]]
+// CK2-DAG: [[TOELSE]]
+// CK2-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
+// CK2-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
+// CK2-DAG: [[FROM]]
+// CK2-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
+// CK2-DAG: br label %[[TYEND]]
+// CK2-DAG: [[TYEND]]
+// CK2-DAG: [[TYPE0:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
+// CK2-64: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR0BC]], i8* [[PTRADDR0BC]], i64 [[BUSIZE]], i64 [[TYPE0]])
+// CK2-DAG: [[BPTRADDR1BC:%.+]] = bitcast %class.C* [[OBJ]] to i8*
+// CK2-DAG: [[PTRADDR1BC:%.+]] = bitcast %class.B* [[BBEGIN]] to i8*
+// CK2-DAG: br label %[[MEMBER:[^,]+]]
+// CK2-DAG: [[MEMBER]]
+// CK2-DAG: br i1 false, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]]
+// CK2-DAG: [[MEMBERCOM]]
+// 281474976710659 == 0x1,000,000,003
+// CK2-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 281474976710659, [[SHIPRESIZE]]
+// CK2-DAG: br label %[[LTYPE]]
+// CK2-DAG: [[LTYPE]]
+// CK2-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 281474976710659, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ]
+// CK2-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
+// CK2-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
+// CK2-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
+// CK2-DAG: [[ALLOC]]
+// CK2-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
+// CK2-DAG: br label %[[TYEND:[^,]+]]
+// CK2-DAG: [[ALLOCELSE]]
+// CK2-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
+// CK2-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
+// CK2-DAG: [[TO]]
+// CK2-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3
+// CK2-DAG: br label %[[TYEND]]
+// CK2-DAG: [[TOELSE]]
+// CK2-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
+// CK2-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
+// CK2-DAG: [[FROM]]
+// CK2-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
+// CK2-DAG: br label %[[TYEND]]
+// CK2-DAG: [[TYEND]]
+// CK2-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
+// CK2: call void [[BMPRFUNC]](i8* [[HANDLE]], i8* [[BPTRADDR1BC]], i8* [[PTRADDR1BC]], i64 1, i64 [[TYPE1]])
+// CK2: [[PTRNEXT]] = getelementptr %class.C*, %class.C** [[PTR]], i32 1
+// CK2: [[ISDONE:%.+]] = icmp eq %class.C** [[PTRNEXT]], [[PTREND]]
+// CK2: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]]
+
+// CK2: [[LEXIT]]
+// CK2: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1
+// CK2: br i1 [[ISARRAY]], label %[[EVALDEL:[^,]+]], label %[[DONE]]
+// CK2: [[EVALDEL]]
+// CK2: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
+// CK2: [[ISDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0
+// CK2: br i1 [[ISDEL]], label %[[DEL:[^,]+]], label %[[DONE]]
+// CK2: [[DEL]]
+// CK2-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
+// CK2-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4
+// CK2: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE]])
+// CK2: br label %[[DONE]]
+// CK2: [[DONE]]
+// CK2: ret void
+
+#endif // CK2
+
+#endif // HEADER
Index: test/OpenMP/capturing_in_templates.cpp
===================================================================
--- test/OpenMP/capturing_in_templates.cpp
+++ test/OpenMP/capturing_in_templates.cpp
@@ -18,7 +18,7 @@
// CHECK-LABEL: @main
int main(int argc, char **argv) {
-// CHECK: call i32 @__tgt_target(i64 -1, i8* @{{.+}}.region_id, i32 0, i8** null, i8** null, i64* null, i64* null)
+// CHECK: call i32 @__tgt_target_mapper(i64 -1, i8* @{{.+}}.region_id, i32 0, i8** null, i8** null, i64* null, i64* null, i8** null)
#pragma omp target
{
for (int i = 0; i < 64; ++i) {
Index: lib/CodeGen/CodeGenFunction.h
===================================================================
--- lib/CodeGen/CodeGenFunction.h
+++ lib/CodeGen/CodeGenFunction.h
@@ -3114,12 +3114,15 @@
Address BasePointersArray = Address::invalid();
Address PointersArray = Address::invalid();
Address SizesArray = Address::invalid();
+ Address MappersArray = Address::invalid();
unsigned NumberOfTargetItems = 0;
explicit OMPTargetDataInfo() = default;
OMPTargetDataInfo(Address BasePointersArray, Address PointersArray,
- Address SizesArray, unsigned NumberOfTargetItems)
+ Address SizesArray, Address MappersArray,
+ unsigned NumberOfTargetItems)
: BasePointersArray(BasePointersArray), PointersArray(PointersArray),
- SizesArray(SizesArray), NumberOfTargetItems(NumberOfTargetItems) {}
+ SizesArray(SizesArray), MappersArray(MappersArray),
+ NumberOfTargetItems(NumberOfTargetItems) {}
};
void EmitOMPTargetTaskBasedDirective(const OMPExecutableDirective &S,
const RegionCodeGenTy &BodyGen,
Index: lib/CodeGen/CGStmtOpenMP.cpp
===================================================================
--- lib/CodeGen/CGStmtOpenMP.cpp
+++ lib/CodeGen/CGStmtOpenMP.cpp
@@ -3137,36 +3137,41 @@
VarDecl *BPVD = nullptr;
VarDecl *PVD = nullptr;
VarDecl *SVD = nullptr;
+ VarDecl *MVD = nullptr;
if (InputInfo.NumberOfTargetItems > 0) {
auto *CD = CapturedDecl::Create(
getContext(), getContext().getTranslationUnitDecl(), /*NumParams=*/0);
llvm::APInt ArrSize(/*numBits=*/32, InputInfo.NumberOfTargetItems);
- QualType BaseAndPointersType = getContext().getConstantArrayType(
+ QualType BaseAndPointerAndMapperType = getContext().getConstantArrayType(
getContext().VoidPtrTy, ArrSize, ArrayType::Normal,
/*IndexTypeQuals=*/0);
BPVD = createImplicitFirstprivateForType(
- getContext(), Data, BaseAndPointersType, CD, S.getBeginLoc());
+ getContext(), Data, BaseAndPointerAndMapperType, CD, S.getBeginLoc());
PVD = createImplicitFirstprivateForType(
- getContext(), Data, BaseAndPointersType, CD, S.getBeginLoc());
+ getContext(), Data, BaseAndPointerAndMapperType, CD, S.getBeginLoc());
QualType SizesType = getContext().getConstantArrayType(
getContext().getIntTypeForBitwidth(/*DestWidth=*/64, /*Signed=*/1),
ArrSize, ArrayType::Normal,
/*IndexTypeQuals=*/0);
SVD = createImplicitFirstprivateForType(getContext(), Data, SizesType, CD,
S.getBeginLoc());
+ MVD = createImplicitFirstprivateForType(
+ getContext(), Data, BaseAndPointerAndMapperType, CD, S.getBeginLoc());
TargetScope.addPrivate(
BPVD, [&InputInfo]() { return InputInfo.BasePointersArray; });
TargetScope.addPrivate(PVD,
[&InputInfo]() { return InputInfo.PointersArray; });
TargetScope.addPrivate(SVD,
[&InputInfo]() { return InputInfo.SizesArray; });
+ TargetScope.addPrivate(MVD,
+ [&InputInfo]() { return InputInfo.MappersArray; });
}
(void)TargetScope.Privatize();
// Build list of dependences.
for (const auto *C : S.getClausesOfKind<OMPDependClause>())
for (const Expr *IRef : C->varlists())
Data.Dependences.emplace_back(C->getDependencyKind(), IRef);
- auto &&CodeGen = [&Data, &S, CS, &BodyGen, BPVD, PVD, SVD,
+ auto &&CodeGen = [&Data, &S, CS, &BodyGen, BPVD, PVD, SVD, MVD,
&InputInfo](CodeGenFunction &CGF, PrePostActionTy &Action) {
// Set proper addresses for generated private copies.
OMPPrivateScope Scope(CGF);
@@ -3207,6 +3212,8 @@
CGF.GetAddrOfLocalVar(PVD), /*Index=*/0);
InputInfo.SizesArray = CGF.Builder.CreateConstArrayGEP(
CGF.GetAddrOfLocalVar(SVD), /*Index=*/0);
+ InputInfo.MappersArray = CGF.Builder.CreateConstArrayGEP(
+ CGF.GetAddrOfLocalVar(MVD), /*Index=*/0);
}
Action.Enter(CGF);
Index: lib/CodeGen/CGOpenMPRuntime.h
===================================================================
--- lib/CodeGen/CGOpenMPRuntime.h
+++ lib/CodeGen/CGOpenMPRuntime.h
@@ -817,6 +817,9 @@
/// Emit the function for the user defined mapper construct.
void emitUserDefinedMapper(const OMPDeclareMapperDecl *D,
CodeGenFunction *CGF = nullptr);
+ /// Get the function for the specified user-defined mapper, if any.
+ virtual llvm::Function *
+ getUserDefinedMapperFunc(const OMPDeclareMapperDecl *D);
/// Emits outlined function for the specified OpenMP parallel directive
/// \a D. This outlined function has type void(*)(kmp_int32 *ThreadID,
@@ -1510,6 +1513,8 @@
llvm::Value *SizesArray = nullptr;
/// The array of map types passed to the runtime library.
llvm::Value *MapTypesArray = nullptr;
+ /// The array of user-defined mappers passed to the runtime library.
+ llvm::Value *MappersArray = nullptr;
/// The total number of pointers passed to the runtime library.
unsigned NumberOfPtrs = 0u;
/// Map between the a declaration of a capture and the corresponding base
@@ -1525,12 +1530,13 @@
PointersArray = nullptr;
SizesArray = nullptr;
MapTypesArray = nullptr;
+ MappersArray = nullptr;
NumberOfPtrs = 0u;
}
/// Return true if the current target data information has valid arrays.
bool isValid() {
return BasePointersArray && PointersArray && SizesArray &&
- MapTypesArray && NumberOfPtrs;
+ MapTypesArray && MappersArray && NumberOfPtrs;
}
bool requiresDevicePointerInfo() { return RequiresDevicePointerInfo; }
};
Index: lib/CodeGen/CGOpenMPRuntime.cpp
===================================================================
--- lib/CodeGen/CGOpenMPRuntime.cpp
+++ lib/CodeGen/CGOpenMPRuntime.cpp
@@ -27,6 +27,7 @@
#include "llvm/Support/Format.h"
#include "llvm/Support/raw_ostream.h"
#include <cassert>
+#include <iostream>
using namespace clang;
using namespace CodeGen;
@@ -739,7 +740,7 @@
// *arg_types);
OMPRTL__tgt_target_data_begin_nowait,
// Call to void __tgt_target_data_end(int64_t device_id, int32_t arg_num,
- // void** args_base, void **args, size_t *arg_sizes, int64_t *arg_types);
+ // void** args_base, void **args, int64_t *arg_sizes, int64_t *arg_types);
OMPRTL__tgt_target_data_end,
// Call to void __tgt_target_data_end_nowait(int64_t device_id, int32_t
// arg_num, void** args_base, void **args, int64_t *arg_sizes, int64_t
@@ -752,6 +753,48 @@
// arg_num, void** args_base, void **args, int64_t *arg_sizes, int64_t
// *arg_types);
OMPRTL__tgt_target_data_update_nowait,
+ // Call to int32_t __tgt_target_mapper(int64_t device_id, void *host_ptr,
+ // int32_t arg_num, void** args_base, void **args, int64_t *arg_sizes, int64_t
+ // *arg_types, void **arg_mappers);
+ OMPRTL__tgt_target_mapper,
+ // Call to int32_t __tgt_target_nowait_mapper(int64_t device_id, void
+ // *host_ptr, int32_t arg_num, void** args_base, void **args, int64_t
+ // *arg_sizes, int64_t *arg_types, void **arg_mappers);
+ OMPRTL__tgt_target_nowait_mapper,
+ // Call to int32_t __tgt_target_teams_mapper(int64_t device_id, void
+ // *host_ptr, int32_t arg_num, void** args_base, void **args, int64_t
+ // *arg_sizes, int64_t *arg_types, void **arg_mappers, int32_t num_teams,
+ // int32_t thread_limit);
+ OMPRTL__tgt_target_teams_mapper,
+ // Call to int32_t __tgt_target_teams_nowait_mapper(int64_t device_id, void
+ // *host_ptr, int32_t arg_num, void** args_base, void **args, int64_t
+ // *arg_sizes, int64_t *arg_types, void **arg_mappers, int32_t num_teams,
+ // int32_t thread_limit);
+ OMPRTL__tgt_target_teams_nowait_mapper,
+ // Call to void __tgt_target_data_begin_mapper(int64_t device_id, int32_t
+ // arg_num, void** args_base, void **args, int64_t *arg_sizes, int64_t
+ // *arg_types, void **arg_mappers);
+ OMPRTL__tgt_target_data_begin_mapper,
+ // Call to void __tgt_target_data_begin_nowait_mapper(int64_t device_id,
+ // int32_t arg_num, void** args_base, void **args, int64_t *arg_sizes, int64_t
+ // *arg_types, void **arg_mappers);
+ OMPRTL__tgt_target_data_begin_nowait_mapper,
+ // Call to void __tgt_target_data_end_mapper(int64_t device_id, int32_t
+ // arg_num, void** args_base, void **args, int64_t *arg_sizes, int64_t
+ // *arg_types, void **arg_mappers);
+ OMPRTL__tgt_target_data_end_mapper,
+ // Call to void __tgt_target_data_end_nowait_mapper(int64_t device_id, int32_t
+ // arg_num, void** args_base, void **args, int64_t *arg_sizes, int64_t
+ // *arg_types, void **arg_mappers);
+ OMPRTL__tgt_target_data_end_nowait_mapper,
+ // Call to void __tgt_target_data_update_mapper(int64_t device_id, int32_t
+ // arg_num, void** args_base, void **args, int64_t *arg_sizes, int64_t
+ // *arg_types, void **arg_mappers);
+ OMPRTL__tgt_target_data_update_mapper,
+ // Call to void __tgt_target_data_update_nowait_mapper(int64_t device_id,
+ // int32_t arg_num, void** args_base, void **args, int64_t *arg_sizes, int64_t
+ // *arg_types, void **arg_mappers);
+ OMPRTL__tgt_target_data_update_nowait_mapper,
// Call to int64_t __tgt_mapper_num_components(void *rt_mapper_handle);
OMPRTL__tgt_mapper_num_components,
// Call to void __tgt_push_mapper_component(void *rt_mapper_handle, void
@@ -2470,6 +2513,179 @@
RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_target_data_update_nowait");
break;
}
+ case OMPRTL__tgt_target_mapper: {
+ // Build int32_t __tgt_target_mapper(int64_t device_id, void *host_ptr,
+ // int32_t arg_num, void** args_base, void **args, int64_t *arg_sizes,
+ // int64_t *arg_types, void **arg_mappers);
+ llvm::Type *TypeParams[] = {CGM.Int64Ty,
+ CGM.VoidPtrTy,
+ CGM.Int32Ty,
+ CGM.VoidPtrPtrTy,
+ CGM.VoidPtrPtrTy,
+ CGM.Int64Ty->getPointerTo(),
+ CGM.Int64Ty->getPointerTo(),
+ CGM.VoidPtrPtrTy};
+ auto *FnTy =
+ llvm::FunctionType::get(CGM.Int32Ty, TypeParams, /*isVarArg*/ false);
+ RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_target_mapper");
+ break;
+ }
+ case OMPRTL__tgt_target_nowait_mapper: {
+ // Build int32_t __tgt_target_nowait_mapper(int64_t device_id, void
+ // *host_ptr, int32_t arg_num, void** args_base, void **args, int64_t
+ // *arg_sizes, int64_t *arg_types, void **arg_mappers);
+ llvm::Type *TypeParams[] = {CGM.Int64Ty,
+ CGM.VoidPtrTy,
+ CGM.Int32Ty,
+ CGM.VoidPtrPtrTy,
+ CGM.VoidPtrPtrTy,
+ CGM.Int64Ty->getPointerTo(),
+ CGM.Int64Ty->getPointerTo(),
+ CGM.VoidPtrPtrTy};
+ auto *FnTy =
+ llvm::FunctionType::get(CGM.Int32Ty, TypeParams, /*isVarArg*/ false);
+ RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_target_nowait_mapper");
+ break;
+ }
+ case OMPRTL__tgt_target_teams_mapper: {
+ // Build int32_t __tgt_target_teams_mapper(int64_t device_id, void
+ // *host_ptr, int32_t arg_num, void** args_base, void **args, int64_t
+ // *arg_sizes, int64_t *arg_types, void **arg_mappers, int32_t num_teams,
+ // int32_t thread_limit);
+ llvm::Type *TypeParams[] = {CGM.Int64Ty,
+ CGM.VoidPtrTy,
+ CGM.Int32Ty,
+ CGM.VoidPtrPtrTy,
+ CGM.VoidPtrPtrTy,
+ CGM.Int64Ty->getPointerTo(),
+ CGM.Int64Ty->getPointerTo(),
+ CGM.VoidPtrPtrTy,
+ CGM.Int32Ty,
+ CGM.Int32Ty};
+ auto *FnTy =
+ llvm::FunctionType::get(CGM.Int32Ty, TypeParams, /*isVarArg*/ false);
+ RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_target_teams_mapper");
+ break;
+ }
+ case OMPRTL__tgt_target_teams_nowait_mapper: {
+ // Build int32_t __tgt_target_teams_nowait_mapper(int64_t device_id, void
+ // *host_ptr, int32_t arg_num, void** args_base, void **args, int64_t
+ // *arg_sizes, int64_t *arg_types, void **arg_mappers, int32_t num_teams,
+ // int32_t thread_limit);
+ llvm::Type *TypeParams[] = {CGM.Int64Ty,
+ CGM.VoidPtrTy,
+ CGM.Int32Ty,
+ CGM.VoidPtrPtrTy,
+ CGM.VoidPtrPtrTy,
+ CGM.Int64Ty->getPointerTo(),
+ CGM.Int64Ty->getPointerTo(),
+ CGM.VoidPtrPtrTy,
+ CGM.Int32Ty,
+ CGM.Int32Ty};
+ auto *FnTy =
+ llvm::FunctionType::get(CGM.Int32Ty, TypeParams, /*isVarArg*/ false);
+ RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_target_teams_nowait_mapper");
+ break;
+ }
+ case OMPRTL__tgt_target_data_begin_mapper: {
+ // Build void __tgt_target_data_begin_mapper(int64_t device_id, int32_t
+ // arg_num, void **args_base, void **args, int64_t *arg_sizes, int64_t
+ // *arg_types, void **arg_mappers);
+ llvm::Type *TypeParams[] = {CGM.Int64Ty,
+ CGM.Int32Ty,
+ CGM.VoidPtrPtrTy,
+ CGM.VoidPtrPtrTy,
+ CGM.Int64Ty->getPointerTo(),
+ CGM.Int64Ty->getPointerTo(),
+ CGM.VoidPtrPtrTy};
+ auto *FnTy =
+ llvm::FunctionType::get(CGM.VoidTy, TypeParams, /*isVarArg*/ false);
+ RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_target_data_begin_mapper");
+ break;
+ }
+ case OMPRTL__tgt_target_data_begin_nowait_mapper: {
+ // Build void __tgt_target_data_begin_nowait_mapper(int64_t device_id,
+ // int32_t arg_num, void** args_base, void **args, int64_t *arg_sizes,
+ // int64_t *arg_types, void **arg_mappers);
+ llvm::Type *TypeParams[] = {CGM.Int64Ty,
+ CGM.Int32Ty,
+ CGM.VoidPtrPtrTy,
+ CGM.VoidPtrPtrTy,
+ CGM.Int64Ty->getPointerTo(),
+ CGM.Int64Ty->getPointerTo(),
+ CGM.VoidPtrPtrTy};
+ auto *FnTy =
+ llvm::FunctionType::get(CGM.VoidTy, TypeParams, /*isVarArg=*/false);
+ RTLFn = CGM.CreateRuntimeFunction(FnTy,
+ "__tgt_target_data_begin_nowait_mapper");
+ break;
+ }
+ case OMPRTL__tgt_target_data_end_mapper: {
+ // Build void __tgt_target_data_end_mapper(int64_t device_id, int32_t
+ // arg_num, void** args_base, void **args, int64_t *arg_sizes, int64_t
+ // *arg_types, void **arg_mappers);
+ llvm::Type *TypeParams[] = {CGM.Int64Ty,
+ CGM.Int32Ty,
+ CGM.VoidPtrPtrTy,
+ CGM.VoidPtrPtrTy,
+ CGM.Int64Ty->getPointerTo(),
+ CGM.Int64Ty->getPointerTo(),
+ CGM.VoidPtrPtrTy};
+ auto *FnTy =
+ llvm::FunctionType::get(CGM.VoidTy, TypeParams, /*isVarArg*/ false);
+ RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_target_data_end_mapper");
+ break;
+ }
+ case OMPRTL__tgt_target_data_end_nowait_mapper: {
+ // Build void __tgt_target_data_end_nowait_mapper(int64_t device_id, int32_t
+ // arg_num, void** args_base, void **args, int64_t *arg_sizes, int64_t
+ // *arg_types, void **arg_mappers);
+ llvm::Type *TypeParams[] = {CGM.Int64Ty,
+ CGM.Int32Ty,
+ CGM.VoidPtrPtrTy,
+ CGM.VoidPtrPtrTy,
+ CGM.Int64Ty->getPointerTo(),
+ CGM.Int64Ty->getPointerTo(),
+ CGM.VoidPtrPtrTy};
+ auto *FnTy =
+ llvm::FunctionType::get(CGM.VoidTy, TypeParams, /*isVarArg=*/false);
+ RTLFn =
+ CGM.CreateRuntimeFunction(FnTy, "__tgt_target_data_end_nowait_mapper");
+ break;
+ }
+ case OMPRTL__tgt_target_data_update_mapper: {
+ // Build void __tgt_target_data_update_mapper(int64_t device_id, int32_t
+ // arg_num, void** args_base, void **args, int64_t *arg_sizes, int64_t
+ // *arg_types, void **arg_mappers);
+ llvm::Type *TypeParams[] = {CGM.Int64Ty,
+ CGM.Int32Ty,
+ CGM.VoidPtrPtrTy,
+ CGM.VoidPtrPtrTy,
+ CGM.Int64Ty->getPointerTo(),
+ CGM.Int64Ty->getPointerTo(),
+ CGM.VoidPtrPtrTy};
+ auto *FnTy =
+ llvm::FunctionType::get(CGM.VoidTy, TypeParams, /*isVarArg*/ false);
+ RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_target_data_update_mapper");
+ break;
+ }
+ case OMPRTL__tgt_target_data_update_nowait_mapper: {
+ // Build void __tgt_target_data_update_nowait_mapper(int64_t device_id,
+ // int32_t arg_num, void** args_base, void **args, int64_t *arg_sizes,
+ // int64_t *arg_types, void **arg_mappers);
+ llvm::Type *TypeParams[] = {CGM.Int64Ty,
+ CGM.Int32Ty,
+ CGM.VoidPtrPtrTy,
+ CGM.VoidPtrPtrTy,
+ CGM.Int64Ty->getPointerTo(),
+ CGM.Int64Ty->getPointerTo(),
+ CGM.VoidPtrPtrTy};
+ auto *FnTy =
+ llvm::FunctionType::get(CGM.VoidTy, TypeParams, /*isVarArg=*/false);
+ RTLFn = CGM.CreateRuntimeFunction(FnTy,
+ "__tgt_target_data_update_nowait_mapper");
+ break;
+ }
case OMPRTL__tgt_mapper_num_components: {
// Build int64_t __tgt_mapper_num_components(void *rt_mapper_handle);
llvm::Type *TypeParams[] = {CGM.VoidPtrTy};
@@ -7154,6 +7370,7 @@
using MapBaseValuesArrayTy = SmallVector<BasePointerInfo, 4>;
using MapValuesArrayTy = SmallVector<llvm::Value *, 4>;
using MapFlagsArrayTy = SmallVector<OpenMPOffloadMappingFlags, 4>;
+ using MapMappersArrayTy = SmallVector<const ValueDecl *, 4>;
/// Map between a struct and the its lowest & highest elements which have been
/// mapped.
@@ -7175,15 +7392,17 @@
ArrayRef<OpenMPMapModifierKind> MapModifiers;
bool ReturnDevicePointer = false;
bool IsImplicit = false;
+ const ValueDecl *Mapper = nullptr;
MapInfo() = default;
MapInfo(
OMPClauseMappableExprCommon::MappableExprComponentListRef Components,
OpenMPMapClauseKind MapType,
- ArrayRef<OpenMPMapModifierKind> MapModifiers,
- bool ReturnDevicePointer, bool IsImplicit)
+ ArrayRef<OpenMPMapModifierKind> MapModifiers, bool ReturnDevicePointer,
+ bool IsImplicit, const ValueDecl *Mapper = nullptr)
: Components(Components), MapType(MapType), MapModifiers(MapModifiers),
- ReturnDevicePointer(ReturnDevicePointer), IsImplicit(IsImplicit) {}
+ ReturnDevicePointer(ReturnDevicePointer), IsImplicit(IsImplicit),
+ Mapper(Mapper) {}
};
/// If use_device_ptr is used on a pointer which is a struct member and there
@@ -7218,7 +7437,7 @@
SmallVector<OMPClauseMappableExprCommon::MappableExprComponentListRef, 4>>
DevPointersMap;
- llvm::Value *getExprTypeSize(const Expr *E) const {
+ llvm::Value *getExprTypeSize(const Expr *E, bool hasMapper) const {
QualType ExprTy = E->getType().getCanonicalType();
// Reference types are ignored for mapping purposes.
@@ -7235,8 +7454,14 @@
// If there is no length associated with the expression, that means we
// are using the whole length of the base.
- if (!OAE->getLength() && OAE->getColonLoc().isValid())
- return CGF.getTypeSize(BaseTy);
+ if (!OAE->getLength() && OAE->getColonLoc().isValid()) {
+ // In case that a user-defined mapper is attached, its size is the
+ // number of array elements instead of the number of total bytes.
+ if (hasMapper)
+ return CGF.Builder.getInt64(1);
+ else
+ return CGF.getTypeSize(BaseTy);
+ }
llvm::Value *ElemSize;
if (const auto *PTy = BaseTy->getAs<PointerType>()) {
@@ -7249,15 +7474,31 @@
// If we don't have a length at this point, that is because we have an
// array section with a single element.
- if (!OAE->getLength())
- return ElemSize;
+ if (!OAE->getLength()) {
+ // In case that a user-defined mapper is attached, its size is the
+ // number of array elements instead of the number of total bytes.
+ if (hasMapper)
+ return CGF.Builder.getInt64(1);
+ else
+ return ElemSize;
+ }
llvm::Value *LengthVal = CGF.EmitScalarExpr(OAE->getLength());
LengthVal =
CGF.Builder.CreateIntCast(LengthVal, CGF.SizeTy, /*isSigned=*/false);
- return CGF.Builder.CreateNUWMul(LengthVal, ElemSize);
+ // In case that a user-defined mapper is attached, its size is the
+ // number of array elements instead of the number of total bytes.
+ if (hasMapper)
+ return LengthVal;
+ else
+ return CGF.Builder.CreateNUWMul(LengthVal, ElemSize);
}
- return CGF.getTypeSize(ExprTy);
+ // In case that a user-defined mapper is attached, its size is the
+ // number of array elements instead of the number of total bytes.
+ if (hasMapper)
+ return CGF.Builder.getInt64(1);
+ else
+ return CGF.getTypeSize(ExprTy);
}
/// Return the corresponding bits for a given map clause modifier. Add
@@ -7344,18 +7585,18 @@
return ConstLength.getSExtValue() != 1;
}
- /// Generate the base pointers, section pointers, sizes and map type
- /// bits for the provided map type, map modifier, and expression components.
- /// \a IsFirstComponent should be set to true if the provided set of
- /// components is the first associated with a capture.
+ /// Generate the base pointers, section pointers, sizes, map type bits, and
+ /// mappers for the provided map type, map modifier, and expression
+ /// components. \a IsFirstComponent should be set to true if the provided set
+ /// of components is the first associated with a capture.
void generateInfoForComponentList(
- OpenMPMapClauseKind MapType,
- ArrayRef<OpenMPMapModifierKind> MapModifiers,
+ OpenMPMapClauseKind MapType, ArrayRef<OpenMPMapModifierKind> MapModifiers,
OMPClauseMappableExprCommon::MappableExprComponentListRef Components,
MapBaseValuesArrayTy &BasePointers, MapValuesArrayTy &Pointers,
MapValuesArrayTy &Sizes, MapFlagsArrayTy &Types,
- StructRangeInfoTy &PartialStruct, bool IsFirstComponentList,
- bool IsImplicit,
+ MapMappersArrayTy &Mappers, StructRangeInfoTy &PartialStruct,
+ bool IsFirstComponentList, bool IsImplicit,
+ const ValueDecl *Mapper = nullptr,
ArrayRef<OMPClauseMappableExprCommon::MappableExprComponentListRef>
OverlappedElements = llvm::None) const {
// The following summarizes what has to be generated for each map and the
@@ -7699,6 +7940,7 @@
Sizes.push_back(CGF.Builder.CreateIntCast(Size, CGF.Int64Ty,
/*isSigned=*/true));
Types.push_back(Flags);
+ Mappers.push_back(nullptr);
LB = CGF.Builder.CreateConstGEP(ComponentLB, 1);
}
BasePointers.push_back(BP.getPointer());
@@ -7710,14 +7952,29 @@
Sizes.push_back(
CGF.Builder.CreateIntCast(Size, CGF.Int64Ty, /*isSigned=*/true));
Types.push_back(Flags);
+ Mappers.push_back(nullptr);
break;
}
- llvm::Value *Size = getExprTypeSize(I->getAssociatedExpression());
if (!IsMemberPointer) {
BasePointers.push_back(BP.getPointer());
Pointers.push_back(LB.getPointer());
- Sizes.push_back(
- CGF.Builder.CreateIntCast(Size, CGF.Int64Ty, /*isSigned=*/true));
+
+ // If Mapper is valid, the last component inherits the mapper.
+ bool hasMapper = Mapper && Next == CE;
+ llvm::Value *Size =
+ getExprTypeSize(I->getAssociatedExpression(), hasMapper);
+ Sizes.push_back(CGF.Builder.CreateIntCast(Size, CGF.Int64Ty,
+ /*isSigned=*/true));
+ if (hasMapper)
+ Mappers.push_back(Mapper);
+ else
+ Mappers.push_back(nullptr);
+ std::cerr << "HH " << "\n";
+ if (hasMapper) {
+ std::cerr << "AM: " << Mappers.size() << " ";
+ I->getAssociatedExpression()->dump();
+ //Size->dump();
+ }
// We need to add a pointer flag for each map that comes from the
// same expression except for the first one. We also need to signal
@@ -7898,7 +8155,7 @@
// Extract device pointer clause information.
for (const auto *C : Dir.getClausesOfKind<OMPIsDevicePtrClause>())
for (auto L : C->component_lists())
- DevPointersMap[L.first].push_back(L.second);
+ DevPointersMap[std::get<0>(L)].push_back(std::get<1>(L));
}
/// Constructor for the declare mapper directive.
@@ -7910,13 +8167,16 @@
/// individual struct members.
void emitCombinedEntry(MapBaseValuesArrayTy &BasePointers,
MapValuesArrayTy &Pointers, MapValuesArrayTy &Sizes,
- MapFlagsArrayTy &Types, MapFlagsArrayTy &CurTypes,
+ MapFlagsArrayTy &Types, MapMappersArrayTy &Mappers,
+ MapFlagsArrayTy &CurTypes,
const StructRangeInfoTy &PartialStruct) const {
// Base is the base of the struct
BasePointers.push_back(PartialStruct.Base.getPointer());
// Pointer is the address of the lowest element
llvm::Value *LB = PartialStruct.LowestElem.second.getPointer();
Pointers.push_back(LB);
+ // There should not be a mapper for a combined entry.
+ Mappers.push_back(nullptr);
// Size is (addr of {highest+1} element) - (addr of lowest element)
llvm::Value *HB = PartialStruct.HighestElem.second.getPointer();
llvm::Value *HAddr = CGF.Builder.CreateConstGEP1_32(HB, /*Idx0=*/1);
@@ -7940,13 +8200,14 @@
setCorrectMemberOfFlag(M, MemberOfFlag);
}
- /// Generate all the base pointers, section pointers, sizes and map
- /// types for the extracted mappable expressions. Also, for each item that
+ /// Generate all the base pointers, section pointers, sizes, map types, and
+ /// mappers for the extracted mappable expressions. Also, for each item that
/// relates with a device pointer, a pair of the relevant declaration and
/// index where it occurs is appended to the device pointers info array.
void generateAllInfo(MapBaseValuesArrayTy &BasePointers,
MapValuesArrayTy &Pointers, MapValuesArrayTy &Sizes,
- MapFlagsArrayTy &Types) const {
+ MapFlagsArrayTy &Types,
+ MapMappersArrayTy &Mappers) const {
// We have to process the component lists that relate with the same
// declaration in a single chunk so that we can generate the map flags
// correctly. Therefore, we organize all lists in a map.
@@ -7954,35 +8215,37 @@
// Helper function to fill the information map for the different supported
// clauses.
- auto &&InfoGen = [&Info](
- const ValueDecl *D,
- OMPClauseMappableExprCommon::MappableExprComponentListRef L,
- OpenMPMapClauseKind MapType,
- ArrayRef<OpenMPMapModifierKind> MapModifiers,
- bool ReturnDevicePointer, bool IsImplicit) {
- const ValueDecl *VD =
- D ? cast<ValueDecl>(D->getCanonicalDecl()) : nullptr;
- Info[VD].emplace_back(L, MapType, MapModifiers, ReturnDevicePointer,
- IsImplicit);
- };
+ auto &&InfoGen =
+ [&Info](const ValueDecl *D,
+ OMPClauseMappableExprCommon::MappableExprComponentListRef L,
+ OpenMPMapClauseKind MapType,
+ ArrayRef<OpenMPMapModifierKind> MapModifiers,
+ bool ReturnDevicePointer, bool IsImplicit,
+ const ValueDecl *Mapper) {
+ const ValueDecl *VD =
+ D ? cast<ValueDecl>(D->getCanonicalDecl()) : nullptr;
+ Info[VD].emplace_back(L, MapType, MapModifiers, ReturnDevicePointer,
+ IsImplicit, Mapper);
+ };
assert(CurDir.is<const OMPExecutableDirective *>() &&
"Expect a executable directive");
const auto *CurExecDir = CurDir.get<const OMPExecutableDirective *>();
for (const auto *C : CurExecDir->getClausesOfKind<OMPMapClause>())
for (const auto &L : C->component_lists()) {
- InfoGen(L.first, L.second, C->getMapType(), C->getMapTypeModifiers(),
- /*ReturnDevicePointer=*/false, C->isImplicit());
+ InfoGen(std::get<0>(L), std::get<1>(L), C->getMapType(),
+ C->getMapTypeModifiers(), /*ReturnDevicePointer=*/false,
+ C->isImplicit(), std::get<2>(L));
}
for (const auto *C : CurExecDir->getClausesOfKind<OMPToClause>())
for (const auto &L : C->component_lists()) {
- InfoGen(L.first, L.second, OMPC_MAP_to, llvm::None,
- /*ReturnDevicePointer=*/false, C->isImplicit());
+ InfoGen(std::get<0>(L), std::get<1>(L), OMPC_MAP_to, llvm::None,
+ /*ReturnDevicePointer=*/false, C->isImplicit(), std::get<2>(L));
}
for (const auto *C : CurExecDir->getClausesOfKind<OMPFromClause>())
for (const auto &L : C->component_lists()) {
- InfoGen(L.first, L.second, OMPC_MAP_from, llvm::None,
- /*ReturnDevicePointer=*/false, C->isImplicit());
+ InfoGen(std::get<0>(L), std::get<1>(L), OMPC_MAP_from, llvm::None,
+ /*ReturnDevicePointer=*/false, C->isImplicit(), std::get<2>(L));
}
// Look at the use_device_ptr clause information and mark the existing map
@@ -7997,10 +8260,13 @@
for (const auto *C :
CurExecDir->getClausesOfKind<OMPUseDevicePtrClause>()) {
for (const auto &L : C->component_lists()) {
- assert(!L.second.empty() && "Not expecting empty list of components!");
- const ValueDecl *VD = L.second.back().getAssociatedDeclaration();
+ OMPClauseMappableExprCommon::MappableExprComponentListRef Components =
+ std::get<1>(L);
+ assert(!Components.empty() &&
+ "Not expecting empty list of components!");
+ const ValueDecl *VD = Components.back().getAssociatedDeclaration();
VD = cast<ValueDecl>(VD->getCanonicalDecl());
- const Expr *IE = L.second.back().getAssociatedExpression();
+ const Expr *IE = Components.back().getAssociatedExpression();
// If the first component is a member expression, we have to look into
// 'this', which maps to null in the map of map information. Otherwise
// look directly for the information.
@@ -8032,8 +8298,8 @@
// Nonetheless, generateInfoForComponentList must be called to take
// the pointer into account for the calculation of the range of the
// partial struct.
- InfoGen(nullptr, L.second, OMPC_MAP_unknown, llvm::None,
- /*ReturnDevicePointer=*/false, C->isImplicit());
+ InfoGen(nullptr, Components, OMPC_MAP_unknown, llvm::None,
+ /*ReturnDevicePointer=*/false, C->isImplicit(), nullptr);
DeferredInfo[nullptr].emplace_back(IE, VD);
} else {
llvm::Value *Ptr =
@@ -8042,6 +8308,7 @@
Pointers.push_back(Ptr);
Sizes.push_back(llvm::Constant::getNullValue(CGF.Int64Ty));
Types.push_back(OMP_MAP_RETURN_PARAM | OMP_MAP_TARGET_PARAM);
+ Mappers.push_back(nullptr);
}
}
}
@@ -8056,6 +8323,7 @@
MapValuesArrayTy CurPointers;
MapValuesArrayTy CurSizes;
MapFlagsArrayTy CurTypes;
+ MapMappersArrayTy CurMappers;
StructRangeInfoTy PartialStruct;
for (const MapInfo &L : M.second) {
@@ -8064,10 +8332,10 @@
// Remember the current base pointer index.
unsigned CurrentBasePointersIdx = CurBasePointers.size();
- generateInfoForComponentList(L.MapType, L.MapModifiers, L.Components,
- CurBasePointers, CurPointers, CurSizes,
- CurTypes, PartialStruct,
- IsFirstComponentList, L.IsImplicit);
+ generateInfoForComponentList(
+ L.MapType, L.MapModifiers, L.Components, CurBasePointers,
+ CurPointers, CurSizes, CurTypes, CurMappers, PartialStruct,
+ IsFirstComponentList, L.IsImplicit, L.Mapper);
// If this entry relates with a device pointer, set the relevant
// declaration and add the 'return pointer' flag.
@@ -8102,29 +8370,31 @@
// correct value of MEMBER_OF.
CurTypes.push_back(OMP_MAP_PTR_AND_OBJ | OMP_MAP_RETURN_PARAM |
OMP_MAP_MEMBER_OF);
+ CurMappers.push_back(nullptr);
}
}
// If there is an entry in PartialStruct it means we have a struct with
// individual members mapped. Emit an extra combined entry.
if (PartialStruct.Base.isValid())
- emitCombinedEntry(BasePointers, Pointers, Sizes, Types, CurTypes,
- PartialStruct);
+ emitCombinedEntry(BasePointers, Pointers, Sizes, Types, Mappers,
+ CurTypes, PartialStruct);
// We need to append the results of this capture to what we already have.
BasePointers.append(CurBasePointers.begin(), CurBasePointers.end());
Pointers.append(CurPointers.begin(), CurPointers.end());
Sizes.append(CurSizes.begin(), CurSizes.end());
Types.append(CurTypes.begin(), CurTypes.end());
+ Mappers.append(CurMappers.begin(), CurMappers.end());
}
}
- /// Generate all the base pointers, section pointers, sizes and map types for
- /// the extracted map clauses of user-defined mapper.
+ /// Generate all the base pointers, section pointers, sizes, map types, and
+ /// mappers for the extracted map clauses of user-defined mapper.
void generateAllInfoForMapper(MapBaseValuesArrayTy &BasePointers,
MapValuesArrayTy &Pointers,
- MapValuesArrayTy &Sizes,
- MapFlagsArrayTy &Types) const {
+ MapValuesArrayTy &Sizes, MapFlagsArrayTy &Types,
+ MapMappersArrayTy &Mappers) const {
assert(CurDir.is<const OMPDeclareMapperDecl *>() &&
"Expect a declare mapper directive");
const auto *CurMapperDir = CurDir.get<const OMPDeclareMapperDecl *>();
@@ -8133,25 +8403,17 @@
// correctly. Therefore, we organize all lists in a map.
llvm::MapVector<const ValueDecl *, SmallVector<MapInfo, 8>> Info;
- // Helper function to fill the information map for the different supported
- // clauses.
- auto &&InfoGen = [&Info](
- const ValueDecl *D,
- OMPClauseMappableExprCommon::MappableExprComponentListRef L,
- OpenMPMapClauseKind MapType,
- ArrayRef<OpenMPMapModifierKind> MapModifiers,
- bool ReturnDevicePointer, bool IsImplicit) {
- const ValueDecl *VD =
- D ? cast<ValueDecl>(D->getCanonicalDecl()) : nullptr;
- Info[VD].emplace_back(L, MapType, MapModifiers, ReturnDevicePointer,
- IsImplicit);
- };
-
+ // Fill the information map for map clauses.
for (const auto *C : CurMapperDir->clauselists()) {
- const auto *MC = cast<OMPMapClause>(C);
+ const auto *MC = cast<const OMPMapClause>(C);
for (const auto &L : MC->component_lists()) {
- InfoGen(L.first, L.second, MC->getMapType(), MC->getMapTypeModifiers(),
- /*ReturnDevicePointer=*/false, MC->isImplicit());
+ const ValueDecl *VD =
+ std::get<0>(L) ? cast<ValueDecl>(std::get<0>(L)->getCanonicalDecl())
+ : nullptr;
+ // Get the corresponding user-defined mapper.
+ Info[VD].emplace_back(
+ std::get<1>(L), MC->getMapType(), MC->getMapTypeModifiers(),
+ /*ReturnDevicePointer=*/false, MC->isImplicit(), std::get<2>(L));
}
}
@@ -8165,29 +8427,31 @@
MapValuesArrayTy CurPointers;
MapValuesArrayTy CurSizes;
MapFlagsArrayTy CurTypes;
+ MapMappersArrayTy CurMappers;
StructRangeInfoTy PartialStruct;
for (const MapInfo &L : M.second) {
assert(!L.Components.empty() &&
"Not expecting declaration with no component lists.");
- generateInfoForComponentList(L.MapType, L.MapModifiers, L.Components,
- CurBasePointers, CurPointers, CurSizes,
- CurTypes, PartialStruct,
- IsFirstComponentList, L.IsImplicit);
+ generateInfoForComponentList(
+ L.MapType, L.MapModifiers, L.Components, CurBasePointers,
+ CurPointers, CurSizes, CurTypes, CurMappers, PartialStruct,
+ IsFirstComponentList, L.IsImplicit, L.Mapper);
IsFirstComponentList = false;
}
// If there is an entry in PartialStruct it means we have a struct with
// individual members mapped. Emit an extra combined entry.
if (PartialStruct.Base.isValid())
- emitCombinedEntry(BasePointers, Pointers, Sizes, Types, CurTypes,
- PartialStruct);
+ emitCombinedEntry(BasePointers, Pointers, Sizes, Types, Mappers,
+ CurTypes, PartialStruct);
// We need to append the results of this capture to what we already have.
BasePointers.append(CurBasePointers.begin(), CurBasePointers.end());
Pointers.append(CurPointers.begin(), CurPointers.end());
Sizes.append(CurSizes.begin(), CurSizes.end());
Types.append(CurTypes.begin(), CurTypes.end());
+ Mappers.append(CurMappers.begin(), CurMappers.end());
}
}
@@ -8195,7 +8459,7 @@
void generateInfoForLambdaCaptures(
const ValueDecl *VD, llvm::Value *Arg, MapBaseValuesArrayTy &BasePointers,
MapValuesArrayTy &Pointers, MapValuesArrayTy &Sizes,
- MapFlagsArrayTy &Types,
+ MapFlagsArrayTy &Types, MapMappersArrayTy &Mappers,
llvm::DenseMap<llvm::Value *, llvm::Value *> &LambdaPointers) const {
const auto *RD = VD->getType()
.getCanonicalType()
@@ -8221,6 +8485,7 @@
CGF.Int64Ty, /*isSigned=*/true));
Types.push_back(OMP_MAP_PTR_AND_OBJ | OMP_MAP_LITERAL |
OMP_MAP_MEMBER_OF | OMP_MAP_IMPLICIT);
+ Mappers.push_back(nullptr);
}
for (const LambdaCapture &LC : RD->captures()) {
if (!LC.capturesVariable())
@@ -8249,6 +8514,7 @@
}
Types.push_back(OMP_MAP_PTR_AND_OBJ | OMP_MAP_LITERAL |
OMP_MAP_MEMBER_OF | OMP_MAP_IMPLICIT);
+ Mappers.push_back(nullptr);
}
}
@@ -8281,13 +8547,14 @@
}
}
- /// Generate the base pointers, section pointers, sizes and map types
- /// associated to a given capture.
+ /// Generate the base pointers, section pointers, sizes, map types, and
+ /// mappers associated to a given capture.
void generateInfoForCapture(const CapturedStmt::Capture *Cap,
llvm::Value *Arg,
MapBaseValuesArrayTy &BasePointers,
MapValuesArrayTy &Pointers,
MapValuesArrayTy &Sizes, MapFlagsArrayTy &Types,
+ MapMappersArrayTy &Mappers,
StructRangeInfoTy &PartialStruct) const {
assert(!Cap->capturesVariableArrayType() &&
"Not expecting to generate map info for a variable array type!");
@@ -8307,25 +8574,30 @@
CGF.Builder.CreateIntCast(CGF.getTypeSize(CGF.getContext().VoidPtrTy),
CGF.Int64Ty, /*isSigned=*/true));
Types.push_back(OMP_MAP_LITERAL | OMP_MAP_TARGET_PARAM);
+ Mappers.push_back(nullptr);
return;
}
using MapData =
std::tuple<OMPClauseMappableExprCommon::MappableExprComponentListRef,
- OpenMPMapClauseKind, ArrayRef<OpenMPMapModifierKind>, bool>;
+ OpenMPMapClauseKind, ArrayRef<OpenMPMapModifierKind>, bool,
+ const ValueDecl *>;
SmallVector<MapData, 4> DeclComponentLists;
assert(CurDir.is<const OMPExecutableDirective *>() &&
"Expect a executable directive");
const auto *CurExecDir = CurDir.get<const OMPExecutableDirective *>();
+ std::cerr << "C " << "\n";
for (const auto *C : CurExecDir->getClausesOfKind<OMPMapClause>()) {
for (const auto &L : C->decl_component_lists(VD)) {
- assert(L.first == VD &&
- "We got information for the wrong declaration??");
- assert(!L.second.empty() &&
+ const ValueDecl *VDecl, *Mapper;
+ OMPClauseMappableExprCommon::MappableExprComponentListRef Components;
+ std::tie(VDecl, Components, Mapper) = L;
+ assert(VDecl == VD && "We got information for the wrong declaration??");
+ assert(!Components.empty() &&
"Not expecting declaration with no component lists.");
- DeclComponentLists.emplace_back(L.second, C->getMapType(),
+ DeclComponentLists.emplace_back(Components, C->getMapType(),
C->getMapTypeModifiers(),
- C->isImplicit());
+ C->isImplicit(), Mapper);
}
}
@@ -8342,11 +8614,12 @@
OpenMPMapClauseKind MapType;
ArrayRef<OpenMPMapModifierKind> MapModifiers;
bool IsImplicit;
- std::tie(Components, MapType, MapModifiers, IsImplicit) = L;
+ const ValueDecl *Mapper;
+ std::tie(Components, MapType, MapModifiers, IsImplicit, Mapper) = L;
++Count;
for (const MapData &L1 : makeArrayRef(DeclComponentLists).slice(Count)) {
OMPClauseMappableExprCommon::MappableExprComponentListRef Components1;
- std::tie(Components1, MapType, MapModifiers, IsImplicit) = L1;
+ std::tie(Components1, MapType, MapModifiers, IsImplicit, Mapper) = L1;
auto CI = Components.rbegin();
auto CE = Components.rend();
auto SI = Components1.rbegin();
@@ -8432,14 +8705,15 @@
OpenMPMapClauseKind MapType;
ArrayRef<OpenMPMapModifierKind> MapModifiers;
bool IsImplicit;
- std::tie(Components, MapType, MapModifiers, IsImplicit) = L;
+ const ValueDecl *Mapper;
+ std::tie(Components, MapType, MapModifiers, IsImplicit, Mapper) = L;
ArrayRef<OMPClauseMappableExprCommon::MappableExprComponentListRef>
OverlappedComponents = Pair.getSecond();
bool IsFirstComponentList = true;
generateInfoForComponentList(MapType, MapModifiers, Components,
BasePointers, Pointers, Sizes, Types,
- PartialStruct, IsFirstComponentList,
- IsImplicit, OverlappedComponents);
+ Mappers, PartialStruct, IsFirstComponentList,
+ IsImplicit, Mapper, OverlappedComponents);
}
// Go through other elements without overlapped elements.
bool IsFirstComponentList = OverlappedData.empty();
@@ -8448,23 +8722,25 @@
OpenMPMapClauseKind MapType;
ArrayRef<OpenMPMapModifierKind> MapModifiers;
bool IsImplicit;
- std::tie(Components, MapType, MapModifiers, IsImplicit) = L;
+ const ValueDecl *Mapper;
+ std::tie(Components, MapType, MapModifiers, IsImplicit, Mapper) = L;
auto It = OverlappedData.find(&L);
if (It == OverlappedData.end())
generateInfoForComponentList(MapType, MapModifiers, Components,
BasePointers, Pointers, Sizes, Types,
- PartialStruct, IsFirstComponentList,
- IsImplicit);
+ Mappers, PartialStruct,
+ IsFirstComponentList, IsImplicit, Mapper);
IsFirstComponentList = false;
}
}
- /// Generate the base pointers, section pointers, sizes and map types
- /// associated with the declare target link variables.
+ /// Generate the base pointers, section pointers, sizes, map types, and
+ /// mappers associated with the declare target link variables.
void generateInfoForDeclareTargetLink(MapBaseValuesArrayTy &BasePointers,
MapValuesArrayTy &Pointers,
MapValuesArrayTy &Sizes,
- MapFlagsArrayTy &Types) const {
+ MapFlagsArrayTy &Types,
+ MapMappersArrayTy &Mappers) const {
assert(CurDir.is<const OMPExecutableDirective *>() &&
"Expect a executable directive");
const auto *CurExecDir = CurDir.get<const OMPExecutableDirective *>();
@@ -8472,9 +8748,9 @@
// but "declare target link" global variables.
for (const auto *C : CurExecDir->getClausesOfKind<OMPMapClause>()) {
for (const auto &L : C->component_lists()) {
- if (!L.first)
+ if (!std::get<0>(L))
continue;
- const auto *VD = dyn_cast<VarDecl>(L.first);
+ const auto *VD = dyn_cast<VarDecl>(std::get<0>(L));
if (!VD)
continue;
llvm::Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
@@ -8484,8 +8760,8 @@
continue;
StructRangeInfoTy PartialStruct;
generateInfoForComponentList(
- C->getMapType(), C->getMapTypeModifiers(), L.second, BasePointers,
- Pointers, Sizes, Types, PartialStruct,
+ C->getMapType(), C->getMapTypeModifiers(), std::get<1>(L),
+ BasePointers, Pointers, Sizes, Types, Mappers, PartialStruct,
/*IsFirstComponentList=*/true, C->isImplicit());
assert(!PartialStruct.Base.isValid() &&
"No partial structs for declare target link expected.");
@@ -8500,7 +8776,8 @@
MapBaseValuesArrayTy &CurBasePointers,
MapValuesArrayTy &CurPointers,
MapValuesArrayTy &CurSizes,
- MapFlagsArrayTy &CurMapTypes) const {
+ MapFlagsArrayTy &CurMapTypes,
+ MapMappersArrayTy &CurMappers) const {
bool IsImplicit = true;
// Do the default mapping.
if (CI.capturesThis()) {
@@ -8575,6 +8852,9 @@
// Add flag stating this is an implicit map.
if (IsImplicit)
CurMapTypes.back() |= OMP_MAP_IMPLICIT;
+
+ // No user-defined mapper for default mapping.
+ CurMappers.push_back(nullptr);
}
};
} // anonymous namespace
@@ -8588,6 +8868,7 @@
MappableExprsHandler::MapValuesArrayTy &Pointers,
MappableExprsHandler::MapValuesArrayTy &Sizes,
MappableExprsHandler::MapFlagsArrayTy &MapTypes,
+ MappableExprsHandler::MapMappersArrayTy &Mappers,
CGOpenMPRuntime::TargetDataInfo &Info) {
CodeGenModule &CGM = CGF.CGM;
ASTContext &Ctx = CGF.getContext();
@@ -8615,6 +8896,8 @@
CGF.CreateMemTemp(PointerArrayType, ".offload_baseptrs").getPointer();
Info.PointersArray =
CGF.CreateMemTemp(PointerArrayType, ".offload_ptrs").getPointer();
+ Info.MappersArray =
+ CGF.CreateMemTemp(PointerArrayType, ".offload_mappers").getPointer();
// If we don't have any VLA types or other types that require runtime
// evaluation, we can use a constant array for the map sizes, otherwise we
@@ -8694,16 +8977,30 @@
CGF.Builder.CreateIntCast(Sizes[I], CGM.Int64Ty, /*isSigned=*/true),
SAddr);
}
+
+ // Fill up the mapper array.
+ llvm::Value *MFunc = llvm::ConstantPointerNull::get(CGM.VoidPtrTy);
+ if (Mappers[I])
+ MFunc = CGM.getOpenMPRuntime().getUserDefinedMapperFunc(
+ cast<OMPDeclareMapperDecl>(Mappers[I]));
+ llvm::Value *M = CGF.Builder.CreateConstInBoundsGEP2_32(
+ llvm::ArrayType::get(CGM.VoidPtrTy, Info.NumberOfPtrs),
+ Info.MappersArray, 0, I);
+ M = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
+ M, MFunc->getType()->getPointerTo(/*AddrSpace=*/0));
+ Address MAddr(M, Ctx.getTypeAlignInChars(Ctx.VoidPtrTy));
+ CGF.Builder.CreateStore(MFunc, MAddr);
}
}
}
/// Emit the arguments to be passed to the runtime library based on the
-/// arrays of pointers, sizes and map types.
+/// arrays of base pointers, pointers, sizes, map types, and mappers.
static void emitOffloadingArraysArgument(
CodeGenFunction &CGF, llvm::Value *&BasePointersArrayArg,
llvm::Value *&PointersArrayArg, llvm::Value *&SizesArrayArg,
- llvm::Value *&MapTypesArrayArg, CGOpenMPRuntime::TargetDataInfo &Info) {
+ llvm::Value *&MapTypesArrayArg, llvm::Value *&MappersArrayArg,
+ CGOpenMPRuntime::TargetDataInfo &Info) {
CodeGenModule &CGM = CGF.CGM;
if (Info.NumberOfPtrs) {
BasePointersArrayArg = CGF.Builder.CreateConstInBoundsGEP2_32(
@@ -8723,12 +9020,17 @@
Info.MapTypesArray,
/*Idx0=*/0,
/*Idx1=*/0);
+ MappersArrayArg = CGF.Builder.CreateConstInBoundsGEP2_32(
+ llvm::ArrayType::get(CGM.VoidPtrTy, Info.NumberOfPtrs),
+ Info.MappersArray,
+ /*Idx0=*/0, /*Idx1=*/0);
} else {
BasePointersArrayArg = llvm::ConstantPointerNull::get(CGM.VoidPtrPtrTy);
PointersArrayArg = llvm::ConstantPointerNull::get(CGM.VoidPtrPtrTy);
SizesArrayArg = llvm::ConstantPointerNull::get(CGM.Int64Ty->getPointerTo());
MapTypesArrayArg =
llvm::ConstantPointerNull::get(CGM.Int64Ty->getPointerTo());
+ MappersArrayArg = llvm::ConstantPointerNull::get(CGM.VoidPtrPtrTy);
}
}
@@ -8942,6 +9244,7 @@
// Emit the loop body block.
MapperCGF.EmitBlock(BodyBB);
+ llvm::BasicBlock *LastBB = BodyBB;
llvm::PHINode *PtrPHI = MapperCGF.Builder.CreatePHI(
PtrBegin->getType(), 2, "omp.arraymap.ptrcurrent");
PtrPHI->addIncoming(PtrBegin, EntryBB);
@@ -8963,8 +9266,10 @@
MappableExprsHandler::MapValuesArrayTy Pointers;
MappableExprsHandler::MapValuesArrayTy Sizes;
MappableExprsHandler::MapFlagsArrayTy MapTypes;
+ MappableExprsHandler::MapMappersArrayTy Mappers;
MappableExprsHandler MEHandler(*D, MapperCGF);
- MEHandler.generateAllInfoForMapper(BasePointers, Pointers, Sizes, MapTypes);
+ MEHandler.generateAllInfoForMapper(BasePointers, Pointers, Sizes, MapTypes,
+ Mappers);
// Call the runtime API __tgt_mapper_num_components to get the number of
// pre-existing components.
@@ -9062,6 +9367,7 @@
MapperCGF.Builder.getInt64(~MappableExprsHandler::OMP_MAP_TO));
// In case of tofrom, do nothing.
MapperCGF.EmitBlock(EndBB);
+ LastBB = EndBB;
llvm::PHINode *CurMapType =
MapperCGF.Builder.CreatePHI(CGM.Int64Ty, 4, "omp.maptype");
CurMapType->addIncoming(AllocMapType, AllocBB);
@@ -9069,22 +9375,28 @@
CurMapType->addIncoming(FromMapType, FromBB);
CurMapType->addIncoming(MemberMapType, ToElseBB);
- // TODO: call the corresponding mapper function if a user-defined mapper is
- // associated with this map clause.
- // Call the runtime API __tgt_push_mapper_component to fill up the runtime
- // data structure.
llvm::Value *OffloadingArgs[] = {Handle, CurBaseArg, CurBeginArg,
CurSizeArg, CurMapType};
- MapperCGF.EmitRuntimeCall(
- createRuntimeFunction(OMPRTL__tgt_push_mapper_component),
- OffloadingArgs);
+ if (Mappers[I]) {
+ // Call the corresponding mapper function.
+ llvm::Function *MapperFunc =
+ getUserDefinedMapperFunc(cast<OMPDeclareMapperDecl>(Mappers[I]));
+ assert(MapperFunc && "Expect a valid mapper function is available.");
+ MapperCGF.Builder.CreateCall(MapperFunc, OffloadingArgs);
+ } else {
+ // Call the runtime API __tgt_push_mapper_component to fill up the runtime
+ // data structure.
+ MapperCGF.EmitRuntimeCall(
+ createRuntimeFunction(OMPRTL__tgt_push_mapper_component),
+ OffloadingArgs);
+ }
}
// Update the pointer to point to the next element that needs to be mapped,
// and check whether we have mapped all elements.
llvm::Value *PtrNext = MapperCGF.Builder.CreateConstGEP1_32(
PtrPHI, /*Idx0=*/1, "omp.arraymap.next");
- PtrPHI->addIncoming(PtrNext, BodyBB);
+ PtrPHI->addIncoming(PtrNext, LastBB);
llvm::Value *IsDone =
MapperCGF.Builder.CreateICmpEQ(PtrNext, PtrEnd, "omp.arraymap.isdone");
llvm::BasicBlock *ExitBB = MapperCGF.createBasicBlock("omp.arraymap.exit");
@@ -9159,6 +9471,15 @@
createRuntimeFunction(OMPRTL__tgt_push_mapper_component), OffloadingArgs);
}
+llvm::Function *
+CGOpenMPRuntime::getUserDefinedMapperFunc(const OMPDeclareMapperDecl *D) {
+ auto I = UDMMap.find(D);
+ if (I != UDMMap.end())
+ return I->second;
+ emitUserDefinedMapper(D);
+ return UDMMap.lookup(D);
+}
+
void CGOpenMPRuntime::emitTargetNumIterationsCall(
CodeGenFunction &CGF, const OMPExecutableDirective &D, const Expr *Device,
const llvm::function_ref<llvm::Value *(
@@ -9289,11 +9610,13 @@
InputInfo.PointersArray.getPointer(),
InputInfo.SizesArray.getPointer(),
MapTypesArray,
+ InputInfo.MappersArray.getPointer(),
NumTeams,
NumThreads};
Return = CGF.EmitRuntimeCall(
- createRuntimeFunction(HasNowait ? OMPRTL__tgt_target_teams_nowait
- : OMPRTL__tgt_target_teams),
+ createRuntimeFunction(HasNowait
+ ? OMPRTL__tgt_target_teams_nowait_mapper
+ : OMPRTL__tgt_target_teams_mapper),
OffloadingArgs);
} else {
llvm::Value *OffloadingArgs[] = {DeviceID,
@@ -9302,10 +9625,11 @@
InputInfo.BasePointersArray.getPointer(),
InputInfo.PointersArray.getPointer(),
InputInfo.SizesArray.getPointer(),
- MapTypesArray};
+ MapTypesArray,
+ InputInfo.MappersArray.getPointer()};
Return = CGF.EmitRuntimeCall(
- createRuntimeFunction(HasNowait ? OMPRTL__tgt_target_nowait
- : OMPRTL__tgt_target),
+ createRuntimeFunction(HasNowait ? OMPRTL__tgt_target_nowait_mapper
+ : OMPRTL__tgt_target_mapper),
OffloadingArgs);
}
@@ -9347,6 +9671,7 @@
MappableExprsHandler::MapValuesArrayTy Pointers;
MappableExprsHandler::MapValuesArrayTy Sizes;
MappableExprsHandler::MapFlagsArrayTy MapTypes;
+ MappableExprsHandler::MapMappersArrayTy Mappers;
// Get mappable expression information.
MappableExprsHandler MEHandler(D, CGF);
@@ -9361,7 +9686,9 @@
MappableExprsHandler::MapValuesArrayTy CurPointers;
MappableExprsHandler::MapValuesArrayTy CurSizes;
MappableExprsHandler::MapFlagsArrayTy CurMapTypes;
+ MappableExprsHandler::MapMappersArrayTy CurMappers;
MappableExprsHandler::StructRangeInfoTy PartialStruct;
+ (*CV)->dump();
// VLA sizes are passed to the outlined region by copy and do not have map
// information associated.
@@ -9374,20 +9701,23 @@
CurMapTypes.push_back(MappableExprsHandler::OMP_MAP_LITERAL |
MappableExprsHandler::OMP_MAP_TARGET_PARAM |
MappableExprsHandler::OMP_MAP_IMPLICIT);
+ CurMappers.push_back(nullptr);
} else {
// If we have any information in the map clause, we use it, otherwise we
// just do a default mapping.
MEHandler.generateInfoForCapture(CI, *CV, CurBasePointers, CurPointers,
- CurSizes, CurMapTypes, PartialStruct);
+ CurSizes, CurMapTypes, CurMappers,
+ PartialStruct);
if (CurBasePointers.empty())
MEHandler.generateDefaultMapInfo(*CI, **RI, *CV, CurBasePointers,
- CurPointers, CurSizes, CurMapTypes);
+ CurPointers, CurSizes, CurMapTypes,
+ CurMappers);
// Generate correct mapping for variables captured by reference in
// lambdas.
if (CI->capturesVariable())
MEHandler.generateInfoForLambdaCaptures(
CI->getCapturedVar(), *CV, CurBasePointers, CurPointers, CurSizes,
- CurMapTypes, LambdaPointers);
+ CurMapTypes, CurMappers, LambdaPointers);
}
// We expect to have at least an element of information for this capture.
assert(!CurBasePointers.empty() &&
@@ -9395,19 +9725,21 @@
assert(CurBasePointers.size() == CurPointers.size() &&
CurBasePointers.size() == CurSizes.size() &&
CurBasePointers.size() == CurMapTypes.size() &&
+ CurBasePointers.size() == CurMappers.size() &&
"Inconsistent map information sizes!");
// If there is an entry in PartialStruct it means we have a struct with
// individual members mapped. Emit an extra combined entry.
if (PartialStruct.Base.isValid())
MEHandler.emitCombinedEntry(BasePointers, Pointers, Sizes, MapTypes,
- CurMapTypes, PartialStruct);
+ Mappers, CurMapTypes, PartialStruct);
// We need to append the results of this capture to what we already have.
BasePointers.append(CurBasePointers.begin(), CurBasePointers.end());
Pointers.append(CurPointers.begin(), CurPointers.end());
Sizes.append(CurSizes.begin(), CurSizes.end());
MapTypes.append(CurMapTypes.begin(), CurMapTypes.end());
+ Mappers.append(CurMappers.begin(), CurMappers.end());
}
// Adjust MEMBER_OF flags for the lambdas captures.
MEHandler.adjustMemberOfForLambdaCaptures(LambdaPointers, BasePointers,
@@ -9415,20 +9747,22 @@
// Map other list items in the map clause which are not captured variables
// but "declare target link" global variables.
MEHandler.generateInfoForDeclareTargetLink(BasePointers, Pointers, Sizes,
- MapTypes);
+ MapTypes, Mappers);
TargetDataInfo Info;
// Fill up the arrays and create the arguments.
- emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, Info);
+ emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, Mappers,
+ Info);
emitOffloadingArraysArgument(CGF, Info.BasePointersArray,
Info.PointersArray, Info.SizesArray,
- Info.MapTypesArray, Info);
+ Info.MapTypesArray, Info.MappersArray, Info);
InputInfo.NumberOfTargetItems = Info.NumberOfPtrs;
InputInfo.BasePointersArray =
Address(Info.BasePointersArray, CGM.getPointerAlign());
InputInfo.PointersArray =
Address(Info.PointersArray, CGM.getPointerAlign());
InputInfo.SizesArray = Address(Info.SizesArray, CGM.getPointerAlign());
+ InputInfo.MappersArray = Address(Info.MappersArray, CGM.getPointerAlign());
MapTypesArray = Info.MapTypesArray;
if (RequiresOuterTask)
CGF.EmitOMPTargetTaskBasedDirective(D, ThenGen, InputInfo);
@@ -9983,20 +10317,24 @@
MappableExprsHandler::MapValuesArrayTy Pointers;
MappableExprsHandler::MapValuesArrayTy Sizes;
MappableExprsHandler::MapFlagsArrayTy MapTypes;
+ MappableExprsHandler::MapMappersArrayTy Mappers;
// Get map clause information.
- MappableExprsHandler MCHandler(D, CGF);
- MCHandler.generateAllInfo(BasePointers, Pointers, Sizes, MapTypes);
+ MappableExprsHandler MEHandler(D, CGF);
+ MEHandler.generateAllInfo(BasePointers, Pointers, Sizes, MapTypes, Mappers);
// Fill up the arrays and create the arguments.
- emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, Info);
+ emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, Mappers,
+ Info);
llvm::Value *BasePointersArrayArg = nullptr;
llvm::Value *PointersArrayArg = nullptr;
llvm::Value *SizesArrayArg = nullptr;
llvm::Value *MapTypesArrayArg = nullptr;
+ llvm::Value *MappersArrayArg = nullptr;
emitOffloadingArraysArgument(CGF, BasePointersArrayArg, PointersArrayArg,
- SizesArrayArg, MapTypesArrayArg, Info);
+ SizesArrayArg, MapTypesArrayArg,
+ MappersArrayArg, Info);
// Emit device ID if any.
llvm::Value *DeviceID = nullptr;
@@ -10011,10 +10349,11 @@
llvm::Value *PointerNum = CGF.Builder.getInt32(Info.NumberOfPtrs);
llvm::Value *OffloadingArgs[] = {
- DeviceID, PointerNum, BasePointersArrayArg,
- PointersArrayArg, SizesArrayArg, MapTypesArrayArg};
- CGF.EmitRuntimeCall(createRuntimeFunction(OMPRTL__tgt_target_data_begin),
- OffloadingArgs);
+ DeviceID, PointerNum, BasePointersArrayArg, PointersArrayArg,
+ SizesArrayArg, MapTypesArrayArg, MappersArrayArg};
+ CGF.EmitRuntimeCall(
+ createRuntimeFunction(OMPRTL__tgt_target_data_begin_mapper),
+ OffloadingArgs);
// If device pointer privatization is required, emit the body of the region
// here. It will have to be duplicated: with and without privatization.
@@ -10031,8 +10370,10 @@
llvm::Value *PointersArrayArg = nullptr;
llvm::Value *SizesArrayArg = nullptr;
llvm::Value *MapTypesArrayArg = nullptr;
+ llvm::Value *MappersArrayArg = nullptr;
emitOffloadingArraysArgument(CGF, BasePointersArrayArg, PointersArrayArg,
- SizesArrayArg, MapTypesArrayArg, Info);
+ SizesArrayArg, MapTypesArrayArg,
+ MappersArrayArg, Info);
// Emit device ID if any.
llvm::Value *DeviceID = nullptr;
@@ -10047,10 +10388,11 @@
llvm::Value *PointerNum = CGF.Builder.getInt32(Info.NumberOfPtrs);
llvm::Value *OffloadingArgs[] = {
- DeviceID, PointerNum, BasePointersArrayArg,
- PointersArrayArg, SizesArrayArg, MapTypesArrayArg};
- CGF.EmitRuntimeCall(createRuntimeFunction(OMPRTL__tgt_target_data_end),
- OffloadingArgs);
+ DeviceID, PointerNum, BasePointersArrayArg, PointersArrayArg,
+ SizesArrayArg, MapTypesArrayArg, MappersArrayArg};
+ CGF.EmitRuntimeCall(
+ createRuntimeFunction(OMPRTL__tgt_target_data_end_mapper),
+ OffloadingArgs);
};
// If we need device pointer privatization, we need to emit the body of the
@@ -10124,24 +10466,25 @@
InputInfo.BasePointersArray.getPointer(),
InputInfo.PointersArray.getPointer(),
InputInfo.SizesArray.getPointer(),
- MapTypesArray};
+ MapTypesArray,
+ InputInfo.MappersArray.getPointer()};
- // Select the right runtime function call for each expected standalone
+ // Select the right runtime function call for each standalone
// directive.
const bool HasNowait = D.hasClausesOfKind<OMPNowaitClause>();
OpenMPRTLFunction RTLFn;
switch (D.getDirectiveKind()) {
case OMPD_target_enter_data:
- RTLFn = HasNowait ? OMPRTL__tgt_target_data_begin_nowait
- : OMPRTL__tgt_target_data_begin;
+ RTLFn = HasNowait ? OMPRTL__tgt_target_data_begin_nowait_mapper
+ : OMPRTL__tgt_target_data_begin_mapper;
break;
case OMPD_target_exit_data:
- RTLFn = HasNowait ? OMPRTL__tgt_target_data_end_nowait
- : OMPRTL__tgt_target_data_end;
+ RTLFn = HasNowait ? OMPRTL__tgt_target_data_end_nowait_mapper
+ : OMPRTL__tgt_target_data_end_mapper;
break;
case OMPD_target_update:
- RTLFn = HasNowait ? OMPRTL__tgt_target_data_update_nowait
- : OMPRTL__tgt_target_data_update;
+ RTLFn = HasNowait ? OMPRTL__tgt_target_data_update_nowait_mapper
+ : OMPRTL__tgt_target_data_update_mapper;
break;
case OMPD_parallel:
case OMPD_for:
@@ -10209,17 +10552,19 @@
MappableExprsHandler::MapValuesArrayTy Pointers;
MappableExprsHandler::MapValuesArrayTy Sizes;
MappableExprsHandler::MapFlagsArrayTy MapTypes;
+ MappableExprsHandler::MapMappersArrayTy Mappers;
// Get map clause information.
MappableExprsHandler MEHandler(D, CGF);
- MEHandler.generateAllInfo(BasePointers, Pointers, Sizes, MapTypes);
+ MEHandler.generateAllInfo(BasePointers, Pointers, Sizes, MapTypes, Mappers);
TargetDataInfo Info;
// Fill up the arrays and create the arguments.
- emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, Info);
+ emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, Mappers,
+ Info);
emitOffloadingArraysArgument(CGF, Info.BasePointersArray,
Info.PointersArray, Info.SizesArray,
- Info.MapTypesArray, Info);
+ Info.MapTypesArray, Info.MappersArray, Info);
InputInfo.NumberOfTargetItems = Info.NumberOfPtrs;
InputInfo.BasePointersArray =
Address(Info.BasePointersArray, CGM.getPointerAlign());
@@ -10227,6 +10572,7 @@
Address(Info.PointersArray, CGM.getPointerAlign());
InputInfo.SizesArray =
Address(Info.SizesArray, CGM.getPointerAlign());
+ InputInfo.MappersArray = Address(Info.MappersArray, CGM.getPointerAlign());
MapTypesArray = Info.MapTypesArray;
if (D.hasClausesOfKind<OMPDependClause>())
CGF.EmitOMPTargetTaskBasedDirective(D, ThenGen, InputInfo);
Index: include/clang/AST/OpenMPClause.h
===================================================================
--- include/clang/AST/OpenMPClause.h
+++ include/clang/AST/OpenMPClause.h
@@ -4286,6 +4286,11 @@
/// Total number of components in this clause.
unsigned NumComponents;
+ /// Whether this clause is possible to have user-defined mappers associated.
+ /// It should be true for map, to, and from clauses, and false for
+ /// use_device_ptr and is_device_ptr.
+ bool hasMapper;
+
/// C++ nested name specifier for the associated user-defined mapper.
NestedNameSpecifierLoc MapperQualifierLoc;
@@ -4306,19 +4311,21 @@
/// NumUniqueDeclarations: number of unique base declarations in this clause;
/// 3) NumComponentLists: number of component lists in this clause; and 4)
/// NumComponents: total number of expression components in the clause.
+ /// \param hasMapper Indicates whether this clause is possible to have
+ /// user-defined mappers associated.
/// \param MapperQualifierLocPtr C++ nested name specifier for the associated
/// user-defined mapper.
/// \param MapperIdInfoPtr The identifier of associated user-defined mapper.
OMPMappableExprListClause(
OpenMPClauseKind K, const OMPVarListLocTy &Locs,
- const OMPMappableExprListSizeTy &Sizes,
+ const OMPMappableExprListSizeTy &Sizes, bool hasMapper = false,
NestedNameSpecifierLoc *MapperQualifierLocPtr = nullptr,
DeclarationNameInfo *MapperIdInfoPtr = nullptr)
: OMPVarListClause<T>(K, Locs.StartLoc, Locs.LParenLoc, Locs.EndLoc,
Sizes.NumVars),
NumUniqueDeclarations(Sizes.NumUniqueDeclarations),
NumComponentLists(Sizes.NumComponentLists),
- NumComponents(Sizes.NumComponents) {
+ NumComponents(Sizes.NumComponents), hasMapper(hasMapper) {
if (MapperQualifierLocPtr)
MapperQualifierLoc = *MapperQualifierLocPtr;
if (MapperIdInfoPtr)
@@ -4517,6 +4524,8 @@
/// Get the user-defined mapper references that are in the trailing objects of
/// the class.
MutableArrayRef<Expr *> getUDMapperRefs() {
+ assert(hasMapper &&
+ "Must be a clause that is possible to have user-defined mappers");
return llvm::makeMutableArrayRef<Expr *>(
static_cast<T *>(this)->template getTrailingObjects<Expr *>() +
OMPVarListClause<T>::varlist_size(),
@@ -4525,9 +4534,11 @@
/// Get the user-defined mappers references that are in the trailing objects
/// of the class.
- ArrayRef<Expr *> getUDMapperRefs() const {
- return llvm::makeArrayRef<Expr *>(
- static_cast<T *>(this)->template getTrailingObjects<Expr *>() +
+ ArrayRef<const Expr *> getUDMapperRefs() const {
+ assert(hasMapper &&
+ "Must be a clause that is possible to have user-defined mappers");
+ return llvm::makeArrayRef<const Expr *>(
+ static_cast<const T *>(this)->template getTrailingObjects<Expr *>() +
OMPVarListClause<T>::varlist_size(),
OMPVarListClause<T>::varlist_size());
}
@@ -4537,6 +4548,8 @@
void setUDMapperRefs(ArrayRef<Expr *> DMDs) {
assert(DMDs.size() == OMPVarListClause<T>::varlist_size() &&
"Unexpected number of user-defined mappers.");
+ assert(hasMapper &&
+ "Must be a clause that is possible to have user-defined mappers");
std::copy(DMDs.begin(), DMDs.end(), getUDMapperRefs().begin());
}
@@ -4573,6 +4586,12 @@
// The list number associated with the current declaration.
ArrayRef<unsigned>::iterator NumListsCur;
+ // Whether this clause is possible to have user-defined mappers associated.
+ bool hasMapper;
+
+ // The user-defined mapper associated with the current declaration.
+ ArrayRef<const Expr *>::iterator MapperCur;
+
// Remaining lists for the current declaration.
unsigned RemainingLists = 0;
@@ -4593,10 +4612,12 @@
explicit const_component_lists_iterator(
ArrayRef<ValueDecl *> UniqueDecls, ArrayRef<unsigned> DeclsListNum,
ArrayRef<unsigned> CumulativeListSizes,
- MappableExprComponentListRef Components)
+ MappableExprComponentListRef Components, bool hasMapper,
+ ArrayRef<const Expr *> Mappers)
: const_component_lists_iterator::iterator_adaptor_base(
Components.begin()),
DeclCur(UniqueDecls.begin()), NumListsCur(DeclsListNum.begin()),
+ hasMapper(hasMapper), MapperCur(Mappers.begin()),
ListSizeCur(CumulativeListSizes.begin()),
ListSizeEnd(CumulativeListSizes.end()), End(Components.end()) {
assert(UniqueDecls.size() == DeclsListNum.size() &&
@@ -4610,9 +4631,11 @@
explicit const_component_lists_iterator(
const ValueDecl *Declaration, ArrayRef<ValueDecl *> UniqueDecls,
ArrayRef<unsigned> DeclsListNum, ArrayRef<unsigned> CumulativeListSizes,
- MappableExprComponentListRef Components)
+ MappableExprComponentListRef Components, bool hasMapper,
+ ArrayRef<const Expr *> Mappers)
: const_component_lists_iterator(UniqueDecls, DeclsListNum,
- CumulativeListSizes, Components) {
+ CumulativeListSizes, Components,
+ hasMapper, Mappers) {
// Look for the desired declaration. While we are looking for it, we
// update the state so that we know the component where a given list
// starts.
@@ -4627,6 +4650,9 @@
std::advance(ListSizeCur, *NumListsCur - 1);
PrevListSize = *ListSizeCur;
++ListSizeCur;
+
+ if (hasMapper)
+ ++MapperCur;
}
// If we didn't find any declaration, advance the iterator to after the
@@ -4652,14 +4678,20 @@
// Return the array with the current list. The sizes are cumulative, so the
// array size is the difference between the current size and previous one.
- std::pair<const ValueDecl *, MappableExprComponentListRef>
+ std::tuple<const ValueDecl *, MappableExprComponentListRef,
+ const ValueDecl *>
operator*() const {
assert(ListSizeCur != ListSizeEnd && "Invalid iterator!");
- return std::make_pair(
+ const ValueDecl *Mapper = nullptr;
+ if (hasMapper && *MapperCur)
+ Mapper = cast<ValueDecl>(cast<DeclRefExpr>(*MapperCur)->getDecl());
+ return std::make_tuple(
*DeclCur,
- MappableExprComponentListRef(&*this->I, *ListSizeCur - PrevListSize));
+ MappableExprComponentListRef(&*this->I, *ListSizeCur - PrevListSize),
+ Mapper);
}
- std::pair<const ValueDecl *, MappableExprComponentListRef>
+ std::tuple<const ValueDecl *, MappableExprComponentListRef,
+ const ValueDecl *>
operator->() const {
return **this;
}
@@ -4682,6 +4714,8 @@
if (!(--RemainingLists)) {
++DeclCur;
++NumListsCur;
+ if (hasMapper)
+ ++MapperCur;
RemainingLists = *NumListsCur;
assert(RemainingLists && "No lists in the following declaration??");
}
@@ -4699,13 +4733,15 @@
const_component_lists_iterator component_lists_begin() const {
return const_component_lists_iterator(
getUniqueDeclsRef(), getDeclNumListsRef(), getComponentListSizesRef(),
- getComponentsRef());
+ getComponentsRef(), hasMapper,
+ hasMapper ? getUDMapperRefs() : ArrayRef<const Expr *>());
}
const_component_lists_iterator component_lists_end() const {
return const_component_lists_iterator(
ArrayRef<ValueDecl *>(), ArrayRef<unsigned>(), ArrayRef<unsigned>(),
MappableExprComponentListRef(getComponentsRef().end(),
- getComponentsRef().end()));
+ getComponentsRef().end()),
+ hasMapper, ArrayRef<const Expr *>());
}
const_component_lists_range component_lists() const {
return {component_lists_begin(), component_lists_end()};
@@ -4717,7 +4753,8 @@
decl_component_lists_begin(const ValueDecl *VD) const {
return const_component_lists_iterator(
VD, getUniqueDeclsRef(), getDeclNumListsRef(),
- getComponentListSizesRef(), getComponentsRef());
+ getComponentListSizesRef(), getComponentsRef(), hasMapper,
+ hasMapper ? getUDMapperRefs() : ArrayRef<const Expr *>());
}
const_component_lists_iterator decl_component_lists_end() const {
return component_lists_end();
@@ -4869,8 +4906,8 @@
OpenMPMapClauseKind MapType, bool MapTypeIsImplicit,
SourceLocation MapLoc, const OMPVarListLocTy &Locs,
const OMPMappableExprListSizeTy &Sizes)
- : OMPMappableExprListClause(OMPC_map, Locs, Sizes, &MapperQualifierLoc,
- &MapperIdInfo),
+ : OMPMappableExprListClause(OMPC_map, Locs, Sizes, /*hasMapper=*/true,
+ &MapperQualifierLoc, &MapperIdInfo),
MapType(MapType), MapTypeIsImplicit(MapTypeIsImplicit), MapLoc(MapLoc) {
assert(llvm::array_lengthof(MapTypeModifiers) == MapModifiers.size() &&
"Unexpected number of map type modifiers.");
@@ -4890,7 +4927,8 @@
/// 3) NumComponentLists: number of component lists in this clause; and 4)
/// NumComponents: total number of expression components in the clause.
explicit OMPMapClause(const OMPMappableExprListSizeTy &Sizes)
- : OMPMappableExprListClause(OMPC_map, OMPVarListLocTy(), Sizes) {}
+ : OMPMappableExprListClause(OMPC_map, OMPVarListLocTy(), Sizes,
+ /*hasMapper=*/true) {}
/// Set map-type-modifier for the clause.
///
@@ -5744,8 +5782,8 @@
DeclarationNameInfo MapperIdInfo,
const OMPVarListLocTy &Locs,
const OMPMappableExprListSizeTy &Sizes)
- : OMPMappableExprListClause(OMPC_to, Locs, Sizes, &MapperQualifierLoc,
- &MapperIdInfo) {}
+ : OMPMappableExprListClause(OMPC_to, Locs, Sizes, /*hasMapper=*/true,
+ &MapperQualifierLoc, &MapperIdInfo) {}
/// Build an empty clause.
///
@@ -5755,7 +5793,8 @@
/// 3) NumComponentLists: number of component lists in this clause; and 4)
/// NumComponents: total number of expression components in the clause.
explicit OMPToClause(const OMPMappableExprListSizeTy &Sizes)
- : OMPMappableExprListClause(OMPC_to, OMPVarListLocTy(), Sizes) {}
+ : OMPMappableExprListClause(OMPC_to, OMPVarListLocTy(), Sizes,
+ /*hasMapper=*/true) {}
/// Define the sizes of each trailing object array except the last one. This
/// is required for TrailingObjects to work properly.
@@ -5862,8 +5901,8 @@
DeclarationNameInfo MapperIdInfo,
const OMPVarListLocTy &Locs,
const OMPMappableExprListSizeTy &Sizes)
- : OMPMappableExprListClause(OMPC_from, Locs, Sizes, &MapperQualifierLoc,
- &MapperIdInfo) {}
+ : OMPMappableExprListClause(OMPC_from, Locs, Sizes, /*hasMapper=*/true,
+ &MapperQualifierLoc, &MapperIdInfo) {}
/// Build an empty clause.
///
@@ -5873,7 +5912,8 @@
/// 3) NumComponentLists: number of component lists in this clause; and 4)
/// NumComponents: total number of expression components in the clause.
explicit OMPFromClause(const OMPMappableExprListSizeTy &Sizes)
- : OMPMappableExprListClause(OMPC_from, OMPVarListLocTy(), Sizes) {}
+ : OMPMappableExprListClause(OMPC_from, OMPVarListLocTy(), Sizes,
+ /*hasMapper=*/true) {}
/// Define the sizes of each trailing object array except the last one. This
/// is required for TrailingObjects to work properly.
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
