[all-commits] [llvm/llvm-project] 011372: [OpenMP] Fix runtime problem due to wrong map size...

[jyu2-git via All-commits]

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: 0113722d82200c39e59dcfbd2f396dbd84ed022b
      https://github.com/llvm/llvm-project/commit/0113722d82200c39e59dcfbd2f396dbd84ed022b
  Author: jyu2-git <jennifer.yu at intel.com>
  Date:   2023-12-07 (Thu, 07 Dec 2023)

  Changed paths:
    M clang/lib/CodeGen/CGOpenMPRuntime.cpp
    M clang/test/OpenMP/target_data_use_device_addr_codegen.cpp
    A openmp/libomptarget/test/offloading/target_map_for_member_data.cpp

  Log Message:
  -----------
  [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)