diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp --- a/clang/lib/Sema/SemaOpenMP.cpp +++ b/clang/lib/Sema/SemaOpenMP.cpp @@ -38,7 +38,7 @@ static const Expr *checkMapClauseExpressionBase( Sema &SemaRef, Expr *E, OMPClauseMappableExprCommon::MappableExprComponentList &CurComponents, - OpenMPClauseKind CKind, bool NoDiagnose); + OpenMPClauseKind CKind, OpenMPDirectiveKind DKind, bool NoDiagnose); namespace { /// Default data sharing attributes, which can be applied to directive. @@ -2997,7 +2997,7 @@ if (isOpenMPTargetExecutionDirective(DKind)) { OMPClauseMappableExprCommon::MappableExprComponentList CurComponents; if (!checkMapClauseExpressionBase(SemaRef, E, CurComponents, OMPC_map, - /*NoDiagnose=*/true)) + Stack->getCurrentDirective(), /*NoDiagnose=*/true)) return; const auto *VD = cast( CurComponents.back().getAssociatedDeclaration()->getCanonicalDecl()); @@ -14723,7 +14723,7 @@ static const Expr *checkMapClauseExpressionBase( Sema &SemaRef, Expr *E, OMPClauseMappableExprCommon::MappableExprComponentList &CurComponents, - OpenMPClauseKind CKind, bool NoDiagnose) { + OpenMPClauseKind CKind, OpenMPDirectiveKind DKind, bool NoDiagnose) { SourceLocation ELoc = E->getExprLoc(); SourceRange ERange = E->getSourceRange(); @@ -14748,6 +14748,10 @@ const Expr *RelevantExpr = nullptr; + // OpenMP 5.0 + // The target update construct was modified to allow array section that + // specify discontiguous storage. + // // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.2] // If a list item is an array section, it must specify contiguous storage. // @@ -14919,7 +14923,10 @@ // pointer. Otherwise, only unitary sections are accepted. if (NotWhole || IsPointer) AllowWholeSizeArraySection = false; - } else if (AllowUnitySizeArraySection && NotUnity) { + } else if (((SemaRef.getLangOpts().OpenMP >= 50 && + DKind != OMPD_target_update) || + SemaRef.getLangOpts().OpenMP < 50) && + (AllowUnitySizeArraySection && NotUnity)) { // A unity or whole array section is not allowed and that is not // compatible with the properties of the current array section. SemaRef.Diag( @@ -15431,7 +15438,8 @@ // Obtain the array or member expression bases if required. Also, fill the // components array with all the components identified in the process. const Expr *BE = checkMapClauseExpressionBase( - SemaRef, SimpleExpr, CurComponents, CKind, /*NoDiagnose=*/false); + SemaRef, SimpleExpr, CurComponents, CKind, DSAS->getCurrentDirective(), + /*NoDiagnose=*/false); if (!BE) continue; diff --git a/clang/test/OpenMP/target_update_ast_print.cpp b/clang/test/OpenMP/target_update_ast_print.cpp --- a/clang/test/OpenMP/target_update_ast_print.cpp +++ b/clang/test/OpenMP/target_update_ast_print.cpp @@ -5,6 +5,14 @@ // RUN: %clang_cc1 -verify -fopenmp-simd -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 -DOMP5 -verify -fopenmp -fopenmp-version=50 -ast-print %s | FileCheck %s --check-prefix=OMP5 +// RUN: %clang_cc1 -DOMP5 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s --check-prefix=OMP5 + +// RUN: %clang_cc1 -DOMP5 -verify -fopenmp-simd -fopenmp-version=50 -ast-print %s | FileCheck %s --check-prefix=OMP5 +// RUN: %clang_cc1 -DOMP5 -fopenmp-simd -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s +// RUN: %clang_cc1 -DOMP5 -fopenmp-simd -fopenmp-version=50 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s --check-prefix=OMP5 // expected-no-diagnostics #ifndef HEADER @@ -20,8 +28,7 @@ #pragma omp target update to(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) - return a + targ + (T)b; -} + // CHECK: static T a; // CHECK-NEXT: U b; // CHECK-NEXT: int l; @@ -38,6 +45,62 @@ // 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) +#ifdef OMP5 + U marr[10][10][10]; +#pragma omp target update to(marr[2][0:2][0:2]) + +#pragma omp target update from(marr[2][0:2][0:2]) + +#pragma omp target update to(marr[:][0:2][0:2]) + +#pragma omp target update from(marr[:][0:2][0:2]) + +#pragma omp target update to(marr[:][:l][l:]) + +#pragma omp target update from(marr[:][:l][l:]) + +#pragma omp target update to(marr[:2][:1][:]) + +#pragma omp target update from(marr[:2][:1][:]) + +#pragma omp target update to(marr[:2][:][:1]) + +#pragma omp target update from(marr[:2][:][:1]) + +#pragma omp target update to(marr[:2][:][1:]) + +#pragma omp target update from(marr[:2][:][1:]) + +#pragma omp target update to(marr[:1][3:2][:2]) + +#pragma omp target update from(marr[:1][3:2][:2]) + +#pragma omp target update to(marr[:1][:2][0]) + +#pragma omp target update from(marr[:1][:2][0]) + +// OMP5: marr[10][10][10]; +// OMP5-NEXT: #pragma omp target update to(marr[2][0:2][0:2]) +// OMP5-NEXT: #pragma omp target update from(marr[2][0:2][0:2]) +// OMP5-NEXT: #pragma omp target update to(marr[:][0:2][0:2]) +// OMP5-NEXT: #pragma omp target update from(marr[:][0:2][0:2]) +// OMP5-NEXT: #pragma omp target update to(marr[:][:l][l:]) +// OMP5-NEXT: #pragma omp target update from(marr[:][:l][l:]) +// OMP5-NEXT: #pragma omp target update to(marr[:2][:1][:]) +// OMP5-NEXT: #pragma omp target update from(marr[:2][:1][:]) +// OMP5-NEXT: #pragma omp target update to(marr[:2][:][:1]) +// OMP5-NEXT: #pragma omp target update from(marr[:2][:][:1]) +// OMP5-NEXT: #pragma omp target update to(marr[:2][:][1:]) +// OMP5-NEXT: #pragma omp target update from(marr[:2][:][1:]) +// OMP5-NEXT: #pragma omp target update to(marr[:1][3:2][:2]) +// OMP5-NEXT: #pragma omp target update from(marr[:1][3:2][:2]) +// OMP5-NEXT: #pragma omp target update to(marr[:1][:2][0]) +// OMP5-NEXT: #pragma omp target update from(marr[:1][:2][0]) +#endif + + return a + targ + (T)b; +} + int main(int argc, char **argv) { static int a; int n; @@ -50,6 +113,40 @@ // CHECK-NEXT: #pragma omp target update to(a) if(f > 0.) device(n) nowait depend(in : n) #pragma omp target update from(f) if(f<0.0) device(n+1) nowait depend(in:n) // CHECK-NEXT: #pragma omp target update from(f) if(f < 0.) device(n + 1) nowait depend(in : n) + +#ifdef OMP5 + float marr[10][10][10]; +// OMP5: marr[10][10][10]; +#pragma omp target update to(marr[2][0:2][0:2]) +// OMP5-NEXT: #pragma omp target update to(marr[2][0:2][0:2]) +#pragma omp target update from(marr[2][0:2][0:2]) +// OMP5-NEXT: #pragma omp target update from(marr[2][0:2][0:2]) +#pragma omp target update to(marr[:][0:2][0:2]) +// OMP5-NEXT: #pragma omp target update to(marr[:][0:2][0:2]) +#pragma omp target update from(marr[:][0:2][0:2]) +// OMP5-NEXT: #pragma omp target update from(marr[:][0:2][0:2]) +#pragma omp target update to(marr[:][:n][n:]) +// OMP5: #pragma omp target update to(marr[:][:n][n:]) +#pragma omp target update from(marr[:2][:1][:]) +// OMP5-NEXT: #pragma omp target update from(marr[:2][:1][:]) +#pragma omp target update to(marr[:2][:][:1]) +// OMP5-NEXT: #pragma omp target update to(marr[:2][:][:1]) +#pragma omp target update from(marr[:2][:][:1]) +// OMP5-NEXT: #pragma omp target update from(marr[:2][:][:1]) +#pragma omp target update to(marr[:2][:][1:]) +// OMP5-NEXT: #pragma omp target update to(marr[:2][:][1:]) +#pragma omp target update from(marr[:2][:][1:]) +// OMP5-NEXT: #pragma omp target update from(marr[:2][:][1:]) +#pragma omp target update to(marr[:1][3:2][:2]) +// OMP5-NEXT: #pragma omp target update to(marr[:1][3:2][:2]) +#pragma omp target update from(marr[:1][3:2][:2]) +// OMP5-NEXT: #pragma omp target update from(marr[:1][3:2][:2]) +#pragma omp target update to(marr[:1][:2][0]) +// OMP5-NEXT: #pragma omp target update to(marr[:1][:2][0]) +#pragma omp target update from(marr[:1][:2][0]) +// OMP5-NEXT: #pragma omp target update from(marr[:1][:2][0]) +#endif + return foo(argc, f) + foo(argv[0][0], f) + a; } diff --git a/clang/test/OpenMP/target_update_codegen.cpp b/clang/test/OpenMP/target_update_codegen.cpp --- a/clang/test/OpenMP/target_update_codegen.cpp +++ b/clang/test/OpenMP/target_update_codegen.cpp @@ -292,4 +292,39 @@ {++arg;} } #endif +///==========================================================================/// +// RUN: %clang_cc1 -DCK5 -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 CK5 --check-prefix CK5-64 +// RUN: %clang_cc1 -DCK5 -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 CK5 --check-prefix CK5-64 +// RUN: %clang_cc1 -DCK5 -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 CK5 --check-prefix CK5-32 +// RUN: %clang_cc1 -DCK5 -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 CK5 --check-prefix CK5-32 + +// RUN: %clang_cc1 -DCK5 -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-ONLY0 %s +// RUN: %clang_cc1 -DCK5 -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-ONLY0 %s +// RUN: %clang_cc1 -DCK5 -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-ONLY0 %s +// RUN: %clang_cc1 -DCK5 -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-ONLY0 %s +// SIMD-ONLY0-NOT: {{__kmpc|__tgt}} +#ifdef CK5 + +// CK5: [[SIZE:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 80] +// CK5: [[MTYPE:@.+]] = {{.+}}constant [1 x i64] [i64 33] + +void target_update_contiguous() { + int marr[10][10][10]; + + // CK5-DAG: call void @__tgt_target_data_update(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE]]{{.+}}) + // CK5-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK5-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + // CK5-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK5-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK5-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to [10 x [10 x [10 x i32]]]** + // CK5-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to [10 x i32]** + // CK5-DAG: store [10 x [10 x [10 x i32]]]* %marr, [10 x [10 x [10 x i32]]]** [[BPC0]] + // CK5-DAG: store [10 x i32]* {{.+}}, [10 x i32]** [[PC0]] + #pragma omp target update to(marr[2][0:2][0:2]) +} +#endif #endif