From c2aa543237843fa7b7c0191b6685062b3512f245 Mon Sep 17 00:00:00 2001 From: Alexey Bataev Date: Mon, 30 Mar 2020 16:06:01 -0400 Subject: [PATCH] [OPENMP50]Codegen for array shaping expression in map clauses. Added codegen support for array shaping operations in map/to/from clauses. --- clang/lib/CodeGen/CGOpenMPRuntime.cpp | 34 ++++- clang/lib/Sema/SemaOpenMP.cpp | 34 ++++- clang/test/OpenMP/target_data_ast_print.cpp | 14 +- clang/test/OpenMP/target_map_codegen.cpp | 76 +++++++++++ clang/test/OpenMP/target_map_messages.cpp | 190 ++++++++++++++------------ clang/test/OpenMP/target_update_ast_print.cpp | 36 ++--- clang/test/OpenMP/target_update_codegen.cpp | 75 ++++++++++ 7 files changed, 336 insertions(+), 123 deletions(-) diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index 31fdc32..6642851 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -7448,6 +7448,20 @@ private: 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(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()) ExprTy = RefTy->getPointeeType().getCanonicalType(); @@ -7779,6 +7793,7 @@ private: const Expr *AssocExpr = I->getAssociatedExpression(); const auto *AE = dyn_cast(AssocExpr); const auto *OASE = dyn_cast(AssocExpr); + const auto *OAShE = dyn_cast(AssocExpr); if (isa(AssocExpr)) { // The base is the 'this' pointer. The content of the pointer is going @@ -7788,6 +7803,11 @@ private: (OASE && isa(OASE->getBase()->IgnoreParenImpCasts()))) { BP = CGF.EmitOMPSharedLValue(AssocExpr).getAddress(CGF); + } else if (OAShE && + isa(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 @@ private: // types. const auto *OASE = dyn_cast(I->getAssociatedExpression()); + const auto *OAShE = + dyn_cast(I->getAssociatedExpression()); const auto *UO = dyn_cast(I->getAssociatedExpression()); const auto *BO = dyn_cast(I->getAssociatedExpression()); bool IsPointer = + OAShE || (OASE && OMPArraySectionExpr::getBaseOriginalType(OASE) .getCanonicalType() ->isAnyPointerType()) || @@ -7890,8 +7913,15 @@ private: isa(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 f9e8e3d..7d2ae17 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(EI->getAssociatedExpression()) || isa(EI->getAssociatedExpression()) || - isa(EI->getAssociatedExpression())) { + isa(EI->getAssociatedExpression()) || + isa(EI->getAssociatedExpression())) { IsVariableAssociatedWithSection = true; // There is nothing more we need to know about this variable. return true; @@ -3225,7 +3226,7 @@ public: 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 @@ public: nullptr && (isa( MC.getAssociatedExpression()) || + isa( + MC.getAssociatedExpression()) || isa( MC.getAssociatedExpression())); }); @@ -3393,8 +3396,10 @@ public: // Do both expressions have the same kind? if (CCI->getAssociatedExpression()->getStmtClass() != SC.getAssociatedExpression()->getStmtClass()) - if (!(isa( - SC.getAssociatedExpression()) && + if (!((isa( + SC.getAssociatedExpression()) || + isa( + SC.getAssociatedExpression())) && isa( CCI->getAssociatedExpression()))) return false; @@ -16284,6 +16289,15 @@ public: 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(CI->getAssociatedExpression()) || - isa(CI->getAssociatedExpression())) && + isa(CI->getAssociatedExpression()) || + isa(CI->getAssociatedExpression())) && (isa(SI->getAssociatedExpression()) || - isa(SI->getAssociatedExpression()))) { + isa(SI->getAssociatedExpression()) || + isa(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( + 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(VE->IgnoreParens()); auto *OASE = dyn_cast(VE->IgnoreParens()); + auto *OAShE = dyn_cast(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 fa67c18..fcd6e92 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 b9766e8..ecfe50c 100644 --- a/clang/test/OpenMP/target_map_codegen.cpp +++ b/clang/test/OpenMP/target_map_codegen.cpp @@ -5354,4 +5354,80 @@ void explicit_maps_single (int ii){ // 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 96932af..a18590f 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 different 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 different 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 different 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 different 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 e60e081..fb6440b 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 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 479461e..fd5a62a 100644 --- a/clang/test/OpenMP/target_update_codegen.cpp +++ b/clang/test/OpenMP/target_update_codegen.cpp @@ -985,4 +985,79 @@ void lvalue_find_base(float **f, SSA *sa) { } #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 -- 2.7.4