diff --git a/clang/include/clang/AST/OpenMPClause.h b/clang/include/clang/AST/OpenMPClause.h --- a/clang/include/clang/AST/OpenMPClause.h +++ b/clang/include/clang/AST/OpenMPClause.h @@ -28,6 +28,7 @@ #include "clang/Basic/SourceLocation.h" #include "llvm/ADT/ArrayRef.h" #include "llvm/ADT/MapVector.h" +#include "llvm/ADT/PointerIntPair.h" #include "llvm/ADT/SmallVector.h" #include "llvm/ADT/iterator.h" #include "llvm/ADT/iterator_range.h" @@ -4738,8 +4739,9 @@ /// subscript it may not have any associated declaration. In that case the /// associated declaration is set to nullptr. class MappableComponent { - /// Expression associated with the component. - Expr *AssociatedExpression = nullptr; + /// Pair of Expression and Non-contiguous pair associated with the + /// component. + llvm::PointerIntPair AssociatedExpressionNonContiguousPr; /// Declaration associated with the declaration. If the component does /// not have a declaration (e.g. array subscripts or section), this is set @@ -4749,14 +4751,22 @@ public: explicit MappableComponent() = default; explicit MappableComponent(Expr *AssociatedExpression, - ValueDecl *AssociatedDeclaration) - : AssociatedExpression(AssociatedExpression), + ValueDecl *AssociatedDeclaration, + bool IsNonContiguous) + : AssociatedExpressionNonContiguousPr(AssociatedExpression, + IsNonContiguous), AssociatedDeclaration( AssociatedDeclaration ? cast(AssociatedDeclaration->getCanonicalDecl()) : nullptr) {} - Expr *getAssociatedExpression() const { return AssociatedExpression; } + Expr *getAssociatedExpression() const { + return AssociatedExpressionNonContiguousPr.getPointer(); + } + + bool isNonContiguous() const { + return AssociatedExpressionNonContiguousPr.getInt(); + } ValueDecl *getAssociatedDeclaration() const { return AssociatedDeclaration; diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -9938,6 +9938,8 @@ "bit fields cannot be used to specify storage in a '%0' clause">; def err_array_section_does_not_specify_contiguous_storage : Error< "array section does not specify contiguous storage">; +def err_array_section_does_not_specify_length : Error< + "array section does not specify length for outermost dimension">; def err_omp_union_type_not_allowed : Error< "mapping of union members is not allowed">; def err_omp_expected_access_to_data_field : Error< 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 @@ -7036,6 +7036,10 @@ /// 0x800 is reserved for compatibility with XLC. /// Produce a runtime error if the data is not already allocated. OMP_MAP_PRESENT = 0x1000, + /// Signal that the runtime library should use args as an array of + /// descriptor_dim pointers and use args_size as dims. Used when we have + /// non-contiguous list items in target update directive + OMP_MAP_NON_CONTIG = 0x100000000000, /// The 16 MSBs of the flags indicate whether the entry is member of some /// struct/class. OMP_MAP_MEMBER_OF = 0xffff000000000000, @@ -7072,16 +7076,26 @@ using MapValuesArrayTy = SmallVector; using MapFlagsArrayTy = SmallVector; using MapMappersArrayTy = SmallVector; + using MapDimArrayTy = SmallVector; + using MapNonContiguousArrayTy = SmallVector; /// This structure contains combined information generated for mappable - /// clauses, including base pointers, pointers, sizes, map types, and - /// user-defined mappers. + /// clauses, including base pointers, pointers, sizes, map types, user-defined + /// mappers, and non-contiguous information. struct MapCombinedInfoTy { + struct StructNonContiguousInfo { + bool IsNonContiguous = false; + MapDimArrayTy Dims; + MapNonContiguousArrayTy Offsets; + MapNonContiguousArrayTy Counts; + MapNonContiguousArrayTy Strides; + }; MapBaseValuesArrayTy BasePointers; MapValuesArrayTy Pointers; MapValuesArrayTy Sizes; MapFlagsArrayTy Types; MapMappersArrayTy Mappers; + StructNonContiguousInfo NonContigInfo; /// Append arrays in \a CurInfo. void append(MapCombinedInfoTy &CurInfo) { @@ -7091,6 +7105,14 @@ Sizes.append(CurInfo.Sizes.begin(), CurInfo.Sizes.end()); Types.append(CurInfo.Types.begin(), CurInfo.Types.end()); Mappers.append(CurInfo.Mappers.begin(), CurInfo.Mappers.end()); + NonContigInfo.Dims.append(CurInfo.NonContigInfo.Dims.begin(), + CurInfo.NonContigInfo.Dims.end()); + NonContigInfo.Offsets.append(CurInfo.NonContigInfo.Offsets.begin(), + CurInfo.NonContigInfo.Offsets.end()); + NonContigInfo.Counts.append(CurInfo.NonContigInfo.Counts.begin(), + CurInfo.NonContigInfo.Counts.end()); + NonContigInfo.Strides.append(CurInfo.NonContigInfo.Strides.begin(), + CurInfo.NonContigInfo.Strides.end()); } }; @@ -7248,7 +7270,7 @@ OpenMPOffloadMappingFlags getMapTypeBits( OpenMPMapClauseKind MapType, ArrayRef MapModifiers, ArrayRef MotionModifiers, bool IsImplicit, - bool AddPtrFlag, bool AddIsTargetParamFlag) const { + bool AddPtrFlag, bool AddIsTargetParamFlag, bool IsNonContiguous) const { OpenMPOffloadMappingFlags Bits = IsImplicit ? OMP_MAP_IMPLICIT : OMP_MAP_NONE; switch (MapType) { @@ -7290,6 +7312,8 @@ if (llvm::find(MotionModifiers, OMPC_MOTION_MODIFIER_present) != MotionModifiers.end()) Bits |= OMP_MAP_PRESENT; + if (IsNonContiguous) + Bits |= OMP_MAP_NON_CONTIG; return Bits; } @@ -7600,6 +7624,12 @@ // whether we are dealing with a member of a declared struct. const MemberExpr *EncounteredME = nullptr; + // Track for the total number of dimension. Start from one for the dummy + // dimension. + uint64_t DimSize = 1; + + bool IsNonContiguous = CombinedInfo.NonContigInfo.IsNonContiguous; + for (; I != CE; ++I) { // If the current component is member of a struct (parent struct) mark it. if (!EncounteredME) { @@ -7629,7 +7659,10 @@ // becomes the base address for the following components. // A final array section, is one whose length can't be proved to be one. + // If the map item is non-contiguous then we don't treat any array section + // as final array section. bool IsFinalArraySection = + !IsNonContiguous && isFinalArraySectionExpression(I->getAssociatedExpression()); // Get information on whether the element is a pointer. Have to do a @@ -7647,7 +7680,10 @@ .getCanonicalType() ->isAnyPointerType()) || I->getAssociatedExpression()->getType()->isAnyPointerType(); - bool IsNonDerefPointer = IsPointer && !UO && !BO; + bool IsNonDerefPointer = IsPointer && !UO && !BO && !IsNonContiguous; + + if (OASE) + ++DimSize; if (Next == CE || IsNonDerefPointer || IsFinalArraySection) { // If this is not the last component, we expect the pointer to be @@ -7704,7 +7740,7 @@ OMP_MAP_MEMBER_OF | getMapTypeBits(MapType, MapModifiers, MotionModifiers, IsImplicit, /*AddPtrFlag=*/false, - /*AddIsTargetParamFlag=*/false); + /*AddIsTargetParamFlag=*/false, IsNonContiguous); LB = BP; llvm::Value *Size = nullptr; // Do bitcopy of all non-overlapped structure elements. @@ -7730,6 +7766,8 @@ Size, CGF.Int64Ty, /*isSigned=*/true)); CombinedInfo.Types.push_back(Flags); CombinedInfo.Mappers.push_back(nullptr); + CombinedInfo.NonContigInfo.Dims.push_back(IsNonContiguous ? DimSize + : 1); LB = CGF.Builder.CreateConstGEP(ComponentLB, 1); } CombinedInfo.BasePointers.push_back(BP.getPointer()); @@ -7742,6 +7780,8 @@ CGF.Builder.CreateIntCast(Size, CGF.Int64Ty, /*isSigned=*/true)); CombinedInfo.Types.push_back(Flags); CombinedInfo.Mappers.push_back(nullptr); + CombinedInfo.NonContigInfo.Dims.push_back(IsNonContiguous ? DimSize + : 1); break; } llvm::Value *Size = getExprTypeSize(I->getAssociatedExpression()); @@ -7750,6 +7790,8 @@ CombinedInfo.Pointers.push_back(LB.getPointer()); CombinedInfo.Sizes.push_back( CGF.Builder.CreateIntCast(Size, CGF.Int64Ty, /*isSigned=*/true)); + CombinedInfo.NonContigInfo.Dims.push_back(IsNonContiguous ? DimSize + : 1); // If Mapper is valid, the last component inherits the mapper. bool HasMapper = Mapper && Next == CE; @@ -7759,11 +7801,11 @@ // same expression except for the first one. We also need to signal // this map is the first one that relates with the current capture // (there is a set of entries for each capture). - OpenMPOffloadMappingFlags Flags = - getMapTypeBits(MapType, MapModifiers, MotionModifiers, IsImplicit, - !IsExpressionFirstInfo || RequiresReference || - FirstPointerInComplexData, - IsCaptureFirstInfo && !RequiresReference); + OpenMPOffloadMappingFlags Flags = getMapTypeBits( + MapType, MapModifiers, MotionModifiers, IsImplicit, + !IsExpressionFirstInfo || RequiresReference || + FirstPointerInComplexData, + IsCaptureFirstInfo && !RequiresReference, IsNonContiguous); if (!IsExpressionFirstInfo) { // If we have a PTR_AND_OBJ pair where the OBJ is a pointer as well, @@ -7824,6 +7866,182 @@ FirstPointerInComplexData = false; } } + + if (!IsNonContiguous) + return; + + const ASTContext &Context = CGF.getContext(); + + // For supporting stride in array section, we need to initialize the first + // dimension size as 1, first offset as 0, and first count as 1 + MapValuesArrayTy CurOffsets = {llvm::ConstantInt::get(CGF.CGM.Int64Ty, 0)}; + MapValuesArrayTy CurCounts = {llvm::ConstantInt::get(CGF.CGM.Int64Ty, 1)}; + MapValuesArrayTy CurStrides; + MapValuesArrayTy DimSizes{llvm::ConstantInt::get(CGF.CGM.Int64Ty, 1)}; + uint64_t ElementTypeSize; + + // Collect Size information for each dimension and get the element size as + // the first Stride. For example, for `int arr[10][10]`, the DimSizes + // should be [10, 10] and the first stride is 4 btyes. + for (const OMPClauseMappableExprCommon::MappableComponent &Component : + Components) { + const Expr *AssocExpr = Component.getAssociatedExpression(); + const auto *OASE = dyn_cast(AssocExpr); + + if (!OASE) + continue; + + QualType Ty = OMPArraySectionExpr::getBaseOriginalType(OASE->getBase()); + auto *CAT = Context.getAsConstantArrayType(Ty); + auto *VAT = Context.getAsVariableArrayType(Ty); + + // We need all the dimension size except for the last dimension. + assert((VAT || CAT || &Component == &*Components.begin()) && + "Should be either ConstantArray or VariableArray if not the " + "first Component"); + + // Get element size if CurStrides is empty. + if (CurStrides.empty()) { + const Type *ElementType = nullptr; + if (CAT) + ElementType = CAT->getElementType().getTypePtr(); + else if (VAT) + ElementType = VAT->getElementType().getTypePtr(); + else + assert(&Component == &*Components.begin() && + "Only expect pointer (non CAT or VAT) when this is the " + "first Component"); + // If ElementType is null, then it means the base is a pointer + // (neither CAT nor VAT) and we'll attempt to get ElementType again + // for next iteration. + if (ElementType) { + // For the case that having pointer as base, we need to remove one + // level of indirection. + if (&Component != &*Components.begin()) + ElementType = ElementType->getPointeeOrArrayElementType(); + ElementTypeSize = + Context.getTypeSizeInChars(ElementType).getQuantity(); + CurStrides.push_back( + llvm::ConstantInt::get(CGF.Int64Ty, ElementTypeSize)); + } + } + // Get dimension value except for the last dimension since we don't need + // it. + if (DimSizes.size() < Components.size() - 1) { + if (CAT) + DimSizes.push_back(llvm::ConstantInt::get( + CGF.Int64Ty, CAT->getSize().getZExtValue())); + else if (VAT) + DimSizes.push_back(CGF.Builder.CreateIntCast( + CGF.EmitScalarExpr(VAT->getSizeExpr()), CGF.Int64Ty, + /*IsSigned=*/false)); + } + } + + // Skip the dummy dimension since we have already have its information. + auto DI = DimSizes.begin() + 1; + // Product of dimension. + llvm::Value *DimProd = + llvm::ConstantInt::get(CGF.CGM.Int64Ty, ElementTypeSize); + + // Collect info for non-contiguous. Notice that offset, count, and stride + // are only meaningful for array-section, so we insert a null for anything + // other than array-section. + // Also, the size of offset, count, and stride are not the same as + // pointers, base_pointers, sizes, or dims. Instead, the size of offset, + // count, and stride are the same as the number of non-contiguous + // declaration in target update to/from clause. + for (const OMPClauseMappableExprCommon::MappableComponent &Component : + Components) { + const Expr *AssocExpr = Component.getAssociatedExpression(); + + if (const auto *AE = dyn_cast(AssocExpr)) { + llvm::Value *Offset = CGF.Builder.CreateIntCast( + CGF.EmitScalarExpr(AE->getIdx()), CGF.Int64Ty, + /*isSigned=*/false); + CurOffsets.push_back(Offset); + CurCounts.push_back(llvm::ConstantInt::get(CGF.Int64Ty, /*V=*/1)); + CurStrides.push_back(CurStrides.back()); + continue; + } + + const auto *OASE = dyn_cast(AssocExpr); + + if (!OASE) + continue; + + // Offset + const Expr *OffsetExpr = OASE->getLowerBound(); + llvm::Value *Offset = nullptr; + if (!OffsetExpr) { + // If offset is absent, then we just set it to zero. + Offset = llvm::ConstantInt::get(CGF.Int64Ty, 0); + } else { + Offset = CGF.Builder.CreateIntCast(CGF.EmitScalarExpr(OffsetExpr), + CGF.Int64Ty, + /*isSigned=*/false); + } + CurOffsets.push_back(Offset); + + // Count + const Expr *CountExpr = OASE->getLength(); + llvm::Value *Count = nullptr; + if (!CountExpr) { + // In Clang, once a high dimension is an array section, we construct all + // the lower dimension as array section, however, for case like + // arr[0:2][2], Clang construct the inner dimension as an array section + // but it actually is not in an array section form according to spec. + if (!OASE->getColonLocFirst().isValid() && + !OASE->getColonLocSecond().isValid()) { + Count = llvm::ConstantInt::get(CGF.Int64Ty, 1); + } else { + // OpenMP 5.0, 2.1.5 Array Sections, Description. + // When the length is absent it defaults to ⌈(size − + // lower-bound)/stride⌉, where size is the size of the array + // dimension. + const Expr *StrideExpr = OASE->getStride(); + llvm::Value *Stride = + StrideExpr + ? CGF.Builder.CreateIntCast(CGF.EmitScalarExpr(StrideExpr), + CGF.Int64Ty, /*isSigned=*/false) + : nullptr; + if (Stride) + Count = CGF.Builder.CreateUDiv( + CGF.Builder.CreateNUWSub(*DI, Offset), Stride); + else + Count = CGF.Builder.CreateNUWSub(*DI, Offset); + } + } else { + Count = CGF.EmitScalarExpr(CountExpr); + } + Count = CGF.Builder.CreateIntCast(Count, CGF.Int64Ty, /*isSigned=*/false); + CurCounts.push_back(Count); + + // Stride_n' = Stride_n * (D_0 * D_1 ... * D_n-1) * Unit size + // Take `int arr[5][5][5]` and `arr[0:2:2][1:2:1][0:2:2]` as an example: + // Offset Count Stride + // D0 0 1 4 (int) <- dummy dimension + // D1 0 2 8 (2 * (1) * 4) + // D2 1 2 20 (1 * (1 * 5) * 4) + // D3 0 2 200 (2 * (1 * 5 * 4) * 4) + const Expr *StrideExpr = OASE->getStride(); + llvm::Value *Stride = + StrideExpr + ? CGF.Builder.CreateIntCast(CGF.EmitScalarExpr(StrideExpr), + CGF.Int64Ty, /*isSigned=*/false) + : nullptr; + DimProd = CGF.Builder.CreateNUWMul(DimProd, *(DI - 1)); + if (Stride) + CurStrides.push_back(CGF.Builder.CreateNUWMul(DimProd, Stride)); + else + CurStrides.push_back(DimProd); + if (DI != DimSizes.end()) + ++DI; + } + + CombinedInfo.NonContigInfo.Offsets.push_back(CurOffsets); + CombinedInfo.NonContigInfo.Counts.push_back(CurCounts); + CombinedInfo.NonContigInfo.Strides.push_back(CurStrides); } /// Return the adjusted map modifiers if the declaration a capture refers to @@ -8229,6 +8447,8 @@ // Remember the current base pointer index. unsigned CurrentBasePointersIdx = CurInfo.BasePointers.size(); + CurInfo.NonContigInfo.IsNonContiguous = + L.Components.back().isNonContiguous(); generateInfoForComponentList(L.MapType, L.MapModifiers, L.MotionModifiers, L.Components, CurInfo, PartialStruct, IsFirstComponentList, @@ -8349,6 +8569,7 @@ // If there is an entry in PartialStruct it means we have a struct with // individual members mapped. Emit an extra combined entry. if (PartialStruct.Base.isValid()) + CurInfo.NonContigInfo.Dims.push_back(0); emitCombinedEntry(CombinedInfo, CurInfo.Types, PartialStruct); // We need to append the results of this capture to what we already have. @@ -8718,13 +8939,82 @@ }; } // anonymous namespace +static void emitNonContiguousDescriptor( + CodeGenFunction &CGF, MappableExprsHandler::MapCombinedInfoTy &CombinedInfo, + CGOpenMPRuntime::TargetDataInfo &Info) { + CodeGenModule &CGM = CGF.CGM; + MappableExprsHandler::MapCombinedInfoTy::StructNonContiguousInfo + &NonContigInfo = CombinedInfo.NonContigInfo; + + // Build an array of struct descriptor_dim and then assign it to + // offload_args. + // + // struct descriptor_dim { + // uint64_t offset; + // uint64_t count; + // uint64_t stride + // }; + ASTContext &C = CGF.getContext(); + QualType Int64Ty = C.getIntTypeForBitwidth(/*DestWidth=*/64, /*Signed=*/0); + RecordDecl *RD; + RD = C.buildImplicitRecord("descriptor_dim"); + RD->startDefinition(); + addFieldToRecordDecl(C, RD, Int64Ty); + addFieldToRecordDecl(C, RD, Int64Ty); + addFieldToRecordDecl(C, RD, Int64Ty); + RD->completeDefinition(); + QualType DimTy = C.getRecordType(RD); + + enum { OffsetFD = 0, CountFD, StrideFD }; + // We need two index variable here since the size of "Dims" is the same as the + // size of Components, however, the size of offset, count, and stride is equal + // to the size of base declaration that is non-contiguous. + for (unsigned I = 0, L = 0, E = NonContigInfo.Dims.size(); I < E; ++I) { + // Skip emitting ir if dimension size is 1 since it cannot be + // non-contiguous. + if (NonContigInfo.Dims[I] == 1) + continue; + llvm::APInt Size(/*numBits=*/32, NonContigInfo.Dims[I]); + QualType ArrayTy = + C.getConstantArrayType(DimTy, Size, nullptr, ArrayType::Normal, 0); + Address DimsAddr = CGF.CreateMemTemp(ArrayTy, "dims"); + for (unsigned II = 0, EE = NonContigInfo.Dims[I]; II < EE; ++II) { + unsigned RevIdx = EE - II - 1; + LValue DimsLVal = CGF.MakeAddrLValue( + CGF.Builder.CreateConstArrayGEP(DimsAddr, II), DimTy); + // Offset + LValue OffsetLVal = CGF.EmitLValueForField( + DimsLVal, *std::next(RD->field_begin(), OffsetFD)); + CGF.EmitStoreOfScalar(NonContigInfo.Offsets[L][RevIdx], OffsetLVal); + // Count + LValue CountLVal = CGF.EmitLValueForField( + DimsLVal, *std::next(RD->field_begin(), CountFD)); + CGF.EmitStoreOfScalar(NonContigInfo.Counts[L][RevIdx], CountLVal); + // Stride + LValue StrideLVal = CGF.EmitLValueForField( + DimsLVal, *std::next(RD->field_begin(), StrideFD)); + CGF.EmitStoreOfScalar(NonContigInfo.Strides[L][RevIdx], StrideLVal); + } + // args[I] = &dims + Address DAddr = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast( + DimsAddr, CGM.Int8PtrTy); + llvm::Value *P = CGF.Builder.CreateConstInBoundsGEP2_32( + llvm::ArrayType::get(CGM.VoidPtrTy, Info.NumberOfPtrs), + Info.PointersArray, 0, I); + Address PAddr(P, CGF.getPointerAlign()); + CGF.Builder.CreateStore(DAddr.getPointer(), PAddr); + ++L; + } +} + /// Emit the arrays used to pass the captures and map information to the /// offloading runtime library. If there is no map or capture information, /// return nullptr by reference. static void emitOffloadingArrays(CodeGenFunction &CGF, MappableExprsHandler::MapCombinedInfoTy &CombinedInfo, - CGOpenMPRuntime::TargetDataInfo &Info) { + CGOpenMPRuntime::TargetDataInfo &Info, + bool IsNonContiguous = false) { CodeGenModule &CGM = CGF.CGM; ASTContext &Ctx = CGF.getContext(); @@ -8770,8 +9060,15 @@ // We expect all the sizes to be constant, so we collect them to create // a constant array. SmallVector ConstSizes; - for (llvm::Value *S : CombinedInfo.Sizes) - ConstSizes.push_back(cast(S)); + for (unsigned I = 0, E = CombinedInfo.Sizes.size(); I < E; ++I) { + if (IsNonContiguous && + (CombinedInfo.Types[I] & MappableExprsHandler::OMP_MAP_NON_CONTIG)) { + ConstSizes.push_back(llvm::ConstantInt::get( + CGF.Int64Ty, CombinedInfo.NonContigInfo.Dims[I])); + } else { + ConstSizes.push_back(cast(CombinedInfo.Sizes[I])); + } + } auto *SizesArrayInit = llvm::ConstantArray::get( llvm::ArrayType::get(CGM.Int64Ty, ConstSizes.size()), ConstSizes); @@ -8872,6 +9169,12 @@ CGF.Builder.CreateStore(MFunc, MAddr); } } + + if (!IsNonContiguous || CombinedInfo.NonContigInfo.Offsets.empty() || + Info.NumberOfPtrs == 0) + return; + + emitNonContiguousDescriptor(CGF, CombinedInfo, Info); } namespace { @@ -10274,7 +10577,7 @@ MEHandler.generateAllInfo(CombinedInfo); // Fill up the arrays and create the arguments. - emitOffloadingArrays(CGF, CombinedInfo, Info); + emitOffloadingArrays(CGF, CombinedInfo, Info, /*IsNonContiguous=*/true); llvm::Value *BasePointersArrayArg = nullptr; llvm::Value *PointersArrayArg = nullptr; @@ -10520,7 +10823,7 @@ TargetDataInfo Info; // Fill up the arrays and create the arguments. - emitOffloadingArrays(CGF, CombinedInfo, Info); + emitOffloadingArrays(CGF, CombinedInfo, Info, /*IsNonContiguous=*/true); bool RequiresOuterTask = D.hasClausesOfKind() || D.hasClausesOfKind(); emitOffloadingArraysArgument( diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp --- a/clang/lib/Sema/SemaOpenMP.cpp +++ b/clang/lib/Sema/SemaOpenMP.cpp @@ -48,7 +48,7 @@ static const Expr *checkMapClauseExpressionBase( Sema &SemaRef, Expr *E, OMPClauseMappableExprCommon::MappableExprComponentList &CurComponents, - OpenMPClauseKind CKind, bool NoDiagnose); + OpenMPClauseKind CKind, OpenMPDirectiveKind DKind, bool NoDiagnose); namespace { /// Default data sharing attributes, which can be applied to directive. @@ -3637,6 +3637,7 @@ if (isOpenMPTargetExecutionDirective(DKind)) { OMPClauseMappableExprCommon::MappableExprComponentList CurComponents; if (!checkMapClauseExpressionBase(SemaRef, E, CurComponents, OMPC_map, + Stack->getCurrentDirective(), /*NoDiagnose=*/true)) return; const auto *VD = cast( @@ -16808,11 +16809,14 @@ class MapBaseChecker final : public StmtVisitor { Sema &SemaRef; OpenMPClauseKind CKind = OMPC_unknown; + OpenMPDirectiveKind DKind = OMPD_unknown; OMPClauseMappableExprCommon::MappableExprComponentList &Components; + bool IsNonContiguous = false; bool NoDiagnose = false; const Expr *RelevantExpr = nullptr; bool AllowUnitySizeArraySection = true; bool AllowWholeSizeArraySection = true; + bool AllowAnotherPtr = true; SourceLocation ELoc; SourceRange ERange; @@ -16837,7 +16841,7 @@ assert(!RelevantExpr && "RelevantExpr is expected to be nullptr"); RelevantExpr = DRE; // Record the component. - Components.emplace_back(DRE, DRE->getDecl()); + Components.emplace_back(DRE, DRE->getDecl(), IsNonContiguous); return true; } @@ -16909,7 +16913,7 @@ AllowWholeSizeArraySection = false; // Record the component. - Components.emplace_back(ME, FD); + Components.emplace_back(ME, FD, IsNonContiguous); return RelevantExpr || Visit(E); } @@ -16947,7 +16951,7 @@ } // Record the component - we don't have any declaration associated. - Components.emplace_back(AE, nullptr); + Components.emplace_back(AE, nullptr, IsNonContiguous); return RelevantExpr || Visit(E); } @@ -16986,6 +16990,13 @@ // pointer. Otherwise, only unitary sections are accepted. if (NotWhole || IsPointer) AllowWholeSizeArraySection = false; + } else if (DKind == OMPD_target_update && + SemaRef.getLangOpts().OpenMP >= 50) { + if (IsPointer && !AllowAnotherPtr) + SemaRef.Diag(ELoc, diag::err_omp_section_length_undefined) + << /*array of unknown bound */ 1; + else + IsNonContiguous = true; } else if (AllowUnitySizeArraySection && NotUnity) { // A unity or whole array section is not allowed and that is not // compatible with the properties of the current array section. @@ -16995,6 +17006,9 @@ return false; } + if (IsPointer) + AllowAnotherPtr = false; + if (const auto *TE = dyn_cast(E)) { Expr::EvalResult ResultR; Expr::EvalResult ResultL; @@ -17020,14 +17034,14 @@ } // Record the component - we don't have any declaration associated. - Components.emplace_back(OASE, nullptr); + Components.emplace_back(OASE, nullptr, /*IsNonContiguous=*/false); return RelevantExpr || Visit(E); } bool VisitOMPArrayShapingExpr(OMPArrayShapingExpr *E) { Expr *Base = E->getBase(); // Record the component - we don't have any declaration associated. - Components.emplace_back(E, nullptr); + Components.emplace_back(E, nullptr, IsNonContiguous); return Visit(Base->IgnoreParenImpCasts()); } @@ -17040,7 +17054,7 @@ } if (!RelevantExpr) { // Record the component if haven't found base decl. - Components.emplace_back(UO, nullptr); + Components.emplace_back(UO, nullptr, /*IsNonContiguous=*/false); } return RelevantExpr || Visit(UO->getSubExpr()->IgnoreParenImpCasts()); } @@ -17056,7 +17070,7 @@ // know the other subtree is just an offset) Expr *LE = BO->getLHS()->IgnoreParenImpCasts(); Expr *RE = BO->getRHS()->IgnoreParenImpCasts(); - Components.emplace_back(BO, nullptr); + Components.emplace_back(BO, nullptr, false); assert((LE->getType().getTypePtr() == BO->getType().getTypePtr() || RE->getType().getTypePtr() == BO->getType().getTypePtr()) && "Either LHS or RHS have base decl inside"); @@ -17067,12 +17081,12 @@ bool VisitCXXThisExpr(CXXThisExpr *CTE) { assert(!RelevantExpr && "RelevantExpr is expected to be nullptr"); RelevantExpr = CTE; - Components.emplace_back(CTE, nullptr); + Components.emplace_back(CTE, nullptr, IsNonContiguous); return true; } bool VisitCXXOperatorCallExpr(CXXOperatorCallExpr *COCE) { assert(!RelevantExpr && "RelevantExpr is expected to be nullptr"); - Components.emplace_back(COCE, nullptr); + Components.emplace_back(COCE, nullptr, IsNonContiguous); return true; } bool VisitStmt(Stmt *) { @@ -17083,10 +17097,10 @@ return RelevantExpr; } explicit MapBaseChecker( - Sema &SemaRef, OpenMPClauseKind CKind, + Sema &SemaRef, OpenMPClauseKind CKind, OpenMPDirectiveKind DKind, OMPClauseMappableExprCommon::MappableExprComponentList &Components, bool NoDiagnose, SourceLocation &ELoc, SourceRange &ERange) - : SemaRef(SemaRef), CKind(CKind), Components(Components), + : SemaRef(SemaRef), CKind(CKind), DKind(DKind), Components(Components), NoDiagnose(NoDiagnose), ELoc(ELoc), ERange(ERange) {} }; } // namespace @@ -17098,13 +17112,30 @@ static const Expr *checkMapClauseExpressionBase( Sema &SemaRef, Expr *E, OMPClauseMappableExprCommon::MappableExprComponentList &CurComponents, - OpenMPClauseKind CKind, bool NoDiagnose) { + OpenMPClauseKind CKind, OpenMPDirectiveKind DKind, bool NoDiagnose) { SourceLocation ELoc = E->getExprLoc(); SourceRange ERange = E->getSourceRange(); - MapBaseChecker Checker(SemaRef, CKind, CurComponents, NoDiagnose, ELoc, + MapBaseChecker Checker(SemaRef, CKind, DKind, CurComponents, NoDiagnose, ELoc, ERange); - if (Checker.Visit(E->IgnoreParens())) + if (Checker.Visit(E->IgnoreParens())) { + // Check if the highest dimension array section has length specified + if (SemaRef.getLangOpts().OpenMP >= 50 && !CurComponents.empty() && + (CKind == OMPC_to || CKind == OMPC_from)) { + auto CI = CurComponents.rbegin(); + auto CE = CurComponents.rend(); + for (; CI != CE; ++CI) { + const auto *OASE = + dyn_cast(CI->getAssociatedExpression()); + if (!OASE) + continue; + if (OASE && OASE->getLength()) + break; + SemaRef.Diag(ELoc, diag::err_array_section_does_not_specify_length) + << ERange; + } + } return Checker.getFoundBase(); + } return nullptr; } @@ -17581,7 +17612,8 @@ // Obtain the array or member expression bases if required. Also, fill the // components array with all the components identified in the process. const Expr *BE = checkMapClauseExpressionBase( - SemaRef, SimpleExpr, CurComponents, CKind, /*NoDiagnose=*/false); + SemaRef, SimpleExpr, CurComponents, CKind, DSAS->getCurrentDirective(), + /*NoDiagnose=*/false); if (!BE) continue; @@ -18881,8 +18913,8 @@ // only need a component. MVLI.VarBaseDeclarations.push_back(D); MVLI.VarComponents.resize(MVLI.VarComponents.size() + 1); - MVLI.VarComponents.back().push_back( - OMPClauseMappableExprCommon::MappableComponent(SimpleRefExpr, D)); + MVLI.VarComponents.back().emplace_back(SimpleRefExpr, D, + /*IsNonContiguous=*/false); } if (MVLI.ProcessedVarList.empty()) @@ -18933,8 +18965,8 @@ if (VD && (isa(RefExpr->IgnoreParenImpCasts()) || isa(RefExpr->IgnoreParenImpCasts()))) Component = DefaultFunctionArrayLvalueConversion(SimpleRefExpr).get(); - MVLI.VarComponents.back().push_back( - OMPClauseMappableExprCommon::MappableComponent(Component, D)); + MVLI.VarComponents.back().emplace_back(Component, D, + /*IsNonContiguous=*/false); } if (MVLI.ProcessedVarList.empty()) @@ -19000,7 +19032,8 @@ // Store the components in the stack so that they can be used to check // against other clauses later on. - OMPClauseMappableExprCommon::MappableComponent MC(SimpleRefExpr, D); + OMPClauseMappableExprCommon::MappableComponent MC( + SimpleRefExpr, D, /*IsNonContiguous=*/false); DSAStack->addMappableExpressionComponents( D, MC, /*WhereFoundClauseKind=*/OMPC_is_device_ptr); diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp --- a/clang/lib/Serialization/ASTReader.cpp +++ b/clang/lib/Serialization/ASTReader.cpp @@ -12639,10 +12639,10 @@ SmallVector Components; Components.reserve(TotalComponents); for (unsigned i = 0; i < TotalComponents; ++i) { - Expr *AssociatedExpr = Record.readExpr(); + Expr *AssociatedExprPr = Record.readExpr(); auto *AssociatedDecl = Record.readDeclAs(); - Components.push_back(OMPClauseMappableExprCommon::MappableComponent( - AssociatedExpr, AssociatedDecl)); + Components.emplace_back(AssociatedExprPr, AssociatedDecl, + /*IsNonContiguous=*/false); } C->setComponents(Components, ListSizes); } @@ -12762,10 +12762,10 @@ SmallVector Components; Components.reserve(TotalComponents); for (unsigned i = 0; i < TotalComponents; ++i) { - Expr *AssociatedExpr = Record.readSubExpr(); + Expr *AssociatedExprPr = Record.readSubExpr(); + bool IsNonContiguous = Record.readBool(); auto *AssociatedDecl = Record.readDeclAs(); - Components.push_back(OMPClauseMappableExprCommon::MappableComponent( - AssociatedExpr, AssociatedDecl)); + Components.emplace_back(AssociatedExprPr, AssociatedDecl, IsNonContiguous); } C->setComponents(Components, ListSizes); } @@ -12818,10 +12818,10 @@ SmallVector Components; Components.reserve(TotalComponents); for (unsigned i = 0; i < TotalComponents; ++i) { - Expr *AssociatedExpr = Record.readSubExpr(); + Expr *AssociatedExprPr = Record.readSubExpr(); + bool IsNonContiguous = Record.readBool(); auto *AssociatedDecl = Record.readDeclAs(); - Components.push_back(OMPClauseMappableExprCommon::MappableComponent( - AssociatedExpr, AssociatedDecl)); + Components.emplace_back(AssociatedExprPr, AssociatedDecl, IsNonContiguous); } C->setComponents(Components, ListSizes); } @@ -12868,10 +12868,10 @@ SmallVector Components; Components.reserve(TotalComponents); for (unsigned i = 0; i < TotalComponents; ++i) { - Expr *AssociatedExpr = Record.readSubExpr(); + auto *AssociatedExprPr = Record.readSubExpr(); auto *AssociatedDecl = Record.readDeclAs(); - Components.push_back(OMPClauseMappableExprCommon::MappableComponent( - AssociatedExpr, AssociatedDecl)); + Components.emplace_back(AssociatedExprPr, AssociatedDecl, + /*IsNonContiguous=*/false); } C->setComponents(Components, ListSizes); } @@ -12912,8 +12912,8 @@ for (unsigned i = 0; i < TotalComponents; ++i) { Expr *AssociatedExpr = Record.readSubExpr(); auto *AssociatedDecl = Record.readDeclAs(); - Components.push_back(OMPClauseMappableExprCommon::MappableComponent( - AssociatedExpr, AssociatedDecl)); + Components.emplace_back(AssociatedExpr, AssociatedDecl, + /*IsNonContiguous*/ false); } C->setComponents(Components, ListSizes); } @@ -12955,8 +12955,8 @@ for (unsigned i = 0; i < TotalComponents; ++i) { Expr *AssociatedExpr = Record.readSubExpr(); auto *AssociatedDecl = Record.readDeclAs(); - Components.push_back(OMPClauseMappableExprCommon::MappableComponent( - AssociatedExpr, AssociatedDecl)); + Components.emplace_back(AssociatedExpr, AssociatedDecl, + /*IsNonContiguous=*/false); } C->setComponents(Components, ListSizes); } diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp --- a/clang/lib/Serialization/ASTWriter.cpp +++ b/clang/lib/Serialization/ASTWriter.cpp @@ -6695,6 +6695,7 @@ Record.push_back(N); for (auto &M : C->all_components()) { Record.AddStmt(M.getAssociatedExpression()); + Record.writeBool(M.isNonContiguous()); Record.AddDeclRef(M.getAssociatedDeclaration()); } } @@ -6724,6 +6725,7 @@ Record.push_back(N); for (auto &M : C->all_components()) { Record.AddStmt(M.getAssociatedExpression()); + Record.writeBool(M.isNonContiguous()); Record.AddDeclRef(M.getAssociatedDeclaration()); } } diff --git a/clang/test/OpenMP/target_update_ast_print.cpp b/clang/test/OpenMP/target_update_ast_print.cpp --- a/clang/test/OpenMP/target_update_ast_print.cpp +++ b/clang/test/OpenMP/target_update_ast_print.cpp @@ -29,6 +29,35 @@ #pragma omp target update from(b, ([a][targ])p) if(l<5) device(l-1) nowait depend(inout:l) + U marr[10][10][10]; +#pragma omp target update to(marr[2][0:2][0:2]) + +#pragma omp target update from(marr[2][0:2][0:2]) + +#pragma omp target update from(marr[:2][0:2][0:2:1]) + +#pragma omp target update to(marr[:l][:l][l:]) + +#pragma omp target update to(marr[:2][:1][:]) + +#pragma omp target update from(marr[:2][:1][:]) + +#pragma omp target update to(marr[:2][:][:1]) + +#pragma omp target update from(marr[:2][:][:1]) + +#pragma omp target update to(marr[:2][:] [1:]) + +#pragma omp target update from(marr[:2][:][1:]) + +#pragma omp target update to(marr[:1][3:2][:2]) + +#pragma omp target update from(marr[:1][3:2][:2]) + +#pragma omp target update to(marr[:1][:2][0]) + +#pragma omp target update from(marr[:1][:2][0]) + int arr[100][100]; #pragma omp target update to(arr[2][0:1:2]) @@ -57,6 +86,21 @@ // CHECK-NEXT: int l; // CHECK-NEXT: #pragma omp target update to(([a][targ])p,a) if(l > 5) device(l) nowait depend(inout : l) // CHECK-NEXT: #pragma omp target update from(b,([a][targ])p) if(l < 5) device(l - 1) nowait depend(inout : l) +// CHECK: marr[10][10][10]; +// CHECK-NEXT: #pragma omp target update to(marr[2][0:2][0:2]) +// CHECK-NEXT: #pragma omp target update from(marr[2][0:2][0:2]) +// CHECK-NEXT: #pragma omp target update from(marr[:2][0:2][0:2:1]) +// CHECK-NEXT: #pragma omp target update to(marr[:l][:l][l:]) +// CHECK-NEXT: #pragma omp target update to(marr[:2][:1][:]) +// CHECK-NEXT: #pragma omp target update from(marr[:2][:1][:]) +// CHECK-NEXT: #pragma omp target update to(marr[:2][:][:1]) +// CHECK-NEXT: #pragma omp target update from(marr[:2][:][:1]) +// CHECK-NEXT: #pragma omp target update to(marr[:2][:][1:]) +// CHECK-NEXT: #pragma omp target update from(marr[:2][:][1:]) +// CHECK-NEXT: #pragma omp target update to(marr[:1][3:2][:2]) +// CHECK-NEXT: #pragma omp target update from(marr[:1][3:2][:2]) +// CHECK-NEXT: #pragma omp target update to(marr[:1][:2][0]) +// CHECK-NEXT: #pragma omp target update from(marr[:1][:2][0]) // CHECK: int arr[100][100]; // CHECK-NEXT: #pragma omp target update to(arr[2][0:1:2]) // CHECK-NEXT: #pragma omp target update from(arr[2][0:1:2]) @@ -86,6 +130,40 @@ // OMP5-NEXT: #pragma omp target update from(present: arr[2][0:1:2], a) #endif +float marr[10][10][10]; +// CHECK: marr[10][10][10]; +#pragma omp target update to(marr[2][0:2][0:2]) +// CHECK-NEXT: #pragma omp target update to(marr[2][0:2][0:2]) +#pragma omp target update from(marr[2][0:2][0:2]) +// CHECK-NEXT: #pragma omp target update from(marr[2][0:2][0:2]) +#pragma omp target update to(marr[:n][:n][n:]) +// CHECK: #pragma omp target update to(marr[:n][:n][n:]) +#pragma omp target update from(marr[:2][:1][:]) +// CHECK-NEXT: #pragma omp target update from(marr[:2][:1][:]) +#pragma omp target update to(marr[:2][:][:1]) +// CHECK-NEXT: #pragma omp target update to(marr[:2][:][:1]) +#pragma omp target update from(marr[:2][:][:1]) +// CHECK-NEXT: #pragma omp target update from(marr[:2][:][:1]) +#pragma omp target update to(marr[:2][:][1:]) +// CHECK-NEXT: #pragma omp target update to(marr[:2][:][1:]) +#pragma omp target update from(marr[:2][:][1:]) +// CHECK-NEXT: #pragma omp target update from(marr[:2][:][1:]) +#pragma omp target update to(marr[:1][3:2][:2]) +// CHECK-NEXT: #pragma omp target update to(marr[:1][3:2][:2]) +#pragma omp target update from(marr[:1][3:2][:2]) +// CHECK-NEXT: #pragma omp target update from(marr[:1][3:2][:2]) +#pragma omp target update to(marr[:1][:2][0]) +// CHECK-NEXT: #pragma omp target update to(marr[:1][:2][0]) +#pragma omp target update from(marr[:1][:2][0]) +// CHECK-NEXT: #pragma omp target update from(marr[:1][:2][0]) +#pragma omp target update to(marr[:2:][0:2][0:2:1]) +// CHECK-NEXT: #pragma omp target update to(marr[:2:][0:2][0:2:1]) +#pragma omp target update from(marr[:2:][0:2][0:2:1]) +// CHECK-NEXT: #pragma omp target update from(marr[:2:][0:2][0:2:1]) +#pragma omp target update to(marr[:2:][:2:][0:2:1]) +// CHECK-NEXT: #pragma omp target update to(marr[:2:][:2:][0:2:1]) +#pragma omp target update from(marr[:2:][:2:][0:2:1]) +// CHECK-NEXT: #pragma omp target update from(marr[:2:][:2:][0:2:1]) return foo(argc, f) + foo(argv[0][0], f) + a; } diff --git a/clang/test/OpenMP/target_update_codegen.cpp b/clang/test/OpenMP/target_update_codegen.cpp --- a/clang/test/OpenMP/target_update_codegen.cpp +++ b/clang/test/OpenMP/target_update_codegen.cpp @@ -1158,5 +1158,493 @@ #pragma omp target update from(present: lb) ; } +#endif + +///==========================================================================/// +// RUN: %clang_cc1 -DCK20 -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 CK20 --check-prefix CK20-64 +// RUN: %clang_cc1 -DCK20 -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 CK20 --check-prefix CK20-64 +// RUN: %clang_cc1 -DCK20 -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 CK20 --check-prefix CK20-32 +// RUN: %clang_cc1 -DCK20 -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 CK20 --check-prefix CK20-32 + +// RUN: %clang_cc1 -DCK20 -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-ONLY19 %s +// RUN: %clang_cc1 -DCK20 -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-ONLY19 %s +// RUN: %clang_cc1 -DCK20 -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-ONLY19 %s +// RUN: %clang_cc1 -DCK20 -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-ONLY19 %s +// SIMD-ONLY19-NOT: {{__kmpc|__tgt}} +#ifdef CK20 + +struct ST { + int a; + double *b; +}; + +// CK20: [[STRUCT_ST:%.+]] = type { i32, double* } +// CK20: [[STRUCT_DESCRIPTOR:%.+]] = type { i64, i64, i64 } + +// CK20: [[MSIZE:@.+]] = {{.+}}constant [1 x i64] [i64 3] +// CK20: [[MTYPE:@.+]] = {{.+}}constant [1 x i64] [i64 17592186044449] + +// CK20-LABEL: _Z3foo +void foo(int arg) { + ST arr[3][4]; + // CK20: [[DIMS:%.+]] = alloca [3 x [[STRUCT_DESCRIPTOR]]], + // CK20: [[ARRAY_IDX:%.+]] = getelementptr inbounds [3 x [4 x [[STRUCT_ST]]]], [3 x [4 x [[STRUCT_ST]]]]* [[ARR:%.+]], {{.+}} 0, {{.+}} 0 + // CK20: [[ARRAY_DECAY:%.+]] = getelementptr inbounds [4 x [[STRUCT_ST]]], [4 x [[STRUCT_ST]]]* [[ARRAY_IDX]], {{.+}} 0, {{.+}} 0 + // CK20: [[ARRAY_IDX_1:%.+]] = getelementptr inbounds [[STRUCT_ST]], [[STRUCT_ST]]* [[ARRAY_DECAY]], {{.+}} + // CK20: [[BP0:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BP:%.+]], {{.+}} 0, {{.+}} 0 + // CK20: [[BPC:%.+]] = bitcast i8** [[BP0]] to [3 x [4 x [[STRUCT_ST]]]]** + // CK20: store [3 x [4 x [[STRUCT_ST]]]]* [[ARR]], [3 x [4 x [[STRUCT_ST]]]]** [[BPC]], + // CK20: [[P0:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[P:%.+]], {{.+}} 0, {{.+}} 0 + // CK20: [[PC:%.+]] = bitcast i8** [[P0]] to [[STRUCT_ST]]** + // CK20: store [[STRUCT_ST]]* [[ARRAY_IDX_1]], [[STRUCT_ST]]** [[PC]], + // CK20: [[DIM_1:%.+]] = getelementptr inbounds [3 x [[STRUCT_DESCRIPTOR]]], [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 0 + // CK20: [[OFFSET:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 0 + // CK20: store i64 0, i64* [[OFFSET]], + // CK20: [[COUNT:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 1 + // CK20: store i64 2, i64* [[COUNT]], + // CK20: [[STRIDE:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 2 + // CK20: store i64 {{32|64}}, i64* [[STRIDE]], + // CK20: [[DIM_2:%.+]] = getelementptr inbounds [3 x [[STRUCT_DESCRIPTOR]]], [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 1 + // CK20: [[OFFSET_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 0 + // CK20: store i64 1, i64* [[OFFSET_2]], + // CK20: [[COUNT_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 1 + // CK20: store i64 4, i64* [[COUNT_2]], + // CK20: [[STRIDE_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 2 + // CK20: store i64 {{8|16}}, i64* [[STRIDE_2]], + // CK20: [[DIM_3:%.+]] = getelementptr inbounds [3 x [[STRUCT_DESCRIPTOR]]], [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 2 + // CK20: [[OFFSET_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 0 + // CK20: store i64 0, i64* [[OFFSET_3]], + // CK20: [[COUNT_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 1 + // CK20: store i64 1, i64* [[COUNT_3]], + // CK20: [[STRIDE_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 2 + // CK20: store i64 {{8|16}}, i64* [[STRIDE_3]], + // CK20-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MSIZE]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE]]{{.+}}) + // CK20-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]] + // CK20-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + // CK20-DAG: [[PC0:%.+]] = bitcast [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]] to i8* + // CK20-DAG: [[PTRS:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* %.offload_ptrs, i32 0, i32 0 + // CK20-DAG: store i8* [[PC0]], i8** [[PTRS]], + +#pragma omp target update to(arr[0:2][1:4]) + { ++arg; } +} + +#endif +///==========================================================================/// +// RUN: %clang_cc1 -DCK21 -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 CK21 --check-prefix CK21-64 +// RUN: %clang_cc1 -DCK21 -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 CK21 --check-prefix CK21-64 +// RUN: %clang_cc1 -DCK21 -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 CK21 --check-prefix CK21-32 +// RUN: %clang_cc1 -DCK21 -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 CK21 --check-prefix CK21-32 + +// RUN: %clang_cc1 -DCK21 -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-ONLY19 %s +// RUN: %clang_cc1 -DCK21 -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-ONLY19 %s +// RUN: %clang_cc1 -DCK21 -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-ONLY19 %s +// RUN: %clang_cc1 -DCK21 -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-ONLY19 %s +// SIMD-ONLY19-NOT: {{__kmpc|__tgt}} +#ifdef CK21 + +// CK21: [[STRUCT_ST:%.+]] = type { [10 x [10 x [10 x double*]]] } +// CK21: [[STRUCT_DESCRIPTOR:%.+]] = type { i64, i64, i64 } + +// CK21: [[MTYPE:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 299067162755073] + +struct ST { + double *dptr[10][10][10]; + + // CK21: _ZN2ST3fooEv + void foo() { + // CK21: [[DIMS:%.+]] = alloca [4 x [[STRUCT_DESCRIPTOR]]], + // CK21: [[ARRAY_IDX:%.+]] = getelementptr inbounds [10 x [10 x [10 x double*]]], [10 x [10 x [10 x double*]]]* [[DPTR:%.+]], {{.+}} 0, {{.+}} 0 + // CK21: [[ARRAY_DECAY:%.+]] = getelementptr inbounds [10 x [10 x double*]], [10 x [10 x double*]]* [[ARRAY_IDX]], {{.+}} 0, {{.+}} 0 + // CK21: [[ARRAY_IDX_1:%.+]] = getelementptr inbounds [10 x double*], [10 x double*]* [[ARRAY_DECAY]], {{.+}} 1 + // CK21: [[ARRAY_DECAY_2:%.+]] = getelementptr inbounds [10 x double*], [10 x double*]* [[ARRAY_IDX_1]], {{.+}} 0, {{.+}} 0 + // CK21: [[ARRAY_IDX_3:%.+]] = getelementptr inbounds {{.+}}, {{.+}}* [[ARRAY_DECAY_2]], {{.+}} 0 + // CK21: [[BP0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BP:%.+]], {{.+}} 0, {{.+}} 0 + // CK21: [[P0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[P:%.+]], i{{.+}} 0, i{{.+}} 0 + // CK21: [[DIM_1:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 0 + // CK21: [[OFFSET_1:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 0 + // CK21: store i64 0, i64* [[OFFSET_1]], + // CK21: [[COUNT_1:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 1 + // CK21: store i64 2, i64* [[COUNT_1]], + // CK21: [[STRIDE_1:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 2 + // CK21: store i64 {{400|800}}, i64* [[STRIDE_1]], + // CK21: [[DIM_2:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 1 + // CK21: [[OFFSET_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 0 + // CK21: store i64 1, i64* [[OFFSET_2]], + // CK21: [[COUNT_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 1 + // CK21: store i64 3, i64* [[COUNT_2]], + // CK21: [[STRIDE_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 2 + // CK21: store i64 {{40|80}}, i64* [[STRIDE_2]], + // CK21: [[DIM_3:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 2 + // CK21: [[OFFSET_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 0 + // CK21: store i64 0, i64* [[OFFSET_3]], + // CK21: [[COUNT_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 1 + // CK21: store i64 4, i64* [[COUNT_3]], + // CK21: [[STRIDE_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 2 + // CK21: store i64 {{4|8}}, i64* [[STRIDE_3]], + // CK21: [[DIM_4:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 3 + // CK21: [[OFFSET_4:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 0 + // CK21: store i64 0, i64* [[OFFSET_4]], + // CK21: [[COUNT_4:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 1 + // CK21: store i64 1, i64* [[COUNT_4]], + // CK21: [[STRIDE_4:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 2 + // CK21: store i64 {{4|8}}, i64* [[STRIDE_4]], + // CK21-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i{{.+}}* [[GEPSZ:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE]]{{.+}}) + // CK21-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]] + // CK21-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + // CK21-DAG: [[PC0:%.+]] = bitcast [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]] to i8* + // CK21-DAG: [[PTRS:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_ptrs, i32 0, i32 0 + // CK21-DAG: store i8* [[PC0]], i8** [[PTRS]], +#pragma omp target update to(dptr[0:2][1:3][0:4]) + } +}; + +void bar() { + ST st; + st.foo(); +} + +#endif +///==========================================================================/// +// RUN: %clang_cc1 -DCK22 -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 CK22 --check-prefix CK22-64 +// RUN: %clang_cc1 -DCK22 -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 CK22 --check-prefix CK22-64 +// RUN: %clang_cc1 -DCK22 -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 CK22 --check-prefix CK22-32 +// RUN: %clang_cc1 -DCK22 -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 CK22 --check-prefix CK22-32 + +// RUN: %clang_cc1 -DCK22 -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-ONLY19 %s +// RUN: %clang_cc1 -DCK22 -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-ONLY19 %s +// RUN: %clang_cc1 -DCK22 -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-ONLY19 %s +// RUN: %clang_cc1 -DCK22 -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-ONLY19 %s +// SIMD-ONLY19-NOT: {{__kmpc|__tgt}} +#ifdef CK22 + +// CK22: [[STRUCT_DESCRIPTOR:%.+]] = type { i64, i64, i64 } + +// CK22: [[MSIZE:@.+]] = {{.+}}constant [1 x i64] [i64 4] +// CK22: [[MTYPE:@.+]] = {{.+}}constant [1 x i64] [i64 17592186044449] + +struct ST { + // CK22: _ZN2ST3fooEPA10_Pi + void foo(int *arr[5][10]) { + // CK22: [[DIMS:%.+]] = alloca [4 x [[STRUCT_DESCRIPTOR]]], + // CK22: [[ARRAY_IDX:%.+]] = getelementptr inbounds [10 x i32*], [10 x i32*]* [[ARR:%.+]], {{.+}} 0 + // CK22: [[ARRAY_DECAY:%.+]] = getelementptr inbounds [10 x i32*], [10 x i32*]* [[ARRAY_IDX]], {{.+}} 0, {{.+}} 0 + // CK22: [[ARRAY_IDX_2:%.+]] = getelementptr inbounds i32*, i32** [[ARRAY_DECAY:%.+]], {{.+}} 1 + // CK22: [[BP0:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BP:%.+]], {{.+}} 0, {{.+}} 0 + // CK22: [[P0:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[P:%.+]], i{{.+}} 0, i{{.+}} 0 + // CK22: [[DIM_1:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 0 + // CK22: [[OFFSET:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 0 + // CK22: store i64 0, i64* [[OFFSET]], + // CK22: [[COUNT:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 1 + // CK22: store i64 2, i64* [[COUNT]], + // CK22: [[STRIDE:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 2 + // CK22: store i64 200, i64* [[STRIDE]], + // CK22: [[DIM_2:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 1 + // CK22: [[OFFSET:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 0 + // CK22: store i64 1, i64* [[OFFSET]], + // CK22: [[COUNT:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 1 + // CK22: store i64 3, i64* [[COUNT]], + // CK22: [[STRIDE:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 2 + // CK22: store i64 40, i64* [[STRIDE]], + // CK22: [[DIM_3:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 2 + // CK22: [[OFFSET:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 0 + // CK22: store i64 0, i64* [[OFFSET]], + // CK22: [[COUNT:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 1 + // CK22: store i64 4, i64* [[COUNT]], + // CK22: [[STRIDE:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 2 + // CK22: store i64 4, i64* [[STRIDE]], + // CK22: [[DIM_4:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 3 + // CK22: [[OFFSET:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 0 + // CK22: store i64 0, i64* [[OFFSET]], + // CK22: [[COUNT:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 1 + // CK22: store i64 1, i64* [[COUNT]], + // CK22: [[STRIDE:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 2 + // CK22: store i64 4, i64* [[STRIDE]], + // CK22-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MSIZE]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE]]{{.+}}) + // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]] + // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + // CK22-DAG: [[PC0:%.+]] = bitcast [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]] to i8* + // CK22-DAG: [[PTRS:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* %.offload_ptrs, i32 0, i32 0 + // CK22-DAG: store i8* [[PC0]], i8** [[PTRS]], +#pragma omp target update to(arr[0:2][1:3][0:4]) + } +}; + +void bar() { + ST st; + int *arr[5][10]; + st.foo(arr); +} + +#endif +///==========================================================================/// +// RUN: %clang_cc1 -DCK23 -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 CK23 --check-prefix CK23-64 +// RUN: %clang_cc1 -DCK23 -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 CK23 --check-prefix CK23-64 +// RUN: %clang_cc1 -DCK23 -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 CK23 --check-prefix CK23-32 +// RUN: %clang_cc1 -DCK23 -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 CK23 --check-prefix CK23-32 + +// RUN: %clang_cc1 -DCK23 -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-ONLY19 %s +// RUN: %clang_cc1 -DCK23 -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-ONLY19 %s +// RUN: %clang_cc1 -DCK23 -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-ONLY19 %s +// RUN: %clang_cc1 -DCK23 -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-ONLY19 %s +// SIMD-ONLY19-NOT: {{__kmpc|__tgt}} +#ifdef CK23 + +// CK23: [[STRUCT_DESCRIPTOR:%.+]] = type { i64, i64, i64 } + +// CK23: [[MSIZE:@.+]] = {{.+}}constant [1 x i64] [i64 4] +// CK23: [[MTYPE:@.+]] = {{.+}}constant [1 x i64] [i64 17592186044449] + +// CK23: foo +void foo(int arg) { + float farr[5][5][5]; + // CK23: [[ARG_ADDR:%.+]] = alloca i32, + // CK23: [[DIMS:%.+]] = alloca [4 x [[STRUCT_DESCRIPTOR]]], + // CK23: [[ARRAY_IDX:%.+]] = getelementptr inbounds [5 x [5 x [5 x float]]], [5 x [5 x [5 x float]]]* [[ARR:%.+]], {{.+}} 0, {{.+}} 0 + // CK23: [[ARRAY_DECAY:%.+]] = getelementptr inbounds [5 x [5 x float]], [5 x [5 x float]]* [[ARRAY_IDX]], {{.+}} 0, {{.+}} 0 + // CK23: [[ARRAY_IDX_1:%.+]] = getelementptr inbounds [5 x float], [5 x float]* [[ARRAY_DECAY]], {{.+}} + // CK23: [[ARRAY_DECAY_2:%.+]] = getelementptr inbounds [5 x float], [5 x float]* [[ARRAY_IDX_1]], {{.+}} 0, {{.+}} 0 + // CK23: [[ARRAY_IDX_2:%.+]] = getelementptr inbounds float, float* [[ARRAY_DECAY_2]], {{.+}} + // CK23: [[MUL:%.+]] = mul nuw i64 4, + // CK23: [[BP0:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BP:%.+]], {{.+}} 0, {{.+}} 0 + // CK23: [[BPC:%.+]] = bitcast i8** [[BP0]] to [5 x [5 x [5 x float]]]** + // CK23: store [5 x [5 x [5 x float]]]* [[ARR]], [5 x [5 x [5 x float]]]** [[BPC]], + // CK23: [[P0:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[P:%.+]], {{.+}} 0, {{.+}} 0 + // CK23: [[PC:%.+]] = bitcast i8** [[P0]] to float** + // CK23: store float* [[ARRAY_IDX_2]], float** [[PC]], + // CK23: [[DIM_1:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 0 + // CK23: [[OFFSET:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 0 + // CK23: store i64 0, i64* [[OFFSET]], + // CK23: [[COUNT:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 1 + // CK23: store i64 2, i64* [[COUNT]], + // CK23: [[STRIDE:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 2 + // CK23: store i64 200, i64* [[STRIDE]], + // CK23: [[DIM_2:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 1 + // CK23: [[OFFSET_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 0 + // CK23: store i64 1, i64* [[OFFSET_2]], + // CK23: [[COUNT_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 1 + // CK23: store i64 2, i64* [[COUNT_2]], + // CK23: [[STRIDE_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 2 + // CK23: store i64 20, i64* [[STRIDE_2]], + // CK23: [[DIM_3:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 2 + // CK23: [[OFFSET_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 0 + // CK23: store i64 0, i64* [[OFFSET_3]], + // CK23: [[COUNT_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 1 + // CK23: store i64 2, i64* [[COUNT_3]], + // CK23: [[STRIDE_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 2 + // CK23: store i64 [[MUL]], i64* [[STRIDE_3]], + // CK23: [[DIM_4:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 3 + // CK23: [[OFFSET_4:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 0 + // CK23: store i64 0, i64* [[OFFSET_4]], + // CK23: [[COUNT_4:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 1 + // CK23: store i64 1, i64* [[COUNT_4]], + // CK23: [[STRIDE_4:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 2 + // CK23: store i64 4, i64* [[STRIDE_4]], + // CK23-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MSIZE]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE]]{{.+}}) + // CK23-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]] + // CK23-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + // CK23-DAG: [[PC0:%.+]] = bitcast [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]] to i8* + // CK23-DAG: [[PTRS:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* %.offload_ptrs, i32 0, i32 0 + // CK23-DAG: store i8* [[PC0]], i8** [[PTRS]], +#pragma omp target update to(farr[0:2:2][1:2:1][0:2:arg]) +} + +#endif +///==========================================================================/// +// RUN: %clang_cc1 -DCK24 -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 CK24 --check-prefix CK24-64 +// RUN: %clang_cc1 -DCK24 -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 CK24 --check-prefix CK24-64 +// RUN: %clang_cc1 -DCK24 -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 CK24 --check-prefix CK24-32 +// RUN: %clang_cc1 -DCK24 -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 CK24 --check-prefix CK24-32 + +// RUN: %clang_cc1 -DCK24 -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-ONLY19 %s +// RUN: %clang_cc1 -DCK24 -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-ONLY19 %s +// RUN: %clang_cc1 -DCK24 -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-ONLY19 %s +// RUN: %clang_cc1 -DCK24 -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-ONLY19 %s +// SIMD-ONLY19-NOT: {{__kmpc|__tgt}} +#ifdef CK24 + +// CK24: [[STRUCT_DESCRIPTOR:%.+]] = type { i64, i64, i64 } + +// CK24: [[MSIZE:@.+]] = {{.+}}constant [1 x i64] [i64 4] +// CK24: [[MTYPE:@.+]] = {{.+}}constant [1 x i64] [i64 17592186044449] + +// CK24: foo +void foo(int arg) { + double darr[3][4][5]; + // CK24: [[DIMS:%.+]] = alloca [4 x [[STRUCT_DESCRIPTOR]]], + // CK24: [[ARRAY_IDX:%.+]] = getelementptr inbounds [3 x [4 x [5 x double]]], [3 x [4 x [5 x double]]]* [[ARR:%.+]], {{.+}} 0, {{.+}} 0 + // CK24: [[ARRAY_DECAY:%.+]] = getelementptr inbounds [4 x [5 x double]], [4 x [5 x double]]* [[ARRAY_IDX]], {{.+}} 0, {{.+}} 0 + // CK24: [[ARRAY_IDX_1:%.+]] = getelementptr inbounds [5 x double], [5 x double]* [[ARRAY_DECAY]], {{.+}} + // CK24: [[ARRAY_DECAY_2:%.+]] = getelementptr inbounds [5 x double], [5 x double]* [[ARRAY_IDX_1]], {{.+}} 0, {{.+}} 0 + // CK24: [[ARRAY_IDX_2:%.+]] = getelementptr inbounds double, double* [[ARRAY_DECAY_2]], {{.+}} + // CK24: [[MUL:%.+]] = mul nuw i64 8, + // CK24: [[SUB:%.+]] = sub nuw i64 4, [[ARG:%.+]] + // CK24: [[LEN:%.+]] = udiv {{.+}} [[SUB]], 1 + // CK24: [[BP0:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BP:%.+]], {{.+}} 0, {{.+}} 0 + // CK24: [[BPC:%.+]] = bitcast i8** [[BP0]] to [3 x [4 x [5 x double]]]** + // CK24: store [3 x [4 x [5 x double]]]* [[ARR]], [3 x [4 x [5 x double]]]** [[BPC]], + // CK24: [[P0:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[P:%.+]], {{.+}} 0, {{.+}} 0 + // CK24: [[PC:%.+]] = bitcast i8** [[P0]] to double** + // CK24: store double* [[ARRAY_IDX_2]], double** [[PC]], + // CK24: [[DIM_1:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 0 + // CK24: [[OFFSET:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 0 + // CK24: store i64 0, i64* [[OFFSET]], + // CK24: [[COUNT:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 1 + // CK24: store i64 2, i64* [[COUNT]], + // CK24: [[STRIDE:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 2 + // CK24: store i64 320, i64* [[STRIDE]], + // CK24: [[DIM_2:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 1 + // CK24: [[OFFSET_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 0 + // CK24: store i64 [[ARG]], i64* [[OFFSET_2]], + // CK24: [[COUNT_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 1 + // CK24: store i64 [[LEN]], i64* [[COUNT_2]], + // CK24: [[STRIDE_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 2 + // CK24: store i64 40, i64* [[STRIDE_2]], + // CK24: [[DIM_3:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 2 + // CK24: [[OFFSET_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 0 + // CK24: store i64 0, i64* [[OFFSET_3]], + // CK24: [[COUNT_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 1 + // CK24: store i64 2, i64* [[COUNT_3]], + // CK24: [[STRIDE_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 2 + // CK24: store i64 [[MUL]], i64* [[STRIDE_3]], + // CK24: [[DIM_4:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 3 + // CK24: [[OFFSET_4:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 0 + // CK24: store i64 0, i64* [[OFFSET_4]], + // CK24: [[COUNT_4:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 1 + // CK24: store i64 1, i64* [[COUNT_4]], + // CK24: [[STRIDE_4:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 2 + // CK24: store i64 8, i64* [[STRIDE_4]], + // CK24-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MSIZE]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE]]{{.+}}) + // CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]] + // CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + // CK24-DAG: [[PC0:%.+]] = bitcast [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]] to i8* + // CK24-DAG: [[PTRS:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* %.offload_ptrs, i32 0, i32 0 + // CK24-DAG: store i8* [[PC0]], i8** [[PTRS]], +#pragma omp target update to(darr[0:2:2][arg: :1][:2:arg]) +} +#endif +///==========================================================================/// +// RUN: %clang_cc1 -DCK25 -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 CK25 --check-prefix CK25-64 +// RUN: %clang_cc1 -DCK25 -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 CK25 --check-prefix CK25-64 +// RUN: %clang_cc1 -DCK25 -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 CK25 --check-prefix CK25-32 +// RUN: %clang_cc1 -DCK25 -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 CK25 --check-prefix CK25-32 + +// RUN: %clang_cc1 -DCK25 -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-ONLY19 %s +// RUN: %clang_cc1 -DCK25 -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-ONLY19 %s +// RUN: %clang_cc1 -DCK25 -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-ONLY19 %s +// RUN: %clang_cc1 -DCK25 -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-ONLY19 %s +// SIMD-ONLY19-NOT: {{__kmpc|__tgt}} +#ifdef CK25 + +// CK25: [[STRUCT_DESCRIPTOR:%.+]] = type { i64, i64, i64 } + +// CK25: [[MSIZE:@.+]] = {{.+}}constant [3 x i64] [i64 4, i64 4, i64 3] +// CK25: [[MTYPE:@.+]] = {{.+}}constant [3 x i64] [i64 17592186044449, i64 33, i64 17592186044449] + +// CK25-LABEL: _Z3foo +void foo(int arg) { + int arr[3][4][5], x; + float farr[4][3]; + + // CK25: [[DIMS:%.+]] = alloca [4 x [[STRUCT_DESCRIPTOR]]], + // CK25: [[DIMS_2:%.+]] = alloca [3 x [[STRUCT_DESCRIPTOR]]], + // CK25: [[ARRAY_IDX:%.+]] = getelementptr inbounds [3 x [4 x [5 x i32]]], [3 x [4 x [5 x i32]]]* [[ARR:%.+]], {{.+}} 0, {{.+}} 0 + // CK25: [[ARRAY_DECAY:%.+]] = getelementptr inbounds [4 x [5 x i32]], [4 x [5 x i32]]* [[ARRAY_IDX]], {{.+}} 0, {{.+}} 0 + // CK25: [[ARRAY_IDX_1:%.+]] = getelementptr inbounds [5 x i32], [5 x i32]* [[ARRAY_DECAY]], {{.+}} + // CK25: [[ARRAY_DECAY_2:%.+]] = getelementptr inbounds [5 x i32], [5 x i32]* [[ARRAY_IDX_1]], {{.+}} 0, {{.+}} 0 + // CK25: [[ARRAY_IDX_3:%.+]] = getelementptr inbounds {{.+}}, {{.+}}* [[ARRAY_DECAY_2]], {{.+}} 1 + // CK25: [[LEN:%.+]] = sub nuw i64 4, [[ARG_ADDR:%.+]] + // CK25: [[ARRAY_IDX_4:%.+]] = getelementptr inbounds [4 x [3 x float]], [4 x [3 x float]]* [[FARR:%.+]], {{.+}} 0, {{.+}} 0 + // CK25: [[ARRAY_DECAY_5:%.+]] = getelementptr inbounds [3 x float], [3 x float]* [[ARRAY_IDX_4]], {{.+}} 0, {{.+}} 0 + // CK25: [[ARRAY_IDX_6:%.+]] = getelementptr inbounds float, float* [[ARRAY_DECAY_5:%.+]], {{.+}} 1 + // CK25: [[BP0:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BP:%.+]], i{{.+}} 0, i{{.+}} 0 + // CK25: [[P0:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[P:%.+]], i{{.+}} 0, i{{.+}} 0 + // CK25: [[DIM_1:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 0 + // CK25: [[OFFSET:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 0 + // CK25: store i64 0, i64* [[OFFSET]], + // CK25: [[COUNT:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 1 + // CK25: store i64 2, i64* [[COUNT]], + // CK25: [[STRIDE:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 2 + // CK25: store i64 80, i64* [[STRIDE]], + // CK25: [[DIM_2:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 1 + // CK25: [[OFFSET_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 0 + // CK25: store i64 [[ARG:%.+]], i64* [[OFFSET_2]], + // CK25: [[COUNT_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 1 + // CK25: store i64 [[LEN]], i64* [[COUNT_2]], + // CK25: [[STRIDE_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 2 + // CK25: store i64 20, i64* [[STRIDE_2]], + // CK25: [[DIM_3:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 2 + // CK25: [[OFFSET_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 0 + // CK25: store i64 1, i64* [[OFFSET_3]], + // CK25: [[COUNT_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 1 + // CK25: store i64 4, i64* [[COUNT_3]], + // CK25: [[STRIDE_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 2 + // CK25: store i64 4, i64* [[STRIDE_3]], + // CK25: [[DIM_4:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 3 + // CK25: [[OFFSET_4:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 0 + // CK25: store i64 0, i64* [[OFFSET_4]], + // CK25: [[COUNT_4:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 1 + // CK25: store i64 1, i64* [[COUNT_4]], + // CK25: [[STRIDE_4:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 2 + // CK25: store i64 4, i64* [[STRIDE_4]], + // CK25: [[PC0:%.+]] = bitcast [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]] to i8* + // CK25: [[PTRS:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* %.offload_ptrs, i32 0, i32 0 + // CK25: store i8* [[PC0]], i8** [[PTRS]], + // CK25: [[DIM_5:%.+]] = getelementptr inbounds [3 x [[STRUCT_DESCRIPTOR]]], [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS_2]], {{.+}} 0, {{.+}} 0 + // CK25: [[OFFSET_2_1:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_5]], {{.+}} 0, {{.+}} 0 + // CK25: store i64 0, i64* [[OFFSET_2_1]], + // CK25: [[COUNT_2_1:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_5]], {{.+}} 0, {{.+}} 1 + // CK25: store i64 2, i64* [[COUNT_2_1]], + // CK25: [[STRIDE_2_1:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_5]], {{.+}} 0, {{.+}} 2 + // CK25: store i64 12, i64* [[STRIDE_2_1]], + // CK25: [[DIM_6:%.+]] = getelementptr inbounds [3 x [[STRUCT_DESCRIPTOR]]], [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS_2]], {{.+}} 0, {{.+}} 1 + // CK25: [[OFFSET_2_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_6]], {{.+}} 0, {{.+}} 0 + // CK25: store i64 1, i64* [[OFFSET_2_2]], + // CK25: [[COUNT_2_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_6]], {{.+}} 0, {{.+}} 1 + // CK25: store i64 2, i64* [[COUNT_2_2]], + // CK25: [[STRIDE_2_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_6]], {{.+}} 0, {{.+}} 2 + // CK25: store i64 4, i64* [[STRIDE_2_2]], + // CK25: [[DIM_7:%.+]] = getelementptr inbounds [3 x [[STRUCT_DESCRIPTOR]]], [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS_2]], {{.+}} 0, {{.+}} 2 + // CK25: [[OFFSET_2_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_7]], {{.+}} 0, {{.+}} 0 + // CK25: store i64 0, i64* [[OFFSET_2_3]], + // CK25: [[COUNT_2_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_7]], {{.+}} 0, {{.+}} 1 + // CK25: store i64 1, i64* [[COUNT_2_3]], + // CK25: [[STRIDE_2_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_7]], {{.+}} 0, {{.+}} 2 + // CK25: store i64 4, i64* [[STRIDE_2_3]], + // CK25: [[PC1:%.+]] = bitcast [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS_2]] to i8* + // CK25: [[PTRS_2:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* %.offload_ptrs, i32 0, i32 2 + // CK25: store i8* [[PC1]], i8** [[PTRS_2]], + // CK25-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[MSIZE]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[MTYPE]]{{.+}}) + // CK25-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]] + // CK25-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + +#pragma omp target update to(arr[0:2][arg:][1:4], x, farr[0:2][1:2]) + { ++arg; } +} + #endif #endif diff --git a/clang/test/OpenMP/target_update_messages.cpp b/clang/test/OpenMP/target_update_messages.cpp --- a/clang/test/OpenMP/target_update_messages.cpp +++ b/clang/test/OpenMP/target_update_messages.cpp @@ -136,6 +136,25 @@ foo(); } + double marr[10][5][10]; +#pragma omp target update to(marr[0:2][2:4][1:2]) // lt50-error {{array section does not specify contiguous storage}} lt50-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + {} +#pragma omp target update from(marr[0:2][2:4][1:2]) // lt50-error {{array section does not specify contiguous storage}} lt50-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + +#pragma omp target update to(marr[0:][1:2:2][1:2]) // ge50-error {{array section does not specify length for outermost dimension}} lt50-error {{expected ']'}} lt50-note {{to match this '['}} lt50-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + {} +#pragma omp target update from(marr[0:][1:2:2][1:2]) // ge50-error {{array section does not specify length for outermost dimension}} lt50-error {{expected ']'}} lt50-note {{to match this '['}} lt50-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + int arr[4][3][2][1]; +#pragma omp target update to(arr[0:2][2:4][:2][1]) // lt50-error {{array section does not specify contiguous storage}} lt50-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + {} +#pragma omp target update from(arr[0:2][2:4][:2][1]) // lt50-error {{array section does not specify contiguous storage}} lt50-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + double ***dptr; +#pragma omp target update to(dptr[0:2][2:4][1:2]) // lt50-error {{array section does not specify contiguous storage}} ge50-error 2 {{section length is unspecified and cannot be inferred because subscripted value is an array of unknown bound}} lt50-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + {} +#pragma omp target update from(dptr[0:2][2:4][1:2]) // lt50-error {{array section does not specify contiguous storage}} ge50-error 2 {{section length is unspecified and cannot be inferred because subscripted value is an array of unknown bound}} lt50-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + int iarr[5][5]; // ge50-error@+4 {{section stride is evaluated to a non-positive value -1}} // lt50-error@+3 {{expected ']'}} @@ -148,6 +167,21 @@ // lt50-note@+2 {{to match this '['}} // expected-error@+1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} #pragma omp target update from(iarr[0:][1:2:-1]) + {} +// lt50-error@+5 {{expected expression}} +// ge50-error@+4 {{array section does not specify length for outermost dimension}} +// lt50-error@+3 {{expected ']'}} +// lt50-note@+2 {{to match this '['}} +// lt50-error@+1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(iarr[0: :2][1:2]) + {} +// lt50-error@+5 {{expected expression}} +// ge50-error@+4 {{array section does not specify length for outermost dimension}} +// lt50-error@+3 {{expected ']'}} +// lt50-note@+2 {{to match this '['}} +// lt50-error@+1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(iarr[0: :2][1:2]) + {} return tmain(argc, argv); } diff --git a/clang/test/OpenMP/target_update_to_messages.cpp b/clang/test/OpenMP/target_update_to_messages.cpp --- a/clang/test/OpenMP/target_update_to_messages.cpp +++ b/clang/test/OpenMP/target_update_to_messages.cpp @@ -79,6 +79,10 @@ #pragma omp target update to(*(*(this->ptr)+a+this->ptr)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} #pragma omp target update to(*(this+this)) // expected-error {{invalid operands to binary expression ('S8 *' and 'S8 *')}} {} + + double marr[10][5][10]; +#pragma omp target update to(marr [0:1][2:4][1:2]) // le45-error {{array section does not specify contiguous storage}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + {} } };