jhuber6 created this revision.
jhuber6 added a reviewer: jdoerfert.
Herald added subscribers: guansong, yaxunl.
jhuber6 requested review of this revision.
Herald added subscribers: llvm-commits, openmp-commits, cfe-commits, sstefan1.
Herald added projects: clang, OpenMP, LLVM.
The custom mapper API did not previously support the mapping names added
previously. This means they were not present if a user requested debugging
information while using the mapper functions. This adds basic support for
passing the mapped names to the runtime library.
Repository:
rG LLVM Github Monorepo
https://reviews.llvm.org/D94806
Files:
clang/lib/CodeGen/CGOpenMPRuntime.cpp
clang/test/OpenMP/declare_mapper_codegen.cpp
clang/test/OpenMP/target_depend_codegen.cpp
llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
llvm/test/Transforms/OpenMP/add_attributes.ll
openmp/libomptarget/src/interface.cpp
openmp/libomptarget/src/omptarget.cpp
openmp/libomptarget/src/private.h
openmp/libomptarget/test/mapping/declare_mapper_api.cpp
Index: openmp/libomptarget/test/mapping/declare_mapper_api.cpp
===================================================================
--- openmp/libomptarget/test/mapping/declare_mapper_api.cpp
+++ openmp/libomptarget/test/mapping/declare_mapper_api.cpp
@@ -15,9 +15,10 @@
void *Begin;
int64_t Size;
int64_t Type;
+ void *Name;
MapComponentInfoTy() = default;
- MapComponentInfoTy(void *Base, void *Begin, int64_t Size, int64_t Type)
- : Base(Base), Begin(Begin), Size(Size), Type(Type) {}
+ MapComponentInfoTy(void *Base, void *Begin, int64_t Size, int64_t Type, void *Name)
+ : Base(Base), Begin(Begin), Size(Size), Type(Type), Name(Name) {}
};
struct MapperComponentsTy {
@@ -30,7 +31,8 @@
#endif
int64_t __tgt_mapper_num_components(void *rt_mapper_handle);
void __tgt_push_mapper_component(void *rt_mapper_handle, void *base,
- void *begin, int64_t size, int64_t type);
+ void *begin, int64_t size, int64_t type,
+ void *name);
#ifdef __cplusplus
}
#endif
@@ -40,8 +42,8 @@
void *base, *begin;
int64_t size, type;
// Push 2 elements into MC.
- __tgt_push_mapper_component((void *)&MC, base, begin, size, type);
- __tgt_push_mapper_component((void *)&MC, base, begin, size, type);
+ __tgt_push_mapper_component((void *)&MC, base, begin, size, type, nullptr);
+ __tgt_push_mapper_component((void *)&MC, base, begin, size, type, nullptr);
int64_t num = __tgt_mapper_num_components((void *)&MC);
// CHECK: num=2
printf("num=%" PRId64 "\n", num);
Index: openmp/libomptarget/src/private.h
===================================================================
--- openmp/libomptarget/src/private.h
+++ openmp/libomptarget/src/private.h
@@ -48,9 +48,11 @@
void *Begin;
int64_t Size;
int64_t Type;
+ void *Name;
MapComponentInfoTy() = default;
- MapComponentInfoTy(void *Base, void *Begin, int64_t Size, int64_t Type)
- : Base(Base), Begin(Begin), Size(Size), Type(Type) {}
+ MapComponentInfoTy(void *Base, void *Begin, int64_t Size, int64_t Type,
+ void *Name)
+ : Base(Base), Begin(Begin), Size(Size), Type(Type), Name(Name) {}
};
// This structure stores all components of a user-defined mapper. The number of
@@ -64,8 +66,10 @@
// The mapper function pointer type. It follows the signature below:
// void .omp_mapper.<type_name>.<mapper_id>.(void *rt_mapper_handle,
// void *base, void *begin,
-// size_t size, int64_t type);
-typedef void (*MapperFuncPtrTy)(void *, void *, void *, int64_t, int64_t);
+// size_t size, int64_t type,
+// void * name);
+typedef void (*MapperFuncPtrTy)(void *, void *, void *, int64_t, int64_t,
+ void *);
// Function pointer type for target_data_* functions (targetDataBegin,
// targetDataEnd and targetDataUpdate).
Index: openmp/libomptarget/src/omptarget.cpp
===================================================================
--- openmp/libomptarget/src/omptarget.cpp
+++ openmp/libomptarget/src/omptarget.cpp
@@ -209,15 +209,16 @@
/// Call the user-defined mapper function followed by the appropriate
// target_data_* function (target_data_{begin,end,update}).
int targetDataMapper(DeviceTy &Device, void *arg_base, void *arg,
- int64_t arg_size, int64_t arg_type, void *arg_mapper,
+ int64_t arg_size, int64_t arg_type,
+ map_var_info_t arg_names, void *arg_mapper,
TargetDataFuncPtrTy target_data_function) {
DP("Calling the mapper function " DPxMOD "\n", DPxPTR(arg_mapper));
// The mapper function fills up Components.
MapperComponentsTy MapperComponents;
MapperFuncPtrTy MapperFuncPtr = (MapperFuncPtrTy)(arg_mapper);
- (*MapperFuncPtr)((void *)&MapperComponents, arg_base, arg, arg_size,
- arg_type);
+ (*MapperFuncPtr)((void *)&MapperComponents, arg_base, arg, arg_size, arg_type,
+ arg_names);
// Construct new arrays for args_base, args, arg_sizes and arg_types
// using the information in MapperComponents and call the corresponding
@@ -226,6 +227,7 @@
std::vector<void *> MapperArgs(MapperComponents.Components.size());
std::vector<int64_t> MapperArgSizes(MapperComponents.Components.size());
std::vector<int64_t> MapperArgTypes(MapperComponents.Components.size());
+ std::vector<void *> MapperArgNames(MapperComponents.Components.size());
for (unsigned I = 0, E = MapperComponents.Components.size(); I < E; ++I) {
auto &C =
@@ -235,12 +237,13 @@
MapperArgs[I] = C.Begin;
MapperArgSizes[I] = C.Size;
MapperArgTypes[I] = C.Type;
+ MapperArgNames[I] = C.Name;
}
int rc = target_data_function(Device, MapperComponents.Components.size(),
MapperArgsBase.data(), MapperArgs.data(),
MapperArgSizes.data(), MapperArgTypes.data(),
- /*arg_names*/ nullptr, /*arg_mappers*/ nullptr,
+ MapperArgNames.data(), /*arg_mappers*/ nullptr,
/*__tgt_async_info*/ nullptr);
return rc;
@@ -264,8 +267,10 @@
// with new arguments.
DP("Calling targetDataMapper for the %dth argument\n", i);
+ map_var_info_t arg_name = (!arg_names) ? nullptr : arg_names[i];
int rc = targetDataMapper(Device, args_base[i], args[i], arg_sizes[i],
- arg_types[i], arg_mappers[i], targetDataBegin);
+ arg_types[i], arg_name, arg_mappers[i],
+ targetDataBegin);
if (rc != OFFLOAD_SUCCESS) {
REPORT("Call to targetDataBegin via targetDataMapper for custom mapper"
@@ -464,8 +469,10 @@
// with new arguments.
DP("Calling targetDataMapper for the %dth argument\n", I);
- Ret = targetDataMapper(Device, ArgBases[I], Args[I], ArgSizes[I],
- ArgTypes[I], ArgMappers[I], targetDataEnd);
+ map_var_info_t ArgName = (!ArgNames) ? nullptr : ArgNames[I];
+ Ret =
+ targetDataMapper(Device, ArgBases[I], Args[I], ArgSizes[I],
+ ArgTypes[I], ArgName, ArgMappers[I], targetDataEnd);
if (Ret != OFFLOAD_SUCCESS) {
REPORT("Call to targetDataEnd via targetDataMapper for custom mapper"
@@ -785,8 +792,10 @@
// with new arguments.
DP("Calling targetDataMapper for the %dth argument\n", I);
+ map_var_info_t ArgName = (!ArgNames) ? nullptr : ArgNames[I];
int Ret = targetDataMapper(Device, ArgsBase[I], Args[I], ArgSizes[I],
- ArgTypes[I], ArgMappers[I], targetDataUpdate);
+ ArgTypes[I], ArgName, ArgMappers[I],
+ targetDataUpdate);
if (Ret != OFFLOAD_SUCCESS) {
REPORT("Call to targetDataUpdate via targetDataMapper for custom mapper"
Index: openmp/libomptarget/src/interface.cpp
===================================================================
--- openmp/libomptarget/src/interface.cpp
+++ openmp/libomptarget/src/interface.cpp
@@ -487,16 +487,17 @@
// Push back one component for a user-defined mapper.
EXTERN void __tgt_push_mapper_component(void *rt_mapper_handle, void *base,
- void *begin, int64_t size,
- int64_t type) {
+ void *begin, int64_t size, int64_t type,
+ void *name) {
TIMESCOPE();
DP("__tgt_push_mapper_component(Handle=" DPxMOD
") adds an entry (Base=" DPxMOD ", Begin=" DPxMOD ", Size=%" PRId64
- ", Type=0x%" PRIx64 ").\n",
- DPxPTR(rt_mapper_handle), DPxPTR(base), DPxPTR(begin), size, type);
+ ", Type=0x%" PRIx64 ", Name=%s).\n",
+ DPxPTR(rt_mapper_handle), DPxPTR(base), DPxPTR(begin), size, type,
+ (name) ? getNameFromMapping(name).c_str() : "unknown");
auto *MapperComponentsPtr = (struct MapperComponentsTy *)rt_mapper_handle;
MapperComponentsPtr->Components.push_back(
- MapComponentInfoTy(base, begin, size, type));
+ MapComponentInfoTy(base, begin, size, type, name));
}
EXTERN void __kmpc_push_target_tripcount(ident_t *loc, int64_t device_id,
Index: llvm/test/Transforms/OpenMP/add_attributes.ll
===================================================================
--- llvm/test/Transforms/OpenMP/add_attributes.ll
+++ llvm/test/Transforms/OpenMP/add_attributes.ll
@@ -657,7 +657,7 @@
declare i64 @__tgt_mapper_num_components(i8*)
-declare void @__tgt_push_mapper_component(i8*, i8*, i8*, i64, i64)
+declare void @__tgt_push_mapper_component(i8*, i8*, i8*, i64, i64, i8*)
declare i8* @__kmpc_task_allow_completion_event(%struct.ident_t*, i32, i8*)
Index: llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
===================================================================
--- llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
+++ llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
@@ -402,7 +402,7 @@
VoidPtrPtr, VoidPtrPtr, Int64Ptr, Int64Ptr, VoidPtrPtr, VoidPtrPtr)
__OMP_RTL(__tgt_mapper_num_components, false, Int64, VoidPtr)
__OMP_RTL(__tgt_push_mapper_component, false, Void, VoidPtr, VoidPtr, VoidPtr,
- Int64, Int64)
+ Int64, Int64, VoidPtr)
__OMP_RTL(__kmpc_task_allow_completion_event, false, VoidPtr, IdentPtr,
/* Int */ Int32, /* kmp_task_t */ VoidPtr)
Index: clang/test/OpenMP/target_depend_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_depend_codegen.cpp
+++ clang/test/OpenMP/target_depend_codegen.cpp
@@ -130,7 +130,7 @@
// CHECK-DAG: [[CPADDR2:%.+]] = bitcast i8** [[PADDR2]] to [[STRUCT_TT]]**
// CHECK-DAG: store [[STRUCT_TT]]* [[D_ADDR:%.+]], [[STRUCT_TT]]** [[CBPADDR2]]
// CHECK-DAG: store [[STRUCT_TT]]* [[D_ADDR]], [[STRUCT_TT]]** [[CPADDR2]]
- // CHECK-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MAPPER_ID:@.+]] to i8*), i8** [[MADDR2]],
+ // CHECK-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64, i8*)* [[MAPPER_ID:@.+]] to i8*), i8** [[MADDR2]],
// CHECK-DAG: [[BP_START:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BP]], i32 0, i32 0
// CHECK-DAG: [[P_START:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[P]], i32 0, i32 0
Index: clang/test/OpenMP/declare_mapper_codegen.cpp
===================================================================
--- clang/test/OpenMP/declare_mapper_codegen.cpp
+++ clang/test/OpenMP/declare_mapper_codegen.cpp
@@ -86,7 +86,7 @@
#pragma omp declare mapper(id: C s) map(s.a, s.b[0:2])
-// CK0: define {{.*}}void [[MPRFUNC:@[.]omp_mapper[.].*C[.]id]](i8*{{.*}}, i8*{{.*}}, i8*{{.*}}, i64{{.*}}, i64{{.*}})
+// CK0: define {{.*}}void [[MPRFUNC:@[.]omp_mapper[.].*C[.]id]](i8*{{.*}}, i8*{{.*}}, i8*{{.*}}, i64{{.*}}, i64{{.*}}, i8*{{.*}})
// CK0: store i8* %{{[^,]+}}, i8** [[HANDLEADDR:%[^,]+]]
// CK0: store i8* %{{[^,]+}}, i8** [[BPTRADDR:%[^,]+]]
// CK0: store i8* %{{[^,]+}}, i8** [[VPTRADDR:%[^,]+]]
@@ -112,7 +112,7 @@
// CK0-64-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
// CK0-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8
// CK0-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4
-// CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]])
+// CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]], {{.*}})
// CK0: br label %[[LHEAD:[^,]+]]
// CK0: [[LHEAD]]
@@ -165,7 +165,7 @@
// CK0-DAG: br label %[[TYEND]]
// CK0-DAG: [[TYEND]]
// CK0-DAG: [[PHITYPE0:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
-// CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR0BC]], i8* [[PTRADDR0BC]], i64 [[CUSIZE]], i64 [[PHITYPE0]])
+// CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR0BC]], i8* [[PTRADDR0BC]], i64 [[CUSIZE]], i64 [[PHITYPE0]], {{.*}})
// CK0-DAG: [[BPTRADDR1BC:%.+]] = bitcast %class.C* [[OBJ]] to i8*
// CK0-DAG: [[PTRADDR1BC:%.+]] = bitcast i32* [[ABEGIN]] to i8*
// CK0-DAG: br label %[[MEMBER:[^,]+]]
@@ -197,7 +197,7 @@
// CK0-DAG: br label %[[TYEND]]
// CK0-DAG: [[TYEND]]
// CK0-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
-// CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR1BC]], i8* [[PTRADDR1BC]], i64 4, i64 [[TYPE1]])
+// CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR1BC]], i8* [[PTRADDR1BC]], i64 4, i64 [[TYPE1]], {{.*}})
// CK0-DAG: [[BPTRADDR2BC:%.+]] = bitcast double** [[BBEGIN]] to i8*
// CK0-DAG: [[PTRADDR2BC:%.+]] = bitcast double* [[BARRBEGINGEP]] to i8*
// CK0-DAG: br label %[[MEMBER:[^,]+]]
@@ -229,7 +229,7 @@
// CK0-DAG: br label %[[TYEND]]
// CK0-DAG: [[TYEND]]
// CK0-DAG: [[TYPE2:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
-// CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR2BC]], i8* [[PTRADDR2BC]], i64 16, i64 [[TYPE2]])
+// CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR2BC]], i8* [[PTRADDR2BC]], i64 16, i64 [[TYPE2]], {{.*}})
// CK0: [[PTRNEXT]] = getelementptr %class.C*, %class.C** [[PTR]], i32 1
// CK0: [[ISDONE:%.+]] = icmp eq %class.C** [[PTRNEXT]], [[PTREND]]
// CK0: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]]
@@ -245,7 +245,7 @@
// CK0-64-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
// CK0-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8
// CK0-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4
-// CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE]])
+// CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE]], {{.*}})
// CK0: br label %[[DONE]]
// CK0: [[DONE]]
// CK0: ret void
@@ -268,7 +268,7 @@
// CK0-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to %class.C**
// CK0-DAG: store %class.C* [[VAL:%[^,]+]], %class.C** [[CBP1]]
// CK0-DAG: store %class.C* [[VAL]], %class.C** [[CP1]]
- // CK0-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[MPR1]]
+ // CK0-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64, i8*)* [[MPRFUNC]] to i8*), i8** [[MPR1]]
// CK0: call void [[KERNEL_1:@.+]](%class.C* [[VAL]])
#pragma omp target map(mapper(id),tofrom: c)
{
@@ -282,7 +282,7 @@
// CK0: [[P2CAST:%.+]] = bitcast i8** [[P2GEP]] to %class.C**
// CK0: store %class.C* [[CADDR]], %class.C** [[P2CAST]], align
// CK0: [[MAPPER2GEP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[OFFLOAD_MAPPER2:%[^,]+]], i[[SZ]] 0, i[[SZ]] 0
- // CK0: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[MAPPER2GEP]], align
+ // CK0: store i8* bitcast (void (i8*, i8*, i8*, i64, i64, i8*)* [[MPRFUNC]] to i8*), i8** [[MAPPER2GEP]], align
// CK0: [[BP2:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[OFFLOAD_BP2]], i32 0, i32 0
// CK0: [[P2:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[OFFLOAD_P2]], i32 0, i32 0
// CK0: [[MAPPER2:%.+]] = bitcast [1 x i8*]* [[OFFLOAD_MAPPER2]] to i8**
@@ -307,7 +307,7 @@
// CK0-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to %class.C**
// CK0-DAG: store %class.C* [[VAL:%[^,]+]], %class.C** [[CBP1]]
// CK0-DAG: store %class.C* [[VAL]], %class.C** [[CP1]]
- // CK0-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[MPR1]]
+ // CK0-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64, i8*)* [[MPRFUNC]] to i8*), i8** [[MPR1]]
// CK0: call void [[KERNEL_3:@.+]](%class.C* [[VAL]])
#pragma omp target teams map(mapper(id),to: c)
{
@@ -336,7 +336,7 @@
// CK0-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to %class.C**
// CK0-DAG: store %class.C* [[VAL:%[^,]+]], %class.C** [[CBP1]]
// CK0-DAG: store %class.C* [[VAL]], %class.C** [[CP1]]
- // CK0-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[MPR1]]
+ // CK0-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64, i8*)* [[MPRFUNC]] to i8*), i8** [[MPR1]]
#pragma omp target enter data map(mapper(id),to: c)
// CK0-DAG: call i32 @__kmpc_omp_task([[IDENT_T]]* @{{[^,]+}}, i32 %{{[^,]+}}, i8* [[TASK_2:%.+]])
@@ -371,7 +371,7 @@
// CK0-DAG: call void @llvm.memcpy.p0i8.p0i8.i[[sz]](i8* align {{4|8}} [[FPMPRADDR]], i8* align {{4|8}} [[MPRADDR]], i[[sz]] {{4|8}}, i1 false)
// CK0-DAG: [[MPRGEP]] = bitcast [1 x i8*]* [[MPR:%.+]] to i8**
// CK0-DAG: [[MPRGEP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[MPR]], i[[sz]] 0, i[[sz]] 0
- // CK0-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[MPRGEP]], align
+ // CK0-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64, i8*)* [[MPRFUNC]] to i8*), i8** [[MPRGEP]], align
#pragma omp target enter data map(mapper(id),to: c) nowait
// CK0-DAG: call void @__tgt_target_data_end_mapper(%struct.ident_t* @{{.+}}, i64 {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[EXDSIZES]]{{.+}}, {{.+}}[[EXDTYPES]]{{.+}}, i8** null, i8** [[MPRGEP:%.+]])
@@ -385,7 +385,7 @@
// CK0-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to %class.C**
// CK0-DAG: store %class.C* [[VAL:%[^,]+]], %class.C** [[CBP1]]
// CK0-DAG: store %class.C* [[VAL]], %class.C** [[CP1]]
- // CK0-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[MPR1]]
+ // CK0-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64, i8*)* [[MPRFUNC]] to i8*), i8** [[MPR1]]
#pragma omp target exit data map(mapper(id),from: c)
// CK0-DAG: call i32 @__kmpc_omp_task([[IDENT_T]]* @{{[^,]+}}, i32 %{{[^,]+}}, i8* [[TASK_3:%.+]])
@@ -420,7 +420,7 @@
// CK0-DAG: call void @llvm.memcpy.p0i8.p0i8.i[[sz]](i8* align {{4|8}} [[FPMPRADDR]], i8* align {{4|8}} [[MPRADDR]], i[[sz]] {{4|8}}, i1 false)
// CK0-DAG: [[MPRGEP]] = bitcast [1 x i8*]* [[MPR:%.+]] to i8**
// CK0-DAG: [[MPRGEP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[MPR]], i[[sz]] 0, i[[sz]] 0
- // CK0-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[MPRGEP]], align
+ // CK0-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64, i8*)* [[MPRFUNC]] to i8*), i8** [[MPRGEP]], align
#pragma omp target exit data map(mapper(id),from: c) nowait
// CK0-DAG: call void @__tgt_target_data_update_mapper(%struct.ident_t* @{{.+}}, i64 -1, i32 1, i8** [[TGEPBP:%.+]], i8** [[TGEPP:%.+]], i64* getelementptr {{.+}}[1 x i64]* [[TSIZES]], i32 0, i32 0), {{.+}}getelementptr {{.+}}[1 x i64]* [[TTYPES]]{{.+}}, i8** null, i8** [[TMPRGEP:%.+]])
@@ -434,7 +434,7 @@
// CK0-DAG: [[TCP0:%.+]] = bitcast i8** [[TP0]] to %class.C**
// CK0-DAG: store %class.C* [[VAL]], %class.C** [[TCBP0]]
// CK0-DAG: store %class.C* [[VAL]], %class.C** [[TCP0]]
- // CK0-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[TMPR1]]
+ // CK0-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64, i8*)* [[MPRFUNC]] to i8*), i8** [[TMPR1]]
#pragma omp target update to(mapper(id): c)
// CK0-DAG: call void @__tgt_target_data_update_mapper(%struct.ident_t* @{{.+}}, i64 -1, i32 1, i8** [[FGEPBP:%.+]], i8** [[FGEPP:%.+]], i64* getelementptr {{.+}}[1 x i64]* [[FSIZES]], i32 0, i32 0), {{.+}}getelementptr {{.+}}[1 x i64]* [[FTYPES]]{{.+}}, i8** null, i8** [[FMPRGEP:%.+]])
@@ -448,7 +448,7 @@
// CK0-DAG: [[FCP0:%.+]] = bitcast i8** [[FP0]] to %class.C**
// CK0-DAG: store %class.C* [[VAL]], %class.C** [[FCBP0]]
// CK0-DAG: store %class.C* [[VAL]], %class.C** [[FCP0]]
- // CK0-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[FMPR1]]
+ // CK0-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64, i8*)* [[MPRFUNC]] to i8*), i8** [[FMPR1]]
#pragma omp target update from(mapper(id): c)
// CK0-DAG: call i32 @__kmpc_omp_task([[IDENT_T]]* @{{[^,]+}}, i32 %{{[^,]+}}, i8* [[TASK_4:%.+]])
@@ -483,7 +483,7 @@
// CK0-DAG: call void @llvm.memcpy.p0i8.p0i8.i[[sz]](i8* align {{4|8}} [[FPMPRADDR]], i8* align {{4|8}} [[MPRADDR]], i[[sz]] {{4|8}}, i1 false)
// CK0-DAG: [[MPRGEP]] = bitcast [1 x i8*]* [[MPR:%.+]] to i8**
// CK0-DAG: [[MPRGEP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[MPR]], i[[sz]] 0, i[[sz]] 0
- // CK0-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[MPRGEP]], align
+ // CK0-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64, i8*)* [[MPRFUNC]] to i8*), i8** [[MPRGEP]], align
#pragma omp target update from(mapper(id): c) nowait
}
@@ -652,7 +652,7 @@
#pragma omp declare mapper(id: C<int> s) map(s.a)
-// CK1-LABEL: define {{.*}}void @.omp_mapper.{{.*}}C{{.*}}.id{{.*}}(i8*{{.*}}, i8*{{.*}}, i8*{{.*}}, i64{{.*}}, i64{{.*}})
+// CK1-LABEL: define {{.*}}void @.omp_mapper.{{.*}}C{{.*}}.id{{.*}}(i8*{{.*}}, i8*{{.*}}, i8*{{.*}}, i64{{.*}}, i64{{.*}}, i8*{{.*}})
// CK1: store i8* %{{[^,]+}}, i8** [[HANDLEADDR:%[^,]+]]
// CK1: store i8* %{{[^,]+}}, i8** [[BPTRADDR:%[^,]+]]
// CK1: store i8* %{{[^,]+}}, i8** [[VPTRADDR:%[^,]+]]
@@ -676,7 +676,7 @@
// CK1: [[INIT]]
// CK1-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 4
// CK1-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4
-// CK1: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]])
+// CK1: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]], {{.*}})
// CK1: br label %[[LHEAD:[^,]+]]
// CK1: [[LHEAD]]
@@ -718,7 +718,7 @@
// CK1-DAG: br label %[[TYEND]]
// CK1-DAG: [[TYEND]]
// CK1-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
-// CK1: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR1BC]], i8* [[PTRADDR1BC]], i64 4, i64 [[TYPE1]])
+// CK1: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR1BC]], i8* [[PTRADDR1BC]], i64 4, i64 [[TYPE1]], {{.*}})
// CK1: [[PTRNEXT]] = getelementptr %class.C*, %class.C** [[PTR]], i32 1
// CK1: [[ISDONE:%.+]] = icmp eq %class.C** [[PTRNEXT]], [[PTREND]]
// CK1: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]]
@@ -733,7 +733,7 @@
// CK1: [[DEL]]
// CK1-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 4
// CK1-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4
-// CK1: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE]])
+// CK1: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE]], {{.*}})
// CK1: br label %[[DONE]]
// CK1: [[DONE]]
// CK1: ret void
@@ -774,9 +774,9 @@
#pragma omp declare mapper(id: C s) map(s.b)
-// CK2: define {{.*}}void [[BMPRFUNC:@[.]omp_mapper[.].*B[.]default]](i8*{{.*}}, i8*{{.*}}, i8*{{.*}}, i64{{.*}}, i64{{.*}})
+// CK2: define {{.*}}void [[BMPRFUNC:@[.]omp_mapper[.].*B[.]default]](i8*{{.*}}, i8*{{.*}}, i8*{{.*}}, i64{{.*}}, i64{{.*}}, i8*{{.*}})
-// CK2-LABEL: define {{.*}}void @.omp_mapper.{{.*}}C{{.*}}.id(i8*{{.*}}, i8*{{.*}}, i8*{{.*}}, i64{{.*}}, i64{{.*}})
+// CK2-LABEL: define {{.*}}void @.omp_mapper.{{.*}}C{{.*}}.id(i8*{{.*}}, i8*{{.*}}, i8*{{.*}}, i64{{.*}}, i64{{.*}}, i8*{{.*}})
// CK2: store i8* %{{[^,]+}}, i8** [[HANDLEADDR:%[^,]+]]
// CK2: store i8* %{{[^,]+}}, i8** [[BPTRADDR:%[^,]+]]
// CK2: store i8* %{{[^,]+}}, i8** [[VPTRADDR:%[^,]+]]
@@ -800,7 +800,7 @@
// 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: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]], {{.*}})
// CK2: br label %[[LHEAD:[^,]+]]
// CK2: [[LHEAD]]
@@ -842,7 +842,7 @@
// 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 8, i64 [[TYPE1]])
+// CK2: call void [[BMPRFUNC]](i8* [[HANDLE]], i8* [[BPTRADDR1BC]], i8* [[PTRADDR1BC]], i64 8, 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]]
@@ -857,7 +857,7 @@
// 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: 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
@@ -900,7 +900,7 @@
#pragma omp declare mapper(id: C s) map(s.a, s.b[0:2])
-// CK3: define {{.*}}void [[MPRFUNC:@[.]omp_mapper[.].*C[.]id]](i8*{{.*}}, i8*{{.*}}, i8*{{.*}}, i64{{.*}}, i64{{.*}})
+// CK3: define {{.*}}void [[MPRFUNC:@[.]omp_mapper[.].*C[.]id]](i8*{{.*}}, i8*{{.*}}, i8*{{.*}}, i64{{.*}}, i64{{.*}}, i8*{{.*}})
// CK3-LABEL: define {{.*}}void @{{.*}}foo{{.*}}
void foo(int a){
@@ -922,7 +922,7 @@
// CK3-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to %class.C**
// CK3-DAG: store %class.B* [[BVAL]], %class.B** [[CBP1]]
// CK3-DAG: store %class.C* [[BC]], %class.C** [[CP1]]
- // CK3-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[MPR1]]
+ // CK3-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64, i8*)* [[MPRFUNC]] to i8*), i8** [[MPR1]]
// CK3-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 1
// CK3-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 1
// CK3-DAG: [[MPR2:%.+]] = getelementptr inbounds {{.+}}[[MPR]], i{{64|32}} 0, i{{64|32}} 1
@@ -931,7 +931,7 @@
// CK3-DAG: store [10 x %class.C]* [[CVAL]], [10 x %class.C]** [[CBP2]]
// CK3-DAG: [[CVALGEP:%.+]] = getelementptr inbounds {{.+}}[[CVAL]], i{{64|32}} 0, i{{64|32}} 0
// CK3-DAG: store %class.C* [[CVALGEP]], %class.C** [[CP2]]
- // CK3-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[MPR2]]
+ // CK3-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64, i8*)* [[MPRFUNC]] to i8*), i8** [[MPR2]]
// CK3: call void [[KERNEL:@.+]](%class.B* [[BVAL]], [10 x %class.C]* [[CVAL]])
#pragma omp target map(mapper(id),tofrom: c[0:10], b.c)
for (int i = 0; i < 10; i++) {
@@ -980,7 +980,7 @@
#pragma omp declare mapper(id: C s) map(s.a, s.b[0:2])
-// CK4: define {{.*}}void [[MPRFUNC:@[.]omp_mapper[.].*C[.]id]](i8*{{.*}}, i8*{{.*}}, i8*{{.*}}, i64{{.*}}, i64{{.*}})
+// CK4: define {{.*}}void [[MPRFUNC:@[.]omp_mapper[.].*C[.]id]](i8*{{.*}}, i8*{{.*}}, i8*{{.*}}, i64{{.*}}, i64{{.*}}, i8*{{.*}})
// CK4: store i8* %{{[^,]+}}, i8** [[HANDLEADDR:%[^,]+]]
// CK4: store i8* %{{[^,]+}}, i8** [[BPTRADDR:%[^,]+]]
// CK4: store i8* %{{[^,]+}}, i8** [[VPTRADDR:%[^,]+]]
@@ -1006,7 +1006,7 @@
// CK4-64-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
// CK4-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8
// CK4-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4
-// CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]])
+// CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]], {{.*}})
// CK4: br label %[[LHEAD:[^,]+]]
// CK4: [[LHEAD]]
@@ -1059,7 +1059,7 @@
// CK4-DAG: br label %[[TYEND]]
// CK4-DAG: [[TYEND]]
// CK4-DAG: [[PHITYPE0:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
-// CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR0BC]], i8* [[PTRADDR0BC]], i64 [[CUSIZE]], i64 [[PHITYPE0]])
+// CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR0BC]], i8* [[PTRADDR0BC]], i64 [[CUSIZE]], i64 [[PHITYPE0]], {{.*}})
// CK4-DAG: [[BPTRADDR1BC:%.+]] = bitcast %class.C* [[OBJ]] to i8*
// CK4-DAG: [[PTRADDR1BC:%.+]] = bitcast i32* [[ABEGIN]] to i8*
// CK4-DAG: br label %[[MEMBER:[^,]+]]
@@ -1091,7 +1091,7 @@
// CK4-DAG: br label %[[TYEND]]
// CK4-DAG: [[TYEND]]
// CK4-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
-// CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR1BC]], i8* [[PTRADDR1BC]], i64 4, i64 [[TYPE1]])
+// CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR1BC]], i8* [[PTRADDR1BC]], i64 4, i64 [[TYPE1]], {{.*}})
// CK4-DAG: [[BPTRADDR2BC:%.+]] = bitcast double** [[BBEGIN]] to i8*
// CK4-DAG: [[PTRADDR2BC:%.+]] = bitcast double* [[BARRBEGINGEP]] to i8*
// CK4-DAG: br label %[[MEMBER:[^,]+]]
@@ -1123,7 +1123,7 @@
// CK4-DAG: br label %[[TYEND]]
// CK4-DAG: [[TYEND]]
// CK4-DAG: [[TYPE2:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
-// CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR2BC]], i8* [[PTRADDR2BC]], i64 16, i64 [[TYPE2]])
+// CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR2BC]], i8* [[PTRADDR2BC]], i64 16, i64 [[TYPE2]], {{.*}})
// CK4: [[PTRNEXT]] = getelementptr %class.C*, %class.C** [[PTR]], i32 1
// CK4: [[ISDONE:%.+]] = icmp eq %class.C** [[PTRNEXT]], [[PTREND]]
// CK4: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]]
@@ -1139,7 +1139,7 @@
// CK4-64-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
// CK4-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8
// CK4-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4
-// CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE]])
+// CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE]], {{.*}})
// CK4: br label %[[DONE]]
// CK4: [[DONE]]
// CK4: ret void
@@ -1162,7 +1162,7 @@
// CK4-DAG: [[TCP0:%.+]] = bitcast i8** [[TP0]] to %class.C**
// CK4-DAG: store %class.C* [[VAL:%[^,]+]], %class.C** [[TCBP0]]
// CK4-DAG: store %class.C* [[VAL]], %class.C** [[TCP0]]
- // CK4-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[TMPR1]]
+ // CK4-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64, i8*)* [[MPRFUNC]] to i8*), i8** [[TMPR1]]
#pragma omp target update to(present, mapper(id): c)
// CK4-DAG: call void @__tgt_target_data_update_mapper(%struct.ident_t* @{{.+}}, i64 -1, i32 1, i8** [[FGEPBP:%.+]], i8** [[FGEPP:%.+]], i64* getelementptr {{.+}}[1 x i64]* [[FSIZES]], i32 0, i32 0), {{.+}}getelementptr {{.+}}[1 x i64]* [[FTYPES]]{{.+}}, i8** null, i8** [[FMPRGEP:%.+]])
@@ -1176,7 +1176,7 @@
// CK4-DAG: [[FCP0:%.+]] = bitcast i8** [[FP0]] to %class.C**
// CK4-DAG: store %class.C* [[VAL]], %class.C** [[FCBP0]]
// CK4-DAG: store %class.C* [[VAL]], %class.C** [[FCP0]]
- // CK4-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[FMPR1]]
+ // CK4-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64, i8*)* [[MPRFUNC]] to i8*), i8** [[FMPR1]]
#pragma omp target update from(mapper(id), present: c)
}
Index: clang/lib/CodeGen/CGOpenMPRuntime.cpp
===================================================================
--- clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -9515,7 +9515,8 @@
/// \code
/// void .omp_mapper.<type_name>.<mapper_id>.(void *rt_mapper_handle,
/// void *base, void *begin,
-/// int64_t size, int64_t type) {
+/// int64_t size, int64_t type,
+/// void *name = nullptr) {
/// // Allocate space for an array section first.
/// if (size > 1 && !maptype.IsDelete)
/// __tgt_push_mapper_component(rt_mapper_handle, base, begin,
@@ -9526,10 +9527,11 @@
/// for (auto c : all_components) {
/// if (c.hasMapper())
/// (*c.Mapper())(rt_mapper_handle, c.arg_base, c.arg_begin, c.arg_size,
-/// c.arg_type);
+/// c.arg_type, c.arg_name);
/// else
/// __tgt_push_mapper_component(rt_mapper_handle, c.arg_base,
-/// c.arg_begin, c.arg_size, c.arg_type);
+/// c.arg_begin, c.arg_size, c.arg_type,
+/// c.arg_name);
/// }
/// }
/// // Delete the array section.
@@ -9562,12 +9564,15 @@
ImplicitParamDecl::Other);
ImplicitParamDecl TypeArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr, Int64Ty,
ImplicitParamDecl::Other);
+ ImplicitParamDecl NameArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr, C.VoidPtrTy,
+ ImplicitParamDecl::Other);
FunctionArgList Args;
Args.push_back(&HandleArg);
Args.push_back(&BaseArg);
Args.push_back(&BeginArg);
Args.push_back(&SizeArg);
Args.push_back(&TypeArg);
+ Args.push_back(&NameArg);
const CGFunctionInfo &FnInfo =
CGM.getTypes().arrangeBuiltinFunctionDeclaration(C.VoidTy, Args);
llvm::FunctionType *FnTy = CGM.getTypes().GetFunctionType(FnInfo);
@@ -9667,6 +9672,10 @@
llvm::Value *CurBeginArg = MapperCGF.Builder.CreateBitCast(
Info.Pointers[I], CGM.getTypes().ConvertTypeForMem(C.VoidPtrTy));
llvm::Value *CurSizeArg = Info.Sizes[I];
+ llvm::Value *CurNameArg =
+ (CGM.getCodeGenOpts().getDebugInfo() == codegenoptions::NoDebugInfo)
+ ? llvm::ConstantPointerNull::get(CGM.VoidPtrTy)
+ : emitMappingInformation(MapperCGF, OMPBuilder, Info.Exprs[I]);
// Extract the MEMBER_OF field from the map type.
llvm::BasicBlock *MemberBB = MapperCGF.createBasicBlock("omp.member");
@@ -9755,8 +9764,8 @@
CurMapType->addIncoming(FromMapType, FromBB);
CurMapType->addIncoming(MemberMapType, ToElseBB);
- llvm::Value *OffloadingArgs[] = {Handle, CurBaseArg, CurBeginArg,
- CurSizeArg, CurMapType};
+ llvm::Value *OffloadingArgs[] = {Handle, CurBaseArg, CurBeginArg,
+ CurSizeArg, CurMapType, CurNameArg};
if (Info.Mappers[I]) {
// Call the corresponding mapper function.
llvm::Function *MapperFunc = getOrCreateUserDefinedMapperFunc(
@@ -9846,9 +9855,12 @@
MapType,
MapperCGF.Builder.getInt64(~(MappableExprsHandler::OMP_MAP_TO |
MappableExprsHandler::OMP_MAP_FROM)));
+ llvm::Value *MapNameArg = llvm::ConstantPointerNull::get(CGM.VoidPtrTy);
+
// Call the runtime API __tgt_push_mapper_component to fill up the runtime
// data structure.
- llvm::Value *OffloadingArgs[] = {Handle, Base, Begin, ArraySize, MapTypeArg};
+ llvm::Value *OffloadingArgs[] = {Handle, Base, Begin,
+ ArraySize, MapTypeArg, MapNameArg};
MapperCGF.EmitRuntimeCall(
OMPBuilder.getOrCreateRuntimeFunction(CGM.getModule(),
OMPRTL___tgt_push_mapper_component),
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
