[clang] 80182a7 - [OpenACC][CIR] Implement 'wait' directive lowering
via cfe-commits
cfe-commits at lists.llvm.org
Thu Apr 24 14:26:28 PDT 2025
Author: erichkeane
Date: 2025-04-24T14:26:24-07:00
New Revision: 80182a7d5d66c8dc90bb4623c1f722aba7ebe45b
URL: https://github.com/llvm/llvm-project/commit/80182a7d5d66c8dc90bb4623c1f722aba7ebe45b
DIFF: https://github.com/llvm/llvm-project/commit/80182a7d5d66c8dc90bb4623c1f722aba7ebe45b.diff
LOG: [OpenACC][CIR] Implement 'wait' directive lowering
This construct has a couple of 'intexprs' which are lowered the same way
as clauses, plus has a pair of simple clauses that needed lowering.
This patch does all of that.
Added:
clang/test/CIR/CodeGenOpenACC/wait.c
Modified:
clang/lib/CIR/CodeGen/CIRGenFunction.h
clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
Removed:
################################################################################
diff --git a/clang/lib/CIR/CodeGen/CIRGenFunction.h b/clang/lib/CIR/CodeGen/CIRGenFunction.h
index f533d0ab53cd2..74fcd081dec18 100644
--- a/clang/lib/CIR/CodeGen/CIRGenFunction.h
+++ b/clang/lib/CIR/CodeGen/CIRGenFunction.h
@@ -626,10 +626,9 @@ class CIRGenFunction : public CIRGenTypeCache {
//===--------------------------------------------------------------------===//
private:
template <typename Op>
- mlir::LogicalResult
- emitOpenACCOp(mlir::Location start, OpenACCDirectiveKind dirKind,
- SourceLocation dirLoc,
- llvm::ArrayRef<const OpenACCClause *> clauses);
+ Op emitOpenACCOp(mlir::Location start, OpenACCDirectiveKind dirKind,
+ SourceLocation dirLoc,
+ llvm::ArrayRef<const OpenACCClause *> clauses);
// Function to do the basic implementation of an operation with an Associated
// Statement. Models AssociatedStmtConstruct.
template <typename Op, typename TermOp>
diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
index 57ef06df068b7..688fca1bf2751 100644
--- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
@@ -317,10 +317,18 @@ class OpenACCClauseCIREmitter final
operation.getAsyncOperandsDeviceTypeAttr(),
createIntExpr(clause.getIntExpr()), range));
}
+ } else if constexpr (isOneOfTypes<OpTy, WaitOp>) {
+ // Wait doesn't have a device_type, so its handling here is slightly
+ //
diff erent.
+ if (!clause.hasIntExpr())
+ operation.setAsync(true);
+ else
+ operation.getAsyncOperandMutable().append(
+ createIntExpr(clause.getIntExpr()));
} else {
// TODO: When we've implemented this for everything, switch this to an
// unreachable. Combined constructs remain. Data, enter data, exit data,
- // update, wait, combined constructs remain.
+ // update, combined constructs remain.
return clauseNotImplemented(clause);
}
}
@@ -345,7 +353,7 @@ class OpenACCClauseCIREmitter final
void VisitIfClause(const OpenACCIfClause &clause) {
if constexpr (isOneOfTypes<OpTy, ParallelOp, SerialOp, KernelsOp, InitOp,
- ShutdownOp, SetOp, DataOp>) {
+ ShutdownOp, SetOp, DataOp, WaitOp>) {
operation.getIfCondMutable().append(
createCondition(clause.getConditionExpr()));
} else {
@@ -353,7 +361,7 @@ class OpenACCClauseCIREmitter final
// until we can write tests/know what we're doing with codegen to make
// sure we get it right.
// TODO: When we've implemented this for everything, switch this to an
- // unreachable. Enter data, exit data, host_data, update, wait, combined
+ // unreachable. Enter data, exit data, host_data, update, combined
// constructs remain.
return clauseNotImplemented(clause);
}
@@ -444,11 +452,9 @@ mlir::LogicalResult CIRGenFunction::emitOpenACCOpAssociatedStmt(
}
template <typename Op>
-mlir::LogicalResult CIRGenFunction::emitOpenACCOp(
+Op CIRGenFunction::emitOpenACCOp(
mlir::Location start, OpenACCDirectiveKind dirKind, SourceLocation dirLoc,
llvm::ArrayRef<const OpenACCClause *> clauses) {
- mlir::LogicalResult res = mlir::success();
-
llvm::SmallVector<mlir::Type> retTy;
llvm::SmallVector<mlir::Value> operands;
auto op = builder.create<Op>(start, retTy, operands);
@@ -461,7 +467,7 @@ mlir::LogicalResult CIRGenFunction::emitOpenACCOp(
makeClauseEmitter(op, *this, builder, dirKind, dirLoc)
.VisitClauseList(clauses);
}
- return res;
+ return op;
}
mlir::LogicalResult
@@ -500,22 +506,61 @@ CIRGenFunction::emitOpenACCDataConstruct(const OpenACCDataConstruct &s) {
mlir::LogicalResult
CIRGenFunction::emitOpenACCInitConstruct(const OpenACCInitConstruct &s) {
mlir::Location start = getLoc(s.getSourceRange().getBegin());
- return emitOpenACCOp<InitOp>(start, s.getDirectiveKind(), s.getDirectiveLoc(),
+ emitOpenACCOp<InitOp>(start, s.getDirectiveKind(), s.getDirectiveLoc(),
s.clauses());
+ return mlir::success();
}
mlir::LogicalResult
CIRGenFunction::emitOpenACCSetConstruct(const OpenACCSetConstruct &s) {
mlir::Location start = getLoc(s.getSourceRange().getBegin());
- return emitOpenACCOp<SetOp>(start, s.getDirectiveKind(), s.getDirectiveLoc(),
+ emitOpenACCOp<SetOp>(start, s.getDirectiveKind(), s.getDirectiveLoc(),
s.clauses());
+ return mlir::success();
}
mlir::LogicalResult CIRGenFunction::emitOpenACCShutdownConstruct(
const OpenACCShutdownConstruct &s) {
mlir::Location start = getLoc(s.getSourceRange().getBegin());
- return emitOpenACCOp<ShutdownOp>(start, s.getDirectiveKind(),
+ emitOpenACCOp<ShutdownOp>(start, s.getDirectiveKind(),
s.getDirectiveLoc(), s.clauses());
+ return mlir::success();
+}
+
+mlir::LogicalResult
+CIRGenFunction::emitOpenACCWaitConstruct(const OpenACCWaitConstruct &s) {
+ mlir::Location start = getLoc(s.getSourceRange().getBegin());
+ auto waitOp = emitOpenACCOp<WaitOp>(start, s.getDirectiveKind(),
+ s.getDirectiveLoc(), s.clauses());
+
+ auto createIntExpr = [this](const Expr *intExpr) {
+ mlir::Value expr = emitScalarExpr(intExpr);
+ mlir::Location exprLoc = cgm.getLoc(intExpr->getBeginLoc());
+
+ mlir::IntegerType targetType = mlir::IntegerType::get(
+ &getMLIRContext(), getContext().getIntWidth(intExpr->getType()),
+ intExpr->getType()->isSignedIntegerOrEnumerationType()
+ ? mlir::IntegerType::SignednessSemantics::Signed
+ : mlir::IntegerType::SignednessSemantics::Unsigned);
+
+ auto conversionOp = builder.create<mlir::UnrealizedConversionCastOp>(
+ exprLoc, targetType, expr);
+ return conversionOp.getResult(0);
+ };
+
+ // Emit the correct 'wait' clauses.
+ {
+ mlir::OpBuilder::InsertionGuard guardCase(builder);
+ builder.setInsertionPoint(waitOp);
+
+ if (s.hasDevNumExpr())
+ waitOp.getWaitDevnumMutable().append(createIntExpr(s.getDevNumExpr()));
+
+ for (Expr *QueueExpr : s.getQueueIdExprs())
+ waitOp.getWaitOperandsMutable().append(createIntExpr(QueueExpr));
+ }
+
+ return mlir::success();
}
mlir::LogicalResult
@@ -544,11 +589,6 @@ mlir::LogicalResult CIRGenFunction::emitOpenACCHostDataConstruct(
return mlir::failure();
}
mlir::LogicalResult
-CIRGenFunction::emitOpenACCWaitConstruct(const OpenACCWaitConstruct &s) {
- cgm.errorNYI(s.getSourceRange(), "OpenACC Wait Construct");
- return mlir::failure();
-}
-mlir::LogicalResult
CIRGenFunction::emitOpenACCUpdateConstruct(const OpenACCUpdateConstruct &s) {
cgm.errorNYI(s.getSourceRange(), "OpenACC Update Construct");
return mlir::failure();
diff --git a/clang/test/CIR/CodeGenOpenACC/wait.c b/clang/test/CIR/CodeGenOpenACC/wait.c
new file mode 100644
index 0000000000000..569846a91ab8a
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenACC/wait.c
@@ -0,0 +1,77 @@
+// RUN: %clang_cc1 -fopenacc -emit-cir -fclangir %s -o - | FileCheck %s
+
+void acc_wait(int cond) {
+ // CHECK: cir.func @acc_wait(%[[ARG:.*]]: !s32i{{.*}}) {
+ // CHECK-NEXT: %[[COND:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["cond", init]
+ // CHECK-NEXT: cir.store %[[ARG]], %[[COND]] : !s32i, !cir.ptr<!s32i>
+
+#pragma acc wait
+ // CHECK-NEXT: acc.wait
+
+#pragma acc wait if (cond)
+ // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
+ // CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast(int_to_bool, %[[COND_LOAD]] : !s32i), !cir.bool
+ // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[BOOL_CAST]] : !cir.bool to i1
+ // CHECK-NEXT: acc.wait if(%[[CONV_CAST]])
+
+#pragma acc wait async
+ // CHECK-NEXT: acc.wait attributes {async}
+
+#pragma acc wait async(cond)
+ // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
+ // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32
+ // CHECK-NEXT: acc.wait async(%[[CONV_CAST]] : si32) loc
+
+#pragma acc wait(1)
+ // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE_LITERAL]] : !s32i to si32
+ // CHECK-NEXT: acc.wait(%[[ONE_CAST]] : si32) loc
+
+#pragma acc wait(1, 2) async
+ // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE_LITERAL]] : !s32i to si32
+ // CHECK-NEXT: %[[TWO_LITERAL:.*]] = cir.const #cir.int<2> : !s32i
+ // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO_LITERAL]] : !s32i to si32
+ // CHECK-NEXT: acc.wait(%[[ONE_CAST]], %[[TWO_CAST]] : si32, si32) attributes {async}
+
+
+#pragma acc wait(queues:1) if (cond)
+ // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
+ // CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast(int_to_bool, %[[COND_LOAD]] : !s32i), !cir.bool
+ // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[BOOL_CAST]] : !cir.bool to i1
+ // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE_LITERAL]] : !s32i to si32
+ // CHECK-NEXT: acc.wait(%[[ONE_CAST]] : si32) if(%[[CONV_CAST]])
+
+#pragma acc wait(queues:1, 2) async(cond)
+ // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
+ // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32
+ // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE_LITERAL]] : !s32i to si32
+ // CHECK-NEXT: %[[TWO_LITERAL:.*]] = cir.const #cir.int<2> : !s32i
+ // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO_LITERAL]] : !s32i to si32
+ // CHECK-NEXT: acc.wait(%[[ONE_CAST]], %[[TWO_CAST]] : si32, si32) async(%[[CONV_CAST]] : si32) loc
+
+#pragma acc wait(devnum:1: 2, 3) if (cond)
+ // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
+ // CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast(int_to_bool, %[[COND_LOAD]] : !s32i), !cir.bool
+ // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[BOOL_CAST]] : !cir.bool to i1
+ // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE_LITERAL]] : !s32i to si32
+ // CHECK-NEXT: %[[TWO_LITERAL:.*]] = cir.const #cir.int<2> : !s32i
+ // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO_LITERAL]] : !s32i to si32
+ // CHECK-NEXT: %[[THREE_LITERAL:.*]] = cir.const #cir.int<3> : !s32i
+ // CHECK-NEXT: %[[THREE_CAST:.*]] = builtin.unrealized_conversion_cast %[[THREE_LITERAL]] : !s32i to si32
+ // CHECK-NEXT: acc.wait(%[[TWO_CAST]], %[[THREE_CAST]] : si32, si32) wait_devnum(%[[ONE_CAST]] : si32) if(%[[CONV_CAST]]) loc
+
+#pragma acc wait(devnum:1: queues: 2, 3) async
+ // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE_LITERAL]] : !s32i to si32
+ // CHECK-NEXT: %[[TWO_LITERAL:.*]] = cir.const #cir.int<2> : !s32i
+ // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO_LITERAL]] : !s32i to si32
+ // CHECK-NEXT: %[[THREE_LITERAL:.*]] = cir.const #cir.int<3> : !s32i
+ // CHECK-NEXT: %[[THREE_CAST:.*]] = builtin.unrealized_conversion_cast %[[THREE_LITERAL]] : !s32i to si32
+ // CHECK-NEXT: acc.wait(%[[TWO_CAST]], %[[THREE_CAST]] : si32, si32) wait_devnum(%[[ONE_CAST]] : si32) attributes {async}
+
+ // CHECK-NEXT: cir.return
+}
More information about the cfe-commits
mailing list