Author: Erich Keane Date: 2025-04-15T15:52:04-07:00 New Revision: af63e1b505453de3e6a281d1b72e62fa8d396b23
URL: https://github.com/llvm/llvm-project/commit/af63e1b505453de3e6a281d1b72e62fa8d396b23 DIFF: https://github.com/llvm/llvm-project/commit/af63e1b505453de3e6a281d1b72e62fa8d396b23.diff LOG: [OpenACC][CIR] Implement 'self' lowering on compute constructs (#135851) This is our first attempt at lowering a clause that is an 'operand' in the OpenACC operand, so it does quite a bit of refactoring. My previous plans on how to emit the clauses was not viable, so we instead do 'create the op, then use the visitor to fill in the operands'. This resulted in the 'applyAttributes' function getting removed and a few other functions simplified. Additionally, it requires setting the insertion point a little to make sure we're inserting 'around' the operation correctly. Finally, since the OpenACC dialect only understands the MLIR types, we had to introduce a use of the unrealized-conversion-cast, which we'll probably getting good use out of in the future. Added: Modified: clang/include/clang/AST/OpenACCClause.h clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp clang/test/CIR/CodeGenOpenACC/kernels.c clang/test/CIR/CodeGenOpenACC/parallel.c clang/test/CIR/CodeGenOpenACC/serial.c Removed: ################################################################################ diff --git a/clang/include/clang/AST/OpenACCClause.h b/clang/include/clang/AST/OpenACCClause.h index 3687af76a559f..681567228cbb0 100644 --- a/clang/include/clang/AST/OpenACCClause.h +++ b/clang/include/clang/AST/OpenACCClause.h @@ -430,6 +430,11 @@ class OpenACCSelfClause final } bool isConditionExprClause() const { return HasConditionExpr.has_value(); } + bool isVarListClause() const { return !isConditionExprClause(); } + bool isEmptySelfClause() const { + return (isConditionExprClause() && !hasConditionExpr()) || + (!isConditionExprClause() && getVarList().empty()); + } bool hasConditionExpr() const { assert(HasConditionExpr.has_value() && diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp index 152f996ed0fed..3bcc6f908a841 100644 --- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp @@ -32,46 +32,51 @@ constexpr bool isOneOfTypes = template <typename ToTest, typename T> constexpr bool isOneOfTypes<ToTest, T> = std::is_same_v<ToTest, T>; +template <typename OpTy> class OpenACCClauseCIREmitter final - : public OpenACCClauseVisitor<OpenACCClauseCIREmitter> { - CIRGenModule &cgm; + : public OpenACCClauseVisitor<OpenACCClauseCIREmitter<OpTy>> { + OpTy &operation; + CIRGenFunction &cgf; + CIRGenBuilderTy &builder; + // This is necessary since a few of the clauses emit diff erently based on the // directive kind they are attached to. OpenACCDirectiveKind dirKind; + // TODO(cir): This source location should be able to go away once the NYI + // diagnostics are gone. SourceLocation dirLoc; - struct AttributeData { - // Value of the 'default' attribute, added on 'data' and 'compute'/etc - // constructs as a 'default-attr'. - std::optional<ClauseDefaultValue> defaultVal = std::nullopt; - // For directives that have their device type architectures listed in - // attributes (init/shutdown/etc), the list of architectures to be emitted. - llvm::SmallVector<mlir::acc::DeviceType> deviceTypeArchs{}; - } attrData; - void clauseNotImplemented(const OpenACCClause &c) { - cgm.errorNYI(c.getSourceRange(), "OpenACC Clause", c.getClauseKind()); + cgf.cgm.errorNYI(c.getSourceRange(), "OpenACC Clause", c.getClauseKind()); } public: - OpenACCClauseCIREmitter(CIRGenModule &cgm, OpenACCDirectiveKind dirKind, - SourceLocation dirLoc) - : cgm(cgm), dirKind(dirKind), dirLoc(dirLoc) {} + OpenACCClauseCIREmitter(OpTy &operation, CIRGenFunction &cgf, + CIRGenBuilderTy &builder, + OpenACCDirectiveKind dirKind, SourceLocation dirLoc) + : operation(operation), cgf(cgf), builder(builder), dirKind(dirKind), + dirLoc(dirLoc) {} void VisitClause(const OpenACCClause &clause) { clauseNotImplemented(clause); } void VisitDefaultClause(const OpenACCDefaultClause &clause) { - switch (clause.getDefaultClauseKind()) { - case OpenACCDefaultClauseKind::None: - attrData.defaultVal = ClauseDefaultValue::None; - break; - case OpenACCDefaultClauseKind::Present: - attrData.defaultVal = ClauseDefaultValue::Present; - break; - case OpenACCDefaultClauseKind::Invalid: - break; + // This type-trait checks if 'op'(the first arg) is one of the mlir::acc + // operations listed in the rest of the arguments. + if constexpr (isOneOfTypes<OpTy, ParallelOp, SerialOp, KernelsOp, DataOp>) { + switch (clause.getDefaultClauseKind()) { + case OpenACCDefaultClauseKind::None: + operation.setDefaultAttr(ClauseDefaultValue::None); + break; + case OpenACCDefaultClauseKind::Present: + operation.setDefaultAttr(ClauseDefaultValue::Present); + break; + case OpenACCDefaultClauseKind::Invalid: + break; + } + } else { + return clauseNotImplemented(clause); } } @@ -89,64 +94,70 @@ class OpenACCClauseCIREmitter final } void VisitDeviceTypeClause(const OpenACCDeviceTypeClause &clause) { + if constexpr (isOneOfTypes<OpTy, InitOp, ShutdownOp>) { + llvm::SmallVector<mlir::Attribute> deviceTypes; + std::optional<mlir::ArrayAttr> existingDeviceTypes = + operation.getDeviceTypes(); + + // Ensure we keep the existing ones, and in the correct 'new' order. + if (existingDeviceTypes) { + for (const mlir::Attribute &Attr : *existingDeviceTypes) + deviceTypes.push_back(mlir::acc::DeviceTypeAttr::get( + builder.getContext(), + cast<mlir::acc::DeviceTypeAttr>(Attr).getValue())); + } - switch (dirKind) { - case OpenACCDirectiveKind::Init: - case OpenACCDirectiveKind::Set: - case OpenACCDirectiveKind::Shutdown: { - // Device type has a list that is either a 'star' (emitted as 'star'), - // or an identifer list, all of which get added for attributes. - - for (const DeviceTypeArgument &arg : clause.getArchitectures()) - attrData.deviceTypeArchs.push_back(decodeDeviceType(arg.first)); - break; - } - default: + for (const DeviceTypeArgument &arg : clause.getArchitectures()) { + deviceTypes.push_back(mlir::acc::DeviceTypeAttr::get( + builder.getContext(), decodeDeviceType(arg.first))); + } + operation.removeDeviceTypesAttr(); + operation.setDeviceTypesAttr( + mlir::ArrayAttr::get(builder.getContext(), deviceTypes)); + } else if constexpr (isOneOfTypes<OpTy, SetOp>) { + assert(!operation.getDeviceTypeAttr() && "already have device-type?"); + assert(clause.getArchitectures().size() <= 1); + + if (!clause.getArchitectures().empty()) + operation.setDeviceType( + decodeDeviceType(clause.getArchitectures()[0].first)); + } else { return clauseNotImplemented(clause); } } - // Apply any of the clauses that resulted in an 'attribute'. - template <typename Op> - void applyAttributes(CIRGenBuilderTy &builder, Op &op) { - - if (attrData.defaultVal.has_value()) { - // FIXME: OpenACC: as we implement this for other directive kinds, we have - // to expand this list. - // This type-trait checks if 'op'(the first arg) is one of the mlir::acc - // operations listed in the rest of the arguments. - if constexpr (isOneOfTypes<Op, ParallelOp, SerialOp, KernelsOp, DataOp>) - op.setDefaultAttr(*attrData.defaultVal); - else - cgm.errorNYI(dirLoc, "OpenACC 'default' clause lowering for ", dirKind); - } - - if (!attrData.deviceTypeArchs.empty()) { - // FIXME: OpenACC: as we implement this for other directive kinds, we have - // to expand this list, or more likely, have a 'noop' branch as most other - // uses of this apply to the operands instead. - // This type-trait checks if 'op'(the first arg) is one of the mlir::acc - if constexpr (isOneOfTypes<Op, InitOp, ShutdownOp>) { - llvm::SmallVector<mlir::Attribute> deviceTypes; - for (mlir::acc::DeviceType DT : attrData.deviceTypeArchs) - deviceTypes.push_back( - mlir::acc::DeviceTypeAttr::get(builder.getContext(), DT)); - - op.setDeviceTypesAttr( - mlir::ArrayAttr::get(builder.getContext(), deviceTypes)); - } else if constexpr (isOneOfTypes<Op, SetOp>) { - assert(attrData.deviceTypeArchs.size() <= 1 && - "Set can only have a single architecture"); - if (!attrData.deviceTypeArchs.empty()) - op.setDeviceType(attrData.deviceTypeArchs[0]); + void VisitSelfClause(const OpenACCSelfClause &clause) { + if constexpr (isOneOfTypes<OpTy, ParallelOp, SerialOp, KernelsOp>) { + if (clause.isEmptySelfClause()) { + operation.setSelfAttr(true); + } else if (clause.isConditionExprClause()) { + assert(clause.hasConditionExpr()); + mlir::Value condition = + cgf.evaluateExprAsBool(clause.getConditionExpr()); + + mlir::Location exprLoc = + cgf.cgm.getLoc(clause.getConditionExpr()->getBeginLoc()); + mlir::IntegerType targetType = mlir::IntegerType::get( + &cgf.getMLIRContext(), /*width=*/1, + mlir::IntegerType::SignednessSemantics::Signless); + auto conversionOp = builder.create<mlir::UnrealizedConversionCastOp>( + exprLoc, targetType, condition); + operation.getSelfCondMutable().append(conversionOp.getResult(0)); } else { - cgm.errorNYI(dirLoc, "OpenACC 'device_type' clause lowering for ", - dirKind); + llvm_unreachable("var-list version of self shouldn't get here"); } + } else { + return clauseNotImplemented(clause); } } }; +template <typename OpTy> +auto makeClauseEmitter(OpTy &op, CIRGenFunction &cgf, CIRGenBuilderTy &builder, + OpenACCDirectiveKind dirKind, SourceLocation dirLoc) { + return OpenACCClauseCIREmitter<OpTy>(op, cgf, builder, dirKind, dirLoc); +} + } // namespace template <typename Op, typename TermOp> @@ -158,24 +169,27 @@ mlir::LogicalResult CIRGenFunction::emitOpenACCOpAssociatedStmt( llvm::SmallVector<mlir::Type> retTy; llvm::SmallVector<mlir::Value> operands; - - // Clause-emitter must be here because it might modify operands. - OpenACCClauseCIREmitter clauseEmitter(getCIRGenModule(), dirKind, dirLoc); - clauseEmitter.VisitClauseList(clauses); - auto op = builder.create<Op>(start, retTy, operands); - // Apply the attributes derived from the clauses. - clauseEmitter.applyAttributes(builder, op); + { + mlir::OpBuilder::InsertionGuard guardCase(builder); + // Sets insertion point before the 'op', since every new expression needs to + // be before the operation. + builder.setInsertionPoint(op); + makeClauseEmitter(op, *this, builder, dirKind, dirLoc) + .VisitClauseList(clauses); + } - mlir::Block &block = op.getRegion().emplaceBlock(); - mlir::OpBuilder::InsertionGuard guardCase(builder); - builder.setInsertionPointToEnd(&block); + { + mlir::Block &block = op.getRegion().emplaceBlock(); + mlir::OpBuilder::InsertionGuard guardCase(builder); + builder.setInsertionPointToEnd(&block); - LexicalScope ls{*this, start, builder.getInsertionBlock()}; - res = emitStmt(associatedStmt, /*useCurrentScope=*/true); + LexicalScope ls{*this, start, builder.getInsertionBlock()}; + res = emitStmt(associatedStmt, /*useCurrentScope=*/true); - builder.create<TermOp>(end); + builder.create<TermOp>(end); + } return res; } @@ -187,14 +201,16 @@ mlir::LogicalResult CIRGenFunction::emitOpenACCOp( llvm::SmallVector<mlir::Type> retTy; llvm::SmallVector<mlir::Value> operands; - - // Clause-emitter must be here because it might modify operands. - OpenACCClauseCIREmitter clauseEmitter(getCIRGenModule(), dirKind, dirLoc); - clauseEmitter.VisitClauseList(clauses); - auto op = builder.create<Op>(start, retTy, operands); - // Apply the attributes derived from the clauses. - clauseEmitter.applyAttributes(builder, op); + + { + mlir::OpBuilder::InsertionGuard guardCase(builder); + // Sets insertion point before the 'op', since every new expression needs to + // be before the operation. + builder.setInsertionPoint(op); + makeClauseEmitter(op, *this, builder, dirKind, dirLoc) + .VisitClauseList(clauses); + } return res; } @@ -254,46 +270,46 @@ mlir::LogicalResult CIRGenFunction::emitOpenACCShutdownConstruct( mlir::LogicalResult CIRGenFunction::emitOpenACCLoopConstruct(const OpenACCLoopConstruct &s) { - getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Loop Construct"); + cgm.errorNYI(s.getSourceRange(), "OpenACC Loop Construct"); return mlir::failure(); } mlir::LogicalResult CIRGenFunction::emitOpenACCCombinedConstruct( const OpenACCCombinedConstruct &s) { - getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Combined Construct"); + cgm.errorNYI(s.getSourceRange(), "OpenACC Combined Construct"); return mlir::failure(); } mlir::LogicalResult CIRGenFunction::emitOpenACCEnterDataConstruct( const OpenACCEnterDataConstruct &s) { - getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC EnterData Construct"); + cgm.errorNYI(s.getSourceRange(), "OpenACC EnterData Construct"); return mlir::failure(); } mlir::LogicalResult CIRGenFunction::emitOpenACCExitDataConstruct( const OpenACCExitDataConstruct &s) { - getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC ExitData Construct"); + cgm.errorNYI(s.getSourceRange(), "OpenACC ExitData Construct"); return mlir::failure(); } mlir::LogicalResult CIRGenFunction::emitOpenACCHostDataConstruct( const OpenACCHostDataConstruct &s) { - getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC HostData Construct"); + cgm.errorNYI(s.getSourceRange(), "OpenACC HostData Construct"); return mlir::failure(); } mlir::LogicalResult CIRGenFunction::emitOpenACCWaitConstruct(const OpenACCWaitConstruct &s) { - getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Wait Construct"); + cgm.errorNYI(s.getSourceRange(), "OpenACC Wait Construct"); return mlir::failure(); } mlir::LogicalResult CIRGenFunction::emitOpenACCUpdateConstruct(const OpenACCUpdateConstruct &s) { - getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Update Construct"); + cgm.errorNYI(s.getSourceRange(), "OpenACC Update Construct"); return mlir::failure(); } mlir::LogicalResult CIRGenFunction::emitOpenACCAtomicConstruct(const OpenACCAtomicConstruct &s) { - getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Atomic Construct"); + cgm.errorNYI(s.getSourceRange(), "OpenACC Atomic Construct"); return mlir::failure(); } mlir::LogicalResult CIRGenFunction::emitOpenACCCacheConstruct(const OpenACCCacheConstruct &s) { - getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Cache Construct"); + cgm.errorNYI(s.getSourceRange(), "OpenACC Cache Construct"); return mlir::failure(); } diff --git a/clang/test/CIR/CodeGenOpenACC/kernels.c b/clang/test/CIR/CodeGenOpenACC/kernels.c index 0c950fe3d0f9c..934daf9e8ecc0 100644 --- a/clang/test/CIR/CodeGenOpenACC/kernels.c +++ b/clang/test/CIR/CodeGenOpenACC/kernels.c @@ -1,7 +1,9 @@ // RUN: %clang_cc1 -fopenacc -emit-cir -fclangir %s -o - | FileCheck %s -void acc_kernels(void) { - // CHECK: cir.func @acc_kernels() { +void acc_kernels(int cond) { + // CHECK: cir.func @acc_kernels(%[[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 kernels {} @@ -38,5 +40,29 @@ void acc_kernels(void) { // CHECK-NEXT: acc.terminator // CHECK-NEXT:} +#pragma acc kernels self + {} + // CHECK-NEXT: acc.kernels { + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } attributes {selfAttr} + +#pragma acc kernels self(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.kernels self(%[[CONV_CAST]]) { + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + +#pragma acc kernels self(0) + {} + // CHECK-NEXT: %[[ZERO_LITERAL:.*]] = cir.const #cir.int<0> : !s32i + // CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast(int_to_bool, %[[ZERO_LITERAL]] : !s32i), !cir.bool + // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[BOOL_CAST]] : !cir.bool to i1 + // CHECK-NEXT: acc.kernels self(%[[CONV_CAST]]) { + // CHECK-NEXT: acc.terminator + // CHECK-NEXT: } loc + // CHECK-NEXT: cir.return } diff --git a/clang/test/CIR/CodeGenOpenACC/parallel.c b/clang/test/CIR/CodeGenOpenACC/parallel.c index e18270435460c..c7a4bda6faa74 100644 --- a/clang/test/CIR/CodeGenOpenACC/parallel.c +++ b/clang/test/CIR/CodeGenOpenACC/parallel.c @@ -1,7 +1,9 @@ // RUN: %clang_cc1 -fopenacc -emit-cir -fclangir %s -o - | FileCheck %s -void acc_parallel(void) { - // CHECK: cir.func @acc_parallel() { +void acc_parallel(int cond) { + // CHECK: cir.func @acc_parallel(%[[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 parallel {} // CHECK-NEXT: acc.parallel { @@ -37,5 +39,29 @@ void acc_parallel(void) { // CHECK-NEXT: acc.yield // CHECK-NEXT:} +#pragma acc parallel self + {} + // CHECK-NEXT: acc.parallel { + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } attributes {selfAttr} + +#pragma acc parallel self(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.parallel self(%[[CONV_CAST]]) { + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + +#pragma acc parallel self(0) + {} + // CHECK-NEXT: %[[ZERO_LITERAL:.*]] = cir.const #cir.int<0> : !s32i + // CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast(int_to_bool, %[[ZERO_LITERAL]] : !s32i), !cir.bool + // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[BOOL_CAST]] : !cir.bool to i1 + // CHECK-NEXT: acc.parallel self(%[[CONV_CAST]]) { + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: cir.return } diff --git a/clang/test/CIR/CodeGenOpenACC/serial.c b/clang/test/CIR/CodeGenOpenACC/serial.c index 72a0995549da3..38a38ad6c9514 100644 --- a/clang/test/CIR/CodeGenOpenACC/serial.c +++ b/clang/test/CIR/CodeGenOpenACC/serial.c @@ -1,7 +1,9 @@ // RUN: %clang_cc1 -fopenacc -emit-cir -fclangir %s -o - | FileCheck %s -void acc_serial(void) { - // CHECK: cir.func @acc_serial() { +void acc_serial(int cond) { + // CHECK: cir.func @acc_serial(%[[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 serial {} @@ -38,5 +40,29 @@ void acc_serial(void) { // CHECK-NEXT: acc.yield // CHECK-NEXT:} +#pragma acc serial self + {} + // CHECK-NEXT: acc.serial { + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } attributes {selfAttr} + +#pragma acc serial self(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.serial self(%[[CONV_CAST]]) { + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + +#pragma acc serial self(0) + {} + // CHECK-NEXT: %[[ZERO_LITERAL:.*]] = cir.const #cir.int<0> : !s32i + // CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast(int_to_bool, %[[ZERO_LITERAL]] : !s32i), !cir.bool + // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[BOOL_CAST]] : !cir.bool to i1 + // CHECK-NEXT: acc.serial self(%[[CONV_CAST]]) { + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: cir.return } _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
