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:
> > > 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?
> I need a correction. When outer map(x) exists, actually the address(not
> value) of x should be a device address, and the code cannot even print x.
> Printing &x should be fine.
In the context of the above comment, should &x on the device be an address I
can verify, somehow, to make sure that it's correct or is it a completely new
device address?
So for example, should it be the same as when I do a use_device_ptr but print
the &x in that case? (With the current master those two addresses are not the
same.)
I guess what I need is an example of using use_device_addr that actually does
something meaningful because with the current main branch printing the &x of a
use_device_addr(x) is nil.
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D133694/new/
https://reviews.llvm.org/D133694
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
