From a5ea6760674762cb597cf328dc467f1296633da0 Mon Sep 17 00:00:00 2001 From: Akash Banerjee Date: Fri, 30 Jun 2023 16:03:57 +0100 Subject: [PATCH] Reverting commit 0d8d718171192301f2beb10bd08ce62e70281a5e as it broke libomptarget tests --- clang/lib/CodeGen/CGOpenMPRuntime.cpp | 196 +++++++++++++-------- clang/test/OpenMP/target_data_codegen.cpp | 24 ++- .../OpenMP/target_data_use_device_ptr_codegen.cpp | 1 + llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h | 8 +- llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 10 +- 5 files changed, 142 insertions(+), 97 deletions(-) diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index dfc8f71..bd0169d 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(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 1dd2173..a29d1ed 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 0e9dbd3..745b0ed 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 ed0c923..9a9ed01 100644 --- a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h +++ b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h @@ -2098,10 +2098,6 @@ public: /// \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 @@ public: omp::RuntimeFunction *MapperFunc = nullptr, function_ref - BodyGenCB = nullptr, - function_ref DeviceAddrCB = nullptr, - function_ref CustomMapperCB = nullptr); + BodyGenCB = nullptr); using TargetBodyGenCallbackTy = function_ref; diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp index 8c3ff59..c3eefde 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 GenMapInfoCB, omp::RuntimeFunction *MapperFunc, function_ref - BodyGenCB, - function_ref DeviceAddrCB, - function_ref 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); -- 2.7.4