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

[Johannes Doerfert via Phabricator via cfe-commits]

jdoerfert added a comment.

This looks pretty good. The device vs if clause thing should be fixed while we are here.



================
Comment at: clang/lib/CodeGen/CGOpenMPRuntime.cpp:10399
+    DeviceID = CGF.Builder.getInt64(OMP_DEVICEID_UNDEF);
+  }
 
----------------
Move this behind the if generation and guard the emit scalar with a conditional, all is only needed if both Device is present and IfCond is not known true.


================
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:[^,]+]]
----------------
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 


================
Comment at: clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp:133-134
   }
   // CK1:     [[BEND]]:
-  // CK1:     [[CMP:%.+]] = icmp ne ptr %{{.+}}, null
   // CK1:     br i1 [[CMP]], label %[[BTHEN:.+]], label %[[BELSE:.+]]
----------------
TIFitis wrote:
> I think it is incorrect to recompute the branch condition here.
> 
> Consider the following C code:
> ```
> int a = 1;
> #pragma omp target data map(tofrom : a) if(a > 10) {a = 100;}
> ```
> In this case the if condition is false before executing the region, but becomes true after.
> If the branch condition is recomputed then it would take the `IfThen` branch here for executing the end_mapper call. This is incorrect and the end_mapper call should not be made partially without a begin _mapper call here.
> 
> Using the OMPIRBuilder fixes this.
Agreed. In general, we should always reuse the values we compute at the beginning of a region.


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D150860