[PATCH] D72811: [WIP][OPENMP5.0] allow lvalue for motion clause

Chi Chun Chen via Phabricator via cfe-commits cfe-commits at lists.llvm.org
Tue Feb 11 12:33:22 PST 2020


cchen marked an inline comment as done.
cchen added a comment.

> Quoted Text





================
Comment at: clang/lib/Sema/SemaOpenMP.cpp:15554
+        LocatorChecker Checker;
+        if (Checker.Visit(OrigExpr)) {
+          llvm::copy(Checker.getComponents(),
----------------
ABataev wrote:
> cchen wrote:
> > ABataev wrote:
> > > cchen wrote:
> > > > ABataev wrote:
> > > > > General question about several cases. How we're going to support them?
> > > > > 1. (a ? b : c). 
> > > > > 2. __builtin_choose_expr(a, b, c).
> > > > > 3. a = b.
> > > > > 4. a?:b
> > > > > 5. __builtin_convertvector(x, ty)
> > > > > 6. (int&)a
> > > > > 7. v.xy, where v is an extended vector
> > > > > 8. a.b, where b is a bitfield
> > > > > 9. __builtin_bit_cast(v, ty)
> > > > > 10. const_cast<ty &>(a)
> > > > > 11. dynamic_cast<ty &>(a)
> > > > > 12. reinterpret_cast
> > > > > 13. static_cast
> > > > > 14. typeid() (also is an lvalue).
> > > > > 15. __uuidof(*comPtr)
> > > > > 16. lambda calls
> > > > > 17. User defined literals
> > > > > 
> > > > I think we could first evaluate the lvalue, and then pass the result of `&lvalue` to device? 
> > > What do you mean when you say `evaluate lvalue`? In veneral, it can be evaluated at the runtime. Do you propose to implement dyic translation? It will definetely slow down the execution on the device.
> > I mean evaluate lvalue before sending &lvalue to device. For example:
> > ```
> > int a, b;
> > b = 5;
> > #pragma omp target map(a = b) // assign b to a before sending `&(a=b)` to device
> > {
> >    a++;
> > }
> > ```
> > Therefore, the above example have the same semantics as this:
> > ```
> > int a, b;
> > b = 5;
> > int &c = (a=b)
> > #pragma omp target map(c)
> > {
> >    a++;
> > }
> > ```
> Sure, we do this already, generally speaking. The main problem here is how to map the address and associate it with address on the device. We do it at the compile time, support of general lvalues requires runtime translation + special instrumentation of the device code for address translation.
  - (a ? b : c).
    - Error out for now, not sure what to for this one
  - __builtin_choose_expr(a, b, c).
    - Allow this one since we know what to send to device at compile time
  - a = b.
   - Sema accept this one, but need more work for codegen I think (Clang not emit the a=b ir if "a=b" inside map/motion clause)
  - a?:b
    - What's the name of this kind of expression? I don't know what to do for this one
  - __builtin_convertvector(x, ty)
    - not sure
  - v.xy, where v is an extended vector
    - not sure
  - a.b, where b is a bitfield
    - Error out for this one since bitfield is not addressable
  - __builtin_bit_cast(v, ty)
    - Error out for this one since bitfield is not addressable
  - casts
    - Error out if the expression inside the cast is not lvalue
  - typeid() (also is an lvalue).
    - I guess we can support this just like normal variable?
  - __uuidof(*comPtr)
    - not sure
  - lambda calls
    - OpenMP 5.0 spec has rules for lambda. Now not error out for sema but need more work for codegen
  - User defined literals
     - error out since not addressable



Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D72811





More information about the cfe-commits mailing list