Update copy-gaps tests.

abhinavgaba · abhinavgaba · commit 502dbb4d8088 · 2025-08-14T16:03:55.000-07:00
diff --git a/clang/test/OpenMP/copy-gaps-1.cpp b/clang/test/OpenMP/copy-gaps-1.cpp
@@ -18,6 +18,17 @@ struct T : public S {
 int main() {
   T v;
 
+// &v,       &v,                 sizeof(v),                         ALLOC | PARAM
+// &v,       &v.z + sizeof(int), sizeof(v.p1 to v.p2),              TO | FROM | MEMBER_OF_1
+// &v,       &v.c + sizeof(int), sizeof(T) - (&v.c + sizeof(int))), TO | FROM | MEMBER_OF_1
+// &v,       &v.x,               sizeof(v.x),                       TO | FROM | MEMBER_OF_1
+// &v,       &v.y,               sizeof(v.y),                       TO | FROM | MEMBER_OF_1
+// &v,       &v.z,               sizeof(v.z),                       TO | FROM | MEMBER_OF_1
+// &v,       &v.a,               sizeof(v.a),                       TO | FROM | MEMBER_OF_1
+// &v,       &v.b,               sizeof(v.b),                       TO | FROM | MEMBER_OF_1
+// &v,       &v.c,               sizeof(v.c),                       TO | FROM | MEMBER_OF_1
+// &v.p1[0], &v.p1[0],           8 * sizeof(int),                   TO | FROM
+// &v.p1,    &v.p1[0],           sizeof(v.p1),                      ATTACH
 #pragma omp target map(tofrom: v, v.x, v.y, v.z, v.p1[:8], v.a, v.b, v.c)
   {
     v.x++;
@@ -32,21 +43,21 @@ int main() {
   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: [[CSTSZ:@.+]] = private {{.*}}constant [11 x i64] [i64 0, i64 0, i64 0, i64 4, i64 4, i64 4, i64 4, i64 4, i64 4, i64 32, i64 8]
+// CHECK: [[CSTTY:@.+]] = private {{.*}}constant [11 x i64] [i64 [[#0x20]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x3]], i64 [[#0x4000]]]
 
 // 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-DAG: [[SZBASE]] = getelementptr inbounds [11 x i64], ptr [[SIZES:%[^,]*]], i32 0, i32 0
 
-// Check for filling of four non-constant size elements here: the whole struct
+// Check for filling of three 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: [[STR:%.+]] = getelementptr inbounds [11 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: [[P1P2:%.+]] = getelementptr inbounds [11 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: [[PAD:%.+]] = getelementptr inbounds [11 x i64], ptr [[SIZES]], i32 0, i32 2
 // CHECK-DAG: store i64 %{{.+}}, ptr [[PAD]], align 8
diff --git a/clang/test/OpenMP/copy-gaps-6.cpp b/clang/test/OpenMP/copy-gaps-6.cpp
@@ -11,13 +11,22 @@ struct S {
 int main() {
   S v;
 
+// &v, &v,                 sizeof(v),            ALLOC | PARAM
+// &v, &v.x + sizeof(int), sizeof(v.arr to v.y), TO | FROM | MEMBER_OF_1
+// &v, &v.x,               sizeof(v.x),          TO | FROM | MEMBER_OF_1
+// &v, &v.z,               sizeof(v.z),          TO | FROM | MEMBER_OF_1
 #pragma omp target map(tofrom: v, v.x, v.z)
   {
     v.x++;
     v.y += 2;
     v.z += 3;
   }
 
+// &v,        &v,                 sizeof(v),            ALLOC | PARAM
+// &v,        &v.x + sizeof(int), sizeof(v.arr to v.z), TO | FROM | MEMBER_OF_1
+// &v,        &v.x,               sizeof(v.x),          TO | FROM | MEMBER_OF_1
+// &v.arr[0], &v.arr[0],          4 * sizeof(int),      TO | FROM
+// &v.arr,    &v.arr[0],          sizeof(void*),        ATTACH
 #pragma omp target map(tofrom: v, v.x, v.arr[:1])
   {
     v.x++;
@@ -26,6 +35,9 @@ int main() {
     v.z += 4;
   }
 
+// &v,        &v,        sizeof(v),       TO | FROM | PARAM
+// &v.arr[0], &v.arr[0], 4 * sizeof(int), TO | FROM
+// &v.arr,    &v.arr[0], sizeof(void*),   ATTACH
 #pragma omp target map(tofrom: v, v.arr[:1])
   {
     v.x++;
@@ -40,11 +52,11 @@ int main() {
 // 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: [[CSTSZ1:@.+]] = private {{.*}}constant [5 x i64] [i64 0, i64 0, i64 4, i64 4, i64 8]
+// CHECK: [[CSTTY1:@.+]] = private {{.*}}constant [5 x i64] [i64 [[#0x20]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x3]], i64 [[#0x4000]]]
 
-// 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: [[CSTSZ2:@.+]] = private {{.*}}constant [3 x i64] [i64 24, i64 4, i64 8]
+// CHECK: [[CSTTY2:@.+]] = private {{.*}}constant [3 x i64] [i64 [[#0x23]], i64 [[#0x3]], i64 [[#0x4000]]]
 
 // 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
@@ -64,24 +76,19 @@ int main() {
 // 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
+// CHECK-DAG: [[SZBASE]] = getelementptr inbounds [5 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: [[STR:%.+]] = getelementptr inbounds [5 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: [[ARRYZ:%.+]] = getelementptr inbounds [5 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
+// No overlap, so no non-constant size element here.
+// CHECK-NOT: store ptr [[CSTSZ2]], ptr [[KSIZE]], align 8
diff --git a/clang/test/OpenMP/target_map_pointer_defalut_mapper_codegen.cpp b/clang/test/OpenMP/target_map_pointer_defalut_mapper_codegen.cpp
@@ -50,9 +50,9 @@ void foo() {
 // CHECK: @.offload_maptypes = private unnamed_addr constant [6 x i64] [i64 [[#0x220]], i64 0, i64 [[#0x2000000000003]], i64 [[#0x2000000000003]], i64 [[#0x2000000000203]], i64 [[#0x4000]]]
 
 // &sp[0], &sp[0],   sizeof(sp[0]), ALLOC | PARAM
-// &sp[0], &sp[0].e, sizeof(e),     MEMBER_OF_2 | TO | FROM
-// &sp[0], &sp[0].h, sizeof(h),     MEMBER_OF_2 | TO | FROM
-// &sp[0], &sp[0].f, sizeof(f),     MEMBER_OF_2 | TO | FROM | IMPLICIT, mapper_of_c
+// &sp[0], &sp[0].e, sizeof(e),     MEMBER_OF_1 | TO | FROM
+// &sp[0], &sp[0].h, sizeof(h),     MEMBER_OF_1 | TO | FROM
+// &sp[0], &sp[0].f, sizeof(f),     MEMBER_OF_1 | TO | FROM | IMPLICIT, mapper_of_c
 // &sp,    &sp[0],   sizeof(void*), ATTACH
 // CHECK: @.offload_sizes.1 = private unnamed_addr constant [5 x i64] [i64 0, i64 0, i64 0, i64 4, i64 8]
 // CHECK: @.offload_maptypes.2 = private unnamed_addr constant [5 x i64] [i64 [[#0x20]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000203]], i64 [[#0x4000]]]
