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 +} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
