[clang] 4a49226 - [OpenMP] Change clang emitTargetDataCalls to use OMPIRBuilder
Akash Banerjee via cfe-commits
cfe-commits at lists.llvm.org
Wed Jul 5 06:04:04 PDT 2023
Author: Akash Banerjee
Date: 2023-07-05T14:03:54+01:00
New Revision: 4a49226537abafc9b372ba8ef61d14a817a5226a
URL: https://github.com/llvm/llvm-project/commit/4a49226537abafc9b372ba8ef61d14a817a5226a
DIFF: https://github.com/llvm/llvm-project/commit/4a49226537abafc9b372ba8ef61d14a817a5226a.diff
LOG: [OpenMP] Change clang emitTargetDataCalls to use OMPIRBuilder
This patch changes the emitTargetDataCalls function in clang to make use of the OpenMPIRBuilder::createTargetData function for Target Data directive code gen.
Reapplying commit 0d8d718171192301f2beb10bd08ce62e70281a5e after fixing libomptarget test failure related to debug info.
Depends on D146557
Reviewed By: jdoerfert
Differential Revision: https://reviews.llvm.org/D150860
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
mlir/test/Target/LLVMIR/omptarget-llvm.mlir
Removed:
################################################################################
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 188b21e3a88949..180f3265ab1938 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -9026,14 +9026,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) {
@@ -10379,140 +10379,97 @@ void CGOpenMPRuntime::emitTargetDataCalls(
// off.
PrePostActionTy NoPrivAction;
- // 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;
+ 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);
+ }
+ // 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);
- // Fill up the arrays and create the arguments.
- emitOffloadingArrays(CGF, CombinedInfo, Info, OMPBuilder,
- /*IsNonContiguous=*/true);
-
- 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);
+ 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);
}
- // 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);
+ return CombinedInfo;
};
-
- // 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);
+ 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;
}
-
- // 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);
+ return InsertPointTy(CGF.Builder.GetInsertBlock(),
+ CGF.Builder.GetInsertPoint());
};
- // 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);
+ 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);
}
};
- // 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);
- }
+ 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]));
+ }
+ return MFunc;
+ };
- // 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);
- }
+ // Source location for the ident struct
+ llvm::Value *RTLoc = emitUpdateLocation(CGF, D.getBeginLoc());
- if (IfCond) {
- emitIfClause(CGF, IfCond, EndThenGen, EndElseGen);
- } else {
- RegionCodeGenTy RCG(EndThenGen);
- RCG(CGF);
- }
+ CGF.Builder.restoreIP(OMPBuilder.createTargetData(
+ OmpLoc, AllocaIP, CodeGenIP, DeviceID, IfCondVal, Info, GenMapInfoCB,
+ /*MapperFunc=*/nullptr, BodyCB, DeviceAddrCB, CustomMapperCB, RTLoc));
}
void CGOpenMPRuntime::emitTargetDataStandAloneCall(
diff --git a/clang/test/OpenMP/target_data_codegen.cpp b/clang/test/OpenMP/target_data_codegen.cpp
index a29d1ed68ae952..1dd2173922f4a6 100644
--- a/clang/test/OpenMP/target_data_codegen.cpp
+++ b/clang/test/OpenMP/target_data_codegen.cpp
@@ -63,9 +63,7 @@ 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: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64
- // CK1-DAG: [[DEVi32]] = load i32, ptr %{{[^,]+}},
+ // 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: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
// CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P]]
#pragma omp target data if(1+3-5) device(arg) map(from: gc)
@@ -354,11 +352,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: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64
-// CK2-DAG: [[DEVi32]] = load i32, ptr %{{[^,]+}},
+// 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: [[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:%[^,]+]]
@@ -388,9 +386,7 @@ 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: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64
-// CK2-DAG: [[DEVi32]] = load i32, ptr %{{[^,]+}},
+// 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: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
// CK2-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P]]
// CK2-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[PS]]
@@ -467,11 +463,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: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64
-// CK4-DAG: [[DEVi32]] = load i32, ptr %{{[^,]+}},
+// 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: [[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:%[^,]+]]
@@ -501,9 +497,7 @@ 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: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64
-// CK4-DAG: [[DEVi32]] = load i32, ptr %{{[^,]+}},
+// 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: [[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 745b0edc3460b1..0e9dbd39fd6418 100644
--- a/clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp
+++ b/clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp
@@ -131,7 +131,6 @@ 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 9a9ed016d85a3b..722ccd9048e5da 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
+++ b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
@@ -2098,6 +2098,10 @@ 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,
@@ -2106,7 +2110,10 @@ class OpenMPIRBuilder {
omp::RuntimeFunction *MapperFunc = nullptr,
function_ref<InsertPointTy(InsertPointTy CodeGenIP,
BodyGenTy BodyGenType)>
- BodyGenCB = nullptr);
+ BodyGenCB = nullptr,
+ function_ref<void(unsigned int, Value *, Value *)> DeviceAddrCB = nullptr,
+ function_ref<Value *(unsigned int)> CustomMapperCB = nullptr,
+ Value *SrcLocInfo = 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 c3eefde2fa2c84..af5d03a2484c46 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -4174,31 +4174,37 @@ OpenMPIRBuilder::InsertPointTy OpenMPIRBuilder::createTargetData(
function_ref<MapInfosTy &(InsertPointTy CodeGenIP)> GenMapInfoCB,
omp::RuntimeFunction *MapperFunc,
function_ref<InsertPointTy(InsertPointTy CodeGenIP, BodyGenTy BodyGenType)>
- BodyGenCB) {
+ BodyGenCB,
+ function_ref<void(unsigned int, Value *, Value *)> DeviceAddrCB,
+ function_ref<Value *(unsigned int)> CustomMapperCB, Value *SrcLocInfo) {
if (!updateToLocation(Loc))
return InsertPointTy();
Builder.restoreIP(CodeGenIP);
bool IsStandAlone = !BodyGenCB;
-
+ MapInfosTy *MapInfo;
// 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 = [&](InsertPointTy AllocaIP, InsertPointTy CodeGenIP) {
- emitOffloadingArrays(AllocaIP, Builder.saveIP(),
- GenMapInfoCB(Builder.saveIP()), Info,
- /*IsNonContiguous=*/true);
+ MapInfo = &GenMapInfoCB(Builder.saveIP());
+ emitOffloadingArrays(AllocaIP, Builder.saveIP(), *MapInfo, Info,
+ /*IsNonContiguous=*/true, DeviceAddrCB,
+ CustomMapperCB);
TargetDataRTArgs RTArgs;
- emitOffloadingArraysArgument(Builder, RTArgs, Info);
+ emitOffloadingArraysArgument(Builder, RTArgs, Info,
+ !MapInfo->Names.empty());
// Emit the number of elements in the offloading arrays.
Value *PointerNum = Builder.getInt32(Info.NumberOfPtrs);
// Source location for the ident struct
- uint32_t SrcLocStrSize;
- Constant *SrcLocStr = getOrCreateSrcLocStr(Loc, SrcLocStrSize);
- Value *SrcLocInfo = getOrCreateIdent(SrcLocStr, SrcLocStrSize);
+ if (!SrcLocInfo) {
+ uint32_t SrcLocStrSize;
+ Constant *SrcLocStr = getOrCreateSrcLocStr(Loc, SrcLocStrSize);
+ SrcLocInfo = getOrCreateIdent(SrcLocStr, SrcLocStrSize);
+ }
Value *OffloadingArgs[] = {SrcLocInfo, DeviceID,
PointerNum, RTArgs.BasePointersArray,
@@ -4233,16 +4239,18 @@ OpenMPIRBuilder::InsertPointTy OpenMPIRBuilder::createTargetData(
// Generate code for the closing of the data region.
auto EndThenGen = [&](InsertPointTy AllocaIP, InsertPointTy CodeGenIP) {
TargetDataRTArgs RTArgs;
- emitOffloadingArraysArgument(Builder, RTArgs, Info, /*EmitDebug=*/false,
+ emitOffloadingArraysArgument(Builder, RTArgs, Info, !MapInfo->Names.empty(),
/*ForEndCall=*/true);
// Emit the number of elements in the offloading arrays.
Value *PointerNum = Builder.getInt32(Info.NumberOfPtrs);
// Source location for the ident struct
- uint32_t SrcLocStrSize;
- Constant *SrcLocStr = getOrCreateSrcLocStr(Loc, SrcLocStrSize);
- Value *SrcLocInfo = getOrCreateIdent(SrcLocStr, SrcLocStrSize);
+ if (!SrcLocInfo) {
+ uint32_t SrcLocStrSize;
+ Constant *SrcLocStr = getOrCreateSrcLocStr(Loc, SrcLocStrSize);
+ SrcLocInfo = getOrCreateIdent(SrcLocStr, SrcLocStrSize);
+ }
Value *OffloadingArgs[] = {SrcLocInfo, DeviceID,
PointerNum, RTArgs.BasePointersArray,
@@ -4731,6 +4739,9 @@ void OpenMPIRBuilder::emitOffloadingArrays(
auto *MapNamesArrayGbl =
createOffloadMapnames(CombinedInfo.Names, MapnamesName);
Info.RTArgs.MapNamesArray = MapNamesArrayGbl;
+ } else {
+ Info.RTArgs.MapNamesArray = Constant::getNullValue(
+ Type::getInt8Ty(Builder.getContext())->getPointerTo());
}
// If there's a present map type modifier, it must not be applied to the end
diff --git a/mlir/test/Target/LLVMIR/omptarget-llvm.mlir b/mlir/test/Target/LLVMIR/omptarget-llvm.mlir
index d39741fd1160ae..1573f30d5b391b 100644
--- a/mlir/test/Target/LLVMIR/omptarget-llvm.mlir
+++ b/mlir/test/Target/LLVMIR/omptarget-llvm.mlir
@@ -28,11 +28,11 @@ llvm.func @_QPopenmp_target_data() {
// CHECK: store ptr null, ptr %[[VAL_8]], align 8
// CHECK: %[[VAL_9:.*]] = getelementptr inbounds [1 x ptr], ptr %[[VAL_0]], i32 0, i32 0
// CHECK: %[[VAL_10:.*]] = getelementptr inbounds [1 x ptr], ptr %[[VAL_1]], i32 0, i32 0
-// CHECK: call void @__tgt_target_data_begin_mapper(ptr @2, i64 -1, i32 1, ptr %[[VAL_9]], ptr %[[VAL_10]], ptr @.offload_sizes, ptr @.offload_maptypes, ptr null, ptr null)
+// CHECK: call void @__tgt_target_data_begin_mapper(ptr @2, i64 -1, i32 1, ptr %[[VAL_9]], ptr %[[VAL_10]], ptr @.offload_sizes, ptr @.offload_maptypes, ptr @.offload_mapnames, ptr null)
// CHECK: store i32 99, ptr %[[VAL_3]], align 4
// CHECK: %[[VAL_11:.*]] = getelementptr inbounds [1 x ptr], ptr %[[VAL_0]], i32 0, i32 0
// CHECK: %[[VAL_12:.*]] = getelementptr inbounds [1 x ptr], ptr %[[VAL_1]], i32 0, i32 0
-// CHECK: call void @__tgt_target_data_end_mapper(ptr @2, i64 -1, i32 1, ptr %[[VAL_11]], ptr %[[VAL_12]], ptr @.offload_sizes, ptr @.offload_maptypes, ptr null, ptr null)
+// CHECK: call void @__tgt_target_data_end_mapper(ptr @2, i64 -1, i32 1, ptr %[[VAL_11]], ptr %[[VAL_12]], ptr @.offload_sizes, ptr @.offload_maptypes, ptr @.offload_mapnames, ptr null)
// CHECK: ret void
// -----
@@ -67,12 +67,12 @@ llvm.func @_QPopenmp_target_data_region(%1 : !llvm.ptr<array<1024 x i32>>) {
// CHECK: store ptr null, ptr %[[VAL_8]], align 8
// CHECK: %[[VAL_9:.*]] = getelementptr inbounds [1 x ptr], ptr %[[VAL_0]], i32 0, i32 0
// CHECK: %[[VAL_10:.*]] = getelementptr inbounds [1 x ptr], ptr %[[VAL_1]], i32 0, i32 0
-// CHECK: call void @__tgt_target_data_begin_mapper(ptr @2, i64 -1, i32 1, ptr %[[VAL_9]], ptr %[[VAL_10]], ptr @.offload_sizes, ptr @.offload_maptypes, ptr null, ptr null)
+// CHECK: call void @__tgt_target_data_begin_mapper(ptr @2, i64 -1, i32 1, ptr %[[VAL_9]], ptr %[[VAL_10]], ptr @.offload_sizes, ptr @.offload_maptypes, ptr @.offload_mapnames, ptr null)
// CHECK: %[[VAL_11:.*]] = getelementptr [1024 x i32], ptr %[[VAL_6]], i32 0, i64 0
// CHECK: store i32 99, ptr %[[VAL_11]], align 4
// CHECK: %[[VAL_12:.*]] = getelementptr inbounds [1 x ptr], ptr %[[VAL_0]], i32 0, i32 0
// CHECK: %[[VAL_13:.*]] = getelementptr inbounds [1 x ptr], ptr %[[VAL_1]], i32 0, i32 0
-// CHECK: call void @__tgt_target_data_end_mapper(ptr @2, i64 -1, i32 1, ptr %[[VAL_12]], ptr %[[VAL_13]], ptr @.offload_sizes, ptr @.offload_maptypes, ptr null, ptr null)
+// CHECK: call void @__tgt_target_data_end_mapper(ptr @2, i64 -1, i32 1, ptr %[[VAL_12]], ptr %[[VAL_13]], ptr @.offload_sizes, ptr @.offload_maptypes, ptr @.offload_mapnames, ptr null)
// CHECK: ret void
// -----
@@ -136,7 +136,7 @@ llvm.func @_QPomp_target_enter_exit(%1 : !llvm.ptr<array<1024 x i32>>, %3 : !llv
// CHECK: store ptr null, ptr %[[VAL_22]], align 8
// CHECK: %[[VAL_23:.*]] = getelementptr inbounds [2 x ptr], ptr %[[VAL_3]], i32 0, i32 0
// CHECK: %[[VAL_24:.*]] = getelementptr inbounds [2 x ptr], ptr %[[VAL_4]], i32 0, i32 0
-// CHECK: call void @__tgt_target_data_begin_mapper(ptr @3, i64 -1, i32 2, ptr %[[VAL_23]], ptr %[[VAL_24]], ptr @.offload_sizes, ptr @.offload_maptypes, ptr null, ptr null)
+// CHECK: call void @__tgt_target_data_begin_mapper(ptr @3, i64 -1, i32 2, ptr %[[VAL_23]], ptr %[[VAL_24]], ptr @.offload_sizes, ptr @.offload_maptypes, ptr @.offload_mapnames, ptr null)
// CHECK: br label %[[VAL_25:.*]]
// CHECK: omp_if.else: ; preds = %[[VAL_11]]
// CHECK: br label %[[VAL_25]]
@@ -160,7 +160,7 @@ llvm.func @_QPomp_target_enter_exit(%1 : !llvm.ptr<array<1024 x i32>>, %3 : !llv
// CHECK: store ptr null, ptr %[[VAL_36]], align 8
// CHECK: %[[VAL_37:.*]] = getelementptr inbounds [2 x ptr], ptr %[[VAL_0]], i32 0, i32 0
// CHECK: %[[VAL_38:.*]] = getelementptr inbounds [2 x ptr], ptr %[[VAL_1]], i32 0, i32 0
-// CHECK: call void @__tgt_target_data_end_mapper(ptr @3, i64 -1, i32 2, ptr %[[VAL_37]], ptr %[[VAL_38]], ptr @.offload_sizes.1, ptr @.offload_maptypes.2, ptr null, ptr null)
+// CHECK: call void @__tgt_target_data_end_mapper(ptr @3, i64 -1, i32 2, ptr %[[VAL_37]], ptr %[[VAL_38]], ptr @.offload_sizes.1, ptr @.offload_maptypes.2, ptr @.offload_mapnames.3, ptr null)
// CHECK: br label %[[VAL_39:.*]]
// CHECK: omp_if.else5: ; preds = %[[VAL_25]]
// CHECK: br label %[[VAL_39]]
More information about the cfe-commits
mailing list