diff --git a/clang/include/clang/AST/OpenMPClause.h b/clang/include/clang/AST/OpenMPClause.h --- a/clang/include/clang/AST/OpenMPClause.h +++ b/clang/include/clang/AST/OpenMPClause.h @@ -8547,10 +8547,11 @@ /// Clang specific specialization of the OMPContext to lookup target features. struct TargetOMPContext final : public llvm::omp::OMPContext { - TargetOMPContext(ASTContext &ASTCtx, std::function &&DiagUnknownTrait, - const FunctionDecl *CurrentFunctionDecl); + const FunctionDecl *CurrentFunctionDecl, + ArrayRef ConstructTraits); + virtual ~TargetOMPContext() = default; /// See llvm::omp::OMPContext::matchesISATrait diff --git a/clang/lib/AST/OpenMPClause.cpp b/clang/lib/AST/OpenMPClause.cpp --- a/clang/lib/AST/OpenMPClause.cpp +++ b/clang/lib/AST/OpenMPClause.cpp @@ -19,6 +19,7 @@ #include "clang/Basic/OpenMPKinds.h" #include "clang/Basic/TargetInfo.h" #include "llvm/ADT/SmallPtrSet.h" +#include "llvm/Frontend/OpenMP/OMPContext.h" #include "llvm/Support/Casting.h" #include "llvm/Support/ErrorHandling.h" #include @@ -2342,8 +2343,6 @@ getOpenMPContextTraitPropertyForSelector( Selector.Kind) && "Ill-formed construct selector!"); - - VMI.ConstructTraits.push_back(Selector.Properties.front().Kind); } } } @@ -2474,7 +2473,8 @@ TargetOMPContext::TargetOMPContext( ASTContext &ASTCtx, std::function &&DiagUnknownTrait, - const FunctionDecl *CurrentFunctionDecl) + const FunctionDecl *CurrentFunctionDecl, + ArrayRef ConstructTraits) : OMPContext(ASTCtx.getLangOpts().OpenMPIsDevice, ASTCtx.getTargetInfo().getTriple()), FeatureValidityCheck([&](StringRef FeatureName) { @@ -2482,6 +2482,9 @@ }), DiagUnknownTrait(std::move(DiagUnknownTrait)) { ASTCtx.getFunctionFeatureMap(FeatureMap, CurrentFunctionDecl); + + for (llvm::omp::TraitProperty Property : ConstructTraits) + addTrait(Property); } bool TargetOMPContext::matchesISATrait(StringRef RawString) const { diff --git a/clang/lib/Parse/ParseOpenMP.cpp b/clang/lib/Parse/ParseOpenMP.cpp --- a/clang/lib/Parse/ParseOpenMP.cpp +++ b/clang/lib/Parse/ParseOpenMP.cpp @@ -2046,8 +2046,10 @@ // improve the diagnostic location. Diag(Loc, diag::warn_unknown_begin_declare_variant_isa_trait) << ISATrait; }; - TargetOMPContext OMPCtx(ASTCtx, std::move(DiagUnknownTrait), - /* CurrentFunctionDecl */ nullptr); + TargetOMPContext OMPCtx( + ASTCtx, std::move(DiagUnknownTrait), + /* CurrentFunctionDecl */ nullptr, + /* ConstructTraits */ ArrayRef()); if (isVariantApplicableInContext(VMI, OMPCtx, /* DeviceSetOnly */ true)) { Actions.ActOnOpenMPBeginDeclareVariant(Loc, TI); 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 @@ -310,6 +310,8 @@ /// Vector of previously encountered target directives SmallVector TargetLocations; SourceLocation AtomicLocation; + /// Vector of declare variant construct traits. + SmallVector ConstructTraits; public: explicit DSAStackTy(Sema &S) : SemaRef(S) {} @@ -726,6 +728,24 @@ OMPC_DEFAULTMAP_MODIFIER_unknown; } + ArrayRef getConstructTraits() { + return ConstructTraits; + } + + void handleConstructTrait(llvm::omp::TraitProperty Property, bool Insert) { + if (Insert) + ConstructTraits.emplace_back(Property); + else { + SmallVector Traits( + std::move(ConstructTraits)); + std::reverse(Traits.begin(), Traits.end()); + llvm::omp::TraitProperty FirstTrait = Traits.pop_back_val(); + assert(FirstTrait == Property); + std::reverse_copy(Traits.begin(), Traits.end(), + std::back_inserter(ConstructTraits)); + } + } + DefaultDataSharingAttributes getDefaultDSA(unsigned Level) const { return getStackSize() <= Level ? DSA_unspecified : getStackElemAtLevel(Level).DefaultAttr; @@ -3871,6 +3891,26 @@ }; } // namespace +static void handleDeclareVariantConstructTrait(DSAStackTy *Stack, + OpenMPDirectiveKind DKind, + bool Insert) { + if (isOpenMPTargetExecutionDirective(DKind)) + Stack->handleConstructTrait( + llvm::omp::TraitProperty::construct_target_target, Insert); + if (isOpenMPTeamsDirective(DKind)) + Stack->handleConstructTrait(llvm::omp::TraitProperty::construct_teams_teams, + Insert); + if (isOpenMPParallelDirective(DKind)) + Stack->handleConstructTrait( + llvm::omp::TraitProperty::construct_parallel_parallel, Insert); + if (isOpenMPWorksharingDirective(DKind)) + Stack->handleConstructTrait(llvm::omp::TraitProperty::construct_for_for, + Insert); + if (isOpenMPSimdDirective(DKind)) + Stack->handleConstructTrait(llvm::omp::TraitProperty::construct_simd_simd, + Insert); +} + void Sema::ActOnOpenMPRegionStart(OpenMPDirectiveKind DKind, Scope *CurScope) { switch (DKind) { case OMPD_parallel: @@ -4285,6 +4325,7 @@ llvm_unreachable("Unknown OpenMP directive"); } DSAStack->setContext(CurContext); + handleDeclareVariantConstructTrait(DSAStack, DKind, /* Insert */ true); } int Sema::getNumberOfConstructScopes(unsigned Level) const { @@ -6223,6 +6264,8 @@ ErrorFound = Res.isInvalid() || ErrorFound; + handleDeclareVariantConstructTrait(DSAStack, Kind, /* Insert */ false); + // Check variables in the clauses if default(none) or // default(firstprivate) was specified. if (DSAStack->getDefaultDSA() == DSA_none || @@ -6804,7 +6847,7 @@ << ISATrait; }; TargetOMPContext OMPCtx(Context, std::move(DiagUnknownTrait), - getCurFunctionDecl()); + getCurFunctionDecl(), DSAStack->getConstructTraits()); QualType CalleeFnType = CalleeFnDecl->getType(); diff --git a/clang/test/OpenMP/declare_variant_construct_codegen_1.c b/clang/test/OpenMP/declare_variant_construct_codegen_1.c new file mode 100644 --- /dev/null +++ b/clang/test/OpenMP/declare_variant_construct_codegen_1.c @@ -0,0 +1,87 @@ +// RUN: %clang_cc1 -verify -fopenmp -triple x86_64-unknown-linux -emit-llvm %s -o - | FileCheck %s --check-prefix=CK1 +// RUN: %clang_cc1 -fopenmp -x c -triple x86_64-unknown-linux -emit-pch -o %t -fopenmp-version=45 %s +// RUN: %clang_cc1 -fopenmp -x c -triple x86_64-unknown-linux -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=45 | FileCheck %s --check-prefix=CK1 +// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c -triple x86_64-unknown-linux -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -o - | FileCheck %s --check-prefix=CK1 +// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -o - | FileCheck %s --check-prefix=CK1 +// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix=CK1 + +// RUN: %clang_cc1 -verify -fopenmp-simd -triple x86_64-unknown-linux -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}" +// RUN: %clang_cc1 -fopenmp-simd -x c -triple x86_64-unknown-linux -emit-pch -o %t -fopenmp-version=45 %s +// RUN: %clang_cc1 -fopenmp-simd -x c -triple x86_64-unknown-linux -include-pch %t -verify %s -emit-llvm -o - -fopenmp-version=45 | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}" +// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -x c -triple x86_64-unknown-linux -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}" +// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}" +// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -x c -triple x86_64-unknown-linux -fopenmp-targets=amdgcn-amd-amdhsa -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}" + +// expected-no-diagnostics + +#ifndef HEADER +#define HEADER + +#define N 100 + +void p_vxv(int *v1, int *v2, int *v3, int n); +void t_vxv(int *v1, int *v2, int *v3, int n); + +#pragma omp declare variant(t_vxv) match(construct={target}) +#pragma omp declare variant(p_vxv) match(construct={parallel}) +void vxv(int *v1, int *v2, int *v3, int n) { + for (int i = 0; i < n; i++) v3[i] = v1[i] * v2[i]; +} +// CK1: define dso_local void @vxv + +void p_vxv(int *v1, int *v2, int *v3, int n) { +#pragma omp for + for (int i = 0; i < n; i++) v3[i] = v1[i] * v2[i] * 3; +} +// CK1: define dso_local void @p_vxv + +#pragma omp declare target +void t_vxv(int *v1, int *v2, int *v3, int n) { +#pragma distribute simd + for (int i = 0; i < n; i++) v3[i] = v1[i] * v2[i] * 2; +} +#pragma omp end declare target +// CK1: define dso_local void @t_vxv + + +// CK1-LABEL: define {{[^@]+}}@main +int main() { + int v1[N], v2[N], v3[N]; + + // init + for (int i = 0; i < N; i++) { + v1[i] = (i + 1); + v2[i] = -(i + 1); + v3[i] = 0; + } + +#pragma omp target teams map(to: v1[:N],v2[:N]) map(from: v3[:N]) + { + vxv(v1, v2, v3, N); + } +// CK1: call void @__omp_offloading_[[OFFLOAD:.+]]({{.+}}) + + vxv(v1, v2, v3, N); +// CK1: call void @vxv + +#pragma omp parallel + { + vxv(v1, v2, v3, N); + } +// CK1: call void ({{.+}}) @__kmpc_fork_call(%struct.ident_t* {{.+}}, i32 3, void ({{.+}})* bitcast (void (i32*, i32*, [100 x i32]*, [100 x i32]*, [100 x i32]*)* [[PARALLEL_REGION:@.+]] to void + + return 0; +} + +// CK1: define internal void @__omp_offloading_[[OFFLOAD]]({{.+}}) +// CK1: call void ({{.+}}) @__kmpc_fork_teams(%struct.ident_t* {{.+}}, i32 3, void ({{.+}})* bitcast (void (i32*, i32*, [100 x i32]*, [100 x i32]*, [100 x i32]*)* [[TARGET_REGION:@.+]] to void + +// CK1: define internal void [[TARGET_REGION]]( +// CK1: call void @t_vxv + +// CK1: define internal void [[PARALLEL_REGION]]( +// CK1: call void @p_vxv + +#endif // HEADER