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