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
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
