From 0d8d718171192301f2beb10bd08ce62e70281a5e Mon Sep 17 00:00:00 2001 From: Akash Banerjee Date: Tue, 27 Jun 2023 13:50:17 +0100 Subject: [PATCH] [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. Depends on D146557 Differential Revision: https://reviews.llvm.org/D150860 --- 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, 97 insertions(+), 142 deletions(-) diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index b79539a..619f562 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) { @@ -10377,140 +10377,94 @@ 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); - } - - // 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); - } + 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])); + } + return MFunc; + }; - 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)); } void CGOpenMPRuntime::emitTargetDataStandAloneCall( diff --git a/clang/test/OpenMP/target_data_codegen.cpp b/clang/test/OpenMP/target_data_codegen.cpp index a29d1ed..1dd2173 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 745b0ed..0e9dbd3 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 7657ad3..ac16e82 100644 --- a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h +++ b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h @@ -2032,6 +2032,10 @@ 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, @@ -2040,7 +2044,9 @@ public: omp::RuntimeFunction *MapperFunc = nullptr, function_ref - BodyGenCB = nullptr); + BodyGenCB = nullptr, + function_ref DeviceAddrCB = nullptr, + function_ref CustomMapperCB = nullptr); using TargetBodyGenCallbackTy = function_ref; diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp index e3aa391..90d831c3 100644 --- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp +++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp @@ -4084,7 +4084,9 @@ OpenMPIRBuilder::InsertPointTy OpenMPIRBuilder::createTargetData( function_ref GenMapInfoCB, omp::RuntimeFunction *MapperFunc, function_ref - BodyGenCB) { + BodyGenCB, + function_ref DeviceAddrCB, + function_ref CustomMapperCB) { if (!updateToLocation(Loc)) return InsertPointTy(); @@ -4095,9 +4097,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); + emitOffloadingArrays( + AllocaIP, Builder.saveIP(), GenMapInfoCB(Builder.saveIP()), Info, + /*IsNonContiguous=*/true, DeviceAddrCB, CustomMapperCB); TargetDataRTArgs RTArgs; emitOffloadingArraysArgument(Builder, RTArgs, Info); -- 2.7.4