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
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
