Author: Joseph Huber Date: 2021-01-21T09:26:44-05:00 New Revision: e4eaf9d82064901ff028332d1644eddafac73f69
URL: https://github.com/llvm/llvm-project/commit/e4eaf9d82064901ff028332d1644eddafac73f69 DIFF: https://github.com/llvm/llvm-project/commit/e4eaf9d82064901ff028332d1644eddafac73f69.diff LOG: [OpenMP] Add support for mapping names in mapper API Summary: 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. Reviewers: jdoerfert Differential Revision: https://reviews.llvm.org/D94806 Added: Modified: clang/lib/CodeGen/CGOpenMPRuntime.cpp clang/test/OpenMP/declare_mapper_codegen.cpp clang/test/OpenMP/target_depend_codegen.cpp clang/test/OpenMP/target_map_names.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 Removed: ################################################################################ diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index 22df862db1b5..7a69fe2c013f 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -9502,7 +9502,8 @@ getNestedDistributeDirective(ASTContext &Ctx, const OMPExecutableDirective &D) { /// \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, @@ -9513,10 +9514,11 @@ getNestedDistributeDirective(ASTContext &Ctx, const OMPExecutableDirective &D) { /// 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. @@ -9549,12 +9551,15 @@ void CGOpenMPRuntime::emitUserDefinedMapper(const OMPDeclareMapperDecl *D, 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); @@ -9654,6 +9659,10 @@ void CGOpenMPRuntime::emitUserDefinedMapper(const OMPDeclareMapperDecl *D, 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"); @@ -9742,8 +9751,8 @@ void CGOpenMPRuntime::emitUserDefinedMapper(const OMPDeclareMapperDecl *D, 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( @@ -9833,9 +9842,12 @@ void CGOpenMPRuntime::emitUDMapperArrayInitOrDel( 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), diff --git a/clang/test/OpenMP/declare_mapper_codegen.cpp b/clang/test/OpenMP/declare_mapper_codegen.cpp index 5e5b175cb5e9..4bb9ca975dc2 100644 --- a/clang/test/OpenMP/declare_mapper_codegen.cpp +++ b/clang/test/OpenMP/declare_mapper_codegen.cpp @@ -86,7 +86,7 @@ class C { #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 @@ class C { // 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 @@ class C { // 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 @@ class C { // 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 @@ class C { // 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 @@ class C { // 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 @@ void foo(int a){ // 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 @@ void foo(int a){ // 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 @@ void foo(int a){ // 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 @@ void foo(int a){ // 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 @@ void foo(int a){ // 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 @@ void foo(int a){ // 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 @@ void foo(int a){ // 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 @@ void foo(int a){ // 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 @@ void foo(int a){ // 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 @@ void foo(int a){ // 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 @@ class C { #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 @@ class C { // 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 @@ class C { // 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 @@ class C { // 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 @@ class C { #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 @@ class C { // 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 @@ class C { // 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 @@ class C { // 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 @@ class B { #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 @@ void foo(int a){ // 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 @@ void foo(int a){ // 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 @@ class C { #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 @@ class C { // 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 @@ class C { // 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 @@ class C { // 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 @@ class C { // 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 @@ class C { // 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 @@ void foo(int a){ // 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 @@ void foo(int a){ // 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) } diff --git a/clang/test/OpenMP/target_depend_codegen.cpp b/clang/test/OpenMP/target_depend_codegen.cpp index a7dcbe7cc6ed..45b53f553a11 100644 --- a/clang/test/OpenMP/target_depend_codegen.cpp +++ b/clang/test/OpenMP/target_depend_codegen.cpp @@ -130,7 +130,7 @@ int foo(int n) { // 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 diff --git a/clang/test/OpenMP/target_map_names.cpp b/clang/test/OpenMP/target_map_names.cpp index bc68afd2166d..b7ac4f8bbfb6 100644 --- a/clang/test/OpenMP/target_map_names.cpp +++ b/clang/test/OpenMP/target_map_names.cpp @@ -167,6 +167,20 @@ void baz() { #pragma omp target update to(t) nowait } +struct S3 { + double Z[64]; +}; + +#pragma omp declare mapper(id: S3 s) map(s.Z[0:64]) + +void qux() { + S3 s; +#pragma omp target map(mapper(id), to:s) + { } +} + +// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s.Z[0:64];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00" + // DEBUG: %{{.+}} = call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 -1, i8* @{{.+}}, i32 1, i8** %{{.+}}, i8** %{{.+}}, i64* {{.+}}, i64* {{.+}}, i8** getelementptr inbounds ([{{[0-9]+}} x i8*], [{{[0-9]+}} x i8*]* @.offload_mapnames{{.*}}, i32 0, i32 0), i8** {{.+}}) // DEBUG: %{{.+}} = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, i64 -1, i8* @{{.+}}, i32 1, i8** %{{.+}}, i8** %{{.+}}, i64* {{.+}}, i64* {{.+}}, i8** getelementptr inbounds ([{{[0-9]+}} x i8*], [{{[0-9]+}} x i8*]* @.offload_mapnames{{.*}}, i32 0, i32 0), i8** {{.+}}, i32 {{.+}}, i32 {{.+}}) // DEBUG: call void @__tgt_target_data_begin_mapper(%struct.ident_t* @{{.+}}, i64 -1, i32 1, i8** %{{.+}}, i8** %{{.+}}, i64* {{.+}}, i64* {{.+}}, i8** getelementptr inbounds ([{{[0-9]+}} x i8*], [{{[0-9]+}} x i8*]* @.offload_mapnames{{.*}}, i32 0, i32 0), i8** {{.+}}) diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def b/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def index 38496d3ba983..844046167975 100644 --- a/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def +++ b/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def @@ -402,7 +402,7 @@ __OMP_RTL(__tgt_target_data_update_nowait_mapper, false, Void, IdentPtr, Int64, 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) diff --git a/llvm/test/Transforms/OpenMP/add_attributes.ll b/llvm/test/Transforms/OpenMP/add_attributes.ll index 4ce87160516d..b294542667bd 100644 --- a/llvm/test/Transforms/OpenMP/add_attributes.ll +++ b/llvm/test/Transforms/OpenMP/add_attributes.ll @@ -657,7 +657,7 @@ declare void @__tgt_target_data_update_nowait_mapper(%struct.ident_t*, i64, i32, 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*) @@ -1189,7 +1189,7 @@ declare void @__kmpc_proxy_task_completed_ooo(i8*) ; CHECK-NEXT: declare i64 @__tgt_mapper_num_components(i8*) ; CHECK: ; Function Attrs: nounwind -; CHECK-NEXT: declare void @__tgt_push_mapper_component(i8*, i8*, i8*, i64, i64) +; CHECK-NEXT: declare void @__tgt_push_mapper_component(i8*, i8*, i8*, i64, i64, i8*) ; CHECK: ; Function Attrs: nounwind ; CHECK-NEXT: declare i8* @__kmpc_task_allow_completion_event(%struct.ident_t*, i32, i8*) @@ -1714,7 +1714,7 @@ declare void @__kmpc_proxy_task_completed_ooo(i8*) ; OPTIMISTIC-NEXT: declare i64 @__tgt_mapper_num_components(i8*) ; OPTIMISTIC: ; Function Attrs: nounwind -; OPTIMISTIC-NEXT: declare void @__tgt_push_mapper_component(i8*, i8*, i8*, i64, i64) +; OPTIMISTIC-NEXT: declare void @__tgt_push_mapper_component(i8*, i8*, i8*, i64, i64, i8*) ; OPTIMISTIC: ; Function Attrs: nofree nosync nounwind willreturn ; OPTIMISTIC-NEXT: declare noalias i8* @__kmpc_task_allow_completion_event(%struct.ident_t* nocapture nofree readonly, i32, i8*) diff --git a/openmp/libomptarget/src/interface.cpp b/openmp/libomptarget/src/interface.cpp index b5af0b140585..c773e1fdac8e 100644 --- a/openmp/libomptarget/src/interface.cpp +++ b/openmp/libomptarget/src/interface.cpp @@ -487,16 +487,17 @@ EXTERN int64_t __tgt_mapper_num_components(void *rt_mapper_handle) { // 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, diff --git a/openmp/libomptarget/src/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp index 0846e99461ff..8cb16a489699 100644 --- a/openmp/libomptarget/src/omptarget.cpp +++ b/openmp/libomptarget/src/omptarget.cpp @@ -209,15 +209,16 @@ static int32_t getParentIndex(int64_t type) { /// 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 @@ int targetDataMapper(DeviceTy &Device, void *arg_base, void *arg, 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 @@ int targetDataMapper(DeviceTy &Device, void *arg_base, void *arg, 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 @@ int targetDataBegin(DeviceTy &Device, int32_t arg_num, void **args_base, // 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" @@ -329,7 +334,7 @@ int targetDataBegin(DeviceTy &Device, int32_t arg_num, void **args_base, // PTR_AND_OBJ entry is handled below, and so the allocation might fail // when HasPresentModifier. PointerTgtPtrBegin = Device.getOrAllocTgtPtr( - HstPtrBase, HstPtrBase, sizeof(void *), HstPtrName, Pointer_IsNew, + HstPtrBase, HstPtrBase, sizeof(void *), nullptr, Pointer_IsNew, IsHostPtr, IsImplicit, UpdateRef, HasCloseModifier, HasPresentModifier); if (!PointerTgtPtrBegin) { @@ -464,8 +469,10 @@ int targetDataEnd(DeviceTy &Device, int32_t ArgNum, void **ArgBases, // 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 @@ int targetDataUpdate(DeviceTy &Device, int32_t ArgNum, void **ArgsBase, // 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" diff --git a/openmp/libomptarget/src/private.h b/openmp/libomptarget/src/private.h index 00826ae417dc..22672987ae29 100644 --- a/openmp/libomptarget/src/private.h +++ b/openmp/libomptarget/src/private.h @@ -48,9 +48,11 @@ struct MapComponentInfoTy { 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 @@ struct MapperComponentsTy { // 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). diff --git a/openmp/libomptarget/test/mapping/declare_mapper_api.cpp b/openmp/libomptarget/test/mapping/declare_mapper_api.cpp index 54a5ad61538b..eda0c86ee105 100644 --- a/openmp/libomptarget/test/mapping/declare_mapper_api.cpp +++ b/openmp/libomptarget/test/mapping/declare_mapper_api.cpp @@ -15,9 +15,10 @@ struct MapComponentInfoTy { 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 @@ extern "C" { #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 @@ int main(int argc, char *argv[]) { 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); _______________________________________________ llvm-branch-commits mailing list llvm-branch-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
