diff --git a/clang/test/OpenMP/target_data_use_device_addr_codegen.cpp b/clang/test/OpenMP/target_data_use_device_addr_codegen.cpp
--- a/clang/test/OpenMP/target_data_use_device_addr_codegen.cpp
+++ b/clang/test/OpenMP/target_data_use_device_addr_codegen.cpp
@@ -11,13 +11,13 @@
 #ifndef HEADER
 #define HEADER
 
-// CHECK-DAG: [[SIZES1:@.+]] = private unnamed_addr constant [5 x i64] zeroinitializer
+// CHECK-DAG: [[SIZES1:@.+]] = private unnamed_addr constant [6 x i64] [i64 4, i64 16, i64 4, i64 4, i64 0, i64 4]
 // 64 = 0x40 = OMP_MAP_RETURN_PARAM
-// CHECK-DAG: [[MAPTYPES1:@.+]] = private unnamed_addr constant [5 x i64] [i64 64, i64 64, i64 64, i64 64, i64 64]
-// CHECK-DAG: [[SIZES2:@.+]] = private unnamed_addr constant [5 x i64] zeroinitializer
+// CHECK-DAG: [[MAPTYPES1:@.+]] = private unnamed_addr constant [6 x i64] [i64 67, i64 67, i64 3, i64 67, i64 67, i64 67]
+// CHECK-DAG: [[SIZES2:@.+]] = private unnamed_addr constant [6 x i64] [i64 0, i64 4, i64 16, i64 4, i64 4, i64 0]
 // 0 = OMP_MAP_NONE
 // 281474976710720 = 0x1000000000040 = OMP_MAP_MEMBER_OF | OMP_MAP_RETURN_PARAM
-// CHECK-DAG: [[MAPTYPES2:@.+]] = private unnamed_addr constant [5 x i64] [i64 0, i64 281474976710720, i64 281474976710720, i64 281474976710720, i64 281474976710720]
+// CHECK-DAG: [[MAPTYPES2:@.+]] = private unnamed_addr constant [6 x i64] [i64 0, i64 281474976710723, i64 281474976710739, i64 281474976710739, i64 281474976710675, i64 281474976710723]
 struct S {
   int a = 0;
   int *ptr = &a;
@@ -25,7 +25,7 @@
   int arr[4];
   S() {}
   void foo() {
-#pragma omp target data use_device_addr(a, ptr [3:4], ref, ptr[0], arr[:a])
+#pragma omp target data map(tofrom: a, ptr [3:4], ref, ptr[0], arr[:a]) use_device_addr(a, ptr [3:4], ref, ptr[0], arr[:a])
     ++a, ++*ptr, ++ref, ++arr[0];
   }
 };
@@ -38,7 +38,7 @@
   float vla[(int)a];
   S s;
   s.foo();
-#pragma omp target data use_device_addr(a, ptr [3:4], ref, ptr[0], arr[:(int)a], vla[0])
+#pragma omp target data map(tofrom: a, ptr [3:4], ref, ptr[0], arr[:(int)a], vla[0]) use_device_addr(a, ptr [3:4], ref, ptr[0], arr[:(int)a], vla[0])
   ++a, ++*ptr, ++ref, ++arr[0], ++vla[0];
   return a;
 }
@@ -48,51 +48,74 @@
 // CHECK: [[PTR_ADDR:%.+]] = alloca float*,
 // CHECK: [[REF_ADDR:%.+]] = alloca float*,
 // CHECK: [[ARR_ADDR:%.+]] = alloca [4 x float],
-// CHECK: [[BPTRS:%.+]] = alloca [5 x i8*],
-// CHECK: [[PTRS:%.+]] = alloca [5 x i8*],
+// CHECK: [[BPTRS:%.+]] = alloca [6 x i8*],
+// CHECK: [[PTRS:%.+]] = alloca [6 x i8*],
+// CHECK: [[MAP_PTRS:%.+]] = alloca [6 x i8*],
+// CHECK: [[SIZES:%.+]] = alloca [6 x i64],
 // CHECK: [[VLA_ADDR:%.+]] = alloca float, i64 %{{.+}},
 // CHECK: [[PTR:%.+]] = load float*, float** [[PTR_ADDR]],
-// CHECK: [[REF:%.+]] = load float*, float** [[REF_ADDR]],
-// CHECK: [[ARR:%.+]] = getelementptr inbounds [4 x float], [4 x float]* [[ARR_ADDR]], i64 0, i64 0
-// CHECK: [[BPTR0:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BPTRS]], i32 0, i32 0
+// CHECK-NEXT: [[P4:%.+]] = load float*, float** [[PTR_ADDR]], align 8
+// CHECK-NEXT: [[ARR_IDX:%.+]] = getelementptr inbounds float, float* [[P4]], i64 3
+// CHECK: [[P5:%.+]] = load float*, float** [[PTR_ADDR]], align 8
+// CHECK-NEXT: [[P6:%.+]] = load float*, float** [[PTR_ADDR]], align 8
+// CHECK-NEXT: [[ARR_IDX1:%.+]] = getelementptr inbounds float, float* [[P6]], i64 0
+// CHECK: [[P7:%.+]] = load float*, float** [[REF_ADDR]],
+// CHECK-NEXT: [[REF:%.+]] = load float*, float** [[REF_ADDR]],
+// CHECK-NEXT: [[ARR_IDX2:%.+]] = getelementptr inbounds [4 x float], [4 x float]* [[ARR_ADDR]], i64 0, i64 0
+// CHECK: [[P10:%.+]] = mul nuw i64 {{.+}}, 4
+// CHECK-NEXT: [[ARR_IDX5:%.+]] = getelementptr inbounds float, float* [[VLA_ADDR]], i64 0
+// CHECK-NEXT: [[P11:%.+]] = bitcast [6 x i64]* [[SIZES]] to i8*
+// CHECK-NEXT: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 8 [[P11]], i8* align 8 bitcast ([6 x i64]* [[SIZES1]] to i8*), i64 48, i1 false)
+// CHECK: [[BPTR0:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BPTRS]], i32 0, i32 0
 // CHECK: [[BPTR0_A_ADDR:%.+]] = bitcast i8** [[BPTR0]] to float**
 // CHECK: store float* [[A_ADDR]], float** [[BPTR0_A_ADDR]],
-// CHECK: [[PTR0:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS]], i32 0, i32 0
+// CHECK: [[PTR0:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[PTRS]], i32 0, i32 0
 // CHECK: [[PTR0_A_ADDR:%.+]] = bitcast i8** [[PTR0]] to float**
 // CHECK: store float* [[A_ADDR]], float** [[PTR0_A_ADDR]],
-// CHECK: [[BPTR1:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BPTRS]], i32 0, i32 1
+// CHECK: [[BPTR1:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BPTRS]], i32 0, i32 1
 // CHECK: [[BPTR1_PTR_ADDR:%.+]] = bitcast i8** [[BPTR1]] to float**
 // CHECK: store float* [[PTR]], float** [[BPTR1_PTR_ADDR]],
-// CHECK: [[PTR1:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS]], i32 0, i32 1
+// CHECK: [[PTR1:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[PTRS]], i32 0, i32 1
 // CHECK: [[PTR1_PTR_ADDR:%.+]] = bitcast i8** [[PTR1]] to float**
-// CHECK: store float* [[PTR]], float** [[PTR1_PTR_ADDR]],
-// CHECK: [[BPTR2:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BPTRS]], i32 0, i32 2
+// CHECK: store float* [[ARR_IDX]], float** [[PTR1_PTR_ADDR]],
+// CHECK: [[BPTR2:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BPTRS]], i32 0, i32 2
 // CHECK: [[BPTR2_REF_ADDR:%.+]] = bitcast i8** [[BPTR2]] to float**
-// CHECK: store float* [[REF]], float** [[BPTR2_REF_ADDR]],
-// CHECK: [[PTR2:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS]], i32 0, i32 2
+// CHECK: store float* [[P5]], float** [[BPTR2_REF_ADDR]],
+// CHECK: [[PTR2:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[PTRS]], i32 0, i32 2
 // CHECK: [[PTR2_REF_ADDR:%.+]] = bitcast i8** [[PTR2]] to float**
-// CHECK: store float* [[REF]], float** [[PTR2_REF_ADDR]],
-// CHECK: [[BPTR3:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BPTRS]], i32 0, i32 3
+// CHECK: store float* [[ARR_IDX1]], float** [[PTR2_REF_ADDR]],
+// CHECK: [[BPTR3:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BPTRS]], i32 0, i32 3
 // CHECK: [[BPTR3_ARR_ADDR:%.+]] = bitcast i8** [[BPTR3]] to float**
-// CHECK: store float* [[ARR]], float** [[BPTR3_ARR_ADDR]],
-// CHECK: [[PTR3:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS]], i32 0, i32 3
+// CHECK: store float* [[P7]], float** [[BPTR3_ARR_ADDR]],
+// CHECK: [[PTR3:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[PTRS]], i32 0, i32 3
 // CHECK: [[PTR3_ARR_ADDR:%.+]] = bitcast i8** [[PTR3]] to float**
-// CHECK: store float* [[ARR]], float** [[PTR3_ARR_ADDR]],
-// CHECK: [[BPTR4:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BPTRS]], i32 0, i32 4
-// CHECK: [[BPTR4_VLA_ADDR:%.+]] = bitcast i8** [[BPTR4]] to float**
-// CHECK: store float* [[VLA_ADDR]], float** [[BPTR4_VLA_ADDR]],
-// CHECK: [[PTR4:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS]], i32 0, i32 4
-// CHECK: [[PTR4_VLA_ADDR:%.+]] = bitcast i8** [[PTR4]] to float**
-// CHECK: store float* [[VLA_ADDR]], float** [[PTR4_VLA_ADDR]],
-// CHECK: [[BPTR:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BPTRS]], i32 0, i32 0
-// CHECK: [[PTR:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS]], i32 0, i32 0
-// CHECK: call void @__tgt_target_data_begin_mapper(%struct.ident_t* @{{.+}}, i64 -1, i32 5, i8** [[BPTR]], i8** [[PTR]], i64* getelementptr inbounds ([5 x i64], [5 x i64]* [[SIZES1]], i32 0, i32 0), i64* getelementptr inbounds ([5 x i64], [5 x i64]* [[MAPTYPES1]], i32 0, i32 0), i8** null, i8** null)
+// CHECK: store float* [[REF]], float** [[PTR3_ARR_ADDR]],
+// CHECK: [[BPTR4:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BPTRS]], i32 0, i32 4
+// CHECK: [[BPTR4_ARR_ADDR:%.+]] = bitcast i8** [[BPTR4]] to [4 x float]**
+// CHECK: store [4 x float]* [[ARR_ADDR]], [4 x float]** [[BPTR4_ARR_ADDR]], align 
+// CHECK: [[PTR4:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[PTRS]], i32 0, i32 4
+// CHECK: [[PTR4_ARR_ADDR:%.+]] = bitcast i8** [[PTR4]] to float**
+// CHECK: store float* [[ARR_IDX2]], float** [[PTR4_ARR_ADDR]], align 8
+// CHECK: [[SIZE_PTR:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[SIZES]], i32 0, i32 4
+// CHECK: store i64 [[P10:%.+]], i64* [[SIZE_PTR]], align 8
+// CHECK: [[MAP_PTR:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[MAP_PTRS]], i64 0, i64 4
+// CHECK: store i8* null, i8** [[MAP_PTR]], align 8
+// CHECK: [[BPTR5:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BPTRS]], i32 0, i32 5
+// CHECK: [[BPTR5_VLA_ADDR:%.+]] = bitcast i8** [[BPTR5]] to float**
+// CHECK: store float* [[VLA_ADDR]], float** [[BPTR5_VLA_ADDR]],
+// CHECK: [[PTR5:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[PTRS]], i32 0, i32 5
+// CHECK: [[PTR5_ARR_ADDR:%.+]] = bitcast i8** [[PTR5]] to float**
+// CHECK: store float* [[ARR_IDX5]], float** [[PTR5_ARR_ADDR]],
+
+// CHECK: [[BPTR:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BPTRS]], i32 0, i32 0
+// CHECK: [[PTR:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[PTRS]], i32 0, i32 0
+// CHECK: [[SIZE:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[SIZES]], i32 0, i32 0
+// CHECK: call void @__tgt_target_data_begin_mapper(%struct.ident_t* @{{.+}}, i64 -1, i32 6, i8** [[BPTR]], i8** [[PTR]], i64* [[SIZE]], i64* getelementptr inbounds ([6 x i64], [6 x i64]* [[MAPTYPES1]], i32 0, i32 0), i8** null, i8** null)
 // CHECK: [[A_REF:%.+]] = load float*, float** [[BPTR0_A_ADDR]],
-// CHECK: [[REF_REF:%.+]] = load float*, float** [[BPTR2_REF_ADDR]],
+// CHECK: [[REF_REF:%.+]] = load float*, float** [[BPTR3_ARR_ADDR]],
 // CHECK: store float* [[REF_REF]], float** [[TMP_REF_ADDR:%.+]],
-// CHECK: [[BPTR3_ARR_ADDR_CAST:%.+]] = bitcast float** [[BPTR3_ARR_ADDR]] to [4 x float]**
-// CHECK: [[ARR_REF:%.+]] = load [4 x float]*, [4 x float]** [[BPTR3_ARR_ADDR_CAST]],
-// CHECK: [[VLA_REF:%.+]] = load float*, float** [[BPTR4_VLA_ADDR]],
+// CHECK: [[ARR_REF:%.+]] = load [4 x float]*, [4 x float]** [[BPTR4_ARR_ADDR]],
+// CHECK: [[VLA_REF:%.+]] = load float*, float** [[BPTR5_VLA_ADDR]],
 // CHECK: [[A:%.+]] = load float, float* [[A_REF]],
 // CHECK: [[INC:%.+]] = fadd float [[A]], 1.000000e+00
 // CHECK: store float [[INC]], float* [[A_REF]],
@@ -112,84 +135,92 @@
 // CHECK: [[VLA0:%.+]] = load float, float* [[VLA0_ADDR]],
 // CHECK: [[INC:%.+]] = fadd float [[VLA0]], 1.000000e+00
 // CHECK: store float [[INC]], float* [[VLA0_ADDR]],
-// CHECK: [[BPTR:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BPTRS]], i32 0, i32 0
-// CHECK: [[PTR:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS]], i32 0, i32 0
-// CHECK: call void @__tgt_target_data_end_mapper(%struct.ident_t* @{{.+}}, i64 -1, i32 5, i8** [[BPTR]], i8** [[PTR]], i64* getelementptr inbounds ([5 x i64], [5 x i64]* [[SIZES1]], i32 0, i32 0), i64* getelementptr inbounds ([5 x i64], [5 x i64]* [[MAPTYPES1]], i32 0, i32 0), i8** null, i8** null)
+// CHECK: [[BPTR:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BPTRS]], i32 0, i32 0
+// CHECK: [[PTR:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[PTRS]], i32 0, i32 0
+// CHECK: [[SIZE:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[SIZES]], i32 0, i32 0
+// CHECK: call void @__tgt_target_data_end_mapper(%struct.ident_t* @{{.+}}, i64 -1, i32 6, i8** [[BPTR]], i8** [[PTR]], i64* [[SIZE]], i64* getelementptr inbounds ([6 x i64], [6 x i64]* [[MAPTYPES1]], i32 0, i32 0), i8** null, i8** null)
 
 // CHECK: foo
-// %this.addr = alloca %struct.S*, align 8
-// CHECK: [[BPTRS:%.+]] = alloca [5 x i8*],
-// CHECK: [[PTRS:%.+]] = alloca [5 x i8*],
-// CHECK: [[SIZES:%.+]] = alloca [5 x i64],
-// %tmp = alloca i32*, align 8
-// %tmp6 = alloca i32**, align 8
-// %tmp7 = alloca i32*, align 8
-// %tmp8 = alloca i32**, align 8
-// %tmp9 = alloca [4 x i32]*, align 8
-// store %struct.S* %this, %struct.S** %this.addr, align 8
-// %this1 = load %struct.S*, %struct.S** %this.addr, align 8
+// CHECK: [[BPTRS:%.+]] = alloca [6 x i8*],
+// CHECK: [[PTRS:%.+]] = alloca [6 x i8*],
+// CHECK: [[MAP_PTRS:%.+]] = alloca [6 x i8*],
+// CHECK: [[SIZES:%.+]] = alloca [6 x i64],
 // CHECK: [[A_ADDR:%.+]] = getelementptr inbounds %struct.S, %struct.S* [[THIS:%.+]], i32 0, i32 0
-// %ptr = getelementptr inbounds %struct.S, %struct.S* %this1, i32 0, i32 1
-// %ref = getelementptr inbounds %struct.S, %struct.S* %this1, i32 0, i32 2
-// %0 = load i32*, i32** %ref, align 8
-// CHECK: [[ARR_ADDR:%.+]] = getelementptr inbounds %struct.S, %struct.S* [[THIS]], i32 0, i32 3
-// CHECK: [[A_ADDR2:%.+]] = getelementptr inbounds %struct.S, %struct.S* [[THIS]], i32 0, i32 0
 // CHECK: [[PTR_ADDR:%.+]] = getelementptr inbounds %struct.S, %struct.S* [[THIS]], i32 0, i32 1
+// CHECK: [[ARR_IDX:%.+]] = getelementptr inbounds i32, i32* %{{.+}}, i64 3
 // CHECK: [[REF_REF:%.+]] = getelementptr inbounds %struct.S, %struct.S* [[THIS]], i32 0, i32 2
 // CHECK: [[REF_PTR:%.+]] = load i32*, i32** [[REF_REF]],
-// CHECK: [[ARR_ADDR2:%.+]] = getelementptr inbounds %struct.S, %struct.S* [[THIS]], i32 0, i32 3
-// CHECK: [[ARR_END:%.+]] = getelementptr [4 x i32], [4 x i32]* [[ARR_ADDR]], i32 1
-// CHECK: [[BEGIN:%.+]] = bitcast i32* [[A_ADDR]] to i8*
-// CHECK: [[END:%.+]] = bitcast [4 x i32]* [[ARR_END]] to i8*
+// CHECK-NEXT: [[P3:%.+]] = getelementptr inbounds %struct.S, %struct.S* [[THIS]], i32 0, i32 1
+// CHECK: [[ARR_IDX5:%.+]] = getelementptr inbounds i32, i32* {{.+}}, i64 0
+// CHECK: [[ARR_ADDR:%.+]] = getelementptr inbounds %struct.S, %struct.S* [[THIS]], i32 0, i32 3
+
+// CHECK: [[ARR_IDX6:%.+]] = getelementptr inbounds [4 x i32], [4 x i32]* [[ARR_ADDR]], i64 0, i64 0
+// CHECK: [[A_ADDR2:%.+]] = getelementptr inbounds %struct.S, %struct.S* [[THIS]], i32 0, i32 0
+// CHECK: [[P4:%.+]] = mul nuw i64 [[CONV:%.+]], 4
+// CHECK: [[ARR_END:%.+]] = getelementptr i32, i32* [[ARR_IDX6]], i32 1
+// CHECK: [[BEGIN:%.+]] = bitcast i32* %a to i8*
+// CHECK: [[END:%.+]] = bitcast i32* [[ARR_END]] to i8*
 // CHECK: [[E:%.+]] = ptrtoint i8* [[END]] to i64
 // CHECK: [[B:%.+]] = ptrtoint i8* [[BEGIN]] to i64
 // CHECK: [[DIFF:%.+]] = sub i64 [[E]], [[B]]
 // CHECK: [[SZ:%.+]] = sdiv exact i64 [[DIFF]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
-// CHECK: [[BPTR0:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BPTRS]], i32 0, i32 0
+// CHECK: [[BPTR0:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BPTRS]], i32 0, i32 0
 // CHECK: [[BPTR0_S:%.+]] = bitcast i8** [[BPTR0]] to %struct.S**
 // CHECK: store %struct.S* [[THIS]], %struct.S** [[BPTR0_S]],
-// CHECK: [[PTR0:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS]], i32 0, i32 0
+// CHECK: [[PTR0:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[PTRS]], i32 0, i32 0
 // CHECK: [[PTR0_BEGIN:%.+]] = bitcast i8** [[PTR0]] to i32**
 // CHECK: store i32* [[A_ADDR]], i32** [[PTR0_BEGIN]],
-// CHECK: [[SIZE0:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[SIZES]], i32 0, i32 0
+// CHECK: [[SIZE0:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[SIZES]], i32 0, i32 0
 // CHECK: store i64 [[SZ]], i64* [[SIZE0]],
-// CHECK: [[BPTR1:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BPTRS]], i32 0, i32 1
-// CHECK: [[BPTR1_A_ADDR:%.+]] = bitcast i8** [[BPTR1]] to i32**
-// CHECK: store i32* [[A_ADDR2]], i32** [[BPTR1_A_ADDR]],
-// CHECK: [[PTR1:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS]], i32 0, i32 1
+// CHECK: [[BPTR1:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BPTRS]], i32 0, i32 1
+// CHECK: [[BPTR1_A_ADDR:%.+]] = bitcast i8** [[BPTR1]] to %struct.S**
+// CHECK: store %struct.S* [[THIS]], %struct.S** [[BPTR1_A_ADDR]]
+// CHECK: [[PTR1:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[PTRS]], i32 0, i32 1
 // CHECK: [[PTR1_A_ADDR:%.+]] = bitcast i8** [[PTR1]] to i32**
-// CHECK: store i32* [[A_ADDR2]], i32** [[PTR1_A_ADDR]],
-// CHECK: [[BPTR2:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BPTRS]], i32 0, i32 2
+// CHECK: store i32* [[A_ADDR]], i32** [[PTR1_A_ADDR]],
+// CHECK: [[BPTR2:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BPTRS]], i32 0, i32 2
 // CHECK: [[BPTR2_PTR_ADDR:%.+]] = bitcast i8** [[BPTR2]] to i32***
 // CHECK: store i32** [[PTR_ADDR]], i32*** [[BPTR2_PTR_ADDR]],
-// CHECK: [[PTR2:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS]], i32 0, i32 2
-// CHECK: [[PTR2_PTR_ADDR:%.+]] = bitcast i8** [[PTR2]] to i32***
-// CHECK: store i32** [[PTR_ADDR]], i32*** [[PTR2_PTR_ADDR]],
-// CHECK: [[BPTR3:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BPTRS]], i32 0, i32 3
-// CHECK: [[BPTR3_REF_PTR:%.+]] = bitcast i8** [[BPTR3]] to i32**
-// CHECK: store i32* [[REF_PTR]], i32** [[BPTR3_REF_PTR]],
-// CHECK: [[PTR3:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS]], i32 0, i32 3
+// CHECK: [[PTR2:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[PTRS]], i32 0, i32 2
+// CHECK: [[PTR2_PTR_ADDR:%.+]] = bitcast i8** [[PTR2]] to i32** 
+// CHECK: store i32* [[ARR_IDX]], i32** [[PTR2_PTR_ADDR]],
+// CHECK: [[BPTR3:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BPTRS]], i32 0, i32 3
+// CHECK: [[BPTR3_REF_PTR:%.+]] = bitcast i8** [[BPTR3]] to %struct.S**
+// CHECK: store %struct.S* [[THIS]], %struct.S** [[BPTR3_REF_PTR]]
+// CHECK: [[PTR3:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[PTRS]], i32 0, i32 3
 // CHECK: [[PTR3_REF_PTR:%.+]] = bitcast i8** [[PTR3]] to i32**
 // CHECK: store i32* [[REF_PTR]], i32** [[PTR3_REF_PTR]],
-// CHECK: [[BPTR4:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BPTRS]], i32 0, i32 4
-// CHECK: [[BPTR4_ARR_ADDR:%.+]] = bitcast i8** [[BPTR4]] to [4 x i32]**
-// CHECK: store [4 x i32]* [[ARR_ADDR2]], [4 x i32]** [[BPTR4_ARR_ADDR]],
-// CHECK: [[PTR4:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS]], i32 0, i32 4
-// CHECK: [[PTR4_ARR_ADDR:%.+]] = bitcast i8** [[PTR4]] to [4 x i32]**
-// CHECK: store [4 x i32]* [[ARR_ADDR2]], [4 x i32]** [[PTR4_ARR_ADDR]],
-// CHECK: [[BPTR:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BPTRS]], i32 0, i32 0
-// CHECK: [[PTR:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS]], i32 0, i32 0
-// CHECK: [[SIZE:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[SIZES]], i32 0, i32 0
-// CHECK: call void @__tgt_target_data_begin_mapper(%struct.ident_t* @{{.+}}, i64 -1, i32 5, i8** [[BPTR]], i8** [[PTR]], i64* [[SIZE]], i64* getelementptr inbounds ([5 x i64], [5 x i64]* [[MAPTYPES2]], i32 0, i32 0), i8** null, i8** null)
-// CHECK: [[A_ADDR:%.+]] = load i32*, i32** [[BPTR1_A_ADDR]],
+// CHECK: [[BPTR4:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BPTRS]], i32 0, i32 4
+// CHECK: [[BPTR4_ARR_ADDR:%.+]] = bitcast i8** [[BPTR4]] to i32***
+// CHECK: store i32** [[P3]], i32*** [[BPTR4_ARR_ADDR]],
+// CHECK: [[PTR4:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[PTRS]], i32 0, i32 4
+// CHECK: [[PTR4_ARR_ADDR:%.+]] = bitcast i8** [[PTR4]] to i32**
+// CHECK: store i32* [[ARR_IDX5]], i32** [[PTR4_ARR_ADDR]]
+
+// CHECK: [[BPTR5:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BPTRS]], i32 0, i32 5
+// CHECK: [[BPTR5_ARR_ADDR:%.+]] = bitcast i8** [[BPTR5]] to %struct.S**
+// CHECK: store %struct.S* [[THIS]], %struct.S** [[BPTR5_ARR_ADDR]], align 8
+// CHECK: [[PTR5:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[PTRS]], i32 0, i32 5
+// CHECK: [[PTR5_ARR_ADDR:%.+]] = bitcast i8** [[PTR5]] to i32**
+// CHECK: store i32* [[ARR_IDX6]], i32** [[PTR5_ARR_ADDR]], align 8
+// CHECK: [[SIZE1:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[SIZES]], i32 0, i32 5
+// CHECK: store i64 [[P4]], i64* [[SIZE1]], align 8
+// CHECK: [[BPTR:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BPTRS]], i32 0, i32 0
+// CHECK: [[PTR:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[PTRS]], i32 0, i32 0
+// CHECK: [[SIZE:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[SIZES]], i32 0, i32 0
+// CHECK: call void @__tgt_target_data_begin_mapper(%struct.ident_t* @{{.+}}, i64 -1, i32 6, i8** [[BPTR]], i8** [[PTR]], i64* [[SIZE]], i64* getelementptr inbounds ([6 x i64], [6 x i64]* [[MAPTYPES2]], i32 0, i32 0), i8** null, i8** null)
+// CHECK: [[BPTR1_A_ADDR2:%.+]] = bitcast %struct.S** [[BPTR1_A_ADDR]] to i32**
+// CHECK: [[A_ADDR:%.+]] = load i32*, i32** [[BPTR1_A_ADDR2]],
 // CHECK: store i32* [[A_ADDR]], i32** [[A_REF:%.+]],
 // CHECK: [[PTR_ADDR:%.+]] = load i32**, i32*** [[BPTR2_PTR_ADDR]],
 // CHECK: store i32** [[PTR_ADDR]], i32*** [[PTR_REF:%.+]],
-// CHECK: [[REF_PTR:%.+]] = load i32*, i32** [[BPTR3_REF_PTR]],
+// CHECK: [[BPTR3_REF_PTR2:%.+]] = bitcast %struct.S** [[BPTR3_REF_PTR]] to i32**
+// CHECK: [[REF_PTR:%.+]] = load i32*, i32** [[BPTR3_REF_PTR2]],
 // CHECK: store i32* [[REF_PTR]], i32** [[REF_REF:%.+]],
 // CHECK: [[PTR_ADDR:%.+]] = load i32**, i32*** [[BPTR2_PTR_ADDR]],
 // CHECK: store i32** [[PTR_ADDR]], i32*** [[PTR_REF2:%.+]],
-// CHECK: [[ARR_ADDR:%.+]] = load [4 x i32]*, [4 x i32]** [[BPTR4_ARR_ADDR]],
+// CHECK: [[BPTR5_ARR_ADDR2:%.+]] = bitcast %struct.S** [[BPTR5_ARR_ADDR]] to [4 x i32]**
+// CHECK: [[ARR_ADDR:%.+]] = load [4 x i32]*, [4 x i32]** [[BPTR5_ARR_ADDR2]],
 // CHECK: store [4 x i32]* [[ARR_ADDR]], [4 x i32]** [[ARR_REF:%.+]],
 // CHECK: [[A_ADDR:%.+]] = load i32*, i32** [[A_REF]],
 // CHECK: [[A:%.+]] = load i32, i32* [[A_ADDR]],
@@ -209,9 +240,9 @@
 // CHECK: [[VAL:%.+]] = load i32, i32* [[ARR0_ADDR]],
 // CHECK: [[INC:%.+]] = add nsw i32 [[VAL]], 1
 // CHECK: store i32 [[INC]], i32* [[ARR0_ADDR]],
-// CHECK: [[BPTR:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BPTRS]], i32 0, i32 0
-// CHECK: [[PTR:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS]], i32 0, i32 0
-// CHECK: [[SIZE:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[SIZES]], i32 0, i32 0
-// CHECK: call void @__tgt_target_data_end_mapper(%struct.ident_t* @{{.+}}, i64 -1, i32 5, i8** [[BPTR]], i8** [[PTR]], i64* [[SIZE]], i64* getelementptr inbounds ([5 x i64], [5 x i64]* [[MAPTYPES2]], i32 0, i32 0), i8** null, i8** null)
+// CHECK: [[BPTR:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BPTRS]], i32 0, i32 0
+// CHECK: [[PTR:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[PTRS]], i32 0, i32 0
+// CHECK: [[SIZE:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[SIZES]], i32 0, i32 0
+// CHECK: call void @__tgt_target_data_end_mapper(%struct.ident_t* @{{.+}}, i64 -1, i32 6, i8** [[BPTR]], i8** [[PTR]], i64* [[SIZE]], i64* getelementptr inbounds ([6 x i64], [6 x i64]* [[MAPTYPES2]], i32 0, i32 0), i8** null, i8** null)
 
 #endif