cchen updated this revision to Diff 372531.
cchen added a comment.
Remove braces
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D109635/new/
https://reviews.llvm.org/D109635
Files:
clang/include/clang/AST/OpenMPClause.h
clang/lib/AST/OpenMPClause.cpp
clang/lib/Parse/ParseOpenMP.cpp
clang/lib/Sema/SemaOpenMP.cpp
clang/test/OpenMP/declare_variant_construct_codegen_1.c
Index: clang/test/OpenMP/declare_variant_construct_codegen_1.c
===================================================================
--- /dev/null
+++ 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
Index: clang/lib/Sema/SemaOpenMP.cpp
===================================================================
--- clang/lib/Sema/SemaOpenMP.cpp
+++ clang/lib/Sema/SemaOpenMP.cpp
@@ -310,6 +310,8 @@
/// Vector of previously encountered target directives
SmallVector<SourceLocation, 2> TargetLocations;
SourceLocation AtomicLocation;
+ /// Vector of declare variant construct traits.
+ SmallVector<llvm::omp::TraitProperty, 8> ConstructTraits;
public:
explicit DSAStackTy(Sema &S) : SemaRef(S) {}
@@ -726,6 +728,24 @@
OMPC_DEFAULTMAP_MODIFIER_unknown;
}
+ ArrayRef<llvm::omp::TraitProperty> getConstructTraits() {
+ return ConstructTraits;
+ }
+
+ void handleConstructTrait(llvm::omp::TraitProperty Property, bool Insert) {
+ if (Insert)
+ ConstructTraits.emplace_back(Property);
+ else {
+ SmallVector<llvm::omp::TraitProperty, 8> 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();
Index: clang/lib/Parse/ParseOpenMP.cpp
===================================================================
--- clang/lib/Parse/ParseOpenMP.cpp
+++ 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<llvm::omp::TraitProperty>());
if (isVariantApplicableInContext(VMI, OMPCtx, /* DeviceSetOnly */ true)) {
Actions.ActOnOpenMPBeginDeclareVariant(Loc, TI);
Index: clang/lib/AST/OpenMPClause.cpp
===================================================================
--- clang/lib/AST/OpenMPClause.cpp
+++ 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 <algorithm>
@@ -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<void(StringRef)> &&DiagUnknownTrait,
- const FunctionDecl *CurrentFunctionDecl)
+ const FunctionDecl *CurrentFunctionDecl,
+ ArrayRef<llvm::omp::TraitProperty> 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 {
Index: clang/include/clang/AST/OpenMPClause.h
===================================================================
--- clang/include/clang/AST/OpenMPClause.h
+++ 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<void(StringRef)> &&DiagUnknownTrait,
- const FunctionDecl *CurrentFunctionDecl);
+ const FunctionDecl *CurrentFunctionDecl,
+ ArrayRef<llvm::omp::TraitProperty> ConstructTraits);
+
virtual ~TargetOMPContext() = default;
/// See llvm::omp::OMPContext::matchesISATrait
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
