[clang] a5ea676 - Reverting commit 0d8d718171192301f2beb10bd08ce62e70281a5e as it broke libomptarget tests
Akash Banerjee via cfe-commits
cfe-commits at lists.llvm.org
Fri Jun 30 08:04:52 PDT 2023
Author: Akash Banerjee
Date: 2023-06-30T16:03:57+01:00
New Revision: a5ea6760674762cb597cf328dc467f1296633da0
URL: https://github.com/llvm/llvm-project/commit/a5ea6760674762cb597cf328dc467f1296633da0
DIFF: https://github.com/llvm/llvm-project/commit/a5ea6760674762cb597cf328dc467f1296633da0.diff
LOG: Reverting commit 0d8d718171192301f2beb10bd08ce62e70281a5e as it broke libomptarget tests
Added:
Modified:
clang/lib/CodeGen/CGOpenMPRuntime.cpp
clang/test/OpenMP/target_data_codegen.cpp
clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp
llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
Removed:
################################################################################
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index dfc8f71ef43583..bd0169d2d80b32 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -9032,14 +9032,14 @@ static void emitOffloadingArrays(
InsertPointTy CodeGenIP(CGF.Builder.GetInsertBlock(),
CGF.Builder.GetInsertPoint());
- auto FillInfoMap = [&](MappableExprsHandler::MappingExprInfo &MapExpr) {
+ auto fillInfoMap = [&](MappableExprsHandler::MappingExprInfo &MapExpr) {
return emitMappingInformation(CGF, OMPBuilder, MapExpr);
};
if (CGM.getCodeGenOpts().getDebugInfo() !=
llvm::codegenoptions::NoDebugInfo) {
CombinedInfo.Names.resize(CombinedInfo.Exprs.size());
llvm::transform(CombinedInfo.Exprs, CombinedInfo.Names.begin(),
- FillInfoMap);
+ fillInfoMap);
}
auto DeviceAddrCB = [&](unsigned int I, llvm::Value *BP, llvm::Value *BPVal) {
@@ -10385,94 +10385,140 @@ void CGOpenMPRuntime::emitTargetDataCalls(
// off.
PrePostActionTy NoPrivAction;
- using InsertPointTy = llvm::OpenMPIRBuilder::InsertPointTy;
- InsertPointTy AllocaIP(CGF.AllocaInsertPt->getParent(),
- CGF.AllocaInsertPt->getIterator());
- InsertPointTy CodeGenIP(CGF.Builder.GetInsertBlock(),
- CGF.Builder.GetInsertPoint());
- llvm::OpenMPIRBuilder::LocationDescription OmpLoc(CodeGenIP);
-
- llvm::Value *IfCondVal = nullptr;
- if (IfCond)
- IfCondVal = CGF.EvaluateExprAsBool(IfCond);
-
- // Emit device ID if any.
- llvm::Value *DeviceID = nullptr;
- if (Device) {
- DeviceID = CGF.Builder.CreateIntCast(CGF.EmitScalarExpr(Device),
- CGF.Int64Ty, /*isSigned=*/true);
- } else {
- DeviceID = CGF.Builder.getInt64(OMP_DEVICEID_UNDEF);
- }
+ // Generate the code for the opening of the data environment. Capture all the
+ // arguments of the runtime call by reference because they are used in the
+ // closing of the region.
+ auto &&BeginThenGen = [this, &D, Device, &Info,
+ &CodeGen](CodeGenFunction &CGF, PrePostActionTy &) {
+ // Fill up the arrays with all the mapped variables.
+ MappableExprsHandler::MapCombinedInfoTy CombinedInfo;
- // Fill up the arrays with all the mapped variables.
- MappableExprsHandler::MapCombinedInfoTy CombinedInfo;
- auto GenMapInfoCB =
- [&](InsertPointTy CodeGenIP) -> llvm::OpenMPIRBuilder::MapInfosTy & {
- CGF.Builder.restoreIP(CodeGenIP);
// Get map clause information.
MappableExprsHandler MEHandler(D, CGF);
MEHandler.generateAllInfo(CombinedInfo);
- auto FillInfoMap = [&](MappableExprsHandler::MappingExprInfo &MapExpr) {
- return emitMappingInformation(CGF, OMPBuilder, MapExpr);
- };
- if (CGM.getCodeGenOpts().getDebugInfo() !=
- llvm::codegenoptions::NoDebugInfo) {
- CombinedInfo.Names.resize(CombinedInfo.Exprs.size());
- llvm::transform(CombinedInfo.Exprs, CombinedInfo.Names.begin(),
- FillInfoMap);
- }
+ // Fill up the arrays and create the arguments.
+ emitOffloadingArrays(CGF, CombinedInfo, Info, OMPBuilder,
+ /*IsNonContiguous=*/true);
- return CombinedInfo;
- };
- using BodyGenTy = llvm::OpenMPIRBuilder::BodyGenTy;
- auto BodyCB = [&](InsertPointTy CodeGenIP, BodyGenTy BodyGenType) {
- CGF.Builder.restoreIP(CodeGenIP);
- switch (BodyGenType) {
- case BodyGenTy::Priv:
- if (!Info.CaptureDeviceAddrMap.empty())
- CodeGen(CGF);
- break;
- case BodyGenTy::DupNoPriv:
- if (!Info.CaptureDeviceAddrMap.empty()) {
- CodeGen.setAction(NoPrivAction);
- CodeGen(CGF);
- }
- break;
- case BodyGenTy::NoPriv:
- if (Info.CaptureDeviceAddrMap.empty()) {
- CodeGen.setAction(NoPrivAction);
- CodeGen(CGF);
- }
- break;
+ llvm::OpenMPIRBuilder::TargetDataRTArgs RTArgs;
+ bool EmitDebug = CGF.CGM.getCodeGenOpts().getDebugInfo() !=
+ llvm::codegenoptions::NoDebugInfo;
+ OMPBuilder.emitOffloadingArraysArgument(CGF.Builder, RTArgs, Info,
+ EmitDebug);
+
+ // Emit device ID if any.
+ llvm::Value *DeviceID = nullptr;
+ if (Device) {
+ DeviceID = CGF.Builder.CreateIntCast(CGF.EmitScalarExpr(Device),
+ CGF.Int64Ty, /*isSigned=*/true);
+ } else {
+ DeviceID = CGF.Builder.getInt64(OMP_DEVICEID_UNDEF);
}
- return InsertPointTy(CGF.Builder.GetInsertBlock(),
- CGF.Builder.GetInsertPoint());
+
+ // Emit the number of elements in the offloading arrays.
+ llvm::Value *PointerNum = CGF.Builder.getInt32(Info.NumberOfPtrs);
+ //
+ // Source location for the ident struct
+ llvm::Value *RTLoc = emitUpdateLocation(CGF, D.getBeginLoc());
+
+ llvm::Value *OffloadingArgs[] = {RTLoc,
+ DeviceID,
+ PointerNum,
+ RTArgs.BasePointersArray,
+ RTArgs.PointersArray,
+ RTArgs.SizesArray,
+ RTArgs.MapTypesArray,
+ RTArgs.MapNamesArray,
+ RTArgs.MappersArray};
+ CGF.EmitRuntimeCall(
+ OMPBuilder.getOrCreateRuntimeFunction(
+ CGM.getModule(), OMPRTL___tgt_target_data_begin_mapper),
+ OffloadingArgs);
+
+ // If device pointer privatization is required, emit the body of the region
+ // here. It will have to be duplicated: with and without privatization.
+ if (!Info.CaptureDeviceAddrMap.empty())
+ CodeGen(CGF);
};
- auto DeviceAddrCB = [&](unsigned int I, llvm::Value *BP, llvm::Value *BPVal) {
- if (const ValueDecl *DevVD = CombinedInfo.DevicePtrDecls[I]) {
- ASTContext &Ctx = CGF.getContext();
- Address BPAddr(BP, BPVal->getType(),
- Ctx.getTypeAlignInChars(Ctx.VoidPtrTy));
- Info.CaptureDeviceAddrMap.try_emplace(DevVD, BPAddr);
+ // Generate code for the closing of the data region.
+ auto &&EndThenGen = [this, Device, &Info, &D](CodeGenFunction &CGF,
+ PrePostActionTy &) {
+ assert(Info.isValid() && "Invalid data environment closing arguments.");
+
+ llvm::OpenMPIRBuilder::TargetDataRTArgs RTArgs;
+ bool EmitDebug = CGF.CGM.getCodeGenOpts().getDebugInfo() !=
+ llvm::codegenoptions::NoDebugInfo;
+ OMPBuilder.emitOffloadingArraysArgument(CGF.Builder, RTArgs, Info,
+ EmitDebug,
+ /*ForEndCall=*/true);
+
+ // Emit device ID if any.
+ llvm::Value *DeviceID = nullptr;
+ if (Device) {
+ DeviceID = CGF.Builder.CreateIntCast(CGF.EmitScalarExpr(Device),
+ CGF.Int64Ty, /*isSigned=*/true);
+ } else {
+ DeviceID = CGF.Builder.getInt64(OMP_DEVICEID_UNDEF);
}
+
+ // Emit the number of elements in the offloading arrays.
+ llvm::Value *PointerNum = CGF.Builder.getInt32(Info.NumberOfPtrs);
+
+ // Source location for the ident struct
+ llvm::Value *RTLoc = emitUpdateLocation(CGF, D.getBeginLoc());
+
+ llvm::Value *OffloadingArgs[] = {RTLoc,
+ DeviceID,
+ PointerNum,
+ RTArgs.BasePointersArray,
+ RTArgs.PointersArray,
+ RTArgs.SizesArray,
+ RTArgs.MapTypesArray,
+ RTArgs.MapNamesArray,
+ RTArgs.MappersArray};
+ CGF.EmitRuntimeCall(
+ OMPBuilder.getOrCreateRuntimeFunction(
+ CGM.getModule(), OMPRTL___tgt_target_data_end_mapper),
+ OffloadingArgs);
};
- auto CustomMapperCB = [&](unsigned int I) {
- llvm::Value *MFunc = nullptr;
- if (CombinedInfo.Mappers[I]) {
- Info.HasMapper = true;
- MFunc = CGF.CGM.getOpenMPRuntime().getOrCreateUserDefinedMapperFunc(
- cast<OMPDeclareMapperDecl>(CombinedInfo.Mappers[I]));
+ // If we need device pointer privatization, we need to emit the body of the
+ // region with no privatization in the 'else' branch of the conditional.
+ // Otherwise, we don't have to do anything.
+ auto &&BeginElseGen = [&Info, &CodeGen, &NoPrivAction](CodeGenFunction &CGF,
+ PrePostActionTy &) {
+ if (!Info.CaptureDeviceAddrMap.empty()) {
+ CodeGen.setAction(NoPrivAction);
+ CodeGen(CGF);
}
- return MFunc;
};
- CGF.Builder.restoreIP(OMPBuilder.createTargetData(
- OmpLoc, AllocaIP, CodeGenIP, DeviceID, IfCondVal, Info, GenMapInfoCB,
- /*MapperFunc=*/nullptr, BodyCB, DeviceAddrCB, CustomMapperCB));
+ // We don't have to do anything to close the region if the if clause evaluates
+ // to false.
+ auto &&EndElseGen = [](CodeGenFunction &CGF, PrePostActionTy &) {};
+
+ if (IfCond) {
+ emitIfClause(CGF, IfCond, BeginThenGen, BeginElseGen);
+ } else {
+ RegionCodeGenTy RCG(BeginThenGen);
+ RCG(CGF);
+ }
+
+ // If we don't require privatization of device pointers, we emit the body in
+ // between the runtime calls. This avoids duplicating the body code.
+ if (Info.CaptureDeviceAddrMap.empty()) {
+ CodeGen.setAction(NoPrivAction);
+ CodeGen(CGF);
+ }
+
+ if (IfCond) {
+ emitIfClause(CGF, IfCond, EndThenGen, EndElseGen);
+ } else {
+ RegionCodeGenTy RCG(EndThenGen);
+ RCG(CGF);
+ }
}
void CGOpenMPRuntime::emitTargetDataStandAloneCall(
diff --git a/clang/test/OpenMP/target_data_codegen.cpp b/clang/test/OpenMP/target_data_codegen.cpp
index 1dd2173922f4a6..a29d1ed68ae952 100644
--- a/clang/test/OpenMP/target_data_codegen.cpp
+++ b/clang/test/OpenMP/target_data_codegen.cpp
@@ -63,7 +63,9 @@ void foo(int arg) {
// CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
- // 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: 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]]
// CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P]]
#pragma omp target data if(1+3-5) device(arg) map(from: gc)
@@ -352,11 +354,11 @@ int bar(int arg){
}
// Region 00
-// CK2-DAG: [[DEV:%[^,]+]] = sext i32 [[DEVi32:%[^,]+]] to i64
-// CK2-DAG: [[DEVi32]] = load i32, ptr %{{[^,]+}},
// CK2: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
// CK2: [[IFTHEN]]
-// CK2-DAG: call void @__tgt_target_data_begin_mapper(ptr @{{.+}}, i64 [[DEV]], i32 2, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[GEPS:%[^,]+]], ptr [[MTYPE00]], ptr null, ptr null)
+// CK2-DAG: call void @__tgt_target_data_begin_mapper(ptr @{{.+}}, i64 [[DEV:%[^,]+]], i32 2, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[GEPS:%[^,]+]], ptr [[MTYPE00]], ptr null, ptr null)
+// CK2-DAG: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64
+// CK2-DAG: [[DEVi32]] = load i32, ptr %{{[^,]+}},
// CK2-DAG: [[GEPBP]] = getelementptr inbounds [2 x ptr], ptr [[BP:%[^,]+]]
// CK2-DAG: [[GEPP]] = getelementptr inbounds [2 x ptr], ptr [[P:%[^,]+]]
// CK2-DAG: [[GEPS]] = getelementptr inbounds [2 x i64], ptr [[PS:%[^,]+]]
@@ -386,7 +388,9 @@ int bar(int arg){
// CK2: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
// CK2: [[IFTHEN]]
-// CK2-DAG: call void @__tgt_target_data_end_mapper(ptr @{{.+}}, i64 [[DEV]], i32 2, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[GEPS:%[^,]+]], ptr [[MTYPE00]], ptr null, ptr null)
+// CK2-DAG: call void @__tgt_target_data_end_mapper(ptr @{{.+}}, i64 [[DEV:%[^,]+]], i32 2, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[GEPS:%[^,]+]], ptr [[MTYPE00]], ptr null, ptr null)
+// CK2-DAG: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64
+// CK2-DAG: [[DEVi32]] = load i32, ptr %{{[^,]+}},
// CK2-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
// CK2-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P]]
// CK2-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[PS]]
@@ -463,11 +467,11 @@ int bar(int arg){
}
// Region 00
-// CK4-DAG: [[DEV:%[^,]+]] = sext i32 [[DEVi32:%[^,]+]] to i64
-// CK4-DAG: [[DEVi32]] = load i32, ptr %{{[^,]+}},
// CK4: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
// CK4: [[IFTHEN]]
-// CK4-DAG: call void @__tgt_target_data_begin_mapper(ptr @{{.+}}, i64 [[DEV]], i32 2, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[GEPS:%.+]], ptr [[MTYPE00]], ptr null, ptr null)
+// CK4-DAG: call void @__tgt_target_data_begin_mapper(ptr @{{.+}}, i64 [[DEV:%[^,]+]], i32 2, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[GEPS:%.+]], ptr [[MTYPE00]], ptr null, ptr null)
+// CK4-DAG: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64
+// CK4-DAG: [[DEVi32]] = load i32, ptr %{{[^,]+}},
// CK4-DAG: [[GEPBP]] = getelementptr inbounds [2 x ptr], ptr [[BP:%[^,]+]]
// CK4-DAG: [[GEPP]] = getelementptr inbounds [2 x ptr], ptr [[P:%[^,]+]]
// CK4-DAG: [[GEPS]] = getelementptr inbounds [2 x i64], ptr [[PS:%[^,]+]]
@@ -497,7 +501,9 @@ int bar(int arg){
// CK4: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
// CK4: [[IFTHEN]]
-// CK4-DAG: call void @__tgt_target_data_end_mapper(ptr @{{.+}}, i64 [[DEV]], i32 2, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[GEPS:%.+]], ptr [[MTYPE00]], ptr null, ptr null)
+// CK4-DAG: call void @__tgt_target_data_end_mapper(ptr @{{.+}}, i64 [[DEV:%[^,]+]], i32 2, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[GEPS:%.+]], ptr [[MTYPE00]], ptr null, ptr null)
+// CK4-DAG: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64
+// CK4-DAG: [[DEVi32]] = load i32, ptr %{{[^,]+}},
// CK4-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
// CK4-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P]]
// CK4-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[PS]]
diff --git a/clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp b/clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp
index 0e9dbd39fd6418..745b0edc3460b1 100644
--- a/clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp
+++ b/clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp
@@ -131,6 +131,7 @@ void foo(float *&lr, T *&tr) {
++l;
}
// CK1: [[BEND]]:
+ // CK1: [[CMP:%.+]] = icmp ne ptr %{{.+}}, null
// CK1: br i1 [[CMP]], label %[[BTHEN:.+]], label %[[BELSE:.+]]
// CK1: [[BTHEN]]:
diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
index ed0c923ceaca12..9a9ed016d85a3b 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
+++ b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
@@ -2098,10 +2098,6 @@ class OpenMPIRBuilder {
/// \param Info Stores all information realted to the Target Data directive.
/// \param GenMapInfoCB Callback that populates the MapInfos and returns.
/// \param BodyGenCB Optional Callback to generate the region code.
- /// \param DeviceAddrCB Optional callback to generate code related to
- /// use_device_ptr and use_device_addr.
- /// \param CustomMapperCB Optional callback to generate code related to
- /// custom mappers.
OpenMPIRBuilder::InsertPointTy createTargetData(
const LocationDescription &Loc, InsertPointTy AllocaIP,
InsertPointTy CodeGenIP, Value *DeviceID, Value *IfCond,
@@ -2110,9 +2106,7 @@ class OpenMPIRBuilder {
omp::RuntimeFunction *MapperFunc = nullptr,
function_ref<InsertPointTy(InsertPointTy CodeGenIP,
BodyGenTy BodyGenType)>
- BodyGenCB = nullptr,
- function_ref<void(unsigned int, Value *, Value *)> DeviceAddrCB = nullptr,
- function_ref<Value *(unsigned int)> CustomMapperCB = nullptr);
+ BodyGenCB = nullptr);
using TargetBodyGenCallbackTy = function_ref<InsertPointTy(
InsertPointTy AllocaIP, InsertPointTy CodeGenIP)>;
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index 8c3ff591af1e35..c3eefde2fa2c84 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -4174,9 +4174,7 @@ OpenMPIRBuilder::InsertPointTy OpenMPIRBuilder::createTargetData(
function_ref<MapInfosTy &(InsertPointTy CodeGenIP)> GenMapInfoCB,
omp::RuntimeFunction *MapperFunc,
function_ref<InsertPointTy(InsertPointTy CodeGenIP, BodyGenTy BodyGenType)>
- BodyGenCB,
- function_ref<void(unsigned int, Value *, Value *)> DeviceAddrCB,
- function_ref<Value *(unsigned int)> CustomMapperCB) {
+ BodyGenCB) {
if (!updateToLocation(Loc))
return InsertPointTy();
@@ -4187,9 +4185,9 @@ OpenMPIRBuilder::InsertPointTy OpenMPIRBuilder::createTargetData(
// arguments of the runtime call by reference because they are used in the
// closing of the region.
auto BeginThenGen = [&](InsertPointTy AllocaIP, InsertPointTy CodeGenIP) {
- emitOffloadingArrays(
- AllocaIP, Builder.saveIP(), GenMapInfoCB(Builder.saveIP()), Info,
- /*IsNonContiguous=*/true, DeviceAddrCB, CustomMapperCB);
+ emitOffloadingArrays(AllocaIP, Builder.saveIP(),
+ GenMapInfoCB(Builder.saveIP()), Info,
+ /*IsNonContiguous=*/true);
TargetDataRTArgs RTArgs;
emitOffloadingArraysArgument(Builder, RTArgs, Info);
More information about the cfe-commits
mailing list