[PATCH] D133694: [Clang][OpenMP] Fix use_device_addr

Gheorghe-Teodor Bercea via Phabricator via cfe-commits cfe-commits at lists.llvm.org
Mon Sep 12 11:23:06 PDT 2022


doru1004 added inline comments.


================
Comment at: clang/test/OpenMP/target_data_use_device_addr_codegen_ptr.cpp:14
+    {
+        #pragma omp target data use_device_addr(x)
+        {
----------------
ye-luo wrote:
> doru1004 wrote:
> > ye-luo wrote:
> > > In my understanding of the spec.
> > > `map(tofrom:x[0:256])` only maps the memory segment that x points to. x itself as a pointer scalar is not mapped.
> > > use_device_addr(x) should fail to find the map of x scalar.
> > > 5.2 spec.
> > > If the list item is not a mapped list item, it is assumed to be accessible on the target device.
> > > To me, it seems just keep &x as it was, in this case &x remains a host address.
> > > 
> > > But in your patch description, it seems treating x differently from a scalar.
> > > 
> > > I also applied your patch on main and got segfault because the x has a value of device address and x[0] fails. This should be the behavior of use_device_ptr instead of use_device_addr.
> > > To me, it seems just keep &x as it was, in this case &x remains a host address.
> > 
> > So does this mean that if I do something like this in the target data I should get different addresses for x:
> > 
> > 
> > ```
> >         #pragma omp target data use_device_ptr(x)
> >         {
> >             fprintf(stderr, "x: %p\n", __LINE__, x);
> >         }
> > 
> >         #pragma omp target data use_device_addr(x)
> >         {
> >             fprintf(stderr, "x: %p\n", __LINE__, x);
> >         }
> > ```
> > 
> > 
> > > I also applied your patch on main and got segfault because the x has a value of device address and x[0] fails.
> > 
> > That's my fault x[0] was the wrong thing to use actually.
> > 
> > 
> When you have an outer `target data map(x)`, then two printf differ. If there is no outer `map(x)`, two printf should be identical.
> When you have an outer `target data map(x)`, then two printf differ. If there is no outer `map(x)`, two printf should be identical.

This is super helpful thank you! I'll make sure that happens.

In the case when an outer target data exists, the print of the x which is under use_device_addr should print the same address as printing x on the host? 


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

https://reviews.llvm.org/D133694



More information about the cfe-commits mailing list