diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index 188b21e3a8894..180f3265ab193 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(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 a29d1ed68ae95..1dd2173922f4a 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 745b0edc3460b..0e9dbd39fd641 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 9a9ed016d85a3..722ccd9048e5d 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 - BodyGenCB = nullptr); + BodyGenCB = nullptr, + function_ref DeviceAddrCB = nullptr, + function_ref CustomMapperCB = nullptr, + Value *SrcLocInfo = nullptr); using TargetBodyGenCallbackTy = function_ref; diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp index c3eefde2fa2c8..af5d03a2484c4 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 GenMapInfoCB, omp::RuntimeFunction *MapperFunc, function_ref - BodyGenCB) { + BodyGenCB, + function_ref DeviceAddrCB, + function_ref 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 d39741fd1160a..1573f30d5b391 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>) { // 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>, %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>, %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]]