diff --git a/clang/include/clang/Basic/OpenMPKinds.def b/clang/include/clang/Basic/OpenMPKinds.def --- a/clang/include/clang/Basic/OpenMPKinds.def +++ b/clang/include/clang/Basic/OpenMPKinds.def @@ -89,6 +89,7 @@ OPENMP_DEFAULTMAP_MODIFIER(firstprivate) OPENMP_DEFAULTMAP_MODIFIER(none) OPENMP_DEFAULTMAP_MODIFIER(default) +OPENMP_DEFAULTMAP_MODIFIER(present) // Static attributes for 'depend' clause. OPENMP_DEPEND_KIND(in) 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 @@ -3306,6 +3306,14 @@ case OMPC_DEFAULTMAP_MODIFIER_tofrom: Kind = OMPC_MAP_tofrom; break; + case OMPC_DEFAULTMAP_MODIFIER_present: + // OpenMP 5.1 [2.21.7.3] defaultmap clause, Description] + // If implicit-behavior is present, each variable referenced in the + // construct in the category specified by variable-category is treated as if + // it had been listed in a map clause with the map-type of alloc and + // map-type-modifier of present. + Kind = OMPC_MAP_alloc; + break; case OMPC_DEFAULTMAP_MODIFIER_firstprivate: case OMPC_DEFAULTMAP_MODIFIER_last: llvm_unreachable("Unexpected defaultmap implicit behavior"); @@ -3332,8 +3340,11 @@ bool ErrorFound = false; bool TryCaptureCXXThisMembers = false; CapturedStmt *CS = nullptr; + const static unsigned DefaultmapKindNum = OMPC_DEFAULTMAP_pointer + 1; llvm::SmallVector ImplicitFirstprivate; - llvm::SmallVector ImplicitMap[OMPC_MAP_delete]; + llvm::SmallVector ImplicitMap[DefaultmapKindNum][OMPC_MAP_delete]; + llvm::SmallVector + ImplicitMapModifier[DefaultmapKindNum]; Sema::VarsWithInheritedDSAType VarsWithInheritedDSA; llvm::SmallDenseSet ImplicitDeclarations; @@ -3466,6 +3477,18 @@ } } } + if (SemaRef.getLangOpts().OpenMP > 50) { + bool IsModifierPresent = Stack->getDefaultmapModifier(ClauseKind) == + OMPC_DEFAULTMAP_MODIFIER_present; + if (IsModifierPresent) { + if (llvm::find(ImplicitMapModifier[ClauseKind], + OMPC_MAP_MODIFIER_present) == + std::end(ImplicitMapModifier[ClauseKind])) { + ImplicitMapModifier[ClauseKind].push_back( + OMPC_MAP_MODIFIER_present); + } + } + } if (isOpenMPTargetExecutionDirective(DKind) && !Stack->isLoopControlVariable(VD).first) { @@ -3506,7 +3529,7 @@ Stack->getDefaultmapModifier(ClauseKind); OpenMPMapClauseKind Kind = getMapClauseKindFromModifier( M, ClauseKind == OMPC_DEFAULTMAP_aggregate || Res); - ImplicitMap[Kind].emplace_back(E); + ImplicitMap[ClauseKind][Kind].emplace_back(E); } return; } @@ -3593,9 +3616,11 @@ OpenMPDefaultmapClauseModifier Modifier = Stack->getDefaultmapModifier(OMPC_DEFAULTMAP_aggregate); + OpenMPDefaultmapClauseKind ClauseKind = + getVariableCategoryFromDecl(SemaRef.getLangOpts(), FD); OpenMPMapClauseKind Kind = getMapClauseKindFromModifier( Modifier, /*IsAggregateOrDeclareTarget*/ true); - ImplicitMap[Kind].emplace_back(E); + ImplicitMap[ClauseKind][Kind].emplace_back(E); return; } @@ -3731,8 +3756,13 @@ ArrayRef getImplicitFirstprivate() const { return ImplicitFirstprivate; } - ArrayRef getImplicitMap(OpenMPDefaultmapClauseKind Kind) const { - return ImplicitMap[Kind]; + ArrayRef getImplicitMap(OpenMPDefaultmapClauseKind DK, + OpenMPMapClauseKind MK) const { + return ImplicitMap[DK][MK]; + } + ArrayRef + getImplicitMapModifier(OpenMPDefaultmapClauseKind Kind) const { + return ImplicitMapModifier[Kind]; } const Sema::VarsWithInheritedDSAType &getVarsWithInheritedDSA() const { return VarsWithInheritedDSA; @@ -5082,11 +5112,25 @@ SmallVector ImplicitFirstprivates( DSAChecker.getImplicitFirstprivate().begin(), DSAChecker.getImplicitFirstprivate().end()); - SmallVector ImplicitMaps[OMPC_MAP_delete]; - for (unsigned I = 0; I < OMPC_MAP_delete; ++I) { - ArrayRef ImplicitMap = - DSAChecker.getImplicitMap(static_cast(I)); - ImplicitMaps[I].append(ImplicitMap.begin(), ImplicitMap.end()); + const unsigned DefaultmapKindNum = OMPC_DEFAULTMAP_pointer + 1; + SmallVector ImplicitMaps[DefaultmapKindNum][OMPC_MAP_delete]; + SmallVector + ImplicitMapModifiers[DefaultmapKindNum]; + SmallVector + ImplicitMapModifiersLoc[DefaultmapKindNum]; + for (unsigned VC = 0; VC < DefaultmapKindNum; ++VC) { + auto Kind = static_cast(VC); + for (unsigned I = 0; I < OMPC_MAP_delete; ++I) { + ArrayRef ImplicitMap = DSAChecker.getImplicitMap( + Kind, static_cast(I)); + ImplicitMaps[VC][I].append(ImplicitMap.begin(), ImplicitMap.end()); + } + ArrayRef ImplicitModifier = + DSAChecker.getImplicitMapModifier(Kind); + ImplicitMapModifiers[VC].append(ImplicitModifier.begin(), + ImplicitModifier.end()); + std::fill_n(std::back_inserter(ImplicitMapModifiersLoc[VC]), + ImplicitModifier.size(), SourceLocation()); } // Mark taskgroup task_reduction descriptors as implicitly firstprivate. for (OMPClause *C : Clauses) { @@ -5112,23 +5156,26 @@ ErrorFound = true; } } - int ClauseKindCnt = -1; - for (ArrayRef ImplicitMap : ImplicitMaps) { - ++ClauseKindCnt; - if (ImplicitMap.empty()) - continue; - CXXScopeSpec MapperIdScopeSpec; - DeclarationNameInfo MapperId; - auto Kind = static_cast(ClauseKindCnt); - if (OMPClause *Implicit = ActOnOpenMPMapClause( - llvm::None, llvm::None, MapperIdScopeSpec, MapperId, Kind, - /*IsMapTypeImplicit=*/true, SourceLocation(), SourceLocation(), - ImplicitMap, OMPVarListLocTy())) { - ClausesWithImplicit.emplace_back(Implicit); - ErrorFound |= - cast(Implicit)->varlist_size() != ImplicitMap.size(); - } else { - ErrorFound = true; + for (unsigned I = 0, E = DefaultmapKindNum; I < E; ++I) { + int ClauseKindCnt = -1; + for (ArrayRef ImplicitMap : ImplicitMaps[I]) { + ++ClauseKindCnt; + if (ImplicitMap.empty()) + continue; + CXXScopeSpec MapperIdScopeSpec; + DeclarationNameInfo MapperId; + auto Kind = static_cast(ClauseKindCnt); + if (OMPClause *Implicit = ActOnOpenMPMapClause( + ImplicitMapModifiers[I], ImplicitMapModifiersLoc[I], + MapperIdScopeSpec, MapperId, Kind, /*IsMapTypeImplicit=*/true, + SourceLocation(), SourceLocation(), ImplicitMap, + OMPVarListLocTy())) { + ClausesWithImplicit.emplace_back(Implicit); + ErrorFound |= cast(Implicit)->varlist_size() != + ImplicitMap.size(); + } else { + ErrorFound = true; + } } } } @@ -18534,20 +18581,38 @@ bool isDefaultmapKind = (Kind != OMPC_DEFAULTMAP_unknown) || (LangOpts.OpenMP >= 50 && KindLoc.isInvalid()); if (!isDefaultmapKind || !isDefaultmapModifier) { - std::string ModifierValue = "'alloc', 'from', 'to', 'tofrom', " + StringRef KindValue = "'scalar', 'aggregate', 'pointer'"; + if (LangOpts.OpenMP == 50) { + StringRef ModifierValue = "'alloc', 'from', 'to', 'tofrom', " "'firstprivate', 'none', 'default'"; - std::string KindValue = "'scalar', 'aggregate', 'pointer'"; - if (!isDefaultmapKind && isDefaultmapModifier) { - Diag(KindLoc, diag::err_omp_unexpected_clause_value) - << KindValue << getOpenMPClauseName(OMPC_defaultmap); - } else if (isDefaultmapKind && !isDefaultmapModifier) { - Diag(MLoc, diag::err_omp_unexpected_clause_value) - << ModifierValue << getOpenMPClauseName(OMPC_defaultmap); + if (!isDefaultmapKind && isDefaultmapModifier) { + Diag(KindLoc, diag::err_omp_unexpected_clause_value) + << KindValue << getOpenMPClauseName(OMPC_defaultmap); + } else if (isDefaultmapKind && !isDefaultmapModifier) { + Diag(MLoc, diag::err_omp_unexpected_clause_value) + << ModifierValue << getOpenMPClauseName(OMPC_defaultmap); + } else { + Diag(MLoc, diag::err_omp_unexpected_clause_value) + << ModifierValue << getOpenMPClauseName(OMPC_defaultmap); + Diag(KindLoc, diag::err_omp_unexpected_clause_value) + << KindValue << getOpenMPClauseName(OMPC_defaultmap); + } } else { - Diag(MLoc, diag::err_omp_unexpected_clause_value) - << ModifierValue << getOpenMPClauseName(OMPC_defaultmap); - Diag(KindLoc, diag::err_omp_unexpected_clause_value) - << KindValue << getOpenMPClauseName(OMPC_defaultmap); + StringRef ModifierValue = + "'alloc', 'from', 'to', 'tofrom', " + "'firstprivate', 'none', 'default', 'present'"; + if (!isDefaultmapKind && isDefaultmapModifier) { + Diag(KindLoc, diag::err_omp_unexpected_clause_value) + << KindValue << getOpenMPClauseName(OMPC_defaultmap); + } else if (isDefaultmapKind && !isDefaultmapModifier) { + Diag(MLoc, diag::err_omp_unexpected_clause_value) + << ModifierValue << getOpenMPClauseName(OMPC_defaultmap); + } else { + Diag(MLoc, diag::err_omp_unexpected_clause_value) + << ModifierValue << getOpenMPClauseName(OMPC_defaultmap); + Diag(KindLoc, diag::err_omp_unexpected_clause_value) + << KindValue << getOpenMPClauseName(OMPC_defaultmap); + } } return nullptr; } diff --git a/clang/test/OpenMP/target_ast_print.cpp b/clang/test/OpenMP/target_ast_print.cpp --- a/clang/test/OpenMP/target_ast_print.cpp +++ b/clang/test/OpenMP/target_ast_print.cpp @@ -1073,4 +1073,56 @@ } #endif //OMP5 + +#ifdef OMP51 + +///==========================================================================/// +// RUN: %clang_cc1 -DOMP51 -verify -fopenmp -fopenmp-version=51 -ast-print %s | FileCheck %s --check-prefix OMP51 +// RUN: %clang_cc1 -DOMP51 -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -emit-pch -o %t %s +// RUN: %clang_cc1 -DOMP51 -fopenmp -fopenmp-version=51 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s --check-prefix OMP51 + +// RUN: %clang_cc1 -DOMP51 -verify -fopenmp-simd -fopenmp-version=51 -ast-print %s | FileCheck %s --check-prefix OMP51 +// RUN: %clang_cc1 -DOMP51 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -emit-pch -o %t %s +// RUN: %clang_cc1 -DOMP51 -fopenmp-simd -fopenmp-version=51 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s --check-prefix OMP51 + +void foo() {} + +template +T tmain(T argc, T *argv) { + #pragma omp target defaultmap(present: scalar) + foo(); + #pragma omp target defaultmap(present: aggregate) + foo(); + #pragma omp target defaultmap(present: pointer) + foo(); + + return 0; +} + +// OMP51: template T tmain(T argc, T *argv) { +// OMP51-NEXT: #pragma omp target defaultmap(present: scalar) +// OMP51-NEXT: foo() +// OMP51-NEXT: #pragma omp target defaultmap(present: aggregate) +// OMP51-NEXT: foo() +// OMP51-NEXT: #pragma omp target defaultmap(present: pointer) +// OMP51-NEXT: foo() + +// OMP51-LABEL: int main(int argc, char **argv) { +int main (int argc, char **argv) { +#pragma omp target defaultmap(present: scalar) +// OMP51-NEXT: #pragma omp target defaultmap(present: scalar) + foo(); +// OMP51-NEXT: foo(); +#pragma omp target defaultmap(present: aggregate) +// OMP51-NEXT: #pragma omp target defaultmap(present: aggregate) + foo(); +// OMP51-NEXT: foo(); +#pragma omp target defaultmap(present: pointer) +// OMP51-NEXT: #pragma omp target defaultmap(present: pointer) + foo(); +// OMP51-NEXT: foo(); + + return tmain(argc, &argc) + tmain(argv[0][0], argv[0]); +} +#endif #endif diff --git a/clang/test/OpenMP/target_defaultmap_codegen.cpp b/clang/test/OpenMP/target_defaultmap_codegen_01.cpp rename from clang/test/OpenMP/target_defaultmap_codegen.cpp rename to clang/test/OpenMP/target_defaultmap_codegen_01.cpp --- a/clang/test/OpenMP/target_defaultmap_codegen.cpp +++ b/clang/test/OpenMP/target_defaultmap_codegen_01.cpp @@ -1520,7 +1520,7 @@ // SIMD-ONLY18-NOT: {{__kmpc|__tgt}} #ifdef CK26 -// CK26-DAG: [[SIZES:@.+]] = {{.+}}constant [3 x i64] [i64 4096, i64 4, i64 {{.+}}] +// CK26-DAG: [[SIZES:@.+]] = {{.+}}constant [3 x i64] [i64 4, i64 4096, i64 {{.+}}] // Map types: OMP_MAP_TO | MAP_ALWAYS | OMP_MAP_IMPLICIT = 533 // CK26-DAG: [[TYPES:@.+]] = {{.+}}constant [3 x i64] [i64 531, i64 531, i64 531] diff --git a/clang/test/OpenMP/target_defaultmap_codegen_02.cpp b/clang/test/OpenMP/target_defaultmap_codegen_02.cpp new file mode 100644 --- /dev/null +++ b/clang/test/OpenMP/target_defaultmap_codegen_02.cpp @@ -0,0 +1,191 @@ +// expected-no-diagnostics +#ifndef HEADER +#define HEADER + +///==========================================================================/// +// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK1 --check-prefix CK1-64 +// RUN: %clang_cc1 -DCK1 -fopenmp -fopenmp-version=51 -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=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK1 --check-prefix CK1-64 +// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK1 --check-prefix CK1-32 +// RUN: %clang_cc1 -DCK1 -fopenmp -fopenmp-version=51 -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=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK1 --check-prefix CK1-32 + +// RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY6 %s +// RUN: %clang_cc1 -DCK1 -fopenmp-simd -fopenmp-version=51 -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=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY6 %s +// RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY6 %s +// RUN: %clang_cc1 -DCK1 -fopenmp-simd -fopenmp-version=51 -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=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY6 %s +// SIMD-ONLY6-NOT: {{__kmpc|__tgt}} +#ifdef CK1 + +// CK1-LABEL: @.__omp_offloading_{{.*}}implicit_present_scalar{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 + +// CK1-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 4] +// Map types: OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT | OMP_MAP_PRESENT = 4640 +// CK1-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 4640] + +// CK1-LABEL: implicit_present_scalar{{.*}}( +void implicit_present_scalar(int a) { + // CK1-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}, i8** null, i8** null) + // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 + // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 + // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 + // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 + // CK1-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i32** + // CK1-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32** + // CK1-DAG: store i32* [[DECL:%[^,]+]], i32** [[CBP1]] + // CK1-DAG: store i32* [[DECL]], i32** [[CP1]] + #pragma omp target defaultmap(present: scalar) + { + a += 1.0; + } +} + + +#endif +///==========================================================================/// +// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK2 --check-prefix CK2-64 +// RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-version=51 -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=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK2 --check-prefix CK2-64 +// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK2 --check-prefix CK2-32 +// RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-version=51 -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=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK2 --check-prefix CK2-32 + +// RUN: %clang_cc1 -DCK2 -verify -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY6 %s +// RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-version=51 -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=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY6 %s +// RUN: %clang_cc1 -DCK2 -verify -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY6 %s +// RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-version=51 -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=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY6 %s +// SIMD-ONLY6-NOT: {{__kmpc|__tgt}} +#ifdef CK2 + +// CK2-LABEL: @.__omp_offloading_{{.*}}implicit_present_aggregate{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 + +// CK2-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16] +// Map types: OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT | OMP_MAP_PRESENT = 4640 +// CK2-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 4640] + +// CK2-LABEL: implicit_present_aggregate{{.*}}( +void implicit_present_aggregate(int a) { + double darr[2] = {(double)a, (double)a}; + + // CK2-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}, i8** null, i8** null) + // CK2-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 + // CK2-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 + // CK2-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 + // CK2-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 + // CK2-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [2 x double]** + // CK2-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to [2 x double]** + // CK2-DAG: store [2 x double]* [[DECL:%[^,]+]], [2 x double]** [[CBP1]] + // CK2-DAG: store [2 x double]* [[DECL]], [2 x double]** [[CP1]] + + // CK2: call void [[KERNEL:@.+]]([2 x double]* [[DECL]]) + #pragma omp target defaultmap(present: aggregate) + { + darr[0] += 1.0; + darr[1] += 1.0; + } +} + + +#endif +///==========================================================================/// +// RUN: %clang_cc1 -DCK3 -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK3 +// RUN: %clang_cc1 -DCK3 -fopenmp -fopenmp-version=51 -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=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK3 +// RUN: %clang_cc1 -DCK3 -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK3 +// RUN: %clang_cc1 -DCK3 -fopenmp -fopenmp-version=51 -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=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK3 + +// RUN: %clang_cc1 -DCK3 -verify -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s +// RUN: %clang_cc1 -DCK3 -fopenmp-simd -fopenmp-version=51 -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=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s +// RUN: %clang_cc1 -DCK3 -verify -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s +// RUN: %clang_cc1 -DCK3 -fopenmp-simd -fopenmp-version=51 -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=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s +// SIMD-ONLY8-NOT: {{__kmpc|__tgt}} +#ifdef CK3 + + +// CK3-LABEL: @.__omp_offloading_{{.*}}explicit_present_pointer{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 + +// CK3: [[SIZE:@.+]] = private {{.*}}constant [1 x i64] [i64 {{8|4}}] +// Map types: OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT | OMP_MAP_PRESENT = 4640 +// CK3: [[MTYPE:@.+]] = private {{.*}}constant [1 x i64] [i64 4640] + +// CK3-LABEL: explicit_present_pointer{{.*}}( +void explicit_present_pointer() { + int *pa; + + // CK3-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE]]{{.+}}, i8** null, i8** null) + // CK3-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK3-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK3-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK3-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK3-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32*** + // CK3-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32*** + // CK3-DAG: store i32** [[VAR0:%.+]], i32*** [[CBP0]] + // CK3-DAG: store i32** [[VAR0]], i32*** [[CP0]] + + #pragma omp target defaultmap(present: pointer) + { + pa[50]++; + } +} + +#endif +#ifdef CK4 + +///==========================================================================/// +// RUN: %clang_cc1 -DCK4 -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK4 +// RUN: %clang_cc1 -DCK4 -fopenmp -fopenmp-version=51 -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=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK4 +// RUN: %clang_cc1 -DCK4 -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK4 +// RUN: %clang_cc1 -DCK4 -fopenmp -fopenmp-version=51 -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=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK4 + +// RUN: %clang_cc1 -DCK4 -verify -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY10 %s +// RUN: %clang_cc1 -DCK4 -fopenmp-simd -fopenmp-version=51 -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=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY10 %s +// RUN: %clang_cc1 -DCK4 -verify -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY10 %s +// RUN: %clang_cc1 -DCK4 -fopenmp-simd -fopenmp-version=51 -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=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY10 %s +// SIMD-ONLY10-NOT: {{__kmpc|__tgt}} + +// CK4-LABEL: @.__omp_offloading_{{.*}}implicit_present_double_complex{{.*}}.region_id = weak constant i8 0 + +// CK4-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16] +// Map types: OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT | OMP_MAP_PRESENT = 4640 +// CK4-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 4640] + +// CK4-LABEL: implicit_present_double_complex{{.*}}( +void implicit_present_double_complex (int a){ + double _Complex dc = (double)a; + + // CK4-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}, i8** null, i8** null) + // CK4-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 + // CK4-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 + // CK4-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 + // CK4-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 + // CK4-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to { double, double }** + // CK4-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to { double, double }** + // CK4-DAG: store { double, double }* [[PTR:%[^,]+]], { double, double }** [[CBP1]] + // CK4-DAG: store { double, double }* [[PTR]], { double, double }** [[CP1]] + + // CK4: call void [[KERNEL:@.+]]({ double, double }* [[PTR]]) + #pragma omp target defaultmap(present:scalar) + { + dc *= dc; + } +} + +// CK4: define internal void [[KERNEL]]({ double, double }* {{.*}}[[ARG:%.+]]) +// CK4: [[ADDR:%.+]] = alloca { double, double }*, +// CK4: store { double, double }* [[ARG]], { double, double }** [[ADDR]], +// CK4: [[REF:%.+]] = load { double, double }*, { double, double }** [[ADDR]], +// CK4: {{.+}} = getelementptr inbounds { double, double }, { double, double }* [[REF]], i32 0, i32 0 +#endif +#endif diff --git a/clang/test/OpenMP/target_defaultmap_messages.cpp b/clang/test/OpenMP/target_defaultmap_messages.cpp --- a/clang/test/OpenMP/target_defaultmap_messages.cpp +++ b/clang/test/OpenMP/target_defaultmap_messages.cpp @@ -1,3 +1,6 @@ +// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=51 %s -verify=expected,omp51 -Wuninitialized -DOMP51 +// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=51 %s -verify=expected,omp51 -Wuninitialized -DOMP51 + // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=50 %s -verify=expected,omp5 -Wuninitialized -DOMP5 // RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=50 %s -verify=expected,omp5 -Wuninitialized -DOMP5 @@ -27,23 +30,24 @@ T tmain(T argc, S **argv) { #pragma omp target defaultmap // expected-error {{expected '(' after 'defaultmap'}} foo(); -#pragma omp target defaultmap( // omp5-error {{expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default' in OpenMP clause 'defaultmap'}} expected-error {{expected ')'}} expected-note {{to match this '('}} omp45-error {{expected 'tofrom' in OpenMP clause 'defaultmap'}} +#pragma omp target defaultmap( // omp51-error {{expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default', 'present' in OpenMP clause 'defaultmap'}} omp5-error {{expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default' in OpenMP clause 'defaultmap'}} expected-error {{expected ')'}} expected-note {{to match this '('}} omp45-error {{expected 'tofrom' in OpenMP clause 'defaultmap'}} foo(); -#pragma omp target defaultmap() // omp5-error {{expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default' in OpenMP clause 'defaultmap'}} omp45-error {{expected 'tofrom' in OpenMP clause 'defaultmap'}} +#pragma omp target defaultmap() // omp51-error {{expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default', 'present' in OpenMP clause 'defaultmap'}} omp5-error {{expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default' in OpenMP clause 'defaultmap'}} omp45-error {{expected 'tofrom' in OpenMP clause 'defaultmap'}} foo(); #pragma omp target defaultmap(tofrom // expected-error {{expected ')'}} expected-note {{to match this '('}} omp45-warning {{missing ':' after defaultmap modifier - ignoring}} omp45-error {{expected 'scalar' in OpenMP clause 'defaultmap'}} foo(); - #pragma omp target defaultmap (tofrom: // expected-error {{expected ')'}} expected-note {{to match this '('}} omp5-error {{expected 'scalar', 'aggregate', 'pointer' in OpenMP clause 'defaultmap'}} omp45-error {{expected 'scalar' in OpenMP clause 'defaultmap'}} + #pragma omp target defaultmap (tofrom: // expected-error {{expected ')'}} expected-note {{to match this '('}} omp51-error {{expected 'scalar', 'aggregate', 'pointer' in OpenMP clause 'defaultmap'}} omp5-error {{expected 'scalar', 'aggregate', 'pointer' in OpenMP clause 'defaultmap'}} omp45-error {{expected 'scalar' in OpenMP clause 'defaultmap'}} foo(); #pragma omp target defaultmap(tofrom) // omp45-warning {{missing ':' after defaultmap modifier - ignoring}} omp45-error {{expected 'scalar' in OpenMP clause 'defaultmap'}} foo(); #pragma omp target defaultmap(tofrom, // expected-error {{expected ')'}} omp45-warning {{missing ':' after defaultmap modifier - ignoring}} expected-note {{to match this '('}} omp45-error {{expected 'scalar' in OpenMP clause 'defaultmap'}} foo(); - #pragma omp target defaultmap (scalar: // omp5-error {{expected 'scalar', 'aggregate', 'pointer' in OpenMP clause 'defaultmap'}} expected-error {{expected ')'}} omp5-error {{expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default' in OpenMP clause 'defaultmap'}} expected-note {{to match this '('}} omp45-error {{expected 'tofrom' in OpenMP clause 'defaultmap'}} + #pragma omp target defaultmap (scalar: // omp51-error {{expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default', 'present' in OpenMP clause 'defaultmap'}} omp51-error {{expected 'scalar', 'aggregate', 'pointer' in OpenMP clause 'defaultmap'}} omp5-error {{expected 'scalar', 'aggregate', 'pointer' in OpenMP clause 'defaultmap'}} expected-error {{expected ')'}} omp5-error {{expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default' in OpenMP clause 'defaultmap'}} expected-note {{to match this '('}} omp45-error {{expected 'tofrom' in OpenMP clause 'defaultmap'}} foo(); #pragma omp target defaultmap(tofrom, scalar // expected-error {{expected ')'}} omp45-warning {{missing ':' after defaultmap modifier - ignoring}} expected-note {{to match this '('}} omp45-error {{expected 'scalar' in OpenMP clause 'defaultmap'}} foo(); - #pragma omp target defaultmap(tofrom:scalar) defaultmap(tofrom:scalar) // omp45-error {{directive '#pragma omp target' cannot contain more than one 'defaultmap' clause}} omp5-error {{at most one defaultmap clause for each variable-category can appear on the directive}} + #pragma omp target defaultmap(tofrom:scalar) defaultmap(tofrom:scalar) // omp45-error {{directive '#pragma omp target' cannot contain more than one 'defaultmap' clause}} omp5-error {{at most one defaultmap clause for each variable-category can appear on the directive}} omp51-error {{at most one defaultmap clause for each variable-category can appear on the directive}} + foo(); #if OMP5 @@ -76,29 +80,36 @@ baz.bar.a += argc; // expected-error {{variable 'baz' must have explicitly specified data sharing attributes, data mapping attributes, or in an is_device_ptr clause}} #endif +#if OMP51 +#pragma omp target defaultmap(present: somthing) // omp51-error {{expected 'scalar', 'aggregate', 'pointer' in OpenMP clause 'defaultmap'}}} + foo(); +#pragma omp target defaultmap(present: scalar) defaultmap(present: scalar) // omp51-error {{at most one defaultmap clause for each variable-category can appear on the directive}} + foo(); +#endif + return argc; } int main(int argc, char **argv) { #pragma omp target defaultmap // expected-error {{expected '(' after 'defaultmap'}} foo(); -#pragma omp target defaultmap( // omp5-error {{expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default' in OpenMP clause 'defaultmap'}} expected-error {{expected ')'}} expected-note {{to match this '('}} omp45-error {{expected 'tofrom' in OpenMP clause 'defaultmap'}} +#pragma omp target defaultmap( // omp51-error {{expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default', 'present' in OpenMP clause 'defaultmap'}} omp5-error {{expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default' in OpenMP clause 'defaultmap'}} expected-error {{expected ')'}} expected-note {{to match this '('}} omp45-error {{expected 'tofrom' in OpenMP clause 'defaultmap'}} foo(); -#pragma omp target defaultmap() // omp5-error {{expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default' in OpenMP clause 'defaultmap'}} omp45-error {{expected 'tofrom' in OpenMP clause 'defaultmap'}} +#pragma omp target defaultmap() // omp51-error {{expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default', 'present' in OpenMP clause 'defaultmap'}} omp5-error {{expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default' in OpenMP clause 'defaultmap'}} omp45-error {{expected 'tofrom' in OpenMP clause 'defaultmap'}} foo(); #pragma omp target defaultmap(tofrom // expected-error {{expected ')'}} expected-note {{to match this '('}} omp45-warning {{missing ':' after defaultmap modifier - ignoring}} omp45-error {{expected 'scalar' in OpenMP clause 'defaultmap'}} foo(); -#pragma omp target defaultmap(tofrom: // expected-error {{expected ')'}} expected-note {{to match this '('}} omp5-error {{expected 'scalar', 'aggregate', 'pointer' in OpenMP clause 'defaultmap'}} omp45-error {{expected 'scalar' in OpenMP clause 'defaultmap'}} +#pragma omp target defaultmap(tofrom: // expected-error {{expected ')'}} expected-note {{to match this '('}} omp51-error {{expected 'scalar', 'aggregate', 'pointer' in OpenMP clause 'defaultmap'}} omp5-error {{expected 'scalar', 'aggregate', 'pointer' in OpenMP clause 'defaultmap'}} omp45-error {{expected 'scalar' in OpenMP clause 'defaultmap'}} foo(); #pragma omp target defaultmap(tofrom) // omp45-warning {{missing ':' after defaultmap modifier - ignoring}} omp45-error {{expected 'scalar' in OpenMP clause 'defaultmap'}} foo(); #pragma omp target defaultmap(tofrom, // expected-error {{expected ')'}} omp45-warning {{missing ':' after defaultmap modifier - ignoring}} expected-note {{to match this '('}} omp45-error {{expected 'scalar' in OpenMP clause 'defaultmap'}} foo(); -#pragma omp target defaultmap(scalar: // omp5-error {{expected 'scalar', 'aggregate', 'pointer' in OpenMP clause 'defaultmap'}} expected-error {{expected ')'}} omp5-error {{expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default' in OpenMP clause 'defaultmap'}} expected-note {{to match this '('}} omp45-error {{expected 'tofrom' in OpenMP clause 'defaultmap'}} +#pragma omp target defaultmap(scalar: // omp51-error {{expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default', 'present' in OpenMP clause 'defaultmap'}} omp51-error {{expected 'scalar', 'aggregate', 'pointer' in OpenMP clause 'defaultmap'}} omp5-error {{expected 'scalar', 'aggregate', 'pointer' in OpenMP clause 'defaultmap'}} expected-error {{expected ')'}} omp5-error {{expected 'alloc', 'from', 'to', 'tofrom', 'firstprivate', 'none', 'default' in OpenMP clause 'defaultmap'}} expected-note {{to match this '('}} omp45-error {{expected 'tofrom' in OpenMP clause 'defaultmap'}} foo(); #pragma omp target defaultmap(tofrom, scalar // expected-error {{expected ')'}} omp45-warning {{missing ':' after defaultmap modifier - ignoring}} expected-note {{to match this '('}} omp45-error {{expected 'scalar' in OpenMP clause 'defaultmap'}} foo(); -#pragma omp target defaultmap(tofrom: scalar) defaultmap(tofrom: scalar) // omp45-error {{directive '#pragma omp target' cannot contain more than one 'defaultmap' clause}} omp5-error {{at most one defaultmap clause for each variable-category can appear on the directive}} +#pragma omp target defaultmap(tofrom: scalar) defaultmap(tofrom: scalar) // omp45-error {{directive '#pragma omp target' cannot contain more than one 'defaultmap' clause}} omp5-error {{at most one defaultmap clause for each variable-category can appear on the directive}} omp51-error {{at most one defaultmap clause for each variable-category can appear on the directive}} foo(); #ifdef OMP5