[Openmp-commits] [PATCH] D142508: [OpenMP][libomptarget] Fix alignment calculation for mapping struct members.

[Pavel Kopyl via Phabricator via Openmp-commits]

pavelkopyl created this revision.
Herald added subscribers: guansong, yaxunl.
Herald added a project: All.
pavelkopyl requested review of this revision.
Herald added a reviewer: jdoerfert.
Herald added subscribers: openmp-commits, sstefan1.
Herald added a project: OpenMP.

This patch fixes calculation of 'HstPtrBegin' address to prevent
'explicit extension not allowed: ...' error that may happen when the
structure being mapped has less that 8-byte alignment:

  #include <omp.h>
  int main() {
    struct S {
      int i;
      int j;
    } s;
    s.i = 20;
    s.j = 30;
    #pragma omp target data map(tofrom : s)
    {
      #pragma omp target map(from : s.i, s.j)
      {
        s.i = 21;
        s.j = 31;
      }
    }
    return 0;
  }

In case 's' object is only 4-byte aligned (please, see D135462 <https://reviews.llvm.org/D135462> for details),
we may get the following error error:

  ...
  Libomptarget --> Looking up mapping(HstPtrBegin=0x00007ffc9b2581e4, Size=8)...
  Libomptarget --> Creating new map entry with HstPtrBase=0x00007ffc9b2581e4, HstPtrBegin=0x00007ffc9b2581e4...
  ...
  Libomptarget --> Moving 8 bytes (hst:0x00007ffc9b2581e4) -> (tgt:0x0000561d0739b390)
  ...
  Libomptarget --> Entry  0: Base=0x00007ffc9b2581e4, Begin=0x00007ffc9b2581e4, Size=8, Type=0x20, Name=unknown
  Libomptarget --> Entry  1: Base=0x00007ffc9b2581e4, Begin=0x00007ffc9b2581e4, Size=4, Type=0x1000000000002, Name=s.i
  Libomptarget --> Entry  2: Base=0x00007ffc9b2581e4, Begin=0x00007ffc9b2581e8, Size=4, Type=0x1000000000002, Name=s.j
  ...
  Libomptarget --> Using a padding of 4 bytes for begin address 0x00007ffc9b2581e4
  Libomptarget --> Looking up mapping(HstPtrBegin=0x00007ffc9b2581e0, Size=12)...
  Libomptarget --> WARNING: Pointer is not mapped but section extends into already mapped data
  Libomptarget message: explicit extension not allowed: host address specified is 0x00007ffc9b2581e0 (12 bytes),
   but device allocation maps to host at 0x00007ffc9b2581e4 (8 bytes)

Here we have wrong HstPtrBegin address - 0x00007ffc9b2581e0 that points
to a memory before 's' object itself, which starts at HstPtrBase (0x00007ffc9b2581e4).
This, in turn, leads to inconsistency in mapping sizes created and requested.


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D142508

Files:
  openmp/libomptarget/src/omptarget.cpp


Index: openmp/libomptarget/src/omptarget.cpp
===================================================================
--- openmp/libomptarget/src/omptarget.cpp
+++ openmp/libomptarget/src/omptarget.cpp
@@ -586,7 +586,7 @@
     if (getParentIndex(ArgTypes[I]) < 0 && NextI < ArgNum &&
         getParentIndex(ArgTypes[NextI]) == I) {
       Padding = (int64_t)HstPtrBegin % Alignment;
-      if (Padding) {
+      if (Padding && ((int64_t)HstPtrBegin - (int64_t)HstPtrBase >= Padding)) {
         DP("Using a padding of %" PRId64 " bytes for begin address " DPxMOD
            "\n",
            Padding, DPxPTR(HstPtrBegin));
@@ -940,7 +940,8 @@
     if (getParentIndex(ArgTypes[I]) < 0 && NextI < ArgNum &&
         getParentIndex(ArgTypes[NextI]) == I) {
       int64_t Padding = (int64_t)HstPtrBegin % Alignment;
-      if (Padding) {
+
+      if (Padding && ((int64_t)HstPtrBegin - (int64_t)ArgBases[I] >= Padding)) {
         DP("Using a Padding of %" PRId64 " bytes for begin address " DPxMOD
            "\n",
            Padding, DPxPTR(HstPtrBegin));


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D142508.491915.patch
Type: text/x-patch
Size: 1054 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/openmp-commits/attachments/20230124/9980e5fd/attachment.bin>