[clang] [OpenMP] Add runtime selection for metadirective with non-constant conditions. (PR #192455)
Zahira Ammarguellat via cfe-commits
cfe-commits at lists.llvm.org
Mon Apr 27 07:11:55 PDT 2026
https://github.com/zahiraam updated https://github.com/llvm/llvm-project/pull/192455
>From 34c79fc24c47f7b7c5f8769b48f4d8dc5c007177 Mon Sep 17 00:00:00 2001
From: Ammarguellat <zahira.ammarguellat at intel.com>
Date: Thu, 16 Apr 2026 06:51:31 -0700
Subject: [PATCH 1/4] [OpenMP] Add runtime selection for metadirective with
non-constant user conditions
---
clang/include/clang/AST/StmtOpenMP.h | 24 +++
clang/lib/CodeGen/CGStmtOpenMP.cpp | 148 ++++++++++++++++++
clang/lib/Parse/ParseOpenMP.cpp | 106 +++++++++++++
clang/lib/Sema/SemaOpenMP.cpp | 11 ++
clang/lib/Sema/SemaStmt.cpp | 13 +-
.../metadirective_user_condition_codegen.cpp | 134 ++++++++++++++++
6 files changed, 431 insertions(+), 5 deletions(-)
create mode 100644 clang/test/OpenMP/metadirective_user_condition_codegen.cpp
diff --git a/clang/include/clang/AST/StmtOpenMP.h b/clang/include/clang/AST/StmtOpenMP.h
index dbc76e7df8ecd..435a060f42281 100644
--- a/clang/include/clang/AST/StmtOpenMP.h
+++ b/clang/include/clang/AST/StmtOpenMP.h
@@ -6377,6 +6377,14 @@ class OMPMetaDirective final : public OMPExecutableDirective {
friend class OMPExecutableDirective;
Stmt *IfStmt;
+ // Runtime selection support for non-constant user conditions.
+ // These fields store metadata for generating runtime if-else chains when
+ // user conditions cannot be evaluated at compile time.
+ bool HasNonConstantConditions = false;
+ unsigned NumConditions = 0;
+ Expr **Conditions = nullptr;
+ OpenMPDirectiveKind *DirectiveVariants = nullptr;
+
OMPMetaDirective(SourceLocation StartLoc, SourceLocation EndLoc)
: OMPExecutableDirective(OMPMetaDirectiveClass,
llvm::omp::OMPD_metadirective, StartLoc,
@@ -6397,6 +6405,22 @@ class OMPMetaDirective final : public OMPExecutableDirective {
EmptyShell);
Stmt *getIfStmt() const { return IfStmt; }
+ void setNonConstantUserConditions(ArrayRef<Expr *> Conds,
+ ArrayRef<OpenMPDirectiveKind> Variants) {
+ HasNonConstantConditions = true;
+ NumConditions = Conds.size();
+ Conditions = const_cast<Expr **>(Conds.data());
+ DirectiveVariants = const_cast<OpenMPDirectiveKind *>(Variants.data());
+ }
+
+ bool hasNonConstantConditions() const { return HasNonConstantConditions; }
+ ArrayRef<Expr *> getConditions() const {
+ return llvm::ArrayRef(Conditions, NumConditions);
+ }
+ ArrayRef<OpenMPDirectiveKind> getDirectiveVariants() const {
+ return llvm::ArrayRef(DirectiveVariants, NumConditions);
+ }
+
static bool classof(const Stmt *T) {
return T->getStmtClass() == OMPMetaDirectiveClass;
}
diff --git a/clang/lib/CodeGen/CGStmtOpenMP.cpp b/clang/lib/CodeGen/CGStmtOpenMP.cpp
index 59d0e6825a975..604a1762d05cb 100644
--- a/clang/lib/CodeGen/CGStmtOpenMP.cpp
+++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp
@@ -2158,7 +2158,155 @@ void CodeGenFunction::EmitOMPParallelDirective(const OMPParallelDirective &S) {
checkForLastprivateConditionalUpdate(*this, S);
}
+/// Emit an OpenMP directive with the given kind and body statement.
+/// Used for runtime metadirective selection.
+static void EmitOMPDirectiveWithBody(CodeGenFunction &CGF,
+ OpenMPDirectiveKind DKind,
+ const Stmt *Body,
+ SourceLocation Loc) {
+ auto &RT = CGF.CGM.getOpenMPRuntime();
+
+ switch (DKind) {
+ case OMPD_parallel: {
+ // Outline the body into a separate function and call __kmpc_fork_call.
+ auto &OMPBuilder = RT.getOMPBuilder();
+ llvm::Module &M = CGF.CGM.getModule();
+ llvm::LLVMContext &Ctx = M.getContext();
+ llvm::Type *Int32PtrTy = llvm::PointerType::getUnqual(Ctx);
+ llvm::FunctionType *FnTy = llvm::FunctionType::get(
+ llvm::Type::getVoidTy(Ctx), {Int32PtrTy, Int32PtrTy}, false);
+ llvm::Function *OutlinedFn = llvm::Function::Create(
+ FnTy, llvm::GlobalValue::InternalLinkage, ".omp_outlined.", &M);
+ OutlinedFn->addFnAttr(llvm::Attribute::NoUnwind);
+ QualType Int32PtrQTy =
+ CGF.getContext().getPointerType(CGF.getContext().IntTy);
+ FunctionArgList Args;
+ ImplicitParamDecl GtidArg(CGF.getContext(), Int32PtrQTy,
+ ImplicitParamKind::Other);
+ ImplicitParamDecl BtidArg(CGF.getContext(), Int32PtrQTy,
+ ImplicitParamKind::Other);
+ Args.push_back(&GtidArg);
+ Args.push_back(&BtidArg);
+
+ // Create CGFunctionInfo for the outlined function.
+ const CGFunctionInfo &FnInfo =
+ CGF.CGM.getTypes().arrangeBuiltinFunctionDeclaration(
+ CGF.getContext().VoidTy, Args);
+
+ // Generate code in the outlined function.
+ CodeGenFunction OutlinedCGF(CGF.CGM);
+ OutlinedCGF.StartFunction(GlobalDecl(), CGF.getContext().VoidTy, OutlinedFn,
+ FnInfo, Args, Loc, Loc);
+
+ // Emit the body in the outlined function.
+ OutlinedCGF.EmitStmt(Body);
+ OutlinedCGF.FinishFunction();
+
+ // Emit call to __kmpc_fork_call(loc, 0, outlined_fn).
+ llvm::Value *RTLoc = RT.emitUpdateLocation(CGF, Loc);
+ llvm::Value *ForkCallArgs[] = {
+ RTLoc,
+ CGF.Builder.getInt32(0),
+ CGF.Builder.CreateBitCast(OutlinedFn,
+ llvm::PointerType::getUnqual(Ctx))};
+
+ llvm::FunctionCallee ForkCallFn = OMPBuilder.getOrCreateRuntimeFunction(
+ M, llvm::omp::OMPRTL___kmpc_fork_call);
+ CGF.EmitRuntimeCall(ForkCallFn, ForkCallArgs);
+ break;
+ }
+ case OMPD_single: {
+ auto &&CodeGen = [Body](CodeGenFunction &CGF, PrePostActionTy &Action) {
+ Action.Enter(CGF);
+ CGF.EmitStmt(Body);
+ };
+
+ // Use emitSingleRegion which handles __kmpc_single/__kmpc_end_single.
+ RT.emitSingleRegion(CGF, CodeGen, Loc,
+ /*CopyprivateVars=*/{}, /*DestExprs=*/{},
+ /*SrcExprs=*/{}, /*AssignmentOps=*/{});
+ RT.emitBarrierCall(CGF, Loc, OMPD_single);
+ break;
+ }
+ default:
+ // For unsupported directives, just emit the body.
+ CGF.EmitStmt(Body);
+ break;
+ }
+}
+
void CodeGenFunction::EmitOMPMetaDirective(const OMPMetaDirective &S) {
+ if (S.hasNonConstantConditions()) {
+ // Build runtime if-else chain.
+ ArrayRef<Expr *> Conditions = S.getConditions();
+ ArrayRef<OpenMPDirectiveKind> DirectiveVariants = S.getDirectiveVariants();
+
+ assert(Conditions.size() == DirectiveVariants.size() &&
+ "Mismatch between conditions and directive variants");
+
+ // Build if-else chain from the conditions. Process in reverse order: skip
+ // 'otherwise' clauses and emit conditions from last to first.
+ llvm::BasicBlock *ContBB = createBasicBlock("omp.meta.end");
+
+ // Find the index of the 'otherwise' clause (condition == nullptr).
+ int otherwiseIdx = -1;
+ for (unsigned i = 0; i < Conditions.size(); ++i) {
+ if (!Conditions[i]) {
+ otherwiseIdx = i;
+ break;
+ }
+ }
+
+ // Build if-else chain.
+ for (int i = Conditions.size() - 1; i >= 0; --i) {
+ Expr *Cond = Conditions[i];
+ OpenMPDirectiveKind DKind = DirectiveVariants[i];
+
+ if (!Cond) {
+ // This is the 'otherwise' clause; will be emitted in final else.
+ continue;
+ }
+
+ llvm::BasicBlock *ThenBB = createBasicBlock("omp.meta.then");
+ llvm::BasicBlock *ElseBB = createBasicBlock("omp.meta.else");
+ EmitBranchOnBoolExpr(Cond, ThenBB, ElseBB, 0);
+ EmitBlock(ThenBB);
+ {
+ // Save LocalDeclMap to allow reusing the same AST in multiple branches.
+ llvm::DenseMap<const Decl *, Address> SavedLocalDeclMap = LocalDeclMap;
+
+ LexicalScope Scope(*this, S.getSourceRange());
+ // Emit the OpenMP directive with the body.
+ EmitOMPDirectiveWithBody(*this, DKind, S.getAssociatedStmt(),
+ S.getBeginLoc());
+
+ // Restore LocalDeclMap for next branch.
+ LocalDeclMap = SavedLocalDeclMap;
+ }
+ EmitBranch(ContBB);
+
+ // Continue with else block.
+ EmitBlock(ElseBB);
+ }
+ // Emit the 'otherwise' clause in the final else block.
+ if (otherwiseIdx >= 0) {
+ // Save LocalDeclMap to allow reusing the same AST.
+ llvm::DenseMap<const Decl *, Address> SavedLocalDeclMap = LocalDeclMap;
+
+ LexicalScope Scope(*this, S.getSourceRange());
+ OpenMPDirectiveKind DKind = DirectiveVariants[otherwiseIdx];
+ // Emit the OpenMP directive with the body.
+ EmitOMPDirectiveWithBody(*this, DKind, S.getAssociatedStmt(),
+ S.getBeginLoc());
+
+ // Restore LocalDeclMap.
+ LocalDeclMap = SavedLocalDeclMap;
+ }
+ EmitBranch(ContBB);
+ EmitBlock(ContBB);
+ return;
+ }
+ // Compile-time selection.
EmitStmt(S.getIfStmt());
}
diff --git a/clang/lib/Parse/ParseOpenMP.cpp b/clang/lib/Parse/ParseOpenMP.cpp
index 0e92c3fa1b572..f73fe79237d8f 100644
--- a/clang/lib/Parse/ParseOpenMP.cpp
+++ b/clang/lib/Parse/ParseOpenMP.cpp
@@ -2583,6 +2583,8 @@ StmtResult Parser::ParseOpenMPDeclarativeOrExecutableDirective(
case OMPD_metadirective: {
ConsumeToken();
SmallVector<VariantMatchInfo, 4> VMIs;
+ SmallVector<OMPTraitInfo *, 4> TraitInfos;
+ SmallVector<OpenMPClauseKind, 4> ClauseKinds;
// First iteration of parsing all clauses of metadirective.
// This iteration only parses and collects all context selector ignoring the
@@ -2662,11 +2664,115 @@ StmtResult Parser::ParseOpenMPDeclarativeOrExecutableDirective(
TI.getAsVariantMatchInfo(ASTContext, VMI);
VMIs.push_back(VMI);
+ TraitInfos.push_back(&TI);
+ ClauseKinds.push_back(CKind);
}
TPA.Revert();
// End of the first iteration. Parser is reset to the start of metadirective
+ // Check if we have non-constant user conditions.
+ bool HasNonConstantUserCondition = false;
+ for (unsigned i = 0; i < TraitInfos.size(); ++i) {
+ if (ClauseKinds[i] == OMPC_when) {
+ OMPTraitInfo *TI = TraitInfos[i];
+ TI->anyScoreOrCondition([&](Expr *&E, bool IsScore) {
+ if (!IsScore && E && !E->isIntegerConstantExpr(ASTContext)) {
+ HasNonConstantUserCondition = true;
+ return true;
+ }
+ return false;
+ });
+ if (HasNonConstantUserCondition)
+ break;
+ }
+ }
+ if (HasNonConstantUserCondition) {
+ // Extract directive kinds and conditions.
+ SmallVector<Expr *, 4> Conditions;
+ SmallVector<OpenMPDirectiveKind, 4> DirectiveKinds;
+
+ while (Tok.isNot(tok::annot_pragma_openmp_end)) {
+ OpenMPClauseKind CKind = Tok.isAnnotation()
+ ? OMPC_unknown
+ : getOpenMPClauseKind(PP.getSpelling(Tok));
+ SourceLocation ClauseLoc = ConsumeToken();
+ T.consumeOpen();
+
+ Expr *Condition = nullptr;
+ if (CKind == OMPC_when) {
+ OMPTraitInfo &TI = Actions.getASTContext().getNewOMPTraitInfo();
+ parseOMPContextSelectors(ClauseLoc, TI);
+
+ // Extract the user condition expression.
+ TI.anyScoreOrCondition([&](Expr *&E, bool IsScore) {
+ if (!IsScore && E) {
+ Condition = E;
+ return true;
+ }
+ return false;
+ });
+
+ ConsumeAnyToken(); // consume ':'.
+ }
+
+ OpenMPDirectiveKind DKind = OMPD_unknown;
+ if (!Tok.is(tok::r_paren)) {
+ DKind = parseOpenMPDirectiveKind(*this);
+ }
+
+ Conditions.push_back(Condition);
+ DirectiveKinds.push_back(DKind);
+ // Skip to closing paren.
+ 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--;
+ if (Tok.is(tok::annot_pragma_openmp_end))
+ break;
+ ConsumeAnyToken();
+ }
+ if (Tok.is(tok::r_paren))
+ T.consumeClose();
+ }
+ // Consume end token and parse body once.
+ assert(Tok.is(tok::annot_pragma_openmp_end));
+ SourceLocation EndLoc = ConsumeAnnotationToken();
+
+ StmtResult AssociatedStmt = ParseStatement();
+ if (AssociatedStmt.isInvalid())
+ return StmtError();
+
+ // Allocate arrays in ASTContext.
+ Expr **ConditionsArray = nullptr;
+ OpenMPDirectiveKind *DirectiveKindsArray = nullptr;
+ if (!Conditions.empty()) {
+ ConditionsArray = new (ASTContext) Expr *[Conditions.size()];
+ std::copy(Conditions.begin(), Conditions.end(), ConditionsArray);
+
+ DirectiveKindsArray =
+ new (ASTContext) OpenMPDirectiveKind[DirectiveKinds.size()];
+ std::copy(DirectiveKinds.begin(), DirectiveKinds.end(),
+ DirectiveKindsArray);
+ }
+ // Create OMPMetaDirective with runtime selection metadata.
+ SmallVector<OMPClause *, 1> Clauses;
+ Directive = Actions.OpenMP().ActOnOpenMPMetaDirective(
+ Clauses, AssociatedStmt.get(), Loc, EndLoc);
+
+ if (Directive.isUsable() && !Conditions.empty()) {
+ if (auto *MD = dyn_cast_or_null<OMPMetaDirective>(Directive.get())) {
+ ArrayRef<Expr *> CondsRef(ConditionsArray, Conditions.size());
+ ArrayRef<OpenMPDirectiveKind> KindsRef(DirectiveKindsArray,
+ DirectiveKinds.size());
+ MD->setNonConstantUserConditions(CondsRef, KindsRef);
+ }
+ }
+
+ return Directive;
+ }
std::function<void(StringRef)> DiagUnknownTrait =
[this, Loc](StringRef ISATrait) {
// TODO Track the selector locations in a way that is accessible here
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index ec9b2e68ef6cf..6b890f3e76f2a 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -7953,6 +7953,17 @@ setBranchProtectedScope(Sema &SemaRef, OpenMPDirectiveKind DKind, Stmt *AStmt) {
return CS;
}
+StmtResult SemaOpenMP::ActOnOpenMPMetaDirective(ArrayRef<OMPClause *> Clauses,
+ Stmt *AStmt,
+ SourceLocation StartLoc,
+ SourceLocation EndLoc) {
+ if (!AStmt)
+ return StmtError();
+
+ return OMPMetaDirective::Create(getASTContext(), StartLoc, EndLoc, Clauses,
+ AStmt, /*IfStmt=*/nullptr);
+}
+
StmtResult
SemaOpenMP::ActOnOpenMPParallelDirective(ArrayRef<OMPClause *> Clauses,
Stmt *AStmt, SourceLocation StartLoc,
diff --git a/clang/lib/Sema/SemaStmt.cpp b/clang/lib/Sema/SemaStmt.cpp
index 531147ef35b08..60d4214bb3c8a 100644
--- a/clang/lib/Sema/SemaStmt.cpp
+++ b/clang/lib/Sema/SemaStmt.cpp
@@ -4700,11 +4700,14 @@ buildCapturedStmtCaptureList(Sema &S, CapturedRegionScopeInfo *RSI,
S.OpenMP().setOpenMPCaptureKind(Field, Cap.getVariable(),
RSI->OpenMPLevel);
- Captures.push_back(CapturedStmt::Capture(
- Cap.getLocation(),
- Cap.isReferenceCapture() ? CapturedStmt::VCK_ByRef
- : CapturedStmt::VCK_ByCopy,
- cast<VarDecl>(Cap.getVariable())));
+ ValueDecl *CapVar = Cap.getVariable();
+ if (auto *BD = dyn_cast<BindingDecl>(CapVar))
+ CapVar = cast<VarDecl>(BD->getDecomposedDecl());
+ Captures.push_back(CapturedStmt::Capture(Cap.getLocation(),
+ Cap.isReferenceCapture()
+ ? CapturedStmt::VCK_ByRef
+ : CapturedStmt::VCK_ByCopy,
+ cast<VarDecl>(CapVar)));
}
CaptureInits.push_back(Init.get());
}
diff --git a/clang/test/OpenMP/metadirective_user_condition_codegen.cpp b/clang/test/OpenMP/metadirective_user_condition_codegen.cpp
new file mode 100644
index 0000000000000..f2a087f28d1f3
--- /dev/null
+++ b/clang/test/OpenMP/metadirective_user_condition_codegen.cpp
@@ -0,0 +1,134 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=52 \
+// RUN: -triple x86_64-unknown-linux -emit-llvm %s -o - | FileCheck %s
+
+// expected-no-diagnostics
+
+#ifndef HEADER
+#define HEADER
+
+void bar();
+
+
+// CHECK-LABEL: define dso_local void @_Z27test_runtime_user_conditionb(
+// CHECK-SAME: i1 noundef zeroext [[USE_GPU:%.*]]) #[[ATTR0:[0-9]+]] {
+// CHECK-NEXT: [[ENTRY:.*:]]
+// CHECK-NEXT: [[USE_GPU_ADDR:%.*]] = alloca i8, align 1
+// CHECK-NEXT: [[STOREDV:%.*]] = zext i1 [[USE_GPU]] to i8
+// CHECK-NEXT: store i8 [[STOREDV]], ptr [[USE_GPU_ADDR]], align 1
+// CHECK-NEXT: [[TMP0:%.*]] = load i8, ptr [[USE_GPU_ADDR]], align 1
+// CHECK-NEXT: [[LOADEDV:%.*]] = trunc i8 [[TMP0]] to i1
+// CHECK-NEXT: br i1 [[LOADEDV]], label %[[OMP_META_THEN:.*]], label %[[OMP_META_ELSE:.*]]
+// CHECK: [[OMP_META_THEN]]:
+// CHECK-NEXT: call void @_Z3barv()
+// CHECK-NEXT: br label %[[OMP_META_END:.*]]
+// CHECK: [[OMP_META_ELSE]]:
+// CHECK-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1:[0-9]+]], i32 0, ptr @.omp_outlined.)
+// CHECK-NEXT: br label %[[OMP_META_END]]
+// CHECK: [[OMP_META_END]]:
+// CHECK-NEXT: ret void
+//
+void test_runtime_user_condition(bool use_gpu) {
+ #pragma omp metadirective \
+ when(user={condition(use_gpu)}: target) \
+ otherwise(parallel)
+ {
+ bar();
+ }
+}
+
+
+// CHECK-LABEL: define dso_local void @_Z35test_runtime_user_condition_integeri(
+// CHECK-SAME: i32 noundef [[THRESHOLD:%.*]]) #[[ATTR0]] {
+// CHECK-NEXT: [[ENTRY:.*:]]
+// CHECK-NEXT: [[THRESHOLD_ADDR:%.*]] = alloca i32, align 4
+// CHECK-NEXT: [[I:%.*]] = alloca i32, align 4
+// CHECK-NEXT: store i32 [[THRESHOLD]], ptr [[THRESHOLD_ADDR]], align 4
+// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr [[THRESHOLD_ADDR]], align 4
+// CHECK-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP0]], 100
+// CHECK-NEXT: br i1 [[CMP]], label %[[OMP_META_THEN:.*]], label %[[OMP_META_ELSE:.*]]
+// CHECK: [[OMP_META_THEN]]:
+// CHECK-NEXT: store i32 0, ptr [[I]], align 4
+// CHECK-NEXT: br label %[[FOR_COND:.*]]
+// CHECK: [[FOR_COND]]:
+// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[I]], align 4
+// CHECK-NEXT: [[CMP1:%.*]] = icmp slt i32 [[TMP1]], 10
+// CHECK-NEXT: br i1 [[CMP1]], label %[[FOR_BODY:.*]], label %[[FOR_END:.*]]
+// CHECK: [[FOR_BODY]]:
+// CHECK-NEXT: call void @_Z3barv()
+// CHECK-NEXT: br label %[[FOR_INC:.*]]
+// CHECK: [[FOR_INC]]:
+// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[I]], align 4
+// CHECK-NEXT: [[INC:%.*]] = add nsw i32 [[TMP2]], 1
+// CHECK-NEXT: store i32 [[INC]], ptr [[I]], align 4
+// CHECK-NEXT: br label %[[FOR_COND]], !llvm.loop [[LOOP4:![0-9]+]]
+// CHECK: [[FOR_END]]:
+// CHECK-NEXT: br label %[[OMP_META_END:.*]]
+// CHECK: [[OMP_META_ELSE]]:
+// CHECK-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 0, ptr @.omp_outlined..1)
+// CHECK-NEXT: br label %[[OMP_META_END]]
+// CHECK: [[OMP_META_END]]:
+// CHECK-NEXT: ret void
+//
+void test_runtime_user_condition_integer(int threshold) {
+
+
+
+ #pragma omp metadirective \
+ when(user={condition(threshold > 100)}: parallel for) \
+ otherwise(parallel)
+ {
+ for (int i = 0; i < 10; ++i) {
+ bar();
+ }
+ }
+}
+
+// CHECK-LABEL: define dso_local void @_Z33test_constant_user_condition_truev(
+// CHECK-SAME: ) #[[ATTR0]] {
+// CHECK-NEXT: [[ENTRY:.*:]]
+// CHECK-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 0, ptr @_Z33test_constant_user_condition_truev.omp_outlined)
+// CHECK-NEXT: ret void
+//
+void test_constant_user_condition_true() {
+ // For constant conditions, no runtime branch should be generated
+
+ #pragma omp metadirective \
+ when(user={condition(true)}: parallel) \
+ otherwise(single)
+ {
+ bar();
+ }
+}
+
+// CHECK-LABEL: define dso_local void @_Z34test_constant_user_condition_falsev(
+// CHECK-SAME: ) #[[ATTR0]] {
+// CHECK-NEXT: [[ENTRY:.*:]]
+// CHECK-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])
+// CHECK-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_single(ptr @[[GLOB1]], i32 [[TMP0]])
+// CHECK-NEXT: [[TMP2:%.*]] = icmp ne i32 [[TMP1]], 0
+// CHECK-NEXT: br i1 [[TMP2]], label %[[OMP_IF_THEN:.*]], label %[[OMP_IF_END:.*]]
+// CHECK: [[OMP_IF_THEN]]:
+// CHECK-NEXT: call void @_Z3barv()
+// CHECK-NEXT: call void @__kmpc_end_single(ptr @[[GLOB1]], i32 [[TMP0]])
+// CHECK-NEXT: br label %[[OMP_IF_END]]
+// CHECK: [[OMP_IF_END]]:
+// CHECK-NEXT: call void @__kmpc_barrier(ptr @[[GLOB2:[0-9]+]], i32 [[TMP0]])
+// CHECK-NEXT: ret void
+//
+void test_constant_user_condition_false() {
+ // For constant conditions, no runtime branch should be generated
+
+ #pragma omp metadirective \
+ when(user={condition(false)}: parallel) \
+ otherwise(single)
+ {
+ bar();
+ }
+}
+
+#endif
+//.
+// CHECK: [[LOOP4]] = distinct !{[[LOOP4]], [[META5:![0-9]+]]}
+// CHECK: [[META5]] = !{!"llvm.loop.mustprogress"}
+//.
>From ca332f555f9e943854cabc27156d9d5d5610374a Mon Sep 17 00:00:00 2001
From: Ammarguellat <zahira.ammarguellat at intel.com>
Date: Thu, 16 Apr 2026 06:56:13 -0700
Subject: [PATCH 2/4] Fix format
---
clang/include/clang/AST/StmtOpenMP.h | 2 +-
clang/lib/CodeGen/CGStmtOpenMP.cpp | 6 ++----
2 files changed, 3 insertions(+), 5 deletions(-)
diff --git a/clang/include/clang/AST/StmtOpenMP.h b/clang/include/clang/AST/StmtOpenMP.h
index 435a060f42281..398c484e26029 100644
--- a/clang/include/clang/AST/StmtOpenMP.h
+++ b/clang/include/clang/AST/StmtOpenMP.h
@@ -6406,7 +6406,7 @@ class OMPMetaDirective final : public OMPExecutableDirective {
Stmt *getIfStmt() const { return IfStmt; }
void setNonConstantUserConditions(ArrayRef<Expr *> Conds,
- ArrayRef<OpenMPDirectiveKind> Variants) {
+ ArrayRef<OpenMPDirectiveKind> Variants) {
HasNonConstantConditions = true;
NumConditions = Conds.size();
Conditions = const_cast<Expr **>(Conds.data());
diff --git a/clang/lib/CodeGen/CGStmtOpenMP.cpp b/clang/lib/CodeGen/CGStmtOpenMP.cpp
index 604a1762d05cb..0363728230218 100644
--- a/clang/lib/CodeGen/CGStmtOpenMP.cpp
+++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp
@@ -2162,8 +2162,7 @@ void CodeGenFunction::EmitOMPParallelDirective(const OMPParallelDirective &S) {
/// Used for runtime metadirective selection.
static void EmitOMPDirectiveWithBody(CodeGenFunction &CGF,
OpenMPDirectiveKind DKind,
- const Stmt *Body,
- SourceLocation Loc) {
+ const Stmt *Body, SourceLocation Loc) {
auto &RT = CGF.CGM.getOpenMPRuntime();
switch (DKind) {
@@ -2205,8 +2204,7 @@ static void EmitOMPDirectiveWithBody(CodeGenFunction &CGF,
// Emit call to __kmpc_fork_call(loc, 0, outlined_fn).
llvm::Value *RTLoc = RT.emitUpdateLocation(CGF, Loc);
llvm::Value *ForkCallArgs[] = {
- RTLoc,
- CGF.Builder.getInt32(0),
+ RTLoc, CGF.Builder.getInt32(0),
CGF.Builder.CreateBitCast(OutlinedFn,
llvm::PointerType::getUnqual(Ctx))};
>From 64618efaecc83685c29243d8c794a65f67b033b4 Mon Sep 17 00:00:00 2001
From: Ammarguellat <zahira.ammarguellat at intel.com>
Date: Thu, 23 Apr 2026 10:48:37 -0700
Subject: [PATCH 3/4] Pass 1
---
clang/include/clang/AST/StmtOpenMP.h | 69 +++--
clang/include/clang/Sema/SemaOpenMP.h | 4 +-
clang/lib/AST/StmtOpenMP.cpp | 31 +-
clang/lib/AST/StmtPrinter.cpp | 28 +-
clang/lib/CodeGen/CGStmtOpenMP.cpp | 143 ++-------
clang/lib/Parse/ParseOpenMP.cpp | 289 ++++++++----------
clang/lib/Sema/SemaOpenMP.cpp | 16 +-
clang/lib/Serialization/ASTReaderStmt.cpp | 11 +-
clang/lib/Serialization/ASTWriterStmt.cpp | 1 +
clang/test/OpenMP/metadirective_ast_print.c | 74 ++++-
clang/test/OpenMP/metadirective_ast_print.cpp | 14 +
11 files changed, 333 insertions(+), 347 deletions(-)
create mode 100644 clang/test/OpenMP/metadirective_ast_print.cpp
diff --git a/clang/include/clang/AST/StmtOpenMP.h b/clang/include/clang/AST/StmtOpenMP.h
index 398c484e26029..bc8e63ae13cf2 100644
--- a/clang/include/clang/AST/StmtOpenMP.h
+++ b/clang/include/clang/AST/StmtOpenMP.h
@@ -6375,50 +6375,49 @@ class OMPMaskedDirective final : public OMPExecutableDirective {
class OMPMetaDirective final : public OMPExecutableDirective {
friend class ASTStmtReader;
friend class OMPExecutableDirective;
- Stmt *IfStmt;
- // Runtime selection support for non-constant user conditions.
- // These fields store metadata for generating runtime if-else chains when
- // user conditions cannot be evaluated at compile time.
- bool HasNonConstantConditions = false;
- unsigned NumConditions = 0;
- Expr **Conditions = nullptr;
- OpenMPDirectiveKind *DirectiveVariants = nullptr;
+ // Children layout: [IfStmt, Cond0, Dir0, Cond1, Dir1, ...]
+ // IfStmt: compile-time resolved directive (may be null for runtime-only)
+ // CondI: Expr* condition for variant I (null = 'otherwise')
+ // DirI: OMPExecutableDirective* for variant I
+ unsigned NumVariants = 0;
- OMPMetaDirective(SourceLocation StartLoc, SourceLocation EndLoc)
+ OMPMetaDirective(SourceLocation StartLoc, SourceLocation EndLoc,
+ unsigned NumVariants)
: OMPExecutableDirective(OMPMetaDirectiveClass,
- llvm::omp::OMPD_metadirective, StartLoc,
- EndLoc) {}
- explicit OMPMetaDirective()
- : OMPExecutableDirective(OMPMetaDirectiveClass,
- llvm::omp::OMPD_metadirective, SourceLocation(),
- SourceLocation()) {}
+ llvm::omp::OMPD_metadirective,
+ StartLoc, EndLoc),
+ NumVariants(NumVariants) {}
- void setIfStmt(Stmt *S) { IfStmt = S; }
+ explicit OMPMetaDirective(unsigned NumVariants)
+ : OMPExecutableDirective(OMPMetaDirectiveClass,
+ llvm::omp::OMPD_metadirective,
+ SourceLocation(), SourceLocation()),
+ NumVariants(NumVariants) {}
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);
- Stmt *getIfStmt() const { return IfStmt; }
+ static OMPMetaDirective *
+ Create(const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc,
+ ArrayRef<OMPClause *> Clauses, Stmt *AssociatedStmt, Stmt *IfStmt,
+ ArrayRef<Expr *> Conditions, ArrayRef<Stmt *> Directives);
- void setNonConstantUserConditions(ArrayRef<Expr *> Conds,
- ArrayRef<OpenMPDirectiveKind> Variants) {
- HasNonConstantConditions = true;
- NumConditions = Conds.size();
- Conditions = const_cast<Expr **>(Conds.data());
- DirectiveVariants = const_cast<OpenMPDirectiveKind *>(Variants.data());
- }
+ static OMPMetaDirective *CreateEmpty(const ASTContext &C,
+ unsigned NumClauses,
+ unsigned NumVariants, EmptyShell);
- bool hasNonConstantConditions() const { return HasNonConstantConditions; }
- ArrayRef<Expr *> getConditions() const {
- return llvm::ArrayRef(Conditions, NumConditions);
+ Stmt *getIfStmt() const { return Data->getChildren()[0]; }
+ void setIfStmt(Stmt *S) { Data->getChildren()[0] = S; }
+
+ unsigned getNumVariants() const { return NumVariants; }
+
+ Expr *getVariantCondition(unsigned I) const {
+ assert(I < NumVariants);
+ return cast_or_null<Expr>(Data->getChildren()[1 + 2 * I]);
}
- ArrayRef<OpenMPDirectiveKind> getDirectiveVariants() const {
- return llvm::ArrayRef(DirectiveVariants, NumConditions);
+
+ Stmt *getVariantDirective(unsigned I) const {
+ assert(I < NumVariants);
+ return Data->getChildren()[2 + 2 * I];
}
static bool classof(const Stmt *T) {
diff --git a/clang/include/clang/Sema/SemaOpenMP.h b/clang/include/clang/Sema/SemaOpenMP.h
index 3621ce96b8724..2394a81c5e66c 100644
--- a/clang/include/clang/Sema/SemaOpenMP.h
+++ b/clang/include/clang/Sema/SemaOpenMP.h
@@ -214,7 +214,9 @@ class SemaOpenMP : public SemaBase {
/// of the associated statement.
StmtResult ActOnOpenMPMetaDirective(ArrayRef<OMPClause *> Clauses,
Stmt *AStmt, SourceLocation StartLoc,
- SourceLocation EndLoc);
+ SourceLocation EndLoc,
+ ArrayRef<Expr *> Conditions,
+ ArrayRef<Stmt *> Directives);
// OpenMP directives and clauses.
/// Called on correct id-expression from the '#pragma omp
diff --git a/clang/lib/AST/StmtOpenMP.cpp b/clang/lib/AST/StmtOpenMP.cpp
index 9d6b315effb41..314f8251aba04 100644
--- a/clang/lib/AST/StmtOpenMP.cpp
+++ b/clang/lib/AST/StmtOpenMP.cpp
@@ -260,23 +260,34 @@ void OMPLoopDirective::setFinalsConditions(ArrayRef<Expr *> A) {
llvm::copy(A, getFinalsConditions().begin());
}
-OMPMetaDirective *OMPMetaDirective::Create(const ASTContext &C,
- SourceLocation StartLoc,
- SourceLocation EndLoc,
- ArrayRef<OMPClause *> Clauses,
- Stmt *AssociatedStmt, Stmt *IfStmt) {
+OMPMetaDirective *
+OMPMetaDirective::Create(const ASTContext &C, SourceLocation StartLoc,
+ SourceLocation EndLoc, ArrayRef<OMPClause *> Clauses,
+ Stmt *AssociatedStmt, Stmt *IfStmt,
+ ArrayRef<Expr *> Conditions,
+ ArrayRef<Stmt *> Directives) {
+ assert(Conditions.size() == Directives.size());
+ unsigned NumVariants = Conditions.size();
auto *Dir = createDirective<OMPMetaDirective>(
- C, Clauses, AssociatedStmt, /*NumChildren=*/1, StartLoc, EndLoc);
- Dir->setIfStmt(IfStmt);
+ C, Clauses, AssociatedStmt,
+ /*NumChildren=*/1 + 2 * NumVariants,
+ StartLoc, EndLoc, NumVariants);
+ Dir->Data->getChildren()[0] = IfStmt;
+ for (unsigned I = 0; I < NumVariants; ++I) {
+ Dir->Data->getChildren()[1 + 2 * I] = Conditions[I];
+ Dir->Data->getChildren()[2 + 2 * I] = Directives[I];
+ }
return Dir;
}
OMPMetaDirective *OMPMetaDirective::CreateEmpty(const ASTContext &C,
unsigned NumClauses,
+ unsigned NumVariants,
EmptyShell) {
- return createEmptyDirective<OMPMetaDirective>(C, NumClauses,
- /*HasAssociatedStmt=*/true,
- /*NumChildren=*/1);
+ return createEmptyDirective<OMPMetaDirective>(
+ C, NumClauses,
+ /*HasAssociatedStmt=*/false,
+ /*NumChildren=*/1 + 2 * NumVariants, NumVariants);
}
OMPParallelDirective *OMPParallelDirective::Create(
diff --git a/clang/lib/AST/StmtPrinter.cpp b/clang/lib/AST/StmtPrinter.cpp
index e0b930ba0a21a..52842612b9244 100644
--- a/clang/lib/AST/StmtPrinter.cpp
+++ b/clang/lib/AST/StmtPrinter.cpp
@@ -762,7 +762,33 @@ void StmtPrinter::PrintOMPExecutableDirective(OMPExecutableDirective *S,
void StmtPrinter::VisitOMPMetaDirective(OMPMetaDirective *Node) {
Indent() << "#pragma omp metadirective";
- PrintOMPExecutableDirective(Node);
+ if (Node->getNumVariants() > 0) {
+ for (unsigned I = 0; I < Node->getNumVariants(); ++I) {
+ Expr *Cond = Node->getVariantCondition(I);
+ Stmt *Dir = Node->getVariantDirective(I);
+
+ if (Cond) {
+ OS << " when(";
+ // Print the condition expression.
+ Cond->printPretty(OS, nullptr, Policy);
+ OS << ":";
+ } else {
+ OS << " otherwise(";
+ }
+
+ // Print the sub-directive.
+ if (Dir && !isa<NullStmt>(Dir)) {
+ OS << " ";
+ PrintStmt(Dir);
+ }
+ OS << ")";
+ }
+ OS << NL;
+ } else {
+ OS << NL;
+ if (Stmt *If = Node->getIfStmt())
+ PrintStmt(If);
+ }
}
void StmtPrinter::VisitOMPParallelDirective(OMPParallelDirective *Node) {
diff --git a/clang/lib/CodeGen/CGStmtOpenMP.cpp b/clang/lib/CodeGen/CGStmtOpenMP.cpp
index 0363728230218..1d5a79fac7cda 100644
--- a/clang/lib/CodeGen/CGStmtOpenMP.cpp
+++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp
@@ -2158,153 +2158,48 @@ void CodeGenFunction::EmitOMPParallelDirective(const OMPParallelDirective &S) {
checkForLastprivateConditionalUpdate(*this, S);
}
-/// Emit an OpenMP directive with the given kind and body statement.
-/// Used for runtime metadirective selection.
-static void EmitOMPDirectiveWithBody(CodeGenFunction &CGF,
- OpenMPDirectiveKind DKind,
- const Stmt *Body, SourceLocation Loc) {
- auto &RT = CGF.CGM.getOpenMPRuntime();
-
- switch (DKind) {
- case OMPD_parallel: {
- // Outline the body into a separate function and call __kmpc_fork_call.
- auto &OMPBuilder = RT.getOMPBuilder();
- llvm::Module &M = CGF.CGM.getModule();
- llvm::LLVMContext &Ctx = M.getContext();
- llvm::Type *Int32PtrTy = llvm::PointerType::getUnqual(Ctx);
- llvm::FunctionType *FnTy = llvm::FunctionType::get(
- llvm::Type::getVoidTy(Ctx), {Int32PtrTy, Int32PtrTy}, false);
- llvm::Function *OutlinedFn = llvm::Function::Create(
- FnTy, llvm::GlobalValue::InternalLinkage, ".omp_outlined.", &M);
- OutlinedFn->addFnAttr(llvm::Attribute::NoUnwind);
- QualType Int32PtrQTy =
- CGF.getContext().getPointerType(CGF.getContext().IntTy);
- FunctionArgList Args;
- ImplicitParamDecl GtidArg(CGF.getContext(), Int32PtrQTy,
- ImplicitParamKind::Other);
- ImplicitParamDecl BtidArg(CGF.getContext(), Int32PtrQTy,
- ImplicitParamKind::Other);
- Args.push_back(&GtidArg);
- Args.push_back(&BtidArg);
-
- // Create CGFunctionInfo for the outlined function.
- const CGFunctionInfo &FnInfo =
- CGF.CGM.getTypes().arrangeBuiltinFunctionDeclaration(
- CGF.getContext().VoidTy, Args);
-
- // Generate code in the outlined function.
- CodeGenFunction OutlinedCGF(CGF.CGM);
- OutlinedCGF.StartFunction(GlobalDecl(), CGF.getContext().VoidTy, OutlinedFn,
- FnInfo, Args, Loc, Loc);
-
- // Emit the body in the outlined function.
- OutlinedCGF.EmitStmt(Body);
- OutlinedCGF.FinishFunction();
-
- // Emit call to __kmpc_fork_call(loc, 0, outlined_fn).
- llvm::Value *RTLoc = RT.emitUpdateLocation(CGF, Loc);
- llvm::Value *ForkCallArgs[] = {
- RTLoc, CGF.Builder.getInt32(0),
- CGF.Builder.CreateBitCast(OutlinedFn,
- llvm::PointerType::getUnqual(Ctx))};
-
- llvm::FunctionCallee ForkCallFn = OMPBuilder.getOrCreateRuntimeFunction(
- M, llvm::omp::OMPRTL___kmpc_fork_call);
- CGF.EmitRuntimeCall(ForkCallFn, ForkCallArgs);
- break;
- }
- case OMPD_single: {
- auto &&CodeGen = [Body](CodeGenFunction &CGF, PrePostActionTy &Action) {
- Action.Enter(CGF);
- CGF.EmitStmt(Body);
- };
-
- // Use emitSingleRegion which handles __kmpc_single/__kmpc_end_single.
- RT.emitSingleRegion(CGF, CodeGen, Loc,
- /*CopyprivateVars=*/{}, /*DestExprs=*/{},
- /*SrcExprs=*/{}, /*AssignmentOps=*/{});
- RT.emitBarrierCall(CGF, Loc, OMPD_single);
- break;
- }
- default:
- // For unsupported directives, just emit the body.
- CGF.EmitStmt(Body);
- break;
- }
-}
-
void CodeGenFunction::EmitOMPMetaDirective(const OMPMetaDirective &S) {
- if (S.hasNonConstantConditions()) {
- // Build runtime if-else chain.
- ArrayRef<Expr *> Conditions = S.getConditions();
- ArrayRef<OpenMPDirectiveKind> DirectiveVariants = S.getDirectiveVariants();
-
- assert(Conditions.size() == DirectiveVariants.size() &&
- "Mismatch between conditions and directive variants");
-
- // Build if-else chain from the conditions. Process in reverse order: skip
- // 'otherwise' clauses and emit conditions from last to first.
+#if 0
+ // Runtime path: emit an if-else chain over the variant conditions.
+ if (S.getNumVariants() > 0) {
llvm::BasicBlock *ContBB = createBasicBlock("omp.meta.end");
- // Find the index of the 'otherwise' clause (condition == nullptr).
- int otherwiseIdx = -1;
- for (unsigned i = 0; i < Conditions.size(); ++i) {
- if (!Conditions[i]) {
- otherwiseIdx = i;
+ // Find the 'otherwise' clause (null condition).
+ int OtherwiseIdx = -1;
+ for (unsigned I = 0; I < S.getNumVariants(); ++I) {
+ if (!S.getVariantCondition(I)) {
+ OtherwiseIdx = I;
break;
}
}
- // Build if-else chain.
- for (int i = Conditions.size() - 1; i >= 0; --i) {
- Expr *Cond = Conditions[i];
- OpenMPDirectiveKind DKind = DirectiveVariants[i];
-
- if (!Cond) {
- // This is the 'otherwise' clause; will be emitted in final else.
+ // Emit each 'when' clause as a conditional branch.
+ for (unsigned I = 0; I < S.getNumVariants(); ++I) {
+ Expr *Cond = S.getVariantCondition(I);
+ if (!Cond)
continue;
- }
llvm::BasicBlock *ThenBB = createBasicBlock("omp.meta.then");
llvm::BasicBlock *ElseBB = createBasicBlock("omp.meta.else");
- EmitBranchOnBoolExpr(Cond, ThenBB, ElseBB, 0);
- EmitBlock(ThenBB);
- {
- // Save LocalDeclMap to allow reusing the same AST in multiple branches.
- llvm::DenseMap<const Decl *, Address> SavedLocalDeclMap = LocalDeclMap;
- LexicalScope Scope(*this, S.getSourceRange());
- // Emit the OpenMP directive with the body.
- EmitOMPDirectiveWithBody(*this, DKind, S.getAssociatedStmt(),
- S.getBeginLoc());
+ EmitBranchOnBoolExpr(Cond, ThenBB, ElseBB, 0);
- // Restore LocalDeclMap for next branch.
- LocalDeclMap = SavedLocalDeclMap;
- }
+ EmitBlock(ThenBB);
+ EmitStmt(S.getVariantDirective(I));
EmitBranch(ContBB);
- // Continue with else block.
EmitBlock(ElseBB);
}
- // Emit the 'otherwise' clause in the final else block.
- if (otherwiseIdx >= 0) {
- // Save LocalDeclMap to allow reusing the same AST.
- llvm::DenseMap<const Decl *, Address> SavedLocalDeclMap = LocalDeclMap;
- LexicalScope Scope(*this, S.getSourceRange());
- OpenMPDirectiveKind DKind = DirectiveVariants[otherwiseIdx];
- // Emit the OpenMP directive with the body.
- EmitOMPDirectiveWithBody(*this, DKind, S.getAssociatedStmt(),
- S.getBeginLoc());
+ // Emit the 'otherwise' clause in the final else block.
+ if (OtherwiseIdx >= 0)
+ EmitStmt(S.getVariantDirective(OtherwiseIdx));
- // Restore LocalDeclMap.
- LocalDeclMap = SavedLocalDeclMap;
- }
EmitBranch(ContBB);
EmitBlock(ContBB);
return;
}
- // Compile-time selection.
+#endif
EmitStmt(S.getIfStmt());
}
diff --git a/clang/lib/Parse/ParseOpenMP.cpp b/clang/lib/Parse/ParseOpenMP.cpp
index f73fe79237d8f..f9f114bc02731 100644
--- a/clang/lib/Parse/ParseOpenMP.cpp
+++ b/clang/lib/Parse/ParseOpenMP.cpp
@@ -2671,188 +2671,167 @@ StmtResult Parser::ParseOpenMPDeclarativeOrExecutableDirective(
TPA.Revert();
// End of the first iteration. Parser is reset to the start of metadirective
- // Check if we have non-constant user conditions.
+ // Detect non-constant user conditions.
bool HasNonConstantUserCondition = false;
- for (unsigned i = 0; i < TraitInfos.size(); ++i) {
- if (ClauseKinds[i] == OMPC_when) {
- OMPTraitInfo *TI = TraitInfos[i];
- TI->anyScoreOrCondition([&](Expr *&E, bool IsScore) {
- if (!IsScore && E && !E->isIntegerConstantExpr(ASTContext)) {
- HasNonConstantUserCondition = true;
- return true;
- }
- return false;
- });
- if (HasNonConstantUserCondition)
- break;
- }
+ for (unsigned I = 0; I < TraitInfos.size(); ++I) {
+ if (ClauseKinds[I] != OMPC_when)
+ continue;
+ TraitInfos[I]->anyScoreOrCondition([&](Expr *&E, bool IsScore) {
+ if (!IsScore && E && !E->isIntegerConstantExpr(ASTContext)) {
+ HasNonConstantUserCondition = true;
+ return true;
+ }
+ return false;
+ });
+ if (HasNonConstantUserCondition)
+ break;
}
- if (HasNonConstantUserCondition) {
- // Extract directive kinds and conditions.
- SmallVector<Expr *, 4> Conditions;
- SmallVector<OpenMPDirectiveKind, 4> DirectiveKinds;
- while (Tok.isNot(tok::annot_pragma_openmp_end)) {
- OpenMPClauseKind CKind = Tok.isAnnotation()
- ? OMPC_unknown
- : getOpenMPClauseKind(PP.getSpelling(Tok));
- SourceLocation ClauseLoc = ConsumeToken();
- T.consumeOpen();
-
- Expr *Condition = nullptr;
- if (CKind == OMPC_when) {
- OMPTraitInfo &TI = Actions.getASTContext().getNewOMPTraitInfo();
- parseOMPContextSelectors(ClauseLoc, TI);
-
- // Extract the user condition expression.
- TI.anyScoreOrCondition([&](Expr *&E, bool IsScore) {
- if (!IsScore && E) {
- Condition = E;
- return true;
- }
- return false;
- });
-
- ConsumeAnyToken(); // consume ':'.
- }
+ if (!HasNonConstantUserCondition) {
+ std::function<void(StringRef)> DiagUnknownTrait =
+ [this, Loc](StringRef ISATrait) {
+ Diag(Loc, diag::warn_unknown_declare_variant_isa_trait) << ISATrait;
+ };
+ TargetOMPContext OMPCtx(ASTContext, std::move(DiagUnknownTrait),
+ /*CurrentFunctionDecl=*/nullptr,
+ ArrayRef<llvm::omp::TraitProperty>(),
+ Actions.OpenMP().getOpenMPDeviceNum());
- OpenMPDirectiveKind DKind = OMPD_unknown;
- if (!Tok.is(tok::r_paren)) {
- DKind = parseOpenMPDirectiveKind(*this);
- }
+ int BestIdx = getBestVariantMatchForContext(VMIs, OMPCtx);
- Conditions.push_back(Condition);
- DirectiveKinds.push_back(DKind);
- // Skip to closing paren.
- 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--;
- if (Tok.is(tok::annot_pragma_openmp_end))
- break;
- ConsumeAnyToken();
- }
- if (Tok.is(tok::r_paren))
- T.consumeClose();
+ if (BestIdx == -1) {
+ SkipUntil(tok::annot_pragma_openmp_end);
+ if (Tok.is(tok::annot_pragma_openmp_end))
+ ConsumeAnnotationToken();
+ return StmtEmpty();
}
- // Consume end token and parse body once.
- assert(Tok.is(tok::annot_pragma_openmp_end));
- SourceLocation EndLoc = ConsumeAnnotationToken();
- StmtResult AssociatedStmt = ParseStatement();
- if (AssociatedStmt.isInvalid())
- return StmtError();
+ BalancedDelimiterTracker T2(*this, tok::l_paren,
+ tok::annot_pragma_openmp_end);
+ int Idx = 0;
+ Stmt *IfStmt = nullptr;
- // Allocate arrays in ASTContext.
- Expr **ConditionsArray = nullptr;
- OpenMPDirectiveKind *DirectiveKindsArray = nullptr;
- if (!Conditions.empty()) {
- ConditionsArray = new (ASTContext) Expr *[Conditions.size()];
- std::copy(Conditions.begin(), Conditions.end(), ConditionsArray);
-
- DirectiveKindsArray =
- new (ASTContext) OpenMPDirectiveKind[DirectiveKinds.size()];
- std::copy(DirectiveKinds.begin(), DirectiveKinds.end(),
- DirectiveKindsArray);
- }
- // Create OMPMetaDirective with runtime selection metadata.
- SmallVector<OMPClause *, 1> Clauses;
- Directive = Actions.OpenMP().ActOnOpenMPMetaDirective(
- Clauses, AssociatedStmt.get(), Loc, EndLoc);
-
- if (Directive.isUsable() && !Conditions.empty()) {
- if (auto *MD = dyn_cast_or_null<OMPMetaDirective>(Directive.get())) {
- ArrayRef<Expr *> CondsRef(ConditionsArray, Conditions.size());
- ArrayRef<OpenMPDirectiveKind> KindsRef(DirectiveKindsArray,
- DirectiveKinds.size());
- MD->setNonConstantUserConditions(CondsRef, KindsRef);
+ while (Tok.isNot(tok::annot_pragma_openmp_end)) {
+ if (Tok.isAnnotation())
+ break;
+ OpenMPClauseKind CKind = getOpenMPClauseKind(PP.getSpelling(Tok));
+ if (CKind == OMPC_unknown)
+ break;
+ SourceLocation ClauseLoc = ConsumeToken();
+ T2.consumeOpen();
+
+ if (Idx != BestIdx) {
+ 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--;
+ if (Tok.is(tok::annot_pragma_openmp_end))
+ break;
+ ConsumeAnyToken();
+ }
+ if (Tok.is(tok::r_paren))
+ T2.consumeClose();
+ ++Idx;
+ continue;
+ }
+ if (CKind == OMPC_when) {
+ OMPTraitInfo &TI = Actions.getASTContext().getNewOMPTraitInfo();
+ parseOMPContextSelectors(ClauseLoc, TI);
+ if (Tok.is(tok::colon))
+ ConsumeAnyToken();
+ }
+ if (Tok.is(tok::r_paren)) {
+ T2.consumeClose();
+ IfStmt = nullptr;
+ } else {
+ StmtResult SR = ParseOpenMPDeclarativeOrExecutableDirective(
+ StmtCtx, /*ReadDirectiveWithinMetadirective=*/true);
+ if (SR.isInvalid()) {
+ SkipUntil(tok::annot_pragma_openmp_end);
+ return StmtError();
}
+ IfStmt = SR.get();
+ }
+ break;
}
+ // Consume pragma end if still present.
+ if (Tok.is(tok::annot_pragma_openmp_end))
+ ConsumeAnnotationToken();
+ if (!IfStmt)
+ return StmtEmpty();
+ SourceLocation EndLoc = Tok.getLocation();
+ Directive = Actions.OpenMP().ActOnOpenMPMetaDirective(
+ /*Clauses=*/{}, IfStmt, Loc, EndLoc,
+ /*Conditions=*/{}, /*Directives=*/{});
return Directive;
}
- std::function<void(StringRef)> DiagUnknownTrait =
- [this, Loc](StringRef ISATrait) {
- // TODO Track the selector locations in a way that is accessible here
- // to improve the diagnostic location.
- Diag(Loc, diag::warn_unknown_declare_variant_isa_trait) << ISATrait;
- };
- TargetOMPContext OMPCtx(ASTContext, std::move(DiagUnknownTrait),
- /* CurrentFunctionDecl */ nullptr,
- ArrayRef<llvm::omp::TraitProperty>(),
- Actions.OpenMP().getOpenMPDeviceNum());
-
- // A single match is returned for OpenMP 5.0
- int BestIdx = getBestVariantMatchForContext(VMIs, OMPCtx);
-
- int Idx = 0;
- // In OpenMP 5.0 metadirective is either replaced by another directive or
- // ignored.
- // TODO: In OpenMP 5.1 generate multiple directives based upon the matches
- // found by getBestWhenMatchForContext.
+
+ SmallVector<Expr *, 4> Conditions;
+ SmallVector<Stmt *, 4> SubDirectives;
+ BalancedDelimiterTracker T2(*this, tok::l_paren,
+ tok::annot_pragma_openmp_end);
+
while (Tok.isNot(tok::annot_pragma_openmp_end)) {
- // OpenMP 5.0 implementation - Skip to the best index found.
- if (Idx++ != BestIdx) {
- 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;
+ if (Tok.isAnnotation()) {
+ break;
}
- OpenMPClauseKind CKind = Tok.isAnnotation()
- ? OMPC_unknown
- : getOpenMPClauseKind(PP.getSpelling(Tok));
- SourceLocation Loc = ConsumeToken();
+ OpenMPClauseKind CKind = getOpenMPClauseKind(PP.getSpelling(Tok));
+ if (CKind == OMPC_unknown)
+ break;
- // Parse '('.
- T.consumeOpen();
+ SourceLocation ClauseLoc = ConsumeToken();
+ T2.consumeOpen();
- // Skip ContextSelectors for when clause
+ Expr *Condition = nullptr;
if (CKind == OMPC_when) {
OMPTraitInfo &TI = Actions.getASTContext().getNewOMPTraitInfo();
- // parse and skip the ContextSelectors
- parseOMPContextSelectors(Loc, TI);
-
- // Parse ':'
- ConsumeAnyToken();
+ parseOMPContextSelectors(ClauseLoc, TI);
+ TI.anyScoreOrCondition([&](Expr *&E, bool IsScore) {
+ if (!IsScore && E) {
+ Condition = E;
+ return true;
+ }
+ return false;
+ });
+ if (Tok.is(tok::colon))
+ ConsumeAnyToken();
}
- // If no directive is passed, skip in OpenMP 5.0.
- // TODO: Generate nothing directive from OpenMP 5.1.
+ StmtResult SubDirective;
if (Tok.is(tok::r_paren)) {
+ SubDirective = StmtEmpty();
+ T2.consumeClose();
+ } else {
+ SubDirective = ParseOpenMPDeclarativeOrExecutableDirective(
+ StmtCtx, /*ReadDirectiveWithinMetadirective=*/true);
+ }
+
+ if (SubDirective.isInvalid()) {
SkipUntil(tok::annot_pragma_openmp_end);
- break;
+ return StmtError();
}
- // Parse Directive
- Directive = ParseOpenMPDeclarativeOrExecutableDirective(
- StmtCtx,
- /*ReadDirectiveWithinMetadirective=*/true);
- break;
+ Conditions.push_back(Condition);
+ Stmt *SubStmt = SubDirective.get();
+ if (isa_and_nonnull<NullStmt>(SubStmt))
+ SubStmt = nullptr;
+ SubDirectives.push_back(SubStmt);
}
- // If no match is found and no otherwise clause is present, skip
- // OMP5.2 Chapter 7.4: If no otherwise clause is specified the effect is as
- // if one was specified without an associated directive variant.
- if (BestIdx == -1 && Idx > 0) {
- assert(Tok.is(tok::annot_pragma_openmp_end) &&
- "Expecting the end of the pragma here");
- ConsumeAnnotationToken();
- return StmtEmpty();
- }
- break;
+
+ SourceLocation EndLoc;
+ if (Tok.is(tok::annot_pragma_openmp_end))
+ EndLoc = ConsumeAnnotationToken();
+ else
+ EndLoc = Loc;
+
+ Directive = Actions.OpenMP().ActOnOpenMPMetaDirective(
+ /*Clauses=*/{}, /*AStmt=*/nullptr, Loc, EndLoc, Conditions,
+ SubDirectives);
+ return Directive;
}
case OMPD_threadprivate: {
// FIXME: Should this be permitted in C++?
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index 6b890f3e76f2a..662346f07d7dc 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -7956,12 +7956,20 @@ setBranchProtectedScope(Sema &SemaRef, OpenMPDirectiveKind DKind, Stmt *AStmt) {
StmtResult SemaOpenMP::ActOnOpenMPMetaDirective(ArrayRef<OMPClause *> Clauses,
Stmt *AStmt,
SourceLocation StartLoc,
- SourceLocation EndLoc) {
- if (!AStmt)
- return StmtError();
+ SourceLocation EndLoc,
+ ArrayRef<Expr *> Conditions,
+ ArrayRef<Stmt *> Directives) {
+ assert(Conditions.size() == Directives.size() &&
+ "Conditions and Directives must have the same size");
+
+ // Determine if any condition is a non-constant runtime expression.
+ // If all conditions are constant (or this is a compile-time-only
+ // metadirective), IfStmt will already have been resolved by the caller
+ // and Conditions/Directives may be empty.
+ Stmt *IfStmt = Conditions.empty() ? AStmt : nullptr;
return OMPMetaDirective::Create(getASTContext(), StartLoc, EndLoc, Clauses,
- AStmt, /*IfStmt=*/nullptr);
+ AStmt, IfStmt, Conditions, Directives);
}
StmtResult
diff --git a/clang/lib/Serialization/ASTReaderStmt.cpp b/clang/lib/Serialization/ASTReaderStmt.cpp
index fb81e4fefdebb..06c3157259df6 100644
--- a/clang/lib/Serialization/ASTReaderStmt.cpp
+++ b/clang/lib/Serialization/ASTReaderStmt.cpp
@@ -2482,7 +2482,7 @@ void ASTStmtReader::VisitOMPLoopDirective(OMPLoopDirective *D) {
void ASTStmtReader::VisitOMPMetaDirective(OMPMetaDirective *D) {
VisitStmt(D);
// The NumClauses field was read in ReadStmtFromStream.
- Record.skipInts(1);
+ Record.skipInts(2);
VisitOMPExecutableDirective(D);
}
@@ -3642,10 +3642,13 @@ Stmt *ASTReader::ReadStmtFromStream(ModuleFile &F) {
S = OMPCanonicalLoop::createEmpty(Context);
break;
- case STMT_OMP_META_DIRECTIVE:
- S = OMPMetaDirective::CreateEmpty(
- Context, Record[ASTStmtReader::NumStmtFields], Empty);
+ case STMT_OMP_META_DIRECTIVE: {
+ unsigned NumClauses = Record[ASTStmtReader::NumStmtFields];
+ unsigned NumVariants = Record[ASTStmtReader::NumStmtFields + 1];
+ S = OMPMetaDirective::CreateEmpty(Context, NumClauses, NumVariants,
+ Empty);
break;
+ }
case STMT_OMP_PARALLEL_DIRECTIVE:
S =
diff --git a/clang/lib/Serialization/ASTWriterStmt.cpp b/clang/lib/Serialization/ASTWriterStmt.cpp
index 4612cd2a7944d..699555a53a463 100644
--- a/clang/lib/Serialization/ASTWriterStmt.cpp
+++ b/clang/lib/Serialization/ASTWriterStmt.cpp
@@ -2499,6 +2499,7 @@ void ASTStmtWriter::VisitOMPLoopDirective(OMPLoopDirective *D) {
void ASTStmtWriter::VisitOMPMetaDirective(OMPMetaDirective *D) {
VisitStmt(D);
Record.push_back(D->getNumClauses());
+ Record.push_back(D->getNumVariants());
VisitOMPExecutableDirective(D);
Code = serialization::STMT_OMP_META_DIRECTIVE;
}
diff --git a/clang/test/OpenMP/metadirective_ast_print.c b/clang/test/OpenMP/metadirective_ast_print.c
index 75ef5fa26827c..d9e2867cb374e 100644
--- a/clang/test/OpenMP/metadirective_ast_print.c
+++ b/clang/test/OpenMP/metadirective_ast_print.c
@@ -126,32 +126,55 @@ void foo1(void) {
// OMP52: void bar(void);
// OMP52: void foo1(void)
+// OMP52-NEXT: #pragma omp metadirective
// OMP52-NEXT: #pragma omp parallel
// OMP52-NEXT: bar()
+// OMP52-NEXT: #pragma omp metadirective
// OMP52-NEXT: #pragma omp parallel
// OMP52-NEXT: bar()
+// OMP52-NEXT: #pragma omp metadirective
// OMP52-NEXT: #pragma omp parallel
-// OMP52-NEXT: bar()
+// OMP52-NEXT: bar();
+// OMP52-NEXT: #pragma omp metadirective
// OMP52-NEXT: #pragma omp parallel
-// OMP52-NEXT: bar()
+// OMP52-NEXT: bar();
+// OMP52-NEXT: #pragma omp metadirective
// OMP52-NEXT: #pragma omp parallel
-// OMP52-NEXT: bar()
+// OMP52-NEXT: bar();
+// OMP52-NEXT: #pragma omp metadirective
// OMP52-NEXT: #pragma omp parallel for
// OMP52-NEXT: for (int i = 0; i < 100; i++)
-// OMP52: #pragma omp parallel
+// OMP52: #pragma omp metadirective
+// OMP52-NEXT: #pragma omp parallel
// OMP52-NEXT: for (int i = 0; i < 100; i++)
-// OMP52: #pragma omp parallel for
+// OMP52: #pragma omp metadirective
+// OMP52-NEXT: #pragma omp parallel
// OMP52-NEXT: for (int i = 0; i < 100; i++)
+// OMP52: #pragma omp metadirective
// OMP52: #pragma omp parallel
// OMP52-NEXT: for (int i = 0; i < 100; i++)
+// OMP52: #pragma omp metadirective
// OMP52: #pragma omp parallel for
// OMP52-NEXT: for (int i = 0; i < 16; i++) {
// OMP52-NEXT: #pragma omp simd
// OMP52-NEXT: for (int j = 0; j < 16; j++)
+// OMP52: #pragma omp metadirective
+// OMP52-NEXT: #pragma omp parallel for
+// OMP52-NEXT: for (int i = 0; i < 100; i++)
// OMP52-GPU: #pragma omp teams distribute parallel for
-// OMP52-GPU-NEXT: for (int i = 0; i < 100; i++)
+// OMP52: #pragma omp metadirective
+// OMP52-NEXT: #pragma omp parallel for
+// OMP52-NEXT: for (int i = 0; i < 100; i++)
// OMP52: for (int i = 0; i < 16; i++)
// OMP52: for (int i = 0; i < 16; i++)
+// OMP52: for (int i = 0; i < 10; i++)
+// OMP52: #pragma omp metadirective
+// OMP52-NEXT: #pragma omp parallel
+// OMP52-NEXT: for (int i = 0; i < 10; i++)
+// OMP52: #pragma omp metadirective
+// OMP52-NEXT: #pragma omp parallel for
+// OMP52-NEXT: for (int i = 0; i < 10; i++)
+// OMP52-GPU-NEXT: for (int i = 0; i < 100; i++)
#else
void bar(void);
@@ -264,33 +287,58 @@ void foo2(void) {
// DEFAULT: void bar(void);
// DEFAULT: void foo2(void)
+// DEFAULT: #pragma omp metadirective
// DEFAULT-NEXT: #pragma omp parallel
// DEFAULT-NEXT: bar()
+// DEFAULT-NEXT: #pragma omp metadirective
// DEFAULT-NEXT: #pragma omp parallel
// DEFAULT-NEXT: bar()
+// DEFAULT-NEXT: #pragma omp metadirective
// DEFAULT-NEXT: #pragma omp parallel
// DEFAULT-NEXT: bar()
+// DEFAULT-NEXT: #pragma omp metadirective
// DEFAULT-NEXT: #pragma omp parallel
// DEFAULT-NEXT: bar()
-// DEFAULT-NEXT: #pragma omp parallel
-// DEFAULT-NEXT: bar()
+// DEFAULT-NEXT: #pragma omp metadirective
+// DEFAULT-NEXT: #pragma omp parallel
+// DEFAULT-NEXT: bar();
+// DEFAULT-NEXT: #pragma omp metadirective
// DEFAULT-NEXT: #pragma omp parallel for
// DEFAULT-NEXT: for (int i = 0; i < 100; i++)
-// DEFAULT: #pragma omp parallel
+// DEFAULT: #pragma omp metadirective
+// DEFAULT-NEXT: #pragma omp parallel
// DEFAULT-NEXT: for (int i = 0; i < 100; i++)
-// DEFAULT: #pragma omp parallel for
+// DEFAULT: #pragma omp metadirective
+// DEFAULT-NEXT: #pragma omp parallel
// DEFAULT-NEXT: for (int i = 0; i < 100; i++)
-// DEFAULT: #pragma omp parallel
+// DEFAULT: #pragma omp metadirective
+// DEFAULT-NEXT: #pragma omp parallel
// DEFAULT-NEXT: for (int i = 0; i < 100; i++)
-// DEFAULT: #pragma omp parallel for
+// DEFAULT: #pragma omp metadirective
+// DEFAULT-NEXT: #pragma omp parallel for
// DEFAULT-NEXT: for (int i = 0; i < 16; i++) {
// DEFAULT-NEXT: #pragma omp simd
// DEFAULT-NEXT: for (int j = 0; j < 16; j++)
// DEFAULT-GPU: #pragma omp teams distribute parallel for
// DEFAULT-GPU-NEXT: for (int i = 0; i < 100; i++)
+// DEFAULT: #pragma omp metadirective
+// DEFAULT-NEXT: #pragma omp parallel for
+// DEFAULT-NEXT: for (int i = 0; i < 100; i++)
+// DEFAULT: #pragma omp metadirective
+// DEFAULT-NEXT: #pragma omp parallel for
+// DEFAULT-NEXT: for (int i = 0; i < 100; i++)
// DEFAULT: for (int i = 0; i < 16; i++)
// DEFAULT: for (int i = 0; i < 16; i++)
-
+// DEFAULT: for (int i = 0; i < 10; i++)
+// DEFAULT: #pragma omp metadirective
+// DEFAULT-NEXT: #pragma omp parallel
+// DEFAULT-NEXT: for (int i = 0; i < 10; i++)
+// DEFAULT: #pragma omp metadirective
+// DEFAULT-NEXT: #pragma omp parallel for
+// DEFAULT-NEXT: for (int i = 0; i < 10; i++)
+// DEFAULT: #pragma omp metadirective
+// DEFAULT-NEXT: #pragma omp parallel for
+// DEFAULT-NEXT: for (int i = 0; i < 10; i++)
#endif
#endif
diff --git a/clang/test/OpenMP/metadirective_ast_print.cpp b/clang/test/OpenMP/metadirective_ast_print.cpp
new file mode 100644
index 0000000000000..d2fa8d0d0f599
--- /dev/null
+++ b/clang/test/OpenMP/metadirective_ast_print.cpp
@@ -0,0 +1,14 @@
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=52 -triple x86_64-unknown-linux-gnu -ast-print %s -o -
+
+// expected-no-diagnostics
+void bar ();
+void test_nonconstant_condition(bool use_gpu) {
+ #pragma omp metadirective \
+ when(user={condition(use_gpu)}: parallel) \
+ otherwise(single)
+ {
+ int x = 0;
+ }
+}
+
+// CHECK: #pragma omp metadirective when(use_gpu: #pragma omp parallel
>From ddc100a473a0c0f81dc7f482c14965907892a9ae Mon Sep 17 00:00:00 2001
From: Ammarguellat <zahira.ammarguellat at intel.com>
Date: Mon, 27 Apr 2026 07:11:39 -0700
Subject: [PATCH 4/4] Removed codegen code and focused on getting the correct
AST
---
clang/lib/AST/StmtPrinter.cpp | 21 +-
clang/lib/Parse/ParseOpenMP.cpp | 256 ++++++++++--------
clang/lib/Sema/SemaOpenMP.cpp | 3 +-
clang/test/OpenMP/metadirective_ast_print.c | 75 +----
clang/test/OpenMP/metadirective_ast_print.cpp | 14 +-
.../metadirective_user_condition_codegen.cpp | 134 ---------
6 files changed, 186 insertions(+), 317 deletions(-)
delete mode 100644 clang/test/OpenMP/metadirective_user_condition_codegen.cpp
diff --git a/clang/lib/AST/StmtPrinter.cpp b/clang/lib/AST/StmtPrinter.cpp
index 52842612b9244..26f5ad063783a 100644
--- a/clang/lib/AST/StmtPrinter.cpp
+++ b/clang/lib/AST/StmtPrinter.cpp
@@ -761,31 +761,34 @@ void StmtPrinter::PrintOMPExecutableDirective(OMPExecutableDirective *S,
}
void StmtPrinter::VisitOMPMetaDirective(OMPMetaDirective *Node) {
- Indent() << "#pragma omp metadirective";
if (Node->getNumVariants() > 0) {
+ // Runtime path. Print full metadirective with all variants.
+ Indent() << "#pragma omp metadirective";
for (unsigned I = 0; I < Node->getNumVariants(); ++I) {
Expr *Cond = Node->getVariantCondition(I);
Stmt *Dir = Node->getVariantDirective(I);
-
if (Cond) {
OS << " when(";
- // Print the condition expression.
Cond->printPretty(OS, nullptr, Policy);
OS << ":";
} else {
OS << " otherwise(";
}
-
- // Print the sub-directive.
- if (Dir && !isa<NullStmt>(Dir)) {
- OS << " ";
- PrintStmt(Dir);
+ if (Dir) {
+ if (auto *OED = dyn_cast<OMPExecutableDirective>(Dir)) {
+ unsigned OpenMPVersion = Context ? Context->getLangOpts().OpenMP
+ : llvm::omp::FallbackVersion;
+ OS << " #pragma omp "
+ << getOpenMPDirectiveName(OED->getDirectiveKind(), OpenMPVersion);
+ }
}
OS << ")";
}
OS << NL;
+ if (Node->hasAssociatedStmt() && Node->getAssociatedStmt())
+ PrintStmt(Node->getAssociatedStmt());
} else {
- OS << NL;
+ Indent() << "#pragma omp metadirective" << NL;
if (Stmt *If = Node->getIfStmt())
PrintStmt(If);
}
diff --git a/clang/lib/Parse/ParseOpenMP.cpp b/clang/lib/Parse/ParseOpenMP.cpp
index f9f114bc02731..08b2f2ca05772 100644
--- a/clang/lib/Parse/ParseOpenMP.cpp
+++ b/clang/lib/Parse/ParseOpenMP.cpp
@@ -2594,6 +2594,7 @@ StmtResult Parser::ParseOpenMPDeclarativeOrExecutableDirective(
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
@@ -2672,44 +2673,121 @@ StmtResult Parser::ParseOpenMPDeclarativeOrExecutableDirective(
// End of the first iteration. Parser is reset to the start of metadirective
// Detect non-constant user conditions.
- bool HasNonConstantUserCondition = false;
+ bool NeedsRuntimeSelection = false;
for (unsigned I = 0; I < TraitInfos.size(); ++I) {
if (ClauseKinds[I] != OMPC_when)
continue;
TraitInfos[I]->anyScoreOrCondition([&](Expr *&E, bool IsScore) {
if (!IsScore && E && !E->isIntegerConstantExpr(ASTContext)) {
- HasNonConstantUserCondition = true;
+ NeedsRuntimeSelection = true;
return true;
}
return false;
});
- if (HasNonConstantUserCondition)
+ if (NeedsRuntimeSelection )
break;
}
- if (!HasNonConstantUserCondition) {
- std::function<void(StringRef)> DiagUnknownTrait =
- [this, Loc](StringRef ISATrait) {
- Diag(Loc, diag::warn_unknown_declare_variant_isa_trait) << ISATrait;
- };
- TargetOMPContext OMPCtx(ASTContext, std::move(DiagUnknownTrait),
- /*CurrentFunctionDecl=*/nullptr,
- ArrayRef<llvm::omp::TraitProperty>(),
- Actions.OpenMP().getOpenMPDeviceNum());
+ if (!NeedsRuntimeSelection) {
+ std::function<void(StringRef)> DiagUnknownTrait =
+ [this, Loc](StringRef ISATrait) {
+ // TODO Track the selector locations in a way that is accessible here
+ // to improve the diagnostic location.
+ Diag(Loc, diag::warn_unknown_declare_variant_isa_trait) << ISATrait;
+ };
+ TargetOMPContext OMPCtx(ASTContext, std::move(DiagUnknownTrait),
+ /* CurrentFunctionDecl */ nullptr,
+ ArrayRef<llvm::omp::TraitProperty>(),
+ Actions.OpenMP().getOpenMPDeviceNum());
+
+ // A single match is returned for OpenMP 5.0
+ int BestIdx = getBestVariantMatchForContext(VMIs, OMPCtx);
+
+ int Idx = 0;
+ // In OpenMP 5.0 metadirective is either replaced by another directive or
+ // ignored.
+ // TODO: In OpenMP 5.1 generate multiple directives based upon the matches
+ // found by getBestWhenMatchForContext.
+ while (Tok.isNot(tok::annot_pragma_openmp_end)) {
+ // OpenMP 5.0 implementation - Skip to the best index found.
+ if (Idx++ != BestIdx) {
+ 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 the ContextSelectors
+ parseOMPContextSelectors(Loc, TI);
- int BestIdx = getBestVariantMatchForContext(VMIs, OMPCtx);
+ // Parse ':'
+ ConsumeAnyToken();
+ }
- if (BestIdx == -1) {
+ // If no directive is passed, skip in OpenMP 5.0.
+ // TODO: Generate nothing directive from OpenMP 5.1.
+ if (Tok.is(tok::r_paren)) {
SkipUntil(tok::annot_pragma_openmp_end);
- if (Tok.is(tok::annot_pragma_openmp_end))
- ConsumeAnnotationToken();
- return StmtEmpty();
+ break;
}
+ // Parse Directive
+ Directive = ParseOpenMPDeclarativeOrExecutableDirective(
+ StmtCtx,
+ /*ReadDirectiveWithinMetadirective=*/true);
+ break;
+ }
+ // If no match is found and no otherwise clause is present, skip
+ // OMP5.2 Chapter 7.4: If no otherwise clause is specified the effect is as
+ // if one was specified without an associated directive variant.
+ if (BestIdx == -1 && Idx > 0) {
+ assert(Tok.is(tok::annot_pragma_openmp_end) &&
+ "Expecting the end of the pragma here");
+ ConsumeAnnotationToken();
+ return StmtEmpty();
+ }
+ break;
+ }
+
+ // Runtime path: NeedsRuntimeSelection == true.
+ // For each clause, do a separate tentative parse to extract its
+ // sub-directive. This is needed because
+ // ParseOpenMPDeclarativeOrExecutableDirective consumes the pragma
+ // end on each call, so we can't parse multiple clauses in one
+ // linear pass.
+ SmallVector<Expr *, 4> Conditions;
+ SmallVector<Stmt *, 4> SubDirectives;
+
+ for (unsigned TargetIdx = 0; TargetIdx < TraitInfos.size(); ++TargetIdx) {
+ // Revert to the start of the clause list for each iteration.
+ TentativeParsingAction TPA2(*this);
BalancedDelimiterTracker T2(*this, tok::l_paren,
tok::annot_pragma_openmp_end);
- int Idx = 0;
- Stmt *IfStmt = nullptr;
+ unsigned Idx = 0;
+ Expr *Condition = nullptr;
+ Stmt *SubStmt = nullptr;
while (Tok.isNot(tok::annot_pragma_openmp_end)) {
if (Tok.isAnnotation())
@@ -2717,18 +2795,19 @@ StmtResult Parser::ParseOpenMPDeclarativeOrExecutableDirective(
OpenMPClauseKind CKind = getOpenMPClauseKind(PP.getSpelling(Tok));
if (CKind == OMPC_unknown)
break;
+
SourceLocation ClauseLoc = ConsumeToken();
T2.consumeOpen();
- if (Idx != BestIdx) {
- 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--;
- if (Tok.is(tok::annot_pragma_openmp_end))
- break;
+ if (Idx != TargetIdx) {
+ 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--;
+ if (Tok.is(tok::annot_pragma_openmp_end))
+ break;
ConsumeAnyToken();
}
if (Tok.is(tok::r_paren))
@@ -2736,100 +2815,63 @@ StmtResult Parser::ParseOpenMPDeclarativeOrExecutableDirective(
++Idx;
continue;
}
- if (CKind == OMPC_when) {
- OMPTraitInfo &TI = Actions.getASTContext().getNewOMPTraitInfo();
- parseOMPContextSelectors(ClauseLoc, TI);
- if (Tok.is(tok::colon))
- ConsumeAnyToken();
- }
- if (Tok.is(tok::r_paren)) {
- T2.consumeClose();
- IfStmt = nullptr;
- } else {
- StmtResult SR = ParseOpenMPDeclarativeOrExecutableDirective(
- StmtCtx, /*ReadDirectiveWithinMetadirective=*/true);
- if (SR.isInvalid()) {
- SkipUntil(tok::annot_pragma_openmp_end);
- return StmtError();
- }
- IfStmt = SR.get();
- }
- break;
- }
-
- // Consume pragma end if still present.
- if (Tok.is(tok::annot_pragma_openmp_end))
- ConsumeAnnotationToken();
- if (!IfStmt)
- return StmtEmpty();
- SourceLocation EndLoc = Tok.getLocation();
- Directive = Actions.OpenMP().ActOnOpenMPMetaDirective(
- /*Clauses=*/{}, IfStmt, Loc, EndLoc,
- /*Conditions=*/{}, /*Directives=*/{});
- return Directive;
- }
-
- SmallVector<Expr *, 4> Conditions;
- SmallVector<Stmt *, 4> SubDirectives;
- BalancedDelimiterTracker T2(*this, tok::l_paren,
- tok::annot_pragma_openmp_end);
- while (Tok.isNot(tok::annot_pragma_openmp_end)) {
- if (Tok.isAnnotation()) {
- break;
- }
-
- OpenMPClauseKind CKind = getOpenMPClauseKind(PP.getSpelling(Tok));
- if (CKind == OMPC_unknown)
- break;
-
- SourceLocation ClauseLoc = ConsumeToken();
- T2.consumeOpen();
+ // This is the target clause. Parse it fully.
+ if (CKind == OMPC_when) {
+ OMPTraitInfo &TI = Actions.getASTContext().getNewOMPTraitInfo();
+ parseOMPContextSelectors(ClauseLoc, TI);
+ TI.anyScoreOrCondition([&](Expr *&E, bool IsScore) {
+ if (!IsScore && E) {
+ Condition = E;
+ return true;
+ }
+ return false;
+ });
+ if (Tok.is(tok::colon))
+ ConsumeAnyToken();
+ }
+ // For 'otherwise'/'default', Condition stays nullptr.
- Expr *Condition = nullptr;
- if (CKind == OMPC_when) {
- OMPTraitInfo &TI = Actions.getASTContext().getNewOMPTraitInfo();
- parseOMPContextSelectors(ClauseLoc, TI);
- TI.anyScoreOrCondition([&](Expr *&E, bool IsScore) {
- if (!IsScore && E) {
- Condition = E;
- return true;
+ if (Tok.is(tok::r_paren)) {
+ T2.consumeClose();
+ SubStmt = nullptr;
+ } else {
+ StmtResult SR = ParseOpenMPDeclarativeOrExecutableDirective(
+ StmtCtx, /*ReadDirectiveWithinMetadirective=*/true);
+ if (SR.isInvalid()) {
+ TPA2.Commit();
+ SkipUntil(tok::annot_pragma_openmp_end);
+ return StmtError();
}
- return false;
- });
- if (Tok.is(tok::colon))
- ConsumeAnyToken();
- }
-
- StmtResult SubDirective;
- if (Tok.is(tok::r_paren)) {
- SubDirective = StmtEmpty();
- T2.consumeClose();
- } else {
- SubDirective = ParseOpenMPDeclarativeOrExecutableDirective(
- StmtCtx, /*ReadDirectiveWithinMetadirective=*/true);
- }
-
- if (SubDirective.isInvalid()) {
- SkipUntil(tok::annot_pragma_openmp_end);
- return StmtError();
+ SubStmt = SR.get();
+ if (isa_and_nonnull<NullStmt>(SubStmt))
+ SubStmt = nullptr;
+ }
+ break;
}
+ // Revert token stream for the next iteration.
+ TPA2.Revert();
Conditions.push_back(Condition);
- Stmt *SubStmt = SubDirective.get();
- if (isa_and_nonnull<NullStmt>(SubStmt))
- SubStmt = nullptr;
- SubDirectives.push_back(SubStmt);
+ SubDirectives.push_back(SubStmt);
}
+ // After all clauses parsed via tentative passes, consume pragma end once.
+ SkipUntil(tok::annot_pragma_openmp_end);
SourceLocation EndLoc;
if (Tok.is(tok::annot_pragma_openmp_end))
EndLoc = ConsumeAnnotationToken();
else
EndLoc = Loc;
+ StmtResult AssocStmt;
+ if (Tok.isNot(tok::r_brace) && Tok.isNot(tok::eof) &&
+ Tok.isNot(tok::annot_pragma_openmp) &&
+ Tok.isNot(tok::annot_attr_openmp)) {
+ AssocStmt = ParseStatement();
+ }
Directive = Actions.OpenMP().ActOnOpenMPMetaDirective(
- /*Clauses=*/{}, /*AStmt=*/nullptr, Loc, EndLoc, Conditions,
+ /*Clauses=*/{}, /*AStmt=*/AssocStmt.get(), Loc, EndLoc, Conditions,
SubDirectives);
return Directive;
}
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index 662346f07d7dc..4655685fcbaf4 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -7967,9 +7967,10 @@ StmtResult SemaOpenMP::ActOnOpenMPMetaDirective(ArrayRef<OMPClause *> Clauses,
// metadirective), IfStmt will already have been resolved by the caller
// and Conditions/Directives may be empty.
Stmt *IfStmt = Conditions.empty() ? AStmt : nullptr;
+ Stmt *AssocStmt = Conditions.empty() ? nullptr : AStmt;
return OMPMetaDirective::Create(getASTContext(), StartLoc, EndLoc, Clauses,
- AStmt, IfStmt, Conditions, Directives);
+ AssocStmt, IfStmt, Conditions, Directives);
}
StmtResult
diff --git a/clang/test/OpenMP/metadirective_ast_print.c b/clang/test/OpenMP/metadirective_ast_print.c
index d9e2867cb374e..42ce2bc4db764 100644
--- a/clang/test/OpenMP/metadirective_ast_print.c
+++ b/clang/test/OpenMP/metadirective_ast_print.c
@@ -126,55 +126,32 @@ void foo1(void) {
// OMP52: void bar(void);
// OMP52: void foo1(void)
-// OMP52-NEXT: #pragma omp metadirective
// OMP52-NEXT: #pragma omp parallel
// OMP52-NEXT: bar()
-// OMP52-NEXT: #pragma omp metadirective
// OMP52-NEXT: #pragma omp parallel
// OMP52-NEXT: bar()
-// OMP52-NEXT: #pragma omp metadirective
// OMP52-NEXT: #pragma omp parallel
-// OMP52-NEXT: bar();
-// OMP52-NEXT: #pragma omp metadirective
+// OMP52-NEXT: bar()
// OMP52-NEXT: #pragma omp parallel
-// OMP52-NEXT: bar();
-// OMP52-NEXT: #pragma omp metadirective
+// OMP52-NEXT: bar()
// OMP52-NEXT: #pragma omp parallel
-// OMP52-NEXT: bar();
-// OMP52-NEXT: #pragma omp metadirective
+// OMP52-NEXT: bar()
// OMP52-NEXT: #pragma omp parallel for
// OMP52-NEXT: for (int i = 0; i < 100; i++)
-// OMP52: #pragma omp metadirective
-// OMP52-NEXT: #pragma omp parallel
+// OMP52: #pragma omp parallel
// OMP52-NEXT: for (int i = 0; i < 100; i++)
-// OMP52: #pragma omp metadirective
-// OMP52-NEXT: #pragma omp parallel
+// OMP52: #pragma omp parallel for
// OMP52-NEXT: for (int i = 0; i < 100; i++)
-// OMP52: #pragma omp metadirective
// OMP52: #pragma omp parallel
// OMP52-NEXT: for (int i = 0; i < 100; i++)
-// OMP52: #pragma omp metadirective
// OMP52: #pragma omp parallel for
// OMP52-NEXT: for (int i = 0; i < 16; i++) {
// OMP52-NEXT: #pragma omp simd
// OMP52-NEXT: for (int j = 0; j < 16; j++)
-// OMP52: #pragma omp metadirective
-// OMP52-NEXT: #pragma omp parallel for
-// OMP52-NEXT: for (int i = 0; i < 100; i++)
// OMP52-GPU: #pragma omp teams distribute parallel for
-// OMP52: #pragma omp metadirective
-// OMP52-NEXT: #pragma omp parallel for
-// OMP52-NEXT: for (int i = 0; i < 100; i++)
+// OMP52-GPU-NEXT: for (int i = 0; i < 100; i++)
// OMP52: for (int i = 0; i < 16; i++)
// OMP52: for (int i = 0; i < 16; i++)
-// OMP52: for (int i = 0; i < 10; i++)
-// OMP52: #pragma omp metadirective
-// OMP52-NEXT: #pragma omp parallel
-// OMP52-NEXT: for (int i = 0; i < 10; i++)
-// OMP52: #pragma omp metadirective
-// OMP52-NEXT: #pragma omp parallel for
-// OMP52-NEXT: for (int i = 0; i < 10; i++)
-// OMP52-GPU-NEXT: for (int i = 0; i < 100; i++)
#else
void bar(void);
@@ -287,59 +264,33 @@ void foo2(void) {
// DEFAULT: void bar(void);
// DEFAULT: void foo2(void)
-// DEFAULT: #pragma omp metadirective
// DEFAULT-NEXT: #pragma omp parallel
// DEFAULT-NEXT: bar()
-// DEFAULT-NEXT: #pragma omp metadirective
// DEFAULT-NEXT: #pragma omp parallel
// DEFAULT-NEXT: bar()
-// DEFAULT-NEXT: #pragma omp metadirective
// DEFAULT-NEXT: #pragma omp parallel
// DEFAULT-NEXT: bar()
-// DEFAULT-NEXT: #pragma omp metadirective
// DEFAULT-NEXT: #pragma omp parallel
// DEFAULT-NEXT: bar()
-// DEFAULT-NEXT: #pragma omp metadirective
-// DEFAULT-NEXT: #pragma omp parallel
-// DEFAULT-NEXT: bar();
-// DEFAULT-NEXT: #pragma omp metadirective
+// DEFAULT-NEXT: #pragma omp parallel
+// DEFAULT-NEXT: bar()
// DEFAULT-NEXT: #pragma omp parallel for
// DEFAULT-NEXT: for (int i = 0; i < 100; i++)
-// DEFAULT: #pragma omp metadirective
-// DEFAULT-NEXT: #pragma omp parallel
+// DEFAULT: #pragma omp parallel
// DEFAULT-NEXT: for (int i = 0; i < 100; i++)
-// DEFAULT: #pragma omp metadirective
-// DEFAULT-NEXT: #pragma omp parallel
+// DEFAULT: #pragma omp parallel for
// DEFAULT-NEXT: for (int i = 0; i < 100; i++)
-// DEFAULT: #pragma omp metadirective
-// DEFAULT-NEXT: #pragma omp parallel
+// DEFAULT: #pragma omp parallel
// DEFAULT-NEXT: for (int i = 0; i < 100; i++)
-// DEFAULT: #pragma omp metadirective
-// DEFAULT-NEXT: #pragma omp parallel for
+// DEFAULT: #pragma omp parallel for
// DEFAULT-NEXT: for (int i = 0; i < 16; i++) {
// DEFAULT-NEXT: #pragma omp simd
// DEFAULT-NEXT: for (int j = 0; j < 16; j++)
// DEFAULT-GPU: #pragma omp teams distribute parallel for
// DEFAULT-GPU-NEXT: for (int i = 0; i < 100; i++)
-// DEFAULT: #pragma omp metadirective
-// DEFAULT-NEXT: #pragma omp parallel for
-// DEFAULT-NEXT: for (int i = 0; i < 100; i++)
-// DEFAULT: #pragma omp metadirective
-// DEFAULT-NEXT: #pragma omp parallel for
-// DEFAULT-NEXT: for (int i = 0; i < 100; i++)
// DEFAULT: for (int i = 0; i < 16; i++)
// DEFAULT: for (int i = 0; i < 16; i++)
-// DEFAULT: for (int i = 0; i < 10; i++)
-// DEFAULT: #pragma omp metadirective
-// DEFAULT-NEXT: #pragma omp parallel
-// DEFAULT-NEXT: for (int i = 0; i < 10; i++)
-// DEFAULT: #pragma omp metadirective
-// DEFAULT-NEXT: #pragma omp parallel for
-// DEFAULT-NEXT: for (int i = 0; i < 10; i++)
-// DEFAULT: #pragma omp metadirective
-// DEFAULT-NEXT: #pragma omp parallel for
-// DEFAULT-NEXT: for (int i = 0; i < 10; i++)
+
#endif
#endif
-
diff --git a/clang/test/OpenMP/metadirective_ast_print.cpp b/clang/test/OpenMP/metadirective_ast_print.cpp
index d2fa8d0d0f599..3c1432838e056 100644
--- a/clang/test/OpenMP/metadirective_ast_print.cpp
+++ b/clang/test/OpenMP/metadirective_ast_print.cpp
@@ -1,14 +1,20 @@
-// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=52 -triple x86_64-unknown-linux-gnu -ast-print %s -o -
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=52 \
+// RUN: -triple x86_64-unknown-linux-gnu -ast-print %s -o - | FileCheck %s
// expected-no-diagnostics
-void bar ();
+
+void bar();
void test_nonconstant_condition(bool use_gpu) {
#pragma omp metadirective \
when(user={condition(use_gpu)}: parallel) \
otherwise(single)
{
- int x = 0;
+ bar();
}
}
-// CHECK: #pragma omp metadirective when(use_gpu: #pragma omp parallel
+// LABEL: void test_nonconstant_condition(bool use_gpu)
+// CHECK: #pragma omp metadirective when(use_gpu: #pragma omp parallel) otherwise( #pragma omp single)
+// CHECK: bar();
+
+
diff --git a/clang/test/OpenMP/metadirective_user_condition_codegen.cpp b/clang/test/OpenMP/metadirective_user_condition_codegen.cpp
deleted file mode 100644
index f2a087f28d1f3..0000000000000
--- a/clang/test/OpenMP/metadirective_user_condition_codegen.cpp
+++ /dev/null
@@ -1,134 +0,0 @@
-// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6
-// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=52 \
-// RUN: -triple x86_64-unknown-linux -emit-llvm %s -o - | FileCheck %s
-
-// expected-no-diagnostics
-
-#ifndef HEADER
-#define HEADER
-
-void bar();
-
-
-// CHECK-LABEL: define dso_local void @_Z27test_runtime_user_conditionb(
-// CHECK-SAME: i1 noundef zeroext [[USE_GPU:%.*]]) #[[ATTR0:[0-9]+]] {
-// CHECK-NEXT: [[ENTRY:.*:]]
-// CHECK-NEXT: [[USE_GPU_ADDR:%.*]] = alloca i8, align 1
-// CHECK-NEXT: [[STOREDV:%.*]] = zext i1 [[USE_GPU]] to i8
-// CHECK-NEXT: store i8 [[STOREDV]], ptr [[USE_GPU_ADDR]], align 1
-// CHECK-NEXT: [[TMP0:%.*]] = load i8, ptr [[USE_GPU_ADDR]], align 1
-// CHECK-NEXT: [[LOADEDV:%.*]] = trunc i8 [[TMP0]] to i1
-// CHECK-NEXT: br i1 [[LOADEDV]], label %[[OMP_META_THEN:.*]], label %[[OMP_META_ELSE:.*]]
-// CHECK: [[OMP_META_THEN]]:
-// CHECK-NEXT: call void @_Z3barv()
-// CHECK-NEXT: br label %[[OMP_META_END:.*]]
-// CHECK: [[OMP_META_ELSE]]:
-// CHECK-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1:[0-9]+]], i32 0, ptr @.omp_outlined.)
-// CHECK-NEXT: br label %[[OMP_META_END]]
-// CHECK: [[OMP_META_END]]:
-// CHECK-NEXT: ret void
-//
-void test_runtime_user_condition(bool use_gpu) {
- #pragma omp metadirective \
- when(user={condition(use_gpu)}: target) \
- otherwise(parallel)
- {
- bar();
- }
-}
-
-
-// CHECK-LABEL: define dso_local void @_Z35test_runtime_user_condition_integeri(
-// CHECK-SAME: i32 noundef [[THRESHOLD:%.*]]) #[[ATTR0]] {
-// CHECK-NEXT: [[ENTRY:.*:]]
-// CHECK-NEXT: [[THRESHOLD_ADDR:%.*]] = alloca i32, align 4
-// CHECK-NEXT: [[I:%.*]] = alloca i32, align 4
-// CHECK-NEXT: store i32 [[THRESHOLD]], ptr [[THRESHOLD_ADDR]], align 4
-// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr [[THRESHOLD_ADDR]], align 4
-// CHECK-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP0]], 100
-// CHECK-NEXT: br i1 [[CMP]], label %[[OMP_META_THEN:.*]], label %[[OMP_META_ELSE:.*]]
-// CHECK: [[OMP_META_THEN]]:
-// CHECK-NEXT: store i32 0, ptr [[I]], align 4
-// CHECK-NEXT: br label %[[FOR_COND:.*]]
-// CHECK: [[FOR_COND]]:
-// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[I]], align 4
-// CHECK-NEXT: [[CMP1:%.*]] = icmp slt i32 [[TMP1]], 10
-// CHECK-NEXT: br i1 [[CMP1]], label %[[FOR_BODY:.*]], label %[[FOR_END:.*]]
-// CHECK: [[FOR_BODY]]:
-// CHECK-NEXT: call void @_Z3barv()
-// CHECK-NEXT: br label %[[FOR_INC:.*]]
-// CHECK: [[FOR_INC]]:
-// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[I]], align 4
-// CHECK-NEXT: [[INC:%.*]] = add nsw i32 [[TMP2]], 1
-// CHECK-NEXT: store i32 [[INC]], ptr [[I]], align 4
-// CHECK-NEXT: br label %[[FOR_COND]], !llvm.loop [[LOOP4:![0-9]+]]
-// CHECK: [[FOR_END]]:
-// CHECK-NEXT: br label %[[OMP_META_END:.*]]
-// CHECK: [[OMP_META_ELSE]]:
-// CHECK-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 0, ptr @.omp_outlined..1)
-// CHECK-NEXT: br label %[[OMP_META_END]]
-// CHECK: [[OMP_META_END]]:
-// CHECK-NEXT: ret void
-//
-void test_runtime_user_condition_integer(int threshold) {
-
-
-
- #pragma omp metadirective \
- when(user={condition(threshold > 100)}: parallel for) \
- otherwise(parallel)
- {
- for (int i = 0; i < 10; ++i) {
- bar();
- }
- }
-}
-
-// CHECK-LABEL: define dso_local void @_Z33test_constant_user_condition_truev(
-// CHECK-SAME: ) #[[ATTR0]] {
-// CHECK-NEXT: [[ENTRY:.*:]]
-// CHECK-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 0, ptr @_Z33test_constant_user_condition_truev.omp_outlined)
-// CHECK-NEXT: ret void
-//
-void test_constant_user_condition_true() {
- // For constant conditions, no runtime branch should be generated
-
- #pragma omp metadirective \
- when(user={condition(true)}: parallel) \
- otherwise(single)
- {
- bar();
- }
-}
-
-// CHECK-LABEL: define dso_local void @_Z34test_constant_user_condition_falsev(
-// CHECK-SAME: ) #[[ATTR0]] {
-// CHECK-NEXT: [[ENTRY:.*:]]
-// CHECK-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])
-// CHECK-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_single(ptr @[[GLOB1]], i32 [[TMP0]])
-// CHECK-NEXT: [[TMP2:%.*]] = icmp ne i32 [[TMP1]], 0
-// CHECK-NEXT: br i1 [[TMP2]], label %[[OMP_IF_THEN:.*]], label %[[OMP_IF_END:.*]]
-// CHECK: [[OMP_IF_THEN]]:
-// CHECK-NEXT: call void @_Z3barv()
-// CHECK-NEXT: call void @__kmpc_end_single(ptr @[[GLOB1]], i32 [[TMP0]])
-// CHECK-NEXT: br label %[[OMP_IF_END]]
-// CHECK: [[OMP_IF_END]]:
-// CHECK-NEXT: call void @__kmpc_barrier(ptr @[[GLOB2:[0-9]+]], i32 [[TMP0]])
-// CHECK-NEXT: ret void
-//
-void test_constant_user_condition_false() {
- // For constant conditions, no runtime branch should be generated
-
- #pragma omp metadirective \
- when(user={condition(false)}: parallel) \
- otherwise(single)
- {
- bar();
- }
-}
-
-#endif
-//.
-// CHECK: [[LOOP4]] = distinct !{[[LOOP4]], [[META5:![0-9]+]]}
-// CHECK: [[META5]] = !{!"llvm.loop.mustprogress"}
-//.
More information about the cfe-commits
mailing list