[Openmp-commits] [openmp] 002d61d - [OpenMP] Fix `present` for exit from `omp target data`
Joel E. Denny via Openmp-commits
openmp-commits at lists.llvm.org
Wed Aug 5 07:04:37 PDT 2020
Author: Joel E. Denny
Date: 2020-08-05T10:03:31-04:00
New Revision: 002d61db2b7790dc884953bf9271878bf0af3a8e
URL: https://github.com/llvm/llvm-project/commit/002d61db2b7790dc884953bf9271878bf0af3a8e
DIFF: https://github.com/llvm/llvm-project/commit/002d61db2b7790dc884953bf9271878bf0af3a8e.diff
LOG: [OpenMP] Fix `present` for exit from `omp target data`
Without this patch, the following example fails but shouldn't
according to OpenMP TR8:
```
#pragma omp target enter data map(alloc:i)
#pragma omp target data map(present, alloc: i)
{
#pragma omp target exit data map(delete:i)
} // fails presence check here
```
OpenMP TR8 sec. 2.22.7.1 "map Clause", p. 321, L23-26 states:
> If the map clause appears on a target, target data, target enter
> data or target exit data construct with a present map-type-modifier
> then on entry to the region if the corresponding list item does not
> appear in the device data environment an error occurs and the
> program terminates.
There is no corresponding statement about the exit from a region.
Thus, the `present` modifier should:
1. Check for presence upon entry into any region, including a `target
exit data` region. This behavior is already implemented correctly.
2. Should not check for presence upon exit from any region, including
a `target` or `target data` region. Without this patch, this
behavior is not implemented correctly, breaking the above example.
In the case of `target data`, this patch fixes the latter behavior by
removing the `present` modifier from the map types Clang generates for
the runtime call at the end of the region.
In the case of `target`, we have not found a valid OpenMP program for
which such a fix would matter. It appears that, if a program can
guarantee that data is present at the beginning of a `target` region
so that there's no error there, that data is also guaranteed to be
present at the end. This patch adds a comment to the runtime to
document this case.
Reviewed By: grokos, RaviNarayanaswamy, ABataev
Differential Revision: https://reviews.llvm.org/D84422
Added:
openmp/libomptarget/test/mapping/present/target_data_at_exit.c
Modified:
clang/lib/CodeGen/CGOpenMPRuntime.cpp
clang/lib/CodeGen/CGOpenMPRuntime.h
clang/lib/CodeGen/CGStmtOpenMP.cpp
clang/test/OpenMP/target_data_codegen.cpp
openmp/libomptarget/src/omptarget.cpp
Removed:
################################################################################
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 60c7081b135b..547a9307dce2 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -8826,6 +8826,30 @@ emitOffloadingArrays(CodeGenFunction &CGF,
MapTypesArrayGbl->setUnnamedAddr(llvm::GlobalValue::UnnamedAddr::Global);
Info.MapTypesArray = MapTypesArrayGbl;
+ // If there's a present map type modifier, it must not be applied to the end
+ // of a region, so generate a separate map type array in that case.
+ if (Info.separateBeginEndCalls()) {
+ bool EndMapTypesDiffer = false;
+ for (uint64_t &Type : Mapping) {
+ if (Type & MappableExprsHandler::OMP_MAP_PRESENT) {
+ Type &= ~MappableExprsHandler::OMP_MAP_PRESENT;
+ EndMapTypesDiffer = true;
+ }
+ }
+ if (EndMapTypesDiffer) {
+ MapTypesArrayInit =
+ llvm::ConstantDataArray::get(CGF.Builder.getContext(), Mapping);
+ MaptypesName = CGM.getOpenMPRuntime().getName({"offload_maptypes"});
+ MapTypesArrayGbl = new llvm::GlobalVariable(
+ CGM.getModule(), MapTypesArrayInit->getType(),
+ /*isConstant=*/true, llvm::GlobalValue::PrivateLinkage,
+ MapTypesArrayInit, MaptypesName);
+ MapTypesArrayGbl->setUnnamedAddr(
+ llvm::GlobalValue::UnnamedAddr::Global);
+ Info.MapTypesArrayEnd = MapTypesArrayGbl;
+ }
+ }
+
for (unsigned I = 0; I < Info.NumberOfPtrs; ++I) {
llvm::Value *BPVal = *CombinedInfo.BasePointers[I];
llvm::Value *BP = CGF.Builder.CreateConstInBoundsGEP2_32(
@@ -8878,12 +8902,16 @@ emitOffloadingArrays(CodeGenFunction &CGF,
}
/// Emit the arguments to be passed to the runtime library based on the
-/// arrays of base pointers, pointers, sizes, map types, and mappers.
+/// arrays of base pointers, pointers, sizes, map types, and mappers. If
+/// ForEndCall, emit map types to be passed for the end of the region instead of
+/// the beginning.
static void emitOffloadingArraysArgument(
CodeGenFunction &CGF, llvm::Value *&BasePointersArrayArg,
llvm::Value *&PointersArrayArg, llvm::Value *&SizesArrayArg,
llvm::Value *&MapTypesArrayArg, llvm::Value *&MappersArrayArg,
- CGOpenMPRuntime::TargetDataInfo &Info) {
+ CGOpenMPRuntime::TargetDataInfo &Info, bool ForEndCall = false) {
+ assert((!ForEndCall || Info.separateBeginEndCalls()) &&
+ "expected region end call to runtime only when end call is separate");
CodeGenModule &CGM = CGF.CGM;
if (Info.NumberOfPtrs) {
BasePointersArrayArg = CGF.Builder.CreateConstInBoundsGEP2_32(
@@ -8900,7 +8928,8 @@ static void emitOffloadingArraysArgument(
/*Idx0=*/0, /*Idx1=*/0);
MapTypesArrayArg = CGF.Builder.CreateConstInBoundsGEP2_32(
llvm::ArrayType::get(CGM.Int64Ty, Info.NumberOfPtrs),
- Info.MapTypesArray,
+ ForEndCall && Info.MapTypesArrayEnd ? Info.MapTypesArrayEnd
+ : Info.MapTypesArray,
/*Idx0=*/0,
/*Idx1=*/0);
MappersArrayArg =
@@ -10267,7 +10296,7 @@ void CGOpenMPRuntime::emitTargetDataCalls(
llvm::Value *MappersArrayArg = nullptr;
emitOffloadingArraysArgument(CGF, BasePointersArrayArg, PointersArrayArg,
SizesArrayArg, MapTypesArrayArg,
- MappersArrayArg, Info);
+ MappersArrayArg, Info, /*ForEndCall=*/false);
// Emit device ID if any.
llvm::Value *DeviceID = nullptr;
@@ -10307,7 +10336,7 @@ void CGOpenMPRuntime::emitTargetDataCalls(
llvm::Value *MappersArrayArg = nullptr;
emitOffloadingArraysArgument(CGF, BasePointersArrayArg, PointersArrayArg,
SizesArrayArg, MapTypesArrayArg,
- MappersArrayArg, Info);
+ MappersArrayArg, Info, /*ForEndCall=*/true);
// Emit device ID if any.
llvm::Value *DeviceID = nullptr;
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.h b/clang/lib/CodeGen/CGOpenMPRuntime.h
index 0b91975343f7..1536181d67c5 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.h
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.h
@@ -1614,6 +1614,9 @@ class CGOpenMPRuntime {
class TargetDataInfo {
/// Set to true if device pointer information have to be obtained.
bool RequiresDevicePointerInfo = false;
+ /// Set to true if Clang emits separate runtime calls for the beginning and
+ /// end of the region. These calls might have separate map type arrays.
+ bool SeparateBeginEndCalls = false;
public:
/// The array of base pointer passed to the runtime library.
@@ -1622,8 +1625,14 @@ class CGOpenMPRuntime {
llvm::Value *PointersArray = nullptr;
/// The array of sizes passed to the runtime library.
llvm::Value *SizesArray = nullptr;
- /// The array of map types passed to the runtime library.
+ /// The array of map types passed to the runtime library for the beginning
+ /// of the region or for the entire region if there are no separate map
+ /// types for the region end.
llvm::Value *MapTypesArray = nullptr;
+ /// The array of map types passed to the runtime library for the end of the
+ /// region, or nullptr if there are no separate map types for the region
+ /// end.
+ llvm::Value *MapTypesArrayEnd = nullptr;
/// The array of user-defined mappers passed to the runtime library.
llvm::Value *MappersArray = nullptr;
/// Indicate whether any user-defined mapper exists.
@@ -1635,14 +1644,17 @@ class CGOpenMPRuntime {
llvm::DenseMap<const ValueDecl *, Address> CaptureDeviceAddrMap;
explicit TargetDataInfo() {}
- explicit TargetDataInfo(bool RequiresDevicePointerInfo)
- : RequiresDevicePointerInfo(RequiresDevicePointerInfo) {}
+ explicit TargetDataInfo(bool RequiresDevicePointerInfo,
+ bool SeparateBeginEndCalls)
+ : RequiresDevicePointerInfo(RequiresDevicePointerInfo),
+ SeparateBeginEndCalls(SeparateBeginEndCalls) {}
/// Clear information about the data arrays.
void clearArrayInfo() {
BasePointersArray = nullptr;
PointersArray = nullptr;
SizesArray = nullptr;
MapTypesArray = nullptr;
+ MapTypesArrayEnd = nullptr;
MappersArray = nullptr;
HasMapper = false;
NumberOfPtrs = 0u;
@@ -1653,6 +1665,7 @@ class CGOpenMPRuntime {
MapTypesArray && (!HasMapper || MappersArray) && NumberOfPtrs;
}
bool requiresDevicePointerInfo() { return RequiresDevicePointerInfo; }
+ bool separateBeginEndCalls() { return SeparateBeginEndCalls; }
};
/// Emit the target data mapping code associated with \a D.
diff --git a/clang/lib/CodeGen/CGStmtOpenMP.cpp b/clang/lib/CodeGen/CGStmtOpenMP.cpp
index df1cc1666de4..0bb57171db44 100644
--- a/clang/lib/CodeGen/CGStmtOpenMP.cpp
+++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp
@@ -6039,7 +6039,8 @@ void CodeGenFunction::EmitOMPUseDeviceAddrClause(
// Generate the instructions for '#pragma omp target data' directive.
void CodeGenFunction::EmitOMPTargetDataDirective(
const OMPTargetDataDirective &S) {
- CGOpenMPRuntime::TargetDataInfo Info(/*RequiresDevicePointerInfo=*/true);
+ CGOpenMPRuntime::TargetDataInfo Info(/*RequiresDevicePointerInfo=*/true,
+ /*SeparateBeginEndCalls=*/true);
// Create a pre/post action to signal the privatization of the device pointer.
// This action can be replaced by the OpenMP runtime code generation to
diff --git a/clang/test/OpenMP/target_data_codegen.cpp b/clang/test/OpenMP/target_data_codegen.cpp
index a2bb8cdf5ba8..67401b38bc32 100644
--- a/clang/test/OpenMP/target_data_codegen.cpp
+++ b/clang/test/OpenMP/target_data_codegen.cpp
@@ -256,10 +256,16 @@ ST<int> gb;
double gc[100];
// PRESENT=0x1000 | TARGET_PARAM=0x20 | TO=0x1 = 0x1021
-// CK1A: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1021]]]
+// CK1A: [[MTYPE00Begin:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1021]]]
+
+// TARGET_PARAM=0x20 | TO=0x1 = 0x21
+// CK1A: [[MTYPE00End:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x21]]]
// PRESENT=0x1000 | CLOSE=0x400 | TARGET_PARAM=0x20 | ALWAYS=0x4 | TO=0x1 = 0x1425
-// CK1A: [[MTYPE01:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1425]]]
+// CK1A: [[MTYPE01Begin:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1425]]]
+
+// CLOSE=0x400 | TARGET_PARAM=0x20 | ALWAYS=0x4 | TO=0x1 = 0x425
+// CK1A: [[MTYPE01End:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x425]]]
// CK1A-LABEL: _Z3fooi
void foo(int arg) {
@@ -267,7 +273,7 @@ void foo(int arg) {
float lb[arg];
// Region 00
- // CK1A-DAG: call void @__tgt_target_data_begin_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz:32|64]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}})
+ // CK1A-DAG: call void @__tgt_target_data_begin_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz:32|64]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00Begin]]{{.+}})
// CK1A-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK1A-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
// CK1A-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
@@ -285,7 +291,7 @@ void foo(int arg) {
// CK1A-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
// CK1A: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
- // CK1A-DAG: call void @__tgt_target_data_end_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}})
+ // CK1A-DAG: call void @__tgt_target_data_end_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00End]]{{.+}})
// CK1A-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
// CK1A-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P]]
// CK1A-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S]]
@@ -293,7 +299,7 @@ void foo(int arg) {
{++arg;}
// Region 01
- // CK1A-DAG: call void @__tgt_target_data_begin_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}})
+ // CK1A-DAG: call void @__tgt_target_data_begin_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01Begin]]{{.+}})
// CK1A-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK1A-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
// CK1A-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
@@ -311,7 +317,7 @@ void foo(int arg) {
// CK1A-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
// CK1A: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
- // CK1A-DAG: call void @__tgt_target_data_end_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}})
+ // CK1A-DAG: call void @__tgt_target_data_end_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01End]]{{.+}})
// CK1A-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
// CK1A-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P]]
// CK1A-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S]]
diff --git a/openmp/libomptarget/src/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp
index f4d79d8064b9..541fe0d52fbb 100644
--- a/openmp/libomptarget/src/omptarget.cpp
+++ b/openmp/libomptarget/src/omptarget.cpp
@@ -506,8 +506,14 @@ int targetDataEnd(DeviceTy &Device, int32_t ArgNum, void **ArgBases,
DP("Mapping does not exist (%s)\n",
(HasPresentModifier ? "'present' map type modifier" : "ignored"));
if (HasPresentModifier) {
- // FIXME: This should not be an error on exit from "omp target data",
- // but it should be an error upon entering an "omp target exit data".
+ // This should be an error upon entering an "omp target exit data". It
+ // should not be an error upon exiting an "omp target data" or "omp
+ // target". For "omp target data", Clang thus doesn't include present
+ // modifiers for end calls. For "omp target", we have not found a valid
+ // OpenMP program for which the error matters: it appears that, if a
+ // program can guarantee that data is present at the beginning of an
+ // "omp target" region so that there's no error there, that data is also
+ // guaranteed to be present at the end.
MESSAGE("device mapping required by 'present' map type modifier does "
"not exist for host address " DPxMOD " (%ld bytes)",
DPxPTR(HstPtrBegin), DataSize);
diff --git a/openmp/libomptarget/test/mapping/present/target_data_at_exit.c b/openmp/libomptarget/test/mapping/present/target_data_at_exit.c
new file mode 100644
index 000000000000..9a258ba57a50
--- /dev/null
+++ b/openmp/libomptarget/test/mapping/present/target_data_at_exit.c
@@ -0,0 +1,37 @@
+// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-aarch64-unknown-linux-gnu 2>&1 \
+// RUN: | %fcheck-aarch64-unknown-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-powerpc64-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-powerpc64le-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64le-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-x86_64-pc-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-x86_64-pc-linux-gnu 2>&1 \
+// RUN: | %fcheck-x86_64-pc-linux-gnu
+
+#include <stdio.h>
+
+int main() {
+ int i;
+
+#pragma omp target enter data map(alloc:i)
+
+ // i isn't present at the end of the target data region, but the "present"
+ // modifier is only checked at the beginning of a region.
+#pragma omp target data map(present, alloc: i)
+ {
+#pragma omp target exit data map(delete:i)
+ }
+
+ // CHECK-NOT: Libomptarget
+ // CHECK: success
+ // CHECK-NOT: Libomptarget
+ fprintf(stderr, "success\n");
+
+ return 0;
+}
More information about the Openmp-commits
mailing list