[openmp] [clang] [OpenMP] Fix runtime problem due wrong map size. (PR #74692)
via cfe-commits
cfe-commits at lists.llvm.org
Wed Dec 6 19:37:31 PST 2023
https://github.com/jyu2-git created https://github.com/llvm/llvm-project/pull/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 %arrayidx is wrong for (D.b[:2]) should be:
%arrayidx = getelementptr inbounds [2 x float], ptr %b, i64 0, i64 1
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)
>From 50c6009e4f4184ed7710a7ee3d8ee0983306edc1 Mon Sep 17 00:00:00 2001
From: Jennifer Yu <jennifer.yu at intel.com>
Date: Wed, 6 Dec 2023 13:53:16 -0800
Subject: [PATCH] [OpenMP] Fix runtime problem due wrong map size.
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 %arrayidx is wrong for (D.b[:2]) should be:
%arrayidx = getelementptr inbounds [2 x float], ptr %b, i64 0, i64 1
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 | 9 +++++++-
.../target_data_use_device_addr_codegen.cpp | 8 ++++++-
.../offloading/target_map_for_member_data.cpp | 23 +++++++++++++++++++
3 files changed, 38 insertions(+), 2 deletions(-)
create mode 100644 openmp/libomptarget/test/offloading/target_map_for_member_data.cpp
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 55648963df36a..7f7e6f5306664 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/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};
+ }
}
}
diff --git a/clang/test/OpenMP/target_data_use_device_addr_codegen.cpp b/clang/test/OpenMP/target_data_use_device_addr_codegen.cpp
index 7e70cdf74ad37..ae0653d0585d4 100644
--- a/clang/test/OpenMP/target_data_use_device_addr_codegen.cpp
+++ b/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]]
diff --git a/openmp/libomptarget/test/offloading/target_map_for_member_data.cpp b/openmp/libomptarget/test/offloading/target_map_for_member_data.cpp
new file mode 100644
index 0000000000000..8c8b4668c32e6
--- /dev/null
+++ 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[2];
+};
+
+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
More information about the cfe-commits
mailing list