7 files changed, 339 insertions, 25 deletions
diff --git a/clang/test/OpenMP/copy-gaps-1.cpp b/clang/test/OpenMP/copy-gaps-1.cpp
new file mode 100644
index 0000000..3d4fae3
--- /dev/null
+++ b/clang/test/OpenMP/copy-gaps-1.cpp
@@ -0,0 +1,52 @@
+// RUN: %clang_cc1 -verify -triple x86_64-pc-linux-gnu -fopenmp-targets=amdgcn-amd-amdhsa -fopenmp -emit-llvm %s -o - | FileCheck %s
+// expected-no-diagnostics
+
+struct S {
+  int x;
+  int y;
+  int z;
+  int *p1;
+  int *p2;
+};
+
+struct T : public S {
+  int a;
+  int b;
+  int c;
+};
+
+int main() {
+  T v;
+
+#pragma omp target map(tofrom: v, v.x, v.y, v.z, v.p1[:8], v.a, v.b, v.c)
+  {
+    v.x++;
+    v.y += 2;
+    v.z += 3;
+    v.p1[0] += 4;
+    v.a += 7;
+    v.b += 5;
+    v.c += 6;
+  }
+
+  return 0;
+}
+
+// CHECK: [[CSTSZ:@.+]] = private {{.*}}constant [10 x i64] [i64 0, i64 0, i64 0, i64 4, i64 4, i64 4, i64 32, i64 4, i64 4, i64 4]
+// CHECK: [[CSTTY:@.+]] = private {{.*}}constant [10 x i64] [i64 [[#0x20]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000013]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]]]
+
+// CHECK-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
+// CHECK-DAG: [[KSIZE:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
+// CHECK-DAG: store ptr [[SZBASE:%.+]], ptr [[KSIZE]], align 8
+// CHECK-DAG: [[SZBASE]] = getelementptr inbounds [10 x i64], ptr [[SIZES:%[^,]*]], i32 0, i32 0
+
+// Check for filling of four non-constant size elements here: the whole struct
+// size, the (padded) region covering p1 & p2, and the padding at the end of
+// struct T.
+
+// CHECK-DAG: [[STR:%.+]] = getelementptr inbounds [10 x i64], ptr [[SIZES]], i32 0, i32 0
+// CHECK-DAG: store i64 %{{.+}}, ptr [[STR]], align 8
+// CHECK-DAG: [[P1P2:%.+]] = getelementptr inbounds [10 x i64], ptr [[SIZES]], i32 0, i32 1
+// CHECK-DAG: store i64 %{{.+}}, ptr [[P1P2]], align 8
+// CHECK-DAG: [[PAD:%.+]] = getelementptr inbounds [10 x i64], ptr [[SIZES]], i32 0, i32 2
+// CHECK-DAG: store i64 %{{.+}}, ptr [[PAD]], align 8
diff --git a/clang/test/OpenMP/copy-gaps-2.cpp b/clang/test/OpenMP/copy-gaps-2.cpp
new file mode 100644
index 0000000..5bf603a
--- /dev/null
+++ b/clang/test/OpenMP/copy-gaps-2.cpp
@@ -0,0 +1,52 @@
+// RUN: %clang_cc1 -verify -triple x86_64-pc-linux-gnu -fopenmp-targets=amdgcn-amd-amdhsa -fopenmp -emit-llvm %s -o - | FileCheck %s
+// expected-no-diagnostics
+
+struct S {
+  int x;
+  int y;
+  int z;
+};
+
+struct M : public S {
+  int mid;
+};
+
+struct T : public M {
+  int a;
+  int b;
+  int c;
+};
+
+int main() {
+  T v;
+
+#pragma omp target map(tofrom: v, v.y, v.z, v.a)
+  {
+    v.y++;
+    v.z += 2;
+    v.a += 3;
+    v.mid += 5;
+  }
+
+  return 0;
+}
+
+// CHECK: [[CSTSZ:@.+]] = private {{.*}}constant [7 x i64] [i64 0, i64 0, i64 0, i64 0, i64 4, i64 4, i64 4]
+// CHECK: [[CSTTY:@.+]] = private {{.*}}constant [7 x i64] [i64 [[#0x20]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]]]
+
+// CHECK-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
+// CHECK-DAG: [[KSIZE:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
+// CHECK-DAG: store ptr [[SZBASE:%.+]], ptr [[KSIZE]], align 8
+// CHECK-DAG: [[SZBASE]] = getelementptr inbounds [7 x i64], ptr [[SIZES:%[^,]*]], i32 0, i32 0
+
+// Fill four non-constant size elements here: the whole struct size, the region
+// covering v.x, the region covering v.mid and the region covering v.b and v.c.
+
+// CHECK-DAG: [[STR:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 0
+// CHECK-DAG: store i64 %{{.+}}, ptr [[STR]], align 8
+// CHECK-DAG: [[X:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 1
+// CHECK-DAG: store i64 %{{.+}}, ptr [[X]], align 8
+// CHECK-DAG: [[MID:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 2
+// CHECK-DAG: store i64 %{{.+}}, ptr [[MID]], align 8
+// CHECK-DAG: [[BC:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 3
+// CHECK-DAG: store i64 %{{.+}}, ptr [[BC]], align 8
diff --git a/clang/test/OpenMP/copy-gaps-3.cpp b/clang/test/OpenMP/copy-gaps-3.cpp
new file mode 100644
index 0000000..5febb18
--- /dev/null
+++ b/clang/test/OpenMP/copy-gaps-3.cpp
@@ -0,0 +1,46 @@
+// RUN: %clang_cc1 -verify -triple x86_64-pc-linux-gnu -fopenmp-targets=amdgcn-amd-amdhsa -fopenmp -emit-llvm %s -o - | FileCheck %s
+// expected-no-diagnostics
+
+struct S {
+  int x;
+  int y;
+  int z;
+};
+
+struct T : public S {
+  int a;
+  int b;
+  int c;
+};
+
+int main() {
+  T v;
+
+  // This one should have no gap between v.z & v.a.
+#pragma omp target map(tofrom: v, v.y, v.z, v.a)
+  {
+    v.y++;
+    v.z += 2;
+    v.a += 3;
+  }
+
+  return 0;
+}
+
+// CHECK: [[CSTSZ:@.+]] = private {{.*}}constant [6 x i64] [i64 0, i64 0, i64 0, i64 4, i64 4, i64 4]
+// CHECK: [[CSTTY:@.+]] = private {{.*}}constant [6 x i64] [i64 [[#0x20]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]]]
+
+// CHECK-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
+// CHECK-DAG: [[KSIZE:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
+// CHECK-DAG: store ptr [[SZBASE:%.+]], ptr [[KSIZE]], align 8
+// CHECK-DAG: [[SZBASE]] = getelementptr inbounds [6 x i64], ptr [[SIZES:%[^,]*]], i32 0, i32 0
+
+// Fill three non-constant size elements here: the whole struct size, the region
+// covering v.x, and the region covering v.b and v.c.
+
+// CHECK-DAG: [[STR:%.+]] = getelementptr inbounds [6 x i64], ptr [[SIZES]], i32 0, i32 0
+// CHECK-DAG: store i64 %{{.+}}, ptr [[STR]], align 8
+// CHECK-DAG: [[X:%.+]] = getelementptr inbounds [6 x i64], ptr [[SIZES]], i32 0, i32 1
+// CHECK-DAG: store i64 %{{.+}}, ptr [[X]], align 8
+// CHECK-DAG: [[BC:%.+]] = getelementptr inbounds [6 x i64], ptr [[SIZES]], i32 0, i32 2
+// CHECK-DAG: store i64 %{{.+}}, ptr [[BC]], align 8
diff --git a/clang/test/OpenMP/copy-gaps-4.cpp b/clang/test/OpenMP/copy-gaps-4.cpp
new file mode 100644
index 0000000..7060fe3
--- /dev/null
+++ b/clang/test/OpenMP/copy-gaps-4.cpp
@@ -0,0 +1,48 @@
+// RUN: %clang_cc1 -verify -triple x86_64-pc-linux-gnu -fopenmp-targets=amdgcn-amd-amdhsa -fopenmp -emit-llvm %s -o - | FileCheck %s
+// expected-no-diagnostics
+
+struct S {
+  int x;
+  int y;
+  char z; // Hidden padding after here...
+};
+
+struct T : public S {
+  int a;
+  int b;
+  int c;
+};
+
+int main() {
+  T v;
+
+#pragma omp target map(tofrom: v, v.y, v.z, v.a)
+  {
+    v.y++;
+    v.z += 2;
+    v.a += 3;
+  }
+
+  return 0;
+}
+
+// CHECK: [[CSTSZ:@.+]] = private {{.*}}constant [7 x i64] [i64 0, i64 0, i64 0, i64 0, i64 4, i64 1, i64 4]
+// CHECK: [[CSTTY:@.+]] = private {{.*}}constant [7 x i64] [i64 [[#0x20]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]]]
+
+// CHECK-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
+// CHECK-DAG: [[KSIZE:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
+// CHECK-DAG: store ptr [[SZBASE:%.+]], ptr [[KSIZE]], align 8
+// CHECK-DAG: [[SZBASE]] = getelementptr inbounds [7 x i64], ptr [[SIZES:%[^,]*]], i32 0, i32 0
+
+// Fill four non-constant size elements here: the whole struct size, the region
+// covering v.x, the region covering padding after v.z and the region covering
+// v.b and v.c.
+
+// CHECK-DAG: [[STR:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 0
+// CHECK-DAG: store i64 %{{.+}}, ptr [[STR]], align 8
+// CHECK-DAG: [[X:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 1
+// CHECK-DAG: store i64 %{{.+}}, ptr [[X]], align 8
+// CHECK-DAG: [[PAD:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 2
+// CHECK-DAG: store i64 %{{.+}}, ptr [[PAD]], align 8
+// CHECK-DAG: [[BC:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 3
+// CHECK-DAG: store i64 %{{.+}}, ptr [[BC]], align 8
diff --git a/clang/test/OpenMP/copy-gaps-5.cpp b/clang/test/OpenMP/copy-gaps-5.cpp
new file mode 100644
index 0000000..fae675d
--- /dev/null
+++ b/clang/test/OpenMP/copy-gaps-5.cpp
@@ -0,0 +1,50 @@
+// RUN: %clang_cc1 -verify -triple x86_64-pc-linux-gnu -fopenmp-targets=amdgcn-amd-amdhsa -fopenmp -emit-llvm %s -o - | FileCheck %s
+// expected-no-diagnostics
+
+template<typename C>
+struct S {
+  C x;
+  C y;
+  char z; // Hidden padding after here...
+};
+
+template<typename C>
+struct T : public S<C> {
+  C a;
+  C b;
+  C c;
+};
+
+int main() {
+  T<int> v;
+
+#pragma omp target map(tofrom: v, v.y, v.z, v.a)
+  {
+    v.y++;
+    v.z += 2;
+    v.a += 3;
+  }
+
+  return 0;
+}
+
+// CHECK: [[CSTSZ:@.+]] = private {{.*}}constant [7 x i64] [i64 0, i64 0, i64 0, i64 0, i64 4, i64 1, i64 4]
+// CHECK: [[CSTTY:@.+]] = private {{.*}}constant [7 x i64] [i64 [[#0x20]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]]]
+
+// CHECK-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
+// CHECK-DAG: [[KSIZE:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
+// CHECK-DAG: store ptr [[SZBASE:%.+]], ptr [[KSIZE]], align 8
+// CHECK-DAG: [[SZBASE]] = getelementptr inbounds [7 x i64], ptr [[SIZES:%[^,]*]], i32 0, i32 0
+
+// Fill four non-constant size elements here: the whole struct size, the region
+// covering v.x, the region covering padding after v.z and the region covering
+// v.b and v.c.
+
+// CHECK-DAG: [[STR:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 0
+// CHECK-DAG: store i64 %{{.+}}, ptr [[STR]], align 8
+// CHECK-DAG: [[X:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 1
+// CHECK-DAG: store i64 %{{.+}}, ptr [[X]], align 8
+// CHECK-DAG: [[PAD:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 2
+// CHECK-DAG: store i64 %{{.+}}, ptr [[PAD]], align 8
+// CHECK-DAG: [[BC:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 3
+// CHECK-DAG: store i64 %{{.+}}, ptr [[BC]], align 8
diff --git a/clang/test/OpenMP/copy-gaps-6.cpp b/clang/test/OpenMP/copy-gaps-6.cpp
new file mode 100644
index 0000000..9c62fde
--- /dev/null
+++ b/clang/test/OpenMP/copy-gaps-6.cpp
@@ -0,0 +1,87 @@
+// RUN: %clang_cc1 -verify -triple x86_64-pc-linux-gnu -fopenmp-targets=amdgcn-amd-amdhsa -fopenmp -emit-llvm %s -o - | FileCheck %s
+// expected-no-diagnostics
+
+struct S {
+  int x;
+  int *arr;
+  int y;
+  int z;
+};
+
+int main() {
+  S v;
+
+#pragma omp target map(tofrom: v, v.x, v.z)
+  {
+    v.x++;
+    v.y += 2;
+    v.z += 3;
+  }
+
+#pragma omp target map(tofrom: v, v.x, v.arr[:1])
+  {
+    v.x++;
+    v.y += 2;
+    v.arr[0] += 2;
+    v.z += 4;
+  }
+
+#pragma omp target map(tofrom: v, v.arr[:1])
+  {
+    v.x++;
+    v.y += 2;
+    v.arr[0] += 2;
+    v.z += 4;
+  }
+
+  return 0;
+}
+
+// CHECK: [[CSTSZ0:@.+]] = private {{.*}}constant [4 x i64] [i64 0, i64 0, i64 4, i64 4]
+// CHECK: [[CSTTY0:@.+]] = private {{.*}}constant [4 x i64] [i64 [[#0x20]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]]]
+
+// CHECK: [[CSTSZ1:@.+]] = private {{.*}}constant [4 x i64] [i64 0, i64 0, i64 4, i64 4]
+// CHECK: [[CSTTY1:@.+]] = private {{.*}}constant [4 x i64] [i64 [[#0x20]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000013]]]
+
+// CHECK: [[CSTSZ2:@.+]] = private {{.*}}constant [3 x i64] [i64 0, i64 24, i64 4]
+// CHECK: [[CSTTY2:@.+]] = private {{.*}}constant [3 x i64] [i64 [[#0x20]], i64 [[#0x1000000000003]], i64 [[#0x1000000000013]]]
+
+// CHECK-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
+// CHECK-DAG: [[KSIZE:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
+// CHECK-DAG: store ptr [[SZBASE:%.+]], ptr [[KSIZE]], align 8
+// CHECK-DAG: [[SZBASE]] = getelementptr inbounds [4 x i64], ptr [[SIZES:%[^,]*]], i32 0, i32 0
+
+// Fill two non-constant size elements here: the whole struct size, and the
+// region covering v.arr and v.y.
+
+// CHECK-DAG: [[STR:%.+]] = getelementptr inbounds [4 x i64], ptr [[SIZES]], i32 0, i32 0
+// CHECK-DAG: store i64 %{{.+}}, ptr [[STR]], align 8
+// CHECK-DAG: [[ARRY:%.+]] = getelementptr inbounds [4 x i64], ptr [[SIZES]], i32 0, i32 1
+// CHECK-DAG: store i64 %{{.+}}, ptr [[ARRY]], align 8
+
+// CHECK: call void
+
+// CHECK-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
+// CHECK-DAG: [[KSIZE:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
+// CHECK-DAG: store ptr [[SZBASE:%.+]], ptr [[KSIZE]], align 8
+// CHECK-DAG: [[SZBASE]] = getelementptr inbounds [4 x i64], ptr [[SIZES:%[^,]*]], i32 0, i32 0
+
+// Fill two non-constant size elements here: the whole struct size, and the
+// region covering v.arr, v.y and v.z.
+
+// CHECK-DAG: [[STR:%.+]] = getelementptr inbounds [4 x i64], ptr [[SIZES]], i32 0, i32 0
+// CHECK-DAG: store i64 %{{.+}}, ptr [[STR]], align 8
+// CHECK-DAG: [[ARRYZ:%.+]] = getelementptr inbounds [4 x i64], ptr [[SIZES]], i32 0, i32 1
+// CHECK-DAG: store i64 %{{.+}}, ptr [[ARRYZ]], align 8
+
+// CHECK: call void
+
+// CHECK-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
+// CHECK-DAG: [[KSIZE:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
+// CHECK-DAG: store ptr [[SZBASE:%.+]], ptr [[KSIZE]], align 8
+// CHECK-DAG: [[SZBASE]] = getelementptr inbounds [3 x i64], ptr [[SIZES:%[^,]*]], i32 0, i32 0
+
+// Fill one non-constant size element here: the whole struct size.
+
+// CHECK-DAG: [[STR:%.+]] = getelementptr inbounds [3 x i64], ptr [[SIZES]], i32 0, i32 0
+// CHECK-DAG: store i64 %{{.+}}, ptr [[STR]], align 8
diff --git a/clang/test/OpenMP/target_map_codegen_35.cpp b/clang/test/OpenMP/target_map_codegen_35.cpp
index afa0965..c4fc49c 100644
--- a/clang/test/OpenMP/target_map_codegen_35.cpp
+++ b/clang/test/OpenMP/target_map_codegen_35.cpp
@@ -27,11 +27,11 @@ public:
   void foo();
 };
 
-// CK35-DAG: [[SIZE_TO:@.+]] = private {{.*}}constant [4 x i64] [i64 0, i64 0, i64 0, i64 8]
+// CK35-DAG: [[SIZE_TO:@.+]] = private {{.*}}constant [3 x i64] [i64 0, i64 0, i64 8]
 // TARGET_PARAM = 0x20
 // MEMBER_OF_1 | TO = 0x1000000000001
 // MEMBER_OF_1 | PTR_AND_OBJ | TO = 0x1000000000011
-// CK35-DAG: [[MTYPE_TO:@.+]] = {{.+}}constant [4 x i64] [i64 [[#0x20]], i64 [[#0x1000000000001]], i64 [[#0x1000000000001]], i64 [[#0x1000000000011]]]
+// CK35-DAG: [[MTYPE_TO:@.+]] = {{.+}}constant [3 x i64] [i64 [[#0x20]], i64 [[#0x1000000000001]], i64 [[#0x1000000000011]]]
 // CK35-DAG: [[SIZE_FROM:@.+]] = private {{.*}}constant [2 x i64] [i64 0, i64 8]
 // TARGET_PARAM = 0x20
 // MEMBER_OF_1 | PTR_AND_OBJ | FROM = 0x1000000000012
@@ -86,35 +86,14 @@ void ref_map() {
   // CK35-DAG: [[B_BEGIN_INTPTR]] = ptrtoint ptr [[B_BEGIN_VOID:%.+]] to i64
   // CK35-DAG: [[B_ADDR:%.+]] = getelementptr inbounds nuw %class.S, ptr [[S_ADDR]], i32 0, i32 1
 
-  // pass MEMBER_OF_1 | TO {&s, &s.b+1, ((ptr)(&s+1)-(ptr)(&s.b+1))} to copy the data of remainder of s.
+  // pass MEMBER_OF_1 | PTR_AND_OBJ | TO {&s, &s.b, 8|4} to copy the data of s.b.
 
   // CK35-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
   // CK35-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
-  // CK35-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2
 
 
   // CK35-DAG: store ptr [[S_ADDR]], ptr [[BP2]],
-  // CK35-DAG: store ptr [[B_END:%.+]], ptr [[P2]],
-  // CK35-DAG: store i64 [[REM_SIZE:%.+]], ptr [[S2]],
-
-  // CK35-DAG: [[B_END]] = getelementptr ptr, ptr [[B_ADDR]], i{{.+}} 1
-
-  // CK35-DAG: [[REM_SIZE]] = sdiv exact i64 [[SZ:%.+]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64)
-  // CK35-DAG: [[SZ]] = sub i64 [[S_END_INTPTR:%.+]], [[B_END_INTPTR:%.+]]
-  // CK35-DAG: [[B_END_INTPTR]] = ptrtoint ptr [[B_END_VOID:%.+]] to i64
-  // CK35-DAG: [[S_END_INTPTR]] = ptrtoint ptr [[S_END_VOID:%.+]] to i64
-  // CK35-DAG: [[S_END_VOID]] = getelementptr i8, ptr [[S_LAST:%.+]], i{{.+}} 1
-  // CK35-64-DAG: [[S_LAST]] = getelementptr i8, ptr [[S_VOIDPTR:%.+]], i64 15
-  // CK35-32-DAG: [[S_LAST]] = getelementptr i8, ptr [[S_VOIDPTR:%.+]], i32 7
-
-  // pass MEMBER_OF_1 | PTR_AND_OBJ | TO {&s, &s.b, 8|4} to copy the data of s.b.
-
-  // CK35-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 3
-  // CK35-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 3
-
-
-  // CK35-DAG: store ptr [[S_ADDR]], ptr [[BP3]],
-  // CK35-DAG: store ptr [[B_ADDR:%.+]], ptr [[P3]],
+  // CK35-DAG: store ptr [[B_ADDR:%.+]], ptr [[P2]],
 
   // CK35-DAG: [[B_ADDR]] = load ptr, ptr [[B_REF:%.+]], align {{[0-9]+}}, !nonnull !{{[0-9]+}}, !align !{{[0-9]+}}
   // CK35-DAG: [[B_REF]] = getelementptr inbounds nuw %class.S, ptr [[S_ADDR]], i32 0, i32 1