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

[Alexey Bataev via Phabricator via cfe-commits]

ABataev added a comment.

In D91370#2392107 <https://reviews.llvm.org/D91370#2392107>, @jhuber6 wrote:

> This stops it from crashing Clang, but I'm not sure if it's fixing the underlying problem. When I compile and run this program it crashes in libomptarget. If you get rid of the `float f[50]` in the struct it works as expected.
>
> `$ clang++ -fopenmp -fopenmp-targets=nvptx64 test.cpp && ./a.out`
>
>   #include <stdio.h>
>   
>   struct S2 {
>       float f[50];
>       double *p; 
>   };
>   
>   int main() {
>     S2 s;
>   #pragma omp target map(s.p)
>     { s.p = nullptr; }
>   }
>
> Output from running with debugging
>
>   Libomptarget --> Device 0 is ready to use.
>   Target CUDA RTL --> Load data from image 0x0000000000400a30
>   Target CUDA RTL --> CUDA module successfully loaded!
>   Target CUDA RTL --> Entry point 0x0000000000000000 maps to __omp_offloading_fd02_267be70_main_l10 (0x000000000174b170)
>   Target CUDA RTL --> Sending global device environment data 4 bytes
>   Libomptarget --> Entry  0: Base=0x00007fffcca22ad8, Begin=0x00007fffcca22ba0, Size=8, Type=0x20
>   Libomptarget --> Entry  1: Base=0x00007fffcca22ad8, Begin=0x00007fffcca22ad8, Size=208, Type=0x1000000000203
>   Libomptarget --> Looking up mapping(HstPtrBegin=0x00007fffcca22ba0, Size=8)...
>   Libomptarget --> MemoryManagerTy::allocate: size 8 with host pointer 0x00007fffcca22ba0.
>   Libomptarget --> findBucket: Size 8 is floored to 8.
>   Libomptarget --> Cannot find a node in the FreeLists. Allocate on device.
>   Libomptarget --> Node address 0x0000000001732d88, target pointer 0x00007f0506400000, size 8
>   Libomptarget --> Creating new map entry: HstBase=0x00007fffcca22ad8, HstBegin=0x00007fffcca22ba0, HstEnd=0x00007fffcca22ba8, TgtBegin=0x00007f0506400000
>   Libomptarget --> There are 8 bytes allocated at target address 0x00007f0506400000 - is new
>   Libomptarget --> Looking up mapping(HstPtrBegin=0x00007fffcca22ad8, Size=208)...
>   Libomptarget --> WARNING: Pointer is not mapped but section extends into already mapped data
>   Libomptarget --> Mapping exists (implicit) with HstPtrBegin=0x00007fffcca22ad8, TgtPtrBegin=0x00007f05063fff38, Size=208, RefCount=1
>   Libomptarget --> There are 208 bytes allocated at target address 0x00007f05063fff38 - is not new
>   Libomptarget --> DeviceTy::getMapEntry: requested entry found
>   Libomptarget --> Moving 208 bytes (hst:0x00007fffcca22ad8) -> (tgt:0x00007f05063fff38)
>   Target CUDA RTL --> Error when copying data from host to device. Pointers: host = 0x00007fffcca22ad8, device = 0x00007f05063fff38, size = 208
>   Target CUDA RTL --> CUDA error is: invalid argument
>   Libomptarget --> Copying data to device failed.
>   Libomptarget --> Call to targetDataBegin failed, abort target.
>   Libomptarget --> Failed to process data before launching the kernel.

It is different problem, which should be addressed in the different patch.


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D91370