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 @@ -27,6 +27,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" @@ -4737,8 +4738,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 @@ -4748,14 +4750,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 @@ -9778,6 +9778,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.h b/clang/lib/CodeGen/CGOpenMPRuntime.h --- a/clang/lib/CodeGen/CGOpenMPRuntime.h +++ b/clang/lib/CodeGen/CGOpenMPRuntime.h @@ -1625,6 +1625,14 @@ /// Map between the a declaration of a capture and the corresponding base /// pointer address where the runtime returns the device pointers. llvm::DenseMap CaptureDeviceAddrMap; + /// The array of dimension size passed to the runtime library. + SmallVector Dims; + /// The array of array of offsets passed to the runtime library. + SmallVector, 4> Offsets; + /// The array of array of counts passed to the runtime library. + SmallVector, 4> Counts; + /// The array of array of strides passed to the runtime library. + SmallVector, 4> Strides; explicit TargetDataInfo() {} explicit TargetDataInfo(bool RequiresDevicePointerInfo) 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 @@ -7040,6 +7040,10 @@ /// Close is a hint to the runtime to allocate memory close to /// the target device. OMP_MAP_CLOSE = 0x400, + /// 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_DESCRIPTOR = 0x100000000000, /// The 16 MSBs of the flags indicate whether the entry is member of some /// struct/class. OMP_MAP_MEMBER_OF = 0xffff000000000000, @@ -7075,6 +7079,8 @@ using MapBaseValuesArrayTy = SmallVector; using MapValuesArrayTy = SmallVector; using MapFlagsArrayTy = SmallVector; + using MapDimArrayTy = SmallVector; + using MapNonContiguousArrayTy = SmallVector; /// Map between a struct and the its lowest & highest elements which have been /// mapped. @@ -7088,6 +7094,14 @@ Address Base = Address::invalid(); }; + struct StructNonContiguousInfo { + bool IsNonContiguous = false; + MapDimArrayTy Dims; + MapNonContiguousArrayTy Offsets; + MapNonContiguousArrayTy Counts; + MapNonContiguousArrayTy Strides; + }; + private: /// Kind that defines how a device pointer has to be returned. struct MapInfo { @@ -7222,9 +7236,11 @@ /// a flag marking the map as a pointer if requested. Add a flag marking the /// map as the first one of a series of maps that relate to the same map /// expression. - OpenMPOffloadMappingFlags getMapTypeBits( - OpenMPMapClauseKind MapType, ArrayRef MapModifiers, - bool IsImplicit, bool AddPtrFlag, bool AddIsTargetParamFlag) const { + OpenMPOffloadMappingFlags + getMapTypeBits(OpenMPMapClauseKind MapType, + ArrayRef MapModifiers, bool IsImplicit, + bool AddPtrFlag, bool AddIsTargetParamFlag, + bool IsNonContiguous) const { OpenMPOffloadMappingFlags Bits = IsImplicit ? OMP_MAP_IMPLICIT : OMP_MAP_NONE; switch (MapType) { @@ -7260,6 +7276,8 @@ if (llvm::find(MapModifiers, OMPC_MAP_MODIFIER_close) != MapModifiers.end()) Bits |= OMP_MAP_CLOSE; + if (IsNonContiguous) + Bits |= OMP_MAP_DESCRIPTOR; return Bits; } @@ -7312,7 +7330,8 @@ MapBaseValuesArrayTy &BasePointers, MapValuesArrayTy &Pointers, MapValuesArrayTy &Sizes, MapFlagsArrayTy &Types, StructRangeInfoTy &PartialStruct, bool IsFirstComponentList, - bool IsImplicit, bool ForDeviceAddr = false, + bool IsImplicit, StructNonContiguousInfo &NonContigInfo, + bool ForDeviceAddr = false, ArrayRef OverlappedElements = llvm::None) const { // The following summarizes what has to be generated for each map and the @@ -7561,6 +7580,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 = NonContigInfo.IsNonContiguous; + for (; I != CE; ++I) { // If the current component is member of a struct (parent struct) mark it. if (!EncounteredME) { @@ -7579,7 +7604,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 @@ -7597,7 +7625,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 @@ -7653,7 +7684,7 @@ OMP_MAP_MEMBER_OF | getMapTypeBits(MapType, MapModifiers, IsImplicit, /*AddPtrFlag=*/false, - /*AddIsTargetParamFlag=*/false); + /*AddIsTargetParamFlag=*/false, IsNonContiguous); LB = BP; llvm::Value *Size = nullptr; // Do bitcopy of all non-overlapped structure elements. @@ -7677,6 +7708,7 @@ Sizes.push_back(CGF.Builder.CreateIntCast(Size, CGF.Int64Ty, /*isSigned=*/true)); Types.push_back(Flags); + NonContigInfo.Dims.push_back(IsNonContiguous ? DimSize : 1); LB = CGF.Builder.CreateConstGEP(ComponentLB, 1); } BasePointers.push_back(BP.getPointer()); @@ -7688,6 +7720,7 @@ Sizes.push_back( CGF.Builder.CreateIntCast(Size, CGF.Int64Ty, /*isSigned=*/true)); Types.push_back(Flags); + NonContigInfo.Dims.push_back(IsNonContiguous ? DimSize : 1); break; } llvm::Value *Size = getExprTypeSize(I->getAssociatedExpression()); @@ -7696,6 +7729,7 @@ Pointers.push_back(LB.getPointer()); Sizes.push_back( CGF.Builder.CreateIntCast(Size, CGF.Int64Ty, /*isSigned=*/true)); + NonContigInfo.Dims.push_back(IsNonContiguous ? DimSize : 1); // We need to add a pointer flag for each map that comes from the // same expression except for the first one. We also need to signal @@ -7704,7 +7738,7 @@ OpenMPOffloadMappingFlags Flags = getMapTypeBits( MapType, MapModifiers, IsImplicit, !IsExpressionFirstInfo || RequiresReference, - IsCaptureFirstInfo && !RequiresReference); + IsCaptureFirstInfo && !RequiresReference, IsNonContiguous); if (!IsExpressionFirstInfo) { // If we have a PTR_AND_OBJ pair where the OBJ is a pointer as well, @@ -7764,6 +7798,180 @@ IsCaptureFirstInfo = 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; + SmallVector 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(); + AssocExpr->dump(); + const auto *OASE = dyn_cast(AssocExpr); + const auto *AE = dyn_cast(AssocExpr); + + if (AE) { + 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; + } + + 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) + : llvm::ConstantInt::get(CGF.Int64Ty, /*V=*/1); + Count = CGF.Builder.CreateUDiv(CGF.Builder.CreateNUWSub(*DI, Offset), + Stride); + } + } 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) + : llvm::ConstantInt::get(CGF.Int64Ty, 1); + DimProd = CGF.Builder.CreateNUWMul(DimProd, *(DI - 1)); + CurStrides.push_back(CGF.Builder.CreateNUWMul(DimProd, Stride)); + if (DI != DimSizes.end()) + ++DI; + } + + NonContigInfo.Offsets.push_back(CurOffsets); + NonContigInfo.Counts.push_back(CurCounts); + NonContigInfo.Strides.push_back(CurStrides); } /// Return the adjusted map modifiers if the declaration a capture refers to @@ -7944,7 +8152,8 @@ /// index where it occurs is appended to the device pointers info array. void generateAllInfo(MapBaseValuesArrayTy &BasePointers, MapValuesArrayTy &Pointers, MapValuesArrayTy &Sizes, - MapFlagsArrayTy &Types) const { + MapFlagsArrayTy &Types, + CGOpenMPRuntime::TargetDataInfo &TargetDataInfo) const { // We have to process the component lists that relate with the same // declaration in a single chunk so that we can generate the map flags // correctly. Therefore, we organize all lists in a map. @@ -8120,6 +8329,7 @@ MapValuesArrayTy CurSizes; MapFlagsArrayTy CurTypes; StructRangeInfoTy PartialStruct; + StructNonContiguousInfo CurNonContigInfo; for (const MapInfo &L : M.second) { assert(!L.Components.empty() && @@ -8127,10 +8337,13 @@ // Remember the current base pointer index. unsigned CurrentBasePointersIdx = CurBasePointers.size(); - generateInfoForComponentList( - L.MapType, L.MapModifiers, L.Components, CurBasePointers, - CurPointers, CurSizes, CurTypes, PartialStruct, - IsFirstComponentList, L.IsImplicit, L.ForDeviceAddr); + CurNonContigInfo.IsNonContiguous = + L.Components.back().isNonContiguous(); + generateInfoForComponentList(L.MapType, L.MapModifiers, L.Components, + CurBasePointers, CurPointers, CurSizes, + CurTypes, PartialStruct, + IsFirstComponentList, L.IsImplicit, + CurNonContigInfo, L.ForDeviceAddr); // If this entry relates with a device pointer, set the relevant // declaration and add the 'return pointer' flag. @@ -8193,6 +8406,14 @@ Pointers.append(CurPointers.begin(), CurPointers.end()); Sizes.append(CurSizes.begin(), CurSizes.end()); Types.append(CurTypes.begin(), CurTypes.end()); + TargetDataInfo.Dims.append(CurNonContigInfo.Dims.begin(), + CurNonContigInfo.Dims.end()); + TargetDataInfo.Offsets.append(CurNonContigInfo.Offsets.begin(), + CurNonContigInfo.Offsets.end()); + TargetDataInfo.Counts.append(CurNonContigInfo.Counts.begin(), + CurNonContigInfo.Counts.end()); + TargetDataInfo.Strides.append(CurNonContigInfo.Strides.begin(), + CurNonContigInfo.Strides.end()); } } @@ -8243,22 +8464,30 @@ MapValuesArrayTy CurSizes; MapFlagsArrayTy CurTypes; StructRangeInfoTy PartialStruct; + StructNonContiguousInfo CurNonContigInfo; for (const MapInfo &L : M.second) { assert(!L.Components.empty() && "Not expecting declaration with no component lists."); - generateInfoForComponentList( - L.MapType, L.MapModifiers, L.Components, CurBasePointers, - CurPointers, CurSizes, CurTypes, PartialStruct, - IsFirstComponentList, L.IsImplicit, L.ForDeviceAddr); + generateInfoForComponentList(L.MapType, L.MapModifiers, L.Components, + CurBasePointers, CurPointers, CurSizes, + CurTypes, PartialStruct, + IsFirstComponentList, L.IsImplicit, + CurNonContigInfo, L.ForDeviceAddr); IsFirstComponentList = false; } // 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()) + if (PartialStruct.Base.isValid()) { + // Make sure Dims have the same size as BP, P, Sizes, and Types. + // Put 0 here to make sure that `emitOffloadingArrays` use it + // to skip processing this one. (OpenMP do not allow non-contigous for + // declare mapper) + CurNonContigInfo.Dims.push_back(0); emitCombinedEntry(BasePointers, Pointers, Sizes, Types, CurTypes, PartialStruct); + } // We need to append the results of this capture to what we already have. BasePointers.append(CurBasePointers.begin(), CurBasePointers.end()); @@ -8503,6 +8732,7 @@ return *It == FD1; }); } + StructNonContiguousInfo NonContigInfo; // Associated with a capture, because the mapping flags depend on it. // Go through all of the elements with the overlapped elements. @@ -8518,7 +8748,7 @@ bool IsFirstComponentList = true; generateInfoForComponentList( MapType, MapModifiers, Components, BasePointers, Pointers, Sizes, - Types, PartialStruct, IsFirstComponentList, IsImplicit, + Types, PartialStruct, IsFirstComponentList, IsImplicit, NonContigInfo, /*ForDeviceAddr=*/false, OverlappedComponents); } // Go through other elements without overlapped elements. @@ -8534,7 +8764,7 @@ generateInfoForComponentList(MapType, MapModifiers, Components, BasePointers, Pointers, Sizes, Types, PartialStruct, IsFirstComponentList, - IsImplicit); + IsImplicit, NonContigInfo); IsFirstComponentList = false; } } @@ -8563,10 +8793,11 @@ !Res || *Res != OMPDeclareTargetDeclAttr::MT_Link) continue; StructRangeInfoTy PartialStruct; + StructNonContiguousInfo NonContigInfo; generateInfoForComponentList( C->getMapType(), C->getMapTypeModifiers(), L.second, BasePointers, Pointers, Sizes, Types, PartialStruct, - /*IsFirstComponentList=*/true, C->isImplicit()); + /*IsFirstComponentList=*/true, C->isImplicit(), NonContigInfo); assert(!PartialStruct.Base.isValid() && "No partial structs for declare target link expected."); } @@ -8659,6 +8890,78 @@ }; } // anonymous namespace +static void emitNonContiguousDescriptor( + CodeGenFunction &CGF, + MappableExprsHandler::MapBaseValuesArrayTy &BasePointers, + MappableExprsHandler::MapValuesArrayTy &Pointers, + MappableExprsHandler::MapValuesArrayTy &Sizes, + MappableExprsHandler::MapFlagsArrayTy &MapTypes, + CGOpenMPRuntime::TargetDataInfo &Info) { + CodeGenModule &CGM = CGF.CGM; + + // 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=*/true); + 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 }; + // The reason we need two index variable here is because 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 = Info.Dims.size(); I < E; ++I) { + // Skip emitting ir if dimension is 1. + if (Info.Dims[I] == 1) + continue; + // For supporting stride, each descriptor has a dummy dimension, hence, + // adding one to the original dimension size. + llvm::APInt Size(/*numBits=*/32, Info.Dims[I]); + QualType ArrayTy = + C.getConstantArrayType(DimTy, Size, nullptr, ArrayType::Normal, 0); + Address DimsAddr = CGF.CreateMemTemp(ArrayTy, "dims"); + for (unsigned II = 0, EE = Info.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(Info.Offsets[L][RevIdx], OffsetLVal); + // Count + LValue CountLVal = CGF.EmitLValueForField( + DimsLVal, *std::next(RD->field_begin(), CountFD)); + CGF.EmitStoreOfScalar(Info.Counts[L][RevIdx], CountLVal); + // Stride + LValue StrideLVal = CGF.EmitLValueForField( + DimsLVal, *std::next(RD->field_begin(), StrideFD)); + CGF.EmitStoreOfScalar(Info.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. @@ -8668,7 +8971,8 @@ MappableExprsHandler::MapValuesArrayTy &Pointers, MappableExprsHandler::MapValuesArrayTy &Sizes, MappableExprsHandler::MapFlagsArrayTy &MapTypes, - CGOpenMPRuntime::TargetDataInfo &Info) { + CGOpenMPRuntime::TargetDataInfo &Info, + bool IsNonContiguous = false) { CodeGenModule &CGM = CGF.CGM; ASTContext &Ctx = CGF.getContext(); @@ -8711,8 +9015,15 @@ // We expect all the sizes to be constant, so we collect them to create // a constant array. SmallVector ConstSizes; - for (llvm::Value *S : Sizes) - ConstSizes.push_back(cast(S)); + for (unsigned I = 0, E = Sizes.size(); I < E; ++I) { + if (IsNonContiguous && + (MapTypes[I] & MappableExprsHandler::OMP_MAP_DESCRIPTOR)) { + ConstSizes.push_back( + llvm::ConstantInt::get(CGF.Int64Ty, Info.Dims[I])); + } else { + ConstSizes.push_back(cast(Sizes[I])); + } + } auto *SizesArrayInit = llvm::ConstantArray::get( llvm::ArrayType::get(CGM.Int64Ty, ConstSizes.size()), ConstSizes); @@ -8776,6 +9087,12 @@ } } } + + if (!IsNonContiguous || Info.Offsets.empty() || Info.NumberOfPtrs == 0) + return; + + emitNonContiguousDescriptor(CGF, BasePointers, Pointers, Sizes, MapTypes, + Info); } /// Emit the arguments to be passed to the runtime library based on the @@ -10138,10 +10455,11 @@ // Get map clause information. MappableExprsHandler MCHandler(D, CGF); - MCHandler.generateAllInfo(BasePointers, Pointers, Sizes, MapTypes); + MCHandler.generateAllInfo(BasePointers, Pointers, Sizes, MapTypes, Info); // Fill up the arrays and create the arguments. - emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, Info); + emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, Info, + /*IsNonContiguous=*/true); llvm::Value *BasePointersArrayArg = nullptr; llvm::Value *PointersArrayArg = nullptr; @@ -10379,11 +10697,12 @@ // Get map clause information. MappableExprsHandler MEHandler(D, CGF); - MEHandler.generateAllInfo(BasePointers, Pointers, Sizes, MapTypes); - TargetDataInfo Info; + MEHandler.generateAllInfo(BasePointers, Pointers, Sizes, MapTypes, Info); + // Fill up the arrays and create the arguments. - emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, Info); + emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, Info, + /*IsNonContiguous=*/true); emitOffloadingArraysArgument(CGF, Info.BasePointersArray, Info.PointersArray, Info.SizesArray, Info.MapTypesArray, Info); 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. @@ -3535,6 +3535,7 @@ if (isOpenMPTargetExecutionDirective(DKind)) { OMPClauseMappableExprCommon::MappableExprComponentList CurComponents; if (!checkMapClauseExpressionBase(SemaRef, E, CurComponents, OMPC_map, + Stack->getCurrentDirective(), /*NoDiagnose=*/true)) return; const auto *VD = cast( @@ -16597,11 +16598,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; @@ -16626,7 +16630,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; } @@ -16698,7 +16702,7 @@ AllowWholeSizeArraySection = false; // Record the component. - Components.emplace_back(ME, FD); + Components.emplace_back(ME, FD, IsNonContiguous); return RelevantExpr || Visit(E); } @@ -16736,7 +16740,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); } @@ -16775,6 +16779,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. @@ -16784,6 +16795,9 @@ return false; } + if (IsPointer) + AllowAnotherPtr = false; + if (const auto *TE = dyn_cast(E)) { Expr::EvalResult ResultR; Expr::EvalResult ResultL; @@ -16809,14 +16823,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()); } @@ -16829,7 +16843,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()); } @@ -16845,7 +16859,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"); @@ -16856,7 +16870,7 @@ 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 VisitStmt(Stmt *) { @@ -16867,10 +16881,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 @@ -16882,13 +16896,31 @@ 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) { + if (!OASE->getLength()) + SemaRef.Diag(ELoc, diag::err_array_section_does_not_specify_length) + << ERange; + else + break; + } + } + } return Checker.getFoundBase(); + } return nullptr; } @@ -17365,7 +17397,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; @@ -18613,8 +18646,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()) @@ -18665,8 +18698,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()) @@ -18732,7 +18765,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 @@ -12531,10 +12531,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); } @@ -12648,10 +12648,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); } @@ -12698,10 +12698,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); } @@ -12748,10 +12748,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); } @@ -12792,8 +12792,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); } @@ -12835,8 +12835,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 @@ -6597,6 +6597,7 @@ Record.push_back(N); for (auto &M : C->all_components()) { Record.AddStmt(M.getAssociatedExpression()); + Record.writeBool(M.isNonContiguous()); Record.AddDeclRef(M.getAssociatedDeclaration()); } } @@ -6621,6 +6622,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 @@ -5,6 +5,14 @@ // RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=50 -ast-print %s | FileCheck %s // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s + +// RUN: %clang_cc1 -DOMP5 -verify -fopenmp -fopenmp-version=50 -ast-print %s | FileCheck %s --check-prefix=OMP5 +// RUN: %clang_cc1 -DOMP5 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s --check-prefix=OMP5 + +// RUN: %clang_cc1 -DOMP5 -verify -fopenmp-simd -fopenmp-version=50 -ast-print %s | FileCheck %s --check-prefix=OMP5 +// RUN: %clang_cc1 -DOMP5 -fopenmp-simd -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s +// RUN: %clang_cc1 -DOMP5 -fopenmp-simd -fopenmp-version=50 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s --check-prefix=OMP5 // expected-no-diagnostics #ifndef HEADER @@ -21,10 +29,61 @@ #pragma omp target update from(b, ([a][targ])p) if(l<5) device(l-1) nowait depend(inout:l) +#ifdef OMP5 + 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]) #pragma omp target update from(arr[2][0:1:2]) + +// OMP5: marr[10][10][10]; +// OMP5-NEXT: #pragma omp target update to(marr[2][0:2][0:2]) +// OMP5-NEXT: #pragma omp target update from(marr[2][0:2][0:2]) +// OMP5-NEXT: #pragma omp target update from(marr[:2][0:2][0:2:1]) +// OMP5-NEXT: #pragma omp target update to(marr[:l][:l][l:]) +// OMP5-NEXT: #pragma omp target update to(marr[:2][:1][:]) +// OMP5-NEXT: #pragma omp target update from(marr[:2][:1][:]) +// OMP5-NEXT: #pragma omp target update to(marr[:2][:][:1]) +// OMP5-NEXT: #pragma omp target update from(marr[:2][:][:1]) +// OMP5-NEXT: #pragma omp target update to(marr[:2][:][1:]) +// OMP5-NEXT: #pragma omp target update from(marr[:2][:][1:]) +// OMP5-NEXT: #pragma omp target update to(marr[:1][3:2][:2]) +// OMP5-NEXT: #pragma omp target update from(marr[:1][3:2][:2]) +// OMP5-NEXT: #pragma omp target update to(marr[:1][:2][0]) +// OMP5-NEXT: #pragma omp target update from(marr[:1][:2][0]) +// OMP5-NEXT: int arr[100][100]; +// OMP5-NEXT: #pragma omp target update to(arr[2][0:1:2]) +// OMP5-NEXT: #pragma omp target update from(arr[2][0:1:2]) +#endif return a + targ + (T)b; } // CHECK: static T a, *p; @@ -42,9 +101,6 @@ // 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: 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]) int main(int argc, char **argv) { static int a; @@ -58,10 +114,47 @@ // CHECK-NEXT: #pragma omp target update to(a) if(f > 0.) device(n) nowait depend(in : n) #pragma omp target update from(f) if(f<0.0) device(n+1) nowait depend(in:n) // CHECK-NEXT: #pragma omp target update from(f) if(f < 0.) device(n + 1) nowait depend(in : n) + +#ifdef OMP5 +float marr[10][10][10]; +// OMP5: marr[10][10][10]; +#pragma omp target update to(marr[2][0:2][0:2]) +// OMP5-NEXT: #pragma omp target update to(marr[2][0:2][0:2]) +#pragma omp target update from(marr[2][0:2][0:2]) +// OMP5-NEXT: #pragma omp target update from(marr[2][0:2][0:2]) +#pragma omp target update to(marr[:n][:n][n:]) +// OMP5: #pragma omp target update to(marr[:n][:n][n:]) +#pragma omp target update from(marr[:2][:1][:]) +// OMP5-NEXT: #pragma omp target update from(marr[:2][:1][:]) +#pragma omp target update to(marr[:2][:][:1]) +// OMP5-NEXT: #pragma omp target update to(marr[:2][:][:1]) +#pragma omp target update from(marr[:2][:][:1]) +// OMP5-NEXT: #pragma omp target update from(marr[:2][:][:1]) +#pragma omp target update to(marr[:2][:][1:]) +// OMP5-NEXT: #pragma omp target update to(marr[:2][:][1:]) +#pragma omp target update from(marr[:2][:][1:]) +// OMP5-NEXT: #pragma omp target update from(marr[:2][:][1:]) +#pragma omp target update to(marr[:1][3:2][:2]) +// OMP5-NEXT: #pragma omp target update to(marr[:1][3:2][:2]) +#pragma omp target update from(marr[:1][3:2][:2]) +// OMP5-NEXT: #pragma omp target update from(marr[:1][3:2][:2]) +#pragma omp target update to(marr[:1][:2][0]) +// OMP5-NEXT: #pragma omp target update to(marr[:1][:2][0]) +#pragma omp target update from(marr[:1][:2][0]) +// OMP5-NEXT: #pragma omp target update from(marr[:1][:2][0]) +#pragma omp target update to(marr[:2:][0:2][0:2:1]) +// OMP5-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]) +// OMP5-NEXT: #pragma omp target update from(marr[:2:][0:2][0:2:1]) +#pragma omp target update to(marr[:2:][:2:][0:2:1]) +// OMP5-NEXT: #pragma omp target update to(marr[:2:][:2:][0:2:1]) +#pragma omp target update from(marr[:2:][:2:][0:2:1]) +// OMP5-NEXT: #pragma omp target update from(marr[:2:][:2:][0:2:1]) #pragma omp target update to(argv[2][0:1:2]) -// CHECK-NEXT: #pragma omp target update to(argv[2][0:1:2]) +// OMP5-NEXT: #pragma omp target update to(argv[2][0:1:2]) #pragma omp target update from(argv[2][0:1:2]) -// CHECK-NEXT: #pragma omp target update from(argv[2][0:1:2]) +// OMP5-NEXT: #pragma omp target update from(argv[2][0:1:2]) +#endif 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 @@ -1059,5 +1059,495 @@ #pragma omp target update from(([sa][5])f) } +#endif +///==========================================================================/// +// RUN: %clang_cc1 -DCK19 -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 CK19 --check-prefix CK19-64 +// RUN: %clang_cc1 -DCK19 -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 CK19 --check-prefix CK19-64 +// RUN: %clang_cc1 -DCK19 -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 CK19 --check-prefix CK19-32 +// RUN: %clang_cc1 -DCK19 -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 CK19 --check-prefix CK19-32 + +// RUN: %clang_cc1 -DCK19 -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 -DCK19 -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 -DCK19 -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 -DCK19 -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 CK19 + +// CK19: [[STRUCT_DESCRIPTOR:%.+]] = type { i64, i64, i64 } + +// CK19: [[MSIZE:@.+]] = {{.+}}constant [3 x i64] [i64 4, i64 4, i64 3] +// CK19: [[MTYPE:@.+]] = {{.+}}constant [3 x i64] [i64 17592186044449, i64 33, i64 17592186044449] + +// CK19-LABEL: _Z3foo +void foo(int arg) { + int arr[3][4][5], x; + float farr[4][3]; + + // CK19: [[DIMS:%.+]] = alloca [4 x [[STRUCT_DESCRIPTOR]]], + // CK19: [[DIMS_2:%.+]] = alloca [3 x [[STRUCT_DESCRIPTOR]]], + // CK19: [[ARRAY_IDX:%.+]] = getelementptr inbounds [3 x [4 x [5 x i32]]], [3 x [4 x [5 x i32]]]* [[ARR:%.+]], {{.+}} 0, {{.+}} 0 + // CK19: [[ARRAY_DECAY:%.+]] = getelementptr inbounds [4 x [5 x i32]], [4 x [5 x i32]]* [[ARRAY_IDX]], {{.+}} 0, {{.+}} 0 + // CK19: [[ARRAY_IDX_1:%.+]] = getelementptr inbounds [5 x i32], [5 x i32]* [[ARRAY_DECAY]], {{.+}} + // CK19: [[ARRAY_DECAY_2:%.+]] = getelementptr inbounds [5 x i32], [5 x i32]* [[ARRAY_IDX_1]], {{.+}} 0, {{.+}} 0 + // CK19: [[ARRAY_IDX_3:%.+]] = getelementptr inbounds {{.+}}, {{.+}}* [[ARRAY_DECAY_2]], {{.+}} 1 + // CK19: [[SUB:%.+]] = sub nuw i64 4, [[ARG_ADDR:%.+]] + // CK19: [[LEN:%.+]] = udiv {{.+}} [[SUB]], 1 + // CK19: [[ARRAY_IDX_4:%.+]] = getelementptr inbounds [4 x [3 x float]], [4 x [3 x float]]* [[FARR:%.+]], {{.+}} 0, {{.+}} 0 + // CK19: [[ARRAY_DECAY_5:%.+]] = getelementptr inbounds [3 x float], [3 x float]* [[ARRAY_IDX_4]], {{.+}} 0, {{.+}} 0 + // CK19: [[ARRAY_IDX_6:%.+]] = getelementptr inbounds float, float* [[ARRAY_DECAY_5:%.+]], {{.+}} 1 + // CK19: [[BP0:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BP:%.+]], i{{.+}} 0, i{{.+}} 0 + // CK19: [[P0:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[P:%.+]], i{{.+}} 0, i{{.+}} 0 + // CK19: [[DIM_1:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 0 + // CK19: [[OFFSET:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 0 + // CK19: store i64 0, i64* [[OFFSET]], + // CK19: [[COUNT:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 1 + // CK19: store i64 2, i64* [[COUNT]], + // CK19: [[STRIDE:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 2 + // CK19: store i64 80, i64* [[STRIDE]], + // CK19: [[DIM_2:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 1 + // CK19: [[OFFSET_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 0 + // CK19: store i64 [[ARG:%.+]], i64* [[OFFSET_2]], + // CK19: [[COUNT_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 1 + // CK19: store i64 [[LEN]], i64* [[COUNT_2]], + // CK19: [[STRIDE_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 2 + // CK19: store i64 20, i64* [[STRIDE_2]], + // CK19: [[DIM_3:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 2 + // CK19: [[OFFSET_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 0 + // CK19: store i64 1, i64* [[OFFSET_3]], + // CK19: [[COUNT_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 1 + // CK19: store i64 4, i64* [[COUNT_3]], + // CK19: [[STRIDE_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 2 + // CK19: store i64 4, i64* [[STRIDE_3]], + // CK19: [[DIM_4:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 3 + // CK19: [[OFFSET_4:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 0 + // CK19: store i64 0, i64* [[OFFSET_4]], + // CK19: [[COUNT_4:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 1 + // CK19: store i64 1, i64* [[COUNT_4]], + // CK19: [[STRIDE_4:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 2 + // CK19: store i64 4, i64* [[STRIDE_4]], + // CK19: [[PC0:%.+]] = bitcast [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]] to i8* + // CK19: [[PTRS:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* %.offload_ptrs, i32 0, i32 0 + // CK19: store i8* [[PC0]], i8** [[PTRS]], + // CK19: [[DIM_5:%.+]] = getelementptr inbounds [3 x [[STRUCT_DESCRIPTOR]]], [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS_2]], {{.+}} 0, {{.+}} 0 + // CK19: [[OFFSET_2_1:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_5]], {{.+}} 0, {{.+}} 0 + // CK19: store i64 0, i64* [[OFFSET_2_1]], + // CK19: [[COUNT_2_1:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_5]], {{.+}} 0, {{.+}} 1 + // CK19: store i64 2, i64* [[COUNT_2_1]], + // CK19: [[STRIDE_2_1:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_5]], {{.+}} 0, {{.+}} 2 + // CK19: store i64 12, i64* [[STRIDE_2_1]], + // CK19: [[DIM_6:%.+]] = getelementptr inbounds [3 x [[STRUCT_DESCRIPTOR]]], [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS_2]], {{.+}} 0, {{.+}} 1 + // CK19: [[OFFSET_2_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_6]], {{.+}} 0, {{.+}} 0 + // CK19: store i64 1, i64* [[OFFSET_2_2]], + // CK19: [[COUNT_2_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_6]], {{.+}} 0, {{.+}} 1 + // CK19: store i64 2, i64* [[COUNT_2_2]], + // CK19: [[STRIDE_2_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_6]], {{.+}} 0, {{.+}} 2 + // CK19: store i64 4, i64* [[STRIDE_2_2]], + // CK19: [[DIM_7:%.+]] = getelementptr inbounds [3 x [[STRUCT_DESCRIPTOR]]], [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS_2]], {{.+}} 0, {{.+}} 2 + // CK19: [[OFFSET_2_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_7]], {{.+}} 0, {{.+}} 0 + // CK19: store i64 0, i64* [[OFFSET_2_3]], + // CK19: [[COUNT_2_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_7]], {{.+}} 0, {{.+}} 1 + // CK19: store i64 1, i64* [[COUNT_2_3]], + // CK19: [[STRIDE_2_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_7]], {{.+}} 0, {{.+}} 2 + // CK19: store i64 4, i64* [[STRIDE_2_3]], + // CK19: [[PC1:%.+]] = bitcast [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS_2]] to i8* + // CK19: [[PTRS_2:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* %.offload_ptrs, i32 0, i32 2 + // CK19: store i8* [[PC1]], i8** [[PTRS_2]], + // CK19-DAG: call void @__tgt_target_data_update(i64 -1, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[MSIZE]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[MTYPE]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + +#pragma omp target update to(arr[0:2][arg:][1:4], x, farr[0:2][1:2]) + { ++arg; } +} + +#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(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(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(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(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(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 #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 @@ -1,8 +1,8 @@ -// RUN: %clang_cc1 -verify=expected,omp45 -fopenmp -fopenmp-version=45 -ferror-limit 100 -o - -std=c++11 %s -Wuninitialized -// RUN: %clang_cc1 -verify=expected,omp50 -fopenmp -fopenmp-version=50 -ferror-limit 100 -o - -std=c++11 %s -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le45 -fopenmp -fopenmp-version=45 -ferror-limit 100 %s -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le50 -fopenmp -fopenmp-version=50 -ferror-limit 100 %s -Wuninitialized -// RUN: %clang_cc1 -verify=expected,omp45 -fopenmp-simd -fopenmp-version=45 -ferror-limit 100 -o - -std=c++11 %s -Wuninitialized -// RUN: %clang_cc1 -verify=expected,omp50 -fopenmp-simd -fopenmp-version=50 -ferror-limit 100 -o - -std=c++11 %s -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd -fopenmp-version=45 -ferror-limit 100 %s -Wuninitialized +// RUN: %clang_cc1 -verify=expected,le50 -fopenmp-simd -fopenmp-version=50 -ferror-limit 100 %s -Wuninitialized void xxx(int argc) { int x; // expected-note {{initialize the variable 'x' to silence this warning}} @@ -39,10 +39,33 @@ foo(); } + double marr[10][5][10]; +#pragma omp target update to(marr[0:2][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'}} + {} +#pragma omp target update from(marr[0:2][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'}} + +#pragma omp target update to(marr[0:][1:2:2][1:2]) // le50-error {{array section does not specify length for outermost dimension}} le45-error {{expected ']'}} le45-note {{to match this '['}} le45-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]) // le50-error {{array section does not specify length for outermost dimension}} le45-error {{expected ']'}} le45-note {{to match this '['}} le45-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]) // 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'}} + {} +#pragma omp target update from(arr[0:2][2:4][:2][1]) // 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'}} + + double ***dptr; +#pragma omp target update to(dptr[0:2][2:4][1:2]) // le45-error {{array section does not specify contiguous storage}} le50-error 2 {{section length is unspecified and cannot be inferred because subscripted value is an array of unknown bound}} le45-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]) // le45-error {{array section does not specify contiguous storage}} le50-error 2 {{section length is unspecified and cannot be inferred because subscripted value is an array of unknown bound}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + int iarr[5][5]; -#pragma omp target update to(iarr[0:][1:2:-1]) // omp50-error {{section stride is evaluated to a non-positive value -1}} omp45-error {{expected ']'}} omp45-note {{to match this '['}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(iarr[0:][1:2:-1]) // le50-error {{section stride is evaluated to a non-positive value -1}} le45-error {{expected ']'}} le45-note {{to match this '['}} expected-error {{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]) // le50-error {{section stride is evaluated to a non-positive value -1}} le45-error {{expected ']'}} le45-note {{to match this '['}} expected-error {{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]) // le50-error {{array section does not specify length for outermost dimension}} le45-error {{expected ']'}} le45-note {{to match this '['}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} le45-error {{expected expression}} {} -#pragma omp target update from(iarr[0:][1:2:-1]) // omp50-error {{section stride is evaluated to a non-positive value -1}} omp45-error {{expected ']'}} omp45-note {{to match this '['}} expected-error {{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]) // le50-error {{array section does not specify length for outermost dimension}} le45-error {{expected ']'}} le45-note {{to match this '['}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} le45-error {{expected expression}} 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'}} + {} } };