https://github.com/jyu2-git updated https://github.com/llvm/llvm-project/pull/74692
>From 50c6009e4f4184ed7710a7ee3d8ee0983306edc1 Mon Sep 17 00:00:00 2001 From: Jennifer Yu <jennifer...@intel.com> Date: Wed, 6 Dec 2023 13:53:16 -0800 Subject: [PATCH 1/2] [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 >From 1027a9d197dfad447cc73e74afdbba525685831c Mon Sep 17 00:00:00 2001 From: Jennifer Yu <jennifer...@intel.com> Date: Thu, 7 Dec 2023 08:27:44 -0800 Subject: [PATCH 2/2] Thanks Alexey for the review. This is address his comment. --- .../libomptarget/test/offloading/target_map_for_member_data.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/openmp/libomptarget/test/offloading/target_map_for_member_data.cpp b/openmp/libomptarget/test/offloading/target_map_for_member_data.cpp index 8c8b4668c32e6..812ede6fc8a26 100644 --- a/openmp/libomptarget/test/offloading/target_map_for_member_data.cpp +++ b/openmp/libomptarget/test/offloading/target_map_for_member_data.cpp @@ -4,7 +4,7 @@ struct DataTy { float a; - float b[2]; + float b[3]; }; int main(int argc, char **argv) { _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits