alokmishra.besu created this revision.
alokmishra.besu added a project: OpenMP.
Herald added subscribers: llvm-commits, cfe-commits, dexonsmith, martong,
arphaman, guansong, hiraditya, yaxunl.
Herald added projects: clang, LLVM.
alokmishra.besu requested review of this revision.
Herald added a reviewer: jdoerfert.
Herald added a subscriber: sstefan1.
This patch supports OpenMP 5.0 metadirective features.
It is implemented keeping the OpenMP 5.1 features like dynamic user condition
in mind.
A new function, getBestWhenMatchForContext, is defined in
llvm/Frontend/OpenMP/OMPContext.h
Currently this function return the index of the when clause with the highest
score from the ones applicable in the Context.
But this function is declared with an array which can be used in OpenMP 5.1
implementation to select all the valid when clauses which can be resolved in
runtime. Currently this array is set to null by default and its implementation
is left for future.
Repository:
rG LLVM Github Monorepo
https://reviews.llvm.org/D91944
Files:
clang/include/clang-c/Index.h
clang/include/clang/AST/RecursiveASTVisitor.h
clang/include/clang/AST/StmtOpenMP.h
clang/include/clang/Basic/DiagnosticSemaKinds.td
clang/include/clang/Basic/StmtNodes.td
clang/include/clang/Sema/Sema.h
clang/include/clang/Serialization/ASTBitCodes.h
clang/lib/AST/OpenMPClause.cpp
clang/lib/AST/StmtOpenMP.cpp
clang/lib/AST/StmtPrinter.cpp
clang/lib/AST/StmtProfile.cpp
clang/lib/Basic/OpenMPKinds.cpp
clang/lib/CodeGen/CGOpenMPRuntime.cpp
clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
clang/lib/CodeGen/CGStmt.cpp
clang/lib/CodeGen/CGStmtOpenMP.cpp
clang/lib/CodeGen/CodeGenFunction.h
clang/lib/Parse/ParseOpenMP.cpp
clang/lib/Sema/SemaExceptionSpec.cpp
clang/lib/Sema/SemaOpenMP.cpp
clang/lib/Sema/TreeTransform.h
clang/lib/Serialization/ASTReaderStmt.cpp
clang/lib/Serialization/ASTWriterStmt.cpp
clang/lib/StaticAnalyzer/Core/ExprEngine.cpp
clang/test/OpenMP/metadirective_ast_print.cpp
clang/test/OpenMP/metadirective_codegen.cpp
clang/test/OpenMP/metadirective_construct.cpp
clang/test/OpenMP/metadirective_empty.cpp
clang/test/OpenMP/metadirective_implementation.cpp
clang/tools/libclang/CIndex.cpp
clang/tools/libclang/CXCursor.cpp
llvm/include/llvm/Frontend/OpenMP/OMP.td
llvm/include/llvm/Frontend/OpenMP/OMPContext.h
llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
llvm/lib/Frontend/OpenMP/OMPContext.cpp
Index: llvm/lib/Frontend/OpenMP/OMPContext.cpp
===================================================================
--- llvm/lib/Frontend/OpenMP/OMPContext.cpp
+++ llvm/lib/Frontend/OpenMP/OMPContext.cpp
@@ -332,6 +332,73 @@
return Score;
}
+int llvm::omp::getBestWhenMatchForContext(
+ const SmallVectorImpl<VariantMatchInfo> &VMIs, const OMPContext &Ctx,
+ SmallVectorImpl<unsigned> *OrderedMatch) {
+
+ APInt BestScore(64, 0);
+ int BestVMIIdx = -1;
+ const VariantMatchInfo *BestVMI = nullptr;
+
+ for (unsigned u = 0, e = VMIs.size(); u < e; ++u) {
+ const VariantMatchInfo &VMI = VMIs[u];
+
+ SmallVector<unsigned, 8> ConstructMatches;
+ // Check if its clearly not the best.
+ APInt Score = getVariantMatchScore(VMI, Ctx, ConstructMatches);
+ if (Score.ult(BestScore))
+ continue;
+ // Equal score need subset checks.
+ if (Score.eq(BestScore)) {
+ // Strict subset are never best.
+ if (isStrictSubset(VMI, *BestVMI))
+ continue;
+ // Same score and the current best is no strict subset so we keep it.
+ if (!isStrictSubset(*BestVMI, VMI))
+ continue;
+ }
+ // New best found.
+ BestVMI = &VMI;
+ BestVMIIdx = u;
+ BestScore = Score;
+ }
+
+ return BestVMIIdx;
+}
+
+/*int llvm::omp::getBestWhenMatchForContext(
+ const SmallVectorImpl<VariantMatchInfo> &VMIs, const OMPContext &Ctx) {
+
+ APInt BestScore(64, 0);
+ int BestVMIIdx = -1;
+ const VariantMatchInfo *BestVMI = nullptr;
+
+ for (unsigned u = 0, e = VMIs.size(); u < e; ++u) {
+ const VariantMatchInfo &VMI = VMIs[u];
+
+ SmallVector<unsigned, 8> ConstructMatches;
+ // Check if its clearly not the best.
+ APInt Score = getVariantMatchScore(VMI, Ctx, ConstructMatches);
+ if (Score.ult(BestScore))
+ continue;
+ // Equal score need subset checks.
+ if (Score.eq(BestScore)) {
+ // Strict subset are never best.
+ if (isStrictSubset(VMI, *BestVMI))
+ continue;
+ // Same score and the current best is no strict subset so we keep it.
+ if (!isStrictSubset(*BestVMI, VMI))
+ continue;
+ }
+ // New best found.
+ BestVMI = &VMI;
+ BestVMIIdx = u;
+ BestScore = Score;
+ }
+
+ return BestVMIIdx;
+}*/
+
int llvm::omp::getBestVariantMatchForContext(
const SmallVectorImpl<VariantMatchInfo> &VMIs, const OMPContext &Ctx) {
Index: llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
===================================================================
--- llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
+++ llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
@@ -115,6 +115,7 @@
__OMP_CLAUSE_NO_CLASS(uniform)
__OMP_CLAUSE_NO_CLASS(device_type)
__OMP_CLAUSE_NO_CLASS(match)
+__OMP_CLAUSE_NO_CLASS(when)
__OMP_IMPLICIT_CLAUSE_CLASS(depobj, "depobj", OMPDepobjClause)
__OMP_IMPLICIT_CLAUSE_CLASS(flush, "flush", OMPFlushClause)
Index: llvm/include/llvm/Frontend/OpenMP/OMPContext.h
===================================================================
--- llvm/include/llvm/Frontend/OpenMP/OMPContext.h
+++ llvm/include/llvm/Frontend/OpenMP/OMPContext.h
@@ -184,6 +184,17 @@
const OMPContext &Ctx,
bool DeviceSetOnly = false);
+/// Return the index (into \p VMIs) of the When clause with the highest score
+/// from the ones applicble in \p Ctx.
+/// In OpenMP 5.1 set OrderedMatch to those conditions which need runtime
+/// resolution.
+int getBestWhenMatchForContext(
+ const SmallVectorImpl<VariantMatchInfo> &VMIs, const OMPContext &Ctx,
+ SmallVectorImpl<unsigned> *OrderedMatch = nullptr);
+/// Return the index (into \p VMIs) of the When clause with the highest score
+/// from the ones applicble in \p Ctx.
+// int getBestWhenMatchForContext(const SmallVectorImpl<VariantMatchInfo> &VMIs,
+// const OMPContext &Ctx);
/// Return the index (into \p VMIs) of the variant with the highest score
/// from the ones applicble in \p Ctx. See llvm::isVariantApplicableInContext.
int getBestVariantMatchForContext(const SmallVectorImpl<VariantMatchInfo> &VMIs,
Index: llvm/include/llvm/Frontend/OpenMP/OMP.td
===================================================================
--- llvm/include/llvm/Frontend/OpenMP/OMP.td
+++ llvm/include/llvm/Frontend/OpenMP/OMP.td
@@ -59,6 +59,7 @@
let clangClass = "OMPCollapseClause";
let flangClassValue = "ScalarIntConstantExpr";
}
+def OMPC_When: Clause<"when"> {}
def OMPC_Default : Clause<"default"> {
let clangClass = "OMPDefaultClause";
let flangClass = "OmpDefaultClause";
@@ -320,6 +321,14 @@
// Definition of OpenMP directives
//===----------------------------------------------------------------------===//
+def OMP_Metadirective : Directive<"metadirective"> {
+ let allowedClauses = [
+ VersionedClause<OMPC_When>
+ ];
+ let allowedOnceClauses = [
+ VersionedClause<OMPC_Default>
+ ];
+}
def OMP_ThreadPrivate : Directive<"threadprivate"> {}
def OMP_Parallel : Directive<"parallel"> {
let allowedClauses = [
Index: clang/tools/libclang/CXCursor.cpp
===================================================================
--- clang/tools/libclang/CXCursor.cpp
+++ clang/tools/libclang/CXCursor.cpp
@@ -639,6 +639,9 @@
case Stmt::MSDependentExistsStmtClass:
K = CXCursor_UnexposedStmt;
break;
+ case Stmt::OMPMetaDirectiveClass:
+ K = CXCursor_OMPMetaDirective;
+ break;
case Stmt::OMPParallelDirectiveClass:
K = CXCursor_OMPParallelDirective;
break;
Index: clang/tools/libclang/CIndex.cpp
===================================================================
--- clang/tools/libclang/CIndex.cpp
+++ clang/tools/libclang/CIndex.cpp
@@ -5525,6 +5525,8 @@
return cxstring::createRef("CXXAccessSpecifier");
case CXCursor_ModuleImportDecl:
return cxstring::createRef("ModuleImport");
+ case CXCursor_OMPMetaDirective:
+ return cxstring::createRef("OMPMetaDirective");
case CXCursor_OMPParallelDirective:
return cxstring::createRef("OMPParallelDirective");
case CXCursor_OMPSimdDirective:
Index: clang/test/OpenMP/metadirective_implementation.cpp
===================================================================
--- /dev/null
+++ clang/test/OpenMP/metadirective_implementation.cpp
@@ -0,0 +1,36 @@
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -o - | FileCheck %s
+// expected-no-diagnostics
+void func() {
+#pragma omp target
+ // Test to check implementation and device.
+ // Since we are building with llvm for nvptx64 device, 512 teams with thread_limit of 32 are generated.
+#pragma omp metadirective \
+ when(implementation = {vendor(llvm)}, device = {arch(nvptx64)} \
+ : teams num_teams(512) thread_limit(32)) \
+ when(implementation = {vendor(amd)}, device = {arch(amdgcn)} \
+ : teams num_teams(512) thread_limit(64)) default(teams)
+#pragma omp distribute parallel for
+ for (int i = 0; i < 10000; i++)
+ ;
+}
+
+// CHECK: define void @_Z4funcv() #0 {
+// CHECK: entry:
+// CHECK: %tmp = alloca i32, align 4
+// CHECK: call void @__kmpc_push_target_tripcount
+// CHECK: %0 = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @3, i64 -1, i8* [[OMP_OUTLINED:@.+]].region_id,
+// CHECK: %1 = icmp ne i32 %0, 0
+// CHECK: br i1 %1, label %omp_offload.failed, label %omp_offload.cont
+// CHECK: omp_offload.failed: ; preds = %entry
+// CHECK: call void [[OMP_OUTLINED:@.+]] #2
+// CHECK: br label %omp_offload.cont
+// CHECK: omp_offload.cont: ; preds = %omp_offload.failed, %entry
+// CHECK: ret void
+// CHECK: }
+// CHECK: define internal void [[OMP_OUTLINED:@.+]] #1 {
+// CHECK: entry:
+// CHECK: %0 = call i32 @__kmpc_global_thread_num(%struct.ident_t* @3)
+// CHECK: call void @__kmpc_push_num_teams(%struct.ident_t* @3, i32 %0, i32 512, i32 32)
+// CHECK: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* @3, i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* [[OMP_OUTLINED:@.+]] to void (i32*, i32*, ...)*))
+// CHECK: ret void
+// CHECK: }
Index: clang/test/OpenMP/metadirective_empty.cpp
===================================================================
--- /dev/null
+++ clang/test/OpenMP/metadirective_empty.cpp
@@ -0,0 +1,31 @@
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -o - | FileCheck %s
+// expected-no-diagnostics
+#define N 1000
+void func() {
+ // Test where a valid when clause contains empty directive.
+ // The directive will be ignored and code for a serial for loop will be generated.
+#pragma omp metadirective when(implementation = {vendor(llvm)} \
+ :) default(parallel for)
+ for (int i = 0; i < N; i++)
+ ;
+}
+
+// CHECK: define void @_Z4funcv() #0 {
+// CHECK: entry:
+// CHECK: %i = alloca i32, align 4
+// CHECK: store i32 0, i32* %i, align 4
+// CHECK: br label %for.cond
+// CHECK: for.cond:
+// CHECK: %0 = load i32, i32* %i, align 4
+// CHECK: %cmp = icmp slt i32 %0, 1000
+// CHECK: br i1 %cmp, label %for.body, label %for.end
+// CHECK: for.body:
+// CHECK: br label %for.inc
+// CHECK: for.inc:
+// CHECK: %1 = load i32, i32* %i, align 4
+// CHECK: %inc = add nsw i32 %1, 1
+// CHECK: store i32 %inc, i32* %i, align 4
+// CHECK: br label %for.cond, !llvm.loop !2
+// CHECK: for.end:
+// CHECK: ret void
+// CHECK: }
Index: clang/test/OpenMP/metadirective_construct.cpp
===================================================================
--- /dev/null
+++ clang/test/OpenMP/metadirective_construct.cpp
@@ -0,0 +1,106 @@
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -o - | FileCheck %s
+// expected-no-diagnostics
+
+#define N 1000
+
+#pragma omp declare target
+void func1() {
+ // Test to check the case where metadirective is called within a target construct.
+ // Here the directive for 'distribute parallel for' will be generated
+#pragma omp metadirective \
+ when(construct = {"target"} \
+ : distribute parallel for) default()
+ for (int i = 0; i < N; i++)
+ ;
+}
+#pragma omp end declare target
+
+void func() {
+#pragma omp target teams
+ func1();
+}
+
+// CHECK: define void @_Z5func1v() #0 {
+// CHECK: entry:
+// CHECK: %.omp.iv = alloca i32, align 4
+// CHECK: %tmp = alloca i32, align 4
+// CHECK: %.omp.comb.lb = alloca i32, align 4
+// CHECK: %.omp.comb.ub = alloca i32, align 4
+// CHECK: %.omp.stride = alloca i32, align 4
+// CHECK: %.omp.is_last = alloca i32, align 4
+// CHECK: %i = alloca i32, align 4
+// CHECK: %{{[0-9]*}} = call i32 @__kmpc_global_thread_num(%struct.ident_t* @2)
+// CHECK: store i32 0, i32* %.omp.comb.lb, align 4
+// CHECK: store i32 999, i32* %.omp.comb.ub, align 4
+// CHECK: store i32 1, i32* %.omp.stride, align 4
+// CHECK: store i32 0, i32* %.omp.is_last, align 4
+// CHECK: call void @__kmpc_for_static_init_4(%struct.ident_t* @1, i32 %{{[0-9]*}}, i32 92, i32* %.omp.is_last, i32* %.omp.comb.lb, i32* %.omp.comb.ub, i32* %.omp.stride, i32 1, i32 1)
+// CHECK: %{{[0-9]*}} = load i32, i32* %.omp.comb.ub, align 4
+// CHECK: %cmp = icmp sgt i32 %{{[0-9]*}}, 999
+// CHECK: br i1 %cmp, label %cond.true, label %cond.false
+// CHECK: cond.true:
+// CHECK: br label %cond.end
+// CHECK: cond.false:
+// CHECK: %{{[0-9]*}} = load i32, i32* %.omp.comb.ub, align 4
+// CHECK: br label %cond.end
+// CHECK: cond.end:
+// CHECK: %cond = phi i32 [ 999, %cond.true ], [ %{{[0-9]*}}, %cond.false ]
+// CHECK: store i32 %cond, i32* %.omp.comb.ub, align 4
+// CHECK: %{{[0-9]*}} = load i32, i32* %.omp.comb.lb, align 4
+// CHECK: store i32 %{{[0-9]*}}, i32* %.omp.iv, align 4
+// CHECK: br label %omp.inner.for.cond
+// CHECK: omp.inner.for.cond:
+// CHECK: %{{[0-9]*}} = load i32, i32* %.omp.iv, align 4
+// CHECK: %{{[0-9]*}} = load i32, i32* %.omp.comb.ub, align 4
+// CHECK: %cmp1 = icmp sle i32 %{{[0-9]*}}, %{{[0-9]*}}
+// CHECK: br i1 %cmp1, label %omp.inner.for.body, label %omp.inner.for.end
+// CHECK: omp.inner.for.body:
+// CHECK: %{{[0-9]*}} = load i32, i32* %.omp.comb.lb, align 4
+// CHECK: %{{[0-9]*}} = zext i32 %{{[0-9]*}} to i64
+// CHECK: %{{[0-9]*}} = load i32, i32* %.omp.comb.ub, align 4
+// CHECK: %{{[0-9]*}} = zext i32 %{{[0-9]*}} to i64
+// CHECK: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* @2, i32 2, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i64, i64)* [[OMP_OUTLINED:@.+]] to void (i32*, i32*, ...)*), i64 %{{[0-9]*}}, i64 %{{[0-9]*}})
+// CHECK: br label %omp.inner.for.inc
+// CHECK: omp.inner.for.inc:
+// CHECK: %{{[0-9]*}} = load i32, i32* %.omp.iv, align 4
+// CHECK: %{{[0-9]*}} = load i32, i32* %.omp.stride, align 4
+// CHECK: %add = add nsw i32 %{{[0-9]*}}, %{{[0-9]*}}
+// CHECK: store i32 %add, i32* %.omp.iv, align 4
+// CHECK: br label %omp.inner.for.cond
+// CHECK: omp.inner.for.end:
+// CHECK: br label %omp.loop.exit
+// CHECK: omp.loop.exit:
+// CHECK: call void @__kmpc_for_static_fini(%struct.ident_t* @1, i32 %{{[0-9]*}})
+// CHECK: ret void
+// CHECK: }
+
+// Declaration of funtion handling distribute parallel for
+// CHECK: define internal void [[OMP_OUTLINED:@.+]](i32* noalias %.global_tid., i32* noalias %.bound_tid., i64 %.previous.lb., i64 %.previous.ub.)
+
+// CHECK: define void @_Z4funcv() #0 {
+// CHECK: entry:
+// CHECK: %{{[0-9]*}} = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @2, i64 -1, i8* [[OMP_OFFLOADING:@.+]].region_id
+// CHECK: %{{[0-9]*}} = icmp ne i32 %{{[0-9]*}}, 0
+// CHECK: br i1 %{{[0-9]*}}, label %omp_offload.failed, label %omp_offload.cont
+// CHECK: omp_offload.failed:
+// CHECK: call void [[OMP_OFFLOADING:@.+]]() #1
+// CHECK: br label %omp_offload.cont
+// CHECK: omp_offload.cont:
+// CHECK: ret void
+// CHECK: }
+
+// CHECK: define internal void [[OMP_OFFLOADING:@.+]]() #2 {
+// CHECK: entry:
+// CHECK: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* @2, i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* @.omp_outlined..1 to void (i32*, i32*, ...)*))
+// CHECK: ret void
+// CHECK: }
+
+// CHECK: define internal void @.omp_outlined..1(i32* noalias %.global_tid., i32* noalias %.bound_tid.) #2 {
+// CHECK: entry:
+// CHECK: %.global_tid..addr = alloca i32*, align 8
+// CHECK: %.bound_tid..addr = alloca i32*, align 8
+// CHECK: store i32* %.global_tid., i32** %.global_tid..addr, align 8
+// CHECK: store i32* %.bound_tid., i32** %.bound_tid..addr, align 8
+// CHECK: call void @_Z5func1v()
+// CHECK: ret void
+// CHECK: }
Index: clang/test/OpenMP/metadirective_codegen.cpp
===================================================================
--- /dev/null
+++ clang/test/OpenMP/metadirective_codegen.cpp
@@ -0,0 +1,58 @@
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -o - | FileCheck %s
+// expected-no-diagnostics
+
+#define N 10000
+int func() {
+ int v1[N], v2[N], v3[N];
+ for (int i = 0; i < N; i++) {
+ v1[i] = (i + 1);
+ v2[i] = -(i + 1);
+ }
+
+#pragma omp metadirective \
+ when(device = {arch("nvptx64")} \
+ : target teams distribute parallel for map(v1 [0:N], v2 [0:N], v3 [0:N])) default()
+ for (int i = 0; i < N; i++) {
+ v3[i] = v1[i] * v2[i];
+ }
+
+ return 0;
+}
+// CHECK: %arrayidx4 = getelementptr inbounds [10000 x i32], [10000 x i32]* %v3, i64 0, i64 0
+// CHECK: %arrayidx5 = getelementptr inbounds [10000 x i32], [10000 x i32]* %v1, i64 0, i64 0
+// CHECK: %arrayidx6 = getelementptr inbounds [10000 x i32], [10000 x i32]* %v2, i64 0, i64 0
+// CHECK: %{{[0-9]*}} = getelementptr inbounds [3 x i8*], [3 x i8*]* %.offload_baseptrs, i32 0, i32 0
+// CHECK: %{{[0-9]*}} = bitcast i8** %{{[0-9]*}} to [10000 x i32]**
+// CHECK: store [10000 x i32]* %v3, [10000 x i32]** %{{[0-9]*}}, align 8
+// CHECK: %{{[0-9]*}} = getelementptr inbounds [3 x i8*], [3 x i8*]* %.offload_ptrs, i32 0, i32 0
+// CHECK: %{{[0-9]*}} = bitcast i8** %{{[0-9]*}} to i32**
+// CHECK: store i32* %arrayidx4, i32** %{{[0-9]*}}, align 8
+// CHECK: %{{[0-9]*}} = getelementptr inbounds [3 x i8*], [3 x i8*]* %.offload_mappers, i64 0, i64 0
+// CHECK: store i8* null, i8** %{{[0-9]*}}, align 8
+// CHECK: %{{[0-9]*}} = getelementptr inbounds [3 x i8*], [3 x i8*]* %.offload_baseptrs, i32 0, i32 1
+// CHECK: %{{[0-9]*}} = bitcast i8** %{{[0-9]*}} to [10000 x i32]**
+// CHECK: store [10000 x i32]* %v1, [10000 x i32]** %{{[0-9]*}}, align 8
+// CHECK: %{{[0-9]*}} = getelementptr inbounds [3 x i8*], [3 x i8*]* %.offload_ptrs, i32 0, i32 1
+// CHECK: %{{[0-9]*}} = bitcast i8** %{{[0-9]*}} to i32**
+// CHECK: store i32* %arrayidx5, i32** %{{[0-9]*}}, align 8
+// CHECK: %{{[0-9]*}} = getelementptr inbounds [3 x i8*], [3 x i8*]* %.offload_mappers, i64 0, i64 1
+// CHECK: store i8* null, i8** %{{[0-9]*}}, align 8
+// CHECK: %{{[0-9]*}} = getelementptr inbounds [3 x i8*], [3 x i8*]* %.offload_baseptrs, i32 0, i32 2
+// CHECK: %{{[0-9]*}} = bitcast i8** %{{[0-9]*}} to [10000 x i32]**
+// CHECK: store [10000 x i32]* %v2, [10000 x i32]** %{{[0-9]*}}, align 8
+// CHECK: %{{[0-9]*}} = getelementptr inbounds [3 x i8*], [3 x i8*]* %.offload_ptrs, i32 0, i32 2
+// CHECK: %{{[0-9]*}} = bitcast i8** %{{[0-9]*}} to i32**
+// CHECK: store i32* %arrayidx6, i32** %{{[0-9]*}}, align 8
+// CHECK: %{{[0-9]*}} = getelementptr inbounds [3 x i8*], [3 x i8*]* %.offload_mappers, i64 0, i64 2
+// CHECK: store i8* null, i8** %{{[0-9]*}}, align 8
+// CHECK: %{{[0-9]*}} = getelementptr inbounds [3 x i8*], [3 x i8*]* %.offload_baseptrs, i32 0, i32 0
+// CHECK: %{{[0-9]*}} = getelementptr inbounds [3 x i8*], [3 x i8*]* %.offload_ptrs, i32 0, i32 0
+// CHECK: call void @__kmpc_push_target_tripcount(%struct.ident_t* @3, i64 -1, i64 10000)
+// CHECK: %{{[0-9]*}} = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @3, i64 -1, i8* [[OMP_OUTLINED:@.+]].region_id, i32 3, i8** %{{[0-9]*}}, i8** %{{[0-9]*}}, i64* getelementptr inbounds ([3 x i64], [3 x i64]* @.offload_sizes, i32 0, i32 0), i64* getelementptr inbounds ([3 x i64], [3 x i64]* @.offload_maptypes, i32 0, i32 0), i8** null, i8** null, i32 0, i32 0)
+// CHECK: %{{[0-9]*}} = icmp ne i32 %{{[0-9]*}}, 0
+// CHECK: br i1 %{{[0-9]*}}, label %omp_offload.failed, label %omp_offload.cont
+// CHECK: omp_offload.failed:
+// CHECK: call void [[OMP_OUTLINED:@.+]]
+// CHECK: br label %omp_offload.cont
+// CHECK: omp_offload.cont:
+// CHECK: ret i32 0
Index: clang/test/OpenMP/metadirective_ast_print.cpp
===================================================================
--- /dev/null
+++ clang/test/OpenMP/metadirective_ast_print.cpp
@@ -0,0 +1,26 @@
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda -ast-print %s -o - | FileCheck %s
+// expected-no-diagnostics
+#define N 10000
+int main() {
+ int v1[N], v2[N], v3[N];
+ for (int i = 0; i < N; i++) {
+ v1[i] = (i + 1);
+ v2[i] = -(i + 1);
+ }
+#pragma omp metadirective \
+ when(device = {arch("nvptx64")} \
+ : target teams distribute parallel for map(v1 [0:N], v2 [0:N], v3 [0:N])) default()
+ for (int i = 0; i < N; i++) {
+ v3[i] = v1[i] * v2[i];
+ }
+ return 0;
+}
+// CHECK: int v1[10000], v2[10000], v3[10000];
+// CHECK: for (int i = 0; i < 10000; i++) {
+// CHECK: v1[i] = (i + 1);
+// CHECK: v2[i] = -(i + 1);
+// CHECK: }
+// CHECK: #pragma omp target teams distribute parallel for map(tofrom: v1[0:10000],v2[0:10000],v3[0:10000])
+// CHECK: for (int i = 0; i < 10000; i++) {
+// CHECK: v3[i] = v1[i] * v2[i];
+// CHECK: }
Index: clang/lib/StaticAnalyzer/Core/ExprEngine.cpp
===================================================================
--- clang/lib/StaticAnalyzer/Core/ExprEngine.cpp
+++ clang/lib/StaticAnalyzer/Core/ExprEngine.cpp
@@ -1292,6 +1292,7 @@
case Stmt::OMPTargetTeamsDistributeParallelForDirectiveClass:
case Stmt::OMPTargetTeamsDistributeParallelForSimdDirectiveClass:
case Stmt::OMPTargetTeamsDistributeSimdDirectiveClass:
+ case Stmt::OMPMetaDirectiveClass:
case Stmt::CapturedStmtClass: {
const ExplodedNode *node = Bldr.generateSink(S, Pred, Pred->getState());
Engine.addAbortedBlock(node, currBldrCtx->getBlock());
Index: clang/lib/Serialization/ASTWriterStmt.cpp
===================================================================
--- clang/lib/Serialization/ASTWriterStmt.cpp
+++ clang/lib/Serialization/ASTWriterStmt.cpp
@@ -2176,6 +2176,13 @@
Record.AddSourceLocation(E->getEndLoc());
}
+void ASTStmtWriter::VisitOMPMetaDirective(OMPMetaDirective *D) {
+ VisitStmt(D);
+ Record.push_back(D->getNumClauses());
+ VisitOMPExecutableDirective(D);
+ Code = serialization::STMT_OMP_META_DIRECTIVE;
+}
+
void ASTStmtWriter::VisitOMPLoopDirective(OMPLoopDirective *D) {
VisitStmt(D);
Record.writeUInt32(D->getCollapsedNumber());
Index: clang/lib/Serialization/ASTReaderStmt.cpp
===================================================================
--- clang/lib/Serialization/ASTReaderStmt.cpp
+++ clang/lib/Serialization/ASTReaderStmt.cpp
@@ -2286,6 +2286,13 @@
VisitOMPExecutableDirective(D);
}
+void ASTStmtReader::VisitOMPMetaDirective(OMPMetaDirective *D) {
+ VisitStmt(D);
+ // The NumClauses field was read in ReadStmtFromStream.
+ Record.skipInts(1);
+ VisitOMPExecutableDirective(D);
+}
+
void ASTStmtReader::VisitOMPParallelDirective(OMPParallelDirective *D) {
VisitStmt(D);
VisitOMPExecutableDirective(D);
@@ -3130,6 +3137,11 @@
nullptr);
break;
+ case STMT_OMP_META_DIRECTIVE:
+ S = OMPMetaDirective::CreateEmpty(
+ Context, Record[ASTStmtReader::NumStmtFields], Empty);
+ break;
+
case STMT_OMP_PARALLEL_DIRECTIVE:
S =
OMPParallelDirective::CreateEmpty(Context,
Index: clang/lib/Sema/TreeTransform.h
===================================================================
--- clang/lib/Sema/TreeTransform.h
+++ clang/lib/Sema/TreeTransform.h
@@ -8379,6 +8379,17 @@
template <typename Derived>
StmtResult
+TreeTransform<Derived>::TransformOMPMetaDirective(OMPMetaDirective *D) {
+ DeclarationNameInfo DirName;
+ getDerived().getSema().StartOpenMPDSABlock(OMPD_metadirective, DirName,
+ nullptr, D->getBeginLoc());
+ StmtResult Res = getDerived().TransformOMPExecutableDirective(D);
+ getDerived().getSema().EndOpenMPDSABlock(Res.get());
+ return Res;
+}
+
+template <typename Derived>
+StmtResult
TreeTransform<Derived>::TransformOMPParallelDirective(OMPParallelDirective *D) {
DeclarationNameInfo DirName;
getDerived().getSema().StartOpenMPDSABlock(OMPD_parallel, DirName, nullptr,
Index: clang/lib/Sema/SemaOpenMP.cpp
===================================================================
--- clang/lib/Sema/SemaOpenMP.cpp
+++ clang/lib/Sema/SemaOpenMP.cpp
@@ -3751,6 +3751,7 @@
void Sema::ActOnOpenMPRegionStart(OpenMPDirectiveKind DKind, Scope *CurScope) {
switch (DKind) {
+ case OMPD_metadirective:
case OMPD_parallel:
case OMPD_parallel_for:
case OMPD_parallel_for_simd:
@@ -5436,6 +5437,7 @@
if (LangOpts.OpenMP >= 50)
AllowedNameModifiers.push_back(OMPD_simd);
break;
+ case OMPD_metadirective:
case OMPD_declare_target:
case OMPD_end_declare_target:
case OMPD_threadprivate:
@@ -5554,6 +5556,7 @@
case OMPC_atomic_default_mem_order:
case OMPC_device_type:
case OMPC_match:
+ case OMPC_when:
default:
llvm_unreachable("Unexpected clause");
}
@@ -11859,6 +11862,7 @@
case OMPC_exclusive:
case OMPC_uses_allocators:
case OMPC_affinity:
+ case OMPC_when:
default:
llvm_unreachable("Clause is not allowed.");
}
@@ -12012,6 +12016,7 @@
case OMPD_atomic:
case OMPD_teams_distribute:
case OMPD_requires:
+ case OMPD_metadirective:
llvm_unreachable("Unexpected OpenMP directive with if-clause");
case OMPD_unknown:
default:
@@ -12091,6 +12096,7 @@
case OMPD_teams_distribute:
case OMPD_teams_distribute_simd:
case OMPD_requires:
+ case OMPD_metadirective:
llvm_unreachable("Unexpected OpenMP directive with num_threads-clause");
case OMPD_unknown:
default:
@@ -12168,6 +12174,7 @@
case OMPD_atomic:
case OMPD_distribute_simd:
case OMPD_requires:
+ case OMPD_metadirective:
llvm_unreachable("Unexpected OpenMP directive with num_teams-clause");
case OMPD_unknown:
default:
@@ -12245,6 +12252,7 @@
case OMPD_atomic:
case OMPD_distribute_simd:
case OMPD_requires:
+ case OMPD_metadirective:
llvm_unreachable("Unexpected OpenMP directive with thread_limit-clause");
case OMPD_unknown:
default:
@@ -12322,6 +12330,7 @@
case OMPD_distribute_simd:
case OMPD_target_teams:
case OMPD_requires:
+ case OMPD_metadirective:
llvm_unreachable("Unexpected OpenMP directive with schedule clause");
case OMPD_unknown:
default:
@@ -12399,6 +12408,7 @@
case OMPD_atomic:
case OMPD_target_teams:
case OMPD_requires:
+ case OMPD_metadirective:
llvm_unreachable("Unexpected OpenMP directive with schedule clause");
case OMPD_unknown:
default:
@@ -12476,6 +12486,7 @@
case OMPD_atomic:
case OMPD_distribute_simd:
case OMPD_requires:
+ case OMPD_metadirective:
llvm_unreachable("Unexpected OpenMP directive with num_teams-clause");
case OMPD_unknown:
default:
@@ -12555,12 +12566,22 @@
case OMPD_atomic:
case OMPD_distribute_simd:
case OMPD_requires:
+ case OMPD_metadirective:
llvm_unreachable("Unexpected OpenMP directive with grainsize-clause");
case OMPD_unknown:
default:
llvm_unreachable("Unknown OpenMP directive");
}
break;
+ case OMPC_when:
+ if (DKind == OMPD_metadirective) {
+ CaptureRegion = OMPD_metadirective;
+ } else if (DKind == OMPD_unknown) {
+ llvm_unreachable("Unknown OpenMP directive");
+ } else {
+ llvm_unreachable("Unexpected OpenMP directive with when clause");
+ }
+ break;
case OMPC_firstprivate:
case OMPC_lastprivate:
case OMPC_reduction:
@@ -13068,6 +13089,7 @@
case OMPC_exclusive:
case OMPC_uses_allocators:
case OMPC_affinity:
+ case OMPC_when:
default:
llvm_unreachable("Clause is not allowed.");
}
@@ -13307,6 +13329,7 @@
case OMPC_exclusive:
case OMPC_uses_allocators:
case OMPC_affinity:
+ case OMPC_when:
default:
llvm_unreachable("Clause is not allowed.");
}
@@ -13545,6 +13568,7 @@
case OMPC_exclusive:
case OMPC_uses_allocators:
case OMPC_affinity:
+ case OMPC_when:
default:
llvm_unreachable("Clause is not allowed.");
}
@@ -13823,6 +13847,7 @@
case OMPC_destroy:
case OMPC_detach:
case OMPC_uses_allocators:
+ case OMPC_when:
default:
llvm_unreachable("Clause is not allowed.");
}
Index: clang/lib/Sema/SemaExceptionSpec.cpp
===================================================================
--- clang/lib/Sema/SemaExceptionSpec.cpp
+++ clang/lib/Sema/SemaExceptionSpec.cpp
@@ -1484,6 +1484,7 @@
case Stmt::OMPTeamsDistributeParallelForDirectiveClass:
case Stmt::OMPTeamsDistributeParallelForSimdDirectiveClass:
case Stmt::OMPTeamsDistributeSimdDirectiveClass:
+ case Stmt::OMPMetaDirectiveClass:
case Stmt::ReturnStmtClass:
case Stmt::SEHExceptStmtClass:
case Stmt::SEHFinallyStmtClass:
Index: clang/lib/Parse/ParseOpenMP.cpp
===================================================================
--- clang/lib/Parse/ParseOpenMP.cpp
+++ clang/lib/Parse/ParseOpenMP.cpp
@@ -2085,6 +2085,7 @@
case OMPD_target_teams_distribute_parallel_for:
case OMPD_target_teams_distribute_parallel_for_simd:
case OMPD_target_teams_distribute_simd:
+ case OMPD_metadirective:
Diag(Tok, diag::err_omp_unexpected_directive)
<< 1 << getOpenMPDirectiveName(DKind);
break;
@@ -2154,9 +2155,177 @@
// Name of critical directive.
DeclarationNameInfo DirName;
StmtResult Directive = StmtError();
+
bool HasAssociatedStatement = true;
switch (DKind) {
+ case OMPD_metadirective: {
+ ConsumeToken();
+ SmallVector<VariantMatchInfo, 4> VMIs;
+ OpenMPDirectiveKind DirKind = OMPD_unknown;
+ StmtResult AssociatedStmt;
+ TentativeParsingAction TPA(*this);
+ ASTContext &ASTContext = Actions.getASTContext();
+
+ BalancedDelimiterTracker T(*this, tok::l_paren,
+ tok::annot_pragma_openmp_end);
+ while (Tok.isNot(tok::annot_pragma_openmp_end)) {
+ OpenMPClauseKind CKind = Tok.isAnnotation()
+ ? OMPC_unknown
+ : getOpenMPClauseKind(PP.getSpelling(Tok));
+ SourceLocation Loc = ConsumeToken();
+
+ // Parse '('.
+ if (T.expectAndConsume(diag::err_expected_lparen_after,
+ getOpenMPClauseName(CKind).data()))
+ return Directive;
+
+ OMPTraitInfo &TI = Actions.getASTContext().getNewOMPTraitInfo();
+ if (CKind == OMPC_when) {
+ // parse and get condition expression to pass to the When clause
+ parseOMPContextSelectors(Loc, TI);
+
+ // Parse ':'
+ if (Tok.is(tok::colon))
+ ConsumeAnyToken();
+ else {
+ Diag(Tok, diag::warn_pragma_expected_colon) << "when clause";
+ return Directive;
+ }
+ }
+ // Skip Directive
+ int paren = 0;
+ while (Tok.isNot(tok::r_paren) || paren != 0) {
+ if (Tok.is(tok::l_paren))
+ paren++;
+ if (Tok.is(tok::r_paren))
+ paren--;
+ ConsumeAnyToken();
+ }
+ // Parse ')'
+ if (Tok.is(tok::r_paren))
+ T.consumeClose();
+
+ VariantMatchInfo VMI;
+ TI.getAsVariantMatchInfo(ASTContext, VMI);
+
+ VMIs.push_back(VMI);
+ }
+
+ TPA.Revert();
+ TargetOMPContext OMPCtx(ASTContext, nullptr, nullptr);
+ int BestIdx = getBestWhenMatchForContext(VMIs, OMPCtx);
+
+ int idx = 0;
+ while (Tok.isNot(tok::annot_pragma_openmp_end)) {
+ if (idx++ != BestIdx) { // Ignore this clause
+ ConsumeToken(); // Consume clause name
+ T.consumeOpen(); // Consume '('
+ int paren = 0;
+ // Skip everything inside the clause
+ while (Tok.isNot(tok::r_paren) || paren != 0) {
+ if (Tok.is(tok::l_paren))
+ paren++;
+ if (Tok.is(tok::r_paren))
+ paren--;
+ ConsumeAnyToken();
+ }
+ // Parse ')'
+ if (Tok.is(tok::r_paren))
+ T.consumeClose();
+ continue;
+ }
+
+ OpenMPClauseKind CKind = Tok.isAnnotation()
+ ? OMPC_unknown
+ : getOpenMPClauseKind(PP.getSpelling(Tok));
+ SourceLocation Loc = ConsumeToken();
+
+ // Parse '('.
+ T.consumeOpen();
+
+ // Skip ContextSelectors for when clause
+ if (CKind == OMPC_when) {
+ OMPTraitInfo &TI = Actions.getASTContext().getNewOMPTraitInfo();
+ // parse and skip ContextSelectors
+ parseOMPContextSelectors(Loc, TI);
+
+ // Parse ':'
+ ConsumeAnyToken();
+ }
+
+ // Parse Directive
+ if (Tok.isNot(tok::r_paren)) {
+ DirKind = parseOpenMPDirectiveKind(*this);
+ ConsumeToken();
+ if (DirKind != OMPD_unknown) {
+ ParseScope OMPDirectiveScope(this, ScopeFlags);
+ Actions.StartOpenMPDSABlock(DirKind, DirName, Actions.getCurScope(),
+ Loc);
+ int paren = 0;
+ while (Tok.isNot(tok::r_paren) || paren != 0) {
+ if (Tok.is(tok::l_paren))
+ paren++;
+ if (Tok.is(tok::r_paren))
+ paren--;
+ OpenMPClauseKind CKind =
+ Tok.isAnnotation() ? OMPC_unknown
+ : getOpenMPClauseKind(PP.getSpelling(Tok));
+ Actions.StartOpenMPClause(CKind);
+ OMPClause *DClause = ParseOpenMPClause(
+ DirKind, CKind, !FirstClauses[(unsigned)CKind].getInt());
+ FirstClauses[(unsigned)CKind].setInt(true);
+ if (DClause) {
+ FirstClauses[(unsigned)CKind].setPointer(DClause);
+ Clauses.push_back(DClause);
+ }
+ // Skip ',' if any.
+ if (Tok.is(tok::comma))
+ ConsumeToken();
+ Actions.EndOpenMPClause();
+ }
+ while (Tok.isNot(tok::annot_pragma_openmp_end))
+ ConsumeAnyToken();
+ // End location of the directive.
+ EndLoc = Tok.getLocation();
+ // Consume final annot_pragma_openmp_end.
+ ConsumeAnnotationToken();
+
+ Actions.ActOnOpenMPRegionStart(DirKind, getCurScope());
+ ParsingOpenMPDirectiveRAII NormalScope(*this, /*Value=*/false);
+ ParseScope InnerStmtScope(this, Scope::DeclScope,
+ getLangOpts().C99 ||
+ getLangOpts().CPlusPlus,
+ Tok.is(tok::l_brace));
+ StmtResult AStmt = ParseStatement();
+ InnerStmtScope.Exit();
+ AssociatedStmt = (Sema::CompoundScopeRAII(Actions), AStmt);
+ AssociatedStmt =
+ Actions.ActOnOpenMPRegionEnd(AssociatedStmt, Clauses);
+ Directive = Actions.ActOnOpenMPExecutableDirective(
+ DirKind, DirName, CancelRegion, Clauses, AssociatedStmt.get(),
+ Loc, EndLoc);
+ // Exit scope.
+ Actions.EndOpenMPDSABlock(Directive.get());
+ OMPDirectiveScope.Exit();
+ } else {
+ // Unknown Directive encountered. Skip directive.
+ Diag(Tok, diag::err_omp_unknown_directive);
+ while (Tok.isNot(tok::annot_pragma_openmp_end))
+ ConsumeAnyToken();
+ ConsumeAnnotationToken();
+ }
+ } else {
+ // Empty Directive. Skip in OpenMP 5.0.
+ // Generate nothing directive from OpenMP 5.1.
+ while (Tok.isNot(tok::annot_pragma_openmp_end))
+ ConsumeAnyToken();
+ ConsumeAnnotationToken();
+ }
+ break;
+ }
+ break;
+ }
case OMPD_threadprivate: {
// FIXME: Should this be permitted in C++?
if ((StmtCtx & ParsedStmtContext::AllowDeclarationsInC) ==
@@ -2759,6 +2928,7 @@
case OMPC_threadprivate:
case OMPC_uniform:
case OMPC_match:
+ case OMPC_when:
if (!WrongDirective)
Diag(Tok, diag::err_omp_unexpected_clause)
<< getOpenMPClauseName(CKind) << getOpenMPDirectiveName(DKind);
Index: clang/lib/CodeGen/CodeGenFunction.h
===================================================================
--- clang/lib/CodeGen/CodeGenFunction.h
+++ clang/lib/CodeGen/CodeGenFunction.h
@@ -3387,6 +3387,7 @@
const RegionCodeGenTy &BodyGen,
OMPTargetDataInfo &InputInfo);
+ void EmitOMPMetaDirective(const OMPMetaDirective &S);
void EmitOMPParallelDirective(const OMPParallelDirective &S);
void EmitOMPSimdDirective(const OMPSimdDirective &S);
void EmitOMPForDirective(const OMPForDirective &S);
Index: clang/lib/CodeGen/CGStmtOpenMP.cpp
===================================================================
--- clang/lib/CodeGen/CGStmtOpenMP.cpp
+++ clang/lib/CodeGen/CGStmtOpenMP.cpp
@@ -5399,6 +5399,7 @@
case OMPC_exclusive:
case OMPC_uses_allocators:
case OMPC_affinity:
+ case OMPC_when:
default:
llvm_unreachable("Clause is not allowed in 'omp atomic'.");
}
@@ -6587,6 +6588,10 @@
CGM.getOpenMPRuntime().emitMasterRegion(*this, CodeGen, S.getBeginLoc());
}
+void CodeGenFunction::EmitOMPMetaDirective(const OMPMetaDirective &S) {
+ EmitStmt(S.getIfStmt());
+}
+
void CodeGenFunction::EmitOMPParallelMasterTaskLoopDirective(
const OMPParallelMasterTaskLoopDirective &S) {
auto &&CodeGen = [this, &S](CodeGenFunction &CGF, PrePostActionTy &Action) {
Index: clang/lib/CodeGen/CGStmt.cpp
===================================================================
--- clang/lib/CodeGen/CGStmt.cpp
+++ clang/lib/CodeGen/CGStmt.cpp
@@ -194,6 +194,9 @@
case Stmt::SEHTryStmtClass:
EmitSEHTryStmt(cast<SEHTryStmt>(*S));
break;
+ case Stmt::OMPMetaDirectiveClass:
+ EmitOMPMetaDirective(cast<OMPMetaDirective>(*S));
+ break;
case Stmt::OMPParallelDirectiveClass:
EmitOMPParallelDirective(cast<OMPParallelDirective>(*S));
break;
Index: clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
===================================================================
--- clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
+++ clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
@@ -709,6 +709,7 @@
case OMPD_parallel_master_taskloop:
case OMPD_parallel_master_taskloop_simd:
case OMPD_requires:
+ case OMPD_metadirective:
case OMPD_unknown:
default:
llvm_unreachable("Unexpected directive.");
@@ -790,6 +791,7 @@
case OMPD_parallel_master_taskloop:
case OMPD_parallel_master_taskloop_simd:
case OMPD_requires:
+ case OMPD_metadirective:
case OMPD_unknown:
default:
break;
@@ -964,6 +966,7 @@
case OMPD_parallel_master_taskloop:
case OMPD_parallel_master_taskloop_simd:
case OMPD_requires:
+ case OMPD_metadirective:
case OMPD_unknown:
default:
llvm_unreachable("Unexpected directive.");
@@ -1051,6 +1054,7 @@
case OMPD_parallel_master_taskloop:
case OMPD_parallel_master_taskloop_simd:
case OMPD_requires:
+ case OMPD_metadirective:
case OMPD_unknown:
default:
break;
Index: clang/lib/CodeGen/CGOpenMPRuntime.cpp
===================================================================
--- clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -6983,6 +6983,7 @@
case OMPD_parallel_master_taskloop_simd:
case OMPD_requires:
case OMPD_unknown:
+ case OMPD_metadirective:
break;
default:
break;
@@ -9499,6 +9500,7 @@
case OMPD_parallel_master_taskloop:
case OMPD_parallel_master_taskloop_simd:
case OMPD_requires:
+ case OMPD_metadirective:
case OMPD_unknown:
default:
llvm_unreachable("Unexpected directive.");
@@ -10202,6 +10204,12 @@
if (!S)
return;
+ if (isa<OMPMetaDirective>(S)) {
+ const auto &M = *cast<OMPMetaDirective>(S);
+ scanForTargetRegionsFunctions(M.getIfStmt(), ParentName);
+ return;
+ }
+
// Codegen OMP target directives that offload compute to the device.
bool RequiresDeviceCodegen =
isa<OMPExecutableDirective>(S) &&
@@ -11010,6 +11018,7 @@
case OMPD_target_parallel_for:
case OMPD_target_parallel_for_simd:
case OMPD_requires:
+ case OMPD_metadirective:
case OMPD_unknown:
default:
llvm_unreachable("Unexpected standalone target data directive.");
Index: clang/lib/Basic/OpenMPKinds.cpp
===================================================================
--- clang/lib/Basic/OpenMPKinds.cpp
+++ clang/lib/Basic/OpenMPKinds.cpp
@@ -180,6 +180,7 @@
case OMPC_exclusive:
case OMPC_uses_allocators:
case OMPC_affinity:
+ case OMPC_when:
break;
default:
break;
@@ -420,6 +421,7 @@
case OMPC_exclusive:
case OMPC_uses_allocators:
case OMPC_affinity:
+ case OMPC_when:
break;
default:
break;
@@ -578,6 +580,9 @@
OpenMPDirectiveKind DKind) {
assert(unsigned(DKind) < llvm::omp::Directive_enumSize);
switch (DKind) {
+ case OMPD_metadirective:
+ CaptureRegions.push_back(OMPD_metadirective);
+ break;
case OMPD_parallel:
case OMPD_parallel_for:
case OMPD_parallel_for_simd:
Index: clang/lib/AST/StmtProfile.cpp
===================================================================
--- clang/lib/AST/StmtProfile.cpp
+++ clang/lib/AST/StmtProfile.cpp
@@ -847,6 +847,10 @@
P.Visit(*I);
}
+void StmtProfiler::VisitOMPMetaDirective(const OMPMetaDirective *S) {
+ VisitOMPExecutableDirective(S);
+}
+
void StmtProfiler::VisitOMPLoopDirective(const OMPLoopDirective *S) {
VisitOMPExecutableDirective(S);
}
Index: clang/lib/AST/StmtPrinter.cpp
===================================================================
--- clang/lib/AST/StmtPrinter.cpp
+++ clang/lib/AST/StmtPrinter.cpp
@@ -650,6 +650,11 @@
PrintStmt(S->getRawStmt());
}
+void StmtPrinter::VisitOMPMetaDirective(OMPMetaDirective *Node) {
+ Indent() << "#pragma omp metadirective";
+ PrintOMPExecutableDirective(Node);
+}
+
void StmtPrinter::VisitOMPParallelDirective(OMPParallelDirective *Node) {
Indent() << "#pragma omp parallel";
PrintOMPExecutableDirective(Node);
Index: clang/lib/AST/StmtOpenMP.cpp
===================================================================
--- clang/lib/AST/StmtOpenMP.cpp
+++ clang/lib/AST/StmtOpenMP.cpp
@@ -191,6 +191,25 @@
llvm::copy(A, getFinalsConditions().begin());
}
+OMPMetaDirective *OMPMetaDirective::Create(const ASTContext &C,
+ SourceLocation StartLoc,
+ SourceLocation EndLoc,
+ ArrayRef<OMPClause *> Clauses,
+ Stmt *AssociatedStmt, Stmt *IfStmt) {
+ auto *Dir = createDirective<OMPMetaDirective>(
+ C, Clauses, AssociatedStmt, /*NumChildren=*/1, StartLoc, EndLoc);
+ Dir->setIfStmt(IfStmt);
+ return Dir;
+}
+
+OMPMetaDirective *OMPMetaDirective::CreateEmpty(const ASTContext &C,
+ unsigned NumClauses,
+ EmptyShell) {
+ return createEmptyDirective<OMPMetaDirective>(C, NumClauses,
+ /*HasAssociatedStmt=*/true,
+ /*NumChildren=*/1);
+}
+
OMPParallelDirective *OMPParallelDirective::Create(
const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc,
ArrayRef<OMPClause *> Clauses, Stmt *AssociatedStmt, Expr *TaskRedRef,
Index: clang/lib/AST/OpenMPClause.cpp
===================================================================
--- clang/lib/AST/OpenMPClause.cpp
+++ clang/lib/AST/OpenMPClause.cpp
@@ -156,6 +156,7 @@
case OMPC_exclusive:
case OMPC_uses_allocators:
case OMPC_affinity:
+ case OMPC_when:
break;
default:
break;
@@ -250,6 +251,7 @@
case OMPC_exclusive:
case OMPC_uses_allocators:
case OMPC_affinity:
+ case OMPC_when:
break;
default:
break;
Index: clang/include/clang/Serialization/ASTBitCodes.h
===================================================================
--- clang/include/clang/Serialization/ASTBitCodes.h
+++ clang/include/clang/Serialization/ASTBitCodes.h
@@ -1830,21 +1830,21 @@
/// A CXXBoolLiteralExpr record.
EXPR_CXX_BOOL_LITERAL,
- EXPR_CXX_NULL_PTR_LITERAL, // CXXNullPtrLiteralExpr
- EXPR_CXX_TYPEID_EXPR, // CXXTypeidExpr (of expr).
- EXPR_CXX_TYPEID_TYPE, // CXXTypeidExpr (of type).
- EXPR_CXX_THIS, // CXXThisExpr
- EXPR_CXX_THROW, // CXXThrowExpr
- EXPR_CXX_DEFAULT_ARG, // CXXDefaultArgExpr
- EXPR_CXX_DEFAULT_INIT, // CXXDefaultInitExpr
- EXPR_CXX_BIND_TEMPORARY, // CXXBindTemporaryExpr
+ EXPR_CXX_NULL_PTR_LITERAL, // CXXNullPtrLiteralExpr
+ EXPR_CXX_TYPEID_EXPR, // CXXTypeidExpr (of expr).
+ EXPR_CXX_TYPEID_TYPE, // CXXTypeidExpr (of type).
+ EXPR_CXX_THIS, // CXXThisExpr
+ EXPR_CXX_THROW, // CXXThrowExpr
+ EXPR_CXX_DEFAULT_ARG, // CXXDefaultArgExpr
+ EXPR_CXX_DEFAULT_INIT, // CXXDefaultInitExpr
+ EXPR_CXX_BIND_TEMPORARY, // CXXBindTemporaryExpr
EXPR_CXX_SCALAR_VALUE_INIT, // CXXScalarValueInitExpr
EXPR_CXX_NEW, // CXXNewExpr
EXPR_CXX_DELETE, // CXXDeleteExpr
EXPR_CXX_PSEUDO_DESTRUCTOR, // CXXPseudoDestructorExpr
- EXPR_EXPR_WITH_CLEANUPS, // ExprWithCleanups
+ EXPR_EXPR_WITH_CLEANUPS, // ExprWithCleanups
EXPR_CXX_DEPENDENT_SCOPE_MEMBER, // CXXDependentScopeMemberExpr
EXPR_CXX_DEPENDENT_SCOPE_DECL_REF, // DependentScopeDeclRefExpr
@@ -1852,41 +1852,42 @@
EXPR_CXX_UNRESOLVED_MEMBER, // UnresolvedMemberExpr
EXPR_CXX_UNRESOLVED_LOOKUP, // UnresolvedLookupExpr
- EXPR_CXX_EXPRESSION_TRAIT, // ExpressionTraitExpr
- EXPR_CXX_NOEXCEPT, // CXXNoexceptExpr
+ EXPR_CXX_EXPRESSION_TRAIT, // ExpressionTraitExpr
+ EXPR_CXX_NOEXCEPT, // CXXNoexceptExpr
- EXPR_OPAQUE_VALUE, // OpaqueValueExpr
- EXPR_BINARY_CONDITIONAL_OPERATOR, // BinaryConditionalOperator
- EXPR_TYPE_TRAIT, // TypeTraitExpr
- EXPR_ARRAY_TYPE_TRAIT, // ArrayTypeTraitIntExpr
+ EXPR_OPAQUE_VALUE, // OpaqueValueExpr
+ EXPR_BINARY_CONDITIONAL_OPERATOR, // BinaryConditionalOperator
+ EXPR_TYPE_TRAIT, // TypeTraitExpr
+ EXPR_ARRAY_TYPE_TRAIT, // ArrayTypeTraitIntExpr
- EXPR_PACK_EXPANSION, // PackExpansionExpr
- EXPR_SIZEOF_PACK, // SizeOfPackExpr
- EXPR_SUBST_NON_TYPE_TEMPLATE_PARM, // SubstNonTypeTemplateParmExpr
- EXPR_SUBST_NON_TYPE_TEMPLATE_PARM_PACK,// SubstNonTypeTemplateParmPackExpr
- EXPR_FUNCTION_PARM_PACK, // FunctionParmPackExpr
- EXPR_MATERIALIZE_TEMPORARY, // MaterializeTemporaryExpr
- EXPR_CXX_FOLD, // CXXFoldExpr
- EXPR_CONCEPT_SPECIALIZATION,// ConceptSpecializationExpr
- EXPR_REQUIRES, // RequiresExpr
+ EXPR_PACK_EXPANSION, // PackExpansionExpr
+ EXPR_SIZEOF_PACK, // SizeOfPackExpr
+ EXPR_SUBST_NON_TYPE_TEMPLATE_PARM, // SubstNonTypeTemplateParmExpr
+ EXPR_SUBST_NON_TYPE_TEMPLATE_PARM_PACK, // SubstNonTypeTemplateParmPackExpr
+ EXPR_FUNCTION_PARM_PACK, // FunctionParmPackExpr
+ EXPR_MATERIALIZE_TEMPORARY, // MaterializeTemporaryExpr
+ EXPR_CXX_FOLD, // CXXFoldExpr
+ EXPR_CONCEPT_SPECIALIZATION, // ConceptSpecializationExpr
+ EXPR_REQUIRES, // RequiresExpr
// CUDA
- EXPR_CUDA_KERNEL_CALL, // CUDAKernelCallExpr
+ EXPR_CUDA_KERNEL_CALL, // CUDAKernelCallExpr
// OpenCL
- EXPR_ASTYPE, // AsTypeExpr
+ EXPR_ASTYPE, // AsTypeExpr
// Microsoft
- EXPR_CXX_PROPERTY_REF_EXPR, // MSPropertyRefExpr
+ EXPR_CXX_PROPERTY_REF_EXPR, // MSPropertyRefExpr
EXPR_CXX_PROPERTY_SUBSCRIPT_EXPR, // MSPropertySubscriptExpr
- EXPR_CXX_UUIDOF_EXPR, // CXXUuidofExpr (of expr).
- EXPR_CXX_UUIDOF_TYPE, // CXXUuidofExpr (of type).
- STMT_SEH_LEAVE, // SEHLeaveStmt
- STMT_SEH_EXCEPT, // SEHExceptStmt
- STMT_SEH_FINALLY, // SEHFinallyStmt
- STMT_SEH_TRY, // SEHTryStmt
+ EXPR_CXX_UUIDOF_EXPR, // CXXUuidofExpr (of expr).
+ EXPR_CXX_UUIDOF_TYPE, // CXXUuidofExpr (of type).
+ STMT_SEH_LEAVE, // SEHLeaveStmt
+ STMT_SEH_EXCEPT, // SEHExceptStmt
+ STMT_SEH_FINALLY, // SEHFinallyStmt
+ STMT_SEH_TRY, // SEHTryStmt
// OpenMP directives
+ STMT_OMP_META_DIRECTIVE,
STMT_OMP_PARALLEL_DIRECTIVE,
STMT_OMP_SIMD_DIRECTIVE,
STMT_OMP_FOR_DIRECTIVE,
@@ -1946,10 +1947,10 @@
EXPR_OMP_ITERATOR,
// ARC
- EXPR_OBJC_BRIDGED_CAST, // ObjCBridgedCastExpr
+ EXPR_OBJC_BRIDGED_CAST, // ObjCBridgedCastExpr
- STMT_MS_DEPENDENT_EXISTS, // MSDependentExistsStmt
- EXPR_LAMBDA, // LambdaExpr
+ STMT_MS_DEPENDENT_EXISTS, // MSDependentExistsStmt
+ EXPR_LAMBDA, // LambdaExpr
STMT_COROUTINE_BODY,
STMT_CORETURN,
EXPR_COAWAIT,
Index: clang/include/clang/Sema/Sema.h
===================================================================
--- clang/include/clang/Sema/Sema.h
+++ clang/include/clang/Sema/Sema.h
@@ -10300,6 +10300,11 @@
void ActOnOpenMPLoopInitialization(SourceLocation ForLoc, Stmt *Init);
// OpenMP directives and clauses.
+ /// Called on well-formed '\#pragma omp metadirective' after parsing
+ /// of the associated statement.
+ StmtResult ActOnOpenMPMetaDirective(ArrayRef<OMPClause *> Clauses,
+ Stmt *AStmt, SourceLocation StartLoc,
+ SourceLocation EndLoc);
/// Called on correct id-expression from the '#pragma omp
/// threadprivate'.
ExprResult ActOnOpenMPIdExpression(Scope *CurScope, CXXScopeSpec &ScopeSpec,
@@ -10795,6 +10800,10 @@
SourceLocation StartLoc,
SourceLocation LParenLoc,
SourceLocation EndLoc);
+ /// Called on well-formed 'when' clause.
+ OMPClause *ActOnOpenMPWhenClause(OMPTraitInfo &TI, SourceLocation StartLoc,
+ SourceLocation LParenLoc,
+ SourceLocation EndLoc);
/// Called on well-formed 'default' clause.
OMPClause *ActOnOpenMPDefaultClause(llvm::omp::DefaultKind Kind,
SourceLocation KindLoc,
Index: clang/include/clang/Basic/StmtNodes.td
===================================================================
--- clang/include/clang/Basic/StmtNodes.td
+++ clang/include/clang/Basic/StmtNodes.td
@@ -217,6 +217,7 @@
// OpenMP Directives.
def OMPExecutableDirective : StmtNode<Stmt, 1>;
+def OMPMetaDirective : StmtNode<OMPExecutableDirective>;
def OMPLoopDirective : StmtNode<OMPExecutableDirective, 1>;
def OMPParallelDirective : StmtNode<OMPExecutableDirective>;
def OMPSimdDirective : StmtNode<OMPLoopDirective>;
Index: clang/include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -10488,6 +10488,9 @@
: Note<"jump bypasses OpenMP structured block">;
def note_omp_exits_structured_block
: Note<"jump exits scope of OpenMP structured block">;
+def err_omp_misplaced_default_clause : Error<
+ "misplaced default clause! Only one default clause is allowed in "
+ "metadirective in the end">;
} // end of OpenMP category
let CategoryName = "Related Result Type Issue" in {
Index: clang/include/clang/AST/StmtOpenMP.h
===================================================================
--- clang/include/clang/AST/StmtOpenMP.h
+++ clang/include/clang/AST/StmtOpenMP.h
@@ -362,6 +362,44 @@
}
};
+/// This represents '#pragma omp metadirective' directive.
+///
+/// \code
+/// #pragma omp metadirective when(user={condition(N>10)}: parallel for)
+/// \endcode
+/// In this example directive '#pragma omp metadirective' has clauses 'when'
+/// with a dynamic user condition to check if a variable 'N > 10'
+///
+class OMPMetaDirective final : public OMPExecutableDirective {
+ friend class ASTStmtReader;
+ friend class OMPExecutableDirective;
+ Stmt *IfStmt;
+
+ OMPMetaDirective(SourceLocation StartLoc, SourceLocation EndLoc)
+ : OMPExecutableDirective(OMPMetaDirectiveClass,
+ llvm::omp::OMPD_metadirective, StartLoc,
+ EndLoc) {}
+ explicit OMPMetaDirective()
+ : OMPExecutableDirective(OMPMetaDirectiveClass,
+ llvm::omp::OMPD_metadirective, SourceLocation(),
+ SourceLocation()) {}
+
+public:
+ static OMPMetaDirective *Create(const ASTContext &C, SourceLocation StartLoc,
+ SourceLocation EndLoc,
+ ArrayRef<OMPClause *> Clauses,
+ Stmt *AssociatedStmt, Stmt *IfStmt);
+ static OMPMetaDirective *CreateEmpty(const ASTContext &C, unsigned NumClauses,
+ EmptyShell);
+
+ void setIfStmt(Stmt *stmt) { IfStmt = stmt; }
+ Stmt *getIfStmt() const { return IfStmt; }
+
+ static bool classof(const Stmt *T) {
+ return T->getStmtClass() == OMPMetaDirectiveClass;
+ }
+};
+
/// This represents '#pragma omp parallel' directive.
///
/// \code
Index: clang/include/clang/AST/RecursiveASTVisitor.h
===================================================================
--- clang/include/clang/AST/RecursiveASTVisitor.h
+++ clang/include/clang/AST/RecursiveASTVisitor.h
@@ -2780,6 +2780,9 @@
return TraverseOMPExecutableDirective(S);
}
+DEF_TRAVERSE_STMT(OMPMetaDirective,
+ { TRY_TO(TraverseOMPExecutableDirective(S)); })
+
DEF_TRAVERSE_STMT(OMPParallelDirective,
{ TRY_TO(TraverseOMPExecutableDirective(S)); })
Index: clang/include/clang-c/Index.h
===================================================================
--- clang/include/clang-c/Index.h
+++ clang/include/clang-c/Index.h
@@ -2568,7 +2568,11 @@
*/
CXCursor_OMPScanDirective = 287,
- CXCursor_LastStmt = CXCursor_OMPScanDirective,
+ /** OpenMP metadirective directive.
+ */
+ CXCursor_OMPMetaDirective = 288,
+
+ CXCursor_LastStmt = CXCursor_OMPMetaDirective,
/**
* Cursor that represents the translation unit itself.
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
