[PATCH] D150860: [OpenMP] Change clang emitTargetDataCalls to use OMPIRBuilder

[Johannes Doerfert via Phabricator via cfe-commits]

jdoerfert added inline comments.


================
Comment at: clang/test/OpenMP/target_data_codegen.cpp:355-356
 // Region 00
+// CK2-DAG: [[DEV:%[^,]+]] = sext i32 [[DEVi32:%[^,]+]] to i64
+// CK2-DAG: [[DEVi32]] = load i32, ptr %{{[^,]+}},
 // CK2: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
----------------
jdoerfert wrote:
> TIFitis wrote:
> > jsjodin wrote:
> > > TIFitis wrote:
> > > > When both if clause and device clause are present, the device clause argument is inadvertently brought outside the `IfThen` region as we emit the `llvm:Value` from the `Clang::Expr` for it before making call to createTargetData.
> > > > 
> > > > I don't think this change would affect any use cases.
> > > > When both if clause and device clause are present, the device clause argument is inadvertently brought outside the `IfThen` region as we emit the `llvm:Value` from the `Clang::Expr` for it before making call to createTargetData.
> > > > 
> > > > I don't think this change would affect any use cases.
> > > 
> > > Is it at all possible that the load could cause an exception if moved outside the if?
> > > 
> > Int pointer is not allowed inside device clause, so I don't think think the load can cause exceptions.
> > 
> > Also this would fix a scenario similar to the if clause, where in the following code `target_begin` would get device as //10// and `target_end` mapper call would get device as //100//.
> > 
> > ```
> > int arg = 10;
> > #pragma omp target data map(to: arg) if(arg < 20) device(arg)
> >   {arg = 100;}
> > ```
> > 
> Resting the value is good, computing it before the if is not. Take:
> 
> `... if(isValidPtr(p)) device(*p)`
> 
> We can generate the good code fairly easily, see above. Won't trigger much and will almost always be folded to straight line code anyway, except in cases like above, which is good. We end up with a phi and reuse the device value, which is also good. Win Win Win 
-Resting +Reusing (apply auto-correction)


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D150860