[clang] c2aa543 - [OPENMP50]Codegen for array shaping expression in map clauses.
Alexey Bataev via cfe-commits
cfe-commits at lists.llvm.org
Tue Mar 31 16:11:34 PDT 2020
Author: Alexey Bataev
Date: 2020-03-31T19:06:49-04:00
New Revision: c2aa543237843fa7b7c0191b6685062b3512f245
URL: https://github.com/llvm/llvm-project/commit/c2aa543237843fa7b7c0191b6685062b3512f245
DIFF: https://github.com/llvm/llvm-project/commit/c2aa543237843fa7b7c0191b6685062b3512f245.diff
LOG: [OPENMP50]Codegen for array shaping expression in map clauses.
Added codegen support for array shaping operations in map/to/from
clauses.
Added:
Modified:
clang/lib/CodeGen/CGOpenMPRuntime.cpp
clang/lib/Sema/SemaOpenMP.cpp
clang/test/OpenMP/target_data_ast_print.cpp
clang/test/OpenMP/target_map_codegen.cpp
clang/test/OpenMP/target_map_messages.cpp
clang/test/OpenMP/target_update_ast_print.cpp
clang/test/OpenMP/target_update_codegen.cpp
Removed:
################################################################################
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 31fdc320d698..6642851a56bc 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -7448,6 +7448,20 @@ class MappableExprsHandler {
llvm::Value *getExprTypeSize(const Expr *E) const {
QualType ExprTy = E->getType().getCanonicalType();
+ // Calculate the size for array shaping expression.
+ if (const auto *OAE = dyn_cast<OMPArrayShapingExpr>(E)) {
+ llvm::Value *Size =
+ CGF.getTypeSize(OAE->getBase()->getType()->getPointeeType());
+ for (const Expr *SE : OAE->getDimensions()) {
+ llvm::Value *Sz = CGF.EmitScalarExpr(SE);
+ Sz = CGF.EmitScalarConversion(Sz, SE->getType(),
+ CGF.getContext().getSizeType(),
+ SE->getExprLoc());
+ Size = CGF.Builder.CreateNUWMul(Size, Sz);
+ }
+ return Size;
+ }
+
// Reference types are ignored for mapping purposes.
if (const auto *RefTy = ExprTy->getAs<ReferenceType>())
ExprTy = RefTy->getPointeeType().getCanonicalType();
@@ -7779,6 +7793,7 @@ class MappableExprsHandler {
const Expr *AssocExpr = I->getAssociatedExpression();
const auto *AE = dyn_cast<ArraySubscriptExpr>(AssocExpr);
const auto *OASE = dyn_cast<OMPArraySectionExpr>(AssocExpr);
+ const auto *OAShE = dyn_cast<OMPArrayShapingExpr>(AssocExpr);
if (isa<MemberExpr>(AssocExpr)) {
// The base is the 'this' pointer. The content of the pointer is going
@@ -7788,6 +7803,11 @@ class MappableExprsHandler {
(OASE &&
isa<CXXThisExpr>(OASE->getBase()->IgnoreParenImpCasts()))) {
BP = CGF.EmitOMPSharedLValue(AssocExpr).getAddress(CGF);
+ } else if (OAShE &&
+ isa<CXXThisExpr>(OAShE->getBase()->IgnoreParenCasts())) {
+ BP = Address(
+ CGF.EmitScalarExpr(OAShE->getBase()),
+ CGF.getContext().getTypeAlignInChars(OAShE->getBase()->getType()));
} else {
// The base is the reference to the variable.
// BP = &Var.
@@ -7870,9 +7890,12 @@ class MappableExprsHandler {
// types.
const auto *OASE =
dyn_cast<OMPArraySectionExpr>(I->getAssociatedExpression());
+ const auto *OAShE =
+ dyn_cast<OMPArrayShapingExpr>(I->getAssociatedExpression());
const auto *UO = dyn_cast<UnaryOperator>(I->getAssociatedExpression());
const auto *BO = dyn_cast<BinaryOperator>(I->getAssociatedExpression());
bool IsPointer =
+ OAShE ||
(OASE && OMPArraySectionExpr::getBaseOriginalType(OASE)
.getCanonicalType()
->isAnyPointerType()) ||
@@ -7890,8 +7913,15 @@ class MappableExprsHandler {
isa<BinaryOperator>(Next->getAssociatedExpression())) &&
"Unexpected expression");
- Address LB = CGF.EmitOMPSharedLValue(I->getAssociatedExpression())
- .getAddress(CGF);
+ Address LB = Address::invalid();
+ if (OAShE) {
+ LB = Address(CGF.EmitScalarExpr(OAShE->getBase()),
+ CGF.getContext().getTypeAlignInChars(
+ OAShE->getBase()->getType()));
+ } else {
+ LB = CGF.EmitOMPSharedLValue(I->getAssociatedExpression())
+ .getAddress(CGF);
+ }
// If this component is a pointer inside the base struct then we don't
// need to create any entry for it - it will be combined with the object
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index f9e8e3d6ccc8..7d2ae172fe4d 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -1943,7 +1943,8 @@ bool Sema::isOpenMPCapturedByRef(const ValueDecl *D, unsigned Level,
if (isa<ArraySubscriptExpr>(EI->getAssociatedExpression()) ||
isa<OMPArraySectionExpr>(EI->getAssociatedExpression()) ||
- isa<MemberExpr>(EI->getAssociatedExpression())) {
+ isa<MemberExpr>(EI->getAssociatedExpression()) ||
+ isa<OMPArrayShapingExpr>(EI->getAssociatedExpression())) {
IsVariableAssociatedWithSection = true;
// There is nothing more we need to know about this variable.
return true;
@@ -3225,7 +3226,7 @@ class DSAAttrChecker final : public StmtVisitor<DSAAttrChecker, void> {
StackComponents,
OpenMPClauseKind) {
// Variable is used if it has been marked as an array, array
- // section or the variable iself.
+ // section, array shaping or the variable iself.
return StackComponents.size() == 1 ||
std::all_of(
std::next(StackComponents.rbegin()),
@@ -3236,6 +3237,8 @@ class DSAAttrChecker final : public StmtVisitor<DSAAttrChecker, void> {
nullptr &&
(isa<OMPArraySectionExpr>(
MC.getAssociatedExpression()) ||
+ isa<OMPArrayShapingExpr>(
+ MC.getAssociatedExpression()) ||
isa<ArraySubscriptExpr>(
MC.getAssociatedExpression()));
});
@@ -3393,8 +3396,10 @@ class DSAAttrChecker final : public StmtVisitor<DSAAttrChecker, void> {
// Do both expressions have the same kind?
if (CCI->getAssociatedExpression()->getStmtClass() !=
SC.getAssociatedExpression()->getStmtClass())
- if (!(isa<OMPArraySectionExpr>(
- SC.getAssociatedExpression()) &&
+ if (!((isa<OMPArraySectionExpr>(
+ SC.getAssociatedExpression()) ||
+ isa<OMPArrayShapingExpr>(
+ SC.getAssociatedExpression())) &&
isa<ArraySubscriptExpr>(
CCI->getAssociatedExpression())))
return false;
@@ -16284,6 +16289,15 @@ class MapBaseChecker final : public StmtVisitor<MapBaseChecker, bool> {
Components.emplace_back(OASE, nullptr);
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);
+
+ return Visit(Base->IgnoreParenImpCasts());
+ }
+
bool VisitUnaryOperator(UnaryOperator *UO) {
if (SemaRef.getLangOpts().OpenMP < 50 || !UO->isLValue() ||
UO->getOpcode() != UO_Deref) {
@@ -16409,9 +16423,11 @@ static bool checkMapConflicts(
// variable in map clauses of the same construct.
if (CurrentRegionOnly &&
(isa<ArraySubscriptExpr>(CI->getAssociatedExpression()) ||
- isa<OMPArraySectionExpr>(CI->getAssociatedExpression())) &&
+ isa<OMPArraySectionExpr>(CI->getAssociatedExpression()) ||
+ isa<OMPArrayShapingExpr>(CI->getAssociatedExpression())) &&
(isa<ArraySubscriptExpr>(SI->getAssociatedExpression()) ||
- isa<OMPArraySectionExpr>(SI->getAssociatedExpression()))) {
+ isa<OMPArraySectionExpr>(SI->getAssociatedExpression()) ||
+ isa<OMPArrayShapingExpr>(SI->getAssociatedExpression()))) {
SemaRef.Diag(CI->getAssociatedExpression()->getExprLoc(),
diag::err_omp_multiple_array_items_in_map_clause)
<< CI->getAssociatedExpression()->getSourceRange();
@@ -16443,6 +16459,9 @@ static bool checkMapConflicts(
const Expr *E = OASE->getBase()->IgnoreParenImpCasts();
Type =
OMPArraySectionExpr::getBaseOriginalType(E).getCanonicalType();
+ } else if (const auto *OASE = dyn_cast<OMPArrayShapingExpr>(
+ SI->getAssociatedExpression())) {
+ Type = OASE->getBase()->getType()->getPointeeType();
}
if (Type.isNull() || Type->isAnyPointerType() ||
checkArrayExpressionDoesNotReferToWholeSize(
@@ -16905,6 +16924,7 @@ static void checkMappableExpressionList(
QualType Type;
auto *ASE = dyn_cast<ArraySubscriptExpr>(VE->IgnoreParens());
auto *OASE = dyn_cast<OMPArraySectionExpr>(VE->IgnoreParens());
+ auto *OAShE = dyn_cast<OMPArrayShapingExpr>(VE->IgnoreParens());
if (ASE) {
Type = ASE->getType().getNonReferenceType();
} else if (OASE) {
@@ -16915,6 +16935,8 @@ static void checkMappableExpressionList(
else
Type = BaseType->getPointeeType();
Type = Type.getNonReferenceType();
+ } else if (OAShE) {
+ Type = OAShE->getBase()->getType()->getPointeeType();
} else {
Type = VE->getType();
}
diff --git a/clang/test/OpenMP/target_data_ast_print.cpp b/clang/test/OpenMP/target_data_ast_print.cpp
index fa67c1834aa4..fcd6e928655c 100644
--- a/clang/test/OpenMP/target_data_ast_print.cpp
+++ b/clang/test/OpenMP/target_data_ast_print.cpp
@@ -1,10 +1,10 @@
-// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -ast-print %s | FileCheck %s
-// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -emit-pch -o %t %s
-// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=50 -ast-print %s | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
-// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -ast-print %s | FileCheck %s
-// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -emit-pch -o %t %s
-// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=50 -ast-print %s | FileCheck %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
// expected-no-diagnostics
#ifndef HEADER
@@ -139,6 +139,8 @@ int main (int argc, char **argv) {
static int a;
// CHECK: static int a;
+#pragma omp target data map(to: ([argc][3][a])argv)
+ // CHECK: #pragma omp target data map(to: ([argc][3][a])argv)
#pragma omp target data map(to: c)
// CHECK: #pragma omp target data map(to: c)
a=2;
diff --git a/clang/test/OpenMP/target_map_codegen.cpp b/clang/test/OpenMP/target_map_codegen.cpp
index b9766e82ce03..ecfe50c01ea6 100644
--- a/clang/test/OpenMP/target_map_codegen.cpp
+++ b/clang/test/OpenMP/target_map_codegen.cpp
@@ -5353,5 +5353,81 @@ void explicit_maps_single (int ii){
// CK31: define {{.+}}[[CALL00]]
// CK31: define {{.+}}[[CALL01]]
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK32 -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 CK32 --check-prefix CK32-64
+// RUN: %clang_cc1 -DCK32 -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 CK32 --check-prefix CK32-64
+// RUN: %clang_cc1 -DCK32 -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 CK32 --check-prefix CK32-32
+// RUN: %clang_cc1 -DCK32 -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 CK32 --check-prefix CK32-32
+
+// RUN: %clang_cc1 -DCK32 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY32 %s
+// RUN: %clang_cc1 -DCK32 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY32 %s
+// RUN: %clang_cc1 -DCK32 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY32 %s
+// RUN: %clang_cc1 -DCK32 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY32 %s
+// SIMD-ONLY32-NOT: {{__kmpc|__tgt}}
+#ifdef CK32
+
+// CK32-DAG: [[MTYPE_TO:@.+]] = {{.+}}constant [1 x i64] [i64 33]
+// CK32-DAG: [[MTYPE_FROM:@.+]] = {{.+}}constant [1 x i64] [i64 34]
+
+void array_shaping(float *f, int sa) {
+
+ // CK32-DAG: call i32 @__tgt_target(i64 -1, i8* @{{.+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE_TO]]{{.+}})
+ // CK32-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+ // CK32-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+ // CK32-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
+
+ // CK32-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+ // CK32-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+ // CK32-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
+
+ // CK32-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to float**
+ // CK32-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to float**
+
+ // CK32-DAG: store float* [[F1:%.+]], float** [[BPC0]],
+ // CK32-DAG: store float* [[F2:%.+]], float** [[PC0]],
+ // CK32-DAG: store i64 [[SIZE:%.+]], i64* [[S0]],
+
+ // CK32-DAG: [[F1]] = load float*, float** [[F_ADDR:%.+]],
+ // CK32-DAG: [[F2]] = load float*, float** [[F_ADDR]],
+ // CK32-64-DAG: [[SIZE]] = mul nuw i64 [[SZ1:%.+]], 4
+ // CK32-64-DAG: [[SZ1]] = mul nuw i64 12, %{{.+}}
+ // CK32-32-DAG: [[SIZE]] = sext i32 [[SZ1:%.+]] to i64
+ // CK32-32-DAG: [[SZ1]] = mul nuw i32 [[SZ2:%.+]], 4
+ // CK32-32-DAG: [[SZ2]] = mul nuw i32 12, %{{.+}}
+ #pragma omp target map(to:([3][sa][4])f)
+ f[0] = 1;
+ sa = 1;
+ // CK32-DAG: call i32 @__tgt_target(i64 -1, i8* @{{.+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE_FROM]]{{.+}})
+ // CK32-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+ // CK32-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+ // CK32-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
+
+ // CK32-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+ // CK32-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+ // CK32-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
+
+ // CK32-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to float**
+ // CK32-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to float**
+
+ // CK32-DAG: store float* [[F1:%.+]], float** [[BPC0]],
+ // CK32-DAG: store float* [[F2:%.+]], float** [[PC0]],
+ // CK32-DAG: store i64 [[SIZE:%.+]], i64* [[S0]],
+
+ // CK32-DAG: [[F1]] = load float*, float** [[F_ADDR:%.+]],
+ // CK32-DAG: [[F2]] = load float*, float** [[F_ADDR]],
+ // CK32-64-DAG: [[SIZE]] = mul nuw i64 [[SZ1:%.+]], 5
+ // CK32-64-DAG: [[SZ1]] = mul nuw i64 4, %{{.+}}
+ // CK32-32-DAG: [[SIZE]] = sext i32 [[SZ1:%.+]] to i64
+ // CK32-32-DAG: [[SZ1]] = mul nuw i32 [[SZ2:%.+]], 5
+ // CK32-32-DAG: [[SZ2]] = mul nuw i32 4, %{{.+}}
+ #pragma omp target map(from: ([sa][5])f)
+ f[0] = 1;
+}
+
#endif
#endif
diff --git a/clang/test/OpenMP/target_map_messages.cpp b/clang/test/OpenMP/target_map_messages.cpp
index 96932af6a04c..a18590fc85fe 100644
--- a/clang/test/OpenMP/target_map_messages.cpp
+++ b/clang/test/OpenMP/target_map_messages.cpp
@@ -140,6 +140,8 @@ struct SA {
{}
#pragma omp target map(close bf: a) // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
{}
+ #pragma omp target map(([b[I]][bf])f) // le45-error {{expected ',' or ']' in lambda capture list}} le45-error {{expected ')'}} le45-note {{to match this '('}}
+ {}
return;
}
};
@@ -189,203 +191,209 @@ void SAclient(int arg) {
SD u;
SC r(p),t(p);
- #pragma omp target map(r)
+#pragma omp target map(r)
{}
- #pragma omp target map(marr[2][0:2][0:2]) // expected-error {{array section does not specify contiguous storage}}
+#pragma omp target map(marr[2] [0:2] [0:2]) // expected-error {{array section does not specify contiguous storage}}
{}
- #pragma omp target map(marr[:][0:2][0:2]) // expected-error {{array section does not specify contiguous storage}}
+#pragma omp target map(marr[:] [0:2] [0:2]) // expected-error {{array section does not specify contiguous storage}}
{}
- #pragma omp target map(marr[2][3][0:2])
+#pragma omp target map(marr[2][3] [0:2])
{}
- #pragma omp target map(marr[:][:][:])
+#pragma omp target map(marr[:][:][:])
{}
- #pragma omp target map(marr[:2][:][:])
+#pragma omp target map(marr[:2][:][:])
{}
- #pragma omp target map(marr[arg:][:][:])
+#pragma omp target map(marr [arg:][:][:])
{}
- #pragma omp target map(marr[arg:])
+#pragma omp target map(marr [arg:])
{}
- #pragma omp target map(marr[arg:][:arg][:]) // correct if arg is the size of dimension 2
+#pragma omp target map(marr [arg:][:arg][:]) // correct if arg is the size of dimension 2
{}
- #pragma omp target map(marr[:arg][:])
+#pragma omp target map(marr[:arg][:])
{}
- #pragma omp target map(marr[:arg][n:])
+#pragma omp target map(marr[:arg] [n:])
{}
- #pragma omp target map(marr[:][:arg][n:]) // correct if arg is the size of dimension 2
+#pragma omp target map(marr[:][:arg] [n:]) // correct if arg is the size of dimension 2
{}
- #pragma omp target map(marr[:][:m][n:]) // expected-error {{array section does not specify contiguous storage}}
+#pragma omp target map(marr[:][:m] [n:]) // expected-error {{array section does not specify contiguous storage}}
{}
- #pragma omp target map(marr[n:m][:arg][n:])
+#pragma omp target map(marr [n:m][:arg] [n:])
{}
- #pragma omp target map(marr[:2][:1][:]) // expected-error {{array section does not specify contiguous storage}}
+#pragma omp target map(marr[:2][:1][:]) // expected-error {{array section does not specify contiguous storage}}
{}
- #pragma omp target map(marr[:2][1:][:]) // expected-error {{array section does not specify contiguous storage}}
+#pragma omp target map(marr[:2] [1:][:]) // expected-error {{array section does not specify contiguous storage}}
{}
- #pragma omp target map(marr[:2][:][:1]) // expected-error {{array section does not specify contiguous storage}}
+#pragma omp target map(marr[:2][:][:1]) // expected-error {{array section does not specify contiguous storage}}
{}
- #pragma omp target map(marr[:2][:][1:]) // expected-error {{array section does not specify contiguous storage}}
+#pragma omp target map(marr[:2][:] [1:]) // expected-error {{array section does not specify contiguous storage}}
{}
- #pragma omp target map(marr[:1][:2][:])
+#pragma omp target map(marr[:1][:2][:])
{}
- #pragma omp target map(marr[:1][0][:])
+#pragma omp target map(marr[:1][0][:])
{}
- #pragma omp target map(marr[:arg][:2][:]) // correct if arg is 1
+#pragma omp target map(marr[:arg][:2][:]) // correct if arg is 1
{}
- #pragma omp target map(marr[:1][3:1][:2])
+#pragma omp target map(marr[:1] [3:1][:2])
{}
- #pragma omp target map(marr[:1][3:arg][:2]) // correct if arg is 1
+#pragma omp target map(marr[:1] [3:arg][:2]) // correct if arg is 1
{}
- #pragma omp target map(marr[:1][3:2][:2]) // expected-error {{array section does not specify contiguous storage}}
+#pragma omp target map(marr[:1] [3:2][:2]) // expected-error {{array section does not specify contiguous storage}}
{}
- #pragma omp target map(marr[:2][:10][:])
+#pragma omp target map(marr[:2][:10][:])
{}
- #pragma omp target map(marr[:2][:][:5+5])
+#pragma omp target map(marr[:2][:][:5 + 5])
{}
- #pragma omp target map(marr[:2][2+2-4:][0:5+5])
+#pragma omp target map(marr[:2] [2 + 2 - 4:] [0:5 + 5])
{}
- #pragma omp target map(marr[:1][:2][0]) // expected-error {{array section does not specify contiguous storage}}
+#pragma omp target map(marr[:1][:2][0]) // expected-error {{array section does not specify contiguous storage}}
{}
- #pragma omp target map(marr2[:1][:2][0])
+#pragma omp target map(marr2[:1][:2][0])
{}
- #pragma omp target map(mvla[:1][:][0]) // correct if the size of dimension 2 is 1.
+#pragma omp target map(mvla[:1][:][0]) // correct if the size of dimension 2 is 1.
{}
- #pragma omp target map(mvla[:2][:arg][:]) // correct if arg is the size of dimension 2.
+#pragma omp target map(mvla[:2][:arg][:]) // correct if arg is the size of dimension 2.
{}
- #pragma omp target map(mvla[:1][:2][0]) // expected-error {{array section does not specify contiguous storage}}
+#pragma omp target map(mvla[:1][:2][0]) // expected-error {{array section does not specify contiguous storage}}
{}
- #pragma omp target map(mvla[1][2:arg][:])
+#pragma omp target map(mvla[1] [2:arg][:])
{}
- #pragma omp target map(mvla[:1][:][:])
+#pragma omp target map(mvla[:1][:][:])
{}
- #pragma omp target map(mvla2[:1][:2][:11])
+#pragma omp target map(mvla2[:1][:2][:11])
{}
- #pragma omp target map(mvla2[:1][:2][:10]) // expected-error {{array section does not specify contiguous storage}}
+#pragma omp target map(mvla2[:1][:2][:10]) // expected-error {{array section does not specify contiguous storage}}
{}
- #pragma omp target map(mptr[:2][2+2-4:1][0:5+5]) // expected-error {{array section does not specify contiguous storage}}
+#pragma omp target map(mptr[:2] [2 + 2 - 4:1] [0:5 + 5]) // expected-error {{array section does not specify contiguous storage}}
{}
- #pragma omp target map(mptr[:1][:2-1][2:4-3])
+#pragma omp target map(mptr[:1][:2 - 1] [2:4 - 3])
{}
- #pragma omp target map(mptr[:1][:arg][2:4-3]) // correct if arg is 1.
+#pragma omp target map(mptr[:1][:arg] [2:4 - 3]) // correct if arg is 1.
{}
- #pragma omp target map(mptr[:1][:2-1][0:2])
+#pragma omp target map(mptr[:1][:2 - 1] [0:2])
{}
- #pragma omp target map(mptr[:1][:2][0:2]) // expected-error {{array section does not specify contiguous storage}}
+#pragma omp target map(mptr[:1][:2] [0:2]) // expected-error {{array section does not specify contiguous storage}}
{}
- #pragma omp target map(mptr[:1][:][0:2]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
+#pragma omp target map(mptr[:1][:] [0:2]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
{}
- #pragma omp target map(mptr[:2][:1][0:2]) // expected-error {{array section does not specify contiguous storage}}
+#pragma omp target map(mptr[:2][:1] [0:2]) // expected-error {{array section does not specify contiguous storage}}
{}
- #pragma omp target map(r.ArrS[0].B)
+#pragma omp target map(r.ArrS[0].B)
{}
- #pragma omp target map(r.ArrS[:1].B) // expected-error {{OpenMP array section is not allowed here}}
+#pragma omp target map(r.ArrS[:1].B) // expected-error {{OpenMP array section is not allowed here}}
{}
- #pragma omp target map(r.ArrS[:arg].B) // expected-error {{OpenMP array section is not allowed here}}
+#pragma omp target map(r.ArrS[:arg].B) // expected-error {{OpenMP array section is not allowed here}}
{}
- #pragma omp target map(r.ArrS[0].Arr[1:23])
+#pragma omp target map(r.ArrS[0].Arr [1:23])
{}
- #pragma omp target map(r.ArrS[0].Arr[1:arg])
+#pragma omp target map(r.ArrS[0].Arr [1:arg])
{}
- #pragma omp target map(r.ArrS[0].Arr[arg:23])
+#pragma omp target map(r.ArrS[0].Arr [arg:23])
{}
- #pragma omp target map(r.ArrS[0].Error) // expected-error {{no member named 'Error' in 'SB'}}
+#pragma omp target map(r.ArrS[0].Error) // expected-error {{no member named 'Error' in 'SB'}}
{}
- #pragma omp target map(r.ArrS[0].A, r.ArrS[1].A) // expected-error {{multiple array elements associated with the same variable are not allowed in map clauses of the same construct}} expected-note {{used here}}
+#pragma omp target map(r.ArrS[0].A, r.ArrS[1].A) // expected-error {{multiple array elements associated with the same variable are not allowed in map clauses of the same construct}} expected-note {{used here}}
{}
- #pragma omp target map(r.ArrS[0].A, t.ArrS[1].A)
+#pragma omp target map(r.ArrS[0].A, t.ArrS[1].A)
{}
- #pragma omp target map(r.PtrS[0], r.PtrS->B) // expected-error {{same pointer dereferenced in multiple
diff erent ways in map clause expressions}} expected-note {{used here}}
+#pragma omp target map(r.PtrS[0], r.PtrS->B) // expected-error {{same pointer dereferenced in multiple
diff erent ways in map clause expressions}} expected-note {{used here}}
{}
- #pragma omp target map(r.PtrS, r.PtrS->B) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}}
+#pragma omp target map(r.PtrS, r.PtrS->B) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}}
{}
- #pragma omp target map(r.PtrS->A, r.PtrS->B)
+#pragma omp target map(r.PtrS->A, r.PtrS->B)
{}
- #pragma omp target map(r.RPtrS[0], r.RPtrS->B) // expected-error {{same pointer dereferenced in multiple
diff erent ways in map clause expressions}} expected-note {{used here}}
+#pragma omp target map(r.RPtrS[0], r.RPtrS->B) // expected-error {{same pointer dereferenced in multiple
diff erent ways in map clause expressions}} expected-note {{used here}}
{}
- #pragma omp target map(r.RPtrS, r.RPtrS->B) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}}
+#pragma omp target map(r.RPtrS, r.RPtrS->B) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}}
{}
- #pragma omp target map(r.RPtrS->A, r.RPtrS->B)
+#pragma omp target map(r.RPtrS->A, r.RPtrS->B)
{}
- #pragma omp target map(r.S.Arr[:12])
+#pragma omp target map(r.S.Arr[:12])
{}
- #pragma omp target map(r.S.foo()[:12]) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le50-error {{expected addressable lvalue in 'map' clause}}
+#pragma omp target map(r.S.foo() [:12]) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le50-error {{expected addressable lvalue in 'map' clause}}
{}
- #pragma omp target map(r.C, r.D)
+#pragma omp target map(r.C, r.D)
{}
- #pragma omp target map(r.C, r.C) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+#pragma omp target map(r.C, r.C) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
{}
- #pragma omp target map(r.C) map(r.C) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+#pragma omp target map(r.C) map(r.C) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
{}
- #pragma omp target map(r.C, r.S) // this would be an error only caught at runtime - Sema would have to make sure there is not way for the missing data between fields to be mapped somewhere else.
+#pragma omp target map(r.C, r.S) // this would be an error only caught at runtime - Sema would have to make sure there is not way for the missing data between fields to be mapped somewhere else.
{}
- #pragma omp target map(r, r.S) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+#pragma omp target map(r, r.S) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
{}
- #pragma omp target map(r.C, t.C)
+#pragma omp target map(r.C, t.C)
{}
- #pragma omp target map(r.A) // expected-error {{bit fields cannot be used to specify storage in a 'map' clause}}
+#pragma omp target map(r.A) // expected-error {{bit fields cannot be used to specify storage in a 'map' clause}}
{}
- #pragma omp target map(r.Arr)
+#pragma omp target map(r.Arr)
{}
- #pragma omp target map(r.Arr[3:5])
+#pragma omp target map(r.Arr [3:5])
{}
- #pragma omp target map(r.Ptr[3:5])
+#pragma omp target map(r.Ptr [3:5])
{}
- #pragma omp target map(r.ArrS[3:5].A) // expected-error {{OpenMP array section is not allowed here}}
+#pragma omp target map(r.ArrS [3:5].A) // expected-error {{OpenMP array section is not allowed here}}
{}
- #pragma omp target map(r.ArrS[3:5].Arr[6:7]) // expected-error {{OpenMP array section is not allowed here}}
+#pragma omp target map(r.ArrS [3:5].Arr [6:7]) // expected-error {{OpenMP array section is not allowed here}}
{}
- #pragma omp target map(r.ArrS[3].Arr[6:7])
+#pragma omp target map(r.ArrS[3].Arr [6:7])
{}
- #pragma omp target map(r.S.Arr[4:5])
+#pragma omp target map(r.S.Arr [4:5])
{}
- #pragma omp target map(r.S.Ptr[4:5])
+#pragma omp target map(r.S.Ptr [4:5])
{}
- #pragma omp target map(r.S.Ptr[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
+#pragma omp target map(r.S.Ptr[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
{}
- #pragma omp target map((p+1)->A) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target map((p + 1)->A) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
{}
- #pragma omp target map(u.B) // expected-error {{mapping of union members is not allowed}}
+#pragma omp target map(u.B) // expected-error {{mapping of union members is not allowed}}
{}
- #pragma omp target
+#pragma omp target
{
u.B = 0;
r.S.foo();
}
- #pragma omp target data map(to: r.C) //expected-note {{used here}}
+#pragma omp target data map(to \
+ : r.C) //expected-note {{used here}}
{
- #pragma omp target map(r.D) // expected-error {{original storage of expression in data environment is shared but data environment do not fully contain mapped expression storage}}
+#pragma omp target map(r.D) // expected-error {{original storage of expression in data environment is shared but data environment do not fully contain mapped expression storage}}
{}
}
- #pragma omp target data map(to: t.Ptr) //expected-note {{used here}}
+#pragma omp target data map(to \
+ : t.Ptr) //expected-note {{used here}}
{
- #pragma omp target map(t.Ptr[:23]) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target map(t.Ptr[:23]) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
{}
}
- #pragma omp target data map(to: t.C, t.D)
+#pragma omp target data map(to \
+ : t.C, t.D)
{
- #pragma omp target data map(to: t.C)
+#pragma omp target data map(to \
+ : t.C)
{
- #pragma omp target map(t.D)
+#pragma omp target map(t.D)
{}
}
}
- #pragma omp target data map(marr[:][:][:])
+#pragma omp target data map(marr[:][:][:])
{
- #pragma omp target data map(marr)
+#pragma omp target data map(marr)
{}
}
- #pragma omp target data map(to: t)
+#pragma omp target data map(to \
+ : t)
{
- #pragma omp target data map(to: t.C)
+#pragma omp target data map(to \
+ : t.C)
{
- #pragma omp target map(t.D)
+#pragma omp target map(t.D)
{}
}
}
diff --git a/clang/test/OpenMP/target_update_ast_print.cpp b/clang/test/OpenMP/target_update_ast_print.cpp
index e60e081b3210..fb6440b87cea 100644
--- a/clang/test/OpenMP/target_update_ast_print.cpp
+++ b/clang/test/OpenMP/target_update_ast_print.cpp
@@ -1,10 +1,10 @@
-// RUN: %clang_cc1 -verify -fopenmp -ast-print %s | FileCheck %s
-// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -emit-pch -o %t %s
-// RUN: %clang_cc1 -fopenmp -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=50 -ast-print %s | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
-// RUN: %clang_cc1 -verify -fopenmp-simd -ast-print %s | FileCheck %s
-// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -emit-pch -o %t %s
-// RUN: %clang_cc1 -fopenmp-simd -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=50 -ast-print %s | FileCheck %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
// expected-no-diagnostics
#ifndef HEADER
@@ -14,29 +14,29 @@ void foo() {}
template <class T, class U>
T foo(T targ, U uarg) {
- static T a;
+ static T a, *p;
U b;
int l;
-#pragma omp target update to(a) if(l>5) device(l) nowait depend(inout:l)
+#pragma omp target update to(([a][targ])p, a) if(l>5) device(l) nowait depend(inout:l)
-#pragma omp target update from(b) if(l<5) device(l-1) nowait depend(inout:l)
+#pragma omp target update from(b, ([a][targ])p) if(l<5) device(l-1) nowait depend(inout:l)
return a + targ + (T)b;
}
-// CHECK: static T a;
+// CHECK: static T a, *p;
// CHECK-NEXT: U b;
// CHECK-NEXT: int l;
-// CHECK-NEXT: #pragma omp target update to(a) if(l > 5) device(l) nowait depend(inout : l){{$}}
-// CHECK-NEXT: #pragma omp target update from(b) if(l < 5) device(l - 1) nowait depend(inout : l)
-// CHECK: static int a;
+// 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: static int a, *p;
// CHECK-NEXT: float b;
// CHECK-NEXT: int l;
-// CHECK-NEXT: #pragma omp target update to(a) if(l > 5) device(l) nowait depend(inout : l)
-// CHECK-NEXT: #pragma omp target update from(b) if(l < 5) device(l - 1) nowait depend(inout : l)
-// CHECK: static char a;
+// 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: static char a, *p;
// CHECK-NEXT: float b;
// CHECK-NEXT: int l;
-// CHECK-NEXT: #pragma omp target update to(a) if(l > 5) device(l) nowait depend(inout : l)
-// CHECK-NEXT: #pragma omp target update from(b) if(l < 5) device(l - 1) nowait depend(inout : 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)
int main(int argc, char **argv) {
static int a;
diff --git a/clang/test/OpenMP/target_update_codegen.cpp b/clang/test/OpenMP/target_update_codegen.cpp
index 479461e7ca80..fd5a62a8067c 100644
--- a/clang/test/OpenMP/target_update_codegen.cpp
+++ b/clang/test/OpenMP/target_update_codegen.cpp
@@ -984,5 +984,80 @@ void lvalue_find_base(float **f, SSA *sa) {
#pragma omp target update from(*(sa->sa->i+*(1+sa->i+f)))
}
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK18 -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 CK18 --check-prefix CK18-64
+// RUN: %clang_cc1 -DCK18 -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 CK18 --check-prefix CK18-64
+// RUN: %clang_cc1 -DCK18 -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 CK18 --check-prefix CK18-32
+// RUN: %clang_cc1 -DCK18 -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 CK18 --check-prefix CK18-32
+
+// RUN: %clang_cc1 -DCK18 -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-ONLY18 %s
+// RUN: %clang_cc1 -DCK18 -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-ONLY18 %s
+// RUN: %clang_cc1 -DCK18 -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-ONLY18 %s
+// RUN: %clang_cc1 -DCK18 -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-ONLY18 %s
+// SIMD-ONLY18-NOT: {{__kmpc|__tgt}}
+#ifdef CK18
+
+// CK18-DAG: [[MTYPE_TO:@.+]] = {{.+}}constant [1 x i64] [i64 33]
+// CK18-DAG: [[MTYPE_FROM:@.+]] = {{.+}}constant [1 x i64] [i64 34]
+
+//CK18-LABEL: array_shaping
+void array_shaping(float *f, int sa) {
+
+ // CK18-DAG: call void @__tgt_target_data_update(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE_TO]]{{.+}})
+ // CK18-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+ // CK18-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+ // CK18-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
+
+ // CK18-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+ // CK18-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+ // CK18-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
+
+ // CK18-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to float**
+ // CK18-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to float**
+
+ // CK18-DAG: store float* [[F1:%.+]], float** [[BPC0]],
+ // CK18-DAG: store float* [[F2:%.+]], float** [[PC0]],
+ // CK18-DAG: store i64 [[SIZE:%.+]], i64* [[S0]],
+
+ // CK18-DAG: [[F1]] = load float*, float** [[F_ADDR:%.+]],
+ // CK18-DAG: [[F2]] = load float*, float** [[F_ADDR]],
+ // CK18-64-DAG: [[SIZE]] = mul nuw i64 [[SZ1:%.+]], 4
+ // CK18-64-DAG: [[SZ1]] = mul nuw i64 12, %{{.+}}
+ // CK18-32-DAG: [[SIZE]] = sext i32 [[SZ1:%.+]] to i64
+ // CK18-32-DAG: [[SZ1]] = mul nuw i32 [[SZ2:%.+]], 4
+ // CK18-32-DAG: [[SZ2]] = mul nuw i32 12, %{{.+}}
+ #pragma omp target update to(([3][sa][4])f)
+ sa = 1;
+ // CK18-DAG: call void @__tgt_target_data_update(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE_FROM]]{{.+}})
+ // CK18-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+ // CK18-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+ // CK18-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
+
+ // CK18-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+ // CK18-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+ // CK18-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
+
+ // CK18-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to float**
+ // CK18-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to float**
+
+ // CK18-DAG: store float* [[F1:%.+]], float** [[BPC0]],
+ // CK18-DAG: store float* [[F2:%.+]], float** [[PC0]],
+ // CK18-DAG: store i64 [[SIZE:%.+]], i64* [[S0]],
+
+ // CK18-DAG: [[F1]] = load float*, float** [[F_ADDR:%.+]],
+ // CK18-DAG: [[F2]] = load float*, float** [[F_ADDR]],
+ // CK18-64-DAG: [[SIZE]] = mul nuw i64 [[SZ1:%.+]], 5
+ // CK18-64-DAG: [[SZ1]] = mul nuw i64 4, %{{.+}}
+ // CK18-32-DAG: [[SIZE]] = sext i32 [[SZ1:%.+]] to i64
+ // CK18-32-DAG: [[SZ1]] = mul nuw i32 [[SZ2:%.+]], 5
+ // CK18-32-DAG: [[SZ2]] = mul nuw i32 4, %{{.+}}
+ #pragma omp target update from(([sa][5])f)
+}
+
#endif
#endif
More information about the cfe-commits
mailing list