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

[Akash Banerjee via Phabricator via llvm-commits]

TIFitis added inline comments.


================
Comment at: clang/test/OpenMP/target_data_codegen.cpp:67-68
-  // CK1-DAG: call void @__tgt_target_data_end_mapper(ptr @{{.+}}, i64 [[DEV:%[^,]+]], i32 1, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[SIZE00]], ptr [[MTYPE00]], ptr null, ptr null)
-  // CK1-DAG: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64
-  // CK1-DAG: [[DEVi32]] = load i32, ptr %{{[^,]+}},
   // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
----------------
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.


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


================
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:.+]]
----------------
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.


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D150860