[OpenMP] Change clang emitTargetDataCalls to use OMPIRBuilder

[TIFitis]

Summary

This patch changes the emitTargetDataCalls function in clang to make use of the OpenMPIRBuilder::createTargetData function for Target Data directive code gen.

Depends on D146557

[TIFitis]

clang/test/OpenMP/target_data_codegen.cpp

67–68

Similar to what I discussed for the if clause, the device clause argument need not be recomputed.
We can reuse the device argument from the begin mapper for the end mapper as well.

355–356

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.

clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp

133–134

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.

[TIFitis]

Rebased

[TIFitis]

Ping for reviews :)

[jplehr]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

10409

I think this can (and should be?) moved into the if in the following line as it is not used outside that scope

10413

capitalize here?

[TIFitis]

Addressed reviewer comments.

[jsjodin]

clang/test/OpenMP/target_data_codegen.cpp

355–356

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?

[TIFitis]

clang/test/OpenMP/target_data_codegen.cpp

355–356

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;}

[TIFitis]

Rebased

[TIFitis]

Rebased

[TIFitis]

Rebased

[jdoerfert]

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

clang/lib/CodeGen/CGOpenMPRuntime.cpp

10415

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.

clang/test/OpenMP/target_data_codegen.cpp

355–356

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

clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp

133–134

Agreed. In general, we should always reuse the values we compute at the beginning of a region.

[jdoerfert]

clang/test/OpenMP/target_data_codegen.cpp

355–356

-Resting +Reusing (apply auto-correction)

[TIFitis]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

10415

I have moved the device generation after the if generation.
But I am not sure what should be the guard here for emitting the scalar.

My idea was to generate something like:

%new_dev = alloca ...
if.then
    %x = load %device
    store %x %new_dev
    ...
    begin_mapper(..., %new_dev, ...)
...
if.then
    ...
    end_mapper(..., %new_dev, ...)

[TIFitis]

Moved device clause emition after if clause.

[TIFitis]

Gentle ping :)

[arsenm]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

10403–10407

Why do you need to do this? I would expect EvaluateExprAsBool to handle the constant case for you

[TIFitis]

Addressed reviewer comments

clang/lib/CodeGen/CGOpenMPRuntime.cpp

10403–10407

CGOpenMPRuntime::emitIfClause has the check for ConstantFoldsToSimpleInteger so I added it here.

Removing it didn't break any test so I've updated it,

[jdoerfert]

Let's move on. I described what this should look like, just for the record.

clang/lib/CodeGen/CGOpenMPRuntime.cpp

10415

What we should do, though I don't think it will really matter, is the following:

device_val = undef;
if_val = eval_if_clause();
if (id_val) {
  device_val = eval_device_clause()
  ...other clauses ...
  begin_then(... Vals ...)
} else {
  begin_else(...)
}

if (if_val) {
  end_then(... Vals ...)
} else {
  end_else(...)
}

[TIFitis]

Rebased to fix build error.