[prev in list] [next in list] [prev in thread] [next in thread]
List: cfe-commits
Subject: [clang] c2aa543 - [OPENMP50]Codegen for array shaping expression in map clauses.
From: Alexey Bataev via cfe-commits <cfe-commits () lists ! llvm ! org>
Date: 2020-03-31 23:11:34
Message-ID: 5e83ce26.1c69fb81.9d3a4.0f7c () mx ! google ! com
[Download RAW message or body]
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
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
[prev in list] [next in list] [prev in thread] [next in thread]
Configure |
About |
News |
Add a list |
Sponsored by KoreLogic