[OpenMP 5.0] Fix user-defined mapper privatization in tasks

[ABataev]

Summary

This patch fixes the problem that user-defined mapper array is not correctly privatized inside a task. This problem causes openmp/libomptarget/test/offloading/target_depend_nowait.cpp fails.

[grokos]

I tried the patch and indeed it fixes the problem with target_depend_nowait.cpp. I'll let someone else do the review though.

[protze.joachim]

I'm not enough invested in the compiler implementation to judge this patch.
I can confirm that this patch fixes the segmentation fault I have seen.

Are these tests supposed to work for offloading to nvidia GPU as well?
I tried and added:

// RUN: %libomptarget-compilexx-run-and-check-nvptx64-nvidia-cuda

which gives me:

ptxas /tmp/target_depend_nowait-37c5c4.s, line 355; error   : Call has wrong number of parameters
ptxas /tmp/target_depend_nowait-37c5c4.s, line 628; error   : Call has wrong number of parameters
ptxas /tmp/target_depend_nowait-37c5c4.s, line 2331; warning : Instruction 'vote' without '.sync' may produce unpredictable results on sm_70 and later architectures
ptxas /tmp/target_depend_nowait-37c5c4.s, line 2331; warning : Instruction 'vote' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
ptxas fatal   : Ptx assembly aborted due to errors

I guess this error is not directly related to your patch?

clang/lib/CodeGen/CGOpenMPRuntime.cpp

10524

Does this mean, that a task is only created if depend clauses are provided?
Does waiting with taskwait work without this wrapper task?

[lildmh]

The problem you had is probably not related to this patch. A lot of libomptarget tests are not working for nvptx now. I suppose someone is going to fix them.

clang/lib/CodeGen/CGOpenMPRuntime.cpp

10524

The case you describe, which is to have a task wrapped around target, is solved with modification on line 9592
This part is to deal with other cases, which is consistent with line 10455 if you look a bit further

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

8915

Do you really need this IsTask flag? The array is created always when Info.NumberOfPtrs > 0. Maybe, just do not check for HasMapper here?

10524

Better to add a flag RequiresOuterTask = D.hasClausesOfKind<OMPDependClause>() and use it here and where the task is generated (line 10455)

[lildmh]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

8915

This is an optimization. The mapper array can be optimized out when there is no actual mappers, so supposed to have less memory accesses in this common case. Not checking HasMapper will disable this optimization.

10524

Sure, please accept this patch after that.

[ABataev]

Add a test

[ye-luo]

What is the current status of this patch?
@lildmh could you update this patch? I'd like to test it against
https://bugs.llvm.org/show_bug.cgi?id=47122

[protze.joachim]

In D84470#2191752, @ABataev wrote:

Add a test

@ABataev did you ask for a codegen test, or is the runtime test openmp/libomptarget/test/offloading/target_depend_nowait.cpp enough?
The test fails without this patch and succeeds with this patch.

[ABataev]

In D84470#2223678, @protze.joachim wrote:

In D84470#2191752, @ABataev wrote:

Add a test

@ABataev did you ask for a codegen test, or is the runtime test openmp/libomptarget/test/offloading/target_depend_nowait.cpp enough?
The test fails without this patch and succeeds with this patch.

Yes, I know that it fails, but each functional change in the frontend must have the corresponding test in clang itself.

[lildmh]

I'll add a test to it. Please give me some time as I am busy with other stuffs.

[jdoerfert]

It looks like @lildmh is not able to provide a test any time soon. @grokos Could you take over this mapper patch as well? I think the current segfault state is rather unfortunate.

[ronlieb]

i applied this patch to our downstream build for amdgcn. it resolves the problem i was seeing with traps on tests with task wait depend tests from sollve.
perhaps rebase it to latest llvm.

[ABataev]

I can update this patch in a week when I'm back from vacation.

[ronlieb]

[AMD Official Use Only - Internal Distribution Only]

Hi Alexey
Any ETA on this one ?

Get Outlook for iOShttps://aka.ms/o0ukef

---

From: Alexey Bataev via Phabricator <[e-mail]>
Sent: Tuesday, September 8, 2020 1:39:06 PM
To: [e-mail] <[e-mail]>; [e-mail] <[e-mail]>; [e-mail] <[e-mail]>; [e-mail] <[e-mail]>
Cc: Lieberman, Ron <[e-mail]>; [e-mail] <[e-mail]>; [e-mail] <[e-mail]>; Liu, Yaxun (Sam) <[e-mail]>; [e-mail] <[e-mail]>; [e-mail] <[e-mail]>
Subject: [PATCH] D84470: [OpenMP 5.0] Fix user-defined mapper privatization in tasks

[CAUTION: External Email]

ABataev added a comment.

I can update this patch in a week when I'm back from vacation.

CHANGES SINCE LAST ACTION

https://nam11.safelinks.protection.outlook.com/?url=https%3A%2F%2Freviews.llvm.org%2FD84470%2Fnew%2F&data=02%7C01%7Cron.lieberman%40amd.com%7C6582a2f2d31d4bbf9b4a08d8541e16cb%7C3dd8961fe4884e608e11a82d994e183d%7C0%7C0%7C637351835496784231&sdata=8%2BJbFULmzVN%2FdafC0QYX%2B9LwN7d1HFZ9Qxt6k4y85%2FA%3D&reserved=0

https://nam11.safelinks.protection.outlook.com/?url=https%3A%2F%2Freviews.llvm.org%2FD84470&amp;data=02%7C01%7Cron.lieberman%40amd.com%7C6582a2f2d31d4bbf9b4a08d8541e16cb%7C3dd8961fe4884e608e11a82d994e183d%7C0%7C0%7C637351835496794187&amp;sdata=gyj3%2B4K3oPlvedNjOzbsnosciYUIxuTOoKq%2BsfnsAOQ%3D&amp;reserved=0

[ABataev]

Hi Ron, hope to prepare it by the end of this week, maybe earlier, if everything goes well. I have some technical issues, but anyway will update the patch and the tests.

[ABataev]

Updated + added clang test.

[ronlieb]

The latest patch applied cleanly to our downstream port.
builds fine, tests very nicely as well. All the failing SOLLVE task wait depend tests now pass.

[ABataev]

In D84470#2277991, @ronlieb wrote:

The latest patch applied cleanly to our downstream port.
builds fine, tests very nicely as well. All the failing SOLLVE task wait depend tests now pass.

Could you accept it then?

[ronlieb]

accepted. not sure if there were any outstanding comments from Johannes. your call

clang/lib/CodeGen/CGOpenMPRuntime.cpp

OpenMPDirectiveKind Kind = isOpenMPTaskLoopDirective(D.getDirectiveKind())
                                  ? OMPD_taskloop
                                  : OMPD_task;
   const CapturedStmt &CS = *D.getCapturedStmt(Kind);
   CodeGenFunction::CGCapturedStmtInfo CapturesInfo(CS);
   LValue SrcBase;
   bool IsTargetTask =
       isOpenMPTargetDataManagementDirective(D.getDirectiveKind()) ||
       isOpenMPTargetExecutionDirective(D.getDirectiveKind());
-  // For target-based directives skip 3 firstprivate arrays BasePointersArray,
-  // PointersArray and SizesArray. The original variables for these arrays are
-  // not captured and we get their addresses explicitly.
+  // For target-based directives skip 4 firstprivate arrays BasePointersArray,
+  // PointersArray, SizesArray, and MappersArray. The original variables for
+  // these arrays are not captured and we get their addresses explicitly.
   if ((!IsTargetTask && !Data.FirstprivateVars.empty() && ForDup) ||
       (IsTargetTask && KmpTaskSharedsPtr.isValid())) {
     SrcBase = CGF.MakeAddrLValue(
         CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
             KmpTaskSharedsPtr, CGF.ConvertTypeForMem(SharedsPtrTy)),
         SharedsTy);
   }
   FI = cast<RecordDecl>(FI->getType()->getAsTagDecl())->field_begin();
   for (const PrivateDataTy &Pair : Privates) {
     // Do not initialize private locals.
     if (Pair.second.isLocalPrivate()) {
       ++FI;
       continue;
     }
     const VarDecl *VD = Pair.second.PrivateCopy;
     const Expr *Init = VD->getAnyInitializer();
     if (Init && (!ForDup || (isa<CXXConstructExpr>(Init) &&
                              !CGF.isTrivialInitializer(Init)))) {
       LValue PrivateLValue = CGF.EmitLValueForField(PrivatesBase, *FI);
       if (const VarDecl *Elem = Pair.second.PrivateElemInit) {
         const VarDecl *OriginalVD = Pair.second.Original;
         // Check if the variable is the target-based BasePointersArray,
-        // PointersArray or SizesArray.
+        // PointersArray, SizesArray, or MappersArray.
         LValue SharedRefLValue;
         QualType Type = PrivateLValue.getType();
         const FieldDecl *SharedField = CapturesInfo.lookup(OriginalVD);
         if (IsTargetTask && !SharedField) {
           assert(isa<ImplicitParamDecl>(OriginalVD) &&
                  isa<CapturedDecl>(OriginalVD->getDeclContext()) &&
                  cast<CapturedDecl>(OriginalVD->getDeclContext())
                          ->getNumParams() == 0 &&
for (unsigned I = 0; I < Info.NumberOfPtrs; ++I) {
         Info.HasMapper = true;
       }
       Address MAddr = CGF.Builder.CreateConstArrayGEP(MappersArray, I);
       CGF.Builder.CreateStore(MFunc, MAddr);
     }
   }
 }
 
+namespace {
+/// Additional arguments for emitOffloadingArraysArgument function.
+struct ArgumentsOptions {
+  bool ForEndCall = false;
+  bool IsTask = false;
+  ArgumentsOptions() = default;
+  ArgumentsOptions(bool ForEndCall, bool IsTask)
+      : ForEndCall(ForEndCall), IsTask(IsTask) {}
+};
+} // namespace
+
 /// Emit the arguments to be passed to the runtime library based on the
 /// arrays of base pointers, pointers, sizes, map types, and mappers.  If
 /// ForEndCall, emit map types to be passed for the end of the region instead of
 /// the beginning.
 static void emitOffloadingArraysArgument(
     CodeGenFunction &CGF, llvm::Value *&BasePointersArrayArg,
     llvm::Value *&PointersArrayArg, llvm::Value *&SizesArrayArg,
     llvm::Value *&MapTypesArrayArg, llvm::Value *&MappersArrayArg,
-    CGOpenMPRuntime::TargetDataInfo &Info, bool ForEndCall = false) {
-  assert((!ForEndCall || Info.separateBeginEndCalls()) &&
+    CGOpenMPRuntime::TargetDataInfo &Info,
+    const ArgumentsOptions &Options = ArgumentsOptions()) {
+  assert((!Options.ForEndCall || Info.separateBeginEndCalls()) &&
          "expected region end call to runtime only when end call is separate");
   CodeGenModule &CGM = CGF.CGM;
   if (Info.NumberOfPtrs) {
     BasePointersArrayArg = CGF.Builder.CreateConstInBoundsGEP2_32(
         llvm::ArrayType::get(CGM.VoidPtrTy, Info.NumberOfPtrs),
         Info.BasePointersArray,
         /*Idx0=*/0, /*Idx1=*/0);
     PointersArrayArg = CGF.Builder.CreateConstInBoundsGEP2_32(
         llvm::ArrayType::get(CGM.VoidPtrTy, Info.NumberOfPtrs),
         Info.PointersArray,
         /*Idx0=*/0,
         /*Idx1=*/0);
     SizesArrayArg = CGF.Builder.CreateConstInBoundsGEP2_32(
         llvm::ArrayType::get(CGM.Int64Ty, Info.NumberOfPtrs), Info.SizesArray,
         /*Idx0=*/0, /*Idx1=*/0);
     MapTypesArrayArg = CGF.Builder.CreateConstInBoundsGEP2_32(
         llvm::ArrayType::get(CGM.Int64Ty, Info.NumberOfPtrs),
-        ForEndCall && Info.MapTypesArrayEnd ? Info.MapTypesArrayEnd
+        Options.ForEndCall && Info.MapTypesArrayEnd ? Info.MapTypesArrayEnd
                                                     : Info.MapTypesArray,
         /*Idx0=*/0,
         /*Idx1=*/0);
+    // Always emit the mapper array address in case of a target task for
+    // privatization.
+    if (!Options.IsTask && !Info.HasMapper)
+      MappersArrayArg = llvm::ConstantPointerNull::get(CGM.VoidPtrPtrTy);
+    else
       MappersArrayArg =
-        Info.HasMapper
-            ? CGF.Builder.CreatePointerCast(Info.MappersArray, CGM.VoidPtrPtrTy)
-            : llvm::ConstantPointerNull::get(CGM.VoidPtrPtrTy);
+          CGF.Builder.CreatePointerCast(Info.MappersArray, CGM.VoidPtrPtrTy);
   } 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);
   }
auto &&TargetThenGen = [this, &ThenGen, &D, &InputInfo, &MapTypesArray,
     // Map any list items in a map clause that were not captures because they
     // weren't referenced within the construct.
     MEHandler.generateAllInfo(CombinedInfo, /*NotTargetParams=*/true,
                               MappedVarSet);
 
     TargetDataInfo Info;
     // Fill up the arrays and create the arguments.
     emitOffloadingArrays(CGF, CombinedInfo, Info);
+    bool HasDependClauses = D.hasClausesOfKind<OMPDependClause>();
     emitOffloadingArraysArgument(CGF, Info.BasePointersArray,
                                  Info.PointersArray, Info.SizesArray,
-                                 Info.MapTypesArray, Info.MappersArray, Info);
+                                 Info.MapTypesArray, Info.MappersArray, Info,
+                                 {/*ForEndTask=*/false, HasDependClauses});
     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;
auto &&BeginThenGen = [this, &D, Device, &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,
-                                 MappersArrayArg, Info, /*ForEndCall=*/false);
+                                 MappersArrayArg, Info);
 
     // Emit device ID if any.
     llvm::Value *DeviceID = nullptr;
     if (Device) {
       DeviceID = CGF.Builder.CreateIntCast(CGF.EmitScalarExpr(Device),
                                            CGF.Int64Ty, /*isSigned=*/true);
     } else {
       DeviceID = CGF.Builder.getInt64(OMP_DEVICEID_UNDEF);
auto &&EndThenGen = [this, Device, &Info](CodeGenFunction &CGF,
 
     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,
-                                 MappersArrayArg, Info, /*ForEndCall=*/true);
+                                 MappersArrayArg, Info,
+                                 {/*ForEndCall=*/true, /*IsTask=*/false});
 
     // Emit device ID if any.
     llvm::Value *DeviceID = nullptr;
     if (Device) {
       DeviceID = CGF.Builder.CreateIntCast(CGF.EmitScalarExpr(Device),
                                            CGF.Int64Ty, /*isSigned=*/true);
     } else {
       DeviceID = CGF.Builder.getInt64(OMP_DEVICEID_UNDEF);
auto &&TargetThenGen = [this, &ThenGen, &D, &InputInfo, &MapTypesArray](
 
     // Get map clause information.
     MappableExprsHandler MEHandler(D, CGF);
     MEHandler.generateAllInfo(CombinedInfo);
 
     TargetDataInfo Info;
     // Fill up the arrays and create the arguments.
     emitOffloadingArrays(CGF, CombinedInfo, Info);
+    bool HasDependClauses = D.hasClausesOfKind<OMPDependClause>();
     emitOffloadingArraysArgument(CGF, Info.BasePointersArray,
                                  Info.PointersArray, Info.SizesArray,
-                                 Info.MapTypesArray, Info.MappersArray, Info);
+                                 Info.MapTypesArray, Info.MappersArray, Info,
+                                 {/*ForEndTask=*/false, HasDependClauses});
     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 (D.hasClausesOfKind<OMPDependClause>())
+    if (HasDependClauses)
       CGF.EmitOMPTargetTaskBasedDirective(D, ThenGen, InputInfo);
     else
       emitInlinedDirective(CGF, D.getDirectiveKind(), ThenGen);
   };
 
   if (IfCond) {
     emitIfClause(CGF, IfCond, TargetThenGen,
                  [](CodeGenFunction &CGF, PrePostActionTy &) {});

clang/test/OpenMP/target_depend_codegen.cpp

 #ifndef HEADER
 #define HEADER
 
 // CHECK-DAG: [[TT:%.+]] = type { i64, i8 }
 // CHECK-DAG: [[ENTTY:%.+]] = type { i8*, i8*, i[[SZ:32|64]], i32, i32 }
 
 // TCHECK: [[ENTTY:%.+]] = type { i8*, i8*, i{{32|64}}, i32, i32 }
 
-// CHECK-DAG: [[SIZET:@.+]] = private unnamed_addr constant [2 x i64] [i64 0, i64 4]
-// CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [2 x i64] [i64 544, i64 800]
+// CHECK-DAG: [[SIZET:@.+]] = private unnamed_addr constant [3 x i64] [i64 0, i64 4, i64 {{16|12}}]
+// CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [3 x i64] [i64 544, i64 800, i64 3]
 // CHECK-DAG: @{{.*}} = weak constant i8 0
 
 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 // TCHECK: @{{.+}} = {{.*}}constant [[ENTTY]]
 // TCHECK-NOT: @{{.+}} = weak constant [[ENTTY]]
 
 // Check target registration is registered as a Ctor.
 // CHECK: appending global [1 x { i32, void ()*, i8* }] [{ i32, void ()*, i8* } { i32 0, void ()* @.omp_offloading.requires_reg, i8* null }]
 
 
 template<typename tx, typename ty>
 struct TT{
   tx X;
   ty Y;
 };
 
+#pragma omp declare mapper(id                     \
+                           : TT <long long, char> \
+                               s) map(s.X, s.Y)
 int global;
 extern int global;
 
 // CHECK: define {{.*}}[[FOO:@.+]](
 int foo(int n) {
   int a = 0;
   short aa = 0;
   float b[10];
#pragma omp target device(global + a) depend(in                   \
   }
 
   // CHECK:       [[ADD:%.+]] = add nsw i32
   // CHECK:       store i32 [[ADD]], i32* [[DEVICE_CAP:%.+]],
 
   // CHECK:       [[BOOL:%.+]] = icmp ne i32 %{{.+}}, 0
   // CHECK:       br i1 [[BOOL]], label %[[THEN:.+]], label %[[ELSE:.+]]
   // CHECK:       [[THEN]]:
-  // CHECK-DAG:   [[BPADDR0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BP:%.+]], i32 0, i32 0
-  // CHECK-DAG:   [[PADDR0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[P:%.+]], i32 0, i32 0
+  // CHECK-DAG:   [[BPADDR0:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BP:%.+]], i32 0, i32 0
+  // CHECK-DAG:   [[PADDR0:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[P:%.+]], i32 0, i32 0
+  // CHECK-DAG:   [[MADDR0:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[M:%.+]], i[[SZ]] 0, i[[SZ]] 0
   // CHECK-DAG:   [[CBPADDR0:%.+]] = bitcast i8** [[BPADDR0]] to i[[SZ]]**
   // CHECK-DAG:   [[CPADDR0:%.+]] = bitcast i8** [[PADDR0]] to i[[SZ]]**
   // CHECK-DAG:   store i[[SZ]]* [[BP0:%[^,]+]], i[[SZ]]** [[CBPADDR0]]
   // CHECK-DAG:   store i[[SZ]]* [[BP0]], i[[SZ]]** [[CPADDR0]]
+  // CHECK-DAG:   store i8* null, i8** [[MADDR0]],
 
-  // CHECK-DAG:   [[BPADDR1:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BP]], i32 0, i32 1
-  // CHECK-DAG:   [[PADDR1:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[P]], i32 0, i32 1
+  // CHECK-DAG:   [[BPADDR1:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BP]], i32 0, i32 1
+  // CHECK-DAG:   [[PADDR1:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[P]], i32 0, i32 1
+  // CHECK-DAG:   [[MADDR1:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[M]], i[[SZ]] 0, i[[SZ]] 1
   // CHECK-DAG:   [[CBPADDR1:%.+]] = bitcast i8** [[BPADDR1]] to i[[SZ]]*
   // CHECK-DAG:   [[CPADDR1:%.+]] = bitcast i8** [[PADDR1]] to i[[SZ]]*
   // CHECK-DAG:   store i[[SZ]] [[BP1:%[^,]+]], i[[SZ]]* [[CBPADDR1]]
   // CHECK-DAG:   store i[[SZ]] [[BP1]], i[[SZ]]* [[CPADDR1]]
-  // CHECK-DAG:   getelementptr inbounds [2 x i8*], [2 x i8*]* [[BP]], i32 0, i32 0
-  // CHECK-DAG:   getelementptr inbounds [2 x i8*], [2 x i8*]* [[P]], i32 0, i32 0
+  // CHECK-DAG:   store i8* null, i8** [[MADDR1]],
+
+  // CHECK-DAG:   [[BPADDR2:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BP]], i32 0, i32 2
+  // CHECK-DAG:   [[PADDR2:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[P]], i32 0, i32 2
+  // CHECK-DAG:   [[MADDR2:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[M]], i[[SZ]] 0, i[[SZ]] 2
+  // CHECK-DAG:   [[CBPADDR2:%.+]] = bitcast i8** [[BPADDR2]] to [[STRUCT_TT:%.+]]**
+  // 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:   [[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
+  // CHECK-DAG:   [[M_START:%.+]] = bitcast [3 x i8*]* [[M]] to i8**
   // CHECK:       [[GEP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2
   // CHECK:       [[DEV:%.+]] = load i32, i32* [[DEVICE_CAP]],
   // CHECK:       store i32 [[DEV]], i32* [[GEP]],
   // CHECK:       [[DEV1:%.+]] = load i32, i32* [[DEVICE_CAP]],
   // CHECK:       [[DEV2:%.+]] = sext i32 [[DEV1]] to i64
 
-  // CHECK:       [[TASK:%.+]] = call i8* @__kmpc_omp_target_task_alloc(%struct.ident_t* @{{.*}}, i32 [[GTID]], i32 1, i[[SZ]] {{120|68}}, i[[SZ]] {{16|12}}, i32 (i32, i8*)* bitcast (i32 (i32, %{{.+}}*)* [[TASK_ENTRY1_:@.+]] to i32 (i32, i8*)*), i64 [[DEV2]])
+  // CHECK:       [[TASK:%.+]] = call i8* @__kmpc_omp_target_task_alloc(%struct.ident_t* @{{.*}}, i32 [[GTID]], i32 1, i[[SZ]] {{152|88}}, i[[SZ]] {{16|12}}, i32 (i32, i8*)* bitcast (i32 (i32, %{{.+}}*)* [[TASK_ENTRY1_:@.+]] to i32 (i32, i8*)*), i64 [[DEV2]])
   // CHECK:       [[BC_TASK:%.+]] = bitcast i8* [[TASK]] to [[TASK_TY1_:%.+]]*
+  // CHECK:       [[BASE:%.+]] = getelementptr inbounds [[TASK_TY1_]], [[TASK_TY1_]]* [[BC_TASK]], i32 0, i32 1
+  // CHECK-64:    [[BP_BASE:%.+]] = getelementptr inbounds [[PRIVS_TY:%.+]], [[PRIVS_TY:%.+]]* [[BASE]], i32 0, i32 1
+  // CHECK-64:    [[BP_CAST:%.+]] = bitcast [3 x i8*]* [[BP_BASE]] to i8*
+  // CHECK-64:    [[BP_SRC:%.+]] = bitcast i8** [[BP_START]] to i8*
+  // CHECK-64:    call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 8 [[BP_CAST]], i8* align 8 [[BP_SRC]], i64 24, i1 false)
+  // CHECK-64:    [[P_BASE:%.+]] = getelementptr inbounds [[PRIVS_TY]], [[PRIVS_TY]]* [[BASE]], i32 0, i32 2
+  // CHECK-64:    [[P_CAST:%.+]] = bitcast [3 x i8*]* [[P_BASE]] to i8*
+  // CHECK-64:    [[P_SRC:%.+]] = bitcast i8** [[P_START]] to i8*
+  // CHECK-64:    call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 8 [[P_CAST]], i8* align 8 [[P_SRC]], i64 24, i1 false)
+  // CHECK-64:    [[SZ_BASE:%.+]] = getelementptr inbounds [[PRIVS_TY]], [[PRIVS_TY]]* [[BASE]], i32 0, i32 3
+  // CHECK-64:    [[SZ_CAST:%.+]] = bitcast [3 x i64]* [[SZ_BASE]] to i8*
+  // CHECK-64:    call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 8 [[SZ_CAST]], i8* align 8 bitcast ([3 x i64]* [[SIZET]] to i8*), i64 24, i1 false)
+  // CHECK-64:    [[M_BASE:%.+]] = getelementptr inbounds [[PRIVS_TY]], [[PRIVS_TY]]* [[BASE]], i32 0, i32 4
+  // CHECK-64:    [[M_CAST:%.+]] = bitcast [3 x i8*]* [[M_BASE]] to i8*
+  // CHECK-64:    [[M_SRC:%.+]] = bitcast i8** [[M_START]] to i8*
+  // CHECK-64:    call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 8 [[M_CAST]], i8* align 8 [[M_SRC]], i64 24, i1 false)
+  // CHECK-32:    [[SZ_BASE:%.+]] = getelementptr inbounds [[PRIVS_TY:%.+]], [[PRIVS_TY:%.+]]* [[BASE]], i32 0, i32 0
+  // CHECK-32:    [[SZ_CAST:%.+]] = bitcast [3 x i64]* [[SZ_BASE]] to i8*
+  // CHECK-32:    call void @llvm.memcpy.p0i8.p0i8.i32(i8* align 4 [[SZ_CAST]], i8* align 4 bitcast ([3 x i64]* [[SIZET]] to i8*), i32 24, i1 false)
+  // CHECK-32:    [[BP_BASE:%.+]] = getelementptr inbounds [[PRIVS_TY]], [[PRIVS_TY]]* [[BASE]], i32 0, i32 3
+  // CHECK-32:    [[BP_CAST:%.+]] = bitcast [3 x i8*]* [[BP_BASE]] to i8*
+  // CHECK-32:    [[BP_SRC:%.+]] = bitcast i8** [[BP_START]] to i8*
+  // CHECK-32:    call void @llvm.memcpy.p0i8.p0i8.i32(i8* align 4 [[BP_CAST]], i8* align 4 [[BP_SRC]], i32 12, i1 false)
+  // CHECK-32:    [[P_BASE:%.+]] = getelementptr inbounds [[PRIVS_TY]], [[PRIVS_TY]]* [[BASE]], i32 0, i32 4
+  // CHECK-32:    [[P_CAST:%.+]] = bitcast [3 x i8*]* [[P_BASE]] to i8*
+  // CHECK-32:    [[P_SRC:%.+]] = bitcast i8** [[P_START]] to i8*
+  // CHECK-32:    call void @llvm.memcpy.p0i8.p0i8.i32(i8* align 4 [[P_CAST]], i8* align 4 [[P_SRC]], i32 12, i1 false)
+  // CHECK-32:    [[M_BASE:%.+]] = getelementptr inbounds [[PRIVS_TY]], [[PRIVS_TY]]* [[BASE]], i32 0, i32 5
+  // CHECK-32:    [[M_CAST:%.+]] = bitcast [3 x i8*]* [[M_BASE]] to i8*
+  // CHECK-32:    [[M_SRC:%.+]] = bitcast i8** [[M_START]] to i8*
+  // CHECK-32:    call void @llvm.memcpy.p0i8.p0i8.i32(i8* align 4 [[M_CAST]], i8* align 4 [[M_SRC]], i32 12, i1 false)
   // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP_START:%.+]], i[[SZ]] 1
   // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP_START]], i[[SZ]] 2
   // CHECK:       [[DEP:%.+]] = bitcast %struct.kmp_depend_info* [[DEP_START]] to i8*
   // CHECK:       call i32 @__kmpc_omp_task_with_deps(%struct.ident_t* @1, i32 [[GTID]], i8* [[TASK]], i32 3, i8* [[DEP]], i32 0, i8* null)
   // CHECK:       br label %[[EXIT:.+]]
 
   // CHECK:       [[ELSE]]:
   // CHECK-NOT:   getelementptr inbounds [2 x i8*], [2 x i8*]*
   // CHECK:       [[GEP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2
   // CHECK:       [[DEV:%.+]] = load i32, i32* [[DEVICE_CAP]],
   // CHECK:       store i32 [[DEV]], i32* [[GEP]],
   // CHECK:       [[DEV1:%.+]] = load i32, i32* [[DEVICE_CAP]],
   // CHECK:       [[DEV2:%.+]] = sext i32 [[DEV1]] to i64
 
   // CHECK:       [[TASK:%.+]] = call i8* @__kmpc_omp_target_task_alloc(%struct.ident_t* @1, i32 [[GTID]], i32 1, i[[SZ]] {{56|28}}, i[[SZ]] {{16|12}}, i32 (i32, i8*)* bitcast (i32 (i32, %{{.+}}*)* [[TASK_ENTRY1__:@.+]] to i32 (i32, i8*)*), i64 [[DEV2]])
   // CHECK:       [[BC_TASK:%.+]] = bitcast i8* [[TASK]] to [[TASK_TY1__:%.+]]*
   // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP_START:%.+]], i[[SZ]] 1
   // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP_START]], i[[SZ]] 2
   // CHECK:       [[DEP:%.+]] = bitcast %struct.kmp_depend_info* [[DEP_START]] to i8*
   // CHECK:       call i32 @__kmpc_omp_task_with_deps(%struct.ident_t* @1, i32 [[GTID]], i8* [[TASK]], i32 3, i8* [[DEP]], i32 0, i8* null)
   // CHECK:       br label %[[EXIT:.+]]
   // CHECK:       [[EXIT]]:
 
-#pragma omp target device(global + a) nowait depend(inout \
-                                                    : global, a, bn) if (a)
+#pragma omp target device(global + a) nowait depend(inout                                          \
+                                                    : global, a, bn) if (a) map(mapper(id), tofrom \
+                                                                                : d)
   {
     static int local1;
     *plocal = global;
     local1 = global;
   }
 
 // CHECK:       [[TASK:%.+]] = call i8* @__kmpc_omp_task_alloc(%struct.ident_t* @1, i32 [[GTID]], i32 1, i[[SZ]] {{48|24}}, i[[SZ]] 4, i32 (i32, i8*)* bitcast (i32 (i32, %{{.+}}*)* [[TASK_ENTRY2:@.+]] to i32 (i32, i8*)*))
 // CHECK:       [[BC_TASK:%.+]] = bitcast i8* [[TASK]] to [[TASK_TY2:%.+]]*

 // CHECK:       [[FAIL]]
 // CHECK:       call void [[HVT0]]()
 // CHECK-NEXT:  br label %[[END]]
 // CHECK:       [[END]]
 // CHECK:       ret i32 0
 
 // CHECK:       define internal void [[HVT1:@.+]](i[[SZ]]* %{{.+}}, i[[SZ]] %{{.+}})
 
-// CHECK:       define internal{{.*}} i32 [[TASK_ENTRY1_]](i32{{.*}}, [[TASK_TY1_]]* noalias %1)
-// CHECK:       call void (i8*, ...) %
-// CHECK:       [[SZT:%.+]] = getelementptr inbounds [2 x i64], [2 x i64]* %{{.+}}, i[[SZ]] 0, i[[SZ]] 0
+// CHECK:       define internal void [[MAPPER_ID]](i8* %{{.+}}, i8* %{{.+}}, i8* %{{.+}}, i64 %{{.+}}, i64 %{{.+}})
+
+// CHECK:       define internal{{.*}} i32 [[TASK_ENTRY1_]](i32{{.*}}, [[TASK_TY1_]]* noalias %{{.+}})
+// CHECK:       call void (i8*, ...) %{{.+}}(i8* %{{.+}}, i[[SZ]]*** %{{.+}}, i32** %{{.+}}, [3 x i8*]** [[BPTR_ADDR:%.+]], [3 x i8*]** [[PTR_ADDR:%.+]], [3 x i64]** [[SZ_ADDR:%.+]], [3 x i8*]** [[M_ADDR:%.+]])
+// CHECK:       [[BPTR_REF:%.+]] = load [3 x i8*]*, [3 x i8*]** [[BPTR_ADDR]],
+// CHECK:       [[PTR_REF:%.+]] = load [3 x i8*]*, [3 x i8*]** [[PTR_ADDR]],
+// CHECK:       [[SZ_REF:%.+]] = load [3 x i64]*, [3 x i64]** [[SZ_ADDR]],
+// CHECK:       [[M_REF:%.+]] = load [3 x i8*]*, [3 x i8*]** [[M_ADDR]],
+// CHECK:       [[BPR:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BPTR_REF]], i[[SZ]] 0, i[[SZ]] 0
+// CHECK:       [[PR:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PTR_REF]], i[[SZ]] 0, i[[SZ]] 0
+// CHECK:       [[SZT:%.+]] = getelementptr inbounds [3 x i64], [3 x i64]* [[SZ_REF]], i[[SZ]] 0, i[[SZ]] 0
+// CHECK:       [[M:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[M_REF]], i[[SZ]] 0, i[[SZ]] 0
 // CHECK:       [[DEVICE_CAP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2
 // CHECK:       [[DEV:%.+]] = load i32, i32* [[DEVICE_CAP]],
 // CHECK:       [[DEVICE:%.+]] = sext i32 [[DEV]] to i64
-// CHECK:       [[RET:%.+]] = call i32 @__tgt_target_nowait_mapper(i64 [[DEVICE]], i8* @{{[^,]+}}, i32 2, i8** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i64* [[SZT]], i64* getelementptr inbounds ([2 x i64], [2 x i64]* [[MAPT]], i32 0, i32 0), i8** [[M:%[^,]+]])
+// CHECK:       [[RET:%.+]] = call i32 @__tgt_target_nowait_mapper(i64 [[DEVICE]], i8* @{{[^,]+}}, i32 3, i8** [[BPR]], i8** [[PR]], i64* [[SZT]], i64* getelementptr inbounds ([3 x i64], [3 x i64]* [[MAPT]], i32 0, i32 0), i8** [[M]])
 
 // CHECK:       [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
 // CHECK-NEXT:  br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
 // CHECK:       [[FAIL]]
 // CHECK:       [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** %
 // CHECK:       [[BP1_I32:%.+]] = load i32, i32* @
 // CHECK-64:    [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
 // CHECK-64:    store i32 [[BP1_I32]], i32* [[BP1_CAST]],