diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -7439,6 +7439,7 @@ // S1 s; // double *p; // struct S2 *ps; + // int &ref; // } // S2 s; // S2 *ps; @@ -7482,6 +7483,14 @@ // optimizes this entry out, same in the examples below) // (***) map the pointee (map: to) // + // map(to: s.ref) + // &s, &(s.ref), sizeof(int*), TARGET_PARAM (*) + // &s, &(s.ref), sizeof(int), MEMBER_OF(1) | PTR_AND_OBJ | TO (***) + // (*) alloc space for struct members, only this is a target parameter + // (**) map the pointer (nothing to be mapped in this example) (the compiler + // optimizes this entry out, same in the examples below) + // (***) map the pointee (map: to) + // // map(s.ps) // &s, &(s.ps), sizeof(S2*), TARGET_PARAM | TO | FROM // @@ -7679,6 +7688,7 @@ uint64_t DimSize = 1; bool IsNonContiguous = CombinedInfo.NonContigInfo.IsNonContiguous; + bool IsPrevMemberReference = false; for (; I != CE; ++I) { // If the current component is member of a struct (parent struct) mark it. @@ -7736,12 +7746,16 @@ .getCanonicalType() ->isAnyPointerType()) || I->getAssociatedExpression()->getType()->isAnyPointerType(); + bool IsMemberReference = isa(I->getAssociatedExpression()) && + MapDecl && + MapDecl->getType()->isLValueReferenceType(); bool IsNonDerefPointer = IsPointer && !UO && !BO && !IsNonContiguous; if (OASE) ++DimSize; - if (Next == CE || IsNonDerefPointer || IsFinalArraySection) { + if (Next == CE || IsMemberReference || IsNonDerefPointer || + IsFinalArraySection) { // If this is not the last component, we expect the pointer to be // associated with an array expression or member expression. assert((Next == CE || @@ -7754,22 +7768,53 @@ "Unexpected expression"); Address LB = Address::invalid(); + Address LowestElem = Address::invalid(); + auto &&EmitMemberExprBase = [](CodeGenFunction &CGF, + const MemberExpr *E) { + const Expr *BaseExpr = E->getBase(); + // If this is s.x, emit s as an lvalue. If it is s->x, emit s as a + // scalar. + LValue BaseLV; + if (E->isArrow()) { + LValueBaseInfo BaseInfo; + TBAAAccessInfo TBAAInfo; + Address Addr = + CGF.EmitPointerWithAlignment(BaseExpr, &BaseInfo, &TBAAInfo); + QualType PtrTy = BaseExpr->getType()->getPointeeType(); + BaseLV = CGF.MakeAddrLValue(Addr, PtrTy, BaseInfo, TBAAInfo); + } else { + BaseLV = CGF.EmitOMPSharedLValue(BaseExpr); + } + return BaseLV; + }; if (OAShE) { - LB = Address(CGF.EmitScalarExpr(OAShE->getBase()), - CGF.getContext().getTypeAlignInChars( - OAShE->getBase()->getType())); - } else { - LB = CGF.EmitOMPSharedLValue(I->getAssociatedExpression()) + LowestElem = LB = Address(CGF.EmitScalarExpr(OAShE->getBase()), + CGF.getContext().getTypeAlignInChars( + OAShE->getBase()->getType())); + } else if (IsMemberReference) { + const auto *ME = cast(I->getAssociatedExpression()); + LValue BaseLVal = EmitMemberExprBase(CGF, ME); + LowestElem = CGF.EmitLValueForFieldInitialization( + BaseLVal, cast(MapDecl)) + .getAddress(CGF); + LB = CGF.EmitLoadOfReferenceLValue(LowestElem, MapDecl->getType()) .getAddress(CGF); + } else { + LowestElem = LB = + CGF.EmitOMPSharedLValue(I->getAssociatedExpression()) + .getAddress(CGF); } // If this component is a pointer inside the base struct then we don't // need to create any entry for it - it will be combined with the object // it is pointing to into a single PTR_AND_OBJ entry. bool IsMemberPointerOrAddr = - (IsPointer || ForDeviceAddr) && EncounteredME && - (dyn_cast(I->getAssociatedExpression()) == - EncounteredME); + EncounteredME && + (((IsPointer || ForDeviceAddr) && + I->getAssociatedExpression() == EncounteredME) || + (IsPrevMemberReference && !IsPointer) || + (IsMemberReference && Next != CE && + !Next->getAssociatedExpression()->getType()->isPointerType())); if (!OverlappedElements.empty() && Next == CE) { // Handle base element with the info for overlapped elements. assert(!PartialStruct.Base.isValid() && "The base element is set."); @@ -7777,11 +7822,11 @@ "Unexpected base element with the pointer type."); // Mark the whole struct as the struct that requires allocation on the // device. - PartialStruct.LowestElem = {0, LB}; + PartialStruct.LowestElem = {0, LowestElem}; CharUnits TypeSize = CGF.getContext().getTypeSizeInChars( I->getAssociatedExpression()->getType()); Address HB = CGF.Builder.CreateConstGEP( - CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(LB, + CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(LowestElem, CGF.VoidPtrTy), TypeSize.getQuantity() - 1); PartialStruct.HighestElem = { @@ -7807,10 +7852,20 @@ Address ComponentLB = Address::invalid(); for (const OMPClauseMappableExprCommon::MappableComponent &MC : Component) { - if (MC.getAssociatedDeclaration()) { - ComponentLB = - CGF.EmitOMPSharedLValue(MC.getAssociatedExpression()) - .getAddress(CGF); + if (const ValueDecl *VD = MC.getAssociatedDeclaration()) { + const auto *FD = dyn_cast(VD); + if (FD && FD->getType()->isLValueReferenceType()) { + const auto *ME = + cast(MC.getAssociatedExpression()); + LValue BaseLVal = EmitMemberExprBase(CGF, ME); + ComponentLB = + CGF.EmitLValueForFieldInitialization(BaseLVal, FD) + .getAddress(CGF); + } else { + ComponentLB = + CGF.EmitOMPSharedLValue(MC.getAssociatedExpression()) + .getAddress(CGF); + } Size = CGF.Builder.CreatePtrDiff( CGF.EmitCastToVoidPtr(ComponentLB.getPointer()), CGF.EmitCastToVoidPtr(LB.getPointer())); @@ -7833,8 +7888,7 @@ CombinedInfo.BasePointers.push_back(BP.getPointer()); CombinedInfo.Pointers.push_back(LB.getPointer()); Size = CGF.Builder.CreatePtrDiff( - CGF.EmitCastToVoidPtr( - CGF.Builder.CreateConstGEP(HB, 1).getPointer()), + CGF.Builder.CreateConstGEP(HB, 1).getPointer(), CGF.EmitCastToVoidPtr(LB.getPointer())); CombinedInfo.Sizes.push_back( CGF.Builder.CreateIntCast(Size, CGF.Int64Ty, /*isSigned=*/true)); @@ -7866,13 +7920,13 @@ OpenMPOffloadMappingFlags Flags = getMapTypeBits( MapType, MapModifiers, MotionModifiers, IsImplicit, !IsExpressionFirstInfo || RequiresReference || - FirstPointerInComplexData, + FirstPointerInComplexData || IsMemberReference, IsCaptureFirstInfo && !RequiresReference, IsNonContiguous); - if (!IsExpressionFirstInfo) { + if (!IsExpressionFirstInfo || IsMemberReference) { // If we have a PTR_AND_OBJ pair where the OBJ is a pointer as well, // then we reset the TO/FROM/ALWAYS/DELETE/CLOSE flags. - if (IsPointer) + if (IsPointer || (IsMemberReference && Next != CE)) Flags &= ~(OMP_MAP_TO | OMP_MAP_FROM | OMP_MAP_ALWAYS | OMP_MAP_DELETE | OMP_MAP_CLOSE); @@ -7898,21 +7952,21 @@ // Update info about the lowest and highest elements for this struct if (!PartialStruct.Base.isValid()) { - PartialStruct.LowestElem = {FieldIndex, LB}; + PartialStruct.LowestElem = {FieldIndex, LowestElem}; if (IsFinalArraySection) { Address HB = CGF.EmitOMPArraySectionExpr(OASE, /*IsLowerBound=*/false) .getAddress(CGF); PartialStruct.HighestElem = {FieldIndex, HB}; } else { - PartialStruct.HighestElem = {FieldIndex, LB}; + PartialStruct.HighestElem = {FieldIndex, LowestElem}; } PartialStruct.Base = BP; PartialStruct.LB = BP; } else if (FieldIndex < PartialStruct.LowestElem.first) { - PartialStruct.LowestElem = {FieldIndex, LB}; + PartialStruct.LowestElem = {FieldIndex, LowestElem}; } else if (FieldIndex > PartialStruct.HighestElem.first) { - PartialStruct.HighestElem = {FieldIndex, LB}; + PartialStruct.HighestElem = {FieldIndex, LowestElem}; } } @@ -7926,11 +7980,12 @@ // The pointer becomes the base for the next element. if (Next != CE) - BP = LB; + BP = IsMemberReference ? LowestElem : LB; IsExpressionFirstInfo = false; IsCaptureFirstInfo = false; FirstPointerInComplexData = false; + IsPrevMemberReference = IsMemberReference; } else if (FirstPointerInComplexData) { QualType Ty = Components.rbegin() ->getAssociatedDeclaration() @@ -8914,7 +8969,6 @@ std::prev(It) ->getAssociatedExpression() ->getType() - .getNonReferenceType() ->isPointerType()) continue; const MapData &BaseData = CI == CE ? L : L1; diff --git a/clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp b/clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp --- a/clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp +++ b/clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp @@ -383,10 +383,10 @@ a++; // CK2: [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 1 - // CK2: [[CBP:%.+]] = bitcast i8** [[BP]] to double*** - // CK2: store double** [[RVAL:%.+]], double*** [[CBP]], + // CK2: [[CBP:%.+]] = bitcast i8** [[BP]] to double**** + // CK2: store double*** [[RVAL:%.+]], double**** [[CBP]], // CK2: call void @__tgt_target_data_begin{{.+}}[[MTYPE01]] - // CK2: [[CBP1:%.+]] = bitcast double*** [[CBP]] to double** + // CK2: [[CBP1:%.+]] = bitcast double**** [[CBP]] to double** // CK2: [[VAL:%.+]] = load double*, double** [[CBP1]], // CK2: store double* [[VAL]], double** [[PVT:%.+]], // CK2: store double** [[PVT]], double*** [[PVT2:%.+]], @@ -428,8 +428,8 @@ la++; // CK2: [[BP1:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* %{{.+}}, i32 0, i32 1 - // CK2: [[CBP1:%.+]] = bitcast i8** [[BP1]] to double*** - // CK2: store double** [[RVAL1:%.+]], double*** [[CBP1]], + // CK2: [[CBP1:%.+]] = bitcast i8** [[BP1]] to double**** + // CK2: store double*** [[RVAL1:%.+]], double**** [[CBP1]], // CK2: [[BP2:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* %{{.+}}, i32 0, i32 2 // CK2: [[CBP2:%.+]] = bitcast i8** [[BP2]] to double*** // CK2: store double** [[RVAL2:%.+]], double*** [[CBP2]], @@ -438,7 +438,7 @@ // CK2: [[VAL2:%.+]] = load double*, double** [[_CBP2]], // CK2: store double* [[VAL2]], double** [[PVT2:%.+]], // CK2: store double** [[PVT2]], double*** [[_PVT2:%.+]], - // CK2: [[_CBP1:%.+]] = bitcast double*** [[CBP1]] to double** + // CK2: [[_CBP1:%.+]] = bitcast double**** [[CBP1]] to double** // CK2: [[VAL1:%.+]] = load double*, double** [[_CBP1]], // CK2: store double* [[VAL1]], double** [[PVT1:%.+]], // CK2: store double** [[PVT1]], double*** [[_PVT1:%.+]], diff --git a/clang/test/OpenMP/target_map_codegen_28.cpp b/clang/test/OpenMP/target_map_codegen_28.cpp --- a/clang/test/OpenMP/target_map_codegen_28.cpp +++ b/clang/test/OpenMP/target_map_codegen_28.cpp @@ -37,13 +37,13 @@ // CK29: [[SSB:%.+]] = type { [[SSA]]*, [[SSA]]** } // CK29-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 -// CK29: [[MTYPE00:@.+]] = private {{.*}}constant [3 x i64] [i64 32, i64 281474976710672, i64 19] +// CK29: [[MTYPE00:@.+]] = private {{.*}}constant [2 x i64] [i64 32, i64 281474976710675] // CK29-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 // CK29: [[MTYPE01:@.+]] = private {{.*}}constant [3 x i64] [i64 32, i64 281474976710672, i64 19] // CK29-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 -// CK29: [[MTYPE02:@.+]] = private {{.*}}constant [3 x i64] [i64 32, i64 281474976710672, i64 19] +// CK29: [[MTYPE02:@.+]] = private {{.*}}constant [2 x i64] [i64 32, i64 281474976710675] struct SSA{ double *p; @@ -60,7 +60,7 @@ void foo() { // Region 00 - // CK29-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{[^,]+}}, i8* {{[^,]+}}, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[Z:64|32]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null, i8** null) + // CK29-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[Z:64|32]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null, i8** null) // CK29-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] // CK29-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] @@ -70,32 +70,22 @@ // CK29-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 // CK29-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 // CK29-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[SSB]]** - // CK29-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[SSA]]** + // CK29-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[SSA]]*** // CK29-DAG: store [[SSB]]* [[VAR0:%.+]], [[SSB]]** [[CBP0]] // CK29-DAG: store [[SSA]]** [[VAR00:%.+]], [[SSA]]*** [[CP0]] // CK29-DAG: store i64 %{{.+}}, i64* [[S0]] // CK29-DAG: [[VAR0]] = load [[SSB]]*, [[SSB]]** % // CK29-DAG: [[VAR00]] = getelementptr inbounds [[SSB]], [[SSB]]* [[VAR0]], i32 0, i32 0 - // CK29-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 - // CK29-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 - // CK29-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1 - // CK29-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[SSA]]*** - // CK29-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to double*** - // CK29-DAG: store [[SSA]]** [[VAR00]], [[SSA]]*** [[CBP1]] - // CK29-DAG: store double** [[VAR1:%.+]], double*** [[CP1]] - // CK29-DAG: store i64 {{8|4}}, i64* [[S1]] - // CK29-DAG: [[VAR1]] = load double**, double*** [[VAR1_REF:%.+]], - // CK29-DAG: [[VAR1_REF]] = getelementptr inbounds [[SSA]], [[SSA]]* %{{.+}}, i32 0, i32 1 - - // CK29-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 - // CK29-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2 - // CK29-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2 - // CK29-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to double*** + // CK29-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK29-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK29-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1 + // CK29-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to double**** // CK29-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to double** - // CK29-DAG: store double** [[VAR1]], double*** [[CBP2]] + // CK29-DAG: store double*** [[VAR1:%.+]], double**** [[CBP2]] // CK29-DAG: store double* [[VAR2:%.+]], double** [[CP2]] // CK29-DAG: store i64 80, i64* [[S2]] + // CK29-DAG: [[VAR1]] = getelementptr inbounds [[SSA]], [[SSA]]* %{{.+}}, i32 0, i32 1 // CK29-DAG: [[VAR2]] = getelementptr inbounds double, double* [[VAR22:%.+]], i{{.+}} 0 // CK29-DAG: [[VAR22]] = load double*, double** %{{.+}}, @@ -116,19 +106,18 @@ // CK29-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 // CK29-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 // CK29-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[SSB]]** - // CK29-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[SSA]]** + // CK29-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[SSA]]**** // CK29-DAG: store [[SSB]]* [[VAR0]], [[SSB]]** [[CBP0]] - // CK29-DAG: store [[SSA]]** [[VAR00:%.+]], [[SSA]]*** [[CP0]] + // CK29-DAG: store [[SSA]]*** [[VAR000:%.+]], [[SSA]]**** [[CP0]] // CK29-DAG: store i64 %{{.+}}, i64* [[S0]] - // CK29-DAG: [[VAR00]] = load [[SSA]]**, [[SSA]]*** [[VAR000:%.+]], // CK29-DAG: [[VAR000]] = getelementptr inbounds [[SSB]], [[SSB]]* [[VAR0]], i32 0, i32 1 // CK29-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 // CK29-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 // CK29-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1 - // CK29-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[SSA]]*** + // CK29-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[SSA]]**** // CK29-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to double*** - // CK29-DAG: store [[SSA]]** [[VAR00]], [[SSA]]*** [[CBP1]] + // CK29-DAG: store [[SSA]]*** [[VAR000]], [[SSA]]**** [[CBP1]] // CK29-DAG: store double** [[VAR1:%.+]], double*** [[CP1]] // CK29-DAG: store i64 {{8|4}}, i64* [[S1]] // CK29-DAG: [[VAR1]] = getelementptr inbounds [[SSA]], [[SSA]]* %{{.+}}, i32 0, i32 0 @@ -151,7 +140,7 @@ } // Region 02 - // CK29-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{[^,]+}}, i8* {{[^,]+}}, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[MTYPE02]]{{.+}}, i8** null, i8** null) + // CK29-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE02]]{{.+}}, i8** null, i8** null) // CK29-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] // CK29-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] @@ -161,32 +150,21 @@ // CK29-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 // CK29-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 // CK29-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[SSB]]** - // CK29-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[SSA]]** + // CK29-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[SSA]]**** // CK29-DAG: store [[SSB]]* [[VAR0]], [[SSB]]** [[CBP0]] - // CK29-DAG: store [[SSA]]** [[VAR00:%.+]], [[SSA]]*** [[CP0]] + // CK29-DAG: store [[SSA]]*** [[VAR000:%.+]], [[SSA]]**** [[CP0]] // CK29-DAG: store i64 %{{.+}}, i64* [[S0]] - // CK29-DAG: [[VAR00]] = load [[SSA]]**, [[SSA]]*** [[VAR000:%.+]], // CK29-DAG: [[VAR000]] = getelementptr inbounds [[SSB]], [[SSB]]* [[VAR0]], i32 0, i32 1 - // CK29-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 - // CK29-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 - // CK29-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1 - // CK29-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[SSA]]*** - // CK29-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to double*** - // CK29-DAG: store [[SSA]]** [[VAR00]], [[SSA]]*** [[CBP1]] - // CK29-DAG: store double** [[VAR1:%.+]], double*** [[CP1]] - // CK29-DAG: store i64 {{8|4}}, i64* [[S1]] - // CK29-DAG: [[VAR1]] = load double**, double*** [[VAR1_REF:%.+]], - // CK29-DAG: [[VAR1_REF]] = getelementptr inbounds [[SSA]], [[SSA]]* %{{.+}}, i32 0, i32 1 - - // CK29-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 - // CK29-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2 - // CK29-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2 - // CK29-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to double*** + // CK29-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK29-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK29-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1 + // CK29-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to double**** // CK29-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to double** - // CK29-DAG: store double** [[VAR1]], double*** [[CBP2]] + // CK29-DAG: store double*** [[VAR1:%.+]], double**** [[CBP2]] // CK29-DAG: store double* [[VAR2:%.+]], double** [[CP2]] // CK29-DAG: store i64 80, i64* [[S2]] + // CK29-DAG: [[VAR1]] = getelementptr inbounds [[SSA]], [[SSA]]* %{{.+}}, i32 0, i32 1 // CK29-DAG: [[VAR2]] = getelementptr inbounds double, double* [[VAR22:%.+]], i{{.+}} 0 // CK29-DAG: [[VAR22]] = load double*, double** %{{.+}}, diff --git a/clang/test/OpenMP/target_map_codegen_35.cpp b/clang/test/OpenMP/target_map_codegen_35.cpp new file mode 100644 --- /dev/null +++ b/clang/test/OpenMP/target_map_codegen_35.cpp @@ -0,0 +1,182 @@ +// expected-no-diagnostics +#ifndef HEADER +#define HEADER + +///==========================================================================/// +// RUN: %clang_cc1 -DCK35 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK35 --check-prefix CK35-64 +// RUN: %clang_cc1 -DCK35 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK35 --check-prefix CK35-64 +// RUN: %clang_cc1 -DCK35 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK35 --check-prefix CK35-32 +// RUN: %clang_cc1 -DCK35 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK35 --check-prefix CK35-32 + +// RUN: %clang_cc1 -DCK35 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY32 %s +// RUN: %clang_cc1 -DCK35 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY32 %s +// RUN: %clang_cc1 -DCK35 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY32 %s +// RUN: %clang_cc1 -DCK35 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY32 %s +// SIMD-ONLY32-NOT: {{__kmpc|__tgt}} +#ifdef CK35 + +class S { +public: + S(double &b) : b(b) {} + int a; + double &b; + void foo(); +}; + +// TARGET_PARAM = 0x20 +// MEMBER_OF_1 | TO = 0x1000000000001 +// MEMBER_OF_1 | PTR_AND_OBJ | TO = 0x1000000000011 +// CK35-DAG: [[MTYPE_TO:@.+]] = {{.+}}constant [4 x i64] [i64 [[#0x20]], i64 [[#0x1000000000001]], i64 [[#0x1000000000001]], i64 [[#0x1000000000011]]] +// TARGET_PARAM = 0x20 +// MEMBER_OF_1 | PTR_AND_OBJ | FROM = 0x1000000000012 +// CK35-DAG: [[MTYPE_FROM:@.+]] = {{.+}}constant [2 x i64] [i64 [[#0x20]], i64 [[#0x1000000000012]]] + +void ref_map() { + double b; + S s(b); + + // CK35-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 -1, i8* @{{.+}}, i32 4, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[4 x i{{.+}}]* [[MTYPE_TO]]{{.+}}, i8** null, i8** null) + // CK35-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK35-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + // CK35-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] + + // pass TARGET_PARAM {&s, &s, ((void*)(&s+1)-(void*)&s)} + + // CK35-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK35-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK35-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 + + // CK35-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to %class.S** + // CK35-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to %class.S** + + // CK35-DAG: store %class.S* [[S_ADDR:%.+]], %class.S** [[BPC0]], + // CK35-DAG: store %class.S* [[S_ADDR]], %class.S** [[PC0]], + // CK35-DAG: store i64 [[S_SIZE:%.+]], i64* [[S0]], + + // CK35-DAG: [[S_SIZE]] = sdiv exact i64 [[SZ:%.+]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64) + // CK35-DAG: [[SZ]] = sub i64 [[S_1_INTPTR:%.+]], [[S_INTPTR:%.+]] + // CK35-DAG: [[S_1_INTPTR]] = ptrtoint i8* [[S_1_VOID:%.+]] to i64 + // CK35-DAG: [[S_INTPTR]] = ptrtoint i8* [[S_VOID:%.+]] to i64 + // CK35-DAG: [[S_1_VOID]] = bitcast %class.S* [[S_1:%.+]] to i8* + // CK35-DAG: [[S_VOID]] = bitcast %class.S* [[S_ADDR]] to i8* + // CK35-DAG: [[S_1]] = getelementptr %class.S, %class.S* [[S_ADDR]], i32 1 + + // pass MEMBER_OF_1 | TO {&s, &s, ((void*)(&s.a+1)-(void*)&s)} to copy the data of s.a. + + // CK35-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK35-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK35-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1 + + // CK35-DAG: [[BPC1:%.+]] = bitcast i8** [[BP1]] to %class.S** + // CK35-DAG: [[PC1:%.+]] = bitcast i8** [[P1]] to %class.S** + + // CK35-DAG: store %class.S* [[S_ADDR]], %class.S** [[BPC1]], + // CK35-DAG: store %class.S* [[S_ADDR]], %class.S** [[PC1]], + // CK35-DAG: store i64 [[A_SIZE:%.+]], i64* [[S1]], + + // CK35-DAG: [[A_SIZE]] = sdiv exact i64 [[SZ:%.+]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64) + // CK35-DAG: [[SZ]] = sub i64 [[B_BEGIN_INTPTR:%.+]], [[S_INTPTR:%.+]] + // CK35-DAG: [[S_INTPTR]] = ptrtoint i8* [[S_VOID:%.+]] to i64 + // CK35-DAG: [[B_BEGIN_INTPTR]] = ptrtoint i8* [[B_BEGIN_VOID:%.+]] to i64 + // CK35-DAG: [[S_VOID]] = bitcast %class.S* [[S_ADDR]] to i8* + // CK35-DAG: [[B_BEGIN_VOID]] = bitcast double** [[B_ADDR:%.+]] to i8* + // CK35-DAG: [[B_ADDR]] = getelementptr inbounds %class.S, %class.S* [[S_ADDR]], i32 0, i32 1 + + // pass MEMBER_OF_1 | TO {&s, &s.b+1, ((void*)(&s+1)-(void*)(&s.b+1))} to copy the data of remainder of s. + + // CK35-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 + // CK35-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2 + // CK35-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2 + + // CK35-DAG: [[BPC2:%.+]] = bitcast i8** [[BP2]] to %class.S** + // CK35-DAG: [[PC2:%.+]] = bitcast i8** [[P2]] to double*** + + // CK35-DAG: store %class.S* [[S_ADDR]], %class.S** [[BPC2]], + // CK35-DAG: store double** [[B_END:%.+]], double*** [[PC2]], + // CK35-DAG: store i64 [[REM_SIZE:%.+]], i64* [[S2]], + + // CK35-DAG: [[B_END]] = getelementptr double*, double** [[B_ADDR]], i{{.+}} 1 + + // CK35-DAG: [[REM_SIZE]] = sdiv exact i64 [[SZ:%.+]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64) + // CK35-DAG: [[SZ]] = sub i64 [[S_END_INTPTR:%.+]], [[B_END_INTPTR:%.+]] + // CK35-DAG: [[B_END_INTPTR]] = ptrtoint i8* [[B_END_VOID:%.+]] to i64 + // CK35-DAG: [[S_END_INTPTR]] = ptrtoint i8* [[S_END_VOID:%.+]] to i64 + // CK35-DAG: [[B_END_VOID]] = bitcast double** [[B_END]] to i8* + // CK35-DAG: [[S_END_VOID]] = getelementptr i8, i8* [[S_LAST:%.+]], i{{.+}} 1 + // CK35-64-DAG: [[S_LAST]] = getelementptr i8, i8* [[S_VOIDPTR:%.+]], i64 15 + // CK35-32-DAG: [[S_LAST]] = getelementptr i8, i8* [[S_VOIDPTR:%.+]], i32 7 + // CK35-DAG: [[S_VOIDPTR]] = bitcast %class.S* [[S_ADDR]] to i8* + + // pass MEMBER_OF_1 | PTR_AND_OBJ | TO {&s, &s.b, 8|4} to copy the data of s.b. + + // CK35-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 3 + // CK35-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 3 + // CK35-DAG: [[S3:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 3 + + // CK35-DAG: [[BPC3:%.+]] = bitcast i8** [[BP3]] to %class.S** + // CK35-DAG: [[PC3:%.+]] = bitcast i8** [[P3]] to double** + + // CK35-DAG: store %class.S* [[S_ADDR]], %class.S** [[BPC3]], + // CK35-DAG: store double* [[B_ADDR:%.+]], double** [[PC3]], + // CK35-DAG: store i64 8, i64* [[S3]], + + // CK35-DAG: [[B_ADDR]] = load double*, double** [[B_REF:%.+]], + // CK35-DAG: [[B_REF]] = getelementptr inbounds %class.S, %class.S* [[S_ADDR]], i32 0, i32 1 + + #pragma omp target map(to: s, s.b) + s.foo(); + + // CK35 : call void + + // CK35-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 -1, i8* @{{.+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE_FROM]]{{.+}}, i8** null, i8** null) + // CK35-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK35-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + // CK35-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] + + // pass TARGET_PARAM {&s, &s.b, ((void*)(&s.b+1)-(void*)&s.b)} + + // CK35-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK35-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK35-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 + + // CK35-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to %class.S** + // CK35-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to double*** + + // CK35-DAG: store %class.S* [[S_ADDR]], %class.S** [[BPC0]], + // CK35-DAG: store double** [[SB_ADDR:%.+]], double*** [[PC0]], + // CK35-DAG: store i64 [[B_SIZE:%.+]], i64* [[S0]], + + // CK35-DAG: [[B_SIZE]] = sdiv exact i64 [[SZ:%.+]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64) + // CK35-DAG: [[SZ]] = sub i64 [[SB_1_INTPTR:%.+]], [[SB_INTPTR:%.+]] + // CK35-DAG: [[SB_1_INTPTR]] = ptrtoint i8* [[SB_1_VOID:%.+]] to i64 + // CK35-DAG: [[SB_INTPTR]] = ptrtoint i8* [[SB_VOID:%.+]] to i64 + // CK35-DAG: [[SB_1_VOID]] = bitcast double** [[SB_1:%.+]] to i8* + // CK35-DAG: [[SB_VOID]] = bitcast double** [[SB_ADDR:%.+]] to i8* + // CK35-DAG: [[SB_ADDR]] = getelementptr inbounds %class.S, %class.S* [[S_ADDR]], i32 0, i32 1 + // CK35-DAG: [[SB_1]] = getelementptr double*, double** [[SB_ADDR]], i{{.+}} 1 + + // pass MEMBER_OF_1 | PTR_AND_OBJ | FROM {&s, &s.b, 8|4} to copy the data of s.c. + + // CK35-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK35-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK35-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1 + + // CK35-DAG: [[BPC1:%.+]] = bitcast i8** [[BP1]] to %class.S** + // CK35-DAG: [[PC1:%.+]] = bitcast i8** [[P1]] to double** + + // CK35-DAG: store %class.S* [[S_ADDR]], %class.S** [[BPC1]], + // CK35-DAG: store double* [[B_ADDR:%.+]], double** [[PC1]], + // CK35-DAG: store i64 8, i64* [[S1]], + + // CK35-DAG: [[B_ADDR]] = load double*, double** [[SB_ADDR]], + + #pragma omp target map(from: s.b) + s.foo(); +} + +#endif // CK35 +#endif diff --git a/openmp/libomptarget/test/mapping/data_member_ref.cpp b/openmp/libomptarget/test/mapping/data_member_ref.cpp new file mode 100644 --- /dev/null +++ b/openmp/libomptarget/test/mapping/data_member_ref.cpp @@ -0,0 +1,69 @@ +// RUN: %libomptarget-compilexx-run-and-check-aarch64-unknown-linux-gnu +// RUN: %libomptarget-compilexx-run-and-check-powerpc64-ibm-linux-gnu +// RUN: %libomptarget-compilexx-run-and-check-powerpc64le-ibm-linux-gnu +// RUN: %libomptarget-compilexx-run-and-check-x86_64-pc-linux-gnu +// RUN: %libomptarget-compilexx-run-and-check-nvptx64-nvidia-cuda + +#include + +struct View { + int Data; +}; + +struct ViewPtr { + int *Data; +}; + +template struct Foo { + Foo(T &V) : VRef(V) {} + T &VRef; +}; + +int main() { + View V; + V.Data = 123456; + Foo Bar(V); + ViewPtr V1; + int Data = 123456; + V1.Data = &Data; + Foo Baz(V1); + + // CHECK: Host 123456. + printf("Host %d.\n", Bar.VRef.Data); +#pragma omp target map(Bar.VRef) + { + // CHECK: Device 123456. + printf("Device %d.\n", Bar.VRef.Data); + V.Data = 654321; + // CHECK: Device 654321. + printf("Device %d.\n", Bar.VRef.Data); + } + // CHECK: Host 654321 654321. + printf("Host %d %d.\n", Bar.VRef.Data, V.Data); + V.Data = 123456; + // CHECK: Host 123456. + printf("Host %d.\n", Bar.VRef.Data); +#pragma omp target map(Bar) map(Bar.VRef) + { + // CHECK: Device 123456. + printf("Device %d.\n", Bar.VRef.Data); + V.Data = 654321; + // CHECK: Device 654321. + printf("Device %d.\n", Bar.VRef.Data); + } + // CHECK: Host 654321 654321. + printf("Host %d %d.\n", Bar.VRef.Data, V.Data); + // CHECK: Host 123456. + printf("Host %d.\n", *Baz.VRef.Data); +#pragma omp target map(*Baz.VRef.Data) + { + // CHECK: Device 123456. + printf("Device %d.\n", *Baz.VRef.Data); + *V1.Data = 654321; + // CHECK: Device 654321. + printf("Device %d.\n", *Baz.VRef.Data); + } + // CHECK: Host 654321 654321 654321. + printf("Host %d %d %d.\n", *Baz.VRef.Data, *V1.Data, Data); + return 0; +}