[PATCH] D91370: [OPENMP]Fix PR48076: Check map types array before accessing its front.

[Joseph Huber via Phabricator via cfe-commits]

jhuber6 added a comment.

In D91370#2392616 <https://reviews.llvm.org/D91370#2392616>, @ABataev wrote:

> Could you try this patch:
>
>   diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
>   index ce8846140d4..854b7f3e830 100644
>   --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
>   +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
>   @@ -9938,7 +9938,7 @@ void CGOpenMPRuntime::emitTargetCall(
>              MappedVarSet.insert(CI->getCapturedVar());
>            else
>              MappedVarSet.insert(nullptr);
>   -        if (CurInfo.BasePointers.empty())
>   +        if (CurInfo.BasePointers.empty() && !PartialStruct.Base.isValid())
>              MEHandler.generateDefaultMapInfo(*CI, **RI, *CV, CurInfo);
>            // Generate correct mapping for variables captured by reference in
>            // lambdas.
>   @@ -9947,7 +9947,7 @@ void CGOpenMPRuntime::emitTargetCall(
>                                                      CurInfo, LambdaPointers);
>          }
>          // We expect to have at least an element of information for this capture.
>   -      assert(!CurInfo.BasePointers.empty() &&
>   +      assert((!CurInfo.BasePointers.empty() || PartialStruct.Base.isValid()) &&
>                 "Non-existing map pointer for capture!");
>          assert(CurInfo.BasePointers.size() == CurInfo.Pointers.size() &&
>                 CurInfo.BasePointers.size() == CurInfo.Sizes.size() &&

It stopped  the crashing but I'm not sure it's working as intended. The following program doesn't crash, but it doesn't change the value of `s.p`. Looking at the debug output it doesn't seem to be mapping the whole struct.

  struct S {
    float f[50];
    double *p;
  };
  
  int main() {
    S s;
    printf(%p\n", s.p);
  #pragma omp target map(tofrom:s.p)
    { s.p = nullptr; }
    printf(%p\n", s.p);
  }

  0x7ffea8eec8f0
  0x7ffea8eec8f0

  Libomptarget --> Entry  0: Base=0x00007fffa3265500, Begin=0x00007fffa3265690, Size=8, Type=0x20
  Libomptarget --> Looking up mapping(HstPtrBegin=0x00007fffa3265690, Size=8)...
  Libomptarget --> MemoryManagerTy::allocate: size 8 with host pointer 0x00007fffa3265690.
  Libomptarget --> findBucket: Size 8 is floored to 8.
  Libomptarget --> Cannot find a node in the FreeLists. Allocate on device.
  Libomptarget --> Node address 0x00000000010511e8, target pointer 0x00007fc846400000, size 8
  Libomptarget --> Creating new map entry: HstBase=0x00007fffa3265500, HstBegin=0x00007fffa3265690, HstEnd=0x00007fffa3265698, TgtBegin=0x00007fc846400000
  Libomptarget --> There are 8 bytes allocated at target address 0x00007fc846400000 - is new
  Libomptarget --> Looking up mapping(HstPtrBegin=0x00007fffa3265690, Size=8)...
  Libomptarget --> Mapping exists with HstPtrBegin=0x00007fffa3265690, TgtPtrBegin=0x00007fc846400000, Size=8, RefCount=1
  Libomptarget --> Obtained target argument 0x00007fc8463ffe70 from host pointer 0x00007fffa3265690

However if it's just a capture with only `#pragma omp target` then it seems to work.

  0x7fff2b234ed0
  (nil)




Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D91370/new/

https://reviews.llvm.org/D91370