Index: lib/CodeGen/CGOpenMPRuntime.cpp =================================================================== --- lib/CodeGen/CGOpenMPRuntime.cpp +++ lib/CodeGen/CGOpenMPRuntime.cpp @@ -4876,6 +4876,9 @@ /// map/privatization results in multiple arguments passed to the runtime /// library. OMP_MAP_FIRST_REF = 0x20, + /// \brief This flag signals that the reference being passed is a pointer to + /// private data. + OMP_MAP_PRIVATE_PTR = 0x80, /// \brief Pass the element to the device by value. OMP_MAP_PRIVATE_VAL = 0x100, }; @@ -4890,6 +4893,9 @@ /// \brief Function the directive is being generated for. CodeGenFunction &CGF; + /// \brief Set of all first private variables in the current directive. + llvm::SmallPtrSet FirstPrivateDecls; + llvm::Value *getExprTypeSize(const Expr *E) const { auto ExprTy = E->getType().getCanonicalType(); @@ -5242,9 +5248,33 @@ } } + /// \brief Return the adjusted map modifiers if the declaration a capture + /// refers to appears in a first-private clause. This is expected to be used + /// only with directives that start with 'target'. + unsigned adjustMapModifiersForPrivateClauses(const CapturedStmt::Capture &Cap, + unsigned CurrentModifiers) { + assert(Cap.capturesVariable() && "Expected capture by reference only!"); + + // A first private variable captured by reference will use only the + // 'private ptr' and 'map to' flag. Return the right flags if the captured + // declaration is known as first-private in this handler. + if (FirstPrivateDecls.count(Cap.getCapturedVar())) + return MappableExprsHandler::OMP_MAP_PRIVATE_PTR | + MappableExprsHandler::OMP_MAP_TO; + + // We didn't modify anything. + return CurrentModifiers; + } + public: MappableExprsHandler(const OMPExecutableDirective &Dir, CodeGenFunction &CGF) - : Directive(Dir), CGF(CGF) {} + : Directive(Dir), CGF(CGF) { + // Extract firstprivate clause information. + for (const auto *C : Dir.getClausesOfKind()) + for (const auto *D : C->varlists()) + FirstPrivateDecls.insert( + cast(cast(D)->getDecl())->getCanonicalDecl()); + } /// \brief Generate all the base pointers, section pointers, sizes and map /// types for the extracted mappable expressions. @@ -5326,6 +5356,86 @@ return; } + + /// \brief Generate the default map information for a given capture \a CI, + /// record field declaration \a RI and captured value \a CV. + void generateDefaultMapInfo( + const CapturedStmt::Capture &CI, const FieldDecl &RI, llvm::Value *CV, + MappableExprsHandler::MapValuesArrayTy &CurBasePointers, + MappableExprsHandler::MapValuesArrayTy &CurPointers, + MappableExprsHandler::MapValuesArrayTy &CurSizes, + MappableExprsHandler::MapFlagsArrayTy &CurMapTypes) { + auto &Ctx = CGF.getContext(); + + // Do the default mapping. + if (CI.capturesThis()) { + CurBasePointers.push_back(CV); + CurPointers.push_back(CV); + const PointerType *PtrTy = cast(RI.getType().getTypePtr()); + CurSizes.push_back(CGF.getTypeSize(PtrTy->getPointeeType())); + // Default map type. + CurMapTypes.push_back(MappableExprsHandler::OMP_MAP_TO | + MappableExprsHandler::OMP_MAP_FROM); + } else if (CI.capturesVariableByCopy()) { + if (!RI.getType()->isAnyPointerType()) { + // If the field is not a pointer, we need to save the actual value + // and load it as a void pointer. + CurMapTypes.push_back(MappableExprsHandler::OMP_MAP_PRIVATE_VAL); + auto DstAddr = CGF.CreateMemTemp(Ctx.getUIntPtrType(), + Twine(CI.getCapturedVar()->getName()) + + ".casted"); + LValue DstLV = CGF.MakeAddrLValue(DstAddr, Ctx.getUIntPtrType()); + + auto *SrcAddrVal = CGF.EmitScalarConversion( + DstAddr.getPointer(), Ctx.getPointerType(Ctx.getUIntPtrType()), + Ctx.getPointerType(RI.getType()), SourceLocation()); + LValue SrcLV = CGF.MakeNaturalAlignAddrLValue(SrcAddrVal, RI.getType()); + + // Store the value using the source type pointer. + CGF.EmitStoreThroughLValue(RValue::get(CV), SrcLV); + + // Load the value using the destination type pointer. + CurBasePointers.push_back( + CGF.EmitLoadOfLValue(DstLV, SourceLocation()).getScalarVal()); + CurPointers.push_back(CurBasePointers.back()); + + // Get the size of the type to be used in the map. + CurSizes.push_back(CGF.getTypeSize(RI.getType())); + } else { + // Pointers are implicitly mapped with a zero size and no flags + // (other than first map that is added for all implicit maps). + CurMapTypes.push_back(0u); + CurBasePointers.push_back(CV); + CurPointers.push_back(CV); + CurSizes.push_back(llvm::Constant::getNullValue(CGF.SizeTy)); + } + } else { + assert(CI.capturesVariable() && "Expected captured reference."); + CurBasePointers.push_back(CV); + CurPointers.push_back(CV); + + const ReferenceType *PtrTy = + cast(RI.getType().getTypePtr()); + QualType ElementType = PtrTy->getPointeeType(); + CurSizes.push_back(CGF.getTypeSize(ElementType)); + // The default map type for a scalar/complex type is 'to' because by + // default the value doesn't have to be retrieved. For an aggregate + // type, the default is 'tofrom'. + CurMapTypes.push_back(ElementType->isAggregateType() + ? (MappableExprsHandler::OMP_MAP_TO | + MappableExprsHandler::OMP_MAP_FROM) + : MappableExprsHandler::OMP_MAP_TO); + + // If we have a capture by reference we may need to add the private + // pointer flag if the base declaration shows in some first-private + // clause. + CurMapTypes.back() = + adjustMapModifiersForPrivateClauses(CI, CurMapTypes.back()); + } + // Every default map produces a single argument, so, it is always the + // first one. + CurMapTypes.back() |= MappableExprsHandler::OMP_MAP_FIRST_REF; + } }; enum OpenMPOffloadingReservedDeviceIDs { @@ -5508,8 +5618,8 @@ MappableExprsHandler::MapValuesArrayTy CurSizes; MappableExprsHandler::MapFlagsArrayTy CurMapTypes; - // Get map clause information. - MappableExprsHandler MCHandler(D, CGF); + // Get mappable expression information. + MappableExprsHandler MEHandler(D, CGF); const CapturedStmt &CS = *cast(D.getAssociatedStmt()); auto RI = CS.getCapturedRecordDecl()->field_begin(); @@ -5537,75 +5647,11 @@ } else { // If we have any information in the map clause, we use it, otherwise we // just do a default mapping. - MCHandler.generateInfoForCapture(CI, CurBasePointers, CurPointers, + MEHandler.generateInfoForCapture(CI, CurBasePointers, CurPointers, CurSizes, CurMapTypes); - - if (CurBasePointers.empty()) { - // Do the default mapping. - if (CI->capturesThis()) { - CurBasePointers.push_back(*CV); - CurPointers.push_back(*CV); - const PointerType *PtrTy = - cast(RI->getType().getTypePtr()); - CurSizes.push_back(CGF.getTypeSize(PtrTy->getPointeeType())); - // Default map type. - CurMapTypes.push_back(MappableExprsHandler::OMP_MAP_TO | - MappableExprsHandler::OMP_MAP_FROM); - } else if (CI->capturesVariableByCopy()) { - if (!RI->getType()->isAnyPointerType()) { - // If the field is not a pointer, we need to save the actual value - // and load it as a void pointer. - CurMapTypes.push_back(MappableExprsHandler::OMP_MAP_PRIVATE_VAL); - auto DstAddr = CGF.CreateMemTemp( - Ctx.getUIntPtrType(), - Twine(CI->getCapturedVar()->getName()) + ".casted"); - LValue DstLV = CGF.MakeAddrLValue(DstAddr, Ctx.getUIntPtrType()); - - auto *SrcAddrVal = CGF.EmitScalarConversion( - DstAddr.getPointer(), Ctx.getPointerType(Ctx.getUIntPtrType()), - Ctx.getPointerType(RI->getType()), SourceLocation()); - LValue SrcLV = - CGF.MakeNaturalAlignAddrLValue(SrcAddrVal, RI->getType()); - - // Store the value using the source type pointer. - CGF.EmitStoreThroughLValue(RValue::get(*CV), SrcLV); - - // Load the value using the destination type pointer. - CurBasePointers.push_back( - CGF.EmitLoadOfLValue(DstLV, SourceLocation()).getScalarVal()); - CurPointers.push_back(CurBasePointers.back()); - - // Get the size of the type to be used in the map. - CurSizes.push_back(CGF.getTypeSize(RI->getType())); - } else { - // Pointers are implicitly mapped with a zero size and no flags - // (other than first map that is added for all implicit maps). - CurMapTypes.push_back(0u); - CurBasePointers.push_back(*CV); - CurPointers.push_back(*CV); - CurSizes.push_back(llvm::Constant::getNullValue(CGM.SizeTy)); - } - } else { - assert(CI->capturesVariable() && "Expected captured reference."); - CurBasePointers.push_back(*CV); - CurPointers.push_back(*CV); - - const ReferenceType *PtrTy = - cast(RI->getType().getTypePtr()); - QualType ElementType = PtrTy->getPointeeType(); - CurSizes.push_back(CGF.getTypeSize(ElementType)); - // The default map type for a scalar/complex type is 'to' because by - // default the value doesn't have to be retrieved. For an aggregate - // type, the default is 'tofrom'. - CurMapTypes.push_back(ElementType->isAggregateType() - ? (MappableExprsHandler::OMP_MAP_TO | - MappableExprsHandler::OMP_MAP_FROM) - : MappableExprsHandler::OMP_MAP_TO); - } - // Every default map produces a single argument, so, it is always the - // first one. - CurMapTypes.back() |= MappableExprsHandler::OMP_MAP_FIRST_REF; - } + if (CurBasePointers.empty()) + MEHandler.generateDefaultMapInfo(*CI, **RI, *CV, CurBasePointers, + CurPointers, CurSizes, CurMapTypes); } // We expect to have at least an element of information for this capture. assert(!CurBasePointers.empty() && "Non-existing map pointer for capture!"); Index: test/OpenMP/target_firstprivate_codegen.cpp =================================================================== --- test/OpenMP/target_firstprivate_codegen.cpp +++ test/OpenMP/target_firstprivate_codegen.cpp @@ -34,14 +34,14 @@ // CHECK-DAG: [[SIZET:@.+]] = private unnamed_addr constant [1 x i{{32|64}}] [i[[SZ:32|64]] 4] // CHECK: [[MAPT:@.+]] = private unnamed_addr constant [1 x i32] [i32 288] -// CHECK-DAG: [[MAPT2:@.+]] = private unnamed_addr constant [9 x i32] [i32 288, i32 35, i32 288, i32 35, i32 35, i32 288, i32 288, i32 35, i32 35] +// CHECK-DAG: [[MAPT2:@.+]] = private unnamed_addr constant [9 x i32] [i32 288, i32 161, i32 288, i32 161, i32 161, i32 288, i32 288, i32 161, i32 161] // CHECK-DAG: [[SIZET3:@.+]] = private unnamed_addr constant [1 x i{{32|64}}] zeroinitializer // CHECK-DAG: [[MAPT3:@.+]] = private unnamed_addr constant [1 x i32] [i32 32] -// CHECK-DAG: [[MAPT4:@.+]] = private unnamed_addr constant [5 x i32] [i32 35, i32 288, i32 288, i32 288, i32 35] +// CHECK-DAG: [[MAPT4:@.+]] = private unnamed_addr constant [5 x i32] [i32 35, i32 288, i32 288, i32 288, i32 161] // CHECK-DAG: [[SIZET5:@.+]] = private unnamed_addr constant [3 x i{{32|64}}] [i[[SZ]] 4, i[[SZ]] 1, i[[SZ]] 40] -// CHECK-DAG: [[MAPT5:@.+]] = private unnamed_addr constant [3 x i32] [i32 288, i32 288, i32 35] +// CHECK-DAG: [[MAPT5:@.+]] = private unnamed_addr constant [3 x i32] [i32 288, i32 288, i32 161] // CHECK-DAG: [[SIZET6:@.+]] = private unnamed_addr constant [2 x i{{32|64}}] [i[[SZ]] 4, i[[SZ]] 40] -// CHECK-DAG: [[MAPT6:@.+]] = private unnamed_addr constant [2 x i32] [i32 288, i32 35] +// CHECK-DAG: [[MAPT6:@.+]] = private unnamed_addr constant [2 x i32] [i32 288, i32 161] // CHECK: define {{.*}}[[FOO:@.+]]( Index: test/OpenMP/target_map_codegen.cpp =================================================================== --- test/OpenMP/target_map_codegen.cpp +++ test/OpenMP/target_map_codegen.cpp @@ -4281,8 +4281,17 @@ // CK27: [[SIZE03:@.+]] = private {{.*}}constant [1 x i[[Z]]] zeroinitializer // CK27: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 35] -// CK27-LABEL: zero_size_section_maps -void zero_size_section_maps (int ii){ +// CK27: [[SIZE05:@.+]] = private {{.*}}constant [1 x i[[Z]]] zeroinitializer +// CK27: [[MTYPE05:@.+]] = private {{.*}}constant [1 x i32] [i32 32] + +// CK27: [[SIZE07:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4] +// CK27: [[MTYPE07:@.+]] = private {{.*}}constant [1 x i32] [i32 288] + +// CK27: [[SIZE09:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 40] +// CK27: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i32] [i32 161] + +// CK27-LABEL: zero_size_section_and_private_maps +void zero_size_section_and_private_maps (int ii){ // Map of a pointer. int *pa; @@ -4367,12 +4376,99 @@ { pa[50]++; } + + int *pvtPtr; + int pvtScl; + int pvtArr[10]; + + // Region 04 + // CK27: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 0, i8** null, i8** null, i{{64|32}}* null, i32* null) + // CK27: call void [[CALL04:@.+]]() + #pragma omp target private(pvtPtr) + { + pvtPtr[5]++; + } + + // Region 05 + // CK27-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE05]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE05]]{{.+}}) + // CK27-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK27-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK27-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK27-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK27-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK27-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK27-DAG: [[CBPVAL0]] = bitcast i32* [[VAR0:%.+]] to i8* + // CK27-DAG: [[CPVAL0]] = bitcast i32* [[VAR0]] to i8* + + // CK27: call void [[CALL05:@.+]](i32* {{[^,]+}}) + #pragma omp target firstprivate(pvtPtr) + { + pvtPtr[5]++; + } + + // Region 06 + // CK27: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 0, i8** null, i8** null, i{{64|32}}* null, i32* null) + // CK27: call void [[CALL06:@.+]]() + #pragma omp target private(pvtScl) + { + pvtScl++; + } + + // Region 07 + // CK27-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZE07]]{{.+}}, {{.+}}[[MTYPE07]]{{.+}}) + // CK27-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 + // CK27-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 + // CK27-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 + // CK27-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 + // CK27-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]], + // CK27-DAG: store i8* [[VALP:%.+]], i8** [[P1]], + // CK27-DAG: [[VALBP]] = inttoptr i[[Z]] [[VAL:%.+]] to i8* + // CK27-DAG: [[VALP]] = inttoptr i[[Z]] [[VAL:%.+]] to i8* + // CK27-DAG: [[VAL]] = load i[[Z]], i[[Z]]* [[ADDR:%.+]], + // CK27-64-DAG: [[CADDR:%.+]] = bitcast i[[Z]]* [[ADDR]] to i32* + // CK27-64-DAG: store i32 {{.+}}, i32* [[CADDR]], + + // CK27: call void [[CALL07:@.+]](i[[Z]] [[VAL]]) + #pragma omp target firstprivate(pvtScl) + { + pvtScl++; + } + + // Region 08 + // CK27: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 0, i8** null, i8** null, i{{64|32}}* null, i32* null) + // CK27: call void [[CALL08:@.+]]() + #pragma omp target private(pvtArr) + { + pvtArr[5]++; + } + + // Region 09 + // CK27-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE09]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE09]]{{.+}}) + // CK27-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK27-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK27-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK27-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK27-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK27-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK27-DAG: [[CBPVAL0]] = bitcast [10 x i32]* [[VAR0:%.+]] to i8* + // CK27-DAG: [[CPVAL0]] = bitcast [10 x i32]* [[VAR0]] to i8* + + // CK27: call void [[CALL09:@.+]]([10 x i32]* {{[^,]+}}) + #pragma omp target firstprivate(pvtArr) + { + pvtArr[5]++; + } } // CK27: define {{.+}}[[CALL00]] // CK27: define {{.+}}[[CALL01]] // CK27: define {{.+}}[[CALL02]] // CK27: define {{.+}}[[CALL03]] - +// CK27: define {{.+}}[[CALL04]] +// CK27: define {{.+}}[[CALL05]] +// CK27: define {{.+}}[[CALL06]] +// CK27: define {{.+}}[[CALL07]] #endif #endif