[OpenMP] Fix runtime problem due to wrong map size. (#74692)

Currently we are missing set up-boundary address for FinalArraySection
as highests elements in partial struct data.

Currently for:
\#pragma omp target map(D.a) map(D.b[:2])
The size is:
  %a = getelementptr inbounds %struct.DataTy, ptr %D, i32 0, i32 0
  %b = getelementptr inbounds %struct.DataTy, ptr %D, i32 0, i32 1
  %arrayidx = getelementptr inbounds [2 x float], ptr %b, i64 0, i64 0
  %2 = getelementptr float, ptr %arrayidx, i32 1
  %3 = ptrtoint ptr %2 to i64
  %4 = ptrtoint ptr %a to i64
  %5 = sub i64 %3, %4
%6 = sdiv exact i64 %5, ptrtoint (ptr getelementptr (i8, ptr null, i32
1) to i64)

Where %2 is wrong for (D.b[:2]) is pointer to first element of array
section. It should pointe to last element of array section.
  
The fix is to emit the pointer to the last element of array section and
use this pointer as the highest element in partial struct data.

After change IR:
  %a = getelementptr inbounds %struct.DataTy, ptr %D, i32 0, i32 0
  %b = getelementptr inbounds %struct.DataTy, ptr %D, i32 0, i32 1
  %arrayidx = getelementptr inbounds [2 x float], ptr %b, i64 0, i64 0
  %b1 = getelementptr inbounds %struct.DataTy, ptr %D, i32 0, i32 1
  %arrayidx2 = getelementptr inbounds [2 x float], ptr %b1, i64 0, i64 1
  %1 = getelementptr float, ptr %arrayidx2, i32 1
  %2 = ptrtoint ptr %1 to i64
  %3 = ptrtoint ptr %a to i64
  %4 = sub i64 %2, %3
%5 = sdiv exact i64 %4, ptrtoint (ptr getelementptr (i8, ptr null, i32
1) to i64)

clang/lib/CodeGen/CGOpenMPRuntime.cpp

@@ -7389,7 +7389,14 @@ class MappableExprsHandler {
           } else if (FieldIndex < PartialStruct.LowestElem.first) {
             PartialStruct.LowestElem = {FieldIndex, LowestElem};
           } else if (FieldIndex > PartialStruct.HighestElem.first) {
-            PartialStruct.HighestElem = {FieldIndex, LowestElem};
+            if (IsFinalArraySection) {
+              Address HB =
+                  CGF.EmitOMPArraySectionExpr(OASE, /*IsLowerBound=*/false)
+                      .getAddress(CGF);
+              PartialStruct.HighestElem = {FieldIndex, HB};
+            } else {
+              PartialStruct.HighestElem = {FieldIndex, LowestElem};
+            }
           }
         }
 

clang/test/OpenMP/target_data_use_device_addr_codegen.cpp

@@ -144,7 +144,13 @@ int main() {
 // CHECK: [[ARR_IDX6:%.+]] = getelementptr inbounds [4 x i32], ptr [[ARR_ADDR]], i64 0, i64 0
 // CHECK: [[A_ADDR2:%.+]] = getelementptr inbounds %struct.S, ptr [[THIS]], i32 0, i32 0
 // CHECK: [[P4:%.+]] = mul nuw i64 [[CONV:%.+]], 4
-// CHECK: [[ARR_END:%.+]] = getelementptr i32, ptr [[ARR_IDX6]], i32 1
+// CHECK: [[A_ADDR3:%.+]] = getelementptr inbounds %struct.S, ptr [[THIS]], i32 0, i32 0
+// CHECK: [[L5:%.+]] = load i32, ptr [[A_ADDR3]]
+// CHECK: [[L6:%.+]] = sext i32 [[L5]] to i64
+// CHECK: [[LB_ADD_LEN:%lb_add_len]] = add nsw i64 -1, [[L6]]
+// CHECK: [[ARR_ADDR9:%.+]] = getelementptr inbounds %struct.S, ptr [[THIS]], i32 0, i32 3
+// CHECK: [[ARR_IDX10:%arrayidx.+]] = getelementptr inbounds [4 x i32], ptr [[ARR_ADDR9]], i64 0, i64 %lb_add_len
+// CHECK: [[ARR_END:%.+]] = getelementptr i32, ptr [[ARR_IDX10]], i32 1
 // CHECK: [[E:%.+]] = ptrtoint ptr [[ARR_END]] to i64
 // CHECK: [[B:%.+]] = ptrtoint ptr [[A_ADDR]] to i64
 // CHECK: [[DIFF:%.+]] = sub i64 [[E]], [[B]]

openmp/libomptarget/test/offloading/target_map_for_member_data.cpp

@@ -0,0 +1,23 @@
+// clang-format off
+// RUN: %libomptarget-compilexx-generic && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-generic 2>&1 | %fcheck-generic
+// clang-format on
+
+struct DataTy {
+  float a;
+  float b[3];
+};
+
+int main(int argc, char **argv) {
+  DataTy D;
+#pragma omp target map(D.a) map(D.b[ : 2])
+  {
+    D.a = 0;
+    D.b[0] = 1;
+  }
+  return 0;
+}
+// clang-format off
+// CHECK: omptarget --> Entry  0: Base=[[DAT_HST_PTR_BASE:0x.*]], Begin=[[DAT_HST_PTR_BASE]], Size=12
+// CHECK: omptarget --> Entry  1: Base=[[DAT_HST_PTR_BASE]], Begin=[[DAT_HST_PTR_BASE]], Size=4,
+// CHECK: omptarget --> Entry  2: Base=[[DAT_HST_PTR_BASE]], Begin=[[DATUM_HST_PTR_BASE:0x.*]], Size=8,
+// clang-format on